From 381dc93b96d90955db42ad5af379503ea852bacb Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Wed, 11 May 2016 12:51:11 +0530
Subject: [PATCH] Support for Atomic inc and dec in HIP

Change-Id: I783e4917cece5cc379894f0d293382315fbfa8b0
---
 include/hcc_detail/hip_runtime.h   |  9 +++++++++
 src/device_util.cpp                | 12 ++++++++++++
 tests/src/hipSimpleAtomicsTest.cpp | 14 +++++---------
 3 files changed, 26 insertions(+), 9 deletions(-)

diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h
index 3288996dce..0d70eaa2a4 100644
--- a/include/hcc_detail/hip_runtime.h
+++ b/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);
diff --git a/src/device_util.cpp b/src/device_util.cpp
index 1177ce8d06..c01ad30ab5 100644
--- a/src/device_util.cpp
+++ b/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)
diff --git a/tests/src/hipSimpleAtomicsTest.cpp b/tests/src/hipSimpleAtomicsTest.cpp
index 1be32f6679..d02252c7ea 100644
--- a/tests/src/hipSimpleAtomicsTest.cpp
+++ b/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);