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 static void overwrite_at (u32 sw[16], const u32 w0, const u32 salt_len)
29 case 1: sw[0] = (sw[0] & 0xff000000) | (w0 >> 8);
30 sw[1] = (sw[1] & 0x00ffffff) | (w0 << 24);
32 case 2: sw[0] = (sw[0] & 0xffff0000) | (w0 >> 16);
33 sw[1] = (sw[1] & 0x0000ffff) | (w0 << 16);
35 case 3: sw[0] = (sw[0] & 0xffffff00) | (w0 >> 24);
36 sw[1] = (sw[1] & 0x000000ff) | (w0 << 8);
40 case 5: sw[1] = (sw[1] & 0xff000000) | (w0 >> 8);
41 sw[2] = (sw[2] & 0x00ffffff) | (w0 << 24);
43 case 6: sw[1] = (sw[1] & 0xffff0000) | (w0 >> 16);
44 sw[2] = (sw[2] & 0x0000ffff) | (w0 << 16);
46 case 7: sw[1] = (sw[1] & 0xffffff00) | (w0 >> 24);
47 sw[2] = (sw[2] & 0x000000ff) | (w0 << 8);
51 case 9: sw[2] = (sw[2] & 0xff000000) | (w0 >> 8);
52 sw[3] = (sw[3] & 0x00ffffff) | (w0 << 24);
54 case 10: sw[2] = (sw[2] & 0xffff0000) | (w0 >> 16);
55 sw[3] = (sw[3] & 0x0000ffff) | (w0 << 16);
57 case 11: sw[2] = (sw[2] & 0xffffff00) | (w0 >> 24);
58 sw[3] = (sw[3] & 0x000000ff) | (w0 << 8);
62 case 13: sw[3] = (sw[3] & 0xff000000) | (w0 >> 8);
63 sw[4] = (sw[4] & 0x00ffffff) | (w0 << 24);
65 case 14: sw[3] = (sw[3] & 0xffff0000) | (w0 >> 16);
66 sw[4] = (sw[4] & 0x0000ffff) | (w0 << 16);
68 case 15: sw[3] = (sw[3] & 0xffffff00) | (w0 >> 24);
69 sw[4] = (sw[4] & 0x000000ff) | (w0 << 8);
73 case 17: sw[4] = (sw[4] & 0xff000000) | (w0 >> 8);
74 sw[5] = (sw[5] & 0x00ffffff) | (w0 << 24);
76 case 18: sw[4] = (sw[4] & 0xffff0000) | (w0 >> 16);
77 sw[5] = (sw[5] & 0x0000ffff) | (w0 << 16);
79 case 19: sw[4] = (sw[4] & 0xffffff00) | (w0 >> 24);
80 sw[5] = (sw[5] & 0x000000ff) | (w0 << 8);
84 case 21: sw[5] = (sw[5] & 0xff000000) | (w0 >> 8);
85 sw[6] = (sw[6] & 0x00ffffff) | (w0 << 24);
87 case 22: sw[5] = (sw[5] & 0xffff0000) | (w0 >> 16);
88 sw[6] = (sw[6] & 0x0000ffff) | (w0 << 16);
90 case 23: sw[5] = (sw[5] & 0xffffff00) | (w0 >> 24);
91 sw[6] = (sw[6] & 0x000000ff) | (w0 << 8);
95 case 25: sw[6] = (sw[6] & 0xff000000) | (w0 >> 8);
96 sw[7] = (sw[7] & 0x00ffffff) | (w0 << 24);
98 case 26: sw[6] = (sw[6] & 0xffff0000) | (w0 >> 16);
99 sw[7] = (sw[7] & 0x0000ffff) | (w0 << 16);
101 case 27: sw[6] = (sw[6] & 0xffffff00) | (w0 >> 24);
102 sw[7] = (sw[7] & 0x000000ff) | (w0 << 8);
106 case 29: sw[7] = (sw[7] & 0xff000000) | (w0 >> 8);
107 sw[8] = (sw[8] & 0x00ffffff) | (w0 << 24);
109 case 30: sw[7] = (sw[7] & 0xffff0000) | (w0 >> 16);
110 sw[8] = (sw[8] & 0x0000ffff) | (w0 << 16);
112 case 31: sw[7] = (sw[7] & 0xffffff00) | (w0 >> 24);
113 sw[8] = (sw[8] & 0x000000ff) | (w0 << 8);
118 static void m00120m (u32 w0[4], u32 w1[4], u32 w2[4], u32 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)
124 const u32 gid = get_global_id (0);
125 const u32 lid = get_local_id (0);
133 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
134 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
135 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
136 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
140 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
141 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
142 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
143 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
159 const u32 salt_len = salt_bufs[salt_pos].salt_len;
161 const u32 pw_salt_len = pw_len + salt_len;
172 w0_t[0] = swap32 (w0[0]);
173 w0_t[1] = swap32 (w0[1]);
174 w0_t[2] = swap32 (w0[2]);
175 w0_t[3] = swap32 (w0[3]);
176 w1_t[0] = swap32 (w1[0]);
177 w1_t[1] = swap32 (w1[1]);
178 w1_t[2] = swap32 (w1[2]);
179 w1_t[3] = swap32 (w1[3]);
180 w2_t[0] = swap32 (w2[0]);
181 w2_t[1] = swap32 (w2[1]);
182 w2_t[2] = swap32 (w2[2]);
183 w2_t[3] = swap32 (w2[3]);
184 w3_t[0] = swap32 (w3[0]);
185 w3_t[1] = swap32 (w3[1]);
186 w3_t[2] = swap32 (w3[2]);
187 w3_t[3] = swap32 (w3[3]);
189 switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
191 w0_t[0] |= salt_buf0[0];
192 w0_t[1] |= salt_buf0[1];
193 w0_t[2] |= salt_buf0[2];
194 w0_t[3] |= salt_buf0[3];
195 w1_t[0] |= salt_buf1[0];
196 w1_t[1] |= salt_buf1[1];
197 w1_t[2] |= salt_buf1[2];
198 w1_t[3] |= salt_buf1[3];
199 w2_t[0] |= salt_buf2[0];
200 w2_t[1] |= salt_buf2[1];
201 w2_t[2] |= salt_buf2[2];
202 w2_t[3] |= salt_buf2[3];
203 w3_t[0] |= salt_buf3[0];
204 w3_t[1] |= salt_buf3[1];
205 w3_t[2] |= salt_buf3[2];
206 w3_t[3] |= salt_buf3[3];
208 w0_t[0] = swap32 (w0_t[0]);
209 w0_t[1] = swap32 (w0_t[1]);
210 w0_t[2] = swap32 (w0_t[2]);
211 w0_t[3] = swap32 (w0_t[3]);
212 w1_t[0] = swap32 (w1_t[0]);
213 w1_t[1] = swap32 (w1_t[1]);
214 w1_t[2] = swap32 (w1_t[2]);
215 w1_t[3] = swap32 (w1_t[3]);
216 w2_t[0] = swap32 (w2_t[0]);
217 w2_t[1] = swap32 (w2_t[1]);
218 w2_t[2] = swap32 (w2_t[2]);
219 w2_t[3] = swap32 (w2_t[3]);
220 w3_t[0] = swap32 (w3_t[0]);
221 w3_t[1] = swap32 (w3_t[1]);
222 w3_t[2] = swap32 (w3_t[2]);
223 w3_t[3] = swap32 (w3_t[3]);
231 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
233 const u32 w0r = bfs_buf[il_pos].i;
235 const u32 w0n = w0l | w0r;
256 overwrite_at (wx, w0n, salt_len);
278 w3_t[3] = pw_salt_len * 8;
293 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t[0]);
294 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t[1]);
295 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w0_t[2]);
296 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w0_t[3]);
297 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w1_t[0]);
298 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w1_t[1]);
299 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t[2]);
300 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t[3]);
301 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t[0]);
302 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w2_t[1]);
303 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w2_t[2]);
304 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w2_t[3]);
305 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w3_t[0]);
306 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t[1]);
307 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t[2]);
308 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w3_t[3]);
309 w0_t[0] = rotl32 ((w3_t[1] ^ w2_t[0] ^ w0_t[2] ^ w0_t[0]), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t[0]);
310 w0_t[1] = rotl32 ((w3_t[2] ^ w2_t[1] ^ w0_t[3] ^ w0_t[1]), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w0_t[1]);
311 w0_t[2] = rotl32 ((w3_t[3] ^ w2_t[2] ^ w1_t[0] ^ w0_t[2]), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w0_t[2]);
312 w0_t[3] = rotl32 ((w0_t[0] ^ w2_t[3] ^ w1_t[1] ^ w0_t[3]), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w0_t[3]);
317 w1_t[0] = rotl32 ((w0_t[1] ^ w3_t[0] ^ w1_t[2] ^ w1_t[0]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t[0]);
318 w1_t[1] = rotl32 ((w0_t[2] ^ w3_t[1] ^ w1_t[3] ^ w1_t[1]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w1_t[1]);
319 w1_t[2] = rotl32 ((w0_t[3] ^ w3_t[2] ^ w2_t[0] ^ w1_t[2]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w1_t[2]);
320 w1_t[3] = rotl32 ((w1_t[0] ^ w3_t[3] ^ w2_t[1] ^ w1_t[3]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t[3]);
321 w2_t[0] = rotl32 ((w1_t[1] ^ w0_t[0] ^ w2_t[2] ^ w2_t[0]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t[0]);
322 w2_t[1] = rotl32 ((w1_t[2] ^ w0_t[1] ^ w2_t[3] ^ w2_t[1]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w2_t[1]);
323 w2_t[2] = rotl32 ((w1_t[3] ^ w0_t[2] ^ w3_t[0] ^ w2_t[2]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t[2]);
324 w2_t[3] = rotl32 ((w2_t[0] ^ w0_t[3] ^ w3_t[1] ^ w2_t[3]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w2_t[3]);
325 w3_t[0] = rotl32 ((w2_t[1] ^ w1_t[0] ^ w3_t[2] ^ w3_t[0]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w3_t[0]);
326 w3_t[1] = rotl32 ((w2_t[2] ^ w1_t[1] ^ w3_t[3] ^ w3_t[1]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w3_t[1]);
327 w3_t[2] = rotl32 ((w2_t[3] ^ w1_t[2] ^ w0_t[0] ^ w3_t[2]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t[2]);
328 w3_t[3] = rotl32 ((w3_t[0] ^ w1_t[3] ^ w0_t[1] ^ w3_t[3]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w3_t[3]);
329 w0_t[0] = rotl32 ((w3_t[1] ^ w2_t[0] ^ w0_t[2] ^ w0_t[0]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t[0]);
330 w0_t[1] = rotl32 ((w3_t[2] ^ w2_t[1] ^ w0_t[3] ^ w0_t[1]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w0_t[1]);
331 w0_t[2] = rotl32 ((w3_t[3] ^ w2_t[2] ^ w1_t[0] ^ w0_t[2]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t[2]);
332 w0_t[3] = rotl32 ((w0_t[0] ^ w2_t[3] ^ w1_t[1] ^ w0_t[3]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w0_t[3]);
333 w1_t[0] = rotl32 ((w0_t[1] ^ w3_t[0] ^ w1_t[2] ^ w1_t[0]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w1_t[0]);
334 w1_t[1] = rotl32 ((w0_t[2] ^ w3_t[1] ^ w1_t[3] ^ w1_t[1]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w1_t[1]);
335 w1_t[2] = rotl32 ((w0_t[3] ^ w3_t[2] ^ w2_t[0] ^ w1_t[2]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t[2]);
336 w1_t[3] = rotl32 ((w1_t[0] ^ w3_t[3] ^ w2_t[1] ^ w1_t[3]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w1_t[3]);
341 w2_t[0] = rotl32 ((w1_t[1] ^ w0_t[0] ^ w2_t[2] ^ w2_t[0]), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t[0]);
342 w2_t[1] = rotl32 ((w1_t[2] ^ w0_t[1] ^ w2_t[3] ^ w2_t[1]), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w2_t[1]);
343 w2_t[2] = rotl32 ((w1_t[3] ^ w0_t[2] ^ w3_t[0] ^ w2_t[2]), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w2_t[2]);
344 w2_t[3] = rotl32 ((w2_t[0] ^ w0_t[3] ^ w3_t[1] ^ w2_t[3]), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w2_t[3]);
345 w3_t[0] = rotl32 ((w2_t[1] ^ w1_t[0] ^ w3_t[2] ^ w3_t[0]), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w3_t[0]);
346 w3_t[1] = rotl32 ((w2_t[2] ^ w1_t[1] ^ w3_t[3] ^ w3_t[1]), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w3_t[1]);
347 w3_t[2] = rotl32 ((w2_t[3] ^ w1_t[2] ^ w0_t[0] ^ w3_t[2]), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t[2]);
348 w3_t[3] = rotl32 ((w3_t[0] ^ w1_t[3] ^ w0_t[1] ^ w3_t[3]), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w3_t[3]);
349 w0_t[0] = rotl32 ((w3_t[1] ^ w2_t[0] ^ w0_t[2] ^ w0_t[0]), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t[0]);
350 w0_t[1] = rotl32 ((w3_t[2] ^ w2_t[1] ^ w0_t[3] ^ w0_t[1]), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w0_t[1]);
351 w0_t[2] = rotl32 ((w3_t[3] ^ w2_t[2] ^ w1_t[0] ^ w0_t[2]), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w0_t[2]);
352 w0_t[3] = rotl32 ((w0_t[0] ^ w2_t[3] ^ w1_t[1] ^ w0_t[3]), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w0_t[3]);
353 w1_t[0] = rotl32 ((w0_t[1] ^ w3_t[0] ^ w1_t[2] ^ w1_t[0]), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w1_t[0]);
354 w1_t[1] = rotl32 ((w0_t[2] ^ w3_t[1] ^ w1_t[3] ^ w1_t[1]), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w1_t[1]);
355 w1_t[2] = rotl32 ((w0_t[3] ^ w3_t[2] ^ w2_t[0] ^ w1_t[2]), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t[2]);
356 w1_t[3] = rotl32 ((w1_t[0] ^ w3_t[3] ^ w2_t[1] ^ w1_t[3]), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w1_t[3]);
357 w2_t[0] = rotl32 ((w1_t[1] ^ w0_t[0] ^ w2_t[2] ^ w2_t[0]), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w2_t[0]);
358 w2_t[1] = rotl32 ((w1_t[2] ^ w0_t[1] ^ w2_t[3] ^ w2_t[1]), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w2_t[1]);
359 w2_t[2] = rotl32 ((w1_t[3] ^ w0_t[2] ^ w3_t[0] ^ w2_t[2]), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w2_t[2]);
360 w2_t[3] = rotl32 ((w2_t[0] ^ w0_t[3] ^ w3_t[1] ^ w2_t[3]), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w2_t[3]);
365 w3_t[0] = rotl32 ((w2_t[1] ^ w1_t[0] ^ w3_t[2] ^ w3_t[0]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t[0]);
366 w3_t[1] = rotl32 ((w2_t[2] ^ w1_t[1] ^ w3_t[3] ^ w3_t[1]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w3_t[1]);
367 w3_t[2] = rotl32 ((w2_t[3] ^ w1_t[2] ^ w0_t[0] ^ w3_t[2]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t[2]);
368 w3_t[3] = rotl32 ((w3_t[0] ^ w1_t[3] ^ w0_t[1] ^ w3_t[3]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w3_t[3]);
369 w0_t[0] = rotl32 ((w3_t[1] ^ w2_t[0] ^ w0_t[2] ^ w0_t[0]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t[0]);
370 w0_t[1] = rotl32 ((w3_t[2] ^ w2_t[1] ^ w0_t[3] ^ w0_t[1]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w0_t[1]);
371 w0_t[2] = rotl32 ((w3_t[3] ^ w2_t[2] ^ w1_t[0] ^ w0_t[2]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w0_t[2]);
372 w0_t[3] = rotl32 ((w0_t[0] ^ w2_t[3] ^ w1_t[1] ^ w0_t[3]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t[3]);
373 w1_t[0] = rotl32 ((w0_t[1] ^ w3_t[0] ^ w1_t[2] ^ w1_t[0]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t[0]);
374 w1_t[1] = rotl32 ((w0_t[2] ^ w3_t[1] ^ w1_t[3] ^ w1_t[1]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w1_t[1]);
375 w1_t[2] = rotl32 ((w0_t[3] ^ w3_t[2] ^ w2_t[0] ^ w1_t[2]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t[2]);
376 w1_t[3] = rotl32 ((w1_t[0] ^ w3_t[3] ^ w2_t[1] ^ w1_t[3]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w1_t[3]);
377 w2_t[0] = rotl32 ((w1_t[1] ^ w0_t[0] ^ w2_t[2] ^ w2_t[0]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w2_t[0]);
378 w2_t[1] = rotl32 ((w1_t[2] ^ w0_t[1] ^ w2_t[3] ^ w2_t[1]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w2_t[1]);
379 w2_t[2] = rotl32 ((w1_t[3] ^ w0_t[2] ^ w3_t[0] ^ w2_t[2]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t[2]);
380 w2_t[3] = rotl32 ((w2_t[0] ^ w0_t[3] ^ w3_t[1] ^ w2_t[3]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w2_t[3]);
381 w3_t[0] = rotl32 ((w2_t[1] ^ w1_t[0] ^ w3_t[2] ^ w3_t[0]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w3_t[0]);
382 w3_t[1] = rotl32 ((w2_t[2] ^ w1_t[1] ^ w3_t[3] ^ w3_t[1]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t[1]);
383 w3_t[2] = rotl32 ((w2_t[3] ^ w1_t[2] ^ w0_t[0] ^ w3_t[2]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w3_t[2]);
384 w3_t[3] = rotl32 ((w3_t[0] ^ w1_t[3] ^ w0_t[1] ^ w3_t[3]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w3_t[3]);
395 static void m00120s (u32 w0[4], u32 w1[4], u32 w2[4], u32 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)
401 const u32 gid = get_global_id (0);
402 const u32 lid = get_local_id (0);
408 const u32 search[4] =
410 digests_buf[digests_offset].digest_buf[DGST_R0],
411 digests_buf[digests_offset].digest_buf[DGST_R1],
412 digests_buf[digests_offset].digest_buf[DGST_R2],
413 digests_buf[digests_offset].digest_buf[DGST_R3]
420 const u32 e_rev = rotl32 (search[1], 2u);
428 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
429 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
430 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
431 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
435 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
436 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
437 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
438 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
454 const u32 salt_len = salt_bufs[salt_pos].salt_len;
456 const u32 pw_salt_len = pw_len + salt_len;
467 w0_t[0] = swap32 (w0[0]);
468 w0_t[1] = swap32 (w0[1]);
469 w0_t[2] = swap32 (w0[2]);
470 w0_t[3] = swap32 (w0[3]);
471 w1_t[0] = swap32 (w1[0]);
472 w1_t[1] = swap32 (w1[1]);
473 w1_t[2] = swap32 (w1[2]);
474 w1_t[3] = swap32 (w1[3]);
475 w2_t[0] = swap32 (w2[0]);
476 w2_t[1] = swap32 (w2[1]);
477 w2_t[2] = swap32 (w2[2]);
478 w2_t[3] = swap32 (w2[3]);
479 w3_t[0] = swap32 (w3[0]);
480 w3_t[1] = swap32 (w3[1]);
481 w3_t[2] = swap32 (w3[2]);
482 w3_t[3] = swap32 (w3[3]);
484 switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
486 w0_t[0] |= salt_buf0[0];
487 w0_t[1] |= salt_buf0[1];
488 w0_t[2] |= salt_buf0[2];
489 w0_t[3] |= salt_buf0[3];
490 w1_t[0] |= salt_buf1[0];
491 w1_t[1] |= salt_buf1[1];
492 w1_t[2] |= salt_buf1[2];
493 w1_t[3] |= salt_buf1[3];
494 w2_t[0] |= salt_buf2[0];
495 w2_t[1] |= salt_buf2[1];
496 w2_t[2] |= salt_buf2[2];
497 w2_t[3] |= salt_buf2[3];
498 w3_t[0] |= salt_buf3[0];
499 w3_t[1] |= salt_buf3[1];
500 w3_t[2] |= salt_buf3[2];
501 w3_t[3] |= salt_buf3[3];
503 w0_t[0] = swap32 (w0_t[0]);
504 w0_t[1] = swap32 (w0_t[1]);
505 w0_t[2] = swap32 (w0_t[2]);
506 w0_t[3] = swap32 (w0_t[3]);
507 w1_t[0] = swap32 (w1_t[0]);
508 w1_t[1] = swap32 (w1_t[1]);
509 w1_t[2] = swap32 (w1_t[2]);
510 w1_t[3] = swap32 (w1_t[3]);
511 w2_t[0] = swap32 (w2_t[0]);
512 w2_t[1] = swap32 (w2_t[1]);
513 w2_t[2] = swap32 (w2_t[2]);
514 w2_t[3] = swap32 (w2_t[3]);
515 w3_t[0] = swap32 (w3_t[0]);
516 w3_t[1] = swap32 (w3_t[1]);
517 w3_t[2] = swap32 (w3_t[2]);
518 w3_t[3] = swap32 (w3_t[3]);
526 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
528 const u32 w0r = bfs_buf[il_pos].i;
530 const u32 w0n = w0l | w0r;
551 overwrite_at (wx, w0n, salt_len);
573 w3_t[3] = pw_salt_len * 8;
588 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t[0]);
589 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t[1]);
590 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w0_t[2]);
591 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w0_t[3]);
592 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w1_t[0]);
593 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w1_t[1]);
594 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t[2]);
595 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t[3]);
596 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t[0]);
597 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w2_t[1]);
598 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w2_t[2]);
599 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w2_t[3]);
600 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w3_t[0]);
601 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t[1]);
602 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t[2]);
603 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w3_t[3]);
604 w0_t[0] = rotl32 ((w3_t[1] ^ w2_t[0] ^ w0_t[2] ^ w0_t[0]), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t[0]);
605 w0_t[1] = rotl32 ((w3_t[2] ^ w2_t[1] ^ w0_t[3] ^ w0_t[1]), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w0_t[1]);
606 w0_t[2] = rotl32 ((w3_t[3] ^ w2_t[2] ^ w1_t[0] ^ w0_t[2]), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w0_t[2]);
607 w0_t[3] = rotl32 ((w0_t[0] ^ w2_t[3] ^ w1_t[1] ^ w0_t[3]), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w0_t[3]);
612 w1_t[0] = rotl32 ((w0_t[1] ^ w3_t[0] ^ w1_t[2] ^ w1_t[0]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t[0]);
613 w1_t[1] = rotl32 ((w0_t[2] ^ w3_t[1] ^ w1_t[3] ^ w1_t[1]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w1_t[1]);
614 w1_t[2] = rotl32 ((w0_t[3] ^ w3_t[2] ^ w2_t[0] ^ w1_t[2]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w1_t[2]);
615 w1_t[3] = rotl32 ((w1_t[0] ^ w3_t[3] ^ w2_t[1] ^ w1_t[3]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t[3]);
616 w2_t[0] = rotl32 ((w1_t[1] ^ w0_t[0] ^ w2_t[2] ^ w2_t[0]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t[0]);
617 w2_t[1] = rotl32 ((w1_t[2] ^ w0_t[1] ^ w2_t[3] ^ w2_t[1]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w2_t[1]);
618 w2_t[2] = rotl32 ((w1_t[3] ^ w0_t[2] ^ w3_t[0] ^ w2_t[2]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t[2]);
619 w2_t[3] = rotl32 ((w2_t[0] ^ w0_t[3] ^ w3_t[1] ^ w2_t[3]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w2_t[3]);
620 w3_t[0] = rotl32 ((w2_t[1] ^ w1_t[0] ^ w3_t[2] ^ w3_t[0]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w3_t[0]);
621 w3_t[1] = rotl32 ((w2_t[2] ^ w1_t[1] ^ w3_t[3] ^ w3_t[1]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w3_t[1]);
622 w3_t[2] = rotl32 ((w2_t[3] ^ w1_t[2] ^ w0_t[0] ^ w3_t[2]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t[2]);
623 w3_t[3] = rotl32 ((w3_t[0] ^ w1_t[3] ^ w0_t[1] ^ w3_t[3]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w3_t[3]);
624 w0_t[0] = rotl32 ((w3_t[1] ^ w2_t[0] ^ w0_t[2] ^ w0_t[0]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t[0]);
625 w0_t[1] = rotl32 ((w3_t[2] ^ w2_t[1] ^ w0_t[3] ^ w0_t[1]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w0_t[1]);
626 w0_t[2] = rotl32 ((w3_t[3] ^ w2_t[2] ^ w1_t[0] ^ w0_t[2]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t[2]);
627 w0_t[3] = rotl32 ((w0_t[0] ^ w2_t[3] ^ w1_t[1] ^ w0_t[3]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w0_t[3]);
628 w1_t[0] = rotl32 ((w0_t[1] ^ w3_t[0] ^ w1_t[2] ^ w1_t[0]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w1_t[0]);
629 w1_t[1] = rotl32 ((w0_t[2] ^ w3_t[1] ^ w1_t[3] ^ w1_t[1]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w1_t[1]);
630 w1_t[2] = rotl32 ((w0_t[3] ^ w3_t[2] ^ w2_t[0] ^ w1_t[2]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t[2]);
631 w1_t[3] = rotl32 ((w1_t[0] ^ w3_t[3] ^ w2_t[1] ^ w1_t[3]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w1_t[3]);
636 w2_t[0] = rotl32 ((w1_t[1] ^ w0_t[0] ^ w2_t[2] ^ w2_t[0]), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t[0]);
637 w2_t[1] = rotl32 ((w1_t[2] ^ w0_t[1] ^ w2_t[3] ^ w2_t[1]), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w2_t[1]);
638 w2_t[2] = rotl32 ((w1_t[3] ^ w0_t[2] ^ w3_t[0] ^ w2_t[2]), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w2_t[2]);
639 w2_t[3] = rotl32 ((w2_t[0] ^ w0_t[3] ^ w3_t[1] ^ w2_t[3]), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w2_t[3]);
640 w3_t[0] = rotl32 ((w2_t[1] ^ w1_t[0] ^ w3_t[2] ^ w3_t[0]), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w3_t[0]);
641 w3_t[1] = rotl32 ((w2_t[2] ^ w1_t[1] ^ w3_t[3] ^ w3_t[1]), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w3_t[1]);
642 w3_t[2] = rotl32 ((w2_t[3] ^ w1_t[2] ^ w0_t[0] ^ w3_t[2]), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t[2]);
643 w3_t[3] = rotl32 ((w3_t[0] ^ w1_t[3] ^ w0_t[1] ^ w3_t[3]), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w3_t[3]);
644 w0_t[0] = rotl32 ((w3_t[1] ^ w2_t[0] ^ w0_t[2] ^ w0_t[0]), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t[0]);
645 w0_t[1] = rotl32 ((w3_t[2] ^ w2_t[1] ^ w0_t[3] ^ w0_t[1]), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w0_t[1]);
646 w0_t[2] = rotl32 ((w3_t[3] ^ w2_t[2] ^ w1_t[0] ^ w0_t[2]), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w0_t[2]);
647 w0_t[3] = rotl32 ((w0_t[0] ^ w2_t[3] ^ w1_t[1] ^ w0_t[3]), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w0_t[3]);
648 w1_t[0] = rotl32 ((w0_t[1] ^ w3_t[0] ^ w1_t[2] ^ w1_t[0]), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w1_t[0]);
649 w1_t[1] = rotl32 ((w0_t[2] ^ w3_t[1] ^ w1_t[3] ^ w1_t[1]), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w1_t[1]);
650 w1_t[2] = rotl32 ((w0_t[3] ^ w3_t[2] ^ w2_t[0] ^ w1_t[2]), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t[2]);
651 w1_t[3] = rotl32 ((w1_t[0] ^ w3_t[3] ^ w2_t[1] ^ w1_t[3]), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w1_t[3]);
652 w2_t[0] = rotl32 ((w1_t[1] ^ w0_t[0] ^ w2_t[2] ^ w2_t[0]), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w2_t[0]);
653 w2_t[1] = rotl32 ((w1_t[2] ^ w0_t[1] ^ w2_t[3] ^ w2_t[1]), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w2_t[1]);
654 w2_t[2] = rotl32 ((w1_t[3] ^ w0_t[2] ^ w3_t[0] ^ w2_t[2]), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w2_t[2]);
655 w2_t[3] = rotl32 ((w2_t[0] ^ w0_t[3] ^ w3_t[1] ^ w2_t[3]), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w2_t[3]);
660 w3_t[0] = rotl32 ((w2_t[1] ^ w1_t[0] ^ w3_t[2] ^ w3_t[0]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t[0]);
661 w3_t[1] = rotl32 ((w2_t[2] ^ w1_t[1] ^ w3_t[3] ^ w3_t[1]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w3_t[1]);
662 w3_t[2] = rotl32 ((w2_t[3] ^ w1_t[2] ^ w0_t[0] ^ w3_t[2]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t[2]);
663 w3_t[3] = rotl32 ((w3_t[0] ^ w1_t[3] ^ w0_t[1] ^ w3_t[3]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w3_t[3]);
664 w0_t[0] = rotl32 ((w3_t[1] ^ w2_t[0] ^ w0_t[2] ^ w0_t[0]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t[0]);
665 w0_t[1] = rotl32 ((w3_t[2] ^ w2_t[1] ^ w0_t[3] ^ w0_t[1]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w0_t[1]);
666 w0_t[2] = rotl32 ((w3_t[3] ^ w2_t[2] ^ w1_t[0] ^ w0_t[2]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w0_t[2]);
667 w0_t[3] = rotl32 ((w0_t[0] ^ w2_t[3] ^ w1_t[1] ^ w0_t[3]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t[3]);
668 w1_t[0] = rotl32 ((w0_t[1] ^ w3_t[0] ^ w1_t[2] ^ w1_t[0]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t[0]);
669 w1_t[1] = rotl32 ((w0_t[2] ^ w3_t[1] ^ w1_t[3] ^ w1_t[1]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w1_t[1]);
670 w1_t[2] = rotl32 ((w0_t[3] ^ w3_t[2] ^ w2_t[0] ^ w1_t[2]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t[2]);
671 w1_t[3] = rotl32 ((w1_t[0] ^ w3_t[3] ^ w2_t[1] ^ w1_t[3]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w1_t[3]);
672 w2_t[0] = rotl32 ((w1_t[1] ^ w0_t[0] ^ w2_t[2] ^ w2_t[0]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w2_t[0]);
673 w2_t[1] = rotl32 ((w1_t[2] ^ w0_t[1] ^ w2_t[3] ^ w2_t[1]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w2_t[1]);
674 w2_t[2] = rotl32 ((w1_t[3] ^ w0_t[2] ^ w3_t[0] ^ w2_t[2]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t[2]);
675 w2_t[3] = rotl32 ((w2_t[0] ^ w0_t[3] ^ w3_t[1] ^ w2_t[3]), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w2_t[3]);
676 w3_t[0] = rotl32 ((w2_t[1] ^ w1_t[0] ^ w3_t[2] ^ w3_t[0]), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w3_t[0]);
678 if (allx (e != e_rev)) continue;
680 w3_t[1] = rotl32 ((w2_t[2] ^ w1_t[1] ^ w3_t[3] ^ w3_t[1]), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t[1]);
681 w3_t[2] = rotl32 ((w2_t[3] ^ w1_t[2] ^ w0_t[0] ^ w3_t[2]), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w3_t[2]);
682 w3_t[3] = rotl32 ((w3_t[0] ^ w1_t[3] ^ w0_t[1] ^ w3_t[3]), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w3_t[3]);
694 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_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)
700 const u32 gid = get_global_id (0);
702 if (gid >= gid_max) return;
706 w0[0] = pws[gid].i[ 0];
707 w0[1] = pws[gid].i[ 1];
708 w0[2] = pws[gid].i[ 2];
709 w0[3] = pws[gid].i[ 3];
732 const u32 pw_len = pws[gid].pw_len;
738 m00120m (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);
741 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_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)
747 const u32 gid = get_global_id (0);
749 if (gid >= gid_max) return;
753 w0[0] = pws[gid].i[ 0];
754 w0[1] = pws[gid].i[ 1];
755 w0[2] = pws[gid].i[ 2];
756 w0[3] = pws[gid].i[ 3];
760 w1[0] = pws[gid].i[ 4];
761 w1[1] = pws[gid].i[ 5];
762 w1[2] = pws[gid].i[ 6];
763 w1[3] = pws[gid].i[ 7];
779 const u32 pw_len = pws[gid].pw_len;
785 m00120m (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);
788 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_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)
794 const u32 gid = get_global_id (0);
796 if (gid >= gid_max) return;
800 w0[0] = pws[gid].i[ 0];
801 w0[1] = pws[gid].i[ 1];
802 w0[2] = pws[gid].i[ 2];
803 w0[3] = pws[gid].i[ 3];
807 w1[0] = pws[gid].i[ 4];
808 w1[1] = pws[gid].i[ 5];
809 w1[2] = pws[gid].i[ 6];
810 w1[3] = pws[gid].i[ 7];
814 w2[0] = pws[gid].i[ 8];
815 w2[1] = pws[gid].i[ 9];
816 w2[2] = pws[gid].i[10];
817 w2[3] = pws[gid].i[11];
821 w3[0] = pws[gid].i[12];
822 w3[1] = pws[gid].i[13];
826 const u32 pw_len = pws[gid].pw_len;
832 m00120m (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);
835 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_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)
841 const u32 gid = get_global_id (0);
843 if (gid >= gid_max) return;
847 w0[0] = pws[gid].i[ 0];
848 w0[1] = pws[gid].i[ 1];
849 w0[2] = pws[gid].i[ 2];
850 w0[3] = pws[gid].i[ 3];
873 const u32 pw_len = pws[gid].pw_len;
879 m00120s (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);
882 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_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)
888 const u32 gid = get_global_id (0);
890 if (gid >= gid_max) return;
894 w0[0] = pws[gid].i[ 0];
895 w0[1] = pws[gid].i[ 1];
896 w0[2] = pws[gid].i[ 2];
897 w0[3] = pws[gid].i[ 3];
901 w1[0] = pws[gid].i[ 4];
902 w1[1] = pws[gid].i[ 5];
903 w1[2] = pws[gid].i[ 6];
904 w1[3] = pws[gid].i[ 7];
920 const u32 pw_len = pws[gid].pw_len;
926 m00120s (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);
929 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_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)
935 const u32 gid = get_global_id (0);
937 if (gid >= gid_max) return;
941 w0[0] = pws[gid].i[ 0];
942 w0[1] = pws[gid].i[ 1];
943 w0[2] = pws[gid].i[ 2];
944 w0[3] = pws[gid].i[ 3];
948 w1[0] = pws[gid].i[ 4];
949 w1[1] = pws[gid].i[ 5];
950 w1[2] = pws[gid].i[ 6];
951 w1[3] = pws[gid].i[ 7];
955 w2[0] = pws[gid].i[ 8];
956 w2[1] = pws[gid].i[ 9];
957 w2[2] = pws[gid].i[10];
958 w2[3] = pws[gid].i[11];
962 w3[0] = pws[gid].i[12];
963 w3[1] = pws[gid].i[13];
967 const u32 pw_len = pws[gid].pw_len;
973 m00120s (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);