2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
16 #include "include/kernel_functions.c"
17 #include "OpenCL/types_ocl.c"
18 #include "OpenCL/common.c"
20 #define COMPARE_S "OpenCL/check_single_comp4.c"
21 #define COMPARE_M "OpenCL/check_multi_comp4.c"
25 #define SBOG_LPSti64 \
26 s_sbob_sl64[0][(t[0] >> (i * 8)) & 0xff] ^ \
27 s_sbob_sl64[1][(t[1] >> (i * 8)) & 0xff] ^ \
28 s_sbob_sl64[2][(t[2] >> (i * 8)) & 0xff] ^ \
29 s_sbob_sl64[3][(t[3] >> (i * 8)) & 0xff] ^ \
30 s_sbob_sl64[4][(t[4] >> (i * 8)) & 0xff] ^ \
31 s_sbob_sl64[5][(t[5] >> (i * 8)) & 0xff] ^ \
32 s_sbob_sl64[6][(t[6] >> (i * 8)) & 0xff] ^ \
33 s_sbob_sl64[7][(t[7] >> (i * 8)) & 0xff]
37 __constant u64 sbob_sl64[8][256] =
2105 __constant u64 sbob_rc64[12][8] =
2229 static void streebog_g (u64 h[8], const u64 m[8], __local u64 s_sbob_sl64[8][256])
2236 for (int i = 0; i < 8; i++)
2241 for (int i = 0; i < 8; i++)
2243 k[i] = SBOG_LPSti64;
2247 for (int i = 0; i < 8; i++)
2252 for (int r = 0; r < 12; r++)
2255 for (int i = 0; i < 8; i++)
2261 for (int i = 0; i < 8; i++)
2263 s[i] = SBOG_LPSti64;
2266 for (int i = 0; i < 8; i++)
2268 t[i] = k[i] ^ sbob_rc64[r][i];
2272 for (int i = 0; i < 8; i++)
2274 k[i] = SBOG_LPSti64;
2279 for (int i = 0; i < 8; i++)
2281 h[i] ^= s[i] ^ k[i] ^ m[i];
2285 static void m11800m (__local u64 s_sbob_sl64[8][256], u32 w[16], 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_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)
2291 const u32 gid = get_global_id (0);
2292 const u32 lid = get_local_id (0);
2300 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
2302 const u32 w0r = bfs_buf[il_pos].i;
2307 * reverse message block
2312 m[0] = hl32_to_64 (w[15], w[14]);
2313 m[1] = hl32_to_64 (w[13], w[12]);
2314 m[2] = hl32_to_64 (w[11], w[10]);
2315 m[3] = hl32_to_64 (w[ 9], w[ 8]);
2316 m[4] = hl32_to_64 (w[ 7], w[ 6]);
2317 m[5] = hl32_to_64 (w[ 5], w[ 4]);
2318 m[6] = hl32_to_64 (w[ 3], w[ 2]);
2319 m[7] = hl32_to_64 (w[ 1], w[ 0]);
2321 m[0] = swap64 (m[0]);
2322 m[1] = swap64 (m[1]);
2323 m[2] = swap64 (m[2]);
2324 m[3] = swap64 (m[3]);
2325 m[4] = swap64 (m[4]);
2326 m[5] = swap64 (m[5]);
2327 m[6] = swap64 (m[6]);
2328 m[7] = swap64 (m[7]);
2330 // state buffer (hash)
2343 streebog_g (h, m, s_sbob_sl64);
2354 z[7] = swap64 ((u64) (pw_len * 8));
2356 streebog_g (h, z, s_sbob_sl64);
2357 streebog_g (h, m, s_sbob_sl64);
2359 const u32 r0 = l32_from_64 (h[0]);
2360 const u32 r1 = h32_from_64 (h[0]);
2361 const u32 r2 = l32_from_64 (h[1]);
2362 const u32 r3 = h32_from_64 (h[1]);
2368 static void m11800s (__local u64 s_sbob_sl64[8][256], u32 w[16], 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_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)
2374 const u32 gid = get_global_id (0);
2375 const u32 lid = get_local_id (0);
2381 const u32 search[4] =
2383 digests_buf[digests_offset].digest_buf[DGST_R0],
2384 digests_buf[digests_offset].digest_buf[DGST_R1],
2385 digests_buf[digests_offset].digest_buf[DGST_R2],
2386 digests_buf[digests_offset].digest_buf[DGST_R3]
2395 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
2397 const u32 w0r = bfs_buf[il_pos].i;
2402 * reverse message block
2407 m[0] = hl32_to_64 (w[15], w[14]);
2408 m[1] = hl32_to_64 (w[13], w[12]);
2409 m[2] = hl32_to_64 (w[11], w[10]);
2410 m[3] = hl32_to_64 (w[ 9], w[ 8]);
2411 m[4] = hl32_to_64 (w[ 7], w[ 6]);
2412 m[5] = hl32_to_64 (w[ 5], w[ 4]);
2413 m[6] = hl32_to_64 (w[ 3], w[ 2]);
2414 m[7] = hl32_to_64 (w[ 1], w[ 0]);
2416 m[0] = swap64 (m[0]);
2417 m[1] = swap64 (m[1]);
2418 m[2] = swap64 (m[2]);
2419 m[3] = swap64 (m[3]);
2420 m[4] = swap64 (m[4]);
2421 m[5] = swap64 (m[5]);
2422 m[6] = swap64 (m[6]);
2423 m[7] = swap64 (m[7]);
2425 // state buffer (hash)
2438 streebog_g (h, m, s_sbob_sl64);
2449 z[7] = swap64 ((u64) (pw_len * 8));
2451 streebog_g (h, z, s_sbob_sl64);
2452 streebog_g (h, m, s_sbob_sl64);
2454 const u32 r0 = l32_from_64 (h[0]);
2455 const u32 r1 = h32_from_64 (h[0]);
2456 const u32 r2 = l32_from_64 (h[1]);
2457 const u32 r3 = h32_from_64 (h[1]);
2463 __kernel void m11800_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_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)
2469 const u32 gid = get_global_id (0);
2470 const u32 lid = get_local_id (0);
2471 const u32 lsz = get_local_size (0);
2474 * shared lookup table
2477 __local u64 s_sbob_sl64[8][256];
2479 for (u32 i = lid; i < 256; i += lsz)
2481 s_sbob_sl64[0][i] = sbob_sl64[0][i];
2482 s_sbob_sl64[1][i] = sbob_sl64[1][i];
2483 s_sbob_sl64[2][i] = sbob_sl64[2][i];
2484 s_sbob_sl64[3][i] = sbob_sl64[3][i];
2485 s_sbob_sl64[4][i] = sbob_sl64[4][i];
2486 s_sbob_sl64[5][i] = sbob_sl64[5][i];
2487 s_sbob_sl64[6][i] = sbob_sl64[6][i];
2488 s_sbob_sl64[7][i] = sbob_sl64[7][i];
2491 barrier (CLK_LOCAL_MEM_FENCE);
2493 if (gid >= gid_max) return;
2501 w[ 0] = pws[gid].i[ 0];
2502 w[ 1] = pws[gid].i[ 1];
2503 w[ 2] = pws[gid].i[ 2];
2504 w[ 3] = pws[gid].i[ 3];
2518 const u32 pw_len = pws[gid].pw_len;
2524 m11800m (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);
2527 __kernel void m11800_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_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)
2533 const u32 gid = get_global_id (0);
2534 const u32 lid = get_local_id (0);
2535 const u32 lsz = get_local_size (0);
2538 * shared lookup table
2541 __local u64 s_sbob_sl64[8][256];
2543 for (u32 i = lid; i < 256; i += lsz)
2545 s_sbob_sl64[0][i] = sbob_sl64[0][i];
2546 s_sbob_sl64[1][i] = sbob_sl64[1][i];
2547 s_sbob_sl64[2][i] = sbob_sl64[2][i];
2548 s_sbob_sl64[3][i] = sbob_sl64[3][i];
2549 s_sbob_sl64[4][i] = sbob_sl64[4][i];
2550 s_sbob_sl64[5][i] = sbob_sl64[5][i];
2551 s_sbob_sl64[6][i] = sbob_sl64[6][i];
2552 s_sbob_sl64[7][i] = sbob_sl64[7][i];
2555 barrier (CLK_LOCAL_MEM_FENCE);
2557 if (gid >= gid_max) return;
2565 w[ 0] = pws[gid].i[ 0];
2566 w[ 1] = pws[gid].i[ 1];
2567 w[ 2] = pws[gid].i[ 2];
2568 w[ 3] = pws[gid].i[ 3];
2569 w[ 4] = pws[gid].i[ 4];
2570 w[ 5] = pws[gid].i[ 5];
2571 w[ 6] = pws[gid].i[ 6];
2572 w[ 7] = pws[gid].i[ 7];
2582 const u32 pw_len = pws[gid].pw_len;
2588 m11800m (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);
2591 __kernel void m11800_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_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)
2597 const u32 gid = get_global_id (0);
2598 const u32 lid = get_local_id (0);
2599 const u32 lsz = get_local_size (0);
2602 * shared lookup table
2605 __local u64 s_sbob_sl64[8][256];
2607 for (u32 i = lid; i < 256; i += lsz)
2609 s_sbob_sl64[0][i] = sbob_sl64[0][i];
2610 s_sbob_sl64[1][i] = sbob_sl64[1][i];
2611 s_sbob_sl64[2][i] = sbob_sl64[2][i];
2612 s_sbob_sl64[3][i] = sbob_sl64[3][i];
2613 s_sbob_sl64[4][i] = sbob_sl64[4][i];
2614 s_sbob_sl64[5][i] = sbob_sl64[5][i];
2615 s_sbob_sl64[6][i] = sbob_sl64[6][i];
2616 s_sbob_sl64[7][i] = sbob_sl64[7][i];
2619 barrier (CLK_LOCAL_MEM_FENCE);
2621 if (gid >= gid_max) return;
2629 w[ 0] = pws[gid].i[ 0];
2630 w[ 1] = pws[gid].i[ 1];
2631 w[ 2] = pws[gid].i[ 2];
2632 w[ 3] = pws[gid].i[ 3];
2633 w[ 4] = pws[gid].i[ 4];
2634 w[ 5] = pws[gid].i[ 5];
2635 w[ 6] = pws[gid].i[ 6];
2636 w[ 7] = pws[gid].i[ 7];
2637 w[ 8] = pws[gid].i[ 8];
2638 w[ 9] = pws[gid].i[ 9];
2639 w[10] = pws[gid].i[10];
2640 w[11] = pws[gid].i[11];
2641 w[12] = pws[gid].i[12];
2642 w[13] = pws[gid].i[13];
2643 w[14] = pws[gid].i[14];
2644 w[15] = pws[gid].i[15];
2646 const u32 pw_len = pws[gid].pw_len;
2652 m11800m (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 __kernel void m11800_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_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 = get_global_id (0);
2662 const u32 lid = get_local_id (0);
2663 const u32 lsz = get_local_size (0);
2666 * shared lookup table
2669 __local u64 s_sbob_sl64[8][256];
2671 for (u32 i = lid; i < 256; i += lsz)
2673 s_sbob_sl64[0][i] = sbob_sl64[0][i];
2674 s_sbob_sl64[1][i] = sbob_sl64[1][i];
2675 s_sbob_sl64[2][i] = sbob_sl64[2][i];
2676 s_sbob_sl64[3][i] = sbob_sl64[3][i];
2677 s_sbob_sl64[4][i] = sbob_sl64[4][i];
2678 s_sbob_sl64[5][i] = sbob_sl64[5][i];
2679 s_sbob_sl64[6][i] = sbob_sl64[6][i];
2680 s_sbob_sl64[7][i] = sbob_sl64[7][i];
2683 barrier (CLK_LOCAL_MEM_FENCE);
2685 if (gid >= gid_max) return;
2693 w[ 0] = pws[gid].i[ 0];
2694 w[ 1] = pws[gid].i[ 1];
2695 w[ 2] = pws[gid].i[ 2];
2696 w[ 3] = pws[gid].i[ 3];
2710 const u32 pw_len = pws[gid].pw_len;
2716 m11800s (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);
2719 __kernel void m11800_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_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)
2725 const u32 gid = get_global_id (0);
2726 const u32 lid = get_local_id (0);
2727 const u32 lsz = get_local_size (0);
2730 * shared lookup table
2733 __local u64 s_sbob_sl64[8][256];
2735 for (u32 i = lid; i < 256; i += lsz)
2737 s_sbob_sl64[0][i] = sbob_sl64[0][i];
2738 s_sbob_sl64[1][i] = sbob_sl64[1][i];
2739 s_sbob_sl64[2][i] = sbob_sl64[2][i];
2740 s_sbob_sl64[3][i] = sbob_sl64[3][i];
2741 s_sbob_sl64[4][i] = sbob_sl64[4][i];
2742 s_sbob_sl64[5][i] = sbob_sl64[5][i];
2743 s_sbob_sl64[6][i] = sbob_sl64[6][i];
2744 s_sbob_sl64[7][i] = sbob_sl64[7][i];
2747 barrier (CLK_LOCAL_MEM_FENCE);
2749 if (gid >= gid_max) return;
2757 w[ 0] = pws[gid].i[ 0];
2758 w[ 1] = pws[gid].i[ 1];
2759 w[ 2] = pws[gid].i[ 2];
2760 w[ 3] = pws[gid].i[ 3];
2761 w[ 4] = pws[gid].i[ 4];
2762 w[ 5] = pws[gid].i[ 5];
2763 w[ 6] = pws[gid].i[ 6];
2764 w[ 7] = pws[gid].i[ 7];
2774 const u32 pw_len = pws[gid].pw_len;
2780 m11800s (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);
2783 __kernel void m11800_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_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)
2789 const u32 gid = get_global_id (0);
2790 const u32 lid = get_local_id (0);
2791 const u32 lsz = get_local_size (0);
2794 * shared lookup table
2797 __local u64 s_sbob_sl64[8][256];
2799 for (u32 i = lid; i < 256; i += lsz)
2801 s_sbob_sl64[0][i] = sbob_sl64[0][i];
2802 s_sbob_sl64[1][i] = sbob_sl64[1][i];
2803 s_sbob_sl64[2][i] = sbob_sl64[2][i];
2804 s_sbob_sl64[3][i] = sbob_sl64[3][i];
2805 s_sbob_sl64[4][i] = sbob_sl64[4][i];
2806 s_sbob_sl64[5][i] = sbob_sl64[5][i];
2807 s_sbob_sl64[6][i] = sbob_sl64[6][i];
2808 s_sbob_sl64[7][i] = sbob_sl64[7][i];
2811 barrier (CLK_LOCAL_MEM_FENCE);
2813 if (gid >= gid_max) return;
2821 w[ 0] = pws[gid].i[ 0];
2822 w[ 1] = pws[gid].i[ 1];
2823 w[ 2] = pws[gid].i[ 2];
2824 w[ 3] = pws[gid].i[ 3];
2825 w[ 4] = pws[gid].i[ 4];
2826 w[ 5] = pws[gid].i[ 5];
2827 w[ 6] = pws[gid].i[ 6];
2828 w[ 7] = pws[gid].i[ 7];
2829 w[ 8] = pws[gid].i[ 8];
2830 w[ 9] = pws[gid].i[ 9];
2831 w[10] = pws[gid].i[10];
2832 w[11] = pws[gid].i[11];
2833 w[12] = pws[gid].i[12];
2834 w[13] = pws[gid].i[13];
2835 w[14] = pws[gid].i[14];
2836 w[15] = pws[gid].i[15];
2838 const u32 pw_len = pws[gid].pw_len;
2844 m11800s (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);