From 8def98b74e4e4e3e1bcb418c9d2d40863fdf3d6b Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 4 Mar 2020 14:17:07 +0530 Subject: [PATCH] Expose support for non-returning atomic FADD Change-Id: If5359488324477315a9bd4f308a75f606c065b39 [ROCm/clr commit: 48f5524fdcca2e15f0475f5f93b5ee987f2c1371] --- .../hipamd/include/hip/hcc_detail/device_library_decls.h | 1 + projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h | 8 ++++++++ 2 files changed, 9 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h index ac35823cd2..182565ad61 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h @@ -72,6 +72,7 @@ extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank( extern "C" __device__ __attribute__((const)) int __ockl_multi_grid_is_valid(void); extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(void); +extern "C" __device__ void __ockl_global_atomic_add_f32(__attribute__((address_space(1))) float*, float); // Introduce local address space #define __local __attribute__((address_space(3))) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h index 263f639e96..7ccfa6b43e 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h @@ -73,6 +73,14 @@ float atomicAdd(float* address, float val) return __uint_as_float(r); } + +__device__ +inline +void atomicAddNoRet(float* address, float val) +{ + __ockl_global_atomic_add_f32((__attribute__((address_space(1))) float*)address, val); +} + __device__ inline double atomicAdd(double* address, double val)