Add support in HMAC for passwords larger than block size of the underlaying hash

OpenCL/inc_hash_md4.cl

@@ -533,7 +533,7 @@ typedef struct md4_hmac_ctx
 
 } md4_hmac_ctx_t;
 
-void md4_hmac_init (md4_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4])
+void md4_hmac_init_64 (md4_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4])
 {
   u32 t0[4];
   u32 t1[4];
@@ -587,6 +587,177 @@ void md4_hmac_init (md4_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const
   md4_update_64 (&ctx->opad, t0, t1, t2, t3, 64);
 }
 
+void md4_hmac_init (md4_hmac_ctx_t *ctx, const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    md4_ctx_t tmp;
+
+    md4_init (&tmp);
+
+    md4_update (&tmp, w, len);
+
+    md4_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = 0;
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  md4_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
+void md4_hmac_init_global (md4_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    md4_ctx_t tmp;
+
+    md4_init (&tmp);
+
+    md4_update_global (&tmp, w, len);
+
+    md4_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = 0;
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  md4_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
+void md4_hmac_init_global_swap (md4_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    md4_ctx_t tmp;
+
+    md4_init (&tmp);
+
+    md4_update_global_swap (&tmp, w, len);
+
+    md4_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = 0;
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = swap32_S (w[ 0]);
+    w0[1] = swap32_S (w[ 1]);
+    w0[2] = swap32_S (w[ 2]);
+    w0[3] = swap32_S (w[ 3]);
+    w1[0] = swap32_S (w[ 4]);
+    w1[1] = swap32_S (w[ 5]);
+    w1[2] = swap32_S (w[ 6]);
+    w1[3] = swap32_S (w[ 7]);
+    w2[0] = swap32_S (w[ 8]);
+    w2[1] = swap32_S (w[ 9]);
+    w2[2] = swap32_S (w[10]);
+    w2[3] = swap32_S (w[11]);
+    w3[0] = swap32_S (w[12]);
+    w3[1] = swap32_S (w[13]);
+    w3[2] = swap32_S (w[14]);
+    w3[3] = swap32_S (w[15]);
+  }
+
+  md4_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
 void md4_hmac_update_64 (md4_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
 {
   md4_update_64 (&ctx->ipad, w0, w1, w2, w3, len);

OpenCL/inc_hash_md5.cl

@@ -735,6 +735,63 @@ void md5_hmac_init_global (md5_hmac_ctx_t *ctx, __global const u32 *w, const int
   md5_hmac_init_64 (ctx, w0, w1, w2, w3);
 }
 
+void md5_hmac_init_global_swap (md5_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    md5_ctx_t tmp;
+
+    md5_init (&tmp);
+
+    md5_update_global_swap (&tmp, w, len);
+
+    md5_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = 0;
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = swap32_S (w[ 0]);
+    w0[1] = swap32_S (w[ 1]);
+    w0[2] = swap32_S (w[ 2]);
+    w0[3] = swap32_S (w[ 3]);
+    w1[0] = swap32_S (w[ 4]);
+    w1[1] = swap32_S (w[ 5]);
+    w1[2] = swap32_S (w[ 6]);
+    w1[3] = swap32_S (w[ 7]);
+    w2[0] = swap32_S (w[ 8]);
+    w2[1] = swap32_S (w[ 9]);
+    w2[2] = swap32_S (w[10]);
+    w2[3] = swap32_S (w[11]);
+    w3[0] = swap32_S (w[12]);
+    w3[1] = swap32_S (w[13]);
+    w3[2] = swap32_S (w[14]);
+    w3[3] = swap32_S (w[15]);
+  }
+
+  md5_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
 void md5_hmac_update_64 (md5_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
 {
   md5_update_64 (&ctx->ipad, w0, w1, w2, w3, len);

OpenCL/inc_hash_ripemd160.cl

@@ -667,7 +667,7 @@ typedef struct ripemd160_hmac_ctx
 
 } ripemd160_hmac_ctx_t;
 
-void ripemd160_hmac_init (ripemd160_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4])
+void ripemd160_hmac_init_64 (ripemd160_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4])
 {
   u32 t0[4];
   u32 t1[4];
@@ -721,6 +721,177 @@ void ripemd160_hmac_init (ripemd160_hmac_ctx_t *ctx, const u32 w0[4], const u32
   ripemd160_update_64 (&ctx->opad, t0, t1, t2, t3, 64);
 }
 
+void ripemd160_hmac_init (ripemd160_hmac_ctx_t *ctx, const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    ripemd160_ctx_t tmp;
+
+    ripemd160_init (&tmp);
+
+    ripemd160_update (&tmp, w, len);
+
+    ripemd160_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  ripemd160_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
+void ripemd160_hmac_init_global (ripemd160_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    ripemd160_ctx_t tmp;
+
+    ripemd160_init (&tmp);
+
+    ripemd160_update_global (&tmp, w, len);
+
+    ripemd160_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  ripemd160_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
+void ripemd160_hmac_init_global_swap (ripemd160_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    ripemd160_ctx_t tmp;
+
+    ripemd160_init (&tmp);
+
+    ripemd160_update_global_swap (&tmp, w, len);
+
+    ripemd160_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = swap32_S (w[ 0]);
+    w0[1] = swap32_S (w[ 1]);
+    w0[2] = swap32_S (w[ 2]);
+    w0[3] = swap32_S (w[ 3]);
+    w1[0] = swap32_S (w[ 4]);
+    w1[1] = swap32_S (w[ 5]);
+    w1[2] = swap32_S (w[ 6]);
+    w1[3] = swap32_S (w[ 7]);
+    w2[0] = swap32_S (w[ 8]);
+    w2[1] = swap32_S (w[ 9]);
+    w2[2] = swap32_S (w[10]);
+    w2[3] = swap32_S (w[11]);
+    w3[0] = swap32_S (w[12]);
+    w3[1] = swap32_S (w[13]);
+    w3[2] = swap32_S (w[14]);
+    w3[3] = swap32_S (w[15]);
+  }
+
+  ripemd160_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
 void ripemd160_hmac_update_64 (ripemd160_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
 {
   ripemd160_update_64 (&ctx->ipad, w0, w1, w2, w3, len);

OpenCL/inc_hash_sha1.cl

@@ -685,7 +685,7 @@ typedef struct sha1_hmac_ctx
 
 } sha1_hmac_ctx_t;
 
-void sha1_hmac_init (sha1_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4])
+void sha1_hmac_init_64 (sha1_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4])
 {
   u32 t0[4];
   u32 t1[4];
@@ -739,6 +739,177 @@ void sha1_hmac_init (sha1_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], con
   sha1_update_64 (&ctx->opad, t0, t1, t2, t3, 64);
 }
 
+void sha1_hmac_init (sha1_hmac_ctx_t *ctx, const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    sha1_ctx_t tmp;
+
+    sha1_init (&tmp);
+
+    sha1_update (&tmp, w, len);
+
+    sha1_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  sha1_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
+void sha1_hmac_init_global (sha1_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    sha1_ctx_t tmp;
+
+    sha1_init (&tmp);
+
+    sha1_update_global (&tmp, w, len);
+
+    sha1_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  sha1_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
+void sha1_hmac_init_global_swap (sha1_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    sha1_ctx_t tmp;
+
+    sha1_init (&tmp);
+
+    sha1_update_global_swap (&tmp, w, len);
+
+    sha1_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = 0;
+    w1[2] = 0;
+    w1[3] = 0;
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = swap32_S (w[ 0]);
+    w0[1] = swap32_S (w[ 1]);
+    w0[2] = swap32_S (w[ 2]);
+    w0[3] = swap32_S (w[ 3]);
+    w1[0] = swap32_S (w[ 4]);
+    w1[1] = swap32_S (w[ 5]);
+    w1[2] = swap32_S (w[ 6]);
+    w1[3] = swap32_S (w[ 7]);
+    w2[0] = swap32_S (w[ 8]);
+    w2[1] = swap32_S (w[ 9]);
+    w2[2] = swap32_S (w[10]);
+    w2[3] = swap32_S (w[11]);
+    w3[0] = swap32_S (w[12]);
+    w3[1] = swap32_S (w[13]);
+    w3[2] = swap32_S (w[14]);
+    w3[3] = swap32_S (w[15]);
+  }
+
+  sha1_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
 void sha1_hmac_update_64 (sha1_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
 {
   sha1_update_64 (&ctx->ipad, w0, w1, w2, w3, len);

OpenCL/inc_hash_sha256.cl

@@ -788,7 +788,7 @@ typedef struct sha256_hmac_ctx
 
 } sha256_hmac_ctx_t;
 
-void sha256_hmac_init (sha256_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4])
+void sha256_hmac_init_64 (sha256_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4])
 {
   u32 t0[4];
   u32 t1[4];
@@ -842,6 +842,177 @@ void sha256_hmac_init (sha256_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4],
   sha256_update_64 (&ctx->opad, t0, t1, t2, t3, 64);
 }
 
+void sha256_hmac_init (sha256_hmac_ctx_t *ctx, const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    sha256_ctx_t tmp;
+
+    sha256_init (&tmp);
+
+    sha256_update (&tmp, w, len);
+
+    sha256_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = tmp.h[5];
+    w1[2] = tmp.h[6];
+    w1[3] = tmp.h[7];
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  sha256_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
+void sha256_hmac_init_global (sha256_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    sha256_ctx_t tmp;
+
+    sha256_init (&tmp);
+
+    sha256_update_global (&tmp, w, len);
+
+    sha256_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = tmp.h[5];
+    w1[2] = tmp.h[6];
+    w1[3] = tmp.h[7];
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  sha256_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
+void sha256_hmac_init_global_swap (sha256_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    sha256_ctx_t tmp;
+
+    sha256_init (&tmp);
+
+    sha256_update_global_swap (&tmp, w, len);
+
+    sha256_final (&tmp);
+
+    w0[0] = tmp.h[0];
+    w0[1] = tmp.h[1];
+    w0[2] = tmp.h[2];
+    w0[3] = tmp.h[3];
+    w1[0] = tmp.h[4];
+    w1[1] = tmp.h[5];
+    w1[2] = tmp.h[6];
+    w1[3] = tmp.h[7];
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+  }
+  else
+  {
+    w0[0] = swap32_S (w[ 0]);
+    w0[1] = swap32_S (w[ 1]);
+    w0[2] = swap32_S (w[ 2]);
+    w0[3] = swap32_S (w[ 3]);
+    w1[0] = swap32_S (w[ 4]);
+    w1[1] = swap32_S (w[ 5]);
+    w1[2] = swap32_S (w[ 6]);
+    w1[3] = swap32_S (w[ 7]);
+    w2[0] = swap32_S (w[ 8]);
+    w2[1] = swap32_S (w[ 9]);
+    w2[2] = swap32_S (w[10]);
+    w2[3] = swap32_S (w[11]);
+    w3[0] = swap32_S (w[12]);
+    w3[1] = swap32_S (w[13]);
+    w3[2] = swap32_S (w[14]);
+    w3[3] = swap32_S (w[15]);
+  }
+
+  sha256_hmac_init_64 (ctx, w0, w1, w2, w3);
+}
+
 void sha256_hmac_update_64 (sha256_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
 {
   sha256_update_64 (&ctx->ipad, w0, w1, w2, w3, len);

OpenCL/inc_hash_sha384.cl

@@ -1050,7 +1050,7 @@ typedef struct sha384_hmac_ctx
 
 } sha384_hmac_ctx_t;
 
-void sha384_hmac_init (sha384_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], const u32 w4[4], const u32 w5[4], const u32 w6[4], const u32 w7[4])
+void sha384_hmac_init_128 (sha384_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], const u32 w4[4], const u32 w5[4], const u32 w6[4], const u32 w7[4])
 {
   u32 t0[4];
   u32 t1[4];
@@ -1140,6 +1140,285 @@ void sha384_hmac_init (sha384_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4],
   sha384_update_128 (&ctx->opad, t0, t1, t2, t3, t4, t5, t6, t7, 128);
 }
 
+void sha384_hmac_init (sha384_hmac_ctx_t *ctx, const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  if (len > 128)
+  {
+    sha384_ctx_t tmp;
+
+    sha384_init (&tmp);
+
+    sha384_update (&tmp, w, len);
+
+    sha384_final (&tmp);
+
+    w0[0] = h32_from_64_S (tmp.h[0]);
+    w0[1] = l32_from_64_S (tmp.h[0]);
+    w0[2] = h32_from_64_S (tmp.h[1]);
+    w0[3] = l32_from_64_S (tmp.h[1]);
+    w1[0] = h32_from_64_S (tmp.h[2]);
+    w1[1] = l32_from_64_S (tmp.h[2]);
+    w1[2] = h32_from_64_S (tmp.h[3]);
+    w1[3] = l32_from_64_S (tmp.h[3]);
+    w2[0] = h32_from_64_S (tmp.h[4]);
+    w2[1] = l32_from_64_S (tmp.h[4]);
+    w2[2] = h32_from_64_S (tmp.h[5]);
+    w2[3] = l32_from_64_S (tmp.h[5]);
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+    w4[0] = 0;
+    w4[1] = 0;
+    w4[2] = 0;
+    w4[3] = 0;
+    w5[0] = 0;
+    w5[1] = 0;
+    w5[2] = 0;
+    w5[3] = 0;
+    w6[0] = 0;
+    w6[1] = 0;
+    w6[2] = 0;
+    w6[3] = 0;
+    w7[0] = 0;
+    w7[1] = 0;
+    w7[2] = 0;
+    w7[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+    w4[0] = w[16];
+    w4[1] = w[17];
+    w4[2] = w[18];
+    w4[3] = w[19];
+    w5[0] = w[20];
+    w5[1] = w[21];
+    w5[2] = w[22];
+    w5[3] = w[23];
+    w6[0] = w[24];
+    w6[1] = w[25];
+    w6[2] = w[26];
+    w6[3] = w[27];
+    w7[0] = w[28];
+    w7[1] = w[29];
+    w7[2] = w[30];
+    w7[3] = w[31];
+  }
+
+  sha384_hmac_init_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+}
+
+void sha384_hmac_init_global (sha384_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  if (len > 128)
+  {
+    sha384_ctx_t tmp;
+
+    sha384_init (&tmp);
+
+    sha384_update_global (&tmp, w, len);
+
+    sha384_final (&tmp);
+
+    w0[0] = h32_from_64_S (tmp.h[0]);
+    w0[1] = l32_from_64_S (tmp.h[0]);
+    w0[2] = h32_from_64_S (tmp.h[1]);
+    w0[3] = l32_from_64_S (tmp.h[1]);
+    w1[0] = h32_from_64_S (tmp.h[2]);
+    w1[1] = l32_from_64_S (tmp.h[2]);
+    w1[2] = h32_from_64_S (tmp.h[3]);
+    w1[3] = l32_from_64_S (tmp.h[3]);
+    w2[0] = h32_from_64_S (tmp.h[4]);
+    w2[1] = l32_from_64_S (tmp.h[4]);
+    w2[2] = h32_from_64_S (tmp.h[5]);
+    w2[3] = l32_from_64_S (tmp.h[5]);
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+    w4[0] = 0;
+    w4[1] = 0;
+    w4[2] = 0;
+    w4[3] = 0;
+    w5[0] = 0;
+    w5[1] = 0;
+    w5[2] = 0;
+    w5[3] = 0;
+    w6[0] = 0;
+    w6[1] = 0;
+    w6[2] = 0;
+    w6[3] = 0;
+    w7[0] = 0;
+    w7[1] = 0;
+    w7[2] = 0;
+    w7[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+    w4[0] = w[16];
+    w4[1] = w[17];
+    w4[2] = w[18];
+    w4[3] = w[19];
+    w5[0] = w[20];
+    w5[1] = w[21];
+    w5[2] = w[22];
+    w5[3] = w[23];
+    w6[0] = w[24];
+    w6[1] = w[25];
+    w6[2] = w[26];
+    w6[3] = w[27];
+    w7[0] = w[28];
+    w7[1] = w[29];
+    w7[2] = w[30];
+    w7[3] = w[31];
+  }
+
+  sha384_hmac_init_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+}
+
+void sha256_hmac_init_global_swap (sha256_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  if (len > 128)
+  {
+    sha384_ctx_t tmp;
+
+    sha384_init (&tmp);
+
+    sha384_update_global_swap (&tmp, w, len);
+
+    sha384_final (&tmp);
+
+    w0[0] = h32_from_64_S (tmp.h[0]);
+    w0[1] = l32_from_64_S (tmp.h[0]);
+    w0[2] = h32_from_64_S (tmp.h[1]);
+    w0[3] = l32_from_64_S (tmp.h[1]);
+    w1[0] = h32_from_64_S (tmp.h[2]);
+    w1[1] = l32_from_64_S (tmp.h[2]);
+    w1[2] = h32_from_64_S (tmp.h[3]);
+    w1[3] = l32_from_64_S (tmp.h[3]);
+    w2[0] = h32_from_64_S (tmp.h[4]);
+    w2[1] = l32_from_64_S (tmp.h[4]);
+    w2[2] = h32_from_64_S (tmp.h[5]);
+    w2[3] = l32_from_64_S (tmp.h[5]);
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+    w4[0] = 0;
+    w4[1] = 0;
+    w4[2] = 0;
+    w4[3] = 0;
+    w5[0] = 0;
+    w5[1] = 0;
+    w5[2] = 0;
+    w5[3] = 0;
+    w6[0] = 0;
+    w6[1] = 0;
+    w6[2] = 0;
+    w6[3] = 0;
+    w7[0] = 0;
+    w7[1] = 0;
+    w7[2] = 0;
+    w7[3] = 0;
+  }
+  else
+  {
+    w0[0] = swap32_S (w[ 0]);
+    w0[1] = swap32_S (w[ 1]);
+    w0[2] = swap32_S (w[ 2]);
+    w0[3] = swap32_S (w[ 3]);
+    w1[0] = swap32_S (w[ 4]);
+    w1[1] = swap32_S (w[ 5]);
+    w1[2] = swap32_S (w[ 6]);
+    w1[3] = swap32_S (w[ 7]);
+    w2[0] = swap32_S (w[ 8]);
+    w2[1] = swap32_S (w[ 9]);
+    w2[2] = swap32_S (w[10]);
+    w2[3] = swap32_S (w[11]);
+    w3[0] = swap32_S (w[12]);
+    w3[1] = swap32_S (w[13]);
+    w3[2] = swap32_S (w[14]);
+    w3[3] = swap32_S (w[15]);
+    w4[0] = swap32_S (w[16]);
+    w4[1] = swap32_S (w[17]);
+    w4[2] = swap32_S (w[18]);
+    w4[3] = swap32_S (w[19]);
+    w5[0] = swap32_S (w[20]);
+    w5[1] = swap32_S (w[21]);
+    w5[2] = swap32_S (w[22]);
+    w5[3] = swap32_S (w[23]);
+    w6[0] = swap32_S (w[24]);
+    w6[1] = swap32_S (w[25]);
+    w6[2] = swap32_S (w[26]);
+    w6[3] = swap32_S (w[27]);
+    w7[0] = swap32_S (w[28]);
+    w7[1] = swap32_S (w[29]);
+    w7[2] = swap32_S (w[30]);
+    w7[3] = swap32_S (w[31]);
+  }
+
+  sha384_hmac_init_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+}
+
 void sha384_hmac_update_128 (sha384_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u32 w5[4], u32 w6[4], u32 w7[4], const int len)
 {
   sha384_update_128 (&ctx->ipad, w0, w1, w2, w3, w4, w5, w6, w7, len);
@@ -1195,10 +1474,10 @@ void sha384_hmac_final (sha384_hmac_ctx_t *ctx)
   t2[1] = l32_from_64_S (ctx->ipad.h[4]);
   t2[2] = h32_from_64_S (ctx->ipad.h[5]);
   t2[3] = l32_from_64_S (ctx->ipad.h[5]);
-  t3[0] = h32_from_64_S (ctx->ipad.h[6]);
-  t3[1] = l32_from_64_S (ctx->ipad.h[6]);
-  t3[2] = h32_from_64_S (ctx->ipad.h[7]);
-  t3[3] = l32_from_64_S (ctx->ipad.h[7]);
+  t3[0] = 0;
+  t3[1] = 0;
+  t3[2] = 0;
+  t3[3] = 0;
   t4[0] = 0;
   t4[1] = 0;
   t4[2] = 0;

OpenCL/inc_hash_sha512.cl

@@ -1050,7 +1050,7 @@ typedef struct sha512_hmac_ctx
 
 } sha512_hmac_ctx_t;
 
-void sha512_hmac_init (sha512_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], const u32 w4[4], const u32 w5[4], const u32 w6[4], const u32 w7[4])
+void sha512_hmac_init_128 (sha512_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], const u32 w4[4], const u32 w5[4], const u32 w6[4], const u32 w7[4])
 {
   u32 t0[4];
   u32 t1[4];
@@ -1140,6 +1140,285 @@ void sha512_hmac_init (sha512_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4],
   sha512_update_128 (&ctx->opad, t0, t1, t2, t3, t4, t5, t6, t7, 128);
 }
 
+void sha512_hmac_init (sha512_hmac_ctx_t *ctx, const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  if (len > 128)
+  {
+    sha512_ctx_t tmp;
+
+    sha512_init (&tmp);
+
+    sha512_update (&tmp, w, len);
+
+    sha512_final (&tmp);
+
+    w0[0] = h32_from_64_S (tmp.h[0]);
+    w0[1] = l32_from_64_S (tmp.h[0]);
+    w0[2] = h32_from_64_S (tmp.h[1]);
+    w0[3] = l32_from_64_S (tmp.h[1]);
+    w1[0] = h32_from_64_S (tmp.h[2]);
+    w1[1] = l32_from_64_S (tmp.h[2]);
+    w1[2] = h32_from_64_S (tmp.h[3]);
+    w1[3] = l32_from_64_S (tmp.h[3]);
+    w2[0] = h32_from_64_S (tmp.h[4]);
+    w2[1] = l32_from_64_S (tmp.h[4]);
+    w2[2] = h32_from_64_S (tmp.h[5]);
+    w2[3] = l32_from_64_S (tmp.h[5]);
+    w3[0] = h32_from_64_S (tmp.h[6]);
+    w3[1] = l32_from_64_S (tmp.h[6]);
+    w3[2] = h32_from_64_S (tmp.h[7]);
+    w3[3] = l32_from_64_S (tmp.h[7]);
+    w4[0] = 0;
+    w4[1] = 0;
+    w4[2] = 0;
+    w4[3] = 0;
+    w5[0] = 0;
+    w5[1] = 0;
+    w5[2] = 0;
+    w5[3] = 0;
+    w6[0] = 0;
+    w6[1] = 0;
+    w6[2] = 0;
+    w6[3] = 0;
+    w7[0] = 0;
+    w7[1] = 0;
+    w7[2] = 0;
+    w7[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+    w4[0] = w[16];
+    w4[1] = w[17];
+    w4[2] = w[18];
+    w4[3] = w[19];
+    w5[0] = w[20];
+    w5[1] = w[21];
+    w5[2] = w[22];
+    w5[3] = w[23];
+    w6[0] = w[24];
+    w6[1] = w[25];
+    w6[2] = w[26];
+    w6[3] = w[27];
+    w7[0] = w[28];
+    w7[1] = w[29];
+    w7[2] = w[30];
+    w7[3] = w[31];
+  }
+
+  sha512_hmac_init_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+}
+
+void sha512_hmac_init_global (sha512_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  if (len > 128)
+  {
+    sha512_ctx_t tmp;
+
+    sha512_init (&tmp);
+
+    sha512_update_global (&tmp, w, len);
+
+    sha512_final (&tmp);
+
+    w0[0] = h32_from_64_S (tmp.h[0]);
+    w0[1] = l32_from_64_S (tmp.h[0]);
+    w0[2] = h32_from_64_S (tmp.h[1]);
+    w0[3] = l32_from_64_S (tmp.h[1]);
+    w1[0] = h32_from_64_S (tmp.h[2]);
+    w1[1] = l32_from_64_S (tmp.h[2]);
+    w1[2] = h32_from_64_S (tmp.h[3]);
+    w1[3] = l32_from_64_S (tmp.h[3]);
+    w2[0] = h32_from_64_S (tmp.h[4]);
+    w2[1] = l32_from_64_S (tmp.h[4]);
+    w2[2] = h32_from_64_S (tmp.h[5]);
+    w2[3] = l32_from_64_S (tmp.h[5]);
+    w3[0] = h32_from_64_S (tmp.h[6]);
+    w3[1] = l32_from_64_S (tmp.h[6]);
+    w3[2] = h32_from_64_S (tmp.h[7]);
+    w3[3] = l32_from_64_S (tmp.h[7]);
+    w4[0] = 0;
+    w4[1] = 0;
+    w4[2] = 0;
+    w4[3] = 0;
+    w5[0] = 0;
+    w5[1] = 0;
+    w5[2] = 0;
+    w5[3] = 0;
+    w6[0] = 0;
+    w6[1] = 0;
+    w6[2] = 0;
+    w6[3] = 0;
+    w7[0] = 0;
+    w7[1] = 0;
+    w7[2] = 0;
+    w7[3] = 0;
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+    w4[0] = w[16];
+    w4[1] = w[17];
+    w4[2] = w[18];
+    w4[3] = w[19];
+    w5[0] = w[20];
+    w5[1] = w[21];
+    w5[2] = w[22];
+    w5[3] = w[23];
+    w6[0] = w[24];
+    w6[1] = w[25];
+    w6[2] = w[26];
+    w6[3] = w[27];
+    w7[0] = w[28];
+    w7[1] = w[29];
+    w7[2] = w[30];
+    w7[3] = w[31];
+  }
+
+  sha512_hmac_init_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+}
+
+void sha512_hmac_init_global_swap (sha512_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  if (len > 128)
+  {
+    sha512_ctx_t tmp;
+
+    sha512_init (&tmp);
+
+    sha512_update_global_swap (&tmp, w, len);
+
+    sha512_final (&tmp);
+
+    w0[0] = h32_from_64_S (tmp.h[0]);
+    w0[1] = l32_from_64_S (tmp.h[0]);
+    w0[2] = h32_from_64_S (tmp.h[1]);
+    w0[3] = l32_from_64_S (tmp.h[1]);
+    w1[0] = h32_from_64_S (tmp.h[2]);
+    w1[1] = l32_from_64_S (tmp.h[2]);
+    w1[2] = h32_from_64_S (tmp.h[3]);
+    w1[3] = l32_from_64_S (tmp.h[3]);
+    w2[0] = h32_from_64_S (tmp.h[4]);
+    w2[1] = l32_from_64_S (tmp.h[4]);
+    w2[2] = h32_from_64_S (tmp.h[5]);
+    w2[3] = l32_from_64_S (tmp.h[5]);
+    w3[0] = h32_from_64_S (tmp.h[6]);
+    w3[1] = l32_from_64_S (tmp.h[6]);
+    w3[2] = h32_from_64_S (tmp.h[7]);
+    w3[3] = l32_from_64_S (tmp.h[7]);
+    w4[0] = 0;
+    w4[1] = 0;
+    w4[2] = 0;
+    w4[3] = 0;
+    w5[0] = 0;
+    w5[1] = 0;
+    w5[2] = 0;
+    w5[3] = 0;
+    w6[0] = 0;
+    w6[1] = 0;
+    w6[2] = 0;
+    w6[3] = 0;
+    w7[0] = 0;
+    w7[1] = 0;
+    w7[2] = 0;
+    w7[3] = 0;
+  }
+  else
+  {
+    w0[0] = swap32_S (w[ 0]);
+    w0[1] = swap32_S (w[ 1]);
+    w0[2] = swap32_S (w[ 2]);
+    w0[3] = swap32_S (w[ 3]);
+    w1[0] = swap32_S (w[ 4]);
+    w1[1] = swap32_S (w[ 5]);
+    w1[2] = swap32_S (w[ 6]);
+    w1[3] = swap32_S (w[ 7]);
+    w2[0] = swap32_S (w[ 8]);
+    w2[1] = swap32_S (w[ 9]);
+    w2[2] = swap32_S (w[10]);
+    w2[3] = swap32_S (w[11]);
+    w3[0] = swap32_S (w[12]);
+    w3[1] = swap32_S (w[13]);
+    w3[2] = swap32_S (w[14]);
+    w3[3] = swap32_S (w[15]);
+    w4[0] = swap32_S (w[16]);
+    w4[1] = swap32_S (w[17]);
+    w4[2] = swap32_S (w[18]);
+    w4[3] = swap32_S (w[19]);
+    w5[0] = swap32_S (w[20]);
+    w5[1] = swap32_S (w[21]);
+    w5[2] = swap32_S (w[22]);
+    w5[3] = swap32_S (w[23]);
+    w6[0] = swap32_S (w[24]);
+    w6[1] = swap32_S (w[25]);
+    w6[2] = swap32_S (w[26]);
+    w6[3] = swap32_S (w[27]);
+    w7[0] = swap32_S (w[28]);
+    w7[1] = swap32_S (w[29]);
+    w7[2] = swap32_S (w[30]);
+    w7[3] = swap32_S (w[31]);
+  }
+
+  sha512_hmac_init_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+}
+
 void sha512_hmac_update_128 (sha512_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u32 w5[4], u32 w6[4], u32 w7[4], const int len)
 {
   sha512_update_128 (&ctx->ipad, w0, w1, w2, w3, w4, w5, w6, w7, len);

OpenCL/inc_hash_whirlpool.cl

@@ -1847,7 +1847,7 @@ typedef struct whirlpool_hmac_ctx
 
 } whirlpool_hmac_ctx_t;
 
-void whirlpool_hmac_init (whirlpool_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
+void whirlpool_hmac_init_64 (whirlpool_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
 {
   u32 t0[4];
   u32 t1[4];
@@ -1901,6 +1901,177 @@ void whirlpool_hmac_init (whirlpool_hmac_ctx_t *ctx, const u32 w0[4], const u32
   whirlpool_update_64 (&ctx->opad, t0, t1, t2, t3, 64, s_Ch, s_Cl);
 }
 
+void whirlpool_hmac_init (whirlpool_hmac_ctx_t *ctx, const u32 *w, const int len, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    whirlpool_ctx_t tmp;
+
+    whirlpool_init (&tmp);
+
+    whirlpool_update (&tmp, w, len, s_Ch, s_Cl);
+
+    whirlpool_final (&tmp, s_Ch, s_Cl);
+
+    w0[0] = tmp.h[ 0];
+    w0[1] = tmp.h[ 1];
+    w0[2] = tmp.h[ 2];
+    w0[3] = tmp.h[ 3];
+    w1[0] = tmp.h[ 4];
+    w1[1] = tmp.h[ 5];
+    w1[2] = tmp.h[ 6];
+    w1[3] = tmp.h[ 7];
+    w2[0] = tmp.h[ 8];
+    w2[1] = tmp.h[ 9];
+    w2[2] = tmp.h[10];
+    w2[3] = tmp.h[11];
+    w3[0] = tmp.h[12];
+    w3[1] = tmp.h[13];
+    w3[2] = tmp.h[14];
+    w3[3] = tmp.h[15];
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  whirlpool_hmac_init_64 (ctx, w0, w1, w2, w3, s_Ch, s_Cl);
+}
+
+void whirlpool_hmac_init_global (whirlpool_hmac_ctx_t *ctx, __global const u32 *w, const int len, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    whirlpool_ctx_t tmp;
+
+    whirlpool_init (&tmp);
+
+    whirlpool_update_global (&tmp, w, len, s_Ch, s_Cl);
+
+    whirlpool_final (&tmp, s_Ch, s_Cl);
+
+    w0[0] = tmp.h[ 0];
+    w0[1] = tmp.h[ 1];
+    w0[2] = tmp.h[ 2];
+    w0[3] = tmp.h[ 3];
+    w1[0] = tmp.h[ 4];
+    w1[1] = tmp.h[ 5];
+    w1[2] = tmp.h[ 6];
+    w1[3] = tmp.h[ 7];
+    w2[0] = tmp.h[ 8];
+    w2[1] = tmp.h[ 9];
+    w2[2] = tmp.h[10];
+    w2[3] = tmp.h[11];
+    w3[0] = tmp.h[12];
+    w3[1] = tmp.h[13];
+    w3[2] = tmp.h[14];
+    w3[3] = tmp.h[15];
+  }
+  else
+  {
+    w0[0] = w[ 0];
+    w0[1] = w[ 1];
+    w0[2] = w[ 2];
+    w0[3] = w[ 3];
+    w1[0] = w[ 4];
+    w1[1] = w[ 5];
+    w1[2] = w[ 6];
+    w1[3] = w[ 7];
+    w2[0] = w[ 8];
+    w2[1] = w[ 9];
+    w2[2] = w[10];
+    w2[3] = w[11];
+    w3[0] = w[12];
+    w3[1] = w[13];
+    w3[2] = w[14];
+    w3[3] = w[15];
+  }
+
+  whirlpool_hmac_init_64 (ctx, w0, w1, w2, w3, s_Ch, s_Cl);
+}
+
+void whirlpool_hmac_init_global_swap (whirlpool_hmac_ctx_t *ctx, __global const u32 *w, const int len, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  if (len > 64)
+  {
+    whirlpool_ctx_t tmp;
+
+    whirlpool_init (&tmp);
+
+    whirlpool_update_global_swap (&tmp, w, len, s_Ch, s_Cl);
+
+    whirlpool_final (&tmp, s_Ch, s_Cl);
+
+    w0[0] = tmp.h[ 0];
+    w0[1] = tmp.h[ 1];
+    w0[2] = tmp.h[ 2];
+    w0[3] = tmp.h[ 3];
+    w1[0] = tmp.h[ 4];
+    w1[1] = tmp.h[ 5];
+    w1[2] = tmp.h[ 6];
+    w1[3] = tmp.h[ 7];
+    w2[0] = tmp.h[ 8];
+    w2[1] = tmp.h[ 9];
+    w2[2] = tmp.h[10];
+    w2[3] = tmp.h[11];
+    w3[0] = tmp.h[12];
+    w3[1] = tmp.h[13];
+    w3[2] = tmp.h[14];
+    w3[3] = tmp.h[15];
+  }
+  else
+  {
+    w0[0] = swap32_S (w[ 0]);
+    w0[1] = swap32_S (w[ 1]);
+    w0[2] = swap32_S (w[ 2]);
+    w0[3] = swap32_S (w[ 3]);
+    w1[0] = swap32_S (w[ 4]);
+    w1[1] = swap32_S (w[ 5]);
+    w1[2] = swap32_S (w[ 6]);
+    w1[3] = swap32_S (w[ 7]);
+    w2[0] = swap32_S (w[ 8]);
+    w2[1] = swap32_S (w[ 9]);
+    w2[2] = swap32_S (w[10]);
+    w2[3] = swap32_S (w[11]);
+    w3[0] = swap32_S (w[12]);
+    w3[1] = swap32_S (w[13]);
+    w3[2] = swap32_S (w[14]);
+    w3[3] = swap32_S (w[15]);
+  }
+
+  whirlpool_hmac_init_64 (ctx, w0, w1, w2, w3, s_Ch, s_Cl);
+}
+
 void whirlpool_hmac_update_64 (whirlpool_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
 {
   whirlpool_update_64 (&ctx->ipad, w0, w1, w2, w3, len, s_Ch, s_Cl);

OpenCL/inc_types.cl

@@ -279,31 +279,12 @@ inline u32 amd_bytealign_S (const u32 a, const u32 b, const u32 c)
 #ifdef IS_NV
 inline u32 swap32_S (const u32 v)
 {
-  u32 r;
-
-  asm ("prmt.b32 %0, %1, 0, 0x0123;" : "=r"(r) : "r"(v));
-
-  return r;
+  return (as_uint (as_uchar4 (v).s3210));
 }
 
 inline u64 swap64_S (const u64 v)
 {
-  u32 il;
-  u32 ir;
-
-  asm ("mov.b64 {%0, %1}, %2;" : "=r"(il), "=r"(ir) : "l"(v));
-
-  u32 tl;
-  u32 tr;
-
-  asm ("prmt.b32 %0, %1, 0, 0x0123;" : "=r"(tl) : "r"(il));
-  asm ("prmt.b32 %0, %1, 0, 0x0123;" : "=r"(tr) : "r"(ir));
-
-  u64 r;
-
-  asm ("mov.b64 %0, {%1, %2};" : "=l"(r) : "r"(tr), "r"(tl));
-
-  return r;
+  return (as_ulong (as_uchar8 (v).s76543210));
 }
 
 inline u32 rotr32_S (const u32 a, const u32 n)

OpenCL/m02100.cl

@@ -119,7 +119,7 @@ __kernel void m02100_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_64 (&sha1_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];

OpenCL/m02500.cl

@@ -63,48 +63,9 @@ __kernel void m02500_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -124,6 +85,11 @@ __kernel void m02500_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;
@@ -365,7 +331,7 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     sha1_hmac_ctx_t ctx1;
 
-    sha1_hmac_init (&ctx1, w0, w1, w2, w3);
+    sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3);
 
     sha1_hmac_update (&ctx1, pke, 100);
 
@@ -441,7 +407,7 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
       sha1_hmac_ctx_t ctx2;
 
-      sha1_hmac_init (&ctx2, t0, t1, t2, t3);
+      sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3);
 
       sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len);
 
@@ -515,7 +481,7 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     sha1_hmac_ctx_t ctx1;
 
-    sha1_hmac_init (&ctx1, w0, w1, w2, w3);
+    sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3);
 
     sha1_hmac_update (&ctx1, pke, 100);
 
@@ -591,7 +557,7 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
       sha1_hmac_ctx_t ctx2;
 
-      sha1_hmac_init (&ctx2, t0, t1, t2, t3);
+      sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3);
 
       sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len);
 

OpenCL/m02501.cl

@@ -206,7 +206,7 @@ __kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     sha1_hmac_ctx_t ctx1;
 
-    sha1_hmac_init (&ctx1, w0, w1, w2, w3);
+    sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3);
 
     sha1_hmac_update (&ctx1, pke, 100);
 
@@ -282,7 +282,7 @@ __kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
       sha1_hmac_ctx_t ctx2;
 
-      sha1_hmac_init (&ctx2, t0, t1, t2, t3);
+      sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3);
 
       sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len);
 
@@ -356,7 +356,7 @@ __kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     sha1_hmac_ctx_t ctx1;
 
-    sha1_hmac_init (&ctx1, w0, w1, w2, w3);
+    sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3);
 
     sha1_hmac_update (&ctx1, pke, 100);
 
@@ -432,7 +432,7 @@ __kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
       sha1_hmac_ctx_t ctx2;
 
-      sha1_hmac_init (&ctx2, t0, t1, t2, t3);
+      sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3);
 
       sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len);
 

