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"
16 #include "inc_simd.cl"
18 #define GETCHAR(a,p) (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
19 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
21 #define SETSHIFTEDINT(a,n,v) \
23 const u32 s = ((n) & 3) * 8; \
24 const u64 x = (u64) (v) << s; \
25 (a)[((n)/4)+0] &= ~(0xff << ((n & 3) * 8)); \
26 (a)[((n)/4)+0] |= x; \
27 (a)[((n)/4)+1] = x >> 32; \
30 __constant u32 sapb_trans_tbl[256] =
32 // first value hack for 0 byte as part of an optimization
33 0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
34 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
35 0x3f, 0x40, 0x41, 0x50, 0x43, 0x44, 0x45, 0x4b, 0x47, 0x48, 0x4d, 0x4e, 0x54, 0x51, 0x53, 0x46,
36 0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x3b, 0x3c, 0x3d, 0x3e, 0x56, 0x55, 0x5c, 0x49, 0x5d, 0x4a,
37 0x42, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
38 0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x58, 0x5b, 0x59, 0xff, 0x52,
39 0x4c, 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, 0x57, 0x5e, 0x5a, 0x4f, 0xff,
41 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
42 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 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
51 __constant u32 bcodeArray[48] =
53 0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
54 0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
55 0xe2, 0xb7, 0x33, 0x71, 0x8b, 0x9f, 0x5d, 0x01, 0x44, 0x70, 0xae, 0x11, 0xef, 0x28, 0xf0, 0x0d
58 u32 sapb_trans (const u32 in)
62 out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
63 out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
64 out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
65 out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
70 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])
89 u32 sum20 = ((a >> 24) & 3)
97 const u32 w[2] = { w0[0], w0[1] };
99 const u32 s[3] = { salt_buf0[0], salt_buf0[1], salt_buf0[2] };
101 u32 saved_key[4] = { a, b, c, d };
107 // we can assume this because the password must be at least 3
108 // and the username must be at least 1 so we can save the if ()
114 t0 |= bcodeArray[47] << 0;
115 t0 |= (w[0] & 0xff) << 8;
116 t0 |= (s[0] & 0xff) << 16;
117 t0 |= bcodeArray[ 1] << 24;
125 t0 |= (w[0] & 0xff) << 0;
126 t0 |= (s[0] & 0xff) << 8;
127 t0 |= bcodeArray[ 0] << 16;
136 // because the following code can increase i2 by a maximum of 5,
137 // there is an overflow potential of 4 before it comes to the next test for i2 >= sum20
138 // we need to truncate in that case
140 while ((i1 < pw_len) && (i3 < salt_len))
146 if (GETCHAR (saved_key, 15 - i1) & 1)
148 x0 |= bcodeArray[48 - 1 - i1] << 0; i2++;
149 x0 |= GETCHAR (w, i1) << 8; i2++; i1++;
150 x0 |= GETCHAR (s, i3) << 16; i2++; i3++;
151 x0 |= bcodeArray[i2 - i1 - i3] << 24; i2++; i2++;
155 x0 |= GETCHAR (w, i1) << 0; i2++; i1++;
156 x0 |= GETCHAR (s, i3) << 8; i2++; i3++;
157 x0 |= bcodeArray[i2 - i1 - i3] << 16; i2++; i2++;
160 SETSHIFTEDINT (t, i2_sav, x0);
168 while ((i1 < pw_len) || (i3 < salt_len))
170 if (i1 < pw_len) // max 8
172 if (GETCHAR (saved_key, 15 - i1) & 1)
174 PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
179 PUTCHAR (t, i2, GETCHAR (w, i1));
186 PUTCHAR (t, i2, GETCHAR (s, i3));
192 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
205 PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
214 void m07700m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __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)
220 const u32 gid = get_global_id (0);
221 const u32 lid = get_local_id (0);
223 w0[0] = sapb_trans (w0[0]);
224 w0[1] = sapb_trans (w0[1]);
232 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
233 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
234 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
236 salt_buf0[0] = sapb_trans (salt_buf0[0]);
237 salt_buf0[1] = sapb_trans (salt_buf0[1]);
238 salt_buf0[2] = sapb_trans (salt_buf0[2]);
240 const u32 salt_len = salt_bufs[salt_pos].salt_len;
247 s0[0] = salt_buf0[0];
248 s0[1] = salt_buf0[1];
249 s0[2] = salt_buf0[2];
264 switch_buffer_by_offset_le (s0, s1, s2, s3, pw_len);
266 const u32 pw_salt_len = pw_len + salt_len;
274 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
276 const u32x w0r = sapb_trans (ix_create_bft (bfs_buf, il_pos));
278 const u32x w0lr = w0l | w0r;
284 t[ 0] = s0[0] | w0[0];
285 t[ 1] = s0[1] | w0[1];
298 t[14] = pw_salt_len * 8;
301 PUTCHAR (t, pw_salt_len, 0x80);
312 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
313 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
314 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
315 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
316 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
317 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
318 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
319 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
320 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
321 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
322 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
323 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
324 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
325 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
326 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
327 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
329 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
330 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
331 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
332 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
333 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
334 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
335 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
336 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
337 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
338 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
339 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
340 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
341 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
342 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
343 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
344 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
346 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
347 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
348 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
349 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
350 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
351 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
352 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
353 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
354 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
355 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
356 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
357 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
358 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
359 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
360 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
361 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
363 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
364 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
365 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
366 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
367 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
368 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
369 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
370 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
371 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
372 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
373 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
374 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
375 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
376 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
377 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
378 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
385 const u32 sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
387 SETSHIFTEDINT (t, sum20, 0x80);
396 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
397 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
398 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
399 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
400 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
401 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
402 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
403 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
404 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
405 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
406 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
407 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
408 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
409 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
410 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
411 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
413 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
414 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
415 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
416 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
417 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
418 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
419 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
420 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
421 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
422 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
423 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
424 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
425 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
426 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
427 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
428 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
430 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
431 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
432 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
433 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
434 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
435 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
436 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
437 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
438 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
439 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
440 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
441 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
442 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
443 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
444 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
445 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
447 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
448 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
449 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
450 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
451 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
452 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
453 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
454 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
455 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
456 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
457 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
458 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
459 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
460 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
461 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
462 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
474 COMPARE_M_SIMD (a, b, c, d);
478 void m07700s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __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)
484 const u32 gid = get_global_id (0);
485 const u32 lid = get_local_id (0);
487 w0[0] = sapb_trans (w0[0]);
488 w0[1] = sapb_trans (w0[1]);
496 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
497 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
498 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
500 salt_buf0[0] = sapb_trans (salt_buf0[0]);
501 salt_buf0[1] = sapb_trans (salt_buf0[1]);
502 salt_buf0[2] = sapb_trans (salt_buf0[2]);
504 const u32 salt_len = salt_bufs[salt_pos].salt_len;
511 s0[0] = salt_buf0[0];
512 s0[1] = salt_buf0[1];
513 s0[2] = salt_buf0[2];
528 switch_buffer_by_offset_le (s0, s1, s2, s3, pw_len);
530 const u32 pw_salt_len = pw_len + salt_len;
536 const u32 search[4] =
538 digests_buf[digests_offset].digest_buf[DGST_R0],
539 digests_buf[digests_offset].digest_buf[DGST_R1],
550 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
552 const u32x w0r = sapb_trans (ix_create_bft (bfs_buf, il_pos));
554 const u32x w0lr = w0l | w0r;
560 t[ 0] = s0[0] | w0[0];
561 t[ 1] = s0[1] | w0[1];
574 t[14] = pw_salt_len * 8;
577 PUTCHAR (t, pw_salt_len, 0x80);
588 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
589 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
590 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
591 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
592 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
593 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
594 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
595 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
596 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
597 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
598 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
599 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
600 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
601 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
602 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
603 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
605 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
606 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
607 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
608 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
609 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
610 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
611 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
612 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
613 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
614 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
615 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
616 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
617 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
618 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
619 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
620 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
622 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
623 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
624 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
625 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
626 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
627 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
628 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
629 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
630 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
631 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
632 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
633 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
634 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
635 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
636 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
637 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
639 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
640 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
641 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
642 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
643 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
644 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
645 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
646 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
647 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
648 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
649 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
650 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
651 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
652 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
653 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
654 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
661 const u32 sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
663 SETSHIFTEDINT (t, sum20, 0x80);
672 MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
673 MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
674 MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
675 MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
676 MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
677 MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
678 MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
679 MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
680 MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
681 MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
682 MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
683 MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
684 MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
685 MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
686 MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
687 MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
689 MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
690 MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
691 MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
692 MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
693 MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
694 MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
695 MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
696 MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
697 MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
698 MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
699 MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
700 MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
701 MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
702 MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
703 MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
704 MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
706 MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
707 MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
708 MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
709 MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
710 MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
711 MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
712 MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
713 MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
714 MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
715 MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
716 MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
717 MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
718 MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
719 MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
720 MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
721 MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
723 MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
724 MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
725 MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
726 MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
727 MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
728 MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
729 MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
730 MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
731 MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
732 MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
733 MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
734 MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
735 MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
736 MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
737 MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
738 MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
750 COMPARE_S_SIMD (a, b, c, d);
754 __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)
760 const u32 gid = get_global_id (0);
762 if (gid >= gid_max) return;
768 //const u32 lid = get_local_id (0);
772 w0[0] = pws[gid].i[ 0];
773 w0[1] = pws[gid].i[ 1];
774 w0[2] = pws[gid].i[ 2];
775 w0[3] = pws[gid].i[ 3];
798 const u32 pw_len = pws[gid].pw_len;
804 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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
807 __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)
813 const u32 gid = get_global_id (0);
815 if (gid >= gid_max) return;
821 //const u32 lid = get_local_id (0);
825 w0[0] = pws[gid].i[ 0];
826 w0[1] = pws[gid].i[ 1];
827 w0[2] = pws[gid].i[ 2];
828 w0[3] = pws[gid].i[ 3];
832 w1[0] = pws[gid].i[ 4];
833 w1[1] = pws[gid].i[ 5];
834 w1[2] = pws[gid].i[ 6];
835 w1[3] = pws[gid].i[ 7];
851 const u32 pw_len = pws[gid].pw_len;
857 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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
860 __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)
864 __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)
870 const u32 gid = get_global_id (0);
872 if (gid >= gid_max) return;
878 //const u32 lid = get_local_id (0);
882 w0[0] = pws[gid].i[ 0];
883 w0[1] = pws[gid].i[ 1];
884 w0[2] = pws[gid].i[ 2];
885 w0[3] = pws[gid].i[ 3];
908 const u32 pw_len = pws[gid].pw_len;
914 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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
917 __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)
923 const u32 gid = get_global_id (0);
925 if (gid >= gid_max) return;
931 //const u32 lid = get_local_id (0);
935 w0[0] = pws[gid].i[ 0];
936 w0[1] = pws[gid].i[ 1];
937 w0[2] = pws[gid].i[ 2];
938 w0[3] = pws[gid].i[ 3];
942 w1[0] = pws[gid].i[ 4];
943 w1[1] = pws[gid].i[ 5];
944 w1[2] = pws[gid].i[ 6];
945 w1[3] = pws[gid].i[ 7];
961 const u32 pw_len = pws[gid].pw_len;
967 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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
970 __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)