diff --git a/OpenCL/m01100_a0-optimized.cl b/OpenCL/m01100_a0-optimized.cl
index bd0b99fe9..f31993bf2 100644
--- a/OpenCL/m01100_a0-optimized.cl
+++ b/OpenCL/m01100_a0-optimized.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_RULES ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -270,7 +270,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_RULES ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m01100_a1-optimized.cl b/OpenCL/m01100_a1-optimized.cl
index 64a69d4b1..a63c340d1 100644
--- a/OpenCL/m01100_a1-optimized.cl
+++ b/OpenCL/m01100_a1-optimized.cl
@@ -49,7 +49,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_BASIC ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -330,7 +330,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_BASIC ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m01100_a3-optimized.cl b/OpenCL/m01100_a3-optimized.cl
index 73d1b826d..57e255107 100644
--- a/OpenCL/m01100_a3-optimized.cl
+++ b/OpenCL/m01100_a3-optimized.cl
@@ -541,7 +541,7 @@ KERNEL_FQ void m01100_m04 (KERN_ATTR_VECTOR ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -595,7 +595,7 @@ KERNEL_FQ void m01100_m08 (KERN_ATTR_VECTOR ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -649,7 +649,7 @@ KERNEL_FQ void m01100_m16 (KERN_ATTR_VECTOR ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -703,7 +703,7 @@ KERNEL_FQ void m01100_s04 (KERN_ATTR_VECTOR ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -757,7 +757,7 @@ KERNEL_FQ void m01100_s08 (KERN_ATTR_VECTOR ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -811,7 +811,7 @@ KERNEL_FQ void m01100_s16 (KERN_ATTR_VECTOR ())
     s_salt_buf[0].salt_buf[10] = (16 + s_salt_buf[0].salt_len) * 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m01500_a0-pure.cl b/OpenCL/m01500_a0-pure.cl
index f3a4d5913..a1d94d749 100644
--- a/OpenCL/m01500_a0-pure.cl
+++ b/OpenCL/m01500_a0-pure.cl
@@ -519,7 +519,7 @@ KERNEL_FQ void m01500_mxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -603,7 +603,7 @@ KERNEL_FQ void m01500_sxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m01500_a1-pure.cl b/OpenCL/m01500_a1-pure.cl
index 88a06b749..81b0a22cb 100644
--- a/OpenCL/m01500_a1-pure.cl
+++ b/OpenCL/m01500_a1-pure.cl
@@ -517,7 +517,7 @@ KERNEL_FQ void m01500_mxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -680,7 +680,7 @@ KERNEL_FQ void m01500_sxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02500-pure.cl b/OpenCL/m02500-pure.cl
index 8066d3298..bb46357f5 100644
--- a/OpenCL/m02500-pure.cl
+++ b/OpenCL/m02500-pure.cl
@@ -682,7 +682,7 @@ KERNEL_FQ void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t)
   #ifdef IS_CUDA
   __syncthreads();
   #else
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
   #endif
 
   #else
diff --git a/OpenCL/m02501-pure.cl b/OpenCL/m02501-pure.cl
index 6da60615e..f51e8f2db 100644
--- a/OpenCL/m02501-pure.cl
+++ b/OpenCL/m02501-pure.cl
@@ -549,7 +549,7 @@ KERNEL_FQ void m02501_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pmk_tmp_t, wpa_eapol_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m02610_a0-optimized.cl b/OpenCL/m02610_a0-optimized.cl
index df8899488..6744a5026 100644
--- a/OpenCL/m02610_a0-optimized.cl
+++ b/OpenCL/m02610_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -336,7 +336,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02610_a0-pure.cl b/OpenCL/m02610_a0-pure.cl
index 5c209f995..092557282 100644
--- a/OpenCL/m02610_a0-pure.cl
+++ b/OpenCL/m02610_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -155,7 +155,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02610_a1-optimized.cl b/OpenCL/m02610_a1-optimized.cl
index ad323f53e..bca78bc86 100644
--- a/OpenCL/m02610_a1-optimized.cl
+++ b/OpenCL/m02610_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -393,7 +393,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02610_a1-pure.cl b/OpenCL/m02610_a1-pure.cl
index 46f3808bb..ba64a81e5 100644
--- a/OpenCL/m02610_a1-pure.cl
+++ b/OpenCL/m02610_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -151,7 +151,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02610_a3-optimized.cl b/OpenCL/m02610_a3-optimized.cl
index 73af945f8..482e34094 100644
--- a/OpenCL/m02610_a3-optimized.cl
+++ b/OpenCL/m02610_a3-optimized.cl
@@ -616,7 +616,7 @@ KERNEL_FQ void m02610_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -686,7 +686,7 @@ KERNEL_FQ void m02610_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -756,7 +756,7 @@ KERNEL_FQ void m02610_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -826,7 +826,7 @@ KERNEL_FQ void m02610_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -896,7 +896,7 @@ KERNEL_FQ void m02610_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -966,7 +966,7 @@ KERNEL_FQ void m02610_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02610_a3-pure.cl b/OpenCL/m02610_a3-pure.cl
index 842897b54..6f992a8a5 100644
--- a/OpenCL/m02610_a3-pure.cl
+++ b/OpenCL/m02610_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m02610_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -164,7 +164,7 @@ KERNEL_FQ void m02610_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02710_a0-optimized.cl b/OpenCL/m02710_a0-optimized.cl
index cfe83ae2d..74c7190f8 100644
--- a/OpenCL/m02710_a0-optimized.cl
+++ b/OpenCL/m02710_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -421,7 +421,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02710_a1-optimized.cl b/OpenCL/m02710_a1-optimized.cl
index 89927efc9..45595e2d6 100644
--- a/OpenCL/m02710_a1-optimized.cl
+++ b/OpenCL/m02710_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -478,7 +478,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02710_a3-optimized.cl b/OpenCL/m02710_a3-optimized.cl
index a02f67260..c37eb27e9 100644
--- a/OpenCL/m02710_a3-optimized.cl
+++ b/OpenCL/m02710_a3-optimized.cl
@@ -785,7 +785,7 @@ KERNEL_FQ void m02710_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -855,7 +855,7 @@ KERNEL_FQ void m02710_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -925,7 +925,7 @@ KERNEL_FQ void m02710_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -995,7 +995,7 @@ KERNEL_FQ void m02710_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1065,7 +1065,7 @@ KERNEL_FQ void m02710_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1135,7 +1135,7 @@ KERNEL_FQ void m02710_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02810_a0-optimized.cl b/OpenCL/m02810_a0-optimized.cl
index 25757ed87..3c5094284 100644
--- a/OpenCL/m02810_a0-optimized.cl
+++ b/OpenCL/m02810_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -420,7 +420,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02810_a0-pure.cl b/OpenCL/m02810_a0-pure.cl
index 5e22e1b3b..c101ba360 100644
--- a/OpenCL/m02810_a0-pure.cl
+++ b/OpenCL/m02810_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -178,7 +178,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02810_a1-optimized.cl b/OpenCL/m02810_a1-optimized.cl
index 081e75270..d173d4832 100644
--- a/OpenCL/m02810_a1-optimized.cl
+++ b/OpenCL/m02810_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -477,7 +477,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02810_a1-pure.cl b/OpenCL/m02810_a1-pure.cl
index 187fe0dfc..31cb8de8c 100644
--- a/OpenCL/m02810_a1-pure.cl
+++ b/OpenCL/m02810_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -174,7 +174,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02810_a3-optimized.cl b/OpenCL/m02810_a3-optimized.cl
index 1f80d06f3..0fb3dd780 100644
--- a/OpenCL/m02810_a3-optimized.cl
+++ b/OpenCL/m02810_a3-optimized.cl
@@ -783,7 +783,7 @@ KERNEL_FQ void m02810_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -853,7 +853,7 @@ KERNEL_FQ void m02810_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -923,7 +923,7 @@ KERNEL_FQ void m02810_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -993,7 +993,7 @@ KERNEL_FQ void m02810_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1063,7 +1063,7 @@ KERNEL_FQ void m02810_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1133,7 +1133,7 @@ KERNEL_FQ void m02810_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m02810_a3-pure.cl b/OpenCL/m02810_a3-pure.cl
index 0af147c6b..903954989 100644
--- a/OpenCL/m02810_a3-pure.cl
+++ b/OpenCL/m02810_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m02810_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -187,7 +187,7 @@ KERNEL_FQ void m02810_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03000_a0-pure.cl b/OpenCL/m03000_a0-pure.cl
index 531d294e9..c35d938de 100644
--- a/OpenCL/m03000_a0-pure.cl
+++ b/OpenCL/m03000_a0-pure.cl
@@ -529,7 +529,7 @@ KERNEL_FQ void m03000_mxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -614,7 +614,7 @@ KERNEL_FQ void m03000_sxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03000_a1-pure.cl b/OpenCL/m03000_a1-pure.cl
index 0e00a37ef..ffadf5214 100644
--- a/OpenCL/m03000_a1-pure.cl
+++ b/OpenCL/m03000_a1-pure.cl
@@ -527,7 +527,7 @@ KERNEL_FQ void m03000_mxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -691,7 +691,7 @@ KERNEL_FQ void m03000_sxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03100_a0-optimized.cl b/OpenCL/m03100_a0-optimized.cl
index 426594c63..d37a1db37 100644
--- a/OpenCL/m03100_a0-optimized.cl
+++ b/OpenCL/m03100_a0-optimized.cl
@@ -56,7 +56,7 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -288,7 +288,7 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m03100_a1-optimized.cl b/OpenCL/m03100_a1-optimized.cl
index d565e1530..601c1c62b 100644
--- a/OpenCL/m03100_a1-optimized.cl
+++ b/OpenCL/m03100_a1-optimized.cl
@@ -54,7 +54,7 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -346,7 +346,7 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m03100_a3-optimized.cl b/OpenCL/m03100_a3-optimized.cl
index 41fe2f1f5..fb8618271 100644
--- a/OpenCL/m03100_a3-optimized.cl
+++ b/OpenCL/m03100_a3-optimized.cl
@@ -452,7 +452,7 @@ KERNEL_FQ void m03100_m04 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -535,7 +535,7 @@ KERNEL_FQ void m03100_m08 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -622,7 +622,7 @@ KERNEL_FQ void m03100_s04 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -705,7 +705,7 @@ KERNEL_FQ void m03100_s08 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m03710_a0-optimized.cl b/OpenCL/m03710_a0-optimized.cl
index 13f534ec4..1956e8eaf 100644
--- a/OpenCL/m03710_a0-optimized.cl
+++ b/OpenCL/m03710_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -367,7 +367,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03710_a0-pure.cl b/OpenCL/m03710_a0-pure.cl
index 600e2bfbd..a5ac06721 100644
--- a/OpenCL/m03710_a0-pure.cl
+++ b/OpenCL/m03710_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -168,7 +168,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03710_a1-optimized.cl b/OpenCL/m03710_a1-optimized.cl
index 6a09adc8d..885e12019 100644
--- a/OpenCL/m03710_a1-optimized.cl
+++ b/OpenCL/m03710_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -424,7 +424,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03710_a1-pure.cl b/OpenCL/m03710_a1-pure.cl
index 334a09bdb..48804bdc0 100644
--- a/OpenCL/m03710_a1-pure.cl
+++ b/OpenCL/m03710_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -164,7 +164,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03710_a3-optimized.cl b/OpenCL/m03710_a3-optimized.cl
index 6feff2427..6a92fe84a 100644
--- a/OpenCL/m03710_a3-optimized.cl
+++ b/OpenCL/m03710_a3-optimized.cl
@@ -643,7 +643,7 @@ KERNEL_FQ void m03710_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -747,7 +747,7 @@ KERNEL_FQ void m03710_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -817,7 +817,7 @@ KERNEL_FQ void m03710_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -887,7 +887,7 @@ KERNEL_FQ void m03710_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -957,7 +957,7 @@ KERNEL_FQ void m03710_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -993,7 +993,7 @@ KERNEL_FQ void m03710_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03710_a3-pure.cl b/OpenCL/m03710_a3-pure.cl
index b9c61f98f..0583ba6e0 100644
--- a/OpenCL/m03710_a3-pure.cl
+++ b/OpenCL/m03710_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m03710_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -177,7 +177,7 @@ KERNEL_FQ void m03710_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03910_a0-optimized.cl b/OpenCL/m03910_a0-optimized.cl
index 38346dea8..38586a366 100644
--- a/OpenCL/m03910_a0-optimized.cl
+++ b/OpenCL/m03910_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -420,7 +420,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03910_a0-pure.cl b/OpenCL/m03910_a0-pure.cl
index 9cebaca44..79fdb2265 100644
--- a/OpenCL/m03910_a0-pure.cl
+++ b/OpenCL/m03910_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -178,7 +178,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03910_a1-optimized.cl b/OpenCL/m03910_a1-optimized.cl
index d27fbdd63..c4b2f5cb2 100644
--- a/OpenCL/m03910_a1-optimized.cl
+++ b/OpenCL/m03910_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -477,7 +477,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03910_a1-pure.cl b/OpenCL/m03910_a1-pure.cl
index 3c3706a99..a27aa1fa4 100644
--- a/OpenCL/m03910_a1-pure.cl
+++ b/OpenCL/m03910_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -174,7 +174,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03910_a3-optimized.cl b/OpenCL/m03910_a3-optimized.cl
index aeacff0e7..3134cd506 100644
--- a/OpenCL/m03910_a3-optimized.cl
+++ b/OpenCL/m03910_a3-optimized.cl
@@ -783,7 +783,7 @@ KERNEL_FQ void m03910_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -853,7 +853,7 @@ KERNEL_FQ void m03910_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -923,7 +923,7 @@ KERNEL_FQ void m03910_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -993,7 +993,7 @@ KERNEL_FQ void m03910_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1063,7 +1063,7 @@ KERNEL_FQ void m03910_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1133,7 +1133,7 @@ KERNEL_FQ void m03910_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m03910_a3-pure.cl b/OpenCL/m03910_a3-pure.cl
index b4c269e4c..73698ba4e 100644
--- a/OpenCL/m03910_a3-pure.cl
+++ b/OpenCL/m03910_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m03910_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -187,7 +187,7 @@ KERNEL_FQ void m03910_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04010_a0-optimized.cl b/OpenCL/m04010_a0-optimized.cl
index efe2e8074..8f675941a 100644
--- a/OpenCL/m04010_a0-optimized.cl
+++ b/OpenCL/m04010_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -393,7 +393,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04010_a0-pure.cl b/OpenCL/m04010_a0-pure.cl
index 0eab0a63b..6557df376 100644
--- a/OpenCL/m04010_a0-pure.cl
+++ b/OpenCL/m04010_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -159,7 +159,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04010_a1-optimized.cl b/OpenCL/m04010_a1-optimized.cl
index 2d9d5ac11..22f5e9830 100644
--- a/OpenCL/m04010_a1-optimized.cl
+++ b/OpenCL/m04010_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -449,7 +449,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04010_a1-pure.cl b/OpenCL/m04010_a1-pure.cl
index 66d5c9f55..935273eb0 100644
--- a/OpenCL/m04010_a1-pure.cl
+++ b/OpenCL/m04010_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -155,7 +155,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04010_a3-optimized.cl b/OpenCL/m04010_a3-optimized.cl
index 34d678766..4d69d9c25 100644
--- a/OpenCL/m04010_a3-optimized.cl
+++ b/OpenCL/m04010_a3-optimized.cl
@@ -683,7 +683,7 @@ KERNEL_FQ void m04010_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -787,7 +787,7 @@ KERNEL_FQ void m04010_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -857,7 +857,7 @@ KERNEL_FQ void m04010_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -927,7 +927,7 @@ KERNEL_FQ void m04010_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -997,7 +997,7 @@ KERNEL_FQ void m04010_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1033,7 +1033,7 @@ KERNEL_FQ void m04010_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04010_a3-pure.cl b/OpenCL/m04010_a3-pure.cl
index 2f6882a17..3bc48a738 100644
--- a/OpenCL/m04010_a3-pure.cl
+++ b/OpenCL/m04010_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04010_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -172,7 +172,7 @@ KERNEL_FQ void m04010_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04110_a0-optimized.cl b/OpenCL/m04110_a0-optimized.cl
index 77d12bc03..972d0113f 100644
--- a/OpenCL/m04110_a0-optimized.cl
+++ b/OpenCL/m04110_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -438,7 +438,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04110_a0-pure.cl b/OpenCL/m04110_a0-pure.cl
index 8347442e3..2426e5a16 100644
--- a/OpenCL/m04110_a0-pure.cl
+++ b/OpenCL/m04110_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -172,7 +172,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04110_a1-optimized.cl b/OpenCL/m04110_a1-optimized.cl
index 36bb2b2db..af3cb9274 100644
--- a/OpenCL/m04110_a1-optimized.cl
+++ b/OpenCL/m04110_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -496,7 +496,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04110_a1-pure.cl b/OpenCL/m04110_a1-pure.cl
index 9d9273f2a..dd230f3fe 100644
--- a/OpenCL/m04110_a1-pure.cl
+++ b/OpenCL/m04110_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -168,7 +168,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04110_a3-optimized.cl b/OpenCL/m04110_a3-optimized.cl
index c9dd1fcc5..3521ecb6d 100644
--- a/OpenCL/m04110_a3-optimized.cl
+++ b/OpenCL/m04110_a3-optimized.cl
@@ -739,7 +739,7 @@ KERNEL_FQ void m04110_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -843,7 +843,7 @@ KERNEL_FQ void m04110_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -913,7 +913,7 @@ KERNEL_FQ void m04110_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -983,7 +983,7 @@ KERNEL_FQ void m04110_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1053,7 +1053,7 @@ KERNEL_FQ void m04110_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1089,7 +1089,7 @@ KERNEL_FQ void m04110_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04110_a3-pure.cl b/OpenCL/m04110_a3-pure.cl
index 4c74ccbd2..e7c4bfbd0 100644
--- a/OpenCL/m04110_a3-pure.cl
+++ b/OpenCL/m04110_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04110_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -183,7 +183,7 @@ KERNEL_FQ void m04110_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04310_a0-optimized.cl b/OpenCL/m04310_a0-optimized.cl
index dc693f19c..33bf2d335 100644
--- a/OpenCL/m04310_a0-optimized.cl
+++ b/OpenCL/m04310_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -336,7 +336,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04310_a0-pure.cl b/OpenCL/m04310_a0-pure.cl
index 0a0230892..cec409c1d 100644
--- a/OpenCL/m04310_a0-pure.cl
+++ b/OpenCL/m04310_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -155,7 +155,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04310_a1-optimized.cl b/OpenCL/m04310_a1-optimized.cl
index cc0b41619..bbd0c9508 100644
--- a/OpenCL/m04310_a1-optimized.cl
+++ b/OpenCL/m04310_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -393,7 +393,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04310_a1-pure.cl b/OpenCL/m04310_a1-pure.cl
index 2ce1d1fa7..b0acaf470 100644
--- a/OpenCL/m04310_a1-pure.cl
+++ b/OpenCL/m04310_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -151,7 +151,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04310_a3-optimized.cl b/OpenCL/m04310_a3-optimized.cl
index 159bf0f4a..6ecf6b1b3 100644
--- a/OpenCL/m04310_a3-optimized.cl
+++ b/OpenCL/m04310_a3-optimized.cl
@@ -616,7 +616,7 @@ KERNEL_FQ void m04310_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -686,7 +686,7 @@ KERNEL_FQ void m04310_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -756,7 +756,7 @@ KERNEL_FQ void m04310_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -826,7 +826,7 @@ KERNEL_FQ void m04310_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -896,7 +896,7 @@ KERNEL_FQ void m04310_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -966,7 +966,7 @@ KERNEL_FQ void m04310_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04310_a3-pure.cl b/OpenCL/m04310_a3-pure.cl
index f291aecb6..d40944557 100644
--- a/OpenCL/m04310_a3-pure.cl
+++ b/OpenCL/m04310_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04310_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -164,7 +164,7 @@ KERNEL_FQ void m04310_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04400_a0-optimized.cl b/OpenCL/m04400_a0-optimized.cl
index ffad3ae09..05e04b25a 100644
--- a/OpenCL/m04400_a0-optimized.cl
+++ b/OpenCL/m04400_a0-optimized.cl
@@ -53,7 +53,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -363,7 +363,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04400_a0-pure.cl b/OpenCL/m04400_a0-pure.cl
index 72718568e..dc63016e7 100644
--- a/OpenCL/m04400_a0-pure.cl
+++ b/OpenCL/m04400_a0-pure.cl
@@ -53,7 +53,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -150,7 +150,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04400_a1-optimized.cl b/OpenCL/m04400_a1-optimized.cl
index 795178874..698ef9d86 100644
--- a/OpenCL/m04400_a1-optimized.cl
+++ b/OpenCL/m04400_a1-optimized.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -419,7 +419,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04400_a1-pure.cl b/OpenCL/m04400_a1-pure.cl
index 9404b95fc..73df68e7b 100644
--- a/OpenCL/m04400_a1-pure.cl
+++ b/OpenCL/m04400_a1-pure.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -146,7 +146,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04400_a3-optimized.cl b/OpenCL/m04400_a3-optimized.cl
index 33224fc61..220db208f 100644
--- a/OpenCL/m04400_a3-optimized.cl
+++ b/OpenCL/m04400_a3-optimized.cl
@@ -588,7 +588,7 @@ KERNEL_FQ void m04400_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -658,7 +658,7 @@ KERNEL_FQ void m04400_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -728,7 +728,7 @@ KERNEL_FQ void m04400_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -798,7 +798,7 @@ KERNEL_FQ void m04400_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -868,7 +868,7 @@ KERNEL_FQ void m04400_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -938,7 +938,7 @@ KERNEL_FQ void m04400_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04400_a3-pure.cl b/OpenCL/m04400_a3-pure.cl
index 213a956ef..5cdee2e1a 100644
--- a/OpenCL/m04400_a3-pure.cl
+++ b/OpenCL/m04400_a3-pure.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m04400_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -159,7 +159,7 @@ KERNEL_FQ void m04400_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04500_a0-optimized.cl b/OpenCL/m04500_a0-optimized.cl
index 573e50223..aa6785d5c 100644
--- a/OpenCL/m04500_a0-optimized.cl
+++ b/OpenCL/m04500_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -389,7 +389,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04500_a0-pure.cl b/OpenCL/m04500_a0-pure.cl
index 1dfa8b061..b8a539961 100644
--- a/OpenCL/m04500_a0-pure.cl
+++ b/OpenCL/m04500_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -149,7 +149,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04500_a1-optimized.cl b/OpenCL/m04500_a1-optimized.cl
index 05d593d27..e9e0c42f9 100644
--- a/OpenCL/m04500_a1-optimized.cl
+++ b/OpenCL/m04500_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -445,7 +445,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04500_a1-pure.cl b/OpenCL/m04500_a1-pure.cl
index 32db391dd..eead59928 100644
--- a/OpenCL/m04500_a1-pure.cl
+++ b/OpenCL/m04500_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -145,7 +145,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04500_a3-optimized.cl b/OpenCL/m04500_a3-optimized.cl
index 1be7a4321..a0f9f632f 100644
--- a/OpenCL/m04500_a3-optimized.cl
+++ b/OpenCL/m04500_a3-optimized.cl
@@ -647,7 +647,7 @@ KERNEL_FQ void m04500_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -717,7 +717,7 @@ KERNEL_FQ void m04500_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -787,7 +787,7 @@ KERNEL_FQ void m04500_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -857,7 +857,7 @@ KERNEL_FQ void m04500_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -927,7 +927,7 @@ KERNEL_FQ void m04500_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -997,7 +997,7 @@ KERNEL_FQ void m04500_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04500_a3-pure.cl b/OpenCL/m04500_a3-pure.cl
index 086e72121..82173d7a4 100644
--- a/OpenCL/m04500_a3-pure.cl
+++ b/OpenCL/m04500_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04500_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -158,7 +158,7 @@ KERNEL_FQ void m04500_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04520_a0-optimized.cl b/OpenCL/m04520_a0-optimized.cl
index a34bca23c..237702cc0 100644
--- a/OpenCL/m04520_a0-optimized.cl
+++ b/OpenCL/m04520_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -620,7 +620,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04520_a0-pure.cl b/OpenCL/m04520_a0-pure.cl
index b53826e38..5cd66b395 100644
--- a/OpenCL/m04520_a0-pure.cl
+++ b/OpenCL/m04520_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -164,7 +164,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04520_a1-optimized.cl b/OpenCL/m04520_a1-optimized.cl
index 955343e5c..2c6cf6a41 100644
--- a/OpenCL/m04520_a1-optimized.cl
+++ b/OpenCL/m04520_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -676,7 +676,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04520_a1-pure.cl b/OpenCL/m04520_a1-pure.cl
index 54179adab..4a34801fc 100644
--- a/OpenCL/m04520_a1-pure.cl
+++ b/OpenCL/m04520_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -160,7 +160,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04520_a3-optimized.cl b/OpenCL/m04520_a3-optimized.cl
index c3d1a96d3..52e229952 100644
--- a/OpenCL/m04520_a3-optimized.cl
+++ b/OpenCL/m04520_a3-optimized.cl
@@ -1100,7 +1100,7 @@ KERNEL_FQ void m04520_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1170,7 +1170,7 @@ KERNEL_FQ void m04520_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1240,7 +1240,7 @@ KERNEL_FQ void m04520_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1310,7 +1310,7 @@ KERNEL_FQ void m04520_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1380,7 +1380,7 @@ KERNEL_FQ void m04520_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1450,7 +1450,7 @@ KERNEL_FQ void m04520_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04520_a3-pure.cl b/OpenCL/m04520_a3-pure.cl
index c1e49123c..8ced12d82 100644
--- a/OpenCL/m04520_a3-pure.cl
+++ b/OpenCL/m04520_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m04520_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -175,7 +175,7 @@ KERNEL_FQ void m04520_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04700_a0-optimized.cl b/OpenCL/m04700_a0-optimized.cl
index 13a2b98c8..19ffea0a4 100644
--- a/OpenCL/m04700_a0-optimized.cl
+++ b/OpenCL/m04700_a0-optimized.cl
@@ -53,7 +53,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -346,7 +346,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04700_a0-pure.cl b/OpenCL/m04700_a0-pure.cl
index 056a4bab3..3645fa853 100644
--- a/OpenCL/m04700_a0-pure.cl
+++ b/OpenCL/m04700_a0-pure.cl
@@ -53,7 +53,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -145,7 +145,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04700_a1-optimized.cl b/OpenCL/m04700_a1-optimized.cl
index 5f590df8c..982e57043 100644
--- a/OpenCL/m04700_a1-optimized.cl
+++ b/OpenCL/m04700_a1-optimized.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -399,7 +399,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04700_a1-pure.cl b/OpenCL/m04700_a1-pure.cl
index 357bf7a78..ddc4286ed 100644
--- a/OpenCL/m04700_a1-pure.cl
+++ b/OpenCL/m04700_a1-pure.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -141,7 +141,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04700_a3-optimized.cl b/OpenCL/m04700_a3-optimized.cl
index d1bf72589..e2c1ab949 100644
--- a/OpenCL/m04700_a3-optimized.cl
+++ b/OpenCL/m04700_a3-optimized.cl
@@ -588,7 +588,7 @@ KERNEL_FQ void m04700_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -658,7 +658,7 @@ KERNEL_FQ void m04700_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -728,7 +728,7 @@ KERNEL_FQ void m04700_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -798,7 +798,7 @@ KERNEL_FQ void m04700_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -868,7 +868,7 @@ KERNEL_FQ void m04700_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -938,7 +938,7 @@ KERNEL_FQ void m04700_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m04700_a3-pure.cl b/OpenCL/m04700_a3-pure.cl
index 29ca21e72..26a7524be 100644
--- a/OpenCL/m04700_a3-pure.cl
+++ b/OpenCL/m04700_a3-pure.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m04700_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -154,7 +154,7 @@ KERNEL_FQ void m04700_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05300_a0-optimized.cl b/OpenCL/m05300_a0-optimized.cl
index 11a175095..f20d0d097 100644
--- a/OpenCL/m05300_a0-optimized.cl
+++ b/OpenCL/m05300_a0-optimized.cl
@@ -138,7 +138,7 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -313,7 +313,7 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05300_a1-optimized.cl b/OpenCL/m05300_a1-optimized.cl
index 290f60a35..fc2f75c04 100644
--- a/OpenCL/m05300_a1-optimized.cl
+++ b/OpenCL/m05300_a1-optimized.cl
@@ -136,7 +136,7 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -371,7 +371,7 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05300_a3-optimized.cl b/OpenCL/m05300_a3-optimized.cl
index c228edc09..1f961d197 100644
--- a/OpenCL/m05300_a3-optimized.cl
+++ b/OpenCL/m05300_a3-optimized.cl
@@ -442,7 +442,7 @@ KERNEL_FQ void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -515,7 +515,7 @@ KERNEL_FQ void m05300_m08 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -588,7 +588,7 @@ KERNEL_FQ void m05300_m16 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -661,7 +661,7 @@ KERNEL_FQ void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -734,7 +734,7 @@ KERNEL_FQ void m05300_s08 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -807,7 +807,7 @@ KERNEL_FQ void m05300_s16 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05400_a0-optimized.cl b/OpenCL/m05400_a0-optimized.cl
index 06cfe4d95..c2332db9e 100644
--- a/OpenCL/m05400_a0-optimized.cl
+++ b/OpenCL/m05400_a0-optimized.cl
@@ -142,7 +142,7 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -326,7 +326,7 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05400_a1-optimized.cl b/OpenCL/m05400_a1-optimized.cl
index 576f280d4..9ca6763e1 100644
--- a/OpenCL/m05400_a1-optimized.cl
+++ b/OpenCL/m05400_a1-optimized.cl
@@ -140,7 +140,7 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -392,7 +392,7 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05400_a3-optimized.cl b/OpenCL/m05400_a3-optimized.cl
index 21d723531..6a1b47117 100644
--- a/OpenCL/m05400_a3-optimized.cl
+++ b/OpenCL/m05400_a3-optimized.cl
@@ -446,7 +446,7 @@ KERNEL_FQ void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -519,7 +519,7 @@ KERNEL_FQ void m05400_m08 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -592,7 +592,7 @@ KERNEL_FQ void m05400_m16 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -665,7 +665,7 @@ KERNEL_FQ void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -738,7 +738,7 @@ KERNEL_FQ void m05400_s08 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -811,7 +811,7 @@ KERNEL_FQ void m05400_s16 (KERN_ATTR_ESALT (ikepsk_t))
     s_msg_buf[i] = hc_swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05500_a0-optimized.cl b/OpenCL/m05500_a0-optimized.cl
index 7db6100e7..7a3a58037 100644
--- a/OpenCL/m05500_a0-optimized.cl
+++ b/OpenCL/m05500_a0-optimized.cl
@@ -539,7 +539,7 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -752,7 +752,7 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05500_a0-pure.cl b/OpenCL/m05500_a0-pure.cl
index f83be3388..f13503698 100644
--- a/OpenCL/m05500_a0-pure.cl
+++ b/OpenCL/m05500_a0-pure.cl
@@ -539,7 +539,7 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -663,7 +663,7 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05500_a1-optimized.cl b/OpenCL/m05500_a1-optimized.cl
index 20de5bc7e..bd80ae13e 100644
--- a/OpenCL/m05500_a1-optimized.cl
+++ b/OpenCL/m05500_a1-optimized.cl
@@ -537,7 +537,7 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -803,7 +803,7 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05500_a1-pure.cl b/OpenCL/m05500_a1-pure.cl
index 14e152681..577117f3a 100644
--- a/OpenCL/m05500_a1-pure.cl
+++ b/OpenCL/m05500_a1-pure.cl
@@ -537,7 +537,7 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -659,7 +659,7 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05500_a3-optimized.cl b/OpenCL/m05500_a3-optimized.cl
index 5f43cd563..b1a34c2c0 100644
--- a/OpenCL/m05500_a3-optimized.cl
+++ b/OpenCL/m05500_a3-optimized.cl
@@ -870,7 +870,7 @@ KERNEL_FQ void m05500_m04 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -944,7 +944,7 @@ KERNEL_FQ void m05500_m08 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1018,7 +1018,7 @@ KERNEL_FQ void m05500_m16 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1092,7 +1092,7 @@ KERNEL_FQ void m05500_s04 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1166,7 +1166,7 @@ KERNEL_FQ void m05500_s08 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1240,7 +1240,7 @@ KERNEL_FQ void m05500_s16 (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05500_a3-pure.cl b/OpenCL/m05500_a3-pure.cl
index c4aa11718..698003a40 100644
--- a/OpenCL/m05500_a3-pure.cl
+++ b/OpenCL/m05500_a3-pure.cl
@@ -537,7 +537,7 @@ KERNEL_FQ void m05500_mxx (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -672,7 +672,7 @@ KERNEL_FQ void m05500_sxx (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05600_a0-optimized.cl b/OpenCL/m05600_a0-optimized.cl
index a1cb6f71d..1ee105135 100644
--- a/OpenCL/m05600_a0-optimized.cl
+++ b/OpenCL/m05600_a0-optimized.cl
@@ -141,7 +141,7 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_RULES_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -380,7 +380,7 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_RULES_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05600_a1-optimized.cl b/OpenCL/m05600_a1-optimized.cl
index 16af9bf75..c8d7bdf11 100644
--- a/OpenCL/m05600_a1-optimized.cl
+++ b/OpenCL/m05600_a1-optimized.cl
@@ -139,7 +139,7 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -436,7 +436,7 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05600_a3-optimized.cl b/OpenCL/m05600_a3-optimized.cl
index ba42483e9..646e07455 100644
--- a/OpenCL/m05600_a3-optimized.cl
+++ b/OpenCL/m05600_a3-optimized.cl
@@ -563,7 +563,7 @@ KERNEL_FQ void m05600_m04 (KERN_ATTR_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -636,7 +636,7 @@ KERNEL_FQ void m05600_m08 (KERN_ATTR_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -709,7 +709,7 @@ KERNEL_FQ void m05600_m16 (KERN_ATTR_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -782,7 +782,7 @@ KERNEL_FQ void m05600_s04 (KERN_ATTR_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -855,7 +855,7 @@ KERNEL_FQ void m05600_s08 (KERN_ATTR_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -928,7 +928,7 @@ KERNEL_FQ void m05600_s16 (KERN_ATTR_ESALT (netntlm_t))
     s_chall_buf[i] = esalt_bufs[digests_offset].chall_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05800-optimized.cl b/OpenCL/m05800-optimized.cl
index 339376659..bedf6625e 100644
--- a/OpenCL/m05800-optimized.cl
+++ b/OpenCL/m05800-optimized.cl
@@ -2308,7 +2308,7 @@ KERNEL_FQ void m05800_loop (KERN_ATTR_TMPS (androidpin_tmp_t))
     s_pc_len[i] = c_pc_len[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m05800-pure.cl b/OpenCL/m05800-pure.cl
index a397ae119..465e47176 100644
--- a/OpenCL/m05800-pure.cl
+++ b/OpenCL/m05800-pure.cl
@@ -2128,7 +2128,7 @@ KERNEL_FQ void m05800_loop (KERN_ATTR_TMPS (androidpin_tmp_t))
     s_pc_len[i] = c_pc_len[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m06100_a0-optimized.cl b/OpenCL/m06100_a0-optimized.cl
index 046d0c4f9..00243058c 100644
--- a/OpenCL/m06100_a0-optimized.cl
+++ b/OpenCL/m06100_a0-optimized.cl
@@ -60,7 +60,7 @@ KERNEL_FQ void m06100_m04 (KERN_ATTR_RULES ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -200,7 +200,7 @@ KERNEL_FQ void m06100_s04 (KERN_ATTR_RULES ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06100_a0-pure.cl b/OpenCL/m06100_a0-pure.cl
index deca7a8ba..4482625ca 100644
--- a/OpenCL/m06100_a0-pure.cl
+++ b/OpenCL/m06100_a0-pure.cl
@@ -55,7 +55,7 @@ KERNEL_FQ void m06100_mxx (KERN_ATTR_RULES ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -139,7 +139,7 @@ KERNEL_FQ void m06100_sxx (KERN_ATTR_RULES ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06100_a1-optimized.cl b/OpenCL/m06100_a1-optimized.cl
index f81455869..e6cb58a41 100644
--- a/OpenCL/m06100_a1-optimized.cl
+++ b/OpenCL/m06100_a1-optimized.cl
@@ -58,7 +58,7 @@ KERNEL_FQ void m06100_m04 (KERN_ATTR_BASIC ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -256,7 +256,7 @@ KERNEL_FQ void m06100_s04 (KERN_ATTR_BASIC ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06100_a1-pure.cl b/OpenCL/m06100_a1-pure.cl
index 611927094..aea3847fb 100644
--- a/OpenCL/m06100_a1-pure.cl
+++ b/OpenCL/m06100_a1-pure.cl
@@ -53,7 +53,7 @@ KERNEL_FQ void m06100_mxx (KERN_ATTR_BASIC ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -135,7 +135,7 @@ KERNEL_FQ void m06100_sxx (KERN_ATTR_BASIC ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06100_a3-optimized.cl b/OpenCL/m06100_a3-optimized.cl
index da2d8745e..b9e10a98f 100644
--- a/OpenCL/m06100_a3-optimized.cl
+++ b/OpenCL/m06100_a3-optimized.cl
@@ -208,7 +208,7 @@ KERNEL_FQ void m06100_m04 (KERN_ATTR_BASIC ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -300,7 +300,7 @@ KERNEL_FQ void m06100_m08 (KERN_ATTR_BASIC ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -396,7 +396,7 @@ KERNEL_FQ void m06100_s04 (KERN_ATTR_BASIC ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -488,7 +488,7 @@ KERNEL_FQ void m06100_s08 (KERN_ATTR_BASIC ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06100_a3-pure.cl b/OpenCL/m06100_a3-pure.cl
index 3f2f65bfd..2f5bedfb1 100644
--- a/OpenCL/m06100_a3-pure.cl
+++ b/OpenCL/m06100_a3-pure.cl
@@ -53,7 +53,7 @@ KERNEL_FQ void m06100_mxx (KERN_ATTR_VECTOR ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -148,7 +148,7 @@ KERNEL_FQ void m06100_sxx (KERN_ATTR_VECTOR ())
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06211-pure.cl b/OpenCL/m06211-pure.cl
index 529049401..7e84978c8 100644
--- a/OpenCL/m06211-pure.cl
+++ b/OpenCL/m06211-pure.cl
@@ -99,7 +99,7 @@ KERNEL_FQ void m06211_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -331,7 +331,7 @@ KERNEL_FQ void m06211_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06212-pure.cl b/OpenCL/m06212-pure.cl
index f6228daea..a60a315ee 100644
--- a/OpenCL/m06212-pure.cl
+++ b/OpenCL/m06212-pure.cl
@@ -99,7 +99,7 @@ KERNEL_FQ void m06212_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -331,7 +331,7 @@ KERNEL_FQ void m06212_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06213-pure.cl b/OpenCL/m06213-pure.cl
index 1bd79a274..45b304de7 100644
--- a/OpenCL/m06213-pure.cl
+++ b/OpenCL/m06213-pure.cl
@@ -99,7 +99,7 @@ KERNEL_FQ void m06213_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -331,7 +331,7 @@ KERNEL_FQ void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06221-pure.cl b/OpenCL/m06221-pure.cl
index 75e6a556e..0e8dc1e5f 100644
--- a/OpenCL/m06221-pure.cl
+++ b/OpenCL/m06221-pure.cl
@@ -121,7 +121,7 @@ KERNEL_FQ void m06221_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -475,7 +475,7 @@ KERNEL_FQ void m06221_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06222-pure.cl b/OpenCL/m06222-pure.cl
index f39eba2a5..2af5c856f 100644
--- a/OpenCL/m06222-pure.cl
+++ b/OpenCL/m06222-pure.cl
@@ -121,7 +121,7 @@ KERNEL_FQ void m06222_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -475,7 +475,7 @@ KERNEL_FQ void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06223-pure.cl b/OpenCL/m06223-pure.cl
index 2c2468f59..3721447bc 100644
--- a/OpenCL/m06223-pure.cl
+++ b/OpenCL/m06223-pure.cl
@@ -121,7 +121,7 @@ KERNEL_FQ void m06223_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -475,7 +475,7 @@ KERNEL_FQ void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06231-pure.cl b/OpenCL/m06231-pure.cl
index 57fc6cdcc..57d84f57c 100644
--- a/OpenCL/m06231-pure.cl
+++ b/OpenCL/m06231-pure.cl
@@ -159,7 +159,7 @@ KERNEL_FQ void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   /**
    * Whirlpool shared
@@ -191,7 +191,7 @@ KERNEL_FQ void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -403,7 +403,7 @@ KERNEL_FQ void m06231_loop (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -609,7 +609,7 @@ KERNEL_FQ void m06231_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -657,7 +657,7 @@ KERNEL_FQ void m06231_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06232-pure.cl b/OpenCL/m06232-pure.cl
index 9634349d5..b369aa7cb 100644
--- a/OpenCL/m06232-pure.cl
+++ b/OpenCL/m06232-pure.cl
@@ -159,7 +159,7 @@ KERNEL_FQ void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   /**
    * Whirlpool shared
@@ -191,7 +191,7 @@ KERNEL_FQ void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -403,7 +403,7 @@ KERNEL_FQ void m06232_loop (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -609,7 +609,7 @@ KERNEL_FQ void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -657,7 +657,7 @@ KERNEL_FQ void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06233-pure.cl b/OpenCL/m06233-pure.cl
index 7e4066daa..a1b94ae89 100644
--- a/OpenCL/m06233-pure.cl
+++ b/OpenCL/m06233-pure.cl
@@ -159,7 +159,7 @@ KERNEL_FQ void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   /**
    * Whirlpool shared
@@ -191,7 +191,7 @@ KERNEL_FQ void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -403,7 +403,7 @@ KERNEL_FQ void m06233_loop (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -609,7 +609,7 @@ KERNEL_FQ void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -657,7 +657,7 @@ KERNEL_FQ void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06600-pure.cl b/OpenCL/m06600-pure.cl
index b84adff7b..245877c4e 100644
--- a/OpenCL/m06600-pure.cl
+++ b/OpenCL/m06600-pure.cl
@@ -267,7 +267,7 @@ KERNEL_FQ void m06600_comp (KERN_ATTR_TMPS (agilekey_tmp_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06800-pure.cl b/OpenCL/m06800-pure.cl
index 2bf55cd0d..5b78bfc50 100644
--- a/OpenCL/m06800-pure.cl
+++ b/OpenCL/m06800-pure.cl
@@ -298,7 +298,7 @@ KERNEL_FQ void m06800_comp (KERN_ATTR_TMPS (lastpass_tmp_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m06900_a0-optimized.cl b/OpenCL/m06900_a0-optimized.cl
index d699856c0..3274d0665 100644
--- a/OpenCL/m06900_a0-optimized.cl
+++ b/OpenCL/m06900_a0-optimized.cl
@@ -721,7 +721,7 @@ KERNEL_FQ void m06900_m04 (KERN_ATTR_RULES ())
     s_tables[3][i] = c_tables[3][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -937,7 +937,7 @@ KERNEL_FQ void m06900_s04 (KERN_ATTR_RULES ())
     s_tables[3][i] = c_tables[3][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m06900_a1-optimized.cl b/OpenCL/m06900_a1-optimized.cl
index a6ca9bb5d..156b85b2f 100644
--- a/OpenCL/m06900_a1-optimized.cl
+++ b/OpenCL/m06900_a1-optimized.cl
@@ -719,7 +719,7 @@ KERNEL_FQ void m06900_m04 (KERN_ATTR_BASIC ())
     s_tables[3][i] = c_tables[3][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -989,7 +989,7 @@ KERNEL_FQ void m06900_s04 (KERN_ATTR_BASIC ())
     s_tables[3][i] = c_tables[3][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m06900_a3-optimized.cl b/OpenCL/m06900_a3-optimized.cl
index 5d9996777..d68d73cc8 100644
--- a/OpenCL/m06900_a3-optimized.cl
+++ b/OpenCL/m06900_a3-optimized.cl
@@ -1079,7 +1079,7 @@ KERNEL_FQ void m06900_m04 (KERN_ATTR_BASIC ())
     s_tables[3][i] = c_tables[3][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1148,7 +1148,7 @@ KERNEL_FQ void m06900_m08 (KERN_ATTR_BASIC ())
     s_tables[3][i] = c_tables[3][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1221,7 +1221,7 @@ KERNEL_FQ void m06900_s04 (KERN_ATTR_BASIC ())
     s_tables[3][i] = c_tables[3][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1290,7 +1290,7 @@ KERNEL_FQ void m06900_s08 (KERN_ATTR_BASIC ())
     s_tables[3][i] = c_tables[3][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08000_a0-optimized.cl b/OpenCL/m08000_a0-optimized.cl
index c3f3e6edf..6057ef561 100644
--- a/OpenCL/m08000_a0-optimized.cl
+++ b/OpenCL/m08000_a0-optimized.cl
@@ -236,7 +236,7 @@ KERNEL_FQ void m08000_m04 (KERN_ATTR_RULES ())
     w_s2[i] = 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (lid == 0)
   {
@@ -264,7 +264,7 @@ KERNEL_FQ void m08000_m04 (KERN_ATTR_RULES ())
     }
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -407,7 +407,7 @@ KERNEL_FQ void m08000_s04 (KERN_ATTR_RULES ())
     w_s2[i] = 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (lid == 0)
   {
@@ -435,7 +435,7 @@ KERNEL_FQ void m08000_s04 (KERN_ATTR_RULES ())
     }
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08000_a1-optimized.cl b/OpenCL/m08000_a1-optimized.cl
index defb4c2f0..3f597c2f2 100644
--- a/OpenCL/m08000_a1-optimized.cl
+++ b/OpenCL/m08000_a1-optimized.cl
@@ -234,7 +234,7 @@ KERNEL_FQ void m08000_m04 (KERN_ATTR_BASIC ())
     w_s2[i] = 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (lid == 0)
   {
@@ -262,7 +262,7 @@ KERNEL_FQ void m08000_m04 (KERN_ATTR_BASIC ())
     }
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -459,7 +459,7 @@ KERNEL_FQ void m08000_s04 (KERN_ATTR_BASIC ())
     w_s2[i] = 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (lid == 0)
   {
@@ -487,7 +487,7 @@ KERNEL_FQ void m08000_s04 (KERN_ATTR_BASIC ())
     }
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08000_a3-optimized.cl b/OpenCL/m08000_a3-optimized.cl
index 1439bce9e..0209e9805 100644
--- a/OpenCL/m08000_a3-optimized.cl
+++ b/OpenCL/m08000_a3-optimized.cl
@@ -231,7 +231,7 @@ DECLSPEC void m08000m (LOCAL_AS u32 *w_s1, LOCAL_AS u32 *w_s2, u32 *w, const u32
     w_s2[i] = 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (lid == 0)
   {
@@ -259,7 +259,7 @@ DECLSPEC void m08000m (LOCAL_AS u32 *w_s1, LOCAL_AS u32 *w_s2, u32 *w, const u32
     }
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -353,7 +353,7 @@ DECLSPEC void m08000s (LOCAL_AS u32 *w_s1, LOCAL_AS u32 *w_s2, u32 *w, const u32
     w_s2[i] = 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (lid == 0)
   {
@@ -381,7 +381,7 @@ DECLSPEC void m08000s (LOCAL_AS u32 *w_s1, LOCAL_AS u32 *w_s2, u32 *w, const u32
     }
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08400_a0-optimized.cl b/OpenCL/m08400_a0-optimized.cl
index 21a973218..6ca791374 100644
--- a/OpenCL/m08400_a0-optimized.cl
+++ b/OpenCL/m08400_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m08400_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -306,7 +306,7 @@ KERNEL_FQ void m08400_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08400_a0-pure.cl b/OpenCL/m08400_a0-pure.cl
index eb3ac3cae..cd55c0d83 100644
--- a/OpenCL/m08400_a0-pure.cl
+++ b/OpenCL/m08400_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m08400_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -203,7 +203,7 @@ KERNEL_FQ void m08400_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08400_a1-optimized.cl b/OpenCL/m08400_a1-optimized.cl
index 77c9a78e9..37a2189cd 100644
--- a/OpenCL/m08400_a1-optimized.cl
+++ b/OpenCL/m08400_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m08400_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -362,7 +362,7 @@ KERNEL_FQ void m08400_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08400_a1-pure.cl b/OpenCL/m08400_a1-pure.cl
index 67e3cb552..5d994aab7 100644
--- a/OpenCL/m08400_a1-pure.cl
+++ b/OpenCL/m08400_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m08400_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -199,7 +199,7 @@ KERNEL_FQ void m08400_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08400_a3-optimized.cl b/OpenCL/m08400_a3-optimized.cl
index ba99799e7..23474987f 100644
--- a/OpenCL/m08400_a3-optimized.cl
+++ b/OpenCL/m08400_a3-optimized.cl
@@ -482,7 +482,7 @@ KERNEL_FQ void m08400_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -552,7 +552,7 @@ KERNEL_FQ void m08400_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -622,7 +622,7 @@ KERNEL_FQ void m08400_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -692,7 +692,7 @@ KERNEL_FQ void m08400_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -762,7 +762,7 @@ KERNEL_FQ void m08400_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -832,7 +832,7 @@ KERNEL_FQ void m08400_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08400_a3-pure.cl b/OpenCL/m08400_a3-pure.cl
index 47847863c..1d495ade1 100644
--- a/OpenCL/m08400_a3-pure.cl
+++ b/OpenCL/m08400_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m08400_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -216,7 +216,7 @@ KERNEL_FQ void m08400_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08500_a0-pure.cl b/OpenCL/m08500_a0-pure.cl
index 04cefa798..fcb84c9b0 100644
--- a/OpenCL/m08500_a0-pure.cl
+++ b/OpenCL/m08500_a0-pure.cl
@@ -559,7 +559,7 @@ KERNEL_FQ void m08500_mxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -657,7 +657,7 @@ KERNEL_FQ void m08500_sxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08500_a1-pure.cl b/OpenCL/m08500_a1-pure.cl
index d7b02244b..6eb3590b9 100644
--- a/OpenCL/m08500_a1-pure.cl
+++ b/OpenCL/m08500_a1-pure.cl
@@ -557,7 +557,7 @@ KERNEL_FQ void m08500_mxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -713,7 +713,7 @@ KERNEL_FQ void m08500_sxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08500_a3-pure.cl b/OpenCL/m08500_a3-pure.cl
index 8ae9721c5..d50e4174d 100644
--- a/OpenCL/m08500_a3-pure.cl
+++ b/OpenCL/m08500_a3-pure.cl
@@ -695,7 +695,7 @@ KERNEL_FQ void m08500_mxx (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -769,7 +769,7 @@ KERNEL_FQ void m08500_sxx (KERN_ATTR_VECTOR ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08600_a0-pure.cl b/OpenCL/m08600_a0-pure.cl
index c88b30787..621f90fc3 100644
--- a/OpenCL/m08600_a0-pure.cl
+++ b/OpenCL/m08600_a0-pure.cl
@@ -250,7 +250,7 @@ KERNEL_FQ void m08600_mxx (KERN_ATTR_RULES ())
     s_lotus_magic_table[i] = lotus_magic_table[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -314,7 +314,7 @@ KERNEL_FQ void m08600_sxx (KERN_ATTR_RULES ())
     s_lotus_magic_table[i] = lotus_magic_table[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08600_a1-pure.cl b/OpenCL/m08600_a1-pure.cl
index 5253d21dd..89ef057c8 100644
--- a/OpenCL/m08600_a1-pure.cl
+++ b/OpenCL/m08600_a1-pure.cl
@@ -248,7 +248,7 @@ KERNEL_FQ void m08600_mxx (KERN_ATTR_BASIC ())
     s_lotus_magic_table[i] = lotus_magic_table[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -372,7 +372,7 @@ KERNEL_FQ void m08600_sxx (KERN_ATTR_BASIC ())
     s_lotus_magic_table[i] = lotus_magic_table[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08600_a3-pure.cl b/OpenCL/m08600_a3-pure.cl
index cd596b5c0..3b579ac55 100644
--- a/OpenCL/m08600_a3-pure.cl
+++ b/OpenCL/m08600_a3-pure.cl
@@ -353,7 +353,7 @@ KERNEL_FQ void m08600_mxx (KERN_ATTR_VECTOR ())
     s_lotus_magic_table[i] = lotus_magic_table[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -410,7 +410,7 @@ KERNEL_FQ void m08600_sxx (KERN_ATTR_VECTOR ())
     s_lotus_magic_table[i] = lotus_magic_table[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08700_a0-optimized.cl b/OpenCL/m08700_a0-optimized.cl
index 6cb6a3b9f..bb18dbed7 100644
--- a/OpenCL/m08700_a0-optimized.cl
+++ b/OpenCL/m08700_a0-optimized.cl
@@ -298,7 +298,7 @@ KERNEL_FQ void m08700_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -485,7 +485,7 @@ KERNEL_FQ void m08700_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08700_a1-optimized.cl b/OpenCL/m08700_a1-optimized.cl
index 2041e3f5d..8f4c3e33a 100644
--- a/OpenCL/m08700_a1-optimized.cl
+++ b/OpenCL/m08700_a1-optimized.cl
@@ -296,7 +296,7 @@ KERNEL_FQ void m08700_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -543,7 +543,7 @@ KERNEL_FQ void m08700_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08700_a3-optimized.cl b/OpenCL/m08700_a3-optimized.cl
index 897bc16c5..bb05bf13f 100644
--- a/OpenCL/m08700_a3-optimized.cl
+++ b/OpenCL/m08700_a3-optimized.cl
@@ -575,7 +575,7 @@ KERNEL_FQ void m08700_m04 (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -643,7 +643,7 @@ KERNEL_FQ void m08700_m08 (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -711,7 +711,7 @@ KERNEL_FQ void m08700_m16 (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -779,7 +779,7 @@ KERNEL_FQ void m08700_s04 (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -847,7 +847,7 @@ KERNEL_FQ void m08700_s08 (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -915,7 +915,7 @@ KERNEL_FQ void m08700_s16 (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m08800-pure.cl b/OpenCL/m08800-pure.cl
index c8e5381e6..b4ef1c066 100644
--- a/OpenCL/m08800-pure.cl
+++ b/OpenCL/m08800-pure.cl
@@ -263,7 +263,7 @@ KERNEL_FQ void m08800_comp (KERN_ATTR_TMPS_ESALT (androidfde_tmp_t, androidfde_t
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m09100-pure.cl b/OpenCL/m09100-pure.cl
index a0a3cd7d4..997b611bd 100644
--- a/OpenCL/m09100-pure.cl
+++ b/OpenCL/m09100-pure.cl
@@ -426,7 +426,7 @@ KERNEL_FQ void m09100_init (KERN_ATTR_TMPS (lotus8_tmp_t))
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m09400-pure.cl b/OpenCL/m09400-pure.cl
index 99747349f..a59d38007 100644
--- a/OpenCL/m09400-pure.cl
+++ b/OpenCL/m09400-pure.cl
@@ -165,7 +165,7 @@ KERNEL_FQ void m09400_comp (KERN_ATTR_TMPS_ESALT (office2007_tmp_t, office2007_t
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m09500-pure.cl b/OpenCL/m09500-pure.cl
index ddcbc5cb9..72bae3d63 100644
--- a/OpenCL/m09500-pure.cl
+++ b/OpenCL/m09500-pure.cl
@@ -163,7 +163,7 @@ KERNEL_FQ void m09500_comp (KERN_ATTR_TMPS_ESALT (office2010_tmp_t, office2010_t
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m09600-pure.cl b/OpenCL/m09600-pure.cl
index c4d0332ee..edbe62eba 100644
--- a/OpenCL/m09600-pure.cl
+++ b/OpenCL/m09600-pure.cl
@@ -209,7 +209,7 @@ KERNEL_FQ void m09600_comp (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m10700-optimized.cl b/OpenCL/m10700-optimized.cl
index 290761db2..8093d2f7a 100644
--- a/OpenCL/m10700-optimized.cl
+++ b/OpenCL/m10700-optimized.cl
@@ -608,7 +608,7 @@ KERNEL_FQ void m10700_loop (KERN_ATTR_TMPS_ESALT (pdf17l8_tmp_t, pdf_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m10700-pure.cl b/OpenCL/m10700-pure.cl
index 6b9dfc686..284a8ca49 100644
--- a/OpenCL/m10700-pure.cl
+++ b/OpenCL/m10700-pure.cl
@@ -1209,7 +1209,7 @@ KERNEL_FQ void m10700_loop (KERN_ATTR_TMPS_ESALT (pdf17l8_tmp_t, pdf_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11100_a0-optimized.cl b/OpenCL/m11100_a0-optimized.cl
index 7b2319f41..7fe39fc4f 100644
--- a/OpenCL/m11100_a0-optimized.cl
+++ b/OpenCL/m11100_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m11100_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -386,7 +386,7 @@ KERNEL_FQ void m11100_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11100_a0-pure.cl b/OpenCL/m11100_a0-pure.cl
index 821f39c07..8e3ca5378 100644
--- a/OpenCL/m11100_a0-pure.cl
+++ b/OpenCL/m11100_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m11100_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -198,7 +198,7 @@ KERNEL_FQ void m11100_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11100_a1-optimized.cl b/OpenCL/m11100_a1-optimized.cl
index d378ff025..bd89d18d9 100644
--- a/OpenCL/m11100_a1-optimized.cl
+++ b/OpenCL/m11100_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m11100_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -444,7 +444,7 @@ KERNEL_FQ void m11100_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11100_a1-pure.cl b/OpenCL/m11100_a1-pure.cl
index 92bdbac02..a870a76f8 100644
--- a/OpenCL/m11100_a1-pure.cl
+++ b/OpenCL/m11100_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m11100_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -194,7 +194,7 @@ KERNEL_FQ void m11100_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11100_a3-optimized.cl b/OpenCL/m11100_a3-optimized.cl
index 76f4dab57..190917c48 100644
--- a/OpenCL/m11100_a3-optimized.cl
+++ b/OpenCL/m11100_a3-optimized.cl
@@ -675,7 +675,7 @@ KERNEL_FQ void m11100_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -745,7 +745,7 @@ KERNEL_FQ void m11100_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -815,7 +815,7 @@ KERNEL_FQ void m11100_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -885,7 +885,7 @@ KERNEL_FQ void m11100_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -955,7 +955,7 @@ KERNEL_FQ void m11100_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1025,7 +1025,7 @@ KERNEL_FQ void m11100_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11100_a3-pure.cl b/OpenCL/m11100_a3-pure.cl
index a9e172740..9b1ef9e5c 100644
--- a/OpenCL/m11100_a3-pure.cl
+++ b/OpenCL/m11100_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m11100_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -237,7 +237,7 @@ KERNEL_FQ void m11100_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11300-pure.cl b/OpenCL/m11300-pure.cl
index cf0b06e7b..8cf78d701 100644
--- a/OpenCL/m11300-pure.cl
+++ b/OpenCL/m11300-pure.cl
@@ -246,7 +246,7 @@ KERNEL_FQ void m11300_comp (KERN_ATTR_TMPS_ESALT (bitcoin_wallet_tmp_t, bitcoin_
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11400_a0-pure.cl b/OpenCL/m11400_a0-pure.cl
index 28d917d9e..b57c7c20b 100644
--- a/OpenCL/m11400_a0-pure.cl
+++ b/OpenCL/m11400_a0-pure.cl
@@ -62,7 +62,7 @@ KERNEL_FQ void m11400_mxx (KERN_ATTR_RULES_ESALT (sip_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -160,7 +160,7 @@ KERNEL_FQ void m11400_sxx (KERN_ATTR_RULES_ESALT (sip_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11400_a1-pure.cl b/OpenCL/m11400_a1-pure.cl
index 92ac1c8fd..b77777ca5 100644
--- a/OpenCL/m11400_a1-pure.cl
+++ b/OpenCL/m11400_a1-pure.cl
@@ -60,7 +60,7 @@ KERNEL_FQ void m11400_mxx (KERN_ATTR_ESALT (sip_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -154,7 +154,7 @@ KERNEL_FQ void m11400_sxx (KERN_ATTR_ESALT (sip_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11400_a3-pure.cl b/OpenCL/m11400_a3-pure.cl
index a98a3f6ce..6b814d216 100644
--- a/OpenCL/m11400_a3-pure.cl
+++ b/OpenCL/m11400_a3-pure.cl
@@ -60,7 +60,7 @@ KERNEL_FQ void m11400_mxx (KERN_ATTR_VECTOR_ESALT (sip_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -180,7 +180,7 @@ KERNEL_FQ void m11400_sxx (KERN_ATTR_VECTOR_ESALT (sip_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11700_a0-optimized.cl b/OpenCL/m11700_a0-optimized.cl
index e42cc57ac..3dc636ebe 100644
--- a/OpenCL/m11700_a0-optimized.cl
+++ b/OpenCL/m11700_a0-optimized.cl
@@ -114,7 +114,7 @@ KERNEL_FQ void m11700_m04 (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -272,7 +272,7 @@ KERNEL_FQ void m11700_s04 (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11700_a0-pure.cl b/OpenCL/m11700_a0-pure.cl
index d651caa92..20461ae49 100644
--- a/OpenCL/m11700_a0-pure.cl
+++ b/OpenCL/m11700_a0-pure.cl
@@ -45,7 +45,7 @@ KERNEL_FQ void m11700_mxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -118,7 +118,7 @@ KERNEL_FQ void m11700_sxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11700_a1-optimized.cl b/OpenCL/m11700_a1-optimized.cl
index 982e53128..846c25d36 100644
--- a/OpenCL/m11700_a1-optimized.cl
+++ b/OpenCL/m11700_a1-optimized.cl
@@ -112,7 +112,7 @@ KERNEL_FQ void m11700_m04 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -328,7 +328,7 @@ KERNEL_FQ void m11700_s04 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11700_a1-pure.cl b/OpenCL/m11700_a1-pure.cl
index 4046115f1..8d37a83ed 100644
--- a/OpenCL/m11700_a1-pure.cl
+++ b/OpenCL/m11700_a1-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11700_mxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -114,7 +114,7 @@ KERNEL_FQ void m11700_sxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11700_a3-optimized.cl b/OpenCL/m11700_a3-optimized.cl
index 7a13ed762..45baeb97a 100644
--- a/OpenCL/m11700_a3-optimized.cl
+++ b/OpenCL/m11700_a3-optimized.cl
@@ -281,7 +281,7 @@ KERNEL_FQ void m11700_m04 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -345,7 +345,7 @@ KERNEL_FQ void m11700_m08 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -409,7 +409,7 @@ KERNEL_FQ void m11700_m16 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -473,7 +473,7 @@ KERNEL_FQ void m11700_s04 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -537,7 +537,7 @@ KERNEL_FQ void m11700_s08 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -601,7 +601,7 @@ KERNEL_FQ void m11700_s16 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11700_a3-pure.cl b/OpenCL/m11700_a3-pure.cl
index 46988ca46..6261703a8 100644
--- a/OpenCL/m11700_a3-pure.cl
+++ b/OpenCL/m11700_a3-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11700_mxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -127,7 +127,7 @@ KERNEL_FQ void m11700_sxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11750_a0-pure.cl b/OpenCL/m11750_a0-pure.cl
index 8f1a4f03e..6e4f071db 100644
--- a/OpenCL/m11750_a0-pure.cl
+++ b/OpenCL/m11750_a0-pure.cl
@@ -45,7 +45,7 @@ KERNEL_FQ void m11750_mxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -127,7 +127,7 @@ KERNEL_FQ void m11750_sxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11750_a1-pure.cl b/OpenCL/m11750_a1-pure.cl
index fef704bb5..98f023f6b 100644
--- a/OpenCL/m11750_a1-pure.cl
+++ b/OpenCL/m11750_a1-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11750_mxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -150,7 +150,7 @@ KERNEL_FQ void m11750_sxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11750_a3-pure.cl b/OpenCL/m11750_a3-pure.cl
index 389818a4f..d1e989c82 100644
--- a/OpenCL/m11750_a3-pure.cl
+++ b/OpenCL/m11750_a3-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11750_mxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -136,7 +136,7 @@ KERNEL_FQ void m11750_sxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11760_a0-pure.cl b/OpenCL/m11760_a0-pure.cl
index 23784cd05..787c21ac7 100644
--- a/OpenCL/m11760_a0-pure.cl
+++ b/OpenCL/m11760_a0-pure.cl
@@ -45,7 +45,7 @@ KERNEL_FQ void m11760_mxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -129,7 +129,7 @@ KERNEL_FQ void m11760_sxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11760_a1-pure.cl b/OpenCL/m11760_a1-pure.cl
index 1f2f7fa41..62bae3ee6 100644
--- a/OpenCL/m11760_a1-pure.cl
+++ b/OpenCL/m11760_a1-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11760_mxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -152,7 +152,7 @@ KERNEL_FQ void m11760_sxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11760_a3-pure.cl b/OpenCL/m11760_a3-pure.cl
index 7096f4616..0ec476ac7 100644
--- a/OpenCL/m11760_a3-pure.cl
+++ b/OpenCL/m11760_a3-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11760_mxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -138,7 +138,7 @@ KERNEL_FQ void m11760_sxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob256_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11800_a0-optimized.cl b/OpenCL/m11800_a0-optimized.cl
index 30b8ffc46..2c4518667 100644
--- a/OpenCL/m11800_a0-optimized.cl
+++ b/OpenCL/m11800_a0-optimized.cl
@@ -114,7 +114,7 @@ KERNEL_FQ void m11800_m04 (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -272,7 +272,7 @@ KERNEL_FQ void m11800_s04 (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11800_a0-pure.cl b/OpenCL/m11800_a0-pure.cl
index e21d9eae6..41315fc64 100644
--- a/OpenCL/m11800_a0-pure.cl
+++ b/OpenCL/m11800_a0-pure.cl
@@ -45,7 +45,7 @@ KERNEL_FQ void m11800_mxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -118,7 +118,7 @@ KERNEL_FQ void m11800_sxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11800_a1-optimized.cl b/OpenCL/m11800_a1-optimized.cl
index 34846963d..ba0843fa2 100644
--- a/OpenCL/m11800_a1-optimized.cl
+++ b/OpenCL/m11800_a1-optimized.cl
@@ -112,7 +112,7 @@ KERNEL_FQ void m11800_m04 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -328,7 +328,7 @@ KERNEL_FQ void m11800_s04 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11800_a1-pure.cl b/OpenCL/m11800_a1-pure.cl
index 52cb7bd02..8f05197ef 100644
--- a/OpenCL/m11800_a1-pure.cl
+++ b/OpenCL/m11800_a1-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11800_mxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -114,7 +114,7 @@ KERNEL_FQ void m11800_sxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11800_a3-optimized.cl b/OpenCL/m11800_a3-optimized.cl
index 32a2746ce..f9538252f 100644
--- a/OpenCL/m11800_a3-optimized.cl
+++ b/OpenCL/m11800_a3-optimized.cl
@@ -281,7 +281,7 @@ KERNEL_FQ void m11800_m04 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -345,7 +345,7 @@ KERNEL_FQ void m11800_m08 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -409,7 +409,7 @@ KERNEL_FQ void m11800_m16 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -473,7 +473,7 @@ KERNEL_FQ void m11800_s04 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -537,7 +537,7 @@ KERNEL_FQ void m11800_s08 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -601,7 +601,7 @@ KERNEL_FQ void m11800_s16 (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m11800_a3-pure.cl b/OpenCL/m11800_a3-pure.cl
index 12d1899c7..963d004db 100644
--- a/OpenCL/m11800_a3-pure.cl
+++ b/OpenCL/m11800_a3-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11800_mxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -127,7 +127,7 @@ KERNEL_FQ void m11800_sxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11850_a0-pure.cl b/OpenCL/m11850_a0-pure.cl
index 703ad9ac9..db98c9529 100644
--- a/OpenCL/m11850_a0-pure.cl
+++ b/OpenCL/m11850_a0-pure.cl
@@ -45,7 +45,7 @@ KERNEL_FQ void m11850_mxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -127,7 +127,7 @@ KERNEL_FQ void m11850_sxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11850_a1-pure.cl b/OpenCL/m11850_a1-pure.cl
index 6f8c93e1b..e7eac3084 100644
--- a/OpenCL/m11850_a1-pure.cl
+++ b/OpenCL/m11850_a1-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11850_mxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -150,7 +150,7 @@ KERNEL_FQ void m11850_sxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11850_a3-pure.cl b/OpenCL/m11850_a3-pure.cl
index b273a6ef4..91e2da87f 100644
--- a/OpenCL/m11850_a3-pure.cl
+++ b/OpenCL/m11850_a3-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11850_mxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -136,7 +136,7 @@ KERNEL_FQ void m11850_sxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11860_a0-pure.cl b/OpenCL/m11860_a0-pure.cl
index 0f20573ea..318c87fdf 100644
--- a/OpenCL/m11860_a0-pure.cl
+++ b/OpenCL/m11860_a0-pure.cl
@@ -45,7 +45,7 @@ KERNEL_FQ void m11860_mxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -129,7 +129,7 @@ KERNEL_FQ void m11860_sxx (KERN_ATTR_RULES ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11860_a1-pure.cl b/OpenCL/m11860_a1-pure.cl
index 272f8a61c..e7880e570 100644
--- a/OpenCL/m11860_a1-pure.cl
+++ b/OpenCL/m11860_a1-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11860_mxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -152,7 +152,7 @@ KERNEL_FQ void m11860_sxx (KERN_ATTR_BASIC ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m11860_a3-pure.cl b/OpenCL/m11860_a3-pure.cl
index 7c54204be..1d920283b 100644
--- a/OpenCL/m11860_a3-pure.cl
+++ b/OpenCL/m11860_a3-pure.cl
@@ -43,7 +43,7 @@ KERNEL_FQ void m11860_mxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -138,7 +138,7 @@ KERNEL_FQ void m11860_sxx (KERN_ATTR_VECTOR ())
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m12400-pure.cl b/OpenCL/m12400-pure.cl
index d0116b0ac..e7bc2e27e 100644
--- a/OpenCL/m12400-pure.cl
+++ b/OpenCL/m12400-pure.cl
@@ -534,7 +534,7 @@ KERNEL_FQ void m12400_init (KERN_ATTR_TMPS (bsdicrypt_tmp_t))
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -670,7 +670,7 @@ KERNEL_FQ void m12400_loop (KERN_ATTR_TMPS (bsdicrypt_tmp_t))
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m12500-pure.cl b/OpenCL/m12500-pure.cl
index 94d525cf8..99920ff71 100644
--- a/OpenCL/m12500-pure.cl
+++ b/OpenCL/m12500-pure.cl
@@ -318,7 +318,7 @@ KERNEL_FQ void m12500_comp (KERN_ATTR_TMPS_ESALT (rar3_tmp_t, pbkdf2_sha1_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m12600_a0-optimized.cl b/OpenCL/m12600_a0-optimized.cl
index 6106316ec..60a9461cd 100644
--- a/OpenCL/m12600_a0-optimized.cl
+++ b/OpenCL/m12600_a0-optimized.cl
@@ -53,7 +53,7 @@ KERNEL_FQ void m12600_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -393,7 +393,7 @@ KERNEL_FQ void m12600_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m12600_a0-pure.cl b/OpenCL/m12600_a0-pure.cl
index d2f8fb7b3..66ca371dc 100644
--- a/OpenCL/m12600_a0-pure.cl
+++ b/OpenCL/m12600_a0-pure.cl
@@ -53,7 +53,7 @@ KERNEL_FQ void m12600_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -189,7 +189,7 @@ KERNEL_FQ void m12600_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m12600_a1-optimized.cl b/OpenCL/m12600_a1-optimized.cl
index 728bc49a5..ea5805154 100644
--- a/OpenCL/m12600_a1-optimized.cl
+++ b/OpenCL/m12600_a1-optimized.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m12600_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -449,7 +449,7 @@ KERNEL_FQ void m12600_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m12600_a1-pure.cl b/OpenCL/m12600_a1-pure.cl
index 6518aaca0..3b2730f66 100644
--- a/OpenCL/m12600_a1-pure.cl
+++ b/OpenCL/m12600_a1-pure.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m12600_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -185,7 +185,7 @@ KERNEL_FQ void m12600_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m12600_a3-optimized.cl b/OpenCL/m12600_a3-optimized.cl
index e42a0754b..15ec957ac 100644
--- a/OpenCL/m12600_a3-optimized.cl
+++ b/OpenCL/m12600_a3-optimized.cl
@@ -648,7 +648,7 @@ KERNEL_FQ void m12600_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -718,7 +718,7 @@ KERNEL_FQ void m12600_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -788,7 +788,7 @@ KERNEL_FQ void m12600_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -858,7 +858,7 @@ KERNEL_FQ void m12600_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -928,7 +928,7 @@ KERNEL_FQ void m12600_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -998,7 +998,7 @@ KERNEL_FQ void m12600_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m12600_a3-pure.cl b/OpenCL/m12600_a3-pure.cl
index d0e80b99f..bf9cfb7f0 100644
--- a/OpenCL/m12600_a3-pure.cl
+++ b/OpenCL/m12600_a3-pure.cl
@@ -51,7 +51,7 @@ KERNEL_FQ void m12600_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -198,7 +198,7 @@ KERNEL_FQ void m12600_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m12700-pure.cl b/OpenCL/m12700-pure.cl
index 3d2638062..4a0c72d11 100644
--- a/OpenCL/m12700-pure.cl
+++ b/OpenCL/m12700-pure.cl
@@ -276,7 +276,7 @@ KERNEL_FQ void m12700_comp (KERN_ATTR_TMPS (mywallet_tmp_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m12800-pure.cl b/OpenCL/m12800-pure.cl
index 0c32d564f..52318e9ce 100644
--- a/OpenCL/m12800-pure.cl
+++ b/OpenCL/m12800-pure.cl
@@ -100,7 +100,7 @@ KERNEL_FQ void m12800_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sh
                  | ((i1 < 10) ? '0' + i1 : 'A' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13200-pure.cl b/OpenCL/m13200-pure.cl
index 41f912b0b..0ade009bd 100644
--- a/OpenCL/m13200-pure.cl
+++ b/OpenCL/m13200-pure.cl
@@ -112,7 +112,7 @@ KERNEL_FQ void m13200_loop (KERN_ATTR_TMPS (axcrypt_tmp_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13400-pure.cl b/OpenCL/m13400-pure.cl
index d77ecb680..5504ac07c 100644
--- a/OpenCL/m13400-pure.cl
+++ b/OpenCL/m13400-pure.cl
@@ -190,7 +190,7 @@ KERNEL_FQ void m13400_loop (KERN_ATTR_TMPS_ESALT (keepass_tmp_t, keepass_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -290,7 +290,7 @@ KERNEL_FQ void m13400_comp (KERN_ATTR_TMPS_ESALT (keepass_tmp_t, keepass_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13711-pure.cl b/OpenCL/m13711-pure.cl
index d3f29166e..e46f64018 100644
--- a/OpenCL/m13711-pure.cl
+++ b/OpenCL/m13711-pure.cl
@@ -140,7 +140,7 @@ KERNEL_FQ void m13711_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -287,7 +287,7 @@ KERNEL_FQ void m13711_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -471,7 +471,7 @@ KERNEL_FQ void m13711_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13712-pure.cl b/OpenCL/m13712-pure.cl
index 67b09a3e8..09214eb6c 100644
--- a/OpenCL/m13712-pure.cl
+++ b/OpenCL/m13712-pure.cl
@@ -191,7 +191,7 @@ KERNEL_FQ void m13712_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -338,7 +338,7 @@ KERNEL_FQ void m13712_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -523,7 +523,7 @@ KERNEL_FQ void m13712_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13713-pure.cl b/OpenCL/m13713-pure.cl
index d6ea2a0e6..18ec22e13 100644
--- a/OpenCL/m13713-pure.cl
+++ b/OpenCL/m13713-pure.cl
@@ -256,7 +256,7 @@ KERNEL_FQ void m13713_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -403,7 +403,7 @@ KERNEL_FQ void m13713_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -589,7 +589,7 @@ KERNEL_FQ void m13713_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13721-pure.cl b/OpenCL/m13721-pure.cl
index 55f3df252..1a58a3feb 100644
--- a/OpenCL/m13721-pure.cl
+++ b/OpenCL/m13721-pure.cl
@@ -162,7 +162,7 @@ KERNEL_FQ void m13721_init (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -390,7 +390,7 @@ KERNEL_FQ void m13721_loop (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -618,7 +618,7 @@ KERNEL_FQ void m13721_comp (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13722-pure.cl b/OpenCL/m13722-pure.cl
index 2e944cb50..ebb5d377a 100644
--- a/OpenCL/m13722-pure.cl
+++ b/OpenCL/m13722-pure.cl
@@ -213,7 +213,7 @@ KERNEL_FQ void m13722_init (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -441,7 +441,7 @@ KERNEL_FQ void m13722_loop (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -670,7 +670,7 @@ KERNEL_FQ void m13722_comp (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13723-pure.cl b/OpenCL/m13723-pure.cl
index c1dd74f80..ce6924637 100644
--- a/OpenCL/m13723-pure.cl
+++ b/OpenCL/m13723-pure.cl
@@ -278,7 +278,7 @@ KERNEL_FQ void m13723_init (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -506,7 +506,7 @@ KERNEL_FQ void m13723_loop (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -736,7 +736,7 @@ KERNEL_FQ void m13723_comp (KERN_ATTR_TMPS_ESALT (vc64_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13731-pure.cl b/OpenCL/m13731-pure.cl
index 24d9cd1de..33fcd4ad8 100644
--- a/OpenCL/m13731-pure.cl
+++ b/OpenCL/m13731-pure.cl
@@ -200,7 +200,7 @@ KERNEL_FQ void m13731_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   /**
    * Whirlpool shared
@@ -232,7 +232,7 @@ KERNEL_FQ void m13731_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -447,7 +447,7 @@ KERNEL_FQ void m13731_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -495,7 +495,7 @@ KERNEL_FQ void m13731_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -758,7 +758,7 @@ KERNEL_FQ void m13731_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -806,7 +806,7 @@ KERNEL_FQ void m13731_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13732-pure.cl b/OpenCL/m13732-pure.cl
index a5e7ee9b9..ac382dcb8 100644
--- a/OpenCL/m13732-pure.cl
+++ b/OpenCL/m13732-pure.cl
@@ -251,7 +251,7 @@ KERNEL_FQ void m13732_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   /**
    * Whirlpool shared
@@ -283,7 +283,7 @@ KERNEL_FQ void m13732_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -498,7 +498,7 @@ KERNEL_FQ void m13732_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -546,7 +546,7 @@ KERNEL_FQ void m13732_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -810,7 +810,7 @@ KERNEL_FQ void m13732_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -858,7 +858,7 @@ KERNEL_FQ void m13732_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13733-pure.cl b/OpenCL/m13733-pure.cl
index 38f8060e3..122fd306a 100644
--- a/OpenCL/m13733-pure.cl
+++ b/OpenCL/m13733-pure.cl
@@ -316,7 +316,7 @@ KERNEL_FQ void m13733_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   /**
    * Whirlpool shared
@@ -348,7 +348,7 @@ KERNEL_FQ void m13733_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -563,7 +563,7 @@ KERNEL_FQ void m13733_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -611,7 +611,7 @@ KERNEL_FQ void m13733_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -876,7 +876,7 @@ KERNEL_FQ void m13733_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -924,7 +924,7 @@ KERNEL_FQ void m13733_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_Cl[7][i] = Cl[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13751-pure.cl b/OpenCL/m13751-pure.cl
index 8ff4e0717..528b608b9 100644
--- a/OpenCL/m13751-pure.cl
+++ b/OpenCL/m13751-pure.cl
@@ -146,7 +146,7 @@ KERNEL_FQ void m13751_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -322,7 +322,7 @@ KERNEL_FQ void m13751_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -559,7 +559,7 @@ KERNEL_FQ void m13751_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13752-pure.cl b/OpenCL/m13752-pure.cl
index 135ead320..421eca8bb 100644
--- a/OpenCL/m13752-pure.cl
+++ b/OpenCL/m13752-pure.cl
@@ -197,7 +197,7 @@ KERNEL_FQ void m13752_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -373,7 +373,7 @@ KERNEL_FQ void m13752_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -582,7 +582,7 @@ KERNEL_FQ void m13752_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13753-pure.cl b/OpenCL/m13753-pure.cl
index 28c56ee1b..ce0021da5 100644
--- a/OpenCL/m13753-pure.cl
+++ b/OpenCL/m13753-pure.cl
@@ -262,7 +262,7 @@ KERNEL_FQ void m13753_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -438,7 +438,7 @@ KERNEL_FQ void m13753_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -648,7 +648,7 @@ KERNEL_FQ void m13753_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13771-pure.cl b/OpenCL/m13771-pure.cl
index b3b76052c..eeb767e63 100644
--- a/OpenCL/m13771-pure.cl
+++ b/OpenCL/m13771-pure.cl
@@ -188,7 +188,7 @@ KERNEL_FQ void m13771_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #ifdef REAL_SHM
 
@@ -206,7 +206,7 @@ KERNEL_FQ void m13771_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -420,7 +420,7 @@ KERNEL_FQ void m13771_loop (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -651,7 +651,7 @@ KERNEL_FQ void m13771_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13772-pure.cl b/OpenCL/m13772-pure.cl
index eb7864d99..80882acd3 100644
--- a/OpenCL/m13772-pure.cl
+++ b/OpenCL/m13772-pure.cl
@@ -239,7 +239,7 @@ KERNEL_FQ void m13772_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #ifdef REAL_SHM
 
@@ -257,7 +257,7 @@ KERNEL_FQ void m13772_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -471,7 +471,7 @@ KERNEL_FQ void m13772_loop (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -703,7 +703,7 @@ KERNEL_FQ void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13773-pure.cl b/OpenCL/m13773-pure.cl
index eea140305..beb700902 100644
--- a/OpenCL/m13773-pure.cl
+++ b/OpenCL/m13773-pure.cl
@@ -304,7 +304,7 @@ KERNEL_FQ void m13773_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #ifdef REAL_SHM
 
@@ -322,7 +322,7 @@ KERNEL_FQ void m13773_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -536,7 +536,7 @@ KERNEL_FQ void m13773_loop (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_sbob_sl64[7][i] = sbob512_sl64[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -769,7 +769,7 @@ KERNEL_FQ void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, vc_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m13800_a0-optimized.cl b/OpenCL/m13800_a0-optimized.cl
index c12745182..f2be74613 100644
--- a/OpenCL/m13800_a0-optimized.cl
+++ b/OpenCL/m13800_a0-optimized.cl
@@ -441,7 +441,7 @@ KERNEL_FQ void m13800_m04 (KERN_ATTR_RULES_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -637,7 +637,7 @@ KERNEL_FQ void m13800_s04 (KERN_ATTR_RULES_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13800_a1-optimized.cl b/OpenCL/m13800_a1-optimized.cl
index 3a432d316..75777a4c2 100644
--- a/OpenCL/m13800_a1-optimized.cl
+++ b/OpenCL/m13800_a1-optimized.cl
@@ -439,7 +439,7 @@ KERNEL_FQ void m13800_m04 (KERN_ATTR_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -691,7 +691,7 @@ KERNEL_FQ void m13800_s04 (KERN_ATTR_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13800_a3-optimized.cl b/OpenCL/m13800_a3-optimized.cl
index 46e422df4..4a30aec4a 100644
--- a/OpenCL/m13800_a3-optimized.cl
+++ b/OpenCL/m13800_a3-optimized.cl
@@ -753,7 +753,7 @@ KERNEL_FQ void m13800_m04 (KERN_ATTR_VECTOR_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -810,7 +810,7 @@ KERNEL_FQ void m13800_m08 (KERN_ATTR_VECTOR_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -867,7 +867,7 @@ KERNEL_FQ void m13800_m16 (KERN_ATTR_VECTOR_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -924,7 +924,7 @@ KERNEL_FQ void m13800_s04 (KERN_ATTR_VECTOR_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -981,7 +981,7 @@ KERNEL_FQ void m13800_s08 (KERN_ATTR_VECTOR_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1038,7 +1038,7 @@ KERNEL_FQ void m13800_s16 (KERN_ATTR_VECTOR_ESALT (win8phone_t))
     s_esalt[i] = esalt_bufs[digests_offset].salt_buf[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13900_a0-optimized.cl b/OpenCL/m13900_a0-optimized.cl
index 4dce764dc..72ce8024b 100644
--- a/OpenCL/m13900_a0-optimized.cl
+++ b/OpenCL/m13900_a0-optimized.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m13900_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -279,7 +279,7 @@ KERNEL_FQ void m13900_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13900_a0-pure.cl b/OpenCL/m13900_a0-pure.cl
index 81ac9fe94..dae69122d 100644
--- a/OpenCL/m13900_a0-pure.cl
+++ b/OpenCL/m13900_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m13900_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -203,7 +203,7 @@ KERNEL_FQ void m13900_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13900_a1-optimized.cl b/OpenCL/m13900_a1-optimized.cl
index 3290633d1..7a5214719 100644
--- a/OpenCL/m13900_a1-optimized.cl
+++ b/OpenCL/m13900_a1-optimized.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m13900_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -335,7 +335,7 @@ KERNEL_FQ void m13900_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13900_a1-pure.cl b/OpenCL/m13900_a1-pure.cl
index 92d5eb003..1b4c8c6ce 100644
--- a/OpenCL/m13900_a1-pure.cl
+++ b/OpenCL/m13900_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m13900_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -199,7 +199,7 @@ KERNEL_FQ void m13900_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13900_a3-optimized.cl b/OpenCL/m13900_a3-optimized.cl
index 10b4b93f7..e446de3db 100644
--- a/OpenCL/m13900_a3-optimized.cl
+++ b/OpenCL/m13900_a3-optimized.cl
@@ -439,7 +439,7 @@ KERNEL_FQ void m13900_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -509,7 +509,7 @@ KERNEL_FQ void m13900_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -579,7 +579,7 @@ KERNEL_FQ void m13900_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -649,7 +649,7 @@ KERNEL_FQ void m13900_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -719,7 +719,7 @@ KERNEL_FQ void m13900_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -789,7 +789,7 @@ KERNEL_FQ void m13900_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m13900_a3-pure.cl b/OpenCL/m13900_a3-pure.cl
index b0e85397d..7f8b5246d 100644
--- a/OpenCL/m13900_a3-pure.cl
+++ b/OpenCL/m13900_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m13900_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -216,7 +216,7 @@ KERNEL_FQ void m13900_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14000_a0-pure.cl b/OpenCL/m14000_a0-pure.cl
index 2be784b5a..d99c0eda2 100644
--- a/OpenCL/m14000_a0-pure.cl
+++ b/OpenCL/m14000_a0-pure.cl
@@ -534,7 +534,7 @@ KERNEL_FQ void m14000_mxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -639,7 +639,7 @@ KERNEL_FQ void m14000_sxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14000_a1-pure.cl b/OpenCL/m14000_a1-pure.cl
index 38367ff6a..4b85567a8 100644
--- a/OpenCL/m14000_a1-pure.cl
+++ b/OpenCL/m14000_a1-pure.cl
@@ -524,7 +524,7 @@ KERNEL_FQ void m14000_mxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -672,7 +672,7 @@ KERNEL_FQ void m14000_sxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14100_a0-pure.cl b/OpenCL/m14100_a0-pure.cl
index 99ffdb703..1efcaee4a 100644
--- a/OpenCL/m14100_a0-pure.cl
+++ b/OpenCL/m14100_a0-pure.cl
@@ -578,7 +578,7 @@ KERNEL_FQ void m14100_mxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -713,7 +713,7 @@ KERNEL_FQ void m14100_sxx (KERN_ATTR_RULES ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14100_a1-pure.cl b/OpenCL/m14100_a1-pure.cl
index aa861b19f..d1a9465c1 100644
--- a/OpenCL/m14100_a1-pure.cl
+++ b/OpenCL/m14100_a1-pure.cl
@@ -568,7 +568,7 @@ KERNEL_FQ void m14100_mxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -751,7 +751,7 @@ KERNEL_FQ void m14100_sxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14100_a3-pure.cl b/OpenCL/m14100_a3-pure.cl
index 1943413a8..4414be3e0 100644
--- a/OpenCL/m14100_a3-pure.cl
+++ b/OpenCL/m14100_a3-pure.cl
@@ -750,7 +750,7 @@ KERNEL_FQ void m14100_mxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -824,7 +824,7 @@ KERNEL_FQ void m14100_sxx (KERN_ATTR_BASIC ())
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14400_a0-optimized.cl b/OpenCL/m14400_a0-optimized.cl
index f6934b1d0..1cfb306b2 100644
--- a/OpenCL/m14400_a0-optimized.cl
+++ b/OpenCL/m14400_a0-optimized.cl
@@ -142,7 +142,7 @@ KERNEL_FQ void m14400_m04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -413,7 +413,7 @@ KERNEL_FQ void m14400_s04 (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14400_a0-pure.cl b/OpenCL/m14400_a0-pure.cl
index 9271e5ae0..fc7a9350a 100644
--- a/OpenCL/m14400_a0-pure.cl
+++ b/OpenCL/m14400_a0-pure.cl
@@ -52,7 +52,7 @@ KERNEL_FQ void m14400_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -299,7 +299,7 @@ KERNEL_FQ void m14400_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14400_a1-optimized.cl b/OpenCL/m14400_a1-optimized.cl
index a9259b61c..d35895338 100644
--- a/OpenCL/m14400_a1-optimized.cl
+++ b/OpenCL/m14400_a1-optimized.cl
@@ -142,7 +142,7 @@ KERNEL_FQ void m14400_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -477,7 +477,7 @@ KERNEL_FQ void m14400_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14400_a1-pure.cl b/OpenCL/m14400_a1-pure.cl
index 88bf46bad..c90002326 100644
--- a/OpenCL/m14400_a1-pure.cl
+++ b/OpenCL/m14400_a1-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m14400_mxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -295,7 +295,7 @@ KERNEL_FQ void m14400_sxx (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14400_a3-optimized.cl b/OpenCL/m14400_a3-optimized.cl
index 07938f501..6222f1214 100644
--- a/OpenCL/m14400_a3-optimized.cl
+++ b/OpenCL/m14400_a3-optimized.cl
@@ -652,7 +652,7 @@ KERNEL_FQ void m14400_m04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -722,7 +722,7 @@ KERNEL_FQ void m14400_m08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -792,7 +792,7 @@ KERNEL_FQ void m14400_m16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -862,7 +862,7 @@ KERNEL_FQ void m14400_s04 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -932,7 +932,7 @@ KERNEL_FQ void m14400_s08 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -1002,7 +1002,7 @@ KERNEL_FQ void m14400_s16 (KERN_ATTR_BASIC ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14400_a3-pure.cl b/OpenCL/m14400_a3-pure.cl
index 7976d7bc5..4dd93ffaf 100644
--- a/OpenCL/m14400_a3-pure.cl
+++ b/OpenCL/m14400_a3-pure.cl
@@ -50,7 +50,7 @@ KERNEL_FQ void m14400_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -320,7 +320,7 @@ KERNEL_FQ void m14400_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14611-pure.cl b/OpenCL/m14611-pure.cl
index 3acddaf75..1785ceb3a 100644
--- a/OpenCL/m14611-pure.cl
+++ b/OpenCL/m14611-pure.cl
@@ -330,7 +330,7 @@ KERNEL_FQ void m14611_comp (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m14621-pure.cl b/OpenCL/m14621-pure.cl
index 058c6e8a2..11ea28861 100644
--- a/OpenCL/m14621-pure.cl
+++ b/OpenCL/m14621-pure.cl
@@ -369,7 +369,7 @@ KERNEL_FQ void m14621_comp (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m14631-pure.cl b/OpenCL/m14631-pure.cl
index 99da8e3fb..2570b5a44 100644
--- a/OpenCL/m14631-pure.cl
+++ b/OpenCL/m14631-pure.cl
@@ -425,7 +425,7 @@ KERNEL_FQ void m14631_comp (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m14641-pure.cl b/OpenCL/m14641-pure.cl
index 4007b8a10..3deef6114 100644
--- a/OpenCL/m14641-pure.cl
+++ b/OpenCL/m14641-pure.cl
@@ -330,7 +330,7 @@ KERNEL_FQ void m14641_comp (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m14700-pure.cl b/OpenCL/m14700-pure.cl
index fb3673d67..df90ba694 100644
--- a/OpenCL/m14700-pure.cl
+++ b/OpenCL/m14700-pure.cl
@@ -266,7 +266,7 @@ KERNEL_FQ void m14700_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, itunes_back
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m14800-pure.cl b/OpenCL/m14800-pure.cl
index 578fe9716..bb5fb88fc 100644
--- a/OpenCL/m14800-pure.cl
+++ b/OpenCL/m14800-pure.cl
@@ -530,7 +530,7 @@ KERNEL_FQ void m14800_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, itunes_ba
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m14900_a0-optimized.cl b/OpenCL/m14900_a0-optimized.cl
index 36ca69d9c..5f015d28f 100644
--- a/OpenCL/m14900_a0-optimized.cl
+++ b/OpenCL/m14900_a0-optimized.cl
@@ -125,7 +125,7 @@ KERNEL_FQ void m14900_m04 (KERN_ATTR_RULES ())
     s_ftable[i] = c_ftable[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -216,7 +216,7 @@ KERNEL_FQ void m14900_s04 (KERN_ATTR_RULES ())
     s_ftable[i] = c_ftable[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14900_a1-optimized.cl b/OpenCL/m14900_a1-optimized.cl
index 42c7dd6c9..c6b740235 100644
--- a/OpenCL/m14900_a1-optimized.cl
+++ b/OpenCL/m14900_a1-optimized.cl
@@ -123,7 +123,7 @@ KERNEL_FQ void m14900_m04 (KERN_ATTR_BASIC ())
     s_ftable[i] = c_ftable[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -278,7 +278,7 @@ KERNEL_FQ void m14900_s04 (KERN_ATTR_BASIC ())
     s_ftable[i] = c_ftable[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m14900_a3-optimized.cl b/OpenCL/m14900_a3-optimized.cl
index b8a7ce4b1..cefd7cb98 100644
--- a/OpenCL/m14900_a3-optimized.cl
+++ b/OpenCL/m14900_a3-optimized.cl
@@ -231,7 +231,7 @@ KERNEL_FQ void m14900_m04 (KERN_ATTR_BASIC ())
     s_ftable[i] = c_ftable[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -305,7 +305,7 @@ KERNEL_FQ void m14900_s04 (KERN_ATTR_BASIC ())
     s_ftable[i] = c_ftable[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m15300-pure.cl b/OpenCL/m15300-pure.cl
index b1c8fb0b4..8d7bdf942 100644
--- a/OpenCL/m15300-pure.cl
+++ b/OpenCL/m15300-pure.cl
@@ -407,7 +407,7 @@ KERNEL_FQ void m15300_comp (KERN_ATTR_TMPS_ESALT (dpapimk_tmp_v1_t, dpapimk_t))
     s_skb[7][i] = c_skb[7][i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m15900-pure.cl b/OpenCL/m15900-pure.cl
index dbb59c80a..51fdfa0eb 100644
--- a/OpenCL/m15900-pure.cl
+++ b/OpenCL/m15900-pure.cl
@@ -539,7 +539,7 @@ KERNEL_FQ void m15900_comp (KERN_ATTR_TMPS_ESALT (dpapimk_tmp_v2_t, dpapimk_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m16000_a0-pure.cl b/OpenCL/m16000_a0-pure.cl
index ffbc3a6e2..cf02c5325 100644
--- a/OpenCL/m16000_a0-pure.cl
+++ b/OpenCL/m16000_a0-pure.cl
@@ -538,7 +538,7 @@ KERNEL_FQ void m16000_mxx (KERN_ATTR_RULES ())
     s_tripcode_salt[i] = c_tripcode_salt[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -631,7 +631,7 @@ KERNEL_FQ void m16000_sxx (KERN_ATTR_RULES ())
     s_tripcode_salt[i] = c_tripcode_salt[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m16000_a1-pure.cl b/OpenCL/m16000_a1-pure.cl
index c5af56fce..775be6c57 100644
--- a/OpenCL/m16000_a1-pure.cl
+++ b/OpenCL/m16000_a1-pure.cl
@@ -536,7 +536,7 @@ KERNEL_FQ void m16000_mxx (KERN_ATTR_BASIC ())
     s_tripcode_salt[i] = c_tripcode_salt[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -708,7 +708,7 @@ KERNEL_FQ void m16000_sxx (KERN_ATTR_BASIC ())
     s_tripcode_salt[i] = c_tripcode_salt[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m16000_a3-pure.cl b/OpenCL/m16000_a3-pure.cl
index c7cb2311f..23979dded 100644
--- a/OpenCL/m16000_a3-pure.cl
+++ b/OpenCL/m16000_a3-pure.cl
@@ -536,7 +536,7 @@ KERNEL_FQ void m16000_mxx (KERN_ATTR_VECTOR ())
     s_tripcode_salt[i] = c_tripcode_salt[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -657,7 +657,7 @@ KERNEL_FQ void m16000_sxx (KERN_ATTR_VECTOR ())
     s_tripcode_salt[i] = c_tripcode_salt[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m16200-pure.cl b/OpenCL/m16200-pure.cl
index ee1906bee..2b28662f7 100644
--- a/OpenCL/m16200-pure.cl
+++ b/OpenCL/m16200-pure.cl
@@ -307,7 +307,7 @@ KERNEL_FQ void m16200_comp (KERN_ATTR_TMPS_ESALT (apple_secure_notes_tmp_t, appl
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m16300-pure.cl b/OpenCL/m16300-pure.cl
index 908adb69f..6a6c2a6a8 100644
--- a/OpenCL/m16300-pure.cl
+++ b/OpenCL/m16300-pure.cl
@@ -441,7 +441,7 @@ KERNEL_FQ void m16300_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, ethereum_
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m16600_a0-optimized.cl b/OpenCL/m16600_a0-optimized.cl
index c8751feb5..4422d0e70 100644
--- a/OpenCL/m16600_a0-optimized.cl
+++ b/OpenCL/m16600_a0-optimized.cl
@@ -67,7 +67,7 @@ KERNEL_FQ void m16600_m04 (KERN_ATTR_RULES_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -452,7 +452,7 @@ KERNEL_FQ void m16600_s04 (KERN_ATTR_RULES_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m16600_a0-pure.cl b/OpenCL/m16600_a0-pure.cl
index 358543031..5c33c9fbd 100644
--- a/OpenCL/m16600_a0-pure.cl
+++ b/OpenCL/m16600_a0-pure.cl
@@ -67,7 +67,7 @@ KERNEL_FQ void m16600_mxx (KERN_ATTR_RULES_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -258,7 +258,7 @@ KERNEL_FQ void m16600_sxx (KERN_ATTR_RULES_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m16600_a1-optimized.cl b/OpenCL/m16600_a1-optimized.cl
index c9c8212b6..c5942f445 100644
--- a/OpenCL/m16600_a1-optimized.cl
+++ b/OpenCL/m16600_a1-optimized.cl
@@ -65,7 +65,7 @@ KERNEL_FQ void m16600_m04 (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -508,7 +508,7 @@ KERNEL_FQ void m16600_s04 (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m16600_a1-pure.cl b/OpenCL/m16600_a1-pure.cl
index a0be77baa..535916cd1 100644
--- a/OpenCL/m16600_a1-pure.cl
+++ b/OpenCL/m16600_a1-pure.cl
@@ -65,7 +65,7 @@ KERNEL_FQ void m16600_mxx (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -254,7 +254,7 @@ KERNEL_FQ void m16600_sxx (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m16600_a3-optimized.cl b/OpenCL/m16600_a3-optimized.cl
index 04a1a7138..94729cbf9 100644
--- a/OpenCL/m16600_a3-optimized.cl
+++ b/OpenCL/m16600_a3-optimized.cl
@@ -362,7 +362,7 @@ KERNEL_FQ void m16600_m04 (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -457,7 +457,7 @@ KERNEL_FQ void m16600_m08 (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -552,7 +552,7 @@ KERNEL_FQ void m16600_m16 (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -647,7 +647,7 @@ KERNEL_FQ void m16600_s04 (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -742,7 +742,7 @@ KERNEL_FQ void m16600_s08 (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -837,7 +837,7 @@ KERNEL_FQ void m16600_s16 (KERN_ATTR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m16600_a3-pure.cl b/OpenCL/m16600_a3-pure.cl
index 7a68e1ff2..a8e3775d8 100644
--- a/OpenCL/m16600_a3-pure.cl
+++ b/OpenCL/m16600_a3-pure.cl
@@ -65,7 +65,7 @@ KERNEL_FQ void m16600_mxx (KERN_ATTR_VECTOR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
@@ -267,7 +267,7 @@ KERNEL_FQ void m16600_sxx (KERN_ATTR_VECTOR_ESALT (electrum_wallet_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m18300-pure.cl b/OpenCL/m18300-pure.cl
index 6ae8fba68..9e1189d75 100644
--- a/OpenCL/m18300-pure.cl
+++ b/OpenCL/m18300-pure.cl
@@ -307,7 +307,7 @@ KERNEL_FQ void m18300_comp (KERN_ATTR_TMPS_ESALT (apple_secure_notes_tmp_t, appl
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m18400-pure.cl b/OpenCL/m18400-pure.cl
index f7abaea8d..3a4a80c27 100644
--- a/OpenCL/m18400-pure.cl
+++ b/OpenCL/m18400-pure.cl
@@ -303,7 +303,7 @@ KERNEL_FQ void m18400_comp (KERN_ATTR_TMPS_ESALT (odf12_tmp_t, odf12_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m18500_a0-pure.cl b/OpenCL/m18500_a0-pure.cl
index b4cf2faf6..7a68132b8 100644
--- a/OpenCL/m18500_a0-pure.cl
+++ b/OpenCL/m18500_a0-pure.cl
@@ -49,7 +49,7 @@ KERNEL_FQ void m18500_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -165,7 +165,7 @@ KERNEL_FQ void m18500_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m18500_a1-pure.cl b/OpenCL/m18500_a1-pure.cl
index cc48a7134..504d2577e 100644
--- a/OpenCL/m18500_a1-pure.cl
+++ b/OpenCL/m18500_a1-pure.cl
@@ -49,7 +49,7 @@ KERNEL_FQ void m18500_mxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -164,7 +164,7 @@ KERNEL_FQ void m18500_sxx (KERN_ATTR_RULES ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m18500_a3-pure.cl b/OpenCL/m18500_a3-pure.cl
index 46f8624f7..a34072250 100644
--- a/OpenCL/m18500_a3-pure.cl
+++ b/OpenCL/m18500_a3-pure.cl
@@ -47,7 +47,7 @@ KERNEL_FQ void m18500_mxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -174,7 +174,7 @@ KERNEL_FQ void m18500_sxx (KERN_ATTR_VECTOR ())
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m18900-pure.cl b/OpenCL/m18900-pure.cl
index 342a203a5..980d0bea9 100644
--- a/OpenCL/m18900-pure.cl
+++ b/OpenCL/m18900-pure.cl
@@ -271,7 +271,7 @@ KERNEL_FQ void m18900_comp (KERN_ATTR_TMPS_ESALT (android_backup_tmp_t, android_
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m19500_a0-pure.cl b/OpenCL/m19500_a0-pure.cl
index 9ec4dafb1..31f2b56db 100644
--- a/OpenCL/m19500_a0-pure.cl
+++ b/OpenCL/m19500_a0-pure.cl
@@ -62,7 +62,7 @@ KERNEL_FQ void m19500_mxx (KERN_ATTR_RULES_ESALT (devise_hash_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -197,7 +197,7 @@ KERNEL_FQ void m19500_sxx (KERN_ATTR_RULES_ESALT (devise_hash_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m19500_a1-pure.cl b/OpenCL/m19500_a1-pure.cl
index 9760a360a..00f6bc9e7 100644
--- a/OpenCL/m19500_a1-pure.cl
+++ b/OpenCL/m19500_a1-pure.cl
@@ -60,7 +60,7 @@ KERNEL_FQ void m19500_mxx (KERN_ATTR_ESALT (devise_hash_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -191,7 +191,7 @@ KERNEL_FQ void m19500_sxx (KERN_ATTR_ESALT (devise_hash_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m19500_a3-pure.cl b/OpenCL/m19500_a3-pure.cl
index 3900c5ff9..61dce2d47 100644
--- a/OpenCL/m19500_a3-pure.cl
+++ b/OpenCL/m19500_a3-pure.cl
@@ -60,7 +60,7 @@ KERNEL_FQ void m19500_mxx (KERN_ATTR_VECTOR_ESALT (devise_hash_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
@@ -206,7 +206,7 @@ KERNEL_FQ void m19500_sxx (KERN_ATTR_VECTOR_ESALT (devise_hash_t))
                  | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 8;
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   if (gid >= gid_max) return;
 
diff --git a/OpenCL/m19600-pure.cl b/OpenCL/m19600-pure.cl
index 265d7aee5..aa19771fe 100644
--- a/OpenCL/m19600-pure.cl
+++ b/OpenCL/m19600-pure.cl
@@ -323,7 +323,7 @@ KERNEL_FQ void m19600_comp (KERN_ATTR_TMPS_ESALT (krb5tgs_17_tmp_t, krb5tgs_17_t
     s_td4[i] = td4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m19700-pure.cl b/OpenCL/m19700-pure.cl
index df9b79e6b..efa5493ec 100644
--- a/OpenCL/m19700-pure.cl
+++ b/OpenCL/m19700-pure.cl
@@ -323,7 +323,7 @@ KERNEL_FQ void m19700_comp (KERN_ATTR_TMPS_ESALT (krb5tgs_18_tmp_t, krb5tgs_18_t
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m19800-pure.cl b/OpenCL/m19800-pure.cl
index a25f22da8..5c46c3edf 100644
--- a/OpenCL/m19800-pure.cl
+++ b/OpenCL/m19800-pure.cl
@@ -323,7 +323,7 @@ KERNEL_FQ void m19800_comp (KERN_ATTR_TMPS_ESALT (krb5pa_17_tmp_t, krb5pa_17_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m19900-pure.cl b/OpenCL/m19900-pure.cl
index 0d8cccc01..b5cfa8bfd 100644
--- a/OpenCL/m19900-pure.cl
+++ b/OpenCL/m19900-pure.cl
@@ -324,7 +324,7 @@ KERNEL_FQ void m19900_comp (KERN_ATTR_TMPS_ESALT (krb5pa_18_tmp_t, krb5pa_18_t))
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m20011-pure.cl b/OpenCL/m20011-pure.cl
index 3fed10939..4733e9f91 100644
--- a/OpenCL/m20011-pure.cl
+++ b/OpenCL/m20011-pure.cl
@@ -360,7 +360,7 @@ KERNEL_FQ void m20011_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, diskcrypt
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m20012-pure.cl b/OpenCL/m20012-pure.cl
index 2a75acd23..dae331981 100644
--- a/OpenCL/m20012-pure.cl
+++ b/OpenCL/m20012-pure.cl
@@ -360,7 +360,7 @@ KERNEL_FQ void m20012_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, diskcrypt
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else
 
diff --git a/OpenCL/m20013-pure.cl b/OpenCL/m20013-pure.cl
index f7f1239a9..fcd2b815a 100644
--- a/OpenCL/m20013-pure.cl
+++ b/OpenCL/m20013-pure.cl
@@ -360,7 +360,7 @@ KERNEL_FQ void m20013_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, diskcrypt
     s_te4[i] = te4[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS ();
 
   #else