Expose support for non-returning atomic FADD (#1909)

Change-Id: If5359488324477315a9bd4f308a75f606c065b39

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)))

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)