New Strategy: Instead of using volatile, mark the mode as unstable. Remove all volatiles

OpenCL/inc_hash_md4.cl

@@ -110,7 +110,7 @@ DECLSPEC void md4_init (md4_ctx_t *ctx)
 
 DECLSPEC void md4_update_64 (md4_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -696,7 +696,7 @@ DECLSPEC void md4_update_global_utf16le_swap (md4_ctx_t *ctx, const __global u32
 
 DECLSPEC void md4_final (md4_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos);
 
@@ -1230,7 +1230,7 @@ DECLSPEC void md4_init_vector_from_scalar (md4_ctx_vector_t *ctx, md4_ctx_t *ctx
 
 DECLSPEC void md4_update_vector_64 (md4_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -1560,7 +1560,7 @@ DECLSPEC void md4_update_vector_utf16le_swap (md4_ctx_vector_t *ctx, const u32x
 
 DECLSPEC void md4_final_vector (md4_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos);
 

OpenCL/inc_hash_md5.cl

@@ -146,7 +146,7 @@ DECLSPEC void md5_init (md5_ctx_t *ctx)
 
 DECLSPEC void md5_update_64 (md5_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -732,7 +732,7 @@ DECLSPEC void md5_update_global_utf16le_swap (md5_ctx_t *ctx, const __global u32
 
 DECLSPEC void md5_final (md5_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos);
 
@@ -1302,7 +1302,7 @@ DECLSPEC void md5_init_vector_from_scalar (md5_ctx_vector_t *ctx, md5_ctx_t *ctx
 
 DECLSPEC void md5_update_vector_64 (md5_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -1632,7 +1632,7 @@ DECLSPEC void md5_update_vector_utf16le_swap (md5_ctx_vector_t *ctx, const u32x
 
 DECLSPEC void md5_final_vector (md5_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos);
 

OpenCL/inc_hash_ripemd160.cl

@@ -244,7 +244,7 @@ DECLSPEC void ripemd160_init (ripemd160_ctx_t *ctx)
 
 DECLSPEC void ripemd160_update_64 (ripemd160_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -830,7 +830,7 @@ DECLSPEC void ripemd160_update_global_utf16le_swap (ripemd160_ctx_t *ctx, const
 
 DECLSPEC void ripemd160_final (ripemd160_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos);
 
@@ -1499,7 +1499,7 @@ DECLSPEC void ripemd160_init_vector_from_scalar (ripemd160_ctx_vector_t *ctx, ri
 
 DECLSPEC void ripemd160_update_vector_64 (ripemd160_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -1829,7 +1829,7 @@ DECLSPEC void ripemd160_update_vector_utf16le_swap (ripemd160_ctx_vector_t *ctx,
 
 DECLSPEC void ripemd160_final_vector (ripemd160_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos);
 

OpenCL/inc_hash_sha1.cl

@@ -176,7 +176,7 @@ DECLSPEC void sha1_init (sha1_ctx_t *ctx)
 
 DECLSPEC void sha1_update_64 (sha1_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -998,7 +998,7 @@ DECLSPEC void sha1_update_global_utf16be_swap (sha1_ctx_t *ctx, const __global u
 
 DECLSPEC void sha1_final (sha1_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 
@@ -1599,7 +1599,7 @@ DECLSPEC void sha1_init_vector_from_scalar (sha1_ctx_vector_t *ctx, sha1_ctx_t *
 
 DECLSPEC void sha1_update_vector_64 (sha1_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -2013,7 +2013,7 @@ DECLSPEC void sha1_update_vector_utf16beN (sha1_ctx_vector_t *ctx, const u32x *w
 
 DECLSPEC void sha1_final_vector (sha1_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 

OpenCL/inc_hash_sha224.cl

@@ -161,7 +161,7 @@ DECLSPEC void sha224_init (sha224_ctx_t *ctx)
 
 DECLSPEC void sha224_update_64 (sha224_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -747,7 +747,7 @@ DECLSPEC void sha224_update_global_utf16le_swap (sha224_ctx_t *ctx, const __glob
 
 DECLSPEC void sha224_final (sha224_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 
@@ -1316,7 +1316,7 @@ DECLSPEC void sha224_init_vector_from_scalar (sha224_ctx_vector_t *ctx, sha224_c
 
 DECLSPEC void sha224_update_vector_64 (sha224_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -1688,7 +1688,7 @@ DECLSPEC void sha224_update_vector_utf16beN (sha224_ctx_vector_t *ctx, const u32
 
 DECLSPEC void sha224_final_vector (sha224_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 

OpenCL/inc_hash_sha256.cl

@@ -161,7 +161,7 @@ DECLSPEC void sha256_init (sha256_ctx_t *ctx)
 
 DECLSPEC void sha256_update_64 (sha256_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -747,7 +747,7 @@ DECLSPEC void sha256_update_global_utf16le_swap (sha256_ctx_t *ctx, const __glob
 
 DECLSPEC void sha256_final (sha256_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 
@@ -1316,7 +1316,7 @@ DECLSPEC void sha256_init_vector_from_scalar (sha256_ctx_vector_t *ctx, sha256_c
 
 DECLSPEC void sha256_update_vector_64 (sha256_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -1688,7 +1688,7 @@ DECLSPEC void sha256_update_vector_utf16beN (sha256_ctx_vector_t *ctx, const u32
 
 DECLSPEC void sha256_final_vector (sha256_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 

OpenCL/inc_hash_sha384.cl

@@ -185,7 +185,7 @@ DECLSPEC void sha384_init (sha384_ctx_t *ctx)
 
 DECLSPEC void sha384_update_128 (sha384_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, u32 *w4, u32 *w5, u32 *w6, u32 *w7, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 127;
+  const int pos = ctx->len & 127;
 
   ctx->len += len;
 
@@ -1191,7 +1191,7 @@ DECLSPEC void sha384_update_global_utf16le_swap (sha384_ctx_t *ctx, const __glob
 
 DECLSPEC void sha384_final (sha384_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 127;
+  const int pos = ctx->len & 127;
 
   append_0x80_8x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, pos ^ 3);
 
@@ -2012,7 +2012,7 @@ DECLSPEC void sha384_init_vector_from_scalar (sha384_ctx_vector_t *ctx, sha384_c
 
 DECLSPEC void sha384_update_vector_128 (sha384_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 127;
+  const int pos = ctx->len & 127;
 
   ctx->len += len;
 
@@ -2644,7 +2644,7 @@ DECLSPEC void sha384_update_vector_utf16beN (sha384_ctx_vector_t *ctx, const u32
 
 DECLSPEC void sha384_final_vector (sha384_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 127;
+  const int pos = ctx->len & 127;
 
   append_0x80_8x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, pos ^ 3);
 

OpenCL/inc_hash_sha512.cl

@@ -185,7 +185,7 @@ DECLSPEC void sha512_init (sha512_ctx_t *ctx)
 
 DECLSPEC void sha512_update_128 (sha512_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, u32 *w4, u32 *w5, u32 *w6, u32 *w7, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 127;
+  const int pos = ctx->len & 127;
 
   ctx->len += len;
 
@@ -1191,7 +1191,7 @@ DECLSPEC void sha512_update_global_utf16le_swap (sha512_ctx_t *ctx, const __glob
 
 DECLSPEC void sha512_final (sha512_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 127;
+  const int pos = ctx->len & 127;
 
   append_0x80_8x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, pos ^ 3);
 
@@ -2012,7 +2012,7 @@ DECLSPEC void sha512_init_vector_from_scalar (sha512_ctx_vector_t *ctx, sha512_c
 
 DECLSPEC void sha512_update_vector_128 (sha512_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 127;
+  const int pos = ctx->len & 127;
 
   ctx->len += len;
 
@@ -2644,7 +2644,7 @@ DECLSPEC void sha512_update_vector_utf16beN (sha512_ctx_vector_t *ctx, const u32
 
 DECLSPEC void sha512_final_vector (sha512_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 127;
+  const int pos = ctx->len & 127;
 
   append_0x80_8x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, pos ^ 3);
 

OpenCL/inc_hash_streebog256.cl

@@ -799,7 +799,7 @@ DECLSPEC void streebog256_transform (streebog256_ctx_t *ctx, const u32 *w0, cons
 
 DECLSPEC void streebog256_update_64 (streebog256_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len;
+  const int pos = ctx->len;
 
   if ((pos + len) < 64)
   {
@@ -1046,7 +1046,7 @@ DECLSPEC void streebog256_update_global_swap (streebog256_ctx_t *ctx, const __gl
 
 DECLSPEC void streebog256_final (streebog256_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x01_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 
@@ -1487,7 +1487,7 @@ DECLSPEC void streebog256_transform_vector (streebog256_ctx_vector_t *ctx, const
 
 DECLSPEC void streebog256_update_vector_64 (streebog256_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len;
+  const int pos = ctx->len;
 
   if ((pos + len) < 64)
   {
@@ -1677,7 +1677,7 @@ DECLSPEC void streebog256_update_vector_swap (streebog256_ctx_vector_t *ctx, con
 
 DECLSPEC void streebog256_final_vector (streebog256_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x01_4x4_VV (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 

OpenCL/inc_hash_streebog512.cl

@@ -799,7 +799,7 @@ DECLSPEC void streebog512_transform (streebog512_ctx_t *ctx, const u32 *w0, cons
 
 DECLSPEC void streebog512_update_64 (streebog512_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len;
+  const int pos = ctx->len;
 
   if ((pos + len) < 64)
   {
@@ -1046,7 +1046,7 @@ DECLSPEC void streebog512_update_global_swap (streebog512_ctx_t *ctx, const __gl
 
 DECLSPEC void streebog512_final (streebog512_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x01_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 
@@ -1487,7 +1487,7 @@ DECLSPEC void streebog512_transform_vector (streebog512_ctx_vector_t *ctx, const
 
 DECLSPEC void streebog512_update_vector_64 (streebog512_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len;
+  const int pos = ctx->len;
 
   if ((pos + len) < 64)
   {
@@ -1677,7 +1677,7 @@ DECLSPEC void streebog512_update_vector_swap (streebog512_ctx_vector_t *ctx, con
 
 DECLSPEC void streebog512_final_vector (streebog512_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x01_4x4_VV (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 

OpenCL/inc_hash_whirlpool.cl

@@ -1344,7 +1344,7 @@ DECLSPEC void whirlpool_init (whirlpool_ctx_t *ctx, __local u32 (*s_Ch)[256], __
 
 DECLSPEC void whirlpool_update_64 (whirlpool_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -1930,7 +1930,7 @@ DECLSPEC void whirlpool_update_global_utf16le_swap (whirlpool_ctx_t *ctx, const
 
 DECLSPEC void whirlpool_final (whirlpool_ctx_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 
@@ -2603,7 +2603,7 @@ DECLSPEC void whirlpool_init_vector_from_scalar (whirlpool_ctx_vector_t *ctx, wh
 
 DECLSPEC void whirlpool_update_vector_64 (whirlpool_ctx_vector_t *ctx, u32x *w0, u32x *w1, u32x *w2, u32x *w3, const int len)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   ctx->len += len;
 
@@ -2933,7 +2933,7 @@ DECLSPEC void whirlpool_update_vector_utf16le_swap (whirlpool_ctx_vector_t *ctx,
 
 DECLSPEC void whirlpool_final_vector (whirlpool_ctx_vector_t *ctx)
 {
-  MAYBE_VOLATILE const int pos = ctx->len & 63;
+  const int pos = ctx->len & 63;
 
   append_0x80_4x4 (ctx->w0, ctx->w1, ctx->w2, ctx->w3, pos ^ 3);
 

OpenCL/inc_vendor.cl

@@ -75,14 +75,6 @@
 #define DECLSPEC
 #endif
 
-// some kernels under special conditions need treatment
-
-#if defined MAYBE_VOLATILE
-// take the one given by the user
-#else
-#define MAYBE_VOLATILE
-#endif
-
 /**
  * AMD specific
  */

OpenCL/m01500_a0-pure.cl

@@ -498,7 +498,7 @@ __kernel void m01500_mxx (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -582,7 +582,7 @@ __kernel void m01500_sxx (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m01500_a1-pure.cl

@@ -496,7 +496,7 @@ __kernel void m01500_mxx (KERN_ATTR_BASIC ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -659,7 +659,7 @@ __kernel void m01500_sxx (KERN_ATTR_BASIC ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m01600-pure.cl

@@ -258,8 +258,8 @@ __kernel void m01600_loop (KERN_ATTR_TMPS (md5crypt_tmp_t))
 
     if (j1)
     {
-      MAYBE_VOLATILE const u32 off = wpc_len[pc] / 4;
-      MAYBE_VOLATILE const u32 mod = wpc_len[pc] % 4;
+      const u32 off = wpc_len[pc] / 4;
+      const u32 mod = wpc_len[pc] % 4;
 
       u32 *ptr = wpc[pc] + off - 4;
 

OpenCL/m02500-pure.cl

@@ -893,7 +893,7 @@ __kernel void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t))
   __local u32 s_te3[256];
   __local u32 s_te4[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     s_td0[i] = td0[i];
     s_td1[i] = td1[i];

OpenCL/m02501-pure.cl

@@ -768,7 +768,7 @@ __kernel void m02501_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pmk_tmp_t, wpa_eapol_t))
   __local u32 s_te3[256];
   __local u32 s_te4[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     s_td0[i] = td0[i];
     s_td1[i] = td1[i];

OpenCL/m02610_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m02610_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -326,7 +326,7 @@ __kernel void m02610_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02610_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m02610_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -146,7 +146,7 @@ __kernel void m02610_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02610_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m02610_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -383,7 +383,7 @@ __kernel void m02610_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02610_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m02610_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -142,7 +142,7 @@ __kernel void m02610_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02610_a3-optimized.cl

@@ -606,7 +606,7 @@ __kernel void m02610_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -676,7 +676,7 @@ __kernel void m02610_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -746,7 +746,7 @@ __kernel void m02610_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -816,7 +816,7 @@ __kernel void m02610_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -886,7 +886,7 @@ __kernel void m02610_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -956,7 +956,7 @@ __kernel void m02610_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02610_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m02610_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -155,7 +155,7 @@ __kernel void m02610_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02710_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m02710_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -411,7 +411,7 @@ __kernel void m02710_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02710_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m02710_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -468,7 +468,7 @@ __kernel void m02710_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02710_a3-optimized.cl

@@ -775,7 +775,7 @@ __kernel void m02710_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -845,7 +845,7 @@ __kernel void m02710_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -915,7 +915,7 @@ __kernel void m02710_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -985,7 +985,7 @@ __kernel void m02710_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1055,7 +1055,7 @@ __kernel void m02710_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1125,7 +1125,7 @@ __kernel void m02710_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02810_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m02810_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -412,7 +412,7 @@ __kernel void m02810_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02810_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m02810_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -169,7 +169,7 @@ __kernel void m02810_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02810_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m02810_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -469,7 +469,7 @@ __kernel void m02810_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02810_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m02810_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -165,7 +165,7 @@ __kernel void m02810_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02810_a3-optimized.cl

@@ -777,7 +777,7 @@ __kernel void m02810_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -847,7 +847,7 @@ __kernel void m02810_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -917,7 +917,7 @@ __kernel void m02810_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -987,7 +987,7 @@ __kernel void m02810_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1057,7 +1057,7 @@ __kernel void m02810_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1127,7 +1127,7 @@ __kernel void m02810_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m02810_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m02810_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -178,7 +178,7 @@ __kernel void m02810_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03000_a0-pure.cl

@@ -508,7 +508,7 @@ __kernel void m03000_mxx (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -593,7 +593,7 @@ __kernel void m03000_sxx (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m03000_a1-pure.cl

@@ -506,7 +506,7 @@ __kernel void m03000_mxx (KERN_ATTR_BASIC ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -670,7 +670,7 @@ __kernel void m03000_sxx (KERN_ATTR_BASIC ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m03100_a0-optimized.cl

@@ -504,7 +504,7 @@ __kernel void m03100_m04 (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -727,7 +727,7 @@ __kernel void m03100_s04 (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m03100_a1-optimized.cl

@@ -502,7 +502,7 @@ __kernel void m03100_m04 (KERN_ATTR_BASIC ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -785,7 +785,7 @@ __kernel void m03100_s04 (KERN_ATTR_BASIC ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m03100_a3-optimized.cl

@@ -900,7 +900,7 @@ __kernel void m03100_m04 (KERN_ATTR_VECTOR ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -974,7 +974,7 @@ __kernel void m03100_m08 (KERN_ATTR_VECTOR ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -1052,7 +1052,7 @@ __kernel void m03100_s04 (KERN_ATTR_VECTOR ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -1126,7 +1126,7 @@ __kernel void m03100_s08 (KERN_ATTR_VECTOR ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m03710_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m03710_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -357,7 +357,7 @@ __kernel void m03710_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03710_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m03710_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -159,7 +159,7 @@ __kernel void m03710_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03710_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m03710_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -414,7 +414,7 @@ __kernel void m03710_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03710_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m03710_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -155,7 +155,7 @@ __kernel void m03710_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03710_a3-optimized.cl

@@ -633,7 +633,7 @@ __kernel void m03710_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -737,7 +737,7 @@ __kernel void m03710_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -807,7 +807,7 @@ __kernel void m03710_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -877,7 +877,7 @@ __kernel void m03710_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -947,7 +947,7 @@ __kernel void m03710_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -983,7 +983,7 @@ __kernel void m03710_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03710_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m03710_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -168,7 +168,7 @@ __kernel void m03710_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03910_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m03910_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -412,7 +412,7 @@ __kernel void m03910_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03910_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m03910_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -169,7 +169,7 @@ __kernel void m03910_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03910_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m03910_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -469,7 +469,7 @@ __kernel void m03910_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03910_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m03910_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -165,7 +165,7 @@ __kernel void m03910_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03910_a3-optimized.cl

@@ -777,7 +777,7 @@ __kernel void m03910_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -847,7 +847,7 @@ __kernel void m03910_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -917,7 +917,7 @@ __kernel void m03910_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -987,7 +987,7 @@ __kernel void m03910_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1057,7 +1057,7 @@ __kernel void m03910_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1127,7 +1127,7 @@ __kernel void m03910_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m03910_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m03910_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -178,7 +178,7 @@ __kernel void m03910_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04010_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m04010_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -383,7 +383,7 @@ __kernel void m04010_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04010_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m04010_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -150,7 +150,7 @@ __kernel void m04010_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04010_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m04010_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -439,7 +439,7 @@ __kernel void m04010_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04010_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04010_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -146,7 +146,7 @@ __kernel void m04010_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04010_a3-optimized.cl

@@ -673,7 +673,7 @@ __kernel void m04010_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -777,7 +777,7 @@ __kernel void m04010_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -847,7 +847,7 @@ __kernel void m04010_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -917,7 +917,7 @@ __kernel void m04010_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -987,7 +987,7 @@ __kernel void m04010_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1023,7 +1023,7 @@ __kernel void m04010_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04010_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04010_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -163,7 +163,7 @@ __kernel void m04010_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04110_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m04110_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -428,7 +428,7 @@ __kernel void m04110_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04110_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m04110_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -163,7 +163,7 @@ __kernel void m04110_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04110_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m04110_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -486,7 +486,7 @@ __kernel void m04110_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04110_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04110_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -159,7 +159,7 @@ __kernel void m04110_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04110_a3-optimized.cl

@@ -729,7 +729,7 @@ __kernel void m04110_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -833,7 +833,7 @@ __kernel void m04110_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -903,7 +903,7 @@ __kernel void m04110_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -973,7 +973,7 @@ __kernel void m04110_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1043,7 +1043,7 @@ __kernel void m04110_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1079,7 +1079,7 @@ __kernel void m04110_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04110_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04110_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -174,7 +174,7 @@ __kernel void m04110_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04310_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m04310_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -326,7 +326,7 @@ __kernel void m04310_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04310_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m04310_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -146,7 +146,7 @@ __kernel void m04310_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04310_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m04310_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -383,7 +383,7 @@ __kernel void m04310_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04310_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04310_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -142,7 +142,7 @@ __kernel void m04310_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04310_a3-optimized.cl

@@ -606,7 +606,7 @@ __kernel void m04310_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -676,7 +676,7 @@ __kernel void m04310_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -746,7 +746,7 @@ __kernel void m04310_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -816,7 +816,7 @@ __kernel void m04310_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -886,7 +886,7 @@ __kernel void m04310_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -956,7 +956,7 @@ __kernel void m04310_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04310_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04310_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -155,7 +155,7 @@ __kernel void m04310_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04400_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m04400_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -352,7 +352,7 @@ __kernel void m04400_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04400_a0-pure.cl

@@ -44,7 +44,7 @@ __kernel void m04400_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -141,7 +141,7 @@ __kernel void m04400_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04400_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m04400_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -408,7 +408,7 @@ __kernel void m04400_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04400_a1-pure.cl

@@ -42,7 +42,7 @@ __kernel void m04400_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -137,7 +137,7 @@ __kernel void m04400_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04400_a3-optimized.cl

@@ -577,7 +577,7 @@ __kernel void m04400_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -647,7 +647,7 @@ __kernel void m04400_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -717,7 +717,7 @@ __kernel void m04400_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -787,7 +787,7 @@ __kernel void m04400_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -857,7 +857,7 @@ __kernel void m04400_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -927,7 +927,7 @@ __kernel void m04400_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04400_a3-pure.cl

@@ -42,7 +42,7 @@ __kernel void m04400_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -150,7 +150,7 @@ __kernel void m04400_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04500_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m04500_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -379,7 +379,7 @@ __kernel void m04500_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04500_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m04500_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -140,7 +140,7 @@ __kernel void m04500_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04500_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m04500_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -435,7 +435,7 @@ __kernel void m04500_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04500_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04500_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -136,7 +136,7 @@ __kernel void m04500_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04500_a3-optimized.cl

@@ -637,7 +637,7 @@ __kernel void m04500_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -707,7 +707,7 @@ __kernel void m04500_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -777,7 +777,7 @@ __kernel void m04500_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -847,7 +847,7 @@ __kernel void m04500_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -917,7 +917,7 @@ __kernel void m04500_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -987,7 +987,7 @@ __kernel void m04500_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04500_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04500_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -149,7 +149,7 @@ __kernel void m04500_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04520_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m04520_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -610,7 +610,7 @@ __kernel void m04520_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04520_a0-pure.cl

@@ -43,7 +43,7 @@ __kernel void m04520_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -155,7 +155,7 @@ __kernel void m04520_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04520_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m04520_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -666,7 +666,7 @@ __kernel void m04520_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04520_a1-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04520_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -151,7 +151,7 @@ __kernel void m04520_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04520_a3-optimized.cl

@@ -1090,7 +1090,7 @@ __kernel void m04520_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1160,7 +1160,7 @@ __kernel void m04520_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1230,7 +1230,7 @@ __kernel void m04520_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1300,7 +1300,7 @@ __kernel void m04520_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1370,7 +1370,7 @@ __kernel void m04520_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -1440,7 +1440,7 @@ __kernel void m04520_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04520_a3-pure.cl

@@ -41,7 +41,7 @@ __kernel void m04520_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -166,7 +166,7 @@ __kernel void m04520_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04700_a0-optimized.cl

@@ -42,7 +42,7 @@ __kernel void m04700_m04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -335,7 +335,7 @@ __kernel void m04700_s04 (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04700_a0-pure.cl

@@ -44,7 +44,7 @@ __kernel void m04700_mxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -136,7 +136,7 @@ __kernel void m04700_sxx (KERN_ATTR_RULES ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04700_a1-optimized.cl

@@ -40,7 +40,7 @@ __kernel void m04700_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -388,7 +388,7 @@ __kernel void m04700_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04700_a1-pure.cl

@@ -42,7 +42,7 @@ __kernel void m04700_mxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -132,7 +132,7 @@ __kernel void m04700_sxx (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04700_a3-optimized.cl

@@ -577,7 +577,7 @@ __kernel void m04700_m04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -647,7 +647,7 @@ __kernel void m04700_m08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -717,7 +717,7 @@ __kernel void m04700_m16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -787,7 +787,7 @@ __kernel void m04700_s04 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -857,7 +857,7 @@ __kernel void m04700_s08 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -927,7 +927,7 @@ __kernel void m04700_s16 (KERN_ATTR_BASIC ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m04700_a3-pure.cl

@@ -42,7 +42,7 @@ __kernel void m04700_mxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;
@@ -145,7 +145,7 @@ __kernel void m04700_sxx (KERN_ATTR_VECTOR ())
 
   __local u32 l_bin2asc[256];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  for (u32 i = lid; i < 256; i += lsz)
   {
     const u32 i0 = (i >> 0) & 15;
     const u32 i1 = (i >> 4) & 15;

OpenCL/m05300_a0-optimized.cl

@@ -116,14 +116,14 @@ __kernel void m05300_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
@@ -291,14 +291,14 @@ __kernel void m05300_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }

OpenCL/m05300_a1-optimized.cl

@@ -114,14 +114,14 @@ __kernel void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
@@ -349,14 +349,14 @@ __kernel void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }

OpenCL/m05300_a3-optimized.cl

@@ -420,14 +420,14 @@ __kernel void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
@@ -493,14 +493,14 @@ __kernel void m05300_m08 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
@@ -566,14 +566,14 @@ __kernel void m05300_m16 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
@@ -639,14 +639,14 @@ __kernel void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
@@ -712,14 +712,14 @@ __kernel void m05300_s08 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }
@@ -785,14 +785,14 @@ __kernel void m05300_s16 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = esalt_bufs[digests_offset].nr_buf[i];
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = esalt_bufs[digests_offset].msg_buf[i];
   }

OpenCL/m05400_a0-optimized.cl

@@ -120,14 +120,14 @@ __kernel void m05400_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
@@ -304,14 +304,14 @@ __kernel void m05400_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }

OpenCL/m05400_a1-optimized.cl

@@ -118,14 +118,14 @@ __kernel void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
@@ -370,14 +370,14 @@ __kernel void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }

OpenCL/m05400_a3-optimized.cl

@@ -424,14 +424,14 @@ __kernel void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
@@ -497,14 +497,14 @@ __kernel void m05400_m08 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
@@ -570,14 +570,14 @@ __kernel void m05400_m16 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
@@ -643,14 +643,14 @@ __kernel void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
@@ -716,14 +716,14 @@ __kernel void m05400_s08 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }
@@ -789,14 +789,14 @@ __kernel void m05400_s16 (KERN_ATTR_ESALT (ikepsk_t))
 
   __local u32 s_nr_buf[16];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 16; i += lsz)
+  for (u32 i = lid; i < 16; i += lsz)
   {
     s_nr_buf[i] = swap32_S (esalt_bufs[digests_offset].nr_buf[i]);
   }
 
   __local u32 s_msg_buf[128];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 128; i += lsz)
+  for (u32 i = lid; i < 128; i += lsz)
   {
     s_msg_buf[i] = swap32_S (esalt_bufs[digests_offset].msg_buf[i]);
   }

OpenCL/m05500_a0-optimized.cl

@@ -505,7 +505,7 @@ __kernel void m05500_m04 (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -718,7 +718,7 @@ __kernel void m05500_s04 (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m05500_a0-pure.cl

@@ -506,7 +506,7 @@ __kernel void m05500_mxx (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -630,7 +630,7 @@ __kernel void m05500_sxx (KERN_ATTR_RULES ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];

OpenCL/m05500_a1-optimized.cl

@@ -503,7 +503,7 @@ __kernel void m05500_m04 (KERN_ATTR_BASIC ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];
@@ -769,7 +769,7 @@ __kernel void m05500_s04 (KERN_ATTR_BASIC ())
   __local u32 s_SPtrans[8][64];
   __local u32 s_skb[8][64];
 
-  for (MAYBE_VOLATILE u32 i = lid; i < 64; i += lsz)
+  for (u32 i = lid; i < 64; i += lsz)
   {
     s_SPtrans[0][i] = c_SPtrans[0][i];
     s_SPtrans[1][i] = c_SPtrans[1][i];