2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
16 #include "include/kernel_functions.c"
17 #include "types_ocl.c"
20 #define COMPARE_S "check_single_comp4.c"
21 #define COMPARE_M "check_multi_comp4.c"
23 #define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
24 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
26 #define SETSHIFTEDINT(a,n,v) \
28 const u32 s = ((n) & 3) * 8; \
29 const u64 x = (u64) (v) << s; \
30 (a)[((n)/4)+0] &= ~(0xff << ((n & 3) * 8)); \
31 (a)[((n)/4)+0] |= x; \
32 (a)[((n)/4)+1] = x >> 32; \
35 __constant u32 sapb_trans_tbl[256] =
37 // first value hack for 0 byte as part of an optimization
38 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
39 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
40 0x3f, 0x40, 0x41, 0x50, 0x43, 0x44, 0x45, 0x4b, 0x47, 0x48, 0x4d, 0x4e, 0x54, 0x51, 0x53, 0x46,
41 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x3b, 0x3c, 0x3d, 0x3e, 0x56, 0x55, 0x5c, 0x49, 0x5d, 0x4a,
42 0x42, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
43 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x58, 0x5b, 0x59, 0xff, 0x52,
44 0x4c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
45 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x57, 0x5e, 0x5a, 0x4f, 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,
49 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
50 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
51 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
52 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
53 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
56 __constant u32 bcodeArray[48] =
58 0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
59 0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
60 0xe2, 0xb7, 0x33, 0x71, 0x8b, 0x9f, 0x5d, 0x01, 0x44, 0x70, 0xae, 0x11, 0xef, 0x28, 0xf0, 0x0d
63 static u32 sapb_trans (const u32 in)
68 out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
69 out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
70 out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
71 out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
77 static 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])
96 u32 sum20 = ((a >> 24) & 3)
104 const u32 w[2] = { w0[0], w0[1] };
106 const u32 s[3] = { salt_buf0[0], salt_buf0[1], salt_buf0[2] };
108 u32 saved_key[4] = { a, b, c, d };
114 // we can assume this because the password must be at least 3
115 // and the username must be at least 1 so we can save the if ()
121 t0 |= bcodeArray[47] << 0;
122 t0 |= (w[0] & 0xff) << 8;
123 t0 |= (s[0] & 0xff) << 16;
124 t0 |= bcodeArray[ 1] << 24;
132 t0 |= (w[0] & 0xff) << 0;
133 t0 |= (s[0] & 0xff) << 8;
134 t0 |= bcodeArray[ 0] << 16;
143 // because the following code can increase i2 by a maximum of 5,
144 // there is an overflow potential of 4 before it comes to the next test for i2 >= sum20
145 // we need to truncate in that case
147 while ((i1 < pw_len) && (i3 < salt_len))
153 if (GETCHAR (saved_key, 15 - i1) & 1)
155 x0 |= bcodeArray[48 - 1 - i1] << 0; i2++;
156 x0 |= GETCHAR (w, i1) << 8; i2++; i1++;
157 x0 |= GETCHAR (s, i3) << 16; i2++; i3++;
158 x0 |= bcodeArray[i2 - i1 - i3] << 24; i2++; i2++;
162 x0 |= GETCHAR (w, i1) << 0; i2++; i1++;
163 x0 |= GETCHAR (s, i3) << 8; i2++; i3++;
164 x0 |= bcodeArray[i2 - i1 - i3] << 16; i2++; i2++;
167 SETSHIFTEDINT (t, i2_sav, x0);
175 while ((i1 < pw_len) || (i3 < salt_len))
177 if (i1 < pw_len) // max 8
179 if (GETCHAR (saved_key, 15 - i1) & 1)
181 PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
186 PUTCHAR (t, i2, GETCHAR (w, i1));
193 PUTCHAR (t, i2, GETCHAR (s, i3));
199 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
212 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
221 static void m07700m (u32 w0[4], u32 w1[4], u32 w2[4], u32 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)
227 const u32 gid = get_global_id (0);
228 const u32 lid = get_local_id (0);
230 w0[0] = sapb_trans (w0[0]);
231 w0[1] = sapb_trans (w0[1]);
239 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
240 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
241 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
243 salt_buf0[0] = sapb_trans (salt_buf0[0]);
244 salt_buf0[1] = sapb_trans (salt_buf0[1]);
245 salt_buf0[2] = sapb_trans (salt_buf0[2]);
247 const u32 salt_len = salt_bufs[salt_pos].salt_len;
251 s0[0] = salt_buf0[0];
252 s0[1] = salt_buf0[1];
253 s0[2] = salt_buf0[2];
277 switch_buffer_by_offset (s0, s1, s2, s3, pw_len);
279 const u32 pw_salt_len = pw_len + salt_len;
287 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
289 const u32 w0r = sapb_trans (bfs_buf[il_pos].i);
295 t[ 0] = s0[0] | w0[0];
296 t[ 1] = s0[1] | w0[1];
309 t[14] = pw_salt_len * 8;
312 PUTCHAR (t, pw_salt_len, 0x80);
323 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
324 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
325 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
326 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
327 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
328 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
329 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
330 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
331 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
332 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
333 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
334 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
335 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
336 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
337 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
338 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
340 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
341 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
342 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
343 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
344 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
345 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
346 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
347 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
348 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
349 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
350 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
351 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
352 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
353 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
354 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
355 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
357 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
358 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
359 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
360 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
361 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
362 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
363 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
364 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
365 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
366 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
367 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
368 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
369 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
370 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
371 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
372 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
374 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
375 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
376 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
377 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
378 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
379 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
380 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
381 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
382 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
383 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
384 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
385 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
386 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
387 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
388 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
389 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
396 const u32 sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
398 SETSHIFTEDINT (t, sum20, 0x80);
407 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
408 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
409 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
410 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
411 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
412 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
413 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
414 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
415 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
416 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
417 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
418 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
419 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
420 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
421 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
422 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
424 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
425 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
426 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
427 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
428 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
429 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
430 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
431 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
432 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
433 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
434 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
435 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
436 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
437 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
438 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
439 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
441 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
442 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
443 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
444 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
445 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
446 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
447 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
448 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
449 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
450 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
451 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
452 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
453 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
454 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
455 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
456 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
458 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
459 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
460 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
461 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
462 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
463 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
464 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
465 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
466 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
467 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
468 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
469 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
470 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
471 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
472 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
473 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
492 static void m07700s (u32 w0[4], u32 w1[4], u32 w2[4], u32 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)
498 const u32 gid = get_global_id (0);
499 const u32 lid = get_local_id (0);
501 w0[0] = sapb_trans (w0[0]);
502 w0[1] = sapb_trans (w0[1]);
510 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
511 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
512 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
514 salt_buf0[0] = sapb_trans (salt_buf0[0]);
515 salt_buf0[1] = sapb_trans (salt_buf0[1]);
516 salt_buf0[2] = sapb_trans (salt_buf0[2]);
518 const u32 salt_len = salt_bufs[salt_pos].salt_len;
522 s0[0] = salt_buf0[0];
523 s0[1] = salt_buf0[1];
524 s0[2] = salt_buf0[2];
548 switch_buffer_by_offset (s0, s1, s2, s3, pw_len);
550 const u32 pw_salt_len = pw_len + salt_len;
556 const u32 search[4] =
558 digests_buf[digests_offset].digest_buf[DGST_R0],
559 digests_buf[digests_offset].digest_buf[DGST_R1],
560 digests_buf[digests_offset].digest_buf[DGST_R2],
561 digests_buf[digests_offset].digest_buf[DGST_R3]
570 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
572 const u32 w0r = sapb_trans (bfs_buf[il_pos].i);
578 t[ 0] = s0[0] | w0[0];
579 t[ 1] = s0[1] | w0[1];
592 t[14] = pw_salt_len * 8;
595 PUTCHAR (t, pw_salt_len, 0x80);
606 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
607 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
608 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
609 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
610 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
611 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
612 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
613 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
614 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
615 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
616 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
617 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
618 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
619 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
620 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
621 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
623 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
624 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
625 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
626 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
627 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
628 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
629 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
630 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
631 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
632 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
633 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
634 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
635 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
636 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
637 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
638 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
640 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
641 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
642 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
643 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
644 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
645 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
646 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
647 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
648 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
649 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
650 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
651 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
652 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
653 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
654 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
655 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
657 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
658 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
659 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
660 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
661 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
662 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
663 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
664 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
665 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
666 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
667 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
668 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
669 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
670 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
671 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
672 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
679 const u32 sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
681 SETSHIFTEDINT (t, sum20, 0x80);
690 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
691 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
692 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
693 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
694 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
695 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
696 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
697 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
698 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
699 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
700 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
701 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
702 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
703 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
704 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
705 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
707 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
708 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
709 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
710 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
711 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
712 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
713 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
714 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
715 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
716 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
717 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
718 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
719 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
720 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
721 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
722 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
724 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
725 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
726 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
727 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
728 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
729 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
730 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
731 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
732 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
733 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
734 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
735 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
736 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
737 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
738 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
739 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
741 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
742 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
743 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
744 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
745 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
746 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
747 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
748 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
749 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
750 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
751 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
752 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
753 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
754 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
755 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
756 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
775 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07700_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)
781 const u32 gid = get_global_id (0);
783 if (gid >= gid_max) return;
789 const u32 lid = get_local_id (0);
793 w0[0] = pws[gid].i[ 0];
794 w0[1] = pws[gid].i[ 1];
795 w0[2] = pws[gid].i[ 2];
796 w0[3] = pws[gid].i[ 3];
819 const u32 pw_len = pws[gid].pw_len;
825 m07700m (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);
828 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07700_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)
834 const u32 gid = get_global_id (0);
836 if (gid >= gid_max) return;
842 const u32 lid = get_local_id (0);
846 w0[0] = pws[gid].i[ 0];
847 w0[1] = pws[gid].i[ 1];
848 w0[2] = pws[gid].i[ 2];
849 w0[3] = pws[gid].i[ 3];
853 w1[0] = pws[gid].i[ 4];
854 w1[1] = pws[gid].i[ 5];
855 w1[2] = pws[gid].i[ 6];
856 w1[3] = pws[gid].i[ 7];
872 const u32 pw_len = pws[gid].pw_len;
878 m07700m (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);
881 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07700_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)
885 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07700_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)
891 const u32 gid = get_global_id (0);
893 if (gid >= gid_max) return;
899 const u32 lid = get_local_id (0);
903 w0[0] = pws[gid].i[ 0];
904 w0[1] = pws[gid].i[ 1];
905 w0[2] = pws[gid].i[ 2];
906 w0[3] = pws[gid].i[ 3];
929 const u32 pw_len = pws[gid].pw_len;
935 m07700s (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);
938 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07700_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)
944 const u32 gid = get_global_id (0);
946 if (gid >= gid_max) return;
952 const u32 lid = get_local_id (0);
956 w0[0] = pws[gid].i[ 0];
957 w0[1] = pws[gid].i[ 1];
958 w0[2] = pws[gid].i[ 2];
959 w0[3] = pws[gid].i[ 3];
963 w1[0] = pws[gid].i[ 4];
964 w1[1] = pws[gid].i[ 5];
965 w1[2] = pws[gid].i[ 6];
966 w1[3] = pws[gid].i[ 7];
982 const u32 pw_len = pws[gid].pw_len;
988 m07700s (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);
991 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07700_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)