Merge pull request #1 from hashcat/master

Merge out to bring fork up to date

.appveyor.yml

@@ -1,7 +1,6 @@
 environment:
   CYG_MIRROR: http://cygwin.mirror.constant.com
   CYG_PACKAGES: make,gcc-core,libiconv-devel
-  DOCUMENT_FOLDER: /share/doc/hashcat
   matrix:
     - CYG_ROOT: C:\cygwin64
       CYG_CACHE: C:\cygwin64\var\cache\setup
@@ -15,12 +14,10 @@ environment:
       CC: gcc
     - MSYSTEM: MINGW64
       MSYS_CACHE: C:\msys64\var\cache\pacman\pkg
-      KERNEL_CACHE: C:\msys64\usr\local\bin\OpenCL\
       BASH: C:\msys64\usr\bin\bash
       CC: gcc
     - MSYSTEM: MINGW32
       MSYS_CACHE: C:\msys64\var\cache\pacman\pkg
-      KERNEL_CACHE: C:\msys64\usr\local\bin\OpenCL\
       BASH: C:\msys64\usr\bin\bash
       CC: gcc
 
@@ -42,23 +39,18 @@ install:
   - if defined MSYSTEM (%BASH% -lc "pacman -Suuy --noconfirm")
 
 build_script:
-  - if defined BASH (%BASH% -lc "cd $(cygpath ${APPVEYOR_BUILD_FOLDER}) && git submodule update --init && make install")
+  - if defined BASH (%BASH% -lc "cd $(cygpath ${APPVEYOR_BUILD_FOLDER}) && git submodule update --init && make")
 
 test_script:
   # some file globbing tests
   #  1. hash file should not exist and therefore hashcat should complain (if it does not there might be a problem)
   #  2. hash file should expand to example0.hash and succeed
   - ps: >-
-      if (Test-Path Env:\MSYSTEM)
+      & $env:BASH -lc "cd '$env:APPVEYOR_BUILD_FOLDER' && ./hashcat.exe -m 0 --show *file_not_found.hash" 2>&1 | out-null
-      {
-        mkdir $env:KERNEL_CACHE 2>&1 | out-null
-      }
-
-      & $env:BASH -lc "hashcat.exe -m 0 --show $env:DOCUMENT_FOLDER/*file_not_found.hash" 2>&1 | out-null
 
       if ($LastExitCode -eq 0)
       {
         throw "test failed"
       }
 
-      & $env:BASH -lc "hashcat.exe -m 0 --show $env:DOCUMENT_FOLDER/*ple0.hash"
+      & $env:BASH -lc "cd '$env:APPVEYOR_BUILD_FOLDER' && ./hashcat.exe -m 0 --show *ple0.hash"

OpenCL/inc_cipher_aes.cl

@@ -3,7 +3,7 @@
  * License.....: MIT
  */
 
-__constant u32a te0[256] =
+__constant static u32a te0[256] =
 {
   0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
   0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
@@ -71,7 +71,7 @@ __constant u32a te0[256] =
   0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
 };
 
-__constant u32a te1[256] =
+__constant static u32a te1[256] =
 {
   0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
   0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
@@ -139,7 +139,7 @@ __constant u32a te1[256] =
   0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
 };
 
-__constant u32a te2[256] =
+__constant static u32a te2[256] =
 {
   0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
   0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
@@ -207,7 +207,7 @@ __constant u32a te2[256] =
   0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
 };
 
-__constant u32a te3[256] =
+__constant static u32a te3[256] =
 {
   0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
   0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
@@ -275,7 +275,7 @@ __constant u32a te3[256] =
   0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
 };
 
-__constant u32a te4[256] =
+__constant static u32a te4[256] =
 {
   0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
   0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
@@ -343,7 +343,7 @@ __constant u32a te4[256] =
   0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
 };
 
-__constant u32a td0[256] =
+__constant static u32a td0[256] =
 {
   0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
   0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
@@ -411,7 +411,7 @@ __constant u32a td0[256] =
   0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
 };
 
-__constant u32a td1[256] =
+__constant static u32a td1[256] =
 {
   0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
   0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
@@ -479,7 +479,7 @@ __constant u32a td1[256] =
   0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
 };
 
-__constant u32a td2[256] =
+__constant static u32a td2[256] =
 {
   0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
   0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
@@ -547,7 +547,7 @@ __constant u32a td2[256] =
   0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
 };
 
-__constant u32a td3[256] =
+__constant static u32a td3[256] =
 {
   0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
   0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
@@ -615,7 +615,7 @@ __constant u32a td3[256] =
   0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
 };
 
-__constant u32a td4[256] =
+__constant static u32a td4[256] =
 {
   0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
   0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
@@ -683,7 +683,7 @@ __constant u32a td4[256] =
   0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
 };
 