OpenCL/m06211.cl

@@ -132,7 +132,7 @@ __kernel void m06211_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   ripemd160_hmac_ctx_t ripemd160_hmac_ctx;
 
-  ripemd160_hmac_init (&ripemd160_hmac_ctx, w0, w1, w2, w3);
+  ripemd160_hmac_init_64 (&ripemd160_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = ripemd160_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = ripemd160_hmac_ctx.ipad.h[1];

OpenCL/m06212.cl

@@ -132,7 +132,7 @@ __kernel void m06212_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   ripemd160_hmac_ctx_t ripemd160_hmac_ctx;
 
-  ripemd160_hmac_init (&ripemd160_hmac_ctx, w0, w1, w2, w3);
+  ripemd160_hmac_init_64 (&ripemd160_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = ripemd160_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = ripemd160_hmac_ctx.ipad.h[1];

OpenCL/m06213.cl

@@ -132,7 +132,7 @@ __kernel void m06213_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   ripemd160_hmac_ctx_t ripemd160_hmac_ctx;
 
-  ripemd160_hmac_init (&ripemd160_hmac_ctx, w0, w1, w2, w3);
+  ripemd160_hmac_init_64 (&ripemd160_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = ripemd160_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = ripemd160_hmac_ctx.ipad.h[1];

OpenCL/m06221.cl

@@ -207,7 +207,7 @@ __kernel void m06221_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_128 (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
 
   tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
@@ -477,25 +477,25 @@ __kernel void m06221_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
   u32 ukey1[8];
 
-  ukey1[0] = swap32_S (h32_from_64 (tmps[gid].out[ 0]));
-  ukey1[1] = swap32_S (l32_from_64 (tmps[gid].out[ 0]));
-  ukey1[2] = swap32_S (h32_from_64 (tmps[gid].out[ 1]));
-  ukey1[3] = swap32_S (l32_from_64 (tmps[gid].out[ 1]));
-  ukey1[4] = swap32_S (h32_from_64 (tmps[gid].out[ 2]));
-  ukey1[5] = swap32_S (l32_from_64 (tmps[gid].out[ 2]));
-  ukey1[6] = swap32_S (h32_from_64 (tmps[gid].out[ 3]));
-  ukey1[7] = swap32_S (l32_from_64 (tmps[gid].out[ 3]));
+  ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0]));
+  ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0]));
+  ukey1[2] = swap32_S (h32_from_64_S (tmps[gid].out[1]));
+  ukey1[3] = swap32_S (l32_from_64_S (tmps[gid].out[1]));
+  ukey1[4] = swap32_S (h32_from_64_S (tmps[gid].out[2]));
+  ukey1[5] = swap32_S (l32_from_64_S (tmps[gid].out[2]));
+  ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3]));
+  ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3]));
 
   u32 ukey2[8];
 
