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 __device__ __constant__ bf_t c_bfs[1024];
45 __device__ static void m05100m (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], 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)
51 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
52 const u32 lid = threadIdx.x;
66 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
68 const u32 w0r = c_bfs[il_pos].i;
77 MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
78 MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
79 MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
80 MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
81 MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
82 MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
83 MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
84 MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
85 MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
86 MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
87 MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
88 MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
89 MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
90 MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
91 MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
92 MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
94 MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
95 MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
96 MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
97 MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
98 MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
99 MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
100 MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
101 MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
102 MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
103 MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
104 MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
105 MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
106 MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
107 MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
108 MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
109 MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
111 MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20);
112 MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21);
113 MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22);
114 MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23);
115 MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20);
116 MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21);
117 MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22);
118 MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23);
119 MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20);
120 MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21);
121 MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22);
122 MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23);
123 MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20);
124 MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21);
125 MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22);
126 MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23);
128 MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
129 MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
130 MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
131 MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
132 MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
133 MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
134 MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
135 MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
136 MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
137 MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
138 MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
139 MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
140 MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
141 MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
142 MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
143 MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
156 #include VECT_COMPARE_M
165 #include VECT_COMPARE_M
174 #include VECT_COMPARE_M
179 __device__ static void m05100s (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], 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)
185 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
186 const u32 lid = threadIdx.x;
198 const u32 search[4] =
200 digests_buf[digests_offset].digest_buf[DGST_R0],
201 digests_buf[digests_offset].digest_buf[DGST_R1],
202 digests_buf[digests_offset].digest_buf[DGST_R2],
203 digests_buf[digests_offset].digest_buf[DGST_R3]
212 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
214 const u32 w0r = c_bfs[il_pos].i;
223 MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
224 MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
225 MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
226 MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
227 MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
228 MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
229 MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
230 MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
231 MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
232 MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
233 MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
234 MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
235 MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
236 MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
237 MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
238 MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
240 MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
241 MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
242 MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
243 MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
244 MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
245 MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
246 MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
247 MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
248 MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
249 MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
250 MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
251 MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
252 MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
253 MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
254 MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
255 MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
257 MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20);
258 MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21);
259 MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22);
260 MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23);
261 MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20);
262 MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21);
263 MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22);
264 MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23);
265 MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20);
266 MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21);
267 MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22);
268 MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23);
269 MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20);
270 MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21);
271 MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22);
272 MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23);
274 MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
275 MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
276 MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
277 MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
278 MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
279 MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
280 MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
281 MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
282 MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
283 MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
284 MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
285 MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
286 MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
287 MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
288 MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
289 MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
302 #include VECT_COMPARE_M
311 #include VECT_COMPARE_M
320 #include VECT_COMPARE_M
325 extern "C" __global__ void __launch_bounds__ (256, 1) m05100_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)
331 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
333 if (gid >= gid_max) return;
337 w0[0] = pws[gid].i[ 0];
338 w0[1] = pws[gid].i[ 1];
339 w0[2] = pws[gid].i[ 2];
340 w0[3] = pws[gid].i[ 3];
363 const u32 pw_len = pws[gid].pw_len;
369 m05100m (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);
372 extern "C" __global__ void __launch_bounds__ (256, 1) m05100_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)
378 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
380 if (gid >= gid_max) return;
384 w0[0] = pws[gid].i[ 0];
385 w0[1] = pws[gid].i[ 1];
386 w0[2] = pws[gid].i[ 2];
387 w0[3] = pws[gid].i[ 3];
391 w1[0] = pws[gid].i[ 4];
392 w1[1] = pws[gid].i[ 5];
393 w1[2] = pws[gid].i[ 6];
394 w1[3] = pws[gid].i[ 7];
410 const u32 pw_len = pws[gid].pw_len;
416 m05100m (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);
419 extern "C" __global__ void __launch_bounds__ (256, 1) m05100_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)
425 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
427 if (gid >= gid_max) return;
431 w0[0] = pws[gid].i[ 0];
432 w0[1] = pws[gid].i[ 1];
433 w0[2] = pws[gid].i[ 2];
434 w0[3] = pws[gid].i[ 3];
438 w1[0] = pws[gid].i[ 4];
439 w1[1] = pws[gid].i[ 5];
440 w1[2] = pws[gid].i[ 6];
441 w1[3] = pws[gid].i[ 7];
445 w2[0] = pws[gid].i[ 8];
446 w2[1] = pws[gid].i[ 9];
447 w2[2] = pws[gid].i[10];
448 w2[3] = pws[gid].i[11];
452 w3[0] = pws[gid].i[12];
453 w3[1] = pws[gid].i[13];
457 const u32 pw_len = pws[gid].pw_len;
463 m05100m (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);
466 extern "C" __global__ void __launch_bounds__ (256, 1) m05100_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)
472 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
474 if (gid >= gid_max) return;
478 w0[0] = pws[gid].i[ 0];
479 w0[1] = pws[gid].i[ 1];
480 w0[2] = pws[gid].i[ 2];
481 w0[3] = pws[gid].i[ 3];
504 const u32 pw_len = pws[gid].pw_len;
510 m05100s (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);
513 extern "C" __global__ void __launch_bounds__ (256, 1) m05100_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)
519 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
521 if (gid >= gid_max) return;
525 w0[0] = pws[gid].i[ 0];
526 w0[1] = pws[gid].i[ 1];
527 w0[2] = pws[gid].i[ 2];
528 w0[3] = pws[gid].i[ 3];
532 w1[0] = pws[gid].i[ 4];
533 w1[1] = pws[gid].i[ 5];
534 w1[2] = pws[gid].i[ 6];
535 w1[3] = pws[gid].i[ 7];
551 const u32 pw_len = pws[gid].pw_len;
557 m05100s (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);
560 extern "C" __global__ void __launch_bounds__ (256, 1) m05100_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)
566 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
568 if (gid >= gid_max) return;
572 w0[0] = pws[gid].i[ 0];
573 w0[1] = pws[gid].i[ 1];
574 w0[2] = pws[gid].i[ 2];
575 w0[3] = pws[gid].i[ 3];
579 w1[0] = pws[gid].i[ 4];
580 w1[1] = pws[gid].i[ 5];
581 w1[2] = pws[gid].i[ 6];
582 w1[3] = pws[gid].i[ 7];
586 w2[0] = pws[gid].i[ 8];
587 w2[1] = pws[gid].i[ 9];
588 w2[2] = pws[gid].i[10];
589 w2[3] = pws[gid].i[11];
593 w3[0] = pws[gid].i[12];
594 w3[1] = pws[gid].i[13];
598 const u32 pw_len = pws[gid].pw_len;
604 m05100s (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);