2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
28 #include "include/kernel_functions.c"
29 #include "types_amd.c"
30 #include "common_amd.c"
33 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
34 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
38 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
39 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
43 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
44 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
47 #define PERM_OP(a,b,tt,n,m) \
57 #define HPERM_OP(a,tt,n,m) \
63 tt = tt >> (16 + n); \
67 __constant u32 c_SPtrans[8][64] =
70 0x02080800, 0x00080000, 0x02000002, 0x02080802,
71 0x02000000, 0x00080802, 0x00080002, 0x02000002,
72 0x00080802, 0x02080800, 0x02080000, 0x00000802,
73 0x02000802, 0x02000000, 0x00000000, 0x00080002,
74 0x00080000, 0x00000002, 0x02000800, 0x00080800,
75 0x02080802, 0x02080000, 0x00000802, 0x02000800,
76 0x00000002, 0x00000800, 0x00080800, 0x02080002,
77 0x00000800, 0x02000802, 0x02080002, 0x00000000,
78 0x00000000, 0x02080802, 0x02000800, 0x00080002,
79 0x02080800, 0x00080000, 0x00000802, 0x02000800,
80 0x02080002, 0x00000800, 0x00080800, 0x02000002,
81 0x00080802, 0x00000002, 0x02000002, 0x02080000,
82 0x02080802, 0x00080800, 0x02080000, 0x02000802,
83 0x02000000, 0x00000802, 0x00080002, 0x00000000,
84 0x00080000, 0x02000000, 0x02000802, 0x02080800,
85 0x00000002, 0x02080002, 0x00000800, 0x00080802,
87 0x40108010, 0x00000000, 0x00108000, 0x40100000,
88 0x40000010, 0x00008010, 0x40008000, 0x00108000,
89 0x00008000, 0x40100010, 0x00000010, 0x40008000,
90 0x00100010, 0x40108000, 0x40100000, 0x00000010,
91 0x00100000, 0x40008010, 0x40100010, 0x00008000,
92 0x00108010, 0x40000000, 0x00000000, 0x00100010,
93 0x40008010, 0x00108010, 0x40108000, 0x40000010,
94 0x40000000, 0x00100000, 0x00008010, 0x40108010,
95 0x00100010, 0x40108000, 0x40008000, 0x00108010,
96 0x40108010, 0x00100010, 0x40000010, 0x00000000,
97 0x40000000, 0x00008010, 0x00100000, 0x40100010,
98 0x00008000, 0x40000000, 0x00108010, 0x40008010,
99 0x40108000, 0x00008000, 0x00000000, 0x40000010,
100 0x00000010, 0x40108010, 0x00108000, 0x40100000,
101 0x40100010, 0x00100000, 0x00008010, 0x40008000,
102 0x40008010, 0x00000010, 0x40100000, 0x00108000,
104 0x04000001, 0x04040100, 0x00000100, 0x04000101,
105 0x00040001, 0x04000000, 0x04000101, 0x00040100,
106 0x04000100, 0x00040000, 0x04040000, 0x00000001,
107 0x04040101, 0x00000101, 0x00000001, 0x04040001,
108 0x00000000, 0x00040001, 0x04040100, 0x00000100,
109 0x00000101, 0x04040101, 0x00040000, 0x04000001,
110 0x04040001, 0x04000100, 0x00040101, 0x04040000,
111 0x00040100, 0x00000000, 0x04000000, 0x00040101,
112 0x04040100, 0x00000100, 0x00000001, 0x00040000,
113 0x00000101, 0x00040001, 0x04040000, 0x04000101,
114 0x00000000, 0x04040100, 0x00040100, 0x04040001,
115 0x00040001, 0x04000000, 0x04040101, 0x00000001,
116 0x00040101, 0x04000001, 0x04000000, 0x04040101,
117 0x00040000, 0x04000100, 0x04000101, 0x00040100,
118 0x04000100, 0x00000000, 0x04040001, 0x00000101,
119 0x04000001, 0x00040101, 0x00000100, 0x04040000,
121 0x00401008, 0x10001000, 0x00000008, 0x10401008,
122 0x00000000, 0x10400000, 0x10001008, 0x00400008,
123 0x10401000, 0x10000008, 0x10000000, 0x00001008,
124 0x10000008, 0x00401008, 0x00400000, 0x10000000,
125 0x10400008, 0x00401000, 0x00001000, 0x00000008,
126 0x00401000, 0x10001008, 0x10400000, 0x00001000,
127 0x00001008, 0x00000000, 0x00400008, 0x10401000,
128 0x10001000, 0x10400008, 0x10401008, 0x00400000,
129 0x10400008, 0x00001008, 0x00400000, 0x10000008,
130 0x00401000, 0x10001000, 0x00000008, 0x10400000,
131 0x10001008, 0x00000000, 0x00001000, 0x00400008,
132 0x00000000, 0x10400008, 0x10401000, 0x00001000,
133 0x10000000, 0x10401008, 0x00401008, 0x00400000,
134 0x10401008, 0x00000008, 0x10001000, 0x00401008,
135 0x00400008, 0x00401000, 0x10400000, 0x10001008,
136 0x00001008, 0x10000000, 0x10000008, 0x10401000,
138 0x08000000, 0x00010000, 0x00000400, 0x08010420,
139 0x08010020, 0x08000400, 0x00010420, 0x08010000,
140 0x00010000, 0x00000020, 0x08000020, 0x00010400,
141 0x08000420, 0x08010020, 0x08010400, 0x00000000,
142 0x00010400, 0x08000000, 0x00010020, 0x00000420,
143 0x08000400, 0x00010420, 0x00000000, 0x08000020,
144 0x00000020, 0x08000420, 0x08010420, 0x00010020,
145 0x08010000, 0x00000400, 0x00000420, 0x08010400,
146 0x08010400, 0x08000420, 0x00010020, 0x08010000,
147 0x00010000, 0x00000020, 0x08000020, 0x08000400,
148 0x08000000, 0x00010400, 0x08010420, 0x00000000,
149 0x00010420, 0x08000000, 0x00000400, 0x00010020,
150 0x08000420, 0x00000400, 0x00000000, 0x08010420,
151 0x08010020, 0x08010400, 0x00000420, 0x00010000,
152 0x00010400, 0x08010020, 0x08000400, 0x00000420,
153 0x00000020, 0x00010420, 0x08010000, 0x08000020,
155 0x80000040, 0x00200040, 0x00000000, 0x80202000,
156 0x00200040, 0x00002000, 0x80002040, 0x00200000,
157 0x00002040, 0x80202040, 0x00202000, 0x80000000,
158 0x80002000, 0x80000040, 0x80200000, 0x00202040,
159 0x00200000, 0x80002040, 0x80200040, 0x00000000,
160 0x00002000, 0x00000040, 0x80202000, 0x80200040,
161 0x80202040, 0x80200000, 0x80000000, 0x00002040,
162 0x00000040, 0x00202000, 0x00202040, 0x80002000,
163 0x00002040, 0x80000000, 0x80002000, 0x00202040,
164 0x80202000, 0x00200040, 0x00000000, 0x80002000,
165 0x80000000, 0x00002000, 0x80200040, 0x00200000,
166 0x00200040, 0x80202040, 0x00202000, 0x00000040,
167 0x80202040, 0x00202000, 0x00200000, 0x80002040,
168 0x80000040, 0x80200000, 0x00202040, 0x00000000,
169 0x00002000, 0x80000040, 0x80002040, 0x80202000,
170 0x80200000, 0x00002040, 0x00000040, 0x80200040,
172 0x00004000, 0x00000200, 0x01000200, 0x01000004,
173 0x01004204, 0x00004004, 0x00004200, 0x00000000,
174 0x01000000, 0x01000204, 0x00000204, 0x01004000,
175 0x00000004, 0x01004200, 0x01004000, 0x00000204,
176 0x01000204, 0x00004000, 0x00004004, 0x01004204,
177 0x00000000, 0x01000200, 0x01000004, 0x00004200,
178 0x01004004, 0x00004204, 0x01004200, 0x00000004,
179 0x00004204, 0x01004004, 0x00000200, 0x01000000,
180 0x00004204, 0x01004000, 0x01004004, 0x00000204,
181 0x00004000, 0x00000200, 0x01000000, 0x01004004,
182 0x01000204, 0x00004204, 0x00004200, 0x00000000,
183 0x00000200, 0x01000004, 0x00000004, 0x01000200,
184 0x00000000, 0x01000204, 0x01000200, 0x00004200,
185 0x00000204, 0x00004000, 0x01004204, 0x01000000,
186 0x01004200, 0x00000004, 0x00004004, 0x01004204,
187 0x01000004, 0x01004200, 0x01004000, 0x00004004,
189 0x20800080, 0x20820000, 0x00020080, 0x00000000,
190 0x20020000, 0x00800080, 0x20800000, 0x20820080,
191 0x00000080, 0x20000000, 0x00820000, 0x00020080,
192 0x00820080, 0x20020080, 0x20000080, 0x20800000,
193 0x00020000, 0x00820080, 0x00800080, 0x20020000,
194 0x20820080, 0x20000080, 0x00000000, 0x00820000,
195 0x20000000, 0x00800000, 0x20020080, 0x20800080,
196 0x00800000, 0x00020000, 0x20820000, 0x00000080,
197 0x00800000, 0x00020000, 0x20000080, 0x20820080,
198 0x00020080, 0x20000000, 0x00000000, 0x00820000,
199 0x20800080, 0x20020080, 0x20020000, 0x00800080,
200 0x20820000, 0x00000080, 0x00800080, 0x20020000,
201 0x20820080, 0x00800000, 0x20800000, 0x20000080,
202 0x00820000, 0x00020080, 0x20020080, 0x20800000,
203 0x00000080, 0x20820000, 0x00820080, 0x00000000,
204 0x20000000, 0x20800080, 0x00020000, 0x00820080,
207 __constant u32 c_skb[8][64] =
209 /* for C bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
210 0x00000000, 0x00000010, 0x20000000, 0x20000010,
211 0x00010000, 0x00010010, 0x20010000, 0x20010010,
212 0x00000800, 0x00000810, 0x20000800, 0x20000810,
213 0x00010800, 0x00010810, 0x20010800, 0x20010810,
214 0x00000020, 0x00000030, 0x20000020, 0x20000030,
215 0x00010020, 0x00010030, 0x20010020, 0x20010030,
216 0x00000820, 0x00000830, 0x20000820, 0x20000830,
217 0x00010820, 0x00010830, 0x20010820, 0x20010830,
218 0x00080000, 0x00080010, 0x20080000, 0x20080010,
219 0x00090000, 0x00090010, 0x20090000, 0x20090010,
220 0x00080800, 0x00080810, 0x20080800, 0x20080810,
221 0x00090800, 0x00090810, 0x20090800, 0x20090810,
222 0x00080020, 0x00080030, 0x20080020, 0x20080030,
223 0x00090020, 0x00090030, 0x20090020, 0x20090030,
224 0x00080820, 0x00080830, 0x20080820, 0x20080830,
225 0x00090820, 0x00090830, 0x20090820, 0x20090830,
226 /* for C bits (numbered as per FIPS 46) 7 8 10 11 12 13 */
227 0x00000000, 0x02000000, 0x00002000, 0x02002000,
228 0x00200000, 0x02200000, 0x00202000, 0x02202000,
229 0x00000004, 0x02000004, 0x00002004, 0x02002004,
230 0x00200004, 0x02200004, 0x00202004, 0x02202004,
231 0x00000400, 0x02000400, 0x00002400, 0x02002400,
232 0x00200400, 0x02200400, 0x00202400, 0x02202400,
233 0x00000404, 0x02000404, 0x00002404, 0x02002404,
234 0x00200404, 0x02200404, 0x00202404, 0x02202404,
235 0x10000000, 0x12000000, 0x10002000, 0x12002000,
236 0x10200000, 0x12200000, 0x10202000, 0x12202000,
237 0x10000004, 0x12000004, 0x10002004, 0x12002004,
238 0x10200004, 0x12200004, 0x10202004, 0x12202004,
239 0x10000400, 0x12000400, 0x10002400, 0x12002400,
240 0x10200400, 0x12200400, 0x10202400, 0x12202400,
241 0x10000404, 0x12000404, 0x10002404, 0x12002404,
242 0x10200404, 0x12200404, 0x10202404, 0x12202404,
243 /* for C bits (numbered as per FIPS 46) 14 15 16 17 19 20 */
244 0x00000000, 0x00000001, 0x00040000, 0x00040001,
245 0x01000000, 0x01000001, 0x01040000, 0x01040001,
246 0x00000002, 0x00000003, 0x00040002, 0x00040003,
247 0x01000002, 0x01000003, 0x01040002, 0x01040003,
248 0x00000200, 0x00000201, 0x00040200, 0x00040201,
249 0x01000200, 0x01000201, 0x01040200, 0x01040201,
250 0x00000202, 0x00000203, 0x00040202, 0x00040203,
251 0x01000202, 0x01000203, 0x01040202, 0x01040203,
252 0x08000000, 0x08000001, 0x08040000, 0x08040001,
253 0x09000000, 0x09000001, 0x09040000, 0x09040001,
254 0x08000002, 0x08000003, 0x08040002, 0x08040003,
255 0x09000002, 0x09000003, 0x09040002, 0x09040003,
256 0x08000200, 0x08000201, 0x08040200, 0x08040201,
257 0x09000200, 0x09000201, 0x09040200, 0x09040201,
258 0x08000202, 0x08000203, 0x08040202, 0x08040203,
259 0x09000202, 0x09000203, 0x09040202, 0x09040203,
260 /* for C bits (numbered as per FIPS 46) 21 23 24 26 27 28 */
261 0x00000000, 0x00100000, 0x00000100, 0x00100100,
262 0x00000008, 0x00100008, 0x00000108, 0x00100108,
263 0x00001000, 0x00101000, 0x00001100, 0x00101100,
264 0x00001008, 0x00101008, 0x00001108, 0x00101108,
265 0x04000000, 0x04100000, 0x04000100, 0x04100100,
266 0x04000008, 0x04100008, 0x04000108, 0x04100108,
267 0x04001000, 0x04101000, 0x04001100, 0x04101100,
268 0x04001008, 0x04101008, 0x04001108, 0x04101108,
269 0x00020000, 0x00120000, 0x00020100, 0x00120100,
270 0x00020008, 0x00120008, 0x00020108, 0x00120108,
271 0x00021000, 0x00121000, 0x00021100, 0x00121100,
272 0x00021008, 0x00121008, 0x00021108, 0x00121108,
273 0x04020000, 0x04120000, 0x04020100, 0x04120100,
274 0x04020008, 0x04120008, 0x04020108, 0x04120108,
275 0x04021000, 0x04121000, 0x04021100, 0x04121100,
276 0x04021008, 0x04121008, 0x04021108, 0x04121108,
277 /* for D bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
278 0x00000000, 0x10000000, 0x00010000, 0x10010000,
279 0x00000004, 0x10000004, 0x00010004, 0x10010004,
280 0x20000000, 0x30000000, 0x20010000, 0x30010000,
281 0x20000004, 0x30000004, 0x20010004, 0x30010004,
282 0x00100000, 0x10100000, 0x00110000, 0x10110000,
283 0x00100004, 0x10100004, 0x00110004, 0x10110004,
284 0x20100000, 0x30100000, 0x20110000, 0x30110000,
285 0x20100004, 0x30100004, 0x20110004, 0x30110004,
286 0x00001000, 0x10001000, 0x00011000, 0x10011000,
287 0x00001004, 0x10001004, 0x00011004, 0x10011004,
288 0x20001000, 0x30001000, 0x20011000, 0x30011000,
289 0x20001004, 0x30001004, 0x20011004, 0x30011004,
290 0x00101000, 0x10101000, 0x00111000, 0x10111000,
291 0x00101004, 0x10101004, 0x00111004, 0x10111004,
292 0x20101000, 0x30101000, 0x20111000, 0x30111000,
293 0x20101004, 0x30101004, 0x20111004, 0x30111004,
294 /* for D bits (numbered as per FIPS 46) 8 9 11 12 13 14 */
295 0x00000000, 0x08000000, 0x00000008, 0x08000008,
296 0x00000400, 0x08000400, 0x00000408, 0x08000408,
297 0x00020000, 0x08020000, 0x00020008, 0x08020008,
298 0x00020400, 0x08020400, 0x00020408, 0x08020408,
299 0x00000001, 0x08000001, 0x00000009, 0x08000009,
300 0x00000401, 0x08000401, 0x00000409, 0x08000409,
301 0x00020001, 0x08020001, 0x00020009, 0x08020009,
302 0x00020401, 0x08020401, 0x00020409, 0x08020409,
303 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
304 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
305 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
306 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
307 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
308 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
309 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
310 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
311 /* for D bits (numbered as per FIPS 46) 16 17 18 19 20 21 */
312 0x00000000, 0x00000100, 0x00080000, 0x00080100,
313 0x01000000, 0x01000100, 0x01080000, 0x01080100,
314 0x00000010, 0x00000110, 0x00080010, 0x00080110,
315 0x01000010, 0x01000110, 0x01080010, 0x01080110,
316 0x00200000, 0x00200100, 0x00280000, 0x00280100,
317 0x01200000, 0x01200100, 0x01280000, 0x01280100,
318 0x00200010, 0x00200110, 0x00280010, 0x00280110,
319 0x01200010, 0x01200110, 0x01280010, 0x01280110,
320 0x00000200, 0x00000300, 0x00080200, 0x00080300,
321 0x01000200, 0x01000300, 0x01080200, 0x01080300,
322 0x00000210, 0x00000310, 0x00080210, 0x00080310,
323 0x01000210, 0x01000310, 0x01080210, 0x01080310,
324 0x00200200, 0x00200300, 0x00280200, 0x00280300,
325 0x01200200, 0x01200300, 0x01280200, 0x01280300,
326 0x00200210, 0x00200310, 0x00280210, 0x00280310,
327 0x01200210, 0x01200310, 0x01280210, 0x01280310,
328 /* for D bits (numbered as per FIPS 46) 22 23 24 25 27 28 */
329 0x00000000, 0x04000000, 0x00040000, 0x04040000,
330 0x00000002, 0x04000002, 0x00040002, 0x04040002,
331 0x00002000, 0x04002000, 0x00042000, 0x04042000,
332 0x00002002, 0x04002002, 0x00042002, 0x04042002,
333 0x00000020, 0x04000020, 0x00040020, 0x04040020,
334 0x00000022, 0x04000022, 0x00040022, 0x04040022,
335 0x00002020, 0x04002020, 0x00042020, 0x04042020,
336 0x00002022, 0x04002022, 0x00042022, 0x04042022,
337 0x00000800, 0x04000800, 0x00040800, 0x04040800,
338 0x00000802, 0x04000802, 0x00040802, 0x04040802,
339 0x00002800, 0x04002800, 0x00042800, 0x04042800,
340 0x00002802, 0x04002802, 0x00042802, 0x04042802,
341 0x00000820, 0x04000820, 0x00040820, 0x04040820,
342 0x00000822, 0x04000822, 0x00040822, 0x04040822,
343 0x00002820, 0x04002820, 0x00042820, 0x04042820,
344 0x00002822, 0x04002822, 0x00042822, 0x04042822
347 __constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
348 __constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
351 #define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
355 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
359 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
362 static void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 s_SPtrans[8][64])
368 for (u32 i = 0; i < 16; i += 2)
373 u = Kc[i + 0] ^ rotl32 (r, 30u);
374 t = Kd[i + 0] ^ rotl32 (r, 26u);
377 ^ BOX (amd_bfe (u, 0, 6), 0, s_SPtrans)
378 ^ BOX (amd_bfe (u, 8, 6), 2, s_SPtrans)
379 ^ BOX (amd_bfe (u, 16, 6), 4, s_SPtrans)
380 ^ BOX (amd_bfe (u, 24, 6), 6, s_SPtrans)
381 ^ BOX (amd_bfe (t, 0, 6), 1, s_SPtrans)
382 ^ BOX (amd_bfe (t, 8, 6), 3, s_SPtrans)
383 ^ BOX (amd_bfe (t, 16, 6), 5, s_SPtrans)
384 ^ BOX (amd_bfe (t, 24, 6), 7, s_SPtrans);
386 u = Kc[i + 1] ^ rotl32 (l, 30u);
387 t = Kd[i + 1] ^ rotl32 (l, 26u);
390 ^ BOX (amd_bfe (u, 0, 6), 0, s_SPtrans)
391 ^ BOX (amd_bfe (u, 8, 6), 2, s_SPtrans)
392 ^ BOX (amd_bfe (u, 16, 6), 4, s_SPtrans)
393 ^ BOX (amd_bfe (u, 24, 6), 6, s_SPtrans)
394 ^ BOX (amd_bfe (t, 0, 6), 1, s_SPtrans)
395 ^ BOX (amd_bfe (t, 8, 6), 3, s_SPtrans)
396 ^ BOX (amd_bfe (t, 16, 6), 5, s_SPtrans)
397 ^ BOX (amd_bfe (t, 24, 6), 7, s_SPtrans);
404 static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 s_skb[8][64])
408 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
409 HPERM_OP (c, tt, 2, 0xcccc0000);
410 HPERM_OP (d, tt, 2, 0xcccc0000);
411 PERM_OP (d, c, tt, 1, 0x55555555);
412 PERM_OP (c, d, tt, 8, 0x00ff00ff);
413 PERM_OP (d, c, tt, 1, 0x55555555);
415 d = ((d & 0x000000ff) << 16)
416 | ((d & 0x0000ff00) << 0)
417 | ((d & 0x00ff0000) >> 16)
418 | ((c & 0xf0000000) >> 4);
423 for (u32 i = 0; i < 16; i++)
425 c = c >> shifts3s0[i] | c << shifts3s1[i];
426 d = d >> shifts3s0[i] | d << shifts3s1[i];
431 const u32x c00 = (c >> 0) & 0x0000003f;
432 const u32x c06 = (c >> 6) & 0x00383003;
433 const u32x c07 = (c >> 7) & 0x0000003c;
434 const u32x c13 = (c >> 13) & 0x0000060f;
435 const u32x c20 = (c >> 20) & 0x00000001;
437 u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
438 | BOX (((c06 >> 0) & 0xff)
439 |((c07 >> 0) & 0xff), 1, s_skb)
440 | BOX (((c13 >> 0) & 0xff)
441 |((c06 >> 8) & 0xff), 2, s_skb)
442 | BOX (((c20 >> 0) & 0xff)
444 |((c06 >> 16) & 0xff), 3, s_skb);
446 const u32x d00 = (d >> 0) & 0x00003c3f;
447 const u32x d07 = (d >> 7) & 0x00003f03;
448 const u32x d21 = (d >> 21) & 0x0000000f;
449 const u32x d22 = (d >> 22) & 0x00000030;
451 u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
452 | BOX (((d07 >> 0) & 0xff)
453 |((d00 >> 8) & 0xff), 5, s_skb)
454 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
455 | BOX (((d21 >> 0) & 0xff)
456 |((d22 >> 0) & 0xff), 7, s_skb);
458 Kc[i] = ((t << 16) | (s & 0x0000ffff));
459 Kd[i] = ((s >> 16) | (t & 0xffff0000));
463 static void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
466 const uchar4 t0 = as_uchar4 (w0);
467 const uchar4 t1 = as_uchar4 (w1);
472 k0.s0 = (t0.s0 >> 0);
473 k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
474 k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
475 k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
476 k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
477 k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
478 k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
479 k1.s3 = (t1.s2 << 1);
481 out[0] = as_uint (k0);
482 out[1] = as_uint (k1);
486 const uchar8 t0 = as_uchar8 (w0);
487 const uchar8 t1 = as_uchar8 (w1);
492 k0.s0 = (t0.s0 >> 0);
493 k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
494 k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
495 k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
496 k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
497 k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
498 k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
499 k1.s3 = (t1.s2 << 1);
501 k0.s4 = (t0.s4 >> 0);
502 k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
503 k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
504 k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
505 k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
506 k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
507 k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
508 k1.s7 = (t1.s6 << 1);
510 out[0] = as_uint2 (k0);
511 out[1] = as_uint2 (k1);
515 const uchar16 t0 = as_uchar16 (w0);
516 const uchar16 t1 = as_uchar16 (w1);
521 k0.s0 = (t0.s0 >> 0);
522 k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
523 k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
524 k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
525 k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
526 k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
527 k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
528 k1.s3 = (t1.s2 << 1);
530 k0.s4 = (t0.s4 >> 0);
531 k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
532 k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
533 k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
534 k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
535 k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
536 k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
537 k1.s7 = (t1.s6 << 1);
539 k0.s8 = (t0.s8 >> 0);
540 k0.s9 = (t0.s8 << 7) | (t0.s9 >> 1);
541 k0.sa = (t0.s9 << 6) | (t0.sa >> 2);
542 k0.sb = (t0.sa << 5) | (t0.sb >> 3);
543 k1.s8 = (t0.sb << 4) | (t1.s8 >> 4);
544 k1.s9 = (t1.s8 << 3) | (t1.s9 >> 5);
545 k1.sa = (t1.s9 << 2) | (t1.sa >> 6);
546 k1.sb = (t1.sa << 1);
548 k0.sc = (t0.sc >> 0);
549 k0.sd = (t0.sc << 7) | (t0.sd >> 1);
550 k0.se = (t0.sd << 6) | (t0.se >> 2);
551 k0.sf = (t0.se << 5) | (t0.sf >> 3);
552 k1.sc = (t0.sf << 4) | (t1.sc >> 4);
553 k1.sd = (t1.sc << 3) | (t1.sd >> 5);
554 k1.se = (t1.sd << 2) | (t1.se >> 6);
555 k1.sf = (t1.se << 1);
557 out[0] = as_uint4 (k0);
558 out[1] = as_uint4 (k1);
562 __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)
568 const u32 lid = get_local_id (0);
574 const u32 gid = get_global_id (0);
578 wordl0[0] = pws[gid].i[ 0];
579 wordl0[1] = pws[gid].i[ 1];
580 wordl0[2] = pws[gid].i[ 2];
581 wordl0[3] = pws[gid].i[ 3];
585 wordl1[0] = pws[gid].i[ 4];
586 wordl1[1] = pws[gid].i[ 5];
587 wordl1[2] = pws[gid].i[ 6];
588 wordl1[3] = pws[gid].i[ 7];
604 const u32 pw_l_len = pws[gid].pw_len;
606 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
608 append_0x80_2 (wordl0, wordl1, pw_l_len);
610 switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
617 __local u32 s_SPtrans[8][64];
618 __local u32 s_skb[8][64];
620 s_SPtrans[0][lid] = c_SPtrans[0][lid];
621 s_SPtrans[1][lid] = c_SPtrans[1][lid];
622 s_SPtrans[2][lid] = c_SPtrans[2][lid];
623 s_SPtrans[3][lid] = c_SPtrans[3][lid];
624 s_SPtrans[4][lid] = c_SPtrans[4][lid];
625 s_SPtrans[5][lid] = c_SPtrans[5][lid];
626 s_SPtrans[6][lid] = c_SPtrans[6][lid];
627 s_SPtrans[7][lid] = c_SPtrans[7][lid];
629 s_skb[0][lid] = c_skb[0][lid];
630 s_skb[1][lid] = c_skb[1][lid];
631 s_skb[2][lid] = c_skb[2][lid];
632 s_skb[3][lid] = c_skb[3][lid];
633 s_skb[4][lid] = c_skb[4][lid];
634 s_skb[5][lid] = c_skb[5][lid];
635 s_skb[6][lid] = c_skb[6][lid];
636 s_skb[7][lid] = c_skb[7][lid];
638 barrier (CLK_LOCAL_MEM_FENCE);
640 if (gid >= gid_max) return;
646 const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
647 const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
648 const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
659 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
661 const u32 pw_r_len = combs_buf[il_pos].pw_len;
663 const u32 pw_len = pw_l_len + pw_r_len;
667 wordr0[0] = combs_buf[il_pos].i[0];
668 wordr0[1] = combs_buf[il_pos].i[1];
669 wordr0[2] = combs_buf[il_pos].i[2];
670 wordr0[3] = combs_buf[il_pos].i[3];
674 wordr1[0] = combs_buf[il_pos].i[4];
675 wordr1[1] = combs_buf[il_pos].i[5];
676 wordr1[2] = combs_buf[il_pos].i[6];
677 wordr1[3] = combs_buf[il_pos].i[7];
693 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
695 switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
700 w0[0] = wordl0[0] | wordr0[0];
701 w0[1] = wordl0[1] | wordr0[1];
702 w0[2] = wordl0[2] | wordr0[2];
703 w0[3] = wordl0[3] | wordr0[3];
707 w1[0] = wordl1[0] | wordr1[0];
708 w1[1] = wordl1[1] | wordr1[1];
709 w1[2] = wordl1[2] | wordr1[2];
710 w1[3] = wordl1[3] | wordr1[3];
731 make_unicode (w0, w0_t, w1_t);
732 make_unicode (w1, w2_t, w3_t);
734 w3_t[2] = pw_len * 8 * 2;
741 MD4_STEP (MD4_Fo, a, b, c, d, w0_t[0], MD4C00, MD4S00);
742 MD4_STEP (MD4_Fo, d, a, b, c, w0_t[1], MD4C00, MD4S01);
743 MD4_STEP (MD4_Fo, c, d, a, b, w0_t[2], MD4C00, MD4S02);
744 MD4_STEP (MD4_Fo, b, c, d, a, w0_t[3], MD4C00, MD4S03);
745 MD4_STEP (MD4_Fo, a, b, c, d, w1_t[0], MD4C00, MD4S00);
746 MD4_STEP (MD4_Fo, d, a, b, c, w1_t[1], MD4C00, MD4S01);
747 MD4_STEP (MD4_Fo, c, d, a, b, w1_t[2], MD4C00, MD4S02);
748 MD4_STEP (MD4_Fo, b, c, d, a, w1_t[3], MD4C00, MD4S03);
749 MD4_STEP (MD4_Fo, a, b, c, d, w2_t[0], MD4C00, MD4S00);
750 MD4_STEP (MD4_Fo, d, a, b, c, w2_t[1], MD4C00, MD4S01);
751 MD4_STEP (MD4_Fo, c, d, a, b, w2_t[2], MD4C00, MD4S02);
752 MD4_STEP (MD4_Fo, b, c, d, a, w2_t[3], MD4C00, MD4S03);
753 MD4_STEP (MD4_Fo, a, b, c, d, w3_t[0], MD4C00, MD4S00);
754 MD4_STEP (MD4_Fo, d, a, b, c, w3_t[1], MD4C00, MD4S01);
755 MD4_STEP (MD4_Fo, c, d, a, b, w3_t[2], MD4C00, MD4S02);
756 MD4_STEP (MD4_Fo, b, c, d, a, w3_t[3], MD4C00, MD4S03);
758 MD4_STEP (MD4_Go, a, b, c, d, w0_t[0], MD4C01, MD4S10);
759 MD4_STEP (MD4_Go, d, a, b, c, w1_t[0], MD4C01, MD4S11);
760 MD4_STEP (MD4_Go, c, d, a, b, w2_t[0], MD4C01, MD4S12);
761 MD4_STEP (MD4_Go, b, c, d, a, w3_t[0], MD4C01, MD4S13);
762 MD4_STEP (MD4_Go, a, b, c, d, w0_t[1], MD4C01, MD4S10);
763 MD4_STEP (MD4_Go, d, a, b, c, w1_t[1], MD4C01, MD4S11);
764 MD4_STEP (MD4_Go, c, d, a, b, w2_t[1], MD4C01, MD4S12);
765 MD4_STEP (MD4_Go, b, c, d, a, w3_t[1], MD4C01, MD4S13);
766 MD4_STEP (MD4_Go, a, b, c, d, w0_t[2], MD4C01, MD4S10);
767 MD4_STEP (MD4_Go, d, a, b, c, w1_t[2], MD4C01, MD4S11);
768 MD4_STEP (MD4_Go, c, d, a, b, w2_t[2], MD4C01, MD4S12);
769 MD4_STEP (MD4_Go, b, c, d, a, w3_t[2], MD4C01, MD4S13);
770 MD4_STEP (MD4_Go, a, b, c, d, w0_t[3], MD4C01, MD4S10);
771 MD4_STEP (MD4_Go, d, a, b, c, w1_t[3], MD4C01, MD4S11);
772 MD4_STEP (MD4_Go, c, d, a, b, w2_t[3], MD4C01, MD4S12);
773 MD4_STEP (MD4_Go, b, c, d, a, w3_t[3], MD4C01, MD4S13);
775 MD4_STEP (MD4_H , a, b, c, d, w0_t[0], MD4C02, MD4S20);
776 MD4_STEP (MD4_H , d, a, b, c, w2_t[0], MD4C02, MD4S21);
777 MD4_STEP (MD4_H , c, d, a, b, w1_t[0], MD4C02, MD4S22);
778 MD4_STEP (MD4_H , b, c, d, a, w3_t[0], MD4C02, MD4S23);
779 MD4_STEP (MD4_H , a, b, c, d, w0_t[2], MD4C02, MD4S20);
780 MD4_STEP (MD4_H , d, a, b, c, w2_t[2], MD4C02, MD4S21);
781 MD4_STEP (MD4_H , c, d, a, b, w1_t[2], MD4C02, MD4S22);
782 MD4_STEP (MD4_H , b, c, d, a, w3_t[2], MD4C02, MD4S23);
783 MD4_STEP (MD4_H , a, b, c, d, w0_t[1], MD4C02, MD4S20);
784 MD4_STEP (MD4_H , d, a, b, c, w2_t[1], MD4C02, MD4S21);
785 MD4_STEP (MD4_H , c, d, a, b, w1_t[1], MD4C02, MD4S22);
786 MD4_STEP (MD4_H , b, c, d, a, w3_t[1], MD4C02, MD4S23);
787 MD4_STEP (MD4_H , a, b, c, d, w0_t[3], MD4C02, MD4S20);
788 MD4_STEP (MD4_H , d, a, b, c, w2_t[3], MD4C02, MD4S21);
790 if (allx (s2 != ((d + MD4M_D) >> 16))) continue;
792 MD4_STEP (MD4_H , c, d, a, b, w1_t[3], MD4C02, MD4S22);
793 MD4_STEP (MD4_H , b, c, d, a, w3_t[3], MD4C02, MD4S23);
806 transform_netntlmv1_key (a, b, key);
811 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
815 _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
821 transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
823 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
827 _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
833 const u32x r0 = iv1[0];
834 const u32x r1 = iv1[1];
835 const u32x r2 = iv2[0];
836 const u32x r3 = iv2[1];
838 #include VECT_COMPARE_M
842 __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)
846 __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)
850 __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)
856 const u32 lid = get_local_id (0);
862 const u32 gid = get_global_id (0);
866 wordl0[0] = pws[gid].i[ 0];
867 wordl0[1] = pws[gid].i[ 1];
868 wordl0[2] = pws[gid].i[ 2];
869 wordl0[3] = pws[gid].i[ 3];
873 wordl1[0] = pws[gid].i[ 4];
874 wordl1[1] = pws[gid].i[ 5];
875 wordl1[2] = pws[gid].i[ 6];
876 wordl1[3] = pws[gid].i[ 7];
892 const u32 pw_l_len = pws[gid].pw_len;
894 if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
896 append_0x80_2 (wordl0, wordl1, pw_l_len);
898 switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
905 __local u32 s_SPtrans[8][64];
906 __local u32 s_skb[8][64];
908 s_SPtrans[0][lid] = c_SPtrans[0][lid];
909 s_SPtrans[1][lid] = c_SPtrans[1][lid];
910 s_SPtrans[2][lid] = c_SPtrans[2][lid];
911 s_SPtrans[3][lid] = c_SPtrans[3][lid];
912 s_SPtrans[4][lid] = c_SPtrans[4][lid];
913 s_SPtrans[5][lid] = c_SPtrans[5][lid];
914 s_SPtrans[6][lid] = c_SPtrans[6][lid];
915 s_SPtrans[7][lid] = c_SPtrans[7][lid];
917 s_skb[0][lid] = c_skb[0][lid];
918 s_skb[1][lid] = c_skb[1][lid];
919 s_skb[2][lid] = c_skb[2][lid];
920 s_skb[3][lid] = c_skb[3][lid];
921 s_skb[4][lid] = c_skb[4][lid];
922 s_skb[5][lid] = c_skb[5][lid];
923 s_skb[6][lid] = c_skb[6][lid];
924 s_skb[7][lid] = c_skb[7][lid];
926 barrier (CLK_LOCAL_MEM_FENCE);
928 if (gid >= gid_max) return;
934 const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
935 const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
936 const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
947 const u32 search[4] =
949 digests_buf[digests_offset].digest_buf[DGST_R0],
950 digests_buf[digests_offset].digest_buf[DGST_R1],
951 digests_buf[digests_offset].digest_buf[DGST_R2],
952 digests_buf[digests_offset].digest_buf[DGST_R3]
959 for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
961 const u32 pw_r_len = combs_buf[il_pos].pw_len;
963 const u32 pw_len = pw_l_len + pw_r_len;
967 wordr0[0] = combs_buf[il_pos].i[0];
968 wordr0[1] = combs_buf[il_pos].i[1];
969 wordr0[2] = combs_buf[il_pos].i[2];
970 wordr0[3] = combs_buf[il_pos].i[3];
974 wordr1[0] = combs_buf[il_pos].i[4];
975 wordr1[1] = combs_buf[il_pos].i[5];
976 wordr1[2] = combs_buf[il_pos].i[6];
977 wordr1[3] = combs_buf[il_pos].i[7];
993 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
995 switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
1000 w0[0] = wordl0[0] | wordr0[0];
1001 w0[1] = wordl0[1] | wordr0[1];
1002 w0[2] = wordl0[2] | wordr0[2];
1003 w0[3] = wordl0[3] | wordr0[3];
1007 w1[0] = wordl1[0] | wordr1[0];
1008 w1[1] = wordl1[1] | wordr1[1];
1009 w1[2] = wordl1[2] | wordr1[2];
1010 w1[3] = wordl1[3] | wordr1[3];
1031 make_unicode (w0, w0_t, w1_t);
1032 make_unicode (w1, w2_t, w3_t);
1034 w3_t[2] = pw_len * 8 * 2;
1041 MD4_STEP (MD4_Fo, a, b, c, d, w0_t[0], MD4C00, MD4S00);
1042 MD4_STEP (MD4_Fo, d, a, b, c, w0_t[1], MD4C00, MD4S01);
1043 MD4_STEP (MD4_Fo, c, d, a, b, w0_t[2], MD4C00, MD4S02);
1044 MD4_STEP (MD4_Fo, b, c, d, a, w0_t[3], MD4C00, MD4S03);
1045 MD4_STEP (MD4_Fo, a, b, c, d, w1_t[0], MD4C00, MD4S00);
1046 MD4_STEP (MD4_Fo, d, a, b, c, w1_t[1], MD4C00, MD4S01);
1047 MD4_STEP (MD4_Fo, c, d, a, b, w1_t[2], MD4C00, MD4S02);
1048 MD4_STEP (MD4_Fo, b, c, d, a, w1_t[3], MD4C00, MD4S03);
1049 MD4_STEP (MD4_Fo, a, b, c, d, w2_t[0], MD4C00, MD4S00);
1050 MD4_STEP (MD4_Fo, d, a, b, c, w2_t[1], MD4C00, MD4S01);
1051 MD4_STEP (MD4_Fo, c, d, a, b, w2_t[2], MD4C00, MD4S02);
1052 MD4_STEP (MD4_Fo, b, c, d, a, w2_t[3], MD4C00, MD4S03);
1053 MD4_STEP (MD4_Fo, a, b, c, d, w3_t[0], MD4C00, MD4S00);
1054 MD4_STEP (MD4_Fo, d, a, b, c, w3_t[1], MD4C00, MD4S01);
1055 MD4_STEP (MD4_Fo, c, d, a, b, w3_t[2], MD4C00, MD4S02);
1056 MD4_STEP (MD4_Fo, b, c, d, a, w3_t[3], MD4C00, MD4S03);
1058 MD4_STEP (MD4_Go, a, b, c, d, w0_t[0], MD4C01, MD4S10);
1059 MD4_STEP (MD4_Go, d, a, b, c, w1_t[0], MD4C01, MD4S11);
1060 MD4_STEP (MD4_Go, c, d, a, b, w2_t[0], MD4C01, MD4S12);
1061 MD4_STEP (MD4_Go, b, c, d, a, w3_t[0], MD4C01, MD4S13);
1062 MD4_STEP (MD4_Go, a, b, c, d, w0_t[1], MD4C01, MD4S10);
1063 MD4_STEP (MD4_Go, d, a, b, c, w1_t[1], MD4C01, MD4S11);
1064 MD4_STEP (MD4_Go, c, d, a, b, w2_t[1], MD4C01, MD4S12);
1065 MD4_STEP (MD4_Go, b, c, d, a, w3_t[1], MD4C01, MD4S13);
1066 MD4_STEP (MD4_Go, a, b, c, d, w0_t[2], MD4C01, MD4S10);
1067 MD4_STEP (MD4_Go, d, a, b, c, w1_t[2], MD4C01, MD4S11);
1068 MD4_STEP (MD4_Go, c, d, a, b, w2_t[2], MD4C01, MD4S12);
1069 MD4_STEP (MD4_Go, b, c, d, a, w3_t[2], MD4C01, MD4S13);
1070 MD4_STEP (MD4_Go, a, b, c, d, w0_t[3], MD4C01, MD4S10);
1071 MD4_STEP (MD4_Go, d, a, b, c, w1_t[3], MD4C01, MD4S11);
1072 MD4_STEP (MD4_Go, c, d, a, b, w2_t[3], MD4C01, MD4S12);
1073 MD4_STEP (MD4_Go, b, c, d, a, w3_t[3], MD4C01, MD4S13);
1075 MD4_STEP (MD4_H , a, b, c, d, w0_t[0], MD4C02, MD4S20);
1076 MD4_STEP (MD4_H , d, a, b, c, w2_t[0], MD4C02, MD4S21);
1077 MD4_STEP (MD4_H , c, d, a, b, w1_t[0], MD4C02, MD4S22);
1078 MD4_STEP (MD4_H , b, c, d, a, w3_t[0], MD4C02, MD4S23);
1079 MD4_STEP (MD4_H , a, b, c, d, w0_t[2], MD4C02, MD4S20);
1080 MD4_STEP (MD4_H , d, a, b, c, w2_t[2], MD4C02, MD4S21);
1081 MD4_STEP (MD4_H , c, d, a, b, w1_t[2], MD4C02, MD4S22);
1082 MD4_STEP (MD4_H , b, c, d, a, w3_t[2], MD4C02, MD4S23);
1083 MD4_STEP (MD4_H , a, b, c, d, w0_t[1], MD4C02, MD4S20);
1084 MD4_STEP (MD4_H , d, a, b, c, w2_t[1], MD4C02, MD4S21);
1085 MD4_STEP (MD4_H , c, d, a, b, w1_t[1], MD4C02, MD4S22);
1086 MD4_STEP (MD4_H , b, c, d, a, w3_t[1], MD4C02, MD4S23);
1087 MD4_STEP (MD4_H , a, b, c, d, w0_t[3], MD4C02, MD4S20);
1088 MD4_STEP (MD4_H , d, a, b, c, w2_t[3], MD4C02, MD4S21);
1090 if (allx (s2 != ((d + MD4M_D) >> 16))) continue;
1092 MD4_STEP (MD4_H , c, d, a, b, w1_t[3], MD4C02, MD4S22);
1093 MD4_STEP (MD4_H , b, c, d, a, w3_t[3], MD4C02, MD4S23);
1106 transform_netntlmv1_key (a, b, key);
1111 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
1115 _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
1122 transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
1124 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
1128 _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
1140 const u32x r0 = iv1[0];
1141 const u32x r1 = iv1[1];
1142 const u32x r2 = iv2[0];
1143 const u32x r3 = iv2[1];
1145 #include VECT_COMPARE_S
1149 __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)
1153 __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)