diff --git a/OpenCL/inc_cipher_aes.cl b/OpenCL/inc_cipher_aes.cl
index 8425414cd..716addc39 100644
--- a/OpenCL/inc_cipher_aes.cl
+++ b/OpenCL/inc_cipher_aes.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_cipher_aes.h"
 
-CONSTANT_AS u32a te0[256] =
+CONSTANT_VK u32a te0[256] =
 {
   0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
   0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@@ -77,7 +77,7 @@ CONSTANT_AS u32a te0[256] =
   0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
 };
 
-CONSTANT_AS u32a te1[256] =
+CONSTANT_VK u32a te1[256] =
 {
   0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
   0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@@ -145,7 +145,7 @@ CONSTANT_AS u32a te1[256] =
   0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
 };
 
-CONSTANT_AS u32a te2[256] =
+CONSTANT_VK u32a te2[256] =
 {
   0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
   0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@@ -213,7 +213,7 @@ CONSTANT_AS u32a te2[256] =
   0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
 };
 
-CONSTANT_AS u32a te3[256] =
+CONSTANT_VK u32a te3[256] =
 {
   0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
   0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@@ -281,7 +281,7 @@ CONSTANT_AS u32a te3[256] =
   0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
 };
 
-CONSTANT_AS u32a te4[256] =
+CONSTANT_VK u32a te4[256] =
 {
   0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
   0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@@ -349,7 +349,7 @@ CONSTANT_AS u32a te4[256] =
   0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
 };
 
-CONSTANT_AS u32a td0[256] =
+CONSTANT_VK u32a td0[256] =
 {
   0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
   0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@@ -417,7 +417,7 @@ CONSTANT_AS u32a td0[256] =
   0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
 };
 
-CONSTANT_AS u32a td1[256] =
+CONSTANT_VK u32a td1[256] =
 {
   0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
   0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@@ -485,7 +485,7 @@ CONSTANT_AS u32a td1[256] =
   0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
 };
 
-CONSTANT_AS u32a td2[256] =
+CONSTANT_VK u32a td2[256] =
 {
   0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
   0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@@ -553,7 +553,7 @@ CONSTANT_AS u32a td2[256] =
   0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
 };
 
-CONSTANT_AS u32a td3[256] =
+CONSTANT_VK u32a td3[256] =
 {
   0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
   0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@@ -621,7 +621,7 @@ CONSTANT_AS u32a td3[256] =
   0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
 };
 
