2 * Author......: Jens Steube <jens.steube@gmail.com>
10 #include "include/constants.h"
11 #include "include/kernel_vendor.h"
18 #include "include/kernel_functions.c"
19 #include "OpenCL/types_ocl.c"
20 #include "OpenCL/common.c"
21 #include "OpenCL/simd.c"
23 __kernel void m01100_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
29 const u32 lid = get_local_id (0);
35 const u32 gid = get_global_id (0);
37 if (gid >= gid_max) return;
42 pw_buf0[0] = pws[gid].i[0];
43 pw_buf0[1] = pws[gid].i[1];
44 pw_buf0[2] = pws[gid].i[2];
45 pw_buf0[3] = pws[gid].i[3];
46 pw_buf1[0] = pws[gid].i[4];
47 pw_buf1[1] = pws[gid].i[5];
48 pw_buf1[2] = pws[gid].i[6];
49 pw_buf1[3] = pws[gid].i[7];
51 const u32 pw_l_len = pws[gid].pw_len;
61 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
62 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
63 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
64 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
65 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
66 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
67 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
68 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
69 salt_buf2[0] = salt_bufs[salt_pos].salt_buf[ 8];
70 salt_buf2[1] = salt_bufs[salt_pos].salt_buf[ 9];
71 salt_buf2[2] = salt_bufs[salt_pos].salt_buf[10];
72 salt_buf2[3] = salt_bufs[salt_pos].salt_buf[11];
74 const u32 salt_len = salt_bufs[salt_pos].salt_len;
80 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
82 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
84 const u32x pw_len = pw_l_len + pw_r_len;
87 * concat password candidate
90 u32x wordl0[4] = { 0 };
91 u32x wordl1[4] = { 0 };
92 u32x wordl2[4] = { 0 };
93 u32x wordl3[4] = { 0 };
95 wordl0[0] = pw_buf0[0];
96 wordl0[1] = pw_buf0[1];
97 wordl0[2] = pw_buf0[2];
98 wordl0[3] = pw_buf0[3];
99 wordl1[0] = pw_buf1[0];
100 wordl1[1] = pw_buf1[1];
101 wordl1[2] = pw_buf1[2];
102 wordl1[3] = pw_buf1[3];
104 u32x wordr0[4] = { 0 };
105 u32x wordr1[4] = { 0 };
106 u32x wordr2[4] = { 0 };
107 u32x wordr3[4] = { 0 };
109 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
110 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
111 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
112 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
113 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
114 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
115 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
116 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
118 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
120 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
124 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
132 w0[0] = wordl0[0] | wordr0[0];
133 w0[1] = wordl0[1] | wordr0[1];
134 w0[2] = wordl0[2] | wordr0[2];
135 w0[3] = wordl0[3] | wordr0[3];
136 w1[0] = wordl1[0] | wordr1[0];
137 w1[1] = wordl1[1] | wordr1[1];
138 w1[2] = wordl1[2] | wordr1[2];
139 w1[3] = wordl1[3] | wordr1[3];
140 w2[0] = wordl2[0] | wordr2[0];
141 w2[1] = wordl2[1] | wordr2[1];
142 w2[2] = wordl2[2] | wordr2[2];
143 w2[3] = wordl2[3] | wordr2[3];
144 w3[0] = wordl3[0] | wordr3[0];
145 w3[1] = wordl3[1] | wordr3[1];
146 w3[2] = wordl3[2] | wordr3[2];
147 w3[3] = wordl3[3] | wordr3[3];
149 make_unicode (w1, w2, w3);
150 make_unicode (w0, w0, w1);
152 w3[2] = pw_len * 2 * 8;
164 MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
165 MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
166 MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
167 MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
168 MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
169 MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
170 MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
171 MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
172 MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
173 MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
174 MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
175 MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
176 MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
177 MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
178 MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
179 MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
181 MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
182 MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
183 MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
184 MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
185 MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
186 MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
187 MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
188 MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
189 MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
190 MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
191 MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
192 MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
193 MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
194 MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
195 MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
196 MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
198 MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
199 MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
200 MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
201 MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
202 MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
203 MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
204 MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
205 MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
206 MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
207 MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
208 MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
209 MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
210 MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
211 MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
212 MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
213 MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
224 w1[0] = salt_buf0[0];
225 w1[1] = salt_buf0[1];
226 w1[2] = salt_buf0[2];
227 w1[3] = salt_buf0[3];
228 w2[0] = salt_buf1[0];
229 w2[1] = salt_buf1[1];
230 w2[2] = salt_buf1[2];
231 w2[3] = salt_buf1[3];
232 w3[0] = salt_buf2[0];
233 w3[1] = salt_buf2[1];
234 w3[2] = (16 + salt_len) * 8;
242 MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
243 MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
244 MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
245 MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
246 MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
247 MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
248 MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
249 MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
250 MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
251 MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
252 MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
253 MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
254 MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
255 MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
256 MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
257 MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
259 MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
260 MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
261 MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
262 MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
263 MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
264 MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
265 MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
266 MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
267 MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
268 MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
269 MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
270 MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
271 MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
272 MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
273 MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
274 MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
276 MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
277 MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
278 MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
279 MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
280 MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
281 MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
282 MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
283 MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
284 MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
285 MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
286 MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
287 MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
288 MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
289 MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
290 MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
291 MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
293 COMPARE_M_SIMD (a, d, c, b);
297 __kernel void m01100_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
301 __kernel void m01100_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
305 __kernel void m01100_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
311 const u32 lid = get_local_id (0);
317 const u32 gid = get_global_id (0);
319 if (gid >= gid_max) return;
324 pw_buf0[0] = pws[gid].i[0];
325 pw_buf0[1] = pws[gid].i[1];
326 pw_buf0[2] = pws[gid].i[2];
327 pw_buf0[3] = pws[gid].i[3];
328 pw_buf1[0] = pws[gid].i[4];
329 pw_buf1[1] = pws[gid].i[5];
330 pw_buf1[2] = pws[gid].i[6];
331 pw_buf1[3] = pws[gid].i[7];
333 const u32 pw_l_len = pws[gid].pw_len;
343 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
344 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
345 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
346 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
347 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
348 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
349 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
350 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
351 salt_buf2[0] = salt_bufs[salt_pos].salt_buf[ 8];
352 salt_buf2[1] = salt_bufs[salt_pos].salt_buf[ 9];
353 salt_buf2[2] = salt_bufs[salt_pos].salt_buf[10];
354 salt_buf2[3] = salt_bufs[salt_pos].salt_buf[11];
356 const u32 salt_len = salt_bufs[salt_pos].salt_len;
362 const u32 search[4] =
364 digests_buf[digests_offset].digest_buf[DGST_R0],
365 digests_buf[digests_offset].digest_buf[DGST_R1],
366 digests_buf[digests_offset].digest_buf[DGST_R2],
367 digests_buf[digests_offset].digest_buf[DGST_R3]
374 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
376 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
378 const u32x pw_len = pw_l_len + pw_r_len;
381 * concat password candidate
384 u32x wordl0[4] = { 0 };
385 u32x wordl1[4] = { 0 };
386 u32x wordl2[4] = { 0 };
387 u32x wordl3[4] = { 0 };
389 wordl0[0] = pw_buf0[0];
390 wordl0[1] = pw_buf0[1];
391 wordl0[2] = pw_buf0[2];
392 wordl0[3] = pw_buf0[3];
393 wordl1[0] = pw_buf1[0];
394 wordl1[1] = pw_buf1[1];
395 wordl1[2] = pw_buf1[2];
396 wordl1[3] = pw_buf1[3];
398 u32x wordr0[4] = { 0 };
399 u32x wordr1[4] = { 0 };
400 u32x wordr2[4] = { 0 };
401 u32x wordr3[4] = { 0 };
403 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
404 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
405 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
406 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
407 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
408 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
409 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
410 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
412 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
414 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
418 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
426 w0[0] = wordl0[0] | wordr0[0];
427 w0[1] = wordl0[1] | wordr0[1];
428 w0[2] = wordl0[2] | wordr0[2];
429 w0[3] = wordl0[3] | wordr0[3];
430 w1[0] = wordl1[0] | wordr1[0];
431 w1[1] = wordl1[1] | wordr1[1];
432 w1[2] = wordl1[2] | wordr1[2];
433 w1[3] = wordl1[3] | wordr1[3];
434 w2[0] = wordl2[0] | wordr2[0];
435 w2[1] = wordl2[1] | wordr2[1];
436 w2[2] = wordl2[2] | wordr2[2];
437 w2[3] = wordl2[3] | wordr2[3];
438 w3[0] = wordl3[0] | wordr3[0];
439 w3[1] = wordl3[1] | wordr3[1];
440 w3[2] = wordl3[2] | wordr3[2];
441 w3[3] = wordl3[3] | wordr3[3];
443 make_unicode (w1, w2, w3);
444 make_unicode (w0, w0, w1);
446 w3[2] = pw_len * 2 * 8;
458 MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
459 MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
460 MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
461 MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
462 MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
463 MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
464 MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
465 MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
466 MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
467 MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
468 MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
469 MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
470 MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
471 MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
472 MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
473 MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
475 MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
476 MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
477 MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
478 MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
479 MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
480 MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
481 MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
482 MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
483 MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
484 MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
485 MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
486 MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
487 MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
488 MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
489 MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
490 MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
492 MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
493 MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
494 MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
495 MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
496 MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
497 MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
498 MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
499 MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
500 MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
501 MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
502 MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
503 MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
504 MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
505 MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
506 MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
507 MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
518 w1[0] = salt_buf0[0];
519 w1[1] = salt_buf0[1];
520 w1[2] = salt_buf0[2];
521 w1[3] = salt_buf0[3];
522 w2[0] = salt_buf1[0];
523 w2[1] = salt_buf1[1];
524 w2[2] = salt_buf1[2];
525 w2[3] = salt_buf1[3];
526 w3[0] = salt_buf2[0];
527 w3[1] = salt_buf2[1];
528 w3[2] = (16 + salt_len) * 8;
536 MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
537 MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
538 MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
539 MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
540 MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
541 MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
542 MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
543 MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
544 MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
545 MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
546 MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
547 MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
548 MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
549 MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
550 MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
551 MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
553 MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
554 MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
555 MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
556 MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
557 MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
558 MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
559 MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
560 MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
561 MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
562 MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
563 MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
564 MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
565 MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
566 MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
567 MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
568 MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
570 MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
571 MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
572 MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
573 MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
574 MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
575 MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
576 MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
577 MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
578 MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
579 MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
580 MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
581 MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
582 MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
584 if (MATCHES_NONE_VS (a, search[0])) continue;
586 MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
587 MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
588 MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
590 COMPARE_S_SIMD (a, d, c, b);
594 __kernel void m01100_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
598 __kernel void m01100_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)