From 0ee021d444550d81ddc0d1ff476fec1a0e44e184 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sat, 28 Mar 2020 17:28:33 +0530 Subject: [PATCH] Remove address_space(1) typecast and use __ockl_atomic_add_noret_f32 (#1956) * Remove address_space(1) typecast for ockl_global_atomic_add_f32 * use __ockl_atomic_add_noret_f32 [ROCm/hip commit: cbc3d1713fbb036c3e1f4534e680103c08f30f81] --- projects/hip/include/hip/hcc_detail/device_library_decls.h | 2 +- projects/hip/include/hip/hcc_detail/hip_atomic.h | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/device_library_decls.h b/projects/hip/include/hip/hcc_detail/device_library_decls.h index 2c4e6929ef..2eadb86774 100644 --- a/projects/hip/include/hip/hcc_detail/device_library_decls.h +++ b/projects/hip/include/hip/hcc_detail/device_library_decls.h @@ -72,7 +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); +extern "C" __device__ void __ockl_atomic_add_noret_f32(float*, float); // Introduce local address space #define __local __attribute__((address_space(3))) diff --git a/projects/hip/include/hip/hcc_detail/hip_atomic.h b/projects/hip/include/hip/hcc_detail/hip_atomic.h index 7ccfa6b43e..d00ebcdabb 100644 --- a/projects/hip/include/hip/hcc_detail/hip_atomic.h +++ b/projects/hip/include/hip/hcc_detail/hip_atomic.h @@ -78,7 +78,7 @@ __device__ inline void atomicAddNoRet(float* address, float val) { - __ockl_global_atomic_add_f32((__attribute__((address_space(1))) float*)address, val); + __ockl_atomic_add_noret_f32(address, val); } __device__