X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm10100_a0.cl;h=117790e330082e80de3480311d6a892f475cc90a;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=833dfa40bd12ebbad1dc2b04037830e915435e2f;hpb=5f7c47b461b7e3ffb727baeee62f41ac537cc0ba;p=hashcat.git diff --git a/OpenCL/m10100_a0.cl b/OpenCL/m10100_a0.cl index 833dfa4..117790e 100644 --- a/OpenCL/m10100_a0.cl +++ b/OpenCL/m10100_a0.cl @@ -5,28 +5,23 @@ #define _SIPHASH_ -#include "include/constants.h" -#include "include/kernel_vendor.h" - -#define DGST_R0 0 -#define DGST_R1 1 -#define DGST_R2 2 -#define DGST_R3 3 - -#include "include/kernel_functions.c" -#include "OpenCL/types_ocl.c" -#include "OpenCL/common.c" -#include "include/rp_gpu.h" -#include "rp.c" - -#define COMPARE_S "OpenCL/check_single_comp4.c" -#define COMPARE_M "OpenCL/check_multi_comp4.c" +//incompatible to simd +//#define NEW_SIMD_CODE + +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" +#include "inc_rp.h" +#include "inc_rp.cl" +#include "inc_simd.cl" #define SIPROUND(v0,v1,v2,v3) \ (v0) += (v1); \ (v1) = rotl64 ((v1), 13); \ (v1) ^= (v0); \ - (v0) = as_ulong (as_uint2 ((v0)).s10); \ + (v0) = rotl64 ((v0), 32); \ (v2) += (v3); \ (v3) = rotl64 ((v3), 16); \ (v3) ^= (v2); \ @@ -36,9 +31,9 @@ (v2) += (v1); \ (v1) = rotl64 ((v1), 17); \ (v1) ^= (v2); \ - (v2) = as_ulong (as_uint2 ((v2)).s10); + (v2) = rotl64 ((v2), 32) -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m04 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m10100_m04 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier @@ -55,14 +50,12 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m04 (__glo if (gid >= gid_max) return; u32 pw_buf0[4]; + u32 pw_buf1[4]; pw_buf0[0] = pws[gid].i[ 0]; pw_buf0[1] = pws[gid].i[ 1]; pw_buf0[2] = pws[gid].i[ 2]; pw_buf0[3] = pws[gid].i[ 3]; - - u32 pw_buf1[4]; - pw_buf1[0] = pws[gid].i[ 4]; pw_buf1[1] = pws[gid].i[ 5]; pw_buf1[2] = pws[gid].i[ 6]; @@ -71,13 +64,13 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m04 (__glo const u32 pw_len = pws[gid].pw_len; /** - * base + * salt */ - u64 v0p = SIPHASHM_0; - u64 v1p = SIPHASHM_1; - u64 v2p = SIPHASHM_2; - u64 v3p = SIPHASHM_3; + u64x v0p = SIPHASHM_0; + u64x v1p = SIPHASHM_1; + u64x v2p = SIPHASHM_2; + u64x v3p = SIPHASHM_3; v0p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]); v1p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]); @@ -88,44 +81,50 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m04 (__glo * loop */ - for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - u32 w[16]; - - w[ 0] = pw_buf0[0]; - w[ 1] = pw_buf0[1]; - w[ 2] = pw_buf0[2]; - w[ 3] = pw_buf0[3]; - w[ 4] = pw_buf1[0]; - w[ 5] = pw_buf1[1]; - w[ 6] = pw_buf1[2]; - w[ 7] = pw_buf1[3]; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 0; - w[15] = 0; - - const u32 out_len = apply_rules (rules_buf[il_pos].cmds, &w[0], &w[4], pw_len); - - u64 *w_ptr = (u64 *) w; - - w_ptr[out_len / 8] |= (u64) out_len << 56; - - u64 v0 = v0p; - u64 v1 = v1p; - u64 v2 = v2p; - u64 v3 = v3p; + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; + + const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); + + /** + * siphash + */ + + switch (out_len / 8) + { + case 0: w0[1] |= out_len << 24; break; + case 1: w0[3] |= out_len << 24; break; + case 2: w1[1] |= out_len << 24; break; + case 3: w1[3] |= out_len << 24; break; + } + + u64x v0 = v0p; + u64x v1 = v1p; + u64x v2 = v2p; + u64x v3 = v3p; int i; int j; - for (i = 0, j = 0; i <= pw_len; i += 8, j += 2) + for (i = 0, j = 0; i <= out_len && i < 16; i += 8, j += 2) { - u64 m = hl32_to_64 (w[j + 1], w[j + 0]); + u64x m = hl32_to_64 (w0[j + 1], w0[j + 0]); + + v3 ^= m; + + SIPROUND (v0, v1, v2, v3); + SIPROUND (v0, v1, v2, v3); + + v0 ^= m; + } + + for ( j = 0; i <= out_len && i < 32; i += 8, j += 2) + { + u64x m = hl32_to_64 (w1[j + 1], w1[j + 0]); v3 ^= m; @@ -142,29 +141,26 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m04 (__glo SIPROUND (v0, v1, v2, v3); SIPROUND (v0, v1, v2, v3); - const u64 v = v0 ^ v1 ^ v2 ^ v3; + const u64x v = v0 ^ v1 ^ v2 ^ v3; - const u32 a = l32_from_64 (v); - const u32 b = h32_from_64 (v); + const u32x a = l32_from_64 (v); + const u32x b = h32_from_64 (v); - const u32 r0 = a; - const u32 r1 = b; - const u32 r2 = 0; - const u32 r3 = 0; + const u32x z = 0; - #include COMPARE_M + COMPARE_M_SIMD (a, b, z, z); } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m08 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m10100_m08 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m16 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m10100_m16 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_s04 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m10100_s04 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier @@ -181,14 +177,12 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_s04 (__glo if (gid >= gid_max) return; u32 pw_buf0[4]; + u32 pw_buf1[4]; pw_buf0[0] = pws[gid].i[ 0]; pw_buf0[1] = pws[gid].i[ 1]; pw_buf0[2] = pws[gid].i[ 2]; pw_buf0[3] = pws[gid].i[ 3]; - - u32 pw_buf1[4]; - pw_buf1[0] = pws[gid].i[ 4]; pw_buf1[1] = pws[gid].i[ 5]; pw_buf1[2] = pws[gid].i[ 6]; @@ -197,13 +191,13 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_s04 (__glo const u32 pw_len = pws[gid].pw_len; /** - * base + * salt */ - u64 v0p = SIPHASHM_0; - u64 v1p = SIPHASHM_1; - u64 v2p = SIPHASHM_2; - u64 v3p = SIPHASHM_3; + u64x v0p = SIPHASHM_0; + u64x v1p = SIPHASHM_1; + u64x v2p = SIPHASHM_2; + u64x v3p = SIPHASHM_3; v0p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]); v1p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]); @@ -218,52 +212,58 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_s04 (__glo { 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] + 0, + 0 }; /** * loop */ - for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - u32 w[16]; - - w[ 0] = pw_buf0[0]; - w[ 1] = pw_buf0[1]; - w[ 2] = pw_buf0[2]; - w[ 3] = pw_buf0[3]; - w[ 4] = pw_buf1[0]; - w[ 5] = pw_buf1[1]; - w[ 6] = pw_buf1[2]; - w[ 7] = pw_buf1[3]; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 0; - w[15] = 0; - - const u32 out_len = apply_rules (rules_buf[il_pos].cmds, &w[0], &w[4], pw_len); - - u64 *w_ptr = (u64 *) w; - - w_ptr[out_len / 8] |= (u64) out_len << 56; - - u64 v0 = v0p; - u64 v1 = v1p; - u64 v2 = v2p; - u64 v3 = v3p; + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; + + const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); + + /** + * siphash + */ + + switch (out_len / 8) + { + case 0: w0[1] |= out_len << 24; break; + case 1: w0[3] |= out_len << 24; break; + case 2: w1[1] |= out_len << 24; break; + case 3: w1[3] |= out_len << 24; break; + } + + u64x v0 = v0p; + u64x v1 = v1p; + u64x v2 = v2p; + u64x v3 = v3p; int i; int j; - for (i = 0, j = 0; i <= pw_len; i += 8, j += 2) + for (i = 0, j = 0; i <= out_len && i < 16; i += 8, j += 2) + { + u64x m = hl32_to_64 (w0[j + 1], w0[j + 0]); + + v3 ^= m; + + SIPROUND (v0, v1, v2, v3); + SIPROUND (v0, v1, v2, v3); + + v0 ^= m; + } + + for ( j = 0; i <= out_len && i < 32; i += 8, j += 2) { - u64 m = hl32_to_64 (w[j + 1], w[j + 0]); + u64x m = hl32_to_64 (w1[j + 1], w1[j + 0]); v3 ^= m; @@ -280,24 +280,21 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_s04 (__glo SIPROUND (v0, v1, v2, v3); SIPROUND (v0, v1, v2, v3); - const u64 v = v0 ^ v1 ^ v2 ^ v3; + const u64x v = v0 ^ v1 ^ v2 ^ v3; - const u32 a = l32_from_64 (v); - const u32 b = h32_from_64 (v); + const u32x a = l32_from_64 (v); + const u32x b = h32_from_64 (v); - const u32 r0 = a; - const u32 r1 = b; - const u32 r2 = 0; - const u32 r3 = 0; + const u32x z = 0; - #include COMPARE_S + COMPARE_S_SIMD (a, b, z, z); } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_s08 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m10100_s08 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_s16 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m10100_s16 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { }