From a2fbf0612995d360ed1216e9557e3700fe1cd1e9 Mon Sep 17 00:00:00 2001
From: gargrahul <rgarg@hcBLR-Fiji2.amd.com>
Date: Mon, 22 Feb 2016 16:21:52 +0530
Subject: [PATCH] Update for shared atomics support

---
 docs/markdown/hip_kernel_language.md |  1 -
 include/hcc_detail/hip_runtime.h     |  4 ++--
 src/hip_hcc.cpp                      | 12 ++++++------
 3 files changed, 8 insertions(+), 9 deletions(-)

diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md
index 9d20fe82a3..4d3e72ce65 100644
--- a/docs/markdown/hip_kernel_language.md
+++ b/docs/markdown/hip_kernel_language.md
@@ -475,7 +475,6 @@ HIP supports the following atomic operations.
 ### Caveats and Features Under-Development:
 
 - HIP enables atomic operations on 32-bit integers. Additionally, it supports an atomic float add. AMD hardware, however, implements the float add using a CAS loop, so this function may not perform efficiently.
-- hcc currently maps `__shared__` atomics to `__device__` atomics. Optimal support is under development.
 - wrapping increment and decrement are under development.
 
 ## Warp Cross-Lane Functions
diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h
index 8474f066df..75bc40aade 100644
--- a/include/hcc_detail/hip_runtime.h
+++ b/include/hcc_detail/hip_runtime.h
@@ -66,8 +66,8 @@ THE SOFTWARE.
     // 32-bit Atomics:
 #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__       (1)
 #define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__   (1)
-#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__       (0)
-#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__   (0)
+#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__       (1)
+#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__   (1)
 #define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__           (0)
 
 // 64-bit Atomics:
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 91c9c7ed55..ee208bcced 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -40,7 +40,7 @@ THE SOFTWARE.
 
 #define USE_PINNED_HOST (__hcc_workweek__ >= 1601)
 
-//#define USE_ASYNC_COPY 
+//#define USE_ASYNC_COPY
 
 #define INLINE static inline
 
@@ -360,11 +360,11 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop)
 
     prop->arch.hasGlobalInt32Atomics       = 1;
     prop->arch.hasGlobalFloatAtomicExch    = 1;
-    prop->arch.hasSharedInt32Atomics       = 0; // TODO-hcc-atomics
-    prop->arch.hasSharedFloatAtomicExch    = 0; // TODO-hcc-atomics
+    prop->arch.hasSharedInt32Atomics       = 1;
+    prop->arch.hasSharedFloatAtomicExch    = 1;
     prop->arch.hasFloatAtomicAdd           = 0;
     prop->arch.hasGlobalInt64Atomics       = 1;
-    prop->arch.hasSharedInt64Atomics       = 0; // TODO-hcc-atomics
+    prop->arch.hasSharedInt64Atomics       = 1;
     prop->arch.hasDoubles                  = 1; // TODO - true for Fiji.
     prop->arch.hasWarpVote                 = 1;
     prop->arch.hasWarpBallot               = 1;
@@ -476,7 +476,7 @@ INLINE ihipDevice_t *ihipGetTlsDefaultDevice()
 {
     // If this is invalid, the TLS state is corrupt.
     // This can fire if called before devices are initialized.
-    // TODO - consider replacing assert with error code 
+    // TODO - consider replacing assert with error code
     assert (ihipIsValidDevice(tls_defaultDevice));
 
     return &g_devices[tls_defaultDevice];
@@ -1428,7 +1428,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
     } else {
         e = hipErrorInvalidResourceHandle;
     }
-        
+
 
 #else
     // TODO-hsart - what synchronization does hsa_copy provide?