X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm05300_a3.cl;h=a1fe855f239f5f50bee4c592b353633e3114d50d;hb=8702d0e3e155e6a61ca1b0dcdc1ec924d92b8436;hp=73dc951b74c69bfb1f4eab850ee02bd0fd77cda1;hpb=0bf4e3c34a6a799ccc34f403bed70119574ca9c8;p=hashcat.git diff --git a/OpenCL/m05300_a3.cl b/OpenCL/m05300_a3.cl index 73dc951..a1fe855 100644 --- a/OpenCL/m05300_a3.cl +++ b/OpenCL/m05300_a3.cl @@ -1,48 +1,44 @@ /** - * Author......: Jens Steube + * Authors.....: Jens Steube + * Gabriele Gristina + * * License.....: MIT */ #define _MD5_ -#include "include/constants.h" -#include "include/kernel_vendor.h" +#define NEW_SIMD_CODE -#define DGST_R0 0 -#define DGST_R1 3 -#define DGST_R2 2 -#define DGST_R3 1 +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" +#include "inc_simd.cl" -#include "include/kernel_functions.c" -#include "types_ocl.c" -#include "common.c" - -#define COMPARE_S "check_single_comp4.c" -#define COMPARE_M "check_multi_comp4.c" - -static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4]) +void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4]) { - u32 a = digest[0]; - u32 b = digest[1]; - u32 c = digest[2]; - u32 d = digest[3]; - - u32 w0_t = w0[0]; - u32 w1_t = w0[1]; - u32 w2_t = w0[2]; - u32 w3_t = w0[3]; - u32 w4_t = w1[0]; - u32 w5_t = w1[1]; - u32 w6_t = w1[2]; - u32 w7_t = w1[3]; - u32 w8_t = w2[0]; - u32 w9_t = w2[1]; - u32 wa_t = w2[2]; - u32 wb_t = w2[3]; - u32 wc_t = w3[0]; - u32 wd_t = w3[1]; - u32 we_t = w3[2]; - u32 wf_t = w3[3]; + u32x a = digest[0]; + u32x b = digest[1]; + u32x c = digest[2]; + u32x d = digest[3]; + + u32x w0_t = w0[0]; + u32x w1_t = w0[1]; + u32x w2_t = w0[2]; + u32x w3_t = w0[3]; + u32x w4_t = w1[0]; + u32x w5_t = w1[1]; + u32x w6_t = w1[2]; + u32x w7_t = w1[3]; + u32x w8_t = w2[0]; + u32x w9_t = w2[1]; + u32x wa_t = w2[2]; + u32x wb_t = w2[3]; + u32x wc_t = w3[0]; + u32x wd_t = w3[1]; + u32x we_t = w3[2]; + u32x wf_t = w3[3]; MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00); MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01); @@ -118,7 +114,7 @@ static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], co digest[3] += d; } -static void hmac_md5_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4]) +void hmac_md5_pad (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[4], u32x opad[4]) { w0[0] = w0[0] ^ 0x36363636; w0[1] = w0[1] ^ 0x36363636; @@ -169,7 +165,7 @@ static void hmac_md5_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4 md5_transform (w0, w1, w2, w3, opad); } -static void hmac_md5_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4], u32 digest[4]) +void hmac_md5_run (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[4], u32x opad[4], u32x digest[4]) { digest[0] = ipad[0]; digest[1] = ipad[1]; @@ -203,7 +199,7 @@ static void hmac_md5_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4 md5_transform (w0, w1, w2, w3, digest); } -static void m05300m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __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 ikepsk_t *ikepsk_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 s_msg_buf[128]) +void m05300m (__local u32 *w_s, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __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 ikepsk_t *ikepsk_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, __local u32 *s_msg_buf) { /** * modifier @@ -219,101 +215,67 @@ static void m05300m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le const u32 nr_len = ikepsk_bufs[salt_pos].nr_len; const u32 msg_len = ikepsk_bufs[salt_pos].msg_len; - u32 salt_buf0[4]; - - salt_buf0[0] = ikepsk_bufs[salt_pos].nr_buf[ 0]; - salt_buf0[1] = ikepsk_bufs[salt_pos].nr_buf[ 1]; - salt_buf0[2] = ikepsk_bufs[salt_pos].nr_buf[ 2]; - salt_buf0[3] = ikepsk_bufs[salt_pos].nr_buf[ 3]; - - u32 salt_buf1[4]; - - salt_buf1[0] = ikepsk_bufs[salt_pos].nr_buf[ 4]; - salt_buf1[1] = ikepsk_bufs[salt_pos].nr_buf[ 5]; - salt_buf1[2] = ikepsk_bufs[salt_pos].nr_buf[ 6]; - salt_buf1[3] = ikepsk_bufs[salt_pos].nr_buf[ 7]; - - u32 salt_buf2[4]; - - salt_buf2[0] = ikepsk_bufs[salt_pos].nr_buf[ 8]; - salt_buf2[1] = ikepsk_bufs[salt_pos].nr_buf[ 9]; - salt_buf2[2] = ikepsk_bufs[salt_pos].nr_buf[10]; - salt_buf2[3] = ikepsk_bufs[salt_pos].nr_buf[11]; - - u32 salt_buf3[4]; - - salt_buf3[0] = ikepsk_bufs[salt_pos].nr_buf[12]; - salt_buf3[1] = ikepsk_bufs[salt_pos].nr_buf[13]; - salt_buf3[2] = 0; - salt_buf3[3] = 0; - /** * loop */ u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 w0r = bfs_buf[il_pos].i; + const u32x w0r = ix_create_bft (bfs_buf, il_pos); - w0[0] = w0l | w0r; - - /** - * pads - */ + const u32x w0lr = w0l | w0r; - u32 w0_t[4]; + u32x w0_t[4]; + u32x w1_t[4]; + u32x w2_t[4]; + u32x w3_t[4]; - w0_t[0] = w0[0]; + w0_t[0] = w0lr; w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - - u32 w1_t[4]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - - u32 w2_t[4]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - - u32 w3_t[4]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; w3_t[3] = w3[3]; - u32 ipad[4]; - u32 opad[4]; + /** + * pads + */ + + u32x ipad[4]; + u32x opad[4]; hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad); - w0_t[0] = salt_buf0[0]; - w0_t[1] = salt_buf0[1]; - w0_t[2] = salt_buf0[2]; - w0_t[3] = salt_buf0[3]; - w1_t[0] = salt_buf1[0]; - w1_t[1] = salt_buf1[1]; - w1_t[2] = salt_buf1[2]; - w1_t[3] = salt_buf1[3]; - w2_t[0] = salt_buf2[0]; - w2_t[1] = salt_buf2[1]; - w2_t[2] = salt_buf2[2]; - w2_t[3] = salt_buf2[3]; - w3_t[0] = salt_buf3[0]; - w3_t[1] = salt_buf3[1]; + w0_t[0] = w_s[ 0]; + w0_t[1] = w_s[ 1]; + w0_t[2] = w_s[ 2]; + w0_t[3] = w_s[ 3]; + w1_t[0] = w_s[ 4]; + w1_t[1] = w_s[ 5]; + w1_t[2] = w_s[ 6]; + w1_t[3] = w_s[ 7]; + w2_t[0] = w_s[ 8]; + w2_t[1] = w_s[ 9]; + w2_t[2] = w_s[10]; + w2_t[3] = w_s[11]; + w3_t[0] = w_s[12]; + w3_t[1] = w_s[13]; w3_t[2] = (64 + nr_len) * 8; w3_t[3] = 0; - u32 digest[4]; + u32x digest[4]; hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest); @@ -380,16 +342,11 @@ static void m05300m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest); - const u32 r0 = digest[0]; - const u32 r1 = digest[3]; - const u32 r2 = digest[2]; - const u32 r3 = digest[1]; - - #include COMPARE_M + COMPARE_M_SIMD (digest[0], digest[3], digest[2], digest[1]); } } -static void m05300s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __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 ikepsk_t *ikepsk_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 s_msg_buf[128]) +void m05300s (__local u32 *w_s, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __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 ikepsk_t *ikepsk_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, __local u32 *s_msg_buf) { /** * modifier @@ -405,34 +362,6 @@ static void m05300s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le const u32 nr_len = ikepsk_bufs[salt_pos].nr_len; const u32 msg_len = ikepsk_bufs[salt_pos].msg_len; - u32 salt_buf0[4]; - - salt_buf0[0] = ikepsk_bufs[salt_pos].nr_buf[ 0]; - salt_buf0[1] = ikepsk_bufs[salt_pos].nr_buf[ 1]; - salt_buf0[2] = ikepsk_bufs[salt_pos].nr_buf[ 2]; - salt_buf0[3] = ikepsk_bufs[salt_pos].nr_buf[ 3]; - - u32 salt_buf1[4]; - - salt_buf1[0] = ikepsk_bufs[salt_pos].nr_buf[ 4]; - salt_buf1[1] = ikepsk_bufs[salt_pos].nr_buf[ 5]; - salt_buf1[2] = ikepsk_bufs[salt_pos].nr_buf[ 6]; - salt_buf1[3] = ikepsk_bufs[salt_pos].nr_buf[ 7]; - - u32 salt_buf2[4]; - - salt_buf2[0] = ikepsk_bufs[salt_pos].nr_buf[ 8]; - salt_buf2[1] = ikepsk_bufs[salt_pos].nr_buf[ 9]; - salt_buf2[2] = ikepsk_bufs[salt_pos].nr_buf[10]; - salt_buf2[3] = ikepsk_bufs[salt_pos].nr_buf[11]; - - u32 salt_buf3[4]; - - salt_buf3[0] = ikepsk_bufs[salt_pos].nr_buf[12]; - salt_buf3[1] = ikepsk_bufs[salt_pos].nr_buf[13]; - salt_buf3[2] = 0; - salt_buf3[3] = 0; - /** * digest */ @@ -451,67 +380,61 @@ static void m05300s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 w0r = bfs_buf[il_pos].i; - - w0[0] = w0l | w0r; + const u32x w0r = ix_create_bft (bfs_buf, il_pos); - /** - * pads - */ + const u32x w0lr = w0l | w0r; - u32 w0_t[4]; + u32x w0_t[4]; + u32x w1_t[4]; + u32x w2_t[4]; + u32x w3_t[4]; - w0_t[0] = w0[0]; + w0_t[0] = w0lr; w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - - u32 w1_t[4]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - - u32 w2_t[4]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - - u32 w3_t[4]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; w3_t[3] = w3[3]; - u32 ipad[4]; - u32 opad[4]; + /** + * pads + */ + + u32x ipad[4]; + u32x opad[4]; hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad); - w0_t[0] = salt_buf0[0]; - w0_t[1] = salt_buf0[1]; - w0_t[2] = salt_buf0[2]; - w0_t[3] = salt_buf0[3]; - w1_t[0] = salt_buf1[0]; - w1_t[1] = salt_buf1[1]; - w1_t[2] = salt_buf1[2]; - w1_t[3] = salt_buf1[3]; - w2_t[0] = salt_buf2[0]; - w2_t[1] = salt_buf2[1]; - w2_t[2] = salt_buf2[2]; - w2_t[3] = salt_buf2[3]; - w3_t[0] = salt_buf3[0]; - w3_t[1] = salt_buf3[1]; + w0_t[0] = w_s[ 0]; + w0_t[1] = w_s[ 1]; + w0_t[2] = w_s[ 2]; + w0_t[3] = w_s[ 3]; + w1_t[0] = w_s[ 4]; + w1_t[1] = w_s[ 5]; + w1_t[2] = w_s[ 6]; + w1_t[3] = w_s[ 7]; + w2_t[0] = w_s[ 8]; + w2_t[1] = w_s[ 9]; + w2_t[2] = w_s[10]; + w2_t[3] = w_s[11]; + w3_t[0] = w_s[12]; + w3_t[1] = w_s[13]; w3_t[2] = (64 + nr_len) * 8; w3_t[3] = 0; - u32 digest[4]; + u32x digest[4]; hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest); @@ -578,28 +501,47 @@ static void m05300s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest); - const u32 r0 = digest[0]; - const u32 r1 = digest[3]; - const u32 r2 = digest[2]; - const u32 r3 = digest[1]; - - #include COMPARE_S + COMPARE_S_SIMD (digest[0], digest[3], digest[2], digest[1]); } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_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 ikepsk_t *ikepsk_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m05300_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 ikepsk_t *ikepsk_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); /** - * base + * s_msg */ - const u32 gid = get_global_id (0); + __local u32 w_s[16]; + + for (u32 i = lid; i < 16; i += lsz) + { + w_s[i] = ikepsk_bufs[salt_pos].nr_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + __local u32 s_msg_buf[128]; + + for (u32 i = lid; i < 128; i += lsz) + { + s_msg_buf[i] = ikepsk_bufs[salt_pos].msg_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + /** + * base + */ u32 w0[4]; @@ -631,41 +573,50 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_m04 (__glo const u32 pw_len = pws[gid].pw_len; - /** - * s_msg - */ - - __local u32 s_msg_buf[128]; - - const u32 lid2 = lid * 2; - - s_msg_buf[lid2 + 0] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 0]; - s_msg_buf[lid2 + 1] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 1]; - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - /** * main */ - m05300m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_msg_buf); + m05300m (w_s, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_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, s_msg_buf); } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_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 ikepsk_t *ikepsk_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m05300_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 ikepsk_t *ikepsk_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); /** - * base + * s_msg */ - const u32 gid = get_global_id (0); + __local u32 w_s[16]; + + for (u32 i = lid; i < 16; i += lsz) + { + w_s[i] = ikepsk_bufs[salt_pos].nr_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + __local u32 s_msg_buf[128]; + + for (u32 i = lid; i < 128; i += lsz) + { + s_msg_buf[i] = ikepsk_bufs[salt_pos].msg_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + /** + * base + */ u32 w0[4]; @@ -697,41 +648,50 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_m08 (__glo const u32 pw_len = pws[gid].pw_len; - /** - * s_msg - */ - - __local u32 s_msg_buf[128]; - - const u32 lid2 = lid * 2; - - s_msg_buf[lid2 + 0] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 0]; - s_msg_buf[lid2 + 1] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 1]; - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - /** * main */ - m05300m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_msg_buf); + m05300m (w_s, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_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, s_msg_buf); } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_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 ikepsk_t *ikepsk_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m05300_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 ikepsk_t *ikepsk_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); /** - * base + * s_msg */ - const u32 gid = get_global_id (0); + __local u32 w_s[16]; + + for (u32 i = lid; i < 16; i += lsz) + { + w_s[i] = ikepsk_bufs[salt_pos].nr_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + __local u32 s_msg_buf[128]; + + for (u32 i = lid; i < 128; i += lsz) + { + s_msg_buf[i] = ikepsk_bufs[salt_pos].msg_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + /** + * base + */ u32 w0[4]; @@ -763,41 +723,50 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_m16 (__glo const u32 pw_len = pws[gid].pw_len; - /** - * s_msg - */ - - __local u32 s_msg_buf[128]; - - const u32 lid2 = lid * 2; - - s_msg_buf[lid2 + 0] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 0]; - s_msg_buf[lid2 + 1] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 1]; - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - /** * main */ - m05300m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_msg_buf); + m05300m (w_s, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_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, s_msg_buf); } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_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 ikepsk_t *ikepsk_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m05300_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 ikepsk_t *ikepsk_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); /** - * base + * s_msg */ - const u32 gid = get_global_id (0); + __local u32 w_s[16]; + + for (u32 i = lid; i < 16; i += lsz) + { + w_s[i] = ikepsk_bufs[salt_pos].nr_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + __local u32 s_msg_buf[128]; + + for (u32 i = lid; i < 128; i += lsz) + { + s_msg_buf[i] = ikepsk_bufs[salt_pos].msg_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + /** + * base + */ u32 w0[4]; @@ -829,41 +798,50 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_s04 (__glo const u32 pw_len = pws[gid].pw_len; - /** - * s_msg - */ - - __local u32 s_msg_buf[128]; - - const u32 lid2 = lid * 2; - - s_msg_buf[lid2 + 0] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 0]; - s_msg_buf[lid2 + 1] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 1]; - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - /** * main */ - m05300s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_msg_buf); + m05300s (w_s, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_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, s_msg_buf); } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_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 ikepsk_t *ikepsk_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m05300_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 ikepsk_t *ikepsk_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); /** - * base + * s_msg */ - const u32 gid = get_global_id (0); + __local u32 w_s[16]; + + for (u32 i = lid; i < 16; i += lsz) + { + w_s[i] = ikepsk_bufs[salt_pos].nr_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + __local u32 s_msg_buf[128]; + + for (u32 i = lid; i < 128; i += lsz) + { + s_msg_buf[i] = ikepsk_bufs[salt_pos].msg_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + /** + * base + */ u32 w0[4]; @@ -895,41 +873,50 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_s08 (__glo const u32 pw_len = pws[gid].pw_len; - /** - * s_msg - */ - - __local u32 s_msg_buf[128]; - - const u32 lid2 = lid * 2; - - s_msg_buf[lid2 + 0] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 0]; - s_msg_buf[lid2 + 1] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 1]; - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - /** * main */ - m05300s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_msg_buf); + m05300s (w_s, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_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, s_msg_buf); } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_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 ikepsk_t *ikepsk_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m05300_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 ikepsk_t *ikepsk_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); /** - * base + * s_msg */ - const u32 gid = get_global_id (0); + __local u32 w_s[16]; + + for (u32 i = lid; i < 16; i += lsz) + { + w_s[i] = ikepsk_bufs[salt_pos].nr_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + __local u32 s_msg_buf[128]; + + for (u32 i = lid; i < 128; i += lsz) + { + s_msg_buf[i] = ikepsk_bufs[salt_pos].msg_buf[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + /** + * base + */ u32 w0[4]; @@ -961,24 +948,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05300_s16 (__glo const u32 pw_len = pws[gid].pw_len; - /** - * s_msg - */ - - __local u32 s_msg_buf[128]; - - const u32 lid2 = lid * 2; - - s_msg_buf[lid2 + 0] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 0]; - s_msg_buf[lid2 + 1] = ikepsk_bufs[salt_pos].msg_buf[lid2 + 1]; - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - /** * main */ - m05300s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_msg_buf); + m05300s (w_s, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, ikepsk_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, s_msg_buf); }