From: Jens Steube Date: Thu, 4 Feb 2016 21:47:02 +0000 (+0100) Subject: Converted to new SIMD: -m 11000 -a 0 X-Git-Tag: v3.00-beta~398 X-Git-Url: https://www.flypig.org.uk/git/?a=commitdiff_plain;h=1bede41eac6216bd57bf7cdf30965715bba6d144;p=hashcat.git Converted to new SIMD: -m 11000 -a 0 --- diff --git a/OpenCL/m11000_a0.cl b/OpenCL/m11000_a0.cl index a46d26a..9e86fa8 100644 --- a/OpenCL/m11000_a0.cl +++ b/OpenCL/m11000_a0.cl @@ -5,6 +5,8 @@ #define _MD5_ +#define NEW_SIMD_CODE + #include "include/constants.h" #include "include/kernel_vendor.h" @@ -18,9 +20,7 @@ #include "OpenCL/common.c" #include "include/rp_kernel.h" #include "OpenCL/rp.c" - -#define COMPARE_S "OpenCL/check_single_comp4.c" -#define COMPARE_M "OpenCL/check_multi_comp4.c" +#include "OpenCL/simd.c" __kernel void m11000_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_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) { @@ -91,35 +91,18 @@ __kernel void m11000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, * loop */ - for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < rules_cnt; il_pos += VECT_SIZE) { - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; - - w0[0] = pw_buf0[0]; - w0[1] = pw_buf0[1]; - w0[2] = pw_buf0[2]; - w0[3] = pw_buf0[3]; - w1[0] = pw_buf1[0]; - w1[1] = pw_buf1[1]; - w1[2] = pw_buf1[2]; - w1[3] = pw_buf1[3]; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len); - - const u32 pw_salt_len = pw_len + salt_len; - - append_0x80_4x4 (w0, w1, w2, w3, out_len); + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; + + const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); + + append_0x80_2x4 (w0, w1, out_len); + + const u32 pw_salt_len = salt_len + out_len; /** * prepend salt @@ -127,10 +110,10 @@ __kernel void m11000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, // first step fixed 56 bytes of salt - u32 w0_t[4]; - u32 w1_t[4]; - u32 w2_t[4]; - u32 w3_t[4]; + u32x w0_t[4]; + u32x w1_t[4]; + u32x w2_t[4]; + u32x w3_t[4]; w0_t[0] = salt_buf0[0]; w0_t[1] = salt_buf0[1]; @@ -158,10 +141,10 @@ __kernel void m11000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, // first transform - u32 a = MD5M_A; - u32 b = MD5M_B; - u32 c = MD5M_C; - u32 d = MD5M_D; + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00); MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01); @@ -236,10 +219,10 @@ __kernel void m11000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += MD5M_C; d += MD5M_D; - u32 r_a = a; - u32 r_b = b; - u32 r_c = c; - u32 r_d = d; + u32x r_a = a; + u32x r_b = b; + u32x r_c = c; + u32x r_d = d; // 2nd transform @@ -333,12 +316,7 @@ __kernel void m11000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += r_c; d += r_d; - const u32 r0 = a; - const u32 r1 = d; - const u32 r2 = c; - const u32 r3 = b; - - #include COMPARE_M + COMPARE_M_SIMD (a, d, c, b); } } @@ -431,35 +409,18 @@ __kernel void m11000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, * loop */ - for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < rules_cnt; il_pos += VECT_SIZE) { - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; - - w0[0] = pw_buf0[0]; - w0[1] = pw_buf0[1]; - w0[2] = pw_buf0[2]; - w0[3] = pw_buf0[3]; - w1[0] = pw_buf1[0]; - w1[1] = pw_buf1[1]; - w1[2] = pw_buf1[2]; - w1[3] = pw_buf1[3]; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len); - - const u32 pw_salt_len = pw_len + salt_len; - - append_0x80_4x4 (w0, w1, w2, w3, out_len); + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; + + const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); + + append_0x80_2x4 (w0, w1, out_len); + + const u32 pw_salt_len = salt_len + out_len; /** * prepend salt @@ -467,10 +428,10 @@ __kernel void m11000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, // first step fixed 56 bytes of salt - u32 w0_t[4]; - u32 w1_t[4]; - u32 w2_t[4]; - u32 w3_t[4]; + u32x w0_t[4]; + u32x w1_t[4]; + u32x w2_t[4]; + u32x w3_t[4]; w0_t[0] = salt_buf0[0]; w0_t[1] = salt_buf0[1]; @@ -498,10 +459,10 @@ __kernel void m11000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, // first transform - u32 a = MD5M_A; - u32 b = MD5M_B; - u32 c = MD5M_C; - u32 d = MD5M_D; + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00); MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01); @@ -576,10 +537,10 @@ __kernel void m11000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += MD5M_C; d += MD5M_D; - u32 r_a = a; - u32 r_b = b; - u32 r_c = c; - u32 r_d = d; + u32x r_a = a; + u32x r_b = b; + u32x r_c = c; + u32x r_d = d; // 2nd transform @@ -673,12 +634,7 @@ __kernel void m11000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += r_c; d += r_d; - const u32 r0 = a; - const u32 r1 = d; - const u32 r2 = c; - const u32 r3 = b; - - #include COMPARE_S + COMPARE_S_SIMD (a, d, c, b); } }