X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm11700_a0.cl;h=a2dcf742f843c06f4cac6dfdafc1892350ab5531;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=875de1ebe8d7a82690b57f6fa323efcbc1bb17a8;hpb=61744662c0c9a3bd31b4224dcc3198217855e391;p=hashcat.git diff --git a/OpenCL/m11700_a0.cl b/OpenCL/m11700_a0.cl index 875de1e..a2dcf74 100644 --- a/OpenCL/m11700_a0.cl +++ b/OpenCL/m11700_a0.cl @@ -1,38 +1,48 @@ /** - * Author......: Jens Steube + * Authors.....: Jens Steube + * Gabriele Gristina + * magnum + * * License.....: MIT */ #define _GOST2012_256_ -#include "include/constants.h" -#include "include/kernel_vendor.h" +//too much register pressure +//#define NEW_SIMD_CODE -#define DGST_R0 0 -#define DGST_R1 1 -#define DGST_R2 2 -#define DGST_R3 3 - -#include "include/kernel_functions.c" -#include "OpenCL/types_ocl.c" -#include "OpenCL/common.c" -#include "include/rp_gpu.h" -#include "OpenCL/rp.c" - -#define COMPARE_S "OpenCL/check_single_comp4.c" -#define COMPARE_M "OpenCL/check_multi_comp4.c" +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" +#include "inc_rp.h" +#include "inc_rp.cl" +#include "inc_simd.cl" #define INITVAL 0x0101010101010101 -#define SBOG_LPSti64 \ - s_sbob_sl64[0][(t[0] >> (i * 8)) & 0xff] ^ \ - s_sbob_sl64[1][(t[1] >> (i * 8)) & 0xff] ^ \ - s_sbob_sl64[2][(t[2] >> (i * 8)) & 0xff] ^ \ - s_sbob_sl64[3][(t[3] >> (i * 8)) & 0xff] ^ \ - s_sbob_sl64[4][(t[4] >> (i * 8)) & 0xff] ^ \ - s_sbob_sl64[5][(t[5] >> (i * 8)) & 0xff] ^ \ - s_sbob_sl64[6][(t[6] >> (i * 8)) & 0xff] ^ \ - s_sbob_sl64[7][(t[7] >> (i * 8)) & 0xff] +#if VECT_SIZE == 1 +#define BOX(S,n,i) (S)[(n)][(i)] +#elif VECT_SIZE == 2 +#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1]) +#elif VECT_SIZE == 4 +#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3]) +#elif VECT_SIZE == 8 +#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7]) +#elif VECT_SIZE == 16 +#define BOX(S,n,i) (u64x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf]) +#endif + +#define SBOG_LPSti64 \ + BOX (s_sbob_sl64, 0, ((t[0] >> (i * 8)) & 0xff)) ^ \ + BOX (s_sbob_sl64, 1, ((t[1] >> (i * 8)) & 0xff)) ^ \ + BOX (s_sbob_sl64, 2, ((t[2] >> (i * 8)) & 0xff)) ^ \ + BOX (s_sbob_sl64, 3, ((t[3] >> (i * 8)) & 0xff)) ^ \ + BOX (s_sbob_sl64, 4, ((t[4] >> (i * 8)) & 0xff)) ^ \ + BOX (s_sbob_sl64, 5, ((t[5] >> (i * 8)) & 0xff)) ^ \ + BOX (s_sbob_sl64, 6, ((t[6] >> (i * 8)) & 0xff)) ^ \ + BOX (s_sbob_sl64, 7, ((t[7] >> (i * 8)) & 0xff)) // constants @@ -2228,13 +2238,15 @@ __constant u64 sbob_rc64[12][8] = }, }; -static void streebog_g (u64 h[8], const u64 m[8], __local u64 s_sbob_sl64[8][256]) +void streebog_g (u64x h[8], const u64x m[8], __local u64 (*s_sbob_sl64)[256]) { - u64 k[8]; - u64 s[8]; - u64 t[8]; + u64x k[8]; + u64x s[8]; + u64x t[8]; + #ifdef _unroll #pragma unroll + #endif for (int i = 0; i < 8; i++) { t[i] = h[i]; @@ -2245,7 +2257,9 @@ static void streebog_g (u64 h[8], const u64 m[8], __local u64 s_sbob_sl64[8][256 k[i] = SBOG_LPSti64; } + #ifdef _unroll #pragma unroll + #endif for (int i = 0; i < 8; i++) { s[i] = m[i]; @@ -2253,13 +2267,17 @@ static void streebog_g (u64 h[8], const u64 m[8], __local u64 s_sbob_sl64[8][256 for (int r = 0; r < 12; r++) { + #ifdef _unroll #pragma unroll + #endif for (int i = 0; i < 8; i++) { t[i] = s[i] ^ k[i]; } + #ifdef _unroll #pragma unroll + #endif for (int i = 0; i < 8; i++) { s[i] = SBOG_LPSti64; @@ -2270,88 +2288,67 @@ static void streebog_g (u64 h[8], const u64 m[8], __local u64 s_sbob_sl64[8][256 t[i] = k[i] ^ sbob_rc64[r][i]; } + #ifdef _unroll #pragma unroll + #endif for (int i = 0; i < 8; i++) { k[i] = SBOG_LPSti64; } } + #ifdef _unroll #pragma unroll + #endif for (int i = 0; i < 8; i++) { h[i] ^= s[i] ^ k[i] ^ m[i]; } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11700_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) { /** * modifier */ + const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); + const u32 lsz = get_local_size (0); /** * shared lookup table */ - const u32 lid4 = lid * 4; - __local u64 s_sbob_sl64[8][256]; - s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0]; - s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1]; - s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2]; - s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3]; - s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0]; - s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1]; - s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2]; - s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3]; - s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0]; - s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1]; - s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2]; - s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3]; - s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0]; - s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1]; - s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2]; - s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3]; - s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0]; - s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1]; - s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2]; - s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3]; - s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0]; - s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1]; - s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2]; - s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3]; - s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0]; - s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1]; - s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2]; - s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3]; - s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0]; - s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1]; - s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2]; - s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3]; + for (u32 i = lid; i < 256; i += lsz) + { + s_sbob_sl64[0][i] = sbob_sl64[0][i]; + s_sbob_sl64[1][i] = sbob_sl64[1][i]; + s_sbob_sl64[2][i] = sbob_sl64[2][i]; + s_sbob_sl64[3][i] = sbob_sl64[3][i]; + s_sbob_sl64[4][i] = sbob_sl64[4][i]; + s_sbob_sl64[5][i] = sbob_sl64[5][i]; + s_sbob_sl64[6][i] = sbob_sl64[6][i]; + s_sbob_sl64[7][i] = sbob_sl64[7][i]; + } barrier (CLK_LOCAL_MEM_FENCE); + if (gid >= gid_max) return; + /** * base */ - const u32 gid = get_global_id (0); - - if (gid >= gid_max) return; - u32 pw_buf0[4]; + u32 pw_buf1[4]; pw_buf0[0] = pws[gid].i[ 0]; pw_buf0[1] = pws[gid].i[ 1]; pw_buf0[2] = pws[gid].i[ 2]; pw_buf0[3] = pws[gid].i[ 3]; - - u32 pw_buf1[4]; - pw_buf1[0] = pws[gid].i[ 4]; pw_buf1[1] = pws[gid].i[ 5]; pw_buf1[2] = pws[gid].i[ 6]; @@ -2363,18 +2360,31 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo * loop */ - for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - u32 w[16]; - - w[ 0] = pw_buf0[0]; - w[ 1] = pw_buf0[1]; - w[ 2] = pw_buf0[2]; - w[ 3] = pw_buf0[3]; - w[ 4] = pw_buf1[0]; - w[ 5] = pw_buf1[1]; - w[ 6] = pw_buf1[2]; - w[ 7] = pw_buf1[3]; + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; + + const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); + + append_0x01_2x4_VV (w0, w1, out_len); + + /** + * GOST + */ + + u32x w[16]; + + w[ 0] = w0[0]; + w[ 1] = w0[1]; + w[ 2] = w0[2]; + w[ 3] = w0[3]; + w[ 4] = w1[0]; + w[ 5] = w1[1]; + w[ 6] = w1[2]; + w[ 7] = w1[3]; w[ 8] = 0; w[ 9] = 0; w[10] = 0; @@ -2384,15 +2394,11 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo w[14] = 0; w[15] = 0; - const u32 out_len = apply_rules (rules_buf[il_pos].cmds, &w[0], &w[1], pw_len); - - append_0x01_2x4 (&w[0], &w[1], out_len); - /** * reverse message block */ - u64 m[8]; + u64x m[8]; m[0] = hl32_to_64 (w[15], w[14]); m[1] = hl32_to_64 (w[13], w[12]); @@ -2414,7 +2420,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo // state buffer (hash) - u64 h[8]; + u64x h[8]; h[0] = INITVAL; h[1] = INITVAL; @@ -2427,7 +2433,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo streebog_g (h, m, s_sbob_sl64); - u64 z[8]; + u64x z[8]; z[0] = 0; z[1] = 0; @@ -2436,97 +2442,71 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap64 ((u64) (out_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); - const u32 r0 = l32_from_64 (h[0]); - const u32 r1 = h32_from_64 (h[0]); - const u32 r2 = l32_from_64 (h[1]); - const u32 r3 = h32_from_64 (h[1]); + const u32x r0 = l32_from_64 (h[0]); + const u32x r1 = h32_from_64 (h[0]); + const u32x r2 = l32_from_64 (h[1]); + const u32x r3 = h32_from_64 (h[1]); - #include COMPARE_M + COMPARE_M_SIMD (r0, r1, r2, r3); } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11700_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) { } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11700_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) { } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11700_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) { /** * modifier */ + const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); - + const u32 lsz = get_local_size (0); /** * shared lookup table */ - const u32 lid4 = lid * 4; - __local u64 s_sbob_sl64[8][256]; - s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0]; - s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1]; - s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2]; - s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3]; - s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0]; - s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1]; - s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2]; - s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3]; - s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0]; - s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1]; - s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2]; - s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3]; - s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0]; - s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1]; - s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2]; - s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3]; - s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0]; - s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1]; - s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2]; - s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3]; - s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0]; - s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1]; - s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2]; - s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3]; - s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0]; - s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1]; - s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2]; - s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3]; - s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0]; - s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1]; - s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2]; - s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3]; + for (u32 i = lid; i < 256; i += lsz) + { + s_sbob_sl64[0][i] = sbob_sl64[0][i]; + s_sbob_sl64[1][i] = sbob_sl64[1][i]; + s_sbob_sl64[2][i] = sbob_sl64[2][i]; + s_sbob_sl64[3][i] = sbob_sl64[3][i]; + s_sbob_sl64[4][i] = sbob_sl64[4][i]; + s_sbob_sl64[5][i] = sbob_sl64[5][i]; + s_sbob_sl64[6][i] = sbob_sl64[6][i]; + s_sbob_sl64[7][i] = sbob_sl64[7][i]; + } barrier (CLK_LOCAL_MEM_FENCE); + if (gid >= gid_max) return; + /** * base */ - const u32 gid = get_global_id (0); - - if (gid >= gid_max) return; - u32 pw_buf0[4]; + u32 pw_buf1[4]; pw_buf0[0] = pws[gid].i[ 0]; pw_buf0[1] = pws[gid].i[ 1]; pw_buf0[2] = pws[gid].i[ 2]; pw_buf0[3] = pws[gid].i[ 3]; - - u32 pw_buf1[4]; - pw_buf1[0] = pws[gid].i[ 4]; pw_buf1[1] = pws[gid].i[ 5]; pw_buf1[2] = pws[gid].i[ 6]; @@ -2550,18 +2530,31 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo * loop */ - for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - u32 w[16]; - - w[ 0] = pw_buf0[0]; - w[ 1] = pw_buf0[1]; - w[ 2] = pw_buf0[2]; - w[ 3] = pw_buf0[3]; - w[ 4] = pw_buf1[0]; - w[ 5] = pw_buf1[1]; - w[ 6] = pw_buf1[2]; - w[ 7] = pw_buf1[3]; + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; + + const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); + + append_0x01_2x4_VV (w0, w1, out_len); + + /** + * GOST + */ + + u32x w[16]; + + w[ 0] = w0[0]; + w[ 1] = w0[1]; + w[ 2] = w0[2]; + w[ 3] = w0[3]; + w[ 4] = w1[0]; + w[ 5] = w1[1]; + w[ 6] = w1[2]; + w[ 7] = w1[3]; w[ 8] = 0; w[ 9] = 0; w[10] = 0; @@ -2571,15 +2564,11 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo w[14] = 0; w[15] = 0; - const u32 out_len = apply_rules (rules_buf[il_pos].cmds, &w[0], &w[1], pw_len); - - append_0x01_2x4 (&w[0], &w[1], out_len); - /** * reverse message block */ - u64 m[8]; + u64x m[8]; m[0] = hl32_to_64 (w[15], w[14]); m[1] = hl32_to_64 (w[13], w[12]); @@ -2601,7 +2590,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo // state buffer (hash) - u64 h[8]; + u64x h[8]; h[0] = INITVAL; h[1] = INITVAL; @@ -2614,7 +2603,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo streebog_g (h, m, s_sbob_sl64); - u64 z[8]; + u64x z[8]; z[0] = 0; z[1] = 0; @@ -2623,24 +2612,24 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap64 ((u64) (out_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); - const u32 r0 = l32_from_64 (h[0]); - const u32 r1 = h32_from_64 (h[0]); - const u32 r2 = l32_from_64 (h[1]); - const u32 r3 = h32_from_64 (h[1]); + const u32x r0 = l32_from_64 (h[0]); + const u32x r1 = h32_from_64 (h[0]); + const u32x r2 = l32_from_64 (h[1]); + const u32x r3 = h32_from_64 (h[1]); - #include COMPARE_S + COMPARE_S_SIMD (r0, r1, r2, r3); } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11700_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) { } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11700_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) { }