2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
28 #include "include/kernel_functions.c"
29 #include "types_amd.c"
30 #include "common_amd.c"
33 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
34 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
38 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
39 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
43 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
44 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
47 static void m04900m (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
53 const u32 gid = get_global_id (0);
54 const u32 lid = get_local_id (0);
62 salt_buf0_t[0] = salt_bufs[salt_pos].salt_buf[0];
63 salt_buf0_t[1] = salt_bufs[salt_pos].salt_buf[1];
64 salt_buf0_t[2] = salt_bufs[salt_pos].salt_buf[2];
65 salt_buf0_t[3] = salt_bufs[salt_pos].salt_buf[3];
69 salt_buf1_t[0] = salt_bufs[salt_pos].salt_buf[4];
70 salt_buf1_t[1] = salt_bufs[salt_pos].salt_buf[5];
71 salt_buf1_t[2] = salt_bufs[salt_pos].salt_buf[6];
72 salt_buf1_t[3] = salt_bufs[salt_pos].salt_buf[7];
88 const u32 salt_len = salt_bufs[salt_pos].salt_len;
90 const u32 pw_salt_len = salt_len + pw_len + salt_len;
92 // first we need to switch the right-hand salt to the correct position (2nd salt)
94 switch_buffer_by_offset (salt_buf0_t, salt_buf1_t, salt_buf2_t, salt_buf3_t, salt_len + pw_len);
98 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
99 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
100 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
101 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
105 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
106 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
107 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
108 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
124 // concatenate the 1st and 2nd instance of the salt
126 salt_buf0[0] |= salt_buf0_t[0];
127 salt_buf0[1] |= salt_buf0_t[1];
128 salt_buf0[2] |= salt_buf0_t[2];
129 salt_buf0[3] |= salt_buf0_t[3];
131 salt_buf1[0] |= salt_buf1_t[0];
132 salt_buf1[1] |= salt_buf1_t[1];
133 salt_buf1[2] |= salt_buf1_t[2];
134 salt_buf1[3] |= salt_buf1_t[3];
136 salt_buf2[0] |= salt_buf2_t[0];
137 salt_buf2[1] |= salt_buf2_t[1];
138 salt_buf2[2] |= salt_buf2_t[2];
139 salt_buf2[3] |= salt_buf2_t[3];
141 salt_buf3[0] |= salt_buf3_t[0];
142 salt_buf3[1] |= salt_buf3_t[1];
143 salt_buf3[2] |= salt_buf3_t[2];
144 salt_buf3[3] |= salt_buf3_t[3];
146 append_0x80_4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_salt_len);
154 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
156 const u32 w0r = bfs_buf[il_pos].i;
189 * put the password after the first salt but before the second salt
192 switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
194 w0_t[0] |= salt_buf0[0];
195 w0_t[1] |= salt_buf0[1];
196 w0_t[2] |= salt_buf0[2];
197 w0_t[3] |= salt_buf0[3];
198 w1_t[0] |= salt_buf1[0];
199 w1_t[1] |= salt_buf1[1];
200 w1_t[2] |= salt_buf1[2];
201 w1_t[3] |= salt_buf1[3];
202 w2_t[0] |= salt_buf2[0];
203 w2_t[1] |= salt_buf2[1];
204 w2_t[2] |= salt_buf2[2];
205 w2_t[3] |= salt_buf2[3];
206 w3_t[0] |= salt_buf3[0];
207 w3_t[1] |= salt_buf3[1];
208 w3_t[2] |= salt_buf3[2];
210 u32x w0 = swap_workaround (w0_t[0]);
211 u32x w1 = swap_workaround (w0_t[1]);
212 u32x w2 = swap_workaround (w0_t[2]);
213 u32x w3 = swap_workaround (w0_t[3]);
214 u32x w4 = swap_workaround (w1_t[0]);
215 u32x w5 = swap_workaround (w1_t[1]);
216 u32x w6 = swap_workaround (w1_t[2]);
217 u32x w7 = swap_workaround (w1_t[3]);
218 u32x w8 = swap_workaround (w2_t[0]);
219 u32x w9 = swap_workaround (w2_t[1]);
220 u32x wa = swap_workaround (w2_t[2]);
221 u32x wb = swap_workaround (w2_t[3]);
222 u32x wc = swap_workaround (w3_t[0]);
223 u32x wd = swap_workaround (w3_t[1]);
224 u32x we = swap_workaround (w3_t[2]);
225 u32x wf = pw_salt_len * 8;
240 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0);
241 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1);
242 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2);
243 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3);
244 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4);
245 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5);
246 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6);
247 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7);
248 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8);
249 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9);
250 SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa);
251 SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb);
252 SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc);
253 SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd);
254 SHA1_STEP (SHA1_F0o, b, c, d, e, a, we);
255 SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf);
256 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0);
257 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1);
258 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2);
259 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3);
264 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4);
265 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5);
266 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6);
267 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7);
268 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8);
269 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9);
270 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa);
271 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb);
272 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc);
273 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd);
274 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we);
275 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf);
276 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0);
277 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1);
278 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2);
279 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3);
280 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4);
281 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5);
282 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6);
283 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7);
288 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8);
289 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9);
290 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa);
291 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb);
292 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc);
293 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd);
294 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we);
295 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf);
296 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0);
297 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1);
298 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2);
299 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3);
300 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4);
301 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5);
302 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6);
303 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7);
304 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8);
305 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9);
306 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa);
307 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb);
312 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc);
313 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd);
314 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we);
315 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf);
316 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0);
317 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1);
318 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2);
319 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3);
320 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4);
321 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5);
322 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6);
323 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7);
324 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8);
325 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9);
326 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa);
327 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb);
328 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc);
329 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd);
330 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we);
331 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf);
338 #include VECT_COMPARE_M
342 static void m04900s (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
348 const u32 gid = get_global_id (0);
349 const u32 lid = get_local_id (0);
355 const u32 search[4] =
357 digests_buf[digests_offset].digest_buf[DGST_R0],
358 digests_buf[digests_offset].digest_buf[DGST_R1],
359 digests_buf[digests_offset].digest_buf[DGST_R2],
360 digests_buf[digests_offset].digest_buf[DGST_R3]
367 const u32 e_rev = rotl32 (search[1], 2u);
375 salt_buf0_t[0] = salt_bufs[salt_pos].salt_buf[0];
376 salt_buf0_t[1] = salt_bufs[salt_pos].salt_buf[1];
377 salt_buf0_t[2] = salt_bufs[salt_pos].salt_buf[2];
378 salt_buf0_t[3] = salt_bufs[salt_pos].salt_buf[3];
382 salt_buf1_t[0] = salt_bufs[salt_pos].salt_buf[4];
383 salt_buf1_t[1] = salt_bufs[salt_pos].salt_buf[5];
384 salt_buf1_t[2] = salt_bufs[salt_pos].salt_buf[6];
385 salt_buf1_t[3] = salt_bufs[salt_pos].salt_buf[7];
401 const u32 salt_len = salt_bufs[salt_pos].salt_len;
403 const u32 pw_salt_len = salt_len + pw_len + salt_len;
405 // first we need to switch the right-hand salt to the correct position (2nd salt)
407 switch_buffer_by_offset (salt_buf0_t, salt_buf1_t, salt_buf2_t, salt_buf3_t, salt_len + pw_len);
411 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
412 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
413 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
414 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
418 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
419 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
420 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
421 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
437 // concatenate the 1st and 2nd instance of the salt
439 salt_buf0[0] |= salt_buf0_t[0];
440 salt_buf0[1] |= salt_buf0_t[1];
441 salt_buf0[2] |= salt_buf0_t[2];
442 salt_buf0[3] |= salt_buf0_t[3];
444 salt_buf1[0] |= salt_buf1_t[0];
445 salt_buf1[1] |= salt_buf1_t[1];
446 salt_buf1[2] |= salt_buf1_t[2];
447 salt_buf1[3] |= salt_buf1_t[3];
449 salt_buf2[0] |= salt_buf2_t[0];
450 salt_buf2[1] |= salt_buf2_t[1];
451 salt_buf2[2] |= salt_buf2_t[2];
452 salt_buf2[3] |= salt_buf2_t[3];
454 salt_buf3[0] |= salt_buf3_t[0];
455 salt_buf3[1] |= salt_buf3_t[1];
456 salt_buf3[2] |= salt_buf3_t[2];
457 salt_buf3[3] |= salt_buf3_t[3];
459 append_0x80_4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_salt_len);
467 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
469 const u32 w0r = bfs_buf[il_pos].i;
502 * put the password after the first salt but before the second salt
505 switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
507 w0_t[0] |= salt_buf0[0];
508 w0_t[1] |= salt_buf0[1];
509 w0_t[2] |= salt_buf0[2];
510 w0_t[3] |= salt_buf0[3];
511 w1_t[0] |= salt_buf1[0];
512 w1_t[1] |= salt_buf1[1];
513 w1_t[2] |= salt_buf1[2];
514 w1_t[3] |= salt_buf1[3];
515 w2_t[0] |= salt_buf2[0];
516 w2_t[1] |= salt_buf2[1];
517 w2_t[2] |= salt_buf2[2];
518 w2_t[3] |= salt_buf2[3];
519 w3_t[0] |= salt_buf3[0];
520 w3_t[1] |= salt_buf3[1];
521 w3_t[2] |= salt_buf3[2];
523 u32x w0 = swap_workaround (w0_t[0]);
524 u32x w1 = swap_workaround (w0_t[1]);
525 u32x w2 = swap_workaround (w0_t[2]);
526 u32x w3 = swap_workaround (w0_t[3]);
527 u32x w4 = swap_workaround (w1_t[0]);
528 u32x w5 = swap_workaround (w1_t[1]);
529 u32x w6 = swap_workaround (w1_t[2]);
530 u32x w7 = swap_workaround (w1_t[3]);
531 u32x w8 = swap_workaround (w2_t[0]);
532 u32x w9 = swap_workaround (w2_t[1]);
533 u32x wa = swap_workaround (w2_t[2]);
534 u32x wb = swap_workaround (w2_t[3]);
535 u32x wc = swap_workaround (w3_t[0]);
536 u32x wd = swap_workaround (w3_t[1]);
537 u32x we = swap_workaround (w3_t[2]);
538 u32x wf = pw_salt_len * 8;
553 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0);
554 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1);
555 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2);
556 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3);
557 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4);
558 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5);
559 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6);
560 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7);
561 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8);
562 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9);
563 SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa);
564 SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb);
565 SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc);
566 SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd);
567 SHA1_STEP (SHA1_F0o, b, c, d, e, a, we);
568 SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf);
569 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0);
570 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1);
571 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2);
572 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3);
577 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4);
578 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5);
579 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6);
580 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7);
581 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8);
582 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9);
583 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa);
584 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb);
585 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc);
586 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd);
587 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we);
588 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf);
589 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0);
590 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1);
591 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2);
592 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3);
593 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4);
594 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5);
595 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6);
596 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7);
601 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8);
602 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9);
603 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa);
604 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb);
605 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc);
606 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd);
607 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we);
608 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf);
609 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0);
610 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1);
611 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2);
612 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3);
613 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4);
614 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5);
615 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6);
616 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7);
617 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8);
618 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9);
619 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa);
620 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb);
625 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc);
626 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd);
627 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we);
628 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf);
629 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0);
630 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1);
631 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2);
632 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3);
633 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4);
634 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5);
635 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6);
636 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7);
637 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8);
638 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9);
639 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa);
640 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb);
642 if (allx (e != e_rev)) continue;
644 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc);
645 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd);
646 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we);
647 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf);
654 #include VECT_COMPARE_S
658 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
664 const u32 gid = get_global_id (0);
666 if (gid >= gid_max) return;
672 const u32 lid = get_local_id (0);
676 w0[0] = pws[gid].i[ 0];
677 w0[1] = pws[gid].i[ 1];
678 w0[2] = pws[gid].i[ 2];
679 w0[3] = pws[gid].i[ 3];
699 w3[2] = pws[gid].i[14];
702 const u32 pw_len = pws[gid].pw_len;
708 m04900m (w0, w1, w2, w3, 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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
711 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
717 const u32 gid = get_global_id (0);
719 if (gid >= gid_max) return;
725 const u32 lid = get_local_id (0);
729 w0[0] = pws[gid].i[ 0];
730 w0[1] = pws[gid].i[ 1];
731 w0[2] = pws[gid].i[ 2];
732 w0[3] = pws[gid].i[ 3];
736 w1[0] = pws[gid].i[ 4];
737 w1[1] = pws[gid].i[ 5];
738 w1[2] = pws[gid].i[ 6];
739 w1[3] = pws[gid].i[ 7];
752 w3[2] = pws[gid].i[14];
755 const u32 pw_len = pws[gid].pw_len;
761 m04900m (w0, w1, w2, w3, 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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
764 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
770 const u32 gid = get_global_id (0);
772 if (gid >= gid_max) return;
778 const u32 lid = get_local_id (0);
782 w0[0] = pws[gid].i[ 0];
783 w0[1] = pws[gid].i[ 1];
784 w0[2] = pws[gid].i[ 2];
785 w0[3] = pws[gid].i[ 3];
789 w1[0] = pws[gid].i[ 4];
790 w1[1] = pws[gid].i[ 5];
791 w1[2] = pws[gid].i[ 6];
792 w1[3] = pws[gid].i[ 7];
796 w2[0] = pws[gid].i[ 8];
797 w2[1] = pws[gid].i[ 9];
798 w2[2] = pws[gid].i[10];
799 w2[3] = pws[gid].i[11];
803 w3[0] = pws[gid].i[12];
804 w3[1] = pws[gid].i[13];
805 w3[2] = pws[gid].i[14];
806 w3[3] = pws[gid].i[15];
808 const u32 pw_len = pws[gid].pw_len;
814 m04900m (w0, w1, w2, w3, 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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
817 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
823 const u32 gid = get_global_id (0);
825 if (gid >= gid_max) return;
831 const u32 lid = get_local_id (0);
835 w0[0] = pws[gid].i[ 0];
836 w0[1] = pws[gid].i[ 1];
837 w0[2] = pws[gid].i[ 2];
838 w0[3] = pws[gid].i[ 3];
858 w3[2] = pws[gid].i[14];
861 const u32 pw_len = pws[gid].pw_len;
867 m04900s (w0, w1, w2, w3, 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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
870 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
876 const u32 gid = get_global_id (0);
878 if (gid >= gid_max) return;
884 const u32 lid = get_local_id (0);
888 w0[0] = pws[gid].i[ 0];
889 w0[1] = pws[gid].i[ 1];
890 w0[2] = pws[gid].i[ 2];
891 w0[3] = pws[gid].i[ 3];
895 w1[0] = pws[gid].i[ 4];
896 w1[1] = pws[gid].i[ 5];
897 w1[2] = pws[gid].i[ 6];
898 w1[3] = pws[gid].i[ 7];
911 w3[2] = pws[gid].i[14];
914 const u32 pw_len = pws[gid].pw_len;
920 m04900s (w0, w1, w2, w3, 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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
923 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
929 const u32 gid = get_global_id (0);
931 if (gid >= gid_max) return;
937 const u32 lid = get_local_id (0);
941 w0[0] = pws[gid].i[ 0];
942 w0[1] = pws[gid].i[ 1];
943 w0[2] = pws[gid].i[ 2];
944 w0[3] = pws[gid].i[ 3];
948 w1[0] = pws[gid].i[ 4];
949 w1[1] = pws[gid].i[ 5];
950 w1[2] = pws[gid].i[ 6];
951 w1[3] = pws[gid].i[ 7];
955 w2[0] = pws[gid].i[ 8];
956 w2[1] = pws[gid].i[ 9];
957 w2[2] = pws[gid].i[10];
958 w2[3] = pws[gid].i[11];
962 w3[0] = pws[gid].i[12];
963 w3[1] = pws[gid].i[13];
964 w3[2] = pws[gid].i[14];
965 w3[3] = pws[gid].i[15];
967 const u32 pw_len = pws[gid].pw_len;
973 m04900s (w0, w1, w2, w3, 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_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);