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 "types_ocl.c"
19 #include "include/rp_gpu.h"
22 #define COMPARE_S "check_single_comp4.c"
23 #define COMPARE_M "check_multi_comp4.c"
26 #define uint_to_hex_lower8(i) l_bin2asc[(i)]
30 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
34 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
37 static u32 memcat32 (u32 block0[16], u32 block1[16], const u32 block_len, const u32 append0[4], const u32 append1[4], const u32 append2[4], const u32 append3[4], const u32 append_len)
39 const u32 mod = block_len & 3;
40 const u32 div = block_len / 4;
42 const int offset_minus_4 = 4 - mod;
46 append0_t[0] = amd_bytealign (append0[0], 0, offset_minus_4);
47 append0_t[1] = amd_bytealign (append0[1], append0[0], offset_minus_4);
48 append0_t[2] = amd_bytealign (append0[2], append0[1], offset_minus_4);
49 append0_t[3] = amd_bytealign (append0[3], append0[2], offset_minus_4);
53 append1_t[0] = amd_bytealign (append1[0], append0[3], offset_minus_4);
54 append1_t[1] = amd_bytealign (append1[1], append1[0], offset_minus_4);
55 append1_t[2] = amd_bytealign (append1[2], append1[1], offset_minus_4);
56 append1_t[3] = amd_bytealign (append1[3], append1[2], offset_minus_4);
60 append2_t[0] = amd_bytealign (append2[0], append1[3], offset_minus_4);
61 append2_t[1] = amd_bytealign (append2[1], append2[0], offset_minus_4);
62 append2_t[2] = amd_bytealign (append2[2], append2[1], offset_minus_4);
63 append2_t[3] = amd_bytealign (append2[3], append2[2], offset_minus_4);
67 append3_t[0] = amd_bytealign (append3[0], append2[3], offset_minus_4);
68 append3_t[1] = amd_bytealign (append3[1], append3[0], offset_minus_4);
69 append3_t[2] = amd_bytealign (append3[2], append3[1], offset_minus_4);
70 append3_t[3] = amd_bytealign (append3[3], append3[2], offset_minus_4);
74 append4_t[0] = amd_bytealign ( 0, append3[3], offset_minus_4);
81 append0_t[0] = append0[0];
82 append0_t[1] = append0[1];
83 append0_t[2] = append0[2];
84 append0_t[3] = append0[3];
86 append1_t[0] = append1[0];
87 append1_t[1] = append1[1];
88 append1_t[2] = append1[2];
89 append1_t[3] = append1[3];
91 append2_t[0] = append2[0];
92 append2_t[1] = append2[1];
93 append2_t[2] = append2[2];
94 append2_t[3] = append2[3];
96 append3_t[0] = append3[0];
97 append3_t[1] = append3[1];
98 append3_t[2] = append3[2];
99 append3_t[3] = append3[3];
109 case 0: block0[ 0] |= append0_t[0];
110 block0[ 1] = append0_t[1];
111 block0[ 2] = append0_t[2];
112 block0[ 3] = append0_t[3];
114 block0[ 4] = append1_t[0];
115 block0[ 5] = append1_t[1];
116 block0[ 6] = append1_t[2];
117 block0[ 7] = append1_t[3];
119 block0[ 8] = append2_t[0];
120 block0[ 9] = append2_t[1];
121 block0[10] = append2_t[2];
122 block0[11] = append2_t[3];
124 block0[12] = append3_t[0];
125 block0[13] = append3_t[1];
126 block0[14] = append3_t[2];
127 block0[15] = append3_t[3];
129 block1[ 0] = append4_t[0];
130 block1[ 1] = append4_t[1];
131 block1[ 2] = append4_t[2];
132 block1[ 3] = append4_t[3];
135 case 1: block0[ 1] |= append0_t[0];
136 block0[ 2] = append0_t[1];
137 block0[ 3] = append0_t[2];
138 block0[ 4] = append0_t[3];
140 block0[ 5] = append1_t[0];
141 block0[ 6] = append1_t[1];
142 block0[ 7] = append1_t[2];
143 block0[ 8] = append1_t[3];
145 block0[ 9] = append2_t[0];
146 block0[10] = append2_t[1];
147 block0[11] = append2_t[2];
148 block0[12] = append2_t[3];
150 block0[13] = append3_t[0];
151 block0[14] = append3_t[1];
152 block0[15] = append3_t[2];
153 block1[ 0] = append3_t[3];
155 block1[ 1] = append4_t[0];
156 block1[ 2] = append4_t[1];
157 block1[ 3] = append4_t[2];
158 block1[ 4] = append4_t[3];
161 case 2: block0[ 2] |= append0_t[0];
162 block0[ 3] = append0_t[1];
163 block0[ 4] = append0_t[2];
164 block0[ 5] = append0_t[3];
166 block0[ 6] = append1_t[0];
167 block0[ 7] = append1_t[1];
168 block0[ 8] = append1_t[2];
169 block0[ 9] = append1_t[3];
171 block0[10] = append2_t[0];
172 block0[11] = append2_t[1];
173 block0[12] = append2_t[2];
174 block0[13] = append2_t[3];
176 block0[14] = append3_t[0];
177 block0[15] = append3_t[1];
178 block1[ 0] = append3_t[2];
179 block1[ 1] = append3_t[3];
181 block1[ 2] = append4_t[0];
182 block1[ 3] = append4_t[1];
183 block1[ 4] = append4_t[2];
184 block1[ 5] = append4_t[3];
187 case 3: block0[ 3] |= append0_t[0];
188 block0[ 4] = append0_t[1];
189 block0[ 5] = append0_t[2];
190 block0[ 6] = append0_t[3];
192 block0[ 7] = append1_t[0];
193 block0[ 8] = append1_t[1];
194 block0[ 9] = append1_t[2];
195 block0[10] = append1_t[3];
197 block0[11] = append2_t[0];
198 block0[12] = append2_t[1];
199 block0[13] = append2_t[2];
200 block0[14] = append2_t[3];
202 block0[15] = append3_t[0];
203 block1[ 0] = append3_t[1];
204 block1[ 1] = append3_t[2];
205 block1[ 2] = append3_t[3];
207 block1[ 3] = append4_t[0];
208 block1[ 4] = append4_t[1];
209 block1[ 5] = append4_t[2];
210 block1[ 6] = append4_t[3];
213 case 4: block0[ 4] |= append0_t[0];
214 block0[ 5] = append0_t[1];
215 block0[ 6] = append0_t[2];
216 block0[ 7] = append0_t[3];
218 block0[ 8] = append1_t[0];
219 block0[ 9] = append1_t[1];
220 block0[10] = append1_t[2];
221 block0[11] = append1_t[3];
223 block0[12] = append2_t[0];
224 block0[13] = append2_t[1];
225 block0[14] = append2_t[2];
226 block0[15] = append2_t[3];
228 block1[ 0] = append3_t[0];
229 block1[ 1] = append3_t[1];
230 block1[ 2] = append3_t[2];
231 block1[ 3] = append3_t[3];
233 block1[ 4] = append4_t[0];
234 block1[ 5] = append4_t[1];
235 block1[ 6] = append4_t[2];
236 block1[ 7] = append4_t[3];
239 case 5: block0[ 5] |= append0_t[0];
240 block0[ 6] = append0_t[1];
241 block0[ 7] = append0_t[2];
242 block0[ 8] = append0_t[3];
244 block0[ 9] = append1_t[0];
245 block0[10] = append1_t[1];
246 block0[11] = append1_t[2];
247 block0[12] = append1_t[3];
249 block0[13] = append2_t[0];
250 block0[14] = append2_t[1];
251 block0[15] = append2_t[2];
252 block1[ 0] = append2_t[3];
254 block1[ 1] = append3_t[0];
255 block1[ 2] = append3_t[1];
256 block1[ 3] = append3_t[2];
257 block1[ 4] = append3_t[3];
259 block1[ 5] = append4_t[0];
260 block1[ 6] = append4_t[1];
261 block1[ 7] = append4_t[2];
262 block1[ 8] = append4_t[3];
265 case 6: block0[ 6] |= append0_t[0];
266 block0[ 7] = append0_t[1];
267 block0[ 8] = append0_t[2];
268 block0[ 9] = append0_t[3];
270 block0[10] = append1_t[0];
271 block0[11] = append1_t[1];
272 block0[12] = append1_t[2];
273 block0[13] = append1_t[3];
275 block0[14] = append2_t[0];
276 block0[15] = append2_t[1];
277 block1[ 0] = append2_t[2];
278 block1[ 1] = append2_t[3];
280 block1[ 2] = append3_t[0];
281 block1[ 3] = append3_t[1];
282 block1[ 4] = append3_t[2];
283 block1[ 5] = append3_t[3];
285 block1[ 6] = append4_t[0];
286 block1[ 7] = append4_t[1];
287 block1[ 8] = append4_t[2];
288 block1[ 9] = append4_t[3];
291 case 7: block0[ 7] |= append0_t[0];
292 block0[ 8] = append0_t[1];
293 block0[ 9] = append0_t[2];
294 block0[10] = append0_t[3];
296 block0[11] = append1_t[0];
297 block0[12] = append1_t[1];
298 block0[13] = append1_t[2];
299 block0[14] = append1_t[3];
301 block0[15] = append2_t[0];
302 block1[ 0] = append2_t[1];
303 block1[ 1] = append2_t[2];
304 block1[ 2] = append2_t[3];
306 block1[ 3] = append3_t[0];
307 block1[ 4] = append3_t[1];
308 block1[ 5] = append3_t[2];
309 block1[ 6] = append3_t[3];
311 block1[ 7] = append4_t[0];
312 block1[ 8] = append4_t[1];
313 block1[ 9] = append4_t[2];
314 block1[10] = append4_t[3];
317 case 8: block0[ 8] |= append0_t[0];
318 block0[ 9] = append0_t[1];
319 block0[10] = append0_t[2];
320 block0[11] = append0_t[3];
322 block0[12] = append1_t[0];
323 block0[13] = append1_t[1];
324 block0[14] = append1_t[2];
325 block0[15] = append1_t[3];
327 block1[ 0] = append2_t[0];
328 block1[ 1] = append2_t[1];
329 block1[ 2] = append2_t[2];
330 block1[ 3] = append2_t[3];
332 block1[ 4] = append3_t[0];
333 block1[ 5] = append3_t[1];
334 block1[ 6] = append3_t[2];
335 block1[ 7] = append3_t[3];
337 block1[ 8] = append4_t[0];
338 block1[ 9] = append4_t[1];
339 block1[10] = append4_t[2];
340 block1[11] = append4_t[3];
343 case 9: block0[ 9] |= append0_t[0];
344 block0[10] = append0_t[1];
345 block0[11] = append0_t[2];
346 block0[12] = append0_t[3];
348 block0[13] = append1_t[0];
349 block0[14] = append1_t[1];
350 block0[15] = append1_t[2];
351 block1[ 0] = append1_t[3];
353 block1[ 1] = append2_t[0];
354 block1[ 2] = append2_t[1];
355 block1[ 3] = append2_t[2];
356 block1[ 4] = append2_t[3];
358 block1[ 5] = append3_t[0];
359 block1[ 6] = append3_t[1];
360 block1[ 7] = append3_t[2];
361 block1[ 8] = append3_t[3];
363 block1[ 9] = append4_t[0];
364 block1[10] = append4_t[1];
365 block1[11] = append4_t[2];
366 block1[12] = append4_t[3];
369 case 10: block0[10] |= append0_t[0];
370 block0[11] = append0_t[1];
371 block0[12] = append0_t[2];
372 block0[13] = append0_t[3];
374 block0[14] = append1_t[0];
375 block0[15] = append1_t[1];
376 block1[ 0] = append1_t[2];
377 block1[ 1] = append1_t[3];
379 block1[ 2] = append2_t[0];
380 block1[ 3] = append2_t[1];
381 block1[ 4] = append2_t[2];
382 block1[ 5] = append2_t[3];
384 block1[ 6] = append3_t[0];
385 block1[ 7] = append3_t[1];
386 block1[ 8] = append3_t[2];
387 block1[ 9] = append3_t[3];
389 block1[10] = append4_t[0];
390 block1[11] = append4_t[1];
391 block1[12] = append4_t[2];
392 block1[13] = append4_t[3];
395 case 11: block0[11] |= append0_t[0];
396 block0[12] = append0_t[1];
397 block0[13] = append0_t[2];
398 block0[14] = append0_t[3];
400 block0[15] = append1_t[0];
401 block1[ 0] = append1_t[1];
402 block1[ 1] = append1_t[2];
403 block1[ 2] = append1_t[3];
405 block1[ 3] = append2_t[0];
406 block1[ 4] = append2_t[1];
407 block1[ 5] = append2_t[2];
408 block1[ 6] = append2_t[3];
410 block1[ 7] = append3_t[0];
411 block1[ 8] = append3_t[1];
412 block1[ 9] = append3_t[2];
413 block1[10] = append3_t[3];
415 block1[11] = append4_t[0];
416 block1[12] = append4_t[1];
417 block1[13] = append4_t[2];
418 block1[14] = append4_t[3];
421 case 12: block0[12] |= append0_t[0];
422 block0[13] = append0_t[1];
423 block0[14] = append0_t[2];
424 block0[15] = append0_t[3];
426 block1[ 0] = append1_t[0];
427 block1[ 1] = append1_t[1];
428 block1[ 2] = append1_t[2];
429 block1[ 3] = append1_t[3];
431 block1[ 4] = append2_t[0];
432 block1[ 5] = append2_t[1];
433 block1[ 6] = append2_t[2];
434 block1[ 7] = append2_t[3];
436 block1[ 8] = append3_t[0];
437 block1[ 9] = append3_t[1];
438 block1[10] = append3_t[2];
439 block1[11] = append3_t[3];
441 block1[12] = append4_t[0];
442 block1[13] = append4_t[1];
443 block1[14] = append4_t[2];
444 block1[15] = append4_t[3];
447 case 13: block0[13] |= append0_t[0];
448 block0[14] = append0_t[1];
449 block0[15] = append0_t[2];
450 block1[ 0] = append0_t[3];
452 block1[ 1] = append1_t[0];
453 block1[ 2] = append1_t[1];
454 block1[ 3] = append1_t[2];
455 block1[ 4] = append1_t[3];
457 block1[ 5] = append2_t[0];
458 block1[ 6] = append2_t[1];
459 block1[ 7] = append2_t[2];
460 block1[ 8] = append2_t[3];
462 block1[ 9] = append3_t[0];
463 block1[10] = append3_t[1];
464 block1[11] = append3_t[2];
465 block1[12] = append3_t[3];
467 block1[13] = append4_t[0];
468 block1[14] = append4_t[1];
469 block1[15] = append4_t[2];
472 case 14: block0[14] |= append0_t[0];
473 block0[15] = append0_t[1];
474 block1[ 0] = append0_t[2];
475 block1[ 1] = append0_t[3];
477 block1[ 2] = append1_t[0];
478 block1[ 3] = append1_t[1];
479 block1[ 4] = append1_t[2];
480 block1[ 5] = append1_t[3];
482 block1[ 6] = append2_t[0];
483 block1[ 7] = append2_t[1];
484 block1[ 8] = append2_t[2];
485 block1[ 9] = append2_t[3];
487 block1[10] = append3_t[0];
488 block1[11] = append3_t[1];
489 block1[12] = append3_t[2];
490 block1[13] = append3_t[3];
492 block1[14] = append4_t[0];
493 block1[15] = append4_t[1];
496 case 15: block0[15] |= append0_t[0];
497 block1[ 0] = append0_t[1];
498 block1[ 1] = append0_t[2];
499 block1[ 2] = append0_t[3];
501 block1[ 3] = append1_t[1];
502 block1[ 4] = append1_t[2];
503 block1[ 5] = append1_t[3];
504 block1[ 6] = append1_t[0];
506 block1[ 7] = append2_t[0];
507 block1[ 8] = append2_t[1];
508 block1[ 9] = append2_t[2];
509 block1[10] = append2_t[3];
511 block1[11] = append3_t[0];
512 block1[12] = append3_t[1];
513 block1[13] = append3_t[2];
514 block1[14] = append3_t[3];
516 block1[15] = append4_t[0];
519 case 16: block1[ 0] |= append0_t[0];
520 block1[ 1] = append0_t[1];
521 block1[ 2] = append0_t[2];
522 block1[ 3] = append0_t[3];
524 block1[ 4] = append1_t[0];
525 block1[ 5] = append1_t[1];
526 block1[ 6] = append1_t[2];
527 block1[ 7] = append1_t[3];
529 block1[ 8] = append2_t[0];
530 block1[ 9] = append2_t[1];
531 block1[10] = append2_t[2];
532 block1[11] = append2_t[3];
534 block1[12] = append3_t[0];
535 block1[13] = append3_t[1];
536 block1[14] = append3_t[2];
537 block1[15] = append3_t[3];
540 case 17: block1[ 1] |= append0_t[0];
541 block1[ 2] = append0_t[1];
542 block1[ 3] = append0_t[2];
543 block1[ 4] = append0_t[3];
545 block1[ 5] = append1_t[0];
546 block1[ 6] = append1_t[1];
547 block1[ 7] = append1_t[2];
548 block1[ 8] = append1_t[3];
550 block1[ 9] = append2_t[0];
551 block1[10] = append2_t[1];
552 block1[11] = append2_t[2];
553 block1[12] = append2_t[3];
555 block1[13] = append3_t[0];
556 block1[14] = append3_t[1];
557 block1[15] = append3_t[2];
560 case 18: block1[ 2] |= append0_t[0];
561 block1[ 3] = append0_t[1];
562 block1[ 4] = append0_t[2];
563 block1[ 5] = append0_t[3];
565 block1[ 6] = append1_t[0];
566 block1[ 7] = append1_t[1];
567 block1[ 8] = append1_t[2];
568 block1[ 9] = append1_t[3];
570 block1[10] = append2_t[0];
571 block1[11] = append2_t[1];
572 block1[12] = append2_t[2];
573 block1[13] = append2_t[3];
575 block1[14] = append3_t[0];
576 block1[15] = append3_t[1];
579 case 19: block1[ 3] |= append0_t[0];
580 block1[ 4] = append0_t[1];
581 block1[ 5] = append0_t[2];
582 block1[ 6] = append0_t[3];
584 block1[ 7] = append1_t[0];
585 block1[ 8] = append1_t[1];
586 block1[ 9] = append1_t[2];
587 block1[10] = append1_t[3];
589 block1[11] = append2_t[0];
590 block1[12] = append2_t[1];
591 block1[13] = append2_t[2];
592 block1[14] = append2_t[3];
594 block1[15] = append3_t[0];
597 case 20: block1[ 4] |= append0_t[0];
598 block1[ 5] = append0_t[1];
599 block1[ 6] = append0_t[2];
600 block1[ 7] = append0_t[3];
602 block1[ 8] = append1_t[0];
603 block1[ 9] = append1_t[1];
604 block1[10] = append1_t[2];
605 block1[11] = append1_t[3];
607 block1[12] = append2_t[0];
608 block1[13] = append2_t[1];
609 block1[14] = append2_t[2];
610 block1[15] = append2_t[3];
613 case 21: block1[ 5] |= append0_t[0];
614 block1[ 6] = append0_t[1];
615 block1[ 7] = append0_t[2];
616 block1[ 8] = append0_t[3];
618 block1[ 9] = append1_t[0];
619 block1[10] = append1_t[1];
620 block1[11] = append1_t[2];
621 block1[12] = append1_t[3];
623 block1[13] = append2_t[0];
624 block1[14] = append2_t[1];
625 block1[15] = append2_t[2];
628 case 22: block1[ 6] |= append0_t[0];
629 block1[ 7] = append0_t[1];
630 block1[ 8] = append0_t[2];
631 block1[ 9] = append0_t[3];
633 block1[10] = append1_t[0];
634 block1[11] = append1_t[1];
635 block1[12] = append1_t[2];
636 block1[13] = append1_t[3];
638 block1[14] = append2_t[0];
639 block1[15] = append2_t[1];
642 case 23: block1[ 7] |= append0_t[0];
643 block1[ 8] = append0_t[1];
644 block1[ 9] = append0_t[2];
645 block1[10] = append0_t[3];
647 block1[11] = append1_t[0];
648 block1[12] = append1_t[1];
649 block1[13] = append1_t[2];
650 block1[14] = append1_t[3];
652 block1[15] = append2_t[0];
655 case 24: block1[ 8] |= append0_t[0];
656 block1[ 9] = append0_t[1];
657 block1[10] = append0_t[2];
658 block1[11] = append0_t[3];
660 block1[12] = append1_t[0];
661 block1[13] = append1_t[1];
662 block1[14] = append1_t[2];
663 block1[15] = append1_t[3];
666 case 25: block1[ 9] |= append0_t[0];
667 block1[10] = append0_t[1];
668 block1[11] = append0_t[2];
669 block1[12] = append0_t[3];
671 block1[13] = append1_t[0];
672 block1[14] = append1_t[1];
673 block1[15] = append1_t[2];
676 case 26: block1[10] |= append0_t[0];
677 block1[11] = append0_t[1];
678 block1[12] = append0_t[2];
679 block1[13] = append0_t[3];
681 block1[14] = append1_t[0];
682 block1[15] = append1_t[1];
685 case 27: block1[11] |= append0_t[0];
686 block1[12] = append0_t[1];
687 block1[13] = append0_t[2];
688 block1[14] = append0_t[3];
690 block1[15] = append1_t[0];
693 case 28: block1[12] |= append0_t[0];
694 block1[13] = append0_t[1];
695 block1[14] = append0_t[2];
696 block1[15] = append0_t[3];
699 case 29: block1[13] |= append0_t[0];
700 block1[14] = append0_t[1];
701 block1[15] = append0_t[2];
704 case 30: block1[14] |= append0_t[0];
705 block1[15] = append0_t[1];
709 u32 new_len = block_len + append_len;
714 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11400_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 sip_t *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)
720 const u32 lid = get_local_id (0);
726 const u32 gid = get_global_id (0);
730 pw_buf0[0] = pws[gid].i[ 0];
731 pw_buf0[1] = pws[gid].i[ 1];
732 pw_buf0[2] = pws[gid].i[ 2];
733 pw_buf0[3] = pws[gid].i[ 3];
737 pw_buf1[0] = pws[gid].i[ 4];
738 pw_buf1[1] = pws[gid].i[ 5];
739 pw_buf1[2] = pws[gid].i[ 6];
740 pw_buf1[3] = pws[gid].i[ 7];
742 const u32 pw_len = pws[gid].pw_len;
748 __local u32 l_bin2asc[256];
750 const u32 lid4 = lid * 4;
752 const u32 lid40 = lid4 + 0;
753 const u32 lid41 = lid4 + 1;
754 const u32 lid42 = lid4 + 2;
755 const u32 lid43 = lid4 + 3;
757 const u32 v400 = (lid40 >> 0) & 15;
758 const u32 v401 = (lid40 >> 4) & 15;
759 const u32 v410 = (lid41 >> 0) & 15;
760 const u32 v411 = (lid41 >> 4) & 15;
761 const u32 v420 = (lid42 >> 0) & 15;
762 const u32 v421 = (lid42 >> 4) & 15;
763 const u32 v430 = (lid43 >> 0) & 15;
764 const u32 v431 = (lid43 >> 4) & 15;
766 l_bin2asc[lid40] = ((v400 < 10) ? '0' + v400 : 'a' - 10 + v400) << 8
767 | ((v401 < 10) ? '0' + v401 : 'a' - 10 + v401) << 0;
768 l_bin2asc[lid41] = ((v410 < 10) ? '0' + v410 : 'a' - 10 + v410) << 8
769 | ((v411 < 10) ? '0' + v411 : 'a' - 10 + v411) << 0;
770 l_bin2asc[lid42] = ((v420 < 10) ? '0' + v420 : 'a' - 10 + v420) << 8
771 | ((v421 < 10) ? '0' + v421 : 'a' - 10 + v421) << 0;
772 l_bin2asc[lid43] = ((v430 < 10) ? '0' + v430 : 'a' - 10 + v430) << 8
773 | ((v431 < 10) ? '0' + v431 : 'a' - 10 + v431) << 0;
775 barrier (CLK_LOCAL_MEM_FENCE);
777 if (gid >= gid_max) return;
783 const u32 salt_len = esalt_bufs[salt_pos].salt_len; // not a bug, we need to get it from the esalt
787 salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0];
788 salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1];
789 salt_buf0[ 2] = esalt_bufs[salt_pos].salt_buf[ 2];
790 salt_buf0[ 3] = esalt_bufs[salt_pos].salt_buf[ 3];
791 salt_buf0[ 4] = esalt_bufs[salt_pos].salt_buf[ 4];
792 salt_buf0[ 5] = esalt_bufs[salt_pos].salt_buf[ 5];
793 salt_buf0[ 6] = esalt_bufs[salt_pos].salt_buf[ 6];
794 salt_buf0[ 7] = esalt_bufs[salt_pos].salt_buf[ 7];
795 salt_buf0[ 8] = esalt_bufs[salt_pos].salt_buf[ 8];
796 salt_buf0[ 9] = esalt_bufs[salt_pos].salt_buf[ 9];
797 salt_buf0[10] = esalt_bufs[salt_pos].salt_buf[10];
798 salt_buf0[11] = esalt_bufs[salt_pos].salt_buf[11];
799 salt_buf0[12] = esalt_bufs[salt_pos].salt_buf[12];
800 salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13];
801 salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14];
802 salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15];
806 salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16];
807 salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17];
808 salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18];
809 salt_buf1[ 3] = esalt_bufs[salt_pos].salt_buf[19];
810 salt_buf1[ 4] = esalt_bufs[salt_pos].salt_buf[20];
811 salt_buf1[ 5] = esalt_bufs[salt_pos].salt_buf[21];
812 salt_buf1[ 6] = esalt_bufs[salt_pos].salt_buf[22];
813 salt_buf1[ 7] = esalt_bufs[salt_pos].salt_buf[23];
814 salt_buf1[ 8] = esalt_bufs[salt_pos].salt_buf[24];
815 salt_buf1[ 9] = esalt_bufs[salt_pos].salt_buf[25];
816 salt_buf1[10] = esalt_bufs[salt_pos].salt_buf[26];
817 salt_buf1[11] = esalt_bufs[salt_pos].salt_buf[27];
818 salt_buf1[12] = esalt_bufs[salt_pos].salt_buf[28];
819 salt_buf1[13] = esalt_bufs[salt_pos].salt_buf[29];
827 const u32 esalt_len = esalt_bufs[salt_pos].esalt_len;
831 esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0];
832 esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1];
833 esalt_buf0[ 2] = esalt_bufs[salt_pos].esalt_buf[ 2];
834 esalt_buf0[ 3] = esalt_bufs[salt_pos].esalt_buf[ 3];
835 esalt_buf0[ 4] = esalt_bufs[salt_pos].esalt_buf[ 4];
836 esalt_buf0[ 5] = esalt_bufs[salt_pos].esalt_buf[ 5];
837 esalt_buf0[ 6] = esalt_bufs[salt_pos].esalt_buf[ 6];
838 esalt_buf0[ 7] = esalt_bufs[salt_pos].esalt_buf[ 7];
839 esalt_buf0[ 8] = esalt_bufs[salt_pos].esalt_buf[ 8];
840 esalt_buf0[ 9] = esalt_bufs[salt_pos].esalt_buf[ 9];
841 esalt_buf0[10] = esalt_bufs[salt_pos].esalt_buf[10];
842 esalt_buf0[11] = esalt_bufs[salt_pos].esalt_buf[11];
843 esalt_buf0[12] = esalt_bufs[salt_pos].esalt_buf[12];
844 esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13];
845 esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14];
846 esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15];
850 esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16];
851 esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17];
852 esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18];
853 esalt_buf1[ 3] = esalt_bufs[salt_pos].esalt_buf[19];
854 esalt_buf1[ 4] = esalt_bufs[salt_pos].esalt_buf[20];
855 esalt_buf1[ 5] = esalt_bufs[salt_pos].esalt_buf[21];
856 esalt_buf1[ 6] = esalt_bufs[salt_pos].esalt_buf[22];
857 esalt_buf1[ 7] = esalt_bufs[salt_pos].esalt_buf[23];
858 esalt_buf1[ 8] = esalt_bufs[salt_pos].esalt_buf[24];
859 esalt_buf1[ 9] = esalt_bufs[salt_pos].esalt_buf[25];
860 esalt_buf1[10] = esalt_bufs[salt_pos].esalt_buf[26];
861 esalt_buf1[11] = esalt_bufs[salt_pos].esalt_buf[27];
862 esalt_buf1[12] = esalt_bufs[salt_pos].esalt_buf[28];
863 esalt_buf1[13] = esalt_bufs[salt_pos].esalt_buf[29];
864 esalt_buf1[14] = esalt_bufs[salt_pos].esalt_buf[30];
865 esalt_buf1[15] = esalt_bufs[salt_pos].esalt_buf[31];
869 esalt_buf2[ 0] = esalt_bufs[salt_pos].esalt_buf[32];
870 esalt_buf2[ 1] = esalt_bufs[salt_pos].esalt_buf[33];
871 esalt_buf2[ 2] = esalt_bufs[salt_pos].esalt_buf[34];
872 esalt_buf2[ 3] = esalt_bufs[salt_pos].esalt_buf[35];
873 esalt_buf2[ 4] = esalt_bufs[salt_pos].esalt_buf[36];
874 esalt_buf2[ 5] = esalt_bufs[salt_pos].esalt_buf[37];
886 const u32 digest_esalt_len = 32 + esalt_len;
887 const u32 remaining_bytes = digest_esalt_len + 1 - 64; // substract previous block
893 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
923 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
925 append_0x80_2x4 (w0, w1, out_len);
927 const u32 pw_salt_len = salt_len + out_len;
930 * HA1 = md5 ($salt . $pass)
933 // append the pass to the salt
937 block0[ 0] = salt_buf0[ 0];
938 block0[ 1] = salt_buf0[ 1];
939 block0[ 2] = salt_buf0[ 2];
940 block0[ 3] = salt_buf0[ 3];
941 block0[ 4] = salt_buf0[ 4];
942 block0[ 5] = salt_buf0[ 5];
943 block0[ 6] = salt_buf0[ 6];
944 block0[ 7] = salt_buf0[ 7];
945 block0[ 8] = salt_buf0[ 8];
946 block0[ 9] = salt_buf0[ 9];
947 block0[10] = salt_buf0[10];
948 block0[11] = salt_buf0[11];
949 block0[12] = salt_buf0[12];
950 block0[13] = salt_buf0[13];
951 block0[14] = salt_buf0[14];
952 block0[15] = salt_buf0[15];
956 block1[ 0] = salt_buf1[ 0];
957 block1[ 1] = salt_buf1[ 1];
958 block1[ 2] = salt_buf1[ 2];
959 block1[ 3] = salt_buf1[ 3];
960 block1[ 4] = salt_buf1[ 4];
961 block1[ 5] = salt_buf1[ 5];
962 block1[ 6] = salt_buf1[ 6];
963 block1[ 7] = salt_buf1[ 7];
964 block1[ 8] = salt_buf1[ 8];
965 block1[ 9] = salt_buf1[ 9];
966 block1[10] = salt_buf1[10];
967 block1[11] = salt_buf1[11];
968 block1[12] = salt_buf1[12];
969 block1[13] = salt_buf1[13];
970 block1[14] = salt_buf1[14];
971 block1[15] = salt_buf1[15];
975 block_len = memcat32 (block0, block1, salt_len, w0, w1, w2, w3, out_len);
979 w0_t[0] = block0[ 0];
980 w0_t[1] = block0[ 1];
981 w0_t[2] = block0[ 2];
982 w0_t[3] = block0[ 3];
986 w1_t[0] = block0[ 4];
987 w1_t[1] = block0[ 5];
988 w1_t[2] = block0[ 6];
989 w1_t[3] = block0[ 7];
993 w2_t[0] = block0[ 8];
994 w2_t[1] = block0[ 9];
995 w2_t[2] = block0[10];
996 w2_t[3] = block0[11];
1000 w3_t[0] = block0[12];
1001 w3_t[1] = block0[13];
1002 w3_t[2] = block0[14];
1003 w3_t[3] = block0[15];
1007 w3_t[2] = pw_salt_len * 8;
1019 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1020 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1021 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1022 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1023 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1024 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1025 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1026 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1027 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1028 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1029 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1030 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1031 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1032 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1033 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1034 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1036 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1037 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1038 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1039 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1040 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1041 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1042 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1043 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1044 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1045 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1046 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1047 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1048 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1049 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1050 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1051 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1053 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1054 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1055 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1056 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1057 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1058 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1059 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1060 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1061 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1062 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1063 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1064 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1065 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1066 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1067 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1068 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1070 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1071 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1072 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1073 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1074 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1075 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1076 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1077 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1078 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1079 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1080 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1081 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1082 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1083 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1084 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1085 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1099 w0_t[0] = block1[ 0];
1100 w0_t[1] = block1[ 1];
1101 w0_t[2] = block1[ 2];
1102 w0_t[3] = block1[ 3];
1104 w1_t[0] = block1[ 4];
1105 w1_t[1] = block1[ 5];
1106 w1_t[2] = block1[ 6];
1107 w1_t[3] = block1[ 7];
1109 w2_t[0] = block1[ 8];
1110 w2_t[1] = block1[ 9];
1111 w2_t[2] = block1[10];
1112 w2_t[3] = block1[11];
1114 w3_t[0] = block1[12];
1115 w3_t[1] = block1[13];
1116 w3_t[2] = pw_salt_len * 8;
1119 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1120 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1121 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1122 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1123 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1124 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1125 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1126 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1127 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1128 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1129 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1130 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1131 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1132 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1133 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1134 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1136 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1137 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1138 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1139 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1140 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1141 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1142 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1143 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1144 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1145 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1146 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1147 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1148 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1149 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1150 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1151 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1153 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1154 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1155 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1156 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1157 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1158 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1159 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1160 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1161 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1162 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1163 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1164 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1165 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1166 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1167 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1168 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1170 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1171 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1172 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1173 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1174 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1175 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1176 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1177 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1178 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1179 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1180 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1181 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1182 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1183 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1184 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1185 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1194 * final = md5 ($HA1 . $esalt)
1195 * we have at least 2 MD5 blocks/transformations, but we might need 3
1198 w0_t[0] = uint_to_hex_lower8 ((a >> 0) & 255) << 0
1199 | uint_to_hex_lower8 ((a >> 8) & 255) << 16;
1200 w0_t[1] = uint_to_hex_lower8 ((a >> 16) & 255) << 0
1201 | uint_to_hex_lower8 ((a >> 24) & 255) << 16;
1202 w0_t[2] = uint_to_hex_lower8 ((b >> 0) & 255) << 0
1203 | uint_to_hex_lower8 ((b >> 8) & 255) << 16;
1204 w0_t[3] = uint_to_hex_lower8 ((b >> 16) & 255) << 0
1205 | uint_to_hex_lower8 ((b >> 24) & 255) << 16;
1206 w1_t[0] = uint_to_hex_lower8 ((c >> 0) & 255) << 0
1207 | uint_to_hex_lower8 ((c >> 8) & 255) << 16;
1208 w1_t[1] = uint_to_hex_lower8 ((c >> 16) & 255) << 0
1209 | uint_to_hex_lower8 ((c >> 24) & 255) << 16;
1210 w1_t[2] = uint_to_hex_lower8 ((d >> 0) & 255) << 0
1211 | uint_to_hex_lower8 ((d >> 8) & 255) << 16;
1212 w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0
1213 | uint_to_hex_lower8 ((d >> 24) & 255) << 16;
1215 w2_t[0] = esalt_buf0[0];
1216 w2_t[1] = esalt_buf0[1];
1217 w2_t[2] = esalt_buf0[2];
1218 w2_t[3] = esalt_buf0[3];
1220 w3_t[0] = esalt_buf0[4];
1221 w3_t[1] = esalt_buf0[5];
1222 w3_t[2] = esalt_buf0[6];
1223 w3_t[3] = esalt_buf0[7];
1233 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1234 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1235 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1236 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1237 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1238 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1239 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1240 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1241 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1242 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1243 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1244 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1245 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1246 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1247 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1248 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1250 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1251 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1252 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1253 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1254 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1255 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1256 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1257 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1258 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1259 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1260 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1261 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1262 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1263 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1264 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1265 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1267 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1268 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1269 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1270 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1271 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1272 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1273 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1274 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1275 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1276 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1277 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1278 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1279 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1280 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1281 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1282 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1284 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1285 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1286 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1287 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1288 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1289 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1290 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1291 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1292 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1293 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1294 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1295 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1296 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1297 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1298 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1299 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1313 w0_t[0] = esalt_buf0[ 8];
1314 w0_t[1] = esalt_buf0[ 9];
1315 w0_t[2] = esalt_buf0[10];
1316 w0_t[3] = esalt_buf0[11];
1318 w1_t[0] = esalt_buf0[12];
1319 w1_t[1] = esalt_buf0[13];
1320 w1_t[2] = esalt_buf0[14];
1321 w1_t[3] = esalt_buf0[15];
1323 w2_t[0] = esalt_buf1[ 0];
1324 w2_t[1] = esalt_buf1[ 1];
1325 w2_t[2] = esalt_buf1[ 2];
1326 w2_t[3] = esalt_buf1[ 3];
1328 w3_t[0] = esalt_buf1[ 4];
1329 w3_t[1] = esalt_buf1[ 5];
1330 w3_t[2] = esalt_buf1[ 6];
1331 w3_t[3] = esalt_buf1[ 7];
1333 // it is the final block when no more than 55 bytes left
1335 if (remaining_bytes < 56)
1337 // it is the last block !
1339 w3_t[2] = digest_esalt_len * 8;
1342 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1343 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1344 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1345 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1346 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1347 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1348 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1349 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1350 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1351 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1352 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1353 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1354 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1355 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1356 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1357 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1359 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1360 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1361 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1362 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1363 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1364 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1365 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1366 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1367 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1368 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1369 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1370 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1371 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1372 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1373 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1374 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1376 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1377 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1378 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1379 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1380 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1381 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1382 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1383 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1384 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1385 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1386 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1387 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1388 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1389 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1390 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1391 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1393 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1394 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1395 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1396 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1397 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1398 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1399 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1400 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1401 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1402 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1403 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1404 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1405 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1406 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1407 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1408 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1410 // sometimes (not rare at all) we need a third block :(
1412 if (remaining_bytes > 55)
1414 // this is for sure the final block
1426 w0_t[0] = esalt_buf1[ 8];
1427 w0_t[1] = esalt_buf1[ 9];
1428 w0_t[2] = esalt_buf1[10];
1429 w0_t[3] = esalt_buf1[11];
1431 w1_t[0] = esalt_buf1[12];
1432 w1_t[1] = esalt_buf1[13];
1433 w1_t[2] = esalt_buf1[14];
1434 w1_t[3] = esalt_buf1[15];
1436 w2_t[0] = esalt_buf2[ 0];
1437 w2_t[1] = esalt_buf2[ 1];
1438 w2_t[2] = esalt_buf2[ 2];
1439 w2_t[3] = esalt_buf2[ 3];
1441 w3_t[0] = esalt_buf2[ 4];
1442 w3_t[1] = esalt_buf2[ 5];
1443 w3_t[2] = digest_esalt_len * 8;
1446 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1447 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1448 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1449 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1450 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1451 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1452 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1453 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1454 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1455 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1456 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1457 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1458 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1459 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1460 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1461 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1463 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1464 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1465 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1466 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1467 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1468 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1469 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1470 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1471 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1472 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1473 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1474 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1475 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1476 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1477 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1478 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1480 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1481 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1482 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1483 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1484 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1485 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1486 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1487 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1488 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1489 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1490 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1491 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1492 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1493 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1494 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1495 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1497 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1498 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1499 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1500 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1501 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1502 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1503 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1504 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1505 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1506 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1507 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1508 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1509 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1510 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1511 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1512 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1529 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11400_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 sip_t *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)
1533 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11400_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 sip_t *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)
1537 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11400_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 sip_t *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)
1543 const u32 lid = get_local_id (0);
1549 const u32 gid = get_global_id (0);
1553 pw_buf0[0] = pws[gid].i[ 0];
1554 pw_buf0[1] = pws[gid].i[ 1];
1555 pw_buf0[2] = pws[gid].i[ 2];
1556 pw_buf0[3] = pws[gid].i[ 3];
1560 pw_buf1[0] = pws[gid].i[ 4];
1561 pw_buf1[1] = pws[gid].i[ 5];
1562 pw_buf1[2] = pws[gid].i[ 6];
1563 pw_buf1[3] = pws[gid].i[ 7];
1565 const u32 pw_len = pws[gid].pw_len;
1571 __local u32 l_bin2asc[256];
1573 const u32 lid4 = lid * 4;
1575 const u32 lid40 = lid4 + 0;
1576 const u32 lid41 = lid4 + 1;
1577 const u32 lid42 = lid4 + 2;
1578 const u32 lid43 = lid4 + 3;
1580 const u32 v400 = (lid40 >> 0) & 15;
1581 const u32 v401 = (lid40 >> 4) & 15;
1582 const u32 v410 = (lid41 >> 0) & 15;
1583 const u32 v411 = (lid41 >> 4) & 15;
1584 const u32 v420 = (lid42 >> 0) & 15;
1585 const u32 v421 = (lid42 >> 4) & 15;
1586 const u32 v430 = (lid43 >> 0) & 15;
1587 const u32 v431 = (lid43 >> 4) & 15;
1589 l_bin2asc[lid40] = ((v400 < 10) ? '0' + v400 : 'a' - 10 + v400) << 8
1590 | ((v401 < 10) ? '0' + v401 : 'a' - 10 + v401) << 0;
1591 l_bin2asc[lid41] = ((v410 < 10) ? '0' + v410 : 'a' - 10 + v410) << 8
1592 | ((v411 < 10) ? '0' + v411 : 'a' - 10 + v411) << 0;
1593 l_bin2asc[lid42] = ((v420 < 10) ? '0' + v420 : 'a' - 10 + v420) << 8
1594 | ((v421 < 10) ? '0' + v421 : 'a' - 10 + v421) << 0;
1595 l_bin2asc[lid43] = ((v430 < 10) ? '0' + v430 : 'a' - 10 + v430) << 8
1596 | ((v431 < 10) ? '0' + v431 : 'a' - 10 + v431) << 0;
1598 barrier (CLK_LOCAL_MEM_FENCE);
1600 if (gid >= gid_max) return;
1606 const u32 salt_len = esalt_bufs[salt_pos].salt_len; // not a bug, we need to get it from the esalt
1610 salt_buf0[ 0] = esalt_bufs[salt_pos].salt_buf[ 0];
1611 salt_buf0[ 1] = esalt_bufs[salt_pos].salt_buf[ 1];
1612 salt_buf0[ 2] = esalt_bufs[salt_pos].salt_buf[ 2];
1613 salt_buf0[ 3] = esalt_bufs[salt_pos].salt_buf[ 3];
1614 salt_buf0[ 4] = esalt_bufs[salt_pos].salt_buf[ 4];
1615 salt_buf0[ 5] = esalt_bufs[salt_pos].salt_buf[ 5];
1616 salt_buf0[ 6] = esalt_bufs[salt_pos].salt_buf[ 6];
1617 salt_buf0[ 7] = esalt_bufs[salt_pos].salt_buf[ 7];
1618 salt_buf0[ 8] = esalt_bufs[salt_pos].salt_buf[ 8];
1619 salt_buf0[ 9] = esalt_bufs[salt_pos].salt_buf[ 9];
1620 salt_buf0[10] = esalt_bufs[salt_pos].salt_buf[10];
1621 salt_buf0[11] = esalt_bufs[salt_pos].salt_buf[11];
1622 salt_buf0[12] = esalt_bufs[salt_pos].salt_buf[12];
1623 salt_buf0[13] = esalt_bufs[salt_pos].salt_buf[13];
1624 salt_buf0[14] = esalt_bufs[salt_pos].salt_buf[14];
1625 salt_buf0[15] = esalt_bufs[salt_pos].salt_buf[15];
1629 salt_buf1[ 0] = esalt_bufs[salt_pos].salt_buf[16];
1630 salt_buf1[ 1] = esalt_bufs[salt_pos].salt_buf[17];
1631 salt_buf1[ 2] = esalt_bufs[salt_pos].salt_buf[18];
1632 salt_buf1[ 3] = esalt_bufs[salt_pos].salt_buf[19];
1633 salt_buf1[ 4] = esalt_bufs[salt_pos].salt_buf[20];
1634 salt_buf1[ 5] = esalt_bufs[salt_pos].salt_buf[21];
1635 salt_buf1[ 6] = esalt_bufs[salt_pos].salt_buf[22];
1636 salt_buf1[ 7] = esalt_bufs[salt_pos].salt_buf[23];
1637 salt_buf1[ 8] = esalt_bufs[salt_pos].salt_buf[24];
1638 salt_buf1[ 9] = esalt_bufs[salt_pos].salt_buf[25];
1639 salt_buf1[10] = esalt_bufs[salt_pos].salt_buf[26];
1640 salt_buf1[11] = esalt_bufs[salt_pos].salt_buf[27];
1641 salt_buf1[12] = esalt_bufs[salt_pos].salt_buf[28];
1642 salt_buf1[13] = esalt_bufs[salt_pos].salt_buf[29];
1650 const u32 esalt_len = esalt_bufs[salt_pos].esalt_len;
1654 esalt_buf0[ 0] = esalt_bufs[salt_pos].esalt_buf[ 0];
1655 esalt_buf0[ 1] = esalt_bufs[salt_pos].esalt_buf[ 1];
1656 esalt_buf0[ 2] = esalt_bufs[salt_pos].esalt_buf[ 2];
1657 esalt_buf0[ 3] = esalt_bufs[salt_pos].esalt_buf[ 3];
1658 esalt_buf0[ 4] = esalt_bufs[salt_pos].esalt_buf[ 4];
1659 esalt_buf0[ 5] = esalt_bufs[salt_pos].esalt_buf[ 5];
1660 esalt_buf0[ 6] = esalt_bufs[salt_pos].esalt_buf[ 6];
1661 esalt_buf0[ 7] = esalt_bufs[salt_pos].esalt_buf[ 7];
1662 esalt_buf0[ 8] = esalt_bufs[salt_pos].esalt_buf[ 8];
1663 esalt_buf0[ 9] = esalt_bufs[salt_pos].esalt_buf[ 9];
1664 esalt_buf0[10] = esalt_bufs[salt_pos].esalt_buf[10];
1665 esalt_buf0[11] = esalt_bufs[salt_pos].esalt_buf[11];
1666 esalt_buf0[12] = esalt_bufs[salt_pos].esalt_buf[12];
1667 esalt_buf0[13] = esalt_bufs[salt_pos].esalt_buf[13];
1668 esalt_buf0[14] = esalt_bufs[salt_pos].esalt_buf[14];
1669 esalt_buf0[15] = esalt_bufs[salt_pos].esalt_buf[15];
1673 esalt_buf1[ 0] = esalt_bufs[salt_pos].esalt_buf[16];
1674 esalt_buf1[ 1] = esalt_bufs[salt_pos].esalt_buf[17];
1675 esalt_buf1[ 2] = esalt_bufs[salt_pos].esalt_buf[18];
1676 esalt_buf1[ 3] = esalt_bufs[salt_pos].esalt_buf[19];
1677 esalt_buf1[ 4] = esalt_bufs[salt_pos].esalt_buf[20];
1678 esalt_buf1[ 5] = esalt_bufs[salt_pos].esalt_buf[21];
1679 esalt_buf1[ 6] = esalt_bufs[salt_pos].esalt_buf[22];
1680 esalt_buf1[ 7] = esalt_bufs[salt_pos].esalt_buf[23];
1681 esalt_buf1[ 8] = esalt_bufs[salt_pos].esalt_buf[24];
1682 esalt_buf1[ 9] = esalt_bufs[salt_pos].esalt_buf[25];
1683 esalt_buf1[10] = esalt_bufs[salt_pos].esalt_buf[26];
1684 esalt_buf1[11] = esalt_bufs[salt_pos].esalt_buf[27];
1685 esalt_buf1[12] = esalt_bufs[salt_pos].esalt_buf[28];
1686 esalt_buf1[13] = esalt_bufs[salt_pos].esalt_buf[29];
1687 esalt_buf1[14] = esalt_bufs[salt_pos].esalt_buf[30];
1688 esalt_buf1[15] = esalt_bufs[salt_pos].esalt_buf[31];
1692 esalt_buf2[ 0] = esalt_bufs[salt_pos].esalt_buf[32];
1693 esalt_buf2[ 1] = esalt_bufs[salt_pos].esalt_buf[33];
1694 esalt_buf2[ 2] = esalt_bufs[salt_pos].esalt_buf[34];
1695 esalt_buf2[ 3] = esalt_bufs[salt_pos].esalt_buf[35];
1696 esalt_buf2[ 4] = esalt_bufs[salt_pos].esalt_buf[36];
1697 esalt_buf2[ 5] = esalt_bufs[salt_pos].esalt_buf[37];
1709 const u32 digest_esalt_len = 32 + esalt_len;
1710 const u32 remaining_bytes = digest_esalt_len + 1 - 64; // substract previous block
1716 const u32 search[4] =
1718 digests_buf[digests_offset].digest_buf[DGST_R0],
1719 digests_buf[digests_offset].digest_buf[DGST_R1],
1720 digests_buf[digests_offset].digest_buf[DGST_R2],
1721 digests_buf[digests_offset].digest_buf[DGST_R3]
1728 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
1758 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
1760 append_0x80_2x4 (w0, w1, out_len);
1762 const u32 pw_salt_len = salt_len + out_len;
1765 * HA1 = md5 ($salt . $pass)
1768 // append the pass to the salt
1772 block0[ 0] = salt_buf0[ 0];
1773 block0[ 1] = salt_buf0[ 1];
1774 block0[ 2] = salt_buf0[ 2];
1775 block0[ 3] = salt_buf0[ 3];
1776 block0[ 4] = salt_buf0[ 4];
1777 block0[ 5] = salt_buf0[ 5];
1778 block0[ 6] = salt_buf0[ 6];
1779 block0[ 7] = salt_buf0[ 7];
1780 block0[ 8] = salt_buf0[ 8];
1781 block0[ 9] = salt_buf0[ 9];
1782 block0[10] = salt_buf0[10];
1783 block0[11] = salt_buf0[11];
1784 block0[12] = salt_buf0[12];
1785 block0[13] = salt_buf0[13];
1786 block0[14] = salt_buf0[14];
1787 block0[15] = salt_buf0[15];
1791 block1[ 0] = salt_buf1[ 0];
1792 block1[ 1] = salt_buf1[ 1];
1793 block1[ 2] = salt_buf1[ 2];
1794 block1[ 3] = salt_buf1[ 3];
1795 block1[ 4] = salt_buf1[ 4];
1796 block1[ 5] = salt_buf1[ 5];
1797 block1[ 6] = salt_buf1[ 6];
1798 block1[ 7] = salt_buf1[ 7];
1799 block1[ 8] = salt_buf1[ 8];
1800 block1[ 9] = salt_buf1[ 9];
1801 block1[10] = salt_buf1[10];
1802 block1[11] = salt_buf1[11];
1803 block1[12] = salt_buf1[12];
1804 block1[13] = salt_buf1[13];
1805 block1[14] = salt_buf1[14];
1806 block1[15] = salt_buf1[15];
1810 block_len = memcat32 (block0, block1, salt_len, w0, w1, w2, w3, out_len);
1814 w0_t[0] = block0[ 0];
1815 w0_t[1] = block0[ 1];
1816 w0_t[2] = block0[ 2];
1817 w0_t[3] = block0[ 3];
1821 w1_t[0] = block0[ 4];
1822 w1_t[1] = block0[ 5];
1823 w1_t[2] = block0[ 6];
1824 w1_t[3] = block0[ 7];
1828 w2_t[0] = block0[ 8];
1829 w2_t[1] = block0[ 9];
1830 w2_t[2] = block0[10];
1831 w2_t[3] = block0[11];
1835 w3_t[0] = block0[12];
1836 w3_t[1] = block0[13];
1837 w3_t[2] = block0[14];
1838 w3_t[3] = block0[15];
1842 w3_t[2] = pw_salt_len * 8;
1854 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1855 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1856 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1857 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1858 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1859 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1860 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1861 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1862 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1863 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1864 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1865 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1866 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1867 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1868 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1869 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1871 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1872 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1873 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1874 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1875 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1876 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1877 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1878 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1879 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1880 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1881 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1882 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1883 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1884 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1885 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1886 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1888 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1889 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1890 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1891 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1892 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1893 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1894 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1895 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1896 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1897 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1898 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1899 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
1900 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
1901 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
1902 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
1903 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
1905 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
1906 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
1907 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
1908 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
1909 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
1910 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
1911 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
1912 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
1913 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
1914 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
1915 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
1916 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
1917 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
1918 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
1919 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
1920 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
1934 w0_t[0] = block1[ 0];
1935 w0_t[1] = block1[ 1];
1936 w0_t[2] = block1[ 2];
1937 w0_t[3] = block1[ 3];
1939 w1_t[0] = block1[ 4];
1940 w1_t[1] = block1[ 5];
1941 w1_t[2] = block1[ 6];
1942 w1_t[3] = block1[ 7];
1944 w2_t[0] = block1[ 8];
1945 w2_t[1] = block1[ 9];
1946 w2_t[2] = block1[10];
1947 w2_t[3] = block1[11];
1949 w3_t[0] = block1[12];
1950 w3_t[1] = block1[13];
1951 w3_t[2] = pw_salt_len * 8;
1954 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
1955 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
1956 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
1957 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
1958 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
1959 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
1960 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
1961 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
1962 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
1963 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
1964 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
1965 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
1966 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
1967 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
1968 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
1969 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
1971 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
1972 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
1973 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
1974 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
1975 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
1976 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
1977 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
1978 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
1979 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
1980 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
1981 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
1982 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
1983 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
1984 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
1985 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
1986 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
1988 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
1989 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
1990 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
1991 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
1992 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
1993 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
1994 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
1995 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
1996 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
1997 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
1998 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
1999 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
2000 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
2001 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
2002 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
2003 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
2005 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
2006 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
2007 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
2008 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
2009 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
2010 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
2011 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
2012 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
2013 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
2014 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
2015 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
2016 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
2017 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
2018 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
2019 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
2020 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
2029 * final = md5 ($HA1 . $esalt)
2030 * we have at least 2 MD5 blocks/transformations, but we might need 3
2033 w0_t[0] = uint_to_hex_lower8 ((a >> 0) & 255) << 0
2034 | uint_to_hex_lower8 ((a >> 8) & 255) << 16;
2035 w0_t[1] = uint_to_hex_lower8 ((a >> 16) & 255) << 0
2036 | uint_to_hex_lower8 ((a >> 24) & 255) << 16;
2037 w0_t[2] = uint_to_hex_lower8 ((b >> 0) & 255) << 0
2038 | uint_to_hex_lower8 ((b >> 8) & 255) << 16;
2039 w0_t[3] = uint_to_hex_lower8 ((b >> 16) & 255) << 0
2040 | uint_to_hex_lower8 ((b >> 24) & 255) << 16;
2041 w1_t[0] = uint_to_hex_lower8 ((c >> 0) & 255) << 0
2042 | uint_to_hex_lower8 ((c >> 8) & 255) << 16;
2043 w1_t[1] = uint_to_hex_lower8 ((c >> 16) & 255) << 0
2044 | uint_to_hex_lower8 ((c >> 24) & 255) << 16;
2045 w1_t[2] = uint_to_hex_lower8 ((d >> 0) & 255) << 0
2046 | uint_to_hex_lower8 ((d >> 8) & 255) << 16;
2047 w1_t[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0
2048 | uint_to_hex_lower8 ((d >> 24) & 255) << 16;
2050 w2_t[0] = esalt_buf0[0];
2051 w2_t[1] = esalt_buf0[1];
2052 w2_t[2] = esalt_buf0[2];
2053 w2_t[3] = esalt_buf0[3];
2055 w3_t[0] = esalt_buf0[4];
2056 w3_t[1] = esalt_buf0[5];
2057 w3_t[2] = esalt_buf0[6];
2058 w3_t[3] = esalt_buf0[7];
2068 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
2069 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
2070 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
2071 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
2072 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
2073 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
2074 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
2075 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
2076 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
2077 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
2078 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
2079 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
2080 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
2081 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
2082 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
2083 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
2085 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
2086 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
2087 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
2088 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
2089 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
2090 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
2091 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
2092 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
2093 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
2094 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
2095 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
2096 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
2097 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
2098 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
2099 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
2100 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
2102 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
2103 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
2104 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
2105 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
2106 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
2107 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
2108 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
2109 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
2110 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
2111 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
2112 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
2113 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
2114 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
2115 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
2116 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
2117 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
2119 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
2120 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
2121 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
2122 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
2123 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
2124 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
2125 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
2126 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
2127 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
2128 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
2129 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
2130 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
2131 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
2132 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
2133 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
2134 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
2148 w0_t[0] = esalt_buf0[ 8];
2149 w0_t[1] = esalt_buf0[ 9];
2150 w0_t[2] = esalt_buf0[10];
2151 w0_t[3] = esalt_buf0[11];
2153 w1_t[0] = esalt_buf0[12];
2154 w1_t[1] = esalt_buf0[13];
2155 w1_t[2] = esalt_buf0[14];
2156 w1_t[3] = esalt_buf0[15];
2158 w2_t[0] = esalt_buf1[ 0];
2159 w2_t[1] = esalt_buf1[ 1];
2160 w2_t[2] = esalt_buf1[ 2];
2161 w2_t[3] = esalt_buf1[ 3];
2163 w3_t[0] = esalt_buf1[ 4];
2164 w3_t[1] = esalt_buf1[ 5];
2165 w3_t[2] = esalt_buf1[ 6];
2166 w3_t[3] = esalt_buf1[ 7];
2168 // it is the final block when no more than 55 bytes left
2170 if (remaining_bytes < 56)
2172 // it is the last block !
2174 w3_t[2] = digest_esalt_len * 8;
2177 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
2178 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
2179 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
2180 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
2181 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
2182 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
2183 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
2184 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
2185 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
2186 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
2187 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
2188 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
2189 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
2190 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
2191 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
2192 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
2194 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
2195 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
2196 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
2197 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
2198 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
2199 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
2200 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
2201 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
2202 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
2203 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
2204 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
2205 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
2206 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
2207 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
2208 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
2209 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
2211 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
2212 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
2213 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
2214 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
2215 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
2216 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
2217 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
2218 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
2219 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
2220 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
2221 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
2222 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
2223 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
2224 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
2225 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
2226 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
2228 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
2229 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
2230 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
2231 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
2232 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
2233 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
2234 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
2235 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
2236 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
2237 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
2238 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
2239 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
2240 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
2241 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
2242 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
2243 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
2245 // sometimes (not rare at all) we need a third block :(
2247 if (remaining_bytes > 55)
2249 // this is for sure the final block
2261 w0_t[0] = esalt_buf1[ 8];
2262 w0_t[1] = esalt_buf1[ 9];
2263 w0_t[2] = esalt_buf1[10];
2264 w0_t[3] = esalt_buf1[11];
2266 w1_t[0] = esalt_buf1[12];
2267 w1_t[1] = esalt_buf1[13];
2268 w1_t[2] = esalt_buf1[14];
2269 w1_t[3] = esalt_buf1[15];
2271 w2_t[0] = esalt_buf2[ 0];
2272 w2_t[1] = esalt_buf2[ 1];
2273 w2_t[2] = esalt_buf2[ 2];
2274 w2_t[3] = esalt_buf2[ 3];
2276 w3_t[0] = esalt_buf2[ 4];
2277 w3_t[1] = esalt_buf2[ 5];
2278 w3_t[2] = digest_esalt_len * 8;
2281 MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
2282 MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
2283 MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
2284 MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
2285 MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
2286 MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
2287 MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
2288 MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
2289 MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
2290 MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
2291 MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
2292 MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
2293 MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
2294 MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
2295 MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
2296 MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
2298 MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
2299 MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
2300 MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
2301 MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
2302 MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
2303 MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
2304 MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
2305 MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
2306 MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
2307 MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
2308 MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
2309 MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
2310 MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
2311 MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
2312 MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
2313 MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
2315 MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
2316 MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
2317 MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
2318 MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
2319 MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
2320 MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
2321 MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
2322 MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
2323 MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
2324 MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
2325 MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
2326 MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
2327 MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
2328 MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
2329 MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
2330 MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
2332 MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
2333 MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
2334 MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
2335 MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
2336 MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
2337 MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
2338 MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
2339 MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
2340 MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
2341 MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
2342 MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
2343 MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
2344 MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
2345 MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
2346 MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
2347 MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
2364 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11400_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 sip_t *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)
2368 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11400_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 sip_t *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)