Add support for device and constant global variables in the SYCL backend

[AuroraPerego]

This tries to fix #2070.
This implementation changes the current API.
The declaration of device memory is done with the macro ALPAKA_STATIC_ACC_MEM_GLOBAL(type, name) that wraps the CUDA/HIP/serial variables in a struct alpaka::DevGlobal and declares a sycl::ext::oneapi::experimental::device_global<type> for SYCL. The inline attribute is used to ensure that only one instance of that variable exists across different translation units.
name is used for the memcpy, while name.get() must be used in the kernel to align with the behavior of the SYCL backend.
The memcpy has been specialized for the device global variables.

The test with the SYCL backend failed with the original KernelExecutionFixture because it creates a new queue instead of using the one used for the memcpy. I added a constructor that takes in input also the queue.

Another issue with the test is that being compiled with the flags to enable the SYCL backend, the macro for the device global variable expands to the SYCL one (sycl::ext::oneapi::experimental::device_global<type>) and therefore it fails when running on the AccCpuSerial. I have disabled this accelerator just for this test.
Note that this happens also with the CUDA backend (with the macro expanded to __device__ alpaka::DevGlobal<type> name), but for some strange reason it works on the serial backend.

Thanks to @fwyzard for the help :)

[fwyzard]

Many thanks for this work !

As it introduces a breaking changes to the API, I would ask for enough time to consider all impacts before merging the changes.

[fwyzard]

why ALPAKA_NO_HOST_ACC_WARNING ?

[AuroraPerego]

