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"
31 #include "include/rp_gpu.h"
35 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
36 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
40 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
41 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
45 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
46 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
50 #define uint_to_hex_lower8(i) l_bin2asc[(i)]
54 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
58 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
61 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
67 const u32 lid = get_local_id (0);
73 const u32 gid = get_global_id (0);
77 pw_buf0[0] = pws[gid].i[ 0];
78 pw_buf0[1] = pws[gid].i[ 1];
79 pw_buf0[2] = pws[gid].i[ 2];
80 pw_buf0[3] = pws[gid].i[ 3];
84 pw_buf1[0] = pws[gid].i[ 4];
85 pw_buf1[1] = pws[gid].i[ 5];
86 pw_buf1[2] = pws[gid].i[ 6];
87 pw_buf1[3] = pws[gid].i[ 7];
89 const u32 pw_len = pws[gid].pw_len;
97 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
98 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
99 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
100 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
104 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
105 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
106 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
107 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
123 const u32 salt_len = salt_bufs[salt_pos].salt_len;
125 const u32 pw_salt_len = 32 + salt_len;
131 __local u32 l_bin2asc[256];
133 const u32 lid4 = lid * 4;
135 const u32 lid40 = lid4 + 0;
136 const u32 lid41 = lid4 + 1;
137 const u32 lid42 = lid4 + 2;
138 const u32 lid43 = lid4 + 3;
140 const u32 v400 = (lid40 >> 0) & 15;
141 const u32 v401 = (lid40 >> 4) & 15;
142 const u32 v410 = (lid41 >> 0) & 15;
143 const u32 v411 = (lid41 >> 4) & 15;
144 const u32 v420 = (lid42 >> 0) & 15;
145 const u32 v421 = (lid42 >> 4) & 15;
146 const u32 v430 = (lid43 >> 0) & 15;
147 const u32 v431 = (lid43 >> 4) & 15;
149 l_bin2asc[lid40] = ((v400 < 10) ? '0' + v400 : 'a' - 10 + v400) << 8
150 | ((v401 < 10) ? '0' + v401 : 'a' - 10 + v401) << 0;
151 l_bin2asc[lid41] = ((v410 < 10) ? '0' + v410 : 'a' - 10 + v410) << 8
152 | ((v411 < 10) ? '0' + v411 : 'a' - 10 + v411) << 0;
153 l_bin2asc[lid42] = ((v420 < 10) ? '0' + v420 : 'a' - 10 + v420) << 8
154 | ((v421 < 10) ? '0' + v421 : 'a' - 10 + v421) << 0;
155 l_bin2asc[lid43] = ((v430 < 10) ? '0' + v430 : 'a' - 10 + v430) << 8
156 | ((v431 < 10) ? '0' + v431 : 'a' - 10 + v431) << 0;
158 barrier (CLK_LOCAL_MEM_FENCE);
160 if (gid >= gid_max) return;
166 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
196 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
198 append_0x80_2 (w0, w1, out_len);
207 MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
208 MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
209 MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
210 MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
211 MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
212 MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
213 MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
214 MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
215 MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
216 MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
217 MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
218 MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
219 MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
220 MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
221 MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
222 MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
224 MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
225 MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
226 MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
227 MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
228 MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
229 MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
230 MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
231 MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
232 MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
233 MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
234 MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
235 MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
236 MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
237 MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
238 MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
239 MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
241 MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20);
242 MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21);
243 MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22);
244 MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23);
245 MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20);
246 MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21);
247 MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22);
248 MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23);
249 MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20);
250 MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21);
251 MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22);
252 MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23);
253 MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20);
254 MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21);
255 MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22);
256 MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23);
258 MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
259 MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
260 MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
261 MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
262 MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
263 MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
264 MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
265 MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
266 MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
267 MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
268 MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
269 MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
270 MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
271 MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
272 MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
273 MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
285 w0_t[0] = uint_to_hex_lower8 ((a >> 0) & 255) << 0
286 | uint_to_hex_lower8 ((a >> 8) & 255) << 16;
287 w0_t[1] = uint_to_hex_lower8 ((a >> 16) & 255) << 0
288 | uint_to_hex_lower8 ((a >> 24) & 255) << 16;
289 w0_t[2] = uint_to_hex_lower8 ((b >> 0) & 255) << 0
290 | uint_to_hex_lower8 ((b >> 8) & 255) << 16;
291 w0_t[3] = uint_to_hex_lower8 ((b >> 16) & 255) << 0
292 | uint_to_hex_lower8 ((b >> 24) & 255) << 16;
293 w1_t[0] = uint_to_hex_lower8 ((c >> 0) & 255) << 0
294 | uint_to_hex_lower8 ((c >> 8) & 255) << 16;
295 w1_t[1] = uint_to_hex_lower8 ((c >> 16) & 255) << 0
296 | uint_to_hex_lower8 ((c >> 24) & 255) << 16;
297 w1_t[2] = uint_to_hex_lower8 ((d >> 0) & 255) << 0
298 | uint_to_hex_lower8 ((d >> 8) & 255) << 16;
299 w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0
300 | uint_to_hex_lower8 ((d >> 24) & 255) << 16;
302 w2_t[0] = 0x00000080;
316 switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
318 w3_t[2] = pw_salt_len * 8;
320 w0_t[0] |= salt_buf0[0];
321 w0_t[1] |= salt_buf0[1];
322 w0_t[2] |= salt_buf0[2];
323 w0_t[3] |= salt_buf0[3];
324 w1_t[0] |= salt_buf1[0];
325 w1_t[1] |= salt_buf1[1];
326 w1_t[2] |= salt_buf1[2];
327 w1_t[3] |= salt_buf1[3];
328 w2_t[0] |= salt_buf2[0];
329 w2_t[1] |= salt_buf2[1];
330 w2_t[2] |= salt_buf2[2];
331 w2_t[3] |= salt_buf2[3];
332 w3_t[0] |= salt_buf3[0];
333 w3_t[1] |= salt_buf3[1];
334 w3_t[2] |= salt_buf3[2];
335 w3_t[3] |= salt_buf3[3];
346 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
347 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
348 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
349 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
350 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
351 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
352 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
353 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
354 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
355 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
356 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
357 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
358 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
359 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
360 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
361 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
363 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
364 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
365 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
366 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
367 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
368 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
369 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
370 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
371 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
372 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
373 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
374 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
375 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
376 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
377 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
378 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
380 MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
381 MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
382 MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
383 MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
384 MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
385 MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
386 MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
387 MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
388 MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
389 MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
390 MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
391 MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
392 MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
393 MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
394 MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
395 MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
397 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
398 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
399 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
400 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
401 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
402 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
403 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
404 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
405 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
406 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
407 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
408 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
409 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
410 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
411 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
412 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
419 #include VECT_COMPARE_M
423 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
427 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
431 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
437 const u32 lid = get_local_id (0);
443 const u32 gid = get_global_id (0);
447 pw_buf0[0] = pws[gid].i[ 0];
448 pw_buf0[1] = pws[gid].i[ 1];
449 pw_buf0[2] = pws[gid].i[ 2];
450 pw_buf0[3] = pws[gid].i[ 3];
454 pw_buf1[0] = pws[gid].i[ 4];
455 pw_buf1[1] = pws[gid].i[ 5];
456 pw_buf1[2] = pws[gid].i[ 6];
457 pw_buf1[3] = pws[gid].i[ 7];
459 const u32 pw_len = pws[gid].pw_len;
467 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
468 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
469 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
470 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
474 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
475 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
476 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
477 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
493 const u32 salt_len = salt_bufs[salt_pos].salt_len;
495 const u32 pw_salt_len = 32 + salt_len;
501 const u32 search[4] =
503 digests_buf[digests_offset].digest_buf[DGST_R0],
504 digests_buf[digests_offset].digest_buf[DGST_R1],
505 digests_buf[digests_offset].digest_buf[DGST_R2],
506 digests_buf[digests_offset].digest_buf[DGST_R3]
513 __local u32 l_bin2asc[256];
515 const u32 lid4 = lid * 4;
517 const u32 lid40 = lid4 + 0;
518 const u32 lid41 = lid4 + 1;
519 const u32 lid42 = lid4 + 2;
520 const u32 lid43 = lid4 + 3;
522 const u32 v400 = (lid40 >> 0) & 15;
523 const u32 v401 = (lid40 >> 4) & 15;
524 const u32 v410 = (lid41 >> 0) & 15;
525 const u32 v411 = (lid41 >> 4) & 15;
526 const u32 v420 = (lid42 >> 0) & 15;
527 const u32 v421 = (lid42 >> 4) & 15;
528 const u32 v430 = (lid43 >> 0) & 15;
529 const u32 v431 = (lid43 >> 4) & 15;
531 l_bin2asc[lid40] = ((v400 < 10) ? '0' + v400 : 'a' - 10 + v400) << 8
532 | ((v401 < 10) ? '0' + v401 : 'a' - 10 + v401) << 0;
533 l_bin2asc[lid41] = ((v410 < 10) ? '0' + v410 : 'a' - 10 + v410) << 8
534 | ((v411 < 10) ? '0' + v411 : 'a' - 10 + v411) << 0;
535 l_bin2asc[lid42] = ((v420 < 10) ? '0' + v420 : 'a' - 10 + v420) << 8
536 | ((v421 < 10) ? '0' + v421 : 'a' - 10 + v421) << 0;
537 l_bin2asc[lid43] = ((v430 < 10) ? '0' + v430 : 'a' - 10 + v430) << 8
538 | ((v431 < 10) ? '0' + v431 : 'a' - 10 + v431) << 0;
540 barrier (CLK_LOCAL_MEM_FENCE);
542 if (gid >= gid_max) return;
548 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
578 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
580 append_0x80_2 (w0, w1, out_len);
589 MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
590 MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
591 MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
592 MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
593 MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
594 MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
595 MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
596 MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
597 MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
598 MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
599 MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
600 MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
601 MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
602 MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
603 MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
604 MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
606 MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
607 MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
608 MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
609 MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
610 MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
611 MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
612 MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
613 MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
614 MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
615 MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
616 MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
617 MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
618 MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
619 MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
620 MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
621 MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
623 MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20);
624 MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21);
625 MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22);
626 MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23);
627 MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20);
628 MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21);
629 MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22);
630 MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23);
631 MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20);
632 MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21);
633 MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22);
634 MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23);
635 MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20);
636 MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21);
637 MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22);
638 MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23);
640 MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
641 MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
642 MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
643 MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
644 MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
645 MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
646 MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
647 MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
648 MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
649 MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
650 MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
651 MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
652 MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
653 MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
654 MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
655 MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
667 w0_t[0] = uint_to_hex_lower8 ((a >> 0) & 255) << 0
668 | uint_to_hex_lower8 ((a >> 8) & 255) << 16;
669 w0_t[1] = uint_to_hex_lower8 ((a >> 16) & 255) << 0
670 | uint_to_hex_lower8 ((a >> 24) & 255) << 16;
671 w0_t[2] = uint_to_hex_lower8 ((b >> 0) & 255) << 0
672 | uint_to_hex_lower8 ((b >> 8) & 255) << 16;
673 w0_t[3] = uint_to_hex_lower8 ((b >> 16) & 255) << 0
674 | uint_to_hex_lower8 ((b >> 24) & 255) << 16;
675 w1_t[0] = uint_to_hex_lower8 ((c >> 0) & 255) << 0
676 | uint_to_hex_lower8 ((c >> 8) & 255) << 16;
677 w1_t[1] = uint_to_hex_lower8 ((c >> 16) & 255) << 0
678 | uint_to_hex_lower8 ((c >> 24) & 255) << 16;
679 w1_t[2] = uint_to_hex_lower8 ((d >> 0) & 255) << 0
680 | uint_to_hex_lower8 ((d >> 8) & 255) << 16;
681 w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0
682 | uint_to_hex_lower8 ((d >> 24) & 255) << 16;
684 w2_t[0] = 0x00000080;
698 switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
700 w3_t[2] = pw_salt_len * 8;
702 w0_t[0] |= salt_buf0[0];
703 w0_t[1] |= salt_buf0[1];
704 w0_t[2] |= salt_buf0[2];
705 w0_t[3] |= salt_buf0[3];
706 w1_t[0] |= salt_buf1[0];
707 w1_t[1] |= salt_buf1[1];
708 w1_t[2] |= salt_buf1[2];
709 w1_t[3] |= salt_buf1[3];
710 w2_t[0] |= salt_buf2[0];
711 w2_t[1] |= salt_buf2[1];
712 w2_t[2] |= salt_buf2[2];
713 w2_t[3] |= salt_buf2[3];
714 w3_t[0] |= salt_buf3[0];
715 w3_t[1] |= salt_buf3[1];
716 w3_t[2] |= salt_buf3[2];
717 w3_t[3] |= salt_buf3[3];
728 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
729 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
730 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
731 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
732 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
733 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
734 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
735 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
736 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
737 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
738 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
739 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
740 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
741 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
742 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
743 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
745 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
746 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
747 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
748 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
749 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
750 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
751 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
752 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
753 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
754 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
755 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
756 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
757 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
758 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
759 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
760 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
762 MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
763 MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
764 MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
765 MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
766 MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
767 MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
768 MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
769 MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
770 MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
771 MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
772 MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
773 MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
774 MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
775 MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
776 MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
777 MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
779 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
780 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
781 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
782 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
783 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
784 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
785 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
786 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
787 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
788 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
789 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
790 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
791 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
792 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
793 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
794 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
801 #include VECT_COMPARE_S
805 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
809 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)