X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm01500_a3.cl;h=a84f922120f9694d3a314548b84520bf62f64cc3;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=0dd3781ab8855190c3e0a9fac810dacd121cca3b;hpb=838b6a3b8bff2db1f4a9fad11ecb19b6fa2b4fec;p=hashcat.git diff --git a/OpenCL/m01500_a3.cl b/OpenCL/m01500_a3.cl index 0dd3781..a84f922 100644 --- a/OpenCL/m01500_a3.cl +++ b/OpenCL/m01500_a3.cl @@ -7,20 +7,14 @@ #define _DES_ -#include "include/constants.h" -#include "include/kernel_vendor.h" +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" -#define DGST_R0 0 -#define DGST_R1 1 -#define DGST_R2 2 -#define DGST_R3 3 - -#include "include/kernel_functions.c" -#include "OpenCL/types_ocl.c" -#include "OpenCL/common.c" - -#define COMPARE_S "OpenCL/check_single_comp4_bs.c" -#define COMPARE_M "OpenCL/check_multi_comp4_bs.c" +#define COMPARE_S "inc_comp_single_bs.cl" +#define COMPARE_M "inc_comp_multi_bs.cl" #define myselx(a,b,c) ((c) ? (b) : (a)) @@ -1569,17 +1563,9 @@ void DESCrypt (const u32 SALT, const u32 K00, const u32 K01, const u32 K02, cons for (u32 ii = 0; ii < 25; ii++) { - #ifdef IS_NV - #if CUDA_ARCH >= 500 - #else + #ifdef _unroll #pragma unroll #endif - #endif - - #ifdef IS_AMD - #pragma unroll - #endif - for (u32 i = 0; i < 2; i++) { if (i) KEYSET10 else KEYSET00 @@ -1705,17 +1691,9 @@ void DESCrypt (const u32 SALT, const u32 K00, const u32 K01, const u32 K02, cons for (u32 ii = 0; ii < 25; ii++) { - #ifdef IS_NV - #if CUDA_ARCH >= 500 - #else + #ifdef _unroll #pragma unroll #endif - #endif - - #ifdef IS_AMD - #pragma unroll - #endif - for (u32 i = 0; i < 2; i++) { if (i) KEYSET10 else KEYSET00 @@ -1906,7 +1884,7 @@ void transpose32c (u32 data[32]) swap (data[30], data[31], 1, 0x55555555); } -void m01500m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) +void m01500m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) { /** * base @@ -1931,62 +1909,62 @@ void m01500m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global co const u32 w0s = (w0 << 1) & 0xfefefefe; const u32 w1s = (w1 << 1) & 0xfefefefe; - #define K00 -((w0s >> ( 0 + 7)) & 1) - #define K01 -((w0s >> ( 0 + 6)) & 1) - #define K02 -((w0s >> ( 0 + 5)) & 1) - #define K03 -((w0s >> ( 0 + 4)) & 1) - #define K04 -((w0s >> ( 0 + 3)) & 1) - #define K05 -((w0s >> ( 0 + 2)) & 1) - #define K06 -((w0s >> ( 0 + 1)) & 1) - #define K07 -((w0s >> ( 8 + 7)) & 1) - #define K08 -((w0s >> ( 8 + 6)) & 1) - #define K09 -((w0s >> ( 8 + 5)) & 1) - #define K10 -((w0s >> ( 8 + 4)) & 1) - #define K11 -((w0s >> ( 8 + 3)) & 1) - #define K12 -((w0s >> ( 8 + 2)) & 1) - #define K13 -((w0s >> ( 8 + 1)) & 1) - #define K14 -((w0s >> (16 + 7)) & 1) - #define K15 -((w0s >> (16 + 6)) & 1) - #define K16 -((w0s >> (16 + 5)) & 1) - #define K17 -((w0s >> (16 + 4)) & 1) - #define K18 -((w0s >> (16 + 3)) & 1) - #define K19 -((w0s >> (16 + 2)) & 1) - #define K20 -((w0s >> (16 + 1)) & 1) - #define K21 -((w0s >> (24 + 7)) & 1) - #define K22 -((w0s >> (24 + 6)) & 1) - #define K23 -((w0s >> (24 + 5)) & 1) - #define K24 -((w0s >> (24 + 4)) & 1) - #define K25 -((w0s >> (24 + 3)) & 1) - #define K26 -((w0s >> (24 + 2)) & 1) - #define K27 -((w0s >> (24 + 1)) & 1) - #define K28 -((w1s >> ( 0 + 7)) & 1) - #define K29 -((w1s >> ( 0 + 6)) & 1) - #define K30 -((w1s >> ( 0 + 5)) & 1) - #define K31 -((w1s >> ( 0 + 4)) & 1) - #define K32 -((w1s >> ( 0 + 3)) & 1) - #define K33 -((w1s >> ( 0 + 2)) & 1) - #define K34 -((w1s >> ( 0 + 1)) & 1) - #define K35 -((w1s >> ( 8 + 7)) & 1) - #define K36 -((w1s >> ( 8 + 6)) & 1) - #define K37 -((w1s >> ( 8 + 5)) & 1) - #define K38 -((w1s >> ( 8 + 4)) & 1) - #define K39 -((w1s >> ( 8 + 3)) & 1) - #define K40 -((w1s >> ( 8 + 2)) & 1) - #define K41 -((w1s >> ( 8 + 1)) & 1) - #define K42 -((w1s >> (16 + 7)) & 1) - #define K43 -((w1s >> (16 + 6)) & 1) - #define K44 -((w1s >> (16 + 5)) & 1) - #define K45 -((w1s >> (16 + 4)) & 1) - #define K46 -((w1s >> (16 + 3)) & 1) - #define K47 -((w1s >> (16 + 2)) & 1) - #define K48 -((w1s >> (16 + 1)) & 1) - #define K49 -((w1s >> (24 + 7)) & 1) - #define K50 -((w1s >> (24 + 6)) & 1) - #define K51 -((w1s >> (24 + 5)) & 1) - #define K52 -((w1s >> (24 + 4)) & 1) - #define K53 -((w1s >> (24 + 3)) & 1) - #define K54 -((w1s >> (24 + 2)) & 1) - #define K55 -((w1s >> (24 + 1)) & 1) + #define K00 (((w0s >> ( 0 + 7)) & 1) ? -1 : 0) + #define K01 (((w0s >> ( 0 + 6)) & 1) ? -1 : 0) + #define K02 (((w0s >> ( 0 + 5)) & 1) ? -1 : 0) + #define K03 (((w0s >> ( 0 + 4)) & 1) ? -1 : 0) + #define K04 (((w0s >> ( 0 + 3)) & 1) ? -1 : 0) + #define K05 (((w0s >> ( 0 + 2)) & 1) ? -1 : 0) + #define K06 (((w0s >> ( 0 + 1)) & 1) ? -1 : 0) + #define K07 (((w0s >> ( 8 + 7)) & 1) ? -1 : 0) + #define K08 (((w0s >> ( 8 + 6)) & 1) ? -1 : 0) + #define K09 (((w0s >> ( 8 + 5)) & 1) ? -1 : 0) + #define K10 (((w0s >> ( 8 + 4)) & 1) ? -1 : 0) + #define K11 (((w0s >> ( 8 + 3)) & 1) ? -1 : 0) + #define K12 (((w0s >> ( 8 + 2)) & 1) ? -1 : 0) + #define K13 (((w0s >> ( 8 + 1)) & 1) ? -1 : 0) + #define K14 (((w0s >> (16 + 7)) & 1) ? -1 : 0) + #define K15 (((w0s >> (16 + 6)) & 1) ? -1 : 0) + #define K16 (((w0s >> (16 + 5)) & 1) ? -1 : 0) + #define K17 (((w0s >> (16 + 4)) & 1) ? -1 : 0) + #define K18 (((w0s >> (16 + 3)) & 1) ? -1 : 0) + #define K19 (((w0s >> (16 + 2)) & 1) ? -1 : 0) + #define K20 (((w0s >> (16 + 1)) & 1) ? -1 : 0) + #define K21 (((w0s >> (24 + 7)) & 1) ? -1 : 0) + #define K22 (((w0s >> (24 + 6)) & 1) ? -1 : 0) + #define K23 (((w0s >> (24 + 5)) & 1) ? -1 : 0) + #define K24 (((w0s >> (24 + 4)) & 1) ? -1 : 0) + #define K25 (((w0s >> (24 + 3)) & 1) ? -1 : 0) + #define K26 (((w0s >> (24 + 2)) & 1) ? -1 : 0) + #define K27 (((w0s >> (24 + 1)) & 1) ? -1 : 0) + #define K28 (((w1s >> ( 0 + 7)) & 1) ? -1 : 0) + #define K29 (((w1s >> ( 0 + 6)) & 1) ? -1 : 0) + #define K30 (((w1s >> ( 0 + 5)) & 1) ? -1 : 0) + #define K31 (((w1s >> ( 0 + 4)) & 1) ? -1 : 0) + #define K32 (((w1s >> ( 0 + 3)) & 1) ? -1 : 0) + #define K33 (((w1s >> ( 0 + 2)) & 1) ? -1 : 0) + #define K34 (((w1s >> ( 0 + 1)) & 1) ? -1 : 0) + #define K35 (((w1s >> ( 8 + 7)) & 1) ? -1 : 0) + #define K36 (((w1s >> ( 8 + 6)) & 1) ? -1 : 0) + #define K37 (((w1s >> ( 8 + 5)) & 1) ? -1 : 0) + #define K38 (((w1s >> ( 8 + 4)) & 1) ? -1 : 0) + #define K39 (((w1s >> ( 8 + 3)) & 1) ? -1 : 0) + #define K40 (((w1s >> ( 8 + 2)) & 1) ? -1 : 0) + #define K41 (((w1s >> ( 8 + 1)) & 1) ? -1 : 0) + #define K42 (((w1s >> (16 + 7)) & 1) ? -1 : 0) + #define K43 (((w1s >> (16 + 6)) & 1) ? -1 : 0) + #define K44 (((w1s >> (16 + 5)) & 1) ? -1 : 0) + #define K45 (((w1s >> (16 + 4)) & 1) ? -1 : 0) + #define K46 (((w1s >> (16 + 3)) & 1) ? -1 : 0) + #define K47 (((w1s >> (16 + 2)) & 1) ? -1 : 0) + #define K48 (((w1s >> (16 + 1)) & 1) ? -1 : 0) + #define K49 (((w1s >> (24 + 7)) & 1) ? -1 : 0) + #define K50 (((w1s >> (24 + 6)) & 1) ? -1 : 0) + #define K51 (((w1s >> (24 + 5)) & 1) ? -1 : 0) + #define K52 (((w1s >> (24 + 4)) & 1) ? -1 : 0) + #define K53 (((w1s >> (24 + 3)) & 1) ? -1 : 0) + #define K54 (((w1s >> (24 + 2)) & 1) ? -1 : 0) + #define K55 (((w1s >> (24 + 1)) & 1) ? -1 : 0) /** * inner loop @@ -2222,7 +2200,9 @@ void m01500m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global co u32 tmpResult = 0; + #ifdef _unroll #pragma unroll + #endif for (int i = 0; i < 32; i++) { const u32 b0 = -((search[0] >> i) & 1); @@ -2249,7 +2229,9 @@ void m01500m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global co u32 out0[32]; u32 out1[32]; + #ifdef _unroll #pragma unroll + #endif for (int i = 0; i < 32; i++) { out0[i] = out[ 0 + 31 - i]; @@ -2259,7 +2241,9 @@ void m01500m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global co transpose32c (out0); transpose32c (out1); + #ifdef _unroll #pragma unroll + #endif for (int slice = 0; slice < 32; slice++) { const u32 r0 = out0[31 - slice]; @@ -2272,7 +2256,7 @@ void m01500m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global co } } -void m01500s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) +void m01500s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) { /** * base @@ -2294,70 +2278,70 @@ void m01500s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global co const u32 s0 = digests_buf[0].digest_buf[0]; const u32 s1 = digests_buf[0].digest_buf[1]; - #define S00 -((s0 >> 0) & 1) - #define S01 -((s0 >> 1) & 1) - #define S02 -((s0 >> 2) & 1) - #define S03 -((s0 >> 3) & 1) - #define S04 -((s0 >> 4) & 1) - #define S05 -((s0 >> 5) & 1) - #define S06 -((s0 >> 6) & 1) - #define S07 -((s0 >> 7) & 1) - #define S08 -((s0 >> 8) & 1) - #define S09 -((s0 >> 9) & 1) - #define S10 -((s0 >> 10) & 1) - #define S11 -((s0 >> 11) & 1) - #define S12 -((s0 >> 12) & 1) - #define S13 -((s0 >> 13) & 1) - #define S14 -((s0 >> 14) & 1) - #define S15 -((s0 >> 15) & 1) - #define S16 -((s0 >> 16) & 1) - #define S17 -((s0 >> 17) & 1) - #define S18 -((s0 >> 18) & 1) - #define S19 -((s0 >> 19) & 1) - #define S20 -((s0 >> 20) & 1) - #define S21 -((s0 >> 21) & 1) - #define S22 -((s0 >> 22) & 1) - #define S23 -((s0 >> 23) & 1) - #define S24 -((s0 >> 24) & 1) - #define S25 -((s0 >> 25) & 1) - #define S26 -((s0 >> 26) & 1) - #define S27 -((s0 >> 27) & 1) - #define S28 -((s0 >> 28) & 1) - #define S29 -((s0 >> 29) & 1) - #define S30 -((s0 >> 30) & 1) - #define S31 -((s0 >> 31) & 1) - #define S32 -((s1 >> 0) & 1) - #define S33 -((s1 >> 1) & 1) - #define S34 -((s1 >> 2) & 1) - #define S35 -((s1 >> 3) & 1) - #define S36 -((s1 >> 4) & 1) - #define S37 -((s1 >> 5) & 1) - #define S38 -((s1 >> 6) & 1) - #define S39 -((s1 >> 7) & 1) - #define S40 -((s1 >> 8) & 1) - #define S41 -((s1 >> 9) & 1) - #define S42 -((s1 >> 10) & 1) - #define S43 -((s1 >> 11) & 1) - #define S44 -((s1 >> 12) & 1) - #define S45 -((s1 >> 13) & 1) - #define S46 -((s1 >> 14) & 1) - #define S47 -((s1 >> 15) & 1) - #define S48 -((s1 >> 16) & 1) - #define S49 -((s1 >> 17) & 1) - #define S50 -((s1 >> 18) & 1) - #define S51 -((s1 >> 19) & 1) - #define S52 -((s1 >> 20) & 1) - #define S53 -((s1 >> 21) & 1) - #define S54 -((s1 >> 22) & 1) - #define S55 -((s1 >> 23) & 1) - #define S56 -((s1 >> 24) & 1) - #define S57 -((s1 >> 25) & 1) - #define S58 -((s1 >> 26) & 1) - #define S59 -((s1 >> 27) & 1) - #define S60 -((s1 >> 28) & 1) - #define S61 -((s1 >> 29) & 1) - #define S62 -((s1 >> 30) & 1) - #define S63 -((s1 >> 31) & 1) + #define S00 (((s0 >> 0) & 1) ? -1 : 0) + #define S01 (((s0 >> 1) & 1) ? -1 : 0) + #define S02 (((s0 >> 2) & 1) ? -1 : 0) + #define S03 (((s0 >> 3) & 1) ? -1 : 0) + #define S04 (((s0 >> 4) & 1) ? -1 : 0) + #define S05 (((s0 >> 5) & 1) ? -1 : 0) + #define S06 (((s0 >> 6) & 1) ? -1 : 0) + #define S07 (((s0 >> 7) & 1) ? -1 : 0) + #define S08 (((s0 >> 8) & 1) ? -1 : 0) + #define S09 (((s0 >> 9) & 1) ? -1 : 0) + #define S10 (((s0 >> 10) & 1) ? -1 : 0) + #define S11 (((s0 >> 11) & 1) ? -1 : 0) + #define S12 (((s0 >> 12) & 1) ? -1 : 0) + #define S13 (((s0 >> 13) & 1) ? -1 : 0) + #define S14 (((s0 >> 14) & 1) ? -1 : 0) + #define S15 (((s0 >> 15) & 1) ? -1 : 0) + #define S16 (((s0 >> 16) & 1) ? -1 : 0) + #define S17 (((s0 >> 17) & 1) ? -1 : 0) + #define S18 (((s0 >> 18) & 1) ? -1 : 0) + #define S19 (((s0 >> 19) & 1) ? -1 : 0) + #define S20 (((s0 >> 20) & 1) ? -1 : 0) + #define S21 (((s0 >> 21) & 1) ? -1 : 0) + #define S22 (((s0 >> 22) & 1) ? -1 : 0) + #define S23 (((s0 >> 23) & 1) ? -1 : 0) + #define S24 (((s0 >> 24) & 1) ? -1 : 0) + #define S25 (((s0 >> 25) & 1) ? -1 : 0) + #define S26 (((s0 >> 26) & 1) ? -1 : 0) + #define S27 (((s0 >> 27) & 1) ? -1 : 0) + #define S28 (((s0 >> 28) & 1) ? -1 : 0) + #define S29 (((s0 >> 29) & 1) ? -1 : 0) + #define S30 (((s0 >> 30) & 1) ? -1 : 0) + #define S31 (((s0 >> 31) & 1) ? -1 : 0) + #define S32 (((s1 >> 0) & 1) ? -1 : 0) + #define S33 (((s1 >> 1) & 1) ? -1 : 0) + #define S34 (((s1 >> 2) & 1) ? -1 : 0) + #define S35 (((s1 >> 3) & 1) ? -1 : 0) + #define S36 (((s1 >> 4) & 1) ? -1 : 0) + #define S37 (((s1 >> 5) & 1) ? -1 : 0) + #define S38 (((s1 >> 6) & 1) ? -1 : 0) + #define S39 (((s1 >> 7) & 1) ? -1 : 0) + #define S40 (((s1 >> 8) & 1) ? -1 : 0) + #define S41 (((s1 >> 9) & 1) ? -1 : 0) + #define S42 (((s1 >> 10) & 1) ? -1 : 0) + #define S43 (((s1 >> 11) & 1) ? -1 : 0) + #define S44 (((s1 >> 12) & 1) ? -1 : 0) + #define S45 (((s1 >> 13) & 1) ? -1 : 0) + #define S46 (((s1 >> 14) & 1) ? -1 : 0) + #define S47 (((s1 >> 15) & 1) ? -1 : 0) + #define S48 (((s1 >> 16) & 1) ? -1 : 0) + #define S49 (((s1 >> 17) & 1) ? -1 : 0) + #define S50 (((s1 >> 18) & 1) ? -1 : 0) + #define S51 (((s1 >> 19) & 1) ? -1 : 0) + #define S52 (((s1 >> 20) & 1) ? -1 : 0) + #define S53 (((s1 >> 21) & 1) ? -1 : 0) + #define S54 (((s1 >> 22) & 1) ? -1 : 0) + #define S55 (((s1 >> 23) & 1) ? -1 : 0) + #define S56 (((s1 >> 24) & 1) ? -1 : 0) + #define S57 (((s1 >> 25) & 1) ? -1 : 0) + #define S58 (((s1 >> 26) & 1) ? -1 : 0) + #define S59 (((s1 >> 27) & 1) ? -1 : 0) + #define S60 (((s1 >> 28) & 1) ? -1 : 0) + #define S61 (((s1 >> 29) & 1) ? -1 : 0) + #define S62 (((s1 >> 30) & 1) ? -1 : 0) + #define S63 (((s1 >> 31) & 1) ? -1 : 0) /** * base @@ -2369,62 +2353,62 @@ void m01500s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global co const u32 w0s = (w0 << 1) & 0xfefefefe; const u32 w1s = (w1 << 1) & 0xfefefefe; - #define K00 -((w0s >> ( 0 + 7)) & 1) - #define K01 -((w0s >> ( 0 + 6)) & 1) - #define K02 -((w0s >> ( 0 + 5)) & 1) - #define K03 -((w0s >> ( 0 + 4)) & 1) - #define K04 -((w0s >> ( 0 + 3)) & 1) - #define K05 -((w0s >> ( 0 + 2)) & 1) - #define K06 -((w0s >> ( 0 + 1)) & 1) - #define K07 -((w0s >> ( 8 + 7)) & 1) - #define K08 -((w0s >> ( 8 + 6)) & 1) - #define K09 -((w0s >> ( 8 + 5)) & 1) - #define K10 -((w0s >> ( 8 + 4)) & 1) - #define K11 -((w0s >> ( 8 + 3)) & 1) - #define K12 -((w0s >> ( 8 + 2)) & 1) - #define K13 -((w0s >> ( 8 + 1)) & 1) - #define K14 -((w0s >> (16 + 7)) & 1) - #define K15 -((w0s >> (16 + 6)) & 1) - #define K16 -((w0s >> (16 + 5)) & 1) - #define K17 -((w0s >> (16 + 4)) & 1) - #define K18 -((w0s >> (16 + 3)) & 1) - #define K19 -((w0s >> (16 + 2)) & 1) - #define K20 -((w0s >> (16 + 1)) & 1) - #define K21 -((w0s >> (24 + 7)) & 1) - #define K22 -((w0s >> (24 + 6)) & 1) - #define K23 -((w0s >> (24 + 5)) & 1) - #define K24 -((w0s >> (24 + 4)) & 1) - #define K25 -((w0s >> (24 + 3)) & 1) - #define K26 -((w0s >> (24 + 2)) & 1) - #define K27 -((w0s >> (24 + 1)) & 1) - #define K28 -((w1s >> ( 0 + 7)) & 1) - #define K29 -((w1s >> ( 0 + 6)) & 1) - #define K30 -((w1s >> ( 0 + 5)) & 1) - #define K31 -((w1s >> ( 0 + 4)) & 1) - #define K32 -((w1s >> ( 0 + 3)) & 1) - #define K33 -((w1s >> ( 0 + 2)) & 1) - #define K34 -((w1s >> ( 0 + 1)) & 1) - #define K35 -((w1s >> ( 8 + 7)) & 1) - #define K36 -((w1s >> ( 8 + 6)) & 1) - #define K37 -((w1s >> ( 8 + 5)) & 1) - #define K38 -((w1s >> ( 8 + 4)) & 1) - #define K39 -((w1s >> ( 8 + 3)) & 1) - #define K40 -((w1s >> ( 8 + 2)) & 1) - #define K41 -((w1s >> ( 8 + 1)) & 1) - #define K42 -((w1s >> (16 + 7)) & 1) - #define K43 -((w1s >> (16 + 6)) & 1) - #define K44 -((w1s >> (16 + 5)) & 1) - #define K45 -((w1s >> (16 + 4)) & 1) - #define K46 -((w1s >> (16 + 3)) & 1) - #define K47 -((w1s >> (16 + 2)) & 1) - #define K48 -((w1s >> (16 + 1)) & 1) - #define K49 -((w1s >> (24 + 7)) & 1) - #define K50 -((w1s >> (24 + 6)) & 1) - #define K51 -((w1s >> (24 + 5)) & 1) - #define K52 -((w1s >> (24 + 4)) & 1) - #define K53 -((w1s >> (24 + 3)) & 1) - #define K54 -((w1s >> (24 + 2)) & 1) - #define K55 -((w1s >> (24 + 1)) & 1) + #define K00 (((w0s >> ( 0 + 7)) & 1) ? -1 : 0) + #define K01 (((w0s >> ( 0 + 6)) & 1) ? -1 : 0) + #define K02 (((w0s >> ( 0 + 5)) & 1) ? -1 : 0) + #define K03 (((w0s >> ( 0 + 4)) & 1) ? -1 : 0) + #define K04 (((w0s >> ( 0 + 3)) & 1) ? -1 : 0) + #define K05 (((w0s >> ( 0 + 2)) & 1) ? -1 : 0) + #define K06 (((w0s >> ( 0 + 1)) & 1) ? -1 : 0) + #define K07 (((w0s >> ( 8 + 7)) & 1) ? -1 : 0) + #define K08 (((w0s >> ( 8 + 6)) & 1) ? -1 : 0) + #define K09 (((w0s >> ( 8 + 5)) & 1) ? -1 : 0) + #define K10 (((w0s >> ( 8 + 4)) & 1) ? -1 : 0) + #define K11 (((w0s >> ( 8 + 3)) & 1) ? -1 : 0) + #define K12 (((w0s >> ( 8 + 2)) & 1) ? -1 : 0) + #define K13 (((w0s >> ( 8 + 1)) & 1) ? -1 : 0) + #define K14 (((w0s >> (16 + 7)) & 1) ? -1 : 0) + #define K15 (((w0s >> (16 + 6)) & 1) ? -1 : 0) + #define K16 (((w0s >> (16 + 5)) & 1) ? -1 : 0) + #define K17 (((w0s >> (16 + 4)) & 1) ? -1 : 0) + #define K18 (((w0s >> (16 + 3)) & 1) ? -1 : 0) + #define K19 (((w0s >> (16 + 2)) & 1) ? -1 : 0) + #define K20 (((w0s >> (16 + 1)) & 1) ? -1 : 0) + #define K21 (((w0s >> (24 + 7)) & 1) ? -1 : 0) + #define K22 (((w0s >> (24 + 6)) & 1) ? -1 : 0) + #define K23 (((w0s >> (24 + 5)) & 1) ? -1 : 0) + #define K24 (((w0s >> (24 + 4)) & 1) ? -1 : 0) + #define K25 (((w0s >> (24 + 3)) & 1) ? -1 : 0) + #define K26 (((w0s >> (24 + 2)) & 1) ? -1 : 0) + #define K27 (((w0s >> (24 + 1)) & 1) ? -1 : 0) + #define K28 (((w1s >> ( 0 + 7)) & 1) ? -1 : 0) + #define K29 (((w1s >> ( 0 + 6)) & 1) ? -1 : 0) + #define K30 (((w1s >> ( 0 + 5)) & 1) ? -1 : 0) + #define K31 (((w1s >> ( 0 + 4)) & 1) ? -1 : 0) + #define K32 (((w1s >> ( 0 + 3)) & 1) ? -1 : 0) + #define K33 (((w1s >> ( 0 + 2)) & 1) ? -1 : 0) + #define K34 (((w1s >> ( 0 + 1)) & 1) ? -1 : 0) + #define K35 (((w1s >> ( 8 + 7)) & 1) ? -1 : 0) + #define K36 (((w1s >> ( 8 + 6)) & 1) ? -1 : 0) + #define K37 (((w1s >> ( 8 + 5)) & 1) ? -1 : 0) + #define K38 (((w1s >> ( 8 + 4)) & 1) ? -1 : 0) + #define K39 (((w1s >> ( 8 + 3)) & 1) ? -1 : 0) + #define K40 (((w1s >> ( 8 + 2)) & 1) ? -1 : 0) + #define K41 (((w1s >> ( 8 + 1)) & 1) ? -1 : 0) + #define K42 (((w1s >> (16 + 7)) & 1) ? -1 : 0) + #define K43 (((w1s >> (16 + 6)) & 1) ? -1 : 0) + #define K44 (((w1s >> (16 + 5)) & 1) ? -1 : 0) + #define K45 (((w1s >> (16 + 4)) & 1) ? -1 : 0) + #define K46 (((w1s >> (16 + 3)) & 1) ? -1 : 0) + #define K47 (((w1s >> (16 + 2)) & 1) ? -1 : 0) + #define K48 (((w1s >> (16 + 1)) & 1) ? -1 : 0) + #define K49 (((w1s >> (24 + 7)) & 1) ? -1 : 0) + #define K50 (((w1s >> (24 + 6)) & 1) ? -1 : 0) + #define K51 (((w1s >> (24 + 5)) & 1) ? -1 : 0) + #define K52 (((w1s >> (24 + 4)) & 1) ? -1 : 0) + #define K53 (((w1s >> (24 + 3)) & 1) ? -1 : 0) + #define K54 (((w1s >> (24 + 2)) & 1) ? -1 : 0) + #define K55 (((w1s >> (24 + 1)) & 1) ? -1 : 0) /** * inner loop @@ -2676,7 +2660,9 @@ __kernel void m01500_tm (__global u32 *mod, __global bs_word_t *words_buf_r) const u32 w0s = (w0 << 1) & 0xfefefefe; + #ifdef _unroll #pragma unroll + #endif for (int i = 0, j = 0; i < 32; i += 8, j += 7) { atomic_or (&words_buf_r[block].b[j + 0], (((w0s >> (i + 7)) & 1) << slice)); @@ -2689,7 +2675,7 @@ __kernel void m01500_tm (__global u32 *mod, __global bs_word_t *words_buf_r) } } -__kernel void m01500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m01500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * base @@ -2704,18 +2690,18 @@ __kernel void m01500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, * main */ - m01500m (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); + m01500m (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); } -__kernel void m01500_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m01500_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { } -__kernel void m01500_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m01500_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { } -__kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * base @@ -2730,13 +2716,13 @@ __kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, * main */ - m01500s (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); + m01500s (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); } -__kernel void m01500_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m01500_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { } -__kernel void m01500_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m01500_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { }