/** * Author......: Jens Steube * License.....: MIT */ #define _KECCAK_ #include "include/constants.h" #include "include/kernel_vendor.h" #ifdef VLIW1 #define VECT_SIZE1 #endif #ifdef VLIW4 #define VECT_SIZE1 #endif #ifdef VLIW5 #define VECT_SIZE1 #endif #define DGST_R0 2 #define DGST_R1 3 #define DGST_R2 4 #define DGST_R3 5 #include "include/kernel_functions.c" #include "types_amd.c" #include "common_amd.c" #ifdef VECT_SIZE1 #define VECT_COMPARE_S "check_single_vect1_comp4.c" #define VECT_COMPARE_M "check_multi_vect1_comp4.c" #endif #ifdef VECT_SIZE2 #define VECT_COMPARE_S "check_single_vect2_comp4.c" #define VECT_COMPARE_M "check_multi_vect2_comp4.c" #endif __constant u64 keccakf_rndc[24] = { 0x0000000000000001, 0x0000000000008082, 0x800000000000808a, 0x8000000080008000, 0x000000000000808b, 0x0000000080000001, 0x8000000080008081, 0x8000000000008009, 0x000000000000008a, 0x0000000000000088, 0x0000000080008009, 0x000000008000000a, 0x000000008000808b, 0x800000000000008b, 0x8000000000008089, 0x8000000000008003, 0x8000000000008002, 0x8000000000000080, 0x000000000000800a, 0x800000008000000a, 0x8000000080008081, 0x8000000000008080, 0x0000000080000001, 0x8000000080008008 }; __constant u32 keccakf_rotc[24] = { 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44 }; __constant u32 keccakf_piln[24] = { 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 }; #ifndef KECCAK_ROUNDS #define KECCAK_ROUNDS 24 #endif #define Theta1(s) (st[0 + s] ^ st[5 + s] ^ st[10 + s] ^ st[15 + s] ^ st[20 + s]) #define Theta2(s) \ { \ st[ 0 + s] ^= t; \ st[ 5 + s] ^= t; \ st[10 + s] ^= t; \ st[15 + s] ^= t; \ st[20 + s] ^= t; \ } #define Rho_Pi(s) \ { \ u32 j = keccakf_piln[s]; \ u32 k = keccakf_rotc[s]; \ bc0 = st[j]; \ st[j] = rotl64 (t, k); \ t = bc0; \ } #define Chi(s) \ { \ bc0 = st[0 + s]; \ bc1 = st[1 + s]; \ bc2 = st[2 + s]; \ bc3 = st[3 + s]; \ bc4 = st[4 + s]; \ st[0 + s] ^= ~bc1 & bc2; \ st[1 + s] ^= ~bc2 & bc3; \ st[2 + s] ^= ~bc3 & bc4; \ st[3 + s] ^= ~bc4 & bc0; \ st[4 + s] ^= ~bc0 & bc1; \ } static void m05000m (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 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) { /** * modifier */ const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); /** * 0x80 keccak, very special */ const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen; const u32 rsiz = 200 - (2 * mdlen); const u32 add80w = (rsiz - 1) / 8; /** * loop */ u32x w0l = w0[0]; for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++) { const u32 w0r = bfs_buf[il_pos].i; w0[0] = w0l | w0r; u64x st[25]; #ifdef VECT_SIZE1 st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32; st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32; st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32; st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32; st[ 4] = (u64x) (w2[0]) | (u64x) (w2[1]) << 32; st[ 5] = (u64x) (w2[2]) | (u64x) (w2[3]) << 32; st[ 6] = (u64x) (w3[0]) | (u64x) (w3[1]) << 32; st[ 7] = (u64x) (w3[2]) | (u64x) (w3[3]) << 32; #endif #ifdef VECT_SIZE2 st[ 0] = (u64x) (w0[0].s0, w0[0].s1) | (u64x) (w0[1].s0, w0[1].s1) << 32; st[ 1] = (u64x) (w0[2].s0, w0[2].s1) | (u64x) (w0[3].s0, w0[3].s1) << 32; st[ 2] = (u64x) (w1[0].s0, w1[0].s1) | (u64x) (w1[1].s0, w1[1].s1) << 32; st[ 3] = (u64x) (w1[2].s0, w1[2].s1) | (u64x) (w1[3].s0, w1[3].s1) << 32; st[ 4] = (u64x) (w2[0].s0, w2[0].s1) | (u64x) (w2[1].s0, w2[1].s1) << 32; st[ 5] = (u64x) (w2[2].s0, w2[2].s1) | (u64x) (w2[3].s0, w2[3].s1) << 32; st[ 6] = (u64x) (w3[0].s0, w3[0].s1) | (u64x) (w3[1].s0, w3[1].s1) << 32; st[ 7] = (u64x) (w3[2].s0, w3[2].s1) | (u64x) (w3[3].s0, w3[3].s1) << 32; #endif st[ 8] = 0; st[ 9] = 0; st[10] = 0; st[11] = 0; st[12] = 0; st[13] = 0; st[14] = 0; st[15] = 0; st[16] = 0; st[17] = 0; st[18] = 0; st[19] = 0; st[20] = 0; st[21] = 0; st[22] = 0; st[23] = 0; st[24] = 0; st[add80w] |= 0x8000000000000000; int round; for (round = 0; round < KECCAK_ROUNDS; round++) { // Theta u64x bc0 = Theta1 (0); u64x bc1 = Theta1 (1); u64x bc2 = Theta1 (2); u64x bc3 = Theta1 (3); u64x bc4 = Theta1 (4); u64x t; t = bc4 ^ rotl64 (bc1, 1); Theta2 (0); t = bc0 ^ rotl64 (bc2, 1); Theta2 (1); t = bc1 ^ rotl64 (bc3, 1); Theta2 (2); t = bc2 ^ rotl64 (bc4, 1); Theta2 (3); t = bc3 ^ rotl64 (bc0, 1); Theta2 (4); // Rho Pi t = st[1]; Rho_Pi (0); Rho_Pi (1); Rho_Pi (2); Rho_Pi (3); Rho_Pi (4); Rho_Pi (5); Rho_Pi (6); Rho_Pi (7); Rho_Pi (8); Rho_Pi (9); Rho_Pi (10); Rho_Pi (11); Rho_Pi (12); Rho_Pi (13); Rho_Pi (14); Rho_Pi (15); Rho_Pi (16); Rho_Pi (17); Rho_Pi (18); Rho_Pi (19); Rho_Pi (20); Rho_Pi (21); Rho_Pi (22); Rho_Pi (23); // Chi Chi (0); Chi (5); Chi (10); Chi (15); Chi (20); // Iota st[0] ^= keccakf_rndc[round]; } const u32x r0 = l32_from_64 (st[1]); const u32x r1 = h32_from_64 (st[1]); const u32x r2 = l32_from_64 (st[2]); const u32x r3 = h32_from_64 (st[2]); #include VECT_COMPARE_M } } static void m05000s (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 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) { /** * modifier */ const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); /** * digest */ const u32 search[4] = { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R2], digests_buf[digests_offset].digest_buf[DGST_R3] }; /** * 0x80 keccak, very special */ const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen; const u32 rsiz = 200 - (2 * mdlen); const u32 add80w = (rsiz - 1) / 8; /** * loop */ u32x w0l = w0[0]; for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++) { const u32 w0r = bfs_buf[il_pos].i; w0[0] = w0l | w0r; u64x st[25]; #ifdef VECT_SIZE1 st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32; st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32; st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32; st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32; st[ 4] = (u64x) (w2[0]) | (u64x) (w2[1]) << 32; st[ 5] = (u64x) (w2[2]) | (u64x) (w2[3]) << 32; st[ 6] = (u64x) (w3[0]) | (u64x) (w3[1]) << 32; st[ 7] = (u64x) (w3[2]) | (u64x) (w3[3]) << 32; #endif #ifdef VECT_SIZE2 st[ 0] = (u64x) (w0[0].s0, w0[0].s1) | (u64x) (w0[1].s0, w0[1].s1) << 32; st[ 1] = (u64x) (w0[2].s0, w0[2].s1) | (u64x) (w0[3].s0, w0[3].s1) << 32; st[ 2] = (u64x) (w1[0].s0, w1[0].s1) | (u64x) (w1[1].s0, w1[1].s1) << 32; st[ 3] = (u64x) (w1[2].s0, w1[2].s1) | (u64x) (w1[3].s0, w1[3].s1) << 32; st[ 4] = (u64x) (w2[0].s0, w2[0].s1) | (u64x) (w2[1].s0, w2[1].s1) << 32; st[ 5] = (u64x) (w2[2].s0, w2[2].s1) | (u64x) (w2[3].s0, w2[3].s1) << 32; st[ 6] = (u64x) (w3[0].s0, w3[0].s1) | (u64x) (w3[1].s0, w3[1].s1) << 32; st[ 7] = (u64x) (w3[2].s0, w3[2].s1) | (u64x) (w3[3].s0, w3[3].s1) << 32; #endif st[ 8] = 0; st[ 9] = 0; st[10] = 0; st[11] = 0; st[12] = 0; st[13] = 0; st[14] = 0; st[15] = 0; st[16] = 0; st[17] = 0; st[18] = 0; st[19] = 0; st[20] = 0; st[21] = 0; st[22] = 0; st[23] = 0; st[24] = 0; st[add80w] |= 0x8000000000000000; int round; for (round = 0; round < KECCAK_ROUNDS; round++) { // Theta u64x bc0 = Theta1 (0); u64x bc1 = Theta1 (1); u64x bc2 = Theta1 (2); u64x bc3 = Theta1 (3); u64x bc4 = Theta1 (4); u64x t; t = bc4 ^ rotl64 (bc1, 1); Theta2 (0); t = bc0 ^ rotl64 (bc2, 1); Theta2 (1); t = bc1 ^ rotl64 (bc3, 1); Theta2 (2); t = bc2 ^ rotl64 (bc4, 1); Theta2 (3); t = bc3 ^ rotl64 (bc0, 1); Theta2 (4); // Rho Pi t = st[1]; Rho_Pi (0); Rho_Pi (1); Rho_Pi (2); Rho_Pi (3); Rho_Pi (4); Rho_Pi (5); Rho_Pi (6); Rho_Pi (7); Rho_Pi (8); Rho_Pi (9); Rho_Pi (10); Rho_Pi (11); Rho_Pi (12); Rho_Pi (13); Rho_Pi (14); Rho_Pi (15); Rho_Pi (16); Rho_Pi (17); Rho_Pi (18); Rho_Pi (19); Rho_Pi (20); Rho_Pi (21); Rho_Pi (22); Rho_Pi (23); // Chi Chi (0); Chi (5); Chi (10); Chi (15); Chi (20); // Iota st[0] ^= keccakf_rndc[round]; } const u32x r0 = l32_from_64 (st[1]); const u32x r1 = h32_from_64 (st[1]); const u32x r2 = l32_from_64 (st[2]); const u32x r3 = h32_from_64 (st[2]); #include VECT_COMPARE_S } } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_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) { /** * base */ const u32 gid = get_global_id (0); if (gid >= gid_max) return; u32x w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32x w1[4]; w1[0] = 0; w1[1] = 0; w1[2] = 0; w1[3] = 0; u32x w2[4]; w2[0] = 0; w2[1] = 0; w2[2] = 0; w2[3] = 0; u32x w3[4]; w3[0] = 0; w3[1] = 0; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * main */ m05000m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_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) { /** * base */ const u32 gid = get_global_id (0); if (gid >= gid_max) return; u32x w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32x w1[4]; w1[0] = pws[gid].i[ 4]; w1[1] = pws[gid].i[ 5]; w1[2] = pws[gid].i[ 6]; w1[3] = pws[gid].i[ 7]; u32x w2[4]; w2[0] = 0; w2[1] = 0; w2[2] = 0; w2[3] = 0; u32x w3[4]; w3[0] = 0; w3[1] = 0; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * main */ m05000m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_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) { /** * base */ const u32 gid = get_global_id (0); if (gid >= gid_max) return; u32x w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32x w1[4]; w1[0] = pws[gid].i[ 4]; w1[1] = pws[gid].i[ 5]; w1[2] = pws[gid].i[ 6]; w1[3] = pws[gid].i[ 7]; u32x w2[4]; w2[0] = pws[gid].i[ 8]; w2[1] = pws[gid].i[ 9]; w2[2] = pws[gid].i[10]; w2[3] = pws[gid].i[11]; u32x w3[4]; w3[0] = pws[gid].i[12]; w3[1] = pws[gid].i[13]; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * main */ m05000m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_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) { /** * base */ const u32 gid = get_global_id (0); if (gid >= gid_max) return; u32x w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32x w1[4]; w1[0] = 0; w1[1] = 0; w1[2] = 0; w1[3] = 0; u32x w2[4]; w2[0] = 0; w2[1] = 0; w2[2] = 0; w2[3] = 0; u32x w3[4]; w3[0] = 0; w3[1] = 0; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * main */ m05000s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_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) { /** * base */ const u32 gid = get_global_id (0); if (gid >= gid_max) return; u32x w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32x w1[4]; w1[0] = pws[gid].i[ 4]; w1[1] = pws[gid].i[ 5]; w1[2] = pws[gid].i[ 6]; w1[3] = pws[gid].i[ 7]; u32x w2[4]; w2[0] = 0; w2[1] = 0; w2[2] = 0; w2[3] = 0; u32x w3[4]; w3[0] = 0; w3[1] = 0; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * main */ m05000s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_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) { /** * base */ const u32 gid = get_global_id (0); if (gid >= gid_max) return; u32x w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32x w1[4]; w1[0] = pws[gid].i[ 4]; w1[1] = pws[gid].i[ 5]; w1[2] = pws[gid].i[ 6]; w1[3] = pws[gid].i[ 7]; u32x w2[4]; w2[0] = pws[gid].i[ 8]; w2[1] = pws[gid].i[ 9]; w2[2] = pws[gid].i[10]; w2[3] = pws[gid].i[11]; u32x w3[4]; w3[0] = pws[gid].i[12]; w3[1] = pws[gid].i[13]; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * main */ m05000s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset); }