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 0x02080800, 0x00080000, 0x02000002, 0x02080802,
47 0x02000000, 0x00080802, 0x00080002, 0x02000002,
48 0x00080802, 0x02080800, 0x02080000, 0x00000802,
49 0x02000802, 0x02000000, 0x00000000, 0x00080002,
50 0x00080000, 0x00000002, 0x02000800, 0x00080800,
51 0x02080802, 0x02080000, 0x00000802, 0x02000800,
52 0x00000002, 0x00000800, 0x00080800, 0x02080002,
53 0x00000800, 0x02000802, 0x02080002, 0x00000000,
54 0x00000000, 0x02080802, 0x02000800, 0x00080002,
55 0x02080800, 0x00080000, 0x00000802, 0x02000800,
56 0x02080002, 0x00000800, 0x00080800, 0x02000002,
57 0x00080802, 0x00000002, 0x02000002, 0x02080000,
58 0x02080802, 0x00080800, 0x02080000, 0x02000802,
59 0x02000000, 0x00000802, 0x00080002, 0x00000000,
60 0x00080000, 0x02000000, 0x02000802, 0x02080800,
61 0x00000002, 0x02080002, 0x00000800, 0x00080802,
64 0x40108010, 0x00000000, 0x00108000, 0x40100000,
65 0x40000010, 0x00008010, 0x40008000, 0x00108000,
66 0x00008000, 0x40100010, 0x00000010, 0x40008000,
67 0x00100010, 0x40108000, 0x40100000, 0x00000010,
68 0x00100000, 0x40008010, 0x40100010, 0x00008000,
69 0x00108010, 0x40000000, 0x00000000, 0x00100010,
70 0x40008010, 0x00108010, 0x40108000, 0x40000010,
71 0x40000000, 0x00100000, 0x00008010, 0x40108010,
72 0x00100010, 0x40108000, 0x40008000, 0x00108010,
73 0x40108010, 0x00100010, 0x40000010, 0x00000000,
74 0x40000000, 0x00008010, 0x00100000, 0x40100010,
75 0x00008000, 0x40000000, 0x00108010, 0x40008010,
76 0x40108000, 0x00008000, 0x00000000, 0x40000010,
77 0x00000010, 0x40108010, 0x00108000, 0x40100000,
78 0x40100010, 0x00100000, 0x00008010, 0x40008000,
79 0x40008010, 0x00000010, 0x40100000, 0x00108000,
82 0x04000001, 0x04040100, 0x00000100, 0x04000101,
83 0x00040001, 0x04000000, 0x04000101, 0x00040100,
84 0x04000100, 0x00040000, 0x04040000, 0x00000001,
85 0x04040101, 0x00000101, 0x00000001, 0x04040001,
86 0x00000000, 0x00040001, 0x04040100, 0x00000100,
87 0x00000101, 0x04040101, 0x00040000, 0x04000001,
88 0x04040001, 0x04000100, 0x00040101, 0x04040000,
89 0x00040100, 0x00000000, 0x04000000, 0x00040101,
90 0x04040100, 0x00000100, 0x00000001, 0x00040000,
91 0x00000101, 0x00040001, 0x04040000, 0x04000101,
92 0x00000000, 0x04040100, 0x00040100, 0x04040001,
93 0x00040001, 0x04000000, 0x04040101, 0x00000001,
94 0x00040101, 0x04000001, 0x04000000, 0x04040101,
95 0x00040000, 0x04000100, 0x04000101, 0x00040100,
96 0x04000100, 0x00000000, 0x04040001, 0x00000101,
97 0x04000001, 0x00040101, 0x00000100, 0x04040000,
100 0x00401008, 0x10001000, 0x00000008, 0x10401008,
101 0x00000000, 0x10400000, 0x10001008, 0x00400008,
102 0x10401000, 0x10000008, 0x10000000, 0x00001008,
103 0x10000008, 0x00401008, 0x00400000, 0x10000000,
104 0x10400008, 0x00401000, 0x00001000, 0x00000008,
105 0x00401000, 0x10001008, 0x10400000, 0x00001000,
106 0x00001008, 0x00000000, 0x00400008, 0x10401000,
107 0x10001000, 0x10400008, 0x10401008, 0x00400000,
108 0x10400008, 0x00001008, 0x00400000, 0x10000008,
109 0x00401000, 0x10001000, 0x00000008, 0x10400000,
110 0x10001008, 0x00000000, 0x00001000, 0x00400008,
111 0x00000000, 0x10400008, 0x10401000, 0x00001000,
112 0x10000000, 0x10401008, 0x00401008, 0x00400000,
113 0x10401008, 0x00000008, 0x10001000, 0x00401008,
114 0x00400008, 0x00401000, 0x10400000, 0x10001008,
115 0x00001008, 0x10000000, 0x10000008, 0x10401000,
118 0x08000000, 0x00010000, 0x00000400, 0x08010420,
119 0x08010020, 0x08000400, 0x00010420, 0x08010000,
120 0x00010000, 0x00000020, 0x08000020, 0x00010400,
121 0x08000420, 0x08010020, 0x08010400, 0x00000000,
122 0x00010400, 0x08000000, 0x00010020, 0x00000420,
123 0x08000400, 0x00010420, 0x00000000, 0x08000020,
124 0x00000020, 0x08000420, 0x08010420, 0x00010020,
125 0x08010000, 0x00000400, 0x00000420, 0x08010400,
126 0x08010400, 0x08000420, 0x00010020, 0x08010000,
127 0x00010000, 0x00000020, 0x08000020, 0x08000400,
128 0x08000000, 0x00010400, 0x08010420, 0x00000000,
129 0x00010420, 0x08000000, 0x00000400, 0x00010020,
130 0x08000420, 0x00000400, 0x00000000, 0x08010420,
131 0x08010020, 0x08010400, 0x00000420, 0x00010000,
132 0x00010400, 0x08010020, 0x08000400, 0x00000420,
133 0x00000020, 0x00010420, 0x08010000, 0x08000020,
136 0x80000040, 0x00200040, 0x00000000, 0x80202000,
137 0x00200040, 0x00002000, 0x80002040, 0x00200000,
138 0x00002040, 0x80202040, 0x00202000, 0x80000000,
139 0x80002000, 0x80000040, 0x80200000, 0x00202040,
140 0x00200000, 0x80002040, 0x80200040, 0x00000000,
141 0x00002000, 0x00000040, 0x80202000, 0x80200040,
142 0x80202040, 0x80200000, 0x80000000, 0x00002040,
143 0x00000040, 0x00202000, 0x00202040, 0x80002000,
144 0x00002040, 0x80000000, 0x80002000, 0x00202040,
145 0x80202000, 0x00200040, 0x00000000, 0x80002000,
146 0x80000000, 0x00002000, 0x80200040, 0x00200000,
147 0x00200040, 0x80202040, 0x00202000, 0x00000040,
148 0x80202040, 0x00202000, 0x00200000, 0x80002040,
149 0x80000040, 0x80200000, 0x00202040, 0x00000000,
150 0x00002000, 0x80000040, 0x80002040, 0x80202000,
151 0x80200000, 0x00002040, 0x00000040, 0x80200040,
154 0x00004000, 0x00000200, 0x01000200, 0x01000004,
155 0x01004204, 0x00004004, 0x00004200, 0x00000000,
156 0x01000000, 0x01000204, 0x00000204, 0x01004000,
157 0x00000004, 0x01004200, 0x01004000, 0x00000204,
158 0x01000204, 0x00004000, 0x00004004, 0x01004204,
159 0x00000000, 0x01000200, 0x01000004, 0x00004200,
160 0x01004004, 0x00004204, 0x01004200, 0x00000004,
161 0x00004204, 0x01004004, 0x00000200, 0x01000000,
162 0x00004204, 0x01004000, 0x01004004, 0x00000204,
163 0x00004000, 0x00000200, 0x01000000, 0x01004004,
164 0x01000204, 0x00004204, 0x00004200, 0x00000000,
165 0x00000200, 0x01000004, 0x00000004, 0x01000200,
166 0x00000000, 0x01000204, 0x01000200, 0x00004200,
167 0x00000204, 0x00004000, 0x01004204, 0x01000000,
168 0x01004200, 0x00000004, 0x00004004, 0x01004204,
169 0x01000004, 0x01004200, 0x01004000, 0x00004004,
172 0x20800080, 0x20820000, 0x00020080, 0x00000000,
173 0x20020000, 0x00800080, 0x20800000, 0x20820080,
174 0x00000080, 0x20000000, 0x00820000, 0x00020080,
175 0x00820080, 0x20020080, 0x20000080, 0x20800000,
176 0x00020000, 0x00820080, 0x00800080, 0x20020000,
177 0x20820080, 0x20000080, 0x00000000, 0x00820000,
178 0x20000000, 0x00800000, 0x20020080, 0x20800080,
179 0x00800000, 0x00020000, 0x20820000, 0x00000080,
180 0x00800000, 0x00020000, 0x20000080, 0x20820080,
181 0x00020080, 0x20000000, 0x00000000, 0x00820000,
182 0x20800080, 0x20020080, 0x20020000, 0x00800080,
183 0x20820000, 0x00000080, 0x00800080, 0x20020000,
184 0x20820080, 0x00800000, 0x20800000, 0x20000080,
185 0x00820000, 0x00020080, 0x20020080, 0x20800000,
186 0x00000080, 0x20820000, 0x00820080, 0x00000000,
187 0x20000000, 0x20800080, 0x00020000, 0x00820080,
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 LM_IV_0_IP_RR3 0x2400b807
340 #define LM_IV_1_IP_RR3 0xaa190747
342 #define BOX(i,n,S) (S)[(n)][(i)]
344 static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
350 for (u32 i = 0; i < 16; i += 2)
355 u = Kc[i + 0] ^ rotl32 (r, 30u);
356 t = Kd[i + 0] ^ rotl32 (r, 26u);
358 l ^= BOX (((u >> 0) & 0x3f), 0, s_SPtrans)
359 | BOX (((u >> 8) & 0x3f), 2, s_SPtrans)
360 | BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
361 | BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
362 | BOX (((t >> 0) & 0x3f), 1, s_SPtrans)
363 | BOX (((t >> 8) & 0x3f), 3, s_SPtrans)
364 | BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
365 | BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
367 u = Kc[i + 1] ^ rotl32 (l, 30u);
368 t = Kd[i + 1] ^ rotl32 (l, 26u);
370 r ^= BOX (((u >> 0) & 0x3f), 0, s_SPtrans)
371 | BOX (((u >> 8) & 0x3f), 2, s_SPtrans)
372 | BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
373 | BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
374 | BOX (((t >> 0) & 0x3f), 1, s_SPtrans)
375 | BOX (((t >> 8) & 0x3f), 3, s_SPtrans)
376 | BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
377 | BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
380 iv[0] = rotl32 (l, 29);
381 iv[1] = rotl32 (r, 29);
384 static void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 s_skb[8][64])
388 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
389 HPERM_OP (c, tt, 2, 0xcccc0000);
390 HPERM_OP (d, tt, 2, 0xcccc0000);
391 PERM_OP (d, c, tt, 1, 0x55555555);
392 PERM_OP (c, d, tt, 8, 0x00ff00ff);
393 PERM_OP (d, c, tt, 1, 0x55555555);
395 d = ((d & 0x000000ff) << 16)
396 | ((d & 0x0000ff00) << 0)
397 | ((d & 0x00ff0000) >> 16)
398 | ((c & 0xf0000000) >> 4);
403 for (u32 i = 0; i < 16; i++)
405 if ((i < 2) || (i == 8) || (i == 15))
407 c = ((c >> 1) | (c << 27));
408 d = ((d >> 1) | (d << 27));
412 c = ((c >> 2) | (c << 26));
413 d = ((d >> 2) | (d << 26));
419 const u32 c00 = (c >> 0) & 0x0000003f;
420 const u32 c06 = (c >> 6) & 0x00383003;
421 const u32 c07 = (c >> 7) & 0x0000003c;
422 const u32 c13 = (c >> 13) & 0x0000060f;
423 const u32 c20 = (c >> 20) & 0x00000001;
425 u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
426 | BOX (((c06 >> 0) & 0xff)
427 |((c07 >> 0) & 0xff), 1, s_skb)
428 | BOX (((c13 >> 0) & 0xff)
429 |((c06 >> 8) & 0xff), 2, s_skb)
430 | BOX (((c20 >> 0) & 0xff)
432 |((c06 >> 16) & 0xff), 3, s_skb);
434 const u32 d00 = (d >> 0) & 0x00003c3f;
435 const u32 d07 = (d >> 7) & 0x00003f03;
436 const u32 d21 = (d >> 21) & 0x0000000f;
437 const u32 d22 = (d >> 22) & 0x00000030;
439 u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
440 | BOX (((d07 >> 0) & 0xff)
441 |((d00 >> 8) & 0xff), 5, s_skb)
442 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
443 | BOX (((d21 >> 0) & 0xff)
444 |((d22 >> 0) & 0xff), 7, s_skb);
446 Kc[i] = ((t << 16) | (s & 0x0000ffff));
447 Kd[i] = ((s >> 16) | (t & 0xffff0000));
451 static void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
453 const uchar4 t0 = as_uchar4 (w0);
454 const uchar4 t1 = as_uchar4 (w1);
459 k0.s0 = (t0.s0 >> 0);
460 k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
461 k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
462 k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
463 k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
464 k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
465 k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
466 k1.s3 = (t1.s2 << 1);
468 out[0] = as_uint (k0);
469 out[1] = as_uint (k1);
472 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03000_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
478 const u32 lid = get_local_id (0);
484 const u32 gid = get_global_id (0);
488 wordl0[0] = pws[gid].i[ 0];
489 wordl0[1] = pws[gid].i[ 1];
514 const u32 pw_l_len = pws[gid].pw_len;
516 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
518 switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
525 __local u32 s_SPtrans[8][64];
527 __local u32 s_skb[8][64];
529 s_SPtrans[0][lid] = c_SPtrans[0][lid];
530 s_SPtrans[1][lid] = c_SPtrans[1][lid];
531 s_SPtrans[2][lid] = c_SPtrans[2][lid];
532 s_SPtrans[3][lid] = c_SPtrans[3][lid];
533 s_SPtrans[4][lid] = c_SPtrans[4][lid];
534 s_SPtrans[5][lid] = c_SPtrans[5][lid];
535 s_SPtrans[6][lid] = c_SPtrans[6][lid];
536 s_SPtrans[7][lid] = c_SPtrans[7][lid];
538 s_skb[0][lid] = c_skb[0][lid];
539 s_skb[1][lid] = c_skb[1][lid];
540 s_skb[2][lid] = c_skb[2][lid];
541 s_skb[3][lid] = c_skb[3][lid];
542 s_skb[4][lid] = c_skb[4][lid];
543 s_skb[5][lid] = c_skb[5][lid];
544 s_skb[6][lid] = c_skb[6][lid];
545 s_skb[7][lid] = c_skb[7][lid];
547 barrier (CLK_LOCAL_MEM_FENCE);
549 if (gid >= gid_max) return;
555 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
557 const u32 pw_r_len = combs_buf[il_pos].pw_len;
559 u32 pw_len = pw_l_len + pw_r_len;
561 pw_len = (pw_len >= 7) ? 7 : pw_len;
565 wordr0[0] = combs_buf[il_pos].i[0];
566 wordr0[1] = combs_buf[il_pos].i[1];
591 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
593 switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
598 w0[0] = wordl0[0] | wordr0[0];
599 w0[1] = wordl0[1] | wordr0[1];
626 transform_netntlmv1_key (w0[0], w0[1], key);
628 const u32 c = key[0];
629 const u32 d = key[1];
634 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
638 data[0] = LM_IV_0_IP_RR3;
639 data[1] = LM_IV_1_IP_RR3;
643 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
645 const u32 r0 = iv[0];
646 const u32 r1 = iv[1];
654 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03000_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)
658 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03000_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)
662 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03000_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
668 const u32 lid = get_local_id (0);
674 const u32 gid = get_global_id (0);
678 wordl0[0] = pws[gid].i[ 0];
679 wordl0[1] = pws[gid].i[ 1];
704 const u32 pw_l_len = pws[gid].pw_len;
706 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
708 switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
715 __local u32 s_SPtrans[8][64];
717 __local u32 s_skb[8][64];
719 s_SPtrans[0][lid] = c_SPtrans[0][lid];
720 s_SPtrans[1][lid] = c_SPtrans[1][lid];
721 s_SPtrans[2][lid] = c_SPtrans[2][lid];
722 s_SPtrans[3][lid] = c_SPtrans[3][lid];
723 s_SPtrans[4][lid] = c_SPtrans[4][lid];
724 s_SPtrans[5][lid] = c_SPtrans[5][lid];
725 s_SPtrans[6][lid] = c_SPtrans[6][lid];
726 s_SPtrans[7][lid] = c_SPtrans[7][lid];
728 s_skb[0][lid] = c_skb[0][lid];
729 s_skb[1][lid] = c_skb[1][lid];
730 s_skb[2][lid] = c_skb[2][lid];
731 s_skb[3][lid] = c_skb[3][lid];
732 s_skb[4][lid] = c_skb[4][lid];
733 s_skb[5][lid] = c_skb[5][lid];
734 s_skb[6][lid] = c_skb[6][lid];
735 s_skb[7][lid] = c_skb[7][lid];
737 barrier (CLK_LOCAL_MEM_FENCE);
739 if (gid >= gid_max) return;
745 const u32 search[4] =
747 digests_buf[digests_offset].digest_buf[DGST_R0],
748 digests_buf[digests_offset].digest_buf[DGST_R1],
749 digests_buf[digests_offset].digest_buf[DGST_R2],
750 digests_buf[digests_offset].digest_buf[DGST_R3]
757 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
759 const u32 pw_r_len = combs_buf[il_pos].pw_len;
761 u32 pw_len = pw_l_len + pw_r_len;
763 pw_len = (pw_len >= 7) ? 7 : pw_len;
767 wordr0[0] = combs_buf[il_pos].i[0];
768 wordr0[1] = combs_buf[il_pos].i[1];
793 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
795 switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
800 w0[0] = wordl0[0] | wordr0[0];
801 w0[1] = wordl0[1] | wordr0[1];
828 transform_netntlmv1_key (w0[0], w0[1], key);
830 const u32 c = key[0];
831 const u32 d = key[1];
836 _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
840 data[0] = LM_IV_0_IP_RR3;
841 data[1] = LM_IV_1_IP_RR3;
845 _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
847 const u32 r0 = iv[0];
848 const u32 r1 = iv[1];
856 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03000_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)
860 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03000_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)