2 * Author......: Jens Steube <jens.steube@gmail.com>
9 #include "include/constants.h"
10 #include "include/kernel_vendor.h"
29 #include "include/kernel_functions.c"
30 #include "types_amd.c"
31 #include "common_amd.c"
34 #define VECT_COMPARE_S "check_single_vect1_comp4_warp.c"
35 #define VECT_COMPARE_M "check_multi_vect1_comp4_warp.c"
39 #define VECT_COMPARE_S "check_single_vect2_comp4_warp.c"
40 #define VECT_COMPARE_M "check_multi_vect2_comp4_warp.c"
44 #define VECT_COMPARE_S "check_single_vect4_comp4_warp.c"
45 #define VECT_COMPARE_M "check_multi_vect4_comp4_warp.c"
48 __constant u32 crc32tab[0x100] =
50 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
51 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
52 0x0edb8832, 0x79dcb8a4, 0xe0d5e91e, 0x97d2d988,
53 0x09b64c2b, 0x7eb17cbd, 0xe7b82d07, 0x90bf1d91,
54 0x1db71064, 0x6ab020f2, 0xf3b97148, 0x84be41de,
55 0x1adad47d, 0x6ddde4eb, 0xf4d4b551, 0x83d385c7,
56 0x136c9856, 0x646ba8c0, 0xfd62f97a, 0x8a65c9ec,
57 0x14015c4f, 0x63066cd9, 0xfa0f3d63, 0x8d080df5,
58 0x3b6e20c8, 0x4c69105e, 0xd56041e4, 0xa2677172,
59 0x3c03e4d1, 0x4b04d447, 0xd20d85fd, 0xa50ab56b,
60 0x35b5a8fa, 0x42b2986c, 0xdbbbc9d6, 0xacbcf940,
61 0x32d86ce3, 0x45df5c75, 0xdcd60dcf, 0xabd13d59,
62 0x26d930ac, 0x51de003a, 0xc8d75180, 0xbfd06116,
63 0x21b4f4b5, 0x56b3c423, 0xcfba9599, 0xb8bda50f,
64 0x2802b89e, 0x5f058808, 0xc60cd9b2, 0xb10be924,
65 0x2f6f7c87, 0x58684c11, 0xc1611dab, 0xb6662d3d,
66 0x76dc4190, 0x01db7106, 0x98d220bc, 0xefd5102a,
67 0x71b18589, 0x06b6b51f, 0x9fbfe4a5, 0xe8b8d433,
68 0x7807c9a2, 0x0f00f934, 0x9609a88e, 0xe10e9818,
69 0x7f6a0dbb, 0x086d3d2d, 0x91646c97, 0xe6635c01,
70 0x6b6b51f4, 0x1c6c6162, 0x856530d8, 0xf262004e,
71 0x6c0695ed, 0x1b01a57b, 0x8208f4c1, 0xf50fc457,
72 0x65b0d9c6, 0x12b7e950, 0x8bbeb8ea, 0xfcb9887c,
73 0x62dd1ddf, 0x15da2d49, 0x8cd37cf3, 0xfbd44c65,
74 0x4db26158, 0x3ab551ce, 0xa3bc0074, 0xd4bb30e2,
75 0x4adfa541, 0x3dd895d7, 0xa4d1c46d, 0xd3d6f4fb,
76 0x4369e96a, 0x346ed9fc, 0xad678846, 0xda60b8d0,
77 0x44042d73, 0x33031de5, 0xaa0a4c5f, 0xdd0d7cc9,
78 0x5005713c, 0x270241aa, 0xbe0b1010, 0xc90c2086,
79 0x5768b525, 0x206f85b3, 0xb966d409, 0xce61e49f,
80 0x5edef90e, 0x29d9c998, 0xb0d09822, 0xc7d7a8b4,
81 0x59b33d17, 0x2eb40d81, 0xb7bd5c3b, 0xc0ba6cad,
82 0xedb88320, 0x9abfb3b6, 0x03b6e20c, 0x74b1d29a,
83 0xead54739, 0x9dd277af, 0x04db2615, 0x73dc1683,
84 0xe3630b12, 0x94643b84, 0x0d6d6a3e, 0x7a6a5aa8,
85 0xe40ecf0b, 0x9309ff9d, 0x0a00ae27, 0x7d079eb1,
86 0xf00f9344, 0x8708a3d2, 0x1e01f268, 0x6906c2fe,
87 0xf762575d, 0x806567cb, 0x196c3671, 0x6e6b06e7,
88 0xfed41b76, 0x89d32be0, 0x10da7a5a, 0x67dd4acc,
89 0xf9b9df6f, 0x8ebeeff9, 0x17b7be43, 0x60b08ed5,
90 0xd6d6a3e8, 0xa1d1937e, 0x38d8c2c4, 0x4fdff252,
91 0xd1bb67f1, 0xa6bc5767, 0x3fb506dd, 0x48b2364b,
92 0xd80d2bda, 0xaf0a1b4c, 0x36034af6, 0x41047a60,
93 0xdf60efc3, 0xa867df55, 0x316e8eef, 0x4669be79,
94 0xcb61b38c, 0xbc66831a, 0x256fd2a0, 0x5268e236,
95 0xcc0c7795, 0xbb0b4703, 0x220216b9, 0x5505262f,
96 0xc5ba3bbe, 0xb2bd0b28, 0x2bb45a92, 0x5cb36a04,
97 0xc2d7ffa7, 0xb5d0cf31, 0x2cd99e8b, 0x5bdeae1d,
98 0x9b64c2b0, 0xec63f226, 0x756aa39c, 0x026d930a,
99 0x9c0906a9, 0xeb0e363f, 0x72076785, 0x05005713,
100 0x95bf4a82, 0xe2b87a14, 0x7bb12bae, 0x0cb61b38,
101 0x92d28e9b, 0xe5d5be0d, 0x7cdcefb7, 0x0bdbdf21,
102 0x86d3d2d4, 0xf1d4e242, 0x68ddb3f8, 0x1fda836e,
103 0x81be16cd, 0xf6b9265b, 0x6fb077e1, 0x18b74777,
104 0x88085ae6, 0xff0f6a70, 0x66063bca, 0x11010b5c,
105 0x8f659eff, 0xf862ae69, 0x616bffd3, 0x166ccf45,
106 0xa00ae278, 0xd70dd2ee, 0x4e048354, 0x3903b3c2,
107 0xa7672661, 0xd06016f7, 0x4969474d, 0x3e6e77db,
108 0xaed16a4a, 0xd9d65adc, 0x40df0b66, 0x37d83bf0,
109 0xa9bcae53, 0xdebb9ec5, 0x47b2cf7f, 0x30b5ffe9,
110 0xbdbdf21c, 0xcabac28a, 0x53b39330, 0x24b4a3a6,
111 0xbad03605, 0xcdd70693, 0x54de5729, 0x23d967bf,
112 0xb3667a2e, 0xc4614ab8, 0x5d681b02, 0x2a6f2b94,
113 0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d
116 static u32x round_crc32 (u32x a, const u32x v)
118 const u32x k = (a ^ v) & 0xff;
120 const u32x s = a >> 8;
127 a.s0 = crc32tab[k.s0];
128 a.s1 = crc32tab[k.s1];
132 a.s0 = crc32tab[k.s0];
133 a.s1 = crc32tab[k.s1];
134 a.s2 = crc32tab[k.s2];
135 a.s3 = crc32tab[k.s3];
143 static u32x crc32 (const u32x w[16], const u32 pw_len, const u32 iv)
147 if (pw_len >= 1) a = round_crc32 (a, w[0] >> 0);
148 if (pw_len >= 2) a = round_crc32 (a, w[0] >> 8);
149 if (pw_len >= 3) a = round_crc32 (a, w[0] >> 16);
150 if (pw_len >= 4) a = round_crc32 (a, w[0] >> 24);
152 for (u32 i = 4, j = 1; i < pw_len; i += 4, j += 1)
154 if (pw_len >= (i + 1)) a = round_crc32 (a, w[j] >> 0);
155 if (pw_len >= (i + 2)) a = round_crc32 (a, w[j] >> 8);
156 if (pw_len >= (i + 3)) a = round_crc32 (a, w[j] >> 16);
157 if (pw_len >= (i + 4)) a = round_crc32 (a, w[j] >> 24);
163 static void m11500m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
169 const u32 gid = get_global_id (0);
170 const u32 lid = get_local_id (0);
176 const u32 iv = salt_bufs[salt_pos].salt_buf[0];
182 const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
186 for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
188 const u32x w0r = words_buf_r[il_pos];
190 const u32x w0 = w0l | w0r;
211 u32x a = crc32 (w_t, pw_len, iv);
219 #include VECT_COMPARE_M
223 static void m11500s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
229 const u32 gid = get_global_id (0);
230 const u32 lid = get_local_id (0);
236 const u32 iv = salt_bufs[salt_pos].salt_buf[0];
238 const u32 search[4] =
240 digests_buf[digests_offset].digest_buf[DGST_R0],
241 digests_buf[digests_offset].digest_buf[DGST_R1],
242 digests_buf[digests_offset].digest_buf[DGST_R2],
243 digests_buf[digests_offset].digest_buf[DGST_R3]
250 const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
254 for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
256 const u32x w0r = words_buf_r[il_pos];
258 const u32x w0 = w0l | w0r;
279 u32x a = crc32 (w_t, pw_len, iv);
287 #include VECT_COMPARE_S
291 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11500_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
297 const u32 gid = get_global_id (0);
299 if (gid >= gid_max) return;
303 w[ 0] = pws[gid].i[ 0];
304 w[ 1] = pws[gid].i[ 1];
305 w[ 2] = pws[gid].i[ 2];
306 w[ 3] = pws[gid].i[ 3];
320 const u32 pw_len = pws[gid].pw_len;
326 m11500m (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
329 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11500_m08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
335 const u32 gid = get_global_id (0);
337 if (gid >= gid_max) return;
341 w[ 0] = pws[gid].i[ 0];
342 w[ 1] = pws[gid].i[ 1];
343 w[ 2] = pws[gid].i[ 2];
344 w[ 3] = pws[gid].i[ 3];
345 w[ 4] = pws[gid].i[ 4];
346 w[ 5] = pws[gid].i[ 5];
347 w[ 6] = pws[gid].i[ 6];
348 w[ 7] = pws[gid].i[ 7];
358 const u32 pw_len = pws[gid].pw_len;
364 m11500m (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
367 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11500_m16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
373 const u32 gid = get_global_id (0);
375 if (gid >= gid_max) return;
379 w[ 0] = pws[gid].i[ 0];
380 w[ 1] = pws[gid].i[ 1];
381 w[ 2] = pws[gid].i[ 2];
382 w[ 3] = pws[gid].i[ 3];
383 w[ 4] = pws[gid].i[ 4];
384 w[ 5] = pws[gid].i[ 5];
385 w[ 6] = pws[gid].i[ 6];
386 w[ 7] = pws[gid].i[ 7];
387 w[ 8] = pws[gid].i[ 8];
388 w[ 9] = pws[gid].i[ 9];
389 w[10] = pws[gid].i[10];
390 w[11] = pws[gid].i[11];
391 w[12] = pws[gid].i[12];
392 w[13] = pws[gid].i[13];
393 w[14] = pws[gid].i[14];
394 w[15] = pws[gid].i[15];
396 const u32 pw_len = pws[gid].pw_len;
402 m11500m (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
405 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11500_s04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
411 const u32 gid = get_global_id (0);
413 if (gid >= gid_max) return;
417 w[ 0] = pws[gid].i[ 0];
418 w[ 1] = pws[gid].i[ 1];
419 w[ 2] = pws[gid].i[ 2];
420 w[ 3] = pws[gid].i[ 3];
434 const u32 pw_len = pws[gid].pw_len;
440 m11500s (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
443 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11500_s08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
449 const u32 gid = get_global_id (0);
451 if (gid >= gid_max) return;
455 w[ 0] = pws[gid].i[ 0];
456 w[ 1] = pws[gid].i[ 1];
457 w[ 2] = pws[gid].i[ 2];
458 w[ 3] = pws[gid].i[ 3];
459 w[ 4] = pws[gid].i[ 4];
460 w[ 5] = pws[gid].i[ 5];
461 w[ 6] = pws[gid].i[ 6];
462 w[ 7] = pws[gid].i[ 7];
472 const u32 pw_len = pws[gid].pw_len;
478 m11500s (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
481 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11500_s16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
487 const u32 gid = get_global_id (0);
489 if (gid >= gid_max) return;
493 w[ 0] = pws[gid].i[ 0];
494 w[ 1] = pws[gid].i[ 1];
495 w[ 2] = pws[gid].i[ 2];
496 w[ 3] = pws[gid].i[ 3];
497 w[ 4] = pws[gid].i[ 4];
498 w[ 5] = pws[gid].i[ 5];
499 w[ 6] = pws[gid].i[ 6];
500 w[ 7] = pws[gid].i[ 7];
501 w[ 8] = pws[gid].i[ 8];
502 w[ 9] = pws[gid].i[ 9];
503 w[10] = pws[gid].i[10];
504 w[11] = pws[gid].i[11];
505 w[12] = pws[gid].i[12];
506 w[13] = pws[gid].i[13];
507 w[14] = pws[gid].i[14];
508 w[15] = pws[gid].i[15];
510 const u32 pw_len = pws[gid].pw_len;
516 m11500s (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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);