2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
4 * magnum <john.magnum@hushmail.com>
13 #include "inc_vendor.cl"
14 #include "inc_hash_constants.h"
15 #include "inc_hash_functions.cl"
16 #include "inc_types.cl"
17 #include "inc_common.cl"
18 #include "inc_simd.cl"
20 #define PERM_OP(a,b,tt,n,m) \
30 #define HPERM_OP(a,tt,n,m) \
36 tt = tt >> (16 + n); \
42 PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \
43 PERM_OP (l, r, tt, 16, 0x0000ffff); \
44 PERM_OP (r, l, tt, 2, 0x33333333); \
45 PERM_OP (l, r, tt, 8, 0x00ff00ff); \
46 PERM_OP (r, l, tt, 1, 0x55555555); \
51 PERM_OP (l, r, tt, 1, 0x55555555); \
52 PERM_OP (r, l, tt, 8, 0x00ff00ff); \
53 PERM_OP (l, r, tt, 2, 0x33333333); \
54 PERM_OP (r, l, tt, 16, 0x0000ffff); \
55 PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
58 __constant u8 ascii_to_ebcdic_pc[256] =
60 // little hack, can't crack 0-bytes in password, but who cares
61 // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
62 0x2a, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
63 0x8a, 0x89, 0x8f, 0x8c, 0xd3, 0xd0, 0xce, 0xe6, 0x9b, 0x98, 0xd5, 0xe5, 0x92, 0x91, 0x97, 0x94,
64 0x2a, 0x34, 0x54, 0x5d, 0x1c, 0x73, 0x0b, 0x51, 0x31, 0x10, 0x13, 0x37, 0x7c, 0x6b, 0x3d, 0x68,
65 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40, 0x46, 0x45, 0x5b, 0x58, 0x5e, 0x16, 0x32, 0x57, 0x76, 0x75,
66 0x52, 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07,
67 0x04, 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x3e, 0x6b, 0x1f, 0x15, 0x70,
68 0x58, 0xa8, 0xae, 0xad, 0xa2, 0xa1, 0xa7, 0xa4, 0xba, 0xb9, 0x89, 0x8f, 0x8c, 0x83, 0x80, 0x86,
69 0x85, 0x9b, 0x98, 0xef, 0xec, 0xe3, 0xe0, 0xe6, 0xe5, 0xfb, 0xf8, 0x2a, 0x7f, 0x0b, 0xe9, 0xa4,
70 0xea, 0xe9, 0xef, 0xec, 0xe3, 0x80, 0xa7, 0x85, 0xfb, 0xf8, 0xfe, 0xfd, 0xf2, 0xb9, 0xbf, 0x9d,
71 0xcb, 0xc8, 0x9e, 0xcd, 0xc2, 0xc1, 0xc7, 0xba, 0xda, 0xd9, 0xdf, 0xdc, 0xa2, 0x83, 0xd6, 0x68,
72 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07, 0x04,
73 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40,
74 0x46, 0x45, 0x5b, 0xab, 0xbf, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, 0x8a, 0x9e, 0x9d, 0x92, 0x91, 0x97,
75 0x94, 0xea, 0xfe, 0xfd, 0xf2, 0xf1, 0xf7, 0xf4, 0xcb, 0xc8, 0xce, 0xcd, 0xc2, 0xc1, 0xc7, 0xc4,
76 0xda, 0xd9, 0xdf, 0xdc, 0xd3, 0xd0, 0xd6, 0xd5, 0x3e, 0x3d, 0x32, 0x31, 0x37, 0x34, 0x1f, 0x1c,
77 0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
80 __constant u32 c_SPtrans[8][64] =
83 0x02080800, 0x00080000, 0x02000002, 0x02080802,
84 0x02000000, 0x00080802, 0x00080002, 0x02000002,
85 0x00080802, 0x02080800, 0x02080000, 0x00000802,
86 0x02000802, 0x02000000, 0x00000000, 0x00080002,
87 0x00080000, 0x00000002, 0x02000800, 0x00080800,
88 0x02080802, 0x02080000, 0x00000802, 0x02000800,
89 0x00000002, 0x00000800, 0x00080800, 0x02080002,
90 0x00000800, 0x02000802, 0x02080002, 0x00000000,
91 0x00000000, 0x02080802, 0x02000800, 0x00080002,
92 0x02080800, 0x00080000, 0x00000802, 0x02000800,
93 0x02080002, 0x00000800, 0x00080800, 0x02000002,
94 0x00080802, 0x00000002, 0x02000002, 0x02080000,
95 0x02080802, 0x00080800, 0x02080000, 0x02000802,
96 0x02000000, 0x00000802, 0x00080002, 0x00000000,
97 0x00080000, 0x02000000, 0x02000802, 0x02080800,
98 0x00000002, 0x02080002, 0x00000800, 0x00080802,
101 0x40108010, 0x00000000, 0x00108000, 0x40100000,
102 0x40000010, 0x00008010, 0x40008000, 0x00108000,
103 0x00008000, 0x40100010, 0x00000010, 0x40008000,
104 0x00100010, 0x40108000, 0x40100000, 0x00000010,
105 0x00100000, 0x40008010, 0x40100010, 0x00008000,
106 0x00108010, 0x40000000, 0x00000000, 0x00100010,
107 0x40008010, 0x00108010, 0x40108000, 0x40000010,
108 0x40000000, 0x00100000, 0x00008010, 0x40108010,
109 0x00100010, 0x40108000, 0x40008000, 0x00108010,
110 0x40108010, 0x00100010, 0x40000010, 0x00000000,
111 0x40000000, 0x00008010, 0x00100000, 0x40100010,
112 0x00008000, 0x40000000, 0x00108010, 0x40008010,
113 0x40108000, 0x00008000, 0x00000000, 0x40000010,
114 0x00000010, 0x40108010, 0x00108000, 0x40100000,
115 0x40100010, 0x00100000, 0x00008010, 0x40008000,
116 0x40008010, 0x00000010, 0x40100000, 0x00108000,
119 0x04000001, 0x04040100, 0x00000100, 0x04000101,
120 0x00040001, 0x04000000, 0x04000101, 0x00040100,
121 0x04000100, 0x00040000, 0x04040000, 0x00000001,
122 0x04040101, 0x00000101, 0x00000001, 0x04040001,
123 0x00000000, 0x00040001, 0x04040100, 0x00000100,
124 0x00000101, 0x04040101, 0x00040000, 0x04000001,
125 0x04040001, 0x04000100, 0x00040101, 0x04040000,
126 0x00040100, 0x00000000, 0x04000000, 0x00040101,
127 0x04040100, 0x00000100, 0x00000001, 0x00040000,
128 0x00000101, 0x00040001, 0x04040000, 0x04000101,
129 0x00000000, 0x04040100, 0x00040100, 0x04040001,
130 0x00040001, 0x04000000, 0x04040101, 0x00000001,
131 0x00040101, 0x04000001, 0x04000000, 0x04040101,
132 0x00040000, 0x04000100, 0x04000101, 0x00040100,
133 0x04000100, 0x00000000, 0x04040001, 0x00000101,
134 0x04000001, 0x00040101, 0x00000100, 0x04040000,
137 0x00401008, 0x10001000, 0x00000008, 0x10401008,
138 0x00000000, 0x10400000, 0x10001008, 0x00400008,
139 0x10401000, 0x10000008, 0x10000000, 0x00001008,
140 0x10000008, 0x00401008, 0x00400000, 0x10000000,
141 0x10400008, 0x00401000, 0x00001000, 0x00000008,
142 0x00401000, 0x10001008, 0x10400000, 0x00001000,
143 0x00001008, 0x00000000, 0x00400008, 0x10401000,
144 0x10001000, 0x10400008, 0x10401008, 0x00400000,
145 0x10400008, 0x00001008, 0x00400000, 0x10000008,
146 0x00401000, 0x10001000, 0x00000008, 0x10400000,
147 0x10001008, 0x00000000, 0x00001000, 0x00400008,
148 0x00000000, 0x10400008, 0x10401000, 0x00001000,
149 0x10000000, 0x10401008, 0x00401008, 0x00400000,
150 0x10401008, 0x00000008, 0x10001000, 0x00401008,
151 0x00400008, 0x00401000, 0x10400000, 0x10001008,
152 0x00001008, 0x10000000, 0x10000008, 0x10401000,
155 0x08000000, 0x00010000, 0x00000400, 0x08010420,
156 0x08010020, 0x08000400, 0x00010420, 0x08010000,
157 0x00010000, 0x00000020, 0x08000020, 0x00010400,
158 0x08000420, 0x08010020, 0x08010400, 0x00000000,
159 0x00010400, 0x08000000, 0x00010020, 0x00000420,
160 0x08000400, 0x00010420, 0x00000000, 0x08000020,
161 0x00000020, 0x08000420, 0x08010420, 0x00010020,
162 0x08010000, 0x00000400, 0x00000420, 0x08010400,
163 0x08010400, 0x08000420, 0x00010020, 0x08010000,
164 0x00010000, 0x00000020, 0x08000020, 0x08000400,
165 0x08000000, 0x00010400, 0x08010420, 0x00000000,
166 0x00010420, 0x08000000, 0x00000400, 0x00010020,
167 0x08000420, 0x00000400, 0x00000000, 0x08010420,
168 0x08010020, 0x08010400, 0x00000420, 0x00010000,
169 0x00010400, 0x08010020, 0x08000400, 0x00000420,
170 0x00000020, 0x00010420, 0x08010000, 0x08000020,
173 0x80000040, 0x00200040, 0x00000000, 0x80202000,
174 0x00200040, 0x00002000, 0x80002040, 0x00200000,
175 0x00002040, 0x80202040, 0x00202000, 0x80000000,
176 0x80002000, 0x80000040, 0x80200000, 0x00202040,
177 0x00200000, 0x80002040, 0x80200040, 0x00000000,
178 0x00002000, 0x00000040, 0x80202000, 0x80200040,
179 0x80202040, 0x80200000, 0x80000000, 0x00002040,
180 0x00000040, 0x00202000, 0x00202040, 0x80002000,
181 0x00002040, 0x80000000, 0x80002000, 0x00202040,
182 0x80202000, 0x00200040, 0x00000000, 0x80002000,
183 0x80000000, 0x00002000, 0x80200040, 0x00200000,
184 0x00200040, 0x80202040, 0x00202000, 0x00000040,
185 0x80202040, 0x00202000, 0x00200000, 0x80002040,
186 0x80000040, 0x80200000, 0x00202040, 0x00000000,
187 0x00002000, 0x80000040, 0x80002040, 0x80202000,
188 0x80200000, 0x00002040, 0x00000040, 0x80200040,
191 0x00004000, 0x00000200, 0x01000200, 0x01000004,
192 0x01004204, 0x00004004, 0x00004200, 0x00000000,
193 0x01000000, 0x01000204, 0x00000204, 0x01004000,
194 0x00000004, 0x01004200, 0x01004000, 0x00000204,
195 0x01000204, 0x00004000, 0x00004004, 0x01004204,
196 0x00000000, 0x01000200, 0x01000004, 0x00004200,
197 0x01004004, 0x00004204, 0x01004200, 0x00000004,
198 0x00004204, 0x01004004, 0x00000200, 0x01000000,
199 0x00004204, 0x01004000, 0x01004004, 0x00000204,
200 0x00004000, 0x00000200, 0x01000000, 0x01004004,
201 0x01000204, 0x00004204, 0x00004200, 0x00000000,
202 0x00000200, 0x01000004, 0x00000004, 0x01000200,
203 0x00000000, 0x01000204, 0x01000200, 0x00004200,
204 0x00000204, 0x00004000, 0x01004204, 0x01000000,
205 0x01004200, 0x00000004, 0x00004004, 0x01004204,
206 0x01000004, 0x01004200, 0x01004000, 0x00004004,
209 0x20800080, 0x20820000, 0x00020080, 0x00000000,
210 0x20020000, 0x00800080, 0x20800000, 0x20820080,
211 0x00000080, 0x20000000, 0x00820000, 0x00020080,
212 0x00820080, 0x20020080, 0x20000080, 0x20800000,
213 0x00020000, 0x00820080, 0x00800080, 0x20020000,
214 0x20820080, 0x20000080, 0x00000000, 0x00820000,
215 0x20000000, 0x00800000, 0x20020080, 0x20800080,
216 0x00800000, 0x00020000, 0x20820000, 0x00000080,
217 0x00800000, 0x00020000, 0x20000080, 0x20820080,
218 0x00020080, 0x20000000, 0x00000000, 0x00820000,
219 0x20800080, 0x20020080, 0x20020000, 0x00800080,
220 0x20820000, 0x00000080, 0x00800080, 0x20020000,
221 0x20820080, 0x00800000, 0x20800000, 0x20000080,
222 0x00820000, 0x00020080, 0x20020080, 0x20800000,
223 0x00000080, 0x20820000, 0x00820080, 0x00000000,
224 0x20000000, 0x20800080, 0x00020000, 0x00820080,
228 __constant u32 c_skb[8][64] =
231 0x00000000, 0x00000010, 0x20000000, 0x20000010,
232 0x00010000, 0x00010010, 0x20010000, 0x20010010,
233 0x00000800, 0x00000810, 0x20000800, 0x20000810,
234 0x00010800, 0x00010810, 0x20010800, 0x20010810,
235 0x00000020, 0x00000030, 0x20000020, 0x20000030,
236 0x00010020, 0x00010030, 0x20010020, 0x20010030,
237 0x00000820, 0x00000830, 0x20000820, 0x20000830,
238 0x00010820, 0x00010830, 0x20010820, 0x20010830,
239 0x00080000, 0x00080010, 0x20080000, 0x20080010,
240 0x00090000, 0x00090010, 0x20090000, 0x20090010,
241 0x00080800, 0x00080810, 0x20080800, 0x20080810,
242 0x00090800, 0x00090810, 0x20090800, 0x20090810,
243 0x00080020, 0x00080030, 0x20080020, 0x20080030,
244 0x00090020, 0x00090030, 0x20090020, 0x20090030,
245 0x00080820, 0x00080830, 0x20080820, 0x20080830,
246 0x00090820, 0x00090830, 0x20090820, 0x20090830,
249 0x00000000, 0x02000000, 0x00002000, 0x02002000,
250 0x00200000, 0x02200000, 0x00202000, 0x02202000,
251 0x00000004, 0x02000004, 0x00002004, 0x02002004,
252 0x00200004, 0x02200004, 0x00202004, 0x02202004,
253 0x00000400, 0x02000400, 0x00002400, 0x02002400,
254 0x00200400, 0x02200400, 0x00202400, 0x02202400,
255 0x00000404, 0x02000404, 0x00002404, 0x02002404,
256 0x00200404, 0x02200404, 0x00202404, 0x02202404,
257 0x10000000, 0x12000000, 0x10002000, 0x12002000,
258 0x10200000, 0x12200000, 0x10202000, 0x12202000,
259 0x10000004, 0x12000004, 0x10002004, 0x12002004,
260 0x10200004, 0x12200004, 0x10202004, 0x12202004,
261 0x10000400, 0x12000400, 0x10002400, 0x12002400,
262 0x10200400, 0x12200400, 0x10202400, 0x12202400,
263 0x10000404, 0x12000404, 0x10002404, 0x12002404,
264 0x10200404, 0x12200404, 0x10202404, 0x12202404,
267 0x00000000, 0x00000001, 0x00040000, 0x00040001,
268 0x01000000, 0x01000001, 0x01040000, 0x01040001,
269 0x00000002, 0x00000003, 0x00040002, 0x00040003,
270 0x01000002, 0x01000003, 0x01040002, 0x01040003,
271 0x00000200, 0x00000201, 0x00040200, 0x00040201,
272 0x01000200, 0x01000201, 0x01040200, 0x01040201,
273 0x00000202, 0x00000203, 0x00040202, 0x00040203,
274 0x01000202, 0x01000203, 0x01040202, 0x01040203,
275 0x08000000, 0x08000001, 0x08040000, 0x08040001,
276 0x09000000, 0x09000001, 0x09040000, 0x09040001,
277 0x08000002, 0x08000003, 0x08040002, 0x08040003,
278 0x09000002, 0x09000003, 0x09040002, 0x09040003,
279 0x08000200, 0x08000201, 0x08040200, 0x08040201,
280 0x09000200, 0x09000201, 0x09040200, 0x09040201,
281 0x08000202, 0x08000203, 0x08040202, 0x08040203,
282 0x09000202, 0x09000203, 0x09040202, 0x09040203,
285 0x00000000, 0x00100000, 0x00000100, 0x00100100,
286 0x00000008, 0x00100008, 0x00000108, 0x00100108,
287 0x00001000, 0x00101000, 0x00001100, 0x00101100,
288 0x00001008, 0x00101008, 0x00001108, 0x00101108,
289 0x04000000, 0x04100000, 0x04000100, 0x04100100,
290 0x04000008, 0x04100008, 0x04000108, 0x04100108,
291 0x04001000, 0x04101000, 0x04001100, 0x04101100,
292 0x04001008, 0x04101008, 0x04001108, 0x04101108,
293 0x00020000, 0x00120000, 0x00020100, 0x00120100,
294 0x00020008, 0x00120008, 0x00020108, 0x00120108,
295 0x00021000, 0x00121000, 0x00021100, 0x00121100,
296 0x00021008, 0x00121008, 0x00021108, 0x00121108,
297 0x04020000, 0x04120000, 0x04020100, 0x04120100,
298 0x04020008, 0x04120008, 0x04020108, 0x04120108,
299 0x04021000, 0x04121000, 0x04021100, 0x04121100,
300 0x04021008, 0x04121008, 0x04021108, 0x04121108,
303 0x00000000, 0x10000000, 0x00010000, 0x10010000,
304 0x00000004, 0x10000004, 0x00010004, 0x10010004,
305 0x20000000, 0x30000000, 0x20010000, 0x30010000,
306 0x20000004, 0x30000004, 0x20010004, 0x30010004,
307 0x00100000, 0x10100000, 0x00110000, 0x10110000,
308 0x00100004, 0x10100004, 0x00110004, 0x10110004,
309 0x20100000, 0x30100000, 0x20110000, 0x30110000,
310 0x20100004, 0x30100004, 0x20110004, 0x30110004,
311 0x00001000, 0x10001000, 0x00011000, 0x10011000,
312 0x00001004, 0x10001004, 0x00011004, 0x10011004,
313 0x20001000, 0x30001000, 0x20011000, 0x30011000,
314 0x20001004, 0x30001004, 0x20011004, 0x30011004,
315 0x00101000, 0x10101000, 0x00111000, 0x10111000,
316 0x00101004, 0x10101004, 0x00111004, 0x10111004,
317 0x20101000, 0x30101000, 0x20111000, 0x30111000,
318 0x20101004, 0x30101004, 0x20111004, 0x30111004,
321 0x00000000, 0x08000000, 0x00000008, 0x08000008,
322 0x00000400, 0x08000400, 0x00000408, 0x08000408,
323 0x00020000, 0x08020000, 0x00020008, 0x08020008,
324 0x00020400, 0x08020400, 0x00020408, 0x08020408,
325 0x00000001, 0x08000001, 0x00000009, 0x08000009,
326 0x00000401, 0x08000401, 0x00000409, 0x08000409,
327 0x00020001, 0x08020001, 0x00020009, 0x08020009,
328 0x00020401, 0x08020401, 0x00020409, 0x08020409,
329 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
330 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
331 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
332 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
333 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
334 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
335 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
336 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
339 0x00000000, 0x00000100, 0x00080000, 0x00080100,
340 0x01000000, 0x01000100, 0x01080000, 0x01080100,
341 0x00000010, 0x00000110, 0x00080010, 0x00080110,
342 0x01000010, 0x01000110, 0x01080010, 0x01080110,
343 0x00200000, 0x00200100, 0x00280000, 0x00280100,
344 0x01200000, 0x01200100, 0x01280000, 0x01280100,
345 0x00200010, 0x00200110, 0x00280010, 0x00280110,
346 0x01200010, 0x01200110, 0x01280010, 0x01280110,
347 0x00000200, 0x00000300, 0x00080200, 0x00080300,
348 0x01000200, 0x01000300, 0x01080200, 0x01080300,
349 0x00000210, 0x00000310, 0x00080210, 0x00080310,
350 0x01000210, 0x01000310, 0x01080210, 0x01080310,
351 0x00200200, 0x00200300, 0x00280200, 0x00280300,
352 0x01200200, 0x01200300, 0x01280200, 0x01280300,
353 0x00200210, 0x00200310, 0x00280210, 0x00280310,
354 0x01200210, 0x01200310, 0x01280210, 0x01280310,
357 0x00000000, 0x04000000, 0x00040000, 0x04040000,
358 0x00000002, 0x04000002, 0x00040002, 0x04040002,
359 0x00002000, 0x04002000, 0x00042000, 0x04042000,
360 0x00002002, 0x04002002, 0x00042002, 0x04042002,
361 0x00000020, 0x04000020, 0x00040020, 0x04040020,
362 0x00000022, 0x04000022, 0x00040022, 0x04040022,
363 0x00002020, 0x04002020, 0x00042020, 0x04042020,
364 0x00002022, 0x04002022, 0x00042022, 0x04042022,
365 0x00000800, 0x04000800, 0x00040800, 0x04040800,
366 0x00000802, 0x04000802, 0x00040802, 0x04040802,
367 0x00002800, 0x04002800, 0x00042800, 0x04042800,
368 0x00002802, 0x04002802, 0x00042802, 0x04042802,
369 0x00000820, 0x04000820, 0x00040820, 0x04040820,
370 0x00000822, 0x04000822, 0x00040822, 0x04040822,
371 0x00002820, 0x04002820, 0x00042820, 0x04042820,
372 0x00002822, 0x04002822, 0x00042822, 0x04042822
377 #define BOX(i,n,S) (S)[(n)][(i)]
379 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
381 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
383 #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])
384 #elif VECT_SIZE == 16
385 #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])
389 #define BOX1(i,S) (S)[(i)]
391 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
393 #define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
395 #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])
396 #elif VECT_SIZE == 16
397 #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])
400 void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
408 for (u32 i = 0; i < 16; i += 2)
414 t = Kd[i + 0] ^ rotl32 (r, 28u);
416 l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
417 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
418 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
419 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
420 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
421 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
422 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
423 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
426 t = Kd[i + 1] ^ rotl32 (l, 28u);
428 r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
429 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
430 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
431 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
432 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
433 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
434 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
435 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
442 void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
446 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
447 HPERM_OP (c, tt, 2, 0xcccc0000);
448 HPERM_OP (d, tt, 2, 0xcccc0000);
449 PERM_OP (d, c, tt, 1, 0x55555555);
450 PERM_OP (c, d, tt, 8, 0x00ff00ff);
451 PERM_OP (d, c, tt, 1, 0x55555555);
453 d = ((d & 0x000000ff) << 16)
454 | ((d & 0x0000ff00) << 0)
455 | ((d & 0x00ff0000) >> 16)
456 | ((c & 0xf0000000) >> 4);
463 for (u32 i = 0; i < 16; i++)
465 if ((i < 2) || (i == 8) || (i == 15))
467 c = ((c >> 1) | (c << 27));
468 d = ((d >> 1) | (d << 27));
472 c = ((c >> 2) | (c << 26));
473 d = ((d >> 2) | (d << 26));
479 const u32x c00 = (c >> 0) & 0x0000003f;
480 const u32x c06 = (c >> 6) & 0x00383003;
481 const u32x c07 = (c >> 7) & 0x0000003c;
482 const u32x c13 = (c >> 13) & 0x0000060f;
483 const u32x c20 = (c >> 20) & 0x00000001;
485 u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
486 | BOX (((c06 >> 0) & 0xff)
487 |((c07 >> 0) & 0xff), 1, s_skb)
488 | BOX (((c13 >> 0) & 0xff)
489 |((c06 >> 8) & 0xff), 2, s_skb)
490 | BOX (((c20 >> 0) & 0xff)
492 |((c06 >> 16) & 0xff), 3, s_skb);
494 const u32x d00 = (d >> 0) & 0x00003c3f;
495 const u32x d07 = (d >> 7) & 0x00003f03;
496 const u32x d21 = (d >> 21) & 0x0000000f;
497 const u32x d22 = (d >> 22) & 0x00000030;
499 u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
500 | BOX (((d07 >> 0) & 0xff)
501 |((d00 >> 8) & 0xff), 5, s_skb)
502 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
503 | BOX (((d21 >> 0) & 0xff)
504 |((d22 >> 0) & 0xff), 7, s_skb);
506 Kc[i] = ((t << 16) | (s & 0x0000ffff));
507 Kd[i] = ((s >> 16) | (t & 0xffff0000));
509 Kc[i] = rotl32 (Kc[i], 2u);
510 Kd[i] = rotl32 (Kd[i], 2u);
514 void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
516 key[0] = BOX1 (((w0 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0
517 | BOX1 (((w0 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8
518 | BOX1 (((w0 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16
519 | BOX1 (((w0 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24;
521 key[1] = BOX1 (((w1 >> 0) & 0xff), ascii_to_ebcdic_pc) << 0
522 | BOX1 (((w1 >> 8) & 0xff), ascii_to_ebcdic_pc) << 8
523 | BOX1 (((w1 >> 16) & 0xff), ascii_to_ebcdic_pc) << 16
524 | BOX1 (((w1 >> 24) & 0xff), ascii_to_ebcdic_pc) << 24;
527 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_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)
533 const u32 gid = get_global_id (0);
534 const u32 lid = get_local_id (0);
542 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
543 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
553 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
555 const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
557 const u32x w0 = w0l | w0r;
565 transform_racf_key (w0, w1, key);
567 const u32x c = key[0];
568 const u32x d = key[1];
573 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
577 data[0] = salt_buf0[0];
578 data[1] = salt_buf0[1];
582 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
586 COMPARE_M_SIMD (iv[0], iv[1], z, z);
590 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_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)
596 const u32 gid = get_global_id (0);
597 const u32 lid = get_local_id (0);
605 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
606 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
612 const u32 search[4] =
614 digests_buf[digests_offset].digest_buf[DGST_R0],
615 digests_buf[digests_offset].digest_buf[DGST_R1],
628 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
630 const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
632 const u32x w0 = w0l | w0r;
640 transform_racf_key (w0, w1, key);
642 const u32x c = key[0];
643 const u32x d = key[1];
648 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
652 data[0] = salt_buf0[0];
653 data[1] = salt_buf0[1];
657 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
661 COMPARE_S_SIMD (iv[0], iv[1], z, z);
665 __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_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)
671 const u32 gid = get_global_id (0);
672 const u32 lid = get_local_id (0);
673 const u32 lsz = get_local_size (0);
679 __local u32 s_SPtrans[8][64];
680 __local u32 s_skb[8][64];
682 for (u32 i = lid; i < 64; i += lsz)
684 s_SPtrans[0][i] = c_SPtrans[0][i];
685 s_SPtrans[1][i] = c_SPtrans[1][i];
686 s_SPtrans[2][i] = c_SPtrans[2][i];
687 s_SPtrans[3][i] = c_SPtrans[3][i];
688 s_SPtrans[4][i] = c_SPtrans[4][i];
689 s_SPtrans[5][i] = c_SPtrans[5][i];
690 s_SPtrans[6][i] = c_SPtrans[6][i];
691 s_SPtrans[7][i] = c_SPtrans[7][i];
693 s_skb[0][i] = c_skb[0][i];
694 s_skb[1][i] = c_skb[1][i];
695 s_skb[2][i] = c_skb[2][i];
696 s_skb[3][i] = c_skb[3][i];
697 s_skb[4][i] = c_skb[4][i];
698 s_skb[5][i] = c_skb[5][i];
699 s_skb[6][i] = c_skb[6][i];
700 s_skb[7][i] = c_skb[7][i];
703 barrier (CLK_LOCAL_MEM_FENCE);
705 if (gid >= gid_max) return;
713 w[ 0] = pws[gid].i[ 0];
714 w[ 1] = pws[gid].i[ 1];
730 const u32 pw_len = pws[gid].pw_len;
736 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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
739 __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_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)
743 __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_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)
747 __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_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)
753 const u32 gid = get_global_id (0);
754 const u32 lid = get_local_id (0);
755 const u32 lsz = get_local_size (0);
761 __local u32 s_SPtrans[8][64];
762 __local u32 s_skb[8][64];
764 for (u32 i = lid; i < 64; i += lsz)
766 s_SPtrans[0][i] = c_SPtrans[0][i];
767 s_SPtrans[1][i] = c_SPtrans[1][i];
768 s_SPtrans[2][i] = c_SPtrans[2][i];
769 s_SPtrans[3][i] = c_SPtrans[3][i];
770 s_SPtrans[4][i] = c_SPtrans[4][i];
771 s_SPtrans[5][i] = c_SPtrans[5][i];
772 s_SPtrans[6][i] = c_SPtrans[6][i];
773 s_SPtrans[7][i] = c_SPtrans[7][i];
775 s_skb[0][i] = c_skb[0][i];
776 s_skb[1][i] = c_skb[1][i];
777 s_skb[2][i] = c_skb[2][i];
778 s_skb[3][i] = c_skb[3][i];
779 s_skb[4][i] = c_skb[4][i];
780 s_skb[5][i] = c_skb[5][i];
781 s_skb[6][i] = c_skb[6][i];
782 s_skb[7][i] = c_skb[7][i];
785 barrier (CLK_LOCAL_MEM_FENCE);
787 if (gid >= gid_max) return;
795 w[ 0] = pws[gid].i[ 0];
796 w[ 1] = pws[gid].i[ 1];
812 const u32 pw_len = pws[gid].pw_len;
818 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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
821 __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_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)
825 __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_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)