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"
38 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
39 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
50 static void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
55 rc4_key->S[i] = rc4_key->S[j];
59 static void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
64 __local u32 *ptr = (__local u32 *) rc4_key->S;
67 for (u32 i = 0; i < 64; i++)
75 for (u32 i = 0; i < 16; i++)
83 j += rc4_key->S[idx] + (v >> 0); swap (rc4_key, idx, j); idx++;
84 j += rc4_key->S[idx] + (v >> 8); swap (rc4_key, idx, j); idx++;
85 j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
86 j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
90 j += rc4_key->S[idx] + (v >> 0); swap (rc4_key, idx, j); idx++;
91 j += rc4_key->S[idx] + (v >> 8); swap (rc4_key, idx, j); idx++;
92 j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
93 j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
97 j += rc4_key->S[idx] + (v >> 0); swap (rc4_key, idx, j); idx++;
98 j += rc4_key->S[idx] + (v >> 8); swap (rc4_key, idx, j); idx++;
99 j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
100 j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
104 j += rc4_key->S[idx] + (v >> 0); swap (rc4_key, idx, j); idx++;
105 j += rc4_key->S[idx] + (v >> 8); swap (rc4_key, idx, j); idx++;
106 j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
107 j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
111 static u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
114 for (u32 k = 0; k < 4; k++)
123 swap (rc4_key, i, j);
125 idx = rc4_key->S[i] + rc4_key->S[j];
127 xor4 |= rc4_key->S[idx] << 0;
132 swap (rc4_key, i, j);
134 idx = rc4_key->S[i] + rc4_key->S[j];
136 xor4 |= rc4_key->S[idx] << 8;
141 swap (rc4_key, i, j);
143 idx = rc4_key->S[i] + rc4_key->S[j];
145 xor4 |= rc4_key->S[idx] << 16;
150 swap (rc4_key, i, j);
152 idx = rc4_key->S[i] + rc4_key->S[j];
154 xor4 |= rc4_key->S[idx] << 24;
156 out[k] = in[k] ^ xor4;
162 static void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
186 MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
187 MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
188 MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
189 MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
190 MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
191 MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
192 MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
193 MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
194 MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
195 MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
196 MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
197 MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
198 MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
199 MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
200 MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
201 MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
203 MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
204 MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
205 MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
206 MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
207 MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
208 MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
209 MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
210 MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
211 MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
212 MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
213 MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
214 MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
215 MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
216 MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
217 MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
218 MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
220 MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
221 MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
222 MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
223 MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
224 MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
225 MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
226 MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
227 MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
228 MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
229 MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
230 MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
231 MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
232 MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
233 MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
234 MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
235 MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
237 MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
238 MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
239 MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
240 MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
241 MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
242 MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
243 MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
244 MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
245 MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
246 MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
247 MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
248 MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
249 MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
250 MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
251 MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
252 MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
260 static void m09700m (__local RC4_KEY rc4_keys[64], 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 oldoffice01_t *oldoffice01_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)
266 const u32 gid = get_global_id (0);
267 const u32 lid = get_local_id (0);
269 __local RC4_KEY *rc4_key = &rc4_keys[lid];
280 salt_buf_t0[0] = salt_bufs[salt_pos].salt_buf[0];
281 salt_buf_t0[1] = salt_bufs[salt_pos].salt_buf[1];
282 salt_buf_t0[2] = salt_bufs[salt_pos].salt_buf[2];
283 salt_buf_t0[3] = salt_bufs[salt_pos].salt_buf[3];
285 salt_buf_t1[0] = salt_buf_t0[0] << 8;
286 salt_buf_t1[1] = salt_buf_t0[0] >> 24 | salt_buf_t0[1] << 8;
287 salt_buf_t1[2] = salt_buf_t0[1] >> 24 | salt_buf_t0[2] << 8;
288 salt_buf_t1[3] = salt_buf_t0[2] >> 24 | salt_buf_t0[3] << 8;
289 salt_buf_t1[4] = salt_buf_t0[3] >> 24;
291 salt_buf_t2[0] = salt_buf_t0[0] << 16;
292 salt_buf_t2[1] = salt_buf_t0[0] >> 16 | salt_buf_t0[1] << 16;
293 salt_buf_t2[2] = salt_buf_t0[1] >> 16 | salt_buf_t0[2] << 16;
294 salt_buf_t2[3] = salt_buf_t0[2] >> 16 | salt_buf_t0[3] << 16;
295 salt_buf_t2[4] = salt_buf_t0[3] >> 16;
297 salt_buf_t3[0] = salt_buf_t0[0] << 24;
298 salt_buf_t3[1] = salt_buf_t0[0] >> 8 | salt_buf_t0[1] << 24;
299 salt_buf_t3[2] = salt_buf_t0[1] >> 8 | salt_buf_t0[2] << 24;
300 salt_buf_t3[3] = salt_buf_t0[2] >> 8 | salt_buf_t0[3] << 24;
301 salt_buf_t3[4] = salt_buf_t0[3] >> 8;
303 const u32 salt_len = 16;
309 const u32 version = oldoffice01_bufs[salt_pos].version;
311 u32 encryptedVerifier[4];
313 encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
314 encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
315 encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
316 encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
324 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
326 const u32 w0r = bfs_buf[il_pos].i;
349 w3_t[2] = pw_len * 8;
353 u32x digest_t1[2]; // need only first 5 byte
357 digest_t0[0] = MD5M_A;
358 digest_t0[1] = MD5M_B;
359 digest_t0[2] = MD5M_C;
360 digest_t0[3] = MD5M_D;
362 md5_transform (w0_t, w1_t, w2_t, w3_t, digest_t0);
364 // prepare 16 * 21 buffer stuff
375 digest_t0[0] &= 0xffffffff;
376 digest_t0[1] &= 0x000000ff;
377 digest_t0[2] &= 0x00000000;
378 digest_t0[3] &= 0x00000000;
380 digest_t1[0] = digest_t0[0] << 8;
381 digest_t1[1] = digest_t0[0] >> 24 | digest_t0[1] << 8;
383 digest_t2[0] = digest_t0[0] << 16;
384 digest_t2[1] = digest_t0[0] >> 16 | digest_t0[1] << 16;
386 digest_t3[0] = digest_t0[0] << 24;
387 digest_t3[1] = digest_t0[0] >> 8 | digest_t0[1] << 24;
389 // generate the 16 * 21 buffer
409 w0_t[0] = digest_t0[0];
410 w0_t[1] = digest_t0[1];
413 w0_t[1] |= salt_buf_t1[0];
414 w0_t[2] = salt_buf_t1[1];
415 w0_t[3] = salt_buf_t1[2];
416 w1_t[0] = salt_buf_t1[3];
417 w1_t[1] = salt_buf_t1[4];
420 w1_t[1] |= digest_t1[0];
421 w1_t[2] = digest_t1[1];
424 w1_t[2] |= salt_buf_t2[0];
425 w1_t[3] = salt_buf_t2[1];
426 w2_t[0] = salt_buf_t2[2];
427 w2_t[1] = salt_buf_t2[3];
428 w2_t[2] = salt_buf_t2[4];
431 w2_t[2] |= digest_t2[0];
432 w2_t[3] = digest_t2[1];
435 w2_t[3] |= salt_buf_t3[0];
436 w3_t[0] = salt_buf_t3[1];
437 w3_t[1] = salt_buf_t3[2];
438 w3_t[2] = salt_buf_t3[3];
439 w3_t[3] = salt_buf_t3[4];
443 w3_t[3] |= digest_t3[0];
445 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
465 w0_t[0] = digest_t3[1];
468 w0_t[1] = salt_buf_t0[0];
469 w0_t[2] = salt_buf_t0[1];
470 w0_t[3] = salt_buf_t0[2];
471 w1_t[0] = salt_buf_t0[3];
474 w1_t[1] = digest_t0[0];
475 w1_t[2] = digest_t0[1];
478 w1_t[2] |= salt_buf_t1[0];
479 w1_t[3] = salt_buf_t1[1];
480 w2_t[0] = salt_buf_t1[2];
481 w2_t[1] = salt_buf_t1[3];
482 w2_t[2] = salt_buf_t1[4];
485 w2_t[2] |= digest_t1[0];
486 w2_t[3] = digest_t1[1];
489 w2_t[3] |= salt_buf_t2[0];
490 w3_t[0] = salt_buf_t2[1];
491 w3_t[1] = salt_buf_t2[2];
492 w3_t[2] = salt_buf_t2[3];
493 w3_t[3] = salt_buf_t2[4];
496 w3_t[3] |= digest_t2[0];
498 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
518 w0_t[0] = digest_t2[1];
521 w0_t[0] |= salt_buf_t3[0];
522 w0_t[1] = salt_buf_t3[1];
523 w0_t[2] = salt_buf_t3[2];
524 w0_t[3] = salt_buf_t3[3];
525 w1_t[0] = salt_buf_t3[4];
528 w1_t[0] |= digest_t3[0];
529 w1_t[1] = digest_t3[1];
532 w1_t[2] = salt_buf_t0[0];
533 w1_t[3] = salt_buf_t0[1];
534 w2_t[0] = salt_buf_t0[2];
535 w2_t[1] = salt_buf_t0[3];
538 w2_t[2] = digest_t0[0];
539 w2_t[3] = digest_t0[1];
542 w2_t[3] |= salt_buf_t1[0];
543 w3_t[0] = salt_buf_t1[1];
544 w3_t[1] = salt_buf_t1[2];
545 w3_t[2] = salt_buf_t1[3];
546 w3_t[3] = salt_buf_t1[4];
549 w3_t[3] |= digest_t1[0];
551 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
571 w0_t[0] = digest_t1[1];
574 w0_t[0] |= salt_buf_t2[0];
575 w0_t[1] = salt_buf_t2[1];
576 w0_t[2] = salt_buf_t2[2];
577 w0_t[3] = salt_buf_t2[3];
578 w1_t[0] = salt_buf_t2[4];
581 w1_t[0] |= digest_t2[0];
582 w1_t[1] = digest_t2[1];
585 w1_t[1] |= salt_buf_t3[0];
586 w1_t[2] = salt_buf_t3[1];
587 w1_t[3] = salt_buf_t3[2];
588 w2_t[0] = salt_buf_t3[3];
589 w2_t[1] = salt_buf_t3[4];
592 w2_t[1] |= digest_t3[0];
593 w2_t[2] = digest_t3[1];
596 w2_t[3] = salt_buf_t0[0];
597 w3_t[0] = salt_buf_t0[1];
598 w3_t[1] = salt_buf_t0[2];
599 w3_t[2] = salt_buf_t0[3];
602 w3_t[3] = digest_t0[0];
604 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
624 w0_t[0] = digest_t0[1];
627 w0_t[0] |= salt_buf_t1[0];
628 w0_t[1] = salt_buf_t1[1];
629 w0_t[2] = salt_buf_t1[2];
630 w0_t[3] = salt_buf_t1[3];
631 w1_t[0] = salt_buf_t1[4];
634 w1_t[0] |= digest_t1[0];
635 w1_t[1] = digest_t1[1];
638 w1_t[1] |= salt_buf_t2[0];
639 w1_t[2] = salt_buf_t2[1];
640 w1_t[3] = salt_buf_t2[2];
641 w2_t[0] = salt_buf_t2[3];
642 w2_t[1] = salt_buf_t2[4];
645 w2_t[1] |= digest_t2[0];
646 w2_t[2] = digest_t2[1];
649 w2_t[2] |= salt_buf_t3[0];
650 w2_t[3] = salt_buf_t3[1];
651 w3_t[0] = salt_buf_t3[2];
652 w3_t[1] = salt_buf_t3[3];
653 w3_t[2] = salt_buf_t3[4];
656 w3_t[2] |= digest_t3[0];
657 w3_t[3] = digest_t3[1];
659 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
661 w0_t[0] = salt_buf_t0[0];
662 w0_t[1] = salt_buf_t0[1];
663 w0_t[2] = salt_buf_t0[2];
664 w0_t[3] = salt_buf_t0[3];
675 w3_t[2] = 21 * 16 * 8;
678 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
680 // now the 40 bit input for the MD5 which then will generate the RC4 key, so it's precomputable!
683 w0_t[1] = digest[1] & 0xff;
704 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
715 rc4_init_16 (rc4_key, key);
719 u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
743 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
745 rc4_next_16 (rc4_key, 16, j, digest, out);
747 const u32x r0 = out[0];
748 const u32x r1 = out[1];
749 const u32x r2 = out[2];
750 const u32x r3 = out[3];
752 #include VECT_COMPARE_M
756 static void m09700s (__local RC4_KEY rc4_keys[64], 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 oldoffice01_t *oldoffice01_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)
762 const u32 gid = get_global_id (0);
763 const u32 lid = get_local_id (0);
765 __local RC4_KEY *rc4_key = &rc4_keys[lid];
771 const u32 search[4] =
773 digests_buf[digests_offset].digest_buf[DGST_R0],
774 digests_buf[digests_offset].digest_buf[DGST_R1],
775 digests_buf[digests_offset].digest_buf[DGST_R2],
776 digests_buf[digests_offset].digest_buf[DGST_R3]
788 salt_buf_t0[0] = salt_bufs[salt_pos].salt_buf[0];
789 salt_buf_t0[1] = salt_bufs[salt_pos].salt_buf[1];
790 salt_buf_t0[2] = salt_bufs[salt_pos].salt_buf[2];
791 salt_buf_t0[3] = salt_bufs[salt_pos].salt_buf[3];
793 salt_buf_t1[0] = salt_buf_t0[0] << 8;
794 salt_buf_t1[1] = salt_buf_t0[0] >> 24 | salt_buf_t0[1] << 8;
795 salt_buf_t1[2] = salt_buf_t0[1] >> 24 | salt_buf_t0[2] << 8;
796 salt_buf_t1[3] = salt_buf_t0[2] >> 24 | salt_buf_t0[3] << 8;
797 salt_buf_t1[4] = salt_buf_t0[3] >> 24;
799 salt_buf_t2[0] = salt_buf_t0[0] << 16;
800 salt_buf_t2[1] = salt_buf_t0[0] >> 16 | salt_buf_t0[1] << 16;
801 salt_buf_t2[2] = salt_buf_t0[1] >> 16 | salt_buf_t0[2] << 16;
802 salt_buf_t2[3] = salt_buf_t0[2] >> 16 | salt_buf_t0[3] << 16;
803 salt_buf_t2[4] = salt_buf_t0[3] >> 16;
805 salt_buf_t3[0] = salt_buf_t0[0] << 24;
806 salt_buf_t3[1] = salt_buf_t0[0] >> 8 | salt_buf_t0[1] << 24;
807 salt_buf_t3[2] = salt_buf_t0[1] >> 8 | salt_buf_t0[2] << 24;
808 salt_buf_t3[3] = salt_buf_t0[2] >> 8 | salt_buf_t0[3] << 24;
809 salt_buf_t3[4] = salt_buf_t0[3] >> 8;
811 const u32 salt_len = 16;
817 const u32 version = oldoffice01_bufs[salt_pos].version;
819 u32 encryptedVerifier[4];
821 encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
822 encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
823 encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
824 encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
832 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
834 const u32 w0r = bfs_buf[il_pos].i;
857 w3_t[2] = pw_len * 8;
861 u32x digest_t1[2]; // need only first 5 byte
865 digest_t0[0] = MD5M_A;
866 digest_t0[1] = MD5M_B;
867 digest_t0[2] = MD5M_C;
868 digest_t0[3] = MD5M_D;
870 md5_transform (w0_t, w1_t, w2_t, w3_t, digest_t0);
872 // prepare 16 * 21 buffer stuff
883 digest_t0[0] &= 0xffffffff;
884 digest_t0[1] &= 0x000000ff;
885 digest_t0[2] &= 0x00000000;
886 digest_t0[3] &= 0x00000000;
888 digest_t1[0] = digest_t0[0] << 8;
889 digest_t1[1] = digest_t0[0] >> 24 | digest_t0[1] << 8;
891 digest_t2[0] = digest_t0[0] << 16;
892 digest_t2[1] = digest_t0[0] >> 16 | digest_t0[1] << 16;
894 digest_t3[0] = digest_t0[0] << 24;
895 digest_t3[1] = digest_t0[0] >> 8 | digest_t0[1] << 24;
897 // generate the 16 * 21 buffer
917 w0_t[0] = digest_t0[0];
918 w0_t[1] = digest_t0[1];
921 w0_t[1] |= salt_buf_t1[0];
922 w0_t[2] = salt_buf_t1[1];
923 w0_t[3] = salt_buf_t1[2];
924 w1_t[0] = salt_buf_t1[3];
925 w1_t[1] = salt_buf_t1[4];
928 w1_t[1] |= digest_t1[0];
929 w1_t[2] = digest_t1[1];
932 w1_t[2] |= salt_buf_t2[0];
933 w1_t[3] = salt_buf_t2[1];
934 w2_t[0] = salt_buf_t2[2];
935 w2_t[1] = salt_buf_t2[3];
936 w2_t[2] = salt_buf_t2[4];
939 w2_t[2] |= digest_t2[0];
940 w2_t[3] = digest_t2[1];
943 w2_t[3] |= salt_buf_t3[0];
944 w3_t[0] = salt_buf_t3[1];
945 w3_t[1] = salt_buf_t3[2];
946 w3_t[2] = salt_buf_t3[3];
947 w3_t[3] = salt_buf_t3[4];
951 w3_t[3] |= digest_t3[0];
953 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
973 w0_t[0] = digest_t3[1];
976 w0_t[1] = salt_buf_t0[0];
977 w0_t[2] = salt_buf_t0[1];
978 w0_t[3] = salt_buf_t0[2];
979 w1_t[0] = salt_buf_t0[3];
982 w1_t[1] = digest_t0[0];
983 w1_t[2] = digest_t0[1];
986 w1_t[2] |= salt_buf_t1[0];
987 w1_t[3] = salt_buf_t1[1];
988 w2_t[0] = salt_buf_t1[2];
989 w2_t[1] = salt_buf_t1[3];
990 w2_t[2] = salt_buf_t1[4];
993 w2_t[2] |= digest_t1[0];
994 w2_t[3] = digest_t1[1];
997 w2_t[3] |= salt_buf_t2[0];
998 w3_t[0] = salt_buf_t2[1];
999 w3_t[1] = salt_buf_t2[2];
1000 w3_t[2] = salt_buf_t2[3];
1001 w3_t[3] = salt_buf_t2[4];
1004 w3_t[3] |= digest_t2[0];
1006 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1026 w0_t[0] = digest_t2[1];
1029 w0_t[0] |= salt_buf_t3[0];
1030 w0_t[1] = salt_buf_t3[1];
1031 w0_t[2] = salt_buf_t3[2];
1032 w0_t[3] = salt_buf_t3[3];
1033 w1_t[0] = salt_buf_t3[4];
1036 w1_t[0] |= digest_t3[0];
1037 w1_t[1] = digest_t3[1];
1040 w1_t[2] = salt_buf_t0[0];
1041 w1_t[3] = salt_buf_t0[1];
1042 w2_t[0] = salt_buf_t0[2];
1043 w2_t[1] = salt_buf_t0[3];
1046 w2_t[2] = digest_t0[0];
1047 w2_t[3] = digest_t0[1];
1050 w2_t[3] |= salt_buf_t1[0];
1051 w3_t[0] = salt_buf_t1[1];
1052 w3_t[1] = salt_buf_t1[2];
1053 w3_t[2] = salt_buf_t1[3];
1054 w3_t[3] = salt_buf_t1[4];
1057 w3_t[3] |= digest_t1[0];
1059 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1079 w0_t[0] = digest_t1[1];
1082 w0_t[0] |= salt_buf_t2[0];
1083 w0_t[1] = salt_buf_t2[1];
1084 w0_t[2] = salt_buf_t2[2];
1085 w0_t[3] = salt_buf_t2[3];
1086 w1_t[0] = salt_buf_t2[4];
1089 w1_t[0] |= digest_t2[0];
1090 w1_t[1] = digest_t2[1];
1093 w1_t[1] |= salt_buf_t3[0];
1094 w1_t[2] = salt_buf_t3[1];
1095 w1_t[3] = salt_buf_t3[2];
1096 w2_t[0] = salt_buf_t3[3];
1097 w2_t[1] = salt_buf_t3[4];
1100 w2_t[1] |= digest_t3[0];
1101 w2_t[2] = digest_t3[1];
1104 w2_t[3] = salt_buf_t0[0];
1105 w3_t[0] = salt_buf_t0[1];
1106 w3_t[1] = salt_buf_t0[2];
1107 w3_t[2] = salt_buf_t0[3];
1110 w3_t[3] = digest_t0[0];
1112 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1132 w0_t[0] = digest_t0[1];
1135 w0_t[0] |= salt_buf_t1[0];
1136 w0_t[1] = salt_buf_t1[1];
1137 w0_t[2] = salt_buf_t1[2];
1138 w0_t[3] = salt_buf_t1[3];
1139 w1_t[0] = salt_buf_t1[4];
1142 w1_t[0] |= digest_t1[0];
1143 w1_t[1] = digest_t1[1];
1146 w1_t[1] |= salt_buf_t2[0];
1147 w1_t[2] = salt_buf_t2[1];
1148 w1_t[3] = salt_buf_t2[2];
1149 w2_t[0] = salt_buf_t2[3];
1150 w2_t[1] = salt_buf_t2[4];
1153 w2_t[1] |= digest_t2[0];
1154 w2_t[2] = digest_t2[1];
1157 w2_t[2] |= salt_buf_t3[0];
1158 w2_t[3] = salt_buf_t3[1];
1159 w3_t[0] = salt_buf_t3[2];
1160 w3_t[1] = salt_buf_t3[3];
1161 w3_t[2] = salt_buf_t3[4];
1164 w3_t[2] |= digest_t3[0];
1165 w3_t[3] = digest_t3[1];
1167 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1169 w0_t[0] = salt_buf_t0[0];
1170 w0_t[1] = salt_buf_t0[1];
1171 w0_t[2] = salt_buf_t0[2];
1172 w0_t[3] = salt_buf_t0[3];
1183 w3_t[2] = 21 * 16 * 8;
1186 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1188 // now the 40 bit input for the MD5 which then will generate the RC4 key, so it's precomputable!
1190 w0_t[0] = digest[0];
1191 w0_t[1] = digest[1] & 0xff;
1212 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1223 rc4_init_16 (rc4_key, key);
1227 u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
1251 md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1253 rc4_next_16 (rc4_key, 16, j, digest, out);
1255 const u32x r0 = out[0];
1256 const u32x r1 = out[1];
1257 const u32x r2 = out[2];
1258 const u32x r3 = out[3];
1260 #include VECT_COMPARE_S
1264 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m09700_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 oldoffice01_t *oldoffice01_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)
1270 const u32 gid = get_global_id (0);
1272 if (gid >= gid_max) return;
1276 w0[0] = pws[gid].i[ 0];
1277 w0[1] = pws[gid].i[ 1];
1278 w0[2] = pws[gid].i[ 2];
1279 w0[3] = pws[gid].i[ 3];
1302 const u32 pw_len = pws[gid].pw_len;
1308 __local RC4_KEY rc4_keys[64];
1310 m09700m (rc4_keys, 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, oldoffice01_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);
1313 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m09700_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 oldoffice01_t *oldoffice01_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)
1319 const u32 gid = get_global_id (0);
1321 if (gid >= gid_max) return;
1325 w0[0] = pws[gid].i[ 0];
1326 w0[1] = pws[gid].i[ 1];
1327 w0[2] = pws[gid].i[ 2];
1328 w0[3] = pws[gid].i[ 3];
1332 w1[0] = pws[gid].i[ 4];
1333 w1[1] = pws[gid].i[ 5];
1334 w1[2] = pws[gid].i[ 6];
1335 w1[3] = pws[gid].i[ 7];
1351 const u32 pw_len = pws[gid].pw_len;
1357 __local RC4_KEY rc4_keys[64];
1359 m09700m (rc4_keys, 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, oldoffice01_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);
1362 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m09700_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 oldoffice01_t *oldoffice01_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)
1368 const u32 gid = get_global_id (0);
1370 if (gid >= gid_max) return;
1374 w0[0] = pws[gid].i[ 0];
1375 w0[1] = pws[gid].i[ 1];
1376 w0[2] = pws[gid].i[ 2];
1377 w0[3] = pws[gid].i[ 3];
1381 w1[0] = pws[gid].i[ 4];
1382 w1[1] = pws[gid].i[ 5];
1383 w1[2] = pws[gid].i[ 6];
1384 w1[3] = pws[gid].i[ 7];
1388 w2[0] = pws[gid].i[ 8];
1389 w2[1] = pws[gid].i[ 9];
1390 w2[2] = pws[gid].i[10];
1391 w2[3] = pws[gid].i[11];
1395 w3[0] = pws[gid].i[12];
1396 w3[1] = pws[gid].i[13];
1400 const u32 pw_len = pws[gid].pw_len;
1406 __local RC4_KEY rc4_keys[64];
1408 m09700m (rc4_keys, 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, oldoffice01_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);
1411 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m09700_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 oldoffice01_t *oldoffice01_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)
1417 const u32 gid = get_global_id (0);
1419 if (gid >= gid_max) return;
1423 w0[0] = pws[gid].i[ 0];
1424 w0[1] = pws[gid].i[ 1];
1425 w0[2] = pws[gid].i[ 2];
1426 w0[3] = pws[gid].i[ 3];
1449 const u32 pw_len = pws[gid].pw_len;
1455 __local RC4_KEY rc4_keys[64];
1457 m09700s (rc4_keys, 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, oldoffice01_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);
1460 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m09700_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 oldoffice01_t *oldoffice01_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)
1466 const u32 gid = get_global_id (0);
1468 if (gid >= gid_max) return;
1472 w0[0] = pws[gid].i[ 0];
1473 w0[1] = pws[gid].i[ 1];
1474 w0[2] = pws[gid].i[ 2];
1475 w0[3] = pws[gid].i[ 3];
1479 w1[0] = pws[gid].i[ 4];
1480 w1[1] = pws[gid].i[ 5];
1481 w1[2] = pws[gid].i[ 6];
1482 w1[3] = pws[gid].i[ 7];
1498 const u32 pw_len = pws[gid].pw_len;
1504 __local RC4_KEY rc4_keys[64];
1506 m09700s (rc4_keys, 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, oldoffice01_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);
1509 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m09700_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 oldoffice01_t *oldoffice01_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)
1515 const u32 gid = get_global_id (0);
1517 if (gid >= gid_max) return;
1521 w0[0] = pws[gid].i[ 0];
1522 w0[1] = pws[gid].i[ 1];
1523 w0[2] = pws[gid].i[ 2];
1524 w0[3] = pws[gid].i[ 3];
1528 w1[0] = pws[gid].i[ 4];
1529 w1[1] = pws[gid].i[ 5];
1530 w1[2] = pws[gid].i[ 6];
1531 w1[3] = pws[gid].i[ 7];
1535 w2[0] = pws[gid].i[ 8];
1536 w2[1] = pws[gid].i[ 9];
1537 w2[2] = pws[gid].i[10];
1538 w2[3] = pws[gid].i[11];
1542 w3[0] = pws[gid].i[12];
1543 w3[1] = pws[gid].i[13];
1547 const u32 pw_len = pws[gid].pw_len;
1553 __local RC4_KEY rc4_keys[64];
1555 m09700s (rc4_keys, 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, oldoffice01_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);