2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
4 * magnum <john.magnum@hushmail.com>
13 #include "inc_vendor.cl"
14 #include "inc_hash_constants.h"
15 #include "inc_hash_functions.cl"
16 #include "inc_types.cl"
17 #include "inc_common.cl"
18 #include "inc_simd.cl"
20 __constant u32 c_tables[4][256] =
23 0x00072000, 0x00075000, 0x00074800, 0x00071000,
24 0x00076800, 0x00074000, 0x00070000, 0x00077000,
25 0x00073000, 0x00075800, 0x00070800, 0x00076000,
26 0x00073800, 0x00077800, 0x00072800, 0x00071800,
27 0x0005a000, 0x0005d000, 0x0005c800, 0x00059000,
28 0x0005e800, 0x0005c000, 0x00058000, 0x0005f000,
29 0x0005b000, 0x0005d800, 0x00058800, 0x0005e000,
30 0x0005b800, 0x0005f800, 0x0005a800, 0x00059800,
31 0x00022000, 0x00025000, 0x00024800, 0x00021000,
32 0x00026800, 0x00024000, 0x00020000, 0x00027000,
33 0x00023000, 0x00025800, 0x00020800, 0x00026000,
34 0x00023800, 0x00027800, 0x00022800, 0x00021800,
35 0x00062000, 0x00065000, 0x00064800, 0x00061000,
36 0x00066800, 0x00064000, 0x00060000, 0x00067000,
37 0x00063000, 0x00065800, 0x00060800, 0x00066000,
38 0x00063800, 0x00067800, 0x00062800, 0x00061800,
39 0x00032000, 0x00035000, 0x00034800, 0x00031000,
40 0x00036800, 0x00034000, 0x00030000, 0x00037000,
41 0x00033000, 0x00035800, 0x00030800, 0x00036000,
42 0x00033800, 0x00037800, 0x00032800, 0x00031800,
43 0x0006a000, 0x0006d000, 0x0006c800, 0x00069000,
44 0x0006e800, 0x0006c000, 0x00068000, 0x0006f000,
45 0x0006b000, 0x0006d800, 0x00068800, 0x0006e000,
46 0x0006b800, 0x0006f800, 0x0006a800, 0x00069800,
47 0x0007a000, 0x0007d000, 0x0007c800, 0x00079000,
48 0x0007e800, 0x0007c000, 0x00078000, 0x0007f000,
49 0x0007b000, 0x0007d800, 0x00078800, 0x0007e000,
50 0x0007b800, 0x0007f800, 0x0007a800, 0x00079800,
51 0x00052000, 0x00055000, 0x00054800, 0x00051000,
52 0x00056800, 0x00054000, 0x00050000, 0x00057000,
53 0x00053000, 0x00055800, 0x00050800, 0x00056000,
54 0x00053800, 0x00057800, 0x00052800, 0x00051800,
55 0x00012000, 0x00015000, 0x00014800, 0x00011000,
56 0x00016800, 0x00014000, 0x00010000, 0x00017000,
57 0x00013000, 0x00015800, 0x00010800, 0x00016000,
58 0x00013800, 0x00017800, 0x00012800, 0x00011800,
59 0x0001a000, 0x0001d000, 0x0001c800, 0x00019000,
60 0x0001e800, 0x0001c000, 0x00018000, 0x0001f000,
61 0x0001b000, 0x0001d800, 0x00018800, 0x0001e000,
62 0x0001b800, 0x0001f800, 0x0001a800, 0x00019800,
63 0x00042000, 0x00045000, 0x00044800, 0x00041000,
64 0x00046800, 0x00044000, 0x00040000, 0x00047000,
65 0x00043000, 0x00045800, 0x00040800, 0x00046000,
66 0x00043800, 0x00047800, 0x00042800, 0x00041800,
67 0x0000a000, 0x0000d000, 0x0000c800, 0x00009000,
68 0x0000e800, 0x0000c000, 0x00008000, 0x0000f000,
69 0x0000b000, 0x0000d800, 0x00008800, 0x0000e000,
70 0x0000b800, 0x0000f800, 0x0000a800, 0x00009800,
71 0x00002000, 0x00005000, 0x00004800, 0x00001000,
72 0x00006800, 0x00004000, 0x00000000, 0x00007000,
73 0x00003000, 0x00005800, 0x00000800, 0x00006000,
74 0x00003800, 0x00007800, 0x00002800, 0x00001800,
75 0x0003a000, 0x0003d000, 0x0003c800, 0x00039000,
76 0x0003e800, 0x0003c000, 0x00038000, 0x0003f000,
77 0x0003b000, 0x0003d800, 0x00038800, 0x0003e000,
78 0x0003b800, 0x0003f800, 0x0003a800, 0x00039800,
79 0x0002a000, 0x0002d000, 0x0002c800, 0x00029000,
80 0x0002e800, 0x0002c000, 0x00028000, 0x0002f000,
81 0x0002b000, 0x0002d800, 0x00028800, 0x0002e000,
82 0x0002b800, 0x0002f800, 0x0002a800, 0x00029800,
83 0x0004a000, 0x0004d000, 0x0004c800, 0x00049000,
84 0x0004e800, 0x0004c000, 0x00048000, 0x0004f000,
85 0x0004b000, 0x0004d800, 0x00048800, 0x0004e000,
86 0x0004b800, 0x0004f800, 0x0004a800, 0x00049800,
89 0x03a80000, 0x03c00000, 0x03880000, 0x03e80000,
90 0x03d00000, 0x03980000, 0x03a00000, 0x03900000,
91 0x03f00000, 0x03f80000, 0x03e00000, 0x03b80000,
92 0x03b00000, 0x03800000, 0x03c80000, 0x03d80000,
93 0x06a80000, 0x06c00000, 0x06880000, 0x06e80000,
94 0x06d00000, 0x06980000, 0x06a00000, 0x06900000,
95 0x06f00000, 0x06f80000, 0x06e00000, 0x06b80000,
96 0x06b00000, 0x06800000, 0x06c80000, 0x06d80000,
97 0x05280000, 0x05400000, 0x05080000, 0x05680000,
98 0x05500000, 0x05180000, 0x05200000, 0x05100000,
99 0x05700000, 0x05780000, 0x05600000, 0x05380000,
100 0x05300000, 0x05000000, 0x05480000, 0x05580000,
101 0x00a80000, 0x00c00000, 0x00880000, 0x00e80000,
102 0x00d00000, 0x00980000, 0x00a00000, 0x00900000,
103 0x00f00000, 0x00f80000, 0x00e00000, 0x00b80000,
104 0x00b00000, 0x00800000, 0x00c80000, 0x00d80000,
105 0x00280000, 0x00400000, 0x00080000, 0x00680000,
106 0x00500000, 0x00180000, 0x00200000, 0x00100000,
107 0x00700000, 0x00780000, 0x00600000, 0x00380000,
108 0x00300000, 0x00000000, 0x00480000, 0x00580000,
109 0x04280000, 0x04400000, 0x04080000, 0x04680000,
110 0x04500000, 0x04180000, 0x04200000, 0x04100000,
111 0x04700000, 0x04780000, 0x04600000, 0x04380000,
112 0x04300000, 0x04000000, 0x04480000, 0x04580000,
113 0x04a80000, 0x04c00000, 0x04880000, 0x04e80000,
114 0x04d00000, 0x04980000, 0x04a00000, 0x04900000,
115 0x04f00000, 0x04f80000, 0x04e00000, 0x04b80000,
116 0x04b00000, 0x04800000, 0x04c80000, 0x04d80000,
117 0x07a80000, 0x07c00000, 0x07880000, 0x07e80000,
118 0x07d00000, 0x07980000, 0x07a00000, 0x07900000,
119 0x07f00000, 0x07f80000, 0x07e00000, 0x07b80000,
120 0x07b00000, 0x07800000, 0x07c80000, 0x07d80000,
121 0x07280000, 0x07400000, 0x07080000, 0x07680000,
122 0x07500000, 0x07180000, 0x07200000, 0x07100000,
123 0x07700000, 0x07780000, 0x07600000, 0x07380000,
124 0x07300000, 0x07000000, 0x07480000, 0x07580000,
125 0x02280000, 0x02400000, 0x02080000, 0x02680000,
126 0x02500000, 0x02180000, 0x02200000, 0x02100000,
127 0x02700000, 0x02780000, 0x02600000, 0x02380000,
128 0x02300000, 0x02000000, 0x02480000, 0x02580000,
129 0x03280000, 0x03400000, 0x03080000, 0x03680000,
130 0x03500000, 0x03180000, 0x03200000, 0x03100000,
131 0x03700000, 0x03780000, 0x03600000, 0x03380000,
132 0x03300000, 0x03000000, 0x03480000, 0x03580000,
133 0x06280000, 0x06400000, 0x06080000, 0x06680000,
134 0x06500000, 0x06180000, 0x06200000, 0x06100000,
135 0x06700000, 0x06780000, 0x06600000, 0x06380000,
136 0x06300000, 0x06000000, 0x06480000, 0x06580000,
137 0x05a80000, 0x05c00000, 0x05880000, 0x05e80000,
138 0x05d00000, 0x05980000, 0x05a00000, 0x05900000,
139 0x05f00000, 0x05f80000, 0x05e00000, 0x05b80000,
140 0x05b00000, 0x05800000, 0x05c80000, 0x05d80000,
141 0x01280000, 0x01400000, 0x01080000, 0x01680000,
142 0x01500000, 0x01180000, 0x01200000, 0x01100000,
143 0x01700000, 0x01780000, 0x01600000, 0x01380000,
144 0x01300000, 0x01000000, 0x01480000, 0x01580000,
145 0x02a80000, 0x02c00000, 0x02880000, 0x02e80000,
146 0x02d00000, 0x02980000, 0x02a00000, 0x02900000,
147 0x02f00000, 0x02f80000, 0x02e00000, 0x02b80000,
148 0x02b00000, 0x02800000, 0x02c80000, 0x02d80000,
149 0x01a80000, 0x01c00000, 0x01880000, 0x01e80000,
150 0x01d00000, 0x01980000, 0x01a00000, 0x01900000,
151 0x01f00000, 0x01f80000, 0x01e00000, 0x01b80000,
152 0x01b00000, 0x01800000, 0x01c80000, 0x01d80000,
155 0x30000002, 0x60000002, 0x38000002, 0x08000002,
156 0x28000002, 0x78000002, 0x68000002, 0x40000002,
157 0x20000002, 0x50000002, 0x48000002, 0x70000002,
158 0x00000002, 0x18000002, 0x58000002, 0x10000002,
159 0xb0000005, 0xe0000005, 0xb8000005, 0x88000005,
160 0xa8000005, 0xf8000005, 0xe8000005, 0xc0000005,
161 0xa0000005, 0xd0000005, 0xc8000005, 0xf0000005,
162 0x80000005, 0x98000005, 0xd8000005, 0x90000005,
163 0x30000005, 0x60000005, 0x38000005, 0x08000005,
164 0x28000005, 0x78000005, 0x68000005, 0x40000005,
165 0x20000005, 0x50000005, 0x48000005, 0x70000005,
166 0x00000005, 0x18000005, 0x58000005, 0x10000005,
167 0x30000000, 0x60000000, 0x38000000, 0x08000000,
168 0x28000000, 0x78000000, 0x68000000, 0x40000000,
169 0x20000000, 0x50000000, 0x48000000, 0x70000000,
170 0x00000000, 0x18000000, 0x58000000, 0x10000000,
171 0xb0000003, 0xe0000003, 0xb8000003, 0x88000003,
172 0xa8000003, 0xf8000003, 0xe8000003, 0xc0000003,
173 0xa0000003, 0xd0000003, 0xc8000003, 0xf0000003,
174 0x80000003, 0x98000003, 0xd8000003, 0x90000003,
175 0x30000001, 0x60000001, 0x38000001, 0x08000001,
176 0x28000001, 0x78000001, 0x68000001, 0x40000001,
177 0x20000001, 0x50000001, 0x48000001, 0x70000001,
178 0x00000001, 0x18000001, 0x58000001, 0x10000001,
179 0xb0000000, 0xe0000000, 0xb8000000, 0x88000000,
180 0xa8000000, 0xf8000000, 0xe8000000, 0xc0000000,
181 0xa0000000, 0xd0000000, 0xc8000000, 0xf0000000,
182 0x80000000, 0x98000000, 0xd8000000, 0x90000000,
183 0xb0000006, 0xe0000006, 0xb8000006, 0x88000006,
184 0xa8000006, 0xf8000006, 0xe8000006, 0xc0000006,
185 0xa0000006, 0xd0000006, 0xc8000006, 0xf0000006,
186 0x80000006, 0x98000006, 0xd8000006, 0x90000006,
187 0xb0000001, 0xe0000001, 0xb8000001, 0x88000001,
188 0xa8000001, 0xf8000001, 0xe8000001, 0xc0000001,
189 0xa0000001, 0xd0000001, 0xc8000001, 0xf0000001,
190 0x80000001, 0x98000001, 0xd8000001, 0x90000001,
191 0x30000003, 0x60000003, 0x38000003, 0x08000003,
192 0x28000003, 0x78000003, 0x68000003, 0x40000003,
193 0x20000003, 0x50000003, 0x48000003, 0x70000003,
194 0x00000003, 0x18000003, 0x58000003, 0x10000003,
195 0x30000004, 0x60000004, 0x38000004, 0x08000004,
196 0x28000004, 0x78000004, 0x68000004, 0x40000004,
197 0x20000004, 0x50000004, 0x48000004, 0x70000004,
198 0x00000004, 0x18000004, 0x58000004, 0x10000004,
199 0xb0000002, 0xe0000002, 0xb8000002, 0x88000002,
200 0xa8000002, 0xf8000002, 0xe8000002, 0xc0000002,
201 0xa0000002, 0xd0000002, 0xc8000002, 0xf0000002,
202 0x80000002, 0x98000002, 0xd8000002, 0x90000002,
203 0xb0000004, 0xe0000004, 0xb8000004, 0x88000004,
204 0xa8000004, 0xf8000004, 0xe8000004, 0xc0000004,
205 0xa0000004, 0xd0000004, 0xc8000004, 0xf0000004,
206 0x80000004, 0x98000004, 0xd8000004, 0x90000004,
207 0x30000006, 0x60000006, 0x38000006, 0x08000006,
208 0x28000006, 0x78000006, 0x68000006, 0x40000006,
209 0x20000006, 0x50000006, 0x48000006, 0x70000006,
210 0x00000006, 0x18000006, 0x58000006, 0x10000006,
211 0xb0000007, 0xe0000007, 0xb8000007, 0x88000007,
212 0xa8000007, 0xf8000007, 0xe8000007, 0xc0000007,
213 0xa0000007, 0xd0000007, 0xc8000007, 0xf0000007,
214 0x80000007, 0x98000007, 0xd8000007, 0x90000007,
215 0x30000007, 0x60000007, 0x38000007, 0x08000007,
216 0x28000007, 0x78000007, 0x68000007, 0x40000007,
217 0x20000007, 0x50000007, 0x48000007, 0x70000007,
218 0x00000007, 0x18000007, 0x58000007, 0x10000007,
221 0x000000e8, 0x000000d8, 0x000000a0, 0x00000088,
222 0x00000098, 0x000000f8, 0x000000a8, 0x000000c8,
223 0x00000080, 0x000000d0, 0x000000f0, 0x000000b8,
224 0x000000b0, 0x000000c0, 0x00000090, 0x000000e0,
225 0x000007e8, 0x000007d8, 0x000007a0, 0x00000788,
226 0x00000798, 0x000007f8, 0x000007a8, 0x000007c8,
227 0x00000780, 0x000007d0, 0x000007f0, 0x000007b8,
228 0x000007b0, 0x000007c0, 0x00000790, 0x000007e0,
229 0x000006e8, 0x000006d8, 0x000006a0, 0x00000688,
230 0x00000698, 0x000006f8, 0x000006a8, 0x000006c8,
231 0x00000680, 0x000006d0, 0x000006f0, 0x000006b8,
232 0x000006b0, 0x000006c0, 0x00000690, 0x000006e0,
233 0x00000068, 0x00000058, 0x00000020, 0x00000008,
234 0x00000018, 0x00000078, 0x00000028, 0x00000048,
235 0x00000000, 0x00000050, 0x00000070, 0x00000038,
236 0x00000030, 0x00000040, 0x00000010, 0x00000060,
237 0x000002e8, 0x000002d8, 0x000002a0, 0x00000288,
238 0x00000298, 0x000002f8, 0x000002a8, 0x000002c8,
239 0x00000280, 0x000002d0, 0x000002f0, 0x000002b8,
240 0x000002b0, 0x000002c0, 0x00000290, 0x000002e0,
241 0x000003e8, 0x000003d8, 0x000003a0, 0x00000388,
242 0x00000398, 0x000003f8, 0x000003a8, 0x000003c8,
243 0x00000380, 0x000003d0, 0x000003f0, 0x000003b8,
244 0x000003b0, 0x000003c0, 0x00000390, 0x000003e0,
245 0x00000568, 0x00000558, 0x00000520, 0x00000508,
246 0x00000518, 0x00000578, 0x00000528, 0x00000548,
247 0x00000500, 0x00000550, 0x00000570, 0x00000538,
248 0x00000530, 0x00000540, 0x00000510, 0x00000560,
249 0x00000268, 0x00000258, 0x00000220, 0x00000208,
250 0x00000218, 0x00000278, 0x00000228, 0x00000248,
251 0x00000200, 0x00000250, 0x00000270, 0x00000238,
252 0x00000230, 0x00000240, 0x00000210, 0x00000260,
253 0x000004e8, 0x000004d8, 0x000004a0, 0x00000488,
254 0x00000498, 0x000004f8, 0x000004a8, 0x000004c8,
255 0x00000480, 0x000004d0, 0x000004f0, 0x000004b8,
256 0x000004b0, 0x000004c0, 0x00000490, 0x000004e0,
257 0x00000168, 0x00000158, 0x00000120, 0x00000108,
258 0x00000118, 0x00000178, 0x00000128, 0x00000148,
259 0x00000100, 0x00000150, 0x00000170, 0x00000138,
260 0x00000130, 0x00000140, 0x00000110, 0x00000160,
261 0x000001e8, 0x000001d8, 0x000001a0, 0x00000188,
262 0x00000198, 0x000001f8, 0x000001a8, 0x000001c8,
263 0x00000180, 0x000001d0, 0x000001f0, 0x000001b8,
264 0x000001b0, 0x000001c0, 0x00000190, 0x000001e0,
265 0x00000768, 0x00000758, 0x00000720, 0x00000708,
266 0x00000718, 0x00000778, 0x00000728, 0x00000748,
267 0x00000700, 0x00000750, 0x00000770, 0x00000738,
268 0x00000730, 0x00000740, 0x00000710, 0x00000760,
269 0x00000368, 0x00000358, 0x00000320, 0x00000308,
270 0x00000318, 0x00000378, 0x00000328, 0x00000348,
271 0x00000300, 0x00000350, 0x00000370, 0x00000338,
272 0x00000330, 0x00000340, 0x00000310, 0x00000360,
273 0x000005e8, 0x000005d8, 0x000005a0, 0x00000588,
274 0x00000598, 0x000005f8, 0x000005a8, 0x000005c8,
275 0x00000580, 0x000005d0, 0x000005f0, 0x000005b8,
276 0x000005b0, 0x000005c0, 0x00000590, 0x000005e0,
277 0x00000468, 0x00000458, 0x00000420, 0x00000408,
278 0x00000418, 0x00000478, 0x00000428, 0x00000448,
279 0x00000400, 0x00000450, 0x00000470, 0x00000438,
280 0x00000430, 0x00000440, 0x00000410, 0x00000460,
281 0x00000668, 0x00000658, 0x00000620, 0x00000608,
282 0x00000618, 0x00000678, 0x00000628, 0x00000648,
283 0x00000600, 0x00000650, 0x00000670, 0x00000638,
284 0x00000630, 0x00000640, 0x00000610, 0x00000660,
289 #define BOX(i,n,S) (S)[(n)][(i)]
291 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
293 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
295 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
296 #elif VECT_SIZE == 16
297 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
300 #define _round(k1,k2,tbl) \
304 l ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
305 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
306 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
307 BOX (((t >> 24) & 0xff), 3, tbl); \
309 r ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
310 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
311 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
312 BOX (((t >> 24) & 0xff), 3, tbl); \
315 #define R(k,h,s,i,t) \
321 _round (k[0], k[1], t); \
322 _round (k[2], k[3], t); \
323 _round (k[4], k[5], t); \
324 _round (k[6], k[7], t); \
325 _round (k[0], k[1], t); \
326 _round (k[2], k[3], t); \
327 _round (k[4], k[5], t); \
328 _round (k[6], k[7], t); \
329 _round (k[0], k[1], t); \
330 _round (k[2], k[3], t); \
331 _round (k[4], k[5], t); \
332 _round (k[6], k[7], t); \
333 _round (k[7], k[6], t); \
334 _round (k[5], k[4], t); \
335 _round (k[3], k[2], t); \
336 _round (k[1], k[0], t); \
342 w[0] = u[0] ^ v[0]; \
343 w[1] = u[1] ^ v[1]; \
344 w[2] = u[2] ^ v[2]; \
345 w[3] = u[3] ^ v[3]; \
346 w[4] = u[4] ^ v[4]; \
347 w[5] = u[5] ^ v[5]; \
348 w[6] = u[6] ^ v[6]; \
352 k[0] = ((w[0] & 0x000000ff) << 0) \
353 | ((w[2] & 0x000000ff) << 8) \
354 | ((w[4] & 0x000000ff) << 16) \
355 | ((w[6] & 0x000000ff) << 24); \
356 k[1] = ((w[0] & 0x0000ff00) >> 8) \
357 | ((w[2] & 0x0000ff00) >> 0) \
358 | ((w[4] & 0x0000ff00) << 8) \
359 | ((w[6] & 0x0000ff00) << 16); \
360 k[2] = ((w[0] & 0x00ff0000) >> 16) \
361 | ((w[2] & 0x00ff0000) >> 8) \
362 | ((w[4] & 0x00ff0000) << 0) \
363 | ((w[6] & 0x00ff0000) << 8); \
364 k[3] = ((w[0] & 0xff000000) >> 24) \
365 | ((w[2] & 0xff000000) >> 16) \
366 | ((w[4] & 0xff000000) >> 8) \
367 | ((w[6] & 0xff000000) >> 0); \
368 k[4] = ((w[1] & 0x000000ff) << 0) \
369 | ((w[3] & 0x000000ff) << 8) \
370 | ((w[5] & 0x000000ff) << 16) \
371 | ((w[7] & 0x000000ff) << 24); \
372 k[5] = ((w[1] & 0x0000ff00) >> 8) \
373 | ((w[3] & 0x0000ff00) >> 0) \
374 | ((w[5] & 0x0000ff00) << 8) \
375 | ((w[7] & 0x0000ff00) << 16); \
376 k[6] = ((w[1] & 0x00ff0000) >> 16) \
377 | ((w[3] & 0x00ff0000) >> 8) \
378 | ((w[5] & 0x00ff0000) << 0) \
379 | ((w[7] & 0x00ff0000) << 8); \
380 k[7] = ((w[1] & 0xff000000) >> 24) \
381 | ((w[3] & 0xff000000) >> 16) \
382 | ((w[5] & 0xff000000) >> 8) \
383 | ((w[7] & 0xff000000) >> 0);
420 x[0] ^= 0xff00ff00; \
421 x[1] ^= 0xff00ff00; \
422 x[2] ^= 0x00ff00ff; \
423 x[3] ^= 0x00ff00ff; \
424 x[4] ^= 0x00ffff00; \
425 x[5] ^= 0xff0000ff; \
426 x[6] ^= 0x000000ff; \
429 #define SHIFT12(u,m,s) \
430 u[0] = m[0] ^ s[6]; \
431 u[1] = m[1] ^ s[7]; \
432 u[2] = m[2] ^ (s[0] << 16) \
434 ^ (s[0] & 0x0000ffff) \
435 ^ (s[1] & 0x0000ffff) \
440 ^ (s[7] & 0xffff0000) \
442 u[3] = m[3] ^ (s[0] & 0x0000ffff) \
444 ^ (s[1] & 0x0000ffff) \
453 ^ (s[7] & 0x0000ffff) \
456 u[4] = m[4] ^ (s[0] & 0xffff0000) \
459 ^ (s[1] & 0xffff0000) \
468 ^ (s[7] & 0x0000ffff) \
471 u[5] = m[5] ^ (s[0] << 16) \
473 ^ (s[0] & 0xffff0000) \
474 ^ (s[1] & 0x0000ffff) \
484 ^ (s[7] & 0xffff0000) \
500 u[7] = m[7] ^ (s[0] & 0xffff0000) \
502 ^ (s[1] & 0x0000ffff) \
511 ^ (s[7] & 0x0000ffff) \
515 #define SHIFT16(h,v,u) \
516 v[0] = h[0] ^ (u[1] << 16) \
518 v[1] = h[1] ^ (u[2] << 16) \
520 v[2] = h[2] ^ (u[3] << 16) \
522 v[3] = h[3] ^ (u[4] << 16) \
524 v[4] = h[4] ^ (u[5] << 16) \
526 v[5] = h[5] ^ (u[6] << 16) \
528 v[6] = h[6] ^ (u[7] << 16) \
530 v[7] = h[7] ^ (u[0] & 0xffff0000) \
533 ^ (u[1] & 0xffff0000) \
536 ^ (u[7] & 0xffff0000);
538 #define SHIFT61(h,v) \
539 h[0] = (v[0] & 0xffff0000) \
543 ^ (v[1] & 0xffff0000) \
552 ^ (v[7] & 0x0000ffff); \
553 h[1] = (v[0] << 16) \
555 ^ (v[0] & 0xffff0000) \
556 ^ (v[1] & 0x0000ffff) \
564 ^ (v[7] & 0xffff0000) \
566 h[2] = (v[0] & 0x0000ffff) \
570 ^ (v[1] & 0xffff0000) \
578 ^ (v[7] & 0x0000ffff) \
581 h[3] = (v[0] << 16) \
583 ^ (v[0] & 0xffff0000) \
584 ^ (v[1] & 0xffff0000) \
594 ^ (v[7] & 0x0000ffff) \
596 h[4] = (v[0] >> 16) \
610 h[5] = (v[0] << 16) \
611 ^ (v[0] & 0xffff0000) \
614 ^ (v[1] & 0xffff0000) \
628 ^ (v[7] & 0xffff0000); \
660 #define PASS0(h,s,u,v,t) \
671 #define PASS2(h,s,u,v,t) \
683 #define PASS4(h,s,u,v,t) \
694 #define PASS6(h,s,u,v,t) \
703 void m06900m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 (*s_tables)[256])
709 const u32 gid = get_global_id (0);
710 const u32 lid = get_local_id (0);
718 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
720 const u32x w0r = ix_create_bft (bfs_buf, il_pos);
722 const u32x w0lr = w0l | w0r;
763 state_m[0] = state[0];
764 state_m[1] = state[1];
765 state_m[2] = state[2];
766 state_m[3] = state[3];
767 state_m[4] = state[4];
768 state_m[5] = state[5];
769 state_m[6] = state[6];
770 state_m[7] = state[7];
783 //if (pw_len > 0) // not really SIMD compatible
785 PASS0 (state, tmp, state_m, data_m, s_tables);
786 PASS2 (state, tmp, state_m, data_m, s_tables);
787 PASS4 (state, tmp, state_m, data_m, s_tables);
788 PASS6 (state, tmp, state_m, data_m, s_tables);
790 SHIFT12 (state_m, data, tmp);
791 SHIFT16 (state, data_m, state_m);
792 SHIFT61 (state, data_m);
795 data[0] = pw_len * 8;
806 state_m[0] = state[0];
807 state_m[1] = state[1];
808 state_m[2] = state[2];
809 state_m[3] = state[3];
810 state_m[4] = state[4];
811 state_m[5] = state[5];
812 state_m[6] = state[6];
813 state_m[7] = state[7];
824 PASS0 (state, tmp, state_m, data_m, s_tables);
825 PASS2 (state, tmp, state_m, data_m, s_tables);
826 PASS4 (state, tmp, state_m, data_m, s_tables);
827 PASS6 (state, tmp, state_m, data_m, s_tables);
829 SHIFT12 (state_m, data, tmp);
830 SHIFT16 (state, data_m, state_m);
831 SHIFT61 (state, data_m);
844 state_m[0] = state[0];
845 state_m[1] = state[1];
846 state_m[2] = state[2];
847 state_m[3] = state[3];
848 state_m[4] = state[4];
849 state_m[5] = state[5];
850 state_m[6] = state[6];
851 state_m[7] = state[7];
862 PASS0 (state, tmp, state_m, data_m, s_tables);
863 PASS2 (state, tmp, state_m, data_m, s_tables);
864 PASS4 (state, tmp, state_m, data_m, s_tables);
865 PASS6 (state, tmp, state_m, data_m, s_tables);
867 SHIFT12 (state_m, data, tmp);
868 SHIFT16 (state, data_m, state_m);
869 SHIFT61 (state, data_m);
873 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
877 void m06900s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 (*s_tables)[256])
883 const u32 gid = get_global_id (0);
884 const u32 lid = get_local_id (0);
890 const u32 search[4] =
892 digests_buf[digests_offset].digest_buf[DGST_R0],
893 digests_buf[digests_offset].digest_buf[DGST_R1],
894 digests_buf[digests_offset].digest_buf[DGST_R2],
895 digests_buf[digests_offset].digest_buf[DGST_R3]
904 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
906 const u32x w0r = ix_create_bft (bfs_buf, il_pos);
908 const u32x w0lr = w0l | w0r;
949 state_m[0] = state[0];
950 state_m[1] = state[1];
951 state_m[2] = state[2];
952 state_m[3] = state[3];
953 state_m[4] = state[4];
954 state_m[5] = state[5];
955 state_m[6] = state[6];
956 state_m[7] = state[7];
969 //if (pw_len > 0) // not really SIMD compatible
971 PASS0 (state, tmp, state_m, data_m, s_tables);
972 PASS2 (state, tmp, state_m, data_m, s_tables);
973 PASS4 (state, tmp, state_m, data_m, s_tables);
974 PASS6 (state, tmp, state_m, data_m, s_tables);
976 SHIFT12 (state_m, data, tmp);
977 SHIFT16 (state, data_m, state_m);
978 SHIFT61 (state, data_m);
981 data[0] = pw_len * 8;
992 state_m[0] = state[0];
993 state_m[1] = state[1];
994 state_m[2] = state[2];
995 state_m[3] = state[3];
996 state_m[4] = state[4];
997 state_m[5] = state[5];
998 state_m[6] = state[6];
999 state_m[7] = state[7];
1001 data_m[0] = data[0];
1002 data_m[1] = data[1];
1003 data_m[2] = data[2];
1004 data_m[3] = data[3];
1005 data_m[4] = data[4];
1006 data_m[5] = data[5];
1007 data_m[6] = data[6];
1008 data_m[7] = data[7];
1010 PASS0 (state, tmp, state_m, data_m, s_tables);
1011 PASS2 (state, tmp, state_m, data_m, s_tables);
1012 PASS4 (state, tmp, state_m, data_m, s_tables);
1013 PASS6 (state, tmp, state_m, data_m, s_tables);
1015 SHIFT12 (state_m, data, tmp);
1016 SHIFT16 (state, data_m, state_m);
1017 SHIFT61 (state, data_m);
1021 data[0] = state[ 8];
1022 data[1] = state[ 9];
1023 data[2] = state[10];
1024 data[3] = state[11];
1025 data[4] = state[12];
1026 data[5] = state[13];
1027 data[6] = state[14];
1028 data[7] = state[15];
1030 state_m[0] = state[0];
1031 state_m[1] = state[1];
1032 state_m[2] = state[2];
1033 state_m[3] = state[3];
1034 state_m[4] = state[4];
1035 state_m[5] = state[5];
1036 state_m[6] = state[6];
1037 state_m[7] = state[7];
1039 data_m[0] = data[0];
1040 data_m[1] = data[1];
1041 data_m[2] = data[2];
1042 data_m[3] = data[3];
1043 data_m[4] = data[4];
1044 data_m[5] = data[5];
1045 data_m[6] = data[6];
1046 data_m[7] = data[7];
1048 PASS0 (state, tmp, state_m, data_m, s_tables);
1049 PASS2 (state, tmp, state_m, data_m, s_tables);
1050 PASS4 (state, tmp, state_m, data_m, s_tables);
1051 PASS6 (state, tmp, state_m, data_m, s_tables);
1053 SHIFT12 (state_m, data, tmp);
1054 SHIFT16 (state, data_m, state_m);
1055 SHIFT61 (state, data_m);
1059 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
1063 __kernel void m06900_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1069 const u32 gid = get_global_id (0);
1070 const u32 lid = get_local_id (0);
1071 const u32 lsz = get_local_size (0);
1077 __local u32 s_tables[4][256];
1079 for (u32 i = lid; i < 256; i += lsz)
1081 s_tables[0][i] = c_tables[0][i];
1082 s_tables[1][i] = c_tables[1][i];
1083 s_tables[2][i] = c_tables[2][i];
1084 s_tables[3][i] = c_tables[3][i];
1087 barrier (CLK_LOCAL_MEM_FENCE);
1089 if (gid >= gid_max) return;
1097 w0[0] = pws[gid].i[ 0];
1098 w0[1] = pws[gid].i[ 1];
1099 w0[2] = pws[gid].i[ 2];
1100 w0[3] = pws[gid].i[ 3];
1123 const u32 pw_len = pws[gid].pw_len;
1129 m06900m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, s_tables);
1132 __kernel void m06900_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1138 const u32 gid = get_global_id (0);
1139 const u32 lid = get_local_id (0);
1140 const u32 lsz = get_local_size (0);
1146 __local u32 s_tables[4][256];
1148 for (u32 i = lid; i < 256; i += lsz)
1150 s_tables[0][i] = c_tables[0][i];
1151 s_tables[1][i] = c_tables[1][i];
1152 s_tables[2][i] = c_tables[2][i];
1153 s_tables[3][i] = c_tables[3][i];
1156 barrier (CLK_LOCAL_MEM_FENCE);
1158 if (gid >= gid_max) return;
1166 w0[0] = pws[gid].i[ 0];
1167 w0[1] = pws[gid].i[ 1];
1168 w0[2] = pws[gid].i[ 2];
1169 w0[3] = pws[gid].i[ 3];
1173 w1[0] = pws[gid].i[ 4];
1174 w1[1] = pws[gid].i[ 5];
1175 w1[2] = pws[gid].i[ 6];
1176 w1[3] = pws[gid].i[ 7];
1192 const u32 pw_len = pws[gid].pw_len;
1198 m06900m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, s_tables);
1201 __kernel void m06900_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1205 __kernel void m06900_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1211 const u32 gid = get_global_id (0);
1212 const u32 lid = get_local_id (0);
1213 const u32 lsz = get_local_size (0);
1219 __local u32 s_tables[4][256];
1221 for (u32 i = lid; i < 256; i += lsz)
1223 s_tables[0][i] = c_tables[0][i];
1224 s_tables[1][i] = c_tables[1][i];
1225 s_tables[2][i] = c_tables[2][i];
1226 s_tables[3][i] = c_tables[3][i];
1229 barrier (CLK_LOCAL_MEM_FENCE);
1231 if (gid >= gid_max) return;
1239 w0[0] = pws[gid].i[ 0];
1240 w0[1] = pws[gid].i[ 1];
1241 w0[2] = pws[gid].i[ 2];
1242 w0[3] = pws[gid].i[ 3];
1265 const u32 pw_len = pws[gid].pw_len;
1271 m06900s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, s_tables);
1274 __kernel void m06900_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1280 const u32 gid = get_global_id (0);
1281 const u32 lid = get_local_id (0);
1282 const u32 lsz = get_local_size (0);
1288 __local u32 s_tables[4][256];
1290 for (u32 i = lid; i < 256; i += lsz)
1292 s_tables[0][i] = c_tables[0][i];
1293 s_tables[1][i] = c_tables[1][i];
1294 s_tables[2][i] = c_tables[2][i];
1295 s_tables[3][i] = c_tables[3][i];
1298 barrier (CLK_LOCAL_MEM_FENCE);
1300 if (gid >= gid_max) return;
1308 w0[0] = pws[gid].i[ 0];
1309 w0[1] = pws[gid].i[ 1];
1310 w0[2] = pws[gid].i[ 2];
1311 w0[3] = pws[gid].i[ 3];
1315 w1[0] = pws[gid].i[ 4];
1316 w1[1] = pws[gid].i[ 5];
1317 w1[2] = pws[gid].i[ 6];
1318 w1[3] = pws[gid].i[ 7];
1334 const u32 pw_len = pws[gid].pw_len;
1340 m06900s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, s_tables);
1343 __kernel void m06900_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)