Update for shared atomics support

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

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:

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?