diff --git a/OpenCL/m06213.cl b/OpenCL/m06213.cl
index c2857ec0c..5c7f9f6b6 100644
--- a/OpenCL/m06213.cl
+++ b/OpenCL/m06213.cl
@@ -707,7 +707,11 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey3[6] = tmps[gid].out[22];
   ukey3[7] = tmps[gid].out[23];
 
+  #if defined (IS_APPLE) && defined (IS_GPU)
+  volatile u32 ukey4[8];
+  #else
   u32 ukey4[8];
+  #endif
 
   ukey4[0] = tmps[gid].out[24];
   ukey4[1] = tmps[gid].out[25];
@@ -763,7 +767,11 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     }
   }
 
+  #if defined (IS_APPLE) && defined (IS_GPU)
+  volatile u32 ukey5[8];
+  #else
   u32 ukey5[8];
+  #endif
 
   ukey5[0] = tmps[gid].out[32];
   ukey5[1] = tmps[gid].out[33];
@@ -774,7 +782,11 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey5[6] = tmps[gid].out[38];
   ukey5[7] = tmps[gid].out[39];
 
+  #if defined (IS_APPLE) && defined (IS_GPU)
+  volatile u32 ukey6[8];
+  #else
   u32 ukey6[8];
+  #endif
 
   ukey6[0] = tmps[gid].out[40];
   ukey6[1] = tmps[gid].out[41];
diff --git a/OpenCL/m06223.cl b/OpenCL/m06223.cl
index 676194cac..3bc57d908 100644
--- a/OpenCL/m06223.cl
+++ b/OpenCL/m06223.cl
@@ -615,7 +615,11 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey3[6] = swap32 (h32_from_64 (tmps[gid].out[11]));
   ukey3[7] = swap32 (l32_from_64 (tmps[gid].out[11]));
 
+  #if defined (IS_APPLE) && defined (IS_GPU)
+  volatile u32 ukey4[8];
+  #else
   u32 ukey4[8];
+  #endif
 
   ukey4[0] = swap32 (h32_from_64 (tmps[gid].out[12]));
   ukey4[1] = swap32 (l32_from_64 (tmps[gid].out[12]));
diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl
index bdfb6c142..a0f076e68 100644
--- a/OpenCL/m06233.cl
+++ b/OpenCL/m06233.cl
@@ -2019,7 +2019,11 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey3[6] = swap32 (tmps[gid].out[22]);
   ukey3[7] = swap32 (tmps[gid].out[23]);
 
+  #if defined (IS_APPLE) && defined (IS_GPU)
+  volatile u32 ukey4[8];
+  #else
   u32 ukey4[8];
+  #endif
 
   ukey4[0] = swap32 (tmps[gid].out[24]);
   ukey4[1] = swap32 (tmps[gid].out[25]);
@@ -2075,7 +2079,11 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     }
   }
 
+  #if defined (IS_APPLE) && defined (IS_GPU)
+  volatile u32 ukey5[8];
+  #else
   u32 ukey5[8];
+  #endif
 
   ukey5[0] = swap32 (tmps[gid].out[32]);
   ukey5[1] = swap32 (tmps[gid].out[33]);
@@ -2086,7 +2094,11 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   ukey5[6] = swap32 (tmps[gid].out[38]);
   ukey5[7] = swap32 (tmps[gid].out[39]);
 
+  #if defined (IS_APPLE) && defined (IS_GPU)
+  volatile u32 ukey6[8];
+  #else
   u32 ukey6[8];
+  #endif
 
   ukey6[0] = swap32 (tmps[gid].out[40]);
   ukey6[1] = swap32 (tmps[gid].out[41]);