[CUDA] Add support for the generic address space

[ghost]

The SYCL 2020 specification requires that the sycl::atomic_ref class template uses the sycl::access::address_space::generic_space enumeration as the default value for the AddressSpace template parameter. The enumeration and its support have been added to the project recently.

On CUDA, attempts of instantiate sycl::atomic_ref with the generic address space lead to linker errors. When the end-to-end tests for sycl::atomic_ref from https://github.com/intel/llvm-test-suite/tree/intel/SYCL/AtomicRef with generic address space (all the tests with suffix _generic.cpp) are compiled with CUDA support, the following linker errors occur:

ptxas fatal   : Unresolved extern function '_Z19__spirv_AtomicStorePiN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagEi'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z19__spirv_AtomicStorePmN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagEm'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z18__spirv_AtomicLoadPKjN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z18__spirv_AtomicLoadPKmN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z18__spirv_AtomicLoadPKiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z29__spirv_AtomicCompareExchangePiN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_ii'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z29__spirv_AtomicCompareExchangePmN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_mm'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z22__spirv_AtomicExchangePiN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagEi'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z22__spirv_AtomicExchangePmN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagEm'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z18__spirv_AtomicSMaxPiN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagEi'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)
...

ptxas fatal   : Unresolved extern function '_Z18__spirv_AtomicSMinPiN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagEi'
llvm-foreach: 
clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation)

It looks like the libclc library doesn't provide the required functions. The functions should be available for the linker.

[tomflinda]

@t4c1, Could you provide the time when the feature of support for the generic address space is added to the CUDA backend?

[t4c1]

There is no exact timeline planned. See #8807.

[npmiller]

Closing as this was fixed in: #7391

The tests have been re-enabled for CUDA in intel/llvm-test-suite#1446 and have been passing.

Please feel free to re-open if the issue persists.