diff --git a/OpenCL/inc_platform.cl b/OpenCL/inc_platform.cl
index f9b88a313..a7c5c7572 100644
--- a/OpenCL/inc_platform.cl
+++ b/OpenCL/inc_platform.cl
@@ -309,8 +309,9 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
   return out.v64;
 }
 
-#define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0)
-#define SYNC_THREADS() __syncthreads ()
+#define FIXED_THREAD_COUNT(n) __attribute__((amdgpu_flat_work_group_size (1, (n))))
+
+#define SYNC_THREADS() __builtin_amdgcn_s_barrier ()
 #endif
 
 #ifdef IS_OPENCL
diff --git a/OpenCL/m02500-pure.cl b/OpenCL/m02500-pure.cl
index 5bcd35d63..f6e20fa5b 100644
--- a/OpenCL/m02500-pure.cl
+++ b/OpenCL/m02500-pure.cl
@@ -775,11 +775,7 @@ KERNEL_FQ void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t)
     s_te4[i] = te4[i];
   }
 
-  #if defined IS_CUDA || defined IS_HIP
-  __syncthreads();
-  #else
   SYNC_THREADS ();
-  #endif
 
   #else
 
diff --git a/OpenCL/m22000-pure.cl b/OpenCL/m22000-pure.cl
index cfe645bc7..6117d2e31 100644
--- a/OpenCL/m22000-pure.cl
+++ b/OpenCL/m22000-pure.cl
@@ -797,11 +797,7 @@ KERNEL_FQ void m22000_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_t))
     s_te4[i] = te4[i];
   }
 
-  #if defined IS_CUDA || defined IS_HIP
-  __syncthreads();
-  #else
   SYNC_THREADS ();
-  #endif
 
   #else
 
diff --git a/OpenCL/m22001-pure.cl b/OpenCL/m22001-pure.cl
index 95431e08f..23015b863 100644
--- a/OpenCL/m22001-pure.cl
+++ b/OpenCL/m22001-pure.cl
@@ -610,11 +610,7 @@ KERNEL_FQ void m22001_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pmk_tmp_t, wpa_t))
     s_te4[i] = te4[i];
   }
 
-  #if defined IS_CUDA || defined IS_HIP
-  __syncthreads();
-  #else
   SYNC_THREADS ();
-  #endif
 
   #else