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"
19 #include "include/rp_kernel.h"
20 #include "OpenCL/rp.c"
22 #define COMPARE_S "OpenCL/check_single_comp4.c"
23 #define COMPARE_M "OpenCL/check_multi_comp4.c"
25 __kernel void m04900_m04 (__global pw_t *pws, __global kernel_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
31 const u32 lid = get_local_id (0);
37 const u32 gid = get_global_id (0);
39 if (gid >= gid_max) return;
43 pw_buf0[0] = pws[gid].i[ 0];
44 pw_buf0[1] = pws[gid].i[ 1];
45 pw_buf0[2] = pws[gid].i[ 2];
46 pw_buf0[3] = pws[gid].i[ 3];
50 pw_buf1[0] = pws[gid].i[ 4];
51 pw_buf1[1] = pws[gid].i[ 5];
52 pw_buf1[2] = pws[gid].i[ 6];
53 pw_buf1[3] = pws[gid].i[ 7];
55 const u32 pw_len = pws[gid].pw_len;
63 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
64 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
65 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
66 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
70 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
71 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
72 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
73 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
89 const u32 salt_len = salt_bufs[salt_pos].salt_len;
95 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
100 w0_t[1] = pw_buf0[1];
101 w0_t[2] = pw_buf0[2];
102 w0_t[3] = pw_buf0[3];
106 w1_t[0] = pw_buf1[0];
107 w1_t[1] = pw_buf1[1];
108 w1_t[2] = pw_buf1[2];
109 w1_t[3] = pw_buf1[3];
125 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0_t, w1_t, pw_len);
131 switch_buffer_by_offset_le (w0_t, w1_t, w2_t, w3_t, salt_len);
133 w0_t[0] |= salt_buf0[0];
134 w0_t[1] |= salt_buf0[1];
135 w0_t[2] |= salt_buf0[2];
136 w0_t[3] |= salt_buf0[3];
137 w1_t[0] |= salt_buf1[0];
138 w1_t[1] |= salt_buf1[1];
139 w1_t[2] |= salt_buf1[2];
140 w1_t[3] |= salt_buf1[3];
141 w2_t[0] |= salt_buf2[0];
142 w2_t[1] |= salt_buf2[1];
143 w2_t[2] |= salt_buf2[2];
144 w2_t[3] |= salt_buf2[3];
145 w3_t[0] |= salt_buf3[0];
146 w3_t[1] |= salt_buf3[1];
147 w3_t[2] |= salt_buf3[2];
148 w3_t[3] |= salt_buf3[3];
156 s0[0] = salt_buf0[0];
157 s0[1] = salt_buf0[1];
158 s0[2] = salt_buf0[2];
159 s0[3] = salt_buf0[3];
163 s1[0] = salt_buf1[0];
164 s1[1] = salt_buf1[1];
165 s1[2] = salt_buf1[2];
166 s1[3] = salt_buf1[3];
182 switch_buffer_by_offset_le (s0, s1, s2, s3, salt_len + out_len);
201 const u32 pw_salt_len = salt_len + out_len + salt_len;
203 append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
205 u32 w0 = swap32 (w0_t[0]);
206 u32 w1 = swap32 (w0_t[1]);
207 u32 w2 = swap32 (w0_t[2]);
208 u32 w3 = swap32 (w0_t[3]);
209 u32 w4 = swap32 (w1_t[0]);
210 u32 w5 = swap32 (w1_t[1]);
211 u32 w6 = swap32 (w1_t[2]);
212 u32 w7 = swap32 (w1_t[3]);
213 u32 w8 = swap32 (w2_t[0]);
214 u32 w9 = swap32 (w2_t[1]);
215 u32 wa = swap32 (w2_t[2]);
216 u32 wb = swap32 (w2_t[3]);
217 u32 wc = swap32 (w3_t[0]);
218 u32 wd = swap32 (w3_t[1]);
220 u32 wf = pw_salt_len * 8;
235 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0);
236 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1);
237 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2);
238 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3);
239 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4);
240 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5);
241 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6);
242 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7);
243 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8);
244 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9);
245 SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa);
246 SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb);
247 SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc);
248 SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd);
249 SHA1_STEP (SHA1_F0o, b, c, d, e, a, we);
250 SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf);
251 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0);
252 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1);
253 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2);
254 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3);
259 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4);
260 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5);
261 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6);
262 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7);
263 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8);
264 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9);
265 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa);
266 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb);
267 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc);
268 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd);
269 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we);
270 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf);
271 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0);
272 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1);
273 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2);
274 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3);
275 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4);
276 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5);
277 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6);
278 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7);
283 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8);
284 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9);
285 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa);
286 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb);
287 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc);
288 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd);
289 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we);
290 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf);
291 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0);
292 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1);
293 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2);
294 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3);
295 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4);
296 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5);
297 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6);
298 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7);
299 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8);
300 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9);
301 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa);
302 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb);
307 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc);
308 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd);
309 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we);
310 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf);
311 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0);
312 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1);
313 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2);
314 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3);
315 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4);
316 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5);
317 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6);
318 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7);
319 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8);
320 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9);
321 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa);
322 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb);
323 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc);
324 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd);
325 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we);
326 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf);
337 __kernel void m04900_m08 (__global pw_t *pws, __global kernel_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
341 __kernel void m04900_m16 (__global pw_t *pws, __global kernel_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
345 __kernel void m04900_s04 (__global pw_t *pws, __global kernel_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
351 const u32 lid = get_local_id (0);
357 const u32 gid = get_global_id (0);
359 if (gid >= gid_max) return;
363 pw_buf0[0] = pws[gid].i[ 0];
364 pw_buf0[1] = pws[gid].i[ 1];
365 pw_buf0[2] = pws[gid].i[ 2];
366 pw_buf0[3] = pws[gid].i[ 3];
370 pw_buf1[0] = pws[gid].i[ 4];
371 pw_buf1[1] = pws[gid].i[ 5];
372 pw_buf1[2] = pws[gid].i[ 6];
373 pw_buf1[3] = pws[gid].i[ 7];
375 const u32 pw_len = pws[gid].pw_len;
383 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
384 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
385 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
386 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
390 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
391 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
392 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
393 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
409 const u32 salt_len = salt_bufs[salt_pos].salt_len;
415 const u32 search[4] =
417 digests_buf[digests_offset].digest_buf[DGST_R0],
418 digests_buf[digests_offset].digest_buf[DGST_R1],
419 digests_buf[digests_offset].digest_buf[DGST_R2],
420 digests_buf[digests_offset].digest_buf[DGST_R3]
427 const u32 e_rev = rotl32 (search[1], 2u);
433 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
437 w0_t[0] = pw_buf0[0];
438 w0_t[1] = pw_buf0[1];
439 w0_t[2] = pw_buf0[2];
440 w0_t[3] = pw_buf0[3];
444 w1_t[0] = pw_buf1[0];
445 w1_t[1] = pw_buf1[1];
446 w1_t[2] = pw_buf1[2];
447 w1_t[3] = pw_buf1[3];
463 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0_t, w1_t, pw_len);
469 switch_buffer_by_offset_le (w0_t, w1_t, w2_t, w3_t, salt_len);
471 w0_t[0] |= salt_buf0[0];
472 w0_t[1] |= salt_buf0[1];
473 w0_t[2] |= salt_buf0[2];
474 w0_t[3] |= salt_buf0[3];
475 w1_t[0] |= salt_buf1[0];
476 w1_t[1] |= salt_buf1[1];
477 w1_t[2] |= salt_buf1[2];
478 w1_t[3] |= salt_buf1[3];
479 w2_t[0] |= salt_buf2[0];
480 w2_t[1] |= salt_buf2[1];
481 w2_t[2] |= salt_buf2[2];
482 w2_t[3] |= salt_buf2[3];
483 w3_t[0] |= salt_buf3[0];
484 w3_t[1] |= salt_buf3[1];
485 w3_t[2] |= salt_buf3[2];
486 w3_t[3] |= salt_buf3[3];
494 s0[0] = salt_buf0[0];
495 s0[1] = salt_buf0[1];
496 s0[2] = salt_buf0[2];
497 s0[3] = salt_buf0[3];
501 s1[0] = salt_buf1[0];
502 s1[1] = salt_buf1[1];
503 s1[2] = salt_buf1[2];
504 s1[3] = salt_buf1[3];
520 switch_buffer_by_offset_le (s0, s1, s2, s3, salt_len + out_len);
539 const u32 pw_salt_len = salt_len + out_len + salt_len;
541 append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
543 u32 w0 = swap32 (w0_t[0]);
544 u32 w1 = swap32 (w0_t[1]);
545 u32 w2 = swap32 (w0_t[2]);
546 u32 w3 = swap32 (w0_t[3]);
547 u32 w4 = swap32 (w1_t[0]);
548 u32 w5 = swap32 (w1_t[1]);
549 u32 w6 = swap32 (w1_t[2]);
550 u32 w7 = swap32 (w1_t[3]);
551 u32 w8 = swap32 (w2_t[0]);
552 u32 w9 = swap32 (w2_t[1]);
553 u32 wa = swap32 (w2_t[2]);
554 u32 wb = swap32 (w2_t[3]);
555 u32 wc = swap32 (w3_t[0]);
556 u32 wd = swap32 (w3_t[1]);
558 u32 wf = pw_salt_len * 8;
573 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0);
574 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1);
575 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2);
576 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3);
577 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4);
578 SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5);
579 SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6);
580 SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7);
581 SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8);
582 SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9);
583 SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa);
584 SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb);
585 SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc);
586 SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd);
587 SHA1_STEP (SHA1_F0o, b, c, d, e, a, we);
588 SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf);
589 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0);
590 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1);
591 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2);
592 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3);
597 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4);
598 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5);
599 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6);
600 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7);
601 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8);
602 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9);
603 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa);
604 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb);
605 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc);
606 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd);
607 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we);
608 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf);
609 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0);
610 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1);
611 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2);
612 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3);
613 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4);
614 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5);
615 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6);
616 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7);
621 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8);
622 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9);
623 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa);
624 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb);
625 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc);
626 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd);
627 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we);
628 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf);
629 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0);
630 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1);
631 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2);
632 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3);
633 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4);
634 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5);
635 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6);
636 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7);
637 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8);
638 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9);
639 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa);
640 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb);
645 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc);
646 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd);
647 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we);
648 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf);
649 w0 = rotl32 ((wd ^ w8 ^ w2 ^ w0), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0);
650 w1 = rotl32 ((we ^ w9 ^ w3 ^ w1), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1);
651 w2 = rotl32 ((wf ^ wa ^ w4 ^ w2), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2);
652 w3 = rotl32 ((w0 ^ wb ^ w5 ^ w3), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3);
653 w4 = rotl32 ((w1 ^ wc ^ w6 ^ w4), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4);
654 w5 = rotl32 ((w2 ^ wd ^ w7 ^ w5), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5);
655 w6 = rotl32 ((w3 ^ we ^ w8 ^ w6), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6);
656 w7 = rotl32 ((w4 ^ wf ^ w9 ^ w7), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7);
657 w8 = rotl32 ((w5 ^ w0 ^ wa ^ w8), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8);
658 w9 = rotl32 ((w6 ^ w1 ^ wb ^ w9), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9);
659 wa = rotl32 ((w7 ^ w2 ^ wc ^ wa), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa);
660 wb = rotl32 ((w8 ^ w3 ^ wd ^ wb), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb);
662 if (allx (e != e_rev)) continue;
664 wc = rotl32 ((w9 ^ w4 ^ we ^ wc), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc);
665 wd = rotl32 ((wa ^ w5 ^ wf ^ wd), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd);
666 we = rotl32 ((wb ^ w6 ^ w0 ^ we), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we);
667 wf = rotl32 ((wc ^ w7 ^ w1 ^ wf), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf);
678 __kernel void m04900_s08 (__global pw_t *pws, __global kernel_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
682 __kernel void m04900_s16 (__global pw_t *pws, __global kernel_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)