copied from the test (https://github.com/alpaka-group/alpaka/blob/develop/test/unit/mem/view/src/ViewStaticAccMem.cpp#L31)

[fwyzard]

why ALPAKA_NO_HOST_ACC_WARNING ?

[AuroraPerego]

same as above

[fwyzard]

Suggested change

	/* Copyright 2023 Axel Hübl, Benjamin Worpitz, Matthias Werner, Andrea Bocci, Jan Stephan, Bernhard Manfred Gruber,
	* Aurora Perego
	/* Copyright 2024 Axel Hübl, Benjamin Worpitz, Matthias Werner, Andrea Bocci, Jan Stephan, Bernhard Manfred Gruber,
	* Aurora Perego

[fwyzard]

This makes device_global visible to all in the alpaka namespace, which may not be a good idea.

Could you move it at least inside the detail namespace (and use it as detail::device_global below) ?

Or, just use it fully expanded everywhere.

[AuroraPerego]

I've used it fully expanded

[fwyzard]

Suggested change

	// CUDA implementation
	// HIP/ROCm implementation

[fwyzard]

This approach seems too generic, as in principle it can match also for a CUDA DevGlobalImplGeneric.

Also, TAcc should be TTag.

[fwyzard]

Can you add a constrain that TAcc (actually TTag) is one of TagCpuOmp2Blocks, TagCpuOmp2Threads, TagCpuSerial, TagCpuTbbBlocks, TagCpuThreads ?

There may be a smarter way, but at least this should avoid the wrong matches.

[fwyzard]

By the way, IIUC TViewDstFwd here is not really a View, but the underlying type of the global variable.
If that's the case, could you rename it to TType ?

[AuroraPerego]

Can you add a constrain that TAcc (actually TTag) is one of TagCpuOmp2Blocks, TagCpuOmp2Threads, TagCpuSerial, TagCpuTbbBlocks, TagCpuThreads ?

ok, I'll try (and I'll probably also have to prevent the test from running on AccCpuSerial when it is compiled with the flags to enable CUDA/HIP/SYCL)

By the way, IIUC TViewDstFwd here is not really a View, but the underlying type of the global variable. If that's the case, could you rename it to TType ?

You are right, I'll change it

[fwyzard]

TAcc should be TTag, it should somehoe match the TApi (we don't want to use CUDA on a HIP/ROCm global variable).

[AuroraPerego]

I don't know if there is a map between TTag and TApi, but I'll look into it

[AuroraPerego]

Device global variables in oneAPI versions 2023.* require explicitly the device image scope in the constructor, while from 2024.0 this is not needed anymore. For this reason, the required compiler version has been changed to 2024.0.

Now the only failures in the tests are some warnings (that become errors when compiling the tests) in the SYCL headers / footers.

[fwyzard]

Now the only failures in the tests are some warnings (that become errors when compiling the tests) in the SYCL headers / footers.

@psychocoderHPC can we remove -Wreserved-identifier -Wold-style-cast from the flags passed to icpx ?
The offending code is autogenerated, and at least for -Wreserved-identifier is included before any user or alpaka code, so we cannot disable the warnings with a #pragma in the code.

[psychocoderHPC]

Now the only failures in the tests are some warnings (that become errors when compiling the tests) in the SYCL headers / footers.

@psychocoderHPC can we remove -Wreserved-identifier -Wold-style-cast from the flags passed to icpx ? The offending code is autogenerated, and at least for -Wreserved-identifier is included before any user or alpaka code, so we cannot disable the warnings with a #pragma in the code.

offline discussed: yes we can remove the options if nessesary

[SimeonEhrig]

I would prefer passing the Acc type directly instead using the type trait alpaka::trait::AccToTag. I think the common case is, that we have the Acc type available and not the tag type, therefore we would have a lot of copy/past code.

Depending on if it makes sense to provide the possibility to use tags in the user API, you can ether move typename alpaka::trait::AccToTag<Acc>::type to the implementation (no support for tags in the user API) or do some kind of overload and check if the given type is a alpaka accelerator type or tag.

[AuroraPerego]

I've moved it to the implementation

template<typename TAcc, typename T>
    using DevGlobal = typename detail::DevGlobalTrait<typename alpaka::trait::AccToTag<TAcc>::type, T>::Type;

[psychocoderHPC]

@alpaka-group/alpaka-maintainers IMO we can merge this PR. Are there any voices against merging it?

[fwyzard]

OK for me.

Any further developments can happen in follow up PRs.

[psychocoderHPC]

@AuroraPerego We need to review this change again because of the bug shown #2259

In the original code we had constant global memory only and in the new code we have global memory and constant memory.
With your changes is there a difference between ALPAKA_STATIC_ACC_MEM_GLOBAL and ALPAKA_STATIC_ACC_MEM_CONSTANT?

[psychocoderHPC]

IMO in this section we should only use ALPAKA_STATIC_ACC_MEM_CONSTANT and later we have a test for ALPAKA_STATIC_ACC_MEM_GLOBAL

[AuroraPerego]

In the original code we had constant global memory only and in the new code we have global memory and constant memory.

They were different also in the original code, with ALPAKA_STATIC_ACC_MEM_GLOBAL expanding to __device__ and ALPAKA_STATIC_ACC_MEM_CONSTANT expanding to __constant__.

With your changes is there a difference between ALPAKA_STATIC_ACC_MEM_GLOBAL and ALPAKA_STATIC_ACC_MEM_CONSTANT?

Yes, they are using the __device__ and __constant__ attributes respectively (plus the inline attribute)

IMO in this section we should only use ALPAKA_STATIC_ACC_MEM_CONSTANT and later we have a test for ALPAKA_STATIC_ACC_MEM_GLOBAL

ok, but why? we were testing both here also before this PR

[psychocoderHPC]

ok, but why? we were testing both here also before this PR

This is a good question.

What if we define ALPAKA_STATIC_ACC_MEM_GLOBAL static instead of inline this should create a single instance too.

[AuroraPerego]

it compiles and also solves the issue in #2259, I can make a PR with the change if you think it's correct (I don't know enough about static, inline and extern to be sure of that)

[fwyzard]

In general, what macros should be used to support all use cases ?

file1.cu

static __device__ int i = 0;

and

file2.h

extern __device__ int i;

file2.cu

#include "file2.h"

__device__ int i = 0;

?

[AuroraPerego]

Thinking about it, maybe static is not what we want. The ideal would be to make the extern keyword work also with the new implementation to allow the file2.h/file2.cu case, which IIUC is not possible with static.