2 * Author......: Jens Steube <jens.steube@gmail.com>
9 #include "include/constants.h"
10 #include "include/kernel_vendor.h"
29 #include "include/kernel_functions.c"
30 #include "types_amd.c"
31 #include "common_amd.c"
34 #define VECT_COMPARE_S "check_single_vect1_comp4_warp.c"
35 #define VECT_COMPARE_M "check_multi_vect1_comp4_warp.c"
39 #define VECT_COMPARE_S "check_single_vect2_comp4_warp.c"
40 #define VECT_COMPARE_M "check_multi_vect2_comp4_warp.c"
44 #define VECT_COMPARE_S "check_single_vect4_comp4_warp.c"
45 #define VECT_COMPARE_M "check_multi_vect4_comp4_warp.c"
48 #define PERM_OP(a,b,tt,n,m) \
58 #define HPERM_OP(a,tt,n,m) \
64 tt = tt >> (16 + n); \
70 PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \
71 PERM_OP (l, r, tt, 16, 0x0000ffff); \
72 PERM_OP (r, l, tt, 2, 0x33333333); \
73 PERM_OP (l, r, tt, 8, 0x00ff00ff); \
74 PERM_OP (r, l, tt, 1, 0x55555555); \
79 PERM_OP (l, r, tt, 1, 0x55555555); \
80 PERM_OP (r, l, tt, 8, 0x00ff00ff); \
81 PERM_OP (l, r, tt, 2, 0x33333333); \
82 PERM_OP (r, l, tt, 16, 0x0000ffff); \
83 PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
86 __constant u8 ascii_to_ebcdic_pc[256] =
88 // little hack, can't crack 0-bytes in password, but who cares
89 // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
90 0x2a, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
91 0x8a, 0x89, 0x8f, 0x8c, 0xd3, 0xd0, 0xce, 0xe6, 0x9b, 0x98, 0xd5, 0xe5, 0x92, 0x91, 0x97, 0x94,
92 0x2a, 0x34, 0x54, 0x5d, 0x1c, 0x73, 0x0b, 0x51, 0x31, 0x10, 0x13, 0x37, 0x7c, 0x6b, 0x3d, 0x68,
93 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40, 0x46, 0x45, 0x5b, 0x58, 0x5e, 0x16, 0x32, 0x57, 0x76, 0x75,
94 0x52, 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07,
95 0x04, 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x3e, 0x6b, 0x1f, 0x15, 0x70,
96 0x58, 0xa8, 0xae, 0xad, 0xa2, 0xa1, 0xa7, 0xa4, 0xba, 0xb9, 0x89, 0x8f, 0x8c, 0x83, 0x80, 0x86,
97 0x85, 0x9b, 0x98, 0xef, 0xec, 0xe3, 0xe0, 0xe6, 0xe5, 0xfb, 0xf8, 0x2a, 0x7f, 0x0b, 0xe9, 0xa4,
98 0xea, 0xe9, 0xef, 0xec, 0xe3, 0x80, 0xa7, 0x85, 0xfb, 0xf8, 0xfe, 0xfd, 0xf2, 0xb9, 0xbf, 0x9d,
99 0xcb, 0xc8, 0x9e, 0xcd, 0xc2, 0xc1, 0xc7, 0xba, 0xda, 0xd9, 0xdf, 0xdc, 0xa2, 0x83, 0xd6, 0x68,
100 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07, 0x04,
101 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40,
102 0x46, 0x45, 0x5b, 0xab, 0xbf, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, 0x8a, 0x9e, 0x9d, 0x92, 0x91, 0x97,
103 0x94, 0xea, 0xfe, 0xfd, 0xf2, 0xf1, 0xf7, 0xf4, 0xcb, 0xc8, 0xce, 0xcd, 0xc2, 0xc1, 0xc7, 0xc4,
104 0xda, 0xd9, 0xdf, 0xdc, 0xd3, 0xd0, 0xd6, 0xd5, 0x3e, 0x3d, 0x32, 0x31, 0x37, 0x34, 0x1f, 0x1c,
105 0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
108 __constant u32 c_SPtrans[8][64] =
111 0x02080800, 0x00080000, 0x02000002, 0x02080802,
112 0x02000000, 0x00080802, 0x00080002, 0x02000002,
113 0x00080802, 0x02080800, 0x02080000, 0x00000802,
114 0x02000802, 0x02000000, 0x00000000, 0x00080002,
115 0x00080000, 0x00000002, 0x02000800, 0x00080800,
116 0x02080802, 0x02080000, 0x00000802, 0x02000800,
117 0x00000002, 0x00000800, 0x00080800, 0x02080002,
118 0x00000800, 0x02000802, 0x02080002, 0x00000000,
119 0x00000000, 0x02080802, 0x02000800, 0x00080002,
120 0x02080800, 0x00080000, 0x00000802, 0x02000800,
121 0x02080002, 0x00000800, 0x00080800, 0x02000002,
122 0x00080802, 0x00000002, 0x02000002, 0x02080000,
123 0x02080802, 0x00080800, 0x02080000, 0x02000802,
124 0x02000000, 0x00000802, 0x00080002, 0x00000000,
125 0x00080000, 0x02000000, 0x02000802, 0x02080800,
126 0x00000002, 0x02080002, 0x00000800, 0x00080802,
128 0x40108010, 0x00000000, 0x00108000, 0x40100000,
129 0x40000010, 0x00008010, 0x40008000, 0x00108000,
130 0x00008000, 0x40100010, 0x00000010, 0x40008000,
131 0x00100010, 0x40108000, 0x40100000, 0x00000010,
132 0x00100000, 0x40008010, 0x40100010, 0x00008000,
133 0x00108010, 0x40000000, 0x00000000, 0x00100010,
134 0x40008010, 0x00108010, 0x40108000, 0x40000010,
135 0x40000000, 0x00100000, 0x00008010, 0x40108010,
136 0x00100010, 0x40108000, 0x40008000, 0x00108010,
137 0x40108010, 0x00100010, 0x40000010, 0x00000000,
138 0x40000000, 0x00008010, 0x00100000, 0x40100010,
139 0x00008000, 0x40000000, 0x00108010, 0x40008010,
140 0x40108000, 0x00008000, 0x00000000, 0x40000010,
141 0x00000010, 0x40108010, 0x00108000, 0x40100000,
142 0x40100010, 0x00100000, 0x00008010, 0x40008000,
143 0x40008010, 0x00000010, 0x40100000, 0x00108000,
145 0x04000001, 0x04040100, 0x00000100, 0x04000101,
146 0x00040001, 0x04000000, 0x04000101, 0x00040100,
147 0x04000100, 0x00040000, 0x04040000, 0x00000001,
148 0x04040101, 0x00000101, 0x00000001, 0x04040001,
149 0x00000000, 0x00040001, 0x04040100, 0x00000100,
150 0x00000101, 0x04040101, 0x00040000, 0x04000001,
151 0x04040001, 0x04000100, 0x00040101, 0x04040000,
152 0x00040100, 0x00000000, 0x04000000, 0x00040101,
153 0x04040100, 0x00000100, 0x00000001, 0x00040000,
154 0x00000101, 0x00040001, 0x04040000, 0x04000101,
155 0x00000000, 0x04040100, 0x00040100, 0x04040001,
156 0x00040001, 0x04000000, 0x04040101, 0x00000001,
157 0x00040101, 0x04000001, 0x04000000, 0x04040101,
158 0x00040000, 0x04000100, 0x04000101, 0x00040100,
159 0x04000100, 0x00000000, 0x04040001, 0x00000101,
160 0x04000001, 0x00040101, 0x00000100, 0x04040000,
162 0x00401008, 0x10001000, 0x00000008, 0x10401008,
163 0x00000000, 0x10400000, 0x10001008, 0x00400008,
164 0x10401000, 0x10000008, 0x10000000, 0x00001008,
165 0x10000008, 0x00401008, 0x00400000, 0x10000000,
166 0x10400008, 0x00401000, 0x00001000, 0x00000008,
167 0x00401000, 0x10001008, 0x10400000, 0x00001000,
168 0x00001008, 0x00000000, 0x00400008, 0x10401000,
169 0x10001000, 0x10400008, 0x10401008, 0x00400000,
170 0x10400008, 0x00001008, 0x00400000, 0x10000008,
171 0x00401000, 0x10001000, 0x00000008, 0x10400000,
172 0x10001008, 0x00000000, 0x00001000, 0x00400008,
173 0x00000000, 0x10400008, 0x10401000, 0x00001000,
174 0x10000000, 0x10401008, 0x00401008, 0x00400000,
175 0x10401008, 0x00000008, 0x10001000, 0x00401008,
176 0x00400008, 0x00401000, 0x10400000, 0x10001008,
177 0x00001008, 0x10000000, 0x10000008, 0x10401000,
179 0x08000000, 0x00010000, 0x00000400, 0x08010420,
180 0x08010020, 0x08000400, 0x00010420, 0x08010000,
181 0x00010000, 0x00000020, 0x08000020, 0x00010400,
182 0x08000420, 0x08010020, 0x08010400, 0x00000000,
183 0x00010400, 0x08000000, 0x00010020, 0x00000420,
184 0x08000400, 0x00010420, 0x00000000, 0x08000020,
185 0x00000020, 0x08000420, 0x08010420, 0x00010020,
186 0x08010000, 0x00000400, 0x00000420, 0x08010400,
187 0x08010400, 0x08000420, 0x00010020, 0x08010000,
188 0x00010000, 0x00000020, 0x08000020, 0x08000400,
189 0x08000000, 0x00010400, 0x08010420, 0x00000000,
190 0x00010420, 0x08000000, 0x00000400, 0x00010020,
191 0x08000420, 0x00000400, 0x00000000, 0x08010420,
192 0x08010020, 0x08010400, 0x00000420, 0x00010000,
193 0x00010400, 0x08010020, 0x08000400, 0x00000420,
194 0x00000020, 0x00010420, 0x08010000, 0x08000020,
196 0x80000040, 0x00200040, 0x00000000, 0x80202000,
197 0x00200040, 0x00002000, 0x80002040, 0x00200000,
198 0x00002040, 0x80202040, 0x00202000, 0x80000000,
199 0x80002000, 0x80000040, 0x80200000, 0x00202040,
200 0x00200000, 0x80002040, 0x80200040, 0x00000000,
201 0x00002000, 0x00000040, 0x80202000, 0x80200040,
202 0x80202040, 0x80200000, 0x80000000, 0x00002040,
203 0x00000040, 0x00202000, 0x00202040, 0x80002000,
204 0x00002040, 0x80000000, 0x80002000, 0x00202040,
205 0x80202000, 0x00200040, 0x00000000, 0x80002000,
206 0x80000000, 0x00002000, 0x80200040, 0x00200000,
207 0x00200040, 0x80202040, 0x00202000, 0x00000040,
208 0x80202040, 0x00202000, 0x00200000, 0x80002040,
209 0x80000040, 0x80200000, 0x00202040, 0x00000000,
210 0x00002000, 0x80000040, 0x80002040, 0x80202000,
211 0x80200000, 0x00002040, 0x00000040, 0x80200040,
213 0x00004000, 0x00000200, 0x01000200, 0x01000004,
214 0x01004204, 0x00004004, 0x00004200, 0x00000000,
215 0x01000000, 0x01000204, 0x00000204, 0x01004000,
216 0x00000004, 0x01004200, 0x01004000, 0x00000204,
217 0x01000204, 0x00004000, 0x00004004, 0x01004204,
218 0x00000000, 0x01000200, 0x01000004, 0x00004200,
219 0x01004004, 0x00004204, 0x01004200, 0x00000004,
220 0x00004204, 0x01004004, 0x00000200, 0x01000000,
221 0x00004204, 0x01004000, 0x01004004, 0x00000204,
222 0x00004000, 0x00000200, 0x01000000, 0x01004004,
223 0x01000204, 0x00004204, 0x00004200, 0x00000000,
224 0x00000200, 0x01000004, 0x00000004, 0x01000200,
225 0x00000000, 0x01000204, 0x01000200, 0x00004200,
226 0x00000204, 0x00004000, 0x01004204, 0x01000000,
227 0x01004200, 0x00000004, 0x00004004, 0x01004204,
228 0x01000004, 0x01004200, 0x01004000, 0x00004004,
230 0x20800080, 0x20820000, 0x00020080, 0x00000000,
231 0x20020000, 0x00800080, 0x20800000, 0x20820080,
232 0x00000080, 0x20000000, 0x00820000, 0x00020080,
233 0x00820080, 0x20020080, 0x20000080, 0x20800000,
234 0x00020000, 0x00820080, 0x00800080, 0x20020000,
235 0x20820080, 0x20000080, 0x00000000, 0x00820000,
236 0x20000000, 0x00800000, 0x20020080, 0x20800080,
237 0x00800000, 0x00020000, 0x20820000, 0x00000080,
238 0x00800000, 0x00020000, 0x20000080, 0x20820080,
239 0x00020080, 0x20000000, 0x00000000, 0x00820000,
240 0x20800080, 0x20020080, 0x20020000, 0x00800080,
241 0x20820000, 0x00000080, 0x00800080, 0x20020000,
242 0x20820080, 0x00800000, 0x20800000, 0x20000080,
243 0x00820000, 0x00020080, 0x20020080, 0x20800000,
244 0x00000080, 0x20820000, 0x00820080, 0x00000000,
245 0x20000000, 0x20800080, 0x00020000, 0x00820080,
248 __constant u32 c_skb[8][64] =
250 /* for C bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
251 0x00000000, 0x00000010, 0x20000000, 0x20000010,
252 0x00010000, 0x00010010, 0x20010000, 0x20010010,
253 0x00000800, 0x00000810, 0x20000800, 0x20000810,
254 0x00010800, 0x00010810, 0x20010800, 0x20010810,
255 0x00000020, 0x00000030, 0x20000020, 0x20000030,
256 0x00010020, 0x00010030, 0x20010020, 0x20010030,
257 0x00000820, 0x00000830, 0x20000820, 0x20000830,
258 0x00010820, 0x00010830, 0x20010820, 0x20010830,
259 0x00080000, 0x00080010, 0x20080000, 0x20080010,
260 0x00090000, 0x00090010, 0x20090000, 0x20090010,
261 0x00080800, 0x00080810, 0x20080800, 0x20080810,
262 0x00090800, 0x00090810, 0x20090800, 0x20090810,
263 0x00080020, 0x00080030, 0x20080020, 0x20080030,
264 0x00090020, 0x00090030, 0x20090020, 0x20090030,
265 0x00080820, 0x00080830, 0x20080820, 0x20080830,
266 0x00090820, 0x00090830, 0x20090820, 0x20090830,
267 /* for C bits (numbered as per FIPS 46) 7 8 10 11 12 13 */
268 0x00000000, 0x02000000, 0x00002000, 0x02002000,
269 0x00200000, 0x02200000, 0x00202000, 0x02202000,
270 0x00000004, 0x02000004, 0x00002004, 0x02002004,
271 0x00200004, 0x02200004, 0x00202004, 0x02202004,
272 0x00000400, 0x02000400, 0x00002400, 0x02002400,
273 0x00200400, 0x02200400, 0x00202400, 0x02202400,
274 0x00000404, 0x02000404, 0x00002404, 0x02002404,
275 0x00200404, 0x02200404, 0x00202404, 0x02202404,
276 0x10000000, 0x12000000, 0x10002000, 0x12002000,
277 0x10200000, 0x12200000, 0x10202000, 0x12202000,
278 0x10000004, 0x12000004, 0x10002004, 0x12002004,
279 0x10200004, 0x12200004, 0x10202004, 0x12202004,
280 0x10000400, 0x12000400, 0x10002400, 0x12002400,
281 0x10200400, 0x12200400, 0x10202400, 0x12202400,
282 0x10000404, 0x12000404, 0x10002404, 0x12002404,
283 0x10200404, 0x12200404, 0x10202404, 0x12202404,
284 /* for C bits (numbered as per FIPS 46) 14 15 16 17 19 20 */
285 0x00000000, 0x00000001, 0x00040000, 0x00040001,
286 0x01000000, 0x01000001, 0x01040000, 0x01040001,
287 0x00000002, 0x00000003, 0x00040002, 0x00040003,
288 0x01000002, 0x01000003, 0x01040002, 0x01040003,
289 0x00000200, 0x00000201, 0x00040200, 0x00040201,
290 0x01000200, 0x01000201, 0x01040200, 0x01040201,
291 0x00000202, 0x00000203, 0x00040202, 0x00040203,
292 0x01000202, 0x01000203, 0x01040202, 0x01040203,
293 0x08000000, 0x08000001, 0x08040000, 0x08040001,
294 0x09000000, 0x09000001, 0x09040000, 0x09040001,
295 0x08000002, 0x08000003, 0x08040002, 0x08040003,
296 0x09000002, 0x09000003, 0x09040002, 0x09040003,
297 0x08000200, 0x08000201, 0x08040200, 0x08040201,
298 0x09000200, 0x09000201, 0x09040200, 0x09040201,
299 0x08000202, 0x08000203, 0x08040202, 0x08040203,
300 0x09000202, 0x09000203, 0x09040202, 0x09040203,
301 /* for C bits (numbered as per FIPS 46) 21 23 24 26 27 28 */
302 0x00000000, 0x00100000, 0x00000100, 0x00100100,
303 0x00000008, 0x00100008, 0x00000108, 0x00100108,
304 0x00001000, 0x00101000, 0x00001100, 0x00101100,
305 0x00001008, 0x00101008, 0x00001108, 0x00101108,
306 0x04000000, 0x04100000, 0x04000100, 0x04100100,
307 0x04000008, 0x04100008, 0x04000108, 0x04100108,
308 0x04001000, 0x04101000, 0x04001100, 0x04101100,
309 0x04001008, 0x04101008, 0x04001108, 0x04101108,
310 0x00020000, 0x00120000, 0x00020100, 0x00120100,
311 0x00020008, 0x00120008, 0x00020108, 0x00120108,
312 0x00021000, 0x00121000, 0x00021100, 0x00121100,
313 0x00021008, 0x00121008, 0x00021108, 0x00121108,
314 0x04020000, 0x04120000, 0x04020100, 0x04120100,
315 0x04020008, 0x04120008, 0x04020108, 0x04120108,
316 0x04021000, 0x04121000, 0x04021100, 0x04121100,
317 0x04021008, 0x04121008, 0x04021108, 0x04121108,
318 /* for D bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
319 0x00000000, 0x10000000, 0x00010000, 0x10010000,
320 0x00000004, 0x10000004, 0x00010004, 0x10010004,
321 0x20000000, 0x30000000, 0x20010000, 0x30010000,
322 0x20000004, 0x30000004, 0x20010004, 0x30010004,
323 0x00100000, 0x10100000, 0x00110000, 0x10110000,
324 0x00100004, 0x10100004, 0x00110004, 0x10110004,
325 0x20100000, 0x30100000, 0x20110000, 0x30110000,
326 0x20100004, 0x30100004, 0x20110004, 0x30110004,
327 0x00001000, 0x10001000, 0x00011000, 0x10011000,
328 0x00001004, 0x10001004, 0x00011004, 0x10011004,
329 0x20001000, 0x30001000, 0x20011000, 0x30011000,
330 0x20001004, 0x30001004, 0x20011004, 0x30011004,
331 0x00101000, 0x10101000, 0x00111000, 0x10111000,
332 0x00101004, 0x10101004, 0x00111004, 0x10111004,
333 0x20101000, 0x30101000, 0x20111000, 0x30111000,
334 0x20101004, 0x30101004, 0x20111004, 0x30111004,
335 /* for D bits (numbered as per FIPS 46) 8 9 11 12 13 14 */
336 0x00000000, 0x08000000, 0x00000008, 0x08000008,
337 0x00000400, 0x08000400, 0x00000408, 0x08000408,
338 0x00020000, 0x08020000, 0x00020008, 0x08020008,
339 0x00020400, 0x08020400, 0x00020408, 0x08020408,
340 0x00000001, 0x08000001, 0x00000009, 0x08000009,
341 0x00000401, 0x08000401, 0x00000409, 0x08000409,
342 0x00020001, 0x08020001, 0x00020009, 0x08020009,
343 0x00020401, 0x08020401, 0x00020409, 0x08020409,
344 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
345 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
346 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
347 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
348 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
349 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
350 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
351 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
352 /* for D bits (numbered as per FIPS 46) 16 17 18 19 20 21 */
353 0x00000000, 0x00000100, 0x00080000, 0x00080100,
354 0x01000000, 0x01000100, 0x01080000, 0x01080100,
355 0x00000010, 0x00000110, 0x00080010, 0x00080110,
356 0x01000010, 0x01000110, 0x01080010, 0x01080110,
357 0x00200000, 0x00200100, 0x00280000, 0x00280100,
358 0x01200000, 0x01200100, 0x01280000, 0x01280100,
359 0x00200010, 0x00200110, 0x00280010, 0x00280110,
360 0x01200010, 0x01200110, 0x01280010, 0x01280110,
361 0x00000200, 0x00000300, 0x00080200, 0x00080300,
362 0x01000200, 0x01000300, 0x01080200, 0x01080300,
363 0x00000210, 0x00000310, 0x00080210, 0x00080310,
364 0x01000210, 0x01000310, 0x01080210, 0x01080310,
365 0x00200200, 0x00200300, 0x00280200, 0x00280300,
366 0x01200200, 0x01200300, 0x01280200, 0x01280300,
367 0x00200210, 0x00200310, 0x00280210, 0x00280310,
368 0x01200210, 0x01200310, 0x01280210, 0x01280310,
369 /* for D bits (numbered as per FIPS 46) 22 23 24 25 27 28 */
370 0x00000000, 0x04000000, 0x00040000, 0x04040000,
371 0x00000002, 0x04000002, 0x00040002, 0x04040002,
372 0x00002000, 0x04002000, 0x00042000, 0x04042000,
373 0x00002002, 0x04002002, 0x00042002, 0x04042002,
374 0x00000020, 0x04000020, 0x00040020, 0x04040020,
375 0x00000022, 0x04000022, 0x00040022, 0x04040022,
376 0x00002020, 0x04002020, 0x00042020, 0x04042020,
377 0x00002022, 0x04002022, 0x00042022, 0x04042022,
378 0x00000800, 0x04000800, 0x00040800, 0x04040800,
379 0x00000802, 0x04000802, 0x00040802, 0x04040802,
380 0x00002800, 0x04002800, 0x00042800, 0x04042800,
381 0x00002802, 0x04002802, 0x00042802, 0x04042802,
382 0x00000820, 0x04000820, 0x00040820, 0x04040820,
383 0x00000822, 0x04000822, 0x00040822, 0x04040822,
384 0x00002820, 0x04002820, 0x00042820, 0x04042820,
385 0x00002822, 0x04002822, 0x00042822, 0x04042822,
388 #define NBOX(i,n,S) (S)[(n)][(i)]
390 static void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 s_SPtrans[8][64])
398 for (u32 i = 0; i < 16; i++)
401 u32x t = Kd[i] ^ rotl32 (r, 28u);
404 l ^= NBOX (((u >> 2) & 0x3f), 0, s_SPtrans)
405 | NBOX (((u >> 10) & 0x3f), 2, s_SPtrans)
406 | NBOX (((u >> 18) & 0x3f), 4, s_SPtrans)
407 | NBOX (((u >> 26) & 0x3f), 6, s_SPtrans)
408 | NBOX (((t >> 2) & 0x3f), 1, s_SPtrans)
409 | NBOX (((t >> 10) & 0x3f), 3, s_SPtrans)
410 | NBOX (((t >> 18) & 0x3f), 5, s_SPtrans)
411 | NBOX (((t >> 26) & 0x3f), 7, s_SPtrans);
415 l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
416 | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
417 | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
418 | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
419 | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
420 | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
421 | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
422 | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
424 l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
425 | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
426 | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
427 | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
428 | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
429 | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
430 | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
431 | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
435 l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
436 | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
437 | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
438 | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
439 | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
440 | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
441 | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
442 | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
444 l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
445 | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
446 | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
447 | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
448 | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
449 | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
450 | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
451 | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
453 l.s2 ^= NBOX (((u.s2 >> 2) & 0x3f), 0, s_SPtrans)
454 | NBOX (((u.s2 >> 10) & 0x3f), 2, s_SPtrans)
455 | NBOX (((u.s2 >> 18) & 0x3f), 4, s_SPtrans)
456 | NBOX (((u.s2 >> 26) & 0x3f), 6, s_SPtrans)
457 | NBOX (((t.s2 >> 2) & 0x3f), 1, s_SPtrans)
458 | NBOX (((t.s2 >> 10) & 0x3f), 3, s_SPtrans)
459 | NBOX (((t.s2 >> 18) & 0x3f), 5, s_SPtrans)
460 | NBOX (((t.s2 >> 26) & 0x3f), 7, s_SPtrans);
462 l.s3 ^= NBOX (((u.s3 >> 2) & 0x3f), 0, s_SPtrans)
463 | NBOX (((u.s3 >> 10) & 0x3f), 2, s_SPtrans)
464 | NBOX (((u.s3 >> 18) & 0x3f), 4, s_SPtrans)
465 | NBOX (((u.s3 >> 26) & 0x3f), 6, s_SPtrans)
466 | NBOX (((t.s3 >> 2) & 0x3f), 1, s_SPtrans)
467 | NBOX (((t.s3 >> 10) & 0x3f), 3, s_SPtrans)
468 | NBOX (((t.s3 >> 18) & 0x3f), 5, s_SPtrans)
469 | NBOX (((t.s3 >> 26) & 0x3f), 7, s_SPtrans);
481 static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 s_skb[8][64])
485 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
486 HPERM_OP (c, tt, 2, 0xcccc0000);
487 HPERM_OP (d, tt, 2, 0xcccc0000);
488 PERM_OP (d, c, tt, 1, 0x55555555);
489 PERM_OP (c, d, tt, 8, 0x00ff00ff);
490 PERM_OP (d, c, tt, 1, 0x55555555);
492 d = ((d & 0x000000ff) << 16)
493 | ((d & 0x0000ff00) << 0)
494 | ((d & 0x00ff0000) >> 16)
495 | ((c & 0xf0000000) >> 4);
500 for (u32 i = 0; i < 16; i++)
502 const u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
503 const u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
505 c = c >> shifts3s0[i] | c << shifts3s1[i];
506 d = d >> shifts3s0[i] | d << shifts3s1[i];
515 s = NBOX ((( c >> 0) & 0x3f), 0, s_skb)
516 | NBOX ((((c >> 6) & 0x03)
517 | ((c >> 7) & 0x3c)), 1, s_skb)
518 | NBOX ((((c >> 13) & 0x0f)
519 | ((c >> 14) & 0x30)), 2, s_skb)
520 | NBOX ((((c >> 20) & 0x01)
522 | ((c >> 22) & 0x38)), 3, s_skb);
524 t = NBOX ((( d >> 0) & 0x3f), 4, s_skb)
525 | NBOX ((((d >> 7) & 0x03)
526 | ((d >> 8) & 0x3c)), 5, s_skb)
527 | NBOX ((((d >> 15) & 0x3f)), 6, s_skb)
528 | NBOX ((((d >> 21) & 0x0f)
529 | ((d >> 22) & 0x30)), 7, s_skb);
533 s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
534 | NBOX ((((c.s0 >> 6) & 0x03)
535 | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
536 | NBOX ((((c.s0 >> 13) & 0x0f)
537 | ((c.s0 >> 14) & 0x30)), 2, s_skb)
538 | NBOX ((((c.s0 >> 20) & 0x01)
539 | ((c.s0 >> 21) & 0x06)
540 | ((c.s0 >> 22) & 0x38)), 3, s_skb);
542 t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
543 | NBOX ((((d.s0 >> 7) & 0x03)
544 | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
545 | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
546 | NBOX ((((d.s0 >> 21) & 0x0f)
547 | ((d.s0 >> 22) & 0x30)), 7, s_skb);
549 s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
550 | NBOX ((((c.s1 >> 6) & 0x03)
551 | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
552 | NBOX ((((c.s1 >> 13) & 0x0f)
553 | ((c.s1 >> 14) & 0x30)), 2, s_skb)
554 | NBOX ((((c.s1 >> 20) & 0x01)
555 | ((c.s1 >> 21) & 0x06)
556 | ((c.s1 >> 22) & 0x38)), 3, s_skb);
558 t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
559 | NBOX ((((d.s1 >> 7) & 0x03)
560 | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
561 | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
562 | NBOX ((((d.s1 >> 21) & 0x0f)
563 | ((d.s1 >> 22) & 0x30)), 7, s_skb);
567 s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
568 | NBOX ((((c.s0 >> 6) & 0x03)
569 | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
570 | NBOX ((((c.s0 >> 13) & 0x0f)
571 | ((c.s0 >> 14) & 0x30)), 2, s_skb)
572 | NBOX ((((c.s0 >> 20) & 0x01)
573 | ((c.s0 >> 21) & 0x06)
574 | ((c.s0 >> 22) & 0x38)), 3, s_skb);
576 t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
577 | NBOX ((((d.s0 >> 7) & 0x03)
578 | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
579 | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
580 | NBOX ((((d.s0 >> 21) & 0x0f)
581 | ((d.s0 >> 22) & 0x30)), 7, s_skb);
583 s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
584 | NBOX ((((c.s1 >> 6) & 0x03)
585 | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
586 | NBOX ((((c.s1 >> 13) & 0x0f)
587 | ((c.s1 >> 14) & 0x30)), 2, s_skb)
588 | NBOX ((((c.s1 >> 20) & 0x01)
589 | ((c.s1 >> 21) & 0x06)
590 | ((c.s1 >> 22) & 0x38)), 3, s_skb);
592 t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
593 | NBOX ((((d.s1 >> 7) & 0x03)
594 | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
595 | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
596 | NBOX ((((d.s1 >> 21) & 0x0f)
597 | ((d.s1 >> 22) & 0x30)), 7, s_skb);
599 s.s2 = NBOX ((( c.s2 >> 0) & 0x3f), 0, s_skb)
600 | NBOX ((((c.s2 >> 6) & 0x03)
601 | ((c.s2 >> 7) & 0x3c)), 1, s_skb)
602 | NBOX ((((c.s2 >> 13) & 0x0f)
603 | ((c.s2 >> 14) & 0x30)), 2, s_skb)
604 | NBOX ((((c.s2 >> 20) & 0x01)
605 | ((c.s2 >> 21) & 0x06)
606 | ((c.s2 >> 22) & 0x38)), 3, s_skb);
608 t.s2 = NBOX ((( d.s2 >> 0) & 0x3f), 4, s_skb)
609 | NBOX ((((d.s2 >> 7) & 0x03)
610 | ((d.s2 >> 8) & 0x3c)), 5, s_skb)
611 | NBOX ((((d.s2 >> 15) & 0x3f)), 6, s_skb)
612 | NBOX ((((d.s2 >> 21) & 0x0f)
613 | ((d.s2 >> 22) & 0x30)), 7, s_skb);
615 s.s3 = NBOX ((( c.s3 >> 0) & 0x3f), 0, s_skb)
616 | NBOX ((((c.s3 >> 6) & 0x03)
617 | ((c.s3 >> 7) & 0x3c)), 1, s_skb)
618 | NBOX ((((c.s3 >> 13) & 0x0f)
619 | ((c.s3 >> 14) & 0x30)), 2, s_skb)
620 | NBOX ((((c.s3 >> 20) & 0x01)
621 | ((c.s3 >> 21) & 0x06)
622 | ((c.s3 >> 22) & 0x38)), 3, s_skb);
624 t.s3 = NBOX ((( d.s3 >> 0) & 0x3f), 4, s_skb)
625 | NBOX ((((d.s3 >> 7) & 0x03)
626 | ((d.s3 >> 8) & 0x3c)), 5, s_skb)
627 | NBOX ((((d.s3 >> 15) & 0x3f)), 6, s_skb)
628 | NBOX ((((d.s3 >> 21) & 0x0f)
629 | ((d.s3 >> 22) & 0x30)), 7, s_skb);
632 #if defined cl_amd_media_ops
633 Kc[i] = amd_bytealign (t, s << 16, 2);
634 Kd[i] = amd_bytealign (t >> 16, s, 2);
636 Kc[i] = ((t << 16) | (s & 0x0000ffff));
637 Kd[i] = ((s >> 16) | (t & 0xffff0000));
640 Kc[i] = rotl32 (Kc[i], 2u);
641 Kd[i] = rotl32 (Kd[i], 2u);
645 static void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
649 key[0] = (ascii_to_ebcdic_pc[(w0 >> 0) & 0xff]) << 0
650 | (ascii_to_ebcdic_pc[(w0 >> 8) & 0xff]) << 8
651 | (ascii_to_ebcdic_pc[(w0 >> 16) & 0xff]) << 16
652 | (ascii_to_ebcdic_pc[(w0 >> 24) & 0xff]) << 24;
654 key[1] = (ascii_to_ebcdic_pc[(w1 >> 0) & 0xff]) << 0
655 | (ascii_to_ebcdic_pc[(w1 >> 8) & 0xff]) << 8
656 | (ascii_to_ebcdic_pc[(w1 >> 16) & 0xff]) << 16
657 | (ascii_to_ebcdic_pc[(w1 >> 24) & 0xff]) << 24;
662 key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
663 | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
664 | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
665 | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
667 key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
668 | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
669 | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
670 | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
672 key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
673 | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
674 | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
675 | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
677 key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
678 | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
679 | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
680 | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
684 key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
685 | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
686 | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
687 | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
689 key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
690 | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
691 | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
692 | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
694 key[0].s2 = (ascii_to_ebcdic_pc[(w0.s2 >> 0) & 0xff]) << 0
695 | (ascii_to_ebcdic_pc[(w0.s2 >> 8) & 0xff]) << 8
696 | (ascii_to_ebcdic_pc[(w0.s2 >> 16) & 0xff]) << 16
697 | (ascii_to_ebcdic_pc[(w0.s2 >> 24) & 0xff]) << 24;
699 key[0].s3 = (ascii_to_ebcdic_pc[(w0.s3 >> 0) & 0xff]) << 0
700 | (ascii_to_ebcdic_pc[(w0.s3 >> 8) & 0xff]) << 8
701 | (ascii_to_ebcdic_pc[(w0.s3 >> 16) & 0xff]) << 16
702 | (ascii_to_ebcdic_pc[(w0.s3 >> 24) & 0xff]) << 24;
704 key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
705 | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
706 | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
707 | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
709 key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
710 | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
711 | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
712 | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
714 key[1].s2 = (ascii_to_ebcdic_pc[(w1.s2 >> 0) & 0xff]) << 0
715 | (ascii_to_ebcdic_pc[(w1.s2 >> 8) & 0xff]) << 8
716 | (ascii_to_ebcdic_pc[(w1.s2 >> 16) & 0xff]) << 16
717 | (ascii_to_ebcdic_pc[(w1.s2 >> 24) & 0xff]) << 24;
719 key[1].s3 = (ascii_to_ebcdic_pc[(w1.s3 >> 0) & 0xff]) << 0
720 | (ascii_to_ebcdic_pc[(w1.s3 >> 8) & 0xff]) << 8
721 | (ascii_to_ebcdic_pc[(w1.s3 >> 16) & 0xff]) << 16
722 | (ascii_to_ebcdic_pc[(w1.s3 >> 24) & 0xff]) << 24;
726 static void m08500m (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
732 const u32 gid = get_global_id (0);
733 const u32 lid = get_local_id (0);
741 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
742 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
748 const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
754 for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
756 const u32x w0r = words_buf_r[il_pos];
758 const u32x w0 = w0l | w0r;
762 transform_racf_key (w0, w1, key);
764 const u32x c = key[0];
765 const u32x d = key[1];
770 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
774 data[0] = salt_buf0[0];
775 data[1] = salt_buf0[1];
779 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
781 const u32x r0 = iv[0];
782 const u32x r1 = iv[1];
786 #include VECT_COMPARE_M
790 static void m08500s (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
796 const u32 gid = get_global_id (0);
797 const u32 lid = get_local_id (0);
805 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
806 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
812 const u32 search[4] =
814 digests_buf[digests_offset].digest_buf[DGST_R0],
815 digests_buf[digests_offset].digest_buf[DGST_R1],
816 digests_buf[digests_offset].digest_buf[DGST_R2],
817 digests_buf[digests_offset].digest_buf[DGST_R3]
824 const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
830 for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
832 const u32x w0r = words_buf_r[il_pos];
834 const u32x w0 = w0l | w0r;
838 transform_racf_key (w0, w1, key);
840 const u32x c = key[0];
841 const u32x d = key[1];
846 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
850 data[0] = salt_buf0[0];
851 data[1] = salt_buf0[1];
855 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
857 const u32x r0 = iv[0];
858 const u32x r1 = iv[1];
862 #include VECT_COMPARE_S
866 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m08500_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
868 __local u32 s_SPtrans[8][64];
870 __local u32 s_skb[8][64];
876 const u32 gid = get_global_id (0);
877 const u32 lid = get_local_id (0);
881 w[ 0] = pws[gid].i[ 0];
882 w[ 1] = pws[gid].i[ 1];
898 const u32 pw_len = pws[gid].pw_len;
904 s_SPtrans[0][lid] = c_SPtrans[0][lid];
905 s_SPtrans[1][lid] = c_SPtrans[1][lid];
906 s_SPtrans[2][lid] = c_SPtrans[2][lid];
907 s_SPtrans[3][lid] = c_SPtrans[3][lid];
908 s_SPtrans[4][lid] = c_SPtrans[4][lid];
909 s_SPtrans[5][lid] = c_SPtrans[5][lid];
910 s_SPtrans[6][lid] = c_SPtrans[6][lid];
911 s_SPtrans[7][lid] = c_SPtrans[7][lid];
913 s_skb[0][lid] = c_skb[0][lid];
914 s_skb[1][lid] = c_skb[1][lid];
915 s_skb[2][lid] = c_skb[2][lid];
916 s_skb[3][lid] = c_skb[3][lid];
917 s_skb[4][lid] = c_skb[4][lid];
918 s_skb[5][lid] = c_skb[5][lid];
919 s_skb[6][lid] = c_skb[6][lid];
920 s_skb[7][lid] = c_skb[7][lid];
922 barrier (CLK_LOCAL_MEM_FENCE);
924 if (gid >= gid_max) return;
930 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, bfs_cnt, digests_cnt, digests_offset);
933 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m08500_m08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
937 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m08500_m16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
941 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m08500_s04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
943 __local u32 s_SPtrans[8][64];
945 __local u32 s_skb[8][64];
951 const u32 gid = get_global_id (0);
952 const u32 lid = get_local_id (0);
956 w[ 0] = pws[gid].i[ 0];
957 w[ 1] = pws[gid].i[ 1];
973 const u32 pw_len = pws[gid].pw_len;
979 s_SPtrans[0][lid] = c_SPtrans[0][lid];
980 s_SPtrans[1][lid] = c_SPtrans[1][lid];
981 s_SPtrans[2][lid] = c_SPtrans[2][lid];
982 s_SPtrans[3][lid] = c_SPtrans[3][lid];
983 s_SPtrans[4][lid] = c_SPtrans[4][lid];
984 s_SPtrans[5][lid] = c_SPtrans[5][lid];
985 s_SPtrans[6][lid] = c_SPtrans[6][lid];
986 s_SPtrans[7][lid] = c_SPtrans[7][lid];
988 s_skb[0][lid] = c_skb[0][lid];
989 s_skb[1][lid] = c_skb[1][lid];
990 s_skb[2][lid] = c_skb[2][lid];
991 s_skb[3][lid] = c_skb[3][lid];
992 s_skb[4][lid] = c_skb[4][lid];
993 s_skb[5][lid] = c_skb[5][lid];
994 s_skb[6][lid] = c_skb[6][lid];
995 s_skb[7][lid] = c_skb[7][lid];
997 barrier (CLK_LOCAL_MEM_FENCE);
999 if (gid >= gid_max) return;
1005 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, bfs_cnt, digests_cnt, digests_offset);
1008 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m08500_s08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1012 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m08500_s16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)