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"
27 #include "include/rp_gpu.h"
31 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
32 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
36 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
37 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
40 __device__ __constant__ u32 padding[8] =
60 __device__ static void swap (RC4_KEY *rc4_key, const u8 i, const u8 j)
65 rc4_key->S[i] = rc4_key->S[j];
69 __device__ static void rc4_init_16 (RC4_KEY *rc4_key, const u32 data[4])
74 u32 *ptr = (u32 *) rc4_key->S;
77 for (u32 i = 0; i < 64; i++)
82 const u32 d0 = data[0] >> 0;
83 const u32 d1 = data[0] >> 8;
84 const u32 d2 = data[0] >> 16;
85 const u32 d3 = data[0] >> 24;
86 const u32 d4 = data[1] >> 0;
92 for (i = 0; i < 255; i += 5)
94 j += rc4_key->S[i + 0]; j += d0; swap (rc4_key, i + 0, j);
95 j += rc4_key->S[i + 1]; j += d1; swap (rc4_key, i + 1, j);
96 j += rc4_key->S[i + 2]; j += d2; swap (rc4_key, i + 2, j);
97 j += rc4_key->S[i + 3]; j += d3; swap (rc4_key, i + 3, j);
98 j += rc4_key->S[i + 4]; j += d4; swap (rc4_key, i + 4, j);
101 j += rc4_key->S[i + 0]; j += d0; swap (rc4_key, i + 0, j);
104 __device__ static u8 rc4_next_16 (RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
106 for (u32 k = 0; k < 4; k++)
115 swap (rc4_key, i, j);
117 idx = rc4_key->S[i] + rc4_key->S[j];
119 xor4 |= rc4_key->S[idx] << 0;
124 swap (rc4_key, i, j);
126 idx = rc4_key->S[i] + rc4_key->S[j];
128 xor4 |= rc4_key->S[idx] << 8;
133 swap (rc4_key, i, j);
135 idx = rc4_key->S[i] + rc4_key->S[j];
137 xor4 |= rc4_key->S[idx] << 16;
142 swap (rc4_key, i, j);
144 idx = rc4_key->S[i] + rc4_key->S[j];
146 xor4 |= rc4_key->S[idx] << 24;
148 out[k] = in[k] ^ xor4;
154 __device__ __constant__ gpu_rule_t c_rules[1024];
156 extern "C" __global__ void __launch_bounds__ (64, 1) m10410_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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
162 const u32 lid = threadIdx.x;
168 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
170 if (gid >= gid_max) return;
174 pw_buf0[0] = pws[gid].i[ 0];
175 pw_buf0[1] = pws[gid].i[ 1];
176 pw_buf0[2] = pws[gid].i[ 2];
177 pw_buf0[3] = pws[gid].i[ 3];
181 pw_buf1[0] = pws[gid].i[ 4];
182 pw_buf1[1] = pws[gid].i[ 5];
183 pw_buf1[2] = pws[gid].i[ 6];
184 pw_buf1[3] = pws[gid].i[ 7];
186 const u32 pw_len = pws[gid].pw_len;
192 __shared__ RC4_KEY rc4_keys[64];
193 RC4_KEY *rc4_key = &rc4_keys[lid];
199 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
229 const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
240 rc4_init_16 (rc4_key, key);
244 rc4_next_16 (rc4_key, 0, 0, padding, out);
246 const u32x r0 = out[0];
247 const u32x r1 = out[1];
248 const u32x r2 = out[2];
249 const u32x r3 = out[3];
251 #include VECT_COMPARE_M
255 extern "C" __global__ void __launch_bounds__ (64, 1) m10410_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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
259 extern "C" __global__ void __launch_bounds__ (64, 1) m10410_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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
263 extern "C" __global__ void __launch_bounds__ (64, 1) m10410_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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
269 const u32 lid = threadIdx.x;
275 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
277 if (gid >= gid_max) return;
281 pw_buf0[0] = pws[gid].i[ 0];
282 pw_buf0[1] = pws[gid].i[ 1];
283 pw_buf0[2] = pws[gid].i[ 2];
284 pw_buf0[3] = pws[gid].i[ 3];
288 pw_buf1[0] = pws[gid].i[ 4];
289 pw_buf1[1] = pws[gid].i[ 5];
290 pw_buf1[2] = pws[gid].i[ 6];
291 pw_buf1[3] = pws[gid].i[ 7];
293 const u32 pw_len = pws[gid].pw_len;
299 const u32 search[4] =
301 digests_buf[digests_offset].digest_buf[DGST_R0],
302 digests_buf[digests_offset].digest_buf[DGST_R1],
303 digests_buf[digests_offset].digest_buf[DGST_R2],
304 digests_buf[digests_offset].digest_buf[DGST_R3]
311 __shared__ RC4_KEY rc4_keys[64];
312 RC4_KEY *rc4_key = &rc4_keys[lid];
318 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
348 const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
359 rc4_init_16 (rc4_key, key);
363 rc4_next_16 (rc4_key, 0, 0, padding, out);
365 const u32x r0 = out[0];
366 const u32x r1 = out[1];
367 const u32x r2 = out[2];
368 const u32x r3 = out[3];
370 #include VECT_COMPARE_S
374 extern "C" __global__ void __launch_bounds__ (64, 1) m10410_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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
378 extern "C" __global__ void __launch_bounds__ (64, 1) m10410_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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)