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"
23 #define INITVAL 0x0101010101010101
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 m11700m (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
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 m11700s (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
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 __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2469 const u32 gid = get_global_id (0);
2470 const u32 lid = get_local_id (0);
2474 w[ 0] = pws[gid].i[ 0];
2475 w[ 1] = pws[gid].i[ 1];
2476 w[ 2] = pws[gid].i[ 2];
2477 w[ 3] = pws[gid].i[ 3];
2491 const u32 pw_len = pws[gid].pw_len;
2494 * shared lookup table
2497 const u32 lid4 = lid * 4;
2499 __local u64 s_sbob_sl64[8][256];
2501 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2502 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2503 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2504 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2505 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2506 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2507 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2508 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2509 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2510 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2511 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2512 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2513 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2514 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2515 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2516 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2517 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2518 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2519 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2520 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2521 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2522 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2523 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2524 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2525 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2526 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2527 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2528 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2529 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2530 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2531 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2532 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2534 barrier (CLK_LOCAL_MEM_FENCE);
2536 if (gid >= gid_max) return;
2542 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);
2545 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2551 const u32 gid = get_global_id (0);
2552 const u32 lid = get_local_id (0);
2556 w[ 0] = pws[gid].i[ 0];
2557 w[ 1] = pws[gid].i[ 1];
2558 w[ 2] = pws[gid].i[ 2];
2559 w[ 3] = pws[gid].i[ 3];
2560 w[ 4] = pws[gid].i[ 4];
2561 w[ 5] = pws[gid].i[ 5];
2562 w[ 6] = pws[gid].i[ 6];
2563 w[ 7] = pws[gid].i[ 7];
2573 const u32 pw_len = pws[gid].pw_len;
2576 * shared lookup table
2579 const u32 lid4 = lid * 4;
2581 __local u64 s_sbob_sl64[8][256];
2583 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2584 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2585 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2586 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2587 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2588 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2589 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2590 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2591 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2592 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2593 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2594 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2595 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2596 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2597 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2598 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2599 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2600 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2601 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2602 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2603 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2604 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2605 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2606 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2607 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2608 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2609 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2610 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2611 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2612 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2613 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2614 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2616 barrier (CLK_LOCAL_MEM_FENCE);
2618 if (gid >= gid_max) return;
2624 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);
2627 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2633 const u32 gid = get_global_id (0);
2634 const u32 lid = get_local_id (0);
2638 w[ 0] = pws[gid].i[ 0];
2639 w[ 1] = pws[gid].i[ 1];
2640 w[ 2] = pws[gid].i[ 2];
2641 w[ 3] = pws[gid].i[ 3];
2642 w[ 4] = pws[gid].i[ 4];
2643 w[ 5] = pws[gid].i[ 5];
2644 w[ 6] = pws[gid].i[ 6];
2645 w[ 7] = pws[gid].i[ 7];
2646 w[ 8] = pws[gid].i[ 8];
2647 w[ 9] = pws[gid].i[ 9];
2648 w[10] = pws[gid].i[10];
2649 w[11] = pws[gid].i[11];
2650 w[12] = pws[gid].i[12];
2651 w[13] = pws[gid].i[13];
2652 w[14] = pws[gid].i[14];
2653 w[15] = pws[gid].i[15];
2655 const u32 pw_len = pws[gid].pw_len;
2658 * shared lookup table
2661 const u32 lid4 = lid * 4;
2663 __local u64 s_sbob_sl64[8][256];
2665 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2666 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2667 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2668 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2669 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2670 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2671 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2672 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2673 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2674 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2675 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2676 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2677 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2678 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2679 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2680 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2681 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2682 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2683 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2684 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2685 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2686 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2687 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2688 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2689 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2690 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2691 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2692 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2693 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2694 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2695 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2696 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2698 barrier (CLK_LOCAL_MEM_FENCE);
2700 if (gid >= gid_max) return;
2706 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);
2709 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2715 const u32 gid = get_global_id (0);
2716 const u32 lid = get_local_id (0);
2720 w[ 0] = pws[gid].i[ 0];
2721 w[ 1] = pws[gid].i[ 1];
2722 w[ 2] = pws[gid].i[ 2];
2723 w[ 3] = pws[gid].i[ 3];
2737 const u32 pw_len = pws[gid].pw_len;
2740 * shared lookup table
2743 const u32 lid4 = lid * 4;
2745 __local u64 s_sbob_sl64[8][256];
2747 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2748 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2749 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2750 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2751 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2752 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2753 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2754 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2755 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2756 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2757 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2758 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2759 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2760 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2761 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2762 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2763 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2764 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2765 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2766 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2767 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2768 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2769 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2770 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2771 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2772 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2773 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2774 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2775 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2776 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2777 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2778 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2780 barrier (CLK_LOCAL_MEM_FENCE);
2782 if (gid >= gid_max) return;
2788 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);
2791 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2797 const u32 gid = get_global_id (0);
2798 const u32 lid = get_local_id (0);
2802 w[ 0] = pws[gid].i[ 0];
2803 w[ 1] = pws[gid].i[ 1];
2804 w[ 2] = pws[gid].i[ 2];
2805 w[ 3] = pws[gid].i[ 3];
2806 w[ 4] = pws[gid].i[ 4];
2807 w[ 5] = pws[gid].i[ 5];
2808 w[ 6] = pws[gid].i[ 6];
2809 w[ 7] = pws[gid].i[ 7];
2819 const u32 pw_len = pws[gid].pw_len;
2822 * shared lookup table
2825 const u32 lid4 = lid * 4;
2827 __local u64 s_sbob_sl64[8][256];
2829 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2830 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2831 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2832 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2833 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2834 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2835 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2836 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2837 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2838 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2839 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2840 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2841 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2842 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2843 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2844 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2845 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2846 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2847 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2848 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2849 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2850 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2851 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2852 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2853 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2854 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2855 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2856 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2857 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2858 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2859 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2860 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2862 barrier (CLK_LOCAL_MEM_FENCE);
2864 if (gid >= gid_max) return;
2870 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);
2873 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2879 const u32 gid = get_global_id (0);
2880 const u32 lid = get_local_id (0);
2884 w[ 0] = pws[gid].i[ 0];
2885 w[ 1] = pws[gid].i[ 1];
2886 w[ 2] = pws[gid].i[ 2];
2887 w[ 3] = pws[gid].i[ 3];
2888 w[ 4] = pws[gid].i[ 4];
2889 w[ 5] = pws[gid].i[ 5];
2890 w[ 6] = pws[gid].i[ 6];
2891 w[ 7] = pws[gid].i[ 7];
2892 w[ 8] = pws[gid].i[ 8];
2893 w[ 9] = pws[gid].i[ 9];
2894 w[10] = pws[gid].i[10];
2895 w[11] = pws[gid].i[11];
2896 w[12] = pws[gid].i[12];
2897 w[13] = pws[gid].i[13];
2898 w[14] = pws[gid].i[14];
2899 w[15] = pws[gid].i[15];
2901 const u32 pw_len = pws[gid].pw_len;
2904 * shared lookup table
2907 const u32 lid4 = lid * 4;
2909 __local u64 s_sbob_sl64[8][256];
2911 s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2912 s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2913 s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2914 s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2915 s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2916 s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2917 s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2918 s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2919 s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2920 s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2921 s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2922 s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2923 s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2924 s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2925 s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2926 s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2927 s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2928 s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2929 s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2930 s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2931 s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2932 s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2933 s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2934 s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2935 s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2936 s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2937 s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2938 s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2939 s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2940 s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2941 s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2942 s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2944 barrier (CLK_LOCAL_MEM_FENCE);
2946 if (gid >= gid_max) return;
2952 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);