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