diff --git a/hipamd/docs/markdown/hip_kernel_language.md b/hipamd/docs/markdown/hip_kernel_language.md
index e7a6baa1a9..582cb8788a 100644
--- a/hipamd/docs/markdown/hip_kernel_language.md
+++ b/hipamd/docs/markdown/hip_kernel_language.md
@@ -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) | ✓ | ✓ |
diff --git a/hipamd/include/hcc_detail/hip_runtime.h b/hipamd/include/hcc_detail/hip_runtime.h
index d7e4cba328..8474f066df 100644
--- a/hipamd/include/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hcc_detail/hip_runtime.h
@@ -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);
 }