2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
24 #include "include/kernel_functions.c"
26 #include "common_nv.c"
29 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
30 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
34 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
35 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
39 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
40 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
43 #define INITVAL 0x0101010101010101
45 #define SBOG_LPSti64 \
46 s_sbob_sl64[0][(t[0] >> (i * 8)) & 0xff] ^ \
47 s_sbob_sl64[1][(t[1] >> (i * 8)) & 0xff] ^ \
48 s_sbob_sl64[2][(t[2] >> (i * 8)) & 0xff] ^ \
49 s_sbob_sl64[3][(t[3] >> (i * 8)) & 0xff] ^ \
50 s_sbob_sl64[4][(t[4] >> (i * 8)) & 0xff] ^ \
51 s_sbob_sl64[5][(t[5] >> (i * 8)) & 0xff] ^ \
52 s_sbob_sl64[6][(t[6] >> (i * 8)) & 0xff] ^ \
53 s_sbob_sl64[7][(t[7] >> (i * 8)) & 0xff]
57 __device__ __constant__ u64 sbob_sl64[8][256] =
2125 __device__ __constant__ u64 sbob_rc64[12][8] =
2249 __device__ static void streebog_g (u64 h[8], const u64 m[8], u64 s_sbob_sl64[8][256])
2256 for (int i = 0; i < 8; i++)
2261 for (int i = 0; i < 8; i++)
2263 k[i] = SBOG_LPSti64;
2267 for (int i = 0; i < 8; i++)
2272 for (int r = 0; r < 12; r++)
2275 for (int i = 0; i < 8; i++)
2281 for (int i = 0; i < 8; i++)
2283 s[i] = SBOG_LPSti64;
2286 for (int i = 0; i < 8; i++)
2288 t[i] = k[i] ^ sbob_rc64[r][i];
2292 for (int i = 0; i < 8; i++)
2294 k[i] = SBOG_LPSti64;
2299 for (int i = 0; i < 8; i++)
2301 h[i] ^= s[i] ^ k[i] ^ m[i];
2305 __device__ __constant__ bf_t c_bfs[1024];
2307 __device__ static void m11700m (u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
2313 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2314 const u32 lid = threadIdx.x;
2322 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
2324 const u32 w0r = c_bfs[il_pos].i;
2329 * reverse message block
2334 m[0] = hl32_to_64 (w[15], w[14]);
2335 m[1] = hl32_to_64 (w[13], w[12]);
2336 m[2] = hl32_to_64 (w[11], w[10]);
2337 m[3] = hl32_to_64 (w[ 9], w[ 8]);
2338 m[4] = hl32_to_64 (w[ 7], w[ 6]);
2339 m[5] = hl32_to_64 (w[ 5], w[ 4]);
2340 m[6] = hl32_to_64 (w[ 3], w[ 2]);
2341 m[7] = hl32_to_64 (w[ 1], w[ 0]);
2343 m[0] = swap_workaround (m[0]);
2344 m[1] = swap_workaround (m[1]);
2345 m[2] = swap_workaround (m[2]);
2346 m[3] = swap_workaround (m[3]);
2347 m[4] = swap_workaround (m[4]);
2348 m[5] = swap_workaround (m[5]);
2349 m[6] = swap_workaround (m[6]);
2350 m[7] = swap_workaround (m[7]);
2352 // state buffer (hash)
2365 streebog_g (h, m, s_sbob_sl64);
2376 z[7] = swap_workaround ((u64) (pw_len * 8));
2378 streebog_g (h, z, s_sbob_sl64);
2379 streebog_g (h, m, s_sbob_sl64);
2381 const u32 r0 = l32_from_64 (h[0]);
2382 const u32 r1 = h32_from_64 (h[0]);
2383 const u32 r2 = l32_from_64 (h[1]);
2384 const u32 r3 = h32_from_64 (h[1]);
2386 #include VECT_COMPARE_M
2390 __device__ static void m11700s (u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
2396 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2397 const u32 lid = threadIdx.x;
2403 const u32 search[4] =
2405 digests_buf[digests_offset].digest_buf[DGST_R0],
2406 digests_buf[digests_offset].digest_buf[DGST_R1],
2407 digests_buf[digests_offset].digest_buf[DGST_R2],
2408 digests_buf[digests_offset].digest_buf[DGST_R3]
2417 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
2419 const u32 w0r = c_bfs[il_pos].i;
2424 * reverse message block
2429 m[0] = hl32_to_64 (w[15], w[14]);
2430 m[1] = hl32_to_64 (w[13], w[12]);
2431 m[2] = hl32_to_64 (w[11], w[10]);
2432 m[3] = hl32_to_64 (w[ 9], w[ 8]);
2433 m[4] = hl32_to_64 (w[ 7], w[ 6]);
2434 m[5] = hl32_to_64 (w[ 5], w[ 4]);
2435 m[6] = hl32_to_64 (w[ 3], w[ 2]);
2436 m[7] = hl32_to_64 (w[ 1], w[ 0]);
2438 m[0] = swap_workaround (m[0]);
2439 m[1] = swap_workaround (m[1]);
2440 m[2] = swap_workaround (m[2]);
2441 m[3] = swap_workaround (m[3]);
2442 m[4] = swap_workaround (m[4]);
2443 m[5] = swap_workaround (m[5]);
2444 m[6] = swap_workaround (m[6]);
2445 m[7] = swap_workaround (m[7]);
2447 // state buffer (hash)
2460 streebog_g (h, m, s_sbob_sl64);
2471 z[7] = swap_workaround ((u64) (pw_len * 8));
2473 streebog_g (h, z, s_sbob_sl64);
2474 streebog_g (h, m, s_sbob_sl64);
2476 const u32 r0 = l32_from_64 (h[0]);
2477 const u32 r1 = h32_from_64 (h[0]);
2478 const u32 r2 = l32_from_64 (h[1]);
2479 const u32 r3 = h32_from_64 (h[1]);
2481 #include VECT_COMPARE_S
2485 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
2491 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2492 const u32 lid = threadIdx.x;
2496 w[ 0] = pws[gid].i[ 0];
2497 w[ 1] = pws[gid].i[ 1];
2498 w[ 2] = pws[gid].i[ 2];
2499 w[ 3] = pws[gid].i[ 3];
2513 const u32 pw_len = pws[gid].pw_len;
2516 * shared lookup table
2519 __shared__ u64 s_sbob_sl64[8][256];
2523 const u32 lid4 = lid * 4;
2525 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2526 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2527 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2528 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2529 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2530 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2531 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2532 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2533 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2534 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2535 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2536 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2537 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2538 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2539 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2540 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2541 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2542 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2543 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2544 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2545 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2546 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2547 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2548 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2549 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2550 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2551 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2552 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2553 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2554 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2555 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2556 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2561 if (gid >= gid_max) return;
2567 m11700m (s_sbob_sl64, w, 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);
2570 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
2576 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2577 const u32 lid = threadIdx.x;
2581 w[ 0] = pws[gid].i[ 0];
2582 w[ 1] = pws[gid].i[ 1];
2583 w[ 2] = pws[gid].i[ 2];
2584 w[ 3] = pws[gid].i[ 3];
2585 w[ 4] = pws[gid].i[ 4];
2586 w[ 5] = pws[gid].i[ 5];
2587 w[ 6] = pws[gid].i[ 6];
2588 w[ 7] = pws[gid].i[ 7];
2598 const u32 pw_len = pws[gid].pw_len;
2601 * shared lookup table
2604 __shared__ u64 s_sbob_sl64[8][256];
2608 const u32 lid4 = lid * 4;
2610 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2611 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2612 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2613 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2614 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2615 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2616 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2617 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2618 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2619 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2620 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2621 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2622 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2623 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2624 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2625 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2626 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2627 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2628 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2629 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2630 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2631 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2632 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2633 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2634 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2635 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2636 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2637 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2638 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2639 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2640 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2641 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2646 if (gid >= gid_max) return;
2652 m11700m (s_sbob_sl64, w, 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);
2655 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
2661 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2662 const u32 lid = threadIdx.x;
2666 w[ 0] = pws[gid].i[ 0];
2667 w[ 1] = pws[gid].i[ 1];
2668 w[ 2] = pws[gid].i[ 2];
2669 w[ 3] = pws[gid].i[ 3];
2670 w[ 4] = pws[gid].i[ 4];
2671 w[ 5] = pws[gid].i[ 5];
2672 w[ 6] = pws[gid].i[ 6];
2673 w[ 7] = pws[gid].i[ 7];
2674 w[ 8] = pws[gid].i[ 8];
2675 w[ 9] = pws[gid].i[ 9];
2676 w[10] = pws[gid].i[10];
2677 w[11] = pws[gid].i[11];
2678 w[12] = pws[gid].i[12];
2679 w[13] = pws[gid].i[13];
2680 w[14] = pws[gid].i[14];
2681 w[15] = pws[gid].i[15];
2683 const u32 pw_len = pws[gid].pw_len;
2686 * shared lookup table
2689 __shared__ u64 s_sbob_sl64[8][256];
2693 const u32 lid4 = lid * 4;
2695 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2696 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2697 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2698 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2699 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2700 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2701 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2702 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2703 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2704 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2705 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2706 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2707 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2708 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2709 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2710 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2711 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2712 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2713 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2714 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2715 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2716 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2717 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2718 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2719 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2720 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2721 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2722 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2723 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2724 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2725 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2726 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2731 if (gid >= gid_max) return;
2737 m11700m (s_sbob_sl64, w, 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);
2740 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
2746 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2747 const u32 lid = threadIdx.x;
2751 w[ 0] = pws[gid].i[ 0];
2752 w[ 1] = pws[gid].i[ 1];
2753 w[ 2] = pws[gid].i[ 2];
2754 w[ 3] = pws[gid].i[ 3];
2768 const u32 pw_len = pws[gid].pw_len;
2771 * shared lookup table
2774 __shared__ u64 s_sbob_sl64[8][256];
2778 const u32 lid4 = lid * 4;
2780 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2781 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2782 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2783 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2784 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2785 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2786 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2787 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2788 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2789 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2790 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2791 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2792 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2793 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2794 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2795 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2796 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2797 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2798 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2799 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2800 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2801 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2802 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2803 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2804 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2805 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2806 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2807 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2808 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2809 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2810 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2811 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2816 if (gid >= gid_max) return;
2822 m11700s (s_sbob_sl64, w, 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);
2825 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
2831 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2832 const u32 lid = threadIdx.x;
2836 w[ 0] = pws[gid].i[ 0];
2837 w[ 1] = pws[gid].i[ 1];
2838 w[ 2] = pws[gid].i[ 2];
2839 w[ 3] = pws[gid].i[ 3];
2840 w[ 4] = pws[gid].i[ 4];
2841 w[ 5] = pws[gid].i[ 5];
2842 w[ 6] = pws[gid].i[ 6];
2843 w[ 7] = pws[gid].i[ 7];
2853 const u32 pw_len = pws[gid].pw_len;
2856 * shared lookup table
2859 __shared__ u64 s_sbob_sl64[8][256];
2863 const u32 lid4 = lid * 4;
2865 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2866 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2867 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2868 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2869 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2870 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2871 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2872 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2873 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2874 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2875 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2876 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2877 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2878 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2879 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2880 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2881 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2882 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2883 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2884 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2885 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2886 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2887 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2888 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2889 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2890 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2891 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2892 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2893 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2894 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2895 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2896 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2901 if (gid >= gid_max) return;
2907 m11700s (s_sbob_sl64, w, 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);
2910 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
2916 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2917 const u32 lid = threadIdx.x;
2921 w[ 0] = pws[gid].i[ 0];
2922 w[ 1] = pws[gid].i[ 1];
2923 w[ 2] = pws[gid].i[ 2];
2924 w[ 3] = pws[gid].i[ 3];
2925 w[ 4] = pws[gid].i[ 4];
2926 w[ 5] = pws[gid].i[ 5];
2927 w[ 6] = pws[gid].i[ 6];
2928 w[ 7] = pws[gid].i[ 7];
2929 w[ 8] = pws[gid].i[ 8];
2930 w[ 9] = pws[gid].i[ 9];
2931 w[10] = pws[gid].i[10];
2932 w[11] = pws[gid].i[11];
2933 w[12] = pws[gid].i[12];
2934 w[13] = pws[gid].i[13];
2935 w[14] = pws[gid].i[14];
2936 w[15] = pws[gid].i[15];
2938 const u32 pw_len = pws[gid].pw_len;
2941 * shared lookup table
2944 __shared__ u64 s_sbob_sl64[8][256];
2948 const u32 lid4 = lid * 4;
2950 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2951 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2952 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2953 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2954 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2955 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2956 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2957 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2958 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2959 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2960 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2961 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2962 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2963 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2964 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2965 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2966 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2967 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2968 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2969 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2970 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2971 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2972 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2973 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2974 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2975 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2976 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2977 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2978 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2979 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2980 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2981 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2986 if (gid >= gid_max) return;
2992 m11700s (s_sbob_sl64, w, 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);