From 8a6e3a5275efedbdb51c78e827646c8fadae8734 Mon Sep 17 00:00:00 2001 From: jsteube Date: Mon, 10 Jul 2017 11:15:15 +0200 Subject: [PATCH] Add support in HMAC for passwords larger than block size of the underlaying hash --- OpenCL/inc_hash_md4.cl | 173 ++++++++++++++++++- OpenCL/inc_hash_md5.cl | 57 +++++++ OpenCL/inc_hash_ripemd160.cl | 173 ++++++++++++++++++- OpenCL/inc_hash_sha1.cl | 173 ++++++++++++++++++- OpenCL/inc_hash_sha256.cl | 173 ++++++++++++++++++- OpenCL/inc_hash_sha384.cl | 289 +++++++++++++++++++++++++++++++- OpenCL/inc_hash_sha512.cl | 281 ++++++++++++++++++++++++++++++- OpenCL/inc_hash_whirlpool.cl | 173 ++++++++++++++++++- OpenCL/inc_types.cl | 23 +-- OpenCL/m02100.cl | 2 +- OpenCL/m02500.cl | 54 ++---- OpenCL/m02501.cl | 8 +- OpenCL/m06211.cl | 2 +- OpenCL/m06212.cl | 2 +- OpenCL/m06213.cl | 2 +- OpenCL/m06221.cl | 34 ++-- OpenCL/m06222.cl | 34 ++-- OpenCL/m06223.cl | 34 ++-- OpenCL/m06231.cl | 2 +- OpenCL/m06232.cl | 2 +- OpenCL/m06233.cl | 2 +- OpenCL/m06400.cl | 46 +----- OpenCL/m06500.cl | 86 ++-------- OpenCL/m06600.cl | 46 +----- OpenCL/m06700.cl | 46 +----- OpenCL/m06800.cl | 54 ++---- OpenCL/m07100.cl | 86 ++-------- OpenCL/m08200.cl | 104 +++--------- OpenCL/m08800.cl | 46 +----- OpenCL/m08900.cl | 90 ++-------- OpenCL/m09100.cl | 2 +- OpenCL/m10900.cl | 46 +----- OpenCL/m11300.cl | 32 ++-- OpenCL/m12000.cl | 46 +----- OpenCL/m12300.cl | 86 ++-------- OpenCL/m12700.cl | 46 +----- OpenCL/m12800.cl | 2 +- OpenCL/m12900.cl | 48 +----- OpenCL/m13000.cl | 46 +----- OpenCL/m13600.cl | 48 +----- OpenCL/m13751.cl | 2 +- OpenCL/m13752.cl | 2 +- OpenCL/m13753.cl | 2 +- OpenCL/m14611.cl | 46 +----- OpenCL/m14612.cl | 46 +----- OpenCL/m14613.cl | 46 +----- OpenCL/m14621.cl | 46 +----- OpenCL/m14622.cl | 46 +----- OpenCL/m14623.cl | 46 +----- OpenCL/m14631.cl | 86 ++-------- OpenCL/m14632.cl | 86 ++-------- OpenCL/m14633.cl | 86 ++-------- OpenCL/m14641.cl | 29 +--- OpenCL/m14642.cl | 29 +--- OpenCL/m14643.cl | 29 +--- OpenCL/m14700.cl | 46 +----- OpenCL/m14800.cl | 48 +----- OpenCL/m15100.cl | 41 +---- OpenCL/m15300.cl | 14 +- OpenCL/m15600.cl | 46 +----- OpenCL/m15700.cl | 90 ++-------- src/interface.c | 311 ++++++++++++----------------------- 62 files changed, 1936 insertions(+), 1986 deletions(-) diff --git a/OpenCL/inc_hash_md4.cl b/OpenCL/inc_hash_md4.cl index 8f40860b7..5cee3bb75 100644 --- a/OpenCL/inc_hash_md4.cl +++ b/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); diff --git a/OpenCL/inc_hash_md5.cl b/OpenCL/inc_hash_md5.cl index 36a6fe6e4..dc432e364 100644 --- a/OpenCL/inc_hash_md5.cl +++ b/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); diff --git a/OpenCL/inc_hash_ripemd160.cl b/OpenCL/inc_hash_ripemd160.cl index a121fb717..b3f0fce6b 100644 --- a/OpenCL/inc_hash_ripemd160.cl +++ b/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); diff --git a/OpenCL/inc_hash_sha1.cl b/OpenCL/inc_hash_sha1.cl index a4ca265c4..b1944eea0 100644 --- a/OpenCL/inc_hash_sha1.cl +++ b/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); diff --git a/OpenCL/inc_hash_sha256.cl b/OpenCL/inc_hash_sha256.cl index f4a0e0ce2..54aac336a 100644 --- a/OpenCL/inc_hash_sha256.cl +++ b/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); diff --git a/OpenCL/inc_hash_sha384.cl b/OpenCL/inc_hash_sha384.cl index 1b2f88acd..5f7e0e64c 100644 --- a/OpenCL/inc_hash_sha384.cl +++ b/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; diff --git a/OpenCL/inc_hash_sha512.cl b/OpenCL/inc_hash_sha512.cl index 610da0f34..490933503 100644 --- a/OpenCL/inc_hash_sha512.cl +++ b/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); diff --git a/OpenCL/inc_hash_whirlpool.cl b/OpenCL/inc_hash_whirlpool.cl index 5bcef30aa..43cd1efa2 100644 --- a/OpenCL/inc_hash_whirlpool.cl +++ b/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); diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index 39920dce5..56f1e74a0 100644 --- a/OpenCL/inc_types.cl +++ b/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) diff --git a/OpenCL/m02100.cl b/OpenCL/m02100.cl index 4eef455eb..1d906c978 100644 --- a/OpenCL/m02100.cl +++ b/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]; diff --git a/OpenCL/m02500.cl b/OpenCL/m02500.cl index b979aa871..5c1794bc8 100644 --- a/OpenCL/m02500.cl +++ b/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); diff --git a/OpenCL/m02501.cl b/OpenCL/m02501.cl index eead090fa..d506bc389 100644 --- a/OpenCL/m02501.cl +++ b/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); diff --git a/OpenCL/m06211.cl b/OpenCL/m06211.cl index 244f93e7b..5cbdf3fd7 100644 --- a/OpenCL/m06211.cl +++ b/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]; diff --git a/OpenCL/m06212.cl b/OpenCL/m06212.cl index 9ad76473e..0a54123e4 100644 --- a/OpenCL/m06212.cl +++ b/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]; diff --git a/OpenCL/m06213.cl b/OpenCL/m06213.cl index 1a0990ce4..33c376a2c 100644 --- a/OpenCL/m06213.cl +++ b/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]; diff --git a/OpenCL/m06221.cl b/OpenCL/m06221.cl index f382e9f4c..9bb3ea3b6 100644 --- a/OpenCL/m06221.cl +++ b/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) { diff --git a/OpenCL/m06222.cl b/OpenCL/m06222.cl index 542ba9f76..4666ff978 100644 --- a/OpenCL/m06222.cl +++ b/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) { diff --git a/OpenCL/m06223.cl b/OpenCL/m06223.cl index 9267d691f..9ffd2c2ed 100644 --- a/OpenCL/m06223.cl +++ b/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) { diff --git a/OpenCL/m06231.cl b/OpenCL/m06231.cl index ec580cc73..311ba2679 100644 --- a/OpenCL/m06231.cl +++ b/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]; diff --git a/OpenCL/m06232.cl b/OpenCL/m06232.cl index 8c2bfb2e3..48d10b8f0 100644 --- a/OpenCL/m06232.cl +++ b/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]; diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl index ee3b54a91..cafd8137d 100644 --- a/OpenCL/m06233.cl +++ b/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]; diff --git a/OpenCL/m06400.cl b/OpenCL/m06400.cl index 62bec2243..8c6840b3d 100644 --- a/OpenCL/m06400.cl +++ b/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; diff --git a/OpenCL/m06500.cl b/OpenCL/m06500.cl index b01448154..97552d3e0 100644 --- a/OpenCL/m06500.cl +++ b/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; diff --git a/OpenCL/m06600.cl b/OpenCL/m06600.cl index 9642bfc71..1b699907e 100644 --- a/OpenCL/m06600.cl +++ b/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; diff --git a/OpenCL/m06700.cl b/OpenCL/m06700.cl index 8a3915a32..ef9c62e37 100644 --- a/OpenCL/m06700.cl +++ b/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; diff --git a/OpenCL/m06800.cl b/OpenCL/m06800.cl index 9ce9e738c..b9d3fe782 100644 --- a/OpenCL/m06800.cl +++ b/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); diff --git a/OpenCL/m07100.cl b/OpenCL/m07100.cl index c535f10c6..2066d3cb1 100644 --- a/OpenCL/m07100.cl +++ b/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; diff --git a/OpenCL/m08200.cl b/OpenCL/m08200.cl index 7a7a25897..48d62c2bd 100644 --- a/OpenCL/m08200.cl +++ b/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); diff --git a/OpenCL/m08800.cl b/OpenCL/m08800.cl index 226079b16..8e6584aaf 100644 --- a/OpenCL/m08800.cl +++ b/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; diff --git a/OpenCL/m08900.cl b/OpenCL/m08900.cl index 6f846b54d..402c2e0c0 100644 --- a/OpenCL/m08900.cl +++ b/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) { diff --git a/OpenCL/m09100.cl b/OpenCL/m09100.cl index 0430b545d..0ec5ac3c5 100644 --- a/OpenCL/m09100.cl +++ b/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]; diff --git a/OpenCL/m10900.cl b/OpenCL/m10900.cl index 1a467f7c2..116e358e6 100644 --- a/OpenCL/m10900.cl +++ b/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; diff --git a/OpenCL/m11300.cl b/OpenCL/m11300.cl index ef403e77e..13565c2ed 100644 --- a/OpenCL/m11300.cl +++ b/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); diff --git a/OpenCL/m12000.cl b/OpenCL/m12000.cl index 468f0a2b9..98c30a600 100644 --- a/OpenCL/m12000.cl +++ b/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; diff --git a/OpenCL/m12300.cl b/OpenCL/m12300.cl index 4049a8406..c74b6beac 100644 --- a/OpenCL/m12300.cl +++ b/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; diff --git a/OpenCL/m12700.cl b/OpenCL/m12700.cl index fc9bdb9c1..3eb4ed541 100644 --- a/OpenCL/m12700.cl +++ b/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]; diff --git a/OpenCL/m12800.cl b/OpenCL/m12800.cl index 673558d99..089bb463f 100644 --- a/OpenCL/m12800.cl +++ b/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]; diff --git a/OpenCL/m12900.cl b/OpenCL/m12900.cl index 921cb3419..18ac1c4e4 100644 --- a/OpenCL/m12900.cl +++ b/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]; diff --git a/OpenCL/m13000.cl b/OpenCL/m13000.cl index d8caf5dd0..321509bc7 100644 --- a/OpenCL/m13000.cl +++ b/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; diff --git a/OpenCL/m13600.cl b/OpenCL/m13600.cl index 1d2898fe0..ad26d6794 100644 --- a/OpenCL/m13600.cl +++ b/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); diff --git a/OpenCL/m13751.cl b/OpenCL/m13751.cl index 680c61463..7c39478fc 100644 --- a/OpenCL/m13751.cl +++ b/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]; diff --git a/OpenCL/m13752.cl b/OpenCL/m13752.cl index 81330bc39..7069eb84b 100644 --- a/OpenCL/m13752.cl +++ b/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]; diff --git a/OpenCL/m13753.cl b/OpenCL/m13753.cl index c8acde476..f3df8e293 100644 --- a/OpenCL/m13753.cl +++ b/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]; diff --git a/OpenCL/m14611.cl b/OpenCL/m14611.cl index 927aaa778..634d7bc27 100644 --- a/OpenCL/m14611.cl +++ b/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; diff --git a/OpenCL/m14612.cl b/OpenCL/m14612.cl index 8ff7bffd4..7d08ed9e9 100644 --- a/OpenCL/m14612.cl +++ b/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; diff --git a/OpenCL/m14613.cl b/OpenCL/m14613.cl index 24d83dd40..494687e27 100644 --- a/OpenCL/m14613.cl +++ b/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; diff --git a/OpenCL/m14621.cl b/OpenCL/m14621.cl index c05b2ad87..0db30ff10 100644 --- a/OpenCL/m14621.cl +++ b/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; diff --git a/OpenCL/m14622.cl b/OpenCL/m14622.cl index f2d587e52..a735e2001 100644 --- a/OpenCL/m14622.cl +++ b/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; diff --git a/OpenCL/m14623.cl b/OpenCL/m14623.cl index 53ef24446..4a8e6c65e 100644 --- a/OpenCL/m14623.cl +++ b/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; diff --git a/OpenCL/m14631.cl b/OpenCL/m14631.cl index de2ea0432..55114e9bd 100644 --- a/OpenCL/m14631.cl +++ b/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; diff --git a/OpenCL/m14632.cl b/OpenCL/m14632.cl index 92de5164f..8408e6a49 100644 --- a/OpenCL/m14632.cl +++ b/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; diff --git a/OpenCL/m14633.cl b/OpenCL/m14633.cl index 66650881f..9be0de7ea 100644 --- a/OpenCL/m14633.cl +++ b/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; diff --git a/OpenCL/m14641.cl b/OpenCL/m14641.cl index 8a215c550..e0a3c68ba 100644 --- a/OpenCL/m14641.cl +++ b/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; diff --git a/OpenCL/m14642.cl b/OpenCL/m14642.cl index c6e1bec33..95d538a19 100644 --- a/OpenCL/m14642.cl +++ b/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; diff --git a/OpenCL/m14643.cl b/OpenCL/m14643.cl index ab6a14c3c..490869f09 100644 --- a/OpenCL/m14643.cl +++ b/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; diff --git a/OpenCL/m14700.cl b/OpenCL/m14700.cl index b533a827d..b06202ff9 100644 --- a/OpenCL/m14700.cl +++ b/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; diff --git a/OpenCL/m14800.cl b/OpenCL/m14800.cl index 01ac28944..0c672531c 100644 --- a/OpenCL/m14800.cl +++ b/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]; diff --git a/OpenCL/m15100.cl b/OpenCL/m15100.cl index 0781f27b7..3dce36e93 100644 --- a/OpenCL/m15100.cl +++ b/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]; diff --git a/OpenCL/m15300.cl b/OpenCL/m15300.cl index dc7181973..d87161ced 100644 --- a/OpenCL/m15300.cl +++ b/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]; diff --git a/OpenCL/m15600.cl b/OpenCL/m15600.cl index b478e5ff2..5d65e9c83 100644 --- a/OpenCL/m15600.cl +++ b/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; diff --git a/OpenCL/m15700.cl b/OpenCL/m15700.cl index 7b6f6577c..57a33dc1b 100644 --- a/OpenCL/m15700.cl +++ b/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) { diff --git a/src/interface.c b/src/interface.c index b7058832b..bc60e1f65 100644 --- a/src/interface.c +++ b/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;