2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
28 #include "include/kernel_functions.c"
29 #include "types_amd.c"
30 #include "common_amd.c"
33 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
34 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
38 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
39 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
43 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
44 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
47 static void overwrite_at (u32x sw[16], const u32x w0, const u32 salt_len)
53 case 1: sw[0] = (sw[0] & 0xff000000) | (w0 >> 8);
54 sw[1] = (sw[1] & 0x00ffffff) | (w0 << 24);
56 case 2: sw[0] = (sw[0] & 0xffff0000) | (w0 >> 16);
57 sw[1] = (sw[1] & 0x0000ffff) | (w0 << 16);
59 case 3: sw[0] = (sw[0] & 0xffffff00) | (w0 >> 24);
60 sw[1] = (sw[1] & 0x000000ff) | (w0 << 8);
64 case 5: sw[1] = (sw[1] & 0xff000000) | (w0 >> 8);
65 sw[2] = (sw[2] & 0x00ffffff) | (w0 << 24);
67 case 6: sw[1] = (sw[1] & 0xffff0000) | (w0 >> 16);
68 sw[2] = (sw[2] & 0x0000ffff) | (w0 << 16);
70 case 7: sw[1] = (sw[1] & 0xffffff00) | (w0 >> 24);
71 sw[2] = (sw[2] & 0x000000ff) | (w0 << 8);
75 case 9: sw[2] = (sw[2] & 0xff000000) | (w0 >> 8);
76 sw[3] = (sw[3] & 0x00ffffff) | (w0 << 24);
78 case 10: sw[2] = (sw[2] & 0xffff0000) | (w0 >> 16);
79 sw[3] = (sw[3] & 0x0000ffff) | (w0 << 16);
81 case 11: sw[2] = (sw[2] & 0xffffff00) | (w0 >> 24);
82 sw[3] = (sw[3] & 0x000000ff) | (w0 << 8);
86 case 13: sw[3] = (sw[3] & 0xff000000) | (w0 >> 8);
87 sw[4] = (sw[4] & 0x00ffffff) | (w0 << 24);
89 case 14: sw[3] = (sw[3] & 0xffff0000) | (w0 >> 16);
90 sw[4] = (sw[4] & 0x0000ffff) | (w0 << 16);
92 case 15: sw[3] = (sw[3] & 0xffffff00) | (w0 >> 24);
93 sw[4] = (sw[4] & 0x000000ff) | (w0 << 8);
97 case 17: sw[4] = (sw[4] & 0xff000000) | (w0 >> 8);
98 sw[5] = (sw[5] & 0x00ffffff) | (w0 << 24);
100 case 18: sw[4] = (sw[4] & 0xffff0000) | (w0 >> 16);
101 sw[5] = (sw[5] & 0x0000ffff) | (w0 << 16);
103 case 19: sw[4] = (sw[4] & 0xffffff00) | (w0 >> 24);
104 sw[5] = (sw[5] & 0x000000ff) | (w0 << 8);
108 case 21: sw[5] = (sw[5] & 0xff000000) | (w0 >> 8);
109 sw[6] = (sw[6] & 0x00ffffff) | (w0 << 24);
111 case 22: sw[5] = (sw[5] & 0xffff0000) | (w0 >> 16);
112 sw[6] = (sw[6] & 0x0000ffff) | (w0 << 16);
114 case 23: sw[5] = (sw[5] & 0xffffff00) | (w0 >> 24);
115 sw[6] = (sw[6] & 0x000000ff) | (w0 << 8);
119 case 25: sw[6] = (sw[6] & 0xff000000) | (w0 >> 8);
120 sw[7] = (sw[7] & 0x00ffffff) | (w0 << 24);
122 case 26: sw[6] = (sw[6] & 0xffff0000) | (w0 >> 16);
123 sw[7] = (sw[7] & 0x0000ffff) | (w0 << 16);
125 case 27: sw[6] = (sw[6] & 0xffffff00) | (w0 >> 24);
126 sw[7] = (sw[7] & 0x000000ff) | (w0 << 8);
130 case 29: sw[7] = (sw[7] & 0xff000000) | (w0 >> 8);
131 sw[8] = (sw[8] & 0x00ffffff) | (w0 << 24);
133 case 30: sw[7] = (sw[7] & 0xffff0000) | (w0 >> 16);
134 sw[8] = (sw[8] & 0x0000ffff) | (w0 << 16);
136 case 31: sw[7] = (sw[7] & 0xffffff00) | (w0 >> 24);
137 sw[8] = (sw[8] & 0x000000ff) | (w0 << 8);
142 static void m00140m (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)
148 const u32 gid = get_global_id (0);
149 const u32 lid = get_local_id (0);
157 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
158 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
159 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
160 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
164 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
165 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
166 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
167 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
183 const u32 salt_len = salt_bufs[salt_pos].salt_len;
185 const u32 pw_salt_len = pw_len + salt_len;
196 w0_t[0] = swap_workaround (w0[0]);
197 w0_t[1] = swap_workaround (w0[1]);
198 w0_t[2] = swap_workaround (w0[2]);
199 w0_t[3] = swap_workaround (w0[3]);
200 w1_t[0] = swap_workaround (w1[0]);
201 w1_t[1] = swap_workaround (w1[1]);
202 w1_t[2] = swap_workaround (w1[2]);
203 w1_t[3] = swap_workaround (w1[3]);
204 w2_t[0] = swap_workaround (w2[0]);
205 w2_t[1] = swap_workaround (w2[1]);
206 w2_t[2] = swap_workaround (w2[2]);
207 w2_t[3] = swap_workaround (w2[3]);
208 w3_t[0] = swap_workaround (w3[0]);
209 w3_t[1] = swap_workaround (w3[1]);
210 w3_t[2] = swap_workaround (w3[2]);
211 w3_t[3] = swap_workaround (w3[3]);
213 switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
215 w0_t[0] |= salt_buf0[0];
216 w0_t[1] |= salt_buf0[1];
217 w0_t[2] |= salt_buf0[2];
218 w0_t[3] |= salt_buf0[3];
219 w1_t[0] |= salt_buf1[0];
220 w1_t[1] |= salt_buf1[1];
221 w1_t[2] |= salt_buf1[2];
222 w1_t[3] |= salt_buf1[3];
223 w2_t[0] |= salt_buf2[0];
224 w2_t[1] |= salt_buf2[1];
225 w2_t[2] |= salt_buf2[2];
226 w2_t[3] |= salt_buf2[3];
227 w3_t[0] |= salt_buf3[0];
228 w3_t[1] |= salt_buf3[1];
229 w3_t[2] |= salt_buf3[2];
230 w3_t[3] |= salt_buf3[3];
232 w0_t[0] = swap_workaround (w0_t[0]);
233 w0_t[1] = swap_workaround (w0_t[1]);
234 w0_t[2] = swap_workaround (w0_t[2]);
235 w0_t[3] = swap_workaround (w0_t[3]);
236 w1_t[0] = swap_workaround (w1_t[0]);
237 w1_t[1] = swap_workaround (w1_t[1]);
238 w1_t[2] = swap_workaround (w1_t[2]);
239 w1_t[3] = swap_workaround (w1_t[3]);
240 w2_t[0] = swap_workaround (w2_t[0]);
241 w2_t[1] = swap_workaround (w2_t[1]);
242 w2_t[2] = swap_workaround (w2_t[2]);
243 w2_t[3] = swap_workaround (w2_t[3]);
244 w3_t[0] = swap_workaround (w3_t[0]);
245 w3_t[1] = swap_workaround (w3_t[1]);
246 w3_t[2] = swap_workaround (w3_t[2]);
247 w3_t[3] = swap_workaround (w3_t[3]);
255 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
257 const u32 w0r = bfs_buf[il_pos].i;
259 const u32x w0n = w0l | w0r;
280 overwrite_at (wx, w0n, salt_len);
302 w3_t[3] = pw_salt_len * 8;
317 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t[0]);
318 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t[1]);
319 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w0_t[2]);
320 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w0_t[3]);
321 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w1_t[0]);
322 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w1_t[1]);
323 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t[2]);
324 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t[3]);
325 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t[0]);
326 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w2_t[1]);
327 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w2_t[2]);
328 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w2_t[3]);
329 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w3_t[0]);
330 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t[1]);
331 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t[2]);
332 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w3_t[3]);
333 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]);
334 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]);
335 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]);
336 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]);
341 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]);
342 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]);
343 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]);
344 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]);
345 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]);
346 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]);
347 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]);
348 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]);
349 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]);
350 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]);
351 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]);
352 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]);
353 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]);
354 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]);
355 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]);
356 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]);
357 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]);
358 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]);
359 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]);
360 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]);
365 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]);
366 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]);
367 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]);
368 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]);
369 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]);
370 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]);
371 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]);
372 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]);
373 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]);
374 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]);
375 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]);
376 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]);
377 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]);
378 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]);
379 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]);
380 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]);
381 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]);
382 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]);
383 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]);
384 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]);
389 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]);
390 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]);
391 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]);
392 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]);
393 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]);
394 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]);
395 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]);
396 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]);
397 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]);
398 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]);
399 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]);
400 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]);
401 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]);
402 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]);
403 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]);
404 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]);
405 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]);
406 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]);
407 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]);
408 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]);
415 #include VECT_COMPARE_M
419 static void m00140s (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)
425 const u32 gid = get_global_id (0);
426 const u32 lid = get_local_id (0);
432 const u32 search[4] =
434 digests_buf[digests_offset].digest_buf[DGST_R0],
435 digests_buf[digests_offset].digest_buf[DGST_R1],
436 digests_buf[digests_offset].digest_buf[DGST_R2],
437 digests_buf[digests_offset].digest_buf[DGST_R3]
444 const u32 e_rev = rotl32 (search[1], 2u);
452 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
453 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
454 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
455 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
459 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
460 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
461 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
462 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
478 const u32 salt_len = salt_bufs[salt_pos].salt_len;
480 const u32 pw_salt_len = pw_len + salt_len;
491 w0_t[0] = swap_workaround (w0[0]);
492 w0_t[1] = swap_workaround (w0[1]);
493 w0_t[2] = swap_workaround (w0[2]);
494 w0_t[3] = swap_workaround (w0[3]);
495 w1_t[0] = swap_workaround (w1[0]);
496 w1_t[1] = swap_workaround (w1[1]);
497 w1_t[2] = swap_workaround (w1[2]);
498 w1_t[3] = swap_workaround (w1[3]);
499 w2_t[0] = swap_workaround (w2[0]);
500 w2_t[1] = swap_workaround (w2[1]);
501 w2_t[2] = swap_workaround (w2[2]);
502 w2_t[3] = swap_workaround (w2[3]);
503 w3_t[0] = swap_workaround (w3[0]);
504 w3_t[1] = swap_workaround (w3[1]);
505 w3_t[2] = swap_workaround (w3[2]);
506 w3_t[3] = swap_workaround (w3[3]);
508 switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
510 w0_t[0] |= salt_buf0[0];
511 w0_t[1] |= salt_buf0[1];
512 w0_t[2] |= salt_buf0[2];
513 w0_t[3] |= salt_buf0[3];
514 w1_t[0] |= salt_buf1[0];
515 w1_t[1] |= salt_buf1[1];
516 w1_t[2] |= salt_buf1[2];
517 w1_t[3] |= salt_buf1[3];
518 w2_t[0] |= salt_buf2[0];
519 w2_t[1] |= salt_buf2[1];
520 w2_t[2] |= salt_buf2[2];
521 w2_t[3] |= salt_buf2[3];
522 w3_t[0] |= salt_buf3[0];
523 w3_t[1] |= salt_buf3[1];
524 w3_t[2] |= salt_buf3[2];
525 w3_t[3] |= salt_buf3[3];
527 w0_t[0] = swap_workaround (w0_t[0]);
528 w0_t[1] = swap_workaround (w0_t[1]);
529 w0_t[2] = swap_workaround (w0_t[2]);
530 w0_t[3] = swap_workaround (w0_t[3]);
531 w1_t[0] = swap_workaround (w1_t[0]);
532 w1_t[1] = swap_workaround (w1_t[1]);
533 w1_t[2] = swap_workaround (w1_t[2]);
534 w1_t[3] = swap_workaround (w1_t[3]);
535 w2_t[0] = swap_workaround (w2_t[0]);
536 w2_t[1] = swap_workaround (w2_t[1]);
537 w2_t[2] = swap_workaround (w2_t[2]);
538 w2_t[3] = swap_workaround (w2_t[3]);
539 w3_t[0] = swap_workaround (w3_t[0]);
540 w3_t[1] = swap_workaround (w3_t[1]);
541 w3_t[2] = swap_workaround (w3_t[2]);
542 w3_t[3] = swap_workaround (w3_t[3]);
550 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
552 const u32 w0r = bfs_buf[il_pos].i;
554 const u32x w0n = w0l | w0r;
575 overwrite_at (wx, w0n, salt_len);
597 w3_t[3] = pw_salt_len * 8;
612 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t[0]);
613 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t[1]);
614 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w0_t[2]);
615 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w0_t[3]);
616 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w1_t[0]);
617 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w1_t[1]);
618 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t[2]);
619 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t[3]);
620 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t[0]);
621 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w2_t[1]);
622 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w2_t[2]);
623 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w2_t[3]);
624 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w3_t[0]);
625 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t[1]);
626 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t[2]);
627 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w3_t[3]);
628 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]);
629 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]);
630 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]);
631 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]);
636 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]);
637 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]);
638 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]);
639 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]);
640 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]);
641 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]);
642 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]);
643 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]);
644 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]);
645 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]);
646 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]);
647 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]);
648 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]);
649 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]);
650 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]);
651 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]);
652 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]);
653 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]);
654 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]);
655 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]);
660 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]);
661 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]);
662 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]);
663 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]);
664 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]);
665 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]);
666 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]);
667 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]);
668 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]);
669 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]);
670 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]);
671 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]);
672 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]);
673 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]);
674 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]);
675 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]);
676 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]);
677 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]);
678 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]);
679 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]);
684 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]);
685 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]);
686 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]);
687 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]);
688 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]);
689 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]);
690 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]);
691 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]);
692 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]);
693 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]);
694 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]);
695 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]);
696 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]);
697 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]);
698 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]);
699 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]);
700 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]);
702 if (allx (e != e_rev)) continue;
704 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]);
705 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]);
706 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]);
714 #include VECT_COMPARE_S
718 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_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)
724 const u32 gid = get_global_id (0);
726 if (gid >= gid_max) return;
730 w0[0] = pws[gid].i[ 0];
731 w0[1] = pws[gid].i[ 1];
732 w0[2] = pws[gid].i[ 2];
733 w0[3] = pws[gid].i[ 3];
756 const u32 pw_len = pws[gid].pw_len;
762 m00140m (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);
765 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_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)
771 const u32 gid = get_global_id (0);
773 if (gid >= gid_max) return;
777 w0[0] = pws[gid].i[ 0];
778 w0[1] = pws[gid].i[ 1];
779 w0[2] = pws[gid].i[ 2];
780 w0[3] = pws[gid].i[ 3];
784 w1[0] = pws[gid].i[ 4];
785 w1[1] = pws[gid].i[ 5];
786 w1[2] = pws[gid].i[ 6];
787 w1[3] = pws[gid].i[ 7];
803 const u32 pw_len = pws[gid].pw_len;
809 m00140m (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);
812 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_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)
818 const u32 gid = get_global_id (0);
820 if (gid >= gid_max) return;
824 w0[0] = pws[gid].i[ 0];
825 w0[1] = pws[gid].i[ 1];
826 w0[2] = pws[gid].i[ 2];
827 w0[3] = pws[gid].i[ 3];
831 w1[0] = pws[gid].i[ 4];
832 w1[1] = pws[gid].i[ 5];
833 w1[2] = pws[gid].i[ 6];
834 w1[3] = pws[gid].i[ 7];
838 w2[0] = pws[gid].i[ 8];
839 w2[1] = pws[gid].i[ 9];
840 w2[2] = pws[gid].i[10];
841 w2[3] = pws[gid].i[11];
845 w3[0] = pws[gid].i[12];
846 w3[1] = pws[gid].i[13];
850 const u32 pw_len = pws[gid].pw_len;
856 m00140m (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);
859 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_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)
865 const u32 gid = get_global_id (0);
867 if (gid >= gid_max) return;
871 w0[0] = pws[gid].i[ 0];
872 w0[1] = pws[gid].i[ 1];
873 w0[2] = pws[gid].i[ 2];
874 w0[3] = pws[gid].i[ 3];
897 const u32 pw_len = pws[gid].pw_len;
903 m00140s (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);
906 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_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)
912 const u32 gid = get_global_id (0);
914 if (gid >= gid_max) return;
918 w0[0] = pws[gid].i[ 0];
919 w0[1] = pws[gid].i[ 1];
920 w0[2] = pws[gid].i[ 2];
921 w0[3] = pws[gid].i[ 3];
925 w1[0] = pws[gid].i[ 4];
926 w1[1] = pws[gid].i[ 5];
927 w1[2] = pws[gid].i[ 6];
928 w1[3] = pws[gid].i[ 7];
944 const u32 pw_len = pws[gid].pw_len;
950 m00140s (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);
953 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_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)
959 const u32 gid = get_global_id (0);
961 if (gid >= gid_max) return;
965 w0[0] = pws[gid].i[ 0];
966 w0[1] = pws[gid].i[ 1];
967 w0[2] = pws[gid].i[ 2];
968 w0[3] = pws[gid].i[ 3];
972 w1[0] = pws[gid].i[ 4];
973 w1[1] = pws[gid].i[ 5];
974 w1[2] = pws[gid].i[ 6];
975 w1[3] = pws[gid].i[ 7];
979 w2[0] = pws[gid].i[ 8];
980 w2[1] = pws[gid].i[ 9];
981 w2[2] = pws[gid].i[10];
982 w2[3] = pws[gid].i[11];
986 w3[0] = pws[gid].i[12];
987 w3[1] = pws[gid].i[13];
991 const u32 pw_len = pws[gid].pw_len;
997 m00140s (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);