AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Expose support for non-returning atomic FADD

Browse Source 
Change-Id: If5359488324477315a9bd4f308a75f606c065b39


[ROCm/clr commit: 48f5524fdc]
This commit is contained in:
Maneesh Gupta
2020-03-04 14:17:07 +05:30
parent fbf9b7429d
commit 8def98b74e
2 changed files with 9 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h
+1
View File
@@ -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__((const)) int __ockl_multi_grid_is_valid(void);
extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(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 // Introduce local address space
#define __local __attribute__((address_space(3))) #define __local __attribute__((address_space(3)))
projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h
+8
View File
@@ -73,6 +73,14 @@ float atomicAdd(float* address, float val)
return __uint_as_float(r); 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__ __device__
inline inline
double atomicAdd(double* address, double val) double atomicAdd(double* address, double val)