2 * Author......: Jens Steube <jens.steube@gmail.com>
9 #include "include/constants.h"
10 #include "include/kernel_vendor.h"
25 #include "include/kernel_functions.c"
27 #include "common_nv.c"
30 #define VECT_COMPARE_S "check_single_vect1_comp4_warp.c"
31 #define VECT_COMPARE_M "check_multi_vect1_comp4_warp.c"
35 #define VECT_COMPARE_S "check_single_vect2_comp4_warp.c"
36 #define VECT_COMPARE_M "check_multi_vect2_comp4_warp.c"
39 #define PERM_OP(a,b,tt,n,m) \
49 #define HPERM_OP(a,tt,n,m) \
55 tt = tt >> (16 + n); \
59 __device__ __constant__ u32 c_SPtrans[8][64] =
62 0x02080800, 0x00080000, 0x02000002, 0x02080802,
63 0x02000000, 0x00080802, 0x00080002, 0x02000002,
64 0x00080802, 0x02080800, 0x02080000, 0x00000802,
65 0x02000802, 0x02000000, 0x00000000, 0x00080002,
66 0x00080000, 0x00000002, 0x02000800, 0x00080800,
67 0x02080802, 0x02080000, 0x00000802, 0x02000800,
68 0x00000002, 0x00000800, 0x00080800, 0x02080002,
69 0x00000800, 0x02000802, 0x02080002, 0x00000000,
70 0x00000000, 0x02080802, 0x02000800, 0x00080002,
71 0x02080800, 0x00080000, 0x00000802, 0x02000800,
72 0x02080002, 0x00000800, 0x00080800, 0x02000002,
73 0x00080802, 0x00000002, 0x02000002, 0x02080000,
74 0x02080802, 0x00080800, 0x02080000, 0x02000802,
75 0x02000000, 0x00000802, 0x00080002, 0x00000000,
76 0x00080000, 0x02000000, 0x02000802, 0x02080800,
77 0x00000002, 0x02080002, 0x00000800, 0x00080802,
79 0x40108010, 0x00000000, 0x00108000, 0x40100000,
80 0x40000010, 0x00008010, 0x40008000, 0x00108000,
81 0x00008000, 0x40100010, 0x00000010, 0x40008000,
82 0x00100010, 0x40108000, 0x40100000, 0x00000010,
83 0x00100000, 0x40008010, 0x40100010, 0x00008000,
84 0x00108010, 0x40000000, 0x00000000, 0x00100010,
85 0x40008010, 0x00108010, 0x40108000, 0x40000010,
86 0x40000000, 0x00100000, 0x00008010, 0x40108010,
87 0x00100010, 0x40108000, 0x40008000, 0x00108010,
88 0x40108010, 0x00100010, 0x40000010, 0x00000000,
89 0x40000000, 0x00008010, 0x00100000, 0x40100010,
90 0x00008000, 0x40000000, 0x00108010, 0x40008010,
91 0x40108000, 0x00008000, 0x00000000, 0x40000010,
92 0x00000010, 0x40108010, 0x00108000, 0x40100000,
93 0x40100010, 0x00100000, 0x00008010, 0x40008000,
94 0x40008010, 0x00000010, 0x40100000, 0x00108000,
96 0x04000001, 0x04040100, 0x00000100, 0x04000101,
97 0x00040001, 0x04000000, 0x04000101, 0x00040100,
98 0x04000100, 0x00040000, 0x04040000, 0x00000001,
99 0x04040101, 0x00000101, 0x00000001, 0x04040001,
100 0x00000000, 0x00040001, 0x04040100, 0x00000100,
101 0x00000101, 0x04040101, 0x00040000, 0x04000001,
102 0x04040001, 0x04000100, 0x00040101, 0x04040000,
103 0x00040100, 0x00000000, 0x04000000, 0x00040101,
104 0x04040100, 0x00000100, 0x00000001, 0x00040000,
105 0x00000101, 0x00040001, 0x04040000, 0x04000101,
106 0x00000000, 0x04040100, 0x00040100, 0x04040001,
107 0x00040001, 0x04000000, 0x04040101, 0x00000001,
108 0x00040101, 0x04000001, 0x04000000, 0x04040101,
109 0x00040000, 0x04000100, 0x04000101, 0x00040100,
110 0x04000100, 0x00000000, 0x04040001, 0x00000101,
111 0x04000001, 0x00040101, 0x00000100, 0x04040000,
113 0x00401008, 0x10001000, 0x00000008, 0x10401008,
114 0x00000000, 0x10400000, 0x10001008, 0x00400008,
115 0x10401000, 0x10000008, 0x10000000, 0x00001008,
116 0x10000008, 0x00401008, 0x00400000, 0x10000000,
117 0x10400008, 0x00401000, 0x00001000, 0x00000008,
118 0x00401000, 0x10001008, 0x10400000, 0x00001000,
119 0x00001008, 0x00000000, 0x00400008, 0x10401000,
120 0x10001000, 0x10400008, 0x10401008, 0x00400000,
121 0x10400008, 0x00001008, 0x00400000, 0x10000008,
122 0x00401000, 0x10001000, 0x00000008, 0x10400000,
123 0x10001008, 0x00000000, 0x00001000, 0x00400008,
124 0x00000000, 0x10400008, 0x10401000, 0x00001000,
125 0x10000000, 0x10401008, 0x00401008, 0x00400000,
126 0x10401008, 0x00000008, 0x10001000, 0x00401008,
127 0x00400008, 0x00401000, 0x10400000, 0x10001008,
128 0x00001008, 0x10000000, 0x10000008, 0x10401000,
130 0x08000000, 0x00010000, 0x00000400, 0x08010420,
131 0x08010020, 0x08000400, 0x00010420, 0x08010000,
132 0x00010000, 0x00000020, 0x08000020, 0x00010400,
133 0x08000420, 0x08010020, 0x08010400, 0x00000000,
134 0x00010400, 0x08000000, 0x00010020, 0x00000420,
135 0x08000400, 0x00010420, 0x00000000, 0x08000020,
136 0x00000020, 0x08000420, 0x08010420, 0x00010020,
137 0x08010000, 0x00000400, 0x00000420, 0x08010400,
138 0x08010400, 0x08000420, 0x00010020, 0x08010000,
139 0x00010000, 0x00000020, 0x08000020, 0x08000400,
140 0x08000000, 0x00010400, 0x08010420, 0x00000000,
141 0x00010420, 0x08000000, 0x00000400, 0x00010020,
142 0x08000420, 0x00000400, 0x00000000, 0x08010420,
143 0x08010020, 0x08010400, 0x00000420, 0x00010000,
144 0x00010400, 0x08010020, 0x08000400, 0x00000420,
145 0x00000020, 0x00010420, 0x08010000, 0x08000020,
147 0x80000040, 0x00200040, 0x00000000, 0x80202000,
148 0x00200040, 0x00002000, 0x80002040, 0x00200000,
149 0x00002040, 0x80202040, 0x00202000, 0x80000000,
150 0x80002000, 0x80000040, 0x80200000, 0x00202040,
151 0x00200000, 0x80002040, 0x80200040, 0x00000000,
152 0x00002000, 0x00000040, 0x80202000, 0x80200040,
153 0x80202040, 0x80200000, 0x80000000, 0x00002040,
154 0x00000040, 0x00202000, 0x00202040, 0x80002000,
155 0x00002040, 0x80000000, 0x80002000, 0x00202040,
156 0x80202000, 0x00200040, 0x00000000, 0x80002000,
157 0x80000000, 0x00002000, 0x80200040, 0x00200000,
158 0x00200040, 0x80202040, 0x00202000, 0x00000040,
159 0x80202040, 0x00202000, 0x00200000, 0x80002040,
160 0x80000040, 0x80200000, 0x00202040, 0x00000000,
161 0x00002000, 0x80000040, 0x80002040, 0x80202000,
162 0x80200000, 0x00002040, 0x00000040, 0x80200040,
164 0x00004000, 0x00000200, 0x01000200, 0x01000004,
165 0x01004204, 0x00004004, 0x00004200, 0x00000000,
166 0x01000000, 0x01000204, 0x00000204, 0x01004000,
167 0x00000004, 0x01004200, 0x01004000, 0x00000204,
168 0x01000204, 0x00004000, 0x00004004, 0x01004204,
169 0x00000000, 0x01000200, 0x01000004, 0x00004200,
170 0x01004004, 0x00004204, 0x01004200, 0x00000004,
171 0x00004204, 0x01004004, 0x00000200, 0x01000000,
172 0x00004204, 0x01004000, 0x01004004, 0x00000204,
173 0x00004000, 0x00000200, 0x01000000, 0x01004004,
174 0x01000204, 0x00004204, 0x00004200, 0x00000000,
175 0x00000200, 0x01000004, 0x00000004, 0x01000200,
176 0x00000000, 0x01000204, 0x01000200, 0x00004200,
177 0x00000204, 0x00004000, 0x01004204, 0x01000000,
178 0x01004200, 0x00000004, 0x00004004, 0x01004204,
179 0x01000004, 0x01004200, 0x01004000, 0x00004004,
181 0x20800080, 0x20820000, 0x00020080, 0x00000000,
182 0x20020000, 0x00800080, 0x20800000, 0x20820080,
183 0x00000080, 0x20000000, 0x00820000, 0x00020080,
184 0x00820080, 0x20020080, 0x20000080, 0x20800000,
185 0x00020000, 0x00820080, 0x00800080, 0x20020000,
186 0x20820080, 0x20000080, 0x00000000, 0x00820000,
187 0x20000000, 0x00800000, 0x20020080, 0x20800080,
188 0x00800000, 0x00020000, 0x20820000, 0x00000080,
189 0x00800000, 0x00020000, 0x20000080, 0x20820080,
190 0x00020080, 0x20000000, 0x00000000, 0x00820000,
191 0x20800080, 0x20020080, 0x20020000, 0x00800080,
192 0x20820000, 0x00000080, 0x00800080, 0x20020000,
193 0x20820080, 0x00800000, 0x20800000, 0x20000080,
194 0x00820000, 0x00020080, 0x20020080, 0x20800000,
195 0x00000080, 0x20820000, 0x00820080, 0x00000000,
196 0x20000000, 0x20800080, 0x00020000, 0x00820080,
199 __device__ __constant__ u32 c_skb[8][64] =
201 /* for C bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
202 0x00000000, 0x00000010, 0x20000000, 0x20000010,
203 0x00010000, 0x00010010, 0x20010000, 0x20010010,
204 0x00000800, 0x00000810, 0x20000800, 0x20000810,
205 0x00010800, 0x00010810, 0x20010800, 0x20010810,
206 0x00000020, 0x00000030, 0x20000020, 0x20000030,
207 0x00010020, 0x00010030, 0x20010020, 0x20010030,
208 0x00000820, 0x00000830, 0x20000820, 0x20000830,
209 0x00010820, 0x00010830, 0x20010820, 0x20010830,
210 0x00080000, 0x00080010, 0x20080000, 0x20080010,
211 0x00090000, 0x00090010, 0x20090000, 0x20090010,
212 0x00080800, 0x00080810, 0x20080800, 0x20080810,
213 0x00090800, 0x00090810, 0x20090800, 0x20090810,
214 0x00080020, 0x00080030, 0x20080020, 0x20080030,
215 0x00090020, 0x00090030, 0x20090020, 0x20090030,
216 0x00080820, 0x00080830, 0x20080820, 0x20080830,
217 0x00090820, 0x00090830, 0x20090820, 0x20090830,
218 /* for C bits (numbered as per FIPS 46) 7 8 10 11 12 13 */
219 0x00000000, 0x02000000, 0x00002000, 0x02002000,
220 0x00200000, 0x02200000, 0x00202000, 0x02202000,
221 0x00000004, 0x02000004, 0x00002004, 0x02002004,
222 0x00200004, 0x02200004, 0x00202004, 0x02202004,
223 0x00000400, 0x02000400, 0x00002400, 0x02002400,
224 0x00200400, 0x02200400, 0x00202400, 0x02202400,
225 0x00000404, 0x02000404, 0x00002404, 0x02002404,
226 0x00200404, 0x02200404, 0x00202404, 0x02202404,
227 0x10000000, 0x12000000, 0x10002000, 0x12002000,
228 0x10200000, 0x12200000, 0x10202000, 0x12202000,
229 0x10000004, 0x12000004, 0x10002004, 0x12002004,
230 0x10200004, 0x12200004, 0x10202004, 0x12202004,
231 0x10000400, 0x12000400, 0x10002400, 0x12002400,
232 0x10200400, 0x12200400, 0x10202400, 0x12202400,
233 0x10000404, 0x12000404, 0x10002404, 0x12002404,
234 0x10200404, 0x12200404, 0x10202404, 0x12202404,
235 /* for C bits (numbered as per FIPS 46) 14 15 16 17 19 20 */
236 0x00000000, 0x00000001, 0x00040000, 0x00040001,
237 0x01000000, 0x01000001, 0x01040000, 0x01040001,
238 0x00000002, 0x00000003, 0x00040002, 0x00040003,
239 0x01000002, 0x01000003, 0x01040002, 0x01040003,
240 0x00000200, 0x00000201, 0x00040200, 0x00040201,
241 0x01000200, 0x01000201, 0x01040200, 0x01040201,
242 0x00000202, 0x00000203, 0x00040202, 0x00040203,
243 0x01000202, 0x01000203, 0x01040202, 0x01040203,
244 0x08000000, 0x08000001, 0x08040000, 0x08040001,
245 0x09000000, 0x09000001, 0x09040000, 0x09040001,
246 0x08000002, 0x08000003, 0x08040002, 0x08040003,
247 0x09000002, 0x09000003, 0x09040002, 0x09040003,
248 0x08000200, 0x08000201, 0x08040200, 0x08040201,
249 0x09000200, 0x09000201, 0x09040200, 0x09040201,
250 0x08000202, 0x08000203, 0x08040202, 0x08040203,
251 0x09000202, 0x09000203, 0x09040202, 0x09040203,
252 /* for C bits (numbered as per FIPS 46) 21 23 24 26 27 28 */
253 0x00000000, 0x00100000, 0x00000100, 0x00100100,
254 0x00000008, 0x00100008, 0x00000108, 0x00100108,
255 0x00001000, 0x00101000, 0x00001100, 0x00101100,
256 0x00001008, 0x00101008, 0x00001108, 0x00101108,
257 0x04000000, 0x04100000, 0x04000100, 0x04100100,
258 0x04000008, 0x04100008, 0x04000108, 0x04100108,
259 0x04001000, 0x04101000, 0x04001100, 0x04101100,
260 0x04001008, 0x04101008, 0x04001108, 0x04101108,
261 0x00020000, 0x00120000, 0x00020100, 0x00120100,
262 0x00020008, 0x00120008, 0x00020108, 0x00120108,
263 0x00021000, 0x00121000, 0x00021100, 0x00121100,
264 0x00021008, 0x00121008, 0x00021108, 0x00121108,
265 0x04020000, 0x04120000, 0x04020100, 0x04120100,
266 0x04020008, 0x04120008, 0x04020108, 0x04120108,
267 0x04021000, 0x04121000, 0x04021100, 0x04121100,
268 0x04021008, 0x04121008, 0x04021108, 0x04121108,
269 /* for D bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
270 0x00000000, 0x10000000, 0x00010000, 0x10010000,
271 0x00000004, 0x10000004, 0x00010004, 0x10010004,
272 0x20000000, 0x30000000, 0x20010000, 0x30010000,
273 0x20000004, 0x30000004, 0x20010004, 0x30010004,
274 0x00100000, 0x10100000, 0x00110000, 0x10110000,
275 0x00100004, 0x10100004, 0x00110004, 0x10110004,
276 0x20100000, 0x30100000, 0x20110000, 0x30110000,
277 0x20100004, 0x30100004, 0x20110004, 0x30110004,
278 0x00001000, 0x10001000, 0x00011000, 0x10011000,
279 0x00001004, 0x10001004, 0x00011004, 0x10011004,
280 0x20001000, 0x30001000, 0x20011000, 0x30011000,
281 0x20001004, 0x30001004, 0x20011004, 0x30011004,
282 0x00101000, 0x10101000, 0x00111000, 0x10111000,
283 0x00101004, 0x10101004, 0x00111004, 0x10111004,
284 0x20101000, 0x30101000, 0x20111000, 0x30111000,
285 0x20101004, 0x30101004, 0x20111004, 0x30111004,
286 /* for D bits (numbered as per FIPS 46) 8 9 11 12 13 14 */
287 0x00000000, 0x08000000, 0x00000008, 0x08000008,
288 0x00000400, 0x08000400, 0x00000408, 0x08000408,
289 0x00020000, 0x08020000, 0x00020008, 0x08020008,
290 0x00020400, 0x08020400, 0x00020408, 0x08020408,
291 0x00000001, 0x08000001, 0x00000009, 0x08000009,
292 0x00000401, 0x08000401, 0x00000409, 0x08000409,
293 0x00020001, 0x08020001, 0x00020009, 0x08020009,
294 0x00020401, 0x08020401, 0x00020409, 0x08020409,
295 0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
296 0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
297 0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
298 0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
299 0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
300 0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
301 0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
302 0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
303 /* for D bits (numbered as per FIPS 46) 16 17 18 19 20 21 */
304 0x00000000, 0x00000100, 0x00080000, 0x00080100,
305 0x01000000, 0x01000100, 0x01080000, 0x01080100,
306 0x00000010, 0x00000110, 0x00080010, 0x00080110,
307 0x01000010, 0x01000110, 0x01080010, 0x01080110,
308 0x00200000, 0x00200100, 0x00280000, 0x00280100,
309 0x01200000, 0x01200100, 0x01280000, 0x01280100,
310 0x00200010, 0x00200110, 0x00280010, 0x00280110,
311 0x01200010, 0x01200110, 0x01280010, 0x01280110,
312 0x00000200, 0x00000300, 0x00080200, 0x00080300,
313 0x01000200, 0x01000300, 0x01080200, 0x01080300,
314 0x00000210, 0x00000310, 0x00080210, 0x00080310,
315 0x01000210, 0x01000310, 0x01080210, 0x01080310,
316 0x00200200, 0x00200300, 0x00280200, 0x00280300,
317 0x01200200, 0x01200300, 0x01280200, 0x01280300,
318 0x00200210, 0x00200310, 0x00280210, 0x00280310,
319 0x01200210, 0x01200310, 0x01280210, 0x01280310,
320 /* for D bits (numbered as per FIPS 46) 22 23 24 25 27 28 */
321 0x00000000, 0x04000000, 0x00040000, 0x04040000,
322 0x00000002, 0x04000002, 0x00040002, 0x04040002,
323 0x00002000, 0x04002000, 0x00042000, 0x04042000,
324 0x00002002, 0x04002002, 0x00042002, 0x04042002,
325 0x00000020, 0x04000020, 0x00040020, 0x04040020,
326 0x00000022, 0x04000022, 0x00040022, 0x04040022,
327 0x00002020, 0x04002020, 0x00042020, 0x04042020,
328 0x00002022, 0x04002022, 0x00042022, 0x04042022,
329 0x00000800, 0x04000800, 0x00040800, 0x04040800,
330 0x00000802, 0x04000802, 0x00040802, 0x04040802,
331 0x00002800, 0x04002800, 0x00042800, 0x04042800,
332 0x00002802, 0x04002802, 0x00042802, 0x04042802,
333 0x00000820, 0x04000820, 0x00040820, 0x04040820,
334 0x00000822, 0x04000822, 0x00040822, 0x04040822,
335 0x00002820, 0x04002820, 0x00042820, 0x04042820,
336 0x00002822, 0x04002822, 0x00042822, 0x04042822
340 #define BOX(i,n,S) u32x ((S)[(n)][(i)])
344 #define BOX(i,n,S) u32x ((S)[(n)][(i).x], (S)[(n)][(i).y])
347 __device__ static void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], u32 s_SPtrans[8][64])
355 for (u32 i = 0; i < 16; i++)
358 u32x t = Kd[i] ^ rotl32 (r, 28u);
360 l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
361 | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
362 | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
363 | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
364 | BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
365 | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
366 | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
367 | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
378 __device__ static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], u32 s_skb[8][64])
382 PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
383 HPERM_OP (c, tt, 2, 0xcccc0000);
384 HPERM_OP (d, tt, 2, 0xcccc0000);
385 PERM_OP (d, c, tt, 1, 0x55555555);
386 PERM_OP (c, d, tt, 8, 0x00ff00ff);
387 PERM_OP (d, c, tt, 1, 0x55555555);
389 d = ((d & 0x000000ff) << 16)
390 | ((d & 0x0000ff00) << 0)
391 | ((d & 0x00ff0000) >> 16)
392 | ((c & 0xf0000000) >> 4);
397 for (u32 i = 0; i < 16; i++)
399 const u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
400 const u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
402 c = c >> shifts3s0[i] | c << shifts3s1[i];
403 d = d >> shifts3s0[i] | d << shifts3s1[i];
408 u32x s = BOX ((( c >> 0) & 0x3f), 0, s_skb)
409 | BOX ((((c >> 6) & 0x03)
410 | ((c >> 7) & 0x3c)), 1, s_skb)
411 | BOX ((((c >> 13) & 0x0f)
412 | ((c >> 14) & 0x30)), 2, s_skb)
413 | BOX ((((c >> 20) & 0x01)
415 | ((c >> 22) & 0x38)), 3, s_skb);
417 u32x t = BOX ((( d >> 0) & 0x3f), 4, s_skb)
418 | BOX ((((d >> 7) & 0x03)
419 | ((d >> 8) & 0x3c)), 5, s_skb)
420 | BOX ((((d >> 15) & 0x3f)), 6, s_skb)
421 | BOX ((((d >> 21) & 0x0f)
422 | ((d >> 22) & 0x30)), 7, s_skb);
424 #if __CUDA_ARCH__ >= 200
425 Kc[i] = __byte_perm (s, t, 0x5410);
426 Kd[i] = __byte_perm (s, t, 0x7632);
428 Kc[i] = ((t << 16) | (s & 0x0000ffff));
429 Kd[i] = ((s >> 16) | (t & 0xffff0000));
432 Kc[i] = rotl32 (Kc[i], 2u);
433 Kd[i] = rotl32 (Kd[i], 2u);
437 __device__ static void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
441 t[0] = (w0 >> 0) & 0xff;
442 t[1] = (w0 >> 8) & 0xff;
443 t[2] = (w0 >> 16) & 0xff;
444 t[3] = (w0 >> 24) & 0xff;
445 t[4] = (w1 >> 0) & 0xff;
446 t[5] = (w1 >> 8) & 0xff;
447 t[6] = (w1 >> 16) & 0xff;
448 t[7] = (w1 >> 24) & 0xff;
453 k[1] = (t[0] << 7) | (t[1] >> 1);
454 k[2] = (t[1] << 6) | (t[2] >> 2);
455 k[3] = (t[2] << 5) | (t[3] >> 3);
456 k[4] = (t[3] << 4) | (t[4] >> 4);
457 k[5] = (t[4] << 3) | (t[5] >> 5);
458 k[6] = (t[5] << 2) | (t[6] >> 6);
461 out[0] = ((k[0] & 0xff) << 0)
462 | ((k[1] & 0xff) << 8)
463 | ((k[2] & 0xff) << 16)
464 | ((k[3] & 0xff) << 24);
466 out[1] = ((k[4] & 0xff) << 0)
467 | ((k[5] & 0xff) << 8)
468 | ((k[6] & 0xff) << 16)
469 | ((k[7] & 0xff) << 24);
472 __device__ __shared__ u32 s_skb[8][64];
473 __device__ __shared__ u32 s_SPtrans[8][64];
475 __device__ __constant__ u32x c_bfs[1024];
477 __device__ static void m05500m (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
483 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
484 const u32 lid = threadIdx.x;
490 const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
491 const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
492 const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
503 const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
507 for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
509 const u32x w0r = c_bfs[il_pos];
511 const u32x w0 = w0l | w0r;
535 MD4_STEP (MD4_Fo, a, b, c, d, w0_t, MD4C00, MD4S00);
536 MD4_STEP (MD4_Fo, d, a, b, c, w1_t, MD4C00, MD4S01);
537 MD4_STEP (MD4_Fo, c, d, a, b, w2_t, MD4C00, MD4S02);
538 MD4_STEP (MD4_Fo, b, c, d, a, w3_t, MD4C00, MD4S03);
539 MD4_STEP (MD4_Fo, a, b, c, d, w4_t, MD4C00, MD4S00);
540 MD4_STEP (MD4_Fo, d, a, b, c, w5_t, MD4C00, MD4S01);
541 MD4_STEP (MD4_Fo, c, d, a, b, w6_t, MD4C00, MD4S02);
542 MD4_STEP (MD4_Fo, b, c, d, a, w7_t, MD4C00, MD4S03);
543 MD4_STEP (MD4_Fo, a, b, c, d, w8_t, MD4C00, MD4S00);
544 MD4_STEP (MD4_Fo, d, a, b, c, w9_t, MD4C00, MD4S01);
545 MD4_STEP (MD4_Fo, c, d, a, b, wa_t, MD4C00, MD4S02);
546 MD4_STEP (MD4_Fo, b, c, d, a, wb_t, MD4C00, MD4S03);
547 MD4_STEP (MD4_Fo, a, b, c, d, wc_t, MD4C00, MD4S00);
548 MD4_STEP (MD4_Fo, d, a, b, c, wd_t, MD4C00, MD4S01);
549 MD4_STEP (MD4_Fo, c, d, a, b, we_t, MD4C00, MD4S02);
550 MD4_STEP (MD4_Fo, b, c, d, a, wf_t, MD4C00, MD4S03);
552 MD4_STEP (MD4_Go, a, b, c, d, w0_t, MD4C01, MD4S10);
553 MD4_STEP (MD4_Go, d, a, b, c, w4_t, MD4C01, MD4S11);
554 MD4_STEP (MD4_Go, c, d, a, b, w8_t, MD4C01, MD4S12);
555 MD4_STEP (MD4_Go, b, c, d, a, wc_t, MD4C01, MD4S13);
556 MD4_STEP (MD4_Go, a, b, c, d, w1_t, MD4C01, MD4S10);
557 MD4_STEP (MD4_Go, d, a, b, c, w5_t, MD4C01, MD4S11);
558 MD4_STEP (MD4_Go, c, d, a, b, w9_t, MD4C01, MD4S12);
559 MD4_STEP (MD4_Go, b, c, d, a, wd_t, MD4C01, MD4S13);
560 MD4_STEP (MD4_Go, a, b, c, d, w2_t, MD4C01, MD4S10);
561 MD4_STEP (MD4_Go, d, a, b, c, w6_t, MD4C01, MD4S11);
562 MD4_STEP (MD4_Go, c, d, a, b, wa_t, MD4C01, MD4S12);
563 MD4_STEP (MD4_Go, b, c, d, a, we_t, MD4C01, MD4S13);
564 MD4_STEP (MD4_Go, a, b, c, d, w3_t, MD4C01, MD4S10);
565 MD4_STEP (MD4_Go, d, a, b, c, w7_t, MD4C01, MD4S11);
566 MD4_STEP (MD4_Go, c, d, a, b, wb_t, MD4C01, MD4S12);
567 MD4_STEP (MD4_Go, b, c, d, a, wf_t, MD4C01, MD4S13);
569 MD4_STEP (MD4_H , a, b, c, d, w0_t, MD4C02, MD4S20);
570 MD4_STEP (MD4_H , d, a, b, c, w8_t, MD4C02, MD4S21);
571 MD4_STEP (MD4_H , c, d, a, b, w4_t, MD4C02, MD4S22);
572 MD4_STEP (MD4_H , b, c, d, a, wc_t, MD4C02, MD4S23);
573 MD4_STEP (MD4_H , a, b, c, d, w2_t, MD4C02, MD4S20);
574 MD4_STEP (MD4_H , d, a, b, c, wa_t, MD4C02, MD4S21);
575 MD4_STEP (MD4_H , c, d, a, b, w6_t, MD4C02, MD4S22);
576 MD4_STEP (MD4_H , b, c, d, a, we_t, MD4C02, MD4S23);
577 MD4_STEP (MD4_H , a, b, c, d, w1_t, MD4C02, MD4S20);
578 MD4_STEP (MD4_H , d, a, b, c, w9_t, MD4C02, MD4S21);
579 MD4_STEP (MD4_H , c, d, a, b, w5_t, MD4C02, MD4S22);
580 MD4_STEP (MD4_H , b, c, d, a, wd_t, MD4C02, MD4S23);
581 MD4_STEP (MD4_H , a, b, c, d, w3_t, MD4C02, MD4S20);
582 MD4_STEP (MD4_H , d, a, b, c, wb_t, MD4C02, MD4S21);
584 if (s2 != ((d + MD4M_D) >> 16)) continue;
586 MD4_STEP (MD4_H , c, d, a, b, w7_t, MD4C02, MD4S22);
587 MD4_STEP (MD4_H , b, c, d, a, wf_t, MD4C02, MD4S23);
600 transform_netntlmv1_key (a, b, key);
605 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
609 _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
615 transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
617 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
621 _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
627 const u32x r0 = iv1[0];
628 const u32x r1 = iv1[1];
629 const u32x r2 = iv2[0];
630 const u32x r3 = iv2[1];
632 #include VECT_COMPARE_M
636 __device__ static void m05500s (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
642 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
643 const u32 lid = threadIdx.x;
649 const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
650 const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
651 const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
662 const u32 search[4] =
664 digests_buf[digests_offset].digest_buf[DGST_R0],
665 digests_buf[digests_offset].digest_buf[DGST_R1],
666 digests_buf[digests_offset].digest_buf[DGST_R2],
667 digests_buf[digests_offset].digest_buf[DGST_R3]
674 const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
678 for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
680 const u32x w0r = c_bfs[il_pos];
682 const u32x w0 = w0l | w0r;
706 MD4_STEP (MD4_Fo, a, b, c, d, w0_t, MD4C00, MD4S00);
707 MD4_STEP (MD4_Fo, d, a, b, c, w1_t, MD4C00, MD4S01);
708 MD4_STEP (MD4_Fo, c, d, a, b, w2_t, MD4C00, MD4S02);
709 MD4_STEP (MD4_Fo, b, c, d, a, w3_t, MD4C00, MD4S03);
710 MD4_STEP (MD4_Fo, a, b, c, d, w4_t, MD4C00, MD4S00);
711 MD4_STEP (MD4_Fo, d, a, b, c, w5_t, MD4C00, MD4S01);
712 MD4_STEP (MD4_Fo, c, d, a, b, w6_t, MD4C00, MD4S02);
713 MD4_STEP (MD4_Fo, b, c, d, a, w7_t, MD4C00, MD4S03);
714 MD4_STEP (MD4_Fo, a, b, c, d, w8_t, MD4C00, MD4S00);
715 MD4_STEP (MD4_Fo, d, a, b, c, w9_t, MD4C00, MD4S01);
716 MD4_STEP (MD4_Fo, c, d, a, b, wa_t, MD4C00, MD4S02);
717 MD4_STEP (MD4_Fo, b, c, d, a, wb_t, MD4C00, MD4S03);
718 MD4_STEP (MD4_Fo, a, b, c, d, wc_t, MD4C00, MD4S00);
719 MD4_STEP (MD4_Fo, d, a, b, c, wd_t, MD4C00, MD4S01);
720 MD4_STEP (MD4_Fo, c, d, a, b, we_t, MD4C00, MD4S02);
721 MD4_STEP (MD4_Fo, b, c, d, a, wf_t, MD4C00, MD4S03);
723 MD4_STEP (MD4_Go, a, b, c, d, w0_t, MD4C01, MD4S10);
724 MD4_STEP (MD4_Go, d, a, b, c, w4_t, MD4C01, MD4S11);
725 MD4_STEP (MD4_Go, c, d, a, b, w8_t, MD4C01, MD4S12);
726 MD4_STEP (MD4_Go, b, c, d, a, wc_t, MD4C01, MD4S13);
727 MD4_STEP (MD4_Go, a, b, c, d, w1_t, MD4C01, MD4S10);
728 MD4_STEP (MD4_Go, d, a, b, c, w5_t, MD4C01, MD4S11);
729 MD4_STEP (MD4_Go, c, d, a, b, w9_t, MD4C01, MD4S12);
730 MD4_STEP (MD4_Go, b, c, d, a, wd_t, MD4C01, MD4S13);
731 MD4_STEP (MD4_Go, a, b, c, d, w2_t, MD4C01, MD4S10);
732 MD4_STEP (MD4_Go, d, a, b, c, w6_t, MD4C01, MD4S11);
733 MD4_STEP (MD4_Go, c, d, a, b, wa_t, MD4C01, MD4S12);
734 MD4_STEP (MD4_Go, b, c, d, a, we_t, MD4C01, MD4S13);
735 MD4_STEP (MD4_Go, a, b, c, d, w3_t, MD4C01, MD4S10);
736 MD4_STEP (MD4_Go, d, a, b, c, w7_t, MD4C01, MD4S11);
737 MD4_STEP (MD4_Go, c, d, a, b, wb_t, MD4C01, MD4S12);
738 MD4_STEP (MD4_Go, b, c, d, a, wf_t, MD4C01, MD4S13);
740 MD4_STEP (MD4_H , a, b, c, d, w0_t, MD4C02, MD4S20);
741 MD4_STEP (MD4_H , d, a, b, c, w8_t, MD4C02, MD4S21);
742 MD4_STEP (MD4_H , c, d, a, b, w4_t, MD4C02, MD4S22);
743 MD4_STEP (MD4_H , b, c, d, a, wc_t, MD4C02, MD4S23);
744 MD4_STEP (MD4_H , a, b, c, d, w2_t, MD4C02, MD4S20);
745 MD4_STEP (MD4_H , d, a, b, c, wa_t, MD4C02, MD4S21);
746 MD4_STEP (MD4_H , c, d, a, b, w6_t, MD4C02, MD4S22);
747 MD4_STEP (MD4_H , b, c, d, a, we_t, MD4C02, MD4S23);
748 MD4_STEP (MD4_H , a, b, c, d, w1_t, MD4C02, MD4S20);
749 MD4_STEP (MD4_H , d, a, b, c, w9_t, MD4C02, MD4S21);
750 MD4_STEP (MD4_H , c, d, a, b, w5_t, MD4C02, MD4S22);
751 MD4_STEP (MD4_H , b, c, d, a, wd_t, MD4C02, MD4S23);
752 MD4_STEP (MD4_H , a, b, c, d, w3_t, MD4C02, MD4S20);
753 MD4_STEP (MD4_H , d, a, b, c, wb_t, MD4C02, MD4S21);
755 if (s2 != ((d + MD4M_D) >> 16)) continue;
757 MD4_STEP (MD4_H , c, d, a, b, w7_t, MD4C02, MD4S22);
758 MD4_STEP (MD4_H , b, c, d, a, wf_t, MD4C02, MD4S23);
771 transform_netntlmv1_key (a, b, key);
776 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
780 _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
786 transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
788 _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
792 _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
798 const u32x r0 = iv1[0];
799 const u32x r1 = iv1[1];
800 const u32x r2 = iv2[0];
801 const u32x r3 = iv2[1];
803 #include VECT_COMPARE_S
807 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
813 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
814 const u32 lid = threadIdx.x;
818 w[ 0] = pws[gid].i[ 0];
819 w[ 1] = pws[gid].i[ 1];
820 w[ 2] = pws[gid].i[ 2];
821 w[ 3] = pws[gid].i[ 3];
832 w[14] = pws[gid].i[14];
835 const u32 pw_len = pws[gid].pw_len;
843 s_skb[0][lid] = c_skb[0][lid];
844 s_skb[1][lid] = c_skb[1][lid];
845 s_skb[2][lid] = c_skb[2][lid];
846 s_skb[3][lid] = c_skb[3][lid];
847 s_skb[4][lid] = c_skb[4][lid];
848 s_skb[5][lid] = c_skb[5][lid];
849 s_skb[6][lid] = c_skb[6][lid];
850 s_skb[7][lid] = c_skb[7][lid];
852 s_SPtrans[0][lid] = c_SPtrans[0][lid];
853 s_SPtrans[1][lid] = c_SPtrans[1][lid];
854 s_SPtrans[2][lid] = c_SPtrans[2][lid];
855 s_SPtrans[3][lid] = c_SPtrans[3][lid];
856 s_SPtrans[4][lid] = c_SPtrans[4][lid];
857 s_SPtrans[5][lid] = c_SPtrans[5][lid];
858 s_SPtrans[6][lid] = c_SPtrans[6][lid];
859 s_SPtrans[7][lid] = c_SPtrans[7][lid];
864 if (gid >= gid_max) return;
870 m05500m (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);
873 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
879 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
880 const u32 lid = threadIdx.x;
884 w[ 0] = pws[gid].i[ 0];
885 w[ 1] = pws[gid].i[ 1];
886 w[ 2] = pws[gid].i[ 2];
887 w[ 3] = pws[gid].i[ 3];
888 w[ 4] = pws[gid].i[ 4];
889 w[ 5] = pws[gid].i[ 5];
890 w[ 6] = pws[gid].i[ 6];
891 w[ 7] = pws[gid].i[ 7];
898 w[14] = pws[gid].i[14];
901 const u32 pw_len = pws[gid].pw_len;
909 s_skb[0][lid] = c_skb[0][lid];
910 s_skb[1][lid] = c_skb[1][lid];
911 s_skb[2][lid] = c_skb[2][lid];
912 s_skb[3][lid] = c_skb[3][lid];
913 s_skb[4][lid] = c_skb[4][lid];
914 s_skb[5][lid] = c_skb[5][lid];
915 s_skb[6][lid] = c_skb[6][lid];
916 s_skb[7][lid] = c_skb[7][lid];
918 s_SPtrans[0][lid] = c_SPtrans[0][lid];
919 s_SPtrans[1][lid] = c_SPtrans[1][lid];
920 s_SPtrans[2][lid] = c_SPtrans[2][lid];
921 s_SPtrans[3][lid] = c_SPtrans[3][lid];
922 s_SPtrans[4][lid] = c_SPtrans[4][lid];
923 s_SPtrans[5][lid] = c_SPtrans[5][lid];
924 s_SPtrans[6][lid] = c_SPtrans[6][lid];
925 s_SPtrans[7][lid] = c_SPtrans[7][lid];
930 if (gid >= gid_max) return;
936 m05500m (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);
939 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
943 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
949 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
950 const u32 lid = threadIdx.x;
954 w[ 0] = pws[gid].i[ 0];
955 w[ 1] = pws[gid].i[ 1];
956 w[ 2] = pws[gid].i[ 2];
957 w[ 3] = pws[gid].i[ 3];
968 w[14] = pws[gid].i[14];
971 const u32 pw_len = pws[gid].pw_len;
979 s_skb[0][lid] = c_skb[0][lid];
980 s_skb[1][lid] = c_skb[1][lid];
981 s_skb[2][lid] = c_skb[2][lid];
982 s_skb[3][lid] = c_skb[3][lid];
983 s_skb[4][lid] = c_skb[4][lid];
984 s_skb[5][lid] = c_skb[5][lid];
985 s_skb[6][lid] = c_skb[6][lid];
986 s_skb[7][lid] = c_skb[7][lid];
988 s_SPtrans[0][lid] = c_SPtrans[0][lid];
989 s_SPtrans[1][lid] = c_SPtrans[1][lid];
990 s_SPtrans[2][lid] = c_SPtrans[2][lid];
991 s_SPtrans[3][lid] = c_SPtrans[3][lid];
992 s_SPtrans[4][lid] = c_SPtrans[4][lid];
993 s_SPtrans[5][lid] = c_SPtrans[5][lid];
994 s_SPtrans[6][lid] = c_SPtrans[6][lid];
995 s_SPtrans[7][lid] = c_SPtrans[7][lid];
1000 if (gid >= gid_max) return;
1006 m05500s (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);
1009 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)
1015 const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1016 const u32 lid = threadIdx.x;
1020 w[ 0] = pws[gid].i[ 0];
1021 w[ 1] = pws[gid].i[ 1];
1022 w[ 2] = pws[gid].i[ 2];
1023 w[ 3] = pws[gid].i[ 3];
1024 w[ 4] = pws[gid].i[ 4];
1025 w[ 5] = pws[gid].i[ 5];
1026 w[ 6] = pws[gid].i[ 6];
1027 w[ 7] = pws[gid].i[ 7];
1034 w[14] = pws[gid].i[14];
1037 const u32 pw_len = pws[gid].pw_len;
1045 s_skb[0][lid] = c_skb[0][lid];
1046 s_skb[1][lid] = c_skb[1][lid];
1047 s_skb[2][lid] = c_skb[2][lid];
1048 s_skb[3][lid] = c_skb[3][lid];
1049 s_skb[4][lid] = c_skb[4][lid];
1050 s_skb[5][lid] = c_skb[5][lid];
1051 s_skb[6][lid] = c_skb[6][lid];
1052 s_skb[7][lid] = c_skb[7][lid];
1054 s_SPtrans[0][lid] = c_SPtrans[0][lid];
1055 s_SPtrans[1][lid] = c_SPtrans[1][lid];
1056 s_SPtrans[2][lid] = c_SPtrans[2][lid];
1057 s_SPtrans[3][lid] = c_SPtrans[3][lid];
1058 s_SPtrans[4][lid] = c_SPtrans[4][lid];
1059 s_SPtrans[5][lid] = c_SPtrans[5][lid];
1060 s_SPtrans[6][lid] = c_SPtrans[6][lid];
1061 s_SPtrans[7][lid] = c_SPtrans[7][lid];
1066 if (gid >= gid_max) return;
1072 m05500s (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);
1075 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, 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)