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 "include/rp_kernel.h"
25 #include "OpenCL/rp.c"
26 #include "OpenCL/simd.c"
28 #define PERM_OP(a,b,tt,n,m) \
38 #define HPERM_OP(a,tt,n,m) \
44 tt = tt >> (16 + n); \
50 PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \
51 PERM_OP (l, r, tt, 16, 0x0000ffff); \
52 PERM_OP (r, l, tt, 2, 0x33333333); \
53 PERM_OP (l, r, tt, 8, 0x00ff00ff); \
54 PERM_OP (r, l, tt, 1, 0x55555555); \
59 PERM_OP (l, r, tt, 1, 0x55555555); \
60 PERM_OP (r, l, tt, 8, 0x00ff00ff); \
61 PERM_OP (l, r, tt, 2, 0x33333333); \
62 PERM_OP (r, l, tt, 16, 0x0000ffff); \
63 PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
66 __constant u8 ascii_to_ebcdic_pc[256] =
68 // little hack, can't crack 0-bytes in password, but who cares
69 // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
70 0x2a, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
71 0x8a, 0x89, 0x8f, 0x8c, 0xd3, 0xd0, 0xce, 0xe6, 0x9b, 0x98, 0xd5, 0xe5, 0x92, 0x91, 0x97, 0x94,
72 0x2a, 0x34, 0x54, 0x5d, 0x1c, 0x73, 0x0b, 0x51, 0x31, 0x10, 0x13, 0x37, 0x7c, 0x6b, 0x3d, 0x68,
73 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40, 0x46, 0x45, 0x5b, 0x58, 0x5e, 0x16, 0x32, 0x57, 0x76, 0x75,
74 0x52, 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07,
75 0x04, 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x3e, 0x6b, 0x1f, 0x15, 0x70,
76 0x58, 0xa8, 0xae, 0xad, 0xa2, 0xa1, 0xa7, 0xa4, 0xba, 0xb9, 0x89, 0x8f, 0x8c, 0x83, 0x80, 0x86,
77 0x85, 0x9b, 0x98, 0xef, 0xec, 0xe3, 0xe0, 0xe6, 0xe5, 0xfb, 0xf8, 0x2a, 0x7f, 0x0b, 0xe9, 0xa4,
78 0xea, 0xe9, 0xef, 0xec, 0xe3, 0x80, 0xa7, 0x85, 0xfb, 0xf8, 0xfe, 0xfd, 0xf2, 0xb9, 0xbf, 0x9d,
79 0xcb, 0xc8, 0x9e, 0xcd, 0xc2, 0xc1, 0xc7, 0xba, 0xda, 0xd9, 0xdf, 0xdc, 0xa2, 0x83, 0xd6, 0x68,
80 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07, 0x04,
81 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40,
82 0x46, 0x45, 0x5b, 0xab, 0xbf, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, 0x8a, 0x9e, 0x9d, 0x92, 0x91, 0x97,
83 0x94, 0xea, 0xfe, 0xfd, 0xf2, 0xf1, 0xf7, 0xf4, 0xcb, 0xc8, 0xce, 0xcd, 0xc2, 0xc1, 0xc7, 0xc4,
84 0xda, 0xd9, 0xdf, 0xdc, 0xd3, 0xd0, 0xd6, 0xd5, 0x3e, 0x3d, 0x32, 0x31, 0x37, 0x34, 0x1f, 0x1c,
85 0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
88 __constant u32 c_SPtrans[8][64] =
91 0x02080800, 0x00080000, 0x02000002, 0x02080802,
92 0x02000000, 0x00080802, 0x00080002, 0x02000002,
93 0x00080802, 0x02080800, 0x02080000, 0x00000802,
94 0x02000802, 0x02000000, 0x00000000, 0x00080002,
95 0x00080000, 0x00000002, 0x02000800, 0x00080800,
96 0x02080802, 0x02080000, 0x00000802, 0x02000800,
97 0x00000002, 0x00000800, 0x00080800, 0x02080002,
98 0x00000800, 0x02000802, 0x02080002, 0x00000000,
99 0x00000000, 0x02080802, 0x02000800, 0x00080002,
100 0x02080800, 0x00080000, 0x00000802, 0x02000800,
101 0x02080002, 0x00000800, 0x00080800, 0x02000002,
102 0x00080802, 0x00000002, 0x02000002, 0x02080000,
103 0x02080802, 0x00080800, 0x02080000, 0x02000802,
104 0x02000000, 0x00000802, 0x00080002, 0x00000000,
105 0x00080000, 0x02000000, 0x02000802, 0x02080800,
106 0x00000002, 0x02080002, 0x00000800, 0x00080802,
109 0x40108010, 0x00000000, 0x00108000, 0x40100000,
110 0x40000010, 0x00008010, 0x40008000, 0x00108000,
111 0x00008000, 0x40100010, 0x00000010, 0x40008000,
112 0x00100010, 0x40108000, 0x40100000, 0x00000010,
113 0x00100000, 0x40008010, 0x40100010, 0x00008000,
114 0x00108010, 0x40000000, 0x00000000, 0x00100010,
115 0x40008010, 0x00108010, 0x40108000, 0x40000010,
116 0x40000000, 0x00100000, 0x00008010, 0x40108010,
117 0x00100010, 0x40108000, 0x40008000, 0x00108010,
118 0x40108010, 0x00100010, 0x40000010, 0x00000000,
119 0x40000000, 0x00008010, 0x00100000, 0x40100010,
120 0x00008000, 0x40000000, 0x00108010, 0x40008010,
121 0x40108000, 0x00008000, 0x00000000, 0x40000010,
122 0x00000010, 0x40108010, 0x00108000, 0x40100000,
123 0x40100010, 0x00100000, 0x00008010, 0x40008000,
124 0x40008010, 0x00000010, 0x40100000, 0x00108000,
127 0x04000001, 0x04040100, 0x00000100, 0x04000101,
128 0x00040001, 0x04000000, 0x04000101, 0x00040100,
129 0x04000100, 0x00040000, 0x04040000, 0x00000001,
130 0x04040101, 0x00000101, 0x00000001, 0x04040001,
131 0x00000000, 0x00040001, 0x04040100, 0x00000100,
132 0x00000101, 0x04040101, 0x00040000, 0x04000001,
133 0x04040001, 0x04000100, 0x00040101, 0x04040000,
134 0x00040100, 0x00000000, 0x04000000, 0x00040101,
135 0x04040100, 0x00000100, 0x00000001, 0x00040000,
136 0x00000101, 0x00040001, 0x04040000, 0x04000101,
137 0x00000000, 0x04040100, 0x00040100, 0x04040001,
138 0x00040001, 0x04000000, 0x04040101, 0x00000001,
139 0x00040101, 0x04000001, 0x04000000, 0x04040101,
140 0x00040000, 0x04000100, 0x04000101, 0x00040100,
141 0x04000100, 0x00000000, 0x04040001, 0x00000101,
142 0x04000001, 0x00040101, 0x00000100, 0x04040000,
145 0x00401008, 0x10001000, 0x00000008, 0x10401008,
146 0x00000000, 0x10400000, 0x10001008, 0x00400008,
147 0x10401000, 0x10000008, 0x10000000, 0x00001008,
148 0x10000008, 0x00401008, 0x00400000, 0x10000000,
149 0x10400008, 0x00401000, 0x00001000, 0x00000008,
150 0x00401000, 0x10001008, 0x10400000, 0x00001000,
151 0x00001008, 0x00000000, 0x00400008, 0x10401000,
152 0x10001000, 0x10400008, 0x10401008, 0x00400000,
153 0x10400008, 0x00001008, 0x00400000, 0x10000008,
154 0x00401000, 0x10001000, 0x00000008, 0x10400000,
155 0x10001008, 0x00000000, 0x00001000, 0x00400008,
156 0x00000000, 0x10400008, 0x10401000, 0x00001000,
157 0x10000000, 0x10401008, 0x00401008, 0x00400000,
158 0x10401008, 0x00000008, 0x10001000, 0x00401008,
159 0x00400008, 0x00401000, 0x10400000, 0x10001008,
160 0x00001008, 0x10000000, 0x10000008, 0x10401000,
163 0x08000000, 0x00010000, 0x00000400, 0x08010420,
164 0x08010020, 0x08000400, 0x00010420, 0x08010000,
165 0x00010000, 0x00000020, 0x08000020, 0x00010400,
166 0x08000420, 0x08010020, 0x08010400, 0x00000000,
167 0x00010400, 0x08000000, 0x00010020, 0x00000420,
168 0x08000400, 0x00010420, 0x00000000, 0x08000020,
169 0x00000020, 0x08000420, 0x08010420, 0x00010020,
170 0x08010000, 0x00000400, 0x00000420, 0x08010400,
171 0x08010400, 0x08000420, 0x00010020, 0x08010000,
172 0x00010000, 0x00000020, 0x08000020, 0x08000400,
173 0x08000000, 0x00010400, 0x08010420, 0x00000000,
174 0x00010420, 0x08000000, 0x00000400, 0x00010020,
175 0x08000420, 0x00000400, 0x00000000, 0x08010420,
176 0x08010020, 0x08010400, 0x00000420, 0x00010000,
177 0x00010400, 0x08010020, 0x08000400, 0x00000420,
178 0x00000020, 0x00010420, 0x08010000, 0x08000020,
181 0x80000040, 0x00200040, 0x00000000, 0x80202000,
182 0x00200040, 0x00002000, 0x80002040, 0x00200000,
183 0x00002040, 0x80202040, 0x00202000, 0x80000000,
184 0x80002000, 0x80000040, 0x80200000, 0x00202040,
185 0x00200000, 0x80002040, 0x80200040, 0x00000000,
186 0x00002000, 0x00000040, 0x80202000, 0x80200040,
187 0x80202040, 0x80200000, 0x80000000, 0x00002040,
188 0x00000040, 0x00202000, 0x00202040, 0x80002000,
189 0x00002040, 0x80000000, 0x80002000, 0x00202040,
190 0x80202000, 0x00200040, 0x00000000, 0x80002000,
191 0x80000000, 0x00002000, 0x80200040, 0x00200000,
192 0x00200040, 0x80202040, 0x00202000, 0x00000040,
193 0x80202040, 0x00202000, 0x00200000, 0x80002040,
194 0x80000040, 0x80200000, 0x00202040, 0x00000000,
195 0x00002000, 0x80000040, 0x80002040, 0x80202000,
196 0x80200000, 0x00002040, 0x00000040, 0x80200040,
199 0x00004000, 0x00000200, 0x01000200, 0x01000004,
200 0x01004204, 0x00004004, 0x00004200, 0x00000000,
201 0x01000000, 0x01000204, 0x00000204, 0x01004000,
202 0x00000004, 0x01004200, 0x01004000, 0x00000204,
203 0x01000204, 0x00004000, 0x00004004, 0x01004204,
204 0x00000000, 0x01000200, 0x01000004, 0x00004200,
205 0x01004004, 0x00004204, 0x01004200, 0x00000004,
206 0x00004204, 0x01004004, 0x00000200, 0x01000000,
207 0x00004204, 0x01004000, 0x01004004, 0x00000204,
208 0x00004000, 0x00000200, 0x01000000, 0x01004004,
209 0x01000204, 0x00004204, 0x00004200, 0x00000000,
210 0x00000200, 0x01000004, 0x00000004, 0x01000200,
211 0x00000000, 0x01000204, 0x01000200, 0x00004200,
212 0x00000204, 0x00004000, 0x01004204, 0x01000000,
213 0x01004200, 0x00000004, 0x00004004, 0x01004204,
214 0x01000004, 0x01004200, 0x01004000, 0x00004004,
217 0x20800080, 0x20820000, 0x00020080, 0x00000000,
218 0x20020000, 0x00800080, 0x20800000, 0x20820080,
219 0x00000080, 0x20000000, 0x00820000, 0x00020080,
220 0x00820080, 0x20020080, 0x20000080, 0x20800000,
221 0x00020000, 0x00820080, 0x00800080, 0x20020000,
222 0x20820080, 0x20000080, 0x00000000, 0x00820000,
223 0x20000000, 0x00800000, 0x20020080, 0x20800080,
224 0x00800000, 0x00020000, 0x20820000, 0x00000080,
225 0x00800000, 0x00020000, 0x20000080, 0x20820080,
226 0x00020080, 0x20000000, 0x00000000, 0x00820000,
227 0x20800080, 0x20020080, 0x20020000, 0x00800080,
228 0x20820000, 0x00000080, 0x00800080, 0x20020000,
229 0x20820080, 0x00800000, 0x20800000, 0x20000080,
230 0x00820000, 0x00020080, 0x20020080, 0x20800000,
231 0x00000080, 0x20820000, 0x00820080, 0x00000000,
232 0x20000000, 0x20800080, 0x00020000, 0x00820080,
236 __constant u32 c_skb[8][64] =
239 0x00000000, 0x00000010, 0x20000000, 0x20000010,
240 0x00010000, 0x00010010, 0x20010000, 0x20010010,
241 0x00000800, 0x00000810, 0x20000800, 0x20000810,
242 0x00010800, 0x00010810, 0x20010800, 0x20010810,
243 0x00000020, 0x00000030, 0x20000020, 0x20000030,
244 0x00010020, 0x00010030, 0x20010020, 0x20010030,
245 0x00000820, 0x00000830, 0x20000820, 0x20000830,
246 0x00010820, 0x00010830, 0x20010820, 0x20010830,
247 0x00080000, 0x00080010, 0x20080000, 0x20080010,
248 0x00090000, 0x00090010, 0x20090000, 0x20090010,
249 0x00080800, 0x00080810, 0x20080800, 0x20080810,
250 0x00090800, 0x00090810, 0x20090800, 0x20090810,
251 0x00080020, 0x00080030, 0x20080020, 0x20080030,
252 0x00090020, 0x00090030, 0x20090020, 0x20090030,
253 0x00080820, 0x00080830, 0x20080820, 0x20080830,
254 0x00090820, 0x00090830, 0x20090820, 0x20090830,
257 0x00000000, 0x02000000, 0x00002000, 0x02002000,
258 0x00200000, 0x02200000, 0x00202000, 0x02202000,
259 0x00000004, 0x02000004, 0x00002004, 0x02002004,
260 0x00200004, 0x02200004, 0x00202004, 0x02202004,
261 0x00000400, 0x02000400, 0x00002400, 0x02002400,
262 0x00200400, 0x02200400, 0x00202400, 0x02202400,
263 0x00000404, 0x02000404, 0x00002404, 0x02002404,
264 0x00200404, 0x02200404, 0x00202404, 0x02202404,
265 0x10000000, 0x12000000, 0x10002000, 0x12002000,
266 0x10200000, 0x12200000, 0x10202000, 0x12202000,
267 0x10000004, 0x12000004, 0x10002004, 0x12002004,
268 0x10200004, 0x12200004, 0x10202004, 0x12202004,
269 0x10000400, 0x12000400, 0x10002400, 0x12002400,
270 0x10200400, 0x12200400, 0x10202400, 0x12202400,
271 0x10000404, 0x12000404, 0x10002404, 0x12002404,
272 0x10200404, 0x12200404, 0x10202404, 0x12202404,
275 0x00000000, 0x00000001, 0x00040000, 0x00040001,
276 0x01000000, 0x01000001, 0x01040000, 0x01040001,
277 0x00000002, 0x00000003, 0x00040002, 0x00040003,
278 0x01000002, 0x01000003, 0x01040002, 0x01040003,
279 0x00000200, 0x00000201, 0x00040200, 0x00040201,
280 0x01000200, 0x01000201, 0x01040200, 0x01040201,
281 0x00000202, 0x00000203, 0x00040202, 0x00040203,
282 0x01000202, 0x01000203, 0x01040202, 0x01040203,
283 0x08000000, 0x08000001, 0x08040000, 0x08040001,
284 0x09000000, 0x09000001, 0x09040000, 0x09040001,
285 0x08000002, 0x08000003, 0x08040002, 0x08040003,
286 0x09000002, 0x09000003, 0x09040002, 0x09040003,
287 0x08000200, 0x08000201, 0x08040200, 0x08040201,
288 0x09000200, 0x09000201, 0x09040200, 0x09040201,
289 0x08000202, 0x08000203, 0x08040202, 0x08040203,
290 0x09000202, 0x09000203, 0x09040202, 0x09040203,
293 0x00000000, 0x00100000, 0x00000100, 0x00100100,
294 0x00000008, 0x00100008, 0x00000108, 0x00100108,
295 0x00001000, 0x00101000, 0x00001100, 0x00101100,
296 0x00001008, 0x00101008, 0x00001108, 0x00101108,
297 0x04000000, 0x04100000, 0x04000100, 0x04100100,
298 0x04000008, 0x04100008, 0x04000108, 0x04100108,
299 0x04001000, 0x04101000, 0x04001100, 0x04101100,
300 0x04001008, 0x04101008, 0x04001108, 0x04101108,
301 0x00020000, 0x00120000, 0x00020100, 0x00120100,
302 0x00020008, 0x00120008, 0x00020108, 0x00120108,
303 0x00021000, 0x00121000, 0x00021100, 0x00121100,
304 0x00021008, 0x00121008, 0x00021108, 0x00121108,
305 0x04020000, 0x04120000, 0x04020100, 0x04120100,
306 0x04020008, 0x04120008, 0x04020108, 0x04120108,
307 0x04021000, 0x04121000, 0x04021100, 0x04121100,
308 0x04021008, 0x04121008, 0x04021108, 0x04121108,
311 0x00000000, 0x10000000, 0x00010000, 0x10010000,
312 0x00000004, 0x10000004, 0x00010004, 0x10010004,
313 0x20000000, 0x30000000, 0x20010000, 0x30010000,
314 0x20000004, 0x30000004, 0x20010004, 0x30010004,
315 0x00100000, 0x10100000, 0x00110000, 0x10110000,
316 0x00100004, 0x10100004, 0x00110004, 0x10110004,
317 0x20100000, 0x30100000, 0x20110000, 0x30110000,
318 0x20100004, 0x30100004, 0x20110004, 0x30110004,
319 0x00001000, 0x10001000, 0x00011000, 0x10011000,
320 0x00001004, 0x10001004, 0x00011004, 0x10011004,
321 0x20001000, 0x30001000, 0x20011000, 0x30011000,
322 0x20001004, 0x30001004, 0x20011004, 0x30011004,
323 0x00101000, 0x10101000, 0x00111000, 0x10111000,
324 0x00101004, 0x10101004, 0x00111004, 0x10111004,
325 0x20101000, 0x30101000, 0x20111000, 0x30111000,
326 0x20101004, 0x30101004, 0x20111004, 0x30111004,
329 0x00000000, 0x08000000, 0x00000008, 0x08000008,
330 0x00000400, 0x08000400, 0x00000408, 0x08000408,
331 0x00020000, 0x08020000, 0x00020008, 0x08020008,
332 0x00020400, 0x08020400, 0x00020408, 0x08020408,
333 0x00000001, 0x08000001, 0x00000009, 0x08000009,
334 0x00000401, 0x08000401, 0x00000409, 0x08000409,
335 0x00020001, 0x08020001, 0x00020009, 0x08020009,
336 0x00020401, 0x08020401, 0x00020409, 0x08020409,
337 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
338 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
339 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
340 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
341 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
342 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
343 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
344 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
347 0x00000000, 0x00000100, 0x00080000, 0x00080100,
348 0x01000000, 0x01000100, 0x01080000, 0x01080100,
349 0x00000010, 0x00000110, 0x00080010, 0x00080110,
350 0x01000010, 0x01000110, 0x01080010, 0x01080110,
351 0x00200000, 0x00200100, 0x00280000, 0x00280100,
352 0x01200000, 0x01200100, 0x01280000, 0x01280100,
353 0x00200010, 0x00200110, 0x00280010, 0x00280110,
354 0x01200010, 0x01200110, 0x01280010, 0x01280110,
355 0x00000200, 0x00000300, 0x00080200, 0x00080300,
356 0x01000200, 0x01000300, 0x01080200, 0x01080300,
357 0x00000210, 0x00000310, 0x00080210, 0x00080310,
358 0x01000210, 0x01000310, 0x01080210, 0x01080310,
359 0x00200200, 0x00200300, 0x00280200, 0x00280300,
360 0x01200200, 0x01200300, 0x01280200, 0x01280300,
361 0x00200210, 0x00200310, 0x00280210, 0x00280310,
362 0x01200210, 0x01200310, 0x01280210, 0x01280310,
365 0x00000000, 0x04000000, 0x00040000, 0x04040000,
366 0x00000002, 0x04000002, 0x00040002, 0x04040002,
367 0x00002000, 0x04002000, 0x00042000, 0x04042000,
368 0x00002002, 0x04002002, 0x00042002, 0x04042002,
369 0x00000020, 0x04000020, 0x00040020, 0x04040020,
370 0x00000022, 0x04000022, 0x00040022, 0x04040022,
371 0x00002020, 0x04002020, 0x00042020, 0x04042020,
372 0x00002022, 0x04002022, 0x00042022, 0x04042022,
373 0x00000800, 0x04000800, 0x00040800, 0x04040800,
374 0x00000802, 0x04000802, 0x00040802, 0x04040802,
375 0x00002800, 0x04002800, 0x00042800, 0x04042800,
376 0x00002802, 0x04002802, 0x00042802, 0x04042802,
377 0x00000820, 0x04000820, 0x00040820, 0x04040820,
378 0x00000822, 0x04000822, 0x00040822, 0x04040822,
379 0x00002820, 0x04002820, 0x00042820, 0x04042820,
380 0x00002822, 0x04002822, 0x00042822, 0x04042822
385 #define BOX(i,n,S) (S)[(n)][(i)]
387 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
389 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
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])
392 #elif VECT_SIZE == 16
393 #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])
397 #define BOX1(i,S) (S)[(i)]
399 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
401 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
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])
404 #elif VECT_SIZE == 16
405 #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])
408 void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
414 for (u32 i = 0; i < 16; i += 2)
420 t = Kd[i + 0] ^ rotl32 (r, 28u);
422 l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
423 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
424 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
425 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
426 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
427 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
428 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
429 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
432 t = Kd[i + 1] ^ rotl32 (l, 28u);
434 r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
435 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
436 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
437 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
438 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
439 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
440 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
441 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
448 void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
452 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
453 HPERM_OP (c, tt, 2, 0xcccc0000);
454 HPERM_OP (d, tt, 2, 0xcccc0000);
455 PERM_OP (d, c, tt, 1, 0x55555555);
456 PERM_OP (c, d, tt, 8, 0x00ff00ff);
457 PERM_OP (d, c, tt, 1, 0x55555555);
459 d = ((d & 0x000000ff) << 16)
460 | ((d & 0x0000ff00) << 0)
461 | ((d & 0x00ff0000) >> 16)
462 | ((c & 0xf0000000) >> 4);
467 for (u32 i = 0; i < 16; i++)
469 if ((i < 2) || (i == 8) || (i == 15))
471 c = ((c >> 1) | (c << 27));
472 d = ((d >> 1) | (d << 27));
476 c = ((c >> 2) | (c << 26));
477 d = ((d >> 2) | (d << 26));
483 const u32x c00 = (c >> 0) & 0x0000003f;
484 const u32x c06 = (c >> 6) & 0x00383003;
485 const u32x c07 = (c >> 7) & 0x0000003c;
486 const u32x c13 = (c >> 13) & 0x0000060f;
487 const u32x c20 = (c >> 20) & 0x00000001;
489 u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
490 | BOX (((c06 >> 0) & 0xff)
491 |((c07 >> 0) & 0xff), 1, s_skb)
492 | BOX (((c13 >> 0) & 0xff)
493 |((c06 >> 8) & 0xff), 2, s_skb)
494 | BOX (((c20 >> 0) & 0xff)
496 |((c06 >> 16) & 0xff), 3, s_skb);
498 const u32x d00 = (d >> 0) & 0x00003c3f;
499 const u32x d07 = (d >> 7) & 0x00003f03;
500 const u32x d21 = (d >> 21) & 0x0000000f;
501 const u32x d22 = (d >> 22) & 0x00000030;
503 u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
504 | BOX (((d07 >> 0) & 0xff)
505 |((d00 >> 8) & 0xff), 5, s_skb)
506 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
507 | BOX (((d21 >> 0) & 0xff)
508 |((d22 >> 0) & 0xff), 7, s_skb);
510 Kc[i] = ((t << 16) | (s & 0x0000ffff));
511 Kd[i] = ((s >> 16) | (t & 0xffff0000));
513 Kc[i] = rotl32 (Kc[i], 2u);
514 Kd[i] = rotl32 (Kd[i], 2u);
518 void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
520 key[0] = BOX1 (((w0 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0
521 | BOX1 (((w0 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8
522 | BOX1 (((w0 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16
523 | BOX1 (((w0 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24;
525 key[1] = BOX1 (((w1 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0
526 | BOX1 (((w1 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8
527 | BOX1 (((w1 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16
528 | BOX1 (((w1 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24;
531 __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_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)
537 const u32 gid = get_global_id (0);
538 const u32 lid = get_local_id (0);
539 const u32 lsz = get_local_size (0);
545 __local u32 s_SPtrans[8][64];
546 __local u32 s_skb[8][64];
548 for (u32 i = lid; i < 64; i += lsz)
550 s_SPtrans[0][i] = c_SPtrans[0][i];
551 s_SPtrans[1][i] = c_SPtrans[1][i];
552 s_SPtrans[2][i] = c_SPtrans[2][i];
553 s_SPtrans[3][i] = c_SPtrans[3][i];
554 s_SPtrans[4][i] = c_SPtrans[4][i];
555 s_SPtrans[5][i] = c_SPtrans[5][i];
556 s_SPtrans[6][i] = c_SPtrans[6][i];
557 s_SPtrans[7][i] = c_SPtrans[7][i];
559 s_skb[0][i] = c_skb[0][i];
560 s_skb[1][i] = c_skb[1][i];
561 s_skb[2][i] = c_skb[2][i];
562 s_skb[3][i] = c_skb[3][i];
563 s_skb[4][i] = c_skb[4][i];
564 s_skb[5][i] = c_skb[5][i];
565 s_skb[6][i] = c_skb[6][i];
566 s_skb[7][i] = c_skb[7][i];
569 barrier (CLK_LOCAL_MEM_FENCE);
571 if (gid >= gid_max) return;
580 pw_buf0[0] = pws[gid].i[ 0];
581 pw_buf0[1] = pws[gid].i[ 1];
589 const u32 pw_len = pws[gid].pw_len;
597 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
598 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
604 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
611 apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
619 transform_racf_key (w0[0], w0[1], key);
621 const u32x c = key[0];
622 const u32x d = key[1];
627 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
631 data[0] = salt_buf0[0];
632 data[1] = salt_buf0[1];
636 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
640 COMPARE_M_SIMD (iv[0], iv[1], z, z);
644 __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_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)
648 __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_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)
652 __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_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)
658 const u32 gid = get_global_id (0);
659 const u32 lid = get_local_id (0);
660 const u32 lsz = get_local_size (0);
666 __local u32 s_SPtrans[8][64];
667 __local u32 s_skb[8][64];
669 for (u32 i = lid; i < 64; i += lsz)
671 s_SPtrans[0][i] = c_SPtrans[0][i];
672 s_SPtrans[1][i] = c_SPtrans[1][i];
673 s_SPtrans[2][i] = c_SPtrans[2][i];
674 s_SPtrans[3][i] = c_SPtrans[3][i];
675 s_SPtrans[4][i] = c_SPtrans[4][i];
676 s_SPtrans[5][i] = c_SPtrans[5][i];
677 s_SPtrans[6][i] = c_SPtrans[6][i];
678 s_SPtrans[7][i] = c_SPtrans[7][i];
680 s_skb[0][i] = c_skb[0][i];
681 s_skb[1][i] = c_skb[1][i];
682 s_skb[2][i] = c_skb[2][i];
683 s_skb[3][i] = c_skb[3][i];
684 s_skb[4][i] = c_skb[4][i];
685 s_skb[5][i] = c_skb[5][i];
686 s_skb[6][i] = c_skb[6][i];
687 s_skb[7][i] = c_skb[7][i];
690 barrier (CLK_LOCAL_MEM_FENCE);
692 if (gid >= gid_max) return;
701 pw_buf0[0] = pws[gid].i[ 0];
702 pw_buf0[1] = pws[gid].i[ 1];
710 const u32 pw_len = pws[gid].pw_len;
718 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
719 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
725 const u32 search[4] =
727 digests_buf[digests_offset].digest_buf[DGST_R0],
728 digests_buf[digests_offset].digest_buf[DGST_R1],
737 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
744 apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
752 transform_racf_key (w0[0], w0[1], key);
754 const u32x c = key[0];
755 const u32x d = key[1];
760 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
764 data[0] = salt_buf0[0];
765 data[1] = salt_buf0[1];
769 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
773 COMPARE_S_SIMD (iv[0], iv[1], z, z);
777 __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_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)
781 __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_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)