2 * Author......: Jens Steube <jens.steube@gmail.com>
9 #include "include/constants.h"
10 #include "include/kernel_vendor.h"
29 #include "include/kernel_functions.c"
30 #include "types_amd.c"
31 #include "common_amd.c"
34 #define VECT_COMPARE_S "check_single_vect1_comp4_warp.c"
35 #define VECT_COMPARE_M "check_multi_vect1_comp4_warp.c"
39 #define VECT_COMPARE_S "check_single_vect2_comp4_warp.c"
40 #define VECT_COMPARE_M "check_multi_vect2_comp4_warp.c"
44 #define VECT_COMPARE_S "check_single_vect4_comp4_warp.c"
45 #define VECT_COMPARE_M "check_multi_vect4_comp4_warp.c"
48 #define PERM_OP(a,b,tt,n,m) \
58 #define HPERM_OP(a,tt,n,m) \
64 tt = tt >> (16 + n); \
68 __constant u32 c_SPtrans[8][64] =
71 0x02080800, 0x00080000, 0x02000002, 0x02080802,
72 0x02000000, 0x00080802, 0x00080002, 0x02000002,
73 0x00080802, 0x02080800, 0x02080000, 0x00000802,
74 0x02000802, 0x02000000, 0x00000000, 0x00080002,
75 0x00080000, 0x00000002, 0x02000800, 0x00080800,
76 0x02080802, 0x02080000, 0x00000802, 0x02000800,
77 0x00000002, 0x00000800, 0x00080800, 0x02080002,
78 0x00000800, 0x02000802, 0x02080002, 0x00000000,
79 0x00000000, 0x02080802, 0x02000800, 0x00080002,
80 0x02080800, 0x00080000, 0x00000802, 0x02000800,
81 0x02080002, 0x00000800, 0x00080800, 0x02000002,
82 0x00080802, 0x00000002, 0x02000002, 0x02080000,
83 0x02080802, 0x00080800, 0x02080000, 0x02000802,
84 0x02000000, 0x00000802, 0x00080002, 0x00000000,
85 0x00080000, 0x02000000, 0x02000802, 0x02080800,
86 0x00000002, 0x02080002, 0x00000800, 0x00080802,
88 0x40108010, 0x00000000, 0x00108000, 0x40100000,
89 0x40000010, 0x00008010, 0x40008000, 0x00108000,
90 0x00008000, 0x40100010, 0x00000010, 0x40008000,
91 0x00100010, 0x40108000, 0x40100000, 0x00000010,
92 0x00100000, 0x40008010, 0x40100010, 0x00008000,
93 0x00108010, 0x40000000, 0x00000000, 0x00100010,
94 0x40008010, 0x00108010, 0x40108000, 0x40000010,
95 0x40000000, 0x00100000, 0x00008010, 0x40108010,
96 0x00100010, 0x40108000, 0x40008000, 0x00108010,
97 0x40108010, 0x00100010, 0x40000010, 0x00000000,
98 0x40000000, 0x00008010, 0x00100000, 0x40100010,
99 0x00008000, 0x40000000, 0x00108010, 0x40008010,
100 0x40108000, 0x00008000, 0x00000000, 0x40000010,
101 0x00000010, 0x40108010, 0x00108000, 0x40100000,
102 0x40100010, 0x00100000, 0x00008010, 0x40008000,
103 0x40008010, 0x00000010, 0x40100000, 0x00108000,
105 0x04000001, 0x04040100, 0x00000100, 0x04000101,
106 0x00040001, 0x04000000, 0x04000101, 0x00040100,
107 0x04000100, 0x00040000, 0x04040000, 0x00000001,
108 0x04040101, 0x00000101, 0x00000001, 0x04040001,
109 0x00000000, 0x00040001, 0x04040100, 0x00000100,
110 0x00000101, 0x04040101, 0x00040000, 0x04000001,
111 0x04040001, 0x04000100, 0x00040101, 0x04040000,
112 0x00040100, 0x00000000, 0x04000000, 0x00040101,
113 0x04040100, 0x00000100, 0x00000001, 0x00040000,
114 0x00000101, 0x00040001, 0x04040000, 0x04000101,
115 0x00000000, 0x04040100, 0x00040100, 0x04040001,
116 0x00040001, 0x04000000, 0x04040101, 0x00000001,
117 0x00040101, 0x04000001, 0x04000000, 0x04040101,
118 0x00040000, 0x04000100, 0x04000101, 0x00040100,
119 0x04000100, 0x00000000, 0x04040001, 0x00000101,
120 0x04000001, 0x00040101, 0x00000100, 0x04040000,
122 0x00401008, 0x10001000, 0x00000008, 0x10401008,
123 0x00000000, 0x10400000, 0x10001008, 0x00400008,
124 0x10401000, 0x10000008, 0x10000000, 0x00001008,
125 0x10000008, 0x00401008, 0x00400000, 0x10000000,
126 0x10400008, 0x00401000, 0x00001000, 0x00000008,
127 0x00401000, 0x10001008, 0x10400000, 0x00001000,
128 0x00001008, 0x00000000, 0x00400008, 0x10401000,
129 0x10001000, 0x10400008, 0x10401008, 0x00400000,
130 0x10400008, 0x00001008, 0x00400000, 0x10000008,
131 0x00401000, 0x10001000, 0x00000008, 0x10400000,
132 0x10001008, 0x00000000, 0x00001000, 0x00400008,
133 0x00000000, 0x10400008, 0x10401000, 0x00001000,
134 0x10000000, 0x10401008, 0x00401008, 0x00400000,
135 0x10401008, 0x00000008, 0x10001000, 0x00401008,
136 0x00400008, 0x00401000, 0x10400000, 0x10001008,
137 0x00001008, 0x10000000, 0x10000008, 0x10401000,
139 0x08000000, 0x00010000, 0x00000400, 0x08010420,
140 0x08010020, 0x08000400, 0x00010420, 0x08010000,
141 0x00010000, 0x00000020, 0x08000020, 0x00010400,
142 0x08000420, 0x08010020, 0x08010400, 0x00000000,
143 0x00010400, 0x08000000, 0x00010020, 0x00000420,
144 0x08000400, 0x00010420, 0x00000000, 0x08000020,
145 0x00000020, 0x08000420, 0x08010420, 0x00010020,
146 0x08010000, 0x00000400, 0x00000420, 0x08010400,
147 0x08010400, 0x08000420, 0x00010020, 0x08010000,
148 0x00010000, 0x00000020, 0x08000020, 0x08000400,
149 0x08000000, 0x00010400, 0x08010420, 0x00000000,
150 0x00010420, 0x08000000, 0x00000400, 0x00010020,
151 0x08000420, 0x00000400, 0x00000000, 0x08010420,
152 0x08010020, 0x08010400, 0x00000420, 0x00010000,
153 0x00010400, 0x08010020, 0x08000400, 0x00000420,
154 0x00000020, 0x00010420, 0x08010000, 0x08000020,
156 0x80000040, 0x00200040, 0x00000000, 0x80202000,
157 0x00200040, 0x00002000, 0x80002040, 0x00200000,
158 0x00002040, 0x80202040, 0x00202000, 0x80000000,
159 0x80002000, 0x80000040, 0x80200000, 0x00202040,
160 0x00200000, 0x80002040, 0x80200040, 0x00000000,
161 0x00002000, 0x00000040, 0x80202000, 0x80200040,
162 0x80202040, 0x80200000, 0x80000000, 0x00002040,
163 0x00000040, 0x00202000, 0x00202040, 0x80002000,
164 0x00002040, 0x80000000, 0x80002000, 0x00202040,
165 0x80202000, 0x00200040, 0x00000000, 0x80002000,
166 0x80000000, 0x00002000, 0x80200040, 0x00200000,
167 0x00200040, 0x80202040, 0x00202000, 0x00000040,
168 0x80202040, 0x00202000, 0x00200000, 0x80002040,
169 0x80000040, 0x80200000, 0x00202040, 0x00000000,
170 0x00002000, 0x80000040, 0x80002040, 0x80202000,
171 0x80200000, 0x00002040, 0x00000040, 0x80200040,
173 0x00004000, 0x00000200, 0x01000200, 0x01000004,
174 0x01004204, 0x00004004, 0x00004200, 0x00000000,
175 0x01000000, 0x01000204, 0x00000204, 0x01004000,
176 0x00000004, 0x01004200, 0x01004000, 0x00000204,
177 0x01000204, 0x00004000, 0x00004004, 0x01004204,
178 0x00000000, 0x01000200, 0x01000004, 0x00004200,
179 0x01004004, 0x00004204, 0x01004200, 0x00000004,
180 0x00004204, 0x01004004, 0x00000200, 0x01000000,
181 0x00004204, 0x01004000, 0x01004004, 0x00000204,
182 0x00004000, 0x00000200, 0x01000000, 0x01004004,
183 0x01000204, 0x00004204, 0x00004200, 0x00000000,
184 0x00000200, 0x01000004, 0x00000004, 0x01000200,
185 0x00000000, 0x01000204, 0x01000200, 0x00004200,
186 0x00000204, 0x00004000, 0x01004204, 0x01000000,
187 0x01004200, 0x00000004, 0x00004004, 0x01004204,
188 0x01000004, 0x01004200, 0x01004000, 0x00004004,
190 0x20800080, 0x20820000, 0x00020080, 0x00000000,
191 0x20020000, 0x00800080, 0x20800000, 0x20820080,
192 0x00000080, 0x20000000, 0x00820000, 0x00020080,
193 0x00820080, 0x20020080, 0x20000080, 0x20800000,
194 0x00020000, 0x00820080, 0x00800080, 0x20020000,
195 0x20820080, 0x20000080, 0x00000000, 0x00820000,
196 0x20000000, 0x00800000, 0x20020080, 0x20800080,
197 0x00800000, 0x00020000, 0x20820000, 0x00000080,
198 0x00800000, 0x00020000, 0x20000080, 0x20820080,
199 0x00020080, 0x20000000, 0x00000000, 0x00820000,
200 0x20800080, 0x20020080, 0x20020000, 0x00800080,
201 0x20820000, 0x00000080, 0x00800080, 0x20020000,
202 0x20820080, 0x00800000, 0x20800000, 0x20000080,
203 0x00820000, 0x00020080, 0x20020080, 0x20800000,
204 0x00000080, 0x20820000, 0x00820080, 0x00000000,
205 0x20000000, 0x20800080, 0x00020000, 0x00820080,
208 __constant u32 c_skb[8][64] =
210 /* for C bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
211 0x00000000, 0x00000010, 0x20000000, 0x20000010,
212 0x00010000, 0x00010010, 0x20010000, 0x20010010,
213 0x00000800, 0x00000810, 0x20000800, 0x20000810,
214 0x00010800, 0x00010810, 0x20010800, 0x20010810,
215 0x00000020, 0x00000030, 0x20000020, 0x20000030,
216 0x00010020, 0x00010030, 0x20010020, 0x20010030,
217 0x00000820, 0x00000830, 0x20000820, 0x20000830,
218 0x00010820, 0x00010830, 0x20010820, 0x20010830,
219 0x00080000, 0x00080010, 0x20080000, 0x20080010,
220 0x00090000, 0x00090010, 0x20090000, 0x20090010,
221 0x00080800, 0x00080810, 0x20080800, 0x20080810,
222 0x00090800, 0x00090810, 0x20090800, 0x20090810,
223 0x00080020, 0x00080030, 0x20080020, 0x20080030,
224 0x00090020, 0x00090030, 0x20090020, 0x20090030,
225 0x00080820, 0x00080830, 0x20080820, 0x20080830,
226 0x00090820, 0x00090830, 0x20090820, 0x20090830,
227 /* for C bits (numbered as per FIPS 46) 7 8 10 11 12 13 */
228 0x00000000, 0x02000000, 0x00002000, 0x02002000,
229 0x00200000, 0x02200000, 0x00202000, 0x02202000,
230 0x00000004, 0x02000004, 0x00002004, 0x02002004,
231 0x00200004, 0x02200004, 0x00202004, 0x02202004,
232 0x00000400, 0x02000400, 0x00002400, 0x02002400,
233 0x00200400, 0x02200400, 0x00202400, 0x02202400,
234 0x00000404, 0x02000404, 0x00002404, 0x02002404,
235 0x00200404, 0x02200404, 0x00202404, 0x02202404,
236 0x10000000, 0x12000000, 0x10002000, 0x12002000,
237 0x10200000, 0x12200000, 0x10202000, 0x12202000,
238 0x10000004, 0x12000004, 0x10002004, 0x12002004,
239 0x10200004, 0x12200004, 0x10202004, 0x12202004,
240 0x10000400, 0x12000400, 0x10002400, 0x12002400,
241 0x10200400, 0x12200400, 0x10202400, 0x12202400,
242 0x10000404, 0x12000404, 0x10002404, 0x12002404,
243 0x10200404, 0x12200404, 0x10202404, 0x12202404,
244 /* for C bits (numbered as per FIPS 46) 14 15 16 17 19 20 */
245 0x00000000, 0x00000001, 0x00040000, 0x00040001,
246 0x01000000, 0x01000001, 0x01040000, 0x01040001,
247 0x00000002, 0x00000003, 0x00040002, 0x00040003,
248 0x01000002, 0x01000003, 0x01040002, 0x01040003,
249 0x00000200, 0x00000201, 0x00040200, 0x00040201,
250 0x01000200, 0x01000201, 0x01040200, 0x01040201,
251 0x00000202, 0x00000203, 0x00040202, 0x00040203,
252 0x01000202, 0x01000203, 0x01040202, 0x01040203,
253 0x08000000, 0x08000001, 0x08040000, 0x08040001,
254 0x09000000, 0x09000001, 0x09040000, 0x09040001,
255 0x08000002, 0x08000003, 0x08040002, 0x08040003,
256 0x09000002, 0x09000003, 0x09040002, 0x09040003,
257 0x08000200, 0x08000201, 0x08040200, 0x08040201,
258 0x09000200, 0x09000201, 0x09040200, 0x09040201,
259 0x08000202, 0x08000203, 0x08040202, 0x08040203,
260 0x09000202, 0x09000203, 0x09040202, 0x09040203,
261 /* for C bits (numbered as per FIPS 46) 21 23 24 26 27 28 */
262 0x00000000, 0x00100000, 0x00000100, 0x00100100,
263 0x00000008, 0x00100008, 0x00000108, 0x00100108,
264 0x00001000, 0x00101000, 0x00001100, 0x00101100,
265 0x00001008, 0x00101008, 0x00001108, 0x00101108,
266 0x04000000, 0x04100000, 0x04000100, 0x04100100,
267 0x04000008, 0x04100008, 0x04000108, 0x04100108,
268 0x04001000, 0x04101000, 0x04001100, 0x04101100,
269 0x04001008, 0x04101008, 0x04001108, 0x04101108,
270 0x00020000, 0x00120000, 0x00020100, 0x00120100,
271 0x00020008, 0x00120008, 0x00020108, 0x00120108,
272 0x00021000, 0x00121000, 0x00021100, 0x00121100,
273 0x00021008, 0x00121008, 0x00021108, 0x00121108,
274 0x04020000, 0x04120000, 0x04020100, 0x04120100,
275 0x04020008, 0x04120008, 0x04020108, 0x04120108,
276 0x04021000, 0x04121000, 0x04021100, 0x04121100,
277 0x04021008, 0x04121008, 0x04021108, 0x04121108,
278 /* for D bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
279 0x00000000, 0x10000000, 0x00010000, 0x10010000,
280 0x00000004, 0x10000004, 0x00010004, 0x10010004,
281 0x20000000, 0x30000000, 0x20010000, 0x30010000,
282 0x20000004, 0x30000004, 0x20010004, 0x30010004,
283 0x00100000, 0x10100000, 0x00110000, 0x10110000,
284 0x00100004, 0x10100004, 0x00110004, 0x10110004,
285 0x20100000, 0x30100000, 0x20110000, 0x30110000,
286 0x20100004, 0x30100004, 0x20110004, 0x30110004,
287 0x00001000, 0x10001000, 0x00011000, 0x10011000,
288 0x00001004, 0x10001004, 0x00011004, 0x10011004,
289 0x20001000, 0x30001000, 0x20011000, 0x30011000,
290 0x20001004, 0x30001004, 0x20011004, 0x30011004,
291 0x00101000, 0x10101000, 0x00111000, 0x10111000,
292 0x00101004, 0x10101004, 0x00111004, 0x10111004,
293 0x20101000, 0x30101000, 0x20111000, 0x30111000,
294 0x20101004, 0x30101004, 0x20111004, 0x30111004,
295 /* for D bits (numbered as per FIPS 46) 8 9 11 12 13 14 */
296 0x00000000, 0x08000000, 0x00000008, 0x08000008,
297 0x00000400, 0x08000400, 0x00000408, 0x08000408,
298 0x00020000, 0x08020000, 0x00020008, 0x08020008,
299 0x00020400, 0x08020400, 0x00020408, 0x08020408,
300 0x00000001, 0x08000001, 0x00000009, 0x08000009,
301 0x00000401, 0x08000401, 0x00000409, 0x08000409,
302 0x00020001, 0x08020001, 0x00020009, 0x08020009,
303 0x00020401, 0x08020401, 0x00020409, 0x08020409,
304 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
305 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
306 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
307 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
308 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
309 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
310 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
311 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
312 /* for D bits (numbered as per FIPS 46) 16 17 18 19 20 21 */
313 0x00000000, 0x00000100, 0x00080000, 0x00080100,
314 0x01000000, 0x01000100, 0x01080000, 0x01080100,
315 0x00000010, 0x00000110, 0x00080010, 0x00080110,
316 0x01000010, 0x01000110, 0x01080010, 0x01080110,
317 0x00200000, 0x00200100, 0x00280000, 0x00280100,
318 0x01200000, 0x01200100, 0x01280000, 0x01280100,
319 0x00200010, 0x00200110, 0x00280010, 0x00280110,
320 0x01200010, 0x01200110, 0x01280010, 0x01280110,
321 0x00000200, 0x00000300, 0x00080200, 0x00080300,
322 0x01000200, 0x01000300, 0x01080200, 0x01080300,
323 0x00000210, 0x00000310, 0x00080210, 0x00080310,
324 0x01000210, 0x01000310, 0x01080210, 0x01080310,
325 0x00200200, 0x00200300, 0x00280200, 0x00280300,
326 0x01200200, 0x01200300, 0x01280200, 0x01280300,
327 0x00200210, 0x00200310, 0x00280210, 0x00280310,
328 0x01200210, 0x01200310, 0x01280210, 0x01280310,
329 /* for D bits (numbered as per FIPS 46) 22 23 24 25 27 28 */
330 0x00000000, 0x04000000, 0x00040000, 0x04040000,
331 0x00000002, 0x04000002, 0x00040002, 0x04040002,
332 0x00002000, 0x04002000, 0x00042000, 0x04042000,
333 0x00002002, 0x04002002, 0x00042002, 0x04042002,
334 0x00000020, 0x04000020, 0x00040020, 0x04040020,
335 0x00000022, 0x04000022, 0x00040022, 0x04040022,
336 0x00002020, 0x04002020, 0x00042020, 0x04042020,
337 0x00002022, 0x04002022, 0x00042022, 0x04042022,
338 0x00000800, 0x04000800, 0x00040800, 0x04040800,
339 0x00000802, 0x04000802, 0x00040802, 0x04040802,
340 0x00002800, 0x04002800, 0x00042800, 0x04042800,
341 0x00002802, 0x04002802, 0x00042802, 0x04042802,
342 0x00000820, 0x04000820, 0x00040820, 0x04040820,
343 0x00000822, 0x04000822, 0x00040822, 0x04040822,
344 0x00002820, 0x04002820, 0x00042820, 0x04042820,
345 0x00002822, 0x04002822, 0x00042822, 0x04042822
348 __constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
349 __constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
352 #define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
356 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
360 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
363 static void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 s_SPtrans[8][64])
369 for (u32 i = 0; i < 16; i += 2)
374 u = Kc[i + 0] ^ rotl32 (r, 30u);
375 t = Kd[i + 0] ^ rotl32 (r, 26u);
378 ^ BOX (amd_bfe (u, 0, 6), 0, s_SPtrans)
379 ^ BOX (amd_bfe (u, 8, 6), 2, s_SPtrans)
380 ^ BOX (amd_bfe (u, 16, 6), 4, s_SPtrans)
381 ^ BOX (amd_bfe (u, 24, 6), 6, s_SPtrans)
382 ^ BOX (amd_bfe (t, 0, 6), 1, s_SPtrans)
383 ^ BOX (amd_bfe (t, 8, 6), 3, s_SPtrans)
384 ^ BOX (amd_bfe (t, 16, 6), 5, s_SPtrans)
385 ^ BOX (amd_bfe (t, 24, 6), 7, s_SPtrans);
387 u = Kc[i + 1] ^ rotl32 (l, 30u);
388 t = Kd[i + 1] ^ rotl32 (l, 26u);
391 ^ BOX (amd_bfe (u, 0, 6), 0, s_SPtrans)
392 ^ BOX (amd_bfe (u, 8, 6), 2, s_SPtrans)
393 ^ BOX (amd_bfe (u, 16, 6), 4, s_SPtrans)
394 ^ BOX (amd_bfe (u, 24, 6), 6, s_SPtrans)
395 ^ BOX (amd_bfe (t, 0, 6), 1, s_SPtrans)
396 ^ BOX (amd_bfe (t, 8, 6), 3, s_SPtrans)
397 ^ BOX (amd_bfe (t, 16, 6), 5, s_SPtrans)
398 ^ BOX (amd_bfe (t, 24, 6), 7, s_SPtrans);
405 static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 s_skb[8][64])
409 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
410 HPERM_OP (c, tt, 2, 0xcccc0000);
411 HPERM_OP (d, tt, 2, 0xcccc0000);
412 PERM_OP (d, c, tt, 1, 0x55555555);
413 PERM_OP (c, d, tt, 8, 0x00ff00ff);
414 PERM_OP (d, c, tt, 1, 0x55555555);
416 d = ((d & 0x000000ff) << 16)
417 | ((d & 0x0000ff00) << 0)
418 | ((d & 0x00ff0000) >> 16)
419 | ((c & 0xf0000000) >> 4);
424 for (u32 i = 0; i < 16; i++)
426 c = c >> shifts3s0[i] | c << shifts3s1[i];
427 d = d >> shifts3s0[i] | d << shifts3s1[i];
432 const u32x c00 = (c >> 0) & 0x0000003f;
433 const u32x c06 = (c >> 6) & 0x00383003;
434 const u32x c07 = (c >> 7) & 0x0000003c;
435 const u32x c13 = (c >> 13) & 0x0000060f;
436 const u32x c20 = (c >> 20) & 0x00000001;
438 u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
439 | BOX (((c06 >> 0) & 0xff)
440 |((c07 >> 0) & 0xff), 1, s_skb)
441 | BOX (((c13 >> 0) & 0xff)
442 |((c06 >> 8) & 0xff), 2, s_skb)
443 | BOX (((c20 >> 0) & 0xff)
445 |((c06 >> 16) & 0xff), 3, s_skb);
447 const u32x d00 = (d >> 0) & 0x00003c3f;
448 const u32x d07 = (d >> 7) & 0x00003f03;
449 const u32x d21 = (d >> 21) & 0x0000000f;
450 const u32x d22 = (d >> 22) & 0x00000030;
452 u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
453 | BOX (((d07 >> 0) & 0xff)
454 |((d00 >> 8) & 0xff), 5, s_skb)
455 | BOX (((d07 >> 8) & 0xff), 6, s_skb)
456 | BOX (((d21 >> 0) & 0xff)
457 |((d22 >> 0) & 0xff), 7, s_skb);
459 Kc[i] = ((t << 16) | (s & 0x0000ffff));
460 Kd[i] = ((s >> 16) | (t & 0xffff0000));
464 static void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
467 const uchar4 t0 = as_uchar4 (w0);
468 const uchar4 t1 = as_uchar4 (w1);
473 k0.s0 = (t0.s0 >> 0);
474 k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
475 k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
476 k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
477 k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
478 k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
479 k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
480 k1.s3 = (t1.s2 << 1);
482 out[0] = as_uint (k0);
483 out[1] = as_uint (k1);
487 const uchar8 t0 = as_uchar8 (w0);
488 const uchar8 t1 = as_uchar8 (w1);
493 k0.s0 = (t0.s0 >> 0);
494 k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
495 k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
496 k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
497 k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
498 k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
499 k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
500 k1.s3 = (t1.s2 << 1);
502 k0.s4 = (t0.s4 >> 0);
503 k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
504 k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
505 k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
506 k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
507 k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
508 k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
509 k1.s7 = (t1.s6 << 1);
511 out[0] = as_uint2 (k0);
512 out[1] = as_uint2 (k1);
516 const uchar16 t0 = as_uchar16 (w0);
517 const uchar16 t1 = as_uchar16 (w1);
522 k0.s0 = (t0.s0 >> 0);
523 k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
524 k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
525 k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
526 k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
527 k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
528 k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
529 k1.s3 = (t1.s2 << 1);
531 k0.s4 = (t0.s4 >> 0);
532 k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
533 k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
534 k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
535 k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
536 k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
537 k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
538 k1.s7 = (t1.s6 << 1);
540 k0.s8 = (t0.s8 >> 0);
541 k0.s9 = (t0.s8 << 7) | (t0.s9 >> 1);
542 k0.sa = (t0.s9 << 6) | (t0.sa >> 2);
543 k0.sb = (t0.sa << 5) | (t0.sb >> 3);
544 k1.s8 = (t0.sb << 4) | (t1.s8 >> 4);
545 k1.s9 = (t1.s8 << 3) | (t1.s9 >> 5);
546 k1.sa = (t1.s9 << 2) | (t1.sa >> 6);
547 k1.sb = (t1.sa << 1);
549 k0.sc = (t0.sc >> 0);
550 k0.sd = (t0.sc << 7) | (t0.sd >> 1);
551 k0.se = (t0.sd << 6) | (t0.se >> 2);
552 k0.sf = (t0.se << 5) | (t0.sf >> 3);
553 k1.sc = (t0.sf << 4) | (t1.sc >> 4);
554 k1.sd = (t1.sc << 3) | (t1.sd >> 5);
555 k1.se = (t1.sd << 2) | (t1.se >> 6);
556 k1.sf = (t1.se << 1);
558 out[0] = as_uint4 (k0);
559 out[1] = as_uint4 (k1);
563 static void m05500m (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
569 const u32 gid = get_global_id (0);
570 const u32 lid = get_local_id (0);
576 const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
577 const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
578 const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
589 const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
593 for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
595 const u32x w0r = words_buf_r[il_pos];
597 const u32x w0 = w0l | w0r;
621 MD4_STEP (MD4_Fo, a, b, c, d, w0_t, MD4C00, MD4S00);
622 MD4_STEP (MD4_Fo, d, a, b, c, w1_t, MD4C00, MD4S01);
623 MD4_STEP (MD4_Fo, c, d, a, b, w2_t, MD4C00, MD4S02);
624 MD4_STEP (MD4_Fo, b, c, d, a, w3_t, MD4C00, MD4S03);
625 MD4_STEP (MD4_Fo, a, b, c, d, w4_t, MD4C00, MD4S00);
626 MD4_STEP (MD4_Fo, d, a, b, c, w5_t, MD4C00, MD4S01);
627 MD4_STEP (MD4_Fo, c, d, a, b, w6_t, MD4C00, MD4S02);
628 MD4_STEP (MD4_Fo, b, c, d, a, w7_t, MD4C00, MD4S03);
629 MD4_STEP (MD4_Fo, a, b, c, d, w8_t, MD4C00, MD4S00);
630 MD4_STEP (MD4_Fo, d, a, b, c, w9_t, MD4C00, MD4S01);
631 MD4_STEP (MD4_Fo, c, d, a, b, wa_t, MD4C00, MD4S02);
632 MD4_STEP (MD4_Fo, b, c, d, a, wb_t, MD4C00, MD4S03);
633 MD4_STEP (MD4_Fo, a, b, c, d, wc_t, MD4C00, MD4S00);
634 MD4_STEP (MD4_Fo, d, a, b, c, wd_t, MD4C00, MD4S01);
635 MD4_STEP (MD4_Fo, c, d, a, b, we_t, MD4C00, MD4S02);
636 MD4_STEP (MD4_Fo, b, c, d, a, wf_t, MD4C00, MD4S03);
638 MD4_STEP (MD4_Go, a, b, c, d, w0_t, MD4C01, MD4S10);
639 MD4_STEP (MD4_Go, d, a, b, c, w4_t, MD4C01, MD4S11);
640 MD4_STEP (MD4_Go, c, d, a, b, w8_t, MD4C01, MD4S12);
641 MD4_STEP (MD4_Go, b, c, d, a, wc_t, MD4C01, MD4S13);
642 MD4_STEP (MD4_Go, a, b, c, d, w1_t, MD4C01, MD4S10);
643 MD4_STEP (MD4_Go, d, a, b, c, w5_t, MD4C01, MD4S11);
644 MD4_STEP (MD4_Go, c, d, a, b, w9_t, MD4C01, MD4S12);
645 MD4_STEP (MD4_Go, b, c, d, a, wd_t, MD4C01, MD4S13);
646 MD4_STEP (MD4_Go, a, b, c, d, w2_t, MD4C01, MD4S10);
647 MD4_STEP (MD4_Go, d, a, b, c, w6_t, MD4C01, MD4S11);
648 MD4_STEP (MD4_Go, c, d, a, b, wa_t, MD4C01, MD4S12);
649 MD4_STEP (MD4_Go, b, c, d, a, we_t, MD4C01, MD4S13);
650 MD4_STEP (MD4_Go, a, b, c, d, w3_t, MD4C01, MD4S10);
651 MD4_STEP (MD4_Go, d, a, b, c, w7_t, MD4C01, MD4S11);
652 MD4_STEP (MD4_Go, c, d, a, b, wb_t, MD4C01, MD4S12);
653 MD4_STEP (MD4_Go, b, c, d, a, wf_t, MD4C01, MD4S13);
655 MD4_STEP (MD4_H , a, b, c, d, w0_t, MD4C02, MD4S20);
656 MD4_STEP (MD4_H , d, a, b, c, w8_t, MD4C02, MD4S21);
657 MD4_STEP (MD4_H , c, d, a, b, w4_t, MD4C02, MD4S22);
658 MD4_STEP (MD4_H , b, c, d, a, wc_t, MD4C02, MD4S23);
659 MD4_STEP (MD4_H , a, b, c, d, w2_t, MD4C02, MD4S20);
660 MD4_STEP (MD4_H , d, a, b, c, wa_t, MD4C02, MD4S21);
661 MD4_STEP (MD4_H , c, d, a, b, w6_t, MD4C02, MD4S22);
662 MD4_STEP (MD4_H , b, c, d, a, we_t, MD4C02, MD4S23);
663 MD4_STEP (MD4_H , a, b, c, d, w1_t, MD4C02, MD4S20);
664 MD4_STEP (MD4_H , d, a, b, c, w9_t, MD4C02, MD4S21);
665 MD4_STEP (MD4_H , c, d, a, b, w5_t, MD4C02, MD4S22);
666 MD4_STEP (MD4_H , b, c, d, a, wd_t, MD4C02, MD4S23);
667 MD4_STEP (MD4_H , a, b, c, d, w3_t, MD4C02, MD4S20);
668 MD4_STEP (MD4_H , d, a, b, c, wb_t, MD4C02, MD4S21);
670 if (allx (s2 != ((d + MD4M_D) >> 16))) continue;
672 MD4_STEP (MD4_H , c, d, a, b, w7_t, MD4C02, MD4S22);
673 MD4_STEP (MD4_H , b, c, d, a, wf_t, MD4C02, MD4S23);
686 transform_netntlmv1_key (a, b, key);
691 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
695 _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
701 transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
703 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
707 _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
713 const u32x r0 = iv1[0];
714 const u32x r1 = iv1[1];
715 const u32x r2 = iv2[0];
716 const u32x r3 = iv2[1];
718 #include VECT_COMPARE_M
722 static void m05500s (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
728 const u32 gid = get_global_id (0);
729 const u32 lid = get_local_id (0);
735 const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
736 const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
737 const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
748 const u32 search[4] =
750 digests_buf[digests_offset].digest_buf[DGST_R0],
751 digests_buf[digests_offset].digest_buf[DGST_R1],
752 digests_buf[digests_offset].digest_buf[DGST_R2],
753 digests_buf[digests_offset].digest_buf[DGST_R3]
760 const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
764 for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
766 const u32x w0r = words_buf_r[il_pos];
768 const u32x w0 = w0l | w0r;
792 MD4_STEP (MD4_Fo, a, b, c, d, w0_t, MD4C00, MD4S00);
793 MD4_STEP (MD4_Fo, d, a, b, c, w1_t, MD4C00, MD4S01);
794 MD4_STEP (MD4_Fo, c, d, a, b, w2_t, MD4C00, MD4S02);
795 MD4_STEP (MD4_Fo, b, c, d, a, w3_t, MD4C00, MD4S03);
796 MD4_STEP (MD4_Fo, a, b, c, d, w4_t, MD4C00, MD4S00);
797 MD4_STEP (MD4_Fo, d, a, b, c, w5_t, MD4C00, MD4S01);
798 MD4_STEP (MD4_Fo, c, d, a, b, w6_t, MD4C00, MD4S02);
799 MD4_STEP (MD4_Fo, b, c, d, a, w7_t, MD4C00, MD4S03);
800 MD4_STEP (MD4_Fo, a, b, c, d, w8_t, MD4C00, MD4S00);
801 MD4_STEP (MD4_Fo, d, a, b, c, w9_t, MD4C00, MD4S01);
802 MD4_STEP (MD4_Fo, c, d, a, b, wa_t, MD4C00, MD4S02);
803 MD4_STEP (MD4_Fo, b, c, d, a, wb_t, MD4C00, MD4S03);
804 MD4_STEP (MD4_Fo, a, b, c, d, wc_t, MD4C00, MD4S00);
805 MD4_STEP (MD4_Fo, d, a, b, c, wd_t, MD4C00, MD4S01);
806 MD4_STEP (MD4_Fo, c, d, a, b, we_t, MD4C00, MD4S02);
807 MD4_STEP (MD4_Fo, b, c, d, a, wf_t, MD4C00, MD4S03);
809 MD4_STEP (MD4_Go, a, b, c, d, w0_t, MD4C01, MD4S10);
810 MD4_STEP (MD4_Go, d, a, b, c, w4_t, MD4C01, MD4S11);
811 MD4_STEP (MD4_Go, c, d, a, b, w8_t, MD4C01, MD4S12);
812 MD4_STEP (MD4_Go, b, c, d, a, wc_t, MD4C01, MD4S13);
813 MD4_STEP (MD4_Go, a, b, c, d, w1_t, MD4C01, MD4S10);
814 MD4_STEP (MD4_Go, d, a, b, c, w5_t, MD4C01, MD4S11);
815 MD4_STEP (MD4_Go, c, d, a, b, w9_t, MD4C01, MD4S12);
816 MD4_STEP (MD4_Go, b, c, d, a, wd_t, MD4C01, MD4S13);
817 MD4_STEP (MD4_Go, a, b, c, d, w2_t, MD4C01, MD4S10);
818 MD4_STEP (MD4_Go, d, a, b, c, w6_t, MD4C01, MD4S11);
819 MD4_STEP (MD4_Go, c, d, a, b, wa_t, MD4C01, MD4S12);
820 MD4_STEP (MD4_Go, b, c, d, a, we_t, MD4C01, MD4S13);
821 MD4_STEP (MD4_Go, a, b, c, d, w3_t, MD4C01, MD4S10);
822 MD4_STEP (MD4_Go, d, a, b, c, w7_t, MD4C01, MD4S11);
823 MD4_STEP (MD4_Go, c, d, a, b, wb_t, MD4C01, MD4S12);
824 MD4_STEP (MD4_Go, b, c, d, a, wf_t, MD4C01, MD4S13);
826 MD4_STEP (MD4_H , a, b, c, d, w0_t, MD4C02, MD4S20);
827 MD4_STEP (MD4_H , d, a, b, c, w8_t, MD4C02, MD4S21);
828 MD4_STEP (MD4_H , c, d, a, b, w4_t, MD4C02, MD4S22);
829 MD4_STEP (MD4_H , b, c, d, a, wc_t, MD4C02, MD4S23);
830 MD4_STEP (MD4_H , a, b, c, d, w2_t, MD4C02, MD4S20);
831 MD4_STEP (MD4_H , d, a, b, c, wa_t, MD4C02, MD4S21);
832 MD4_STEP (MD4_H , c, d, a, b, w6_t, MD4C02, MD4S22);
833 MD4_STEP (MD4_H , b, c, d, a, we_t, MD4C02, MD4S23);
834 MD4_STEP (MD4_H , a, b, c, d, w1_t, MD4C02, MD4S20);
835 MD4_STEP (MD4_H , d, a, b, c, w9_t, MD4C02, MD4S21);
836 MD4_STEP (MD4_H , c, d, a, b, w5_t, MD4C02, MD4S22);
837 MD4_STEP (MD4_H , b, c, d, a, wd_t, MD4C02, MD4S23);
838 MD4_STEP (MD4_H , a, b, c, d, w3_t, MD4C02, MD4S20);
839 MD4_STEP (MD4_H , d, a, b, c, wb_t, MD4C02, MD4S21);
841 if (allx (s2 != ((d + MD4M_D) >> 16))) continue;
843 MD4_STEP (MD4_H , c, d, a, b, w7_t, MD4C02, MD4S22);
844 MD4_STEP (MD4_H , b, c, d, a, wf_t, MD4C02, MD4S23);
857 transform_netntlmv1_key (a, b, key);
862 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
866 _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
873 transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
875 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
879 _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
891 const u32x r0 = iv1[0];
892 const u32x r1 = iv1[1];
893 const u32x r2 = iv2[0];
894 const u32x r3 = iv2[1];
896 #include VECT_COMPARE_S
900 __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 u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
906 const u32 gid = get_global_id (0);
907 const u32 lid = get_local_id (0);
911 w[ 0] = pws[gid].i[ 0];
912 w[ 1] = pws[gid].i[ 1];
913 w[ 2] = pws[gid].i[ 2];
914 w[ 3] = pws[gid].i[ 3];
925 w[14] = pws[gid].i[14];
928 const u32 pw_len = pws[gid].pw_len;
934 __local u32 s_SPtrans[8][64];
935 __local u32 s_skb[8][64];
937 s_SPtrans[0][lid] = c_SPtrans[0][lid];
938 s_SPtrans[1][lid] = c_SPtrans[1][lid];
939 s_SPtrans[2][lid] = c_SPtrans[2][lid];
940 s_SPtrans[3][lid] = c_SPtrans[3][lid];
941 s_SPtrans[4][lid] = c_SPtrans[4][lid];
942 s_SPtrans[5][lid] = c_SPtrans[5][lid];
943 s_SPtrans[6][lid] = c_SPtrans[6][lid];
944 s_SPtrans[7][lid] = c_SPtrans[7][lid];
946 s_skb[0][lid] = c_skb[0][lid];
947 s_skb[1][lid] = c_skb[1][lid];
948 s_skb[2][lid] = c_skb[2][lid];
949 s_skb[3][lid] = c_skb[3][lid];
950 s_skb[4][lid] = c_skb[4][lid];
951 s_skb[5][lid] = c_skb[5][lid];
952 s_skb[6][lid] = c_skb[6][lid];
953 s_skb[7][lid] = c_skb[7][lid];
955 barrier (CLK_LOCAL_MEM_FENCE);
957 if (gid >= gid_max) return;
963 m05500m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
966 __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 u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
972 const u32 gid = get_global_id (0);
973 const u32 lid = get_local_id (0);
977 w[ 0] = pws[gid].i[ 0];
978 w[ 1] = pws[gid].i[ 1];
979 w[ 2] = pws[gid].i[ 2];
980 w[ 3] = pws[gid].i[ 3];
981 w[ 4] = pws[gid].i[ 4];
982 w[ 5] = pws[gid].i[ 5];
983 w[ 6] = pws[gid].i[ 6];
984 w[ 7] = pws[gid].i[ 7];
991 w[14] = pws[gid].i[14];
994 const u32 pw_len = pws[gid].pw_len;
1000 __local u32 s_SPtrans[8][64];
1001 __local u32 s_skb[8][64];
1003 s_SPtrans[0][lid] = c_SPtrans[0][lid];
1004 s_SPtrans[1][lid] = c_SPtrans[1][lid];
1005 s_SPtrans[2][lid] = c_SPtrans[2][lid];
1006 s_SPtrans[3][lid] = c_SPtrans[3][lid];
1007 s_SPtrans[4][lid] = c_SPtrans[4][lid];
1008 s_SPtrans[5][lid] = c_SPtrans[5][lid];
1009 s_SPtrans[6][lid] = c_SPtrans[6][lid];
1010 s_SPtrans[7][lid] = c_SPtrans[7][lid];
1012 s_skb[0][lid] = c_skb[0][lid];
1013 s_skb[1][lid] = c_skb[1][lid];
1014 s_skb[2][lid] = c_skb[2][lid];
1015 s_skb[3][lid] = c_skb[3][lid];
1016 s_skb[4][lid] = c_skb[4][lid];
1017 s_skb[5][lid] = c_skb[5][lid];
1018 s_skb[6][lid] = c_skb[6][lid];
1019 s_skb[7][lid] = c_skb[7][lid];
1021 barrier (CLK_LOCAL_MEM_FENCE);
1023 if (gid >= gid_max) return;
1029 m05500m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
1032 __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 u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1036 __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 u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1042 const u32 gid = get_global_id (0);
1043 const u32 lid = get_local_id (0);
1047 w[ 0] = pws[gid].i[ 0];
1048 w[ 1] = pws[gid].i[ 1];
1049 w[ 2] = pws[gid].i[ 2];
1050 w[ 3] = pws[gid].i[ 3];
1061 w[14] = pws[gid].i[14];
1064 const u32 pw_len = pws[gid].pw_len;
1070 __local u32 s_SPtrans[8][64];
1071 __local u32 s_skb[8][64];
1073 s_SPtrans[0][lid] = c_SPtrans[0][lid];
1074 s_SPtrans[1][lid] = c_SPtrans[1][lid];
1075 s_SPtrans[2][lid] = c_SPtrans[2][lid];
1076 s_SPtrans[3][lid] = c_SPtrans[3][lid];
1077 s_SPtrans[4][lid] = c_SPtrans[4][lid];
1078 s_SPtrans[5][lid] = c_SPtrans[5][lid];
1079 s_SPtrans[6][lid] = c_SPtrans[6][lid];
1080 s_SPtrans[7][lid] = c_SPtrans[7][lid];
1082 s_skb[0][lid] = c_skb[0][lid];
1083 s_skb[1][lid] = c_skb[1][lid];
1084 s_skb[2][lid] = c_skb[2][lid];
1085 s_skb[3][lid] = c_skb[3][lid];
1086 s_skb[4][lid] = c_skb[4][lid];
1087 s_skb[5][lid] = c_skb[5][lid];
1088 s_skb[6][lid] = c_skb[6][lid];
1089 s_skb[7][lid] = c_skb[7][lid];
1091 barrier (CLK_LOCAL_MEM_FENCE);
1093 if (gid >= gid_max) return;
1099 m05500s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
1102 __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 u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1108 const u32 gid = get_global_id (0);
1109 const u32 lid = get_local_id (0);
1113 w[ 0] = pws[gid].i[ 0];
1114 w[ 1] = pws[gid].i[ 1];
1115 w[ 2] = pws[gid].i[ 2];
1116 w[ 3] = pws[gid].i[ 3];
1117 w[ 4] = pws[gid].i[ 4];
1118 w[ 5] = pws[gid].i[ 5];
1119 w[ 6] = pws[gid].i[ 6];
1120 w[ 7] = pws[gid].i[ 7];
1127 w[14] = pws[gid].i[14];
1130 const u32 pw_len = pws[gid].pw_len;
1136 __local u32 s_SPtrans[8][64];
1137 __local u32 s_skb[8][64];
1139 s_SPtrans[0][lid] = c_SPtrans[0][lid];
1140 s_SPtrans[1][lid] = c_SPtrans[1][lid];
1141 s_SPtrans[2][lid] = c_SPtrans[2][lid];
1142 s_SPtrans[3][lid] = c_SPtrans[3][lid];
1143 s_SPtrans[4][lid] = c_SPtrans[4][lid];
1144 s_SPtrans[5][lid] = c_SPtrans[5][lid];
1145 s_SPtrans[6][lid] = c_SPtrans[6][lid];
1146 s_SPtrans[7][lid] = c_SPtrans[7][lid];
1148 s_skb[0][lid] = c_skb[0][lid];
1149 s_skb[1][lid] = c_skb[1][lid];
1150 s_skb[2][lid] = c_skb[2][lid];
1151 s_skb[3][lid] = c_skb[3][lid];
1152 s_skb[4][lid] = c_skb[4][lid];
1153 s_skb[5][lid] = c_skb[5][lid];
1154 s_skb[6][lid] = c_skb[6][lid];
1155 s_skb[7][lid] = c_skb[7][lid];
1157 barrier (CLK_LOCAL_MEM_FENCE);
1159 if (gid >= gid_max) return;
1165 m05500s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
1168 __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 u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)