-  ukey2[0] = swap32_S (h32_from_64 (tmps[gid].out[ 4]));
-  ukey2[1] = swap32_S (l32_from_64 (tmps[gid].out[ 4]));
-  ukey2[2] = swap32_S (h32_from_64 (tmps[gid].out[ 5]));
-  ukey2[3] = swap32_S (l32_from_64 (tmps[gid].out[ 5]));
-  ukey2[4] = swap32_S (h32_from_64 (tmps[gid].out[ 6]));
-  ukey2[5] = swap32_S (l32_from_64 (tmps[gid].out[ 6]));
-  ukey2[6] = swap32_S (h32_from_64 (tmps[gid].out[ 7]));
-  ukey2[7] = swap32_S (l32_from_64 (tmps[gid].out[ 7]));
+  ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4]));
+  ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4]));
+  ukey2[2] = swap32_S (h32_from_64_S (tmps[gid].out[5]));
+  ukey2[3] = swap32_S (l32_from_64_S (tmps[gid].out[5]));
+  ukey2[4] = swap32_S (h32_from_64_S (tmps[gid].out[6]));
+  ukey2[5] = swap32_S (l32_from_64_S (tmps[gid].out[6]));
+  ukey2[6] = swap32_S (h32_from_64_S (tmps[gid].out[7]));
+  ukey2[7] = swap32_S (l32_from_64_S (tmps[gid].out[7]));
 
   if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
   {

OpenCL/m06222.cl

@@ -207,7 +207,7 @@ __kernel void m06222_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_128 (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
 
   tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
@@ -481,14 +481,14 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   u32 ukey1[8];
   #endif
 
-  ukey1[0] = swap32_S (h32_from_64 (tmps[gid].out[ 0]));
-  ukey1[1] = swap32_S (l32_from_64 (tmps[gid].out[ 0]));
-  ukey1[2] = swap32_S (h32_from_64 (tmps[gid].out[ 1]));
-  ukey1[3] = swap32_S (l32_from_64 (tmps[gid].out[ 1]));
-  ukey1[4] = swap32_S (h32_from_64 (tmps[gid].out[ 2]));
-  ukey1[5] = swap32_S (l32_from_64 (tmps[gid].out[ 2]));
-  ukey1[6] = swap32_S (h32_from_64 (tmps[gid].out[ 3]));
-  ukey1[7] = swap32_S (l32_from_64 (tmps[gid].out[ 3]));
+  ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0]));
+  ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0]));
+  ukey1[2] = swap32_S (h32_from_64_S (tmps[gid].out[1]));
+  ukey1[3] = swap32_S (l32_from_64_S (tmps[gid].out[1]));
+  ukey1[4] = swap32_S (h32_from_64_S (tmps[gid].out[2]));
+  ukey1[5] = swap32_S (l32_from_64_S (tmps[gid].out[2]));
+  ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3]));
+  ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3]));
 
   #if defined (IS_APPLE) && defined (IS_GPU)
   volatile u32 ukey2[8];
