2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
28 #include "include/kernel_functions.c"
29 #include "types_amd.c"
30 #include "common_amd.c"
33 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
34 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
37 #define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
38 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
40 #define SETSHIFTEDINT(a,n,v) \
42 const u32 s = ((n) & 3) * 8; \
43 const u64 x = (u64) (v) << s; \
44 (a)[((n)/4)+0] &= ~(0xff << ((n & 3) * 8)); \
45 (a)[((n)/4)+0] |= x; \
46 (a)[((n)/4)+1] = x >> 32; \
49 __constant u32 sapb_trans_tbl[256] =
51 // first value hack for 0 byte as part of an optimization
52 0x00, 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 0x3f, 0x40, 0x41, 0x50, 0x43, 0x44, 0x45, 0x4b, 0x47, 0x48, 0x4d, 0x4e, 0x54, 0x51, 0x53, 0x46,
55 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x3b, 0x3c, 0x3d, 0x3e, 0x56, 0x55, 0x5c, 0x49, 0x5d, 0x4a,
56 0x42, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
57 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x58, 0x5b, 0x59, 0xff, 0x52,
58 0x4c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
59 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x57, 0x5e, 0x5a, 0x4f, 0xff,
60 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
61 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
62 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
63 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
64 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
65 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
66 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
67 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
70 __constant u32 bcodeArray[48] =
72 0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
73 0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
74 0xe2, 0xb7, 0x33, 0x71, 0x8b, 0x9f, 0x5d, 0x01, 0x44, 0x70, 0xae, 0x11, 0xef, 0x28, 0xf0, 0x0d
77 static u32x sapb_trans (const u32x in)
82 out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
83 out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
84 out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
85 out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
91 static u32x walld0rf_magic (const u32x w0[4], const u32 pw_len, const u32x salt_buf0[4], const u32 salt_len, const u32x a, const u32x b, const u32x c, const u32x d, u32x t[16])
110 u32 sum20 = ((a >> 24) & 3)
118 const u32 w[2] = { w0[0], w0[1] };
120 const u32 s[3] = { salt_buf0[0], salt_buf0[1], salt_buf0[2] };
122 u32 saved_key[4] = { a, b, c, d };
128 // we can assume this because the password must be at least 3
129 // and the username must be at least 1 so we can save the if ()
135 t0 |= bcodeArray[47] << 0;
136 t0 |= (w[0] & 0xff) << 8;
137 t0 |= (s[0] & 0xff) << 16;
138 t0 |= bcodeArray[ 1] << 24;
146 t0 |= (w[0] & 0xff) << 0;
147 t0 |= (s[0] & 0xff) << 8;
148 t0 |= bcodeArray[ 0] << 16;
157 // because the following code can increase i2 by a maximum of 5,
158 // there is an overflow potential of 4 before it comes to the next test for i2 >= sum20
159 // we need to truncate in that case
161 while ((i1 < pw_len) && (i3 < salt_len))
167 if (GETCHAR (saved_key, 15 - i1) & 1)
169 x0 |= bcodeArray[48 - 1 - i1] << 0; i2++;
170 x0 |= GETCHAR (w, i1) << 8; i2++; i1++;
171 x0 |= GETCHAR (s, i3) << 16; i2++; i3++;
172 x0 |= bcodeArray[i2 - i1 - i3] << 24; i2++; i2++;
176 x0 |= GETCHAR (w, i1) << 0; i2++; i1++;
177 x0 |= GETCHAR (s, i3) << 8; i2++; i3++;
178 x0 |= bcodeArray[i2 - i1 - i3] << 16; i2++; i2++;
181 SETSHIFTEDINT (t, i2_sav, x0);
189 while ((i1 < pw_len) || (i3 < salt_len))
191 if (i1 < pw_len) // max 8
193 if (GETCHAR (saved_key, 15 - i1) & 1)
195 PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
200 PUTCHAR (t, i2, GETCHAR (w, i1));
207 PUTCHAR (t, i2, GETCHAR (s, i3));
213 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
226 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
235 static void m07700m (u32x w0[4], u32x w1[4], u32x w2[4], u32x 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)
241 const u32 gid = get_global_id (0);
242 const u32 lid = get_local_id (0);
244 w0[0] = sapb_trans (w0[0]);
245 w0[1] = sapb_trans (w0[1]);
253 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
254 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
255 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
257 salt_buf0[0] = sapb_trans (salt_buf0[0]);
258 salt_buf0[1] = sapb_trans (salt_buf0[1]);
259 salt_buf0[2] = sapb_trans (salt_buf0[2]);
261 const u32 salt_len = salt_bufs[salt_pos].salt_len;
265 s0[0] = salt_buf0[0];
266 s0[1] = salt_buf0[1];
267 s0[2] = salt_buf0[2];
291 switch_buffer_by_offset (s0, s1, s2, s3, pw_len);
293 const u32 pw_salt_len = pw_len + salt_len;
301 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
303 const u32 w0r = sapb_trans (bfs_buf[il_pos].i);
309 t[ 0] = s0[0] | w0[0];
310 t[ 1] = s0[1] | w0[1];
323 t[14] = pw_salt_len * 8;
326 PUTCHAR (t, pw_salt_len, 0x80);
337 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
338 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
339 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
340 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
341 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
342 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
343 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
344 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
345 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
346 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
347 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
348 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
349 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
350 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
351 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
352 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
354 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
355 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
356 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
357 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
358 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
359 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
360 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
361 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
362 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
363 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
364 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
365 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
366 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
367 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
368 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
369 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
371 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
372 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
373 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
374 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
375 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
376 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
377 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
378 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
379 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
380 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
381 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
382 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
383 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
384 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
385 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
386 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
388 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
389 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
390 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
391 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
392 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
393 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
394 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
395 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
396 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
397 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
398 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
399 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
400 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
401 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
402 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
403 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
410 const u32x sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
412 SETSHIFTEDINT (t, sum20, 0x80);
421 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
422 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
423 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
424 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
425 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
426 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
427 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
428 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
429 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
430 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
431 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
432 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
433 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
434 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
435 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
436 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
438 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
439 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
440 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
441 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
442 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
443 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
444 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
445 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
446 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
447 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
448 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
449 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
450 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
451 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
452 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
453 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
455 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
456 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
457 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
458 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
459 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
460 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
461 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
462 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
463 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
464 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
465 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
466 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
467 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
468 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
469 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
470 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
472 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
473 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
474 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
475 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
476 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
477 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
478 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
479 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
480 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
481 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
482 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
483 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
484 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
485 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
486 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
487 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
502 #include VECT_COMPARE_M
506 static void m07700s (u32x w0[4], u32x w1[4], u32x w2[4], u32x 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)
512 const u32 gid = get_global_id (0);
513 const u32 lid = get_local_id (0);
515 w0[0] = sapb_trans (w0[0]);
516 w0[1] = sapb_trans (w0[1]);
524 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
525 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
526 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
528 salt_buf0[0] = sapb_trans (salt_buf0[0]);
529 salt_buf0[1] = sapb_trans (salt_buf0[1]);
530 salt_buf0[2] = sapb_trans (salt_buf0[2]);
532 const u32 salt_len = salt_bufs[salt_pos].salt_len;
536 s0[0] = salt_buf0[0];
537 s0[1] = salt_buf0[1];
538 s0[2] = salt_buf0[2];
562 switch_buffer_by_offset (s0, s1, s2, s3, pw_len);
564 const u32 pw_salt_len = pw_len + salt_len;
570 const u32 search[4] =
572 digests_buf[digests_offset].digest_buf[DGST_R0],
573 digests_buf[digests_offset].digest_buf[DGST_R1],
574 digests_buf[digests_offset].digest_buf[DGST_R2],
575 digests_buf[digests_offset].digest_buf[DGST_R3]
584 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
586 const u32 w0r = sapb_trans (bfs_buf[il_pos].i);
592 t[ 0] = s0[0] | w0[0];
593 t[ 1] = s0[1] | w0[1];
606 t[14] = pw_salt_len * 8;
609 PUTCHAR (t, pw_salt_len, 0x80);
620 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
621 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
622 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
623 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
624 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
625 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
626 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
627 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
628 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
629 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
630 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
631 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
632 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
633 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
634 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
635 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
637 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
638 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
639 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
640 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
641 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
642 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
643 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
644 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
645 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
646 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
647 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
648 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
649 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
650 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
651 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
652 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
654 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
655 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
656 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
657 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
658 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
659 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
660 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
661 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
662 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
663 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
664 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
665 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
666 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
667 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
668 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
669 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
671 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
672 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
673 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
674 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
675 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
676 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
677 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
678 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
679 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
680 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
681 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
682 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
683 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
684 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
685 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
686 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
693 const u32x sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
695 SETSHIFTEDINT (t, sum20, 0x80);
704 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
705 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
706 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
707 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
708 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
709 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
710 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
711 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
712 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
713 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
714 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
715 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
716 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
717 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
718 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
719 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
721 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
722 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
723 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
724 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
725 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
726 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
727 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
728 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
729 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
730 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
731 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
732 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
733 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
734 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
735 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
736 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
738 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
739 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
740 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
741 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
742 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
743 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
744 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
745 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
746 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
747 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
748 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
749 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
750 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
751 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
752 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
753 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
755 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
756 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
757 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
758 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
759 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
760 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
761 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
762 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
763 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
764 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
765 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
766 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
767 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
768 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
769 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
770 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
785 #include VECT_COMPARE_S
789 __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)
795 const u32 gid = get_global_id (0);
797 if (gid >= gid_max) return;
803 const u32 lid = get_local_id (0);
807 w0[0] = pws[gid].i[ 0];
808 w0[1] = pws[gid].i[ 1];
809 w0[2] = pws[gid].i[ 2];
810 w0[3] = pws[gid].i[ 3];
833 const u32 pw_len = pws[gid].pw_len;
839 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);
842 __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)
848 const u32 gid = get_global_id (0);
850 if (gid >= gid_max) return;
856 const u32 lid = get_local_id (0);
860 w0[0] = pws[gid].i[ 0];
861 w0[1] = pws[gid].i[ 1];
862 w0[2] = pws[gid].i[ 2];
863 w0[3] = pws[gid].i[ 3];
867 w1[0] = pws[gid].i[ 4];
868 w1[1] = pws[gid].i[ 5];
869 w1[2] = pws[gid].i[ 6];
870 w1[3] = pws[gid].i[ 7];
886 const u32 pw_len = pws[gid].pw_len;
892 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);
895 __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)
899 __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)
905 const u32 gid = get_global_id (0);
907 if (gid >= gid_max) return;
913 const u32 lid = get_local_id (0);
917 w0[0] = pws[gid].i[ 0];
918 w0[1] = pws[gid].i[ 1];
919 w0[2] = pws[gid].i[ 2];
920 w0[3] = pws[gid].i[ 3];
943 const u32 pw_len = pws[gid].pw_len;
949 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);
952 __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)
958 const u32 gid = get_global_id (0);
960 if (gid >= gid_max) return;
966 const u32 lid = get_local_id (0);
970 w0[0] = pws[gid].i[ 0];
971 w0[1] = pws[gid].i[ 1];
972 w0[2] = pws[gid].i[ 2];
973 w0[3] = pws[gid].i[ 3];
977 w1[0] = pws[gid].i[ 4];
978 w1[1] = pws[gid].i[ 5];
979 w1[2] = pws[gid].i[ 6];
980 w1[3] = pws[gid].i[ 7];
996 const u32 pw_len = pws[gid].pw_len;
1002 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);
1005 __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)