question about atomic add over the address

[zjin-lcf]

Can you please suggest a solution (a SYCL program) to the following Cuda program ? In the Cuda program, the atomic operations are performed over the casted address rather than the elements pointed to by "a".

Thank you.

#include <cuda.h>
#include <stdio.h>

__managed__ char* a;

__global__ void test() {
  size_t r = atomicAdd((unsigned long long*)&a, (size_t)1);
}

int main() {

 cudaMallocManaged((void**)&a, 64);
 for (int i = 0; i < 64; i++) a[i] = i;

 test<<<1, 8>>>();

 cudaDeviceSynchronize();
 cudaFree(a);

 return 0;
}

[keryell]

What about using sycl::atomic_ref?
https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:atomic-references

[zjin-lcf]

Thank you for your reply. I have trouble migrating the example to SYCL with atomic_ref.

[zjin-lcf]

DPCT (https://github.com/oneapi-src/SYCLomatic) can convert the example successfully.
Compiling the example with the CUDA support shows the following error:

ptxas fatal : Unresolved extern function '_Z18__spirv_AtomicIAddPyN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagEy'

[AlexeySachkov]

@zjin-lcf, as @keryell pointed out, sycl::atomic_ref is what you should be using. Regarding unresolved extern function error, I would suggest to open an issue, because it looks like a bug. Perhaps we have some pieces missing for some backends