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)