X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm11400_a3.cl;h=f7fee1c31078b0c4cf3bb469206460962bd65cc8;hb=51e1d11b2d6dff4ae331cff6098169455733c412;hp=87d9511ee9cefc388fcdc6da41df6a839b314b8a;hpb=a5be8a75ed3ae5bb297f2551133b7a67af2991fe;p=hashcat.git diff --git a/OpenCL/m11400_a3.cl b/OpenCL/m11400_a3.cl index 87d9511..f7fee1c 100644 --- a/OpenCL/m11400_a3.cl +++ b/OpenCL/m11400_a3.cl @@ -8,20 +8,15 @@ #define _MD5_ -#define NEW_SIMD_CODE +//incompatible because of brances +//#define NEW_SIMD_CODE -#include "include/constants.h" -#include "include/kernel_vendor.h" - -#define DGST_R0 0 -#define DGST_R1 3 -#define DGST_R2 2 -#define DGST_R3 1 - -#include "include/kernel_functions.c" -#include "OpenCL/types_ocl.c" -#include "OpenCL/common.c" -#include "OpenCL/simd.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_simd.cl" #if VECT_SIZE == 1 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)]) @@ -35,7 +30,7 @@ #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7], l_bin2asc[(i).s8], l_bin2asc[(i).s9], l_bin2asc[(i).sa], l_bin2asc[(i).sb], l_bin2asc[(i).sc], l_bin2asc[(i).sd], l_bin2asc[(i).se], l_bin2asc[(i).sf]) #endif -static u32 memcat32 (u32x block0[16], u32x block1[16], const u32 block_len, const u32x append0[4], const u32x append1[4], const u32x append2[4], const u32x append3[4], const u32 append_len) +u32 memcat32 (u32x block0[16], u32x block1[16], const u32 block_len, const u32x append0[4], const u32x append1[4], const u32x append2[4], const u32x append3[4], const u32 append_len) { const u32 mod = block_len & 3; const u32 div = block_len / 4; @@ -756,7 +751,7 @@ static u32 memcat32 (u32x block0[16], u32x block1[16], const u32 block_len, cons return new_len; } -static void m11400m_0_0 (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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc) +void m11400m_0_0 (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 sip_t *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, __local u32 *l_bin2asc) { /** * modifier @@ -765,18 +760,6 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); - /** - * digest - */ - - const u32 search[4] = - { - digests_buf[digests_offset].digest_buf[DGST_R0], - digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] - }; - /** * salt */ @@ -786,6 +769,7 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 pw_salt_len = salt_len + pw_len; u32 salt_buf0[16]; + u32 salt_buf1[16]; salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0]; salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1]; @@ -803,9 +787,6 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13]; salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14]; salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15]; - - u32 salt_buf1[16]; - salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16]; salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17]; salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18]; @@ -830,6 +811,7 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 esalt_len = esalt_bufs[salt_pos].esalt_len; u32 esalt_buf0[16]; + u32 esalt_buf1[16]; esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0]; esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1]; @@ -847,9 +829,6 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13]; esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14]; esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15]; - - u32 esalt_buf1[16]; - esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16]; esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17]; esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18]; @@ -875,9 +854,9 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = w0r_create_bft (bfs_buf, il_pos); + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0lr = w0l | w0r; @@ -888,6 +867,7 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p // append the pass to the salt u32x block0[16]; + u32x block1[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -905,9 +885,6 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p block0[13] = salt_buf0[13]; block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - - u32x block1[16]; - block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; block1[ 2] = salt_buf1[ 2]; @@ -934,17 +911,14 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; @@ -956,17 +930,14 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - w3_t[0] = block0[12]; w3_t[1] = block0[13]; w3_t[2] = pw_salt_len * 8; @@ -1073,12 +1044,10 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p | uint_to_hex_lower8 ((d >> 8) & 255) << 16; w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0 | uint_to_hex_lower8 ((d >> 24) & 255) << 16; - w2_t[0] = esalt_buf0[0]; w2_t[1] = esalt_buf0[1]; w2_t[2] = esalt_buf0[2]; w2_t[3] = esalt_buf0[3]; - w3_t[0] = esalt_buf0[4]; w3_t[1] = esalt_buf0[5]; w3_t[2] = esalt_buf0[6]; @@ -1176,17 +1145,14 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf0[ 9]; w0_t[2] = esalt_buf0[10]; w0_t[3] = esalt_buf0[11]; - w1_t[0] = esalt_buf0[12]; w1_t[1] = esalt_buf0[13]; w1_t[2] = esalt_buf0[14]; w1_t[3] = esalt_buf0[15]; - w2_t[0] = esalt_buf1[ 0]; w2_t[1] = esalt_buf1[ 1]; w2_t[2] = esalt_buf1[ 2]; w2_t[3] = esalt_buf1[ 3]; - w3_t[0] = esalt_buf1[ 4]; w3_t[1] = esalt_buf1[ 5]; w3_t[2] = digest_esalt_len * 8; @@ -1265,11 +1231,11 @@ static void m11400m_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p c += r_c; d += r_d; - COMPARE_S_SIMD (a, d, c, b); + COMPARE_M_SIMD (a, d, c, b); } } -static void m11400m_0_1 (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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc) +void m11400m_0_1 (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 sip_t *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, __local u32 *l_bin2asc) { /** * modifier @@ -1278,18 +1244,6 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); - /** - * digest - */ - - const u32 search[4] = - { - digests_buf[digests_offset].digest_buf[DGST_R0], - digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] - }; - /** * salt */ @@ -1299,6 +1253,7 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 pw_salt_len = salt_len + pw_len; u32 salt_buf0[16]; + u32 salt_buf1[16]; salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0]; salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1]; @@ -1316,9 +1271,6 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13]; salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14]; salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15]; - - u32 salt_buf1[16]; - salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16]; salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17]; salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18]; @@ -1343,6 +1295,8 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 esalt_len = esalt_bufs[salt_pos].esalt_len; u32 esalt_buf0[16]; + u32 esalt_buf1[16]; + u32 esalt_buf2[16]; esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0]; esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1]; @@ -1360,9 +1314,6 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13]; esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14]; esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15]; - - u32 esalt_buf1[16]; - esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16]; esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17]; esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18]; @@ -1379,9 +1330,6 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf1[13] = esalt_bufs[salt_pos].esalt_buf[29]; esalt_buf1[14] = esalt_bufs[salt_pos].esalt_buf[30]; esalt_buf1[15] = esalt_bufs[salt_pos].esalt_buf[31]; - - u32 esalt_buf2[16]; - esalt_buf2[ 0] = esalt_bufs[salt_pos].esalt_buf[32]; esalt_buf2[ 1] = esalt_bufs[salt_pos].esalt_buf[33]; esalt_buf2[ 2] = esalt_bufs[salt_pos].esalt_buf[34]; @@ -1407,9 +1355,9 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = w0r_create_bft (bfs_buf, il_pos); + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0lr = w0l | w0r; @@ -1420,6 +1368,7 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p // append the pass to the salt u32x block0[16]; + u32x block1[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -1437,9 +1386,6 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p block0[13] = salt_buf0[13]; block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - - u32x block1[16]; - block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; block1[ 2] = salt_buf1[ 2]; @@ -1466,17 +1412,14 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; @@ -1488,17 +1431,14 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - w3_t[0] = block0[12]; w3_t[1] = block0[13]; w3_t[2] = pw_salt_len * 8; @@ -1605,12 +1545,10 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p | uint_to_hex_lower8 ((d >> 8) & 255) << 16; w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0 | uint_to_hex_lower8 ((d >> 24) & 255) << 16; - w2_t[0] = esalt_buf0[0]; w2_t[1] = esalt_buf0[1]; w2_t[2] = esalt_buf0[2]; w2_t[3] = esalt_buf0[3]; - w3_t[0] = esalt_buf0[4]; w3_t[1] = esalt_buf0[5]; w3_t[2] = esalt_buf0[6]; @@ -1708,17 +1646,14 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf0[ 9]; w0_t[2] = esalt_buf0[10]; w0_t[3] = esalt_buf0[11]; - w1_t[0] = esalt_buf0[12]; w1_t[1] = esalt_buf0[13]; w1_t[2] = esalt_buf0[14]; w1_t[3] = esalt_buf0[15]; - w2_t[0] = esalt_buf1[ 0]; w2_t[1] = esalt_buf1[ 1]; w2_t[2] = esalt_buf1[ 2]; w2_t[3] = esalt_buf1[ 3]; - w3_t[0] = esalt_buf1[ 4]; w3_t[1] = esalt_buf1[ 5]; w3_t[2] = esalt_buf1[ 6]; @@ -1808,17 +1743,14 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf1[ 9]; w0_t[2] = esalt_buf1[10]; w0_t[3] = esalt_buf1[11]; - w1_t[0] = esalt_buf1[12]; w1_t[1] = esalt_buf1[13]; w1_t[2] = esalt_buf1[14]; w1_t[3] = esalt_buf1[15]; - w2_t[0] = esalt_buf2[ 0]; w2_t[1] = esalt_buf2[ 1]; w2_t[2] = esalt_buf2[ 2]; w2_t[3] = esalt_buf2[ 3]; - w3_t[0] = esalt_buf2[ 4]; w3_t[1] = esalt_buf2[ 5]; w3_t[2] = digest_esalt_len * 8; @@ -1897,11 +1829,11 @@ static void m11400m_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p c += r_c; d += r_d; - COMPARE_S_SIMD (a, d, c, b); + COMPARE_M_SIMD (a, d, c, b); } } -static void m11400m_1_0 (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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc) +void m11400m_1_0 (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 sip_t *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, __local u32 *l_bin2asc) { /** * modifier @@ -1910,18 +1842,6 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); - /** - * digest - */ - - const u32 search[4] = - { - digests_buf[digests_offset].digest_buf[DGST_R0], - digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] - }; - /** * salt */ @@ -1931,6 +1851,7 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 pw_salt_len = salt_len + pw_len; u32 salt_buf0[16]; + u32 salt_buf1[16]; salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0]; salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1]; @@ -1948,9 +1869,6 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13]; salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14]; salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15]; - - u32 salt_buf1[16]; - salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16]; salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17]; salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18]; @@ -1975,6 +1893,7 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 esalt_len = esalt_bufs[salt_pos].esalt_len; u32 esalt_buf0[16]; + u32 esalt_buf1[16]; esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0]; esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1]; @@ -1992,9 +1911,6 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13]; esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14]; esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15]; - - u32 esalt_buf1[16]; - esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16]; esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17]; esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18]; @@ -2020,9 +1936,9 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = w0r_create_bft (bfs_buf, il_pos); + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0lr = w0l | w0r; @@ -2033,6 +1949,7 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p // append the pass to the salt u32x block0[16]; + u32x block1[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -2050,9 +1967,6 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p block0[13] = salt_buf0[13]; block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - - u32x block1[16]; - block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; block1[ 2] = salt_buf1[ 2]; @@ -2079,17 +1993,14 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; @@ -2101,17 +2012,14 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - w3_t[0] = block0[12]; w3_t[1] = block0[13]; w3_t[2] = block0[14]; @@ -2206,17 +2114,14 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block1[ 1]; w0_t[2] = block1[ 2]; w0_t[3] = block1[ 3]; - w1_t[0] = block1[ 4]; w1_t[1] = block1[ 5]; w1_t[2] = block1[ 6]; w1_t[3] = block1[ 7]; - w2_t[0] = block1[ 8]; w2_t[1] = block1[ 9]; w2_t[2] = block1[10]; w2_t[3] = block1[11]; - w3_t[0] = block1[12]; w3_t[1] = block1[13]; w3_t[2] = pw_salt_len * 8; @@ -2316,12 +2221,10 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p | uint_to_hex_lower8 ((d >> 8) & 255) << 16; w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0 | uint_to_hex_lower8 ((d >> 24) & 255) << 16; - w2_t[0] = esalt_buf0[0]; w2_t[1] = esalt_buf0[1]; w2_t[2] = esalt_buf0[2]; w2_t[3] = esalt_buf0[3]; - w3_t[0] = esalt_buf0[4]; w3_t[1] = esalt_buf0[5]; w3_t[2] = esalt_buf0[6]; @@ -2419,17 +2322,14 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf0[ 9]; w0_t[2] = esalt_buf0[10]; w0_t[3] = esalt_buf0[11]; - w1_t[0] = esalt_buf0[12]; w1_t[1] = esalt_buf0[13]; w1_t[2] = esalt_buf0[14]; w1_t[3] = esalt_buf0[15]; - w2_t[0] = esalt_buf1[ 0]; w2_t[1] = esalt_buf1[ 1]; w2_t[2] = esalt_buf1[ 2]; w2_t[3] = esalt_buf1[ 3]; - w3_t[0] = esalt_buf1[ 4]; w3_t[1] = esalt_buf1[ 5]; w3_t[2] = digest_esalt_len * 8; @@ -2508,11 +2408,11 @@ static void m11400m_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p c += r_c; d += r_d; - COMPARE_S_SIMD (a, d, c, b); + COMPARE_M_SIMD (a, d, c, b); } } -static void m11400m_1_1 (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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc) +void m11400m_1_1 (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 sip_t *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, __local u32 *l_bin2asc) { /** * modifier @@ -2521,18 +2421,6 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); - /** - * digest - */ - - const u32 search[4] = - { - digests_buf[digests_offset].digest_buf[DGST_R0], - digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] - }; - /** * salt */ @@ -2542,6 +2430,7 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 pw_salt_len = salt_len + pw_len; u32 salt_buf0[16]; + u32 salt_buf1[16]; salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0]; salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1]; @@ -2559,9 +2448,6 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13]; salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14]; salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15]; - - u32 salt_buf1[16]; - salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16]; salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17]; salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18]; @@ -2586,6 +2472,8 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 esalt_len = esalt_bufs[salt_pos].esalt_len; u32 esalt_buf0[16]; + u32 esalt_buf1[16]; + u32 esalt_buf2[16]; esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0]; esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1]; @@ -2603,9 +2491,6 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13]; esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14]; esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15]; - - u32 esalt_buf1[16]; - esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16]; esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17]; esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18]; @@ -2622,9 +2507,6 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf1[13] = esalt_bufs[salt_pos].esalt_buf[29]; esalt_buf1[14] = esalt_bufs[salt_pos].esalt_buf[30]; esalt_buf1[15] = esalt_bufs[salt_pos].esalt_buf[31]; - - u32 esalt_buf2[16]; - esalt_buf2[ 0] = esalt_bufs[salt_pos].esalt_buf[32]; esalt_buf2[ 1] = esalt_bufs[salt_pos].esalt_buf[33]; esalt_buf2[ 2] = esalt_bufs[salt_pos].esalt_buf[34]; @@ -2650,9 +2532,9 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = w0r_create_bft (bfs_buf, il_pos); + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0lr = w0l | w0r; @@ -2663,6 +2545,7 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p // append the pass to the salt u32x block0[16]; + u32x block1[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -2680,9 +2563,6 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p block0[13] = salt_buf0[13]; block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - - u32 block1[16]; - block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; block1[ 2] = salt_buf1[ 2]; @@ -2709,17 +2589,14 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; @@ -2731,17 +2608,14 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - w3_t[0] = block0[12]; w3_t[1] = block0[13]; w3_t[2] = block0[14]; @@ -2836,17 +2710,14 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block1[ 1]; w0_t[2] = block1[ 2]; w0_t[3] = block1[ 3]; - w1_t[0] = block1[ 4]; w1_t[1] = block1[ 5]; w1_t[2] = block1[ 6]; w1_t[3] = block1[ 7]; - w2_t[0] = block1[ 8]; w2_t[1] = block1[ 9]; w2_t[2] = block1[10]; w2_t[3] = block1[11]; - w3_t[0] = block1[12]; w3_t[1] = block1[13]; w3_t[2] = pw_salt_len * 8; @@ -2946,12 +2817,10 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p | uint_to_hex_lower8 ((d >> 8) & 255) << 16; w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0 | uint_to_hex_lower8 ((d >> 24) & 255) << 16; - w2_t[0] = esalt_buf0[0]; w2_t[1] = esalt_buf0[1]; w2_t[2] = esalt_buf0[2]; w2_t[3] = esalt_buf0[3]; - w3_t[0] = esalt_buf0[4]; w3_t[1] = esalt_buf0[5]; w3_t[2] = esalt_buf0[6]; @@ -3049,17 +2918,14 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf0[ 9]; w0_t[2] = esalt_buf0[10]; w0_t[3] = esalt_buf0[11]; - w1_t[0] = esalt_buf0[12]; w1_t[1] = esalt_buf0[13]; w1_t[2] = esalt_buf0[14]; w1_t[3] = esalt_buf0[15]; - w2_t[0] = esalt_buf1[ 0]; w2_t[1] = esalt_buf1[ 1]; w2_t[2] = esalt_buf1[ 2]; w2_t[3] = esalt_buf1[ 3]; - w3_t[0] = esalt_buf1[ 4]; w3_t[1] = esalt_buf1[ 5]; w3_t[2] = esalt_buf1[ 6]; @@ -3149,17 +3015,14 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf1[ 9]; w0_t[2] = esalt_buf1[10]; w0_t[3] = esalt_buf1[11]; - w1_t[0] = esalt_buf1[12]; w1_t[1] = esalt_buf1[13]; w1_t[2] = esalt_buf1[14]; w1_t[3] = esalt_buf1[15]; - w2_t[0] = esalt_buf2[ 0]; w2_t[1] = esalt_buf2[ 1]; w2_t[2] = esalt_buf2[ 2]; w2_t[3] = esalt_buf2[ 3]; - w3_t[0] = esalt_buf2[ 4]; w3_t[1] = esalt_buf2[ 5]; w3_t[2] = digest_esalt_len * 8; @@ -3238,11 +3101,11 @@ static void m11400m_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p c += r_c; d += r_d; - COMPARE_S_SIMD (a, d, c, b); + COMPARE_M_SIMD (a, d, c, b); } } -static void m11400s_0_0 (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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc) +void m11400s_0_0 (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 sip_t *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, __local u32 *l_bin2asc) { /** * modifier @@ -3272,6 +3135,7 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 pw_salt_len = salt_len + pw_len; u32 salt_buf0[16]; + u32 salt_buf1[16]; salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0]; salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1]; @@ -3289,9 +3153,6 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13]; salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14]; salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15]; - - u32 salt_buf1[16]; - salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16]; salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17]; salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18]; @@ -3316,6 +3177,7 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 esalt_len = esalt_bufs[salt_pos].esalt_len; u32 esalt_buf0[16]; + u32 esalt_buf1[16]; esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0]; esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1]; @@ -3333,9 +3195,6 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13]; esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14]; esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15]; - - u32 esalt_buf1[16]; - esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16]; esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17]; esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18]; @@ -3361,9 +3220,9 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = w0r_create_bft (bfs_buf, il_pos); + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0lr = w0l | w0r; @@ -3374,6 +3233,7 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p // append the pass to the salt u32x block0[16]; + u32x block1[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -3391,9 +3251,6 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p block0[13] = salt_buf0[13]; block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - - u32x block1[16]; - block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; block1[ 2] = salt_buf1[ 2]; @@ -3420,17 +3277,14 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; @@ -3442,17 +3296,14 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - w3_t[0] = block0[12]; w3_t[1] = block0[13]; w3_t[2] = pw_salt_len * 8; @@ -3559,12 +3410,10 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p | uint_to_hex_lower8 ((d >> 8) & 255) << 16; w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0 | uint_to_hex_lower8 ((d >> 24) & 255) << 16; - w2_t[0] = esalt_buf0[0]; w2_t[1] = esalt_buf0[1]; w2_t[2] = esalt_buf0[2]; w2_t[3] = esalt_buf0[3]; - w3_t[0] = esalt_buf0[4]; w3_t[1] = esalt_buf0[5]; w3_t[2] = esalt_buf0[6]; @@ -3662,17 +3511,14 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf0[ 9]; w0_t[2] = esalt_buf0[10]; w0_t[3] = esalt_buf0[11]; - w1_t[0] = esalt_buf0[12]; w1_t[1] = esalt_buf0[13]; w1_t[2] = esalt_buf0[14]; w1_t[3] = esalt_buf0[15]; - w2_t[0] = esalt_buf1[ 0]; w2_t[1] = esalt_buf1[ 1]; w2_t[2] = esalt_buf1[ 2]; w2_t[3] = esalt_buf1[ 3]; - w3_t[0] = esalt_buf1[ 4]; w3_t[1] = esalt_buf1[ 5]; w3_t[2] = digest_esalt_len * 8; @@ -3755,7 +3601,7 @@ static void m11400s_0_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p } } -static void m11400s_0_1 (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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc) +void m11400s_0_1 (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 sip_t *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, __local u32 *l_bin2asc) { /** * modifier @@ -3785,6 +3631,7 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 pw_salt_len = salt_len + pw_len; u32 salt_buf0[16]; + u32 salt_buf1[16]; salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0]; salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1]; @@ -3802,9 +3649,6 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13]; salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14]; salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15]; - - u32 salt_buf1[16]; - salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16]; salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17]; salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18]; @@ -3829,6 +3673,8 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 esalt_len = esalt_bufs[salt_pos].esalt_len; u32 esalt_buf0[16]; + u32 esalt_buf1[16]; + u32 esalt_buf2[16]; esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0]; esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1]; @@ -3846,9 +3692,6 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13]; esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14]; esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15]; - - u32 esalt_buf1[16]; - esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16]; esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17]; esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18]; @@ -3865,9 +3708,6 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf1[13] = esalt_bufs[salt_pos].esalt_buf[29]; esalt_buf1[14] = esalt_bufs[salt_pos].esalt_buf[30]; esalt_buf1[15] = esalt_bufs[salt_pos].esalt_buf[31]; - - u32 esalt_buf2[16]; - esalt_buf2[ 0] = esalt_bufs[salt_pos].esalt_buf[32]; esalt_buf2[ 1] = esalt_bufs[salt_pos].esalt_buf[33]; esalt_buf2[ 2] = esalt_bufs[salt_pos].esalt_buf[34]; @@ -3893,9 +3733,9 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = w0r_create_bft (bfs_buf, il_pos); + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0lr = w0l | w0r; @@ -3906,6 +3746,7 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p // append the pass to the salt u32x block0[16]; + u32x block1[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -3923,9 +3764,6 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p block0[13] = salt_buf0[13]; block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - - u32x block1[16]; - block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; block1[ 2] = salt_buf1[ 2]; @@ -3952,17 +3790,14 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; @@ -3974,17 +3809,14 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - w3_t[0] = block0[12]; w3_t[1] = block0[13]; w3_t[2] = pw_salt_len * 8; @@ -4091,12 +3923,10 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p | uint_to_hex_lower8 ((d >> 8) & 255) << 16; w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0 | uint_to_hex_lower8 ((d >> 24) & 255) << 16; - w2_t[0] = esalt_buf0[0]; w2_t[1] = esalt_buf0[1]; w2_t[2] = esalt_buf0[2]; w2_t[3] = esalt_buf0[3]; - w3_t[0] = esalt_buf0[4]; w3_t[1] = esalt_buf0[5]; w3_t[2] = esalt_buf0[6]; @@ -4194,17 +4024,14 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf0[ 9]; w0_t[2] = esalt_buf0[10]; w0_t[3] = esalt_buf0[11]; - w1_t[0] = esalt_buf0[12]; w1_t[1] = esalt_buf0[13]; w1_t[2] = esalt_buf0[14]; w1_t[3] = esalt_buf0[15]; - w2_t[0] = esalt_buf1[ 0]; w2_t[1] = esalt_buf1[ 1]; w2_t[2] = esalt_buf1[ 2]; w2_t[3] = esalt_buf1[ 3]; - w3_t[0] = esalt_buf1[ 4]; w3_t[1] = esalt_buf1[ 5]; w3_t[2] = esalt_buf1[ 6]; @@ -4294,17 +4121,14 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf1[ 9]; w0_t[2] = esalt_buf1[10]; w0_t[3] = esalt_buf1[11]; - w1_t[0] = esalt_buf1[12]; w1_t[1] = esalt_buf1[13]; w1_t[2] = esalt_buf1[14]; w1_t[3] = esalt_buf1[15]; - w2_t[0] = esalt_buf2[ 0]; w2_t[1] = esalt_buf2[ 1]; w2_t[2] = esalt_buf2[ 2]; w2_t[3] = esalt_buf2[ 3]; - w3_t[0] = esalt_buf2[ 4]; w3_t[1] = esalt_buf2[ 5]; w3_t[2] = digest_esalt_len * 8; @@ -4387,7 +4211,7 @@ static void m11400s_0_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p } } -static void m11400s_1_0 (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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc) +void m11400s_1_0 (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 sip_t *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, __local u32 *l_bin2asc) { /** * modifier @@ -4417,6 +4241,7 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 pw_salt_len = salt_len + pw_len; u32 salt_buf0[16]; + u32 salt_buf1[16]; salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0]; salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1]; @@ -4434,9 +4259,6 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13]; salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14]; salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15]; - - u32 salt_buf1[16]; - salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16]; salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17]; salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18]; @@ -4461,6 +4283,7 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 esalt_len = esalt_bufs[salt_pos].esalt_len; u32 esalt_buf0[16]; + u32 esalt_buf1[16]; esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0]; esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1]; @@ -4478,9 +4301,6 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13]; esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14]; esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15]; - - u32 esalt_buf1[16]; - esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16]; esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17]; esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18]; @@ -4506,9 +4326,9 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = w0r_create_bft (bfs_buf, il_pos); + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0lr = w0l | w0r; @@ -4519,6 +4339,7 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p // append the pass to the salt u32x block0[16]; + u32x block1[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -4536,9 +4357,6 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p block0[13] = salt_buf0[13]; block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - - u32x block1[16]; - block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; block1[ 2] = salt_buf1[ 2]; @@ -4565,17 +4383,14 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; @@ -4587,17 +4402,14 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - w3_t[0] = block0[12]; w3_t[1] = block0[13]; w3_t[2] = block0[14]; @@ -4692,17 +4504,14 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block1[ 1]; w0_t[2] = block1[ 2]; w0_t[3] = block1[ 3]; - w1_t[0] = block1[ 4]; w1_t[1] = block1[ 5]; w1_t[2] = block1[ 6]; w1_t[3] = block1[ 7]; - w2_t[0] = block1[ 8]; w2_t[1] = block1[ 9]; w2_t[2] = block1[10]; w2_t[3] = block1[11]; - w3_t[0] = block1[12]; w3_t[1] = block1[13]; w3_t[2] = pw_salt_len * 8; @@ -4802,12 +4611,10 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p | uint_to_hex_lower8 ((d >> 8) & 255) << 16; w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0 | uint_to_hex_lower8 ((d >> 24) & 255) << 16; - w2_t[0] = esalt_buf0[0]; w2_t[1] = esalt_buf0[1]; w2_t[2] = esalt_buf0[2]; w2_t[3] = esalt_buf0[3]; - w3_t[0] = esalt_buf0[4]; w3_t[1] = esalt_buf0[5]; w3_t[2] = esalt_buf0[6]; @@ -4905,17 +4712,14 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf0[ 9]; w0_t[2] = esalt_buf0[10]; w0_t[3] = esalt_buf0[11]; - w1_t[0] = esalt_buf0[12]; w1_t[1] = esalt_buf0[13]; w1_t[2] = esalt_buf0[14]; w1_t[3] = esalt_buf0[15]; - w2_t[0] = esalt_buf1[ 0]; w2_t[1] = esalt_buf1[ 1]; w2_t[2] = esalt_buf1[ 2]; w2_t[3] = esalt_buf1[ 3]; - w3_t[0] = esalt_buf1[ 4]; w3_t[1] = esalt_buf1[ 5]; w3_t[2] = digest_esalt_len * 8; @@ -4998,7 +4802,7 @@ static void m11400s_1_0 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p } } -static void m11400s_1_1 (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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc) +void m11400s_1_1 (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 sip_t *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, __local u32 *l_bin2asc) { /** * modifier @@ -5028,6 +4832,7 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 pw_salt_len = salt_len + pw_len; u32 salt_buf0[16]; + u32 salt_buf1[16]; salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0]; salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1]; @@ -5045,9 +4850,6 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13]; salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14]; salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15]; - - u32 salt_buf1[16]; - salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16]; salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17]; salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18]; @@ -5072,6 +4874,8 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p const u32 esalt_len = esalt_bufs[salt_pos].esalt_len; u32 esalt_buf0[16]; + u32 esalt_buf1[16]; + u32 esalt_buf2[16]; esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0]; esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1]; @@ -5089,9 +4893,6 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13]; esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14]; esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15]; - - u32 esalt_buf1[16]; - esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16]; esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17]; esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18]; @@ -5108,9 +4909,6 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p esalt_buf1[13] = esalt_bufs[salt_pos].esalt_buf[29]; esalt_buf1[14] = esalt_bufs[salt_pos].esalt_buf[30]; esalt_buf1[15] = esalt_bufs[salt_pos].esalt_buf[31]; - - u32 esalt_buf2[16]; - esalt_buf2[ 0] = esalt_bufs[salt_pos].esalt_buf[32]; esalt_buf2[ 1] = esalt_bufs[salt_pos].esalt_buf[33]; esalt_buf2[ 2] = esalt_bufs[salt_pos].esalt_buf[34]; @@ -5136,9 +4934,9 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p u32 w0l = w0[0]; - for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = w0r_create_bft (bfs_buf, il_pos); + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0lr = w0l | w0r; @@ -5149,6 +4947,7 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p // append the pass to the salt u32x block0[16]; + u32x block1[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -5166,9 +4965,6 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p block0[13] = salt_buf0[13]; block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - - u32x block1[16]; - block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; block1[ 2] = salt_buf1[ 2]; @@ -5195,17 +4991,14 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = w0[1]; w0_t[2] = w0[2]; w0_t[3] = w0[3]; - w1_t[0] = w1[0]; w1_t[1] = w1[1]; w1_t[2] = w1[2]; w1_t[3] = w1[3]; - w2_t[0] = w2[0]; w2_t[1] = w2[1]; w2_t[2] = w2[2]; w2_t[3] = w2[3]; - w3_t[0] = w3[0]; w3_t[1] = w3[1]; w3_t[2] = w3[2]; @@ -5217,17 +5010,14 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - w3_t[0] = block0[12]; w3_t[1] = block0[13]; w3_t[2] = block0[14]; @@ -5322,17 +5112,14 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = block1[ 1]; w0_t[2] = block1[ 2]; w0_t[3] = block1[ 3]; - w1_t[0] = block1[ 4]; w1_t[1] = block1[ 5]; w1_t[2] = block1[ 6]; w1_t[3] = block1[ 7]; - w2_t[0] = block1[ 8]; w2_t[1] = block1[ 9]; w2_t[2] = block1[10]; w2_t[3] = block1[11]; - w3_t[0] = block1[12]; w3_t[1] = block1[13]; w3_t[2] = pw_salt_len * 8; @@ -5432,12 +5219,10 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p | uint_to_hex_lower8 ((d >> 8) & 255) << 16; w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0 | uint_to_hex_lower8 ((d >> 24) & 255) << 16; - w2_t[0] = esalt_buf0[0]; w2_t[1] = esalt_buf0[1]; w2_t[2] = esalt_buf0[2]; w2_t[3] = esalt_buf0[3]; - w3_t[0] = esalt_buf0[4]; w3_t[1] = esalt_buf0[5]; w3_t[2] = esalt_buf0[6]; @@ -5535,17 +5320,14 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf0[ 9]; w0_t[2] = esalt_buf0[10]; w0_t[3] = esalt_buf0[11]; - w1_t[0] = esalt_buf0[12]; w1_t[1] = esalt_buf0[13]; w1_t[2] = esalt_buf0[14]; w1_t[3] = esalt_buf0[15]; - w2_t[0] = esalt_buf1[ 0]; w2_t[1] = esalt_buf1[ 1]; w2_t[2] = esalt_buf1[ 2]; w2_t[3] = esalt_buf1[ 3]; - w3_t[0] = esalt_buf1[ 4]; w3_t[1] = esalt_buf1[ 5]; w3_t[2] = esalt_buf1[ 6]; @@ -5635,17 +5417,14 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p w0_t[1] = esalt_buf1[ 9]; w0_t[2] = esalt_buf1[10]; w0_t[3] = esalt_buf1[11]; - w1_t[0] = esalt_buf1[12]; w1_t[1] = esalt_buf1[13]; w1_t[2] = esalt_buf1[14]; w1_t[3] = esalt_buf1[15]; - w2_t[0] = esalt_buf2[ 0]; w2_t[1] = esalt_buf2[ 1]; w2_t[2] = esalt_buf2[ 2]; w2_t[3] = esalt_buf2[ 3]; - w3_t[0] = esalt_buf2[ 4]; w3_t[1] = esalt_buf2[ 5]; w3_t[2] = digest_esalt_len * 8; @@ -5728,7 +5507,7 @@ static void m11400s_1_1 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 p } } -__kernel void m11400_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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11400_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 sip_t *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 @@ -5804,21 +5583,21 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, switch (sw_1 | sw_2) { case 0: - m11400m_0_0 (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, esalt_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, l_bin2asc); + m11400m_0_0 (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, 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, l_bin2asc); break; case 1: - m11400m_0_1 (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, esalt_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, l_bin2asc); + m11400m_0_1 (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, 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, l_bin2asc); break; case 2: - m11400m_1_0 (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, esalt_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, l_bin2asc); + m11400m_1_0 (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, 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, l_bin2asc); break; case 3: - m11400m_1_1 (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, esalt_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, l_bin2asc); + m11400m_1_1 (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, 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, l_bin2asc); break; } } -__kernel void m11400_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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11400_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 sip_t *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 @@ -5894,21 +5673,21 @@ __kernel void m11400_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, switch (sw_1 | sw_2) { case 0: - m11400m_0_0 (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, esalt_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, l_bin2asc); + m11400m_0_0 (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, 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, l_bin2asc); break; case 1: - m11400m_0_1 (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, esalt_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, l_bin2asc); + m11400m_0_1 (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, 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, l_bin2asc); break; case 2: - m11400m_1_0 (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, esalt_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, l_bin2asc); + m11400m_1_0 (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, 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, l_bin2asc); break; case 3: - m11400m_1_1 (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, esalt_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, l_bin2asc); + m11400m_1_1 (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, 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, l_bin2asc); break; } } -__kernel void m11400_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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11400_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 sip_t *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 @@ -5984,21 +5763,21 @@ __kernel void m11400_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, switch (sw_1 | sw_2) { case 0: - m11400m_0_0 (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, esalt_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, l_bin2asc); + m11400m_0_0 (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, 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, l_bin2asc); break; case 1: - m11400m_0_1 (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, esalt_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, l_bin2asc); + m11400m_0_1 (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, 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, l_bin2asc); break; case 2: - m11400m_1_0 (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, esalt_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, l_bin2asc); + m11400m_1_0 (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, 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, l_bin2asc); break; case 3: - m11400m_1_1 (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, esalt_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, l_bin2asc); + m11400m_1_1 (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, 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, l_bin2asc); break; } } -__kernel void m11400_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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11400_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 sip_t *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 @@ -6074,21 +5853,21 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, switch (sw_1 | sw_2) { case 0: - m11400s_0_0 (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, esalt_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, l_bin2asc); + m11400s_0_0 (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, 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, l_bin2asc); break; case 1: - m11400s_0_1 (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, esalt_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, l_bin2asc); + m11400s_0_1 (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, 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, l_bin2asc); break; case 2: - m11400s_1_0 (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, esalt_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, l_bin2asc); + m11400s_1_0 (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, 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, l_bin2asc); break; case 3: - m11400s_1_1 (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, esalt_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, l_bin2asc); + m11400s_1_1 (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, 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, l_bin2asc); break; } } -__kernel void m11400_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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11400_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 sip_t *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 @@ -6164,21 +5943,21 @@ __kernel void m11400_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, switch (sw_1 | sw_2) { case 0: - m11400s_0_0 (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, esalt_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, l_bin2asc); + m11400s_0_0 (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, 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, l_bin2asc); break; case 1: - m11400s_0_1 (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, esalt_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, l_bin2asc); + m11400s_0_1 (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, 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, l_bin2asc); break; case 2: - m11400s_1_0 (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, esalt_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, l_bin2asc); + m11400s_1_0 (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, 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, l_bin2asc); break; case 3: - m11400s_1_1 (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, esalt_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, l_bin2asc); + m11400s_1_1 (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, 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, l_bin2asc); break; } } -__kernel void m11400_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 sip_t *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m11400_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 sip_t *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 @@ -6254,16 +6033,16 @@ __kernel void m11400_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, switch (sw_1 | sw_2) { case 0: - m11400s_0_0 (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, esalt_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, l_bin2asc); + m11400s_0_0 (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, 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, l_bin2asc); break; case 1: - m11400s_0_1 (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, esalt_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, l_bin2asc); + m11400s_0_1 (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, 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, l_bin2asc); break; case 2: - m11400s_1_0 (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, esalt_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, l_bin2asc); + m11400s_1_0 (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, 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, l_bin2asc); break; case 3: - m11400s_1_1 (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, esalt_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, l_bin2asc); + m11400s_1_1 (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, 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, l_bin2asc); break; } }