2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
12 #include "include/constants.h"
13 #include "include/kernel_vendor.h"
20 #include "include/kernel_functions.c"
21 #include "OpenCL/types_ocl.c"
22 #include "OpenCL/common.c"
23 #include "OpenCL/simd.c"
25 #define PERM_OP(a,b,tt,n,m) \
35 #define HPERM_OP(a,tt,n,m) \
41 tt = tt >> (16 + n); \
47 PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \
48 PERM_OP (l, r, tt, 16, 0x0000ffff); \
49 PERM_OP (r, l, tt, 2, 0x33333333); \
50 PERM_OP (l, r, tt, 8, 0x00ff00ff); \
51 PERM_OP (r, l, tt, 1, 0x55555555); \
56 PERM_OP (l, r, tt, 1, 0x55555555); \
57 PERM_OP (r, l, tt, 8, 0x00ff00ff); \
58 PERM_OP (l, r, tt, 2, 0x33333333); \
59 PERM_OP (r, l, tt, 16, 0x0000ffff); \
60 PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
63 __constant u32 c_SPtrans[8][64] =
66 0x02080800, 0x00080000, 0x02000002, 0x02080802,
67 0x02000000, 0x00080802, 0x00080002, 0x02000002,
68 0x00080802, 0x02080800, 0x02080000, 0x00000802,
69 0x02000802, 0x02000000, 0x00000000, 0x00080002,
70 0x00080000, 0x00000002, 0x02000800, 0x00080800,
71 0x02080802, 0x02080000, 0x00000802, 0x02000800,
72 0x00000002, 0x00000800, 0x00080800, 0x02080002,
73 0x00000800, 0x02000802, 0x02080002, 0x00000000,
74 0x00000000, 0x02080802, 0x02000800, 0x00080002,
75 0x02080800, 0x00080000, 0x00000802, 0x02000800,
76 0x02080002, 0x00000800, 0x00080800, 0x02000002,
77 0x00080802, 0x00000002, 0x02000002, 0x02080000,
78 0x02080802, 0x00080800, 0x02080000, 0x02000802,
79 0x02000000, 0x00000802, 0x00080002, 0x00000000,
80 0x00080000, 0x02000000, 0x02000802, 0x02080800,
81 0x00000002, 0x02080002, 0x00000800, 0x00080802,
84 0x40108010, 0x00000000, 0x00108000, 0x40100000,
85 0x40000010, 0x00008010, 0x40008000, 0x00108000,
86 0x00008000, 0x40100010, 0x00000010, 0x40008000,
87 0x00100010, 0x40108000, 0x40100000, 0x00000010,
88 0x00100000, 0x40008010, 0x40100010, 0x00008000,
89 0x00108010, 0x40000000, 0x00000000, 0x00100010,
90 0x40008010, 0x00108010, 0x40108000, 0x40000010,
91 0x40000000, 0x00100000, 0x00008010, 0x40108010,
92 0x00100010, 0x40108000, 0x40008000, 0x00108010,
93 0x40108010, 0x00100010, 0x40000010, 0x00000000,
94 0x40000000, 0x00008010, 0x00100000, 0x40100010,
95 0x00008000, 0x40000000, 0x00108010, 0x40008010,
96 0x40108000, 0x00008000, 0x00000000, 0x40000010,
97 0x00000010, 0x40108010, 0x00108000, 0x40100000,
98 0x40100010, 0x00100000, 0x00008010, 0x40008000,
99 0x40008010, 0x00000010, 0x40100000, 0x00108000,
102 0x04000001, 0x04040100, 0x00000100, 0x04000101,
103 0x00040001, 0x04000000, 0x04000101, 0x00040100,
104 0x04000100, 0x00040000, 0x04040000, 0x00000001,
105 0x04040101, 0x00000101, 0x00000001, 0x04040001,
106 0x00000000, 0x00040001, 0x04040100, 0x00000100,
107 0x00000101, 0x04040101, 0x00040000, 0x04000001,
108 0x04040001, 0x04000100, 0x00040101, 0x04040000,
109 0x00040100, 0x00000000, 0x04000000, 0x00040101,
110 0x04040100, 0x00000100, 0x00000001, 0x00040000,
111 0x00000101, 0x00040001, 0x04040000, 0x04000101,
112 0x00000000, 0x04040100, 0x00040100, 0x04040001,
113 0x00040001, 0x04000000, 0x04040101, 0x00000001,
114 0x00040101, 0x04000001, 0x04000000, 0x04040101,
115 0x00040000, 0x04000100, 0x04000101, 0x00040100,
116 0x04000100, 0x00000000, 0x04040001, 0x00000101,
117 0x04000001, 0x00040101, 0x00000100, 0x04040000,
120 0x00401008, 0x10001000, 0x00000008, 0x10401008,
121 0x00000000, 0x10400000, 0x10001008, 0x00400008,
122 0x10401000, 0x10000008, 0x10000000, 0x00001008,
123 0x10000008, 0x00401008, 0x00400000, 0x10000000,
124 0x10400008, 0x00401000, 0x00001000, 0x00000008,
125 0x00401000, 0x10001008, 0x10400000, 0x00001000,
126 0x00001008, 0x00000000, 0x00400008, 0x10401000,
127 0x10001000, 0x10400008, 0x10401008, 0x00400000,
128 0x10400008, 0x00001008, 0x00400000, 0x10000008,
129 0x00401000, 0x10001000, 0x00000008, 0x10400000,
130 0x10001008, 0x00000000, 0x00001000, 0x00400008,
131 0x00000000, 0x10400008, 0x10401000, 0x00001000,
132 0x10000000, 0x10401008, 0x00401008, 0x00400000,
133 0x10401008, 0x00000008, 0x10001000, 0x00401008,
134 0x00400008, 0x00401000, 0x10400000, 0x10001008,
135 0x00001008, 0x10000000, 0x10000008, 0x10401000,
138 0x08000000, 0x00010000, 0x00000400, 0x08010420,
139 0x08010020, 0x08000400, 0x00010420, 0x08010000,
140 0x00010000, 0x00000020, 0x08000020, 0x00010400,
141 0x08000420, 0x08010020, 0x08010400, 0x00000000,
142 0x00010400, 0x08000000, 0x00010020, 0x00000420,
143 0x08000400, 0x00010420, 0x00000000, 0x08000020,
144 0x00000020, 0x08000420, 0x08010420, 0x00010020,
145 0x08010000, 0x00000400, 0x00000420, 0x08010400,
146 0x08010400, 0x08000420, 0x00010020, 0x08010000,
147 0x00010000, 0x00000020, 0x08000020, 0x08000400,
148 0x08000000, 0x00010400, 0x08010420, 0x00000000,
149 0x00010420, 0x08000000, 0x00000400, 0x00010020,
150 0x08000420, 0x00000400, 0x00000000, 0x08010420,
151 0x08010020, 0x08010400, 0x00000420, 0x00010000,
152 0x00010400, 0x08010020, 0x08000400, 0x00000420,
153 0x00000020, 0x00010420, 0x08010000, 0x08000020,
156 0x80000040, 0x00200040, 0x00000000, 0x80202000,
157 0x00200040, 0x00002000, 0x80002040, 0x00200000,
158 0x00002040, 0x80202040, 0x00202000, 0x80000000,
159 0x80002000, 0x80000040, 0x80200000, 0x00202040,
160 0x00200000, 0x80002040, 0x80200040, 0x00000000,
161 0x00002000, 0x00000040, 0x80202000, 0x80200040,
162 0x80202040, 0x80200000, 0x80000000, 0x00002040,
163 0x00000040, 0x00202000, 0x00202040, 0x80002000,
164 0x00002040, 0x80000000, 0x80002000, 0x00202040,
165 0x80202000, 0x00200040, 0x00000000, 0x80002000,
166 0x80000000, 0x00002000, 0x80200040, 0x00200000,
167 0x00200040, 0x80202040, 0x00202000, 0x00000040,
168 0x80202040, 0x00202000, 0x00200000, 0x80002040,
169 0x80000040, 0x80200000, 0x00202040, 0x00000000,
170 0x00002000, 0x80000040, 0x80002040, 0x80202000,
171 0x80200000, 0x00002040, 0x00000040, 0x80200040,
174 0x00004000, 0x00000200, 0x01000200, 0x01000004,
175 0x01004204, 0x00004004, 0x00004200, 0x00000000,
176 0x01000000, 0x01000204, 0x00000204, 0x01004000,
177 0x00000004, 0x01004200, 0x01004000, 0x00000204,
178 0x01000204, 0x00004000, 0x00004004, 0x01004204,
179 0x00000000, 0x01000200, 0x01000004, 0x00004200,
180 0x01004004, 0x00004204, 0x01004200, 0x00000004,
181 0x00004204, 0x01004004, 0x00000200, 0x01000000,
182 0x00004204, 0x01004000, 0x01004004, 0x00000204,
183 0x00004000, 0x00000200, 0x01000000, 0x01004004,
184 0x01000204, 0x00004204, 0x00004200, 0x00000000,
185 0x00000200, 0x01000004, 0x00000004, 0x01000200,
186 0x00000000, 0x01000204, 0x01000200, 0x00004200,
187 0x00000204, 0x00004000, 0x01004204, 0x01000000,
188 0x01004200, 0x00000004, 0x00004004, 0x01004204,
189 0x01000004, 0x01004200, 0x01004000, 0x00004004,
192 0x20800080, 0x20820000, 0x00020080, 0x00000000,
193 0x20020000, 0x00800080, 0x20800000, 0x20820080,
194 0x00000080, 0x20000000, 0x00820000, 0x00020080,
195 0x00820080, 0x20020080, 0x20000080, 0x20800000,
196 0x00020000, 0x00820080, 0x00800080, 0x20020000,
197 0x20820080, 0x20000080, 0x00000000, 0x00820000,
198 0x20000000, 0x00800000, 0x20020080, 0x20800080,
199 0x00800000, 0x00020000, 0x20820000, 0x00000080,
200 0x00800000, 0x00020000, 0x20000080, 0x20820080,
201 0x00020080, 0x20000000, 0x00000000, 0x00820000,
202 0x20800080, 0x20020080, 0x20020000, 0x00800080,
203 0x20820000, 0x00000080, 0x00800080, 0x20020000,
204 0x20820080, 0x00800000, 0x20800000, 0x20000080,
205 0x00820000, 0x00020080, 0x20020080, 0x20800000,
206 0x00000080, 0x20820000, 0x00820080, 0x00000000,
207 0x20000000, 0x20800080, 0x00020000, 0x00820080,
211 __constant u32 c_skb[8][64] =
214 0x00000000, 0x00000010, 0x20000000, 0x20000010,
215 0x00010000, 0x00010010, 0x20010000, 0x20010010,
216 0x00000800, 0x00000810, 0x20000800, 0x20000810,
217 0x00010800, 0x00010810, 0x20010800, 0x20010810,
218 0x00000020, 0x00000030, 0x20000020, 0x20000030,
219 0x00010020, 0x00010030, 0x20010020, 0x20010030,
220 0x00000820, 0x00000830, 0x20000820, 0x20000830,
221 0x00010820, 0x00010830, 0x20010820, 0x20010830,
222 0x00080000, 0x00080010, 0x20080000, 0x20080010,
223 0x00090000, 0x00090010, 0x20090000, 0x20090010,
224 0x00080800, 0x00080810, 0x20080800, 0x20080810,
225 0x00090800, 0x00090810, 0x20090800, 0x20090810,
226 0x00080020, 0x00080030, 0x20080020, 0x20080030,
227 0x00090020, 0x00090030, 0x20090020, 0x20090030,
228 0x00080820, 0x00080830, 0x20080820, 0x20080830,
229 0x00090820, 0x00090830, 0x20090820, 0x20090830,
232 0x00000000, 0x02000000, 0x00002000, 0x02002000,
233 0x00200000, 0x02200000, 0x00202000, 0x02202000,
234 0x00000004, 0x02000004, 0x00002004, 0x02002004,
235 0x00200004, 0x02200004, 0x00202004, 0x02202004,
236 0x00000400, 0x02000400, 0x00002400, 0x02002400,
237 0x00200400, 0x02200400, 0x00202400, 0x02202400,
238 0x00000404, 0x02000404, 0x00002404, 0x02002404,
239 0x00200404, 0x02200404, 0x00202404, 0x02202404,
240 0x10000000, 0x12000000, 0x10002000, 0x12002000,
241 0x10200000, 0x12200000, 0x10202000, 0x12202000,
242 0x10000004, 0x12000004, 0x10002004, 0x12002004,
243 0x10200004, 0x12200004, 0x10202004, 0x12202004,
244 0x10000400, 0x12000400, 0x10002400, 0x12002400,
245 0x10200400, 0x12200400, 0x10202400, 0x12202400,
246 0x10000404, 0x12000404, 0x10002404, 0x12002404,
247 0x10200404, 0x12200404, 0x10202404, 0x12202404,
250 0x00000000, 0x00000001, 0x00040000, 0x00040001,
251 0x01000000, 0x01000001, 0x01040000, 0x01040001,
252 0x00000002, 0x00000003, 0x00040002, 0x00040003,
253 0x01000002, 0x01000003, 0x01040002, 0x01040003,
254 0x00000200, 0x00000201, 0x00040200, 0x00040201,
255 0x01000200, 0x01000201, 0x01040200, 0x01040201,
256 0x00000202, 0x00000203, 0x00040202, 0x00040203,
257 0x01000202, 0x01000203, 0x01040202, 0x01040203,
258 0x08000000, 0x08000001, 0x08040000, 0x08040001,
259 0x09000000, 0x09000001, 0x09040000, 0x09040001,
260 0x08000002, 0x08000003, 0x08040002, 0x08040003,
261 0x09000002, 0x09000003, 0x09040002, 0x09040003,
262 0x08000200, 0x08000201, 0x08040200, 0x08040201,
263 0x09000200, 0x09000201, 0x09040200, 0x09040201,
264 0x08000202, 0x08000203, 0x08040202, 0x08040203,
265 0x09000202, 0x09000203, 0x09040202, 0x09040203,
268 0x00000000, 0x00100000, 0x00000100, 0x00100100,
269 0x00000008, 0x00100008, 0x00000108, 0x00100108,
270 0x00001000, 0x00101000, 0x00001100, 0x00101100,
271 0x00001008, 0x00101008, 0x00001108, 0x00101108,
272 0x04000000, 0x04100000, 0x04000100, 0x04100100,
273 0x04000008, 0x04100008, 0x04000108, 0x04100108,
274 0x04001000, 0x04101000, 0x04001100, 0x04101100,
275 0x04001008, 0x04101008, 0x04001108, 0x04101108,
276 0x00020000, 0x00120000, 0x00020100, 0x00120100,
277 0x00020008, 0x00120008, 0x00020108, 0x00120108,
278 0x00021000, 0x00121000, 0x00021100, 0x00121100,
279 0x00021008, 0x00121008, 0x00021108, 0x00121108,
280 0x04020000, 0x04120000, 0x04020100, 0x04120100,
281 0x04020008, 0x04120008, 0x04020108, 0x04120108,
282 0x04021000, 0x04121000, 0x04021100, 0x04121100,
283 0x04021008, 0x04121008, 0x04021108, 0x04121108,
286 0x00000000, 0x10000000, 0x00010000, 0x10010000,
287 0x00000004, 0x10000004, 0x00010004, 0x10010004,
288 0x20000000, 0x30000000, 0x20010000, 0x30010000,
289 0x20000004, 0x30000004, 0x20010004, 0x30010004,
290 0x00100000, 0x10100000, 0x00110000, 0x10110000,
291 0x00100004, 0x10100004, 0x00110004, 0x10110004,
292 0x20100000, 0x30100000, 0x20110000, 0x30110000,
293 0x20100004, 0x30100004, 0x20110004, 0x30110004,
294 0x00001000, 0x10001000, 0x00011000, 0x10011000,
295 0x00001004, 0x10001004, 0x00011004, 0x10011004,
296 0x20001000, 0x30001000, 0x20011000, 0x30011000,
297 0x20001004, 0x30001004, 0x20011004, 0x30011004,
298 0x00101000, 0x10101000, 0x00111000, 0x10111000,
299 0x00101004, 0x10101004, 0x00111004, 0x10111004,
300 0x20101000, 0x30101000, 0x20111000, 0x30111000,
301 0x20101004, 0x30101004, 0x20111004, 0x30111004,
304 0x00000000, 0x08000000, 0x00000008, 0x08000008,
305 0x00000400, 0x08000400, 0x00000408, 0x08000408,
306 0x00020000, 0x08020000, 0x00020008, 0x08020008,
307 0x00020400, 0x08020400, 0x00020408, 0x08020408,
308 0x00000001, 0x08000001, 0x00000009, 0x08000009,
309 0x00000401, 0x08000401, 0x00000409, 0x08000409,
310 0x00020001, 0x08020001, 0x00020009, 0x08020009,
311 0x00020401, 0x08020401, 0x00020409, 0x08020409,
312 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
313 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
314 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
315 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
316 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
317 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
318 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
319 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
322 0x00000000, 0x00000100, 0x00080000, 0x00080100,
323 0x01000000, 0x01000100, 0x01080000, 0x01080100,
324 0x00000010, 0x00000110, 0x00080010, 0x00080110,
325 0x01000010, 0x01000110, 0x01080010, 0x01080110,
326 0x00200000, 0x00200100, 0x00280000, 0x00280100,
327 0x01200000, 0x01200100, 0x01280000, 0x01280100,
328 0x00200010, 0x00200110, 0x00280010, 0x00280110,
329 0x01200010, 0x01200110, 0x01280010, 0x01280110,
330 0x00000200, 0x00000300, 0x00080200, 0x00080300,
331 0x01000200, 0x01000300, 0x01080200, 0x01080300,
332 0x00000210, 0x00000310, 0x00080210, 0x00080310,
333 0x01000210, 0x01000310, 0x01080210, 0x01080310,
334 0x00200200, 0x00200300, 0x00280200, 0x00280300,
335 0x01200200, 0x01200300, 0x01280200, 0x01280300,
336 0x00200210, 0x00200310, 0x00280210, 0x00280310,
337 0x01200210, 0x01200310, 0x01280210, 0x01280310,
340 0x00000000, 0x04000000, 0x00040000, 0x04040000,
341 0x00000002, 0x04000002, 0x00040002, 0x04040002,
342 0x00002000, 0x04002000, 0x00042000, 0x04042000,
343 0x00002002, 0x04002002, 0x00042002, 0x04042002,
344 0x00000020, 0x04000020, 0x00040020, 0x04040020,
345 0x00000022, 0x04000022, 0x00040022, 0x04040022,
346 0x00002020, 0x04002020, 0x00042020, 0x04042020,
347 0x00002022, 0x04002022, 0x00042022, 0x04042022,
348 0x00000800, 0x04000800, 0x00040800, 0x04040800,
349 0x00000802, 0x04000802, 0x00040802, 0x04040802,
350 0x00002800, 0x04002800, 0x00042800, 0x04042800,
351 0x00002802, 0x04002802, 0x00042802, 0x04042802,
352 0x00000820, 0x04000820, 0x00040820, 0x04040820,
353 0x00000822, 0x04000822, 0x00040822, 0x04040822,
354 0x00002820, 0x04002820, 0x00042820, 0x04042820,
355 0x00002822, 0x04002822, 0x00042822, 0x04042822
360 #define BOX(i,n,S) (S)[(n)][(i)]
362 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
364 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
366 #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])
369 static void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
382 for (u32 i = 0; i < 16; i += 2)
388 t = Kd[i + 0] ^ rotl32 (r, 28u);
390 l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
391 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
392 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
393 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
394 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
395 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
396 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
397 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
400 t = Kd[i + 1] ^ rotl32 (l, 28u);
402 r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
403 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
404 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
405 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
406 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
407 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
408 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
409 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
421 static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
425 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
426 HPERM_OP (c, tt, 2, 0xcccc0000);
427 HPERM_OP (d, tt, 2, 0xcccc0000);
428 PERM_OP (d, c, tt, 1, 0x55555555);
429 PERM_OP (c, d, tt, 8, 0x00ff00ff);
430 PERM_OP (d, c, tt, 1, 0x55555555);
432 d = ((d & 0x000000ff) << 16)
433 | ((d & 0x0000ff00) << 0)
434 | ((d & 0x00ff0000) >> 16)
435 | ((c & 0xf0000000) >> 4);
440 for (u32 i = 0; i < 16; i++)
442 if ((i < 2) || (i == 8) || (i == 15))
444 c = ((c >> 1) | (c << 27));
445 d = ((d >> 1) | (d << 27));
449 c = ((c >> 2) | (c << 26));
450 d = ((d >> 2) | (d << 26));
456 const u32x c00 = (c >> 0) & 0x0000003f;
457 const u32x c06 = (c >> 6) & 0x00383003;
458 const u32x c07 = (c >> 7) & 0x0000003c;
459 const u32x c13 = (c >> 13) & 0x0000060f;
460 const u32x c20 = (c >> 20) & 0x00000001;
462 u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
463 | BOX (((c06 >> 0) & 0xff)
464 |((c07 >> 0) & 0xff), 1, s_skb)
465 | BOX (((c13 >> 0) & 0xff)
466 |((c06 >> 8) & 0xff), 2, s_skb)
467 | BOX (((c20 >> 0) & 0xff)
469 |((c06 >> 16) & 0xff), 3, s_skb);
471 const u32x d00 = (d >> 0) & 0x00003c3f;
472 const u32x d07 = (d >> 7) & 0x00003f03;
473 const u32x d21 = (d >> 21) & 0x0000000f;
474 const u32x d22 = (d >> 22) & 0x00000030;
476 u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
477 | BOX (((d07 >> 0) & 0xff)
478 |((d00 >> 8) & 0xff), 5, s_skb)
479 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
480 | BOX (((d21 >> 0) & 0xff)
481 |((d22 >> 0) & 0xff), 7, s_skb);
483 Kc[i] = ((t << 16) | (s & 0x0000ffff));
484 Kd[i] = ((s >> 16) | (t & 0xffff0000));
486 Kc[i] = rotl32 (Kc[i], 2u);
487 Kd[i] = rotl32 (Kd[i], 2u);
491 static void m03100m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
497 const u32 gid = get_global_id (0);
498 const u32 lid = get_local_id (0);
506 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
507 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
508 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
509 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
513 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
514 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
515 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
516 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
525 const u32 salt_len = salt_bufs[salt_pos].salt_len;
527 const u32 salt_word_len = (salt_len + pw_len) * 2;
555 switch_buffer_by_offset_le_S (w0_t, w1_t, w2_t, w3_t, salt_len);
557 w0_t[0] |= salt_buf0[0];
558 w0_t[1] |= salt_buf0[1];
559 w0_t[2] |= salt_buf0[2];
560 w0_t[3] |= salt_buf0[3];
561 w1_t[0] |= salt_buf1[0];
562 w1_t[1] |= salt_buf1[1];
563 w1_t[2] |= salt_buf1[2];
564 w1_t[3] |= salt_buf1[3];
565 w2_t[0] |= salt_buf2[0];
566 w2_t[1] |= salt_buf2[1];
567 w2_t[2] |= salt_buf2[2];
568 w2_t[3] |= salt_buf2[3];
599 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE)
601 const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
603 const u32x w0 = w0l | w0r;
605 overwrite_at_le (dst, w0, salt_len);
608 * precompute key1 since key is static: 0x0123456789abcdef
609 * plus LEFT_ROTATE by 2
651 * key1 (generate key)
659 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
663 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
664 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
669 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
673 * key2 (generate hash)
676 _des_crypt_keysetup (iv[0], iv[1], Kc, Kd, s_skb);
681 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
685 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
686 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
691 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
701 COMPARE_M_SIMD (iv[0], iv[1], c, d);
705 static void m03100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
711 const u32 gid = get_global_id (0);
712 const u32 lid = get_local_id (0);
720 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
721 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
722 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
723 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
727 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
728 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
729 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
730 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
739 const u32 salt_len = salt_bufs[salt_pos].salt_len;
741 const u32 salt_word_len = (salt_len + pw_len) * 2;
769 switch_buffer_by_offset_le_S (w0_t, w1_t, w2_t, w3_t, salt_len);
771 w0_t[0] |= salt_buf0[0];
772 w0_t[1] |= salt_buf0[1];
773 w0_t[2] |= salt_buf0[2];
774 w0_t[3] |= salt_buf0[3];
775 w1_t[0] |= salt_buf1[0];
776 w1_t[1] |= salt_buf1[1];
777 w1_t[2] |= salt_buf1[2];
778 w1_t[3] |= salt_buf1[3];
779 w2_t[0] |= salt_buf2[0];
780 w2_t[1] |= salt_buf2[1];
781 w2_t[2] |= salt_buf2[2];
782 w2_t[3] |= salt_buf2[3];
811 const u32 search[4] =
813 digests_buf[digests_offset].digest_buf[DGST_R0],
814 digests_buf[digests_offset].digest_buf[DGST_R1],
815 digests_buf[digests_offset].digest_buf[DGST_R2],
816 digests_buf[digests_offset].digest_buf[DGST_R3]
825 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE)
827 const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
829 const u32x w0 = w0l | w0r;
831 overwrite_at_le (dst, w0, salt_len);
834 * precompute key1 since key is static: 0x0123456789abcdef
835 * plus LEFT_ROTATE by 2
877 * key1 (generate key)
885 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
889 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
890 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
895 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
899 * key2 (generate hash)
902 _des_crypt_keysetup (iv[0], iv[1], Kc, Kd, s_skb);
907 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
911 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
912 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
917 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
927 COMPARE_S_SIMD (iv[0], iv[1], c, d);
931 __kernel void m03100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
937 const u32 gid = get_global_id (0);
938 const u32 lid = get_local_id (0);
939 const u32 lsz = get_local_size (0);
947 w[ 0] = pws[gid].i[ 0];
948 w[ 1] = pws[gid].i[ 1];
949 w[ 2] = pws[gid].i[ 2];
950 w[ 3] = pws[gid].i[ 3];
964 const u32 pw_len = pws[gid].pw_len;
970 __local u32 s_SPtrans[8][64];
971 __local u32 s_skb[8][64];
973 for (u32 i = lid; i < 64; i += lsz)
975 s_SPtrans[0][i] = c_SPtrans[0][i];
976 s_SPtrans[1][i] = c_SPtrans[1][i];
977 s_SPtrans[2][i] = c_SPtrans[2][i];
978 s_SPtrans[3][i] = c_SPtrans[3][i];
979 s_SPtrans[4][i] = c_SPtrans[4][i];
980 s_SPtrans[5][i] = c_SPtrans[5][i];
981 s_SPtrans[6][i] = c_SPtrans[6][i];
982 s_SPtrans[7][i] = c_SPtrans[7][i];
984 s_skb[0][i] = c_skb[0][i];
985 s_skb[1][i] = c_skb[1][i];
986 s_skb[2][i] = c_skb[2][i];
987 s_skb[3][i] = c_skb[3][i];
988 s_skb[4][i] = c_skb[4][i];
989 s_skb[5][i] = c_skb[5][i];
990 s_skb[6][i] = c_skb[6][i];
991 s_skb[7][i] = c_skb[7][i];
994 barrier (CLK_LOCAL_MEM_FENCE);
996 if (gid >= gid_max) return;
1002 m03100m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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, bfs_cnt, digests_cnt, digests_offset);
1005 __kernel void m03100_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1011 const u32 gid = get_global_id (0);
1012 const u32 lid = get_local_id (0);
1013 const u32 lsz = get_local_size (0);
1021 w[ 0] = pws[gid].i[ 0];
1022 w[ 1] = pws[gid].i[ 1];
1023 w[ 2] = pws[gid].i[ 2];
1024 w[ 3] = pws[gid].i[ 3];
1025 w[ 4] = pws[gid].i[ 4];
1026 w[ 5] = pws[gid].i[ 5];
1027 w[ 6] = pws[gid].i[ 6];
1028 w[ 7] = pws[gid].i[ 7];
1038 const u32 pw_len = pws[gid].pw_len;
1044 __local u32 s_SPtrans[8][64];
1045 __local u32 s_skb[8][64];
1047 for (u32 i = lid; i < 64; i += lsz)
1049 s_SPtrans[0][i] = c_SPtrans[0][i];
1050 s_SPtrans[1][i] = c_SPtrans[1][i];
1051 s_SPtrans[2][i] = c_SPtrans[2][i];
1052 s_SPtrans[3][i] = c_SPtrans[3][i];
1053 s_SPtrans[4][i] = c_SPtrans[4][i];
1054 s_SPtrans[5][i] = c_SPtrans[5][i];
1055 s_SPtrans[6][i] = c_SPtrans[6][i];
1056 s_SPtrans[7][i] = c_SPtrans[7][i];
1058 s_skb[0][i] = c_skb[0][i];
1059 s_skb[1][i] = c_skb[1][i];
1060 s_skb[2][i] = c_skb[2][i];
1061 s_skb[3][i] = c_skb[3][i];
1062 s_skb[4][i] = c_skb[4][i];
1063 s_skb[5][i] = c_skb[5][i];
1064 s_skb[6][i] = c_skb[6][i];
1065 s_skb[7][i] = c_skb[7][i];
1068 barrier (CLK_LOCAL_MEM_FENCE);
1070 if (gid >= gid_max) return;
1076 m03100m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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, bfs_cnt, digests_cnt, digests_offset);
1079 __kernel void m03100_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1083 __kernel void m03100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1089 const u32 gid = get_global_id (0);
1090 const u32 lid = get_local_id (0);
1091 const u32 lsz = get_local_size (0);
1099 w[ 0] = pws[gid].i[ 0];
1100 w[ 1] = pws[gid].i[ 1];
1101 w[ 2] = pws[gid].i[ 2];
1102 w[ 3] = pws[gid].i[ 3];
1116 const u32 pw_len = pws[gid].pw_len;
1122 __local u32 s_SPtrans[8][64];
1123 __local u32 s_skb[8][64];
1125 for (u32 i = lid; i < 64; i += lsz)
1127 s_SPtrans[0][i] = c_SPtrans[0][i];
1128 s_SPtrans[1][i] = c_SPtrans[1][i];
1129 s_SPtrans[2][i] = c_SPtrans[2][i];
1130 s_SPtrans[3][i] = c_SPtrans[3][i];
1131 s_SPtrans[4][i] = c_SPtrans[4][i];
1132 s_SPtrans[5][i] = c_SPtrans[5][i];
1133 s_SPtrans[6][i] = c_SPtrans[6][i];
1134 s_SPtrans[7][i] = c_SPtrans[7][i];
1136 s_skb[0][i] = c_skb[0][i];
1137 s_skb[1][i] = c_skb[1][i];
1138 s_skb[2][i] = c_skb[2][i];
1139 s_skb[3][i] = c_skb[3][i];
1140 s_skb[4][i] = c_skb[4][i];
1141 s_skb[5][i] = c_skb[5][i];
1142 s_skb[6][i] = c_skb[6][i];
1143 s_skb[7][i] = c_skb[7][i];
1146 barrier (CLK_LOCAL_MEM_FENCE);
1148 if (gid >= gid_max) return;
1154 m03100s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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, bfs_cnt, digests_cnt, digests_offset);
1157 __kernel void m03100_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1163 const u32 gid = get_global_id (0);
1164 const u32 lid = get_local_id (0);
1165 const u32 lsz = get_local_size (0);
1173 w[ 0] = pws[gid].i[ 0];
1174 w[ 1] = pws[gid].i[ 1];
1175 w[ 2] = pws[gid].i[ 2];
1176 w[ 3] = pws[gid].i[ 3];
1177 w[ 4] = pws[gid].i[ 4];
1178 w[ 5] = pws[gid].i[ 5];
1179 w[ 6] = pws[gid].i[ 6];
1180 w[ 7] = pws[gid].i[ 7];
1190 const u32 pw_len = pws[gid].pw_len;
1196 __local u32 s_SPtrans[8][64];
1197 __local u32 s_skb[8][64];
1199 for (u32 i = lid; i < 64; i += lsz)
1201 s_SPtrans[0][i] = c_SPtrans[0][i];
1202 s_SPtrans[1][i] = c_SPtrans[1][i];
1203 s_SPtrans[2][i] = c_SPtrans[2][i];
1204 s_SPtrans[3][i] = c_SPtrans[3][i];
1205 s_SPtrans[4][i] = c_SPtrans[4][i];
1206 s_SPtrans[5][i] = c_SPtrans[5][i];
1207 s_SPtrans[6][i] = c_SPtrans[6][i];
1208 s_SPtrans[7][i] = c_SPtrans[7][i];
1210 s_skb[0][i] = c_skb[0][i];
1211 s_skb[1][i] = c_skb[1][i];
1212 s_skb[2][i] = c_skb[2][i];
1213 s_skb[3][i] = c_skb[3][i];
1214 s_skb[4][i] = c_skb[4][i];
1215 s_skb[5][i] = c_skb[5][i];
1216 s_skb[6][i] = c_skb[6][i];
1217 s_skb[7][i] = c_skb[7][i];
1220 barrier (CLK_LOCAL_MEM_FENCE);
1222 if (gid >= gid_max) return;
1228 m03100s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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, bfs_cnt, digests_cnt, digests_offset);
1231 __kernel void m03100_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)