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 "OpenCL/types_ocl.c"
18 #include "OpenCL/common.c"
20 #define COMPARE_S "OpenCL/check_single_comp4.c"
21 #define COMPARE_M "OpenCL/check_multi_comp4.c"
23 #define PERM_OP(a,b,tt,n,m) \
33 #define HPERM_OP(a,tt,n,m) \
39 tt = tt >> (16 + n); \
43 __constant u32 c_SPtrans[8][64] =
46 0x00820200, 0x00020000, 0x80800000, 0x80820200,
47 0x00800000, 0x80020200, 0x80020000, 0x80800000,
48 0x80020200, 0x00820200, 0x00820000, 0x80000200,
49 0x80800200, 0x00800000, 0x00000000, 0x80020000,
50 0x00020000, 0x80000000, 0x00800200, 0x00020200,
51 0x80820200, 0x00820000, 0x80000200, 0x00800200,
52 0x80000000, 0x00000200, 0x00020200, 0x80820000,
53 0x00000200, 0x80800200, 0x80820000, 0x00000000,
54 0x00000000, 0x80820200, 0x00800200, 0x80020000,
55 0x00820200, 0x00020000, 0x80000200, 0x00800200,
56 0x80820000, 0x00000200, 0x00020200, 0x80800000,
57 0x80020200, 0x80000000, 0x80800000, 0x00820000,
58 0x80820200, 0x00020200, 0x00820000, 0x80800200,
59 0x00800000, 0x80000200, 0x80020000, 0x00000000,
60 0x00020000, 0x00800000, 0x80800200, 0x00820200,
61 0x80000000, 0x80820000, 0x00000200, 0x80020200,
64 0x10042004, 0x00000000, 0x00042000, 0x10040000,
65 0x10000004, 0x00002004, 0x10002000, 0x00042000,
66 0x00002000, 0x10040004, 0x00000004, 0x10002000,
67 0x00040004, 0x10042000, 0x10040000, 0x00000004,
68 0x00040000, 0x10002004, 0x10040004, 0x00002000,
69 0x00042004, 0x10000000, 0x00000000, 0x00040004,
70 0x10002004, 0x00042004, 0x10042000, 0x10000004,
71 0x10000000, 0x00040000, 0x00002004, 0x10042004,
72 0x00040004, 0x10042000, 0x10002000, 0x00042004,
73 0x10042004, 0x00040004, 0x10000004, 0x00000000,
74 0x10000000, 0x00002004, 0x00040000, 0x10040004,
75 0x00002000, 0x10000000, 0x00042004, 0x10002004,
76 0x10042000, 0x00002000, 0x00000000, 0x10000004,
77 0x00000004, 0x10042004, 0x00042000, 0x10040000,
78 0x10040004, 0x00040000, 0x00002004, 0x10002000,
79 0x10002004, 0x00000004, 0x10040000, 0x00042000,
82 0x41000000, 0x01010040, 0x00000040, 0x41000040,
83 0x40010000, 0x01000000, 0x41000040, 0x00010040,
84 0x01000040, 0x00010000, 0x01010000, 0x40000000,
85 0x41010040, 0x40000040, 0x40000000, 0x41010000,
86 0x00000000, 0x40010000, 0x01010040, 0x00000040,
87 0x40000040, 0x41010040, 0x00010000, 0x41000000,
88 0x41010000, 0x01000040, 0x40010040, 0x01010000,
89 0x00010040, 0x00000000, 0x01000000, 0x40010040,
90 0x01010040, 0x00000040, 0x40000000, 0x00010000,
91 0x40000040, 0x40010000, 0x01010000, 0x41000040,
92 0x00000000, 0x01010040, 0x00010040, 0x41010000,
93 0x40010000, 0x01000000, 0x41010040, 0x40000000,
94 0x40010040, 0x41000000, 0x01000000, 0x41010040,
95 0x00010000, 0x01000040, 0x41000040, 0x00010040,
96 0x01000040, 0x00000000, 0x41010000, 0x40000040,
97 0x41000000, 0x40010040, 0x00000040, 0x01010000,
100 0x00100402, 0x04000400, 0x00000002, 0x04100402,
101 0x00000000, 0x04100000, 0x04000402, 0x00100002,
102 0x04100400, 0x04000002, 0x04000000, 0x00000402,
103 0x04000002, 0x00100402, 0x00100000, 0x04000000,
104 0x04100002, 0x00100400, 0x00000400, 0x00000002,
105 0x00100400, 0x04000402, 0x04100000, 0x00000400,
106 0x00000402, 0x00000000, 0x00100002, 0x04100400,
107 0x04000400, 0x04100002, 0x04100402, 0x00100000,
108 0x04100002, 0x00000402, 0x00100000, 0x04000002,
109 0x00100400, 0x04000400, 0x00000002, 0x04100000,
110 0x04000402, 0x00000000, 0x00000400, 0x00100002,
111 0x00000000, 0x04100002, 0x04100400, 0x00000400,
112 0x04000000, 0x04100402, 0x00100402, 0x00100000,
113 0x04100402, 0x00000002, 0x04000400, 0x00100402,
114 0x00100002, 0x00100400, 0x04100000, 0x04000402,
115 0x00000402, 0x04000000, 0x04000002, 0x04100400,
118 0x02000000, 0x00004000, 0x00000100, 0x02004108,
119 0x02004008, 0x02000100, 0x00004108, 0x02004000,
120 0x00004000, 0x00000008, 0x02000008, 0x00004100,
121 0x02000108, 0x02004008, 0x02004100, 0x00000000,
122 0x00004100, 0x02000000, 0x00004008, 0x00000108,
123 0x02000100, 0x00004108, 0x00000000, 0x02000008,
124 0x00000008, 0x02000108, 0x02004108, 0x00004008,
125 0x02004000, 0x00000100, 0x00000108, 0x02004100,
126 0x02004100, 0x02000108, 0x00004008, 0x02004000,
127 0x00004000, 0x00000008, 0x02000008, 0x02000100,
128 0x02000000, 0x00004100, 0x02004108, 0x00000000,
129 0x00004108, 0x02000000, 0x00000100, 0x00004008,
130 0x02000108, 0x00000100, 0x00000000, 0x02004108,
131 0x02004008, 0x02004100, 0x00000108, 0x00004000,
132 0x00004100, 0x02004008, 0x02000100, 0x00000108,
133 0x00000008, 0x00004108, 0x02004000, 0x02000008,
136 0x20000010, 0x00080010, 0x00000000, 0x20080800,
137 0x00080010, 0x00000800, 0x20000810, 0x00080000,
138 0x00000810, 0x20080810, 0x00080800, 0x20000000,
139 0x20000800, 0x20000010, 0x20080000, 0x00080810,
140 0x00080000, 0x20000810, 0x20080010, 0x00000000,
141 0x00000800, 0x00000010, 0x20080800, 0x20080010,
142 0x20080810, 0x20080000, 0x20000000, 0x00000810,
143 0x00000010, 0x00080800, 0x00080810, 0x20000800,
144 0x00000810, 0x20000000, 0x20000800, 0x00080810,
145 0x20080800, 0x00080010, 0x00000000, 0x20000800,
146 0x20000000, 0x00000800, 0x20080010, 0x00080000,
147 0x00080010, 0x20080810, 0x00080800, 0x00000010,
148 0x20080810, 0x00080800, 0x00080000, 0x20000810,
149 0x20000010, 0x20080000, 0x00080810, 0x00000000,
150 0x00000800, 0x20000010, 0x20000810, 0x20080800,
151 0x20080000, 0x00000810, 0x00000010, 0x20080010,
154 0x00001000, 0x00000080, 0x00400080, 0x00400001,
155 0x00401081, 0x00001001, 0x00001080, 0x00000000,
156 0x00400000, 0x00400081, 0x00000081, 0x00401000,
157 0x00000001, 0x00401080, 0x00401000, 0x00000081,
158 0x00400081, 0x00001000, 0x00001001, 0x00401081,
159 0x00000000, 0x00400080, 0x00400001, 0x00001080,
160 0x00401001, 0x00001081, 0x00401080, 0x00000001,
161 0x00001081, 0x00401001, 0x00000080, 0x00400000,
162 0x00001081, 0x00401000, 0x00401001, 0x00000081,
163 0x00001000, 0x00000080, 0x00400000, 0x00401001,
164 0x00400081, 0x00001081, 0x00001080, 0x00000000,
165 0x00000080, 0x00400001, 0x00000001, 0x00400080,
166 0x00000000, 0x00400081, 0x00400080, 0x00001080,
167 0x00000081, 0x00001000, 0x00401081, 0x00400000,
168 0x00401080, 0x00000001, 0x00001001, 0x00401081,
169 0x00400001, 0x00401080, 0x00401000, 0x00001001,
172 0x08200020, 0x08208000, 0x00008020, 0x00000000,
173 0x08008000, 0x00200020, 0x08200000, 0x08208020,
174 0x00000020, 0x08000000, 0x00208000, 0x00008020,
175 0x00208020, 0x08008020, 0x08000020, 0x08200000,
176 0x00008000, 0x00208020, 0x00200020, 0x08008000,
177 0x08208020, 0x08000020, 0x00000000, 0x00208000,
178 0x08000000, 0x00200000, 0x08008020, 0x08200020,
179 0x00200000, 0x00008000, 0x08208000, 0x00000020,
180 0x00200000, 0x00008000, 0x08000020, 0x08208020,
181 0x00008020, 0x08000000, 0x00000000, 0x00208000,
182 0x08200020, 0x08008020, 0x08008000, 0x00200020,
183 0x08208000, 0x00000020, 0x00200020, 0x08008000,
184 0x08208020, 0x00200000, 0x08200000, 0x08000020,
185 0x00208000, 0x00008020, 0x08008020, 0x08200000,
186 0x00000020, 0x08208000, 0x00208020, 0x00000000,
187 0x08000000, 0x08200020, 0x00008000, 0x00208020
191 __constant u32 c_skb[8][64] =
194 0x00000000, 0x00000010, 0x20000000, 0x20000010,
195 0x00010000, 0x00010010, 0x20010000, 0x20010010,
196 0x00000800, 0x00000810, 0x20000800, 0x20000810,
197 0x00010800, 0x00010810, 0x20010800, 0x20010810,
198 0x00000020, 0x00000030, 0x20000020, 0x20000030,
199 0x00010020, 0x00010030, 0x20010020, 0x20010030,
200 0x00000820, 0x00000830, 0x20000820, 0x20000830,
201 0x00010820, 0x00010830, 0x20010820, 0x20010830,
202 0x00080000, 0x00080010, 0x20080000, 0x20080010,
203 0x00090000, 0x00090010, 0x20090000, 0x20090010,
204 0x00080800, 0x00080810, 0x20080800, 0x20080810,
205 0x00090800, 0x00090810, 0x20090800, 0x20090810,
206 0x00080020, 0x00080030, 0x20080020, 0x20080030,
207 0x00090020, 0x00090030, 0x20090020, 0x20090030,
208 0x00080820, 0x00080830, 0x20080820, 0x20080830,
209 0x00090820, 0x00090830, 0x20090820, 0x20090830,
212 0x00000000, 0x02000000, 0x00002000, 0x02002000,
213 0x00200000, 0x02200000, 0x00202000, 0x02202000,
214 0x00000004, 0x02000004, 0x00002004, 0x02002004,
215 0x00200004, 0x02200004, 0x00202004, 0x02202004,
216 0x00000400, 0x02000400, 0x00002400, 0x02002400,
217 0x00200400, 0x02200400, 0x00202400, 0x02202400,
218 0x00000404, 0x02000404, 0x00002404, 0x02002404,
219 0x00200404, 0x02200404, 0x00202404, 0x02202404,
220 0x10000000, 0x12000000, 0x10002000, 0x12002000,
221 0x10200000, 0x12200000, 0x10202000, 0x12202000,
222 0x10000004, 0x12000004, 0x10002004, 0x12002004,
223 0x10200004, 0x12200004, 0x10202004, 0x12202004,
224 0x10000400, 0x12000400, 0x10002400, 0x12002400,
225 0x10200400, 0x12200400, 0x10202400, 0x12202400,
226 0x10000404, 0x12000404, 0x10002404, 0x12002404,
227 0x10200404, 0x12200404, 0x10202404, 0x12202404,
230 0x00000000, 0x00000001, 0x00040000, 0x00040001,
231 0x01000000, 0x01000001, 0x01040000, 0x01040001,
232 0x00000002, 0x00000003, 0x00040002, 0x00040003,
233 0x01000002, 0x01000003, 0x01040002, 0x01040003,
234 0x00000200, 0x00000201, 0x00040200, 0x00040201,
235 0x01000200, 0x01000201, 0x01040200, 0x01040201,
236 0x00000202, 0x00000203, 0x00040202, 0x00040203,
237 0x01000202, 0x01000203, 0x01040202, 0x01040203,
238 0x08000000, 0x08000001, 0x08040000, 0x08040001,
239 0x09000000, 0x09000001, 0x09040000, 0x09040001,
240 0x08000002, 0x08000003, 0x08040002, 0x08040003,
241 0x09000002, 0x09000003, 0x09040002, 0x09040003,
242 0x08000200, 0x08000201, 0x08040200, 0x08040201,
243 0x09000200, 0x09000201, 0x09040200, 0x09040201,
244 0x08000202, 0x08000203, 0x08040202, 0x08040203,
245 0x09000202, 0x09000203, 0x09040202, 0x09040203,
248 0x00000000, 0x00100000, 0x00000100, 0x00100100,
249 0x00000008, 0x00100008, 0x00000108, 0x00100108,
250 0x00001000, 0x00101000, 0x00001100, 0x00101100,
251 0x00001008, 0x00101008, 0x00001108, 0x00101108,
252 0x04000000, 0x04100000, 0x04000100, 0x04100100,
253 0x04000008, 0x04100008, 0x04000108, 0x04100108,
254 0x04001000, 0x04101000, 0x04001100, 0x04101100,
255 0x04001008, 0x04101008, 0x04001108, 0x04101108,
256 0x00020000, 0x00120000, 0x00020100, 0x00120100,
257 0x00020008, 0x00120008, 0x00020108, 0x00120108,
258 0x00021000, 0x00121000, 0x00021100, 0x00121100,
259 0x00021008, 0x00121008, 0x00021108, 0x00121108,
260 0x04020000, 0x04120000, 0x04020100, 0x04120100,
261 0x04020008, 0x04120008, 0x04020108, 0x04120108,
262 0x04021000, 0x04121000, 0x04021100, 0x04121100,
263 0x04021008, 0x04121008, 0x04021108, 0x04121108,
266 0x00000000, 0x10000000, 0x00010000, 0x10010000,
267 0x00000004, 0x10000004, 0x00010004, 0x10010004,
268 0x20000000, 0x30000000, 0x20010000, 0x30010000,
269 0x20000004, 0x30000004, 0x20010004, 0x30010004,
270 0x00100000, 0x10100000, 0x00110000, 0x10110000,
271 0x00100004, 0x10100004, 0x00110004, 0x10110004,
272 0x20100000, 0x30100000, 0x20110000, 0x30110000,
273 0x20100004, 0x30100004, 0x20110004, 0x30110004,
274 0x00001000, 0x10001000, 0x00011000, 0x10011000,
275 0x00001004, 0x10001004, 0x00011004, 0x10011004,
276 0x20001000, 0x30001000, 0x20011000, 0x30011000,
277 0x20001004, 0x30001004, 0x20011004, 0x30011004,
278 0x00101000, 0x10101000, 0x00111000, 0x10111000,
279 0x00101004, 0x10101004, 0x00111004, 0x10111004,
280 0x20101000, 0x30101000, 0x20111000, 0x30111000,
281 0x20101004, 0x30101004, 0x20111004, 0x30111004,
284 0x00000000, 0x08000000, 0x00000008, 0x08000008,
285 0x00000400, 0x08000400, 0x00000408, 0x08000408,
286 0x00020000, 0x08020000, 0x00020008, 0x08020008,
287 0x00020400, 0x08020400, 0x00020408, 0x08020408,
288 0x00000001, 0x08000001, 0x00000009, 0x08000009,
289 0x00000401, 0x08000401, 0x00000409, 0x08000409,
290 0x00020001, 0x08020001, 0x00020009, 0x08020009,
291 0x00020401, 0x08020401, 0x00020409, 0x08020409,
292 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
293 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
294 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
295 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
296 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
297 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
298 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
299 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
302 0x00000000, 0x00000100, 0x00080000, 0x00080100,
303 0x01000000, 0x01000100, 0x01080000, 0x01080100,
304 0x00000010, 0x00000110, 0x00080010, 0x00080110,
305 0x01000010, 0x01000110, 0x01080010, 0x01080110,
306 0x00200000, 0x00200100, 0x00280000, 0x00280100,
307 0x01200000, 0x01200100, 0x01280000, 0x01280100,
308 0x00200010, 0x00200110, 0x00280010, 0x00280110,
309 0x01200010, 0x01200110, 0x01280010, 0x01280110,
310 0x00000200, 0x00000300, 0x00080200, 0x00080300,
311 0x01000200, 0x01000300, 0x01080200, 0x01080300,
312 0x00000210, 0x00000310, 0x00080210, 0x00080310,
313 0x01000210, 0x01000310, 0x01080210, 0x01080310,
314 0x00200200, 0x00200300, 0x00280200, 0x00280300,
315 0x01200200, 0x01200300, 0x01280200, 0x01280300,
316 0x00200210, 0x00200310, 0x00280210, 0x00280310,
317 0x01200210, 0x01200310, 0x01280210, 0x01280310,
320 0x00000000, 0x04000000, 0x00040000, 0x04040000,
321 0x00000002, 0x04000002, 0x00040002, 0x04040002,
322 0x00002000, 0x04002000, 0x00042000, 0x04042000,
323 0x00002002, 0x04002002, 0x00042002, 0x04042002,
324 0x00000020, 0x04000020, 0x00040020, 0x04040020,
325 0x00000022, 0x04000022, 0x00040022, 0x04040022,
326 0x00002020, 0x04002020, 0x00042020, 0x04042020,
327 0x00002022, 0x04002022, 0x00042022, 0x04042022,
328 0x00000800, 0x04000800, 0x00040800, 0x04040800,
329 0x00000802, 0x04000802, 0x00040802, 0x04040802,
330 0x00002800, 0x04002800, 0x00042800, 0x04042800,
331 0x00002802, 0x04002802, 0x00042802, 0x04042802,
332 0x00000820, 0x04000820, 0x00040820, 0x04040820,
333 0x00000822, 0x04000822, 0x00040822, 0x04040822,
334 0x00002820, 0x04002820, 0x00042820, 0x04042820,
335 0x00002822, 0x04002822, 0x00042822, 0x04042822
339 #define BOX(i,n,S) (S)[(n)][(i)]
341 static void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 s_skb[8][64])
345 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
346 HPERM_OP (c, tt, 2, 0xcccc0000);
347 HPERM_OP (d, tt, 2, 0xcccc0000);
348 PERM_OP (d, c, tt, 1, 0x55555555);
349 PERM_OP (c, d, tt, 8, 0x00ff00ff);
350 PERM_OP (d, c, tt, 1, 0x55555555);
352 d = ((d & 0x000000ff) << 16)
353 | ((d & 0x0000ff00) << 0)
354 | ((d & 0x00ff0000) >> 16)
355 | ((c & 0xf0000000) >> 4);
360 for (u32 i = 0; i < 16; i++)
362 if ((i < 2) || (i == 8) || (i == 15))
364 c = ((c >> 1) | (c << 27));
365 d = ((d >> 1) | (d << 27));
369 c = ((c >> 2) | (c << 26));
370 d = ((d >> 2) | (d << 26));
376 const u32 c00 = (c >> 0) & 0x0000003f;
377 const u32 c06 = (c >> 6) & 0x00383003;
378 const u32 c07 = (c >> 7) & 0x0000003c;
379 const u32 c13 = (c >> 13) & 0x0000060f;
380 const u32 c20 = (c >> 20) & 0x00000001;
382 u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
383 | BOX (((c06 >> 0) & 0xff)
384 |((c07 >> 0) & 0xff), 1, s_skb)
385 | BOX (((c13 >> 0) & 0xff)
386 |((c06 >> 8) & 0xff), 2, s_skb)
387 | BOX (((c20 >> 0) & 0xff)
389 |((c06 >> 16) & 0xff), 3, s_skb);
391 const u32 d00 = (d >> 0) & 0x00003c3f;
392 const u32 d07 = (d >> 7) & 0x00003f03;
393 const u32 d21 = (d >> 21) & 0x0000000f;
394 const u32 d22 = (d >> 22) & 0x00000030;
396 u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
397 | BOX (((d07 >> 0) & 0xff)
398 |((d00 >> 8) & 0xff), 5, s_skb)
399 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
400 | BOX (((d21 >> 0) & 0xff)
401 |((d22 >> 0) & 0xff), 7, s_skb);
403 Kc[i] = ((t << 16) | (s & 0x0000ffff));
404 Kd[i] = ((s >> 16) | (t & 0xffff0000));
408 static void _des_crypt_encrypt (u32 iv[2], u32 mask, u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
410 const u32 E1 = (mask >> 2) & 0x3f0;
411 const u32 E0 = mask & 0x3f;
416 for (u32 i = 0; i < 25; i++)
418 for (u32 j = 0; j < 16; j += 2)
434 l ^= BOX (((u >> 0) & 0x3f), 0, s_SPtrans)
435 | BOX (((u >> 8) & 0x3f), 2, s_SPtrans)
436 | BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
437 | BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
438 | BOX (((t >> 0) & 0x3f), 1, s_SPtrans)
439 | BOX (((t >> 8) & 0x3f), 3, s_SPtrans)
440 | BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
441 | BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
454 r ^= BOX (((u >> 0) & 0x3f), 0, s_SPtrans)
455 | BOX (((u >> 8) & 0x3f), 2, s_SPtrans)
456 | BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
457 | BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
458 | BOX (((t >> 0) & 0x3f), 1, s_SPtrans)
459 | BOX (((t >> 8) & 0x3f), 3, s_SPtrans)
460 | BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
461 | BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
471 iv[0] = rotl32 (r, 31);
472 iv[1] = rotl32 (l, 31);
475 __kernel void m01500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
481 const u32 gid = get_global_id (0);
482 const u32 lid = get_local_id (0);
483 const u32 lsz = get_local_size (0);
491 wordl0[0] = pws[gid].i[ 0];
492 wordl0[1] = pws[gid].i[ 1];
517 const u32 pw_l_len = pws[gid].pw_len;
519 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
521 switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
528 __local u32 s_SPtrans[8][64];
529 __local u32 s_skb[8][64];
531 for (u32 i = lid; i < 64; i += lsz)
533 s_SPtrans[0][i] = c_SPtrans[0][i];
534 s_SPtrans[1][i] = c_SPtrans[1][i];
535 s_SPtrans[2][i] = c_SPtrans[2][i];
536 s_SPtrans[3][i] = c_SPtrans[3][i];
537 s_SPtrans[4][i] = c_SPtrans[4][i];
538 s_SPtrans[5][i] = c_SPtrans[5][i];
539 s_SPtrans[6][i] = c_SPtrans[6][i];
540 s_SPtrans[7][i] = c_SPtrans[7][i];
542 s_skb[0][i] = c_skb[0][i];
543 s_skb[1][i] = c_skb[1][i];
544 s_skb[2][i] = c_skb[2][i];
545 s_skb[3][i] = c_skb[3][i];
546 s_skb[4][i] = c_skb[4][i];
547 s_skb[5][i] = c_skb[5][i];
548 s_skb[6][i] = c_skb[6][i];
549 s_skb[7][i] = c_skb[7][i];
552 barrier (CLK_LOCAL_MEM_FENCE);
554 if (gid >= gid_max) return;
560 const u32 mask = salt_bufs[salt_pos].salt_buf[0];
566 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
568 const u32 pw_r_len = combs_buf[il_pos].pw_len;
570 u32 pw_len = pw_l_len + pw_r_len;
572 pw_len = (pw_len >= 8) ? 8 : pw_len;
576 wordr0[0] = combs_buf[il_pos].i[0];
577 wordr0[1] = combs_buf[il_pos].i[1];
602 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
604 switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
609 w0[0] = wordl0[0] | wordr0[0];
610 w0[1] = wordl0[1] | wordr0[1];
637 data[0] = (w0[0] << 1) & 0xfefefefe;
638 data[1] = (w0[1] << 1) & 0xfefefefe;
643 _des_crypt_keysetup (data[0], data[1], Kc, Kd, s_skb);
647 _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
649 const u32 r0 = iv[0];
650 const u32 r1 = iv[1];
658 __kernel void m01500_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
662 __kernel void m01500_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
666 __kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
672 const u32 gid = get_global_id (0);
673 const u32 lid = get_local_id (0);
674 const u32 lsz = get_local_size (0);
682 wordl0[0] = pws[gid].i[ 0];
683 wordl0[1] = pws[gid].i[ 1];
708 const u32 pw_l_len = pws[gid].pw_len;
710 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
712 switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
719 __local u32 s_SPtrans[8][64];
720 __local u32 s_skb[8][64];
722 for (u32 i = lid; i < 64; i += lsz)
724 s_SPtrans[0][i] = c_SPtrans[0][i];
725 s_SPtrans[1][i] = c_SPtrans[1][i];
726 s_SPtrans[2][i] = c_SPtrans[2][i];
727 s_SPtrans[3][i] = c_SPtrans[3][i];
728 s_SPtrans[4][i] = c_SPtrans[4][i];
729 s_SPtrans[5][i] = c_SPtrans[5][i];
730 s_SPtrans[6][i] = c_SPtrans[6][i];
731 s_SPtrans[7][i] = c_SPtrans[7][i];
733 s_skb[0][i] = c_skb[0][i];
734 s_skb[1][i] = c_skb[1][i];
735 s_skb[2][i] = c_skb[2][i];
736 s_skb[3][i] = c_skb[3][i];
737 s_skb[4][i] = c_skb[4][i];
738 s_skb[5][i] = c_skb[5][i];
739 s_skb[6][i] = c_skb[6][i];
740 s_skb[7][i] = c_skb[7][i];
743 barrier (CLK_LOCAL_MEM_FENCE);
745 if (gid >= gid_max) return;
751 const u32 mask = salt_bufs[salt_pos].salt_buf[0];
757 const u32 search[4] =
759 digests_buf[digests_offset].digest_buf[DGST_R0],
760 digests_buf[digests_offset].digest_buf[DGST_R1],
761 digests_buf[digests_offset].digest_buf[DGST_R2],
762 digests_buf[digests_offset].digest_buf[DGST_R3]
769 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
771 const u32 pw_r_len = combs_buf[il_pos].pw_len;
773 u32 pw_len = pw_l_len + pw_r_len;
775 pw_len = (pw_len >= 8) ? 8 : pw_len;
779 wordr0[0] = combs_buf[il_pos].i[0];
780 wordr0[1] = combs_buf[il_pos].i[1];
805 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
807 switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
812 w0[0] = wordl0[0] | wordr0[0];
813 w0[1] = wordl0[1] | wordr0[1];
840 data[0] = (w0[0] << 1) & 0xfefefefe;
841 data[1] = (w0[1] << 1) & 0xfefefefe;
846 _des_crypt_keysetup (data[0], data[1], Kc, Kd, s_skb);
850 _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
852 const u32 r0 = iv[0];
853 const u32 r1 = iv[1];
861 __kernel void m01500_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
865 __kernel void m01500_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)