Description
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.