From 240f6298beb47635e8543dc32b2ce7d2b0b9c889 Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 17 Nov 2018 23:14:12 +0100 Subject: [PATCH] Fix some leftovers from switching kernel parameters to macros --- OpenCL/m14100_a3-pure.cl | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/OpenCL/m14100_a3-pure.cl b/OpenCL/m14100_a3-pure.cl index 29db2b70b..034f7416d 100644 --- a/OpenCL/m14100_a3-pure.cl +++ b/OpenCL/m14100_a3-pure.cl @@ -530,7 +530,7 @@ DECLSPEC void _des_crypt_keysetup (u32x c, u32x d, u32x *Kc, u32x *Kd, __local u } } -DECLSPEC void m14100m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ()) +DECLSPEC void m14100m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 *w, const u32 pw_len, KERN_ATTR_BASIC ()) { /** * modifier @@ -558,7 +558,7 @@ DECLSPEC void m14100m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0 = w0l | w0r; @@ -615,7 +615,7 @@ DECLSPEC void m14100m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u } } -DECLSPEC void m14100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ()) +DECLSPEC void m14100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 *w, const u32 pw_len, KERN_ATTR_BASIC ()) { /** * modifier @@ -655,7 +655,7 @@ DECLSPEC void m14100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + const u32x w0r = ix_create_bft (bfs_buf, il_pos); const u32x w0 = w0l | w0r; @@ -712,7 +712,7 @@ DECLSPEC void m14100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u } } -__kernel void m14100_mxx (KERN_ATTR_VECTOR ()) +__kernel void m14100_mxx (KERN_ATTR_BASIC ()) { /** * base @@ -783,10 +783,10 @@ __kernel void m14100_mxx (KERN_ATTR_VECTOR ()) * main */ - m14100m (s_SPtrans, s_skb, w, pw_len, 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); + m14100m (s_SPtrans, s_skb, w, 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, combs_mode, gid_max); } -__kernel void m14100_sxx (KERN_ATTR_VECTOR ()) +__kernel void m14100_sxx (KERN_ATTR_BASIC ()) { /** * base @@ -857,5 +857,5 @@ __kernel void m14100_sxx (KERN_ATTR_VECTOR ()) * main */ - m14100s (s_SPtrans, s_skb, w, pw_len, 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); + m14100s (s_SPtrans, s_skb, w, 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, combs_mode, gid_max); }