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 16131ab920..cf33122570 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
@@ -687,10 +687,18 @@ __device__ inline static void __threadfence_system() {
   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
 }
 __device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags) {
-  if (flags) {
+  if (flags == (__CLK_GLOBAL_MEM_FENCE | __CLK_LOCAL_MEM_FENCE)) {
     __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
     __builtin_amdgcn_s_barrier();
     __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
+  } else if (flags & (__CLK_GLOBAL_MEM_FENCE)) {
+    __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup", "global");
+    __builtin_amdgcn_s_barrier();
+    __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup", "global");
+  } else if (flags & (__CLK_LOCAL_MEM_FENCE)) {
+    __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup", "local");
+    __builtin_amdgcn_s_barrier();
+    __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup", "local");
   } else {
     __builtin_amdgcn_s_barrier();
   }
@@ -698,8 +706,12 @@ __device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags) {
 
 __device__ inline static void __barrier(int n) { __work_group_barrier((__cl_mem_fence_flags)n); }
 
-__device__ inline __attribute__((convergent)) void __syncthreads() {
-  __barrier(__CLK_LOCAL_MEM_FENCE);
+__device__
+inline
+__attribute__((convergent))
+void __syncthreads()
+{
+  __barrier(__CLK_GLOBAL_MEM_FENCE | __CLK_LOCAL_MEM_FENCE);
 }
 
 __device__ inline __attribute__((convergent)) int __syncthreads_count(int predicate) {
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 00a98f5750..33623f5881 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,7 +128,8 @@ __device__ inline static __local void* __to_local(unsigned x) { return (__local
 #endif  //__HIP_DEVICE_COMPILE__
 
 // Using hip.amdgcn.bc - sync threads
-#define __CLK_LOCAL_MEM_FENCE 0x01
+#define __CLK_LOCAL_MEM_FENCE  0x01
+#define __CLK_GLOBAL_MEM_FENCE 0x02
 typedef unsigned __cl_mem_fence_flags;
 
 #endif