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

Removed atomicInc and atomicDec support from HIP

Browse Source
This commit is contained in:
gargrahul
2016-02-10 04:29:55 +05:30
parent 4d4ca3ef3f
commit 51f46d9ddf
2 changed files with 16 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/docs/markdown/hip_kernel_language.md
+2 -2
View File
@@ -435,8 +435,8 @@ HIP supports the following atomic operations.
| int atomicMax(int* address, int val) | ✓ | ✓ |
| unsigned int atomicMax(unsigned int* address,unsigned int val) | ✓ | ✓ |
| unsigned long long int atomicMax(unsigned long long int* address,unsigned long long int val) | ✓ | ✓ |
| unsigned int atomicInc(unsigned int* address)| ✓ <br><sub>Takes one argument.</sub> | ✓ <br><sub> Wrapping increment,takes two arguments.</sub> |
| unsigned int atomicDec(unsigned int* address)| ✓ <br><sub>Takes one argument.</sub> | ✓ <br><sub> Wrapping decrement,takes two arguments.</sub> |
| unsigned int atomicInc(unsigned int* address)| ✗ | ✓ |
| unsigned int atomicDec(unsigned int* address)| ✗ | ✓ |
| int atomicCAS(int* address, int compare, int val) | ✓ | ✓ |
| unsigned int atomicCAS(unsigned int* address,unsigned int compare,unsigned int val) | ✓ | ✓ |
| unsigned long long int atomicCAS(unsigned long long int* address,unsigned long long int compare,unsigned long long int val) | ✓ | ✓ |
hipamd/include/hcc_detail/hip_runtime.h
+14 -26
View File
@@ -195,18 +195,6 @@ __device__ inline unsigned long long int atomicMax(unsigned long long int* addre
return (long long int)hc::atomic_fetch_max((uint64_t*)address,(uint64_t)val);
}
//atomicInc()
__device__ inline unsigned int atomicInc(unsigned int* address)
{
return hc::atomic_fetch_inc(address);
}
//atomicDec()
__device__ inline unsigned int atomicDec(unsigned int* address)
{
return hc::atomic_fetch_dec(address);
}
//atomicCAS()
__device__ inline int atomicCAS(int* address, int compare, int val)
{
@@ -318,17 +306,17 @@ __device__ inline unsigned int __ffsll(unsigned long long int input)
return hc::__lastbit_u32_u64( input)+1;
}
__device__ inline unsigned int __ffs(int input)
__device__ inline unsigned int __ffs(int input)
{
return hc::__lastbit_u32_s32( input)+1;
}
__device__ inline unsigned int __ffsll(long long int input)
__device__ inline unsigned int __ffsll(long long int input)
{
return hc::__lastbit_u32_s64( input)+1;
}
__device__ inline unsigned int __brev( unsigned int input)
__device__ inline unsigned int __brev( unsigned int input)
{
return hc::__bitrev_b32( input);
}
@@ -339,59 +327,59 @@ __device__ inline unsigned long long int __brevll( unsigned long long int input)
}
// warp vote function __all __any __ballot
__device__ inline int __all( int input)
__device__ inline int __all( int input)
{
return hc::__all( input);
}
__device__ inline int __any( int input)
__device__ inline int __any( int input)
{
if( hc::__any( input)!=0) return 1;
else return 0;
}
__device__ inline unsigned long long int __ballot( int input)
__device__ inline unsigned long long int __ballot( int input)
{
return hc::__ballot( input);
}
// warp shuffle functions
__device__ inline int __shfl(int input, int lane, int width)
__device__ inline int __shfl(int input, int lane, int width)
{
return hc::__shfl(input,lane,width);
}
__device__ inline int __shfl_up(int input, unsigned int lane_delta, int width)
__device__ inline int __shfl_up(int input, unsigned int lane_delta, int width)
{
return hc::__shfl_up(input,lane_delta,width);
}
__device__ inline int __shfl_down(int input, unsigned int lane_delta, int width)
__device__ inline int __shfl_down(int input, unsigned int lane_delta, int width)
{
return hc::__shfl_down(input,lane_delta,width);
}
__device__ inline int __shfl_xor(int input, int lane_mask, int width)
__device__ inline int __shfl_xor(int input, int lane_mask, int width)
{
return hc::__shfl_xor(input,lane_mask,width);
}
__device__ inline float __shfl(float input, int lane, int width)
__device__ inline float __shfl(float input, int lane, int width)
{
return hc::__shfl(input,lane,width);
}
__device__ inline float __shfl_up(float input, unsigned int lane_delta, int width)
__device__ inline float __shfl_up(float input, unsigned int lane_delta, int width)
{
return hc::__shfl_up(input,lane_delta,width);
}
__device__ inline float __shfl_down(float input, unsigned int lane_delta, int width)
__device__ inline float __shfl_down(float input, unsigned int lane_delta, int width)
{
return hc::__shfl_down(input,lane_delta,width);
}
__device__ inline float __shfl_xor(float input, int lane_mask, int width)
__device__ inline float __shfl_xor(float input, int lane_mask, int width)
{
return hc::__shfl_xor(input,lane_mask,width);
}