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
359 #define BOX(i,n,S) (S)[(n)][(i)]
361 static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
374 for (u32 i = 0; i < 16; i += 2)
380 t = Kd[i + 0] ^ rotl32 (r, 28u);
382 l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
383 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
384 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
385 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
386 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
387 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
388 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
389 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
392 t = Kd[i + 1] ^ rotl32 (l, 28u);
394 r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
395 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
396 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
397 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
398 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
399 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
400 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
401 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
413 static void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64])
417 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
418 HPERM_OP (c, tt, 2, 0xcccc0000);
419 HPERM_OP (d, tt, 2, 0xcccc0000);
420 PERM_OP (d, c, tt, 1, 0x55555555);
421 PERM_OP (c, d, tt, 8, 0x00ff00ff);
422 PERM_OP (d, c, tt, 1, 0x55555555);
424 d = ((d & 0x000000ff) << 16)
425 | ((d & 0x0000ff00) << 0)
426 | ((d & 0x00ff0000) >> 16)
427 | ((c & 0xf0000000) >> 4);
432 for (u32 i = 0; i < 16; i++)
434 if ((i < 2) || (i == 8) || (i == 15))
436 c = ((c >> 1) | (c << 27));
437 d = ((d >> 1) | (d << 27));
441 c = ((c >> 2) | (c << 26));
442 d = ((d >> 2) | (d << 26));
448 const u32 c00 = (c >> 0) & 0x0000003f;
449 const u32 c06 = (c >> 6) & 0x00383003;
450 const u32 c07 = (c >> 7) & 0x0000003c;
451 const u32 c13 = (c >> 13) & 0x0000060f;
452 const u32 c20 = (c >> 20) & 0x00000001;
454 u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
455 | BOX (((c06 >> 0) & 0xff)
456 |((c07 >> 0) & 0xff), 1, s_skb)
457 | BOX (((c13 >> 0) & 0xff)
458 |((c06 >> 8) & 0xff), 2, s_skb)
459 | BOX (((c20 >> 0) & 0xff)
461 |((c06 >> 16) & 0xff), 3, s_skb);
463 const u32 d00 = (d >> 0) & 0x00003c3f;
464 const u32 d07 = (d >> 7) & 0x00003f03;
465 const u32 d21 = (d >> 21) & 0x0000000f;
466 const u32 d22 = (d >> 22) & 0x00000030;
468 u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
469 | BOX (((d07 >> 0) & 0xff)
470 |((d00 >> 8) & 0xff), 5, s_skb)
471 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
472 | BOX (((d21 >> 0) & 0xff)
473 |((d22 >> 0) & 0xff), 7, s_skb);
475 Kc[i] = ((t << 16) | (s & 0x0000ffff));
476 Kd[i] = ((s >> 16) | (t & 0xffff0000));
478 Kc[i] = rotl32 (Kc[i], 2u);
479 Kd[i] = rotl32 (Kd[i], 2u);
483 __kernel void m03100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
489 const u32 gid = get_global_id (0);
490 const u32 lid = get_local_id (0);
491 const u32 lsz = get_local_size (0);
497 __local u32 s_SPtrans[8][64];
498 __local u32 s_skb[8][64];
500 for (u32 i = lid; i < 64; i += lsz)
502 s_SPtrans[0][i] = c_SPtrans[0][i];
503 s_SPtrans[1][i] = c_SPtrans[1][i];
504 s_SPtrans[2][i] = c_SPtrans[2][i];
505 s_SPtrans[3][i] = c_SPtrans[3][i];
506 s_SPtrans[4][i] = c_SPtrans[4][i];
507 s_SPtrans[5][i] = c_SPtrans[5][i];
508 s_SPtrans[6][i] = c_SPtrans[6][i];
509 s_SPtrans[7][i] = c_SPtrans[7][i];
511 s_skb[0][i] = c_skb[0][i];
512 s_skb[1][i] = c_skb[1][i];
513 s_skb[2][i] = c_skb[2][i];
514 s_skb[3][i] = c_skb[3][i];
515 s_skb[4][i] = c_skb[4][i];
516 s_skb[5][i] = c_skb[5][i];
517 s_skb[6][i] = c_skb[6][i];
518 s_skb[7][i] = c_skb[7][i];
521 barrier (CLK_LOCAL_MEM_FENCE);
523 if (gid >= gid_max) return;
531 wordl0[0] = pws[gid].i[ 0];
532 wordl0[1] = pws[gid].i[ 1];
533 wordl0[2] = pws[gid].i[ 2];
534 wordl0[3] = pws[gid].i[ 3];
538 wordl1[0] = pws[gid].i[ 4];
539 wordl1[1] = pws[gid].i[ 5];
540 wordl1[2] = pws[gid].i[ 6];
541 wordl1[3] = pws[gid].i[ 7];
557 const u32 pw_l_len = pws[gid].pw_len;
559 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
561 switch_buffer_by_offset_le_S (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
571 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
572 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
573 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
574 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
575 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
576 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
577 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
578 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
580 const u32 salt_len = salt_bufs[salt_pos].salt_len;
586 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos += VECT_SIZE)
588 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
590 const u32x pw_len = pw_l_len + pw_r_len;
592 const u32 salt_word_len = (salt_len + pw_len) * 2;
594 u32x wordr0[4] = { 0 };
595 u32x wordr1[4] = { 0 };
596 u32x wordr2[4] = { 0 };
597 u32x wordr3[4] = { 0 };
599 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
600 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
601 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
602 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
603 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
604 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
605 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
606 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
608 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
610 switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
615 w0[0] = wordl0[0] | wordr0[0];
616 w0[1] = wordl0[1] | wordr0[1];
617 w0[2] = wordl0[2] | wordr0[2];
618 w0[3] = wordl0[3] | wordr0[3];
622 w1[0] = wordl1[0] | wordr1[0];
623 w1[1] = wordl1[1] | wordr1[1];
624 w1[2] = wordl1[2] | wordr1[2];
625 w1[3] = wordl1[3] | wordr1[3];
629 w2[0] = wordl2[0] | wordr2[0];
630 w2[1] = wordl2[1] | wordr2[1];
631 w2[2] = wordl2[2] | wordr2[2];
632 w2[3] = wordl2[3] | wordr2[3];
636 w3[0] = wordl3[0] | wordr3[0];
637 w3[1] = wordl3[1] | wordr3[1];
638 w3[2] = wordl3[2] | wordr3[2];
639 w3[3] = wordl3[3] | wordr3[3];
667 switch_buffer_by_offset_le (w0_t, w1_t, w2_t, w3_t, salt_len);
669 w0_t[0] |= salt_buf0[0];
670 w0_t[1] |= salt_buf0[1];
671 w0_t[2] |= salt_buf0[2];
672 w0_t[3] |= salt_buf0[3];
673 w1_t[0] |= salt_buf1[0];
674 w1_t[1] |= salt_buf1[1];
675 w1_t[2] |= salt_buf1[2];
676 w1_t[3] |= salt_buf1[3];
698 * precompute key1 since key is static: 0x0123456789abcdef
699 * plus LEFT_ROTATE by 2
741 * key1 (generate key)
749 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
753 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
754 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
759 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
763 * key2 (generate hash)
766 _des_crypt_keysetup (iv[0], iv[1], Kc, Kd, s_skb);
771 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
775 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
776 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
781 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
791 COMPARE_M_SIMD (iv[0], iv[1], c, d);
795 __kernel void m03100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
799 __kernel void m03100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
803 __kernel void m03100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
809 const u32 gid = get_global_id (0);
810 const u32 lid = get_local_id (0);
811 const u32 lsz = get_local_size (0);
817 __local u32 s_SPtrans[8][64];
818 __local u32 s_skb[8][64];
820 for (u32 i = lid; i < 64; i += lsz)
822 s_SPtrans[0][i] = c_SPtrans[0][i];
823 s_SPtrans[1][i] = c_SPtrans[1][i];
824 s_SPtrans[2][i] = c_SPtrans[2][i];
825 s_SPtrans[3][i] = c_SPtrans[3][i];
826 s_SPtrans[4][i] = c_SPtrans[4][i];
827 s_SPtrans[5][i] = c_SPtrans[5][i];
828 s_SPtrans[6][i] = c_SPtrans[6][i];
829 s_SPtrans[7][i] = c_SPtrans[7][i];
831 s_skb[0][i] = c_skb[0][i];
832 s_skb[1][i] = c_skb[1][i];
833 s_skb[2][i] = c_skb[2][i];
834 s_skb[3][i] = c_skb[3][i];
835 s_skb[4][i] = c_skb[4][i];
836 s_skb[5][i] = c_skb[5][i];
837 s_skb[6][i] = c_skb[6][i];
838 s_skb[7][i] = c_skb[7][i];
841 barrier (CLK_LOCAL_MEM_FENCE);
843 if (gid >= gid_max) return;
852 pws0[0] = pws[gid].i[0];
853 pws0[1] = pws[gid].i[1];
854 pws0[2] = pws[gid].i[2];
855 pws0[3] = pws[gid].i[3];
856 pws1[0] = pws[gid].i[4];
857 pws1[1] = pws[gid].i[5];
858 pws1[2] = pws[gid].i[6];
859 pws1[3] = pws[gid].i[7];
861 const u32 pw_l_len = pws[gid].pw_len;
870 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
871 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
872 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
873 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
874 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
875 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
876 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
877 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
879 const u32 salt_len = salt_bufs[salt_pos].salt_len;
885 const u32 search[4] =
887 digests_buf[digests_offset].digest_buf[DGST_R0],
888 digests_buf[digests_offset].digest_buf[DGST_R1],
889 digests_buf[digests_offset].digest_buf[DGST_R2],
890 digests_buf[digests_offset].digest_buf[DGST_R3]
897 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos += VECT_SIZE)
899 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
901 const u32x pw_len = pw_l_len + pw_r_len;
903 const u32 salt_word_len = (salt_len + pw_len) * 2;
905 u32x wordr0[4] = { 0 };
906 u32x wordr1[4] = { 0 };
907 u32x wordr2[4] = { 0 };
908 u32x wordr3[4] = { 0 };
910 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
911 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
912 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
913 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
914 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
915 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
916 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
917 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
919 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
921 switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
926 w0[0] = wordl0[0] | wordr0[0];
927 w0[1] = wordl0[1] | wordr0[1];
928 w0[2] = wordl0[2] | wordr0[2];
929 w0[3] = wordl0[3] | wordr0[3];
933 w1[0] = wordl1[0] | wordr1[0];
934 w1[1] = wordl1[1] | wordr1[1];
935 w1[2] = wordl1[2] | wordr1[2];
936 w1[3] = wordl1[3] | wordr1[3];
940 w2[0] = wordl2[0] | wordr2[0];
941 w2[1] = wordl2[1] | wordr2[1];
942 w2[2] = wordl2[2] | wordr2[2];
943 w2[3] = wordl2[3] | wordr2[3];
947 w3[0] = wordl3[0] | wordr3[0];
948 w3[1] = wordl3[1] | wordr3[1];
949 w3[2] = wordl3[2] | wordr3[2];
950 w3[3] = wordl3[3] | wordr3[3];
978 switch_buffer_by_offset_le (w0_t, w1_t, w2_t, w3_t, salt_len);
980 w0_t[0] |= salt_buf0[0];
981 w0_t[1] |= salt_buf0[1];
982 w0_t[2] |= salt_buf0[2];
983 w0_t[3] |= salt_buf0[3];
984 w1_t[0] |= salt_buf1[0];
985 w1_t[1] |= salt_buf1[1];
986 w1_t[2] |= salt_buf1[2];
987 w1_t[3] |= salt_buf1[3];
1009 * precompute key1 since key is static: 0x0123456789abcdef
1010 * plus LEFT_ROTATE by 2
1015 Kc[ 0] = 0x64649040;
1016 Kc[ 1] = 0x14909858;
1017 Kc[ 2] = 0xc4b44888;
1018 Kc[ 3] = 0x9094e438;
1019 Kc[ 4] = 0xd8a004f0;
1020 Kc[ 5] = 0xa8f02810;
1021 Kc[ 6] = 0xc84048d8;
1022 Kc[ 7] = 0x68d804a8;
1023 Kc[ 8] = 0x0490e40c;
1024 Kc[ 9] = 0xac183024;
1025 Kc[10] = 0x24c07c10;
1026 Kc[11] = 0x8c88c038;
1027 Kc[12] = 0xc048c824;
1028 Kc[13] = 0x4c0470a8;
1029 Kc[14] = 0x584020b4;
1030 Kc[15] = 0x00742c4c;
1034 Kd[ 0] = 0xa42ce40c;
1035 Kd[ 1] = 0x64689858;
1036 Kd[ 2] = 0x484050b8;
1037 Kd[ 3] = 0xe8184814;
1038 Kd[ 4] = 0x405cc070;
1039 Kd[ 5] = 0xa010784c;
1040 Kd[ 6] = 0x6074a800;
1041 Kd[ 7] = 0x80701c1c;
1042 Kd[ 8] = 0x9cd49430;
1043 Kd[ 9] = 0x4c8ce078;
1044 Kd[10] = 0x5c18c088;
1045 Kd[11] = 0x28a8a4c8;
1046 Kd[12] = 0x3c180838;
1047 Kd[13] = 0xb0b86c20;
1048 Kd[14] = 0xac84a094;
1049 Kd[15] = 0x4ce0c0c4;
1052 * key1 (generate key)
1060 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
1064 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
1065 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
1070 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
1074 * key2 (generate hash)
1077 _des_crypt_keysetup (iv[0], iv[1], Kc, Kd, s_skb);
1082 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
1086 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
1087 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
1092 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
1102 COMPARE_S_SIMD (iv[0], iv[1], c, d);
1106 __kernel void m03100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1110 __kernel void m03100_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)