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"
19 #include "include/rp_gpu.h"
22 #define COMPARE_S "check_single_comp4.c"
23 #define COMPARE_M "check_multi_comp4.c"
25 #define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
26 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
28 #define SETSHIFTEDINT(a,n,v) \
30 const u32 s = ((n) & 3) * 8; \
31 const u64 x = (u64) (v) << s; \
32 (a)[((n)/4)+0] &= ~(0xff << ((n & 3) * 8)); \
33 (a)[((n)/4)+0] |= x; \
34 (a)[((n)/4)+1] = x >> 32; \
37 __constant u32 sapb_trans_tbl[256] =
39 // first value hack for 0 byte as part of an optimization
40 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
41 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
42 0x3f, 0x40, 0x41, 0x50, 0x43, 0x44, 0x45, 0x4b, 0x47, 0x48, 0x4d, 0x4e, 0x54, 0x51, 0x53, 0x46,
43 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x3b, 0x3c, 0x3d, 0x3e, 0x56, 0x55, 0x5c, 0x49, 0x5d, 0x4a,
44 0x42, 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, 0x58, 0x5b, 0x59, 0xff, 0x52,
46 0x4c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
47 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x57, 0x5e, 0x5a, 0x4f, 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,
54 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
55 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
58 __constant u32 bcodeArray[48] =
60 0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
61 0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
62 0xe2, 0xb7, 0x33, 0x71, 0x8b, 0x9f, 0x5d, 0x01, 0x44, 0x70, 0xae, 0x11, 0xef, 0x28, 0xf0, 0x0d
65 static u32 sapb_trans (const u32 in)
70 out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
71 out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
72 out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
73 out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
79 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])
98 u32 sum20 = ((a >> 24) & 3)
106 const u32 w[2] = { w0[0], w0[1] };
108 const u32 s[3] = { salt_buf0[0], salt_buf0[1], salt_buf0[2] };
110 u32 saved_key[4] = { a, b, c, d };
116 // we can assume this because the password must be at least 3
117 // and the username must be at least 1 so we can save the if ()
123 t0 |= bcodeArray[47] << 0;
124 t0 |= (w[0] & 0xff) << 8;
125 t0 |= (s[0] & 0xff) << 16;
126 t0 |= bcodeArray[ 1] << 24;
134 t0 |= (w[0] & 0xff) << 0;
135 t0 |= (s[0] & 0xff) << 8;
136 t0 |= bcodeArray[ 0] << 16;
145 // because the following code can increase i2 by a maximum of 5,
146 // there is an overflow potential of 4 before it comes to the next test for i2 >= sum20
147 // we need to truncate in that case
149 while ((i1 < pw_len) && (i3 < salt_len))
155 if (GETCHAR (saved_key, 15 - i1) & 1)
157 x0 |= bcodeArray[48 - 1 - i1] << 0; i2++;
158 x0 |= GETCHAR (w, i1) << 8; i2++; i1++;
159 x0 |= GETCHAR (s, i3) << 16; i2++; i3++;
160 x0 |= bcodeArray[i2 - i1 - i3] << 24; i2++; i2++;
164 x0 |= GETCHAR (w, i1) << 0; i2++; i1++;
165 x0 |= GETCHAR (s, i3) << 8; i2++; i3++;
166 x0 |= bcodeArray[i2 - i1 - i3] << 16; i2++; i2++;
169 SETSHIFTEDINT (t, i2_sav, x0);
177 while ((i1 < pw_len) || (i3 < salt_len))
179 if (i1 < pw_len) // max 8
181 if (GETCHAR (saved_key, 15 - i1) & 1)
183 PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
188 PUTCHAR (t, i2, GETCHAR (w, i1));
195 PUTCHAR (t, i2, GETCHAR (s, i3));
201 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
214 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
223 __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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
229 const u32 lid = get_local_id (0);
235 const u32 gid = get_global_id (0);
237 if (gid >= gid_max) return;
241 pw_buf0[0] = pws[gid].i[ 0];
242 pw_buf0[1] = pws[gid].i[ 1];
253 const u32 pw_len = pws[gid].pw_len;
261 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
262 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
263 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
266 const u32 salt_len = salt_bufs[salt_pos].salt_len;
268 salt_buf0[0] = sapb_trans (salt_buf0[0]);
269 salt_buf0[1] = sapb_trans (salt_buf0[1]);
270 salt_buf0[2] = sapb_trans (salt_buf0[2]);
276 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
306 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
308 if (out_len > 8) continue; // otherwise it overflows in waldorf function
310 w0[0] = sapb_trans (w0[0]);
311 w0[1] = sapb_trans (w0[1]);
319 s0[0] = salt_buf0[0];
320 s0[1] = salt_buf0[1];
321 s0[2] = salt_buf0[2];
345 switch_buffer_by_offset (s0, s1, s2, s3, out_len);
347 const u32 pw_salt_len = out_len + salt_len;
351 t[ 0] = s0[0] | w0[0];
352 t[ 1] = s0[1] | w0[1];
365 t[14] = pw_salt_len * 8;
368 PUTCHAR (t, pw_salt_len, 0x80);
379 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
380 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
381 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
382 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
383 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
384 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
385 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
386 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
387 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
388 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
389 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
390 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
391 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
392 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
393 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
394 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
396 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
397 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
398 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
399 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
400 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
401 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
402 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
403 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
404 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
405 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
406 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
407 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
408 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
409 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
410 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
411 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
413 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
414 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
415 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
416 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
417 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
418 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
419 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
420 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
421 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
422 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
423 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
424 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
425 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
426 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
427 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
428 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
430 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
431 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
432 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
433 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
434 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
435 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
436 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
437 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
438 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
439 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
440 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
441 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
442 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
443 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
444 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
445 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
452 const u32 sum20 = walld0rf_magic (w0, out_len, salt_buf0, salt_len, a, b, c, d, t);
454 SETSHIFTEDINT (t, sum20, 0x80);
463 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
464 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
465 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
466 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
467 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
468 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
469 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
470 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
471 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
472 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
473 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
474 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
475 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
476 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
477 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
478 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
480 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
481 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
482 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
483 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
484 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
485 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
486 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
487 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
488 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
489 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
490 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
491 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
492 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
493 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
494 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
495 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
497 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
498 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
499 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
500 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
501 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
502 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
503 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
504 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
505 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
506 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
507 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
508 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
509 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
510 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
511 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
512 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
514 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
515 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
516 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
517 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
518 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
519 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
520 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
521 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
522 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
523 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
524 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
525 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
526 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
527 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
528 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
529 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
548 __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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
552 __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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
556 __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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
562 const u32 lid = get_local_id (0);
568 const u32 gid = get_global_id (0);
570 if (gid >= gid_max) return;
574 pw_buf0[0] = pws[gid].i[ 0];
575 pw_buf0[1] = pws[gid].i[ 1];
586 const u32 pw_len = pws[gid].pw_len;
594 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
595 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
596 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
599 const u32 salt_len = salt_bufs[salt_pos].salt_len;
601 salt_buf0[0] = sapb_trans (salt_buf0[0]);
602 salt_buf0[1] = sapb_trans (salt_buf0[1]);
603 salt_buf0[2] = sapb_trans (salt_buf0[2]);
609 const u32 search[4] =
611 digests_buf[digests_offset].digest_buf[DGST_R0],
612 digests_buf[digests_offset].digest_buf[DGST_R1],
613 digests_buf[digests_offset].digest_buf[DGST_R2],
614 digests_buf[digests_offset].digest_buf[DGST_R3]
621 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
651 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
653 if (out_len > 8) continue; // otherwise it overflows in waldorf function
655 w0[0] = sapb_trans (w0[0]);
656 w0[1] = sapb_trans (w0[1]);
664 s0[0] = salt_buf0[0];
665 s0[1] = salt_buf0[1];
666 s0[2] = salt_buf0[2];
690 switch_buffer_by_offset (s0, s1, s2, s3, out_len);
692 const u32 pw_salt_len = out_len + salt_len;
696 t[ 0] = s0[0] | w0[0];
697 t[ 1] = s0[1] | w0[1];
710 t[14] = pw_salt_len * 8;
713 PUTCHAR (t, pw_salt_len, 0x80);
724 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
725 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
726 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
727 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
728 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
729 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
730 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
731 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
732 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
733 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
734 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
735 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
736 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
737 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
738 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
739 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
741 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
742 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
743 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
744 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
745 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
746 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
747 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
748 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
749 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
750 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
751 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
752 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
753 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
754 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
755 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
756 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
758 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
759 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
760 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
761 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
762 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
763 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
764 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
765 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
766 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
767 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
768 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
769 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
770 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
771 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
772 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
773 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
775 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
776 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
777 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
778 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
779 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
780 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
781 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
782 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
783 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
784 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
785 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
786 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
787 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
788 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
789 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
790 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
797 const u32 sum20 = walld0rf_magic (w0, out_len, salt_buf0, salt_len, a, b, c, d, t);
799 SETSHIFTEDINT (t, sum20, 0x80);
808 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
809 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
810 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
811 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
812 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
813 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
814 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
815 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
816 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
817 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
818 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
819 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
820 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
821 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
822 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
823 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
825 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
826 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
827 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
828 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
829 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
830 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
831 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
832 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
833 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
834 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
835 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
836 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
837 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
838 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
839 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
840 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
842 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
843 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
844 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
845 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
846 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
847 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
848 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
849 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
850 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
851 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
852 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
853 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
854 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
855 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
856 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
857 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
859 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
860 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
861 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
862 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
863 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
864 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
865 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
866 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
867 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
868 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
869 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
870 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
871 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
872 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
873 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
874 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
893 __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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
897 __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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)