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 m00200m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32 * words_buf_r, __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)
29 const u32 gid = get_global_id (0);
30 const u32 lid = get_local_id (0);
38 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
40 const u32 w0r = words_buf_r[il_pos];
42 const u32 w0 = w0l | w0r;
51 a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
58 ROUND ((w0 >> 0) & 0xff);
59 ROUND ((w0 >> 8) & 0xff);
60 ROUND ((w0 >> 16) & 0xff);
61 ROUND ((w0 >> 24) & 0xff);
65 ROUND ((w0 >> 0) & 0xff);
66 ROUND ((w0 >> 8) & 0xff);
67 ROUND ((w0 >> 16) & 0xff);
71 ROUND ((w0 >> 0) & 0xff);
72 ROUND ((w0 >> 8) & 0xff);
76 ROUND ((w0 >> 0) & 0xff);
82 for (i = 4, j = 1; i <= (int) pw_len - 4; i += 4, j += 1)
86 ROUND ((wj >> 0) & 0xff);
87 ROUND ((wj >> 8) & 0xff);
88 ROUND ((wj >> 16) & 0xff);
89 ROUND ((wj >> 24) & 0xff);
94 const u32 left = pw_len - i;
98 ROUND ((wj >> 0) & 0xff);
99 ROUND ((wj >> 8) & 0xff);
100 ROUND ((wj >> 16) & 0xff);
104 ROUND ((wj >> 0) & 0xff);
105 ROUND ((wj >> 8) & 0xff);
109 ROUND ((wj >> 0) & 0xff);
124 static void m00200s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32 * words_buf_r, __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)
130 const u32 gid = get_global_id (0);
131 const u32 lid = get_local_id (0);
137 const u32 search[4] =
139 digests_buf[digests_offset].digest_buf[DGST_R0],
140 digests_buf[digests_offset].digest_buf[DGST_R1],
141 digests_buf[digests_offset].digest_buf[DGST_R2],
142 digests_buf[digests_offset].digest_buf[DGST_R3]
151 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
153 const u32 w0r = words_buf_r[il_pos];
155 const u32 w0 = w0l | w0r;
164 a ^= (((a & 0x3f) + add) * (v)) + (a << 8); \
171 ROUND ((w0 >> 0) & 0xff);
172 ROUND ((w0 >> 8) & 0xff);
173 ROUND ((w0 >> 16) & 0xff);
174 ROUND ((w0 >> 24) & 0xff);
176 else if (pw_len == 3)
178 ROUND ((w0 >> 0) & 0xff);
179 ROUND ((w0 >> 8) & 0xff);
180 ROUND ((w0 >> 16) & 0xff);
182 else if (pw_len == 2)
184 ROUND ((w0 >> 0) & 0xff);
185 ROUND ((w0 >> 8) & 0xff);
187 else if (pw_len == 1)
189 ROUND ((w0 >> 0) & 0xff);
195 for (i = 4, j = 1; i <= (int) pw_len - 4; i += 4, j += 1)
199 ROUND ((wj >> 0) & 0xff);
200 ROUND ((wj >> 8) & 0xff);
201 ROUND ((wj >> 16) & 0xff);
202 ROUND ((wj >> 24) & 0xff);
207 const u32 left = pw_len - i;
211 ROUND ((wj >> 0) & 0xff);
212 ROUND ((wj >> 8) & 0xff);
213 ROUND ((wj >> 16) & 0xff);
217 ROUND ((wj >> 0) & 0xff);
218 ROUND ((wj >> 8) & 0xff);
222 ROUND ((wj >> 0) & 0xff);
237 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00200_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32 * words_buf_r, __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)
243 const u32 gid = get_global_id (0);
245 if (gid >= gid_max) return;
249 w[ 0] = pws[gid].i[ 0];
250 w[ 1] = pws[gid].i[ 1];
251 w[ 2] = pws[gid].i[ 2];
252 w[ 3] = pws[gid].i[ 3];
266 const u32 pw_len = pws[gid].pw_len;
272 m00200m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
275 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00200_m08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32 * words_buf_r, __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)
281 const u32 gid = get_global_id (0);
283 if (gid >= gid_max) return;
287 w[ 0] = pws[gid].i[ 0];
288 w[ 1] = pws[gid].i[ 1];
289 w[ 2] = pws[gid].i[ 2];
290 w[ 3] = pws[gid].i[ 3];
291 w[ 4] = pws[gid].i[ 4];
292 w[ 5] = pws[gid].i[ 5];
293 w[ 6] = pws[gid].i[ 6];
294 w[ 7] = pws[gid].i[ 7];
304 const u32 pw_len = pws[gid].pw_len;
310 m00200m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
313 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00200_m16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32 * words_buf_r, __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)
319 const u32 gid = get_global_id (0);
321 if (gid >= gid_max) return;
325 w[ 0] = pws[gid].i[ 0];
326 w[ 1] = pws[gid].i[ 1];
327 w[ 2] = pws[gid].i[ 2];
328 w[ 3] = pws[gid].i[ 3];
329 w[ 4] = pws[gid].i[ 4];
330 w[ 5] = pws[gid].i[ 5];
331 w[ 6] = pws[gid].i[ 6];
332 w[ 7] = pws[gid].i[ 7];
333 w[ 8] = pws[gid].i[ 8];
334 w[ 9] = pws[gid].i[ 9];
335 w[10] = pws[gid].i[10];
336 w[11] = pws[gid].i[11];
337 w[12] = pws[gid].i[12];
338 w[13] = pws[gid].i[13];
339 w[14] = pws[gid].i[14];
340 w[15] = pws[gid].i[15];
342 const u32 pw_len = pws[gid].pw_len;
348 m00200m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
351 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00200_s04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32 * words_buf_r, __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)
357 const u32 gid = get_global_id (0);
359 if (gid >= gid_max) return;
363 w[ 0] = pws[gid].i[ 0];
364 w[ 1] = pws[gid].i[ 1];
365 w[ 2] = pws[gid].i[ 2];
366 w[ 3] = pws[gid].i[ 3];
380 const u32 pw_len = pws[gid].pw_len;
386 m00200s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
389 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00200_s08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32 * words_buf_r, __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)
395 const u32 gid = get_global_id (0);
397 if (gid >= gid_max) return;
401 w[ 0] = pws[gid].i[ 0];
402 w[ 1] = pws[gid].i[ 1];
403 w[ 2] = pws[gid].i[ 2];
404 w[ 3] = pws[gid].i[ 3];
405 w[ 4] = pws[gid].i[ 4];
406 w[ 5] = pws[gid].i[ 5];
407 w[ 6] = pws[gid].i[ 6];
408 w[ 7] = pws[gid].i[ 7];
418 const u32 pw_len = pws[gid].pw_len;
424 m00200s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
427 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00200_s16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32 * words_buf_r, __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)
433 const u32 gid = get_global_id (0);
435 if (gid >= gid_max) return;
439 w[ 0] = pws[gid].i[ 0];
440 w[ 1] = pws[gid].i[ 1];
441 w[ 2] = pws[gid].i[ 2];
442 w[ 3] = pws[gid].i[ 3];
443 w[ 4] = pws[gid].i[ 4];
444 w[ 5] = pws[gid].i[ 5];
445 w[ 6] = pws[gid].i[ 6];
446 w[ 7] = pws[gid].i[ 7];
447 w[ 8] = pws[gid].i[ 8];
448 w[ 9] = pws[gid].i[ 9];
449 w[10] = pws[gid].i[10];
450 w[11] = pws[gid].i[11];
451 w[12] = pws[gid].i[12];
452 w[13] = pws[gid].i[13];
453 w[14] = pws[gid].i[14];
454 w[15] = pws[gid].i[15];
456 const u32 pw_len = pws[gid].pw_len;
462 m00200s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);