@@ -496,14 +496,14 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   u32 ukey2[8];
   #endif
 
-  ukey2[0] = swap32_S (h32_from_64 (tmps[gid].out[ 4]));
-  ukey2[1] = swap32_S (l32_from_64 (tmps[gid].out[ 4]));
-  ukey2[2] = swap32_S (h32_from_64 (tmps[gid].out[ 5]));
-  ukey2[3] = swap32_S (l32_from_64 (tmps[gid].out[ 5]));
-  ukey2[4] = swap32_S (h32_from_64 (tmps[gid].out[ 6]));
-  ukey2[5] = swap32_S (l32_from_64 (tmps[gid].out[ 6]));
-  ukey2[6] = swap32_S (h32_from_64 (tmps[gid].out[ 7]));
-  ukey2[7] = swap32_S (l32_from_64 (tmps[gid].out[ 7]));
+  ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4]));
+  ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4]));
+  ukey2[2] = swap32_S (h32_from_64_S (tmps[gid].out[5]));
+  ukey2[3] = swap32_S (l32_from_64_S (tmps[gid].out[5]));
+  ukey2[4] = swap32_S (h32_from_64_S (tmps[gid].out[6]));
+  ukey2[5] = swap32_S (l32_from_64_S (tmps[gid].out[6]));
+  ukey2[6] = swap32_S (h32_from_64_S (tmps[gid].out[7]));
+  ukey2[7] = swap32_S (l32_from_64_S (tmps[gid].out[7]));
 
   if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
   {

OpenCL/m06223.cl

@@ -207,7 +207,7 @@ __kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_128 (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
 
   tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
@@ -481,14 +481,14 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   u32 ukey1[8];
   #endif
 
-  ukey1[0] = swap32_S (h32_from_64 (tmps[gid].out[ 0]));
-  ukey1[1] = swap32_S (l32_from_64 (tmps[gid].out[ 0]));
-  ukey1[2] = swap32_S (h32_from_64 (tmps[gid].out[ 1]));
-  ukey1[3] = swap32_S (l32_from_64 (tmps[gid].out[ 1]));
-  ukey1[4] = swap32_S (h32_from_64 (tmps[gid].out[ 2]));
-  ukey1[5] = swap32_S (l32_from_64 (tmps[gid].out[ 2]));
-  ukey1[6] = swap32_S (h32_from_64 (tmps[gid].out[ 3]));
-  ukey1[7] = swap32_S (l32_from_64 (tmps[gid].out[ 3]));
+  ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0]));
+  ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0]));
+  ukey1[2] = swap32_S (h32_from_64_S (tmps[gid].out[1]));
+  ukey1[3] = swap32_S (l32_from_64_S (tmps[gid].out[1]));
+  ukey1[4] = swap32_S (h32_from_64_S (tmps[gid].out[2]));
+  ukey1[5] = swap32_S (l32_from_64_S (tmps[gid].out[2]));
+  ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3]));
+  ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3]));
 
   #if defined (IS_APPLE) && defined (IS_GPU)
   volatile u32 ukey2[8];
