From 5efebb7b4818375e224c5013b8fcad492a88ed3b Mon Sep 17 00:00:00 2001 From: "R. Yushaev" <44146334+Naufragous@users.noreply.github.com> Date: Thu, 29 Nov 2018 12:49:03 +0100 Subject: [PATCH] Cleanup VeraCrypt related code Remove unnecessary constant variables by hardcoring values instead of looking up. Precalculate swaps that are known at compile time. Hardcode hashes_shown offset as zero for all TC / VC kernels. --- OpenCL/inc_cipher_kuznyechik.cl | 96 ++++++++++++++++----------------- OpenCL/inc_hash_streebog256.cl | 4 +- OpenCL/inc_hash_streebog512.cl | 4 +- OpenCL/m06212-pure.cl | 4 +- OpenCL/m06213-pure.cl | 6 +-- OpenCL/m06222-pure.cl | 4 +- OpenCL/m06223-pure.cl | 6 +-- OpenCL/m06232-pure.cl | 4 +- OpenCL/m06233-pure.cl | 6 +-- OpenCL/m13751-pure.cl | 6 +-- OpenCL/m13752-pure.cl | 16 +++--- OpenCL/m13753-pure.cl | 22 ++++---- OpenCL/m13771-pure.cl | 10 ++-- OpenCL/m13772-pure.cl | 14 ++--- OpenCL/m13773-pure.cl | 16 +++--- 15 files changed, 107 insertions(+), 111 deletions(-) diff --git a/OpenCL/inc_cipher_kuznyechik.cl b/OpenCL/inc_cipher_kuznyechik.cl index 83a5afc0c..7de25e101 100644 --- a/OpenCL/inc_cipher_kuznyechik.cl +++ b/OpenCL/inc_cipher_kuznyechik.cl @@ -12,7 +12,7 @@ * * */ -__constant u32a k_sbox[256] = +__constant const u32a k_sbox[256] = { 0xfc, 0xee, 0xdd, 0x11, 0xcf, 0x6e, 0x31, 0x16, 0xfb, 0xc4, 0xfa, 0xda, 0x23, 0xc5, 0x04, 0x4d, @@ -48,7 +48,7 @@ __constant u32a k_sbox[256] = 0xd1, 0x66, 0xaf, 0xc2, 0x39, 0x4b, 0x63, 0xb6 }; -__constant u32a k_sbox_inv[256] = +__constant const u32a k_sbox_inv[256] = { 0xa5, 0x2d, 0x32, 0x8f, 0x0e, 0x30, 0x38, 0xc0, 0x54, 0xe6, 0x9e, 0x39, 0x55, 0x7e, 0x52, 0x91, @@ -84,31 +84,26 @@ __constant u32a k_sbox_inv[256] = 0xd6, 0x20, 0x0a, 0x08, 0x00, 0x4c, 0xd7, 0x74 }; -__constant int k_vec[16] = -{ - 0x94, 0x20, 0x85, 0x10, 0xc2, 0xc0, 0x01, 0xfb, - 0x01, 0xc0, 0xc2, 0x10, 0x85, 0x20, 0x94, 0x01 -}; - #define extract_byte(x,n) (((x) >> (8 * (n))) & 0xff) #define k_lookup(w,sbox) \ for (int i = 0; i < 4; i++) \ - w[i] = (sbox[extract_byte(w[i],0)] << 0) \ - | (sbox[extract_byte(w[i],1)] << 8) \ - | (sbox[extract_byte(w[i],2)] << 16) \ - | (sbox[extract_byte(w[i],3)] << 24) + w[i] = sbox[extract_byte (w[i], 0)] << 0 \ + | sbox[extract_byte (w[i], 1)] << 8 \ + | sbox[extract_byte (w[i], 2)] << 16 \ + | sbox[extract_byte (w[i], 3)] << 24 -#define k_vec_xor(n) \ - for (int y = k_vec[(n)]; y > 0; y >>= 1) \ - { \ - z ^= x * (y & 1); \ - x = ((x << 1) ^ ((x >> 7) * 0xc3)) & 0xff; \ +#define k_xor(n) \ + for (int i = (n); i > 0; i /= 2) \ + { \ + z ^= x * (i % 2); \ + x = (x << 1) ^ ((x >> 7) * 0xc3); \ + x &= 0xff; \ } DECLSPEC void kuznyechik_linear (u32 *w) { - // used inside k_vec_xor macro + // used in k_xor macro u32 x; u32 z; @@ -116,22 +111,23 @@ DECLSPEC void kuznyechik_linear (u32 *w) { z = 0; - x = extract_byte (w[3], 3); k_vec_xor (15); - x = extract_byte (w[3], 2); k_vec_xor (14); - x = extract_byte (w[3], 1); k_vec_xor (13); - x = extract_byte (w[3], 0); k_vec_xor (12); - x = extract_byte (w[2], 3); k_vec_xor (11); - x = extract_byte (w[2], 2); k_vec_xor (10); - x = extract_byte (w[2], 1); k_vec_xor ( 9); - x = extract_byte (w[2], 0); k_vec_xor ( 8); - x = extract_byte (w[1], 3); k_vec_xor ( 7); - x = extract_byte (w[1], 2); k_vec_xor ( 6); - x = extract_byte (w[1], 1); k_vec_xor ( 5); - x = extract_byte (w[1], 0); k_vec_xor ( 4); - x = extract_byte (w[0], 3); k_vec_xor ( 3); - x = extract_byte (w[0], 2); k_vec_xor ( 2); - x = extract_byte (w[0], 1); k_vec_xor ( 1); - x = extract_byte (w[0], 0); k_vec_xor ( 0); + // k_xor (1) yields the same result as a simple xor + x = extract_byte (w[3], 3); z ^= x; + x = extract_byte (w[3], 2); k_xor (148); + x = extract_byte (w[3], 1); k_xor (32); + x = extract_byte (w[3], 0); k_xor (133); + x = extract_byte (w[2], 3); k_xor (16); + x = extract_byte (w[2], 2); k_xor (194); + x = extract_byte (w[2], 1); k_xor (192); + x = extract_byte (w[2], 0); z ^= x; + x = extract_byte (w[1], 3); k_xor (251); + x = extract_byte (w[1], 2); z ^= x; + x = extract_byte (w[1], 1); k_xor (192); + x = extract_byte (w[1], 0); k_xor (194); + x = extract_byte (w[0], 3); k_xor (16); + x = extract_byte (w[0], 2); k_xor (133); + x = extract_byte (w[0], 1); k_xor (32); + x = extract_byte (w[0], 0); k_xor (148); // right-shift data block, prepend calculated byte w[3] = (w[3] << 8) | (w[2] >> 24); @@ -143,7 +139,7 @@ DECLSPEC void kuznyechik_linear (u32 *w) DECLSPEC void kuznyechik_linear_inv (u32 *w) { - // used inside k_vec_xor macro + // used in k_xor macro u32 x; u32 z; @@ -157,21 +153,21 @@ DECLSPEC void kuznyechik_linear_inv (u32 *w) w[2] = (w[2] >> 8) | (w[3] << 24); w[3] = (w[3] >> 8); - x = extract_byte (w[0], 0); k_vec_xor ( 0); - x = extract_byte (w[0], 1); k_vec_xor ( 1); - x = extract_byte (w[0], 2); k_vec_xor ( 2); - x = extract_byte (w[0], 3); k_vec_xor ( 3); - x = extract_byte (w[1], 0); k_vec_xor ( 4); - x = extract_byte (w[1], 1); k_vec_xor ( 5); - x = extract_byte (w[1], 2); k_vec_xor ( 6); - x = extract_byte (w[1], 3); k_vec_xor ( 7); - x = extract_byte (w[2], 0); k_vec_xor ( 8); - x = extract_byte (w[2], 1); k_vec_xor ( 9); - x = extract_byte (w[2], 2); k_vec_xor (10); - x = extract_byte (w[2], 3); k_vec_xor (11); - x = extract_byte (w[3], 0); k_vec_xor (12); - x = extract_byte (w[3], 1); k_vec_xor (13); - x = extract_byte (w[3], 2); k_vec_xor (14); + x = extract_byte (w[0], 0); k_xor (148); + x = extract_byte (w[0], 1); k_xor (32); + x = extract_byte (w[0], 2); k_xor (133); + x = extract_byte (w[0], 3); k_xor (16); + x = extract_byte (w[1], 0); k_xor (194); + x = extract_byte (w[1], 1); k_xor (192); + x = extract_byte (w[1], 2); z ^= x; + x = extract_byte (w[1], 3); k_xor (251); + x = extract_byte (w[2], 0); z ^= x; + x = extract_byte (w[2], 1); k_xor (192); + x = extract_byte (w[2], 2); k_xor (194); + x = extract_byte (w[2], 3); k_xor (16); + x = extract_byte (w[3], 0); k_xor (133); + x = extract_byte (w[3], 1); k_xor (32); + x = extract_byte (w[3], 2); k_xor (148); //append calculated byte w[3] |= (z << 24); diff --git a/OpenCL/inc_hash_streebog256.cl b/OpenCL/inc_hash_streebog256.cl index 4eb098543..a6758023d 100644 --- a/OpenCL/inc_hash_streebog256.cl +++ b/OpenCL/inc_hash_streebog256.cl @@ -791,7 +791,7 @@ DECLSPEC void streebog256_transform (streebog256_ctx_t *ctx, const u32 *w0, cons streebog256_g (ctx->h, ctx->n, m, ctx->s_sbob_sl64); u64 counterbuf[8] = { 0 }; - counterbuf[7] = swap64_S ((u64) 0x200); + counterbuf[7] = 0x0002000000000000; streebog256_add (ctx->n, counterbuf); streebog256_add (ctx->s, m); @@ -1479,7 +1479,7 @@ DECLSPEC void streebog256_transform_vector (streebog256_ctx_vector_t *ctx, const streebog256_g_vector (ctx->h, ctx->n, m, ctx->s_sbob_sl64); u64x counterbuf[8] = { 0 }; - counterbuf[7] = swap64 ((u64x) 0x200); + counterbuf[7] = 0x0002000000000000; streebog256_add_vector (ctx->n, counterbuf); streebog256_add_vector (ctx->s, m); diff --git a/OpenCL/inc_hash_streebog512.cl b/OpenCL/inc_hash_streebog512.cl index 194413373..e62a23b78 100644 --- a/OpenCL/inc_hash_streebog512.cl +++ b/OpenCL/inc_hash_streebog512.cl @@ -791,7 +791,7 @@ DECLSPEC void streebog512_transform (streebog512_ctx_t *ctx, const u32 *w0, cons streebog512_g (ctx->h, ctx->n, m, ctx->s_sbob_sl64); u64 counterbuf[8] = { 0 }; - counterbuf[7] = swap64_S ((u64) 0x200); + counterbuf[7] = 0x0002000000000000; streebog512_add (ctx->n, counterbuf); streebog512_add (ctx->s, m); @@ -1479,7 +1479,7 @@ DECLSPEC void streebog512_transform_vector (streebog512_ctx_vector_t *ctx, const streebog512_g_vector (ctx->h, ctx->n, m, ctx->s_sbob_sl64); u64x counterbuf[8] = { 0 }; - counterbuf[7] = swap64 ((u64x) 0x200); + counterbuf[7] = 0x0002000000000000; streebog512_add_vector (ctx->n, counterbuf); streebog512_add_vector (ctx->s, m); diff --git a/OpenCL/m06212-pure.cl b/OpenCL/m06212-pure.cl index 58e0bcf4b..530b58ed5 100644 --- a/OpenCL/m06212-pure.cl +++ b/OpenCL/m06212-pure.cl @@ -441,7 +441,7 @@ __kernel void m06212_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -449,7 +449,7 @@ __kernel void m06212_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m06213-pure.cl b/OpenCL/m06213-pure.cl index efd911330..94929138c 100644 --- a/OpenCL/m06213-pure.cl +++ b/OpenCL/m06213-pure.cl @@ -441,7 +441,7 @@ __kernel void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -449,7 +449,7 @@ __kernel void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -511,7 +511,7 @@ __kernel void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m06222-pure.cl b/OpenCL/m06222-pure.cl index b77fce863..b1c04bf58 100644 --- a/OpenCL/m06222-pure.cl +++ b/OpenCL/m06222-pure.cl @@ -585,7 +585,7 @@ __kernel void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -593,7 +593,7 @@ __kernel void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m06223-pure.cl b/OpenCL/m06223-pure.cl index 011ab917a..6193775cd 100644 --- a/OpenCL/m06223-pure.cl +++ b/OpenCL/m06223-pure.cl @@ -585,7 +585,7 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -593,7 +593,7 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -655,7 +655,7 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t)) if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m06232-pure.cl b/OpenCL/m06232-pure.cl index 663ca3fe3..533cccb68 100644 --- a/OpenCL/m06232-pure.cl +++ b/OpenCL/m06232-pure.cl @@ -699,7 +699,7 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -707,7 +707,7 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m06233-pure.cl b/OpenCL/m06233-pure.cl index cd3bbfb33..7643f733b 100644 --- a/OpenCL/m06233-pure.cl +++ b/OpenCL/m06233-pure.cl @@ -699,7 +699,7 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -707,7 +707,7 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -769,7 +769,7 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m13751-pure.cl b/OpenCL/m13751-pure.cl index eb38def4e..2c659cc4f 100644 --- a/OpenCL/m13751-pure.cl +++ b/OpenCL/m13751-pure.cl @@ -411,7 +411,7 @@ __kernel void m13751_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) 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) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -419,7 +419,7 @@ __kernel void m13751_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -427,7 +427,7 @@ __kernel void m13751_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m13752-pure.cl b/OpenCL/m13752-pure.cl index 2e074f3f6..19c7f2b53 100644 --- a/OpenCL/m13752-pure.cl +++ b/OpenCL/m13752-pure.cl @@ -411,7 +411,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) 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) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -419,7 +419,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -427,7 +427,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -473,7 +473,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -481,7 +481,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -489,7 +489,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -497,7 +497,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -505,7 +505,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m13753-pure.cl b/OpenCL/m13753-pure.cl index 62056f78c..f4c959023 100644 --- a/OpenCL/m13753-pure.cl +++ b/OpenCL/m13753-pure.cl @@ -411,7 +411,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) 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) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -419,7 +419,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -427,7 +427,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -473,7 +473,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -481,7 +481,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -489,7 +489,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -497,7 +497,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -505,7 +505,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -551,7 +551,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -559,7 +559,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -567,7 +567,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m13771-pure.cl b/OpenCL/m13771-pure.cl index f52eeefeb..1a216d96a 100644 --- a/OpenCL/m13771-pure.cl +++ b/OpenCL/m13771-pure.cl @@ -31,12 +31,12 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 u64x padding[8] = { 0 }; u64x message[8]; - padding[7] = swap64 ((u64x) 0x01); + padding[7] = 0x0100000000000000; //inner HMAC: ipad + message //first transform: precalculated ipad hash - counterbuf[7] = swap64 ((u64x) 0x200); + counterbuf[7] = 0x0002000000000000; //second transform: message = previous HMAC digest message[7] = hl32_to_64 (w3[2], w3[3]); @@ -59,7 +59,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64); - counterbuf[7] = swap64 ((u64x) 0x400); + counterbuf[7] = 0x0004000000000000; //final: padding byte streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); @@ -74,7 +74,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 //outer HMAC: opad + digest //first transform: precalculated opad hash - counterbuf[7] = swap64 ((u64x) 0x200); + counterbuf[7] = 0x0002000000000000; //second transform: message = inner HMAC digest message[0] = digest[0]; @@ -97,7 +97,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64); - counterbuf[7] = swap64 ((u64x) 0x400); + counterbuf[7] = 0x0004000000000000; streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); diff --git a/OpenCL/m13772-pure.cl b/OpenCL/m13772-pure.cl index 2c3b0973f..eae06f757 100644 --- a/OpenCL/m13772-pure.cl +++ b/OpenCL/m13772-pure.cl @@ -31,12 +31,12 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 u64x padding[8] = { 0 }; u64x message[8]; - padding[7] = swap64 ((u64x) 0x01); + padding[7] = 0x0100000000000000; //inner HMAC: ipad + message //first transform: precalculated ipad hash - counterbuf[7] = swap64 ((u64x) 0x200); + counterbuf[7] = 0x0002000000000000; //second transform: message = previous HMAC digest message[7] = hl32_to_64 (w3[2], w3[3]); @@ -59,7 +59,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64); - counterbuf[7] = swap64 ((u64x) 0x400); + counterbuf[7] = 0x0004000000000000; //final: padding byte streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); @@ -74,7 +74,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 //outer HMAC: opad + digest //first transform: precalculated opad hash - counterbuf[7] = swap64 ((u64x) 0x200); + counterbuf[7] = 0x0002000000000000; //second transform: message = inner HMAC digest message[0] = digest[0]; @@ -97,7 +97,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64); - counterbuf[7] = swap64 ((u64x) 0x400); + counterbuf[7] = 0x0004000000000000; streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); @@ -627,7 +627,7 @@ __kernel void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -635,7 +635,7 @@ __kernel void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } diff --git a/OpenCL/m13773-pure.cl b/OpenCL/m13773-pure.cl index eae41cf23..dcba9b96b 100644 --- a/OpenCL/m13773-pure.cl +++ b/OpenCL/m13773-pure.cl @@ -31,12 +31,12 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 u64x padding[8] = { 0 }; u64x message[8]; - padding[7] = swap64 ((u64x) 0x01); + padding[7] = 0x0100000000000000; //inner HMAC: ipad + message //first transform: precalculated ipad hash - counterbuf[7] = swap64 ((u64x) 0x200); + counterbuf[7] = 0x0002000000000000; //second transform: message = previous HMAC digest message[7] = hl32_to_64 (w3[2], w3[3]); @@ -59,7 +59,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64); - counterbuf[7] = swap64 ((u64x) 0x400); + counterbuf[7] = 0x0004000000000000; //final: padding byte streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); @@ -74,7 +74,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 //outer HMAC: opad + digest //first transform: precalculated opad hash - counterbuf[7] = swap64 ((u64x) 0x200); + counterbuf[7] = 0x0002000000000000; //second transform: message = inner HMAC digest message[0] = digest[0]; @@ -97,7 +97,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6 streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64); - counterbuf[7] = swap64 ((u64x) 0x400); + counterbuf[7] = 0x0004000000000000; streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); @@ -627,7 +627,7 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -635,7 +635,7 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); } @@ -697,7 +697,7 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t)) if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1) { - if (atomic_inc (&hashes_shown[digests_offset]) == 0) + if (atomic_inc (&hashes_shown[0]) == 0) { mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0); }