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"
18 #include "inc_simd.cl"
20 #define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
21 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
23 #define SETSHIFTEDINT(a,n,v) \
25 const u32 s = ((n) & 3) * 8; \
26 const u64 x = (u64) (v) << s; \
27 (a)[((n)/4)+0] &= ~(0xff << ((n & 3) * 8)); \
28 (a)[((n)/4)+0] |= x; \
29 (a)[((n)/4)+1] = x >> 32; \
32 __constant u32 sapb_trans_tbl[256] =
34 // first value hack for 0 byte as part of an optimization
35 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
36 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
37 0x3f, 0x40, 0x41, 0x50, 0x43, 0x44, 0x45, 0x4b, 0x47, 0x48, 0x4d, 0x4e, 0x54, 0x51, 0x53, 0x46,
38 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x3b, 0x3c, 0x3d, 0x3e, 0x56, 0x55, 0x5c, 0x49, 0x5d, 0x4a,
39 0x42, 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, 0x58, 0x5b, 0x59, 0xff, 0x52,
41 0x4c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
42 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x57, 0x5e, 0x5a, 0x4f, 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,
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
53 __constant u32 bcodeArray[48] =
55 0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
56 0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
57 0xe2, 0xb7, 0x33, 0x71, 0x8b, 0x9f, 0x5d, 0x01, 0x44, 0x70, 0xae, 0x11, 0xef, 0x28, 0xf0, 0x0d
60 u32 sapb_trans (const u32 in)
64 out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
65 out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
66 out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
67 out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
72 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])
91 u32 sum20 = ((a >> 24) & 3)
99 const u32 w[2] = { w0[0], w0[1] };
101 const u32 s[3] = { salt_buf0[0], salt_buf0[1], salt_buf0[2] };
103 u32 saved_key[4] = { a, b, c, d };
109 // we can assume this because the password must be at least 3
110 // and the username must be at least 1 so we can save the if ()
116 t0 |= bcodeArray[47] << 0;
117 t0 |= (w[0] & 0xff) << 8;
118 t0 |= (s[0] & 0xff) << 16;
119 t0 |= bcodeArray[ 1] << 24;
127 t0 |= (w[0] & 0xff) << 0;
128 t0 |= (s[0] & 0xff) << 8;
129 t0 |= bcodeArray[ 0] << 16;
138 // because the following code can increase i2 by a maximum of 5,
139 // there is an overflow potential of 4 before it comes to the next test for i2 >= sum20
140 // we need to truncate in that case
142 while ((i1 < pw_len) && (i3 < salt_len))
148 if (GETCHAR (saved_key, 15 - i1) & 1)
150 x0 |= bcodeArray[48 - 1 - i1] << 0; i2++;
151 x0 |= GETCHAR (w, i1) << 8; i2++; i1++;
152 x0 |= GETCHAR (s, i3) << 16; i2++; i3++;
153 x0 |= bcodeArray[i2 - i1 - i3] << 24; i2++; i2++;
157 x0 |= GETCHAR (w, i1) << 0; i2++; i1++;
158 x0 |= GETCHAR (s, i3) << 8; i2++; i3++;
159 x0 |= bcodeArray[i2 - i1 - i3] << 16; i2++; i2++;
162 SETSHIFTEDINT (t, i2_sav, x0);
170 while ((i1 < pw_len) || (i3 < salt_len))
172 if (i1 < pw_len) // max 8
174 if (GETCHAR (saved_key, 15 - i1) & 1)
176 PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
181 PUTCHAR (t, i2, GETCHAR (w, i1));
188 PUTCHAR (t, i2, GETCHAR (s, i3));
194 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
207 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
216 __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)
222 const u32 lid = get_local_id (0);
228 const u32 gid = get_global_id (0);
230 if (gid >= gid_max) return;
235 pw_buf0[0] = pws[gid].i[0];
236 pw_buf0[1] = pws[gid].i[1];
237 pw_buf0[2] = pws[gid].i[2];
238 pw_buf0[3] = pws[gid].i[3];
239 pw_buf1[0] = pws[gid].i[4];
240 pw_buf1[1] = pws[gid].i[5];
241 pw_buf1[2] = pws[gid].i[6];
242 pw_buf1[3] = pws[gid].i[7];
244 const u32 pw_len = pws[gid].pw_len;
252 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
253 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
254 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
257 const u32 salt_len = salt_bufs[salt_pos].salt_len;
259 salt_buf0[0] = sapb_trans (salt_buf0[0]);
260 salt_buf0[1] = sapb_trans (salt_buf0[1]);
261 salt_buf0[2] = sapb_trans (salt_buf0[2]);
267 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
274 const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
276 if (out_len > 8) continue; // otherwise it overflows in waldorf function
282 w0[0] = sapb_trans (w0[0]);
283 w0[1] = sapb_trans (w0[1]);
294 s0[0] = salt_buf0[0];
295 s0[1] = salt_buf0[1];
296 s0[2] = salt_buf0[2];
311 switch_buffer_by_offset_le (s0, s1, s2, s3, out_len);
313 const u32 pw_salt_len = out_len + salt_len;
317 t[ 0] = s0[0] | w0[0];
318 t[ 1] = s0[1] | w0[1];
331 t[14] = pw_salt_len * 8;
334 PUTCHAR (t, pw_salt_len, 0x80);
345 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
346 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
347 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
348 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
349 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
350 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
351 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
352 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
353 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
354 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
355 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
356 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
357 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
358 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
359 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
360 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
362 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
363 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
364 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
365 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
366 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
367 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
368 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
369 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
370 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
371 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
372 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
373 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
374 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
375 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
376 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
377 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
379 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
380 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
381 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
382 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
383 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
384 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
385 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
386 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
387 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
388 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
389 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
390 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
391 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
392 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
393 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
394 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
396 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
397 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
398 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
399 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
400 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
401 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
402 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
403 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
404 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
405 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
406 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
407 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
408 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
409 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
410 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
411 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
418 const u32 sum20 = walld0rf_magic (w0, out_len, salt_buf0, salt_len, a, b, c, d, t);
420 SETSHIFTEDINT (t, sum20, 0x80);
429 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
430 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
431 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
432 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
433 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
434 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
435 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
436 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
437 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
438 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
439 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
440 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
441 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
442 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
443 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
444 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
446 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
447 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
448 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
449 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
450 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
451 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
452 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
453 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
454 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
455 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
456 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
457 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
458 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
459 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
460 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
461 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
463 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
464 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
465 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
466 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
467 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
468 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
469 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
470 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
471 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
472 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
473 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
474 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
475 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
476 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
477 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
478 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
480 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
481 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
482 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
483 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
484 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
485 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
486 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
487 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
488 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
489 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
490 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
491 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
492 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
493 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
494 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
495 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
507 COMPARE_M_SIMD (a, b, c, d);
511 __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)
515 __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)
519 __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)
525 const u32 lid = get_local_id (0);
531 const u32 gid = get_global_id (0);
533 if (gid >= gid_max) return;
538 pw_buf0[0] = pws[gid].i[0];
539 pw_buf0[1] = pws[gid].i[1];
540 pw_buf0[2] = pws[gid].i[2];
541 pw_buf0[3] = pws[gid].i[3];
542 pw_buf1[0] = pws[gid].i[4];
543 pw_buf1[1] = pws[gid].i[5];
544 pw_buf1[2] = pws[gid].i[6];
545 pw_buf1[3] = pws[gid].i[7];
547 const u32 pw_len = pws[gid].pw_len;
555 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
556 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
557 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
560 const u32 salt_len = salt_bufs[salt_pos].salt_len;
562 salt_buf0[0] = sapb_trans (salt_buf0[0]);
563 salt_buf0[1] = sapb_trans (salt_buf0[1]);
564 salt_buf0[2] = sapb_trans (salt_buf0[2]);
570 const u32 search[4] =
572 digests_buf[digests_offset].digest_buf[DGST_R0],
573 digests_buf[digests_offset].digest_buf[DGST_R1],
582 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
589 const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
591 if (out_len > 8) continue; // otherwise it overflows in waldorf function
597 w0[0] = sapb_trans (w0[0]);
598 w0[1] = sapb_trans (w0[1]);
609 s0[0] = salt_buf0[0];
610 s0[1] = salt_buf0[1];
611 s0[2] = salt_buf0[2];
626 switch_buffer_by_offset_le (s0, s1, s2, s3, out_len);
628 const u32 pw_salt_len = out_len + salt_len;
632 t[ 0] = s0[0] | w0[0];
633 t[ 1] = s0[1] | w0[1];
646 t[14] = pw_salt_len * 8;
649 PUTCHAR (t, pw_salt_len, 0x80);
660 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
661 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
662 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
663 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
664 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
665 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
666 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
667 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
668 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
669 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
670 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
671 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
672 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
673 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
674 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
675 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
677 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
678 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
679 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
680 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
681 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
682 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
683 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
684 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
685 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
686 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
687 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
688 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
689 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
690 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
691 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
692 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
694 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
695 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
696 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
697 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
698 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
699 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
700 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
701 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
702 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
703 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
704 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
705 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
706 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
707 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
708 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
709 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
711 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
712 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
713 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
714 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
715 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
716 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
717 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
718 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
719 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
720 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
721 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
722 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
723 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
724 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
725 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
726 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
733 const u32 sum20 = walld0rf_magic (w0, out_len, salt_buf0, salt_len, a, b, c, d, t);
735 SETSHIFTEDINT (t, sum20, 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);
822 COMPARE_S_SIMD (a, b, c, d);
826 __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)
830 __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)