Merge "Support for Atomic inc and dec in HIP" into amd-master

include/hcc_detail/hip_runtime.h

@@ -376,6 +376,15 @@ __device__ unsigned int atomicXor(unsigned int* address,
 __device__ unsigned long long int atomicXor(unsigned long long int* address,
                                  unsigned long long int val);
 
+//atomicInc()
+__device__ unsigned int atomicInc(unsigned int* address,
+                       unsigned int val);
+
+
+//atomicDec()
+__device__ unsigned int atomicDec(unsigned int* address,
+                       unsigned int val);
+
 
 // integer intrinsic function __poc __clz __ffs __brev
 __device__ unsigned int __popc( unsigned int input);

src/device_util.cpp

@@ -657,7 +657,19 @@ __device__  unsigned long long int atomicXor(unsigned long long int* address,
 	return (long long int)hc::atomic_fetch_xor((uint64_t*)address,(uint64_t)val);
 }
 
+//atomicInc
+__device__  int atomicInc(unsigned int* address,
+                       unsigned int val)
+{
+	return hc::__atomic_wrapinc(address,val);
+}
 
+//atomicDec
+__device__  int atomicDec(unsigned int* address,
+                       unsigned int val)
+{
+	return hc::__atomic_wrapdec(address,val);
+}
 
 
 __device__ unsigned int test__popc(unsigned int input)

tests/src/hipSimpleAtomicsTest.cpp

@@ -121,8 +121,7 @@ int computeGold(int *gpuData, const int len)
 
     for (int i = 0; i < len; ++i)
     {
-        //val = (val >= limit) ? 0 : val+1;
-        val = val+1;
+        val = (val >= limit) ? 0 : val+1;
     }
 
     if (val != gpuData[5])
@@ -136,8 +135,7 @@ int computeGold(int *gpuData, const int len)
 
     for (int i = 0; i < len; ++i)
     {
-        //val = ((val == 0) || (val > limit)) ? limit : val-1;
-        val = val-1;
+        val = ((val == 0) || (val > limit)) ? limit : val-1;
     }
 
     if (val != gpuData[6])
@@ -234,12 +232,10 @@ __global__ void testKernel(hipLaunchParm lp,int *g_odata)
     atomicMin(&g_odata[4], tid);
 
     // Atomic increment (modulo 17+1)
-    //atomicInc((unsigned int *)&g_odata[5], 17);
-    //atomicInc((unsigned int *)&g_odata[5]);
-
+    atomicInc((unsigned int *)&g_odata[5], 17);
+ 
     // Atomic decrement
-   // atomicDec((unsigned int *)&g_odata[6], 137);
-    //atomicDec((unsigned int *)&g_odata[6]);
+    atomicDec((unsigned int *)&g_odata[6], 137);
 
     // Atomic compare-and-swap
     atomicCAS(&g_odata[7], tid-1, tid);