@@ -496,14 +496,14 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   u32 ukey2[8];
   #endif
 
-  ukey2[0] = swap32_S (h32_from_64 (tmps[gid].out[ 4]));
-  ukey2[1] = swap32_S (l32_from_64 (tmps[gid].out[ 4]));
-  ukey2[2] = swap32_S (h32_from_64 (tmps[gid].out[ 5]));
-  ukey2[3] = swap32_S (l32_from_64 (tmps[gid].out[ 5]));
-  ukey2[4] = swap32_S (h32_from_64 (tmps[gid].out[ 6]));
-  ukey2[5] = swap32_S (l32_from_64 (tmps[gid].out[ 6]));
-  ukey2[6] = swap32_S (h32_from_64 (tmps[gid].out[ 7]));
-  ukey2[7] = swap32_S (l32_from_64 (tmps[gid].out[ 7]));
+  ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4]));
+  ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4]));
+  ukey2[2] = swap32_S (h32_from_64_S (tmps[gid].out[5]));
+  ukey2[3] = swap32_S (l32_from_64_S (tmps[gid].out[5]));
+  ukey2[4] = swap32_S (h32_from_64_S (tmps[gid].out[6]));
+  ukey2[5] = swap32_S (l32_from_64_S (tmps[gid].out[6]));
+  ukey2[6] = swap32_S (h32_from_64_S (tmps[gid].out[7]));
+  ukey2[7] = swap32_S (l32_from_64_S (tmps[gid].out[7]));
 
   if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
   {

OpenCL/m06231.cl

@@ -241,7 +241,7 @@ __kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   whirlpool_hmac_ctx_t whirlpool_hmac_ctx;
 
-  whirlpool_hmac_init (&whirlpool_hmac_ctx, w0, w1, w2, w3, s_Ch, s_Cl);
+  whirlpool_hmac_init_64 (&whirlpool_hmac_ctx, w0, w1, w2, w3, s_Ch, s_Cl);
 
   tmps[gid].ipad[ 0] = whirlpool_hmac_ctx.ipad.h[ 0];
   tmps[gid].ipad[ 1] = whirlpool_hmac_ctx.ipad.h[ 1];

OpenCL/m06232.cl

@@ -241,7 +241,7 @@ __kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   whirlpool_hmac_ctx_t whirlpool_hmac_ctx;
 
-  whirlpool_hmac_init (&whirlpool_hmac_ctx, w0, w1, w2, w3, s_Ch, s_Cl);
+  whirlpool_hmac_init_64 (&whirlpool_hmac_ctx, w0, w1, w2, w3, s_Ch, s_Cl);
 
   tmps[gid].ipad[ 0] = whirlpool_hmac_ctx.ipad.h[ 0];
   tmps[gid].ipad[ 1] = whirlpool_hmac_ctx.ipad.h[ 1];

OpenCL/m06233.cl

@@ -241,7 +241,7 @@ __kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   whirlpool_hmac_ctx_t whirlpool_hmac_ctx;
 
-  whirlpool_hmac_init (&whirlpool_hmac_ctx, w0, w1, w2, w3, s_Ch, s_Cl);
+  whirlpool_hmac_init_64 (&whirlpool_hmac_ctx, w0, w1, w2, w3, s_Ch, s_Cl);
 
   tmps[gid].ipad[ 0] = whirlpool_hmac_ctx.ipad.h[ 0];
   tmps[gid].ipad[ 1] = whirlpool_hmac_ctx.ipad.h[ 1];

OpenCL/m06400.cl

@@ -68,48 +68,9 @@ __kernel void m06400_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
@@ -135,6 +96,11 @@ __kernel void m06400_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m06500.cl

@@ -84,84 +84,9 @@ __kernel void m06500_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-  u32 w4[4];
-  u32 w5[4];
-  u32 w6[4];
-  u32 w7[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-  w4[0] = pws[gid].i[16];
-  w4[1] = pws[gid].i[17];
-  w4[2] = pws[gid].i[18];
-  w4[3] = pws[gid].i[19];
-  w5[0] = pws[gid].i[20];
-  w5[1] = pws[gid].i[21];
-  w5[2] = pws[gid].i[22];
-  w5[3] = pws[gid].i[23];
-  w6[0] = pws[gid].i[24];
-  w6[1] = pws[gid].i[25];
-  w6[2] = pws[gid].i[26];
-  w6[3] = pws[gid].i[27];
-  w7[0] = pws[gid].i[28];
-  w7[1] = pws[gid].i[29];
-  w7[2] = pws[gid].i[30];
-  w7[3] = pws[gid].i[31];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-  w4[0] = swap32_S (w4[0]);
-  w4[1] = swap32_S (w4[1]);
-  w4[2] = swap32_S (w4[2]);
-  w4[3] = swap32_S (w4[3]);
-  w5[0] = swap32_S (w5[0]);
-  w5[1] = swap32_S (w5[1]);
-  w5[2] = swap32_S (w5[2]);
-  w5[3] = swap32_S (w5[3]);
-  w6[0] = swap32_S (w6[0]);
-  w6[1] = swap32_S (w6[1]);
-  w6[2] = swap32_S (w6[2]);
-  w6[3] = swap32_S (w6[3]);
-  w7[0] = swap32_S (w7[0]);
-  w7[1] = swap32_S (w7[1]);
-  w7[2] = swap32_S (w7[2]);
-  w7[3] = swap32_S (w7[3]);
-
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
@@ -187,6 +112,15 @@ __kernel void m06500_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha512_hmac_ctx_t sha512_hmac_ctx2 = sha512_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+    u32 w4[4];
+    u32 w5[4];
+    u32 w6[4];
+    u32 w7[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m06600.cl

@@ -61,48 +61,9 @@ __kernel void m06600_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -130,6 +91,11 @@ __kernel void m06600_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m06700.cl

@@ -62,48 +62,9 @@ __kernel void m06700_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -123,6 +84,11 @@ __kernel void m06700_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m06800.cl

@@ -69,48 +69,9 @@ __kernel void m06800_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
@@ -136,6 +97,11 @@ __kernel void m06800_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;
@@ -385,10 +351,10 @@ __kernel void m06800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
     salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 
-    out[0] = swap32 (out[0]);
-    out[1] = swap32 (out[1]);
-    out[2] = swap32 (out[2]);
-    out[3] = swap32 (out[3]);
+    out[0] = swap32_S (out[0]);
+    out[1] = swap32_S (out[1]);
+    out[2] = swap32_S (out[2]);
+    out[3] = swap32_S (out[3]);
 
     truncate_block_4x4_le (out, salt_len);
 

OpenCL/m07100.cl

@@ -84,84 +84,9 @@ __kernel void m07100_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-  u32 w4[4];
-  u32 w5[4];
-  u32 w6[4];
-  u32 w7[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-  w4[0] = pws[gid].i[16];
-  w4[1] = pws[gid].i[17];
-  w4[2] = pws[gid].i[18];
-  w4[3] = pws[gid].i[19];
-  w5[0] = pws[gid].i[20];
-  w5[1] = pws[gid].i[21];
-  w5[2] = pws[gid].i[22];
-  w5[3] = pws[gid].i[23];
-  w6[0] = pws[gid].i[24];
-  w6[1] = pws[gid].i[25];
-  w6[2] = pws[gid].i[26];
-  w6[3] = pws[gid].i[27];
-  w7[0] = pws[gid].i[28];
-  w7[1] = pws[gid].i[29];
-  w7[2] = pws[gid].i[30];
-  w7[3] = pws[gid].i[31];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-  w4[0] = swap32_S (w4[0]);
-  w4[1] = swap32_S (w4[1]);
-  w4[2] = swap32_S (w4[2]);
-  w4[3] = swap32_S (w4[3]);
-  w5[0] = swap32_S (w5[0]);
-  w5[1] = swap32_S (w5[1]);
-  w5[2] = swap32_S (w5[2]);
-  w5[3] = swap32_S (w5[3]);
-  w6[0] = swap32_S (w6[0]);
-  w6[1] = swap32_S (w6[1]);
-  w6[2] = swap32_S (w6[2]);
-  w6[3] = swap32_S (w6[3]);
-  w7[0] = swap32_S (w7[0]);
-  w7[1] = swap32_S (w7[1]);
-  w7[2] = swap32_S (w7[2]);
-  w7[3] = swap32_S (w7[3]);
-
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
@@ -187,6 +112,15 @@ __kernel void m07100_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha512_hmac_ctx_t sha512_hmac_ctx2 = sha512_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+    u32 w4[4];
+    u32 w5[4];
+    u32 w6[4];
+    u32 w7[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m08200.cl

@@ -85,84 +85,9 @@ __kernel void m08200_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-  u32 w4[4];
-  u32 w5[4];
-  u32 w6[4];
-  u32 w7[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-  w4[0] = pws[gid].i[16];
-  w4[1] = pws[gid].i[17];
-  w4[2] = pws[gid].i[18];
-  w4[3] = pws[gid].i[19];
-  w5[0] = pws[gid].i[20];
-  w5[1] = pws[gid].i[21];
-  w5[2] = pws[gid].i[22];
-  w5[3] = pws[gid].i[23];
-  w6[0] = pws[gid].i[24];
-  w6[1] = pws[gid].i[25];
-  w6[2] = pws[gid].i[26];
-  w6[3] = pws[gid].i[27];
-  w7[0] = pws[gid].i[28];
-  w7[1] = pws[gid].i[29];
-  w7[2] = pws[gid].i[30];
-  w7[3] = pws[gid].i[31];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-  w4[0] = swap32_S (w4[0]);
-  w4[1] = swap32_S (w4[1]);
-  w4[2] = swap32_S (w4[2]);
-  w4[3] = swap32_S (w4[3]);
-  w5[0] = swap32_S (w5[0]);
-  w5[1] = swap32_S (w5[1]);
-  w5[2] = swap32_S (w5[2]);
-  w5[3] = swap32_S (w5[3]);
-  w6[0] = swap32_S (w6[0]);
-  w6[1] = swap32_S (w6[1]);
-  w6[2] = swap32_S (w6[2]);
-  w6[3] = swap32_S (w6[3]);
-  w7[0] = swap32_S (w7[0]);
-  w7[1] = swap32_S (w7[1]);
-  w7[2] = swap32_S (w7[2]);
-  w7[3] = swap32_S (w7[3]);
-
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
@@ -188,6 +113,15 @@ __kernel void m08200_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha512_hmac_ctx_t sha512_hmac_ctx2 = sha512_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+    u32 w4[4];
+    u32 w5[4];
+    u32 w6[4];
+    u32 w7[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;
@@ -388,14 +322,14 @@ __kernel void m08200_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   u32 w2[4];
   u32 w3[4];
 
-  w0[0] = tmps[gid].out[4] >> 32;
-  w0[1] = tmps[gid].out[4] & 0xffffffff;
-  w0[2] = tmps[gid].out[5] >> 32;
-  w0[3] = tmps[gid].out[5] & 0xffffffff;
-  w1[0] = tmps[gid].out[6] >> 32;
-  w1[1] = tmps[gid].out[6] & 0xffffffff;
-  w1[2] = tmps[gid].out[7] >> 32;
-  w1[3] = tmps[gid].out[7] & 0xffffffff;
+  w0[0] = h32_from_64_S (tmps[gid].out[4]);
+  w0[1] = l32_from_64_S (tmps[gid].out[4]);
+  w0[2] = h32_from_64_S (tmps[gid].out[5]);
+  w0[3] = l32_from_64_S (tmps[gid].out[5]);
+  w1[0] = h32_from_64_S (tmps[gid].out[6]);
+  w1[1] = l32_from_64_S (tmps[gid].out[6]);
+  w1[2] = h32_from_64_S (tmps[gid].out[7]);
+  w1[3] = l32_from_64_S (tmps[gid].out[7]);
   w2[0] = 0;
   w2[1] = 0;
   w2[2] = 0;
@@ -407,7 +341,7 @@ __kernel void m08200_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha256_hmac_ctx_t ctx;
 
-  sha256_hmac_init (&ctx, w0, w1, w2, w3);
+  sha256_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
   sha256_hmac_update_global (&ctx, esalt_bufs[digests_offset].data_buf, esalt_bufs[digests_offset].data_len);
 

OpenCL/m08800.cl

@@ -61,48 +61,9 @@ __kernel void m08800_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -122,6 +83,11 @@ __kernel void m08800_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m08900.cl

@@ -217,48 +217,9 @@ __kernel void m08900_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
 
@@ -266,6 +227,11 @@ __kernel void m08900_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;
@@ -352,52 +318,18 @@ __kernel void m08900_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
+  /**
+   * 2nd pbkdf2, creates B
+   */
+
   u32 w0[4];
   u32 w1[4];
   u32 w2[4];
   u32 w3[4];
 
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
-  /**
-   * 2nd pbkdf2, creates B
-   */
-
   sha256_hmac_ctx_t ctx;
 
-  sha256_hmac_init (&ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
 
   for (u32 l = 0; l < SCRYPT_CNT4; l += 4)
   {

OpenCL/m09100.cl

@@ -597,7 +597,7 @@ __kernel void m09100_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_64 (&sha1_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];

OpenCL/m10900.cl

@@ -68,48 +68,9 @@ __kernel void m10900_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
@@ -135,6 +96,11 @@ __kernel void m10900_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m11300.cl

@@ -287,21 +287,21 @@ __kernel void m11300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
   u32 key[8];
 
-  key[0] = h32_from_64 (dgst[0]);
-  key[1] = l32_from_64 (dgst[0]);
-  key[2] = h32_from_64 (dgst[1]);
-  key[3] = l32_from_64 (dgst[1]);
-  key[4] = h32_from_64 (dgst[2]);
-  key[5] = l32_from_64 (dgst[2]);
-  key[6] = h32_from_64 (dgst[3]);
-  key[7] = l32_from_64 (dgst[3]);
+  key[0] = h32_from_64_S (dgst[0]);
+  key[1] = l32_from_64_S (dgst[0]);
+  key[2] = h32_from_64_S (dgst[1]);
+  key[3] = l32_from_64_S (dgst[1]);
+  key[4] = h32_from_64_S (dgst[2]);
+  key[5] = l32_from_64_S (dgst[2]);
+  key[6] = h32_from_64_S (dgst[3]);
+  key[7] = l32_from_64_S (dgst[3]);
 
   u32 iv[4];
 
-  iv[0] = h32_from_64 (dgst[4]);
-  iv[1] = l32_from_64 (dgst[4]);
-  iv[2] = h32_from_64 (dgst[5]);
-  iv[3] = l32_from_64 (dgst[5]);
+  iv[0] = h32_from_64_S (dgst[4]);
+  iv[1] = l32_from_64_S (dgst[4]);
+  iv[2] = h32_from_64_S (dgst[5]);
+  iv[3] = l32_from_64_S (dgst[5]);
 
   #define KEYLEN 60
 
@@ -315,10 +315,10 @@ __kernel void m11300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     u32 data[4];
 
-    data[0] = swap32 (esalt_bufs[digests_offset].cry_master_buf[(i / 4) + 0]);
-    data[1] = swap32 (esalt_bufs[digests_offset].cry_master_buf[(i / 4) + 1]);
-    data[2] = swap32 (esalt_bufs[digests_offset].cry_master_buf[(i / 4) + 2]);
-    data[3] = swap32 (esalt_bufs[digests_offset].cry_master_buf[(i / 4) + 3]);
+    data[0] = swap32_S (esalt_bufs[digests_offset].cry_master_buf[(i / 4) + 0]);
+    data[1] = swap32_S (esalt_bufs[digests_offset].cry_master_buf[(i / 4) + 1]);
+    data[2] = swap32_S (esalt_bufs[digests_offset].cry_master_buf[(i / 4) + 2]);
+    data[3] = swap32_S (esalt_bufs[digests_offset].cry_master_buf[(i / 4) + 3]);
 
     AES256_decrypt (ks, data, out, s_td0, s_td1, s_td2, s_td3, s_td4);
 

OpenCL/m12000.cl

@@ -62,48 +62,9 @@ __kernel void m12000_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -123,6 +84,11 @@ __kernel void m12000_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m12300.cl

@@ -84,84 +84,9 @@ __kernel void m12300_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-  u32 w4[4];
-  u32 w5[4];
-  u32 w6[4];
-  u32 w7[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-  w4[0] = pws[gid].i[16];
-  w4[1] = pws[gid].i[17];
-  w4[2] = pws[gid].i[18];
-  w4[3] = pws[gid].i[19];
-  w5[0] = pws[gid].i[20];
-  w5[1] = pws[gid].i[21];
-  w5[2] = pws[gid].i[22];
-  w5[3] = pws[gid].i[23];
-  w6[0] = pws[gid].i[24];
-  w6[1] = pws[gid].i[25];
-  w6[2] = pws[gid].i[26];
-  w6[3] = pws[gid].i[27];
-  w7[0] = pws[gid].i[28];
-  w7[1] = pws[gid].i[29];
-  w7[2] = pws[gid].i[30];
-  w7[3] = pws[gid].i[31];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-  w4[0] = swap32_S (w4[0]);
-  w4[1] = swap32_S (w4[1]);
-  w4[2] = swap32_S (w4[2]);
-  w4[3] = swap32_S (w4[3]);
-  w5[0] = swap32_S (w5[0]);
-  w5[1] = swap32_S (w5[1]);
-  w5[2] = swap32_S (w5[2]);
-  w5[3] = swap32_S (w5[3]);
-  w6[0] = swap32_S (w6[0]);
-  w6[1] = swap32_S (w6[1]);
-  w6[2] = swap32_S (w6[2]);
-  w6[3] = swap32_S (w6[3]);
-  w7[0] = swap32_S (w7[0]);
-  w7[1] = swap32_S (w7[1]);
-  w7[2] = swap32_S (w7[2]);
-  w7[3] = swap32_S (w7[3]);
-
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
@@ -183,6 +108,15 @@ __kernel void m12300_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha512_hmac_update_global (&sha512_hmac_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
 
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
   w0[0] = 0x41555448;
   w0[1] = 0x5f50424b;
   w0[2] = 0x4446325f;

OpenCL/m12700.cl

@@ -63,48 +63,9 @@ __kernel void m12700_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -118,6 +79,11 @@ __kernel void m12700_init (__global pw_t *pws, __global const kernel_rule_t *rul
   tmps[gid].opad[3] = sha1_hmac_ctx.opad.h[3];
   tmps[gid].opad[4] = sha1_hmac_ctx.opad.h[4];
 
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
   w0[0] = salt_bufs[salt_pos].salt_buf[0];
   w0[1] = salt_bufs[salt_pos].salt_buf[1];
   w0[2] = salt_bufs[salt_pos].salt_buf[2];

OpenCL/m12800.cl

@@ -159,7 +159,7 @@ __kernel void m12800_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_64 (&sha256_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];

OpenCL/m12900.cl

@@ -68,48 +68,9 @@ __kernel void m12900_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
@@ -129,6 +90,11 @@ __kernel void m12900_init (__global pw_t *pws, __global const kernel_rule_t *rul
   tmps[gid].opad[6] = sha256_hmac_ctx.opad.h[6];
   tmps[gid].opad[7] = sha256_hmac_ctx.opad.h[7];
 
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
   w0[0] = salt_bufs[salt_pos].salt_buf[0];
   w0[1] = salt_bufs[salt_pos].salt_buf[1];
   w0[2] = salt_bufs[salt_pos].salt_buf[2];
@@ -335,7 +301,7 @@ __kernel void m12900_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha256_hmac_ctx_t ctx;
 
-  sha256_hmac_init (&ctx, w0, w1, w2, w3);
+  sha256_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
   w0[0] = salt_bufs[salt_pos].salt_buf[ 4];
   w0[1] = salt_bufs[salt_pos].salt_buf[ 5];

OpenCL/m13000.cl

@@ -68,48 +68,9 @@ __kernel void m13000_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
@@ -135,6 +96,11 @@ __kernel void m13000_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m13600.cl

@@ -62,48 +62,9 @@ __kernel void m13600_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -117,6 +78,11 @@ __kernel void m13600_init (__global pw_t *pws, __global const kernel_rule_t *rul
   tmps[gid].opad[3] = sha1_hmac_ctx.opad.h[3];
   tmps[gid].opad[4] = sha1_hmac_ctx.opad.h[4];
 
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
   w0[0] = swap32_S (esalt_bufs[digests_offset].salt_buf[0]);
   w0[1] = swap32_S (esalt_bufs[digests_offset].salt_buf[1]);
   w0[2] = swap32_S (esalt_bufs[digests_offset].salt_buf[2]);
@@ -369,7 +335,7 @@ __kernel void m13600_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha1_hmac_ctx_t ctx;
 
-  sha1_hmac_init (&ctx, w0, w1, w2, w3);
+  sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
   sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].data_buf, esalt_bufs[digests_offset].data_len);
 

OpenCL/m13751.cl

@@ -155,7 +155,7 @@ __kernel void m13751_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_64 (&sha256_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];

OpenCL/m13752.cl

@@ -155,7 +155,7 @@ __kernel void m13752_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_64 (&sha256_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];

OpenCL/m13753.cl

@@ -155,7 +155,7 @@ __kernel void m13753_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_64 (&sha256_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];

OpenCL/m14611.cl

@@ -71,48 +71,9 @@ __kernel void m14611_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = sha1_hmac_ctx.ipad.h[1];
@@ -134,6 +95,11 @@ __kernel void m14611_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14612.cl

@@ -71,48 +71,9 @@ __kernel void m14612_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = sha1_hmac_ctx.ipad.h[1];
@@ -134,6 +95,11 @@ __kernel void m14612_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14613.cl

@@ -71,48 +71,9 @@ __kernel void m14613_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = sha1_hmac_ctx.ipad.h[1];
@@ -134,6 +95,11 @@ __kernel void m14613_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14621.cl

@@ -77,48 +77,9 @@ __kernel void m14621_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = sha256_hmac_ctx.ipad.h[1];
@@ -146,6 +107,11 @@ __kernel void m14621_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14622.cl

@@ -78,48 +78,9 @@ __kernel void m14622_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = sha256_hmac_ctx.ipad.h[1];
@@ -147,6 +108,11 @@ __kernel void m14622_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14623.cl

@@ -77,48 +77,9 @@ __kernel void m14623_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = sha256_hmac_ctx.ipad.h[1];
@@ -146,6 +107,11 @@ __kernel void m14623_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14631.cl

@@ -93,84 +93,9 @@ __kernel void m14631_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-  u32 w4[4];
-  u32 w5[4];
-  u32 w6[4];
-  u32 w7[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-  w4[0] = pws[gid].i[16];
-  w4[1] = pws[gid].i[17];
-  w4[2] = pws[gid].i[18];
-  w4[3] = pws[gid].i[19];
-  w5[0] = pws[gid].i[20];
-  w5[1] = pws[gid].i[21];
-  w5[2] = pws[gid].i[22];
-  w5[3] = pws[gid].i[23];
-  w6[0] = pws[gid].i[24];
-  w6[1] = pws[gid].i[25];
-  w6[2] = pws[gid].i[26];
-  w6[3] = pws[gid].i[27];
-  w7[0] = pws[gid].i[28];
-  w7[1] = pws[gid].i[29];
-  w7[2] = pws[gid].i[30];
-  w7[3] = pws[gid].i[31];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-  w4[0] = swap32_S (w4[0]);
-  w4[1] = swap32_S (w4[1]);
-  w4[2] = swap32_S (w4[2]);
-  w4[3] = swap32_S (w4[3]);
-  w5[0] = swap32_S (w5[0]);
-  w5[1] = swap32_S (w5[1]);
-  w5[2] = swap32_S (w5[2]);
-  w5[3] = swap32_S (w5[3]);
-  w6[0] = swap32_S (w6[0]);
-  w6[1] = swap32_S (w6[1]);
-  w6[2] = swap32_S (w6[2]);
-  w6[3] = swap32_S (w6[3]);
-  w7[0] = swap32_S (w7[0]);
-  w7[1] = swap32_S (w7[1]);
-  w7[2] = swap32_S (w7[2]);
-  w7[3] = swap32_S (w7[3]);
-
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad64[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad64[1] = sha512_hmac_ctx.ipad.h[1];
@@ -198,6 +123,15 @@ __kernel void m14631_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha512_hmac_ctx_t sha512_hmac_ctx2 = sha512_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+    u32 w4[4];
+    u32 w5[4];
+    u32 w6[4];
+    u32 w7[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14632.cl

@@ -93,84 +93,9 @@ __kernel void m14632_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-  u32 w4[4];
-  u32 w5[4];
-  u32 w6[4];
-  u32 w7[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-  w4[0] = pws[gid].i[16];
-  w4[1] = pws[gid].i[17];
-  w4[2] = pws[gid].i[18];
-  w4[3] = pws[gid].i[19];
-  w5[0] = pws[gid].i[20];
-  w5[1] = pws[gid].i[21];
-  w5[2] = pws[gid].i[22];
-  w5[3] = pws[gid].i[23];
-  w6[0] = pws[gid].i[24];
-  w6[1] = pws[gid].i[25];
-  w6[2] = pws[gid].i[26];
-  w6[3] = pws[gid].i[27];
-  w7[0] = pws[gid].i[28];
-  w7[1] = pws[gid].i[29];
-  w7[2] = pws[gid].i[30];
-  w7[3] = pws[gid].i[31];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-  w4[0] = swap32_S (w4[0]);
-  w4[1] = swap32_S (w4[1]);
-  w4[2] = swap32_S (w4[2]);
-  w4[3] = swap32_S (w4[3]);
-  w5[0] = swap32_S (w5[0]);
-  w5[1] = swap32_S (w5[1]);
-  w5[2] = swap32_S (w5[2]);
-  w5[3] = swap32_S (w5[3]);
-  w6[0] = swap32_S (w6[0]);
-  w6[1] = swap32_S (w6[1]);
-  w6[2] = swap32_S (w6[2]);
-  w6[3] = swap32_S (w6[3]);
-  w7[0] = swap32_S (w7[0]);
-  w7[1] = swap32_S (w7[1]);
-  w7[2] = swap32_S (w7[2]);
-  w7[3] = swap32_S (w7[3]);
-
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad64[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad64[1] = sha512_hmac_ctx.ipad.h[1];
@@ -198,6 +123,15 @@ __kernel void m14632_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha512_hmac_ctx_t sha512_hmac_ctx2 = sha512_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+    u32 w4[4];
+    u32 w5[4];
+    u32 w6[4];
+    u32 w7[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14633.cl

@@ -93,84 +93,9 @@ __kernel void m14633_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-  u32 w4[4];
-  u32 w5[4];
-  u32 w6[4];
-  u32 w7[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-  w4[0] = pws[gid].i[16];
-  w4[1] = pws[gid].i[17];
-  w4[2] = pws[gid].i[18];
-  w4[3] = pws[gid].i[19];
-  w5[0] = pws[gid].i[20];
-  w5[1] = pws[gid].i[21];
-  w5[2] = pws[gid].i[22];
-  w5[3] = pws[gid].i[23];
-  w6[0] = pws[gid].i[24];
-  w6[1] = pws[gid].i[25];
-  w6[2] = pws[gid].i[26];
-  w6[3] = pws[gid].i[27];
-  w7[0] = pws[gid].i[28];
-  w7[1] = pws[gid].i[29];
-  w7[2] = pws[gid].i[30];
-  w7[3] = pws[gid].i[31];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-  w4[0] = swap32_S (w4[0]);
-  w4[1] = swap32_S (w4[1]);
-  w4[2] = swap32_S (w4[2]);
-  w4[3] = swap32_S (w4[3]);
-  w5[0] = swap32_S (w5[0]);
-  w5[1] = swap32_S (w5[1]);
-  w5[2] = swap32_S (w5[2]);
-  w5[3] = swap32_S (w5[3]);
-  w6[0] = swap32_S (w6[0]);
-  w6[1] = swap32_S (w6[1]);
-  w6[2] = swap32_S (w6[2]);
-  w6[3] = swap32_S (w6[3]);
-  w7[0] = swap32_S (w7[0]);
-  w7[1] = swap32_S (w7[1]);
-  w7[2] = swap32_S (w7[2]);
-  w7[3] = swap32_S (w7[3]);
-
   sha512_hmac_ctx_t sha512_hmac_ctx;
 
-  sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+  sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad64[0] = sha512_hmac_ctx.ipad.h[0];
   tmps[gid].ipad64[1] = sha512_hmac_ctx.ipad.h[1];
@@ -198,6 +123,15 @@ __kernel void m14633_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha512_hmac_ctx_t sha512_hmac_ctx2 = sha512_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+    u32 w4[4];
+    u32 w5[4];
+    u32 w6[4];
+    u32 w7[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14641.cl

@@ -71,31 +71,9 @@ __kernel void m14641_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
   ripemd160_hmac_ctx_t ripemd160_hmac_ctx;
 
-  ripemd160_hmac_init (&ripemd160_hmac_ctx, w0, w1, w2, w3);
+  ripemd160_hmac_init_global (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = ripemd160_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = ripemd160_hmac_ctx.ipad.h[1];
@@ -117,6 +95,11 @@ __kernel void m14641_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     ripemd160_hmac_ctx_t ripemd160_hmac_ctx2 = ripemd160_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j << 24;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14642.cl

@@ -71,31 +71,9 @@ __kernel void m14642_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
   ripemd160_hmac_ctx_t ripemd160_hmac_ctx;
 
-  ripemd160_hmac_init (&ripemd160_hmac_ctx, w0, w1, w2, w3);
+  ripemd160_hmac_init_global (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = ripemd160_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = ripemd160_hmac_ctx.ipad.h[1];
@@ -117,6 +95,11 @@ __kernel void m14642_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     ripemd160_hmac_ctx_t ripemd160_hmac_ctx2 = ripemd160_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j << 24;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14643.cl

@@ -71,31 +71,9 @@ __kernel void m14643_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
   ripemd160_hmac_ctx_t ripemd160_hmac_ctx;
 
-  ripemd160_hmac_init (&ripemd160_hmac_ctx, w0, w1, w2, w3);
+  ripemd160_hmac_init_global (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad32[0] = ripemd160_hmac_ctx.ipad.h[0];
   tmps[gid].ipad32[1] = ripemd160_hmac_ctx.ipad.h[1];
@@ -117,6 +95,11 @@ __kernel void m14643_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     ripemd160_hmac_ctx_t ripemd160_hmac_ctx2 = ripemd160_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j << 24;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14700.cl

@@ -63,48 +63,9 @@ __kernel void m14700_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -124,6 +85,11 @@ __kernel void m14700_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m14800.cl

@@ -103,48 +103,9 @@ __kernel void m14800_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
@@ -164,6 +125,11 @@ __kernel void m14800_init (__global pw_t *pws, __global const kernel_rule_t *rul
   tmps[gid].opad[6] = sha256_hmac_ctx.opad.h[6];
   tmps[gid].opad[7] = sha256_hmac_ctx.opad.h[7];
 
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
   w0[0] = esalt_bufs[digests_offset].dpsl[0];
   w0[1] = esalt_bufs[digests_offset].dpsl[1];
   w0[2] = esalt_bufs[digests_offset].dpsl[2];
@@ -368,7 +334,7 @@ __kernel void m14800_init2 (__global pw_t *pws, __global const kernel_rule_t *ru
 
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_64 (&sha1_hmac_ctx, w0, w1, w2, w3);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];

OpenCL/m15100.cl

@@ -62,48 +62,9 @@ __kernel void m15100_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha1_hmac_ctx_t sha1_hmac_ctx;
 
-  sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+  sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];

OpenCL/m15300.cl

@@ -722,7 +722,7 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   sha1_hmac_ctx_t ctx;
 
-  sha1_hmac_init (&ctx, w0, w1, w2, w3);
+  sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
   sha1_hmac_update_global (&ctx, esalt_bufs[digests_offset].SID, esalt_bufs[digests_offset].SID_len);
 
@@ -767,7 +767,7 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
     sha1_hmac_ctx_t sha1_hmac_ctx;
 
-    sha1_hmac_init (&sha1_hmac_ctx, w0, w1, w2, w3);
+    sha1_hmac_init_64 (&sha1_hmac_ctx, w0, w1, w2, w3);
 
     tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
     tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
@@ -886,7 +886,7 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
     sha512_hmac_ctx_t sha512_hmac_ctx;
 
-    sha512_hmac_init (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
+    sha512_hmac_init_128 (&sha512_hmac_ctx, w0, w1, w2, w3, w5, w5, w6, w7);
 
     tmps[gid].ipad64[0] = sha512_hmac_ctx.ipad.h[0];
     tmps[gid].ipad64[1] = sha512_hmac_ctx.ipad.h[1];
@@ -1438,7 +1438,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     sha1_hmac_ctx_t ctx;
 
-    sha1_hmac_init (&ctx, w0, w1, w2, w3);
+    sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
     w0[0] = hmacSalt[0];
     w0[1] = hmacSalt[1];
@@ -1478,7 +1478,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     w3[2] = 0;
     w3[3] = 0;
 
-    sha1_hmac_init (&ctx, w0, w1, w2, w3);
+    sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
     w0[0] = swap32_S (lastKey[ 0]);
     w0[1] = swap32_S (lastKey[ 1]);
@@ -1637,7 +1637,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     sha512_hmac_ctx_t ctx;
 
-    sha512_hmac_init (&ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+    sha512_hmac_init_128 (&ctx, w0, w1, w2, w3, w4, w5, w6, w7);
 
     w0[0] = hmacSalt[0];
     w0[1] = hmacSalt[1];
@@ -1709,7 +1709,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     w7[2] = 0;
     w7[3] = 0;
 
-    sha512_hmac_init (&ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+    sha512_hmac_init_128 (&ctx, w0, w1, w2, w3, w4, w5, w6, w7);
 
     w0[0] = lastKey[ 0];
     w0[1] = lastKey[ 1];

OpenCL/m15600.cl

@@ -200,48 +200,9 @@ __kernel void m15600_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
   tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
@@ -267,6 +228,11 @@ __kernel void m15600_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;

OpenCL/m15700.cl

@@ -348,48 +348,9 @@ __kernel void m15700_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
   sha256_hmac_ctx_t sha256_hmac_ctx;
 
-  sha256_hmac_init (&sha256_hmac_ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
 
   sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
 
@@ -397,6 +358,11 @@ __kernel void m15700_init (__global pw_t *pws, __global const kernel_rule_t *rul
   {
     sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx;
 
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+
     w0[0] = j;
     w0[1] = 0;
     w0[2] = 0;
@@ -483,52 +449,18 @@ __kernel void m15700_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
   if (gid >= gid_max) return;
 
+  /**
+   * 2nd pbkdf2, creates B
+   */
+
   u32 w0[4];
   u32 w1[4];
   u32 w2[4];
   u32 w3[4];
 
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
-  w0[0] = swap32_S (w0[0]);
-  w0[1] = swap32_S (w0[1]);
-  w0[2] = swap32_S (w0[2]);
-  w0[3] = swap32_S (w0[3]);
-  w1[0] = swap32_S (w1[0]);
-  w1[1] = swap32_S (w1[1]);
-  w1[2] = swap32_S (w1[2]);
-  w1[3] = swap32_S (w1[3]);
-  w2[0] = swap32_S (w2[0]);
-  w2[1] = swap32_S (w2[1]);
-  w2[2] = swap32_S (w2[2]);
-  w2[3] = swap32_S (w2[3]);
-  w3[0] = swap32_S (w3[0]);
-  w3[1] = swap32_S (w3[1]);
-  w3[2] = swap32_S (w3[2]);
-  w3[3] = swap32_S (w3[3]);
-
-  /**
-   * 2nd pbkdf2, creates B
-   */
-
   sha256_hmac_ctx_t ctx;
 
-  sha256_hmac_init (&ctx, w0, w1, w2, w3);
+  sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
 
   for (u32 l = 0; l < SCRYPT_CNT4; l += 4)
   {

src/interface.c

@@ -24621,7 +24621,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
                 break;
   }
 
-  // pw_max
+  // pw_max : some algo suffer from support for long passwords, the user need to add -L to enable support for them
 
   if (user_options->length_limit_disable == true)
   {
@@ -24640,19 +24640,19 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
     {
       case   125: hashconfig->pw_max = 32;
                   break;
-      case   500: hashconfig->pw_max = 15; // pure kernel available
+      case   500: hashconfig->pw_max = 15; // -L available
                   break;
-      case  1600: hashconfig->pw_max = 15; // pure kernel available
+      case  1600: hashconfig->pw_max = 15; // -L available
                   break;
-      case  1800: hashconfig->pw_max = 16; // pure kernel available
+      case  1800: hashconfig->pw_max = 16; // -L available
                   break;
-      case  5800: hashconfig->pw_max = 16; // pure kernel available
+      case  5800: hashconfig->pw_max = 16; // -L available
                   break;
-      case  6300: hashconfig->pw_max = 15; // pure kernel available
+      case  6300: hashconfig->pw_max = 15; // -L available
                   break;
       case  7000: hashconfig->pw_max = 19;
                   break;
-      case  7400: hashconfig->pw_max = 15; // pure kernel available
+      case  7400: hashconfig->pw_max = 15; // -L available
                   break;
       case  7700: hashconfig->pw_max = 8;
                   break;
@@ -24671,214 +24671,115 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
     }
   }
 
+  // pw_max : all modes listed in the follow 2 switch cases are at maximum possible
   // pw_max : some algo are converted to long password support but without dropping performance
 
   switch (hashconfig->hash_mode)
   {
-    case  2100: hashconfig->pw_max = PW_MAX;
-                break;
-    case  5200: hashconfig->pw_max = PW_MAX;
-                break;
-    case  7900: hashconfig->pw_max = PW_MAX;
-                break;
-    case  9400: hashconfig->pw_max = PW_MAX;
-                break;
-    case  9500: hashconfig->pw_max = PW_MAX;
-                break;
-    case  9600: hashconfig->pw_max = PW_MAX;
-                break;
-    case 10300: hashconfig->pw_max = PW_MAX;
-                break;
-    case 11300: hashconfig->pw_max = PW_MAX;
-                break;
-    case 11600: hashconfig->pw_max = PW_MAX;
-                break;
-    case 11900: hashconfig->pw_max = PW_MAX;
-                break;
-    case 12200: hashconfig->pw_max = PW_MAX;
-                break;
-    case 12400: hashconfig->pw_max = PW_MAX;
-                break;
-    case 12800: hashconfig->pw_max = PW_MAX;
-                break;
-    case 13200: hashconfig->pw_max = PW_MAX;
-                break;
-    case 13400: hashconfig->pw_max = PW_MAX;
-                break;
+    case  2100: hashconfig->pw_max = PW_MAX; break;
+    case  5200: hashconfig->pw_max = PW_MAX; break;
+    case  6400: hashconfig->pw_max = PW_MAX; break;
+    case  6500: hashconfig->pw_max = PW_MAX; break;
+    case  6600: hashconfig->pw_max = PW_MAX; break;
+    case  6700: hashconfig->pw_max = PW_MAX; break;
+    case  6800: hashconfig->pw_max = PW_MAX; break;
+    case  7100: hashconfig->pw_max = PW_MAX; break;
+    case  7200: hashconfig->pw_max = PW_MAX; break;
+    case  7900: hashconfig->pw_max = PW_MAX; break;
+    case  8200: hashconfig->pw_max = PW_MAX; break;
+    case  8800: hashconfig->pw_max = PW_MAX; break;
+    case  8900: hashconfig->pw_max = PW_MAX; break;
+    case  9200: hashconfig->pw_max = PW_MAX; break;
+    case  9300: hashconfig->pw_max = PW_MAX; break;
+    case  9400: hashconfig->pw_max = PW_MAX; break;
+    case  9500: hashconfig->pw_max = PW_MAX; break;
+    case  9600: hashconfig->pw_max = PW_MAX; break;
+    case 10000: hashconfig->pw_max = PW_MAX; break;
+    case 10300: hashconfig->pw_max = PW_MAX; break;
+    case 10900: hashconfig->pw_max = PW_MAX; break;
+    case 11300: hashconfig->pw_max = PW_MAX; break;
+    case 11600: hashconfig->pw_max = PW_MAX; break;
+    case 11900: hashconfig->pw_max = PW_MAX; break;
+    case 12000: hashconfig->pw_max = PW_MAX; break;
+    case 12001: hashconfig->pw_max = PW_MAX; break;
+    case 12200: hashconfig->pw_max = PW_MAX; break;
+    case 12300: hashconfig->pw_max = PW_MAX; break;
+    case 12400: hashconfig->pw_max = PW_MAX; break;
+    case 12700: hashconfig->pw_max = PW_MAX; break;
+    case 12800: hashconfig->pw_max = PW_MAX; break;
+    case 12900: hashconfig->pw_max = PW_MAX; break;
+    case 13000: hashconfig->pw_max = PW_MAX; break;
+    case 13200: hashconfig->pw_max = PW_MAX; break;
+    case 13400: hashconfig->pw_max = PW_MAX; break;
+    case 13600: hashconfig->pw_max = PW_MAX; break;
+    case 14611: hashconfig->pw_max = PW_MAX; break;
+    case 14612: hashconfig->pw_max = PW_MAX; break;
+    case 14613: hashconfig->pw_max = PW_MAX; break;
+    case 14621: hashconfig->pw_max = PW_MAX; break;
+    case 14622: hashconfig->pw_max = PW_MAX; break;
+    case 14623: hashconfig->pw_max = PW_MAX; break;
+    case 14631: hashconfig->pw_max = PW_MAX; break;
+    case 14632: hashconfig->pw_max = PW_MAX; break;
+    case 14633: hashconfig->pw_max = PW_MAX; break;
+    case 14641: hashconfig->pw_max = PW_MAX; break;
+    case 14642: hashconfig->pw_max = PW_MAX; break;
+    case 14643: hashconfig->pw_max = PW_MAX; break;
+    case 14700: hashconfig->pw_max = PW_MAX; break;
+    case 14800: hashconfig->pw_max = PW_MAX; break;
+    case 15100: hashconfig->pw_max = PW_MAX; break;
+    case 15300: hashconfig->pw_max = PW_MAX; break;
+    case 15600: hashconfig->pw_max = PW_MAX; break;
+    case 15700: hashconfig->pw_max = PW_MAX; break;
   }
 
   // pw_max : algo specific hard max length
 
   switch (hashconfig->hash_mode)
   {
-    case  1500: hashconfig->pw_max = 8;   // DES max
-                break;
-    case  2100: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case  2500: hashconfig->pw_max = 63;  // WPA max
-                break;
-    case  2501: hashconfig->pw_max = 64;  // WPA PMK max
-                break;
-    case  3000: hashconfig->pw_max = 7;   // LM half max
-                break;
-    case  3200: hashconfig->pw_max = 72;  // bcrypt max
-                break;
-    case  6211: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case  6212: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case  6213: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case  6221: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case  6222: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case  6223: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case  6231: hashconfig->pw_max = 64;  // PBKDF2-HMAC-Whirlpool max
-                break;
-    case  6232: hashconfig->pw_max = 64;  // PBKDF2-HMAC-Whirlpool max
-                break;
-    case  6233: hashconfig->pw_max = 64;  // PBKDF2-HMAC-Whirlpool max
-                break;
-    case  6241: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case  6242: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case  6243: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case  6400: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case  6500: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case  6600: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case  6700: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case  6800: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case  7100: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case  7200: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case  8200: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case  8500: hashconfig->pw_max = 8;   // DES max
-                break;
-    case  8800: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case  8900: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case  9100: hashconfig->pw_max = 64;  // Lotus Notes/Domino 8 max
-                break;
-    case  9200: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case  9300: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case  9710: hashconfig->pw_max = 5;   // RC4-40 max
-                break;
-    case  9810: hashconfig->pw_max = 5;   // RC4-40 max
-                break;
-    case 10000: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 10410: hashconfig->pw_max = 5;   // RC4-40 max
-                break;
-    case 10500: hashconfig->pw_max = 32;  // PDF 1.4 - 1.6 (Acrobat 5 - 8) max
-                break;
-    case 10900: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 12000: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 12001: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 12300: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case 12700: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 12900: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 13000: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 13600: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 13711: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 13712: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 13713: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 13721: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case 13722: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case 13723: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case 13731: hashconfig->pw_max = 64;  // PBKDF2-HMAC-Whirlpool max
-                break;
-    case 13732: hashconfig->pw_max = 64;  // PBKDF2-HMAC-Whirlpool max
-                break;
-    case 13733: hashconfig->pw_max = 64;  // PBKDF2-HMAC-Whirlpool max
-                break;
-    case 13741: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 13742: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 13743: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 13751: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 13752: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 13753: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 13761: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 13762: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 13763: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 14000: hashconfig->pw_max = 8;   // DES max
-                break;
-    case 14100: hashconfig->pw_max = 24;  // 3DES max
-                break;
-    case 14611: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 14612: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 14613: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 14621: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 14622: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 14623: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 14631: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case 14632: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case 14633: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case 14641: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 14642: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 14643: hashconfig->pw_max = 64;  // PBKDF2-HMAC-RipeMD160 max
-                break;
-    case 14700: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 14800: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 14900: hashconfig->pw_max = 10;  // Skip32 max
-                break;
-    case 15100: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA1 max
-                break;
-    case 15300: hashconfig->pw_max = 128; // PBKDF2-HMAC-SHA512 max
-                break;
-    case 15600: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
-    case 15700: hashconfig->pw_max = 64;  // PBKDF2-HMAC-SHA256 max
-                break;
+    case  1500: hashconfig->pw_max = 8;  break; // DES max
+    case  2500: hashconfig->pw_max = 63; break; // WPA/WPA2 limits itself to 63
+    case  2501: hashconfig->pw_max = 64; break; // WPA/WPA2 PMK max
+    case  3000: hashconfig->pw_max = 7;  break; // LM max
+    case  3200: hashconfig->pw_max = 72; break; // blowfish max
+    case  6211: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6212: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6213: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6221: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6222: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6223: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6231: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6232: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6233: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6241: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6242: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  6243: hashconfig->pw_max = 64; break; // TC limits itself to 64
+    case  8500: hashconfig->pw_max = 8;  break; // DES max
+    case  9100: hashconfig->pw_max = 64; break; // Lotus Notes/Domino limits itself to 8
+    case  9710: hashconfig->pw_max = 5;  break; // RC4-40 max
+    case  9810: hashconfig->pw_max = 5;  break; // RC4-40 max
+    case 10410: hashconfig->pw_max = 5;  break; // RC4-40 max
+    case 10500: hashconfig->pw_max = 32; break; // PDF 1.4 - 1.6 (Acrobat 5 - 8) max
+    case 13711: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13712: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13713: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13721: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13722: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13723: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13731: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13732: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13733: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13741: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13742: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13743: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13751: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13752: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13753: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13761: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13762: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 13763: hashconfig->pw_max = 64; break; // VC limits itself to 64
+    case 14000: hashconfig->pw_max = 8;  break; // DES max
+    case 14100: hashconfig->pw_max = 24; break; // 3DES max
+    case 14900: hashconfig->pw_max = 10; break; // Skip32 max
   }
 
   return 0;