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"
20 #define COMPARE_S "check_single_comp4.c"
21 #define COMPARE_M "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,
63 0x40108010, 0x00000000, 0x00108000, 0x40100000,
64 0x40000010, 0x00008010, 0x40008000, 0x00108000,
65 0x00008000, 0x40100010, 0x00000010, 0x40008000,
66 0x00100010, 0x40108000, 0x40100000, 0x00000010,
67 0x00100000, 0x40008010, 0x40100010, 0x00008000,
68 0x00108010, 0x40000000, 0x00000000, 0x00100010,
69 0x40008010, 0x00108010, 0x40108000, 0x40000010,
70 0x40000000, 0x00100000, 0x00008010, 0x40108010,
71 0x00100010, 0x40108000, 0x40008000, 0x00108010,
72 0x40108010, 0x00100010, 0x40000010, 0x00000000,
73 0x40000000, 0x00008010, 0x00100000, 0x40100010,
74 0x00008000, 0x40000000, 0x00108010, 0x40008010,
75 0x40108000, 0x00008000, 0x00000000, 0x40000010,
76 0x00000010, 0x40108010, 0x00108000, 0x40100000,
77 0x40100010, 0x00100000, 0x00008010, 0x40008000,
78 0x40008010, 0x00000010, 0x40100000, 0x00108000,
80 0x04000001, 0x04040100, 0x00000100, 0x04000101,
81 0x00040001, 0x04000000, 0x04000101, 0x00040100,
82 0x04000100, 0x00040000, 0x04040000, 0x00000001,
83 0x04040101, 0x00000101, 0x00000001, 0x04040001,
84 0x00000000, 0x00040001, 0x04040100, 0x00000100,
85 0x00000101, 0x04040101, 0x00040000, 0x04000001,
86 0x04040001, 0x04000100, 0x00040101, 0x04040000,
87 0x00040100, 0x00000000, 0x04000000, 0x00040101,
88 0x04040100, 0x00000100, 0x00000001, 0x00040000,
89 0x00000101, 0x00040001, 0x04040000, 0x04000101,
90 0x00000000, 0x04040100, 0x00040100, 0x04040001,
91 0x00040001, 0x04000000, 0x04040101, 0x00000001,
92 0x00040101, 0x04000001, 0x04000000, 0x04040101,
93 0x00040000, 0x04000100, 0x04000101, 0x00040100,
94 0x04000100, 0x00000000, 0x04040001, 0x00000101,
95 0x04000001, 0x00040101, 0x00000100, 0x04040000,
97 0x00401008, 0x10001000, 0x00000008, 0x10401008,
98 0x00000000, 0x10400000, 0x10001008, 0x00400008,
99 0x10401000, 0x10000008, 0x10000000, 0x00001008,
100 0x10000008, 0x00401008, 0x00400000, 0x10000000,
101 0x10400008, 0x00401000, 0x00001000, 0x00000008,
102 0x00401000, 0x10001008, 0x10400000, 0x00001000,
103 0x00001008, 0x00000000, 0x00400008, 0x10401000,
104 0x10001000, 0x10400008, 0x10401008, 0x00400000,
105 0x10400008, 0x00001008, 0x00400000, 0x10000008,
106 0x00401000, 0x10001000, 0x00000008, 0x10400000,
107 0x10001008, 0x00000000, 0x00001000, 0x00400008,
108 0x00000000, 0x10400008, 0x10401000, 0x00001000,
109 0x10000000, 0x10401008, 0x00401008, 0x00400000,
110 0x10401008, 0x00000008, 0x10001000, 0x00401008,
111 0x00400008, 0x00401000, 0x10400000, 0x10001008,
112 0x00001008, 0x10000000, 0x10000008, 0x10401000,
114 0x08000000, 0x00010000, 0x00000400, 0x08010420,
115 0x08010020, 0x08000400, 0x00010420, 0x08010000,
116 0x00010000, 0x00000020, 0x08000020, 0x00010400,
117 0x08000420, 0x08010020, 0x08010400, 0x00000000,
118 0x00010400, 0x08000000, 0x00010020, 0x00000420,
119 0x08000400, 0x00010420, 0x00000000, 0x08000020,
120 0x00000020, 0x08000420, 0x08010420, 0x00010020,
121 0x08010000, 0x00000400, 0x00000420, 0x08010400,
122 0x08010400, 0x08000420, 0x00010020, 0x08010000,
123 0x00010000, 0x00000020, 0x08000020, 0x08000400,
124 0x08000000, 0x00010400, 0x08010420, 0x00000000,
125 0x00010420, 0x08000000, 0x00000400, 0x00010020,
126 0x08000420, 0x00000400, 0x00000000, 0x08010420,
127 0x08010020, 0x08010400, 0x00000420, 0x00010000,
128 0x00010400, 0x08010020, 0x08000400, 0x00000420,
129 0x00000020, 0x00010420, 0x08010000, 0x08000020,
131 0x80000040, 0x00200040, 0x00000000, 0x80202000,
132 0x00200040, 0x00002000, 0x80002040, 0x00200000,
133 0x00002040, 0x80202040, 0x00202000, 0x80000000,
134 0x80002000, 0x80000040, 0x80200000, 0x00202040,
135 0x00200000, 0x80002040, 0x80200040, 0x00000000,
136 0x00002000, 0x00000040, 0x80202000, 0x80200040,
137 0x80202040, 0x80200000, 0x80000000, 0x00002040,
138 0x00000040, 0x00202000, 0x00202040, 0x80002000,
139 0x00002040, 0x80000000, 0x80002000, 0x00202040,
140 0x80202000, 0x00200040, 0x00000000, 0x80002000,
141 0x80000000, 0x00002000, 0x80200040, 0x00200000,
142 0x00200040, 0x80202040, 0x00202000, 0x00000040,
143 0x80202040, 0x00202000, 0x00200000, 0x80002040,
144 0x80000040, 0x80200000, 0x00202040, 0x00000000,
145 0x00002000, 0x80000040, 0x80002040, 0x80202000,
146 0x80200000, 0x00002040, 0x00000040, 0x80200040,
148 0x00004000, 0x00000200, 0x01000200, 0x01000004,
149 0x01004204, 0x00004004, 0x00004200, 0x00000000,
150 0x01000000, 0x01000204, 0x00000204, 0x01004000,
151 0x00000004, 0x01004200, 0x01004000, 0x00000204,
152 0x01000204, 0x00004000, 0x00004004, 0x01004204,
153 0x00000000, 0x01000200, 0x01000004, 0x00004200,
154 0x01004004, 0x00004204, 0x01004200, 0x00000004,
155 0x00004204, 0x01004004, 0x00000200, 0x01000000,
156 0x00004204, 0x01004000, 0x01004004, 0x00000204,
157 0x00004000, 0x00000200, 0x01000000, 0x01004004,
158 0x01000204, 0x00004204, 0x00004200, 0x00000000,
159 0x00000200, 0x01000004, 0x00000004, 0x01000200,
160 0x00000000, 0x01000204, 0x01000200, 0x00004200,
161 0x00000204, 0x00004000, 0x01004204, 0x01000000,
162 0x01004200, 0x00000004, 0x00004004, 0x01004204,
163 0x01000004, 0x01004200, 0x01004000, 0x00004004,
165 0x20800080, 0x20820000, 0x00020080, 0x00000000,
166 0x20020000, 0x00800080, 0x20800000, 0x20820080,
167 0x00000080, 0x20000000, 0x00820000, 0x00020080,
168 0x00820080, 0x20020080, 0x20000080, 0x20800000,
169 0x00020000, 0x00820080, 0x00800080, 0x20020000,
170 0x20820080, 0x20000080, 0x00000000, 0x00820000,
171 0x20000000, 0x00800000, 0x20020080, 0x20800080,
172 0x00800000, 0x00020000, 0x20820000, 0x00000080,
173 0x00800000, 0x00020000, 0x20000080, 0x20820080,
174 0x00020080, 0x20000000, 0x00000000, 0x00820000,
175 0x20800080, 0x20020080, 0x20020000, 0x00800080,
176 0x20820000, 0x00000080, 0x00800080, 0x20020000,
177 0x20820080, 0x00800000, 0x20800000, 0x20000080,
178 0x00820000, 0x00020080, 0x20020080, 0x20800000,
179 0x00000080, 0x20820000, 0x00820080, 0x00000000,
180 0x20000000, 0x20800080, 0x00020000, 0x00820080,
183 __constant u32 c_skb[8][64] =
185 /* for C bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
186 0x00000000, 0x00000010, 0x20000000, 0x20000010,
187 0x00010000, 0x00010010, 0x20010000, 0x20010010,
188 0x00000800, 0x00000810, 0x20000800, 0x20000810,
189 0x00010800, 0x00010810, 0x20010800, 0x20010810,
190 0x00000020, 0x00000030, 0x20000020, 0x20000030,
191 0x00010020, 0x00010030, 0x20010020, 0x20010030,
192 0x00000820, 0x00000830, 0x20000820, 0x20000830,
193 0x00010820, 0x00010830, 0x20010820, 0x20010830,
194 0x00080000, 0x00080010, 0x20080000, 0x20080010,
195 0x00090000, 0x00090010, 0x20090000, 0x20090010,
196 0x00080800, 0x00080810, 0x20080800, 0x20080810,
197 0x00090800, 0x00090810, 0x20090800, 0x20090810,
198 0x00080020, 0x00080030, 0x20080020, 0x20080030,
199 0x00090020, 0x00090030, 0x20090020, 0x20090030,
200 0x00080820, 0x00080830, 0x20080820, 0x20080830,
201 0x00090820, 0x00090830, 0x20090820, 0x20090830,
202 /* for C bits (numbered as per FIPS 46) 7 8 10 11 12 13 */
203 0x00000000, 0x02000000, 0x00002000, 0x02002000,
204 0x00200000, 0x02200000, 0x00202000, 0x02202000,
205 0x00000004, 0x02000004, 0x00002004, 0x02002004,
206 0x00200004, 0x02200004, 0x00202004, 0x02202004,
207 0x00000400, 0x02000400, 0x00002400, 0x02002400,
208 0x00200400, 0x02200400, 0x00202400, 0x02202400,
209 0x00000404, 0x02000404, 0x00002404, 0x02002404,
210 0x00200404, 0x02200404, 0x00202404, 0x02202404,
211 0x10000000, 0x12000000, 0x10002000, 0x12002000,
212 0x10200000, 0x12200000, 0x10202000, 0x12202000,
213 0x10000004, 0x12000004, 0x10002004, 0x12002004,
214 0x10200004, 0x12200004, 0x10202004, 0x12202004,
215 0x10000400, 0x12000400, 0x10002400, 0x12002400,
216 0x10200400, 0x12200400, 0x10202400, 0x12202400,
217 0x10000404, 0x12000404, 0x10002404, 0x12002404,
218 0x10200404, 0x12200404, 0x10202404, 0x12202404,
219 /* for C bits (numbered as per FIPS 46) 14 15 16 17 19 20 */
220 0x00000000, 0x00000001, 0x00040000, 0x00040001,
221 0x01000000, 0x01000001, 0x01040000, 0x01040001,
222 0x00000002, 0x00000003, 0x00040002, 0x00040003,
223 0x01000002, 0x01000003, 0x01040002, 0x01040003,
224 0x00000200, 0x00000201, 0x00040200, 0x00040201,
225 0x01000200, 0x01000201, 0x01040200, 0x01040201,
226 0x00000202, 0x00000203, 0x00040202, 0x00040203,
227 0x01000202, 0x01000203, 0x01040202, 0x01040203,
228 0x08000000, 0x08000001, 0x08040000, 0x08040001,
229 0x09000000, 0x09000001, 0x09040000, 0x09040001,
230 0x08000002, 0x08000003, 0x08040002, 0x08040003,
231 0x09000002, 0x09000003, 0x09040002, 0x09040003,
232 0x08000200, 0x08000201, 0x08040200, 0x08040201,
233 0x09000200, 0x09000201, 0x09040200, 0x09040201,
234 0x08000202, 0x08000203, 0x08040202, 0x08040203,
235 0x09000202, 0x09000203, 0x09040202, 0x09040203,
236 /* for C bits (numbered as per FIPS 46) 21 23 24 26 27 28 */
237 0x00000000, 0x00100000, 0x00000100, 0x00100100,
238 0x00000008, 0x00100008, 0x00000108, 0x00100108,
239 0x00001000, 0x00101000, 0x00001100, 0x00101100,
240 0x00001008, 0x00101008, 0x00001108, 0x00101108,
241 0x04000000, 0x04100000, 0x04000100, 0x04100100,
242 0x04000008, 0x04100008, 0x04000108, 0x04100108,
243 0x04001000, 0x04101000, 0x04001100, 0x04101100,
244 0x04001008, 0x04101008, 0x04001108, 0x04101108,
245 0x00020000, 0x00120000, 0x00020100, 0x00120100,
246 0x00020008, 0x00120008, 0x00020108, 0x00120108,
247 0x00021000, 0x00121000, 0x00021100, 0x00121100,
248 0x00021008, 0x00121008, 0x00021108, 0x00121108,
249 0x04020000, 0x04120000, 0x04020100, 0x04120100,
250 0x04020008, 0x04120008, 0x04020108, 0x04120108,
251 0x04021000, 0x04121000, 0x04021100, 0x04121100,
252 0x04021008, 0x04121008, 0x04021108, 0x04121108,
253 /* for D bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
254 0x00000000, 0x10000000, 0x00010000, 0x10010000,
255 0x00000004, 0x10000004, 0x00010004, 0x10010004,
256 0x20000000, 0x30000000, 0x20010000, 0x30010000,
257 0x20000004, 0x30000004, 0x20010004, 0x30010004,
258 0x00100000, 0x10100000, 0x00110000, 0x10110000,
259 0x00100004, 0x10100004, 0x00110004, 0x10110004,
260 0x20100000, 0x30100000, 0x20110000, 0x30110000,
261 0x20100004, 0x30100004, 0x20110004, 0x30110004,
262 0x00001000, 0x10001000, 0x00011000, 0x10011000,
263 0x00001004, 0x10001004, 0x00011004, 0x10011004,
264 0x20001000, 0x30001000, 0x20011000, 0x30011000,
265 0x20001004, 0x30001004, 0x20011004, 0x30011004,
266 0x00101000, 0x10101000, 0x00111000, 0x10111000,
267 0x00101004, 0x10101004, 0x00111004, 0x10111004,
268 0x20101000, 0x30101000, 0x20111000, 0x30111000,
269 0x20101004, 0x30101004, 0x20111004, 0x30111004,
270 /* for D bits (numbered as per FIPS 46) 8 9 11 12 13 14 */
271 0x00000000, 0x08000000, 0x00000008, 0x08000008,
272 0x00000400, 0x08000400, 0x00000408, 0x08000408,
273 0x00020000, 0x08020000, 0x00020008, 0x08020008,
274 0x00020400, 0x08020400, 0x00020408, 0x08020408,
275 0x00000001, 0x08000001, 0x00000009, 0x08000009,
276 0x00000401, 0x08000401, 0x00000409, 0x08000409,
277 0x00020001, 0x08020001, 0x00020009, 0x08020009,
278 0x00020401, 0x08020401, 0x00020409, 0x08020409,
279 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
280 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
281 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
282 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
283 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
284 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
285 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
286 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
287 /* for D bits (numbered as per FIPS 46) 16 17 18 19 20 21 */
288 0x00000000, 0x00000100, 0x00080000, 0x00080100,
289 0x01000000, 0x01000100, 0x01080000, 0x01080100,
290 0x00000010, 0x00000110, 0x00080010, 0x00080110,
291 0x01000010, 0x01000110, 0x01080010, 0x01080110,
292 0x00200000, 0x00200100, 0x00280000, 0x00280100,
293 0x01200000, 0x01200100, 0x01280000, 0x01280100,
294 0x00200010, 0x00200110, 0x00280010, 0x00280110,
295 0x01200010, 0x01200110, 0x01280010, 0x01280110,
296 0x00000200, 0x00000300, 0x00080200, 0x00080300,
297 0x01000200, 0x01000300, 0x01080200, 0x01080300,
298 0x00000210, 0x00000310, 0x00080210, 0x00080310,
299 0x01000210, 0x01000310, 0x01080210, 0x01080310,
300 0x00200200, 0x00200300, 0x00280200, 0x00280300,
301 0x01200200, 0x01200300, 0x01280200, 0x01280300,
302 0x00200210, 0x00200310, 0x00280210, 0x00280310,
303 0x01200210, 0x01200310, 0x01280210, 0x01280310,
304 /* for D bits (numbered as per FIPS 46) 22 23 24 25 27 28 */
305 0x00000000, 0x04000000, 0x00040000, 0x04040000,
306 0x00000002, 0x04000002, 0x00040002, 0x04040002,
307 0x00002000, 0x04002000, 0x00042000, 0x04042000,
308 0x00002002, 0x04002002, 0x00042002, 0x04042002,
309 0x00000020, 0x04000020, 0x00040020, 0x04040020,
310 0x00000022, 0x04000022, 0x00040022, 0x04040022,
311 0x00002020, 0x04002020, 0x00042020, 0x04042020,
312 0x00002022, 0x04002022, 0x00042022, 0x04042022,
313 0x00000800, 0x04000800, 0x00040800, 0x04040800,
314 0x00000802, 0x04000802, 0x00040802, 0x04040802,
315 0x00002800, 0x04002800, 0x00042800, 0x04042800,
316 0x00002802, 0x04002802, 0x00042802, 0x04042802,
317 0x00000820, 0x04000820, 0x00040820, 0x04040820,
318 0x00000822, 0x04000822, 0x00040822, 0x04040822,
319 0x00002820, 0x04002820, 0x00042820, 0x04042820,
320 0x00002822, 0x04002822, 0x00042822, 0x04042822
323 __constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
324 __constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
326 #define BOX(i,n,S) (S)[(n)][(i)]
328 static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
334 for (u32 i = 0; i < 16; i += 2)
339 u = Kc[i + 0] ^ rotl32 (r, 30u);
340 t = Kd[i + 0] ^ rotl32 (r, 26u);
343 ^ BOX (amd_bfe (u, 0, 6), 0, s_SPtrans)
344 ^ BOX (amd_bfe (u, 8, 6), 2, s_SPtrans)
345 ^ BOX (amd_bfe (u, 16, 6), 4, s_SPtrans)
346 ^ BOX (amd_bfe (u, 24, 6), 6, s_SPtrans)
347 ^ BOX (amd_bfe (t, 0, 6), 1, s_SPtrans)
348 ^ BOX (amd_bfe (t, 8, 6), 3, s_SPtrans)
349 ^ BOX (amd_bfe (t, 16, 6), 5, s_SPtrans)
350 ^ BOX (amd_bfe (t, 24, 6), 7, s_SPtrans);
352 u = Kc[i + 1] ^ rotl32 (l, 30u);
353 t = Kd[i + 1] ^ rotl32 (l, 26u);
356 ^ BOX (amd_bfe (u, 0, 6), 0, s_SPtrans)
357 ^ BOX (amd_bfe (u, 8, 6), 2, s_SPtrans)
358 ^ BOX (amd_bfe (u, 16, 6), 4, s_SPtrans)
359 ^ BOX (amd_bfe (u, 24, 6), 6, s_SPtrans)
360 ^ BOX (amd_bfe (t, 0, 6), 1, s_SPtrans)
361 ^ BOX (amd_bfe (t, 8, 6), 3, s_SPtrans)
362 ^ BOX (amd_bfe (t, 16, 6), 5, s_SPtrans)
363 ^ BOX (amd_bfe (t, 24, 6), 7, s_SPtrans);
370 static void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 s_skb[8][64])
374 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
375 HPERM_OP (c, tt, 2, 0xcccc0000);
376 HPERM_OP (d, tt, 2, 0xcccc0000);
377 PERM_OP (d, c, tt, 1, 0x55555555);
378 PERM_OP (c, d, tt, 8, 0x00ff00ff);
379 PERM_OP (d, c, tt, 1, 0x55555555);
381 d = ((d & 0x000000ff) << 16)
382 | ((d & 0x0000ff00) << 0)
383 | ((d & 0x00ff0000) >> 16)
384 | ((c & 0xf0000000) >> 4);
389 for (u32 i = 0; i < 16; i++)
391 c = c >> shifts3s0[i] | c << shifts3s1[i];
392 d = d >> shifts3s0[i] | d << shifts3s1[i];
397 const u32 c00 = (c >> 0) & 0x0000003f;
398 const u32 c06 = (c >> 6) & 0x00383003;
399 const u32 c07 = (c >> 7) & 0x0000003c;
400 const u32 c13 = (c >> 13) & 0x0000060f;
401 const u32 c20 = (c >> 20) & 0x00000001;
403 u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
404 | BOX (((c06 >> 0) & 0xff)
405 |((c07 >> 0) & 0xff), 1, s_skb)
406 | BOX (((c13 >> 0) & 0xff)
407 |((c06 >> 8) & 0xff), 2, s_skb)
408 | BOX (((c20 >> 0) & 0xff)
410 |((c06 >> 16) & 0xff), 3, s_skb);
412 const u32 d00 = (d >> 0) & 0x00003c3f;
413 const u32 d07 = (d >> 7) & 0x00003f03;
414 const u32 d21 = (d >> 21) & 0x0000000f;
415 const u32 d22 = (d >> 22) & 0x00000030;
417 u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
418 | BOX (((d07 >> 0) & 0xff)
419 |((d00 >> 8) & 0xff), 5, s_skb)
420 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
421 | BOX (((d21 >> 0) & 0xff)
422 |((d22 >> 0) & 0xff), 7, s_skb);
424 Kc[i] = ((t << 16) | (s & 0x0000ffff));
425 Kd[i] = ((s >> 16) | (t & 0xffff0000));
429 static void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
431 const uchar4 t0 = as_uchar4 (w0);
432 const uchar4 t1 = as_uchar4 (w1);
437 k0.s0 = (t0.s0 >> 0);
438 k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
439 k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
440 k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
441 k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
442 k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
443 k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
444 k1.s3 = (t1.s2 << 1);
446 out[0] = as_uint (k0);
447 out[1] = as_uint (k1);
450 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_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)
456 const u32 lid = get_local_id (0);
462 const u32 gid = get_global_id (0);
466 wordl0[0] = pws[gid].i[ 0];
467 wordl0[1] = pws[gid].i[ 1];
468 wordl0[2] = pws[gid].i[ 2];
469 wordl0[3] = pws[gid].i[ 3];
473 wordl1[0] = pws[gid].i[ 4];
474 wordl1[1] = pws[gid].i[ 5];
475 wordl1[2] = pws[gid].i[ 6];
476 wordl1[3] = pws[gid].i[ 7];
492 const u32 pw_l_len = pws[gid].pw_len;
494 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
496 append_0x80_2x4 (wordl0, wordl1, pw_l_len);
498 switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
505 __local u32 s_SPtrans[8][64];
506 __local u32 s_skb[8][64];
508 s_SPtrans[0][lid] = c_SPtrans[0][lid];
509 s_SPtrans[1][lid] = c_SPtrans[1][lid];
510 s_SPtrans[2][lid] = c_SPtrans[2][lid];
511 s_SPtrans[3][lid] = c_SPtrans[3][lid];
512 s_SPtrans[4][lid] = c_SPtrans[4][lid];
513 s_SPtrans[5][lid] = c_SPtrans[5][lid];
514 s_SPtrans[6][lid] = c_SPtrans[6][lid];
515 s_SPtrans[7][lid] = c_SPtrans[7][lid];
517 s_skb[0][lid] = c_skb[0][lid];
518 s_skb[1][lid] = c_skb[1][lid];
519 s_skb[2][lid] = c_skb[2][lid];
520 s_skb[3][lid] = c_skb[3][lid];
521 s_skb[4][lid] = c_skb[4][lid];
522 s_skb[5][lid] = c_skb[5][lid];
523 s_skb[6][lid] = c_skb[6][lid];
524 s_skb[7][lid] = c_skb[7][lid];
526 barrier (CLK_LOCAL_MEM_FENCE);
528 if (gid >= gid_max) return;
534 const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
535 const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
536 const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
547 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
549 const u32 pw_r_len = combs_buf[il_pos].pw_len;
551 const u32 pw_len = pw_l_len + pw_r_len;
555 wordr0[0] = combs_buf[il_pos].i[0];
556 wordr0[1] = combs_buf[il_pos].i[1];
557 wordr0[2] = combs_buf[il_pos].i[2];
558 wordr0[3] = combs_buf[il_pos].i[3];
562 wordr1[0] = combs_buf[il_pos].i[4];
563 wordr1[1] = combs_buf[il_pos].i[5];
564 wordr1[2] = combs_buf[il_pos].i[6];
565 wordr1[3] = combs_buf[il_pos].i[7];
581 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
583 switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
588 w0[0] = wordl0[0] | wordr0[0];
589 w0[1] = wordl0[1] | wordr0[1];
590 w0[2] = wordl0[2] | wordr0[2];
591 w0[3] = wordl0[3] | wordr0[3];
595 w1[0] = wordl1[0] | wordr1[0];
596 w1[1] = wordl1[1] | wordr1[1];
597 w1[2] = wordl1[2] | wordr1[2];
598 w1[3] = wordl1[3] | wordr1[3];
619 make_unicode (w0, w0_t, w1_t);
620 make_unicode (w1, w2_t, w3_t);
622 w3_t[2] = pw_len * 8 * 2;
629 MD4_STEP (MD4_Fo, a, b, c, d, w0_t[0], MD4C00, MD4S00);
630 MD4_STEP (MD4_Fo, d, a, b, c, w0_t[1], MD4C00, MD4S01);
631 MD4_STEP (MD4_Fo, c, d, a, b, w0_t[2], MD4C00, MD4S02);
632 MD4_STEP (MD4_Fo, b, c, d, a, w0_t[3], MD4C00, MD4S03);
633 MD4_STEP (MD4_Fo, a, b, c, d, w1_t[0], MD4C00, MD4S00);
634 MD4_STEP (MD4_Fo, d, a, b, c, w1_t[1], MD4C00, MD4S01);
635 MD4_STEP (MD4_Fo, c, d, a, b, w1_t[2], MD4C00, MD4S02);
636 MD4_STEP (MD4_Fo, b, c, d, a, w1_t[3], MD4C00, MD4S03);
637 MD4_STEP (MD4_Fo, a, b, c, d, w2_t[0], MD4C00, MD4S00);
638 MD4_STEP (MD4_Fo, d, a, b, c, w2_t[1], MD4C00, MD4S01);
639 MD4_STEP (MD4_Fo, c, d, a, b, w2_t[2], MD4C00, MD4S02);
640 MD4_STEP (MD4_Fo, b, c, d, a, w2_t[3], MD4C00, MD4S03);
641 MD4_STEP (MD4_Fo, a, b, c, d, w3_t[0], MD4C00, MD4S00);
642 MD4_STEP (MD4_Fo, d, a, b, c, w3_t[1], MD4C00, MD4S01);
643 MD4_STEP (MD4_Fo, c, d, a, b, w3_t[2], MD4C00, MD4S02);
644 MD4_STEP (MD4_Fo, b, c, d, a, w3_t[3], MD4C00, MD4S03);
646 MD4_STEP (MD4_Go, a, b, c, d, w0_t[0], MD4C01, MD4S10);
647 MD4_STEP (MD4_Go, d, a, b, c, w1_t[0], MD4C01, MD4S11);
648 MD4_STEP (MD4_Go, c, d, a, b, w2_t[0], MD4C01, MD4S12);
649 MD4_STEP (MD4_Go, b, c, d, a, w3_t[0], MD4C01, MD4S13);
650 MD4_STEP (MD4_Go, a, b, c, d, w0_t[1], MD4C01, MD4S10);
651 MD4_STEP (MD4_Go, d, a, b, c, w1_t[1], MD4C01, MD4S11);
652 MD4_STEP (MD4_Go, c, d, a, b, w2_t[1], MD4C01, MD4S12);
653 MD4_STEP (MD4_Go, b, c, d, a, w3_t[1], MD4C01, MD4S13);
654 MD4_STEP (MD4_Go, a, b, c, d, w0_t[2], MD4C01, MD4S10);
655 MD4_STEP (MD4_Go, d, a, b, c, w1_t[2], MD4C01, MD4S11);
656 MD4_STEP (MD4_Go, c, d, a, b, w2_t[2], MD4C01, MD4S12);
657 MD4_STEP (MD4_Go, b, c, d, a, w3_t[2], MD4C01, MD4S13);
658 MD4_STEP (MD4_Go, a, b, c, d, w0_t[3], MD4C01, MD4S10);
659 MD4_STEP (MD4_Go, d, a, b, c, w1_t[3], MD4C01, MD4S11);
660 MD4_STEP (MD4_Go, c, d, a, b, w2_t[3], MD4C01, MD4S12);
661 MD4_STEP (MD4_Go, b, c, d, a, w3_t[3], MD4C01, MD4S13);
663 MD4_STEP (MD4_H , a, b, c, d, w0_t[0], MD4C02, MD4S20);
664 MD4_STEP (MD4_H , d, a, b, c, w2_t[0], MD4C02, MD4S21);
665 MD4_STEP (MD4_H , c, d, a, b, w1_t[0], MD4C02, MD4S22);
666 MD4_STEP (MD4_H , b, c, d, a, w3_t[0], MD4C02, MD4S23);
667 MD4_STEP (MD4_H , a, b, c, d, w0_t[2], MD4C02, MD4S20);
668 MD4_STEP (MD4_H , d, a, b, c, w2_t[2], MD4C02, MD4S21);
669 MD4_STEP (MD4_H , c, d, a, b, w1_t[2], MD4C02, MD4S22);
670 MD4_STEP (MD4_H , b, c, d, a, w3_t[2], MD4C02, MD4S23);
671 MD4_STEP (MD4_H , a, b, c, d, w0_t[1], MD4C02, MD4S20);
672 MD4_STEP (MD4_H , d, a, b, c, w2_t[1], MD4C02, MD4S21);
673 MD4_STEP (MD4_H , c, d, a, b, w1_t[1], MD4C02, MD4S22);
674 MD4_STEP (MD4_H , b, c, d, a, w3_t[1], MD4C02, MD4S23);
675 MD4_STEP (MD4_H , a, b, c, d, w0_t[3], MD4C02, MD4S20);
676 MD4_STEP (MD4_H , d, a, b, c, w2_t[3], MD4C02, MD4S21);
678 if (allx (s2 != ((d + MD4M_D) >> 16))) continue;
680 MD4_STEP (MD4_H , c, d, a, b, w1_t[3], MD4C02, MD4S22);
681 MD4_STEP (MD4_H , b, c, d, a, w3_t[3], MD4C02, MD4S23);
694 transform_netntlmv1_key (a, b, key);
699 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
703 _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
709 transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
711 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
715 _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
721 const u32 r0 = iv1[0];
722 const u32 r1 = iv1[1];
723 const u32 r2 = iv2[0];
724 const u32 r3 = iv2[1];
730 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_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)
734 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_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)
738 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_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)
744 const u32 lid = get_local_id (0);
750 const u32 gid = get_global_id (0);
754 wordl0[0] = pws[gid].i[ 0];
755 wordl0[1] = pws[gid].i[ 1];
756 wordl0[2] = pws[gid].i[ 2];
757 wordl0[3] = pws[gid].i[ 3];
761 wordl1[0] = pws[gid].i[ 4];
762 wordl1[1] = pws[gid].i[ 5];
763 wordl1[2] = pws[gid].i[ 6];
764 wordl1[3] = pws[gid].i[ 7];
780 const u32 pw_l_len = pws[gid].pw_len;
782 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
784 append_0x80_2x4 (wordl0, wordl1, pw_l_len);
786 switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
793 __local u32 s_SPtrans[8][64];
794 __local u32 s_skb[8][64];
796 s_SPtrans[0][lid] = c_SPtrans[0][lid];
797 s_SPtrans[1][lid] = c_SPtrans[1][lid];
798 s_SPtrans[2][lid] = c_SPtrans[2][lid];
799 s_SPtrans[3][lid] = c_SPtrans[3][lid];
800 s_SPtrans[4][lid] = c_SPtrans[4][lid];
801 s_SPtrans[5][lid] = c_SPtrans[5][lid];
802 s_SPtrans[6][lid] = c_SPtrans[6][lid];
803 s_SPtrans[7][lid] = c_SPtrans[7][lid];
805 s_skb[0][lid] = c_skb[0][lid];
806 s_skb[1][lid] = c_skb[1][lid];
807 s_skb[2][lid] = c_skb[2][lid];
808 s_skb[3][lid] = c_skb[3][lid];
809 s_skb[4][lid] = c_skb[4][lid];
810 s_skb[5][lid] = c_skb[5][lid];
811 s_skb[6][lid] = c_skb[6][lid];
812 s_skb[7][lid] = c_skb[7][lid];
814 barrier (CLK_LOCAL_MEM_FENCE);
816 if (gid >= gid_max) return;
822 const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
823 const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
824 const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
835 const u32 search[4] =
837 digests_buf[digests_offset].digest_buf[DGST_R0],
838 digests_buf[digests_offset].digest_buf[DGST_R1],
839 digests_buf[digests_offset].digest_buf[DGST_R2],
840 digests_buf[digests_offset].digest_buf[DGST_R3]
847 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
849 const u32 pw_r_len = combs_buf[il_pos].pw_len;
851 const u32 pw_len = pw_l_len + pw_r_len;
855 wordr0[0] = combs_buf[il_pos].i[0];
856 wordr0[1] = combs_buf[il_pos].i[1];
857 wordr0[2] = combs_buf[il_pos].i[2];
858 wordr0[3] = combs_buf[il_pos].i[3];
862 wordr1[0] = combs_buf[il_pos].i[4];
863 wordr1[1] = combs_buf[il_pos].i[5];
864 wordr1[2] = combs_buf[il_pos].i[6];
865 wordr1[3] = combs_buf[il_pos].i[7];
881 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
883 switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
888 w0[0] = wordl0[0] | wordr0[0];
889 w0[1] = wordl0[1] | wordr0[1];
890 w0[2] = wordl0[2] | wordr0[2];
891 w0[3] = wordl0[3] | wordr0[3];
895 w1[0] = wordl1[0] | wordr1[0];
896 w1[1] = wordl1[1] | wordr1[1];
897 w1[2] = wordl1[2] | wordr1[2];
898 w1[3] = wordl1[3] | wordr1[3];
919 make_unicode (w0, w0_t, w1_t);
920 make_unicode (w1, w2_t, w3_t);
922 w3_t[2] = pw_len * 8 * 2;
929 MD4_STEP (MD4_Fo, a, b, c, d, w0_t[0], MD4C00, MD4S00);
930 MD4_STEP (MD4_Fo, d, a, b, c, w0_t[1], MD4C00, MD4S01);
931 MD4_STEP (MD4_Fo, c, d, a, b, w0_t[2], MD4C00, MD4S02);
932 MD4_STEP (MD4_Fo, b, c, d, a, w0_t[3], MD4C00, MD4S03);
933 MD4_STEP (MD4_Fo, a, b, c, d, w1_t[0], MD4C00, MD4S00);
934 MD4_STEP (MD4_Fo, d, a, b, c, w1_t[1], MD4C00, MD4S01);
935 MD4_STEP (MD4_Fo, c, d, a, b, w1_t[2], MD4C00, MD4S02);
936 MD4_STEP (MD4_Fo, b, c, d, a, w1_t[3], MD4C00, MD4S03);
937 MD4_STEP (MD4_Fo, a, b, c, d, w2_t[0], MD4C00, MD4S00);
938 MD4_STEP (MD4_Fo, d, a, b, c, w2_t[1], MD4C00, MD4S01);
939 MD4_STEP (MD4_Fo, c, d, a, b, w2_t[2], MD4C00, MD4S02);
940 MD4_STEP (MD4_Fo, b, c, d, a, w2_t[3], MD4C00, MD4S03);
941 MD4_STEP (MD4_Fo, a, b, c, d, w3_t[0], MD4C00, MD4S00);
942 MD4_STEP (MD4_Fo, d, a, b, c, w3_t[1], MD4C00, MD4S01);
943 MD4_STEP (MD4_Fo, c, d, a, b, w3_t[2], MD4C00, MD4S02);
944 MD4_STEP (MD4_Fo, b, c, d, a, w3_t[3], MD4C00, MD4S03);
946 MD4_STEP (MD4_Go, a, b, c, d, w0_t[0], MD4C01, MD4S10);
947 MD4_STEP (MD4_Go, d, a, b, c, w1_t[0], MD4C01, MD4S11);
948 MD4_STEP (MD4_Go, c, d, a, b, w2_t[0], MD4C01, MD4S12);
949 MD4_STEP (MD4_Go, b, c, d, a, w3_t[0], MD4C01, MD4S13);
950 MD4_STEP (MD4_Go, a, b, c, d, w0_t[1], MD4C01, MD4S10);
951 MD4_STEP (MD4_Go, d, a, b, c, w1_t[1], MD4C01, MD4S11);
952 MD4_STEP (MD4_Go, c, d, a, b, w2_t[1], MD4C01, MD4S12);
953 MD4_STEP (MD4_Go, b, c, d, a, w3_t[1], MD4C01, MD4S13);
954 MD4_STEP (MD4_Go, a, b, c, d, w0_t[2], MD4C01, MD4S10);
955 MD4_STEP (MD4_Go, d, a, b, c, w1_t[2], MD4C01, MD4S11);
956 MD4_STEP (MD4_Go, c, d, a, b, w2_t[2], MD4C01, MD4S12);
957 MD4_STEP (MD4_Go, b, c, d, a, w3_t[2], MD4C01, MD4S13);
958 MD4_STEP (MD4_Go, a, b, c, d, w0_t[3], MD4C01, MD4S10);
959 MD4_STEP (MD4_Go, d, a, b, c, w1_t[3], MD4C01, MD4S11);
960 MD4_STEP (MD4_Go, c, d, a, b, w2_t[3], MD4C01, MD4S12);
961 MD4_STEP (MD4_Go, b, c, d, a, w3_t[3], MD4C01, MD4S13);
963 MD4_STEP (MD4_H , a, b, c, d, w0_t[0], MD4C02, MD4S20);
964 MD4_STEP (MD4_H , d, a, b, c, w2_t[0], MD4C02, MD4S21);
965 MD4_STEP (MD4_H , c, d, a, b, w1_t[0], MD4C02, MD4S22);
966 MD4_STEP (MD4_H , b, c, d, a, w3_t[0], MD4C02, MD4S23);
967 MD4_STEP (MD4_H , a, b, c, d, w0_t[2], MD4C02, MD4S20);
968 MD4_STEP (MD4_H , d, a, b, c, w2_t[2], MD4C02, MD4S21);
969 MD4_STEP (MD4_H , c, d, a, b, w1_t[2], MD4C02, MD4S22);
970 MD4_STEP (MD4_H , b, c, d, a, w3_t[2], MD4C02, MD4S23);
971 MD4_STEP (MD4_H , a, b, c, d, w0_t[1], MD4C02, MD4S20);
972 MD4_STEP (MD4_H , d, a, b, c, w2_t[1], MD4C02, MD4S21);
973 MD4_STEP (MD4_H , c, d, a, b, w1_t[1], MD4C02, MD4S22);
974 MD4_STEP (MD4_H , b, c, d, a, w3_t[1], MD4C02, MD4S23);
975 MD4_STEP (MD4_H , a, b, c, d, w0_t[3], MD4C02, MD4S20);
976 MD4_STEP (MD4_H , d, a, b, c, w2_t[3], MD4C02, MD4S21);
978 if (allx (s2 != ((d + MD4M_D) >> 16))) continue;
980 MD4_STEP (MD4_H , c, d, a, b, w1_t[3], MD4C02, MD4S22);
981 MD4_STEP (MD4_H , b, c, d, a, w3_t[3], MD4C02, MD4S23);
994 transform_netntlmv1_key (a, b, key);
999 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
1003 _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
1010 transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
1012 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
1016 _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
1028 const u32 r0 = iv1[0];
1029 const u32 r1 = iv1[1];
1030 const u32 r2 = iv2[0];
1031 const u32 r3 = iv2[1];
1037 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_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)
1041 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_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)