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 "OpenCL/types_ocl.c"
18 #include "OpenCL/common.c"
20 #define COMPARE_S "OpenCL/check_single_comp4.c"
21 #define COMPARE_M "OpenCL/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)
67 out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
68 out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
69 out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
70 out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
75 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])
94 u32 sum20 = ((a >> 24) & 3)
102 const u32 w[2] = { w0[0], w0[1] };
104 const u32 s[3] = { salt_buf0[0], salt_buf0[1], salt_buf0[2] };
106 u32 saved_key[4] = { a, b, c, d };
112 // we can assume this because the password must be at least 3
113 // and the username must be at least 1 so we can save the if ()
119 t0 |= bcodeArray[47] << 0;
120 t0 |= (w[0] & 0xff) << 8;
121 t0 |= (s[0] & 0xff) << 16;
122 t0 |= bcodeArray[ 1] << 24;
130 t0 |= (w[0] & 0xff) << 0;
131 t0 |= (s[0] & 0xff) << 8;
132 t0 |= bcodeArray[ 0] << 16;
141 // because the following code can increase i2 by a maximum of 5,
142 // there is an overflow potential of 4 before it comes to the next test for i2 >= sum20
143 // we need to truncate in that case
145 while ((i1 < pw_len) && (i3 < salt_len))
151 if (GETCHAR (saved_key, 15 - i1) & 1)
153 x0 |= bcodeArray[48 - 1 - i1] << 0; i2++;
154 x0 |= GETCHAR (w, i1) << 8; i2++; i1++;
155 x0 |= GETCHAR (s, i3) << 16; i2++; i3++;
156 x0 |= bcodeArray[i2 - i1 - i3] << 24; i2++; i2++;
160 x0 |= GETCHAR (w, i1) << 0; i2++; i1++;
161 x0 |= GETCHAR (s, i3) << 8; i2++; i3++;
162 x0 |= bcodeArray[i2 - i1 - i3] << 16; i2++; i2++;
165 SETSHIFTEDINT (t, i2_sav, x0);
173 while ((i1 < pw_len) || (i3 < salt_len))
175 if (i1 < pw_len) // max 8
177 if (GETCHAR (saved_key, 15 - i1) & 1)
179 PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
184 PUTCHAR (t, i2, GETCHAR (w, i1));
191 PUTCHAR (t, i2, GETCHAR (s, i3));
197 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
210 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
219 __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_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
225 const u32 lid = get_local_id (0);
231 const u32 gid = get_global_id (0);
233 if (gid >= gid_max) return;
237 wordl0[0] = pws[gid].i[ 0];
238 wordl0[1] = pws[gid].i[ 1];
263 const u32 pw_l_len = pws[gid].pw_len;
265 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
267 switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
276 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
277 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
278 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
281 const u32 salt_len = salt_bufs[salt_pos].salt_len;
283 salt_buf0[0] = sapb_trans (salt_buf0[0]);
284 salt_buf0[1] = sapb_trans (salt_buf0[1]);
285 salt_buf0[2] = sapb_trans (salt_buf0[2]);
295 for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
297 const u32 pw_r_len = combs_buf[il_pos].pw_len;
299 const u32 pw_len = pw_l_len + pw_r_len;
303 wordr0[0] = combs_buf[il_pos].i[0];
304 wordr0[1] = combs_buf[il_pos].i[1];
329 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
331 switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
336 w0[0] = sapb_trans (wordl0[0] | wordr0[0]);
337 w0[1] = sapb_trans (wordl0[1] | wordr0[1]);
347 s0[0] = salt_buf0[0];
348 s0[1] = salt_buf0[1];
349 s0[2] = salt_buf0[2];
373 switch_buffer_by_offset_le (s0, s1, s2, s3, pw_len);
375 const u32 pw_salt_len = pw_len + salt_len;
379 t[ 0] = s0[0] | w0[0];
380 t[ 1] = s0[1] | w0[1];
393 t[14] = pw_salt_len * 8;
396 PUTCHAR (t, pw_salt_len, 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);
480 const u32 sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
482 SETSHIFTEDINT (t, sum20, 0x80);
491 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
492 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
493 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
494 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
495 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
496 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
497 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
498 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
499 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
500 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
501 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
502 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
503 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
504 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
505 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
506 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
508 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
509 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
510 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
511 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
512 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
513 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
514 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
515 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
516 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
517 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
518 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
519 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
520 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
521 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
522 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
523 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
525 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
526 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
527 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
528 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
529 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
530 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
531 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
532 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
533 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
534 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
535 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
536 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
537 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
538 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
539 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
540 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
542 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
543 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
544 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
545 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
546 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
547 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
548 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
549 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
550 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
551 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
552 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
553 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
554 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
555 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
556 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
557 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
576 __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_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
580 __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_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
584 __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_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
590 const u32 lid = get_local_id (0);
596 const u32 gid = get_global_id (0);
598 if (gid >= gid_max) return;
602 wordl0[0] = pws[gid].i[ 0];
603 wordl0[1] = pws[gid].i[ 1];
628 const u32 pw_l_len = pws[gid].pw_len;
630 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
632 switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
641 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
642 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
643 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
646 const u32 salt_len = salt_bufs[salt_pos].salt_len;
648 salt_buf0[0] = sapb_trans (salt_buf0[0]);
649 salt_buf0[1] = sapb_trans (salt_buf0[1]);
650 salt_buf0[2] = sapb_trans (salt_buf0[2]);
656 const u32 search[4] =
658 digests_buf[digests_offset].digest_buf[DGST_R0],
659 digests_buf[digests_offset].digest_buf[DGST_R1],
660 digests_buf[digests_offset].digest_buf[DGST_R2],
661 digests_buf[digests_offset].digest_buf[DGST_R3]
668 for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
670 const u32 pw_r_len = combs_buf[il_pos].pw_len;
672 const u32 pw_len = pw_l_len + pw_r_len;
676 wordr0[0] = combs_buf[il_pos].i[0];
677 wordr0[1] = combs_buf[il_pos].i[1];
702 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
704 switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
709 w0[0] = sapb_trans (wordl0[0] | wordr0[0]);
710 w0[1] = sapb_trans (wordl0[1] | wordr0[1]);
720 s0[0] = salt_buf0[0];
721 s0[1] = salt_buf0[1];
722 s0[2] = salt_buf0[2];
746 switch_buffer_by_offset_le (s0, s1, s2, s3, pw_len);
748 const u32 pw_salt_len = pw_len + salt_len;
752 t[ 0] = s0[0] | w0[0];
753 t[ 1] = s0[1] | w0[1];
766 t[14] = pw_salt_len * 8;
769 PUTCHAR (t, pw_salt_len, 0x80);
780 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
781 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
782 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
783 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
784 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
785 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
786 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
787 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
788 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
789 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
790 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
791 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
792 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
793 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
794 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
795 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
797 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
798 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
799 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
800 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
801 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
802 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
803 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
804 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
805 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
806 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
807 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
808 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
809 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
810 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
811 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
812 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
814 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
815 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
816 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
817 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
818 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
819 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
820 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
821 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
822 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
823 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
824 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
825 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
826 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
827 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
828 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
829 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
831 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
832 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
833 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
834 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
835 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
836 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
837 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
838 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
839 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
840 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
841 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
842 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
843 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
844 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
845 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
846 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
853 const u32 sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
855 SETSHIFTEDINT (t, sum20, 0x80);
864 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
865 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
866 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
867 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
868 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
869 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
870 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
871 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
872 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
873 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
874 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
875 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
876 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
877 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
878 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
879 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
881 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
882 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
883 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
884 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
885 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
886 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
887 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
888 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
889 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
890 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
891 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
892 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
893 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
894 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
895 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
896 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
898 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
899 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
900 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
901 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
902 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
903 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
904 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
905 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
906 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
907 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
908 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
909 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
910 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
911 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
912 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
913 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
915 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
916 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
917 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
918 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
919 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
920 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
921 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
922 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
923 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
924 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
925 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
926 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
927 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
928 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
929 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
930 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
949 __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_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
953 __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_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)