2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
16 #include "include/kernel_functions.c"
17 #include "types_ocl.c"
19 #include "include/rp_gpu.h"
22 #define COMPARE_S "check_single_comp4.c"
23 #define COMPARE_M "check_multi_comp4.c"
25 #define PERM_OP(a,b,tt,n,m) \
35 #define HPERM_OP(a,tt,n,m) \
41 tt = tt >> (16 + n); \
47 PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \
48 PERM_OP (l, r, tt, 16, 0x0000ffff); \
49 PERM_OP (r, l, tt, 2, 0x33333333); \
50 PERM_OP (l, r, tt, 8, 0x00ff00ff); \
51 PERM_OP (r, l, tt, 1, 0x55555555); \
56 PERM_OP (l, r, tt, 1, 0x55555555); \
57 PERM_OP (r, l, tt, 8, 0x00ff00ff); \
58 PERM_OP (l, r, tt, 2, 0x33333333); \
59 PERM_OP (r, l, tt, 16, 0x0000ffff); \
60 PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \
63 __constant u8 ascii_to_ebcdic_pc[256] =
65 // little hack, can't crack 0-bytes in password, but who cares
66 // 0xab, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
67 0x2a, 0xa8, 0xae, 0xad, 0xc4, 0xf1, 0xf7, 0xf4, 0x86, 0xa1, 0xe0, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5,
68 0x8a, 0x89, 0x8f, 0x8c, 0xd3, 0xd0, 0xce, 0xe6, 0x9b, 0x98, 0xd5, 0xe5, 0x92, 0x91, 0x97, 0x94,
69 0x2a, 0x34, 0x54, 0x5d, 0x1c, 0x73, 0x0b, 0x51, 0x31, 0x10, 0x13, 0x37, 0x7c, 0x6b, 0x3d, 0x68,
70 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40, 0x46, 0x45, 0x5b, 0x58, 0x5e, 0x16, 0x32, 0x57, 0x76, 0x75,
71 0x52, 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07,
72 0x04, 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x3e, 0x6b, 0x1f, 0x15, 0x70,
73 0x58, 0xa8, 0xae, 0xad, 0xa2, 0xa1, 0xa7, 0xa4, 0xba, 0xb9, 0x89, 0x8f, 0x8c, 0x83, 0x80, 0x86,
74 0x85, 0x9b, 0x98, 0xef, 0xec, 0xe3, 0xe0, 0xe6, 0xe5, 0xfb, 0xf8, 0x2a, 0x7f, 0x0b, 0xe9, 0xa4,
75 0xea, 0xe9, 0xef, 0xec, 0xe3, 0x80, 0xa7, 0x85, 0xfb, 0xf8, 0xfe, 0xfd, 0xf2, 0xb9, 0xbf, 0x9d,
76 0xcb, 0xc8, 0x9e, 0xcd, 0xc2, 0xc1, 0xc7, 0xba, 0xda, 0xd9, 0xdf, 0xdc, 0xa2, 0x83, 0xd6, 0x68,
77 0x29, 0x2f, 0x2c, 0x23, 0x20, 0x26, 0x25, 0x3b, 0x38, 0x08, 0x0e, 0x0d, 0x02, 0x01, 0x07, 0x04,
78 0x1a, 0x19, 0x6e, 0x6d, 0x62, 0x61, 0x67, 0x64, 0x7a, 0x79, 0x4a, 0x49, 0x4f, 0x4c, 0x43, 0x40,
79 0x46, 0x45, 0x5b, 0xab, 0xbf, 0xbc, 0xb3, 0xb0, 0xb6, 0xb5, 0x8a, 0x9e, 0x9d, 0x92, 0x91, 0x97,
80 0x94, 0xea, 0xfe, 0xfd, 0xf2, 0xf1, 0xf7, 0xf4, 0xcb, 0xc8, 0xce, 0xcd, 0xc2, 0xc1, 0xc7, 0xc4,
81 0xda, 0xd9, 0xdf, 0xdc, 0xd3, 0xd0, 0xd6, 0xd5, 0x3e, 0x3d, 0x32, 0x31, 0x37, 0x34, 0x1f, 0x1c,
82 0x13, 0x10, 0x16, 0x15, 0x7f, 0x7c, 0x73, 0x70, 0x76, 0x75, 0x5e, 0x5d, 0x52, 0x51, 0x57, 0x54,
85 __constant u32 c_SPtrans[8][64] =
88 0x02080800, 0x00080000, 0x02000002, 0x02080802,
89 0x02000000, 0x00080802, 0x00080002, 0x02000002,
90 0x00080802, 0x02080800, 0x02080000, 0x00000802,
91 0x02000802, 0x02000000, 0x00000000, 0x00080002,
92 0x00080000, 0x00000002, 0x02000800, 0x00080800,
93 0x02080802, 0x02080000, 0x00000802, 0x02000800,
94 0x00000002, 0x00000800, 0x00080800, 0x02080002,
95 0x00000800, 0x02000802, 0x02080002, 0x00000000,
96 0x00000000, 0x02080802, 0x02000800, 0x00080002,
97 0x02080800, 0x00080000, 0x00000802, 0x02000800,
98 0x02080002, 0x00000800, 0x00080800, 0x02000002,
99 0x00080802, 0x00000002, 0x02000002, 0x02080000,
100 0x02080802, 0x00080800, 0x02080000, 0x02000802,
101 0x02000000, 0x00000802, 0x00080002, 0x00000000,
102 0x00080000, 0x02000000, 0x02000802, 0x02080800,
103 0x00000002, 0x02080002, 0x00000800, 0x00080802,
105 0x40108010, 0x00000000, 0x00108000, 0x40100000,
106 0x40000010, 0x00008010, 0x40008000, 0x00108000,
107 0x00008000, 0x40100010, 0x00000010, 0x40008000,
108 0x00100010, 0x40108000, 0x40100000, 0x00000010,
109 0x00100000, 0x40008010, 0x40100010, 0x00008000,
110 0x00108010, 0x40000000, 0x00000000, 0x00100010,
111 0x40008010, 0x00108010, 0x40108000, 0x40000010,
112 0x40000000, 0x00100000, 0x00008010, 0x40108010,
113 0x00100010, 0x40108000, 0x40008000, 0x00108010,
114 0x40108010, 0x00100010, 0x40000010, 0x00000000,
115 0x40000000, 0x00008010, 0x00100000, 0x40100010,
116 0x00008000, 0x40000000, 0x00108010, 0x40008010,
117 0x40108000, 0x00008000, 0x00000000, 0x40000010,
118 0x00000010, 0x40108010, 0x00108000, 0x40100000,
119 0x40100010, 0x00100000, 0x00008010, 0x40008000,
120 0x40008010, 0x00000010, 0x40100000, 0x00108000,
122 0x04000001, 0x04040100, 0x00000100, 0x04000101,
123 0x00040001, 0x04000000, 0x04000101, 0x00040100,
124 0x04000100, 0x00040000, 0x04040000, 0x00000001,
125 0x04040101, 0x00000101, 0x00000001, 0x04040001,
126 0x00000000, 0x00040001, 0x04040100, 0x00000100,
127 0x00000101, 0x04040101, 0x00040000, 0x04000001,
128 0x04040001, 0x04000100, 0x00040101, 0x04040000,
129 0x00040100, 0x00000000, 0x04000000, 0x00040101,
130 0x04040100, 0x00000100, 0x00000001, 0x00040000,
131 0x00000101, 0x00040001, 0x04040000, 0x04000101,
132 0x00000000, 0x04040100, 0x00040100, 0x04040001,
133 0x00040001, 0x04000000, 0x04040101, 0x00000001,
134 0x00040101, 0x04000001, 0x04000000, 0x04040101,
135 0x00040000, 0x04000100, 0x04000101, 0x00040100,
136 0x04000100, 0x00000000, 0x04040001, 0x00000101,
137 0x04000001, 0x00040101, 0x00000100, 0x04040000,
139 0x00401008, 0x10001000, 0x00000008, 0x10401008,
140 0x00000000, 0x10400000, 0x10001008, 0x00400008,
141 0x10401000, 0x10000008, 0x10000000, 0x00001008,
142 0x10000008, 0x00401008, 0x00400000, 0x10000000,
143 0x10400008, 0x00401000, 0x00001000, 0x00000008,
144 0x00401000, 0x10001008, 0x10400000, 0x00001000,
145 0x00001008, 0x00000000, 0x00400008, 0x10401000,
146 0x10001000, 0x10400008, 0x10401008, 0x00400000,
147 0x10400008, 0x00001008, 0x00400000, 0x10000008,
148 0x00401000, 0x10001000, 0x00000008, 0x10400000,
149 0x10001008, 0x00000000, 0x00001000, 0x00400008,
150 0x00000000, 0x10400008, 0x10401000, 0x00001000,
151 0x10000000, 0x10401008, 0x00401008, 0x00400000,
152 0x10401008, 0x00000008, 0x10001000, 0x00401008,
153 0x00400008, 0x00401000, 0x10400000, 0x10001008,
154 0x00001008, 0x10000000, 0x10000008, 0x10401000,
156 0x08000000, 0x00010000, 0x00000400, 0x08010420,
157 0x08010020, 0x08000400, 0x00010420, 0x08010000,
158 0x00010000, 0x00000020, 0x08000020, 0x00010400,
159 0x08000420, 0x08010020, 0x08010400, 0x00000000,
160 0x00010400, 0x08000000, 0x00010020, 0x00000420,
161 0x08000400, 0x00010420, 0x00000000, 0x08000020,
162 0x00000020, 0x08000420, 0x08010420, 0x00010020,
163 0x08010000, 0x00000400, 0x00000420, 0x08010400,
164 0x08010400, 0x08000420, 0x00010020, 0x08010000,
165 0x00010000, 0x00000020, 0x08000020, 0x08000400,
166 0x08000000, 0x00010400, 0x08010420, 0x00000000,
167 0x00010420, 0x08000000, 0x00000400, 0x00010020,
168 0x08000420, 0x00000400, 0x00000000, 0x08010420,
169 0x08010020, 0x08010400, 0x00000420, 0x00010000,
170 0x00010400, 0x08010020, 0x08000400, 0x00000420,
171 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,
190 0x00004000, 0x00000200, 0x01000200, 0x01000004,
191 0x01004204, 0x00004004, 0x00004200, 0x00000000,
192 0x01000000, 0x01000204, 0x00000204, 0x01004000,
193 0x00000004, 0x01004200, 0x01004000, 0x00000204,
194 0x01000204, 0x00004000, 0x00004004, 0x01004204,
195 0x00000000, 0x01000200, 0x01000004, 0x00004200,
196 0x01004004, 0x00004204, 0x01004200, 0x00000004,
197 0x00004204, 0x01004004, 0x00000200, 0x01000000,
198 0x00004204, 0x01004000, 0x01004004, 0x00000204,
199 0x00004000, 0x00000200, 0x01000000, 0x01004004,
200 0x01000204, 0x00004204, 0x00004200, 0x00000000,
201 0x00000200, 0x01000004, 0x00000004, 0x01000200,
202 0x00000000, 0x01000204, 0x01000200, 0x00004200,
203 0x00000204, 0x00004000, 0x01004204, 0x01000000,
204 0x01004200, 0x00000004, 0x00004004, 0x01004204,
205 0x01000004, 0x01004200, 0x01004000, 0x00004004,
207 0x20800080, 0x20820000, 0x00020080, 0x00000000,
208 0x20020000, 0x00800080, 0x20800000, 0x20820080,
209 0x00000080, 0x20000000, 0x00820000, 0x00020080,
210 0x00820080, 0x20020080, 0x20000080, 0x20800000,
211 0x00020000, 0x00820080, 0x00800080, 0x20020000,
212 0x20820080, 0x20000080, 0x00000000, 0x00820000,
213 0x20000000, 0x00800000, 0x20020080, 0x20800080,
214 0x00800000, 0x00020000, 0x20820000, 0x00000080,
215 0x00800000, 0x00020000, 0x20000080, 0x20820080,
216 0x00020080, 0x20000000, 0x00000000, 0x00820000,
217 0x20800080, 0x20020080, 0x20020000, 0x00800080,
218 0x20820000, 0x00000080, 0x00800080, 0x20020000,
219 0x20820080, 0x00800000, 0x20800000, 0x20000080,
220 0x00820000, 0x00020080, 0x20020080, 0x20800000,
221 0x00000080, 0x20820000, 0x00820080, 0x00000000,
222 0x20000000, 0x20800080, 0x00020000, 0x00820080,
225 __constant u32 c_skb[8][64] =
227 /* for C bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
228 0x00000000, 0x00000010, 0x20000000, 0x20000010,
229 0x00010000, 0x00010010, 0x20010000, 0x20010010,
230 0x00000800, 0x00000810, 0x20000800, 0x20000810,
231 0x00010800, 0x00010810, 0x20010800, 0x20010810,
232 0x00000020, 0x00000030, 0x20000020, 0x20000030,
233 0x00010020, 0x00010030, 0x20010020, 0x20010030,
234 0x00000820, 0x00000830, 0x20000820, 0x20000830,
235 0x00010820, 0x00010830, 0x20010820, 0x20010830,
236 0x00080000, 0x00080010, 0x20080000, 0x20080010,
237 0x00090000, 0x00090010, 0x20090000, 0x20090010,
238 0x00080800, 0x00080810, 0x20080800, 0x20080810,
239 0x00090800, 0x00090810, 0x20090800, 0x20090810,
240 0x00080020, 0x00080030, 0x20080020, 0x20080030,
241 0x00090020, 0x00090030, 0x20090020, 0x20090030,
242 0x00080820, 0x00080830, 0x20080820, 0x20080830,
243 0x00090820, 0x00090830, 0x20090820, 0x20090830,
244 /* for C bits (numbered as per FIPS 46) 7 8 10 11 12 13 */
245 0x00000000, 0x02000000, 0x00002000, 0x02002000,
246 0x00200000, 0x02200000, 0x00202000, 0x02202000,
247 0x00000004, 0x02000004, 0x00002004, 0x02002004,
248 0x00200004, 0x02200004, 0x00202004, 0x02202004,
249 0x00000400, 0x02000400, 0x00002400, 0x02002400,
250 0x00200400, 0x02200400, 0x00202400, 0x02202400,
251 0x00000404, 0x02000404, 0x00002404, 0x02002404,
252 0x00200404, 0x02200404, 0x00202404, 0x02202404,
253 0x10000000, 0x12000000, 0x10002000, 0x12002000,
254 0x10200000, 0x12200000, 0x10202000, 0x12202000,
255 0x10000004, 0x12000004, 0x10002004, 0x12002004,
256 0x10200004, 0x12200004, 0x10202004, 0x12202004,
257 0x10000400, 0x12000400, 0x10002400, 0x12002400,
258 0x10200400, 0x12200400, 0x10202400, 0x12202400,
259 0x10000404, 0x12000404, 0x10002404, 0x12002404,
260 0x10200404, 0x12200404, 0x10202404, 0x12202404,
261 /* for C bits (numbered as per FIPS 46) 14 15 16 17 19 20 */
262 0x00000000, 0x00000001, 0x00040000, 0x00040001,
263 0x01000000, 0x01000001, 0x01040000, 0x01040001,
264 0x00000002, 0x00000003, 0x00040002, 0x00040003,
265 0x01000002, 0x01000003, 0x01040002, 0x01040003,
266 0x00000200, 0x00000201, 0x00040200, 0x00040201,
267 0x01000200, 0x01000201, 0x01040200, 0x01040201,
268 0x00000202, 0x00000203, 0x00040202, 0x00040203,
269 0x01000202, 0x01000203, 0x01040202, 0x01040203,
270 0x08000000, 0x08000001, 0x08040000, 0x08040001,
271 0x09000000, 0x09000001, 0x09040000, 0x09040001,
272 0x08000002, 0x08000003, 0x08040002, 0x08040003,
273 0x09000002, 0x09000003, 0x09040002, 0x09040003,
274 0x08000200, 0x08000201, 0x08040200, 0x08040201,
275 0x09000200, 0x09000201, 0x09040200, 0x09040201,
276 0x08000202, 0x08000203, 0x08040202, 0x08040203,
277 0x09000202, 0x09000203, 0x09040202, 0x09040203,
278 /* for C bits (numbered as per FIPS 46) 21 23 24 26 27 28 */
279 0x00000000, 0x00100000, 0x00000100, 0x00100100,
280 0x00000008, 0x00100008, 0x00000108, 0x00100108,
281 0x00001000, 0x00101000, 0x00001100, 0x00101100,
282 0x00001008, 0x00101008, 0x00001108, 0x00101108,
283 0x04000000, 0x04100000, 0x04000100, 0x04100100,
284 0x04000008, 0x04100008, 0x04000108, 0x04100108,
285 0x04001000, 0x04101000, 0x04001100, 0x04101100,
286 0x04001008, 0x04101008, 0x04001108, 0x04101108,
287 0x00020000, 0x00120000, 0x00020100, 0x00120100,
288 0x00020008, 0x00120008, 0x00020108, 0x00120108,
289 0x00021000, 0x00121000, 0x00021100, 0x00121100,
290 0x00021008, 0x00121008, 0x00021108, 0x00121108,
291 0x04020000, 0x04120000, 0x04020100, 0x04120100,
292 0x04020008, 0x04120008, 0x04020108, 0x04120108,
293 0x04021000, 0x04121000, 0x04021100, 0x04121100,
294 0x04021008, 0x04121008, 0x04021108, 0x04121108,
295 /* for D bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
296 0x00000000, 0x10000000, 0x00010000, 0x10010000,
297 0x00000004, 0x10000004, 0x00010004, 0x10010004,
298 0x20000000, 0x30000000, 0x20010000, 0x30010000,
299 0x20000004, 0x30000004, 0x20010004, 0x30010004,
300 0x00100000, 0x10100000, 0x00110000, 0x10110000,
301 0x00100004, 0x10100004, 0x00110004, 0x10110004,
302 0x20100000, 0x30100000, 0x20110000, 0x30110000,
303 0x20100004, 0x30100004, 0x20110004, 0x30110004,
304 0x00001000, 0x10001000, 0x00011000, 0x10011000,
305 0x00001004, 0x10001004, 0x00011004, 0x10011004,
306 0x20001000, 0x30001000, 0x20011000, 0x30011000,
307 0x20001004, 0x30001004, 0x20011004, 0x30011004,
308 0x00101000, 0x10101000, 0x00111000, 0x10111000,
309 0x00101004, 0x10101004, 0x00111004, 0x10111004,
310 0x20101000, 0x30101000, 0x20111000, 0x30111000,
311 0x20101004, 0x30101004, 0x20111004, 0x30111004,
312 /* for D bits (numbered as per FIPS 46) 8 9 11 12 13 14 */
313 0x00000000, 0x08000000, 0x00000008, 0x08000008,
314 0x00000400, 0x08000400, 0x00000408, 0x08000408,
315 0x00020000, 0x08020000, 0x00020008, 0x08020008,
316 0x00020400, 0x08020400, 0x00020408, 0x08020408,
317 0x00000001, 0x08000001, 0x00000009, 0x08000009,
318 0x00000401, 0x08000401, 0x00000409, 0x08000409,
319 0x00020001, 0x08020001, 0x00020009, 0x08020009,
320 0x00020401, 0x08020401, 0x00020409, 0x08020409,
321 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
322 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
323 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
324 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
325 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
326 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
327 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
328 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
329 /* for D bits (numbered as per FIPS 46) 16 17 18 19 20 21 */
330 0x00000000, 0x00000100, 0x00080000, 0x00080100,
331 0x01000000, 0x01000100, 0x01080000, 0x01080100,
332 0x00000010, 0x00000110, 0x00080010, 0x00080110,
333 0x01000010, 0x01000110, 0x01080010, 0x01080110,
334 0x00200000, 0x00200100, 0x00280000, 0x00280100,
335 0x01200000, 0x01200100, 0x01280000, 0x01280100,
336 0x00200010, 0x00200110, 0x00280010, 0x00280110,
337 0x01200010, 0x01200110, 0x01280010, 0x01280110,
338 0x00000200, 0x00000300, 0x00080200, 0x00080300,
339 0x01000200, 0x01000300, 0x01080200, 0x01080300,
340 0x00000210, 0x00000310, 0x00080210, 0x00080310,
341 0x01000210, 0x01000310, 0x01080210, 0x01080310,
342 0x00200200, 0x00200300, 0x00280200, 0x00280300,
343 0x01200200, 0x01200300, 0x01280200, 0x01280300,
344 0x00200210, 0x00200310, 0x00280210, 0x00280310,
345 0x01200210, 0x01200310, 0x01280210, 0x01280310,
346 /* for D bits (numbered as per FIPS 46) 22 23 24 25 27 28 */
347 0x00000000, 0x04000000, 0x00040000, 0x04040000,
348 0x00000002, 0x04000002, 0x00040002, 0x04040002,
349 0x00002000, 0x04002000, 0x00042000, 0x04042000,
350 0x00002002, 0x04002002, 0x00042002, 0x04042002,
351 0x00000020, 0x04000020, 0x00040020, 0x04040020,
352 0x00000022, 0x04000022, 0x00040022, 0x04040022,
353 0x00002020, 0x04002020, 0x00042020, 0x04042020,
354 0x00002022, 0x04002022, 0x00042022, 0x04042022,
355 0x00000800, 0x04000800, 0x00040800, 0x04040800,
356 0x00000802, 0x04000802, 0x00040802, 0x04040802,
357 0x00002800, 0x04002800, 0x00042800, 0x04042800,
358 0x00002802, 0x04002802, 0x00042802, 0x04042802,
359 0x00000820, 0x04000820, 0x00040820, 0x04040820,
360 0x00000822, 0x04000822, 0x00040822, 0x04040822,
361 0x00002820, 0x04002820, 0x00042820, 0x04042820,
362 0x00002822, 0x04002822, 0x00042822, 0x04042822,
365 #define NBOX(i,n,S) (S)[(n)][(i)]
367 static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
375 for (u32 i = 0; i < 16; i++)
378 u32 t = Kd[i] ^ rotl32 (r, 28u);
381 l ^= NBOX (((u >> 2) & 0x3f), 0, s_SPtrans)
382 | NBOX (((u >> 10) & 0x3f), 2, s_SPtrans)
383 | NBOX (((u >> 18) & 0x3f), 4, s_SPtrans)
384 | NBOX (((u >> 26) & 0x3f), 6, s_SPtrans)
385 | NBOX (((t >> 2) & 0x3f), 1, s_SPtrans)
386 | NBOX (((t >> 10) & 0x3f), 3, s_SPtrans)
387 | NBOX (((t >> 18) & 0x3f), 5, s_SPtrans)
388 | NBOX (((t >> 26) & 0x3f), 7, s_SPtrans);
392 l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
393 | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
394 | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
395 | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
396 | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
397 | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
398 | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
399 | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
401 l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
402 | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
403 | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
404 | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
405 | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
406 | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
407 | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
408 | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
412 l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
413 | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
414 | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
415 | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
416 | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
417 | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
418 | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
419 | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
421 l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
422 | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
423 | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
424 | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
425 | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
426 | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
427 | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
428 | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
430 l.s2 ^= NBOX (((u.s2 >> 2) & 0x3f), 0, s_SPtrans)
431 | NBOX (((u.s2 >> 10) & 0x3f), 2, s_SPtrans)
432 | NBOX (((u.s2 >> 18) & 0x3f), 4, s_SPtrans)
433 | NBOX (((u.s2 >> 26) & 0x3f), 6, s_SPtrans)
434 | NBOX (((t.s2 >> 2) & 0x3f), 1, s_SPtrans)
435 | NBOX (((t.s2 >> 10) & 0x3f), 3, s_SPtrans)
436 | NBOX (((t.s2 >> 18) & 0x3f), 5, s_SPtrans)
437 | NBOX (((t.s2 >> 26) & 0x3f), 7, s_SPtrans);
439 l.s3 ^= NBOX (((u.s3 >> 2) & 0x3f), 0, s_SPtrans)
440 | NBOX (((u.s3 >> 10) & 0x3f), 2, s_SPtrans)
441 | NBOX (((u.s3 >> 18) & 0x3f), 4, s_SPtrans)
442 | NBOX (((u.s3 >> 26) & 0x3f), 6, s_SPtrans)
443 | NBOX (((t.s3 >> 2) & 0x3f), 1, s_SPtrans)
444 | NBOX (((t.s3 >> 10) & 0x3f), 3, s_SPtrans)
445 | NBOX (((t.s3 >> 18) & 0x3f), 5, s_SPtrans)
446 | NBOX (((t.s3 >> 26) & 0x3f), 7, s_SPtrans);
458 static void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 s_skb[8][64])
462 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
463 HPERM_OP (c, tt, 2, 0xcccc0000);
464 HPERM_OP (d, tt, 2, 0xcccc0000);
465 PERM_OP (d, c, tt, 1, 0x55555555);
466 PERM_OP (c, d, tt, 8, 0x00ff00ff);
467 PERM_OP (d, c, tt, 1, 0x55555555);
469 d = ((d & 0x000000ff) << 16)
470 | ((d & 0x0000ff00) << 0)
471 | ((d & 0x00ff0000) >> 16)
472 | ((c & 0xf0000000) >> 4);
477 for (u32 i = 0; i < 16; i++)
479 const u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
480 const u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
482 c = c >> shifts3s0[i] | c << shifts3s1[i];
483 d = d >> shifts3s0[i] | d << shifts3s1[i];
492 s = NBOX ((( c >> 0) & 0x3f), 0, s_skb)
493 | NBOX ((((c >> 6) & 0x03)
494 | ((c >> 7) & 0x3c)), 1, s_skb)
495 | NBOX ((((c >> 13) & 0x0f)
496 | ((c >> 14) & 0x30)), 2, s_skb)
497 | NBOX ((((c >> 20) & 0x01)
499 | ((c >> 22) & 0x38)), 3, s_skb);
501 t = NBOX ((( d >> 0) & 0x3f), 4, s_skb)
502 | NBOX ((((d >> 7) & 0x03)
503 | ((d >> 8) & 0x3c)), 5, s_skb)
504 | NBOX ((((d >> 15) & 0x3f)), 6, s_skb)
505 | NBOX ((((d >> 21) & 0x0f)
506 | ((d >> 22) & 0x30)), 7, s_skb);
510 s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
511 | NBOX ((((c.s0 >> 6) & 0x03)
512 | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
513 | NBOX ((((c.s0 >> 13) & 0x0f)
514 | ((c.s0 >> 14) & 0x30)), 2, s_skb)
515 | NBOX ((((c.s0 >> 20) & 0x01)
516 | ((c.s0 >> 21) & 0x06)
517 | ((c.s0 >> 22) & 0x38)), 3, s_skb);
519 t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
520 | NBOX ((((d.s0 >> 7) & 0x03)
521 | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
522 | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
523 | NBOX ((((d.s0 >> 21) & 0x0f)
524 | ((d.s0 >> 22) & 0x30)), 7, s_skb);
526 s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
527 | NBOX ((((c.s1 >> 6) & 0x03)
528 | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
529 | NBOX ((((c.s1 >> 13) & 0x0f)
530 | ((c.s1 >> 14) & 0x30)), 2, s_skb)
531 | NBOX ((((c.s1 >> 20) & 0x01)
532 | ((c.s1 >> 21) & 0x06)
533 | ((c.s1 >> 22) & 0x38)), 3, s_skb);
535 t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
536 | NBOX ((((d.s1 >> 7) & 0x03)
537 | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
538 | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
539 | NBOX ((((d.s1 >> 21) & 0x0f)
540 | ((d.s1 >> 22) & 0x30)), 7, s_skb);
544 s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
545 | NBOX ((((c.s0 >> 6) & 0x03)
546 | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
547 | NBOX ((((c.s0 >> 13) & 0x0f)
548 | ((c.s0 >> 14) & 0x30)), 2, s_skb)
549 | NBOX ((((c.s0 >> 20) & 0x01)
550 | ((c.s0 >> 21) & 0x06)
551 | ((c.s0 >> 22) & 0x38)), 3, s_skb);
553 t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
554 | NBOX ((((d.s0 >> 7) & 0x03)
555 | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
556 | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
557 | NBOX ((((d.s0 >> 21) & 0x0f)
558 | ((d.s0 >> 22) & 0x30)), 7, s_skb);
560 s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
561 | NBOX ((((c.s1 >> 6) & 0x03)
562 | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
563 | NBOX ((((c.s1 >> 13) & 0x0f)
564 | ((c.s1 >> 14) & 0x30)), 2, s_skb)
565 | NBOX ((((c.s1 >> 20) & 0x01)
566 | ((c.s1 >> 21) & 0x06)
567 | ((c.s1 >> 22) & 0x38)), 3, s_skb);
569 t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
570 | NBOX ((((d.s1 >> 7) & 0x03)
571 | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
572 | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
573 | NBOX ((((d.s1 >> 21) & 0x0f)
574 | ((d.s1 >> 22) & 0x30)), 7, s_skb);
576 s.s2 = NBOX ((( c.s2 >> 0) & 0x3f), 0, s_skb)
577 | NBOX ((((c.s2 >> 6) & 0x03)
578 | ((c.s2 >> 7) & 0x3c)), 1, s_skb)
579 | NBOX ((((c.s2 >> 13) & 0x0f)
580 | ((c.s2 >> 14) & 0x30)), 2, s_skb)
581 | NBOX ((((c.s2 >> 20) & 0x01)
582 | ((c.s2 >> 21) & 0x06)
583 | ((c.s2 >> 22) & 0x38)), 3, s_skb);
585 t.s2 = NBOX ((( d.s2 >> 0) & 0x3f), 4, s_skb)
586 | NBOX ((((d.s2 >> 7) & 0x03)
587 | ((d.s2 >> 8) & 0x3c)), 5, s_skb)
588 | NBOX ((((d.s2 >> 15) & 0x3f)), 6, s_skb)
589 | NBOX ((((d.s2 >> 21) & 0x0f)
590 | ((d.s2 >> 22) & 0x30)), 7, s_skb);
592 s.s3 = NBOX ((( c.s3 >> 0) & 0x3f), 0, s_skb)
593 | NBOX ((((c.s3 >> 6) & 0x03)
594 | ((c.s3 >> 7) & 0x3c)), 1, s_skb)
595 | NBOX ((((c.s3 >> 13) & 0x0f)
596 | ((c.s3 >> 14) & 0x30)), 2, s_skb)
597 | NBOX ((((c.s3 >> 20) & 0x01)
598 | ((c.s3 >> 21) & 0x06)
599 | ((c.s3 >> 22) & 0x38)), 3, s_skb);
601 t.s3 = NBOX ((( d.s3 >> 0) & 0x3f), 4, s_skb)
602 | NBOX ((((d.s3 >> 7) & 0x03)
603 | ((d.s3 >> 8) & 0x3c)), 5, s_skb)
604 | NBOX ((((d.s3 >> 15) & 0x3f)), 6, s_skb)
605 | NBOX ((((d.s3 >> 21) & 0x0f)
606 | ((d.s3 >> 22) & 0x30)), 7, s_skb);
609 #if defined cl_amd_media_ops
610 Kc[i] = amd_bytealign (t, s << 16, 2);
611 Kd[i] = amd_bytealign (t >> 16, s, 2);
613 Kc[i] = ((t << 16) | (s & 0x0000ffff));
614 Kd[i] = ((s >> 16) | (t & 0xffff0000));
617 Kc[i] = rotl32 (Kc[i], 2u);
618 Kd[i] = rotl32 (Kd[i], 2u);
622 static void transform_racf_key (const u32 w0, const u32 w1, u32 key[2])
626 key[0] = (ascii_to_ebcdic_pc[(w0 >> 0) & 0xff]) << 0
627 | (ascii_to_ebcdic_pc[(w0 >> 8) & 0xff]) << 8
628 | (ascii_to_ebcdic_pc[(w0 >> 16) & 0xff]) << 16
629 | (ascii_to_ebcdic_pc[(w0 >> 24) & 0xff]) << 24;
631 key[1] = (ascii_to_ebcdic_pc[(w1 >> 0) & 0xff]) << 0
632 | (ascii_to_ebcdic_pc[(w1 >> 8) & 0xff]) << 8
633 | (ascii_to_ebcdic_pc[(w1 >> 16) & 0xff]) << 16
634 | (ascii_to_ebcdic_pc[(w1 >> 24) & 0xff]) << 24;
639 key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
640 | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
641 | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
642 | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
644 key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
645 | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
646 | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
647 | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
649 key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
650 | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
651 | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
652 | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
654 key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
655 | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
656 | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
657 | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
661 key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
662 | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
663 | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
664 | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
666 key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
667 | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
668 | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
669 | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
671 key[0].s2 = (ascii_to_ebcdic_pc[(w0.s2 >> 0) & 0xff]) << 0
672 | (ascii_to_ebcdic_pc[(w0.s2 >> 8) & 0xff]) << 8
673 | (ascii_to_ebcdic_pc[(w0.s2 >> 16) & 0xff]) << 16
674 | (ascii_to_ebcdic_pc[(w0.s2 >> 24) & 0xff]) << 24;
676 key[0].s3 = (ascii_to_ebcdic_pc[(w0.s3 >> 0) & 0xff]) << 0
677 | (ascii_to_ebcdic_pc[(w0.s3 >> 8) & 0xff]) << 8
678 | (ascii_to_ebcdic_pc[(w0.s3 >> 16) & 0xff]) << 16
679 | (ascii_to_ebcdic_pc[(w0.s3 >> 24) & 0xff]) << 24;
681 key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
682 | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
683 | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
684 | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
686 key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
687 | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
688 | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
689 | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
691 key[1].s2 = (ascii_to_ebcdic_pc[(w1.s2 >> 0) & 0xff]) << 0
692 | (ascii_to_ebcdic_pc[(w1.s2 >> 8) & 0xff]) << 8
693 | (ascii_to_ebcdic_pc[(w1.s2 >> 16) & 0xff]) << 16
694 | (ascii_to_ebcdic_pc[(w1.s2 >> 24) & 0xff]) << 24;
696 key[1].s3 = (ascii_to_ebcdic_pc[(w1.s3 >> 0) & 0xff]) << 0
697 | (ascii_to_ebcdic_pc[(w1.s3 >> 8) & 0xff]) << 8
698 | (ascii_to_ebcdic_pc[(w1.s3 >> 16) & 0xff]) << 16
699 | (ascii_to_ebcdic_pc[(w1.s3 >> 24) & 0xff]) << 24;
703 __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)
709 const u32 lid = get_local_id (0);
715 const u32 gid = get_global_id (0);
719 pw_buf[0] = pws[gid].i[ 0];
720 pw_buf[1] = pws[gid].i[ 1];
724 const u32 pw_len = pws[gid].pw_len;
732 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
733 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
739 __local u32 s_SPtrans[8][64];
740 __local u32 s_skb[8][64];
742 s_SPtrans[0][lid] = c_SPtrans[0][lid];
743 s_SPtrans[1][lid] = c_SPtrans[1][lid];
744 s_SPtrans[2][lid] = c_SPtrans[2][lid];
745 s_SPtrans[3][lid] = c_SPtrans[3][lid];
746 s_SPtrans[4][lid] = c_SPtrans[4][lid];
747 s_SPtrans[5][lid] = c_SPtrans[5][lid];
748 s_SPtrans[6][lid] = c_SPtrans[6][lid];
749 s_SPtrans[7][lid] = c_SPtrans[7][lid];
751 s_skb[0][lid] = c_skb[0][lid];
752 s_skb[1][lid] = c_skb[1][lid];
753 s_skb[2][lid] = c_skb[2][lid];
754 s_skb[3][lid] = c_skb[3][lid];
755 s_skb[4][lid] = c_skb[4][lid];
756 s_skb[5][lid] = c_skb[5][lid];
757 s_skb[6][lid] = c_skb[6][lid];
758 s_skb[7][lid] = c_skb[7][lid];
760 barrier (CLK_LOCAL_MEM_FENCE);
762 if (gid >= gid_max) return;
768 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
798 u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
800 out_len = (out_len >= 8) ? 8 : out_len;
804 transform_racf_key (w0[0], w0[1], key);
806 const u32 c = key[0];
807 const u32 d = key[1];
812 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
816 data[0] = salt_buf0[0];
817 data[1] = salt_buf0[1];
821 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
823 const u32 r0 = iv[0];
824 const u32 r1 = iv[1];
832 __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)
836 __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)
840 __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)
846 const u32 lid = get_local_id (0);
852 const u32 gid = get_global_id (0);
856 pw_buf[0] = pws[gid].i[ 0];
857 pw_buf[1] = pws[gid].i[ 1];
861 const u32 pw_len = pws[gid].pw_len;
869 salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
870 salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
876 __local u32 s_SPtrans[8][64];
877 __local u32 s_skb[8][64];
879 s_SPtrans[0][lid] = c_SPtrans[0][lid];
880 s_SPtrans[1][lid] = c_SPtrans[1][lid];
881 s_SPtrans[2][lid] = c_SPtrans[2][lid];
882 s_SPtrans[3][lid] = c_SPtrans[3][lid];
883 s_SPtrans[4][lid] = c_SPtrans[4][lid];
884 s_SPtrans[5][lid] = c_SPtrans[5][lid];
885 s_SPtrans[6][lid] = c_SPtrans[6][lid];
886 s_SPtrans[7][lid] = c_SPtrans[7][lid];
888 s_skb[0][lid] = c_skb[0][lid];
889 s_skb[1][lid] = c_skb[1][lid];
890 s_skb[2][lid] = c_skb[2][lid];
891 s_skb[3][lid] = c_skb[3][lid];
892 s_skb[4][lid] = c_skb[4][lid];
893 s_skb[5][lid] = c_skb[5][lid];
894 s_skb[6][lid] = c_skb[6][lid];
895 s_skb[7][lid] = c_skb[7][lid];
897 barrier (CLK_LOCAL_MEM_FENCE);
899 if (gid >= gid_max) return;
905 const u32 search[4] =
907 digests_buf[digests_offset].digest_buf[DGST_R0],
908 digests_buf[digests_offset].digest_buf[DGST_R1],
909 digests_buf[digests_offset].digest_buf[DGST_R2],
910 digests_buf[digests_offset].digest_buf[DGST_R3]
917 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
947 u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
949 out_len = (out_len >= 8) ? 8 : out_len;
953 transform_racf_key (w0[0], w0[1], key);
955 const u32 c = key[0];
956 const u32 d = key[1];
961 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
965 data[0] = salt_buf0[0];
966 data[1] = salt_buf0[1];
970 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
972 const u32 r0 = iv[0];
973 const u32 r1 = iv[1];
981 __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)
985 __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)