-CONSTANT_AS u32a td4[256] =
+CONSTANT_VK u32a td4[256] =
 {
   0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
   0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
diff --git a/OpenCL/inc_cipher_camellia.cl b/OpenCL/inc_cipher_camellia.cl
index 881e541bf..08c73e6f4 100644
--- a/OpenCL/inc_cipher_camellia.cl
+++ b/OpenCL/inc_cipher_camellia.cl
@@ -21,7 +21,7 @@
 #include "inc_common.h"
 #include "inc_cipher_camellia.h"
 
-CONSTANT_AS u32a c_sbox[256] =
+CONSTANT_VK u32a c_sbox[256] =
 {
   0x70, 0x82, 0x2c, 0xec, 0xb3, 0x27, 0xc0, 0xe5,
   0xe4, 0x85, 0x57, 0x35, 0xea, 0x0c, 0xae, 0x41,
diff --git a/OpenCL/inc_cipher_des.cl b/OpenCL/inc_cipher_des.cl
index a90d5788b..b247e1e22 100644
--- a/OpenCL/inc_cipher_des.cl
+++ b/OpenCL/inc_cipher_des.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_cipher_des.h"
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     /* nibble 0 */
@@ -165,7 +165,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/inc_cipher_kuznyechik.cl b/OpenCL/inc_cipher_kuznyechik.cl
index 2c3792b96..a5768f70f 100644
--- a/OpenCL/inc_cipher_kuznyechik.cl
+++ b/OpenCL/inc_cipher_kuznyechik.cl
@@ -18,7 +18,7 @@
 #include "inc_common.h"
 #include "inc_cipher_kuznyechik.h"
 
-CONSTANT_AS u32a k_sbox[256] =
+CONSTANT_VK u32a k_sbox[256] =
 {
   0xfc, 0xee, 0xdd, 0x11, 0xcf, 0x6e, 0x31, 0x16,
   0xfb, 0xc4, 0xfa, 0xda, 0x23, 0xc5, 0x04, 0x4d,
@@ -54,7 +54,7 @@ CONSTANT_AS u32a k_sbox[256] =
   0xd1, 0x66, 0xaf, 0xc2, 0x39, 0x4b, 0x63, 0xb6
 };
 
-CONSTANT_AS u32a k_sbox_inv[256] =
+CONSTANT_VK u32a k_sbox_inv[256] =
 {
   0xa5, 0x2d, 0x32, 0x8f, 0x0e, 0x30, 0x38, 0xc0,
   0x54, 0xe6, 0x9e, 0x39, 0x55, 0x7e, 0x52, 0x91,
diff --git a/OpenCL/inc_cipher_twofish.cl b/OpenCL/inc_cipher_twofish.cl
index 90b6cf600..dc0461b24 100644
--- a/OpenCL/inc_cipher_twofish.cl
+++ b/OpenCL/inc_cipher_twofish.cl
@@ -25,7 +25,7 @@
 #include "inc_common.h"
 #include "inc_cipher_twofish.h"
 
-CONSTANT_AS u32a q_tab[2][256] =
+CONSTANT_VK u32a q_tab[2][256] =
 {
   {
     0xA9, 0x67, 0xB3, 0xE8, 0x04, 0xFD, 0xA3, 0x76, 0x9A, 0x92, 0x80, 0x78,
@@ -77,7 +77,7 @@ CONSTANT_AS u32a q_tab[2][256] =
   }
 };
 
-CONSTANT_AS u32a  m_tab[4][256] =
+CONSTANT_VK u32a  m_tab[4][256] =
 {
   { 0xBCBC3275, 0xECEC21F3, 0x202043C6, 0xB3B3C9F4, 0xDADA03DB, 0x02028B7B,
     0xE2E22BFB, 0x9E9EFAC8, 0xC9C9EC4A, 0xD4D409D3, 0x18186BE6, 0x1E1E9F6B,
diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl
index 844e4ba12..68fdf0b88 100644
--- a/OpenCL/inc_common.cl
+++ b/OpenCL/inc_common.cl
@@ -305,7 +305,9 @@ DECLSPEC u64x hl32_to_64 (const u32x a, const u32x b)
 
 DECLSPEC u32x hc_rotl32 (const u32x a, const int n)
 {
-  #ifdef _CPU_OPENCL_EMU_H
+  #if   defined _CPU_OPENCL_EMU_H
+  return rotl32 (a, n);
+  #elif defined IS_CUDA
   return rotl32 (a, n);
   #else
   return rotate (a, (u32x) (n));
@@ -314,7 +316,9 @@ DECLSPEC u32x hc_rotl32 (const u32x a, const int n)
 
 DECLSPEC u32x hc_rotr32 (const u32x a, const int n)
 {
-  #ifdef _CPU_OPENCL_EMU_H
+  #if   defined _CPU_OPENCL_EMU_H
+  return rotr32 (a, n);
+  #elif defined IS_CUDA
   return rotr32 (a, n);
   #else
   return rotate (a, (u32x) (32 - n));
@@ -323,8 +327,10 @@ DECLSPEC u32x hc_rotr32 (const u32x a, const int n)
 
 DECLSPEC u32 hc_rotl32_S (const u32 a, const int n)
 {
-  #ifdef _CPU_OPENCL_EMU_H
+  #if   defined _CPU_OPENCL_EMU_H
   return rotl32 (a, n);
+  #elif defined IS_CUDA
+  return rotl32_S (a, n);
   #else
   return rotate (a, (u32) (n));
   #endif
@@ -332,8 +338,10 @@ DECLSPEC u32 hc_rotl32_S (const u32 a, const int n)
 
 DECLSPEC u32 hc_rotr32_S (const u32 a, const int n)
 {
-  #ifdef _CPU_OPENCL_EMU_H
+  #if   defined _CPU_OPENCL_EMU_H
   return rotr32 (a, n);
+  #elif defined IS_CUDA
+  return rotr32_S (a, n);
   #else
   return rotate (a, (u32) (32 - n));
   #endif
@@ -341,7 +349,9 @@ DECLSPEC u32 hc_rotr32_S (const u32 a, const int n)
 
 DECLSPEC u64x hc_rotl64 (const u64x a, const int n)
 {
-  #ifdef _CPU_OPENCL_EMU_H
+  #if   defined _CPU_OPENCL_EMU_H
+  return rotl64 (a, n);
+  #elif defined IS_CUDA
   return rotl64 (a, n);
   #else
   return rotate (a, (u64x) (n));
@@ -350,7 +360,9 @@ DECLSPEC u64x hc_rotl64 (const u64x a, const int n)
 
 DECLSPEC u64x hc_rotr64 (const u64x a, const int n)
 {
-  #ifdef _CPU_OPENCL_EMU_H
+  #if   defined _CPU_OPENCL_EMU_H
+  return rotr64 (a, n);
+  #elif defined IS_CUDA
   return rotr64 (a, n);
   #else
   return rotate (a, (u64x) (64 - n));
@@ -359,8 +371,10 @@ DECLSPEC u64x hc_rotr64 (const u64x a, const int n)
 
 DECLSPEC u64 hc_rotl64_S (const u64 a, const int n)
 {
-  #ifdef _CPU_OPENCL_EMU_H
+  #if   defined _CPU_OPENCL_EMU_H
   return rotl64 (a, n);
+  #elif defined IS_CUDA
+  return rotl64_S (a, n);
   #else
   return rotate (a, (u64) (n));
   #endif
@@ -368,8 +382,10 @@ DECLSPEC u64 hc_rotl64_S (const u64 a, const int n)
 
 DECLSPEC u64 hc_rotr64_S (const u64 a, const int n)
 {
-  #ifdef _CPU_OPENCL_EMU_H
+  #if   defined _CPU_OPENCL_EMU_H
   return rotr64 (a, n);
+  #elif defined IS_CUDA
+  return rotr64_S (a, n);
   #else
   return rotate (a, (u64) (64 - n));
   #endif
diff --git a/OpenCL/inc_hash_sha224.cl b/OpenCL/inc_hash_sha224.cl
index 0758b1f54..e93206e44 100644
--- a/OpenCL/inc_hash_sha224.cl
+++ b/OpenCL/inc_hash_sha224.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_hash_sha224.h"
 
-CONSTANT_AS u32a k_sha224[64] =
+CONSTANT_VK u32a k_sha224[64] =
 {
   SHA224C00, SHA224C01, SHA224C02, SHA224C03,
   SHA224C04, SHA224C05, SHA224C06, SHA224C07,
diff --git a/OpenCL/inc_hash_sha256.cl b/OpenCL/inc_hash_sha256.cl
index f52e551da..de2bd5897 100644
--- a/OpenCL/inc_hash_sha256.cl
+++ b/OpenCL/inc_hash_sha256.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_hash_sha256.h"
 
-CONSTANT_AS u32a k_sha256[64] =
+CONSTANT_VK u32a k_sha256[64] =
 {
   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
   SHA256C04, SHA256C05, SHA256C06, SHA256C07,
diff --git a/OpenCL/inc_hash_sha384.cl b/OpenCL/inc_hash_sha384.cl
index cdfa357ca..ea26ec734 100644
--- a/OpenCL/inc_hash_sha384.cl
+++ b/OpenCL/inc_hash_sha384.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_hash_sha384.h"
 
-CONSTANT_AS u64a k_sha384[80] =
+CONSTANT_VK u64a k_sha384[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,
diff --git a/OpenCL/inc_hash_sha512.cl b/OpenCL/inc_hash_sha512.cl
index 9ea5463d8..783a66fbe 100644
--- a/OpenCL/inc_hash_sha512.cl
+++ b/OpenCL/inc_hash_sha512.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_hash_sha512.h"
 
-CONSTANT_AS u64a k_sha512[80] =
+CONSTANT_VK u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,
diff --git a/OpenCL/inc_hash_streebog256.cl b/OpenCL/inc_hash_streebog256.cl
index 58900e7d7..d97926000 100644
--- a/OpenCL/inc_hash_streebog256.cl
+++ b/OpenCL/inc_hash_streebog256.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_hash_streebog256.h"
 
-CONSTANT_AS u64a sbob256_sl64[8][256] =
+CONSTANT_VK u64a sbob256_sl64[8][256] =
 {
   {
     0xd031c397ce553fe6, 0x16ba5b01b006b525, 0xa89bade6296e70c8, 0x6a1f525d77d3435b,
@@ -541,7 +541,7 @@ CONSTANT_AS u64a sbob256_sl64[8][256] =
   },
 };
 
-CONSTANT_AS u64a sbob256_rc64[12][8] =
+CONSTANT_VK u64a sbob256_rc64[12][8] =
 {
   {
     0xe9daca1eda5b08b1, 0x1f7c65c0812fcbeb, 0x16d0452e43766a2f, 0xfcc485758db84e71,
diff --git a/OpenCL/inc_hash_streebog512.cl b/OpenCL/inc_hash_streebog512.cl
index 7ad416ade..fe86b1822 100644
--- a/OpenCL/inc_hash_streebog512.cl
+++ b/OpenCL/inc_hash_streebog512.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_hash_streebog512.h"
 
-CONSTANT_AS u64a sbob512_sl64[8][256] =
+CONSTANT_VK u64a sbob512_sl64[8][256] =
 {
   {
     0xd031c397ce553fe6, 0x16ba5b01b006b525, 0xa89bade6296e70c8, 0x6a1f525d77d3435b,
@@ -541,7 +541,7 @@ CONSTANT_AS u64a sbob512_sl64[8][256] =
   },
 };
 
-CONSTANT_AS u64a sbob512_rc64[12][8] =
+CONSTANT_VK u64a sbob512_rc64[12][8] =
 {
   {
     0xe9daca1eda5b08b1, 0x1f7c65c0812fcbeb, 0x16d0452e43766a2f, 0xfcc485758db84e71,
diff --git a/OpenCL/inc_hash_whirlpool.cl b/OpenCL/inc_hash_whirlpool.cl
index d4e1206c4..b4933e6f7 100644
--- a/OpenCL/inc_hash_whirlpool.cl
+++ b/OpenCL/inc_hash_whirlpool.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_hash_whirlpool.h"
 
-CONSTANT_AS u32a Ch[8][256] =
+CONSTANT_VK u32a Ch[8][256] =
 {
   {
     0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@@ -541,7 +541,7 @@ CONSTANT_AS u32a Ch[8][256] =
   }
 };
 
-CONSTANT_AS u32a Cl[8][256] =
+CONSTANT_VK u32a Cl[8][256] =
 {
   {
     0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@@ -1073,7 +1073,7 @@ CONSTANT_AS u32a Cl[8][256] =
   },
 };
 
-CONSTANT_AS u32a rch[R + 1] =
+CONSTANT_VK u32a rch[R + 1] =
 {
   0x00000000,
   0x1823c6e8,
@@ -1088,7 +1088,7 @@ CONSTANT_AS u32a rch[R + 1] =
   0xca2dbf07,
 };
 
-CONSTANT_AS u32a rcl[R + 1] =
+CONSTANT_VK u32a rcl[R + 1] =
 {
   0x00000000,
   0x87b8014f,
diff --git a/OpenCL/inc_platform.cl b/OpenCL/inc_platform.cl
index 1dc643173..e5924dd13 100644
--- a/OpenCL/inc_platform.cl
+++ b/OpenCL/inc_platform.cl
@@ -44,6 +44,46 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
   return blockDim.x;
 }
 
+DECLSPEC u32x rotl32 (const u32x a, const int n)
+{
+  return ((a << n) | ((a >> (32 - n))));
+}
+
+DECLSPEC u32x rotr32 (const u32x a, const int n)
+{
+  return ((a >> n) | ((a << (32 - n))));
+}
+
+DECLSPEC u32 rotl32_S (const u32 a, const int n)
+{
+  return ((a << n) | ((a >> (32 - n))));
+}
+
+DECLSPEC u32 rotr32_S (const u32 a, const int n)
+{
+  return ((a >> n) | ((a << (32 - n))));
+}
+
+DECLSPEC u64x rotl64 (const u64x a, const int n)
+{
+  return ((a << n) | ((a >> (64 - n))));
+}
+
+DECLSPEC u64x rotr64 (const u64x a, const int n)
+{
+  return ((a >> n) | ((a << (64 - n))));
+}
+
+DECLSPEC u64 rotl64_S (const u64 a, const int n)
+{
+  return ((a << n) | ((a >> (64 - n))));
+}
+
+DECLSPEC u64 rotr64_S (const u64 a, const int n)
+{
+  return ((a >> n) | ((a << (64 - n))));
+}
+
 #define SYNC_THREADS() __syncthreads ()
 #endif
 
diff --git a/OpenCL/inc_platform.h b/OpenCL/inc_platform.h
index a8ce27fef..1055838c0 100644
--- a/OpenCL/inc_platform.h
+++ b/OpenCL/inc_platform.h
@@ -14,7 +14,16 @@ DECLSPEC size_t get_global_id   (const u32 dimindx __attribute__((unused)));
 DECLSPEC size_t get_local_id    (const u32 dimindx __attribute__((unused)));
 DECLSPEC size_t get_local_size  (const u32 dimindx __attribute__((unused)));
 
-#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n))))
+DECLSPEC u32x hc_rotl32   (const u32x a, const int n);
+DECLSPEC u32x hc_rotr32   (const u32x a, const int n);
+DECLSPEC u32  hc_rotl32_S (const u32  a, const int n);
+DECLSPEC u32  hc_rotr32_S (const u32  a, const int n);
+DECLSPEC u64x hc_rotl64   (const u64x a, const int n);
+DECLSPEC u64x hc_rotr64   (const u64x a, const int n);
+DECLSPEC u64  hc_rotl64_S (const u64  a, const int n);
+DECLSPEC u64  hc_rotr64_S (const u64  a, const int n);
+
+//#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n))))
 #define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a))))
 #endif
 
diff --git a/OpenCL/inc_truecrypt_crc32.cl b/OpenCL/inc_truecrypt_crc32.cl
index 391ec91d0..2bc30bffa 100644
--- a/OpenCL/inc_truecrypt_crc32.cl
+++ b/OpenCL/inc_truecrypt_crc32.cl
@@ -9,7 +9,7 @@
 #include "inc_common.h"
 #include "inc_truecrypt_crc32.h"
 
-CONSTANT_AS u32a crc32tab[0x100] =
+CONSTANT_VK u32a crc32tab[0x100] =
 {
   0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
   0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h
index 77997dd97..d1c1f4498 100644
--- a/OpenCL/inc_types.h
+++ b/OpenCL/inc_types.h
@@ -368,7 +368,6 @@ inline __device__ u64x operator ~  (const u64x a) { return u64x (~a.s0, ~a.s1, ~
 
 #if VECT_SIZE == 8
 
-
 class u8x
 {
   private:
@@ -497,6 +496,50 @@ inline __device__ u32x operator *  (const u32x a, const u32x b) { return u32x ((
 
 inline __device__ u32x operator ~  (const u32x a) { return u32x (~a.s0, ~a.s1, ~a.s2, ~a.s3, ~a.s4, ~a.s5, ~a.s6, ~a.s7); }
 
+inline __device__ bool operator != (const u64x a, const u64  b) { return ((a.s0 != b)    && (a.s1 != b)    && (a.s2 != b)    && (a.s3 != b)    && (a.s4 != b)    && (a.s5 != b)    && (a.s6 != b)    && (a.s7 != b)   ); }
+inline __device__ bool operator != (const u64x a, const u64x b) { return ((a.s0 != b.s0) && (a.s1 != b.s1) && (a.s2 != b.s2) && (a.s3 != b.s3) && (a.s4 != b.s4) && (a.s5 != b.s5) && (a.s6 != b.s6) && (a.s7 != b.s7)); }
+
+inline __device__ void operator ^= (u64x &a, const u64  b) { a.s0 ^= b;    a.s1 ^= b;    a.s2 ^= b;    a.s3 ^= b;    a.s4 ^= b;    a.s5 ^= b;    a.s6 ^= b;    a.s7 ^= b;     }
+inline __device__ void operator ^= (u64x &a, const u64x b) { a.s0 ^= b.s0; a.s1 ^= b.s1; a.s2 ^= b.s2; a.s3 ^= b.s3; a.s4 ^= b.s4; a.s5 ^= b.s5; a.s6 ^= b.s6; a.s7 ^= b.s7;  }
+
+inline __device__ void operator |= (u64x &a, const u64  b) { a.s0 |= b;    a.s1 |= b;    a.s2 |= b;    a.s3 |= b;    a.s4 |= b;    a.s5 |= b;    a.s6 |= b;    a.s7 |= b;     }
+inline __device__ void operator |= (u64x &a, const u64x b) { a.s0 |= b.s0; a.s1 |= b.s1; a.s2 |= b.s2; a.s3 |= b.s3; a.s4 |= b.s4; a.s5 |= b.s5; a.s6 |= b.s6; a.s7 |= b.s7;  }
+
+inline __device__ void operator &= (u64x &a, const u64  b) { a.s0 &= b;    a.s1 &= b;    a.s2 &= b;    a.s3 &= b;    a.s4 &= b;    a.s5 &= b;    a.s6 &= b;    a.s7 &= b;     }
+inline __device__ void operator &= (u64x &a, const u64x b) { a.s0 &= b.s0; a.s1 &= b.s1; a.s2 &= b.s2; a.s3 &= b.s3; a.s4 &= b.s4; a.s5 &= b.s5; a.s6 &= b.s6; a.s7 &= b.s7;  }
+
+inline __device__ void operator += (u64x &a, const u64  b) { a.s0 += b;    a.s1 += b;    a.s2 += b;    a.s3 += b;    a.s4 += b;    a.s5 += b;    a.s6 += b;    a.s7 += b;     }
+inline __device__ void operator += (u64x &a, const u64x b) { a.s0 += b.s0; a.s1 += b.s1; a.s2 += b.s2; a.s3 += b.s3; a.s4 += b.s4; a.s5 += b.s5; a.s6 += b.s6; a.s7 += b.s7;  }
+
+inline __device__ void operator -= (u64x &a, const u64  b) { a.s0 -= b;    a.s1 -= b;    a.s2 -= b;    a.s3 -= b;    a.s4 -= b;    a.s5 -= b;    a.s6 -= b;    a.s7 -= b;     }
+inline __device__ void operator -= (u64x &a, const u64x b) { a.s0 -= b.s0; a.s1 -= b.s1; a.s2 -= b.s2; a.s3 -= b.s3; a.s4 -= b.s4; a.s5 -= b.s5; a.s6 -= b.s6; a.s7 -= b.s7;  }
+
+inline __device__ u64x operator << (const u64x a, const u64  b) { return u64x ((a.s0 << b),    (a.s1 << b)   , (a.s2 << b),    (a.s3 << b)   , (a.s4 << b),    (a.s5 << b)   , (a.s6 << b),    (a.s7 << b)   );  }
+inline __device__ u64x operator << (const u64x a, const u64x b) { return u64x ((a.s0 << b.s0), (a.s1 << b.s1), (a.s2 << b.s2), (a.s3 << b.s3), (a.s4 << b.s4), (a.s5 << b.s5), (a.s6 << b.s6), (a.s7 << b.s7));  }
+
+inline __device__ u64x operator >> (const u64x a, const u64  b) { return u64x ((a.s0 >> b),    (a.s1 >> b)   , (a.s2 >> b),    (a.s3 >> b)   , (a.s4 >> b),    (a.s5 >> b)   , (a.s6 >> b),    (a.s7 >> b)   );  }
+inline __device__ u64x operator >> (const u64x a, const u64x b) { return u64x ((a.s0 >> b.s0), (a.s1 >> b.s1), (a.s2 >> b.s2), (a.s3 >> b.s3), (a.s4 >> b.s4), (a.s5 >> b.s5), (a.s6 >> b.s6), (a.s7 >> b.s7));  }
+
+inline __device__ u64x operator ^  (const u64x a, const u64  b) { return u64x ((a.s0 ^  b),    (a.s1 ^  b)   , (a.s2 ^  b),    (a.s3 ^  b)   , (a.s4 ^  b),    (a.s5 ^  b)   , (a.s6 ^  b),    (a.s7 ^  b)   );  }
+inline __device__ u64x operator ^  (const u64x a, const u64x b) { return u64x ((a.s0 ^  b.s0), (a.s1 ^  b.s1), (a.s2 ^  b.s2), (a.s3 ^  b.s3), (a.s4 ^  b.s4), (a.s5 ^  b.s5), (a.s6 ^  b.s6), (a.s7 ^  b.s7));  }
+
+inline __device__ u64x operator |  (const u64x a, const u64  b) { return u64x ((a.s0 |  b),    (a.s1 |  b)   , (a.s2 |  b),    (a.s3 |  b)   , (a.s4 |  b),    (a.s5 |  b)   , (a.s6 |  b),    (a.s7 |  b)   );  }
+inline __device__ u64x operator |  (const u64x a, const u64x b) { return u64x ((a.s0 |  b.s0), (a.s1 |  b.s1), (a.s2 |  b.s2), (a.s3 |  b.s3), (a.s4 |  b.s4), (a.s5 |  b.s5), (a.s6 |  b.s6), (a.s7 |  b.s7));  }
+
+inline __device__ u64x operator &  (const u64x a, const u64  b) { return u64x ((a.s0 &  b),    (a.s1 &  b)   , (a.s2 &  b),    (a.s3 &  b)   , (a.s4 &  b),    (a.s5 &  b)   , (a.s6 &  b),    (a.s7 &  b)   );  }
+inline __device__ u64x operator &  (const u64x a, const u64x b) { return u64x ((a.s0 &  b.s0), (a.s1 &  b.s1), (a.s2 &  b.s2), (a.s3 &  b.s3), (a.s4 &  b.s4), (a.s5 &  b.s5), (a.s6 &  b.s6), (a.s7 &  b.s7));  }
+
+inline __device__ u64x operator +  (const u64x a, const u64  b) { return u64x ((a.s0 +  b),    (a.s1 +  b)   , (a.s2 +  b),    (a.s3 +  b)   , (a.s4 +  b),    (a.s5 +  b)   , (a.s6 +  b),    (a.s7 +  b)   );  }
+inline __device__ u64x operator +  (const u64x a, const u64x b) { return u64x ((a.s0 +  b.s0), (a.s1 +  b.s1), (a.s2 +  b.s2), (a.s3 +  b.s3), (a.s4 +  b.s4), (a.s5 +  b.s5), (a.s6 +  b.s6), (a.s7 +  b.s7));  }
+
+inline __device__ u64x operator -  (const u64x a, const u64  b) { return u64x ((a.s0 -  b),    (a.s1 -  b)   , (a.s2 -  b),    (a.s3 -  b)   , (a.s4 -  b),    (a.s5 -  b)   , (a.s6 -  b),    (a.s7 -  b)   );  }
+inline __device__ u64x operator -  (const u64x a, const u64x b) { return u64x ((a.s0 -  b.s0), (a.s1 -  b.s1), (a.s2 -  b.s2), (a.s3 -  b.s3), (a.s4 -  b.s4), (a.s5 -  b.s5), (a.s6 -  b.s6), (a.s7 -  b.s7));  }
+
+inline __device__ u64x operator *  (const u64x a, const u64  b) { return u64x ((a.s0 *  b),    (a.s1 *  b)   , (a.s2 *  b),    (a.s3 *  b)   , (a.s4 *  b),    (a.s5 *  b)   , (a.s6 *  b),    (a.s7 *  b)   );  }
+inline __device__ u64x operator *  (const u64x a, const u64x b) { return u64x ((a.s0 *  b.s0), (a.s1 *  b.s1), (a.s2 *  b.s2), (a.s3 *  b.s3), (a.s4 *  b.s4), (a.s5 *  b.s5), (a.s6 *  b.s6), (a.s7 *  b.s7));  }
+
+inline __device__ u64x operator ~  (const u64x a) { return u64x (~a.s0, ~a.s1, ~a.s2, ~a.s3, ~a.s4, ~a.s5, ~a.s6, ~a.s7); }
+
 #endif
 
 #if VECT_SIZE == 16
diff --git a/OpenCL/inc_vendor.h b/OpenCL/inc_vendor.h
index f69e0573b..279c87c02 100644
--- a/OpenCL/inc_vendor.h
+++ b/OpenCL/inc_vendor.h
@@ -15,16 +15,19 @@
 #endif
 
 #if defined IS_NATIVE
+#define CONSTANT_VK
 #define CONSTANT_AS
 #define GLOBAL_AS
 #define LOCAL_AS
 #define KERNEL_FQ
 #elif defined IS_CUDA
+#define CONSTANT_VK __constant__
 #define CONSTANT_AS
 #define GLOBAL_AS
 #define LOCAL_AS
 #define KERNEL_FQ   extern "C" __global__
 #elif defined IS_OPENCL
+#define CONSTANT_VK __constant
 #define CONSTANT_AS __constant
 #define GLOBAL_AS   __global
 #define LOCAL_AS    __local
diff --git a/OpenCL/m01500_a0-pure.cl b/OpenCL/m01500_a0-pure.cl
index beabe09f9..944f04de0 100644
--- a/OpenCL/m01500_a0-pure.cl
+++ b/OpenCL/m01500_a0-pure.cl
@@ -35,7 +35,7 @@
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x00820200, 0x00020000, 0x80800000, 0x80820200,
@@ -183,7 +183,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m01500_a1-pure.cl b/OpenCL/m01500_a1-pure.cl
index 3d9f06e01..3bcf6344f 100644
--- a/OpenCL/m01500_a1-pure.cl
+++ b/OpenCL/m01500_a1-pure.cl
@@ -33,7 +33,7 @@
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x00820200, 0x00020000, 0x80800000, 0x80820200,
@@ -181,7 +181,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m03000_a0-pure.cl b/OpenCL/m03000_a0-pure.cl
index 96ef208e2..a2eb1ad60 100644
--- a/OpenCL/m03000_a0-pure.cl
+++ b/OpenCL/m03000_a0-pure.cl
@@ -35,7 +35,7 @@
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -183,7 +183,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m03000_a1-pure.cl b/OpenCL/m03000_a1-pure.cl
index d0e7f6b9f..92665dda4 100644
--- a/OpenCL/m03000_a1-pure.cl
+++ b/OpenCL/m03000_a1-pure.cl
@@ -33,7 +33,7 @@
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -181,7 +181,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m03200-pure.cl b/OpenCL/m03200-pure.cl
index 4b7c6c79d..938cb1f48 100644
--- a/OpenCL/m03200-pure.cl
+++ b/OpenCL/m03200-pure.cl
@@ -28,7 +28,7 @@ typedef struct bcrypt_tmp
 
 // http://www.schneier.com/code/constants.txt
 
-CONSTANT_AS u32a c_sbox0[256] =
+CONSTANT_VK u32a c_sbox0[256] =
 {
   0xd1310ba6, 0x98dfb5ac, 0x2ffd72db, 0xd01adfb7,
   0xb8e1afed, 0x6a267e96, 0xba7c9045, 0xf12c7f99,
@@ -96,7 +96,7 @@ CONSTANT_AS u32a c_sbox0[256] =
   0x53b02d5d, 0xa99f8fa1, 0x08ba4799, 0x6e85076a
 };
 
-CONSTANT_AS u32a c_sbox1[256] =
+CONSTANT_VK u32a c_sbox1[256] =
 {
   0x4b7a70e9, 0xb5b32944, 0xdb75092e, 0xc4192623,
   0xad6ea6b0, 0x49a7df7d, 0x9cee60b8, 0x8fedb266,
@@ -164,7 +164,7 @@ CONSTANT_AS u32a c_sbox1[256] =
   0x153e21e7, 0x8fb03d4a, 0xe6e39f2b, 0xdb83adf7
 };
 
-CONSTANT_AS u32a c_sbox2[256] =
+CONSTANT_VK u32a c_sbox2[256] =
 {
   0xe93d5a68, 0x948140f7, 0xf64c261c, 0x94692934,
   0x411520f7, 0x7602d4f7, 0xbcf46b2e, 0xd4a20068,
@@ -232,7 +232,7 @@ CONSTANT_AS u32a c_sbox2[256] =
   0xd79a3234, 0x92638212, 0x670efa8e, 0x406000e0
 };
 
-CONSTANT_AS u32a c_sbox3[256] =
+CONSTANT_VK u32a c_sbox3[256] =
 {
   0x3a39ce37, 0xd3faf5cf, 0xabc27737, 0x5ac52d1b,
   0x5cb0679e, 0x4fa33742, 0xd3822740, 0x99bc9bbe,
diff --git a/OpenCL/m05500_a0-optimized.cl b/OpenCL/m05500_a0-optimized.cl
index f9fd1e11a..893cbc2c1 100644
--- a/OpenCL/m05500_a0-optimized.cl
+++ b/OpenCL/m05500_a0-optimized.cl
@@ -48,7 +48,7 @@ typedef struct netntlm
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -196,7 +196,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m05500_a0-pure.cl b/OpenCL/m05500_a0-pure.cl
index aaef0fbf9..e2087f918 100644
--- a/OpenCL/m05500_a0-pure.cl
+++ b/OpenCL/m05500_a0-pure.cl
@@ -48,7 +48,7 @@ typedef struct netntlm
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -196,7 +196,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m05500_a1-optimized.cl b/OpenCL/m05500_a1-optimized.cl
index a1f12ff86..2f4a8a912 100644
--- a/OpenCL/m05500_a1-optimized.cl
+++ b/OpenCL/m05500_a1-optimized.cl
@@ -46,7 +46,7 @@ typedef struct netntlm
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -194,7 +194,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m05500_a1-pure.cl b/OpenCL/m05500_a1-pure.cl
index 3e482266c..6872de6cc 100644
--- a/OpenCL/m05500_a1-pure.cl
+++ b/OpenCL/m05500_a1-pure.cl
@@ -46,7 +46,7 @@ typedef struct netntlm
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -194,7 +194,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m05500_a3-optimized.cl b/OpenCL/m05500_a3-optimized.cl
index ec25450db..9901e818d 100644
--- a/OpenCL/m05500_a3-optimized.cl
+++ b/OpenCL/m05500_a3-optimized.cl
@@ -46,7 +46,7 @@ typedef struct netntlm
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -194,7 +194,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m05500_a3-pure.cl b/OpenCL/m05500_a3-pure.cl
index eca392158..467818f9f 100644
--- a/OpenCL/m05500_a3-pure.cl
+++ b/OpenCL/m05500_a3-pure.cl
@@ -46,7 +46,7 @@ typedef struct netntlm
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -194,7 +194,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m05800-optimized.cl b/OpenCL/m05800-optimized.cl
index 658a9f93e..4d013411c 100644
--- a/OpenCL/m05800-optimized.cl
+++ b/OpenCL/m05800-optimized.cl
@@ -20,7 +20,7 @@ typedef struct androidpin_tmp
 
 } androidpin_tmp_t;
 
-CONSTANT_AS u32a c_pc_dec[1024] =
+CONSTANT_VK u32a c_pc_dec[1024] =
 {
   0x00000030,
   0x00000031,
@@ -1048,7 +1048,7 @@ CONSTANT_AS u32a c_pc_dec[1024] =
   0x33323031,
 };
 
-CONSTANT_AS u32a c_pc_len[1024] =
+CONSTANT_VK u32a c_pc_len[1024] =
 {
   1,
   1,
diff --git a/OpenCL/m05800-pure.cl b/OpenCL/m05800-pure.cl
index e7cddac46..1813576b7 100644
--- a/OpenCL/m05800-pure.cl
+++ b/OpenCL/m05800-pure.cl
@@ -20,7 +20,7 @@ typedef struct androidpin_tmp
 
 } androidpin_tmp_t;
 
-CONSTANT_AS u32a c_pc_dec[1024] =
+CONSTANT_VK u32a c_pc_dec[1024] =
 {
   0x00000030,
   0x00000031,
@@ -1048,7 +1048,7 @@ CONSTANT_AS u32a c_pc_dec[1024] =
   0x33323031,
 };
 
-CONSTANT_AS u32a c_pc_len[1024] =
+CONSTANT_VK u32a c_pc_len[1024] =
 {
   1,
   1,
diff --git a/OpenCL/m06900_a0-optimized.cl b/OpenCL/m06900_a0-optimized.cl
index 07939b450..5c3ef3850 100644
--- a/OpenCL/m06900_a0-optimized.cl
+++ b/OpenCL/m06900_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a c_tables[4][256] =
+CONSTANT_VK u32a c_tables[4][256] =
 {
   {
     0x00072000, 0x00075000, 0x00074800, 0x00071000,
diff --git a/OpenCL/m06900_a1-optimized.cl b/OpenCL/m06900_a1-optimized.cl
index 7cf40f6cf..41cc12653 100644
--- a/OpenCL/m06900_a1-optimized.cl
+++ b/OpenCL/m06900_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a c_tables[4][256] =
+CONSTANT_VK u32a c_tables[4][256] =
 {
   {
     0x00072000, 0x00075000, 0x00074800, 0x00071000,
diff --git a/OpenCL/m06900_a3-optimized.cl b/OpenCL/m06900_a3-optimized.cl
index abfcda648..4f201a4c1 100644
--- a/OpenCL/m06900_a3-optimized.cl
+++ b/OpenCL/m06900_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a c_tables[4][256] =
+CONSTANT_VK u32a c_tables[4][256] =
 {
   {
     0x00072000, 0x00075000, 0x00074800, 0x00071000,
diff --git a/OpenCL/m07700_a0-optimized.cl b/OpenCL/m07700_a0-optimized.cl
index 824307804..2c633028b 100644
--- a/OpenCL/m07700_a0-optimized.cl
+++ b/OpenCL/m07700_a0-optimized.cl
@@ -29,7 +29,7 @@
   (a)[((n)/4)+1]  = x >> 32;        \
 }
 
-CONSTANT_AS u32a sapb_trans_tbl[256] =
+CONSTANT_VK u32a sapb_trans_tbl[256] =
 {
   // first value hack for 0 byte as part of an optimization
   0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@@ -50,7 +50,7 @@ CONSTANT_AS u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-CONSTANT_AS u32a bcodeArray[48] =
+CONSTANT_VK u32a bcodeArray[48] =
 {
   0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
   0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
diff --git a/OpenCL/m07700_a1-optimized.cl b/OpenCL/m07700_a1-optimized.cl
index 46877fd70..525d9c257 100644
--- a/OpenCL/m07700_a1-optimized.cl
+++ b/OpenCL/m07700_a1-optimized.cl
@@ -27,7 +27,7 @@
   (a)[((n)/4)+1]  = x >> 32;        \
 }
 
-CONSTANT_AS u32a sapb_trans_tbl[256] =
+CONSTANT_VK u32a sapb_trans_tbl[256] =
 {
   // first value hack for 0 byte as part of an optimization
   0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@@ -48,7 +48,7 @@ CONSTANT_AS u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-CONSTANT_AS u32a bcodeArray[48] =
+CONSTANT_VK u32a bcodeArray[48] =
 {
   0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
   0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
diff --git a/OpenCL/m07700_a3-optimized.cl b/OpenCL/m07700_a3-optimized.cl
index da3865e91..78927041b 100644
--- a/OpenCL/m07700_a3-optimized.cl
+++ b/OpenCL/m07700_a3-optimized.cl
@@ -18,7 +18,7 @@
 #define GETCHAR(a,p)  (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
 
-CONSTANT_AS u32a sapb_trans_tbl[256] =
+CONSTANT_VK u32a sapb_trans_tbl[256] =
 {
   // first value hack for 0 byte as part of an optimization
   0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@@ -39,7 +39,7 @@ CONSTANT_AS u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-CONSTANT_AS u32a bcodeArray[48] =
+CONSTANT_VK u32a bcodeArray[48] =
 {
   0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
   0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
diff --git a/OpenCL/m07701_a0-optimized.cl b/OpenCL/m07701_a0-optimized.cl
index e20c9c505..99fbb8183 100644
--- a/OpenCL/m07701_a0-optimized.cl
+++ b/OpenCL/m07701_a0-optimized.cl
@@ -29,7 +29,7 @@
   (a)[((n)/4)+1]  = x >> 32;        \
 }
 
-CONSTANT_AS u32a sapb_trans_tbl[256] =
+CONSTANT_VK u32a sapb_trans_tbl[256] =
 {
   // first value hack for 0 byte as part of an optimization
   0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@@ -50,7 +50,7 @@ CONSTANT_AS u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-CONSTANT_AS u32a bcodeArray[48] =
+CONSTANT_VK u32a bcodeArray[48] =
 {
   0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
   0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
diff --git a/OpenCL/m07701_a1-optimized.cl b/OpenCL/m07701_a1-optimized.cl
index e403114ab..04609d8e0 100644
--- a/OpenCL/m07701_a1-optimized.cl
+++ b/OpenCL/m07701_a1-optimized.cl
@@ -27,7 +27,7 @@
   (a)[((n)/4)+1]  = x >> 32;        \
 }
 
-CONSTANT_AS u32a sapb_trans_tbl[256] =
+CONSTANT_VK u32a sapb_trans_tbl[256] =
 {
   // first value hack for 0 byte as part of an optimization
   0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@@ -48,7 +48,7 @@ CONSTANT_AS u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-CONSTANT_AS u32a bcodeArray[48] =
+CONSTANT_VK u32a bcodeArray[48] =
 {
   0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
   0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
diff --git a/OpenCL/m07701_a3-optimized.cl b/OpenCL/m07701_a3-optimized.cl
index 5dad95087..506ca491f 100644
--- a/OpenCL/m07701_a3-optimized.cl
+++ b/OpenCL/m07701_a3-optimized.cl
@@ -18,7 +18,7 @@
 #define GETCHAR(a,p)  (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
 
-CONSTANT_AS u32a sapb_trans_tbl[256] =
+CONSTANT_VK u32a sapb_trans_tbl[256] =
 {
   // first value hack for 0 byte as part of an optimization
   0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
@@ -39,7 +39,7 @@ CONSTANT_AS u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-CONSTANT_AS u32a bcodeArray[48] =
+CONSTANT_VK u32a bcodeArray[48] =
 {
   0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
   0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
diff --git a/OpenCL/m07800_a0-optimized.cl b/OpenCL/m07800_a0-optimized.cl
index 720d68ad0..4d86b2858 100644
--- a/OpenCL/m07800_a0-optimized.cl
+++ b/OpenCL/m07800_a0-optimized.cl
@@ -17,7 +17,7 @@
 #include "inc_hash_sha1.cl"
 #endif
 
-CONSTANT_AS u32a theMagicArray[64] =
+CONSTANT_VK u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
diff --git a/OpenCL/m07800_a1-optimized.cl b/OpenCL/m07800_a1-optimized.cl
index b2c8358fb..cc552143f 100644
--- a/OpenCL/m07800_a1-optimized.cl
+++ b/OpenCL/m07800_a1-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_hash_sha1.cl"
 #endif
 
-CONSTANT_AS u32a theMagicArray[64] =
+CONSTANT_VK u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
diff --git a/OpenCL/m07800_a3-optimized.cl b/OpenCL/m07800_a3-optimized.cl
index 55065786d..de22475f0 100644
--- a/OpenCL/m07800_a3-optimized.cl
+++ b/OpenCL/m07800_a3-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_hash_sha1.cl"
 #endif
 
-CONSTANT_AS u32a theMagicArray[64] =
+CONSTANT_VK u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
diff --git a/OpenCL/m07801_a0-optimized.cl b/OpenCL/m07801_a0-optimized.cl
index 0ebdf2e9a..aeb9b95d0 100644
--- a/OpenCL/m07801_a0-optimized.cl
+++ b/OpenCL/m07801_a0-optimized.cl
@@ -17,7 +17,7 @@
 #include "inc_hash_sha1.cl"
 #endif
 
-CONSTANT_AS u32a theMagicArray[64] =
+CONSTANT_VK u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
diff --git a/OpenCL/m07801_a1-optimized.cl b/OpenCL/m07801_a1-optimized.cl
index 60829f3f7..8d055619c 100644
--- a/OpenCL/m07801_a1-optimized.cl
+++ b/OpenCL/m07801_a1-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_hash_sha1.cl"
 #endif
 
-CONSTANT_AS u32a theMagicArray[64] =
+CONSTANT_VK u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
diff --git a/OpenCL/m07801_a3-optimized.cl b/OpenCL/m07801_a3-optimized.cl
index 0a346ca8d..7de9f2a11 100644
--- a/OpenCL/m07801_a3-optimized.cl
+++ b/OpenCL/m07801_a3-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_hash_sha1.cl"
 #endif
 
-CONSTANT_AS u32a theMagicArray[64] =
+CONSTANT_VK u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
diff --git a/OpenCL/m08500_a0-pure.cl b/OpenCL/m08500_a0-pure.cl
index 69c768de3..e64afc6cc 100644
--- a/OpenCL/m08500_a0-pure.cl
+++ b/OpenCL/m08500_a0-pure.cl
@@ -53,7 +53,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_ascii_to_ebcdic_pc[256] =
+CONSTANT_VK u32a c_ascii_to_ebcdic_pc[256] =
 {
   // little hack, can't crack 0-bytes in password, but who cares
   //    0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
@@ -75,7 +75,7 @@ CONSTANT_AS u32a c_ascii_to_ebcdic_pc[256] =
   0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
 };
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -223,7 +223,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m08500_a1-pure.cl b/OpenCL/m08500_a1-pure.cl
index 82df98f55..1a4cf0a23 100644
--- a/OpenCL/m08500_a1-pure.cl
+++ b/OpenCL/m08500_a1-pure.cl
@@ -51,7 +51,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_ascii_to_ebcdic_pc[256] =
+CONSTANT_VK u32a c_ascii_to_ebcdic_pc[256] =
 {
   // little hack, can't crack 0-bytes in password, but who cares
   //    0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
@@ -73,7 +73,7 @@ CONSTANT_AS u32a c_ascii_to_ebcdic_pc[256] =
   0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
 };
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -221,7 +221,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m08500_a3-pure.cl b/OpenCL/m08500_a3-pure.cl
index 277295e7d..2ddb4efc8 100644
--- a/OpenCL/m08500_a3-pure.cl
+++ b/OpenCL/m08500_a3-pure.cl
@@ -51,7 +51,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_ascii_to_ebcdic_pc[256] =
+CONSTANT_VK u32a c_ascii_to_ebcdic_pc[256] =
 {
   // little hack, can't crack 0-bytes in password, but who cares
   //    0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
@@ -73,7 +73,7 @@ CONSTANT_AS u32a c_ascii_to_ebcdic_pc[256] =
   0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
 };
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -221,7 +221,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m08600_a0-pure.cl b/OpenCL/m08600_a0-pure.cl
index 22ac6cf46..04c9a3594 100644
--- a/OpenCL/m08600_a0-pure.cl
+++ b/OpenCL/m08600_a0-pure.cl
@@ -16,7 +16,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a lotus_magic_table[256] =
+CONSTANT_VK u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
diff --git a/OpenCL/m08600_a1-pure.cl b/OpenCL/m08600_a1-pure.cl
index 859278a54..bb081605d 100644
--- a/OpenCL/m08600_a1-pure.cl
+++ b/OpenCL/m08600_a1-pure.cl
@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a lotus_magic_table[256] =
+CONSTANT_VK u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
diff --git a/OpenCL/m08600_a3-pure.cl b/OpenCL/m08600_a3-pure.cl
index da94d4a3c..b7b2f29df 100644
--- a/OpenCL/m08600_a3-pure.cl
+++ b/OpenCL/m08600_a3-pure.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a lotus_magic_table[256] =
+CONSTANT_VK u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
diff --git a/OpenCL/m08700_a0-optimized.cl b/OpenCL/m08700_a0-optimized.cl
index 022281919..add1bba63 100644
--- a/OpenCL/m08700_a0-optimized.cl
+++ b/OpenCL/m08700_a0-optimized.cl
@@ -16,7 +16,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a lotus_magic_table[256] =
+CONSTANT_VK u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
diff --git a/OpenCL/m08700_a1-optimized.cl b/OpenCL/m08700_a1-optimized.cl
index 94968476f..a8fc8adf3 100644
--- a/OpenCL/m08700_a1-optimized.cl
+++ b/OpenCL/m08700_a1-optimized.cl
@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a lotus_magic_table[256] =
+CONSTANT_VK u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
diff --git a/OpenCL/m08700_a3-optimized.cl b/OpenCL/m08700_a3-optimized.cl
index 694a58886..07eaf035c 100644
--- a/OpenCL/m08700_a3-optimized.cl
+++ b/OpenCL/m08700_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a lotus_magic_table[256] =
+CONSTANT_VK u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
diff --git a/OpenCL/m09000-pure.cl b/OpenCL/m09000-pure.cl
index ec34faab2..e04a8594b 100644
--- a/OpenCL/m09000-pure.cl
+++ b/OpenCL/m09000-pure.cl
@@ -29,7 +29,7 @@ typedef struct pwsafe2_tmp
 
 // http://www.schneier.com/code/constants.txt
 
-CONSTANT_AS u32a c_sbox0[256] =
+CONSTANT_VK u32a c_sbox0[256] =
 {
   0xd1310ba6, 0x98dfb5ac, 0x2ffd72db, 0xd01adfb7,
   0xb8e1afed, 0x6a267e96, 0xba7c9045, 0xf12c7f99,
@@ -97,7 +97,7 @@ CONSTANT_AS u32a c_sbox0[256] =
   0x53b02d5d, 0xa99f8fa1, 0x08ba4799, 0x6e85076a
 };
 
-CONSTANT_AS u32a c_sbox1[256] =
+CONSTANT_VK u32a c_sbox1[256] =
 {
   0x4b7a70e9, 0xb5b32944, 0xdb75092e, 0xc4192623,
   0xad6ea6b0, 0x49a7df7d, 0x9cee60b8, 0x8fedb266,
@@ -165,7 +165,7 @@ CONSTANT_AS u32a c_sbox1[256] =
   0x153e21e7, 0x8fb03d4a, 0xe6e39f2b, 0xdb83adf7
 };
 
-CONSTANT_AS u32a c_sbox2[256] =
+CONSTANT_VK u32a c_sbox2[256] =
 {
   0xe93d5a68, 0x948140f7, 0xf64c261c, 0x94692934,
   0x411520f7, 0x7602d4f7, 0xbcf46b2e, 0xd4a20068,
@@ -233,7 +233,7 @@ CONSTANT_AS u32a c_sbox2[256] =
   0xd79a3234, 0x92638212, 0x670efa8e, 0x406000e0
 };
 
-CONSTANT_AS u32a c_sbox3[256] =
+CONSTANT_VK u32a c_sbox3[256] =
 {
   0x3a39ce37, 0xd3faf5cf, 0xabc27737, 0x5ac52d1b,
   0x5cb0679e, 0x4fa33742, 0xd3822740, 0x99bc9bbe,
@@ -301,7 +301,7 @@ CONSTANT_AS u32a c_sbox3[256] =
   0xb74e6132, 0xce77e25b, 0x578fdfe3, 0x3ac372e6
 };
 
-CONSTANT_AS u32a c_pbox[18] =
+CONSTANT_VK u32a c_pbox[18] =
 {
   0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344,
   0xa4093822, 0x299f31d0, 0x082efa98, 0xec4e6c89,
diff --git a/OpenCL/m09100-pure.cl b/OpenCL/m09100-pure.cl
index 51ae77536..004bb3e85 100644
--- a/OpenCL/m09100-pure.cl
+++ b/OpenCL/m09100-pure.cl
@@ -27,7 +27,7 @@ typedef struct lotus8_tmp
 
 } lotus8_tmp_t;
 
-CONSTANT_AS u32a lotus64_table[64] =
+CONSTANT_VK u32a lotus64_table[64] =
 {
   '0', '1', '2', '3', '4', '5', '6', '7',
   '8', '9', 'A', 'B', 'C', 'D', 'E', 'F',
@@ -39,7 +39,7 @@ CONSTANT_AS u32a lotus64_table[64] =
   'u', 'v', 'w', 'x', 'y', 'z', '+', '/',
 };
 
-CONSTANT_AS u32a lotus_magic_table[256] =
+CONSTANT_VK u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
diff --git a/OpenCL/m10400_a0-optimized.cl b/OpenCL/m10400_a0-optimized.cl
index 4da62a07d..fbd378c0b 100644
--- a/OpenCL/m10400_a0-optimized.cl
+++ b/OpenCL/m10400_a0-optimized.cl
@@ -17,7 +17,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10400_a1-optimized.cl b/OpenCL/m10400_a1-optimized.cl
index f1662f380..5ad30d180 100644
--- a/OpenCL/m10400_a1-optimized.cl
+++ b/OpenCL/m10400_a1-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10400_a3-optimized.cl b/OpenCL/m10400_a3-optimized.cl
index 1182a6067..5f134e4d5 100644
--- a/OpenCL/m10400_a3-optimized.cl
+++ b/OpenCL/m10400_a3-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10410_a0-optimized.cl b/OpenCL/m10410_a0-optimized.cl
index 5bbc1c9d7..ede8b87ea 100644
--- a/OpenCL/m10410_a0-optimized.cl
+++ b/OpenCL/m10410_a0-optimized.cl
@@ -17,7 +17,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10410_a1-optimized.cl b/OpenCL/m10410_a1-optimized.cl
index b480d6321..493638f0b 100644
--- a/OpenCL/m10410_a1-optimized.cl
+++ b/OpenCL/m10410_a1-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10410_a3-optimized.cl b/OpenCL/m10410_a3-optimized.cl
index 6bd44c809..0513d4053 100644
--- a/OpenCL/m10410_a3-optimized.cl
+++ b/OpenCL/m10410_a3-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10420_a0-optimized.cl b/OpenCL/m10420_a0-optimized.cl
index 2766e028a..799be19d9 100644
--- a/OpenCL/m10420_a0-optimized.cl
+++ b/OpenCL/m10420_a0-optimized.cl
@@ -16,7 +16,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10420_a1-optimized.cl b/OpenCL/m10420_a1-optimized.cl
index 3fe2fba4e..9f9f26a7f 100644
--- a/OpenCL/m10420_a1-optimized.cl
+++ b/OpenCL/m10420_a1-optimized.cl
@@ -14,7 +14,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10420_a3-optimized.cl b/OpenCL/m10420_a3-optimized.cl
index 47d0f1252..7162048c0 100644
--- a/OpenCL/m10420_a3-optimized.cl
+++ b/OpenCL/m10420_a3-optimized.cl
@@ -14,7 +14,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m10500-pure.cl b/OpenCL/m10500-pure.cl
index ef270689a..844382357 100644
--- a/OpenCL/m10500-pure.cl
+++ b/OpenCL/m10500-pure.cl
@@ -14,7 +14,7 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
-CONSTANT_AS u32a padding[8] =
+CONSTANT_VK u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
diff --git a/OpenCL/m11500_a0-optimized.cl b/OpenCL/m11500_a0-optimized.cl
index 903190d74..9a77848e9 100644
--- a/OpenCL/m11500_a0-optimized.cl
+++ b/OpenCL/m11500_a0-optimized.cl
@@ -16,7 +16,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a crc32tab[0x100] =
+CONSTANT_VK u32a crc32tab[0x100] =
 {
   0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
   0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
diff --git a/OpenCL/m11500_a1-optimized.cl b/OpenCL/m11500_a1-optimized.cl
index 6ee1b5b6c..06a8b1ebb 100644
--- a/OpenCL/m11500_a1-optimized.cl
+++ b/OpenCL/m11500_a1-optimized.cl
@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a crc32tab[0x100] =
+CONSTANT_VK u32a crc32tab[0x100] =
 {
   0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
   0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
diff --git a/OpenCL/m11500_a3-optimized.cl b/OpenCL/m11500_a3-optimized.cl
index 2fd20e7bf..698b93325 100644
--- a/OpenCL/m11500_a3-optimized.cl
+++ b/OpenCL/m11500_a3-optimized.cl
@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u32a crc32tab[0x100] =
+CONSTANT_VK u32a crc32tab[0x100] =
 {
   0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
   0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
diff --git a/OpenCL/m12400-pure.cl b/OpenCL/m12400-pure.cl
index 4c3cf8fe7..43002fca5 100644
--- a/OpenCL/m12400-pure.cl
+++ b/OpenCL/m12400-pure.cl
@@ -60,7 +60,7 @@ typedef struct bsdicrypt_tmp
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x00820200, 0x00020000, 0x80800000, 0x80820200,
@@ -208,7 +208,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m14000_a0-pure.cl b/OpenCL/m14000_a0-pure.cl
index 503e26d28..5a8a3c39c 100644
--- a/OpenCL/m14000_a0-pure.cl
+++ b/OpenCL/m14000_a0-pure.cl
@@ -53,7 +53,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     /* nibble 0 */
@@ -209,7 +209,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m14000_a1-pure.cl b/OpenCL/m14000_a1-pure.cl
index 90954adae..f8f55cf2e 100644
--- a/OpenCL/m14000_a1-pure.cl
+++ b/OpenCL/m14000_a1-pure.cl
@@ -51,7 +51,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -199,7 +199,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m14100_a0-pure.cl b/OpenCL/m14100_a0-pure.cl
index 17d5bed9d..42b06753b 100644
--- a/OpenCL/m14100_a0-pure.cl
+++ b/OpenCL/m14100_a0-pure.cl
@@ -53,7 +53,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     /* nibble 0 */
@@ -209,7 +209,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m14100_a1-pure.cl b/OpenCL/m14100_a1-pure.cl
index 3eba5678c..f44756c2b 100644
--- a/OpenCL/m14100_a1-pure.cl
+++ b/OpenCL/m14100_a1-pure.cl
@@ -51,7 +51,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -199,7 +199,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m14100_a3-pure.cl b/OpenCL/m14100_a3-pure.cl
index 616d944d3..854e1c36b 100644
--- a/OpenCL/m14100_a3-pure.cl
+++ b/OpenCL/m14100_a3-pure.cl
@@ -51,7 +51,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -199,7 +199,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   }
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
diff --git a/OpenCL/m14900_a0-optimized.cl b/OpenCL/m14900_a0-optimized.cl
index f8946baa1..16cf26446 100644
--- a/OpenCL/m14900_a0-optimized.cl
+++ b/OpenCL/m14900_a0-optimized.cl
@@ -16,7 +16,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u8a c_ftable[256] =
+CONSTANT_VK u8a c_ftable[256] =
 {
   0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4,
   0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9,
diff --git a/OpenCL/m14900_a1-optimized.cl b/OpenCL/m14900_a1-optimized.cl
index f672860fe..b9753524c 100644
--- a/OpenCL/m14900_a1-optimized.cl
+++ b/OpenCL/m14900_a1-optimized.cl
@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u8a c_ftable[256] =
+CONSTANT_VK u8a c_ftable[256] =
 {
   0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4,
   0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9,
diff --git a/OpenCL/m14900_a3-optimized.cl b/OpenCL/m14900_a3-optimized.cl
index 7255cfcec..a64f22df7 100644
--- a/OpenCL/m14900_a3-optimized.cl
+++ b/OpenCL/m14900_a3-optimized.cl
@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u8a c_ftable[256] =
+CONSTANT_VK u8a c_ftable[256] =
 {
   0xa3, 0xd7, 0x09, 0x83, 0xf8, 0x48, 0xf6, 0xf4,
   0xb3, 0x21, 0x15, 0x78, 0x99, 0xb1, 0xaf, 0xf9,
diff --git a/OpenCL/m15600-pure.cl b/OpenCL/m15600-pure.cl
index 05652abf7..e6bad07b5 100644
--- a/OpenCL/m15600-pure.cl
+++ b/OpenCL/m15600-pure.cl
@@ -34,7 +34,7 @@ typedef struct ethereum_pbkdf2
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m15700-pure.cl b/OpenCL/m15700-pure.cl
index 15bc64b58..3e12e7100 100644
--- a/OpenCL/m15700-pure.cl
+++ b/OpenCL/m15700-pure.cl
@@ -261,7 +261,7 @@ DECLSPEC void scrypt_smix (uint4 *X, uint4 *T, GLOBAL_AS uint4 *V0, GLOBAL_AS ui
   st[4 + s] ^= ~bc0 & bc1;      \
 }
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m16000_a0-pure.cl b/OpenCL/m16000_a0-pure.cl
index 991e033a0..c78cbff01 100644
--- a/OpenCL/m16000_a0-pure.cl
+++ b/OpenCL/m16000_a0-pure.cl
@@ -35,7 +35,7 @@
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x00820200, 0x00020000, 0x80800000, 0x80820200,
@@ -183,7 +183,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
@@ -331,7 +331,7 @@ CONSTANT_AS u32a c_skb[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_tripcode_salt[128] =
+CONSTANT_VK u32a c_tripcode_salt[128] =
 {
   0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
   0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
diff --git a/OpenCL/m16000_a1-pure.cl b/OpenCL/m16000_a1-pure.cl
index 994e30ace..578021fe8 100644
--- a/OpenCL/m16000_a1-pure.cl
+++ b/OpenCL/m16000_a1-pure.cl
@@ -33,7 +33,7 @@
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x00820200, 0x00020000, 0x80800000, 0x80820200,
@@ -181,7 +181,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
@@ -329,7 +329,7 @@ CONSTANT_AS u32a c_skb[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_tripcode_salt[128] =
+CONSTANT_VK u32a c_tripcode_salt[128] =
 {
   0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
   0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
diff --git a/OpenCL/m16000_a3-pure.cl b/OpenCL/m16000_a3-pure.cl
index 17361374f..12e73ed34 100644
--- a/OpenCL/m16000_a3-pure.cl
+++ b/OpenCL/m16000_a3-pure.cl
@@ -33,7 +33,7 @@
   a  = a ^ tt;              \
 }
 
-CONSTANT_AS u32a c_SPtrans[8][64] =
+CONSTANT_VK u32a c_SPtrans[8][64] =
 {
   {
     0x00820200, 0x00020000, 0x80800000, 0x80820200,
@@ -181,7 +181,7 @@ CONSTANT_AS u32a c_SPtrans[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_skb[8][64] =
+CONSTANT_VK u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,
@@ -329,7 +329,7 @@ CONSTANT_AS u32a c_skb[8][64] =
   },
 };
 
-CONSTANT_AS u32a c_tripcode_salt[128] =
+CONSTANT_VK u32a c_tripcode_salt[128] =
 {
   0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
   0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
diff --git a/OpenCL/m16300-pure.cl b/OpenCL/m16300-pure.cl
index fa3d51042..5201a8ccc 100644
--- a/OpenCL/m16300-pure.cl
+++ b/OpenCL/m16300-pure.cl
@@ -36,7 +36,7 @@ typedef struct pbkdf2_sha256_tmp
 
 } pbkdf2_sha256_tmp_t;
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17300_a0-optimized.cl b/OpenCL/m17300_a0-optimized.cl
index 011de4eb5..a1c7bd237 100644
--- a/OpenCL/m17300_a0-optimized.cl
+++ b/OpenCL/m17300_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17300_a1-optimized.cl b/OpenCL/m17300_a1-optimized.cl
index d110afb2d..cc3c5ac5b 100644
--- a/OpenCL/m17300_a1-optimized.cl
+++ b/OpenCL/m17300_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17300_a3-optimized.cl b/OpenCL/m17300_a3-optimized.cl
index 58898f5cb..cf78c55e6 100644
--- a/OpenCL/m17300_a3-optimized.cl
+++ b/OpenCL/m17300_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17400_a0-optimized.cl b/OpenCL/m17400_a0-optimized.cl
index a157070f0..8dbefbd72 100644
--- a/OpenCL/m17400_a0-optimized.cl
+++ b/OpenCL/m17400_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17400_a1-optimized.cl b/OpenCL/m17400_a1-optimized.cl
index 84b14a2b8..96c99c25b 100644
--- a/OpenCL/m17400_a1-optimized.cl
+++ b/OpenCL/m17400_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17400_a3-optimized.cl b/OpenCL/m17400_a3-optimized.cl
index fd2977b51..1ce91475c 100644
--- a/OpenCL/m17400_a3-optimized.cl
+++ b/OpenCL/m17400_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17500_a0-optimized.cl b/OpenCL/m17500_a0-optimized.cl
index dd62c6956..10a029a81 100644
--- a/OpenCL/m17500_a0-optimized.cl
+++ b/OpenCL/m17500_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17500_a1-optimized.cl b/OpenCL/m17500_a1-optimized.cl
index e6a04db79..3a62ba394 100644
--- a/OpenCL/m17500_a1-optimized.cl
+++ b/OpenCL/m17500_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17500_a3-optimized.cl b/OpenCL/m17500_a3-optimized.cl
index 9c6a3c99b..91705b5c4 100644
--- a/OpenCL/m17500_a3-optimized.cl
+++ b/OpenCL/m17500_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17600_a0-optimized.cl b/OpenCL/m17600_a0-optimized.cl
index bf9c3efb6..0b32f69b5 100644
--- a/OpenCL/m17600_a0-optimized.cl
+++ b/OpenCL/m17600_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17600_a1-optimized.cl b/OpenCL/m17600_a1-optimized.cl
index a605df2a8..5d5117736 100644
--- a/OpenCL/m17600_a1-optimized.cl
+++ b/OpenCL/m17600_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17600_a3-optimized.cl b/OpenCL/m17600_a3-optimized.cl
index aed11f579..7dde3b61a 100644
--- a/OpenCL/m17600_a3-optimized.cl
+++ b/OpenCL/m17600_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17700_a0-optimized.cl b/OpenCL/m17700_a0-optimized.cl
index 3bc5f965f..00355e29a 100644
--- a/OpenCL/m17700_a0-optimized.cl
+++ b/OpenCL/m17700_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17700_a1-optimized.cl b/OpenCL/m17700_a1-optimized.cl
index 5dcf2aff7..88527e9b6 100644
--- a/OpenCL/m17700_a1-optimized.cl
+++ b/OpenCL/m17700_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17700_a3-optimized.cl b/OpenCL/m17700_a3-optimized.cl
index bc912a045..70a3f42d0 100644
--- a/OpenCL/m17700_a3-optimized.cl
+++ b/OpenCL/m17700_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17800_a0-optimized.cl b/OpenCL/m17800_a0-optimized.cl
index dc4cffadd..193f9ce6e 100644
--- a/OpenCL/m17800_a0-optimized.cl
+++ b/OpenCL/m17800_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17800_a1-optimized.cl b/OpenCL/m17800_a1-optimized.cl
index 8733184de..cc1e2ca52 100644
--- a/OpenCL/m17800_a1-optimized.cl
+++ b/OpenCL/m17800_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17800_a3-optimized.cl b/OpenCL/m17800_a3-optimized.cl
index f2d497e36..3de33b43a 100644
--- a/OpenCL/m17800_a3-optimized.cl
+++ b/OpenCL/m17800_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17900_a0-optimized.cl b/OpenCL/m17900_a0-optimized.cl
index 5ed128657..1e9e68e9d 100644
--- a/OpenCL/m17900_a0-optimized.cl
+++ b/OpenCL/m17900_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17900_a1-optimized.cl b/OpenCL/m17900_a1-optimized.cl
index d6518aeb8..e26e19987 100644
--- a/OpenCL/m17900_a1-optimized.cl
+++ b/OpenCL/m17900_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m17900_a3-optimized.cl b/OpenCL/m17900_a3-optimized.cl
index 43d2791d9..db78824ad 100644
--- a/OpenCL/m17900_a3-optimized.cl
+++ b/OpenCL/m17900_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m18000_a0-optimized.cl b/OpenCL/m18000_a0-optimized.cl
index 1d587b1db..6a2f4eeb4 100644
--- a/OpenCL/m18000_a0-optimized.cl
+++ b/OpenCL/m18000_a0-optimized.cl
@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m18000_a1-optimized.cl b/OpenCL/m18000_a1-optimized.cl
index 1964d037a..36caf6052 100644
--- a/OpenCL/m18000_a1-optimized.cl
+++ b/OpenCL/m18000_a1-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m18000_a3-optimized.cl b/OpenCL/m18000_a3-optimized.cl
index a30de8b7e..2dee6304a 100644
--- a/OpenCL/m18000_a3-optimized.cl
+++ b/OpenCL/m18000_a3-optimized.cl
@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #endif
 
-CONSTANT_AS u64a keccakf_rndc[24] =
+CONSTANT_VK u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,
diff --git a/OpenCL/m18600-pure.cl b/OpenCL/m18600-pure.cl
index f68d6f82f..2f8e0ea18 100644
--- a/OpenCL/m18600-pure.cl
+++ b/OpenCL/m18600-pure.cl
@@ -38,7 +38,7 @@ typedef struct odf11
 
 // http://www.schneier.com/code/constants.txt
 
-CONSTANT_AS u32a c_sbox0[256] =
+CONSTANT_VK u32a c_sbox0[256] =
 {
   0xd1310ba6, 0x98dfb5ac, 0x2ffd72db, 0xd01adfb7,
   0xb8e1afed, 0x6a267e96, 0xba7c9045, 0xf12c7f99,
@@ -106,7 +106,7 @@ CONSTANT_AS u32a c_sbox0[256] =
   0x53b02d5d, 0xa99f8fa1, 0x08ba4799, 0x6e85076a
 };
 
-CONSTANT_AS u32a c_sbox1[256] =
+CONSTANT_VK u32a c_sbox1[256] =
 {
   0x4b7a70e9, 0xb5b32944, 0xdb75092e, 0xc4192623,
   0xad6ea6b0, 0x49a7df7d, 0x9cee60b8, 0x8fedb266,
@@ -174,7 +174,7 @@ CONSTANT_AS u32a c_sbox1[256] =
   0x153e21e7, 0x8fb03d4a, 0xe6e39f2b, 0xdb83adf7
 };
 
-CONSTANT_AS u32a c_sbox2[256] =
+CONSTANT_VK u32a c_sbox2[256] =
 {
   0xe93d5a68, 0x948140f7, 0xf64c261c, 0x94692934,
   0x411520f7, 0x7602d4f7, 0xbcf46b2e, 0xd4a20068,
@@ -242,7 +242,7 @@ CONSTANT_AS u32a c_sbox2[256] =
   0xd79a3234, 0x92638212, 0x670efa8e, 0x406000e0
 };
 
-CONSTANT_AS u32a c_sbox3[256] =
+CONSTANT_VK u32a c_sbox3[256] =
 {
   0x3a39ce37, 0xd3faf5cf, 0xabc27737, 0x5ac52d1b,
   0x5cb0679e, 0x4fa33742, 0xd3822740, 0x99bc9bbe,
@@ -310,7 +310,7 @@ CONSTANT_AS u32a c_sbox3[256] =
   0xb74e6132, 0xce77e25b, 0x578fdfe3, 0x3ac372e6
 };
 
-CONSTANT_AS u32a c_pbox[18] =
+CONSTANT_VK u32a c_pbox[18] =
 {
   0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344,
   0xa4093822, 0x299f31d0, 0x082efa98, 0xec4e6c89,
diff --git a/src/backend.c b/src/backend.c
index 9a162c150..eee101e6c 100644
--- a/src/backend.c
+++ b/src/backend.c
@@ -3166,14 +3166,14 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con
       if (rc_cuEventRecord2 == -1) return -1;
     }
 
-    const int rc_cuEventSynchronize = hc_cuEventSynchronize (hashcat_ctx, device_param->cuda_event2);
-
-    if (rc_cuEventSynchronize == -1) return -1;
-
     const int rc_cuStreamSynchronize = hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream);
 
     if (rc_cuStreamSynchronize == -1) return -1;
 
+    const int rc_cuEventSynchronize = hc_cuEventSynchronize (hashcat_ctx, device_param->cuda_event2);
+
+    if (rc_cuEventSynchronize == -1) return -1;
+
     float exec_ms;
 
     const int rc_cuEventElapsedTime = hc_cuEventElapsedTime (hashcat_ctx, &exec_ms, device_param->cuda_event1, device_param->cuda_event2);