From 2b9888486e022ab0b7a30a19061a3397139d6799 Mon Sep 17 00:00:00 2001 From: jsteube Date: Mon, 28 Aug 2017 22:20:40 +0200 Subject: [PATCH 1/2] Make -m 14000 a pure kernel only and add volatile for asm statement --- OpenCL/m01500_a3.cl | 2 +- OpenCL/m03000_a3.cl | 2 +- .../{m14000_a0-optimized.cl => m14000_a0.cl} | 140 +++++++-------- .../{m14000_a1-optimized.cl => m14000_a1.cl} | 168 ++++++++---------- .../{m14000_a3-optimized.cl => m14000_a3.cl} | 22 +-- 5 files changed, 143 insertions(+), 191 deletions(-) rename OpenCL/{m14000_a0-optimized.cl => m14000_a0.cl} (78%) rename OpenCL/{m14000_a1-optimized.cl => m14000_a1.cl} (78%) rename OpenCL/{m14000_a3-optimized.cl => m14000_a3.cl} (96%) diff --git a/OpenCL/m01500_a3.cl b/OpenCL/m01500_a3.cl index 67a3ffb10..4a89ad578 100644 --- a/OpenCL/m01500_a3.cl +++ b/OpenCL/m01500_a3.cl @@ -55,7 +55,7 @@ // The underlying mathematical formulas are NOT copyrighted. // -#define LUT(a,b,c,d,e) u32 a; asm ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); +#define LUT(a,b,c,d,e) u32 a; asm volatile ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4) { diff --git a/OpenCL/m03000_a3.cl b/OpenCL/m03000_a3.cl index 50fd8fff7..f2b01d31b 100644 --- a/OpenCL/m03000_a3.cl +++ b/OpenCL/m03000_a3.cl @@ -52,7 +52,7 @@ // The underlying mathematical formulas are NOT copyrighted. // -#define LUT(a,b,c,d,e) u32 a; asm ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); +#define LUT(a,b,c,d,e) u32 a; asm volatile ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4) { diff --git a/OpenCL/m14000_a0-optimized.cl b/OpenCL/m14000_a0.cl similarity index 78% rename from OpenCL/m14000_a0-optimized.cl rename to OpenCL/m14000_a0.cl index 7f9886b51..afca41a8a 100644 --- a/OpenCL/m14000_a0-optimized.cl +++ b/OpenCL/m14000_a0.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -380,23 +380,23 @@ __constant u32a c_skb[8][64] = #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf]) #endif -void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) +void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64]) { - u32x r = rotl32 (data[0], 3u); - u32x l = rotl32 (data[1], 3u); + u32 r = rotl32_S (data[0], 3u); + u32 l = rotl32_S (data[1], 3u); - u32x tt; + u32 tt; #ifdef _unroll #pragma unroll #endif for (u32 i = 0; i < 16; i += 2) { - u32x u; - u32x t; + u32 u; + u32 t; u = Kc[i + 0] ^ r; - t = Kd[i + 0] ^ rotl32 (r, 28u); + t = Kd[i + 0] ^ rotl32_S (r, 28u); l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -408,7 +408,7 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); u = Kc[i + 1] ^ l; - t = Kd[i + 1] ^ rotl32 (l, 28u); + t = Kd[i + 1] ^ rotl32_S (l, 28u); r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -420,13 +420,13 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); } - iv[0] = rotl32 (l, 29u); - iv[1] = rotl32 (r, 29u); + iv[0] = rotl32_S (l, 29u); + iv[1] = rotl32_S (r, 29u); } -void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64]) +void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64]) { - u32x tt; + u32 tt; PERM_OP (d, c, tt, 4, 0x0f0f0f0f); HPERM_OP (c, tt, 2, 0xcccc0000); @@ -461,42 +461,42 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 c = c & 0x0fffffff; d = d & 0x0fffffff; - const u32x c00 = (c >> 0) & 0x0000003f; - const u32x c06 = (c >> 6) & 0x00383003; - const u32x c07 = (c >> 7) & 0x0000003c; - const u32x c13 = (c >> 13) & 0x0000060f; - const u32x c20 = (c >> 20) & 0x00000001; + const u32 c00 = (c >> 0) & 0x0000003f; + const u32 c06 = (c >> 6) & 0x00383003; + const u32 c07 = (c >> 7) & 0x0000003c; + const u32 c13 = (c >> 13) & 0x0000060f; + const u32 c20 = (c >> 20) & 0x00000001; - u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb) - | BOX (((c06 >> 0) & 0xff) - |((c07 >> 0) & 0xff), 1, s_skb) - | BOX (((c13 >> 0) & 0xff) - |((c06 >> 8) & 0xff), 2, s_skb) - | BOX (((c20 >> 0) & 0xff) - |((c13 >> 8) & 0xff) - |((c06 >> 16) & 0xff), 3, s_skb); + u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb) + | BOX (((c06 >> 0) & 0xff) + |((c07 >> 0) & 0xff), 1, s_skb) + | BOX (((c13 >> 0) & 0xff) + |((c06 >> 8) & 0xff), 2, s_skb) + | BOX (((c20 >> 0) & 0xff) + |((c13 >> 8) & 0xff) + |((c06 >> 16) & 0xff), 3, s_skb); - const u32x d00 = (d >> 0) & 0x00003c3f; - const u32x d07 = (d >> 7) & 0x00003f03; - const u32x d21 = (d >> 21) & 0x0000000f; - const u32x d22 = (d >> 22) & 0x00000030; + const u32 d00 = (d >> 0) & 0x00003c3f; + const u32 d07 = (d >> 7) & 0x00003f03; + const u32 d21 = (d >> 21) & 0x0000000f; + const u32 d22 = (d >> 22) & 0x00000030; - u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb) - | BOX (((d07 >> 0) & 0xff) - |((d00 >> 8) & 0xff), 5, s_skb) - | BOX (((d07 >> 8) & 0xff), 6, s_skb) - | BOX (((d21 >> 0) & 0xff) - |((d22 >> 0) & 0xff), 7, s_skb); + u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb) + | BOX (((d07 >> 0) & 0xff) + |((d00 >> 8) & 0xff), 5, s_skb) + | BOX (((d07 >> 8) & 0xff), 6, s_skb) + | BOX (((d21 >> 0) & 0xff) + |((d22 >> 0) & 0xff), 7, s_skb); Kc[i] = ((t << 16) | (s & 0x0000ffff)); Kd[i] = ((s >> 16) | (t & 0xffff0000)); - Kc[i] = rotl32 (Kc[i], 2u); - Kd[i] = rotl32 (Kd[i], 2u); + Kc[i] = rotl32_S (Kc[i], 2u); + Kd[i] = rotl32_S (Kd[i], 2u); } } -__kernel void m14000_m04 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14000_mxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -571,45 +571,37 @@ __kernel void m14000_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - u32x w0[4] = { 0 }; - u32x w1[4] = { 0 }; - u32x w2[4] = { 0 }; - u32x w3[4] = { 0 }; + u32 w0[4] = { 0 }; + u32 w1[4] = { 0 }; + u32 w2[4] = { 0 }; + u32 w3[4] = { 0 }; apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); - const u32x c = w0[0]; - const u32x d = w0[1]; + const u32 c = w0[0]; + const u32 d = w0[1]; - u32x Kc[16]; - u32x Kd[16]; + u32 Kc[16]; + u32 Kd[16]; _des_crypt_keysetup (c, d, Kc, Kd, s_skb); - u32x data[2]; + u32 data[2]; data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x iv[2]; + u32 iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); - u32x z = 0; + u32 z = 0; COMPARE_M_SIMD (iv[0], iv[1], z, z); } } -__kernel void m14000_m08 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_m16 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_s04 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14000_sxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -696,40 +688,32 @@ __kernel void m14000_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - u32x w0[4] = { 0 }; - u32x w1[4] = { 0 }; - u32x w2[4] = { 0 }; - u32x w3[4] = { 0 }; + u32 w0[4] = { 0 }; + u32 w1[4] = { 0 }; + u32 w2[4] = { 0 }; + u32 w3[4] = { 0 }; apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); - const u32x c = w0[0]; - const u32x d = w0[1]; + const u32 c = w0[0]; + const u32 d = w0[1]; - u32x Kc[16]; - u32x Kd[16]; + u32 Kc[16]; + u32 Kd[16]; _des_crypt_keysetup (c, d, Kc, Kd, s_skb); - u32x data[2]; + u32 data[2]; data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x iv[2]; + u32 iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); - u32x z = 0; + u32 z = 0; COMPARE_S_SIMD (iv[0], iv[1], z, z); } } - -__kernel void m14000_s08 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_s16 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} diff --git a/OpenCL/m14000_a1-optimized.cl b/OpenCL/m14000_a1.cl similarity index 78% rename from OpenCL/m14000_a1-optimized.cl rename to OpenCL/m14000_a1.cl index d3a4392e0..faf95aaf5 100644 --- a/OpenCL/m14000_a1-optimized.cl +++ b/OpenCL/m14000_a1.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -370,23 +370,23 @@ __constant u32a c_skb[8][64] = #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf]) #endif -void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) +void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64]) { - u32x r = rotl32 (data[0], 3u); - u32x l = rotl32 (data[1], 3u); + u32 r = rotl32_S (data[0], 3u); + u32 l = rotl32_S (data[1], 3u); - u32x tt; + u32 tt; #ifdef _unroll #pragma unroll #endif for (u32 i = 0; i < 16; i += 2) { - u32x u; - u32x t; + u32 u; + u32 t; u = Kc[i + 0] ^ r; - t = Kd[i + 0] ^ rotl32 (r, 28u); + t = Kd[i + 0] ^ rotl32_S (r, 28u); l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -398,7 +398,7 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); u = Kc[i + 1] ^ l; - t = Kd[i + 1] ^ rotl32 (l, 28u); + t = Kd[i + 1] ^ rotl32_S (l, 28u); r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -410,13 +410,13 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); } - iv[0] = rotl32 (l, 29u); - iv[1] = rotl32 (r, 29u); + iv[0] = rotl32_S (l, 29u); + iv[1] = rotl32_S (r, 29u); } -void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64]) +void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64]) { - u32x tt; + u32 tt; PERM_OP (d, c, tt, 4, 0x0f0f0f0f); HPERM_OP (c, tt, 2, 0xcccc0000); @@ -451,42 +451,42 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 c = c & 0x0fffffff; d = d & 0x0fffffff; - const u32x c00 = (c >> 0) & 0x0000003f; - const u32x c06 = (c >> 6) & 0x00383003; - const u32x c07 = (c >> 7) & 0x0000003c; - const u32x c13 = (c >> 13) & 0x0000060f; - const u32x c20 = (c >> 20) & 0x00000001; + const u32 c00 = (c >> 0) & 0x0000003f; + const u32 c06 = (c >> 6) & 0x00383003; + const u32 c07 = (c >> 7) & 0x0000003c; + const u32 c13 = (c >> 13) & 0x0000060f; + const u32 c20 = (c >> 20) & 0x00000001; - u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb) - | BOX (((c06 >> 0) & 0xff) - |((c07 >> 0) & 0xff), 1, s_skb) - | BOX (((c13 >> 0) & 0xff) - |((c06 >> 8) & 0xff), 2, s_skb) - | BOX (((c20 >> 0) & 0xff) - |((c13 >> 8) & 0xff) - |((c06 >> 16) & 0xff), 3, s_skb); + u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb) + | BOX (((c06 >> 0) & 0xff) + |((c07 >> 0) & 0xff), 1, s_skb) + | BOX (((c13 >> 0) & 0xff) + |((c06 >> 8) & 0xff), 2, s_skb) + | BOX (((c20 >> 0) & 0xff) + |((c13 >> 8) & 0xff) + |((c06 >> 16) & 0xff), 3, s_skb); - const u32x d00 = (d >> 0) & 0x00003c3f; - const u32x d07 = (d >> 7) & 0x00003f03; - const u32x d21 = (d >> 21) & 0x0000000f; - const u32x d22 = (d >> 22) & 0x00000030; + const u32 d00 = (d >> 0) & 0x00003c3f; + const u32 d07 = (d >> 7) & 0x00003f03; + const u32 d21 = (d >> 21) & 0x0000000f; + const u32 d22 = (d >> 22) & 0x00000030; - u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb) - | BOX (((d07 >> 0) & 0xff) - |((d00 >> 8) & 0xff), 5, s_skb) - | BOX (((d07 >> 8) & 0xff), 6, s_skb) - | BOX (((d21 >> 0) & 0xff) - |((d22 >> 0) & 0xff), 7, s_skb); + u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb) + | BOX (((d07 >> 0) & 0xff) + |((d00 >> 8) & 0xff), 5, s_skb) + | BOX (((d07 >> 8) & 0xff), 6, s_skb) + | BOX (((d21 >> 0) & 0xff) + |((d22 >> 0) & 0xff), 7, s_skb); Kc[i] = ((t << 16) | (s & 0x0000ffff)); Kd[i] = ((s >> 16) | (t & 0xffff0000)); - Kc[i] = rotl32 (Kc[i], 2u); - Kd[i] = rotl32 (Kd[i], 2u); + Kc[i] = rotl32_S (Kc[i], 2u); + Kd[i] = rotl32_S (Kd[i], 2u); } } -__kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14000_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -561,18 +561,18 @@ __kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); - const u32x pw_len = pw_l_len + pw_r_len; + const u32 pw_len = pw_l_len + pw_r_len; /** * concat password candidate */ - u32x wordl0[4] = { 0 }; - u32x wordl1[4] = { 0 }; - u32x wordl2[4] = { 0 }; - u32x wordl3[4] = { 0 }; + u32 wordl0[4] = { 0 }; + u32 wordl1[4] = { 0 }; + u32 wordl2[4] = { 0 }; + u32 wordl3[4] = { 0 }; wordl0[0] = pw_buf0[0]; wordl0[1] = pw_buf0[1]; @@ -583,10 +583,10 @@ __kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule wordl1[2] = pw_buf1[2]; wordl1[3] = pw_buf1[3]; - u32x wordr0[4] = { 0 }; - u32x wordr1[4] = { 0 }; - u32x wordr2[4] = { 0 }; - u32x wordr3[4] = { 0 }; + u32 wordr0[4] = { 0 }; + u32 wordr1[4] = { 0 }; + u32 wordr2[4] = { 0 }; + u32 wordr3[4] = { 0 }; wordr0[0] = ix_create_combt (combs_buf, il_pos, 0); wordr0[1] = ix_create_combt (combs_buf, il_pos, 1); @@ -606,43 +606,35 @@ __kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len); } - u32x w0[2]; + u32 w0[2]; w0[0] = wordl0[0] | wordr0[0]; w0[1] = wordl0[1] | wordr0[1]; - const u32x c = w0[0]; - const u32x d = w0[1]; + const u32 c = w0[0]; + const u32 d = w0[1]; - u32x Kc[16]; - u32x Kd[16]; + u32 Kc[16]; + u32 Kd[16]; _des_crypt_keysetup (c, d, Kc, Kd, s_skb); - u32x data[2]; + u32 data[2]; data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x iv[2]; + u32 iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); - u32x z = 0; + u32 z = 0; COMPARE_M_SIMD (iv[0], iv[1], z, z); } } -__kernel void m14000_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14000_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -729,18 +721,18 @@ __kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); - const u32x pw_len = pw_l_len + pw_r_len; + const u32 pw_len = pw_l_len + pw_r_len; /** * concat password candidate */ - u32x wordl0[4] = { 0 }; - u32x wordl1[4] = { 0 }; - u32x wordl2[4] = { 0 }; - u32x wordl3[4] = { 0 }; + u32 wordl0[4] = { 0 }; + u32 wordl1[4] = { 0 }; + u32 wordl2[4] = { 0 }; + u32 wordl3[4] = { 0 }; wordl0[0] = pw_buf0[0]; wordl0[1] = pw_buf0[1]; @@ -751,10 +743,10 @@ __kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule wordl1[2] = pw_buf1[2]; wordl1[3] = pw_buf1[3]; - u32x wordr0[4] = { 0 }; - u32x wordr1[4] = { 0 }; - u32x wordr2[4] = { 0 }; - u32x wordr3[4] = { 0 }; + u32 wordr0[4] = { 0 }; + u32 wordr1[4] = { 0 }; + u32 wordr2[4] = { 0 }; + u32 wordr3[4] = { 0 }; wordr0[0] = ix_create_combt (combs_buf, il_pos, 0); wordr0[1] = ix_create_combt (combs_buf, il_pos, 1); @@ -774,38 +766,30 @@ __kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len); } - u32x w0[2]; + u32 w0[2]; w0[0] = wordl0[0] | wordr0[0]; w0[1] = wordl0[1] | wordr0[1]; - const u32x c = (w0[0]); - const u32x d = (w0[1]); + const u32 c = (w0[0]); + const u32 d = (w0[1]); - u32x Kc[16]; - u32x Kd[16]; + u32 Kc[16]; + u32 Kd[16]; _des_crypt_keysetup (c, d, Kc, Kd, s_skb); - u32x data[2]; + u32 data[2]; data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x iv[2]; + u32 iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); - u32x z = 0; + u32 z = 0; COMPARE_S_SIMD (iv[0], iv[1], z, z); } } - -__kernel void m14000_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} diff --git a/OpenCL/m14000_a3-optimized.cl b/OpenCL/m14000_a3.cl similarity index 96% rename from OpenCL/m14000_a3-optimized.cl rename to OpenCL/m14000_a3.cl index 58689f7f2..9dcc7fd36 100644 --- a/OpenCL/m14000_a3-optimized.cl +++ b/OpenCL/m14000_a3.cl @@ -52,7 +52,7 @@ // The underlying mathematical formulas are NOT copyrighted. // -#define LUT(a,b,c,d,e) u32 a; asm ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); +#define LUT(a,b,c,d,e) u32 a; asm volatile ("lop3.b32 %0, %1, %2, %3, "#e";" : "=r"(a): "r"(b), "r"(c), "r"(d)); void s1 (const u32 a1, const u32 a2, const u32 a3, const u32 a4, const u32 a5, const u32 a6, u32 *out1, u32 *out2, u32 *out3, u32 *out4) { @@ -2714,7 +2714,7 @@ __kernel void m14000_tm (__global u32 *mod, __global bs_word_t *words_buf_r) } } -__kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14000_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -2732,15 +2732,7 @@ __kernel void m14000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule m14000m (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); } -__kernel void m14000_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14000_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -2757,11 +2749,3 @@ __kernel void m14000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule m14000s (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); } - -__kernel void m14000_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14000_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bs_word_t *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} From 1c169af0ad99a682f2c003560dbef458b9ffeb85 Mon Sep 17 00:00:00 2001 From: jsteube Date: Mon, 28 Aug 2017 22:26:30 +0200 Subject: [PATCH 2/2] Make -m 14100 a pure kernel only --- .../{m14100_a0-optimized.cl => m14100_a0.cl} | 186 +++++------ .../{m14100_a1-optimized.cl => m14100_a1.cl} | 232 +++++++------- .../{m14100_a3-optimized.cl => m14100_a3.cl} | 300 +----------------- 3 files changed, 195 insertions(+), 523 deletions(-) rename OpenCL/{m14100_a0-optimized.cl => m14100_a0.cl} (78%) rename OpenCL/{m14100_a1-optimized.cl => m14100_a1.cl} (78%) rename OpenCL/{m14100_a3-optimized.cl => m14100_a3.cl} (73%) diff --git a/OpenCL/m14100_a0-optimized.cl b/OpenCL/m14100_a0.cl similarity index 78% rename from OpenCL/m14100_a0-optimized.cl rename to OpenCL/m14100_a0.cl index 0a9129c06..2d038630b 100644 --- a/OpenCL/m14100_a0-optimized.cl +++ b/OpenCL/m14100_a0.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -380,23 +380,23 @@ __constant u32a c_skb[8][64] = #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf]) #endif -void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) +void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64]) { - u32x r = rotl32 (data[0], 3u); - u32x l = rotl32 (data[1], 3u); + u32 r = rotl32_S (data[0], 3u); + u32 l = rotl32_S (data[1], 3u); - u32x tt; + u32 tt; #ifdef _unroll #pragma unroll #endif for (u32 i = 0; i < 16; i += 2) { - u32x u; - u32x t; + u32 u; + u32 t; u = Kc[i + 0] ^ r; - t = Kd[i + 0] ^ rotl32 (r, 28u); + t = Kd[i + 0] ^ rotl32_S (r, 28u); l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -408,7 +408,7 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); u = Kc[i + 1] ^ l; - t = Kd[i + 1] ^ rotl32 (l, 28u); + t = Kd[i + 1] ^ rotl32_S (l, 28u); r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -420,27 +420,27 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); } - iv[0] = rotl32 (l, 29u); - iv[1] = rotl32 (r, 29u); + iv[0] = rotl32_S (l, 29u); + iv[1] = rotl32_S (r, 29u); } -void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) +void _des_crypt_decrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64]) { - u32x r = rotl32 (data[0], 3u); - u32x l = rotl32 (data[1], 3u); + u32 r = rotl32_S (data[0], 3u); + u32 l = rotl32_S (data[1], 3u); - u32x tt; + u32 tt; #ifdef _unroll #pragma unroll #endif for (u32 i = 16; i > 0; i -= 2) { - u32x u; - u32x t; + u32 u; + u32 t; u = Kc[i - 1] ^ r; - t = Kd[i - 1] ^ rotl32 (r, 28u); + t = Kd[i - 1] ^ rotl32_S (r, 28u); l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -452,7 +452,7 @@ void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); u = Kc[i - 2] ^ l; - t = Kd[i - 2] ^ rotl32 (l, 28u); + t = Kd[i - 2] ^ rotl32_S (l, 28u); r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -464,13 +464,13 @@ void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); } - iv[0] = rotl32 (l, 29u); - iv[1] = rotl32 (r, 29u); + iv[0] = rotl32_S (l, 29u); + iv[1] = rotl32_S (r, 29u); } -void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64]) +void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64]) { - u32x tt; + u32 tt; PERM_OP (d, c, tt, 4, 0x0f0f0f0f); HPERM_OP (c, tt, 2, 0xcccc0000); @@ -505,13 +505,13 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 c = c & 0x0fffffff; d = d & 0x0fffffff; - const u32x c00 = (c >> 0) & 0x0000003f; - const u32x c06 = (c >> 6) & 0x00383003; - const u32x c07 = (c >> 7) & 0x0000003c; - const u32x c13 = (c >> 13) & 0x0000060f; - const u32x c20 = (c >> 20) & 0x00000001; + const u32 c00 = (c >> 0) & 0x0000003f; + const u32 c06 = (c >> 6) & 0x00383003; + const u32 c07 = (c >> 7) & 0x0000003c; + const u32 c13 = (c >> 13) & 0x0000060f; + const u32 c20 = (c >> 20) & 0x00000001; - u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb) + u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb) | BOX (((c06 >> 0) & 0xff) |((c07 >> 0) & 0xff), 1, s_skb) | BOX (((c13 >> 0) & 0xff) @@ -520,27 +520,27 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 |((c13 >> 8) & 0xff) |((c06 >> 16) & 0xff), 3, s_skb); - const u32x d00 = (d >> 0) & 0x00003c3f; - const u32x d07 = (d >> 7) & 0x00003f03; - const u32x d21 = (d >> 21) & 0x0000000f; - const u32x d22 = (d >> 22) & 0x00000030; + const u32 d00 = (d >> 0) & 0x00003c3f; + const u32 d07 = (d >> 7) & 0x00003f03; + const u32 d21 = (d >> 21) & 0x0000000f; + const u32 d22 = (d >> 22) & 0x00000030; - u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb) - | BOX (((d07 >> 0) & 0xff) - |((d00 >> 8) & 0xff), 5, s_skb) - | BOX (((d07 >> 8) & 0xff), 6, s_skb) - | BOX (((d21 >> 0) & 0xff) - |((d22 >> 0) & 0xff), 7, s_skb); + u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb) + | BOX (((d07 >> 0) & 0xff) + |((d00 >> 8) & 0xff), 5, s_skb) + | BOX (((d07 >> 8) & 0xff), 6, s_skb) + | BOX (((d21 >> 0) & 0xff) + |((d22 >> 0) & 0xff), 7, s_skb); Kc[i] = ((t << 16) | (s & 0x0000ffff)); Kd[i] = ((s >> 16) | (t & 0xffff0000)); - Kc[i] = rotl32 (Kc[i], 2u); - Kd[i] = rotl32 (Kd[i], 2u); + Kc[i] = rotl32_S (Kc[i], 2u); + Kd[i] = rotl32_S (Kd[i], 2u); } } -__kernel void m14100_m04 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14100_mxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -615,75 +615,67 @@ __kernel void m14100_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - u32x w0[4] = { 0 }; - u32x w1[4] = { 0 }; - u32x w2[4] = { 0 }; - u32x w3[4] = { 0 }; + u32 w0[4] = { 0 }; + u32 w1[4] = { 0 }; + u32 w2[4] = { 0 }; + u32 w3[4] = { 0 }; apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); /* First Pass */ - const u32x a = w0[0]; - const u32x b = w0[1]; + const u32 a = w0[0]; + const u32 b = w0[1]; - u32x Ka[16]; - u32x Kb[16]; + u32 Ka[16]; + u32 Kb[16]; _des_crypt_keysetup (a, b, Ka, Kb, s_skb); - u32x data[2]; + u32 data[2]; data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x p1[2]; + u32 p1[2]; _des_crypt_encrypt (p1, data, Ka, Kb, s_SPtrans); /* Second Pass */ - const u32x c = w0[2]; - const u32x d = w0[3]; + const u32 c = w0[2]; + const u32 d = w0[3]; - u32x Kc[16]; - u32x Kd[16]; + u32 Kc[16]; + u32 Kd[16]; _des_crypt_keysetup (c, d, Kc, Kd, s_skb); - u32x p2[2]; + u32 p2[2]; _des_crypt_decrypt (p2, p1, Kc, Kd, s_SPtrans); /* Third Pass */ - const u32x e = w1[0]; - const u32x f = w1[1]; + const u32 e = w1[0]; + const u32 f = w1[1]; - u32x Ke[16]; - u32x Kf[16]; + u32 Ke[16]; + u32 Kf[16]; _des_crypt_keysetup (e, f, Ke, Kf, s_skb); - u32x iv[2]; + u32 iv[2]; _des_crypt_encrypt (iv, p2, Ke, Kf, s_SPtrans); - u32x z = 0; + u32 z = 0; COMPARE_M_SIMD (iv[0], iv[1], z, z); } } -__kernel void m14100_m08 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14100_m16 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14100_s04 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14100_sxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -770,70 +762,62 @@ __kernel void m14100_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - u32x w0[4] = { 0 }; - u32x w1[4] = { 0 }; - u32x w2[4] = { 0 }; - u32x w3[4] = { 0 }; + u32 w0[4] = { 0 }; + u32 w1[4] = { 0 }; + u32 w2[4] = { 0 }; + u32 w3[4] = { 0 }; apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); /* First Pass */ - const u32x a = w0[0]; - const u32x b = w0[1]; + const u32 a = w0[0]; + const u32 b = w0[1]; - u32x Ka[16]; - u32x Kb[16]; + u32 Ka[16]; + u32 Kb[16]; _des_crypt_keysetup (a, b, Ka, Kb, s_skb); - u32x data[2]; + u32 data[2]; data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x p1[2]; + u32 p1[2]; _des_crypt_encrypt (p1, data, Ka, Kb, s_SPtrans); /* Second Pass */ - const u32x c = w0[2]; - const u32x d = w0[3]; + const u32 c = w0[2]; + const u32 d = w0[3]; - u32x Kc[16]; - u32x Kd[16]; + u32 Kc[16]; + u32 Kd[16]; _des_crypt_keysetup (c, d, Kc, Kd, s_skb); - u32x p2[2]; + u32 p2[2]; _des_crypt_decrypt (p2, p1, Kc, Kd, s_SPtrans); /* Third Pass */ - const u32x e = w1[0]; - const u32x f = w1[1]; + const u32 e = w1[0]; + const u32 f = w1[1]; - u32x Ke[16]; - u32x Kf[16]; + u32 Ke[16]; + u32 Kf[16]; _des_crypt_keysetup (e, f, Ke, Kf, s_skb); - u32x iv[2]; + u32 iv[2]; _des_crypt_encrypt (iv, p2, Ke, Kf, s_SPtrans); - u32x z = 0; + u32 z = 0; COMPARE_S_SIMD (iv[0], iv[1], z, z); } } - -__kernel void m14100_s08 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14100_s16 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} diff --git a/OpenCL/m14100_a1-optimized.cl b/OpenCL/m14100_a1.cl similarity index 78% rename from OpenCL/m14100_a1-optimized.cl rename to OpenCL/m14100_a1.cl index 4af156ea3..69cde6c9a 100644 --- a/OpenCL/m14100_a1-optimized.cl +++ b/OpenCL/m14100_a1.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -370,23 +370,23 @@ __constant u32a c_skb[8][64] = #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf]) #endif -void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) +void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64]) { - u32x r = rotl32 (data[0], 3u); - u32x l = rotl32 (data[1], 3u); + u32 r = rotl32_S (data[0], 3u); + u32 l = rotl32_S (data[1], 3u); - u32x tt; + u32 tt; #ifdef _unroll #pragma unroll #endif for (u32 i = 0; i < 16; i += 2) { - u32x u; - u32x t; + u32 u; + u32 t; u = Kc[i + 0] ^ r; - t = Kd[i + 0] ^ rotl32 (r, 28u); + t = Kd[i + 0] ^ rotl32_S (r, 28u); l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -398,7 +398,7 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); u = Kc[i + 1] ^ l; - t = Kd[i + 1] ^ rotl32 (l, 28u); + t = Kd[i + 1] ^ rotl32_S (l, 28u); r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -410,27 +410,27 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); } - iv[0] = rotl32 (l, 29u); - iv[1] = rotl32 (r, 29u); + iv[0] = rotl32_S (l, 29u); + iv[1] = rotl32_S (r, 29u); } -void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) +void _des_crypt_decrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64]) { - u32x r = rotl32 (data[0], 3u); - u32x l = rotl32 (data[1], 3u); + u32 r = rotl32_S (data[0], 3u); + u32 l = rotl32_S (data[1], 3u); - u32x tt; + u32 tt; #ifdef _unroll #pragma unroll #endif for (u32 i = 16; i > 0; i -= 2) { - u32x u; - u32x t; + u32 u; + u32 t; u = Kc[i - 1] ^ r; - t = Kd[i - 1] ^ rotl32 (r, 28u); + t = Kd[i - 1] ^ rotl32_S (r, 28u); l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -442,7 +442,7 @@ void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); u = Kc[i - 2] ^ l; - t = Kd[i - 2] ^ rotl32 (l, 28u); + t = Kd[i - 2] ^ rotl32_S (l, 28u); r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) @@ -454,13 +454,13 @@ void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); } - iv[0] = rotl32 (l, 29u); - iv[1] = rotl32 (r, 29u); + iv[0] = rotl32_S (l, 29u); + iv[1] = rotl32_S (r, 29u); } -void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64]) +void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64]) { - u32x tt; + u32 tt; PERM_OP (d, c, tt, 4, 0x0f0f0f0f); HPERM_OP (c, tt, 2, 0xcccc0000); @@ -495,42 +495,42 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 c = c & 0x0fffffff; d = d & 0x0fffffff; - const u32x c00 = (c >> 0) & 0x0000003f; - const u32x c06 = (c >> 6) & 0x00383003; - const u32x c07 = (c >> 7) & 0x0000003c; - const u32x c13 = (c >> 13) & 0x0000060f; - const u32x c20 = (c >> 20) & 0x00000001; + const u32 c00 = (c >> 0) & 0x0000003f; + const u32 c06 = (c >> 6) & 0x00383003; + const u32 c07 = (c >> 7) & 0x0000003c; + const u32 c13 = (c >> 13) & 0x0000060f; + const u32 c20 = (c >> 20) & 0x00000001; - u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb) - | BOX (((c06 >> 0) & 0xff) - |((c07 >> 0) & 0xff), 1, s_skb) - | BOX (((c13 >> 0) & 0xff) - |((c06 >> 8) & 0xff), 2, s_skb) - | BOX (((c20 >> 0) & 0xff) - |((c13 >> 8) & 0xff) - |((c06 >> 16) & 0xff), 3, s_skb); + u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb) + | BOX (((c06 >> 0) & 0xff) + |((c07 >> 0) & 0xff), 1, s_skb) + | BOX (((c13 >> 0) & 0xff) + |((c06 >> 8) & 0xff), 2, s_skb) + | BOX (((c20 >> 0) & 0xff) + |((c13 >> 8) & 0xff) + |((c06 >> 16) & 0xff), 3, s_skb); - const u32x d00 = (d >> 0) & 0x00003c3f; - const u32x d07 = (d >> 7) & 0x00003f03; - const u32x d21 = (d >> 21) & 0x0000000f; - const u32x d22 = (d >> 22) & 0x00000030; + const u32 d00 = (d >> 0) & 0x00003c3f; + const u32 d07 = (d >> 7) & 0x00003f03; + const u32 d21 = (d >> 21) & 0x0000000f; + const u32 d22 = (d >> 22) & 0x00000030; - u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb) - | BOX (((d07 >> 0) & 0xff) - |((d00 >> 8) & 0xff), 5, s_skb) - | BOX (((d07 >> 8) & 0xff), 6, s_skb) - | BOX (((d21 >> 0) & 0xff) - |((d22 >> 0) & 0xff), 7, s_skb); + u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb) + | BOX (((d07 >> 0) & 0xff) + |((d00 >> 8) & 0xff), 5, s_skb) + | BOX (((d07 >> 8) & 0xff), 6, s_skb) + | BOX (((d21 >> 0) & 0xff) + |((d22 >> 0) & 0xff), 7, s_skb); Kc[i] = ((t << 16) | (s & 0x0000ffff)); Kd[i] = ((s >> 16) | (t & 0xffff0000)); - Kc[i] = rotl32 (Kc[i], 2u); - Kd[i] = rotl32 (Kd[i], 2u); + Kc[i] = rotl32_S (Kc[i], 2u); + Kd[i] = rotl32_S (Kd[i], 2u); } } -__kernel void m14100_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14100_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -605,18 +605,18 @@ __kernel void m14100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); - const u32x pw_len = pw_l_len + pw_r_len; + const u32 pw_len = pw_l_len + pw_r_len; /** * concat password candidate */ - u32x wordl0[4] = { 0 }; - u32x wordl1[4] = { 0 }; - u32x wordl2[4] = { 0 }; - u32x wordl3[4] = { 0 }; + u32 wordl0[4] = { 0 }; + u32 wordl1[4] = { 0 }; + u32 wordl2[4] = { 0 }; + u32 wordl3[4] = { 0 }; wordl0[0] = pw_buf0[0]; wordl0[1] = pw_buf0[1]; @@ -627,10 +627,10 @@ __kernel void m14100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule wordl1[2] = pw_buf1[2]; wordl1[3] = pw_buf1[3]; - u32x wordr0[4] = { 0 }; - u32x wordr1[4] = { 0 }; - u32x wordr2[4] = { 0 }; - u32x wordr3[4] = { 0 }; + u32 wordr0[4] = { 0 }; + u32 wordr1[4] = { 0 }; + u32 wordr2[4] = { 0 }; + u32 wordr3[4] = { 0 }; wordr0[0] = ix_create_combt (combs_buf, il_pos, 0); wordr0[1] = ix_create_combt (combs_buf, il_pos, 1); @@ -650,8 +650,8 @@ __kernel void m14100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len); } - u32x w0[4]; - u32x w1[4]; + u32 w0[4]; + u32 w1[4]; w0[0] = wordl0[0] | wordr0[0]; w0[1] = wordl0[1] | wordr0[1]; @@ -662,66 +662,58 @@ __kernel void m14100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule /* First Pass */ - const u32x a = w0[0]; - const u32x b = w0[1]; + const u32 a = w0[0]; + const u32 b = w0[1]; - u32x Ka[16]; - u32x Kb[16]; + u32 Ka[16]; + u32 Kb[16]; _des_crypt_keysetup (a, b, Ka, Kb, s_skb); - u32x data[2]; + u32 data[2]; data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x p1[2]; + u32 p1[2]; _des_crypt_encrypt (p1, data, Ka, Kb, s_SPtrans); /* Second Pass */ - const u32x c = w0[2]; - const u32x d = w0[3]; + const u32 c = w0[2]; + const u32 d = w0[3]; - u32x Kc[16]; - u32x Kd[16]; + u32 Kc[16]; + u32 Kd[16]; _des_crypt_keysetup (c, d, Kc, Kd, s_skb); - u32x p2[2]; + u32 p2[2]; _des_crypt_decrypt (p2, p1, Kc, Kd, s_SPtrans); /* Third Pass */ - const u32x e = w1[0]; - const u32x f = w1[1]; + const u32 e = w1[0]; + const u32 f = w1[1]; - u32x Ke[16]; - u32x Kf[16]; + u32 Ke[16]; + u32 Kf[16]; _des_crypt_keysetup (e, f, Ke, Kf, s_skb); - u32x iv[2]; + u32 iv[2]; _des_crypt_encrypt (iv, p2, Ke, Kf, s_SPtrans); - u32x z = 0; + u32 z = 0; COMPARE_M_SIMD (iv[0], iv[1], z, z); } } -__kernel void m14100_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14100_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14100_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14100_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -808,18 +800,18 @@ __kernel void m14100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); - const u32x pw_len = pw_l_len + pw_r_len; + const u32 pw_len = pw_l_len + pw_r_len; /** * concat password candidate */ - u32x wordl0[4] = { 0 }; - u32x wordl1[4] = { 0 }; - u32x wordl2[4] = { 0 }; - u32x wordl3[4] = { 0 }; + u32 wordl0[4] = { 0 }; + u32 wordl1[4] = { 0 }; + u32 wordl2[4] = { 0 }; + u32 wordl3[4] = { 0 }; wordl0[0] = pw_buf0[0]; wordl0[1] = pw_buf0[1]; @@ -830,10 +822,10 @@ __kernel void m14100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule wordl1[2] = pw_buf1[2]; wordl1[3] = pw_buf1[3]; - u32x wordr0[4] = { 0 }; - u32x wordr1[4] = { 0 }; - u32x wordr2[4] = { 0 }; - u32x wordr3[4] = { 0 }; + u32 wordr0[4] = { 0 }; + u32 wordr1[4] = { 0 }; + u32 wordr2[4] = { 0 }; + u32 wordr3[4] = { 0 }; wordr0[0] = ix_create_combt (combs_buf, il_pos, 0); wordr0[1] = ix_create_combt (combs_buf, il_pos, 1); @@ -853,8 +845,8 @@ __kernel void m14100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len); } - u32x w0[4]; - u32x w1[4]; + u32 w0[4]; + u32 w1[4]; w0[0] = wordl0[0] | wordr0[0]; w0[1] = wordl0[1] | wordr0[1]; @@ -865,61 +857,53 @@ __kernel void m14100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule /* First Pass */ - const u32x a = w0[0]; - const u32x b = w0[1]; + const u32 a = w0[0]; + const u32 b = w0[1]; - u32x Ka[16]; - u32x Kb[16]; + u32 Ka[16]; + u32 Kb[16]; _des_crypt_keysetup (a, b, Ka, Kb, s_skb); - u32x data[2]; + u32 data[2]; data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x p1[2]; + u32 p1[2]; _des_crypt_encrypt (p1, data, Ka, Kb, s_SPtrans); /* Second Pass */ - const u32x c = w0[2]; - const u32x d = w0[3]; + const u32 c = w0[2]; + const u32 d = w0[3]; - u32x Kc[16]; - u32x Kd[16]; + u32 Kc[16]; + u32 Kd[16]; _des_crypt_keysetup (c, d, Kc, Kd, s_skb); - u32x p2[2]; + u32 p2[2]; _des_crypt_decrypt (p2, p1, Kc, Kd, s_SPtrans); /* Third Pass */ - const u32x e = w1[0]; - const u32x f = w1[1]; + const u32 e = w1[0]; + const u32 f = w1[1]; - u32x Ke[16]; - u32x Kf[16]; + u32 Ke[16]; + u32 Kf[16]; _des_crypt_keysetup (e, f, Ke, Kf, s_skb); - u32x iv[2]; + u32 iv[2]; _des_crypt_encrypt (iv, p2, Ke, Kf, s_SPtrans); - u32x z = 0; + u32 z = 0; COMPARE_S_SIMD (iv[0], iv[1], z, z); } } - -__kernel void m14100_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} - -__kernel void m14100_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ -} diff --git a/OpenCL/m14100_a3-optimized.cl b/OpenCL/m14100_a3.cl similarity index 73% rename from OpenCL/m14100_a3-optimized.cl rename to OpenCL/m14100_a3.cl index 24d5b6db3..deff80f71 100644 --- a/OpenCL/m14100_a3-optimized.cl +++ b/OpenCL/m14100_a3.cl @@ -712,7 +712,7 @@ void m14100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], } } -__kernel void m14100_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14100_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -786,303 +786,7 @@ __kernel void m14100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule m14100m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); } -__kernel void m14100_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ - /** - * base - */ - - const u64 gid = get_global_id (0); - const u64 lid = get_local_id (0); - const u64 lsz = get_local_size (0); - - /** - * shared - */ - - __local u32 s_SPtrans[8][64]; - __local u32 s_skb[8][64]; - - for (u32 i = lid; i < 64; i += lsz) - { - s_SPtrans[0][i] = c_SPtrans[0][i]; - s_SPtrans[1][i] = c_SPtrans[1][i]; - s_SPtrans[2][i] = c_SPtrans[2][i]; - s_SPtrans[3][i] = c_SPtrans[3][i]; - s_SPtrans[4][i] = c_SPtrans[4][i]; - s_SPtrans[5][i] = c_SPtrans[5][i]; - s_SPtrans[6][i] = c_SPtrans[6][i]; - s_SPtrans[7][i] = c_SPtrans[7][i]; - - s_skb[0][i] = c_skb[0][i]; - s_skb[1][i] = c_skb[1][i]; - s_skb[2][i] = c_skb[2][i]; - s_skb[3][i] = c_skb[3][i]; - s_skb[4][i] = c_skb[4][i]; - s_skb[5][i] = c_skb[5][i]; - s_skb[6][i] = c_skb[6][i]; - s_skb[7][i] = c_skb[7][i]; - } - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - - /** - * base - */ - - u32 w[16]; - - w[ 0] = pws[gid].i[ 0]; - w[ 1] = pws[gid].i[ 1]; - w[ 2] = pws[gid].i[ 2]; - w[ 3] = pws[gid].i[ 3]; - w[ 4] = pws[gid].i[ 4]; - w[ 5] = pws[gid].i[ 5]; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 0; - w[15] = 0; - - const u32 pw_len = pws[gid].pw_len; - - /** - * main - */ - - m14100m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); -} - -__kernel void m14100_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ - /** - * base - */ - - const u64 gid = get_global_id (0); - const u64 lid = get_local_id (0); - const u64 lsz = get_local_size (0); - - /** - * shared - */ - - __local u32 s_SPtrans[8][64]; - __local u32 s_skb[8][64]; - - for (u32 i = lid; i < 64; i += lsz) - { - s_SPtrans[0][i] = c_SPtrans[0][i]; - s_SPtrans[1][i] = c_SPtrans[1][i]; - s_SPtrans[2][i] = c_SPtrans[2][i]; - s_SPtrans[3][i] = c_SPtrans[3][i]; - s_SPtrans[4][i] = c_SPtrans[4][i]; - s_SPtrans[5][i] = c_SPtrans[5][i]; - s_SPtrans[6][i] = c_SPtrans[6][i]; - s_SPtrans[7][i] = c_SPtrans[7][i]; - - s_skb[0][i] = c_skb[0][i]; - s_skb[1][i] = c_skb[1][i]; - s_skb[2][i] = c_skb[2][i]; - s_skb[3][i] = c_skb[3][i]; - s_skb[4][i] = c_skb[4][i]; - s_skb[5][i] = c_skb[5][i]; - s_skb[6][i] = c_skb[6][i]; - s_skb[7][i] = c_skb[7][i]; - } - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - - /** - * base - */ - - u32 w[16]; - - w[ 0] = pws[gid].i[ 0]; - w[ 1] = pws[gid].i[ 1]; - w[ 2] = pws[gid].i[ 2]; - w[ 3] = pws[gid].i[ 3]; - w[ 4] = pws[gid].i[ 4]; - w[ 5] = pws[gid].i[ 5]; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 0; - w[15] = 0; - - const u32 pw_len = pws[gid].pw_len; - - /** - * main - */ - - m14100m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); -} - -__kernel void m14100_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ - /** - * base - */ - - const u64 gid = get_global_id (0); - const u64 lid = get_local_id (0); - const u64 lsz = get_local_size (0); - - /** - * shared - */ - - __local u32 s_SPtrans[8][64]; - __local u32 s_skb[8][64]; - - for (u32 i = lid; i < 64; i += lsz) - { - s_SPtrans[0][i] = c_SPtrans[0][i]; - s_SPtrans[1][i] = c_SPtrans[1][i]; - s_SPtrans[2][i] = c_SPtrans[2][i]; - s_SPtrans[3][i] = c_SPtrans[3][i]; - s_SPtrans[4][i] = c_SPtrans[4][i]; - s_SPtrans[5][i] = c_SPtrans[5][i]; - s_SPtrans[6][i] = c_SPtrans[6][i]; - s_SPtrans[7][i] = c_SPtrans[7][i]; - - s_skb[0][i] = c_skb[0][i]; - s_skb[1][i] = c_skb[1][i]; - s_skb[2][i] = c_skb[2][i]; - s_skb[3][i] = c_skb[3][i]; - s_skb[4][i] = c_skb[4][i]; - s_skb[5][i] = c_skb[5][i]; - s_skb[6][i] = c_skb[6][i]; - s_skb[7][i] = c_skb[7][i]; - } - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - - /** - * base - */ - - u32 w[16]; - - w[ 0] = pws[gid].i[ 0]; - w[ 1] = pws[gid].i[ 1]; - w[ 2] = pws[gid].i[ 2]; - w[ 3] = pws[gid].i[ 3]; - w[ 4] = pws[gid].i[ 4]; - w[ 5] = pws[gid].i[ 5]; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 0; - w[15] = 0; - - const u32 pw_len = pws[gid].pw_len; - - /** - * main - */ - - m14100s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); -} - -__kernel void m14100_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ - /** - * base - */ - - const u64 gid = get_global_id (0); - const u64 lid = get_local_id (0); - const u64 lsz = get_local_size (0); - - /** - * shared - */ - - __local u32 s_SPtrans[8][64]; - __local u32 s_skb[8][64]; - - for (u32 i = lid; i < 64; i += lsz) - { - s_SPtrans[0][i] = c_SPtrans[0][i]; - s_SPtrans[1][i] = c_SPtrans[1][i]; - s_SPtrans[2][i] = c_SPtrans[2][i]; - s_SPtrans[3][i] = c_SPtrans[3][i]; - s_SPtrans[4][i] = c_SPtrans[4][i]; - s_SPtrans[5][i] = c_SPtrans[5][i]; - s_SPtrans[6][i] = c_SPtrans[6][i]; - s_SPtrans[7][i] = c_SPtrans[7][i]; - - s_skb[0][i] = c_skb[0][i]; - s_skb[1][i] = c_skb[1][i]; - s_skb[2][i] = c_skb[2][i]; - s_skb[3][i] = c_skb[3][i]; - s_skb[4][i] = c_skb[4][i]; - s_skb[5][i] = c_skb[5][i]; - s_skb[6][i] = c_skb[6][i]; - s_skb[7][i] = c_skb[7][i]; - } - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - - /** - * base - */ - - u32 w[16]; - - w[ 0] = pws[gid].i[ 0]; - w[ 1] = pws[gid].i[ 1]; - w[ 2] = pws[gid].i[ 2]; - w[ 3] = pws[gid].i[ 3]; - w[ 4] = pws[gid].i[ 4]; - w[ 5] = pws[gid].i[ 5]; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 0; - w[15] = 0; - - const u32 pw_len = pws[gid].pw_len; - - /** - * main - */ - - m14100s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); -} - -__kernel void m14100_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14100_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base