2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
4 * magnum <john.magnum@hushmail.com>
13 #include "include/constants.h"
14 #include "include/kernel_vendor.h"
21 #include "include/kernel_functions.c"
22 #include "OpenCL/types_ocl.c"
23 #include "OpenCL/common.c"
24 #include "OpenCL/simd.c"
26 #define PERM_OP(a,b,tt,n,m) \
36 #define HPERM_OP(a,tt,n,m) \
42 tt = tt >> (16 + n); \
48 PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \
49 PERM_OP (l, r, tt, 16, 0x0000ffff); \
50 PERM_OP (r, l, tt, 2, 0x33333333); \
51 PERM_OP (l, r, tt, 8, 0x00ff00ff); \
52 PERM_OP (r, l, tt, 1, 0x55555555); \
57 PERM_OP (l, r, tt, 1, 0x55555555); \
58 PERM_OP (r, l, tt, 8, 0x00ff00ff); \
59 PERM_OP (l, r, tt, 2, 0x33333333); \
60 PERM_OP (r, l, tt, 16, 0x0000ffff); \
61 PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
64 __constant u8 ascii_to_ebcdic_pc[256] =
66 // little hack, can't crack 0-bytes in password, but who cares
67 // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
68 0x2a, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
69 0x8a, 0x89, 0x8f, 0x8c, 0xd3, 0xd0, 0xce, 0xe6, 0x9b, 0x98, 0xd5, 0xe5, 0x92, 0x91, 0x97, 0x94,
70 0x2a, 0x34, 0x54, 0x5d, 0x1c, 0x73, 0x0b, 0x51, 0x31, 0x10, 0x13, 0x37, 0x7c, 0x6b, 0x3d, 0x68,
71 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40, 0x46, 0x45, 0x5b, 0x58, 0x5e, 0x16, 0x32, 0x57, 0x76, 0x75,
72 0x52, 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07,
73 0x04, 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x3e, 0x6b, 0x1f, 0x15, 0x70,
74 0x58, 0xa8, 0xae, 0xad, 0xa2, 0xa1, 0xa7, 0xa4, 0xba, 0xb9, 0x89, 0x8f, 0x8c, 0x83, 0x80, 0x86,
75 0x85, 0x9b, 0x98, 0xef, 0xec, 0xe3, 0xe0, 0xe6, 0xe5, 0xfb, 0xf8, 0x2a, 0x7f, 0x0b, 0xe9, 0xa4,
76 0xea, 0xe9, 0xef, 0xec, 0xe3, 0x80, 0xa7, 0x85, 0xfb, 0xf8, 0xfe, 0xfd, 0xf2, 0xb9, 0xbf, 0x9d,
77 0xcb, 0xc8, 0x9e, 0xcd, 0xc2, 0xc1, 0xc7, 0xba, 0xda, 0xd9, 0xdf, 0xdc, 0xa2, 0x83, 0xd6, 0x68,
78 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07, 0x04,
79 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40,
80 0x46, 0x45, 0x5b, 0xab, 0xbf, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, 0x8a, 0x9e, 0x9d, 0x92, 0x91, 0x97,
81 0x94, 0xea, 0xfe, 0xfd, 0xf2, 0xf1, 0xf7, 0xf4, 0xcb, 0xc8, 0xce, 0xcd, 0xc2, 0xc1, 0xc7, 0xc4,
82 0xda, 0xd9, 0xdf, 0xdc, 0xd3, 0xd0, 0xd6, 0xd5, 0x3e, 0x3d, 0x32, 0x31, 0x37, 0x34, 0x1f, 0x1c,
83 0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
86 __constant u32 c_SPtrans[8][64] =
89 0x02080800, 0x00080000, 0x02000002, 0x02080802,
90 0x02000000, 0x00080802, 0x00080002, 0x02000002,
91 0x00080802, 0x02080800, 0x02080000, 0x00000802,
92 0x02000802, 0x02000000, 0x00000000, 0x00080002,
93 0x00080000, 0x00000002, 0x02000800, 0x00080800,
94 0x02080802, 0x02080000, 0x00000802, 0x02000800,
95 0x00000002, 0x00000800, 0x00080800, 0x02080002,
96 0x00000800, 0x02000802, 0x02080002, 0x00000000,
97 0x00000000, 0x02080802, 0x02000800, 0x00080002,
98 0x02080800, 0x00080000, 0x00000802, 0x02000800,
99 0x02080002, 0x00000800, 0x00080800, 0x02000002,
100 0x00080802, 0x00000002, 0x02000002, 0x02080000,
101 0x02080802, 0x00080800, 0x02080000, 0x02000802,
102 0x02000000, 0x00000802, 0x00080002, 0x00000000,
103 0x00080000, 0x02000000, 0x02000802, 0x02080800,
104 0x00000002, 0x02080002, 0x00000800, 0x00080802,
107 0x40108010, 0x00000000, 0x00108000, 0x40100000,
108 0x40000010, 0x00008010, 0x40008000, 0x00108000,
109 0x00008000, 0x40100010, 0x00000010, 0x40008000,
110 0x00100010, 0x40108000, 0x40100000, 0x00000010,
111 0x00100000, 0x40008010, 0x40100010, 0x00008000,
112 0x00108010, 0x40000000, 0x00000000, 0x00100010,
113 0x40008010, 0x00108010, 0x40108000, 0x40000010,
114 0x40000000, 0x00100000, 0x00008010, 0x40108010,
115 0x00100010, 0x40108000, 0x40008000, 0x00108010,
116 0x40108010, 0x00100010, 0x40000010, 0x00000000,
117 0x40000000, 0x00008010, 0x00100000, 0x40100010,
118 0x00008000, 0x40000000, 0x00108010, 0x40008010,
119 0x40108000, 0x00008000, 0x00000000, 0x40000010,
120 0x00000010, 0x40108010, 0x00108000, 0x40100000,
121 0x40100010, 0x00100000, 0x00008010, 0x40008000,
122 0x40008010, 0x00000010, 0x40100000, 0x00108000,
125 0x04000001, 0x04040100, 0x00000100, 0x04000101,
126 0x00040001, 0x04000000, 0x04000101, 0x00040100,
127 0x04000100, 0x00040000, 0x04040000, 0x00000001,
128 0x04040101, 0x00000101, 0x00000001, 0x04040001,
129 0x00000000, 0x00040001, 0x04040100, 0x00000100,
130 0x00000101, 0x04040101, 0x00040000, 0x04000001,
131 0x04040001, 0x04000100, 0x00040101, 0x04040000,
132 0x00040100, 0x00000000, 0x04000000, 0x00040101,
133 0x04040100, 0x00000100, 0x00000001, 0x00040000,
134 0x00000101, 0x00040001, 0x04040000, 0x04000101,
135 0x00000000, 0x04040100, 0x00040100, 0x04040001,
136 0x00040001, 0x04000000, 0x04040101, 0x00000001,
137 0x00040101, 0x04000001, 0x04000000, 0x04040101,
138 0x00040000, 0x04000100, 0x04000101, 0x00040100,
139 0x04000100, 0x00000000, 0x04040001, 0x00000101,
140 0x04000001, 0x00040101, 0x00000100, 0x04040000,
143 0x00401008, 0x10001000, 0x00000008, 0x10401008,
144 0x00000000, 0x10400000, 0x10001008, 0x00400008,
145 0x10401000, 0x10000008, 0x10000000, 0x00001008,
146 0x10000008, 0x00401008, 0x00400000, 0x10000000,
147 0x10400008, 0x00401000, 0x00001000, 0x00000008,
148 0x00401000, 0x10001008, 0x10400000, 0x00001000,
149 0x00001008, 0x00000000, 0x00400008, 0x10401000,
150 0x10001000, 0x10400008, 0x10401008, 0x00400000,
151 0x10400008, 0x00001008, 0x00400000, 0x10000008,
152 0x00401000, 0x10001000, 0x00000008, 0x10400000,
153 0x10001008, 0x00000000, 0x00001000, 0x00400008,
154 0x00000000, 0x10400008, 0x10401000, 0x00001000,
155 0x10000000, 0x10401008, 0x00401008, 0x00400000,
156 0x10401008, 0x00000008, 0x10001000, 0x00401008,
157 0x00400008, 0x00401000, 0x10400000, 0x10001008,
158 0x00001008, 0x10000000, 0x10000008, 0x10401000,
161 0x08000000, 0x00010000, 0x00000400, 0x08010420,
162 0x08010020, 0x08000400, 0x00010420, 0x08010000,
163 0x00010000, 0x00000020, 0x08000020, 0x00010400,
164 0x08000420, 0x08010020, 0x08010400, 0x00000000,
165 0x00010400, 0x08000000, 0x00010020, 0x00000420,
166 0x08000400, 0x00010420, 0x00000000, 0x08000020,
167 0x00000020, 0x08000420, 0x08010420, 0x00010020,
168 0x08010000, 0x00000400, 0x00000420, 0x08010400,
169 0x08010400, 0x08000420, 0x00010020, 0x08010000,
170 0x00010000, 0x00000020, 0x08000020, 0x08000400,
171 0x08000000, 0x00010400, 0x08010420, 0x00000000,
172 0x00010420, 0x08000000, 0x00000400, 0x00010020,
173 0x08000420, 0x00000400, 0x00000000, 0x08010420,
174 0x08010020, 0x08010400, 0x00000420, 0x00010000,
175 0x00010400, 0x08010020, 0x08000400, 0x00000420,
176 0x00000020, 0x00010420, 0x08010000, 0x08000020,
179 0x80000040, 0x00200040, 0x00000000, 0x80202000,
180 0x00200040, 0x00002000, 0x80002040, 0x00200000,
181 0x00002040, 0x80202040, 0x00202000, 0x80000000,
182 0x80002000, 0x80000040, 0x80200000, 0x00202040,
183 0x00200000, 0x80002040, 0x80200040, 0x00000000,
184 0x00002000, 0x00000040, 0x80202000, 0x80200040,
185 0x80202040, 0x80200000, 0x80000000, 0x00002040,
186 0x00000040, 0x00202000, 0x00202040, 0x80002000,
187 0x00002040, 0x80000000, 0x80002000, 0x00202040,
188 0x80202000, 0x00200040, 0x00000000, 0x80002000,
189 0x80000000, 0x00002000, 0x80200040, 0x00200000,
190 0x00200040, 0x80202040, 0x00202000, 0x00000040,
191 0x80202040, 0x00202000, 0x00200000, 0x80002040,
192 0x80000040, 0x80200000, 0x00202040, 0x00000000,
193 0x00002000, 0x80000040, 0x80002040, 0x80202000,
194 0x80200000, 0x00002040, 0x00000040, 0x80200040,
197 0x00004000, 0x00000200, 0x01000200, 0x01000004,
198 0x01004204, 0x00004004, 0x00004200, 0x00000000,
199 0x01000000, 0x01000204, 0x00000204, 0x01004000,
200 0x00000004, 0x01004200, 0x01004000, 0x00000204,
201 0x01000204, 0x00004000, 0x00004004, 0x01004204,
202 0x00000000, 0x01000200, 0x01000004, 0x00004200,
203 0x01004004, 0x00004204, 0x01004200, 0x00000004,
204 0x00004204, 0x01004004, 0x00000200, 0x01000000,
205 0x00004204, 0x01004000, 0x01004004, 0x00000204,
206 0x00004000, 0x00000200, 0x01000000, 0x01004004,
207 0x01000204, 0x00004204, 0x00004200, 0x00000000,
208 0x00000200, 0x01000004, 0x00000004, 0x01000200,
209 0x00000000, 0x01000204, 0x01000200, 0x00004200,
210 0x00000204, 0x00004000, 0x01004204, 0x01000000,
211 0x01004200, 0x00000004, 0x00004004, 0x01004204,
212 0x01000004, 0x01004200, 0x01004000, 0x00004004,
215 0x20800080, 0x20820000, 0x00020080, 0x00000000,
216 0x20020000, 0x00800080, 0x20800000, 0x20820080,
217 0x00000080, 0x20000000, 0x00820000, 0x00020080,
218 0x00820080, 0x20020080, 0x20000080, 0x20800000,
219 0x00020000, 0x00820080, 0x00800080, 0x20020000,
220 0x20820080, 0x20000080, 0x00000000, 0x00820000,
221 0x20000000, 0x00800000, 0x20020080, 0x20800080,
222 0x00800000, 0x00020000, 0x20820000, 0x00000080,
223 0x00800000, 0x00020000, 0x20000080, 0x20820080,
224 0x00020080, 0x20000000, 0x00000000, 0x00820000,
225 0x20800080, 0x20020080, 0x20020000, 0x00800080,
226 0x20820000, 0x00000080, 0x00800080, 0x20020000,
227 0x20820080, 0x00800000, 0x20800000, 0x20000080,
228 0x00820000, 0x00020080, 0x20020080, 0x20800000,
229 0x00000080, 0x20820000, 0x00820080, 0x00000000,
230 0x20000000, 0x20800080, 0x00020000, 0x00820080,
234 __constant u32 c_skb[8][64] =
237 0x00000000, 0x00000010, 0x20000000, 0x20000010,
238 0x00010000, 0x00010010, 0x20010000, 0x20010010,
239 0x00000800, 0x00000810, 0x20000800, 0x20000810,
240 0x00010800, 0x00010810, 0x20010800, 0x20010810,
241 0x00000020, 0x00000030, 0x20000020, 0x20000030,
242 0x00010020, 0x00010030, 0x20010020, 0x20010030,
243 0x00000820, 0x00000830, 0x20000820, 0x20000830,
244 0x00010820, 0x00010830, 0x20010820, 0x20010830,
245 0x00080000, 0x00080010, 0x20080000, 0x20080010,
246 0x00090000, 0x00090010, 0x20090000, 0x20090010,
247 0x00080800, 0x00080810, 0x20080800, 0x20080810,
248 0x00090800, 0x00090810, 0x20090800, 0x20090810,
249 0x00080020, 0x00080030, 0x20080020, 0x20080030,
250 0x00090020, 0x00090030, 0x20090020, 0x20090030,
251 0x00080820, 0x00080830, 0x20080820, 0x20080830,
252 0x00090820, 0x00090830, 0x20090820, 0x20090830,
255 0x00000000, 0x02000000, 0x00002000, 0x02002000,
256 0x00200000, 0x02200000, 0x00202000, 0x02202000,
257 0x00000004, 0x02000004, 0x00002004, 0x02002004,
258 0x00200004, 0x02200004, 0x00202004, 0x02202004,
259 0x00000400, 0x02000400, 0x00002400, 0x02002400,
260 0x00200400, 0x02200400, 0x00202400, 0x02202400,
261 0x00000404, 0x02000404, 0x00002404, 0x02002404,
262 0x00200404, 0x02200404, 0x00202404, 0x02202404,
263 0x10000000, 0x12000000, 0x10002000, 0x12002000,
264 0x10200000, 0x12200000, 0x10202000, 0x12202000,
265 0x10000004, 0x12000004, 0x10002004, 0x12002004,
266 0x10200004, 0x12200004, 0x10202004, 0x12202004,
267 0x10000400, 0x12000400, 0x10002400, 0x12002400,
268 0x10200400, 0x12200400, 0x10202400, 0x12202400,
269 0x10000404, 0x12000404, 0x10002404, 0x12002404,
270 0x10200404, 0x12200404, 0x10202404, 0x12202404,
273 0x00000000, 0x00000001, 0x00040000, 0x00040001,
274 0x01000000, 0x01000001, 0x01040000, 0x01040001,
275 0x00000002, 0x00000003, 0x00040002, 0x00040003,
276 0x01000002, 0x01000003, 0x01040002, 0x01040003,
277 0x00000200, 0x00000201, 0x00040200, 0x00040201,
278 0x01000200, 0x01000201, 0x01040200, 0x01040201,
279 0x00000202, 0x00000203, 0x00040202, 0x00040203,
280 0x01000202, 0x01000203, 0x01040202, 0x01040203,
281 0x08000000, 0x08000001, 0x08040000, 0x08040001,
282 0x09000000, 0x09000001, 0x09040000, 0x09040001,
283 0x08000002, 0x08000003, 0x08040002, 0x08040003,
284 0x09000002, 0x09000003, 0x09040002, 0x09040003,
285 0x08000200, 0x08000201, 0x08040200, 0x08040201,
286 0x09000200, 0x09000201, 0x09040200, 0x09040201,
287 0x08000202, 0x08000203, 0x08040202, 0x08040203,
288 0x09000202, 0x09000203, 0x09040202, 0x09040203,
291 0x00000000, 0x00100000, 0x00000100, 0x00100100,
292 0x00000008, 0x00100008, 0x00000108, 0x00100108,
293 0x00001000, 0x00101000, 0x00001100, 0x00101100,
294 0x00001008, 0x00101008, 0x00001108, 0x00101108,
295 0x04000000, 0x04100000, 0x04000100, 0x04100100,
296 0x04000008, 0x04100008, 0x04000108, 0x04100108,
297 0x04001000, 0x04101000, 0x04001100, 0x04101100,
298 0x04001008, 0x04101008, 0x04001108, 0x04101108,
299 0x00020000, 0x00120000, 0x00020100, 0x00120100,
300 0x00020008, 0x00120008, 0x00020108, 0x00120108,
301 0x00021000, 0x00121000, 0x00021100, 0x00121100,
302 0x00021008, 0x00121008, 0x00021108, 0x00121108,
303 0x04020000, 0x04120000, 0x04020100, 0x04120100,
304 0x04020008, 0x04120008, 0x04020108, 0x04120108,
305 0x04021000, 0x04121000, 0x04021100, 0x04121100,
306 0x04021008, 0x04121008, 0x04021108, 0x04121108,
309 0x00000000, 0x10000000, 0x00010000, 0x10010000,
310 0x00000004, 0x10000004, 0x00010004, 0x10010004,
311 0x20000000, 0x30000000, 0x20010000, 0x30010000,
312 0x20000004, 0x30000004, 0x20010004, 0x30010004,
313 0x00100000, 0x10100000, 0x00110000, 0x10110000,
314 0x00100004, 0x10100004, 0x00110004, 0x10110004,
315 0x20100000, 0x30100000, 0x20110000, 0x30110000,
316 0x20100004, 0x30100004, 0x20110004, 0x30110004,
317 0x00001000, 0x10001000, 0x00011000, 0x10011000,
318 0x00001004, 0x10001004, 0x00011004, 0x10011004,
319 0x20001000, 0x30001000, 0x20011000, 0x30011000,
320 0x20001004, 0x30001004, 0x20011004, 0x30011004,
321 0x00101000, 0x10101000, 0x00111000, 0x10111000,
322 0x00101004, 0x10101004, 0x00111004, 0x10111004,
323 0x20101000, 0x30101000, 0x20111000, 0x30111000,
324 0x20101004, 0x30101004, 0x20111004, 0x30111004,
327 0x00000000, 0x08000000, 0x00000008, 0x08000008,
328 0x00000400, 0x08000400, 0x00000408, 0x08000408,
329 0x00020000, 0x08020000, 0x00020008, 0x08020008,
330 0x00020400, 0x08020400, 0x00020408, 0x08020408,
331 0x00000001, 0x08000001, 0x00000009, 0x08000009,
332 0x00000401, 0x08000401, 0x00000409, 0x08000409,
333 0x00020001, 0x08020001, 0x00020009, 0x08020009,
334 0x00020401, 0x08020401, 0x00020409, 0x08020409,
335 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
336 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
337 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
338 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
339 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
340 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
341 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
342 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
345 0x00000000, 0x00000100, 0x00080000, 0x00080100,
346 0x01000000, 0x01000100, 0x01080000, 0x01080100,
347 0x00000010, 0x00000110, 0x00080010, 0x00080110,
348 0x01000010, 0x01000110, 0x01080010, 0x01080110,
349 0x00200000, 0x00200100, 0x00280000, 0x00280100,
350 0x01200000, 0x01200100, 0x01280000, 0x01280100,
351 0x00200010, 0x00200110, 0x00280010, 0x00280110,
352 0x01200010, 0x01200110, 0x01280010, 0x01280110,
353 0x00000200, 0x00000300, 0x00080200, 0x00080300,
354 0x01000200, 0x01000300, 0x01080200, 0x01080300,
355 0x00000210, 0x00000310, 0x00080210, 0x00080310,
356 0x01000210, 0x01000310, 0x01080210, 0x01080310,
357 0x00200200, 0x00200300, 0x00280200, 0x00280300,
358 0x01200200, 0x01200300, 0x01280200, 0x01280300,
359 0x00200210, 0x00200310, 0x00280210, 0x00280310,
360 0x01200210, 0x01200310, 0x01280210, 0x01280310,
363 0x00000000, 0x04000000, 0x00040000, 0x04040000,
364 0x00000002, 0x04000002, 0x00040002, 0x04040002,
365 0x00002000, 0x04002000, 0x00042000, 0x04042000,
366 0x00002002, 0x04002002, 0x00042002, 0x04042002,
367 0x00000020, 0x04000020, 0x00040020, 0x04040020,
368 0x00000022, 0x04000022, 0x00040022, 0x04040022,
369 0x00002020, 0x04002020, 0x00042020, 0x04042020,
370 0x00002022, 0x04002022, 0x00042022, 0x04042022,
371 0x00000800, 0x04000800, 0x00040800, 0x04040800,
372 0x00000802, 0x04000802, 0x00040802, 0x04040802,
373 0x00002800, 0x04002800, 0x00042800, 0x04042800,
374 0x00002802, 0x04002802, 0x00042802, 0x04042802,
375 0x00000820, 0x04000820, 0x00040820, 0x04040820,
376 0x00000822, 0x04000822, 0x00040822, 0x04040822,
377 0x00002820, 0x04002820, 0x00042820, 0x04042820,
378 0x00002822, 0x04002822, 0x00042822, 0x04042822
383 #define BOX(i,n,S) (S)[(n)][(i)]
385 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
387 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
389 #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])
390 #elif VECT_SIZE == 16
391 #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])
395 #define BOX1(i,S) (S)[(i)]
397 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
399 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
401 #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])
402 #elif VECT_SIZE == 16
403 #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])
406 static void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
412 for (u32 i = 0; i < 16; i += 2)
418 t = Kd[i + 0] ^ rotl32 (r, 28u);
420 l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
421 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
422 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
423 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
424 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
425 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
426 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
427 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
430 t = Kd[i + 1] ^ rotl32 (l, 28u);
432 r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
433 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
434 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
435 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
436 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
437 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
438 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
439 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
446 static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
450 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
451 HPERM_OP (c, tt, 2, 0xcccc0000);
452 HPERM_OP (d, tt, 2, 0xcccc0000);
453 PERM_OP (d, c, tt, 1, 0x55555555);
454 PERM_OP (c, d, tt, 8, 0x00ff00ff);
455 PERM_OP (d, c, tt, 1, 0x55555555);
457 d = ((d & 0x000000ff) << 16)
458 | ((d & 0x0000ff00) << 0)
459 | ((d & 0x00ff0000) >> 16)
460 | ((c & 0xf0000000) >> 4);
465 for (u32 i = 0; i < 16; i++)
467 if ((i < 2) || (i == 8) || (i == 15))
469 c = ((c >> 1) | (c << 27));
470 d = ((d >> 1) | (d << 27));
474 c = ((c >> 2) | (c << 26));
475 d = ((d >> 2) | (d << 26));
481 const u32x c00 = (c >> 0) & 0x0000003f;
482 const u32x c06 = (c >> 6) & 0x00383003;
483 const u32x c07 = (c >> 7) & 0x0000003c;
484 const u32x c13 = (c >> 13) & 0x0000060f;
485 const u32x c20 = (c >> 20) & 0x00000001;
487 u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
488 | BOX (((c06 >> 0) & 0xff)
489 |((c07 >> 0) & 0xff), 1, s_skb)
490 | BOX (((c13 >> 0) & 0xff)
491 |((c06 >> 8) & 0xff), 2, s_skb)
492 | BOX (((c20 >> 0) & 0xff)
494 |((c06 >> 16) & 0xff), 3, s_skb);
496 const u32x d00 = (d >> 0) & 0x00003c3f;
497 const u32x d07 = (d >> 7) & 0x00003f03;
498 const u32x d21 = (d >> 21) & 0x0000000f;
499 const u32x d22 = (d >> 22) & 0x00000030;
501 u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
502 | BOX (((d07 >> 0) & 0xff)
503 |((d00 >> 8) & 0xff), 5, s_skb)
504 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
505 | BOX (((d21 >> 0) & 0xff)
506 |((d22 >> 0) & 0xff), 7, s_skb);
508 Kc[i] = ((t << 16) | (s & 0x0000ffff));
509 Kd[i] = ((s >> 16) | (t & 0xffff0000));
511 Kc[i] = rotl32 (Kc[i], 2u);
512 Kd[i] = rotl32 (Kd[i], 2u);
516 static void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
518 key[0] = BOX1 (((w0 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0
519 | BOX1 (((w0 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8
520 | BOX1 (((w0 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16
521 | BOX1 (((w0 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24;
523 key[1] = BOX1 (((w1 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0
524 | BOX1 (((w1 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8
525 | BOX1 (((w1 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16
526 | BOX1 (((w1 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24;
529 static void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
535 const u32 gid = get_global_id (0);
536 const u32 lid = get_local_id (0);
544 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
545 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
555 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
557 const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
559 const u32x w0 = w0l | w0r;
567 transform_racf_key (w0, w1, key);
569 const u32x c = key[0];
570 const u32x d = key[1];
575 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
579 data[0] = salt_buf0[0];
580 data[1] = salt_buf0[1];
584 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
589 COMPARE_M_SIMD (iv[0], iv[1], iv2, iv3);
593 static void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
599 const u32 gid = get_global_id (0);
600 const u32 lid = get_local_id (0);
608 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
609 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
615 const u32 search[4] =
617 digests_buf[digests_offset].digest_buf[DGST_R0],
618 digests_buf[digests_offset].digest_buf[DGST_R1],
619 digests_buf[digests_offset].digest_buf[DGST_R2],
620 digests_buf[digests_offset].digest_buf[DGST_R3]
631 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
633 const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
635 const u32x w0 = w0l | w0r;
643 transform_racf_key (w0, w1, key);
645 const u32x c = key[0];
646 const u32x d = key[1];
651 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
655 data[0] = salt_buf0[0];
656 data[1] = salt_buf0[1];
660 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
665 COMPARE_S_SIMD (iv[0], iv[1], iv2, iv3);
669 __kernel void m08500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
675 const u32 gid = get_global_id (0);
676 const u32 lid = get_local_id (0);
677 const u32 lsz = get_local_size (0);
683 __local u32 s_SPtrans[8][64];
684 __local u32 s_skb[8][64];
686 for (u32 i = lid; i < 64; i += lsz)
688 s_SPtrans[0][i] = c_SPtrans[0][i];
689 s_SPtrans[1][i] = c_SPtrans[1][i];
690 s_SPtrans[2][i] = c_SPtrans[2][i];
691 s_SPtrans[3][i] = c_SPtrans[3][i];
692 s_SPtrans[4][i] = c_SPtrans[4][i];
693 s_SPtrans[5][i] = c_SPtrans[5][i];
694 s_SPtrans[6][i] = c_SPtrans[6][i];
695 s_SPtrans[7][i] = c_SPtrans[7][i];
697 s_skb[0][i] = c_skb[0][i];
698 s_skb[1][i] = c_skb[1][i];
699 s_skb[2][i] = c_skb[2][i];
700 s_skb[3][i] = c_skb[3][i];
701 s_skb[4][i] = c_skb[4][i];
702 s_skb[5][i] = c_skb[5][i];
703 s_skb[6][i] = c_skb[6][i];
704 s_skb[7][i] = c_skb[7][i];
707 barrier (CLK_LOCAL_MEM_FENCE);
709 if (gid >= gid_max) return;
717 w[ 0] = pws[gid].i[ 0];
718 w[ 1] = pws[gid].i[ 1];
734 const u32 pw_len = pws[gid].pw_len;
740 m08500m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
743 __kernel void m08500_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
747 __kernel void m08500_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
751 __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
757 const u32 gid = get_global_id (0);
758 const u32 lid = get_local_id (0);
759 const u32 lsz = get_local_size (0);
765 __local u32 s_SPtrans[8][64];
766 __local u32 s_skb[8][64];
768 for (u32 i = lid; i < 64; i += lsz)
770 s_SPtrans[0][i] = c_SPtrans[0][i];
771 s_SPtrans[1][i] = c_SPtrans[1][i];
772 s_SPtrans[2][i] = c_SPtrans[2][i];
773 s_SPtrans[3][i] = c_SPtrans[3][i];
774 s_SPtrans[4][i] = c_SPtrans[4][i];
775 s_SPtrans[5][i] = c_SPtrans[5][i];
776 s_SPtrans[6][i] = c_SPtrans[6][i];
777 s_SPtrans[7][i] = c_SPtrans[7][i];
779 s_skb[0][i] = c_skb[0][i];
780 s_skb[1][i] = c_skb[1][i];
781 s_skb[2][i] = c_skb[2][i];
782 s_skb[3][i] = c_skb[3][i];
783 s_skb[4][i] = c_skb[4][i];
784 s_skb[5][i] = c_skb[5][i];
785 s_skb[6][i] = c_skb[6][i];
786 s_skb[7][i] = c_skb[7][i];
789 barrier (CLK_LOCAL_MEM_FENCE);
791 if (gid >= gid_max) return;
799 w[ 0] = pws[gid].i[ 0];
800 w[ 1] = pws[gid].i[ 1];
816 const u32 pw_len = pws[gid].pw_len;
822 m08500s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
825 __kernel void m08500_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
829 __kernel void m08500_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)