DPCT1007:2: Migration of atomicDec is not supported

[zjin-lcf]

DPCT1007:2: Migration of atomicDec is not supported. However, atomicInc() can be migrated.

I posted a feature request (intel/llvm#8220).

Could you please explain that the implementation of atomicDec() below does not work ? Thanks.

        // atomicDec(&gpuData[8], 137)

        auto ao8 = atomic_ref<T, memory_order::relaxed, memory_scope::device,
                              access::address_space::global_space>(gpuData[8]);
        while (true) {
          T old = ao8.load();
          if (old == 0 || old > 137) {
            if (ao8.compare_exchange_strong(old, 137)) break;
          } else if (ao8.compare_exchange_strong(old, old - 1)) break;
        }

[tomflinda]

@zjin-lcf Could you provide a complete executable reproducer?

[zjin-lcf]

Please see the example

__global__ void f (unsigned int *d) {
   atomicDec(d, 137);
}

[tomflinda]

@zjin-lcf the example you provided is just a device function, I want you to provide a complete executable reproducer, so that I can reproduce the issue you reported,
Or pls try the WA to migrate example your provided

__global__ void f (unsigned int *d) {
   atomicDec(d, 137);
}

with:

template <sycl::access::address_space addressSpace = sycl::access::address_space::global_space,
          sycl::memory_order memoryOrder = sycl::memory_order::relaxed,
          sycl::memory_scope memoryScope = sycl::memory_scope::device>
inline unsigned int atomic_fetch_compare_dec(unsigned int *addr,
                                             unsigned int operand) {
  auto atm = sycl::atomic_ref<unsigned int, memoryOrder, memoryScope,
                                  addressSpace>(addr[0]);
  unsigned int old;
 
	while (true) {
	  old = atm.load();
	  if (old == 0 || old > operand) {
		if (atm.compare_exchange_strong(old, operand)) break;
	  } else if (atm.compare_exchange_strong(old, old - 1)) break;
	}  
  
  return old;
}

void f (unsigned int *d) {
    atomic_fetch_compare_dec<sycl::access::address_space::generic_space>(d, 137);
}

[zjin-lcf]

Sorry. Please see the example in CUDA: https://github.com/zjin-lcf/HeCBench/tree/master/atomicIntrinsics-cuda. It was updated with the addition of atomicInc and atomicDec yesterday.

The SYCL "atomicDec" function hangs, but the SYCL "atomicInc" function works.

Thanks.

[tomflinda]

@zjin-lcf the possible root cause is that the work-group size passed to the SYCL kernel may exceed the HW
limit.

You can try it to do the following changes:

diff --git a/atomicIntrinsics-cuda/main.cu b/atomicIntrinsics-cuda/main.cu
index e72cc587..d433e0d4 100644
--- a/atomicIntrinsics-cuda/main.cu
+++ b/atomicIntrinsics-cuda/main.cu
@@ -26,8 +26,8 @@
 template <class T>
 void testcase(const int repeat)
 {
-  unsigned int len = 1 << 27;
-  unsigned int numThreads = 256;
+  unsigned int len = 256;
+  unsigned int numThreads = 16;
   unsigned int numBlocks = (len + numThreads - 1) / numThreads;
   T gpuData[] = {0, 0, (T)-256, 256, 255, 0, 255, 0, 0};
   unsigned int memSize = sizeof(gpuData);

Then do the migration and apply the WA for "atomicDec((unsigned int *)&g_odata[8], 137)" in the migrated code kernel.h:

template <sycl::access::address_space addressSpace = sycl::access::address_space::global_space,
          sycl::memory_order memoryOrder = sycl::memory_order::relaxed,
          sycl::memory_scope memoryScope = sycl::memory_scope::device>
inline unsigned int atomic_fetch_compare_dec(unsigned int *addr,
                                             unsigned int operand) {
  auto atm = sycl::atomic_ref<unsigned int, memoryOrder, memoryScope,
                                  addressSpace>(addr[0]);
  unsigned int old;
 
	while (true) {
	  old = atm.load();
	  if (old == 0 || old > operand) {
		if (atm.compare_exchange_strong(old, operand)) break;
	  } else if (atm.compare_exchange_strong(old, old - 1)) break;
	}  
  
  return old;
}

template <class T>
void testKernel(T *g_odata, sycl::nd_item<3> item_ct1)
{
   .....
    // Atomic decrement
    /*
    DPCT1007:2: Migration of atomicDec is not supported.
    */
    //atomicDec((unsigned int *)&g_odata[8], 137);
	
    atomic_fetch_compare_dec<sycl::access::address_space::generic_space>(
        (unsigned int *)&g_odata[8], (unsigned int)137);
}

Then the test case should pass.

[zjin-lcf]

Thank you for the answer. I am confused that the Cuda program passes the test and the length and work-group size need the proposed changes in Sycl program. Both programs run on an Nvidia GPU. Could you please explain the hardware limit ?

[tomflinda]

@zjin-lcf it seems a bug of SYCL runtime, we will report this issue to SYCL compiler team.

[zjin-lcf]

Thanks.

[tomflinda]

@zjin-lcf #595 has implemented the migration of atomicDec, for the root cause that why migrated SYCL code without changing the length and work-group size hangs on Nvidia GPU, I have reported it to SYCL compiler team, and waiting for their feedback.

[zjin-lcf]

Thank you. I look forward to understanding the root cause.

[zjin-lcf]

Do you have feedback from the compiler team ?

[tomflinda]

@zjin-lcf the latest feedback said that the issue is under investigation, since this issue is not related to SYComatic, I have closed it.