search

acts-project / vecmem

Vectorised data model base and helper classes.
https://acts-project.github.io/vecmem/
Mozilla Public License 2.0
19 stars 13 forks source link

Introduce vecmem::sycl::local_accessor, main branch (2023.05.31.) #230

Closed krasznaa closed 1 year ago

krasznaa commented 1 year ago

While in traccc I told @guilhermeAlmeida1 that we shouldn't over-think how we use local/shared memory with SYCL, in this project I finally changed my mind...

With modern versions of Intel's compiler there is a warning about using ::sycl::access::target::local on a "normal" ::sycl::accessor. One is rather supposed to use ::sycl::local_accessor nowadays. So to help with using the correct type in our code, I now introduced vecmem::sycl::local_accessor, which would resolve to the correct type with whatever SYCL compiler we use. By checking during CMake configuration whether ::sycl::local_accessor is usable, and if it is, using that. If it's not, it falls back on using ::sycl::accessor.

I thought this would be a more robust setup than to rely on some compiler macros.

As long as people are on board with this, we should eventually update the traccc code to use this type as well.

guilhermeAlmeida1 commented 1 year ago

This works fine by me, but I have a question. If Intel's design choice is to send these kinds of warnings or disfavour backwards compatiblity, should it be up to vecmem to provide a SYCL to "old-SYCL" common interface, when it's no longer about common code between different platforms but a single one?

krasznaa commented 1 year ago

The sycl::local_accessor type is part of the SYCL2020 specification. "SYCL versions" are not handled as rigorously by the compilers as C++ versions just yet, that's where much of this is coming from.

In the future, once all SYCL compilers will properly respect a well defined set of SYCL specification versions, user code will be a bit easier to manage. As at that point it should be possible to tell newer compiler versions to stick to an older SYCL version. But for now checking whether syc::local_accessor is available, and using it if it is, seems to be the best approach to me.