2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
4 * magnum <john.magnum@hushmail.com>
13 #include "include/constants.h"
14 #include "include/kernel_vendor.h"
21 #include "include/kernel_functions.c"
22 #include "OpenCL/types_ocl.c"
23 #include "OpenCL/common.c"
24 #include "OpenCL/simd.c"
26 __constant u32 c_tables[4][256] =
29 0x00072000, 0x00075000, 0x00074800, 0x00071000,
30 0x00076800, 0x00074000, 0x00070000, 0x00077000,
31 0x00073000, 0x00075800, 0x00070800, 0x00076000,
32 0x00073800, 0x00077800, 0x00072800, 0x00071800,
33 0x0005a000, 0x0005d000, 0x0005c800, 0x00059000,
34 0x0005e800, 0x0005c000, 0x00058000, 0x0005f000,
35 0x0005b000, 0x0005d800, 0x00058800, 0x0005e000,
36 0x0005b800, 0x0005f800, 0x0005a800, 0x00059800,
37 0x00022000, 0x00025000, 0x00024800, 0x00021000,
38 0x00026800, 0x00024000, 0x00020000, 0x00027000,
39 0x00023000, 0x00025800, 0x00020800, 0x00026000,
40 0x00023800, 0x00027800, 0x00022800, 0x00021800,
41 0x00062000, 0x00065000, 0x00064800, 0x00061000,
42 0x00066800, 0x00064000, 0x00060000, 0x00067000,
43 0x00063000, 0x00065800, 0x00060800, 0x00066000,
44 0x00063800, 0x00067800, 0x00062800, 0x00061800,
45 0x00032000, 0x00035000, 0x00034800, 0x00031000,
46 0x00036800, 0x00034000, 0x00030000, 0x00037000,
47 0x00033000, 0x00035800, 0x00030800, 0x00036000,
48 0x00033800, 0x00037800, 0x00032800, 0x00031800,
49 0x0006a000, 0x0006d000, 0x0006c800, 0x00069000,
50 0x0006e800, 0x0006c000, 0x00068000, 0x0006f000,
51 0x0006b000, 0x0006d800, 0x00068800, 0x0006e000,
52 0x0006b800, 0x0006f800, 0x0006a800, 0x00069800,
53 0x0007a000, 0x0007d000, 0x0007c800, 0x00079000,
54 0x0007e800, 0x0007c000, 0x00078000, 0x0007f000,
55 0x0007b000, 0x0007d800, 0x00078800, 0x0007e000,
56 0x0007b800, 0x0007f800, 0x0007a800, 0x00079800,
57 0x00052000, 0x00055000, 0x00054800, 0x00051000,
58 0x00056800, 0x00054000, 0x00050000, 0x00057000,
59 0x00053000, 0x00055800, 0x00050800, 0x00056000,
60 0x00053800, 0x00057800, 0x00052800, 0x00051800,
61 0x00012000, 0x00015000, 0x00014800, 0x00011000,
62 0x00016800, 0x00014000, 0x00010000, 0x00017000,
63 0x00013000, 0x00015800, 0x00010800, 0x00016000,
64 0x00013800, 0x00017800, 0x00012800, 0x00011800,
65 0x0001a000, 0x0001d000, 0x0001c800, 0x00019000,
66 0x0001e800, 0x0001c000, 0x00018000, 0x0001f000,
67 0x0001b000, 0x0001d800, 0x00018800, 0x0001e000,
68 0x0001b800, 0x0001f800, 0x0001a800, 0x00019800,
69 0x00042000, 0x00045000, 0x00044800, 0x00041000,
70 0x00046800, 0x00044000, 0x00040000, 0x00047000,
71 0x00043000, 0x00045800, 0x00040800, 0x00046000,
72 0x00043800, 0x00047800, 0x00042800, 0x00041800,
73 0x0000a000, 0x0000d000, 0x0000c800, 0x00009000,
74 0x0000e800, 0x0000c000, 0x00008000, 0x0000f000,
75 0x0000b000, 0x0000d800, 0x00008800, 0x0000e000,
76 0x0000b800, 0x0000f800, 0x0000a800, 0x00009800,
77 0x00002000, 0x00005000, 0x00004800, 0x00001000,
78 0x00006800, 0x00004000, 0x00000000, 0x00007000,
79 0x00003000, 0x00005800, 0x00000800, 0x00006000,
80 0x00003800, 0x00007800, 0x00002800, 0x00001800,
81 0x0003a000, 0x0003d000, 0x0003c800, 0x00039000,
82 0x0003e800, 0x0003c000, 0x00038000, 0x0003f000,
83 0x0003b000, 0x0003d800, 0x00038800, 0x0003e000,
84 0x0003b800, 0x0003f800, 0x0003a800, 0x00039800,
85 0x0002a000, 0x0002d000, 0x0002c800, 0x00029000,
86 0x0002e800, 0x0002c000, 0x00028000, 0x0002f000,
87 0x0002b000, 0x0002d800, 0x00028800, 0x0002e000,
88 0x0002b800, 0x0002f800, 0x0002a800, 0x00029800,
89 0x0004a000, 0x0004d000, 0x0004c800, 0x00049000,
90 0x0004e800, 0x0004c000, 0x00048000, 0x0004f000,
91 0x0004b000, 0x0004d800, 0x00048800, 0x0004e000,
92 0x0004b800, 0x0004f800, 0x0004a800, 0x00049800,
95 0x03a80000, 0x03c00000, 0x03880000, 0x03e80000,
96 0x03d00000, 0x03980000, 0x03a00000, 0x03900000,
97 0x03f00000, 0x03f80000, 0x03e00000, 0x03b80000,
98 0x03b00000, 0x03800000, 0x03c80000, 0x03d80000,
99 0x06a80000, 0x06c00000, 0x06880000, 0x06e80000,
100 0x06d00000, 0x06980000, 0x06a00000, 0x06900000,
101 0x06f00000, 0x06f80000, 0x06e00000, 0x06b80000,
102 0x06b00000, 0x06800000, 0x06c80000, 0x06d80000,
103 0x05280000, 0x05400000, 0x05080000, 0x05680000,
104 0x05500000, 0x05180000, 0x05200000, 0x05100000,
105 0x05700000, 0x05780000, 0x05600000, 0x05380000,
106 0x05300000, 0x05000000, 0x05480000, 0x05580000,
107 0x00a80000, 0x00c00000, 0x00880000, 0x00e80000,
108 0x00d00000, 0x00980000, 0x00a00000, 0x00900000,
109 0x00f00000, 0x00f80000, 0x00e00000, 0x00b80000,
110 0x00b00000, 0x00800000, 0x00c80000, 0x00d80000,
111 0x00280000, 0x00400000, 0x00080000, 0x00680000,
112 0x00500000, 0x00180000, 0x00200000, 0x00100000,
113 0x00700000, 0x00780000, 0x00600000, 0x00380000,
114 0x00300000, 0x00000000, 0x00480000, 0x00580000,
115 0x04280000, 0x04400000, 0x04080000, 0x04680000,
116 0x04500000, 0x04180000, 0x04200000, 0x04100000,
117 0x04700000, 0x04780000, 0x04600000, 0x04380000,
118 0x04300000, 0x04000000, 0x04480000, 0x04580000,
119 0x04a80000, 0x04c00000, 0x04880000, 0x04e80000,
120 0x04d00000, 0x04980000, 0x04a00000, 0x04900000,
121 0x04f00000, 0x04f80000, 0x04e00000, 0x04b80000,
122 0x04b00000, 0x04800000, 0x04c80000, 0x04d80000,
123 0x07a80000, 0x07c00000, 0x07880000, 0x07e80000,
124 0x07d00000, 0x07980000, 0x07a00000, 0x07900000,
125 0x07f00000, 0x07f80000, 0x07e00000, 0x07b80000,
126 0x07b00000, 0x07800000, 0x07c80000, 0x07d80000,
127 0x07280000, 0x07400000, 0x07080000, 0x07680000,
128 0x07500000, 0x07180000, 0x07200000, 0x07100000,
129 0x07700000, 0x07780000, 0x07600000, 0x07380000,
130 0x07300000, 0x07000000, 0x07480000, 0x07580000,
131 0x02280000, 0x02400000, 0x02080000, 0x02680000,
132 0x02500000, 0x02180000, 0x02200000, 0x02100000,
133 0x02700000, 0x02780000, 0x02600000, 0x02380000,
134 0x02300000, 0x02000000, 0x02480000, 0x02580000,
135 0x03280000, 0x03400000, 0x03080000, 0x03680000,
136 0x03500000, 0x03180000, 0x03200000, 0x03100000,
137 0x03700000, 0x03780000, 0x03600000, 0x03380000,
138 0x03300000, 0x03000000, 0x03480000, 0x03580000,
139 0x06280000, 0x06400000, 0x06080000, 0x06680000,
140 0x06500000, 0x06180000, 0x06200000, 0x06100000,
141 0x06700000, 0x06780000, 0x06600000, 0x06380000,
142 0x06300000, 0x06000000, 0x06480000, 0x06580000,
143 0x05a80000, 0x05c00000, 0x05880000, 0x05e80000,
144 0x05d00000, 0x05980000, 0x05a00000, 0x05900000,
145 0x05f00000, 0x05f80000, 0x05e00000, 0x05b80000,
146 0x05b00000, 0x05800000, 0x05c80000, 0x05d80000,
147 0x01280000, 0x01400000, 0x01080000, 0x01680000,
148 0x01500000, 0x01180000, 0x01200000, 0x01100000,
149 0x01700000, 0x01780000, 0x01600000, 0x01380000,
150 0x01300000, 0x01000000, 0x01480000, 0x01580000,
151 0x02a80000, 0x02c00000, 0x02880000, 0x02e80000,
152 0x02d00000, 0x02980000, 0x02a00000, 0x02900000,
153 0x02f00000, 0x02f80000, 0x02e00000, 0x02b80000,
154 0x02b00000, 0x02800000, 0x02c80000, 0x02d80000,
155 0x01a80000, 0x01c00000, 0x01880000, 0x01e80000,
156 0x01d00000, 0x01980000, 0x01a00000, 0x01900000,
157 0x01f00000, 0x01f80000, 0x01e00000, 0x01b80000,
158 0x01b00000, 0x01800000, 0x01c80000, 0x01d80000,
161 0x30000002, 0x60000002, 0x38000002, 0x08000002,
162 0x28000002, 0x78000002, 0x68000002, 0x40000002,
163 0x20000002, 0x50000002, 0x48000002, 0x70000002,
164 0x00000002, 0x18000002, 0x58000002, 0x10000002,
165 0xb0000005, 0xe0000005, 0xb8000005, 0x88000005,
166 0xa8000005, 0xf8000005, 0xe8000005, 0xc0000005,
167 0xa0000005, 0xd0000005, 0xc8000005, 0xf0000005,
168 0x80000005, 0x98000005, 0xd8000005, 0x90000005,
169 0x30000005, 0x60000005, 0x38000005, 0x08000005,
170 0x28000005, 0x78000005, 0x68000005, 0x40000005,
171 0x20000005, 0x50000005, 0x48000005, 0x70000005,
172 0x00000005, 0x18000005, 0x58000005, 0x10000005,
173 0x30000000, 0x60000000, 0x38000000, 0x08000000,
174 0x28000000, 0x78000000, 0x68000000, 0x40000000,
175 0x20000000, 0x50000000, 0x48000000, 0x70000000,
176 0x00000000, 0x18000000, 0x58000000, 0x10000000,
177 0xb0000003, 0xe0000003, 0xb8000003, 0x88000003,
178 0xa8000003, 0xf8000003, 0xe8000003, 0xc0000003,
179 0xa0000003, 0xd0000003, 0xc8000003, 0xf0000003,
180 0x80000003, 0x98000003, 0xd8000003, 0x90000003,
181 0x30000001, 0x60000001, 0x38000001, 0x08000001,
182 0x28000001, 0x78000001, 0x68000001, 0x40000001,
183 0x20000001, 0x50000001, 0x48000001, 0x70000001,
184 0x00000001, 0x18000001, 0x58000001, 0x10000001,
185 0xb0000000, 0xe0000000, 0xb8000000, 0x88000000,
186 0xa8000000, 0xf8000000, 0xe8000000, 0xc0000000,
187 0xa0000000, 0xd0000000, 0xc8000000, 0xf0000000,
188 0x80000000, 0x98000000, 0xd8000000, 0x90000000,
189 0xb0000006, 0xe0000006, 0xb8000006, 0x88000006,
190 0xa8000006, 0xf8000006, 0xe8000006, 0xc0000006,
191 0xa0000006, 0xd0000006, 0xc8000006, 0xf0000006,
192 0x80000006, 0x98000006, 0xd8000006, 0x90000006,
193 0xb0000001, 0xe0000001, 0xb8000001, 0x88000001,
194 0xa8000001, 0xf8000001, 0xe8000001, 0xc0000001,
195 0xa0000001, 0xd0000001, 0xc8000001, 0xf0000001,
196 0x80000001, 0x98000001, 0xd8000001, 0x90000001,
197 0x30000003, 0x60000003, 0x38000003, 0x08000003,
198 0x28000003, 0x78000003, 0x68000003, 0x40000003,
199 0x20000003, 0x50000003, 0x48000003, 0x70000003,
200 0x00000003, 0x18000003, 0x58000003, 0x10000003,
201 0x30000004, 0x60000004, 0x38000004, 0x08000004,
202 0x28000004, 0x78000004, 0x68000004, 0x40000004,
203 0x20000004, 0x50000004, 0x48000004, 0x70000004,
204 0x00000004, 0x18000004, 0x58000004, 0x10000004,
205 0xb0000002, 0xe0000002, 0xb8000002, 0x88000002,
206 0xa8000002, 0xf8000002, 0xe8000002, 0xc0000002,
207 0xa0000002, 0xd0000002, 0xc8000002, 0xf0000002,
208 0x80000002, 0x98000002, 0xd8000002, 0x90000002,
209 0xb0000004, 0xe0000004, 0xb8000004, 0x88000004,
210 0xa8000004, 0xf8000004, 0xe8000004, 0xc0000004,
211 0xa0000004, 0xd0000004, 0xc8000004, 0xf0000004,
212 0x80000004, 0x98000004, 0xd8000004, 0x90000004,
213 0x30000006, 0x60000006, 0x38000006, 0x08000006,
214 0x28000006, 0x78000006, 0x68000006, 0x40000006,
215 0x20000006, 0x50000006, 0x48000006, 0x70000006,
216 0x00000006, 0x18000006, 0x58000006, 0x10000006,
217 0xb0000007, 0xe0000007, 0xb8000007, 0x88000007,
218 0xa8000007, 0xf8000007, 0xe8000007, 0xc0000007,
219 0xa0000007, 0xd0000007, 0xc8000007, 0xf0000007,
220 0x80000007, 0x98000007, 0xd8000007, 0x90000007,
221 0x30000007, 0x60000007, 0x38000007, 0x08000007,
222 0x28000007, 0x78000007, 0x68000007, 0x40000007,
223 0x20000007, 0x50000007, 0x48000007, 0x70000007,
224 0x00000007, 0x18000007, 0x58000007, 0x10000007,
227 0x000000e8, 0x000000d8, 0x000000a0, 0x00000088,
228 0x00000098, 0x000000f8, 0x000000a8, 0x000000c8,
229 0x00000080, 0x000000d0, 0x000000f0, 0x000000b8,
230 0x000000b0, 0x000000c0, 0x00000090, 0x000000e0,
231 0x000007e8, 0x000007d8, 0x000007a0, 0x00000788,
232 0x00000798, 0x000007f8, 0x000007a8, 0x000007c8,
233 0x00000780, 0x000007d0, 0x000007f0, 0x000007b8,
234 0x000007b0, 0x000007c0, 0x00000790, 0x000007e0,
235 0x000006e8, 0x000006d8, 0x000006a0, 0x00000688,
236 0x00000698, 0x000006f8, 0x000006a8, 0x000006c8,
237 0x00000680, 0x000006d0, 0x000006f0, 0x000006b8,
238 0x000006b0, 0x000006c0, 0x00000690, 0x000006e0,
239 0x00000068, 0x00000058, 0x00000020, 0x00000008,
240 0x00000018, 0x00000078, 0x00000028, 0x00000048,
241 0x00000000, 0x00000050, 0x00000070, 0x00000038,
242 0x00000030, 0x00000040, 0x00000010, 0x00000060,
243 0x000002e8, 0x000002d8, 0x000002a0, 0x00000288,
244 0x00000298, 0x000002f8, 0x000002a8, 0x000002c8,
245 0x00000280, 0x000002d0, 0x000002f0, 0x000002b8,
246 0x000002b0, 0x000002c0, 0x00000290, 0x000002e0,
247 0x000003e8, 0x000003d8, 0x000003a0, 0x00000388,
248 0x00000398, 0x000003f8, 0x000003a8, 0x000003c8,
249 0x00000380, 0x000003d0, 0x000003f0, 0x000003b8,
250 0x000003b0, 0x000003c0, 0x00000390, 0x000003e0,
251 0x00000568, 0x00000558, 0x00000520, 0x00000508,
252 0x00000518, 0x00000578, 0x00000528, 0x00000548,
253 0x00000500, 0x00000550, 0x00000570, 0x00000538,
254 0x00000530, 0x00000540, 0x00000510, 0x00000560,
255 0x00000268, 0x00000258, 0x00000220, 0x00000208,
256 0x00000218, 0x00000278, 0x00000228, 0x00000248,
257 0x00000200, 0x00000250, 0x00000270, 0x00000238,
258 0x00000230, 0x00000240, 0x00000210, 0x00000260,
259 0x000004e8, 0x000004d8, 0x000004a0, 0x00000488,
260 0x00000498, 0x000004f8, 0x000004a8, 0x000004c8,
261 0x00000480, 0x000004d0, 0x000004f0, 0x000004b8,
262 0x000004b0, 0x000004c0, 0x00000490, 0x000004e0,
263 0x00000168, 0x00000158, 0x00000120, 0x00000108,
264 0x00000118, 0x00000178, 0x00000128, 0x00000148,
265 0x00000100, 0x00000150, 0x00000170, 0x00000138,
266 0x00000130, 0x00000140, 0x00000110, 0x00000160,
267 0x000001e8, 0x000001d8, 0x000001a0, 0x00000188,
268 0x00000198, 0x000001f8, 0x000001a8, 0x000001c8,
269 0x00000180, 0x000001d0, 0x000001f0, 0x000001b8,
270 0x000001b0, 0x000001c0, 0x00000190, 0x000001e0,
271 0x00000768, 0x00000758, 0x00000720, 0x00000708,
272 0x00000718, 0x00000778, 0x00000728, 0x00000748,
273 0x00000700, 0x00000750, 0x00000770, 0x00000738,
274 0x00000730, 0x00000740, 0x00000710, 0x00000760,
275 0x00000368, 0x00000358, 0x00000320, 0x00000308,
276 0x00000318, 0x00000378, 0x00000328, 0x00000348,
277 0x00000300, 0x00000350, 0x00000370, 0x00000338,
278 0x00000330, 0x00000340, 0x00000310, 0x00000360,
279 0x000005e8, 0x000005d8, 0x000005a0, 0x00000588,
280 0x00000598, 0x000005f8, 0x000005a8, 0x000005c8,
281 0x00000580, 0x000005d0, 0x000005f0, 0x000005b8,
282 0x000005b0, 0x000005c0, 0x00000590, 0x000005e0,
283 0x00000468, 0x00000458, 0x00000420, 0x00000408,
284 0x00000418, 0x00000478, 0x00000428, 0x00000448,
285 0x00000400, 0x00000450, 0x00000470, 0x00000438,
286 0x00000430, 0x00000440, 0x00000410, 0x00000460,
287 0x00000668, 0x00000658, 0x00000620, 0x00000608,
288 0x00000618, 0x00000678, 0x00000628, 0x00000648,
289 0x00000600, 0x00000650, 0x00000670, 0x00000638,
290 0x00000630, 0x00000640, 0x00000610, 0x00000660,
295 #define BOX(i,n,S) (S)[(n)][(i)]
297 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
299 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
301 #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])
302 #elif VECT_SIZE == 16
303 #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])
306 #define _round(k1,k2,tbl) \
310 l ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
311 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
312 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
313 BOX (((t >> 24) & 0xff), 3, tbl); \
315 r ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
316 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
317 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
318 BOX (((t >> 24) & 0xff), 3, tbl); \
321 #define R(k,h,s,i,t) \
327 _round (k[0], k[1], t); \
328 _round (k[2], k[3], t); \
329 _round (k[4], k[5], t); \
330 _round (k[6], k[7], t); \
331 _round (k[0], k[1], t); \
332 _round (k[2], k[3], t); \
333 _round (k[4], k[5], t); \
334 _round (k[6], k[7], t); \
335 _round (k[0], k[1], t); \
336 _round (k[2], k[3], t); \
337 _round (k[4], k[5], t); \
338 _round (k[6], k[7], t); \
339 _round (k[7], k[6], t); \
340 _round (k[5], k[4], t); \
341 _round (k[3], k[2], t); \
342 _round (k[1], k[0], t); \
348 w[0] = u[0] ^ v[0]; \
349 w[1] = u[1] ^ v[1]; \
350 w[2] = u[2] ^ v[2]; \
351 w[3] = u[3] ^ v[3]; \
352 w[4] = u[4] ^ v[4]; \
353 w[5] = u[5] ^ v[5]; \
354 w[6] = u[6] ^ v[6]; \
358 k[0] = ((w[0] & 0x000000ff) << 0) \
359 | ((w[2] & 0x000000ff) << 8) \
360 | ((w[4] & 0x000000ff) << 16) \
361 | ((w[6] & 0x000000ff) << 24); \
362 k[1] = ((w[0] & 0x0000ff00) >> 8) \
363 | ((w[2] & 0x0000ff00) >> 0) \
364 | ((w[4] & 0x0000ff00) << 8) \
365 | ((w[6] & 0x0000ff00) << 16); \
366 k[2] = ((w[0] & 0x00ff0000) >> 16) \
367 | ((w[2] & 0x00ff0000) >> 8) \
368 | ((w[4] & 0x00ff0000) << 0) \
369 | ((w[6] & 0x00ff0000) << 8); \
370 k[3] = ((w[0] & 0xff000000) >> 24) \
371 | ((w[2] & 0xff000000) >> 16) \
372 | ((w[4] & 0xff000000) >> 8) \
373 | ((w[6] & 0xff000000) >> 0); \
374 k[4] = ((w[1] & 0x000000ff) << 0) \
375 | ((w[3] & 0x000000ff) << 8) \
376 | ((w[5] & 0x000000ff) << 16) \
377 | ((w[7] & 0x000000ff) << 24); \
378 k[5] = ((w[1] & 0x0000ff00) >> 8) \
379 | ((w[3] & 0x0000ff00) >> 0) \
380 | ((w[5] & 0x0000ff00) << 8) \
381 | ((w[7] & 0x0000ff00) << 16); \
382 k[6] = ((w[1] & 0x00ff0000) >> 16) \
383 | ((w[3] & 0x00ff0000) >> 8) \
384 | ((w[5] & 0x00ff0000) << 0) \
385 | ((w[7] & 0x00ff0000) << 8); \
386 k[7] = ((w[1] & 0xff000000) >> 24) \
387 | ((w[3] & 0xff000000) >> 16) \
388 | ((w[5] & 0xff000000) >> 8) \
389 | ((w[7] & 0xff000000) >> 0);
426 x[0] ^= 0xff00ff00; \
427 x[1] ^= 0xff00ff00; \
428 x[2] ^= 0x00ff00ff; \
429 x[3] ^= 0x00ff00ff; \
430 x[4] ^= 0x00ffff00; \
431 x[5] ^= 0xff0000ff; \
432 x[6] ^= 0x000000ff; \
435 #define SHIFT12(u,m,s) \
436 u[0] = m[0] ^ s[6]; \
437 u[1] = m[1] ^ s[7]; \
438 u[2] = m[2] ^ (s[0] << 16) \
440 ^ (s[0] & 0x0000ffff) \
441 ^ (s[1] & 0x0000ffff) \
446 ^ (s[7] & 0xffff0000) \
448 u[3] = m[3] ^ (s[0] & 0x0000ffff) \
450 ^ (s[1] & 0x0000ffff) \
459 ^ (s[7] & 0x0000ffff) \
462 u[4] = m[4] ^ (s[0] & 0xffff0000) \
465 ^ (s[1] & 0xffff0000) \
474 ^ (s[7] & 0x0000ffff) \
477 u[5] = m[5] ^ (s[0] << 16) \
479 ^ (s[0] & 0xffff0000) \
480 ^ (s[1] & 0x0000ffff) \
490 ^ (s[7] & 0xffff0000) \
506 u[7] = m[7] ^ (s[0] & 0xffff0000) \
508 ^ (s[1] & 0x0000ffff) \
517 ^ (s[7] & 0x0000ffff) \
521 #define SHIFT16(h,v,u) \
522 v[0] = h[0] ^ (u[1] << 16) \
524 v[1] = h[1] ^ (u[2] << 16) \
526 v[2] = h[2] ^ (u[3] << 16) \
528 v[3] = h[3] ^ (u[4] << 16) \
530 v[4] = h[4] ^ (u[5] << 16) \
532 v[5] = h[5] ^ (u[6] << 16) \
534 v[6] = h[6] ^ (u[7] << 16) \
536 v[7] = h[7] ^ (u[0] & 0xffff0000) \
539 ^ (u[1] & 0xffff0000) \
542 ^ (u[7] & 0xffff0000);
544 #define SHIFT61(h,v) \
545 h[0] = (v[0] & 0xffff0000) \
549 ^ (v[1] & 0xffff0000) \
558 ^ (v[7] & 0x0000ffff); \
559 h[1] = (v[0] << 16) \
561 ^ (v[0] & 0xffff0000) \
562 ^ (v[1] & 0x0000ffff) \
570 ^ (v[7] & 0xffff0000) \
572 h[2] = (v[0] & 0x0000ffff) \
576 ^ (v[1] & 0xffff0000) \
584 ^ (v[7] & 0x0000ffff) \
587 h[3] = (v[0] << 16) \
589 ^ (v[0] & 0xffff0000) \
590 ^ (v[1] & 0xffff0000) \
600 ^ (v[7] & 0x0000ffff) \
602 h[4] = (v[0] >> 16) \
616 h[5] = (v[0] << 16) \
617 ^ (v[0] & 0xffff0000) \
620 ^ (v[1] & 0xffff0000) \
634 ^ (v[7] & 0xffff0000); \
666 #define PASS0(h,s,u,v,t) \
677 #define PASS2(h,s,u,v,t) \
689 #define PASS4(h,s,u,v,t) \
700 #define PASS6(h,s,u,v,t) \
709 static 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])
715 const u32 gid = get_global_id (0);
716 const u32 lid = get_local_id (0);
722 const u32 w14 = pw_len * 8;
730 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
732 const u32x w0r = ix_create_bft (bfs_buf, il_pos);
734 const u32x w0lr = w0l | w0r;
771 state_m[0] = state[0];
772 state_m[1] = state[1];
773 state_m[2] = state[2];
774 state_m[3] = state[3];
775 state_m[4] = state[4];
776 state_m[5] = state[5];
777 state_m[6] = state[6];
778 state_m[7] = state[7];
793 PASS0 (state, tmp, state_m, data_m, s_tables);
794 PASS2 (state, tmp, state_m, data_m, s_tables);
795 PASS4 (state, tmp, state_m, data_m, s_tables);
796 PASS6 (state, tmp, state_m, data_m, s_tables);
798 SHIFT12 (state_m, data, tmp);
799 SHIFT16 (state, data_m, state_m);
800 SHIFT61 (state, data_m);
814 state_m[0] = state[0];
815 state_m[1] = state[1];
816 state_m[2] = state[2];
817 state_m[3] = state[3];
818 state_m[4] = state[4];
819 state_m[5] = state[5];
820 state_m[6] = state[6];
821 state_m[7] = state[7];
832 PASS0 (state, tmp, state_m, data_m, s_tables);
833 PASS2 (state, tmp, state_m, data_m, s_tables);
834 PASS4 (state, tmp, state_m, data_m, s_tables);
835 PASS6 (state, tmp, state_m, data_m, s_tables);
837 SHIFT12 (state_m, data, tmp);
838 SHIFT16 (state, data_m, state_m);
839 SHIFT61 (state, data_m);
852 state_m[0] = state[0];
853 state_m[1] = state[1];
854 state_m[2] = state[2];
855 state_m[3] = state[3];
856 state_m[4] = state[4];
857 state_m[5] = state[5];
858 state_m[6] = state[6];
859 state_m[7] = state[7];
870 PASS0 (state, tmp, state_m, data_m, s_tables);
871 PASS2 (state, tmp, state_m, data_m, s_tables);
872 PASS4 (state, tmp, state_m, data_m, s_tables);
873 PASS6 (state, tmp, state_m, data_m, s_tables);
875 SHIFT12 (state_m, data, tmp);
876 SHIFT16 (state, data_m, state_m);
877 SHIFT61 (state, data_m);
881 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
885 static 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])
891 const u32 gid = get_global_id (0);
892 const u32 lid = get_local_id (0);
898 const u32 w14 = pw_len * 8;
904 const u32 search[4] =
906 digests_buf[digests_offset].digest_buf[DGST_R0],
907 digests_buf[digests_offset].digest_buf[DGST_R1],
908 digests_buf[digests_offset].digest_buf[DGST_R2],
909 digests_buf[digests_offset].digest_buf[DGST_R3]
918 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
920 const u32x w0r = ix_create_bft (bfs_buf, il_pos);
922 const u32x w0lr = w0l | w0r;
959 state_m[0] = state[0];
960 state_m[1] = state[1];
961 state_m[2] = state[2];
962 state_m[3] = state[3];
963 state_m[4] = state[4];
964 state_m[5] = state[5];
965 state_m[6] = state[6];
966 state_m[7] = state[7];
981 PASS0 (state, tmp, state_m, data_m, s_tables);
982 PASS2 (state, tmp, state_m, data_m, s_tables);
983 PASS4 (state, tmp, state_m, data_m, s_tables);
984 PASS6 (state, tmp, state_m, data_m, s_tables);
986 SHIFT12 (state_m, data, tmp);
987 SHIFT16 (state, data_m, state_m);
988 SHIFT61 (state, data_m);
1002 state_m[0] = state[0];
1003 state_m[1] = state[1];
1004 state_m[2] = state[2];
1005 state_m[3] = state[3];
1006 state_m[4] = state[4];
1007 state_m[5] = state[5];
1008 state_m[6] = state[6];
1009 state_m[7] = state[7];
1011 data_m[0] = data[0];
1012 data_m[1] = data[1];
1013 data_m[2] = data[2];
1014 data_m[3] = data[3];
1015 data_m[4] = data[4];
1016 data_m[5] = data[5];
1017 data_m[6] = data[6];
1018 data_m[7] = data[7];
1020 PASS0 (state, tmp, state_m, data_m, s_tables);
1021 PASS2 (state, tmp, state_m, data_m, s_tables);
1022 PASS4 (state, tmp, state_m, data_m, s_tables);
1023 PASS6 (state, tmp, state_m, data_m, s_tables);
1025 SHIFT12 (state_m, data, tmp);
1026 SHIFT16 (state, data_m, state_m);
1027 SHIFT61 (state, data_m);
1031 data[0] = state[ 8];
1032 data[1] = state[ 9];
1033 data[2] = state[10];
1034 data[3] = state[11];
1035 data[4] = state[12];
1036 data[5] = state[13];
1037 data[6] = state[14];
1038 data[7] = state[15];
1040 state_m[0] = state[0];
1041 state_m[1] = state[1];
1042 state_m[2] = state[2];
1043 state_m[3] = state[3];
1044 state_m[4] = state[4];
1045 state_m[5] = state[5];
1046 state_m[6] = state[6];
1047 state_m[7] = state[7];
1049 data_m[0] = data[0];
1050 data_m[1] = data[1];
1051 data_m[2] = data[2];
1052 data_m[3] = data[3];
1053 data_m[4] = data[4];
1054 data_m[5] = data[5];
1055 data_m[6] = data[6];
1056 data_m[7] = data[7];
1058 PASS0 (state, tmp, state_m, data_m, s_tables);
1059 PASS2 (state, tmp, state_m, data_m, s_tables);
1060 PASS4 (state, tmp, state_m, data_m, s_tables);
1061 PASS6 (state, tmp, state_m, data_m, s_tables);
1063 SHIFT12 (state_m, data, tmp);
1064 SHIFT16 (state, data_m, state_m);
1065 SHIFT61 (state, data_m);
1069 COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
1073 __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)
1079 const u32 gid = get_global_id (0);
1080 const u32 lid = get_local_id (0);
1081 const u32 lsz = get_local_size (0);
1087 __local u32 s_tables[4][256];
1089 for (u32 i = lid; i < 256; i += lsz)
1091 s_tables[0][i] = c_tables[0][i];
1092 s_tables[1][i] = c_tables[1][i];
1093 s_tables[2][i] = c_tables[2][i];
1094 s_tables[3][i] = c_tables[3][i];
1097 barrier (CLK_LOCAL_MEM_FENCE);
1099 if (gid >= gid_max) return;
1107 w0[0] = pws[gid].i[ 0];
1108 w0[1] = pws[gid].i[ 1];
1109 w0[2] = pws[gid].i[ 2];
1110 w0[3] = pws[gid].i[ 3];
1133 const u32 pw_len = pws[gid].pw_len;
1139 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);
1142 __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)
1148 const u32 gid = get_global_id (0);
1149 const u32 lid = get_local_id (0);
1150 const u32 lsz = get_local_size (0);
1156 __local u32 s_tables[4][256];
1158 for (u32 i = lid; i < 256; i += lsz)
1160 s_tables[0][i] = c_tables[0][i];
1161 s_tables[1][i] = c_tables[1][i];
1162 s_tables[2][i] = c_tables[2][i];
1163 s_tables[3][i] = c_tables[3][i];
1166 barrier (CLK_LOCAL_MEM_FENCE);
1168 if (gid >= gid_max) return;
1176 w0[0] = pws[gid].i[ 0];
1177 w0[1] = pws[gid].i[ 1];
1178 w0[2] = pws[gid].i[ 2];
1179 w0[3] = pws[gid].i[ 3];
1183 w1[0] = pws[gid].i[ 4];
1184 w1[1] = pws[gid].i[ 5];
1185 w1[2] = pws[gid].i[ 6];
1186 w1[3] = pws[gid].i[ 7];
1202 const u32 pw_len = pws[gid].pw_len;
1208 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);
1211 __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)
1215 __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)
1221 const u32 gid = get_global_id (0);
1222 const u32 lid = get_local_id (0);
1223 const u32 lsz = get_local_size (0);
1229 __local u32 s_tables[4][256];
1231 for (u32 i = lid; i < 256; i += lsz)
1233 s_tables[0][i] = c_tables[0][i];
1234 s_tables[1][i] = c_tables[1][i];
1235 s_tables[2][i] = c_tables[2][i];
1236 s_tables[3][i] = c_tables[3][i];
1239 barrier (CLK_LOCAL_MEM_FENCE);
1241 if (gid >= gid_max) return;
1249 w0[0] = pws[gid].i[ 0];
1250 w0[1] = pws[gid].i[ 1];
1251 w0[2] = pws[gid].i[ 2];
1252 w0[3] = pws[gid].i[ 3];
1275 const u32 pw_len = pws[gid].pw_len;
1281 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);
1284 __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)
1290 const u32 gid = get_global_id (0);
1291 const u32 lid = get_local_id (0);
1292 const u32 lsz = get_local_size (0);
1298 __local u32 s_tables[4][256];
1300 for (u32 i = lid; i < 256; i += lsz)
1302 s_tables[0][i] = c_tables[0][i];
1303 s_tables[1][i] = c_tables[1][i];
1304 s_tables[2][i] = c_tables[2][i];
1305 s_tables[3][i] = c_tables[3][i];
1308 barrier (CLK_LOCAL_MEM_FENCE);
1310 if (gid >= gid_max) return;
1318 w0[0] = pws[gid].i[ 0];
1319 w0[1] = pws[gid].i[ 1];
1320 w0[2] = pws[gid].i[ 2];
1321 w0[3] = pws[gid].i[ 3];
1325 w1[0] = pws[gid].i[ 4];
1326 w1[1] = pws[gid].i[ 5];
1327 w1[2] = pws[gid].i[ 6];
1328 w1[3] = pws[gid].i[ 7];
1344 const u32 pw_len = pws[gid].pw_len;
1350 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);
1353 __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)