diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl index 6cc88efcd..e48829cfa 100644 --- a/OpenCL/inc_common.cl +++ b/OpenCL/inc_common.cl @@ -44,10 +44,10 @@ __global const salt_t *salt_bufs, \ __global const p19 *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, \ + __global uint4 *d_scryptV0_buf, \ + __global uint4 *d_scryptV1_buf, \ + __global uint4 *d_scryptV2_buf, \ + __global uint4 *d_scryptV3_buf, \ const u32 bitmap_mask, \ const u32 bitmap_shift1, \ const u32 bitmap_shift2, \ @@ -75,6 +75,7 @@ #define KERN_ATTR_RULES_ESALT(e) KERN_ATTR (__constant, __global const bf_t *bfs_buf, void, void, e) #define KERN_ATTR_TMPS(t) KERN_ATTR (__global, __global const bf_t *bfs_buf, t, void, void) #define KERN_ATTR_TMPS_ESALT(t,e) KERN_ATTR (__global, __global const bf_t *bfs_buf, t, void, e) +#define KERN_ATTR_TMPS_HOOKS(t,h) KERN_ATTR (__global, __global const bf_t *bfs_buf, t, h, void) #define KERN_ATTR_VECTOR() KERN_ATTR (__global, __constant const u32x *words_buf_r, void, void, void) #define KERN_ATTR_VECTOR_ESALT(e) KERN_ATTR (__global, __constant const u32x *words_buf_r, void, void, e) diff --git a/OpenCL/m02000_a0-pure.cl b/OpenCL/m02000_a0-pure.cl index 1552d30bb..8baa27eeb 100644 --- a/OpenCL/m02000_a0-pure.cl +++ b/OpenCL/m02000_a0-pure.cl @@ -7,10 +7,10 @@ #include "inc_types.cl" #include "inc_common.cl" -__kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *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 u64 gid_max) +__kernel void m02000_mxx (KERN_ATTR_BASIC ()) { } -__kernel void m02000_sxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *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 u64 gid_max) +__kernel void m02000_sxx (KERN_ATTR_BASIC ()) { } diff --git a/OpenCL/m02000_a1-pure.cl b/OpenCL/m02000_a1-pure.cl index 1552d30bb..8baa27eeb 100644 --- a/OpenCL/m02000_a1-pure.cl +++ b/OpenCL/m02000_a1-pure.cl @@ -7,10 +7,10 @@ #include "inc_types.cl" #include "inc_common.cl" -__kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *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 u64 gid_max) +__kernel void m02000_mxx (KERN_ATTR_BASIC ()) { } -__kernel void m02000_sxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *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 u64 gid_max) +__kernel void m02000_sxx (KERN_ATTR_BASIC ()) { } diff --git a/OpenCL/m02000_a3-pure.cl b/OpenCL/m02000_a3-pure.cl index 1552d30bb..8baa27eeb 100644 --- a/OpenCL/m02000_a3-pure.cl +++ b/OpenCL/m02000_a3-pure.cl @@ -7,10 +7,10 @@ #include "inc_types.cl" #include "inc_common.cl" -__kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *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 u64 gid_max) +__kernel void m02000_mxx (KERN_ATTR_BASIC ()) { } -__kernel void m02000_sxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *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 u64 gid_max) +__kernel void m02000_sxx (KERN_ATTR_BASIC ()) { } diff --git a/OpenCL/m08900-pure.cl b/OpenCL/m08900-pure.cl index 61e00a974..fd388c70d 100644 --- a/OpenCL/m08900-pure.cl +++ b/OpenCL/m08900-pure.cl @@ -207,7 +207,7 @@ DECLSPEC void scrypt_smix (uint4 *X, uint4 *T, __global uint4 *V0, __global uint // there can be no __attribute__((reqd_work_group_size(16, 1, 1))) because kernel is used by both -m 8900 and -m 9300 -__kernel void m08900_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global scrypt_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global uint4 *d_scryptV0_buf, __global uint4 *d_scryptV1_buf, __global uint4 *d_scryptV2_buf, __global uint4 *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 u64 gid_max) +__kernel void m08900_init (KERN_ATTR_TMPS (scrypt_tmp_t)) { /** * base @@ -272,7 +272,7 @@ __kernel void m08900_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m08900_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global scrypt_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global uint4 *d_scryptV0_buf, __global uint4 *d_scryptV1_buf, __global uint4 *d_scryptV2_buf, __global uint4 *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 u64 gid_max) +__kernel void m08900_loop (KERN_ATTR_TMPS (scrypt_tmp_t)) { const u64 gid = get_global_id (0); @@ -305,7 +305,7 @@ __kernel void m08900_loop (__global pw_t *pws, __global const kernel_rule_t *rul #endif } -__kernel void m08900_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global scrypt_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global uint4 *d_scryptV0_buf, __global uint4 *d_scryptV1_buf, __global uint4 *d_scryptV2_buf, __global uint4 *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 u64 gid_max) +__kernel void m08900_comp (KERN_ATTR_TMPS (scrypt_tmp_t)) { /** * base diff --git a/OpenCL/m09400-pure.cl b/OpenCL/m09400-pure.cl index 116bf1c5b..c8902b85c 100644 --- a/OpenCL/m09400-pure.cl +++ b/OpenCL/m09400-pure.cl @@ -17,7 +17,7 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" -__kernel void m09400_init (KERN_ATTR_TMPS_ESALT (const office2007_tmp_t, office2007_t)) +__kernel void m09400_init (KERN_ATTR_TMPS_ESALT (office2007_tmp_t, office2007_t)) { /** * base @@ -44,7 +44,7 @@ __kernel void m09400_init (KERN_ATTR_TMPS_ESALT (const office2007_tmp_t, office2 tmps[gid].out[4] = ctx.h[4]; } -__kernel void m09400_loop (KERN_ATTR_TMPS_ESALT (const office2007_tmp_t, office2007_t)) +__kernel void m09400_loop (KERN_ATTR_TMPS_ESALT (office2007_tmp_t, office2007_t)) { const u64 gid = get_global_id (0); @@ -111,7 +111,7 @@ __kernel void m09400_loop (KERN_ATTR_TMPS_ESALT (const office2007_tmp_t, office2 unpackv (tmps, out, gid, 4, t4); } -__kernel void m09400_comp (KERN_ATTR_TMPS_ESALT (const office2007_tmp_t, office2007_t)) +__kernel void m09400_comp (KERN_ATTR_TMPS_ESALT (office2007_tmp_t, office2007_t)) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); diff --git a/OpenCL/m09500-pure.cl b/OpenCL/m09500-pure.cl index 48ab66a42..068f496a0 100644 --- a/OpenCL/m09500-pure.cl +++ b/OpenCL/m09500-pure.cl @@ -17,7 +17,7 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" -__kernel void m09500_init (KERN_ATTR_TMPS_ESALT (const office2010_tmp_t, office2010_t)) +__kernel void m09500_init (KERN_ATTR_TMPS_ESALT (office2010_tmp_t, office2010_t)) { /** * base @@ -44,7 +44,7 @@ __kernel void m09500_init (KERN_ATTR_TMPS_ESALT (const office2010_tmp_t, office2 tmps[gid].out[4] = ctx.h[4]; } -__kernel void m09500_loop (KERN_ATTR_TMPS_ESALT (const office2010_tmp_t, office2010_t)) +__kernel void m09500_loop (KERN_ATTR_TMPS_ESALT (office2010_tmp_t, office2010_t)) { const u64 gid = get_global_id (0); @@ -111,7 +111,7 @@ __kernel void m09500_loop (KERN_ATTR_TMPS_ESALT (const office2010_tmp_t, office2 unpackv (tmps, out, gid, 4, t4); } -__kernel void m09500_comp (KERN_ATTR_TMPS_ESALT (const office2010_tmp_t, office2010_t)) +__kernel void m09500_comp (KERN_ATTR_TMPS_ESALT (office2010_tmp_t, office2010_t)) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); diff --git a/OpenCL/m09600-pure.cl b/OpenCL/m09600-pure.cl index b8cfec7e6..a77319f43 100644 --- a/OpenCL/m09600-pure.cl +++ b/OpenCL/m09600-pure.cl @@ -17,7 +17,7 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" -__kernel void m09600_init (KERN_ATTR_TMPS_ESALT (const office2013_tmp_t, office2013_t)) +__kernel void m09600_init (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t)) { /** * base @@ -47,7 +47,7 @@ __kernel void m09600_init (KERN_ATTR_TMPS_ESALT (const office2013_tmp_t, office2 tmps[gid].out[7] = ctx.h[7]; } -__kernel void m09600_loop (KERN_ATTR_TMPS_ESALT (const office2013_tmp_t, office2013_t)) +__kernel void m09600_loop (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t)) { const u64 gid = get_global_id (0); @@ -157,7 +157,7 @@ __kernel void m09600_loop (KERN_ATTR_TMPS_ESALT (const office2013_tmp_t, office2 unpack64v (tmps, out, gid, 7, t7); } -__kernel void m09600_comp (KERN_ATTR_TMPS_ESALT (const office2013_tmp_t, office2013_t)) +__kernel void m09600_comp (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t)) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); diff --git a/OpenCL/m11300-pure.cl b/OpenCL/m11300-pure.cl index 0b8060177..54ac4c55f 100644 --- a/OpenCL/m11300-pure.cl +++ b/OpenCL/m11300-pure.cl @@ -72,7 +72,7 @@ DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest); } -__kernel void m11300_init (KERN_ATTR_TMPS_ESALT (const bitcoin_wallet_tmp_t, bitcoin_wallet_t)) +__kernel void m11300_init (KERN_ATTR_TMPS_ESALT (bitcoin_wallet_tmp_t, bitcoin_wallet_t)) { /** * base @@ -102,7 +102,7 @@ __kernel void m11300_init (KERN_ATTR_TMPS_ESALT (const bitcoin_wallet_tmp_t, bit tmps[gid].dgst[7] = ctx.h[7]; } -__kernel void m11300_loop (KERN_ATTR_TMPS_ESALT (const bitcoin_wallet_tmp_t, bitcoin_wallet_t)) +__kernel void m11300_loop (KERN_ATTR_TMPS_ESALT (bitcoin_wallet_tmp_t, bitcoin_wallet_t)) { const u64 gid = get_global_id (0); @@ -211,7 +211,7 @@ __kernel void m11300_loop (KERN_ATTR_TMPS_ESALT (const bitcoin_wallet_tmp_t, bit unpack64v (tmps, dgst, gid, 7, t7); } -__kernel void m11300_comp (KERN_ATTR_TMPS_ESALT (const bitcoin_wallet_tmp_t, bitcoin_wallet_t)) +__kernel void m11300_comp (KERN_ATTR_TMPS_ESALT (bitcoin_wallet_tmp_t, bitcoin_wallet_t)) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); diff --git a/OpenCL/m11600-pure.cl b/OpenCL/m11600-pure.cl index c9835e96c..c22df6156 100644 --- a/OpenCL/m11600-pure.cl +++ b/OpenCL/m11600-pure.cl @@ -111,7 +111,7 @@ DECLSPEC void memcat8c_be (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 len, co } } -__kernel void m11600_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global seven_zip_hook_t *seven_zip_hook, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 u64 gid_max) +__kernel void m11600_init (KERN_ATTR_TMPS_HOOKS (seven_zip_tmp_t, seven_zip_hook_t)) { /** * base @@ -158,7 +158,7 @@ __kernel void m11600_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].len = ctx.len; } -__kernel void m11600_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global seven_zip_hook_t *seven_zip_hook, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 u64 gid_max) +__kernel void m11600_loop (KERN_ATTR_TMPS_HOOKS (seven_zip_tmp_t, seven_zip_hook_t)) { /** * base @@ -257,7 +257,7 @@ __kernel void m11600_loop (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].len = ctx.len; } -__kernel void m11600_hook23 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global seven_zip_hook_t *seven_zip_hook, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 u64 gid_max) +__kernel void m11600_hook23 (KERN_ATTR_TMPS_HOOKS (seven_zip_tmp_t, seven_zip_hook_t)) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); @@ -301,17 +301,17 @@ __kernel void m11600_hook23 (__global pw_t *pws, __global const kernel_rule_t *r sha256_final (&ctx); - seven_zip_hook[gid].ukey[0] = swap32_S (ctx.h[0]); - seven_zip_hook[gid].ukey[1] = swap32_S (ctx.h[1]); - seven_zip_hook[gid].ukey[2] = swap32_S (ctx.h[2]); - seven_zip_hook[gid].ukey[3] = swap32_S (ctx.h[3]); - seven_zip_hook[gid].ukey[4] = swap32_S (ctx.h[4]); - seven_zip_hook[gid].ukey[5] = swap32_S (ctx.h[5]); - seven_zip_hook[gid].ukey[6] = swap32_S (ctx.h[6]); - seven_zip_hook[gid].ukey[7] = swap32_S (ctx.h[7]); + hooks[gid].ukey[0] = swap32_S (ctx.h[0]); + hooks[gid].ukey[1] = swap32_S (ctx.h[1]); + hooks[gid].ukey[2] = swap32_S (ctx.h[2]); + hooks[gid].ukey[3] = swap32_S (ctx.h[3]); + hooks[gid].ukey[4] = swap32_S (ctx.h[4]); + hooks[gid].ukey[5] = swap32_S (ctx.h[5]); + hooks[gid].ukey[6] = swap32_S (ctx.h[6]); + hooks[gid].ukey[7] = swap32_S (ctx.h[7]); } -__kernel void m11600_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global seven_zip_hook_t *seven_zip_hook, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 u64 gid_max) +__kernel void m11600_comp (KERN_ATTR_TMPS_HOOKS (seven_zip_tmp_t, seven_zip_hook_t)) { /** * base @@ -321,7 +321,7 @@ __kernel void m11600_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - if (seven_zip_hook[gid].hook_success == 1) + if (hooks[gid].hook_success == 1) { if (atomic_inc (&hashes_shown[digests_offset]) == 0) { diff --git a/OpenCL/m12000-pure.cl b/OpenCL/m12000-pure.cl index 1a02b57ff..e9ac20e59 100644 --- a/OpenCL/m12000-pure.cl +++ b/OpenCL/m12000-pure.cl @@ -52,7 +52,7 @@ DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipa sha1_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m12000_init (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, pbkdf2_sha1_t)) +__kernel void m12000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, pbkdf2_sha1_t)) { /** * base @@ -124,7 +124,7 @@ __kernel void m12000_init (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, pbkdf2 } } -__kernel void m12000_loop (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, pbkdf2_sha1_t)) +__kernel void m12000_loop (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, pbkdf2_sha1_t)) { const u64 gid = get_global_id (0); @@ -209,7 +209,7 @@ __kernel void m12000_loop (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, pbkdf2 } } -__kernel void m12000_comp (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, pbkdf2_sha1_t)) +__kernel void m12000_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, pbkdf2_sha1_t)) { /** * base diff --git a/OpenCL/m13600-pure.cl b/OpenCL/m13600-pure.cl index 8709434cf..cb4f967d0 100644 --- a/OpenCL/m13600-pure.cl +++ b/OpenCL/m13600-pure.cl @@ -52,7 +52,7 @@ DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipa sha1_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m13600_init (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, zip2_t)) +__kernel void m13600_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, zip2_t)) { /** * base @@ -165,7 +165,7 @@ __kernel void m13600_init (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, zip2_t } } -__kernel void m13600_loop (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, zip2_t)) +__kernel void m13600_loop (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, zip2_t)) { const u64 gid = get_global_id (0); @@ -274,7 +274,7 @@ __kernel void m13600_loop (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, zip2_t } } -__kernel void m13600_comp (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, zip2_t)) +__kernel void m13600_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, zip2_t)) { /** * base diff --git a/OpenCL/m14700-pure.cl b/OpenCL/m14700-pure.cl index b24626998..26d386245 100644 --- a/OpenCL/m14700-pure.cl +++ b/OpenCL/m14700-pure.cl @@ -53,7 +53,7 @@ DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipa sha1_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14700_init (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, itunes_backup_t)) +__kernel void m14700_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, itunes_backup_t)) { /** * base @@ -125,7 +125,7 @@ __kernel void m14700_init (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, itunes } } -__kernel void m14700_loop (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, itunes_backup_t)) +__kernel void m14700_loop (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, itunes_backup_t)) { const u64 gid = get_global_id (0); @@ -210,7 +210,7 @@ __kernel void m14700_loop (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, itunes } } -__kernel void m14700_comp (KERN_ATTR_TMPS_ESALT (const pbkdf2_sha1_tmp_t, itunes_backup_t)) +__kernel void m14700_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, itunes_backup_t)) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); diff --git a/OpenCL/m15700-pure.cl b/OpenCL/m15700-pure.cl index 1a34fbd74..8b2ff0b99 100644 --- a/OpenCL/m15700-pure.cl +++ b/OpenCL/m15700-pure.cl @@ -336,7 +336,7 @@ DECLSPEC void keccak_transform_S (u64 *st) } } -__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global scrypt_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const ethereum_scrypt_t *esalt_bufs, __global u32 *d_return_buf, __global uint4 *d_scryptV0_buf, __global uint4 *d_scryptV1_buf, __global uint4 *d_scryptV2_buf, __global uint4 *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 u64 gid_max) +__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_init (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t)) { /** * base @@ -401,7 +401,7 @@ __kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_init (__glob } } -__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global scrypt_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const ethereum_scrypt_t *esalt_bufs, __global u32 *d_return_buf, __global uint4 *d_scryptV0_buf, __global uint4 *d_scryptV1_buf, __global uint4 *d_scryptV2_buf, __global uint4 *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 u64 gid_max) +__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_loop (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t)) { const u64 gid = get_global_id (0); @@ -434,7 +434,7 @@ __kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_loop (__glob #endif } -__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global scrypt_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const ethereum_scrypt_t *esalt_bufs, __global u32 *d_return_buf, __global uint4 *d_scryptV0_buf, __global uint4 *d_scryptV1_buf, __global uint4 *d_scryptV2_buf, __global uint4 *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 u64 gid_max) +__kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_comp (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_t)) { /** * base