2 * Author......: Jens Steube <jens.steube@gmail.com>
8 //too much register pressure
9 //#define NEW_SIMD_CODE
11 #include "inc_vendor.cl"
12 #include "inc_hash_constants.h"
13 #include "inc_hash_functions.cl"
14 #include "inc_types.cl"
15 #include "inc_common.cl"
16 #include "inc_simd.cl"
18 #define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
19 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
21 #define SETSHIFTEDINT(a,n,v) \
23 const u32 s = ((n) & 3) * 8; \
24 const u64 x = (u64) (v) << s; \
25 (a)[((n)/4)+0] &= ~(0xff << ((n & 3) * 8)); \
26 (a)[((n)/4)+0] |= x; \
27 (a)[((n)/4)+1] = x >> 32; \
30 __constant u32 sapb_trans_tbl[256] =
32 // first value hack for 0 byte as part of an optimization
33 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
34 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
35 0x3f, 0x40, 0x41, 0x50, 0x43, 0x44, 0x45, 0x4b, 0x47, 0x48, 0x4d, 0x4e, 0x54, 0x51, 0x53, 0x46,
36 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x3b, 0x3c, 0x3d, 0x3e, 0x56, 0x55, 0x5c, 0x49, 0x5d, 0x4a,
37 0x42, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
38 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x58, 0x5b, 0x59, 0xff, 0x52,
39 0x4c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
40 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x57, 0x5e, 0x5a, 0x4f, 0xff,
41 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
42 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
43 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
44 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
45 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
46 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
47 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
48 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
51 __constant u32 bcodeArray[48] =
53 0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
54 0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
55 0xe2, 0xb7, 0x33, 0x71, 0x8b, 0x9f, 0x5d, 0x01, 0x44, 0x70, 0xae, 0x11, 0xef, 0x28, 0xf0, 0x0d
58 u32 sapb_trans (const u32 in)
62 out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
63 out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
64 out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
65 out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
70 u32 walld0rf_magic (const u32 w0[4], const u32 pw_len, const u32 salt_buf0[4], const u32 salt_len, const u32 a, const u32 b, const u32 c, const u32 d, u32 t[16])
89 u32 sum20 = ((a >> 24) & 3)
97 const u32 w[2] = { w0[0], w0[1] };
99 const u32 s[3] = { salt_buf0[0], salt_buf0[1], salt_buf0[2] };
101 u32 saved_key[4] = { a, b, c, d };
107 // we can assume this because the password must be at least 3
108 // and the username must be at least 1 so we can save the if ()
114 t0 |= bcodeArray[47] << 0;
115 t0 |= (w[0] & 0xff) << 8;
116 t0 |= (s[0] & 0xff) << 16;
117 t0 |= bcodeArray[ 1] << 24;
125 t0 |= (w[0] & 0xff) << 0;
126 t0 |= (s[0] & 0xff) << 8;
127 t0 |= bcodeArray[ 0] << 16;
136 // because the following code can increase i2 by a maximum of 5,
137 // there is an overflow potential of 4 before it comes to the next test for i2 >= sum20
138 // we need to truncate in that case
140 while ((i1 < pw_len) && (i3 < salt_len))
146 if (GETCHAR (saved_key, 15 - i1) & 1)
148 x0 |= bcodeArray[48 - 1 - i1] << 0; i2++;
149 x0 |= GETCHAR (w, i1) << 8; i2++; i1++;
150 x0 |= GETCHAR (s, i3) << 16; i2++; i3++;
151 x0 |= bcodeArray[i2 - i1 - i3] << 24; i2++; i2++;
155 x0 |= GETCHAR (w, i1) << 0; i2++; i1++;
156 x0 |= GETCHAR (s, i3) << 8; i2++; i3++;
157 x0 |= bcodeArray[i2 - i1 - i3] << 16; i2++; i2++;
160 SETSHIFTEDINT (t, i2_sav, x0);
168 while ((i1 < pw_len) || (i3 < salt_len))
170 if (i1 < pw_len) // max 8
172 if (GETCHAR (saved_key, 15 - i1) & 1)
174 PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
179 PUTCHAR (t, i2, GETCHAR (w, i1));
186 PUTCHAR (t, i2, GETCHAR (s, i3));
192 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
205 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
214 __kernel void m07700_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)
220 const u32 lid = get_local_id (0);
226 const u32 gid = get_global_id (0);
228 if (gid >= gid_max) return;
233 pw_buf0[0] = pws[gid].i[0];
234 pw_buf0[1] = pws[gid].i[1];
235 pw_buf0[2] = pws[gid].i[2];
236 pw_buf0[3] = pws[gid].i[3];
237 pw_buf1[0] = pws[gid].i[4];
238 pw_buf1[1] = pws[gid].i[5];
239 pw_buf1[2] = pws[gid].i[6];
240 pw_buf1[3] = pws[gid].i[7];
242 const u32 pw_l_len = pws[gid].pw_len;
250 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
251 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
252 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
255 const u32 salt_len = salt_bufs[salt_pos].salt_len;
257 salt_buf0[0] = sapb_trans (salt_buf0[0]);
258 salt_buf0[1] = sapb_trans (salt_buf0[1]);
259 salt_buf0[2] = sapb_trans (salt_buf0[2]);
265 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
267 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
269 const u32x pw_len = pw_l_len + pw_r_len;
272 * concat password candidate
275 u32x wordl0[4] = { 0 };
276 u32x wordl1[4] = { 0 };
277 u32x wordl2[4] = { 0 };
278 u32x wordl3[4] = { 0 };
280 wordl0[0] = pw_buf0[0];
281 wordl0[1] = pw_buf0[1];
282 wordl0[2] = pw_buf0[2];
283 wordl0[3] = pw_buf0[3];
284 wordl1[0] = pw_buf1[0];
285 wordl1[1] = pw_buf1[1];
286 wordl1[2] = pw_buf1[2];
287 wordl1[3] = pw_buf1[3];
289 u32x wordr0[4] = { 0 };
290 u32x wordr1[4] = { 0 };
291 u32x wordr2[4] = { 0 };
292 u32x wordr3[4] = { 0 };
294 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
295 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
296 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
297 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
298 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
299 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
300 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
301 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
303 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
305 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
309 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
314 w0[0] = wordl0[0] | wordr0[0];
315 w0[1] = wordl0[1] | wordr0[1];
317 if (pw_len > 8) continue; // otherwise it overflows in waldorf function
323 w0[0] = sapb_trans (w0[0]);
324 w0[1] = sapb_trans (w0[1]);
335 s0[0] = salt_buf0[0];
336 s0[1] = salt_buf0[1];
337 s0[2] = salt_buf0[2];
352 switch_buffer_by_offset_le (s0, s1, s2, s3, pw_len);
354 const u32 pw_salt_len = pw_len + salt_len;
358 t[ 0] = s0[0] | w0[0];
359 t[ 1] = s0[1] | w0[1];
372 t[14] = pw_salt_len * 8;
375 PUTCHAR (t, pw_salt_len, 0x80);
386 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
387 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
388 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
389 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
390 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
391 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
392 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
393 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
394 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
395 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
396 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
397 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
398 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
399 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
400 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
401 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
403 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
404 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
405 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
406 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
407 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
408 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
409 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
410 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
411 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
412 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
413 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
414 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
415 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
416 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
417 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
418 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
420 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
421 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
422 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
423 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
424 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
425 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
426 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
427 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
428 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
429 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
430 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
431 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
432 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
433 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
434 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
435 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
437 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
438 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
439 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
440 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
441 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
442 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
443 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
444 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
445 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
446 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
447 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
448 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
449 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
450 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
451 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
452 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
459 const u32 sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
461 SETSHIFTEDINT (t, sum20, 0x80);
470 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
471 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
472 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
473 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
474 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
475 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
476 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
477 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
478 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
479 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
480 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
481 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
482 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
483 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
484 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
485 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
487 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
488 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
489 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
490 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
491 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
492 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
493 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
494 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
495 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
496 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
497 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
498 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
499 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
500 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
501 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
502 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
504 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
505 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
506 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
507 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
508 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
509 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
510 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
511 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
512 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
513 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
514 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
515 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
516 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
517 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
518 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
519 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
521 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
522 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
523 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
524 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
525 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
526 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
527 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
528 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
529 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
530 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
531 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
532 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
533 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
534 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
535 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
536 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
548 COMPARE_M_SIMD (a, b, c, d);
552 __kernel void m07700_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)
556 __kernel void m07700_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)
560 __kernel void m07700_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)
566 const u32 lid = get_local_id (0);
572 const u32 gid = get_global_id (0);
574 if (gid >= gid_max) return;
579 pw_buf0[0] = pws[gid].i[0];
580 pw_buf0[1] = pws[gid].i[1];
581 pw_buf0[2] = pws[gid].i[2];
582 pw_buf0[3] = pws[gid].i[3];
583 pw_buf1[0] = pws[gid].i[4];
584 pw_buf1[1] = pws[gid].i[5];
585 pw_buf1[2] = pws[gid].i[6];
586 pw_buf1[3] = pws[gid].i[7];
588 const u32 pw_l_len = pws[gid].pw_len;
596 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
597 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
598 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
601 const u32 salt_len = salt_bufs[salt_pos].salt_len;
603 salt_buf0[0] = sapb_trans (salt_buf0[0]);
604 salt_buf0[1] = sapb_trans (salt_buf0[1]);
605 salt_buf0[2] = sapb_trans (salt_buf0[2]);
611 const u32 search[4] =
613 digests_buf[digests_offset].digest_buf[DGST_R0],
614 digests_buf[digests_offset].digest_buf[DGST_R1],
623 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
625 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
627 const u32x pw_len = pw_l_len + pw_r_len;
630 * concat password candidate
633 u32x wordl0[4] = { 0 };
634 u32x wordl1[4] = { 0 };
635 u32x wordl2[4] = { 0 };
636 u32x wordl3[4] = { 0 };
638 wordl0[0] = pw_buf0[0];
639 wordl0[1] = pw_buf0[1];
640 wordl0[2] = pw_buf0[2];
641 wordl0[3] = pw_buf0[3];
642 wordl1[0] = pw_buf1[0];
643 wordl1[1] = pw_buf1[1];
644 wordl1[2] = pw_buf1[2];
645 wordl1[3] = pw_buf1[3];
647 u32x wordr0[4] = { 0 };
648 u32x wordr1[4] = { 0 };
649 u32x wordr2[4] = { 0 };
650 u32x wordr3[4] = { 0 };
652 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
653 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
654 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
655 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
656 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
657 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
658 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
659 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
661 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
663 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
667 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
672 w0[0] = wordl0[0] | wordr0[0];
673 w0[1] = wordl0[1] | wordr0[1];
675 if (pw_len > 8) continue; // otherwise it overflows in waldorf function
681 w0[0] = sapb_trans (w0[0]);
682 w0[1] = sapb_trans (w0[1]);
693 s0[0] = salt_buf0[0];
694 s0[1] = salt_buf0[1];
695 s0[2] = salt_buf0[2];
710 switch_buffer_by_offset_le (s0, s1, s2, s3, pw_len);
712 const u32 pw_salt_len = pw_len + salt_len;
716 t[ 0] = s0[0] | w0[0];
717 t[ 1] = s0[1] | w0[1];
730 t[14] = pw_salt_len * 8;
733 PUTCHAR (t, pw_salt_len, 0x80);
744 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
745 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
746 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
747 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
748 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
749 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
750 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
751 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
752 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
753 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
754 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
755 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
756 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
757 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
758 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
759 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
761 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
762 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
763 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
764 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
765 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
766 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
767 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
768 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
769 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
770 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
771 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
772 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
773 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
774 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
775 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
776 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
778 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
779 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
780 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
781 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
782 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
783 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
784 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
785 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
786 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
787 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
788 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
789 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
790 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
791 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
792 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
793 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
795 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
796 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
797 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
798 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
799 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
800 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
801 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
802 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
803 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
804 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
805 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
806 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
807 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
808 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
809 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
810 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
817 const u32 sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
819 SETSHIFTEDINT (t, sum20, 0x80);
828 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
829 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
830 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
831 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
832 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
833 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
834 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
835 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
836 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
837 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
838 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
839 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
840 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
841 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
842 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
843 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
845 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
846 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
847 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
848 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
849 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
850 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
851 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
852 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
853 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
854 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
855 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
856 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
857 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
858 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
859 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
860 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
862 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
863 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
864 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
865 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
866 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
867 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
868 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
869 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
870 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
871 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
872 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
873 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
874 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
875 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
876 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
877 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
879 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
880 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
881 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
882 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
883 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
884 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
885 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
886 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
887 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
888 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
889 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
890 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
891 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
892 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
893 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
894 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
906 COMPARE_S_SIMD (a, b, c, d);
910 __kernel void m07700_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)
914 __kernel void m07700_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)