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 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);
724 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
726 const u32x w0r = ix_create_bft (bfs_buf, il_pos);
728 const u32x w0lr = w0l | w0r;
769 state_m[0] = state[0];
770 state_m[1] = state[1];
771 state_m[2] = state[2];
772 state_m[3] = state[3];
773 state_m[4] = state[4];
774 state_m[5] = state[5];
775 state_m[6] = state[6];
776 state_m[7] = state[7];
789 //if (pw_len > 0) // not really SIMD compatible
791 PASS0 (state, tmp, state_m, data_m, s_tables);
792 PASS2 (state, tmp, state_m, data_m, s_tables);
793 PASS4 (state, tmp, state_m, data_m, s_tables);
794 PASS6 (state, tmp, state_m, data_m, s_tables);
796 SHIFT12 (state_m, data, tmp);
797 SHIFT16 (state, data_m, state_m);
798 SHIFT61 (state, data_m);
801 data[0] = pw_len * 8;
812 state_m[0] = state[0];
813 state_m[1] = state[1];
814 state_m[2] = state[2];
815 state_m[3] = state[3];
816 state_m[4] = state[4];
817 state_m[5] = state[5];
818 state_m[6] = state[6];
819 state_m[7] = state[7];
830 PASS0 (state, tmp, state_m, data_m, s_tables);
831 PASS2 (state, tmp, state_m, data_m, s_tables);
832 PASS4 (state, tmp, state_m, data_m, s_tables);
833 PASS6 (state, tmp, state_m, data_m, s_tables);
835 SHIFT12 (state_m, data, tmp);
836 SHIFT16 (state, data_m, state_m);
837 SHIFT61 (state, data_m);
850 state_m[0] = state[0];
851 state_m[1] = state[1];
852 state_m[2] = state[2];
853 state_m[3] = state[3];
854 state_m[4] = state[4];
855 state_m[5] = state[5];
856 state_m[6] = state[6];
857 state_m[7] = state[7];
868 PASS0 (state, tmp, state_m, data_m, s_tables);
869 PASS2 (state, tmp, state_m, data_m, s_tables);
870 PASS4 (state, tmp, state_m, data_m, s_tables);
871 PASS6 (state, tmp, state_m, data_m, s_tables);
873 SHIFT12 (state_m, data, tmp);
874 SHIFT16 (state, data_m, state_m);
875 SHIFT61 (state, data_m);
879 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
883 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])
889 const u32 gid = get_global_id (0);
890 const u32 lid = get_local_id (0);
896 const u32 search[4] =
898 digests_buf[digests_offset].digest_buf[DGST_R0],
899 digests_buf[digests_offset].digest_buf[DGST_R1],
900 digests_buf[digests_offset].digest_buf[DGST_R2],
901 digests_buf[digests_offset].digest_buf[DGST_R3]
910 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
912 const u32x w0r = ix_create_bft (bfs_buf, il_pos);
914 const u32x w0lr = w0l | w0r;
955 state_m[0] = state[0];
956 state_m[1] = state[1];
957 state_m[2] = state[2];
958 state_m[3] = state[3];
959 state_m[4] = state[4];
960 state_m[5] = state[5];
961 state_m[6] = state[6];
962 state_m[7] = state[7];
975 //if (pw_len > 0) // not really SIMD compatible
977 PASS0 (state, tmp, state_m, data_m, s_tables);
978 PASS2 (state, tmp, state_m, data_m, s_tables);
979 PASS4 (state, tmp, state_m, data_m, s_tables);
980 PASS6 (state, tmp, state_m, data_m, s_tables);
982 SHIFT12 (state_m, data, tmp);
983 SHIFT16 (state, data_m, state_m);
984 SHIFT61 (state, data_m);
987 data[0] = pw_len * 8;
998 state_m[0] = state[0];
999 state_m[1] = state[1];
1000 state_m[2] = state[2];
1001 state_m[3] = state[3];
1002 state_m[4] = state[4];
1003 state_m[5] = state[5];
1004 state_m[6] = state[6];
1005 state_m[7] = state[7];
1007 data_m[0] = data[0];
1008 data_m[1] = data[1];
1009 data_m[2] = data[2];
1010 data_m[3] = data[3];
1011 data_m[4] = data[4];
1012 data_m[5] = data[5];
1013 data_m[6] = data[6];
1014 data_m[7] = data[7];
1016 PASS0 (state, tmp, state_m, data_m, s_tables);
1017 PASS2 (state, tmp, state_m, data_m, s_tables);
1018 PASS4 (state, tmp, state_m, data_m, s_tables);
1019 PASS6 (state, tmp, state_m, data_m, s_tables);
1021 SHIFT12 (state_m, data, tmp);
1022 SHIFT16 (state, data_m, state_m);
1023 SHIFT61 (state, data_m);
1027 data[0] = state[ 8];
1028 data[1] = state[ 9];
1029 data[2] = state[10];
1030 data[3] = state[11];
1031 data[4] = state[12];
1032 data[5] = state[13];
1033 data[6] = state[14];
1034 data[7] = state[15];
1036 state_m[0] = state[0];
1037 state_m[1] = state[1];
1038 state_m[2] = state[2];
1039 state_m[3] = state[3];
1040 state_m[4] = state[4];
1041 state_m[5] = state[5];
1042 state_m[6] = state[6];
1043 state_m[7] = state[7];
1045 data_m[0] = data[0];
1046 data_m[1] = data[1];
1047 data_m[2] = data[2];
1048 data_m[3] = data[3];
1049 data_m[4] = data[4];
1050 data_m[5] = data[5];
1051 data_m[6] = data[6];
1052 data_m[7] = data[7];
1054 PASS0 (state, tmp, state_m, data_m, s_tables);
1055 PASS2 (state, tmp, state_m, data_m, s_tables);
1056 PASS4 (state, tmp, state_m, data_m, s_tables);
1057 PASS6 (state, tmp, state_m, data_m, s_tables);
1059 SHIFT12 (state_m, data, tmp);
1060 SHIFT16 (state, data_m, state_m);
1061 SHIFT61 (state, data_m);
1065 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
1069 __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)
1075 const u32 gid = get_global_id (0);
1076 const u32 lid = get_local_id (0);
1077 const u32 lsz = get_local_size (0);
1083 __local u32 s_tables[4][256];
1085 for (u32 i = lid; i < 256; i += lsz)
1087 s_tables[0][i] = c_tables[0][i];
1088 s_tables[1][i] = c_tables[1][i];
1089 s_tables[2][i] = c_tables[2][i];
1090 s_tables[3][i] = c_tables[3][i];
1093 barrier (CLK_LOCAL_MEM_FENCE);
1095 if (gid >= gid_max) return;
1103 w0[0] = pws[gid].i[ 0];
1104 w0[1] = pws[gid].i[ 1];
1105 w0[2] = pws[gid].i[ 2];
1106 w0[3] = pws[gid].i[ 3];
1129 const u32 pw_len = pws[gid].pw_len;
1135 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);
1138 __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)
1144 const u32 gid = get_global_id (0);
1145 const u32 lid = get_local_id (0);
1146 const u32 lsz = get_local_size (0);
1152 __local u32 s_tables[4][256];
1154 for (u32 i = lid; i < 256; i += lsz)
1156 s_tables[0][i] = c_tables[0][i];
1157 s_tables[1][i] = c_tables[1][i];
1158 s_tables[2][i] = c_tables[2][i];
1159 s_tables[3][i] = c_tables[3][i];
1162 barrier (CLK_LOCAL_MEM_FENCE);
1164 if (gid >= gid_max) return;
1172 w0[0] = pws[gid].i[ 0];
1173 w0[1] = pws[gid].i[ 1];
1174 w0[2] = pws[gid].i[ 2];
1175 w0[3] = pws[gid].i[ 3];
1179 w1[0] = pws[gid].i[ 4];
1180 w1[1] = pws[gid].i[ 5];
1181 w1[2] = pws[gid].i[ 6];
1182 w1[3] = pws[gid].i[ 7];
1198 const u32 pw_len = pws[gid].pw_len;
1204 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);
1207 __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)
1211 __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)
1217 const u32 gid = get_global_id (0);
1218 const u32 lid = get_local_id (0);
1219 const u32 lsz = get_local_size (0);
1225 __local u32 s_tables[4][256];
1227 for (u32 i = lid; i < 256; i += lsz)
1229 s_tables[0][i] = c_tables[0][i];
1230 s_tables[1][i] = c_tables[1][i];
1231 s_tables[2][i] = c_tables[2][i];
1232 s_tables[3][i] = c_tables[3][i];
1235 barrier (CLK_LOCAL_MEM_FENCE);
1237 if (gid >= gid_max) return;
1245 w0[0] = pws[gid].i[ 0];
1246 w0[1] = pws[gid].i[ 1];
1247 w0[2] = pws[gid].i[ 2];
1248 w0[3] = pws[gid].i[ 3];
1271 const u32 pw_len = pws[gid].pw_len;
1277 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);
1280 __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)
1286 const u32 gid = get_global_id (0);
1287 const u32 lid = get_local_id (0);
1288 const u32 lsz = get_local_size (0);
1294 __local u32 s_tables[4][256];
1296 for (u32 i = lid; i < 256; i += lsz)
1298 s_tables[0][i] = c_tables[0][i];
1299 s_tables[1][i] = c_tables[1][i];
1300 s_tables[2][i] = c_tables[2][i];
1301 s_tables[3][i] = c_tables[3][i];
1304 barrier (CLK_LOCAL_MEM_FENCE);
1306 if (gid >= gid_max) return;
1314 w0[0] = pws[gid].i[ 0];
1315 w0[1] = pws[gid].i[ 1];
1316 w0[2] = pws[gid].i[ 2];
1317 w0[3] = pws[gid].i[ 3];
1321 w1[0] = pws[gid].i[ 4];
1322 w1[1] = pws[gid].i[ 5];
1323 w1[2] = pws[gid].i[ 6];
1324 w1[3] = pws[gid].i[ 7];
1340 const u32 pw_len = pws[gid].pw_len;
1346 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);
1349 __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)