diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h
index ab1e1f9e97..664f425396 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h
+++ b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h
@@ -826,21 +826,21 @@ __device__
 inline
 static void __threadfence()
 {
-  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_device);
+    __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
 }
 
 __device__
 inline
 static void __threadfence_block()
 {
-  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_work_group);
+    __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
 }
 
 __device__
 inline
 static void __threadfence_system()
 {
-  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_all_svm_devices);
+    __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
 }
 
 // abort
@@ -913,14 +913,11 @@ void __assertfail()
 }
 #endif /* defined(_WIN32) || defined(_WIN64) */
 
-__device__
-inline
-static void __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope)
-{
+__device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags) {
     if (flags) {
-        __atomic_work_item_fence(flags, __memory_order_release, scope);
+        __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
         __builtin_amdgcn_s_barrier();
-        __atomic_work_item_fence(flags, __memory_order_acquire, scope);
+        __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
     } else {
         __builtin_amdgcn_s_barrier();
     }
@@ -930,7 +927,7 @@ __device__
 inline
 static void __barrier(int n)
 {
-  __work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group);
+  __work_group_barrier((__cl_mem_fence_flags)n);
 }
 
 __device__
diff --git a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h
index 8add4fa279..0222870590 100644
--- a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h
+++ b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h
@@ -128,26 +128,4 @@ __device__ inline static __local void* __to_local(unsigned x) { return (__local
 #define __CLK_LOCAL_MEM_FENCE    0x01
 typedef unsigned __cl_mem_fence_flags;
 
-typedef enum __memory_scope {
-  __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
-  __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
-  __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
-  __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
-  __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
-} __memory_scope;
-
-// enum values aligned with what clang uses in EmitAtomicExpr()
-typedef enum __memory_order
-{
-  __memory_order_relaxed = __ATOMIC_RELAXED,
-  __memory_order_acquire = __ATOMIC_ACQUIRE,
-  __memory_order_release = __ATOMIC_RELEASE,
-  __memory_order_acq_rel = __ATOMIC_ACQ_REL,
-  __memory_order_seq_cst = __ATOMIC_SEQ_CST
-} __memory_order;
-
-// Linked from hip.amdgcn.bc
-extern "C" __device__ void
-__atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope);
-
 #endif