2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
11 //#define NEW_SIMD_CODE
13 #include "inc_vendor.cl"
14 #include "inc_hash_constants.h"
15 #include "inc_hash_functions.cl"
16 #include "inc_types.cl"
17 #include "inc_common.cl"
18 #include "inc_simd.cl"
20 #define PERM_OP(a,b,tt,n,m) \
30 #define HPERM_OP(a,tt,n,m) \
36 tt = tt >> (16 + n); \
42 PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \
43 PERM_OP (l, r, tt, 16, 0x0000ffff); \
44 PERM_OP (r, l, tt, 2, 0x33333333); \
45 PERM_OP (l, r, tt, 8, 0x00ff00ff); \
46 PERM_OP (r, l, tt, 1, 0x55555555); \
51 PERM_OP (l, r, tt, 1, 0x55555555); \
52 PERM_OP (r, l, tt, 8, 0x00ff00ff); \
53 PERM_OP (l, r, tt, 2, 0x33333333); \
54 PERM_OP (r, l, tt, 16, 0x0000ffff); \
55 PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
58 __constant u32 c_SPtrans[8][64] =
61 0x02080800, 0x00080000, 0x02000002, 0x02080802,
62 0x02000000, 0x00080802, 0x00080002, 0x02000002,
63 0x00080802, 0x02080800, 0x02080000, 0x00000802,
64 0x02000802, 0x02000000, 0x00000000, 0x00080002,
65 0x00080000, 0x00000002, 0x02000800, 0x00080800,
66 0x02080802, 0x02080000, 0x00000802, 0x02000800,
67 0x00000002, 0x00000800, 0x00080800, 0x02080002,
68 0x00000800, 0x02000802, 0x02080002, 0x00000000,
69 0x00000000, 0x02080802, 0x02000800, 0x00080002,
70 0x02080800, 0x00080000, 0x00000802, 0x02000800,
71 0x02080002, 0x00000800, 0x00080800, 0x02000002,
72 0x00080802, 0x00000002, 0x02000002, 0x02080000,
73 0x02080802, 0x00080800, 0x02080000, 0x02000802,
74 0x02000000, 0x00000802, 0x00080002, 0x00000000,
75 0x00080000, 0x02000000, 0x02000802, 0x02080800,
76 0x00000002, 0x02080002, 0x00000800, 0x00080802,
79 0x40108010, 0x00000000, 0x00108000, 0x40100000,
80 0x40000010, 0x00008010, 0x40008000, 0x00108000,
81 0x00008000, 0x40100010, 0x00000010, 0x40008000,
82 0x00100010, 0x40108000, 0x40100000, 0x00000010,
83 0x00100000, 0x40008010, 0x40100010, 0x00008000,
84 0x00108010, 0x40000000, 0x00000000, 0x00100010,
85 0x40008010, 0x00108010, 0x40108000, 0x40000010,
86 0x40000000, 0x00100000, 0x00008010, 0x40108010,
87 0x00100010, 0x40108000, 0x40008000, 0x00108010,
88 0x40108010, 0x00100010, 0x40000010, 0x00000000,
89 0x40000000, 0x00008010, 0x00100000, 0x40100010,
90 0x00008000, 0x40000000, 0x00108010, 0x40008010,
91 0x40108000, 0x00008000, 0x00000000, 0x40000010,
92 0x00000010, 0x40108010, 0x00108000, 0x40100000,
93 0x40100010, 0x00100000, 0x00008010, 0x40008000,
94 0x40008010, 0x00000010, 0x40100000, 0x00108000,
97 0x04000001, 0x04040100, 0x00000100, 0x04000101,
98 0x00040001, 0x04000000, 0x04000101, 0x00040100,
99 0x04000100, 0x00040000, 0x04040000, 0x00000001,
100 0x04040101, 0x00000101, 0x00000001, 0x04040001,
101 0x00000000, 0x00040001, 0x04040100, 0x00000100,
102 0x00000101, 0x04040101, 0x00040000, 0x04000001,
103 0x04040001, 0x04000100, 0x00040101, 0x04040000,
104 0x00040100, 0x00000000, 0x04000000, 0x00040101,
105 0x04040100, 0x00000100, 0x00000001, 0x00040000,
106 0x00000101, 0x00040001, 0x04040000, 0x04000101,
107 0x00000000, 0x04040100, 0x00040100, 0x04040001,
108 0x00040001, 0x04000000, 0x04040101, 0x00000001,
109 0x00040101, 0x04000001, 0x04000000, 0x04040101,
110 0x00040000, 0x04000100, 0x04000101, 0x00040100,
111 0x04000100, 0x00000000, 0x04040001, 0x00000101,
112 0x04000001, 0x00040101, 0x00000100, 0x04040000,
115 0x00401008, 0x10001000, 0x00000008, 0x10401008,
116 0x00000000, 0x10400000, 0x10001008, 0x00400008,
117 0x10401000, 0x10000008, 0x10000000, 0x00001008,
118 0x10000008, 0x00401008, 0x00400000, 0x10000000,
119 0x10400008, 0x00401000, 0x00001000, 0x00000008,
120 0x00401000, 0x10001008, 0x10400000, 0x00001000,
121 0x00001008, 0x00000000, 0x00400008, 0x10401000,
122 0x10001000, 0x10400008, 0x10401008, 0x00400000,
123 0x10400008, 0x00001008, 0x00400000, 0x10000008,
124 0x00401000, 0x10001000, 0x00000008, 0x10400000,
125 0x10001008, 0x00000000, 0x00001000, 0x00400008,
126 0x00000000, 0x10400008, 0x10401000, 0x00001000,
127 0x10000000, 0x10401008, 0x00401008, 0x00400000,
128 0x10401008, 0x00000008, 0x10001000, 0x00401008,
129 0x00400008, 0x00401000, 0x10400000, 0x10001008,
130 0x00001008, 0x10000000, 0x10000008, 0x10401000,
133 0x08000000, 0x00010000, 0x00000400, 0x08010420,
134 0x08010020, 0x08000400, 0x00010420, 0x08010000,
135 0x00010000, 0x00000020, 0x08000020, 0x00010400,
136 0x08000420, 0x08010020, 0x08010400, 0x00000000,
137 0x00010400, 0x08000000, 0x00010020, 0x00000420,
138 0x08000400, 0x00010420, 0x00000000, 0x08000020,
139 0x00000020, 0x08000420, 0x08010420, 0x00010020,
140 0x08010000, 0x00000400, 0x00000420, 0x08010400,
141 0x08010400, 0x08000420, 0x00010020, 0x08010000,
142 0x00010000, 0x00000020, 0x08000020, 0x08000400,
143 0x08000000, 0x00010400, 0x08010420, 0x00000000,
144 0x00010420, 0x08000000, 0x00000400, 0x00010020,
145 0x08000420, 0x00000400, 0x00000000, 0x08010420,
146 0x08010020, 0x08010400, 0x00000420, 0x00010000,
147 0x00010400, 0x08010020, 0x08000400, 0x00000420,
148 0x00000020, 0x00010420, 0x08010000, 0x08000020,
151 0x80000040, 0x00200040, 0x00000000, 0x80202000,
152 0x00200040, 0x00002000, 0x80002040, 0x00200000,
153 0x00002040, 0x80202040, 0x00202000, 0x80000000,
154 0x80002000, 0x80000040, 0x80200000, 0x00202040,
155 0x00200000, 0x80002040, 0x80200040, 0x00000000,
156 0x00002000, 0x00000040, 0x80202000, 0x80200040,
157 0x80202040, 0x80200000, 0x80000000, 0x00002040,
158 0x00000040, 0x00202000, 0x00202040, 0x80002000,
159 0x00002040, 0x80000000, 0x80002000, 0x00202040,
160 0x80202000, 0x00200040, 0x00000000, 0x80002000,
161 0x80000000, 0x00002000, 0x80200040, 0x00200000,
162 0x00200040, 0x80202040, 0x00202000, 0x00000040,
163 0x80202040, 0x00202000, 0x00200000, 0x80002040,
164 0x80000040, 0x80200000, 0x00202040, 0x00000000,
165 0x00002000, 0x80000040, 0x80002040, 0x80202000,
166 0x80200000, 0x00002040, 0x00000040, 0x80200040,
169 0x00004000, 0x00000200, 0x01000200, 0x01000004,
170 0x01004204, 0x00004004, 0x00004200, 0x00000000,
171 0x01000000, 0x01000204, 0x00000204, 0x01004000,
172 0x00000004, 0x01004200, 0x01004000, 0x00000204,
173 0x01000204, 0x00004000, 0x00004004, 0x01004204,
174 0x00000000, 0x01000200, 0x01000004, 0x00004200,
175 0x01004004, 0x00004204, 0x01004200, 0x00000004,
176 0x00004204, 0x01004004, 0x00000200, 0x01000000,
177 0x00004204, 0x01004000, 0x01004004, 0x00000204,
178 0x00004000, 0x00000200, 0x01000000, 0x01004004,
179 0x01000204, 0x00004204, 0x00004200, 0x00000000,
180 0x00000200, 0x01000004, 0x00000004, 0x01000200,
181 0x00000000, 0x01000204, 0x01000200, 0x00004200,
182 0x00000204, 0x00004000, 0x01004204, 0x01000000,
183 0x01004200, 0x00000004, 0x00004004, 0x01004204,
184 0x01000004, 0x01004200, 0x01004000, 0x00004004,
187 0x20800080, 0x20820000, 0x00020080, 0x00000000,
188 0x20020000, 0x00800080, 0x20800000, 0x20820080,
189 0x00000080, 0x20000000, 0x00820000, 0x00020080,
190 0x00820080, 0x20020080, 0x20000080, 0x20800000,
191 0x00020000, 0x00820080, 0x00800080, 0x20020000,
192 0x20820080, 0x20000080, 0x00000000, 0x00820000,
193 0x20000000, 0x00800000, 0x20020080, 0x20800080,
194 0x00800000, 0x00020000, 0x20820000, 0x00000080,
195 0x00800000, 0x00020000, 0x20000080, 0x20820080,
196 0x00020080, 0x20000000, 0x00000000, 0x00820000,
197 0x20800080, 0x20020080, 0x20020000, 0x00800080,
198 0x20820000, 0x00000080, 0x00800080, 0x20020000,
199 0x20820080, 0x00800000, 0x20800000, 0x20000080,
200 0x00820000, 0x00020080, 0x20020080, 0x20800000,
201 0x00000080, 0x20820000, 0x00820080, 0x00000000,
202 0x20000000, 0x20800080, 0x00020000, 0x00820080,
206 __constant u32 c_skb[8][64] =
209 0x00000000, 0x00000010, 0x20000000, 0x20000010,
210 0x00010000, 0x00010010, 0x20010000, 0x20010010,
211 0x00000800, 0x00000810, 0x20000800, 0x20000810,
212 0x00010800, 0x00010810, 0x20010800, 0x20010810,
213 0x00000020, 0x00000030, 0x20000020, 0x20000030,
214 0x00010020, 0x00010030, 0x20010020, 0x20010030,
215 0x00000820, 0x00000830, 0x20000820, 0x20000830,
216 0x00010820, 0x00010830, 0x20010820, 0x20010830,
217 0x00080000, 0x00080010, 0x20080000, 0x20080010,
218 0x00090000, 0x00090010, 0x20090000, 0x20090010,
219 0x00080800, 0x00080810, 0x20080800, 0x20080810,
220 0x00090800, 0x00090810, 0x20090800, 0x20090810,
221 0x00080020, 0x00080030, 0x20080020, 0x20080030,
222 0x00090020, 0x00090030, 0x20090020, 0x20090030,
223 0x00080820, 0x00080830, 0x20080820, 0x20080830,
224 0x00090820, 0x00090830, 0x20090820, 0x20090830,
227 0x00000000, 0x02000000, 0x00002000, 0x02002000,
228 0x00200000, 0x02200000, 0x00202000, 0x02202000,
229 0x00000004, 0x02000004, 0x00002004, 0x02002004,
230 0x00200004, 0x02200004, 0x00202004, 0x02202004,
231 0x00000400, 0x02000400, 0x00002400, 0x02002400,
232 0x00200400, 0x02200400, 0x00202400, 0x02202400,
233 0x00000404, 0x02000404, 0x00002404, 0x02002404,
234 0x00200404, 0x02200404, 0x00202404, 0x02202404,
235 0x10000000, 0x12000000, 0x10002000, 0x12002000,
236 0x10200000, 0x12200000, 0x10202000, 0x12202000,
237 0x10000004, 0x12000004, 0x10002004, 0x12002004,
238 0x10200004, 0x12200004, 0x10202004, 0x12202004,
239 0x10000400, 0x12000400, 0x10002400, 0x12002400,
240 0x10200400, 0x12200400, 0x10202400, 0x12202400,
241 0x10000404, 0x12000404, 0x10002404, 0x12002404,
242 0x10200404, 0x12200404, 0x10202404, 0x12202404,
245 0x00000000, 0x00000001, 0x00040000, 0x00040001,
246 0x01000000, 0x01000001, 0x01040000, 0x01040001,
247 0x00000002, 0x00000003, 0x00040002, 0x00040003,
248 0x01000002, 0x01000003, 0x01040002, 0x01040003,
249 0x00000200, 0x00000201, 0x00040200, 0x00040201,
250 0x01000200, 0x01000201, 0x01040200, 0x01040201,
251 0x00000202, 0x00000203, 0x00040202, 0x00040203,
252 0x01000202, 0x01000203, 0x01040202, 0x01040203,
253 0x08000000, 0x08000001, 0x08040000, 0x08040001,
254 0x09000000, 0x09000001, 0x09040000, 0x09040001,
255 0x08000002, 0x08000003, 0x08040002, 0x08040003,
256 0x09000002, 0x09000003, 0x09040002, 0x09040003,
257 0x08000200, 0x08000201, 0x08040200, 0x08040201,
258 0x09000200, 0x09000201, 0x09040200, 0x09040201,
259 0x08000202, 0x08000203, 0x08040202, 0x08040203,
260 0x09000202, 0x09000203, 0x09040202, 0x09040203,
263 0x00000000, 0x00100000, 0x00000100, 0x00100100,
264 0x00000008, 0x00100008, 0x00000108, 0x00100108,
265 0x00001000, 0x00101000, 0x00001100, 0x00101100,
266 0x00001008, 0x00101008, 0x00001108, 0x00101108,
267 0x04000000, 0x04100000, 0x04000100, 0x04100100,
268 0x04000008, 0x04100008, 0x04000108, 0x04100108,
269 0x04001000, 0x04101000, 0x04001100, 0x04101100,
270 0x04001008, 0x04101008, 0x04001108, 0x04101108,
271 0x00020000, 0x00120000, 0x00020100, 0x00120100,
272 0x00020008, 0x00120008, 0x00020108, 0x00120108,
273 0x00021000, 0x00121000, 0x00021100, 0x00121100,
274 0x00021008, 0x00121008, 0x00021108, 0x00121108,
275 0x04020000, 0x04120000, 0x04020100, 0x04120100,
276 0x04020008, 0x04120008, 0x04020108, 0x04120108,
277 0x04021000, 0x04121000, 0x04021100, 0x04121100,
278 0x04021008, 0x04121008, 0x04021108, 0x04121108,
281 0x00000000, 0x10000000, 0x00010000, 0x10010000,
282 0x00000004, 0x10000004, 0x00010004, 0x10010004,
283 0x20000000, 0x30000000, 0x20010000, 0x30010000,
284 0x20000004, 0x30000004, 0x20010004, 0x30010004,
285 0x00100000, 0x10100000, 0x00110000, 0x10110000,
286 0x00100004, 0x10100004, 0x00110004, 0x10110004,
287 0x20100000, 0x30100000, 0x20110000, 0x30110000,
288 0x20100004, 0x30100004, 0x20110004, 0x30110004,
289 0x00001000, 0x10001000, 0x00011000, 0x10011000,
290 0x00001004, 0x10001004, 0x00011004, 0x10011004,
291 0x20001000, 0x30001000, 0x20011000, 0x30011000,
292 0x20001004, 0x30001004, 0x20011004, 0x30011004,
293 0x00101000, 0x10101000, 0x00111000, 0x10111000,
294 0x00101004, 0x10101004, 0x00111004, 0x10111004,
295 0x20101000, 0x30101000, 0x20111000, 0x30111000,
296 0x20101004, 0x30101004, 0x20111004, 0x30111004,
299 0x00000000, 0x08000000, 0x00000008, 0x08000008,
300 0x00000400, 0x08000400, 0x00000408, 0x08000408,
301 0x00020000, 0x08020000, 0x00020008, 0x08020008,
302 0x00020400, 0x08020400, 0x00020408, 0x08020408,
303 0x00000001, 0x08000001, 0x00000009, 0x08000009,
304 0x00000401, 0x08000401, 0x00000409, 0x08000409,
305 0x00020001, 0x08020001, 0x00020009, 0x08020009,
306 0x00020401, 0x08020401, 0x00020409, 0x08020409,
307 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
308 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
309 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
310 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
311 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
312 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
313 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
314 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
317 0x00000000, 0x00000100, 0x00080000, 0x00080100,
318 0x01000000, 0x01000100, 0x01080000, 0x01080100,
319 0x00000010, 0x00000110, 0x00080010, 0x00080110,
320 0x01000010, 0x01000110, 0x01080010, 0x01080110,
321 0x00200000, 0x00200100, 0x00280000, 0x00280100,
322 0x01200000, 0x01200100, 0x01280000, 0x01280100,
323 0x00200010, 0x00200110, 0x00280010, 0x00280110,
324 0x01200010, 0x01200110, 0x01280010, 0x01280110,
325 0x00000200, 0x00000300, 0x00080200, 0x00080300,
326 0x01000200, 0x01000300, 0x01080200, 0x01080300,
327 0x00000210, 0x00000310, 0x00080210, 0x00080310,
328 0x01000210, 0x01000310, 0x01080210, 0x01080310,
329 0x00200200, 0x00200300, 0x00280200, 0x00280300,
330 0x01200200, 0x01200300, 0x01280200, 0x01280300,
331 0x00200210, 0x00200310, 0x00280210, 0x00280310,
332 0x01200210, 0x01200310, 0x01280210, 0x01280310,
335 0x00000000, 0x04000000, 0x00040000, 0x04040000,
336 0x00000002, 0x04000002, 0x00040002, 0x04040002,
337 0x00002000, 0x04002000, 0x00042000, 0x04042000,
338 0x00002002, 0x04002002, 0x00042002, 0x04042002,
339 0x00000020, 0x04000020, 0x00040020, 0x04040020,
340 0x00000022, 0x04000022, 0x00040022, 0x04040022,
341 0x00002020, 0x04002020, 0x00042020, 0x04042020,
342 0x00002022, 0x04002022, 0x00042022, 0x04042022,
343 0x00000800, 0x04000800, 0x00040800, 0x04040800,
344 0x00000802, 0x04000802, 0x00040802, 0x04040802,
345 0x00002800, 0x04002800, 0x00042800, 0x04042800,
346 0x00002802, 0x04002802, 0x00042802, 0x04042802,
347 0x00000820, 0x04000820, 0x00040820, 0x04040820,
348 0x00000822, 0x04000822, 0x00040822, 0x04040822,
349 0x00002820, 0x04002820, 0x00042820, 0x04042820,
350 0x00002822, 0x04002822, 0x00042822, 0x04042822
355 #define BOX(i,n,S) (S)[(n)][(i)]
357 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
359 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
361 #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])
362 #elif VECT_SIZE == 16
363 #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])
366 void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
381 for (u32 i = 0; i < 16; i += 2)
387 t = Kd[i + 0] ^ rotl32 (r, 28u);
389 l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
390 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
391 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
392 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
393 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
394 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
395 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
396 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
399 t = Kd[i + 1] ^ rotl32 (l, 28u);
401 r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
402 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
403 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
404 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
405 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
406 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
407 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
408 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
420 void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
424 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
425 HPERM_OP (c, tt, 2, 0xcccc0000);
426 HPERM_OP (d, tt, 2, 0xcccc0000);
427 PERM_OP (d, c, tt, 1, 0x55555555);
428 PERM_OP (c, d, tt, 8, 0x00ff00ff);
429 PERM_OP (d, c, tt, 1, 0x55555555);
431 d = ((d & 0x000000ff) << 16)
432 | ((d & 0x0000ff00) << 0)
433 | ((d & 0x00ff0000) >> 16)
434 | ((c & 0xf0000000) >> 4);
441 for (u32 i = 0; i < 16; i++)
443 if ((i < 2) || (i == 8) || (i == 15))
445 c = ((c >> 1) | (c << 27));
446 d = ((d >> 1) | (d << 27));
450 c = ((c >> 2) | (c << 26));
451 d = ((d >> 2) | (d << 26));
457 const u32x c00 = (c >> 0) & 0x0000003f;
458 const u32x c06 = (c >> 6) & 0x00383003;
459 const u32x c07 = (c >> 7) & 0x0000003c;
460 const u32x c13 = (c >> 13) & 0x0000060f;
461 const u32x c20 = (c >> 20) & 0x00000001;
463 u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
464 | BOX (((c06 >> 0) & 0xff)
465 |((c07 >> 0) & 0xff), 1, s_skb)
466 | BOX (((c13 >> 0) & 0xff)
467 |((c06 >> 8) & 0xff), 2, s_skb)
468 | BOX (((c20 >> 0) & 0xff)
470 |((c06 >> 16) & 0xff), 3, s_skb);
472 const u32x d00 = (d >> 0) & 0x00003c3f;
473 const u32x d07 = (d >> 7) & 0x00003f03;
474 const u32x d21 = (d >> 21) & 0x0000000f;
475 const u32x d22 = (d >> 22) & 0x00000030;
477 u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
478 | BOX (((d07 >> 0) & 0xff)
479 |((d00 >> 8) & 0xff), 5, s_skb)
480 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
481 | BOX (((d21 >> 0) & 0xff)
482 |((d22 >> 0) & 0xff), 7, s_skb);
484 Kc[i] = ((t << 16) | (s & 0x0000ffff));
485 Kd[i] = ((s >> 16) | (t & 0xffff0000));
487 Kc[i] = rotl32 (Kc[i], 2u);
488 Kd[i] = rotl32 (Kd[i], 2u);
492 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
498 const u32 gid = get_global_id (0);
499 const u32 lid = get_local_id (0);
500 const u32 lsz = get_local_size (0);
506 __local u32 s_SPtrans[8][64];
507 __local u32 s_skb[8][64];
509 for (u32 i = lid; i < 64; i += lsz)
511 s_SPtrans[0][i] = c_SPtrans[0][i];
512 s_SPtrans[1][i] = c_SPtrans[1][i];
513 s_SPtrans[2][i] = c_SPtrans[2][i];
514 s_SPtrans[3][i] = c_SPtrans[3][i];
515 s_SPtrans[4][i] = c_SPtrans[4][i];
516 s_SPtrans[5][i] = c_SPtrans[5][i];
517 s_SPtrans[6][i] = c_SPtrans[6][i];
518 s_SPtrans[7][i] = c_SPtrans[7][i];
520 s_skb[0][i] = c_skb[0][i];
521 s_skb[1][i] = c_skb[1][i];
522 s_skb[2][i] = c_skb[2][i];
523 s_skb[3][i] = c_skb[3][i];
524 s_skb[4][i] = c_skb[4][i];
525 s_skb[5][i] = c_skb[5][i];
526 s_skb[6][i] = c_skb[6][i];
527 s_skb[7][i] = c_skb[7][i];
530 barrier (CLK_LOCAL_MEM_FENCE);
532 if (gid >= gid_max) return;
541 pw_buf0[0] = pws[gid].i[0];
542 pw_buf0[1] = pws[gid].i[1];
543 pw_buf0[2] = pws[gid].i[2];
544 pw_buf0[3] = pws[gid].i[3];
545 pw_buf1[0] = pws[gid].i[4];
546 pw_buf1[1] = pws[gid].i[5];
547 pw_buf1[2] = pws[gid].i[6];
548 pw_buf1[3] = pws[gid].i[7];
550 const u32 pw_l_len = pws[gid].pw_len;
559 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
560 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
561 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
562 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
563 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
564 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
565 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
566 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
568 const u32 salt_len = salt_bufs[salt_pos].salt_len;
574 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
576 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
578 const u32x pw_len = pw_l_len + pw_r_len;
580 const u32x salt_word_len = (salt_len + pw_len) * 2;
583 * concat password candidate
586 u32x wordl0[4] = { 0 };
587 u32x wordl1[4] = { 0 };
588 u32x wordl2[4] = { 0 };
589 u32x wordl3[4] = { 0 };
591 wordl0[0] = pw_buf0[0];
592 wordl0[1] = pw_buf0[1];
593 wordl0[2] = pw_buf0[2];
594 wordl0[3] = pw_buf0[3];
595 wordl1[0] = pw_buf1[0];
596 wordl1[1] = pw_buf1[1];
597 wordl1[2] = pw_buf1[2];
598 wordl1[3] = pw_buf1[3];
600 u32x wordr0[4] = { 0 };
601 u32x wordr1[4] = { 0 };
602 u32x wordr2[4] = { 0 };
603 u32x wordr3[4] = { 0 };
605 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
606 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
607 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
608 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
609 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
610 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
611 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
612 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
614 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
616 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
620 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
628 w0[0] = wordl0[0] | wordr0[0];
629 w0[1] = wordl0[1] | wordr0[1];
630 w0[2] = wordl0[2] | wordr0[2];
631 w0[3] = wordl0[3] | wordr0[3];
632 w1[0] = wordl1[0] | wordr1[0];
633 w1[1] = wordl1[1] | wordr1[1];
634 w1[2] = wordl1[2] | wordr1[2];
635 w1[3] = wordl1[3] | wordr1[3];
636 w2[0] = wordl2[0] | wordr2[0];
637 w2[1] = wordl2[1] | wordr2[1];
638 w2[2] = wordl2[2] | wordr2[2];
639 w2[3] = wordl2[3] | wordr2[3];
640 w3[0] = wordl3[0] | wordr3[0];
641 w3[1] = wordl3[1] | wordr3[1];
642 w3[2] = wordl3[2] | wordr3[2];
643 w3[3] = wordl3[3] | wordr3[3];
649 switch_buffer_by_offset_le (w0, w1, w2, w3, salt_len);
653 dst[ 0] = w0[0] | salt_buf0[0];
654 dst[ 1] = w0[1] | salt_buf0[1];
655 dst[ 2] = w0[2] | salt_buf0[2];
656 dst[ 3] = w0[3] | salt_buf0[3];
657 dst[ 4] = w1[0] | salt_buf1[0];
658 dst[ 5] = w1[1] | salt_buf1[1];
659 dst[ 6] = w1[2] | salt_buf1[2];
660 dst[ 7] = w1[3] | salt_buf1[3];
671 * precompute key1 since key is static: 0x0123456789abcdef
672 * plus LEFT_ROTATE by 2
714 * key1 (generate key)
722 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
726 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
727 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
732 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
736 * key2 (generate hash)
739 _des_crypt_keysetup (iv[0], iv[1], Kc, Kd, s_skb);
744 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
748 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
749 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
754 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
763 COMPARE_M_SIMD (iv[0], iv[1], z, z);
767 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
771 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
775 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
781 const u32 gid = get_global_id (0);
782 const u32 lid = get_local_id (0);
783 const u32 lsz = get_local_size (0);
789 __local u32 s_SPtrans[8][64];
790 __local u32 s_skb[8][64];
792 for (u32 i = lid; i < 64; i += lsz)
794 s_SPtrans[0][i] = c_SPtrans[0][i];
795 s_SPtrans[1][i] = c_SPtrans[1][i];
796 s_SPtrans[2][i] = c_SPtrans[2][i];
797 s_SPtrans[3][i] = c_SPtrans[3][i];
798 s_SPtrans[4][i] = c_SPtrans[4][i];
799 s_SPtrans[5][i] = c_SPtrans[5][i];
800 s_SPtrans[6][i] = c_SPtrans[6][i];
801 s_SPtrans[7][i] = c_SPtrans[7][i];
803 s_skb[0][i] = c_skb[0][i];
804 s_skb[1][i] = c_skb[1][i];
805 s_skb[2][i] = c_skb[2][i];
806 s_skb[3][i] = c_skb[3][i];
807 s_skb[4][i] = c_skb[4][i];
808 s_skb[5][i] = c_skb[5][i];
809 s_skb[6][i] = c_skb[6][i];
810 s_skb[7][i] = c_skb[7][i];
813 barrier (CLK_LOCAL_MEM_FENCE);
815 if (gid >= gid_max) return;
824 pw_buf0[0] = pws[gid].i[0];
825 pw_buf0[1] = pws[gid].i[1];
826 pw_buf0[2] = pws[gid].i[2];
827 pw_buf0[3] = pws[gid].i[3];
828 pw_buf1[0] = pws[gid].i[4];
829 pw_buf1[1] = pws[gid].i[5];
830 pw_buf1[2] = pws[gid].i[6];
831 pw_buf1[3] = pws[gid].i[7];
833 const u32 pw_l_len = pws[gid].pw_len;
842 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
843 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
844 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
845 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
846 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
847 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
848 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
849 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
851 const u32 salt_len = salt_bufs[salt_pos].salt_len;
857 const u32 search[4] =
859 digests_buf[digests_offset].digest_buf[DGST_R0],
860 digests_buf[digests_offset].digest_buf[DGST_R1],
869 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
871 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
873 const u32x pw_len = pw_l_len + pw_r_len;
875 const u32x salt_word_len = (salt_len + pw_len) * 2;
878 * concat password candidate
881 u32x wordl0[4] = { 0 };
882 u32x wordl1[4] = { 0 };
883 u32x wordl2[4] = { 0 };
884 u32x wordl3[4] = { 0 };
886 wordl0[0] = pw_buf0[0];
887 wordl0[1] = pw_buf0[1];
888 wordl0[2] = pw_buf0[2];
889 wordl0[3] = pw_buf0[3];
890 wordl1[0] = pw_buf1[0];
891 wordl1[1] = pw_buf1[1];
892 wordl1[2] = pw_buf1[2];
893 wordl1[3] = pw_buf1[3];
895 u32x wordr0[4] = { 0 };
896 u32x wordr1[4] = { 0 };
897 u32x wordr2[4] = { 0 };
898 u32x wordr3[4] = { 0 };
900 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
901 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
902 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
903 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
904 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
905 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
906 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
907 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
909 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
911 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
915 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
923 w0[0] = wordl0[0] | wordr0[0];
924 w0[1] = wordl0[1] | wordr0[1];
925 w0[2] = wordl0[2] | wordr0[2];
926 w0[3] = wordl0[3] | wordr0[3];
927 w1[0] = wordl1[0] | wordr1[0];
928 w1[1] = wordl1[1] | wordr1[1];
929 w1[2] = wordl1[2] | wordr1[2];
930 w1[3] = wordl1[3] | wordr1[3];
931 w2[0] = wordl2[0] | wordr2[0];
932 w2[1] = wordl2[1] | wordr2[1];
933 w2[2] = wordl2[2] | wordr2[2];
934 w2[3] = wordl2[3] | wordr2[3];
935 w3[0] = wordl3[0] | wordr3[0];
936 w3[1] = wordl3[1] | wordr3[1];
937 w3[2] = wordl3[2] | wordr3[2];
938 w3[3] = wordl3[3] | wordr3[3];
944 switch_buffer_by_offset_le (w0, w1, w2, w3, salt_len);
948 dst[ 0] = w0[0] | salt_buf0[0];
949 dst[ 1] = w0[1] | salt_buf0[1];
950 dst[ 2] = w0[2] | salt_buf0[2];
951 dst[ 3] = w0[3] | salt_buf0[3];
952 dst[ 4] = w1[0] | salt_buf1[0];
953 dst[ 5] = w1[1] | salt_buf1[1];
954 dst[ 6] = w1[2] | salt_buf1[2];
955 dst[ 7] = w1[3] | salt_buf1[3];
966 * precompute key1 since key is static: 0x0123456789abcdef
967 * plus LEFT_ROTATE by 2
1000 Kd[ 9] = 0x4c8ce078;
1001 Kd[10] = 0x5c18c088;
1002 Kd[11] = 0x28a8a4c8;
1003 Kd[12] = 0x3c180838;
1004 Kd[13] = 0xb0b86c20;
1005 Kd[14] = 0xac84a094;
1006 Kd[15] = 0x4ce0c0c4;
1009 * key1 (generate key)
1017 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
1021 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
1022 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
1027 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
1031 * key2 (generate hash)
1034 _des_crypt_keysetup (iv[0], iv[1], Kc, Kd, s_skb);
1039 for (u32 j = 0, k = 0; j < salt_word_len; j += 8, k++)
1043 data[0] = ((dst[k] << 16) & 0xff000000) | ((dst[k] << 8) & 0x0000ff00);
1044 data[1] = ((dst[k] >> 0) & 0xff000000) | ((dst[k] >> 8) & 0x0000ff00);
1049 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
1058 COMPARE_S_SIMD (iv[0], iv[1], z, z);
1062 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
1066 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)