2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
16 #include "include/kernel_functions.c"
17 #include "OpenCL/types_ocl.c"
18 #include "OpenCL/common.c"
20 #define COMPARE_S "OpenCL/check_single_comp4.c"
21 #define COMPARE_M "OpenCL/check_multi_comp4.c"
23 __constant u32 k_sha256[64] =
25 SHA256C00, SHA256C01, SHA256C02, SHA256C03,
26 SHA256C04, SHA256C05, SHA256C06, SHA256C07,
27 SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
28 SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
29 SHA256C10, SHA256C11, SHA256C12, SHA256C13,
30 SHA256C14, SHA256C15, SHA256C16, SHA256C17,
31 SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
32 SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
33 SHA256C20, SHA256C21, SHA256C22, SHA256C23,
34 SHA256C24, SHA256C25, SHA256C26, SHA256C27,
35 SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
36 SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
37 SHA256C30, SHA256C31, SHA256C32, SHA256C33,
38 SHA256C34, SHA256C35, SHA256C36, SHA256C37,
39 SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
40 SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
45 void sha256_transform (const u32 w[16], u32 digest[8])
56 u32 w0_t = swap32 (w[ 0]);
57 u32 w1_t = swap32 (w[ 1]);
58 u32 w2_t = swap32 (w[ 2]);
59 u32 w3_t = swap32 (w[ 3]);
60 u32 w4_t = swap32 (w[ 4]);
61 u32 w5_t = swap32 (w[ 5]);
62 u32 w6_t = swap32 (w[ 6]);
63 u32 w7_t = swap32 (w[ 7]);
64 u32 w8_t = swap32 (w[ 8]);
65 u32 w9_t = swap32 (w[ 9]);
66 u32 wa_t = swap32 (w[10]);
67 u32 wb_t = swap32 (w[11]);
68 u32 wc_t = swap32 (w[12]);
69 u32 wd_t = swap32 (w[13]);
70 u32 we_t = swap32 (w[14]);
71 u32 wf_t = swap32 (w[15]);
73 #define ROUND_EXPAND() \
75 w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); \
76 w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); \
77 w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); \
78 w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); \
79 w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); \
80 w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); \
81 w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); \
82 w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); \
83 w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); \
84 w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); \
85 wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); \
86 wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); \
87 wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); \
88 wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); \
89 we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); \
90 wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); \
93 #define ROUND_STEP(i) \
95 SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i + 0]); \
96 SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i + 1]); \
97 SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i + 2]); \
98 SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i + 3]); \
99 SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i + 4]); \
100 SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i + 5]); \
101 SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i + 6]); \
102 SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i + 7]); \
103 SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i + 8]); \
104 SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i + 9]); \
105 SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
106 SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
107 SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
108 SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
109 SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
110 SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
116 for (int i = 16; i < 64; i += 16)
118 ROUND_EXPAND (); ROUND_STEP (i);
131 void sha256_transform_no14 (const u32 w[16], u32 digest[8])
152 sha256_transform (w_t, digest);
155 void init_ctx (u32 digest[8])
157 digest[0] = SHA256M_A;
158 digest[1] = SHA256M_B;
159 digest[2] = SHA256M_C;
160 digest[3] = SHA256M_D;
161 digest[4] = SHA256M_E;
162 digest[5] = SHA256M_F;
163 digest[6] = SHA256M_G;
164 digest[7] = SHA256M_H;
167 void bzero16 (u32 block[16])
187 void bswap8 (u32 block[16])
189 block[ 0] = swap32 (block[ 0]);
190 block[ 1] = swap32 (block[ 1]);
191 block[ 2] = swap32 (block[ 2]);
192 block[ 3] = swap32 (block[ 3]);
193 block[ 4] = swap32 (block[ 4]);
194 block[ 5] = swap32 (block[ 5]);
195 block[ 6] = swap32 (block[ 6]);
196 block[ 7] = swap32 (block[ 7]);
199 u32 memcat16 (u32 block[16], const u32 block_len, const u32 append[4], const u32 append_len)
201 const u32 mod = block_len & 3;
202 const u32 div = block_len / 4;
210 #if defined IS_AMD || defined IS_GENERIC
211 const int offset_minus_4 = 4 - block_len;
213 tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
214 tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
215 tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
216 tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
217 tmp4 = amd_bytealign ( 0, append[3], offset_minus_4);
230 const int offset_minus_4 = 4 - (block_len & 3);
232 const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
234 tmp0 = __byte_perm ( 0, append[0], selector);
235 tmp1 = __byte_perm (append[0], append[1], selector);
236 tmp2 = __byte_perm (append[1], append[2], selector);
237 tmp3 = __byte_perm (append[2], append[3], selector);
238 tmp4 = __byte_perm (append[3], 0, selector);
243 case 0: block[ 0] |= tmp0;
249 case 1: block[ 1] |= tmp0;
255 case 2: block[ 2] |= tmp0;
261 case 3: block[ 3] |= tmp0;
267 case 4: block[ 4] |= tmp0;
273 case 5: block[ 5] |= tmp0;
279 case 6: block[ 6] |= tmp0;
285 case 7: block[ 7] |= tmp0;
291 case 8: block[ 8] |= tmp0;
297 case 9: block[ 9] |= tmp0;
303 case 10: block[10] |= tmp0;
309 case 11: block[11] |= tmp0;
315 case 12: block[12] |= tmp0;
320 case 13: block[13] |= tmp0;
324 case 14: block[14] |= tmp0;
327 case 15: block[15] |= tmp0;
331 u32 new_len = block_len + append_len;
336 u32 memcat16c (u32 block[16], const u32 block_len, const u32 append[4], const u32 append_len, u32 digest[8])
338 const u32 mod = block_len & 3;
339 const u32 div = block_len / 4;
347 #if defined IS_AMD || defined IS_GENERIC
348 const int offset_minus_4 = 4 - block_len;
350 tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
351 tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
352 tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
353 tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
354 tmp4 = amd_bytealign ( 0, append[3], offset_minus_4);
367 const int offset_minus_4 = 4 - (block_len & 3);
369 const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
371 tmp0 = __byte_perm ( 0, append[0], selector);
372 tmp1 = __byte_perm (append[0], append[1], selector);
373 tmp2 = __byte_perm (append[1], append[2], selector);
374 tmp3 = __byte_perm (append[2], append[3], selector);
375 tmp4 = __byte_perm (append[3], 0, selector);
378 u32 carry[4] = { 0, 0, 0, 0 };
382 case 0: block[ 0] |= tmp0;
388 case 1: block[ 1] |= tmp0;
394 case 2: block[ 2] |= tmp0;
400 case 3: block[ 3] |= tmp0;
406 case 4: block[ 4] |= tmp0;
412 case 5: block[ 5] |= tmp0;
418 case 6: block[ 6] |= tmp0;
424 case 7: block[ 7] |= tmp0;
430 case 8: block[ 8] |= tmp0;
436 case 9: block[ 9] |= tmp0;
442 case 10: block[10] |= tmp0;
448 case 11: block[11] |= tmp0;
454 case 12: block[12] |= tmp0;
460 case 13: block[13] |= tmp0;
466 case 14: block[14] |= tmp0;
472 case 15: block[15] |= tmp0;
480 u32 new_len = block_len + append_len;
486 sha256_transform (block, digest);
499 u32 memcat20 (u32 block[20], const u32 block_len, const u32 append[4], const u32 append_len)
501 const u32 mod = block_len & 3;
502 const u32 div = block_len / 4;
510 #if defined IS_AMD || defined IS_GENERIC
511 const int offset_minus_4 = 4 - block_len;
513 tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
514 tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
515 tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
516 tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
517 tmp4 = amd_bytealign ( 0, append[3], offset_minus_4);
530 const int offset_minus_4 = 4 - (block_len & 3);
532 const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
534 tmp0 = __byte_perm ( 0, append[0], selector);
535 tmp1 = __byte_perm (append[0], append[1], selector);
536 tmp2 = __byte_perm (append[1], append[2], selector);
537 tmp3 = __byte_perm (append[2], append[3], selector);
538 tmp4 = __byte_perm (append[3], 0, selector);
543 case 0: block[ 0] |= tmp0;
549 case 1: block[ 1] |= tmp0;
555 case 2: block[ 2] |= tmp0;
561 case 3: block[ 3] |= tmp0;
567 case 4: block[ 4] |= tmp0;
573 case 5: block[ 5] |= tmp0;
579 case 6: block[ 6] |= tmp0;
585 case 7: block[ 7] |= tmp0;
591 case 8: block[ 8] |= tmp0;
597 case 9: block[ 9] |= tmp0;
603 case 10: block[10] |= tmp0;
609 case 11: block[11] |= tmp0;
615 case 12: block[12] |= tmp0;
621 case 13: block[13] |= tmp0;
627 case 14: block[14] |= tmp0;
633 case 15: block[15] |= tmp0;
641 return block_len + append_len;
644 u32 memcat20_x80 (u32 block[20], const u32 block_len, const u32 append[4], const u32 append_len)
646 const u32 mod = block_len & 3;
647 const u32 div = block_len / 4;
655 #if defined IS_AMD || defined IS_GENERIC
656 const int offset_minus_4 = 4 - block_len;
658 tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
659 tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
660 tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
661 tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
662 tmp4 = amd_bytealign ( 0x80, append[3], offset_minus_4);
675 const int offset_minus_4 = 4 - (block_len & 3);
677 const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
679 tmp0 = __byte_perm ( 0, append[0], selector);
680 tmp1 = __byte_perm (append[0], append[1], selector);
681 tmp2 = __byte_perm (append[1], append[2], selector);
682 tmp3 = __byte_perm (append[2], append[3], selector);
683 tmp4 = __byte_perm (append[3], 0x80, selector);
688 case 0: block[ 0] |= tmp0;
694 case 1: block[ 1] |= tmp0;
700 case 2: block[ 2] |= tmp0;
706 case 3: block[ 3] |= tmp0;
712 case 4: block[ 4] |= tmp0;
718 case 5: block[ 5] |= tmp0;
724 case 6: block[ 6] |= tmp0;
730 case 7: block[ 7] |= tmp0;
736 case 8: block[ 8] |= tmp0;
742 case 9: block[ 9] |= tmp0;
748 case 10: block[10] |= tmp0;
754 case 11: block[11] |= tmp0;
760 case 12: block[12] |= tmp0;
766 case 13: block[13] |= tmp0;
772 case 14: block[14] |= tmp0;
778 case 15: block[15] |= tmp0;
786 return block_len + append_len;
789 __kernel void m07400_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global sha256crypt_tmp_t *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)
795 const u32 gid = get_global_id (0);
797 if (gid >= gid_max) return;
801 w0[0] = pws[gid].i[0];
802 w0[1] = pws[gid].i[1];
803 w0[2] = pws[gid].i[2];
804 w0[3] = pws[gid].i[3];
806 const u32 pw_len = pws[gid].pw_len;
814 salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
815 salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
816 salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
817 salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
819 u32 salt_len = salt_bufs[salt_pos].salt_len;
825 u32 block_len; // never reaches > 64
826 u32 transform_len; // required for w[15] = len * 8
834 /* Prepare for the real work. */
842 block_len = memcat16 (block, block_len, w0, pw_len);
846 block_len = memcat16 (block, block_len, salt_buf, salt_len);
850 block_len = memcat16 (block, block_len, w0, pw_len);
852 append_0x80_1x16 (block, block_len);
854 block[15] = swap32 (block_len * 8);
856 init_ctx (alt_result);
858 sha256_transform (block, alt_result);
866 u32 alt_result_tmp[8];
868 alt_result_tmp[0] = alt_result[0];
869 alt_result_tmp[1] = alt_result[1];
870 alt_result_tmp[2] = alt_result[2];
871 alt_result_tmp[3] = alt_result[3];
872 alt_result_tmp[4] = 0;
873 alt_result_tmp[5] = 0;
874 alt_result_tmp[6] = 0;
875 alt_result_tmp[7] = 0;
877 truncate_block (alt_result_tmp, pw_len);
879 /* Add the key string. */
881 block_len = memcat16 (block, block_len, w0, pw_len);
883 /* The last part is the salt string. This must be at most 8
884 characters and it ends at the first `$' character (for
885 compatibility with existing implementations). */
887 block_len = memcat16 (block, block_len, salt_buf, salt_len);
889 /* Now get result of this (32 bytes) and add it to the other
892 block_len = memcat16 (block, block_len, alt_result_tmp, pw_len);
894 transform_len = block_len;
896 /* Take the binary representation of the length of the key and for every
897 1 add the alternate sum, for every 0 the key. */
899 alt_result_tmp[0] = alt_result[0];
900 alt_result_tmp[1] = alt_result[1];
901 alt_result_tmp[2] = alt_result[2];
902 alt_result_tmp[3] = alt_result[3];
903 alt_result_tmp[4] = alt_result[4];
904 alt_result_tmp[5] = alt_result[5];
905 alt_result_tmp[6] = alt_result[6];
906 alt_result_tmp[7] = alt_result[7];
908 init_ctx (alt_result);
910 for (u32 j = pw_len; j; j >>= 1)
914 block_len = memcat16c (block, block_len, &alt_result_tmp[0], 16, alt_result);
915 block_len = memcat16c (block, block_len, &alt_result_tmp[4], 16, alt_result);
921 block_len = memcat16c (block, block_len, w0, pw_len, alt_result);
923 transform_len += pw_len;
927 append_0x80_1x16 (block, block_len);
931 sha256_transform (block, alt_result);
936 block[15] = swap32 (transform_len * 8);
938 sha256_transform (block, alt_result);
942 tmps[gid].alt_result[0] = alt_result[0];
943 tmps[gid].alt_result[1] = alt_result[1];
944 tmps[gid].alt_result[2] = alt_result[2];
945 tmps[gid].alt_result[3] = alt_result[3];
946 tmps[gid].alt_result[4] = alt_result[4];
947 tmps[gid].alt_result[5] = alt_result[5];
948 tmps[gid].alt_result[6] = alt_result[6];
949 tmps[gid].alt_result[7] = alt_result[7];
951 /* Start computation of P byte sequence. */
959 /* For every character in the password add the entire password. */
963 for (u32 j = 0; j < pw_len; j++)
965 block_len = memcat16c (block, block_len, w0, pw_len, p_bytes);
967 transform_len += pw_len;
970 /* Finish the digest. */
972 append_0x80_1x16 (block, block_len);
976 sha256_transform (block, p_bytes);
981 block[15] = swap32 (transform_len * 8);
983 sha256_transform (block, p_bytes);
987 truncate_block (p_bytes, pw_len);
989 tmps[gid].p_bytes[0] = p_bytes[0];
990 tmps[gid].p_bytes[1] = p_bytes[1];
991 tmps[gid].p_bytes[2] = p_bytes[2];
992 tmps[gid].p_bytes[3] = p_bytes[3];
994 /* Start computation of S byte sequence. */
1002 /* For every character in the password add the entire password. */
1006 for (u32 j = 0; j < 16 + (alt_result[0] & 0xff); j++)
1008 block_len = memcat16c (block, block_len, salt_buf, salt_len, s_bytes);
1010 transform_len += salt_len;
1013 /* Finish the digest. */
1015 append_0x80_1x16 (block, block_len);
1017 if (block_len >= 56)
1019 sha256_transform (block, s_bytes);
1024 block[15] = swap32 (transform_len * 8);
1026 sha256_transform (block, s_bytes);
1030 truncate_block (s_bytes, salt_len);
1032 tmps[gid].s_bytes[0] = s_bytes[0];
1033 tmps[gid].s_bytes[1] = s_bytes[1];
1034 tmps[gid].s_bytes[2] = s_bytes[2];
1035 tmps[gid].s_bytes[3] = s_bytes[3];
1038 __kernel void m07400_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global sha256crypt_tmp_t *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)
1044 const u32 gid = get_global_id (0);
1046 if (gid >= gid_max) return;
1048 const u32 pw_len = pws[gid].pw_len;
1056 p_bytes[0] = tmps[gid].p_bytes[0];
1057 p_bytes[1] = tmps[gid].p_bytes[1];
1058 p_bytes[2] = tmps[gid].p_bytes[2];
1059 p_bytes[3] = tmps[gid].p_bytes[3];
1063 p_bytes_x80[0] = tmps[gid].p_bytes[0];
1064 p_bytes_x80[1] = tmps[gid].p_bytes[1];
1065 p_bytes_x80[2] = tmps[gid].p_bytes[2];
1066 p_bytes_x80[3] = tmps[gid].p_bytes[3];
1068 append_0x80_1x4 (p_bytes_x80, pw_len);
1072 s_bytes[0] = tmps[gid].s_bytes[0];
1073 s_bytes[1] = tmps[gid].s_bytes[1];
1074 s_bytes[2] = tmps[gid].s_bytes[2];
1075 s_bytes[3] = tmps[gid].s_bytes[3];
1079 alt_result[0] = tmps[gid].alt_result[0];
1080 alt_result[1] = tmps[gid].alt_result[1];
1081 alt_result[2] = tmps[gid].alt_result[2];
1082 alt_result[3] = tmps[gid].alt_result[3];
1083 alt_result[4] = tmps[gid].alt_result[4];
1084 alt_result[5] = tmps[gid].alt_result[5];
1085 alt_result[6] = tmps[gid].alt_result[6];
1086 alt_result[7] = tmps[gid].alt_result[7];
1088 u32 salt_len = salt_bufs[salt_pos].salt_len;
1090 /* Repeatedly run the collected hash value through SHA256 to burn
1093 for (u32 i = 0, j = loop_pos; i < loop_cnt; i++, j++)
1101 bzero16 (&block[ 0]);
1102 bzero16 (&block[16]);
1106 const u32 j1 = (j & 1) ? 1 : 0;
1107 const u32 j3 = (j % 3) ? 1 : 0;
1108 const u32 j7 = (j % 7) ? 1 : 0;
1112 block[0] = p_bytes[0];
1113 block[1] = p_bytes[1];
1114 block[2] = p_bytes[2];
1115 block[3] = p_bytes[3];
1121 block[0] = alt_result[0];
1122 block[1] = alt_result[1];
1123 block[2] = alt_result[2];
1124 block[3] = alt_result[3];
1125 block[4] = alt_result[4];
1126 block[5] = alt_result[5];
1127 block[6] = alt_result[6];
1128 block[7] = alt_result[7];
1135 block_len = memcat20 (block, block_len, s_bytes, salt_len);
1140 block_len = memcat20 (block, block_len, p_bytes, pw_len);
1145 block_len = memcat20 (block, block_len, &alt_result[0], 16);
1146 block_len = memcat20_x80 (block, block_len, &alt_result[4], 16);
1150 block_len = memcat20 (block, block_len, p_bytes_x80, pw_len);
1153 if (block_len >= 56)
1155 sha256_transform (block, tmp);
1157 block[ 0] = block[16];
1158 block[ 1] = block[17];
1159 block[ 2] = block[18];
1160 block[ 3] = block[19];
1175 block[15] = swap32 (block_len * 8);
1177 sha256_transform_no14 (block, tmp);
1181 alt_result[0] = tmp[0];
1182 alt_result[1] = tmp[1];
1183 alt_result[2] = tmp[2];
1184 alt_result[3] = tmp[3];
1185 alt_result[4] = tmp[4];
1186 alt_result[5] = tmp[5];
1187 alt_result[6] = tmp[6];
1188 alt_result[7] = tmp[7];
1191 tmps[gid].alt_result[0] = alt_result[0];
1192 tmps[gid].alt_result[1] = alt_result[1];
1193 tmps[gid].alt_result[2] = alt_result[2];
1194 tmps[gid].alt_result[3] = alt_result[3];
1195 tmps[gid].alt_result[4] = alt_result[4];
1196 tmps[gid].alt_result[5] = alt_result[5];
1197 tmps[gid].alt_result[6] = alt_result[6];
1198 tmps[gid].alt_result[7] = alt_result[7];
1201 __kernel void m07400_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global sha256crypt_tmp_t *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)
1207 const u32 gid = get_global_id (0);
1209 if (gid >= gid_max) return;
1211 const u32 lid = get_local_id (0);
1213 const u32 r0 = tmps[gid].alt_result[0];
1214 const u32 r1 = tmps[gid].alt_result[1];
1215 const u32 r2 = tmps[gid].alt_result[2];
1216 const u32 r3 = tmps[gid].alt_result[3];
1225 #define PUTCHAR32_BE(a,p,c) ((u8 *)(a))[(p) ^ 3] = (u8) (c)
1226 #define GETCHAR32_BE(a,p) ((u8 *)(a))[(p) ^ 3]
1236 void sha256_transform (const u32 w[16], u32 digest[8])
1264 #define ROUND_EXPAND() \
1266 w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); \
1267 w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); \
1268 w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); \
1269 w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); \
1270 w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); \
1271 w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); \
1272 w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); \
1273 w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); \
1274 w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); \
1275 w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); \
1276 wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); \
1277 wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); \
1278 wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); \
1279 wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); \
1280 we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); \
1281 wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); \
1284 #define ROUND_STEP(i) \
1286 SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i + 0]); \
1287 SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i + 1]); \
1288 SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i + 2]); \
1289 SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i + 3]); \
1290 SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i + 4]); \
1291 SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i + 5]); \
1292 SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i + 6]); \
1293 SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i + 7]); \
1294 SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i + 8]); \
1295 SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i + 9]); \
1296 SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
1297 SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
1298 SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
1299 SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
1300 SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
1301 SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
1307 for (int i = 16; i < 64; i += 16)
1309 ROUND_EXPAND (); ROUND_STEP (i);
1322 void sha256_init (sha256_ctx_t *sha256_ctx)
1324 sha256_ctx->state[0] = SHA256M_A;
1325 sha256_ctx->state[1] = SHA256M_B;
1326 sha256_ctx->state[2] = SHA256M_C;
1327 sha256_ctx->state[3] = SHA256M_D;
1328 sha256_ctx->state[4] = SHA256M_E;
1329 sha256_ctx->state[5] = SHA256M_F;
1330 sha256_ctx->state[6] = SHA256M_G;
1331 sha256_ctx->state[7] = SHA256M_H;
1333 sha256_ctx->len = 0;
1336 void sha256_update (sha256_ctx_t *sha256_ctx, const u32 *buf, int len)
1338 int pos = sha256_ctx->len & 0x3f;
1340 sha256_ctx->len += len;
1342 if ((pos + len) < 64)
1344 for (int i = 0; i < len; i++)
1346 PUTCHAR32_BE (sha256_ctx->buf, pos++, GETCHAR32_BE (buf, i));
1354 for (int i = 0; i < cnt; i++)
1356 PUTCHAR32_BE (sha256_ctx->buf, pos++, GETCHAR32_BE (buf, i));
1359 sha256_transform (sha256_ctx->buf, sha256_ctx->state);
1363 for (int i = 0; i < len; i++)
1365 PUTCHAR32_BE (sha256_ctx->buf, i, GETCHAR32_BE (buf, cnt + i));
1369 void sha256_final (sha256_ctx_t *sha256_ctx)
1371 int pos = sha256_ctx->len & 0x3f;
1373 for (int i = pos; i < 64; i++)
1375 PUTCHAR32_BE (sha256_ctx->buf, i, 0);
1378 PUTCHAR32_BE (sha256_ctx->buf, pos, 0x80);
1382 sha256_transform (sha256_ctx->buf, sha256_ctx->state);
1384 sha256_ctx->buf[ 0] = 0;
1385 sha256_ctx->buf[ 1] = 0;
1386 sha256_ctx->buf[ 2] = 0;
1387 sha256_ctx->buf[ 3] = 0;
1388 sha256_ctx->buf[ 4] = 0;
1389 sha256_ctx->buf[ 5] = 0;
1390 sha256_ctx->buf[ 6] = 0;
1391 sha256_ctx->buf[ 7] = 0;
1392 sha256_ctx->buf[ 8] = 0;
1393 sha256_ctx->buf[ 9] = 0;
1394 sha256_ctx->buf[10] = 0;
1395 sha256_ctx->buf[11] = 0;
1396 sha256_ctx->buf[12] = 0;
1397 sha256_ctx->buf[13] = 0;
1398 sha256_ctx->buf[14] = 0;
1399 sha256_ctx->buf[15] = 0;
1402 sha256_ctx->buf[15] = sha256_ctx->len * 8;
1404 sha256_transform (sha256_ctx->buf, sha256_ctx->state);
1407 __kernel void m07400_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global sha256crypt_tmp_t *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)
1413 const u32 gid = get_global_id (0);
1415 if (gid >= gid_max) return;
1419 pw[0] = swap32 (pws[gid].i[0]);
1420 pw[1] = swap32 (pws[gid].i[1]);
1421 pw[2] = swap32 (pws[gid].i[2]);
1422 pw[3] = swap32 (pws[gid].i[3]);
1424 const u32 pw_len = pws[gid].pw_len;
1432 salt[0] = swap32 (salt_bufs[salt_pos].salt_buf[0]);
1433 salt[1] = swap32 (salt_bufs[salt_pos].salt_buf[1]);
1434 salt[2] = swap32 (salt_bufs[salt_pos].salt_buf[2]);
1435 salt[3] = swap32 (salt_bufs[salt_pos].salt_buf[3]);
1437 u32 salt_len = salt_bufs[salt_pos].salt_len;
1443 sha256_ctx_t sha256_ctx;
1445 sha256_init (&sha256_ctx);
1447 sha256_update (&sha256_ctx, pw, pw_len);
1448 sha256_update (&sha256_ctx, salt, salt_len);
1449 sha256_update (&sha256_ctx, pw, pw_len);
1451 sha256_final (&sha256_ctx);
1455 tmp[0] = sha256_ctx.state[0];
1456 tmp[1] = sha256_ctx.state[1];
1457 tmp[2] = sha256_ctx.state[2];
1458 tmp[3] = sha256_ctx.state[3];
1459 tmp[4] = sha256_ctx.state[4];
1460 tmp[5] = sha256_ctx.state[5];
1461 tmp[6] = sha256_ctx.state[6];
1462 tmp[7] = sha256_ctx.state[7];
1464 sha256_init (&sha256_ctx);
1466 sha256_update (&sha256_ctx, pw, pw_len);
1467 sha256_update (&sha256_ctx, salt, salt_len);
1468 sha256_update (&sha256_ctx, tmp, pw_len);
1470 for (u32 j = pw_len; j; j >>= 1)
1474 sha256_update (&sha256_ctx, tmp, 32);
1478 sha256_update (&sha256_ctx, pw, pw_len);
1482 sha256_final (&sha256_ctx);
1484 tmps[gid].alt_result[0] = sha256_ctx.state[0];
1485 tmps[gid].alt_result[1] = sha256_ctx.state[1];
1486 tmps[gid].alt_result[2] = sha256_ctx.state[2];
1487 tmps[gid].alt_result[3] = sha256_ctx.state[3];
1488 tmps[gid].alt_result[4] = sha256_ctx.state[4];
1489 tmps[gid].alt_result[5] = sha256_ctx.state[5];
1490 tmps[gid].alt_result[6] = sha256_ctx.state[6];
1491 tmps[gid].alt_result[7] = sha256_ctx.state[7];
1495 sha256_init (&sha256_ctx);
1497 for (u32 j = 0; j < pw_len; j++)
1499 sha256_update (&sha256_ctx, pw, pw_len);
1502 sha256_final (&sha256_ctx);
1504 tmps[gid].p_bytes[0] = sha256_ctx.state[0];
1505 tmps[gid].p_bytes[1] = sha256_ctx.state[1];
1506 tmps[gid].p_bytes[2] = sha256_ctx.state[2];
1507 tmps[gid].p_bytes[3] = sha256_ctx.state[3];
1511 sha256_init (&sha256_ctx);
1513 for (u32 j = 0; j < 16 + ((tmps[gid].alt_result[0] >> 24) & 0xff); j++)
1515 sha256_update (&sha256_ctx, salt, salt_len);
1518 sha256_final (&sha256_ctx);
1520 tmps[gid].s_bytes[0] = sha256_ctx.state[0];
1521 tmps[gid].s_bytes[1] = sha256_ctx.state[1];
1522 tmps[gid].s_bytes[2] = sha256_ctx.state[2];
1523 tmps[gid].s_bytes[3] = sha256_ctx.state[3];
1526 __kernel void m07400_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global sha256crypt_tmp_t *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)
1532 const u32 gid = get_global_id (0);
1534 if (gid >= gid_max) return;
1538 p_bytes0[0] = tmps[gid].p_bytes[0];
1539 p_bytes0[1] = tmps[gid].p_bytes[1];
1540 p_bytes0[2] = tmps[gid].p_bytes[2];
1541 p_bytes0[3] = tmps[gid].p_bytes[3];
1543 const u32 pw_len = pws[gid].pw_len;
1547 s_bytes0[0] = tmps[gid].s_bytes[0];
1548 s_bytes0[1] = tmps[gid].s_bytes[1];
1549 s_bytes0[2] = tmps[gid].s_bytes[2];
1550 s_bytes0[3] = tmps[gid].s_bytes[3];
1552 const u32 salt_len = salt_bufs[salt_pos].salt_len;
1556 wpc_len[0] = 32 + 0 + 0 + pw_len;
1557 wpc_len[1] = pw_len + 0 + 0 + 32;
1558 wpc_len[2] = 32 + salt_len + 0 + pw_len;
1559 wpc_len[3] = pw_len + salt_len + 0 + 32;
1560 wpc_len[4] = 32 + 0 + pw_len + pw_len;
1561 wpc_len[5] = pw_len + 0 + pw_len + 32;
1562 wpc_len[6] = 32 + salt_len + pw_len + pw_len;
1563 wpc_len[7] = pw_len + salt_len + pw_len + 32;
1565 u32 wpc[8][32] = { { 0 } };
1567 for (u32 i = 0; i < 8; i++)
1573 for (u32 j = 0; j < pw_len; j++)
1575 PUTCHAR32_BE (wpc[i], block_len++, GETCHAR32_BE (p_bytes0, j));
1585 for (u32 j = 0; j < salt_len; j++)
1587 PUTCHAR32_BE (wpc[i], block_len++, GETCHAR32_BE (s_bytes0, j));
1593 for (u32 j = 0; j < pw_len; j++)
1595 PUTCHAR32_BE (wpc[i], block_len++, GETCHAR32_BE (p_bytes0, j));
1605 for (u32 j = 0; j < pw_len; j++)
1607 PUTCHAR32_BE (wpc[i], block_len++, GETCHAR32_BE (p_bytes0, j));
1611 PUTCHAR32_BE (wpc[i], block_len, 0x80);
1615 wpc[i][15] = block_len * 8;
1619 wpc[i][31] = block_len * 8;
1629 alt_result[0] = tmps[gid].alt_result[0];
1630 alt_result[1] = tmps[gid].alt_result[1];
1631 alt_result[2] = tmps[gid].alt_result[2];
1632 alt_result[3] = tmps[gid].alt_result[3];
1633 alt_result[4] = tmps[gid].alt_result[4];
1634 alt_result[5] = tmps[gid].alt_result[5];
1635 alt_result[6] = tmps[gid].alt_result[6];
1636 alt_result[7] = tmps[gid].alt_result[7];
1639 /* Repeatedly run the collected hash value through SHA256 to burn
1642 for (u32 i = 0, j = loop_pos; i < loop_cnt; i++, j++)
1644 const u32 j1 = (j & 1) ? 1 : 0;
1645 const u32 j3 = (j % 3) ? 2 : 0;
1646 const u32 j7 = (j % 7) ? 4 : 0;
1648 const u32 pc = j1 + j3 + j7;
1652 block[ 0] = wpc[pc][ 0];
1653 block[ 1] = wpc[pc][ 1];
1654 block[ 2] = wpc[pc][ 2];
1655 block[ 3] = wpc[pc][ 3];
1656 block[ 4] = wpc[pc][ 4];
1657 block[ 5] = wpc[pc][ 5];
1658 block[ 6] = wpc[pc][ 6];
1659 block[ 7] = wpc[pc][ 7];
1660 block[ 8] = wpc[pc][ 8];
1661 block[ 9] = wpc[pc][ 9];
1662 block[10] = wpc[pc][10];
1663 block[11] = wpc[pc][11];
1664 block[12] = wpc[pc][12];
1665 block[13] = wpc[pc][13];
1666 block[14] = wpc[pc][14];
1667 block[15] = wpc[pc][15];
1668 block[16] = wpc[pc][16];
1669 block[17] = wpc[pc][17];
1670 block[18] = wpc[pc][18];
1671 block[19] = wpc[pc][19];
1672 block[20] = wpc[pc][20];
1673 block[21] = wpc[pc][21];
1674 block[22] = wpc[pc][22];
1675 block[23] = wpc[pc][23];
1676 block[24] = wpc[pc][24];
1677 block[25] = wpc[pc][25];
1678 block[26] = wpc[pc][26];
1679 block[27] = wpc[pc][27];
1680 block[28] = wpc[pc][28];
1681 block[29] = wpc[pc][29];
1682 block[30] = wpc[pc][30];
1683 block[31] = wpc[pc][31];
1685 const u32 block_len = wpc_len[pc];
1690 for (u32 k = 0, p = block_len - 32; k < 32; k++, p++)
1692 PUTCHAR32_BE (block, p, GETCHAR32_BE (alt_result, k));
1697 block[0] = alt_result[0];
1698 block[1] = alt_result[1];
1699 block[2] = alt_result[2];
1700 block[3] = alt_result[3];
1701 block[4] = alt_result[4];
1702 block[5] = alt_result[5];
1703 block[6] = alt_result[6];
1704 block[7] = alt_result[7];
1707 alt_result[0] = SHA256M_A;
1708 alt_result[1] = SHA256M_B;
1709 alt_result[2] = SHA256M_C;
1710 alt_result[3] = SHA256M_D;
1711 alt_result[4] = SHA256M_E;
1712 alt_result[5] = SHA256M_F;
1713 alt_result[6] = SHA256M_G;
1714 alt_result[7] = SHA256M_H;
1716 sha256_transform (block, alt_result);
1718 if (block_len >= 56)
1720 sha256_transform (block + 16, alt_result);
1724 tmps[gid].alt_result[0] = alt_result[0];
1725 tmps[gid].alt_result[1] = alt_result[1];
1726 tmps[gid].alt_result[2] = alt_result[2];
1727 tmps[gid].alt_result[3] = alt_result[3];
1728 tmps[gid].alt_result[4] = alt_result[4];
1729 tmps[gid].alt_result[5] = alt_result[5];
1730 tmps[gid].alt_result[6] = alt_result[6];
1731 tmps[gid].alt_result[7] = alt_result[7];
1734 __kernel void m07400_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global sha256crypt_tmp_t *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)
1740 const u32 gid = get_global_id (0);
1742 if (gid >= gid_max) return;
1744 const u32 lid = get_local_id (0);
1746 const u32 r0 = swap32 (tmps[gid].alt_result[0]);
1747 const u32 r1 = swap32 (tmps[gid].alt_result[1]);
1748 const u32 r2 = swap32 (tmps[gid].alt_result[2]);
1749 const u32 r3 = swap32 (tmps[gid].alt_result[3]);