2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
12 #include "inc_vendor.cl"
13 #include "inc_hash_constants.h"
14 #include "inc_hash_functions.cl"
15 #include "inc_types.cl"
16 #include "inc_common.cl"
17 #include "inc_simd.cl"
19 #define PERM_OP(a,b,tt,n,m) \
29 #define HPERM_OP(a,tt,n,m) \
35 tt = tt >> (16 + n); \
41 PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \
42 PERM_OP (l, r, tt, 16, 0x0000ffff); \
43 PERM_OP (r, l, tt, 2, 0x33333333); \
44 PERM_OP (l, r, tt, 8, 0x00ff00ff); \
45 PERM_OP (r, l, tt, 1, 0x55555555); \
50 PERM_OP (l, r, tt, 1, 0x55555555); \
51 PERM_OP (r, l, tt, 8, 0x00ff00ff); \
52 PERM_OP (l, r, tt, 2, 0x33333333); \
53 PERM_OP (r, l, tt, 16, 0x0000ffff); \
54 PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
57 __constant u8 ascii_to_ebcdic_pc[256] =
59 // little hack, can't crack 0-bytes in password, but who cares
60 // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
61 0x2a, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
62 0x8a, 0x89, 0x8f, 0x8c, 0xd3, 0xd0, 0xce, 0xe6, 0x9b, 0x98, 0xd5, 0xe5, 0x92, 0x91, 0x97, 0x94,
63 0x2a, 0x34, 0x54, 0x5d, 0x1c, 0x73, 0x0b, 0x51, 0x31, 0x10, 0x13, 0x37, 0x7c, 0x6b, 0x3d, 0x68,
64 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40, 0x46, 0x45, 0x5b, 0x58, 0x5e, 0x16, 0x32, 0x57, 0x76, 0x75,
65 0x52, 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07,
66 0x04, 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x3e, 0x6b, 0x1f, 0x15, 0x70,
67 0x58, 0xa8, 0xae, 0xad, 0xa2, 0xa1, 0xa7, 0xa4, 0xba, 0xb9, 0x89, 0x8f, 0x8c, 0x83, 0x80, 0x86,
68 0x85, 0x9b, 0x98, 0xef, 0xec, 0xe3, 0xe0, 0xe6, 0xe5, 0xfb, 0xf8, 0x2a, 0x7f, 0x0b, 0xe9, 0xa4,
69 0xea, 0xe9, 0xef, 0xec, 0xe3, 0x80, 0xa7, 0x85, 0xfb, 0xf8, 0xfe, 0xfd, 0xf2, 0xb9, 0xbf, 0x9d,
70 0xcb, 0xc8, 0x9e, 0xcd, 0xc2, 0xc1, 0xc7, 0xba, 0xda, 0xd9, 0xdf, 0xdc, 0xa2, 0x83, 0xd6, 0x68,
71 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07, 0x04,
72 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40,
73 0x46, 0x45, 0x5b, 0xab, 0xbf, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, 0x8a, 0x9e, 0x9d, 0x92, 0x91, 0x97,
74 0x94, 0xea, 0xfe, 0xfd, 0xf2, 0xf1, 0xf7, 0xf4, 0xcb, 0xc8, 0xce, 0xcd, 0xc2, 0xc1, 0xc7, 0xc4,
75 0xda, 0xd9, 0xdf, 0xdc, 0xd3, 0xd0, 0xd6, 0xd5, 0x3e, 0x3d, 0x32, 0x31, 0x37, 0x34, 0x1f, 0x1c,
76 0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
79 __constant u32 c_SPtrans[8][64] =
82 0x02080800, 0x00080000, 0x02000002, 0x02080802,
83 0x02000000, 0x00080802, 0x00080002, 0x02000002,
84 0x00080802, 0x02080800, 0x02080000, 0x00000802,
85 0x02000802, 0x02000000, 0x00000000, 0x00080002,
86 0x00080000, 0x00000002, 0x02000800, 0x00080800,
87 0x02080802, 0x02080000, 0x00000802, 0x02000800,
88 0x00000002, 0x00000800, 0x00080800, 0x02080002,
89 0x00000800, 0x02000802, 0x02080002, 0x00000000,
90 0x00000000, 0x02080802, 0x02000800, 0x00080002,
91 0x02080800, 0x00080000, 0x00000802, 0x02000800,
92 0x02080002, 0x00000800, 0x00080800, 0x02000002,
93 0x00080802, 0x00000002, 0x02000002, 0x02080000,
94 0x02080802, 0x00080800, 0x02080000, 0x02000802,
95 0x02000000, 0x00000802, 0x00080002, 0x00000000,
96 0x00080000, 0x02000000, 0x02000802, 0x02080800,
97 0x00000002, 0x02080002, 0x00000800, 0x00080802,
100 0x40108010, 0x00000000, 0x00108000, 0x40100000,
101 0x40000010, 0x00008010, 0x40008000, 0x00108000,
102 0x00008000, 0x40100010, 0x00000010, 0x40008000,
103 0x00100010, 0x40108000, 0x40100000, 0x00000010,
104 0x00100000, 0x40008010, 0x40100010, 0x00008000,
105 0x00108010, 0x40000000, 0x00000000, 0x00100010,
106 0x40008010, 0x00108010, 0x40108000, 0x40000010,
107 0x40000000, 0x00100000, 0x00008010, 0x40108010,
108 0x00100010, 0x40108000, 0x40008000, 0x00108010,
109 0x40108010, 0x00100010, 0x40000010, 0x00000000,
110 0x40000000, 0x00008010, 0x00100000, 0x40100010,
111 0x00008000, 0x40000000, 0x00108010, 0x40008010,
112 0x40108000, 0x00008000, 0x00000000, 0x40000010,
113 0x00000010, 0x40108010, 0x00108000, 0x40100000,
114 0x40100010, 0x00100000, 0x00008010, 0x40008000,
115 0x40008010, 0x00000010, 0x40100000, 0x00108000,
118 0x04000001, 0x04040100, 0x00000100, 0x04000101,
119 0x00040001, 0x04000000, 0x04000101, 0x00040100,
120 0x04000100, 0x00040000, 0x04040000, 0x00000001,
121 0x04040101, 0x00000101, 0x00000001, 0x04040001,
122 0x00000000, 0x00040001, 0x04040100, 0x00000100,
123 0x00000101, 0x04040101, 0x00040000, 0x04000001,
124 0x04040001, 0x04000100, 0x00040101, 0x04040000,
125 0x00040100, 0x00000000, 0x04000000, 0x00040101,
126 0x04040100, 0x00000100, 0x00000001, 0x00040000,
127 0x00000101, 0x00040001, 0x04040000, 0x04000101,
128 0x00000000, 0x04040100, 0x00040100, 0x04040001,
129 0x00040001, 0x04000000, 0x04040101, 0x00000001,
130 0x00040101, 0x04000001, 0x04000000, 0x04040101,
131 0x00040000, 0x04000100, 0x04000101, 0x00040100,
132 0x04000100, 0x00000000, 0x04040001, 0x00000101,
133 0x04000001, 0x00040101, 0x00000100, 0x04040000,
136 0x00401008, 0x10001000, 0x00000008, 0x10401008,
137 0x00000000, 0x10400000, 0x10001008, 0x00400008,
138 0x10401000, 0x10000008, 0x10000000, 0x00001008,
139 0x10000008, 0x00401008, 0x00400000, 0x10000000,
140 0x10400008, 0x00401000, 0x00001000, 0x00000008,
141 0x00401000, 0x10001008, 0x10400000, 0x00001000,
142 0x00001008, 0x00000000, 0x00400008, 0x10401000,
143 0x10001000, 0x10400008, 0x10401008, 0x00400000,
144 0x10400008, 0x00001008, 0x00400000, 0x10000008,
145 0x00401000, 0x10001000, 0x00000008, 0x10400000,
146 0x10001008, 0x00000000, 0x00001000, 0x00400008,
147 0x00000000, 0x10400008, 0x10401000, 0x00001000,
148 0x10000000, 0x10401008, 0x00401008, 0x00400000,
149 0x10401008, 0x00000008, 0x10001000, 0x00401008,
150 0x00400008, 0x00401000, 0x10400000, 0x10001008,
151 0x00001008, 0x10000000, 0x10000008, 0x10401000,
154 0x08000000, 0x00010000, 0x00000400, 0x08010420,
155 0x08010020, 0x08000400, 0x00010420, 0x08010000,
156 0x00010000, 0x00000020, 0x08000020, 0x00010400,
157 0x08000420, 0x08010020, 0x08010400, 0x00000000,
158 0x00010400, 0x08000000, 0x00010020, 0x00000420,
159 0x08000400, 0x00010420, 0x00000000, 0x08000020,
160 0x00000020, 0x08000420, 0x08010420, 0x00010020,
161 0x08010000, 0x00000400, 0x00000420, 0x08010400,
162 0x08010400, 0x08000420, 0x00010020, 0x08010000,
163 0x00010000, 0x00000020, 0x08000020, 0x08000400,
164 0x08000000, 0x00010400, 0x08010420, 0x00000000,
165 0x00010420, 0x08000000, 0x00000400, 0x00010020,
166 0x08000420, 0x00000400, 0x00000000, 0x08010420,
167 0x08010020, 0x08010400, 0x00000420, 0x00010000,
168 0x00010400, 0x08010020, 0x08000400, 0x00000420,
169 0x00000020, 0x00010420, 0x08010000, 0x08000020,
172 0x80000040, 0x00200040, 0x00000000, 0x80202000,
173 0x00200040, 0x00002000, 0x80002040, 0x00200000,
174 0x00002040, 0x80202040, 0x00202000, 0x80000000,
175 0x80002000, 0x80000040, 0x80200000, 0x00202040,
176 0x00200000, 0x80002040, 0x80200040, 0x00000000,
177 0x00002000, 0x00000040, 0x80202000, 0x80200040,
178 0x80202040, 0x80200000, 0x80000000, 0x00002040,
179 0x00000040, 0x00202000, 0x00202040, 0x80002000,
180 0x00002040, 0x80000000, 0x80002000, 0x00202040,
181 0x80202000, 0x00200040, 0x00000000, 0x80002000,
182 0x80000000, 0x00002000, 0x80200040, 0x00200000,
183 0x00200040, 0x80202040, 0x00202000, 0x00000040,
184 0x80202040, 0x00202000, 0x00200000, 0x80002040,
185 0x80000040, 0x80200000, 0x00202040, 0x00000000,
186 0x00002000, 0x80000040, 0x80002040, 0x80202000,
187 0x80200000, 0x00002040, 0x00000040, 0x80200040,
190 0x00004000, 0x00000200, 0x01000200, 0x01000004,
191 0x01004204, 0x00004004, 0x00004200, 0x00000000,
192 0x01000000, 0x01000204, 0x00000204, 0x01004000,
193 0x00000004, 0x01004200, 0x01004000, 0x00000204,
194 0x01000204, 0x00004000, 0x00004004, 0x01004204,
195 0x00000000, 0x01000200, 0x01000004, 0x00004200,
196 0x01004004, 0x00004204, 0x01004200, 0x00000004,
197 0x00004204, 0x01004004, 0x00000200, 0x01000000,
198 0x00004204, 0x01004000, 0x01004004, 0x00000204,
199 0x00004000, 0x00000200, 0x01000000, 0x01004004,
200 0x01000204, 0x00004204, 0x00004200, 0x00000000,
201 0x00000200, 0x01000004, 0x00000004, 0x01000200,
202 0x00000000, 0x01000204, 0x01000200, 0x00004200,
203 0x00000204, 0x00004000, 0x01004204, 0x01000000,
204 0x01004200, 0x00000004, 0x00004004, 0x01004204,
205 0x01000004, 0x01004200, 0x01004000, 0x00004004,
208 0x20800080, 0x20820000, 0x00020080, 0x00000000,
209 0x20020000, 0x00800080, 0x20800000, 0x20820080,
210 0x00000080, 0x20000000, 0x00820000, 0x00020080,
211 0x00820080, 0x20020080, 0x20000080, 0x20800000,
212 0x00020000, 0x00820080, 0x00800080, 0x20020000,
213 0x20820080, 0x20000080, 0x00000000, 0x00820000,
214 0x20000000, 0x00800000, 0x20020080, 0x20800080,
215 0x00800000, 0x00020000, 0x20820000, 0x00000080,
216 0x00800000, 0x00020000, 0x20000080, 0x20820080,
217 0x00020080, 0x20000000, 0x00000000, 0x00820000,
218 0x20800080, 0x20020080, 0x20020000, 0x00800080,
219 0x20820000, 0x00000080, 0x00800080, 0x20020000,
220 0x20820080, 0x00800000, 0x20800000, 0x20000080,
221 0x00820000, 0x00020080, 0x20020080, 0x20800000,
222 0x00000080, 0x20820000, 0x00820080, 0x00000000,
223 0x20000000, 0x20800080, 0x00020000, 0x00820080,
227 __constant u32 c_skb[8][64] =
230 0x00000000, 0x00000010, 0x20000000, 0x20000010,
231 0x00010000, 0x00010010, 0x20010000, 0x20010010,
232 0x00000800, 0x00000810, 0x20000800, 0x20000810,
233 0x00010800, 0x00010810, 0x20010800, 0x20010810,
234 0x00000020, 0x00000030, 0x20000020, 0x20000030,
235 0x00010020, 0x00010030, 0x20010020, 0x20010030,
236 0x00000820, 0x00000830, 0x20000820, 0x20000830,
237 0x00010820, 0x00010830, 0x20010820, 0x20010830,
238 0x00080000, 0x00080010, 0x20080000, 0x20080010,
239 0x00090000, 0x00090010, 0x20090000, 0x20090010,
240 0x00080800, 0x00080810, 0x20080800, 0x20080810,
241 0x00090800, 0x00090810, 0x20090800, 0x20090810,
242 0x00080020, 0x00080030, 0x20080020, 0x20080030,
243 0x00090020, 0x00090030, 0x20090020, 0x20090030,
244 0x00080820, 0x00080830, 0x20080820, 0x20080830,
245 0x00090820, 0x00090830, 0x20090820, 0x20090830,
248 0x00000000, 0x02000000, 0x00002000, 0x02002000,
249 0x00200000, 0x02200000, 0x00202000, 0x02202000,
250 0x00000004, 0x02000004, 0x00002004, 0x02002004,
251 0x00200004, 0x02200004, 0x00202004, 0x02202004,
252 0x00000400, 0x02000400, 0x00002400, 0x02002400,
253 0x00200400, 0x02200400, 0x00202400, 0x02202400,
254 0x00000404, 0x02000404, 0x00002404, 0x02002404,
255 0x00200404, 0x02200404, 0x00202404, 0x02202404,
256 0x10000000, 0x12000000, 0x10002000, 0x12002000,
257 0x10200000, 0x12200000, 0x10202000, 0x12202000,
258 0x10000004, 0x12000004, 0x10002004, 0x12002004,
259 0x10200004, 0x12200004, 0x10202004, 0x12202004,
260 0x10000400, 0x12000400, 0x10002400, 0x12002400,
261 0x10200400, 0x12200400, 0x10202400, 0x12202400,
262 0x10000404, 0x12000404, 0x10002404, 0x12002404,
263 0x10200404, 0x12200404, 0x10202404, 0x12202404,
266 0x00000000, 0x00000001, 0x00040000, 0x00040001,
267 0x01000000, 0x01000001, 0x01040000, 0x01040001,
268 0x00000002, 0x00000003, 0x00040002, 0x00040003,
269 0x01000002, 0x01000003, 0x01040002, 0x01040003,
270 0x00000200, 0x00000201, 0x00040200, 0x00040201,
271 0x01000200, 0x01000201, 0x01040200, 0x01040201,
272 0x00000202, 0x00000203, 0x00040202, 0x00040203,
273 0x01000202, 0x01000203, 0x01040202, 0x01040203,
274 0x08000000, 0x08000001, 0x08040000, 0x08040001,
275 0x09000000, 0x09000001, 0x09040000, 0x09040001,
276 0x08000002, 0x08000003, 0x08040002, 0x08040003,
277 0x09000002, 0x09000003, 0x09040002, 0x09040003,
278 0x08000200, 0x08000201, 0x08040200, 0x08040201,
279 0x09000200, 0x09000201, 0x09040200, 0x09040201,
280 0x08000202, 0x08000203, 0x08040202, 0x08040203,
281 0x09000202, 0x09000203, 0x09040202, 0x09040203,
284 0x00000000, 0x00100000, 0x00000100, 0x00100100,
285 0x00000008, 0x00100008, 0x00000108, 0x00100108,
286 0x00001000, 0x00101000, 0x00001100, 0x00101100,
287 0x00001008, 0x00101008, 0x00001108, 0x00101108,
288 0x04000000, 0x04100000, 0x04000100, 0x04100100,
289 0x04000008, 0x04100008, 0x04000108, 0x04100108,
290 0x04001000, 0x04101000, 0x04001100, 0x04101100,
291 0x04001008, 0x04101008, 0x04001108, 0x04101108,
292 0x00020000, 0x00120000, 0x00020100, 0x00120100,
293 0x00020008, 0x00120008, 0x00020108, 0x00120108,
294 0x00021000, 0x00121000, 0x00021100, 0x00121100,
295 0x00021008, 0x00121008, 0x00021108, 0x00121108,
296 0x04020000, 0x04120000, 0x04020100, 0x04120100,
297 0x04020008, 0x04120008, 0x04020108, 0x04120108,
298 0x04021000, 0x04121000, 0x04021100, 0x04121100,
299 0x04021008, 0x04121008, 0x04021108, 0x04121108,
302 0x00000000, 0x10000000, 0x00010000, 0x10010000,
303 0x00000004, 0x10000004, 0x00010004, 0x10010004,
304 0x20000000, 0x30000000, 0x20010000, 0x30010000,
305 0x20000004, 0x30000004, 0x20010004, 0x30010004,
306 0x00100000, 0x10100000, 0x00110000, 0x10110000,
307 0x00100004, 0x10100004, 0x00110004, 0x10110004,
308 0x20100000, 0x30100000, 0x20110000, 0x30110000,
309 0x20100004, 0x30100004, 0x20110004, 0x30110004,
310 0x00001000, 0x10001000, 0x00011000, 0x10011000,
311 0x00001004, 0x10001004, 0x00011004, 0x10011004,
312 0x20001000, 0x30001000, 0x20011000, 0x30011000,
313 0x20001004, 0x30001004, 0x20011004, 0x30011004,
314 0x00101000, 0x10101000, 0x00111000, 0x10111000,
315 0x00101004, 0x10101004, 0x00111004, 0x10111004,
316 0x20101000, 0x30101000, 0x20111000, 0x30111000,
317 0x20101004, 0x30101004, 0x20111004, 0x30111004,
320 0x00000000, 0x08000000, 0x00000008, 0x08000008,
321 0x00000400, 0x08000400, 0x00000408, 0x08000408,
322 0x00020000, 0x08020000, 0x00020008, 0x08020008,
323 0x00020400, 0x08020400, 0x00020408, 0x08020408,
324 0x00000001, 0x08000001, 0x00000009, 0x08000009,
325 0x00000401, 0x08000401, 0x00000409, 0x08000409,
326 0x00020001, 0x08020001, 0x00020009, 0x08020009,
327 0x00020401, 0x08020401, 0x00020409, 0x08020409,
328 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
329 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
330 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
331 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
332 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
333 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
334 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
335 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
338 0x00000000, 0x00000100, 0x00080000, 0x00080100,
339 0x01000000, 0x01000100, 0x01080000, 0x01080100,
340 0x00000010, 0x00000110, 0x00080010, 0x00080110,
341 0x01000010, 0x01000110, 0x01080010, 0x01080110,
342 0x00200000, 0x00200100, 0x00280000, 0x00280100,
343 0x01200000, 0x01200100, 0x01280000, 0x01280100,
344 0x00200010, 0x00200110, 0x00280010, 0x00280110,
345 0x01200010, 0x01200110, 0x01280010, 0x01280110,
346 0x00000200, 0x00000300, 0x00080200, 0x00080300,
347 0x01000200, 0x01000300, 0x01080200, 0x01080300,
348 0x00000210, 0x00000310, 0x00080210, 0x00080310,
349 0x01000210, 0x01000310, 0x01080210, 0x01080310,
350 0x00200200, 0x00200300, 0x00280200, 0x00280300,
351 0x01200200, 0x01200300, 0x01280200, 0x01280300,
352 0x00200210, 0x00200310, 0x00280210, 0x00280310,
353 0x01200210, 0x01200310, 0x01280210, 0x01280310,
356 0x00000000, 0x04000000, 0x00040000, 0x04040000,
357 0x00000002, 0x04000002, 0x00040002, 0x04040002,
358 0x00002000, 0x04002000, 0x00042000, 0x04042000,
359 0x00002002, 0x04002002, 0x00042002, 0x04042002,
360 0x00000020, 0x04000020, 0x00040020, 0x04040020,
361 0x00000022, 0x04000022, 0x00040022, 0x04040022,
362 0x00002020, 0x04002020, 0x00042020, 0x04042020,
363 0x00002022, 0x04002022, 0x00042022, 0x04042022,
364 0x00000800, 0x04000800, 0x00040800, 0x04040800,
365 0x00000802, 0x04000802, 0x00040802, 0x04040802,
366 0x00002800, 0x04002800, 0x00042800, 0x04042800,
367 0x00002802, 0x04002802, 0x00042802, 0x04042802,
368 0x00000820, 0x04000820, 0x00040820, 0x04040820,
369 0x00000822, 0x04000822, 0x00040822, 0x04040822,
370 0x00002820, 0x04002820, 0x00042820, 0x04042820,
371 0x00002822, 0x04002822, 0x00042822, 0x04042822
376 #define BOX(i,n,S) (S)[(n)][(i)]
378 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
380 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
382 #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])
383 #elif VECT_SIZE == 16
384 #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])
388 #define BOX1(i,S) (S)[(i)]
390 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
392 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
394 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7])
395 #elif VECT_SIZE == 16
396 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf])
399 void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
407 for (u32 i = 0; i < 16; i += 2)
413 t = Kd[i + 0] ^ rotl32 (r, 28u);
415 l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
416 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
417 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
418 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
419 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
420 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
421 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
422 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
425 t = Kd[i + 1] ^ rotl32 (l, 28u);
427 r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
428 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
429 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
430 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
431 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
432 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
433 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
434 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
441 void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
445 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
446 HPERM_OP (c, tt, 2, 0xcccc0000);
447 HPERM_OP (d, tt, 2, 0xcccc0000);
448 PERM_OP (d, c, tt, 1, 0x55555555);
449 PERM_OP (c, d, tt, 8, 0x00ff00ff);
450 PERM_OP (d, c, tt, 1, 0x55555555);
452 d = ((d & 0x000000ff) << 16)
453 | ((d & 0x0000ff00) << 0)
454 | ((d & 0x00ff0000) >> 16)
455 | ((c & 0xf0000000) >> 4);
462 for (u32 i = 0; i < 16; i++)
464 if ((i < 2) || (i == 8) || (i == 15))
466 c = ((c >> 1) | (c << 27));
467 d = ((d >> 1) | (d << 27));
471 c = ((c >> 2) | (c << 26));
472 d = ((d >> 2) | (d << 26));
478 const u32x c00 = (c >> 0) & 0x0000003f;
479 const u32x c06 = (c >> 6) & 0x00383003;
480 const u32x c07 = (c >> 7) & 0x0000003c;
481 const u32x c13 = (c >> 13) & 0x0000060f;
482 const u32x c20 = (c >> 20) & 0x00000001;
484 u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
485 | BOX (((c06 >> 0) & 0xff)
486 |((c07 >> 0) & 0xff), 1, s_skb)
487 | BOX (((c13 >> 0) & 0xff)
488 |((c06 >> 8) & 0xff), 2, s_skb)
489 | BOX (((c20 >> 0) & 0xff)
491 |((c06 >> 16) & 0xff), 3, s_skb);
493 const u32x d00 = (d >> 0) & 0x00003c3f;
494 const u32x d07 = (d >> 7) & 0x00003f03;
495 const u32x d21 = (d >> 21) & 0x0000000f;
496 const u32x d22 = (d >> 22) & 0x00000030;
498 u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
499 | BOX (((d07 >> 0) & 0xff)
500 |((d00 >> 8) & 0xff), 5, s_skb)
501 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
502 | BOX (((d21 >> 0) & 0xff)
503 |((d22 >> 0) & 0xff), 7, s_skb);
505 Kc[i] = ((t << 16) | (s & 0x0000ffff));
506 Kd[i] = ((s >> 16) | (t & 0xffff0000));
508 Kc[i] = rotl32 (Kc[i], 2u);
509 Kd[i] = rotl32 (Kd[i], 2u);
513 void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
515 key[0] = BOX1 (((w0 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0
516 | BOX1 (((w0 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8
517 | BOX1 (((w0 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16
518 | BOX1 (((w0 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24;
520 key[1] = BOX1 (((w1 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0
521 | BOX1 (((w1 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8
522 | BOX1 (((w1 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16
523 | BOX1 (((w1 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24;
526 __kernel void m08500_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)
532 const u32 gid = get_global_id (0);
533 const u32 lid = get_local_id (0);
534 const u32 lsz = get_local_size (0);
540 __local u32 s_SPtrans[8][64];
541 __local u32 s_skb[8][64];
543 for (u32 i = lid; i < 64; i += lsz)
545 s_SPtrans[0][i] = c_SPtrans[0][i];
546 s_SPtrans[1][i] = c_SPtrans[1][i];
547 s_SPtrans[2][i] = c_SPtrans[2][i];
548 s_SPtrans[3][i] = c_SPtrans[3][i];
549 s_SPtrans[4][i] = c_SPtrans[4][i];
550 s_SPtrans[5][i] = c_SPtrans[5][i];
551 s_SPtrans[6][i] = c_SPtrans[6][i];
552 s_SPtrans[7][i] = c_SPtrans[7][i];
554 s_skb[0][i] = c_skb[0][i];
555 s_skb[1][i] = c_skb[1][i];
556 s_skb[2][i] = c_skb[2][i];
557 s_skb[3][i] = c_skb[3][i];
558 s_skb[4][i] = c_skb[4][i];
559 s_skb[5][i] = c_skb[5][i];
560 s_skb[6][i] = c_skb[6][i];
561 s_skb[7][i] = c_skb[7][i];
564 barrier (CLK_LOCAL_MEM_FENCE);
566 if (gid >= gid_max) return;
575 pw_buf0[0] = pws[gid].i[ 0];
576 pw_buf0[1] = pws[gid].i[ 1];
584 const u32 pw_l_len = pws[gid].pw_len;
592 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
593 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
599 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
601 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
603 const u32x pw_len = pw_l_len + pw_r_len;
606 * concat password candidate
609 u32x wordl0[4] = { 0 };
610 u32x wordl1[4] = { 0 };
611 u32x wordl2[4] = { 0 };
612 u32x wordl3[4] = { 0 };
614 wordl0[0] = pw_buf0[0];
615 wordl0[1] = pw_buf0[1];
616 wordl0[2] = pw_buf0[2];
617 wordl0[3] = pw_buf0[3];
618 wordl1[0] = pw_buf1[0];
619 wordl1[1] = pw_buf1[1];
620 wordl1[2] = pw_buf1[2];
621 wordl1[3] = pw_buf1[3];
623 u32x wordr0[4] = { 0 };
624 u32x wordr1[4] = { 0 };
625 u32x wordr2[4] = { 0 };
626 u32x wordr3[4] = { 0 };
628 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
629 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
630 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
631 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
632 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
633 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
634 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
635 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
637 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
639 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
643 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
648 w0[0] = wordl0[0] | wordr0[0];
649 w0[1] = wordl0[1] | wordr0[1];
657 transform_racf_key (w0[0], w0[1], key);
659 const u32x c = key[0];
660 const u32x d = key[1];
665 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
669 data[0] = salt_buf0[0];
670 data[1] = salt_buf0[1];
674 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
678 COMPARE_M_SIMD (iv[0], iv[1], z, z);
682 __kernel void m08500_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)
686 __kernel void m08500_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)
690 __kernel void m08500_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)
696 const u32 gid = get_global_id (0);
697 const u32 lid = get_local_id (0);
698 const u32 lsz = get_local_size (0);
704 __local u32 s_SPtrans[8][64];
705 __local u32 s_skb[8][64];
707 for (u32 i = lid; i < 64; i += lsz)
709 s_SPtrans[0][i] = c_SPtrans[0][i];
710 s_SPtrans[1][i] = c_SPtrans[1][i];
711 s_SPtrans[2][i] = c_SPtrans[2][i];
712 s_SPtrans[3][i] = c_SPtrans[3][i];
713 s_SPtrans[4][i] = c_SPtrans[4][i];
714 s_SPtrans[5][i] = c_SPtrans[5][i];
715 s_SPtrans[6][i] = c_SPtrans[6][i];
716 s_SPtrans[7][i] = c_SPtrans[7][i];
718 s_skb[0][i] = c_skb[0][i];
719 s_skb[1][i] = c_skb[1][i];
720 s_skb[2][i] = c_skb[2][i];
721 s_skb[3][i] = c_skb[3][i];
722 s_skb[4][i] = c_skb[4][i];
723 s_skb[5][i] = c_skb[5][i];
724 s_skb[6][i] = c_skb[6][i];
725 s_skb[7][i] = c_skb[7][i];
728 barrier (CLK_LOCAL_MEM_FENCE);
730 if (gid >= gid_max) return;
739 pw_buf0[0] = pws[gid].i[ 0];
740 pw_buf0[1] = pws[gid].i[ 1];
748 const u32 pw_l_len = pws[gid].pw_len;
756 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
757 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
763 const u32 search[4] =
765 digests_buf[digests_offset].digest_buf[DGST_R0],
766 digests_buf[digests_offset].digest_buf[DGST_R1],
775 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
777 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
779 const u32x pw_len = pw_l_len + pw_r_len;
782 * concat password candidate
785 u32x wordl0[4] = { 0 };
786 u32x wordl1[4] = { 0 };
787 u32x wordl2[4] = { 0 };
788 u32x wordl3[4] = { 0 };
790 wordl0[0] = pw_buf0[0];
791 wordl0[1] = pw_buf0[1];
792 wordl0[2] = pw_buf0[2];
793 wordl0[3] = pw_buf0[3];
794 wordl1[0] = pw_buf1[0];
795 wordl1[1] = pw_buf1[1];
796 wordl1[2] = pw_buf1[2];
797 wordl1[3] = pw_buf1[3];
799 u32x wordr0[4] = { 0 };
800 u32x wordr1[4] = { 0 };
801 u32x wordr2[4] = { 0 };
802 u32x wordr3[4] = { 0 };
804 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
805 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
806 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
807 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
808 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
809 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
810 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
811 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
813 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
815 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
819 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
824 w0[0] = wordl0[0] | wordr0[0];
825 w0[1] = wordl0[1] | wordr0[1];
833 transform_racf_key (w0[0], w0[1], key);
835 const u32x c = key[0];
836 const u32x d = key[1];
841 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
845 data[0] = salt_buf0[0];
846 data[1] = salt_buf0[1];
850 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
854 COMPARE_S_SIMD (iv[0], iv[1], z, z);
858 __kernel void m08500_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)
862 __kernel void m08500_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)