-__constant u32a rcon[] =
+__constant static u32a rcon[] =
 {
   0x01000000, 0x02000000, 0x04000000, 0x08000000,
   0x10000000, 0x20000000, 0x40000000, 0x80000000,

OpenCL/inc_cipher_twofish.cl

@@ -21,7 +21,7 @@
 
 #define extract_byte(x,n) (((x) >> (8 * (n))) & 0xff)
 
-__constant u32a q_tab[2][256] =
+__constant static u32a q_tab[2][256] =
 {
   {
     0xA9, 0x67, 0xB3, 0xE8, 0x04, 0xFD, 0xA3, 0x76, 0x9A, 0x92, 0x80, 0x78,
@@ -75,7 +75,7 @@ __constant u32a q_tab[2][256] =
 
 #define q(n,x) q_tab[n][x]
 
-__constant u32a  m_tab[4][256] =
+__constant static u32a  m_tab[4][256] =
 {
   { 0xBCBC3275, 0xECEC21F3, 0x202043C6, 0xB3B3C9F4, 0xDADA03DB, 0x02028B7B,
     0xE2E22BFB, 0x9E9EFAC8, 0xC9C9EC4A, 0xD4D409D3, 0x18186BE6, 0x1E1E9F6B,

OpenCL/inc_common.cl

@@ -3,7 +3,7 @@
  * License.....: MIT
  */
 
-__constant u32 c_append_helper[64][16] =
+__constant static u32a c_append_helper[64][16] =
 {
   { 0x000000ff, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000 },
   { 0x0000ff00, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000 },
@@ -208,6 +208,26 @@ float get_entropy (const u32 *buf, const int elems)
   return entropy;
 }
 
+int is_valid_hex_8 (const u8 v)
+{
+  // direct lookup table is slower thanks to CMOV
+
+  if ((v >= '0') && (v <= '9')) return 1;
+  if ((v >= 'a') && (v <= 'f')) return 1;
+
+  return 0;
+}
+
+int is_valid_hex_32 (const u32 v)
+{
+  if (is_valid_hex_8 ((u8) (v >>  0)) == 0) return 0;
+  if (is_valid_hex_8 ((u8) (v >>  8)) == 0) return 0;
+  if (is_valid_hex_8 ((u8) (v >> 16)) == 0) return 0;
+  if (is_valid_hex_8 ((u8) (v >> 24)) == 0) return 0;
+
+  return 1;
+}
+
 /**
  * vector functions
  */
@@ -61594,3 +61614,84 @@ __kernel void gpu_memset (__global uint4 *buf, const u32 value, const u64 gid_ma
 
   buf[gid] = (uint4) (value);
 }
+
+__kernel void gpu_atinit (__global pw_t *buf, const u64 gid_max)
+{
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u32 l32 = l32_from_64_S (gid);
+  const u32 h32 = h32_from_64_S (gid);
+
+  pw_t pw;
+
+  pw.i[ 0] = 0x5c5c5c5c ^ l32;
+  pw.i[ 1] = 0x36363636 ^ h32;
+  pw.i[ 2] = 0;
+  pw.i[ 3] = 0;
+  pw.i[ 4] = 0;
+  pw.i[ 5] = 0;
+  pw.i[ 6] = 0;
+  pw.i[ 7] = 0;
+  pw.i[ 8] = 0;
+  pw.i[ 9] = 0;
+  pw.i[10] = 0;
+  pw.i[11] = 0;
+  pw.i[12] = 0;
+  pw.i[13] = 0;
+  pw.i[14] = 0;
+  pw.i[15] = 0;
+  pw.i[16] = 0;
+  pw.i[17] = 0;
+  pw.i[18] = 0;
+  pw.i[19] = 0;
+  pw.i[20] = 0;
+  pw.i[21] = 0;
+  pw.i[22] = 0;
+  pw.i[23] = 0;
+  pw.i[24] = 0;
+  pw.i[25] = 0;
+  pw.i[26] = 0;
+  pw.i[27] = 0;
+  pw.i[28] = 0;
+  pw.i[29] = 0;
+  pw.i[30] = 0;
+  pw.i[31] = 0;
+  pw.i[32] = 0;
+  pw.i[33] = 0;
+  pw.i[34] = 0;
+  pw.i[35] = 0;
+  pw.i[36] = 0;
+  pw.i[37] = 0;
+  pw.i[38] = 0;
+  pw.i[39] = 0;
+  pw.i[40] = 0;
+  pw.i[41] = 0;
+  pw.i[42] = 0;
+  pw.i[43] = 0;
+  pw.i[44] = 0;
+  pw.i[45] = 0;
+  pw.i[46] = 0;
+  pw.i[47] = 0;
+  pw.i[48] = 0;
+  pw.i[49] = 0;
+  pw.i[50] = 0;
+  pw.i[51] = 0;
+  pw.i[52] = 0;
+  pw.i[53] = 0;
+  pw.i[54] = 0;
+  pw.i[55] = 0;
+  pw.i[56] = 0;
+  pw.i[57] = 0;
+  pw.i[58] = 0;
+  pw.i[59] = 0;
+  pw.i[60] = 0;
+  pw.i[61] = 0;
+  pw.i[62] = 0;
+  pw.i[63] = 0; // yep that's faster
+
+  pw.pw_len = 1 + (l32 & 15);
+
+  buf[gid] = pw;
+}

OpenCL/inc_hash_sha224.cl

@@ -4,7 +4,7 @@
 // input buf needs to be in algorithm native byte order (md5 = LE, sha256 = BE, etc)
 // input buf needs to be 64 byte aligned when using md5_update()
 
-__constant u32a k_sha256[64] =
+__constant static u32a k_sha256[64] =
 {
   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
   SHA256C04, SHA256C05, SHA256C06, SHA256C07,

OpenCL/inc_hash_sha256.cl

@@ -4,7 +4,7 @@
 // input buf needs to be in algorithm native byte order (md5 = LE, sha256 = BE, etc)
 // input buf needs to be 64 byte aligned when using md5_update()
 
-__constant u32a k_sha256[64] =
+__constant static u32a k_sha256[64] =
 {
   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
   SHA256C04, SHA256C05, SHA256C06, SHA256C07,

OpenCL/inc_hash_sha384.cl

@@ -4,7 +4,7 @@
 // input buf needs to be in algorithm native byte order (md5 = LE, sha1 = BE, etc)
 // input buf needs to be 128 byte aligned when using sha512_update()
 
-__constant u64a k_sha384[80] =
+__constant static u64a k_sha384[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/inc_hash_sha512.cl

@@ -4,7 +4,7 @@
 // input buf needs to be in algorithm native byte order (md5 = LE, sha1 = BE, etc)
 // input buf needs to be 128 byte aligned when using sha512_update()
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/inc_hash_whirlpool.cl

@@ -6,7 +6,7 @@
 
 #define R 10
 
-__constant u32a Ch[8][256] =
+__constant static u32a Ch[8][256] =
 {
   {
     0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
@@ -538,7 +538,7 @@ __constant u32a Ch[8][256] =
   }
 };
 
-__constant u32a Cl[8][256] =
+__constant static u32a Cl[8][256] =
 {
   {
     0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
@@ -1070,7 +1070,7 @@ __constant u32a Cl[8][256] =
   },
 };
 
-__constant u32a rch[R + 1] =
+__constant static u32a rch[R + 1] =
 {
   0x00000000,
   0x1823c6e8,
@@ -1085,7 +1085,7 @@ __constant u32a rch[R + 1] =
   0xca2dbf07,
 };
 
-__constant u32a rcl[R + 1] =
+__constant static u32a rcl[R + 1] =
 {
   0x00000000,
   0x87b8014f,

OpenCL/inc_luks_af.cl

@@ -1,7 +1,7 @@
 
 // basically normal XXX_transform() but with a different name to avoid collisions with function nameing
 
-__constant u32a AF_k_sha256[64] =
+__constant static u32a AF_k_sha256[64] =
 {
   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
   SHA256C04, SHA256C05, SHA256C06, SHA256C07,
@@ -21,7 +21,7 @@ __constant u32a AF_k_sha256[64] =
   SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
 };
 
-__constant u64a AF_k_sha512[80] =
+__constant static u64a AF_k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/inc_luks_essiv.cl

@@ -1,4 +1,4 @@
-__constant u32a ESSIV_k_sha256[64] =
+__constant static u32a ESSIV_k_sha256[64] =
 {
   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
   SHA256C04, SHA256C05, SHA256C06, SHA256C07,

OpenCL/inc_truecrypt_crc32.cl

@@ -1,4 +1,4 @@
-__constant u32a crc32tab[0x100] =
+__constant static u32a crc32tab[0x100] =
 {
   0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
   0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,

OpenCL/inc_types.cl

@@ -1467,6 +1467,49 @@ typedef struct ethereum_scrypt
 
 } ethereum_scrypt_t;
 
+typedef struct ethereum_presale
+{
+  u32 iv[4];
+  u32 enc_seed[152];
+  u32 enc_seed_len;
+
+} ethereum_presale_t;
+
+typedef struct tacacs_plus
+{
+  u32 session_buf[16];
+
+  u32 ct_data_buf[64];
+  u32 ct_data_len;
+
+  u32 sequence_buf[16];
+
+} tacacs_plus_t;
+
+typedef struct apple_secure_notes
+{
+  u32 Z_PK;
+  u32 ZCRYPTOITERATIONCOUNT;
+  u32 ZCRYPTOSALT[16];
+  u32 ZCRYPTOWRAPPEDKEY[16];
+
+} apple_secure_notes_t;
+
+typedef struct jwt
+{
+  u32 salt_buf[1024];
+  u32 salt_len;
+
+} jwt_t;
+
+typedef struct electrum_wallet
+{
+  u32 salt_type;
+  u32 iv[4];
+  u32 encrypted[4];
+
+} electrum_wallet_t;
+
 typedef struct pdf14_tmp
 {
   u32 digest[4];
@@ -1855,6 +1898,16 @@ typedef struct dpapimk_tmp_v2
 
 } dpapimk_tmp_v2_t;
 
+typedef struct apple_secure_notes_tmp
+{
+  u32 ipad[8];
+  u32 opad[8];
+
+  u32 dgst[8];
+  u32 out[8];
+
+} apple_secure_notes_tmp_t;
+
 typedef struct bsdicrypt_tmp
 {
   u32 Kc[16];

OpenCL/m00600_a3-optimized.cl

@@ -10,7 +10,7 @@
 #include "inc_hash_functions.cl"
 #include "inc_types.cl"
 #include "inc_common.cl"
-#include "inc_simd.cl",
+#include "inc_simd.cl"
 
 #define BLAKE2B_FINAL   1
 #define BLAKE2B_UPDATE  0

OpenCL/m01500_a0.cl

@@ -34,7 +34,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x00820200, 0x00020000, 0x80800000, 0x80820200,
@@ -182,7 +182,7 @@ __constant u32a c_SPtrans[8][64] =
   },
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m01500_a1.cl

@@ -32,7 +32,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x00820200, 0x00020000, 0x80800000, 0x80820200,
@@ -180,7 +180,7 @@ __constant u32a c_SPtrans[8][64] =
   },
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m01700_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01700_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01700_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01710_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01710_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01710_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01720_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01720_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01720_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01730_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01730_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01730_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01740_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01740_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m01740_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha512[80] =
+__constant static u64a k_sha512[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m02000_a0.cl

@@ -3,8 +3,32 @@
  * License.....: MIT
  */
 
-typedef uint   u32;
+typedef uint  u32;
-typedef ulong  u64;
+typedef ulong u64;
+
+typedef struct pw
+{
+  u32 i[64];
+
+  u32 pw_len;
+
+} pw_t;
+
+static u32 l32_from_64_S (u64 a)
+{
+  const u32 r = (u32) (a);
+
+  return r;
+}
+
+static u32 h32_from_64_S (u64 a)
+{
+  a >>= 32;
+
+  const u32 r = (u32) (a);
+
+  return r;
+}
 
 __kernel void gpu_memset (__global uint4 *buf, const u32 value, const u64 gid_max)
 {
@@ -15,6 +39,87 @@ __kernel void gpu_memset (__global uint4 *buf, const u32 value, const u64 gid_ma
   buf[gid] = (uint4) (value);
 }
 
+__kernel void gpu_atinit (__global pw_t *buf, const u64 gid_max)
+{
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u32 l32 = l32_from_64_S (gid);
+  const u32 h32 = h32_from_64_S (gid);
+
+  pw_t pw;
+
+  pw.i[ 0] = 0x5c5c5c5c ^ l32;
+  pw.i[ 1] = 0x36363636 ^ h32;
+  pw.i[ 2] = 0;
+  pw.i[ 3] = 0;
+  pw.i[ 4] = 0;
+  pw.i[ 5] = 0;
+  pw.i[ 6] = 0;
+  pw.i[ 7] = 0;
+  pw.i[ 8] = 0;
+  pw.i[ 9] = 0;
+  pw.i[10] = 0;
+  pw.i[11] = 0;
+  pw.i[12] = 0;
+  pw.i[13] = 0;
+  pw.i[14] = 0;
+  pw.i[15] = 0;
+  pw.i[16] = 0;
+  pw.i[17] = 0;
+  pw.i[18] = 0;
+  pw.i[19] = 0;
+  pw.i[20] = 0;
+  pw.i[21] = 0;
+  pw.i[22] = 0;
+  pw.i[23] = 0;
+  pw.i[24] = 0;
+  pw.i[25] = 0;
+  pw.i[26] = 0;
+  pw.i[27] = 0;
+  pw.i[28] = 0;
+  pw.i[29] = 0;
+  pw.i[30] = 0;
+  pw.i[31] = 0;
+  pw.i[32] = 0;
+  pw.i[33] = 0;
+  pw.i[34] = 0;
+  pw.i[35] = 0;
+  pw.i[36] = 0;
+  pw.i[37] = 0;
+  pw.i[38] = 0;
+  pw.i[39] = 0;
+  pw.i[40] = 0;
+  pw.i[41] = 0;
+  pw.i[42] = 0;
+  pw.i[43] = 0;
+  pw.i[44] = 0;
+  pw.i[45] = 0;
+  pw.i[46] = 0;
+  pw.i[47] = 0;
+  pw.i[48] = 0;
+  pw.i[49] = 0;
+  pw.i[50] = 0;
+  pw.i[51] = 0;
+  pw.i[52] = 0;
+  pw.i[53] = 0;
+  pw.i[54] = 0;
+  pw.i[55] = 0;
+  pw.i[56] = 0;
+  pw.i[57] = 0;
+  pw.i[58] = 0;
+  pw.i[59] = 0;
+  pw.i[60] = 0;
+  pw.i[61] = 0;
+  pw.i[62] = 0;
+  pw.i[63] = 0; // yep that's faster
+
+  pw.pw_len = 1 + (l32 & 15);
+
+  buf[gid] = pw;
+}
+
 __kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
 }

OpenCL/m02000_a1.cl

@@ -3,8 +3,32 @@
  * License.....: MIT
  */
 
-typedef uint   u32;
+typedef uint  u32;
-typedef ulong  u64;
+typedef ulong u64;
+
+typedef struct pw
+{
+  u32 i[64];
+
+  u32 pw_len;
+
+} pw_t;
+
+static u32 l32_from_64_S (u64 a)
+{
+  const u32 r = (u32) (a);
+
+  return r;
+}
+
+static u32 h32_from_64_S (u64 a)
+{
+  a >>= 32;
+
+  const u32 r = (u32) (a);
+
+  return r;
+}
 
 __kernel void gpu_memset (__global uint4 *buf, const u32 value, const u64 gid_max)
 {
@@ -15,6 +39,87 @@ __kernel void gpu_memset (__global uint4 *buf, const u32 value, const u64 gid_ma
   buf[gid] = (uint4) (value);
 }
 
+__kernel void gpu_atinit (__global pw_t *buf, const u64 gid_max)
+{
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u32 l32 = l32_from_64_S (gid);
+  const u32 h32 = h32_from_64_S (gid);
+
+  pw_t pw;
+
+  pw.i[ 0] = 0x5c5c5c5c ^ l32;
+  pw.i[ 1] = 0x36363636 ^ h32;
+  pw.i[ 2] = 0;
+  pw.i[ 3] = 0;
+  pw.i[ 4] = 0;
+  pw.i[ 5] = 0;
+  pw.i[ 6] = 0;
+  pw.i[ 7] = 0;
+  pw.i[ 8] = 0;
+  pw.i[ 9] = 0;
+  pw.i[10] = 0;
+  pw.i[11] = 0;
+  pw.i[12] = 0;
+  pw.i[13] = 0;
+  pw.i[14] = 0;
+  pw.i[15] = 0;
+  pw.i[16] = 0;
+  pw.i[17] = 0;
+  pw.i[18] = 0;
+  pw.i[19] = 0;
+  pw.i[20] = 0;
+  pw.i[21] = 0;
+  pw.i[22] = 0;
+  pw.i[23] = 0;
+  pw.i[24] = 0;
+  pw.i[25] = 0;
+  pw.i[26] = 0;
+  pw.i[27] = 0;
+  pw.i[28] = 0;
+  pw.i[29] = 0;
+  pw.i[30] = 0;
+  pw.i[31] = 0;
+  pw.i[32] = 0;
+  pw.i[33] = 0;
+  pw.i[34] = 0;
+  pw.i[35] = 0;
+  pw.i[36] = 0;
+  pw.i[37] = 0;
+  pw.i[38] = 0;
+  pw.i[39] = 0;
+  pw.i[40] = 0;
+  pw.i[41] = 0;
+  pw.i[42] = 0;
+  pw.i[43] = 0;
+  pw.i[44] = 0;
+  pw.i[45] = 0;
+  pw.i[46] = 0;
+  pw.i[47] = 0;
+  pw.i[48] = 0;
+  pw.i[49] = 0;
+  pw.i[50] = 0;
+  pw.i[51] = 0;
+  pw.i[52] = 0;
+  pw.i[53] = 0;
+  pw.i[54] = 0;
+  pw.i[55] = 0;
+  pw.i[56] = 0;
+  pw.i[57] = 0;
+  pw.i[58] = 0;
+  pw.i[59] = 0;
+  pw.i[60] = 0;
+  pw.i[61] = 0;
+  pw.i[62] = 0;
+  pw.i[63] = 0; // yep that's faster
+
+  pw.pw_len = 1 + (l32 & 15);
+
+  buf[gid] = pw;
+}
+
 __kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
 }

OpenCL/m02000_a3.cl

@@ -3,8 +3,32 @@
  * License.....: MIT
  */
 
-typedef uint   u32;
+typedef uint  u32;
-typedef ulong  u64;
+typedef ulong u64;
+
+typedef struct pw
+{
+  u32 i[64];
+
+  u32 pw_len;
+
+} pw_t;
+
+static u32 l32_from_64_S (u64 a)
+{
+  const u32 r = (u32) (a);
+
+  return r;
+}
+
+static u32 h32_from_64_S (u64 a)
+{
+  a >>= 32;
+
+  const u32 r = (u32) (a);
+
+  return r;
+}
 
 __kernel void gpu_memset (__global uint4 *buf, const u32 value, const u64 gid_max)
 {
@@ -15,6 +39,87 @@ __kernel void gpu_memset (__global uint4 *buf, const u32 value, const u64 gid_ma
   buf[gid] = (uint4) (value);
 }
 
+__kernel void gpu_atinit (__global pw_t *buf, const u64 gid_max)
+{
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u32 l32 = l32_from_64_S (gid);
+  const u32 h32 = h32_from_64_S (gid);
+
+  pw_t pw;
+
+  pw.i[ 0] = 0x5c5c5c5c ^ l32;
+  pw.i[ 1] = 0x36363636 ^ h32;
+  pw.i[ 2] = 0;
+  pw.i[ 3] = 0;
+  pw.i[ 4] = 0;
+  pw.i[ 5] = 0;
+  pw.i[ 6] = 0;
+  pw.i[ 7] = 0;
+  pw.i[ 8] = 0;
+  pw.i[ 9] = 0;
+  pw.i[10] = 0;
+  pw.i[11] = 0;
+  pw.i[12] = 0;
+  pw.i[13] = 0;
+  pw.i[14] = 0;
+  pw.i[15] = 0;
+  pw.i[16] = 0;
+  pw.i[17] = 0;
+  pw.i[18] = 0;
+  pw.i[19] = 0;
+  pw.i[20] = 0;
+  pw.i[21] = 0;
+  pw.i[22] = 0;
+  pw.i[23] = 0;
+  pw.i[24] = 0;
+  pw.i[25] = 0;
+  pw.i[26] = 0;
+  pw.i[27] = 0;
+  pw.i[28] = 0;
+  pw.i[29] = 0;
+  pw.i[30] = 0;
+  pw.i[31] = 0;
+  pw.i[32] = 0;
+  pw.i[33] = 0;
+  pw.i[34] = 0;
+  pw.i[35] = 0;
+  pw.i[36] = 0;
+  pw.i[37] = 0;
+  pw.i[38] = 0;
+  pw.i[39] = 0;
+  pw.i[40] = 0;
+  pw.i[41] = 0;
+  pw.i[42] = 0;
+  pw.i[43] = 0;
+  pw.i[44] = 0;
+  pw.i[45] = 0;
+  pw.i[46] = 0;
+  pw.i[47] = 0;
+  pw.i[48] = 0;
+  pw.i[49] = 0;
+  pw.i[50] = 0;
+  pw.i[51] = 0;
+  pw.i[52] = 0;
+  pw.i[53] = 0;
+  pw.i[54] = 0;
+  pw.i[55] = 0;
+  pw.i[56] = 0;
+  pw.i[57] = 0;
+  pw.i[58] = 0;
+  pw.i[59] = 0;
+  pw.i[60] = 0;
+  pw.i[61] = 0;
+  pw.i[62] = 0;
+  pw.i[63] = 0; // yep that's faster
+
+  pw.pw_len = 1 + (l32 & 15);
+
+  buf[gid] = pw;
+}
+
 __kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
 }

OpenCL/m02400_a0-optimized.cl

@@ -3,7 +3,7 @@
  * License.....: MIT
  */
 
-#define NEW_SIMD_CODE
+//#define NEW_SIMD_CODE
 
 #include "inc_vendor.cl"
 #include "inc_hash_constants.h"
@@ -57,18 +57,39 @@ __kernel void m02400_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
 
     const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
-    w1[0] = 0x80;
+    if (out_len <= 16)
-    w1[1] = 0;
+    {
-    w1[2] = 0;
+      w1[0] = 0x80;
-    w1[3] = 0;
+      w1[1] = 0;
-    w2[0] = 0;
+      w1[2] = 0;
-    w2[1] = 0;
+      w1[3] = 0;
-    w2[2] = 0;
+      w2[0] = 0;
-    w2[3] = 0;
+      w2[1] = 0;
-    w3[0] = 0;
+      w2[2] = 0;
-    w3[1] = 0;
+      w2[3] = 0;
-    w3[2] = 16 * 8;
+      w3[0] = 0;
-    w3[3] = 0;
+      w3[1] = 0;
+      w3[2] = 16 * 8;
+      w3[3] = 0;
+    }
+    else if (out_len <= 32)
+    {
+      w2[0] = 0x80;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 32 * 8;
+      w3[3] = 0;
+    }
+    else if (out_len <= 48)
+    {
+      w3[0] = 0x80;
+      w3[1] = 0;
+      w3[2] = 48 * 8;
+      w3[3] = 0;
+    }
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;
@@ -217,18 +238,39 @@ __kernel void m02400_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
 
     const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
-    w1[0] = 0x80;
+    if (out_len <= 16)
-    w1[1] = 0;
+    {
-    w1[2] = 0;
+      w1[0] = 0x80;
-    w1[3] = 0;
+      w1[1] = 0;
-    w2[0] = 0;
+      w1[2] = 0;
-    w2[1] = 0;
+      w1[3] = 0;
-    w2[2] = 0;
+      w2[0] = 0;
-    w2[3] = 0;
+      w2[1] = 0;
-    w3[0] = 0;
+      w2[2] = 0;
-    w3[1] = 0;
+      w2[3] = 0;
-    w3[2] = 16 * 8;
+      w3[0] = 0;
-    w3[3] = 0;
+      w3[1] = 0;
+      w3[2] = 16 * 8;
+      w3[3] = 0;
+    }
+    else if (out_len <= 32)
+    {
+      w2[0] = 0x80;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 32 * 8;
+      w3[3] = 0;
+    }
+    else if (out_len <= 48)
+    {
+      w3[0] = 0x80;
+      w3[1] = 0;
+      w3[2] = 48 * 8;
+      w3[3] = 0;
+    }
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;

OpenCL/m02400_a1-optimized.cl

@@ -3,7 +3,7 @@
  * License.....: MIT
  */
 
-#define NEW_SIMD_CODE
+//#define NEW_SIMD_CODE
 
 #include "inc_vendor.cl"
 #include "inc_hash_constants.h"
@@ -119,18 +119,39 @@ __kernel void m02400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
      * md5
      */
 
-    w1[0] = 0x80;
+    if (pw_len <= 16)
-    w1[1] = 0;
+    {
-    w1[2] = 0;
+      w1[0] = 0x80;
-    w1[3] = 0;
+      w1[1] = 0;
-    w2[0] = 0;
+      w1[2] = 0;
-    w2[1] = 0;
+      w1[3] = 0;
-    w2[2] = 0;
+      w2[0] = 0;
-    w2[3] = 0;
+      w2[1] = 0;
-    w3[0] = 0;
+      w2[2] = 0;
-    w3[1] = 0;
+      w2[3] = 0;
-    w3[2] = 16 * 8;
+      w3[0] = 0;
-    w3[3] = 0;
+      w3[1] = 0;
+      w3[2] = 16 * 8;
+      w3[3] = 0;
+    }
+    else if (pw_len <= 32)
+    {
+      w2[0] = 0x80;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 32 * 8;
+      w3[3] = 0;
+    }
+    else if (pw_len <= 48)
+    {
+      w3[0] = 0x80;
+      w3[1] = 0;
+      w3[2] = 48 * 8;
+      w3[3] = 0;
+    }
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;
@@ -343,18 +364,39 @@ __kernel void m02400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
      * md5
      */
 
-    w1[0] = 0x80;
+    if (pw_len <= 16)
-    w1[1] = 0;
+    {
-    w1[2] = 0;
+      w1[0] = 0x80;
-    w1[3] = 0;
+      w1[1] = 0;
-    w2[0] = 0;
+      w1[2] = 0;
-    w2[1] = 0;
+      w1[3] = 0;
-    w2[2] = 0;
+      w2[0] = 0;
-    w2[3] = 0;
+      w2[1] = 0;
-    w3[0] = 0;
+      w2[2] = 0;
-    w3[1] = 0;
+      w2[3] = 0;
-    w3[2] = 16 * 8;
+      w3[0] = 0;
-    w3[3] = 0;
+      w3[1] = 0;
+      w3[2] = 16 * 8;
+      w3[3] = 0;
+    }
+    else if (pw_len <= 32)
+    {
+      w2[0] = 0x80;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 32 * 8;
+      w3[3] = 0;
+    }
+    else if (pw_len <= 48)
+    {
+      w3[0] = 0x80;
+      w3[1] = 0;
+      w3[2] = 48 * 8;
+      w3[3] = 0;
+    }
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;

OpenCL/m02400_a3-optimized.cl

@@ -25,18 +25,39 @@ void m02400m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke
    * algorithm specific
    */
 
-  w[ 4] = 0x80;
+  if (pw_len <= 16)
-  w[ 5] = 0;
+  {
-  w[ 6] = 0;
+    w[ 4] = 0x80;
-  w[ 7] = 0;
+    w[ 5] = 0;
-  w[ 8] = 0;
+    w[ 6] = 0;
-  w[ 9] = 0;
+    w[ 7] = 0;
-  w[10] = 0;
+    w[ 8] = 0;
-  w[11] = 0;
+    w[ 9] = 0;
-  w[12] = 0;
+    w[10] = 0;
-  w[13] = 0;
+    w[11] = 0;
-  w[14] = 16 * 8;
+    w[12] = 0;
-  w[15] = 0;
+    w[13] = 0;
+    w[14] = 16 * 8;
+    w[15] = 0;
+  }
+  else if (pw_len <= 32)
+  {
+    w[ 8] = 0x80;
+    w[ 9] = 0;
+    w[10] = 0;
+    w[11] = 0;
+    w[12] = 0;
+    w[13] = 0;
+    w[14] = 32 * 8;
+    w[15] = 0;
+  }
+  else if (pw_len <= 48)
+  {
+    w[12] = 0x80;
+    w[13] = 0;
+    w[14] = 48 * 8;
+    w[15] = 0;
+  }
 
   /**
    * base
@@ -219,18 +240,39 @@ void m02400s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke
    * algorithm specific
    */
 
-  w[ 4] = 0x80;
+  if (pw_len <= 16)
-  w[ 5] = 0;
+  {
-  w[ 6] = 0;
+    w[ 4] = 0x80;
-  w[ 7] = 0;
+    w[ 5] = 0;
-  w[ 8] = 0;
+    w[ 6] = 0;
-  w[ 9] = 0;
+    w[ 7] = 0;
-  w[10] = 0;
+    w[ 8] = 0;
-  w[11] = 0;
+    w[ 9] = 0;
-  w[12] = 0;
+    w[10] = 0;
-  w[13] = 0;
+    w[11] = 0;
-  w[14] = 16 * 8;
+    w[12] = 0;
-  w[15] = 0;
+    w[13] = 0;
+    w[14] = 16 * 8;
+    w[15] = 0;
+  }
+  else if (pw_len <= 32)
+  {
+    w[ 8] = 0x80;
+    w[ 9] = 0;
+    w[10] = 0;
+    w[11] = 0;
+    w[12] = 0;
+    w[13] = 0;
+    w[14] = 32 * 8;
+    w[15] = 0;
+  }
+  else if (pw_len <= 48)
+  {
+    w[12] = 0x80;
+    w[13] = 0;
+    w[14] = 48 * 8;
+    w[15] = 0;
+  }
 
   /**
    * base

OpenCL/m02410_a0-optimized.cl

@@ -3,7 +3,7 @@
  * License.....: MIT
  */
 
-#define NEW_SIMD_CODE
+//#define NEW_SIMD_CODE
 
 #include "inc_vendor.cl"
 #include "inc_hash_constants.h"
@@ -70,6 +70,8 @@ __kernel void m02410_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
   salt_buf3[2] = 0;
   salt_buf3[3] = 0;
 
+  const u32 salt_len = salt_bufs[salt_pos].salt_len;
+
   /**
    * loop
    */
@@ -132,18 +134,41 @@ __kernel void m02410_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
      * md5
      */
 
-    w1[0] = 0x80;
+    const u32x out_salt_len = out_len + salt_len;
-    w1[1] = 0;
+
-    w1[2] = 0;
+    if (out_salt_len <= 16)
-    w1[3] = 0;
+    {
-    w2[0] = 0;
+      w1[0] = 0x80;
-    w2[1] = 0;
+      w1[1] = 0;
-    w2[2] = 0;
+      w1[2] = 0;
-    w2[3] = 0;
+      w1[3] = 0;
-    w3[0] = 0;
+      w2[0] = 0;
-    w3[1] = 0;
+      w2[1] = 0;
-    w3[2] = 16 * 8;
+      w2[2] = 0;
-    w3[3] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 16 * 8;
+      w3[3] = 0;
+    }
+    else if (out_salt_len <= 32)
+    {
+      w2[0] = 0x80;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 32 * 8;
+      w3[3] = 0;
+    }
+    else if (out_salt_len <= 48)
+    {
+      w3[0] = 0x80;
+      w3[1] = 0;
+      w3[2] = 48 * 8;
+      w3[3] = 0;
+    }
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;
@@ -293,6 +318,8 @@ __kernel void m02410_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
   salt_buf3[2] = 0;
   salt_buf3[3] = 0;
 
+  const u32 salt_len = salt_bufs[salt_pos].salt_len;
+
   /**
    * digest
    */
@@ -367,18 +394,41 @@ __kernel void m02410_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
      * md5
      */
 
-    w1[0] = 0x80;
+    const u32x out_salt_len = out_len + salt_len;
-    w1[1] = 0;
+
-    w1[2] = 0;
+    if (out_salt_len <= 16)
-    w1[3] = 0;
+    {
-    w2[0] = 0;
+      w1[0] = 0x80;
-    w2[1] = 0;
+      w1[1] = 0;
-    w2[2] = 0;
+      w1[2] = 0;
-    w2[3] = 0;
+      w1[3] = 0;
-    w3[0] = 0;
+      w2[0] = 0;
-    w3[1] = 0;
+      w2[1] = 0;
-    w3[2] = 16 * 8;
+      w2[2] = 0;
-    w3[3] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 16 * 8;
+      w3[3] = 0;
+    }
+    else if (out_salt_len <= 32)
+    {
+      w2[0] = 0x80;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 32 * 8;
+      w3[3] = 0;
+    }
+    else if (out_salt_len <= 48)
+    {
+      w3[0] = 0x80;
+      w3[1] = 0;
+      w3[2] = 48 * 8;
+      w3[3] = 0;
+    }
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;

OpenCL/m02410_a1-optimized.cl

@@ -3,7 +3,7 @@
  * License.....: MIT
  */
 
-#define NEW_SIMD_CODE
+//#define NEW_SIMD_CODE
 
 #include "inc_vendor.cl"
 #include "inc_hash_constants.h"
@@ -68,6 +68,8 @@ __kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
   salt_buf3[2] = 0;
   salt_buf3[3] = 0;
 
+  const u32 salt_len = salt_bufs[salt_pos].salt_len;
+
   /**
    * loop
    */
@@ -190,18 +192,41 @@ __kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
      * md5
      */
 
-    w1[0] = 0x80;
+    const u32x pw_salt_len = pw_len + salt_len;
-    w1[1] = 0;
+
-    w1[2] = 0;
+    if (pw_salt_len <= 16)
-    w1[3] = 0;
+    {
-    w2[0] = 0;
+      w1[0] = 0x80;
-    w2[1] = 0;
+      w1[1] = 0;
-    w2[2] = 0;
+      w1[2] = 0;
-    w2[3] = 0;
+      w1[3] = 0;
-    w3[0] = 0;
+      w2[0] = 0;
-    w3[1] = 0;
+      w2[1] = 0;
-    w3[2] = 16 * 8;
+      w2[2] = 0;
-    w3[3] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 16 * 8;
+      w3[3] = 0;
+    }
+    else if (pw_salt_len <= 32)
+    {
+      w2[0] = 0x80;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 32 * 8;
+      w3[3] = 0;
+    }
+    else if (pw_salt_len <= 48)
+    {
+      w3[0] = 0x80;
+      w3[1] = 0;
+      w3[2] = 48 * 8;
+      w3[3] = 0;
+    }
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;
@@ -351,6 +376,8 @@ __kernel void m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   salt_buf3[2] = 0;
   salt_buf3[3] = 0;
 
+  const u32 salt_len = salt_bufs[salt_pos].salt_len;
+
   /**
    * digest
    */
@@ -485,18 +512,41 @@ __kernel void m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
      * md5
      */
 
-    w1[0] = 0x80;
+    const u32x pw_salt_len = pw_len + salt_len;
-    w1[1] = 0;
+
-    w1[2] = 0;
+    if (pw_salt_len <= 16)
-    w1[3] = 0;
+    {
-    w2[0] = 0;
+      w1[0] = 0x80;
-    w2[1] = 0;
+      w1[1] = 0;
-    w2[2] = 0;
+      w1[2] = 0;
-    w2[3] = 0;
+      w1[3] = 0;
-    w3[0] = 0;
+      w2[0] = 0;
-    w3[1] = 0;
+      w2[1] = 0;
-    w3[2] = 16 * 8;
+      w2[2] = 0;
-    w3[3] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 16 * 8;
+      w3[3] = 0;
+    }
+    else if (pw_salt_len <= 32)
+    {
+      w2[0] = 0x80;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 32 * 8;
+      w3[3] = 0;
+    }
+    else if (pw_salt_len <= 48)
+    {
+      w3[0] = 0x80;
+      w3[1] = 0;
+      w3[2] = 48 * 8;
+      w3[3] = 0;
+    }
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;

OpenCL/m02410_a3-optimized.cl

@@ -47,6 +47,8 @@ void m02410m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke
   salt_buf3[2] = 0;
   salt_buf3[3] = 0;
 
+  const u32 salt_len = salt_bufs[salt_pos].salt_len;
+
   switch_buffer_by_offset_le_S (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_len);
 
   w[ 0] |= salt_buf0[0];
@@ -70,18 +72,41 @@ void m02410m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke
    * algorithm specific
    */
 
-  w[ 4] = 0x80;
+  const u32 pw_salt_len = pw_len + salt_len;
-  w[ 5] = 0;
+
-  w[ 6] = 0;
+  if (pw_salt_len <= 16)
-  w[ 7] = 0;
+  {
-  w[ 8] = 0;
+    w[ 4] = 0x80;
-  w[ 9] = 0;
+    w[ 5] = 0;
-  w[10] = 0;
+    w[ 6] = 0;
-  w[11] = 0;
+    w[ 7] = 0;
-  w[12] = 0;
+    w[ 8] = 0;
-  w[13] = 0;
+    w[ 9] = 0;
-  w[14] = 16 * 8;
+    w[10] = 0;
-  w[15] = 0;
+    w[11] = 0;
+    w[12] = 0;
+    w[13] = 0;
+    w[14] = 16 * 8;
+    w[15] = 0;
+  }
+  else if (pw_salt_len <= 32)
+  {
+    w[ 8] = 0x80;
+    w[ 9] = 0;
+    w[10] = 0;
+    w[11] = 0;
+    w[12] = 0;
+    w[13] = 0;
+    w[14] = 32 * 8;
+    w[15] = 0;
+  }
+  else if (pw_salt_len <= 48)
+  {
+    w[12] = 0x80;
+    w[13] = 0;
+    w[14] = 48 * 8;
+    w[15] = 0;
+  }
 
   /**
    * base
@@ -286,6 +311,8 @@ void m02410s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke
   salt_buf3[2] = 0;
   salt_buf3[3] = 0;
 
+  const u32 salt_len = salt_bufs[salt_pos].salt_len;
+
   switch_buffer_by_offset_le_S (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_len);
 
   w[ 0] |= salt_buf0[0];
@@ -309,18 +336,41 @@ void m02410s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke
    * algorithm specific
    */
 
-  w[ 4] = 0x80;
+  const u32 pw_salt_len = pw_len + salt_len;
-  w[ 5] = 0;
+
-  w[ 6] = 0;
+  if (pw_salt_len <= 16)
-  w[ 7] = 0;
+  {
-  w[ 8] = 0;
+    w[ 4] = 0x80;
-  w[ 9] = 0;
+    w[ 5] = 0;
-  w[10] = 0;
+    w[ 6] = 0;
-  w[11] = 0;
+    w[ 7] = 0;
-  w[12] = 0;
+    w[ 8] = 0;
-  w[13] = 0;
+    w[ 9] = 0;
-  w[14] = 16 * 8;
+    w[10] = 0;
-  w[15] = 0;
+    w[11] = 0;
+    w[12] = 0;
+    w[13] = 0;
+    w[14] = 16 * 8;
+    w[15] = 0;
+  }
+  else if (pw_salt_len <= 32)
+  {
+    w[ 8] = 0x80;
+    w[ 9] = 0;
+    w[10] = 0;
+    w[11] = 0;
+    w[12] = 0;
+    w[13] = 0;
+    w[14] = 32 * 8;
+    w[15] = 0;
+  }
+  else if (pw_salt_len <= 48)
+  {
+    w[12] = 0x80;
+    w[13] = 0;
+    w[14] = 48 * 8;
+    w[15] = 0;
+  }
 
   /**
    * base

OpenCL/m03000_a0.cl

@@ -34,7 +34,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -182,7 +182,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m03000_a1.cl

@@ -32,7 +32,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -180,7 +180,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m03100_a0-optimized.cl

@@ -53,7 +53,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -201,7 +201,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m03100_a1-optimized.cl

@@ -51,7 +51,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -199,7 +199,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m03100_a3-optimized.cl

@@ -50,7 +50,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -198,7 +198,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m03200.cl

@@ -14,7 +14,7 @@
 
 // http://www.schneier.com/code/constants.txt
 
-__constant u32a c_sbox0[256] =
+__constant static u32a c_sbox0[256] =
 {
   0xd1310ba6, 0x98dfb5ac, 0x2ffd72db, 0xd01adfb7,
   0xb8e1afed, 0x6a267e96, 0xba7c9045, 0xf12c7f99,
@@ -82,7 +82,7 @@ __constant u32a c_sbox0[256] =
   0x53b02d5d, 0xa99f8fa1, 0x08ba4799, 0x6e85076a
 };
 
-__constant u32a c_sbox1[256] =
+__constant static u32a c_sbox1[256] =
 {
   0x4b7a70e9, 0xb5b32944, 0xdb75092e, 0xc4192623,
   0xad6ea6b0, 0x49a7df7d, 0x9cee60b8, 0x8fedb266,
@@ -150,7 +150,7 @@ __constant u32a c_sbox1[256] =
   0x153e21e7, 0x8fb03d4a, 0xe6e39f2b, 0xdb83adf7
 };
 
-__constant u32a c_sbox2[256] =
+__constant static u32a c_sbox2[256] =
 {
   0xe93d5a68, 0x948140f7, 0xf64c261c, 0x94692934,
   0x411520f7, 0x7602d4f7, 0xbcf46b2e, 0xd4a20068,
@@ -218,7 +218,7 @@ __constant u32a c_sbox2[256] =
   0xd79a3234, 0x92638212, 0x670efa8e, 0x406000e0
 };
 
-__constant u32a c_sbox3[256] =
+__constant static u32a c_sbox3[256] =
 {
   0x3a39ce37, 0xd3faf5cf, 0xabc27737, 0x5ac52d1b,
   0x5cb0679e, 0x4fa33742, 0xd3822740, 0x99bc9bbe,

OpenCL/m05000_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u64a keccakf_rndc[24] =
+__constant static u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

OpenCL/m05000_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a keccakf_rndc[24] =
+__constant static u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

OpenCL/m05000_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a keccakf_rndc[24] =
+__constant static u64a keccakf_rndc[24] =
 {
   0x0000000000000001, 0x0000000000008082, 0x800000000000808a,
   0x8000000080008000, 0x000000000000808b, 0x0000000080000001,

OpenCL/m05500_a0-optimized.cl

@@ -34,7 +34,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -182,7 +182,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m05500_a0.cl

@@ -35,7 +35,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -183,7 +183,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m05500_a1-optimized.cl

@@ -32,7 +32,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -180,7 +180,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m05500_a1.cl

@@ -33,7 +33,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -181,7 +181,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m05500_a3-optimized.cl

@@ -32,7 +32,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -180,7 +180,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m05500_a3.cl

@@ -33,7 +33,7 @@
   a  = a ^ tt;              \
 }
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -181,7 +181,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m05800-optimized.cl

@@ -13,7 +13,7 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
-__constant u32a c_pc_dec[1024] =
+__constant static u32a c_pc_dec[1024] =
 {
   0x00000030,
   0x00000031,
@@ -1041,7 +1041,7 @@ __constant u32a c_pc_dec[1024] =
   0x33323031,
 };
 
-__constant u32a c_pc_len[1024] =
+__constant static u32a c_pc_len[1024] =
 {
   1,
   1,

OpenCL/m05800.cl

@@ -13,7 +13,7 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
-__constant u32a c_pc_dec[1024] =
+__constant static u32a c_pc_dec[1024] =
 {
   0x00000030,
   0x00000031,
@@ -1041,7 +1041,7 @@ __constant u32a c_pc_dec[1024] =
   0x33323031,
 };
 
-__constant u32a c_pc_len[1024] =
+__constant static u32a c_pc_len[1024] =
 {
   1,
   1,

OpenCL/m06900_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u32a c_tables[4][256] =
+__constant static u32a c_tables[4][256] =
 {
   {
     0x00072000, 0x00075000, 0x00074800, 0x00071000,

OpenCL/m06900_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u32a c_tables[4][256] =
+__constant static u32a c_tables[4][256] =
 {
   {
     0x00072000, 0x00075000, 0x00074800, 0x00071000,

OpenCL/m06900_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u32a c_tables[4][256] =
+__constant static u32a c_tables[4][256] =
 {
   {
     0x00072000, 0x00075000, 0x00074800, 0x00071000,

OpenCL/m07700_a0-optimized.cl

@@ -28,7 +28,7 @@
   (a)[((n)/4)+1]  = x >> 32;        \
 }
 
-__constant u32a sapb_trans_tbl[256] =
+__constant static 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,
@@ -49,7 +49,7 @@ __constant u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-__constant u32a bcodeArray[48] =
+__constant static 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,

OpenCL/m07700_a1-optimized.cl

@@ -26,7 +26,7 @@
   (a)[((n)/4)+1]  = x >> 32;        \
 }
 
-__constant u32a sapb_trans_tbl[256] =
+__constant static 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,
@@ -47,7 +47,7 @@ __constant u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-__constant u32a bcodeArray[48] =
+__constant static 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,

OpenCL/m07700_a3-optimized.cl

@@ -17,7 +17,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 u32a sapb_trans_tbl[256] =
+__constant static 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,
@@ -38,7 +38,7 @@ __constant u32a sapb_trans_tbl[256] =
   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
 };
 
-__constant u32a bcodeArray[48] =
+__constant static 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,

OpenCL/m07800_a0-optimized.cl

@@ -16,7 +16,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
-__constant u32 theMagicArray[64] =
+__constant static u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

OpenCL/m07800_a1-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
-__constant u32 theMagicArray[64] =
+__constant static u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

OpenCL/m07800_a3-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
-__constant u32 theMagicArray[64] =
+__constant static u32a theMagicArray[64] =
 {
   0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
   0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,

OpenCL/m08000_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u32a k_sha256[64] =
+__constant static u32a k_sha256[64] =
 {
   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
   SHA256C04, SHA256C05, SHA256C06, SHA256C07,

OpenCL/m08000_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u32a k_sha256[64] =
+__constant static u32a k_sha256[64] =
 {
   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
   SHA256C04, SHA256C05, SHA256C06, SHA256C07,

OpenCL/m08000_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u32a k_sha256[64] =
+__constant static u32a k_sha256[64] =
 {
   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
   SHA256C04, SHA256C05, SHA256C06, SHA256C07,

OpenCL/m08500_a0.cl

@@ -52,7 +52,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-__constant u32a c_ascii_to_ebcdic_pc[256] =
+__constant static 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,
@@ -74,7 +74,7 @@ __constant u32a c_ascii_to_ebcdic_pc[256] =
   0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
 };
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -222,7 +222,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m08500_a1.cl

@@ -50,7 +50,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-__constant u32a c_ascii_to_ebcdic_pc[256] =
+__constant static 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,
@@ -72,7 +72,7 @@ __constant u32a c_ascii_to_ebcdic_pc[256] =
   0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
 };
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -220,7 +220,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m08500_a3.cl

@@ -50,7 +50,7 @@
   PERM_OP (l, r, tt,  4, 0x0f0f0f0f);  \
 }
 
-__constant u32a c_ascii_to_ebcdic_pc[256] =
+__constant static 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,
@@ -72,7 +72,7 @@ __constant u32a c_ascii_to_ebcdic_pc[256] =
   0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
 };
 
-__constant u32a c_SPtrans[8][64] =
+__constant static u32a c_SPtrans[8][64] =
 {
   {
     0x02080800, 0x00080000, 0x02000002, 0x02080802,
@@ -220,7 +220,7 @@ __constant u32a c_SPtrans[8][64] =
   }
 };
 
-__constant u32a c_skb[8][64] =
+__constant static u32a c_skb[8][64] =
 {
   {
     0x00000000, 0x00000010, 0x20000000, 0x20000010,

OpenCL/m08600_a0.cl

@@ -15,7 +15,7 @@
 #include "inc_rp.cl"
 #include "inc_simd.cl"
 
-__constant u32a lotus_magic_table[256] =
+__constant static u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

OpenCL/m08600_a1.cl

@@ -13,7 +13,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u32a lotus_magic_table[256] =
+__constant static u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

OpenCL/m08600_a3.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u32a lotus_magic_table[256] =
+__constant static u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

OpenCL/m08700_a0-optimized.cl

@@ -15,7 +15,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u32a lotus_magic_table[256] =
+__constant static u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

OpenCL/m08700_a1-optimized.cl

@@ -13,7 +13,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u32a lotus_magic_table[256] =
+__constant static u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

OpenCL/m08700_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u32a lotus_magic_table[256] =
+__constant static u32a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

OpenCL/m09000.cl

@@ -14,7 +14,7 @@
 
 // http://www.schneier.com/code/constants.txt
 
-__constant u32a c_sbox0[256] =
+__constant static u32a c_sbox0[256] =
 {
   0xd1310ba6, 0x98dfb5ac, 0x2ffd72db, 0xd01adfb7,
   0xb8e1afed, 0x6a267e96, 0xba7c9045, 0xf12c7f99,
@@ -82,7 +82,7 @@ __constant u32a c_sbox0[256] =
   0x53b02d5d, 0xa99f8fa1, 0x08ba4799, 0x6e85076a
 };
 
-__constant u32a c_sbox1[256] =
+__constant static u32a c_sbox1[256] =
 {
   0x4b7a70e9, 0xb5b32944, 0xdb75092e, 0xc4192623,
   0xad6ea6b0, 0x49a7df7d, 0x9cee60b8, 0x8fedb266,
@@ -150,7 +150,7 @@ __constant u32a c_sbox1[256] =
   0x153e21e7, 0x8fb03d4a, 0xe6e39f2b, 0xdb83adf7
 };
 
-__constant u32a c_sbox2[256] =
+__constant static u32a c_sbox2[256] =
 {
   0xe93d5a68, 0x948140f7, 0xf64c261c, 0x94692934,
   0x411520f7, 0x7602d4f7, 0xbcf46b2e, 0xd4a20068,
@@ -218,7 +218,7 @@ __constant u32a c_sbox2[256] =
   0xd79a3234, 0x92638212, 0x670efa8e, 0x406000e0
 };
 
-__constant u32a c_sbox3[256] =
+__constant static u32a c_sbox3[256] =
 {
   0x3a39ce37, 0xd3faf5cf, 0xabc27737, 0x5ac52d1b,
   0x5cb0679e, 0x4fa33742, 0xd3822740, 0x99bc9bbe,
@@ -286,7 +286,7 @@ __constant u32a c_sbox3[256] =
   0xb74e6132, 0xce77e25b, 0x578fdfe3, 0x3ac372e6
 };
 
-__constant u32a c_pbox[18] =
+__constant static u32a c_pbox[18] =
 {
   0x243f6a88, 0x85a308d3, 0x13198a2e, 0x03707344,
   0xa4093822, 0x299f31d0, 0x082efa98, 0xec4e6c89,

OpenCL/m09100.cl

@@ -17,7 +17,7 @@
 #define COMPARE_M "inc_comp_multi.cl"
 
 // breaks if used with u8a on AMDGPU-PRO
-__constant u8 lotus64_table[64] =
+__constant static u8a lotus64_table[64] =
 {
   '0', '1', '2', '3', '4', '5', '6', '7',
   '8', '9', 'A', 'B', 'C', 'D', 'E', 'F',
@@ -30,7 +30,7 @@ __constant u8 lotus64_table[64] =
 };
 
 // break if used with u8 on NVidia driver 378.x
-__constant u8a lotus_magic_table[256] =
+__constant static u8a lotus_magic_table[256] =
 {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,

OpenCL/m09720_a0-optimized.cl

@@ -514,8 +514,8 @@ __kernel void m09720_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m09720_a1-optimized.cl

@@ -562,8 +562,8 @@ __kernel void m09720_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m09720_a3-optimized.cl

@@ -475,8 +475,8 @@ void m09720s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m09820_a0-optimized.cl

@@ -198,8 +198,8 @@ __kernel void m09820_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m09820_a1-optimized.cl

@@ -246,8 +246,8 @@ __kernel void m09820_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m09820_a3-optimized.cl

@@ -143,8 +143,8 @@ void m09820s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m10400_a0-optimized.cl

@@ -16,7 +16,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32 padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,

OpenCL/m10400_a1-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32 padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,

OpenCL/m10400_a3-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32 padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,

OpenCL/m10410_a0-optimized.cl

@@ -16,7 +16,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32 padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,

OpenCL/m10410_a1-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32 padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,

OpenCL/m10410_a3-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32 padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,

OpenCL/m10420_a0-optimized.cl

@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32a padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
@@ -250,8 +250,8 @@ __kernel void m10420_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m10420_a1-optimized.cl

@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32a padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
@@ -308,8 +308,8 @@ __kernel void m10420_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m10420_a3-optimized.cl

@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-__constant u32a padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,
@@ -260,8 +260,8 @@ void m10420s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
     digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
+    0,
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0
   };
 
   /**

OpenCL/m10500.cl

@@ -13,7 +13,7 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
-__constant u32a padding[8] =
+__constant static u32a padding[8] =
 {
   0x5e4ebf28,
   0x418a754e,

OpenCL/m10800_a0-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha384[80] =
+__constant static u64a k_sha384[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m10800_a1-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha384[80] =
+__constant static u64a k_sha384[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m10800_a3-optimized.cl

@@ -12,7 +12,7 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-__constant u64a k_sha384[80] =
+__constant static u64a k_sha384[80] =
 {
   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
   SHA512C04, SHA512C05, SHA512C06, SHA512C07,

OpenCL/m11500_a0-optimized.cl

@@ -15,7 +15,7 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
-__constant u32a crc32tab[0x100] =
+__constant static u32a crc32tab[0x100] =
 {
   0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
   0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,