2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
12 #include "include/constants.h"
13 #include "include/kernel_vendor.h"
20 #include "include/kernel_functions.c"
21 #include "OpenCL/types_ocl.c"
22 #include "OpenCL/common.c"
23 #include "OpenCL/simd.c"
25 __constant u32 c_tables[4][256] =
28 0x00072000, 0x00075000, 0x00074800, 0x00071000,
29 0x00076800, 0x00074000, 0x00070000, 0x00077000,
30 0x00073000, 0x00075800, 0x00070800, 0x00076000,
31 0x00073800, 0x00077800, 0x00072800, 0x00071800,
32 0x0005a000, 0x0005d000, 0x0005c800, 0x00059000,
33 0x0005e800, 0x0005c000, 0x00058000, 0x0005f000,
34 0x0005b000, 0x0005d800, 0x00058800, 0x0005e000,
35 0x0005b800, 0x0005f800, 0x0005a800, 0x00059800,
36 0x00022000, 0x00025000, 0x00024800, 0x00021000,
37 0x00026800, 0x00024000, 0x00020000, 0x00027000,
38 0x00023000, 0x00025800, 0x00020800, 0x00026000,
39 0x00023800, 0x00027800, 0x00022800, 0x00021800,
40 0x00062000, 0x00065000, 0x00064800, 0x00061000,
41 0x00066800, 0x00064000, 0x00060000, 0x00067000,
42 0x00063000, 0x00065800, 0x00060800, 0x00066000,
43 0x00063800, 0x00067800, 0x00062800, 0x00061800,
44 0x00032000, 0x00035000, 0x00034800, 0x00031000,
45 0x00036800, 0x00034000, 0x00030000, 0x00037000,
46 0x00033000, 0x00035800, 0x00030800, 0x00036000,
47 0x00033800, 0x00037800, 0x00032800, 0x00031800,
48 0x0006a000, 0x0006d000, 0x0006c800, 0x00069000,
49 0x0006e800, 0x0006c000, 0x00068000, 0x0006f000,
50 0x0006b000, 0x0006d800, 0x00068800, 0x0006e000,
51 0x0006b800, 0x0006f800, 0x0006a800, 0x00069800,
52 0x0007a000, 0x0007d000, 0x0007c800, 0x00079000,
53 0x0007e800, 0x0007c000, 0x00078000, 0x0007f000,
54 0x0007b000, 0x0007d800, 0x00078800, 0x0007e000,
55 0x0007b800, 0x0007f800, 0x0007a800, 0x00079800,
56 0x00052000, 0x00055000, 0x00054800, 0x00051000,
57 0x00056800, 0x00054000, 0x00050000, 0x00057000,
58 0x00053000, 0x00055800, 0x00050800, 0x00056000,
59 0x00053800, 0x00057800, 0x00052800, 0x00051800,
60 0x00012000, 0x00015000, 0x00014800, 0x00011000,
61 0x00016800, 0x00014000, 0x00010000, 0x00017000,
62 0x00013000, 0x00015800, 0x00010800, 0x00016000,
63 0x00013800, 0x00017800, 0x00012800, 0x00011800,
64 0x0001a000, 0x0001d000, 0x0001c800, 0x00019000,
65 0x0001e800, 0x0001c000, 0x00018000, 0x0001f000,
66 0x0001b000, 0x0001d800, 0x00018800, 0x0001e000,
67 0x0001b800, 0x0001f800, 0x0001a800, 0x00019800,
68 0x00042000, 0x00045000, 0x00044800, 0x00041000,
69 0x00046800, 0x00044000, 0x00040000, 0x00047000,
70 0x00043000, 0x00045800, 0x00040800, 0x00046000,
71 0x00043800, 0x00047800, 0x00042800, 0x00041800,
72 0x0000a000, 0x0000d000, 0x0000c800, 0x00009000,
73 0x0000e800, 0x0000c000, 0x00008000, 0x0000f000,
74 0x0000b000, 0x0000d800, 0x00008800, 0x0000e000,
75 0x0000b800, 0x0000f800, 0x0000a800, 0x00009800,
76 0x00002000, 0x00005000, 0x00004800, 0x00001000,
77 0x00006800, 0x00004000, 0x00000000, 0x00007000,
78 0x00003000, 0x00005800, 0x00000800, 0x00006000,
79 0x00003800, 0x00007800, 0x00002800, 0x00001800,
80 0x0003a000, 0x0003d000, 0x0003c800, 0x00039000,
81 0x0003e800, 0x0003c000, 0x00038000, 0x0003f000,
82 0x0003b000, 0x0003d800, 0x00038800, 0x0003e000,
83 0x0003b800, 0x0003f800, 0x0003a800, 0x00039800,
84 0x0002a000, 0x0002d000, 0x0002c800, 0x00029000,
85 0x0002e800, 0x0002c000, 0x00028000, 0x0002f000,
86 0x0002b000, 0x0002d800, 0x00028800, 0x0002e000,
87 0x0002b800, 0x0002f800, 0x0002a800, 0x00029800,
88 0x0004a000, 0x0004d000, 0x0004c800, 0x00049000,
89 0x0004e800, 0x0004c000, 0x00048000, 0x0004f000,
90 0x0004b000, 0x0004d800, 0x00048800, 0x0004e000,
91 0x0004b800, 0x0004f800, 0x0004a800, 0x00049800,
94 0x03a80000, 0x03c00000, 0x03880000, 0x03e80000,
95 0x03d00000, 0x03980000, 0x03a00000, 0x03900000,
96 0x03f00000, 0x03f80000, 0x03e00000, 0x03b80000,
97 0x03b00000, 0x03800000, 0x03c80000, 0x03d80000,
98 0x06a80000, 0x06c00000, 0x06880000, 0x06e80000,
99 0x06d00000, 0x06980000, 0x06a00000, 0x06900000,
100 0x06f00000, 0x06f80000, 0x06e00000, 0x06b80000,
101 0x06b00000, 0x06800000, 0x06c80000, 0x06d80000,
102 0x05280000, 0x05400000, 0x05080000, 0x05680000,
103 0x05500000, 0x05180000, 0x05200000, 0x05100000,
104 0x05700000, 0x05780000, 0x05600000, 0x05380000,
105 0x05300000, 0x05000000, 0x05480000, 0x05580000,
106 0x00a80000, 0x00c00000, 0x00880000, 0x00e80000,
107 0x00d00000, 0x00980000, 0x00a00000, 0x00900000,
108 0x00f00000, 0x00f80000, 0x00e00000, 0x00b80000,
109 0x00b00000, 0x00800000, 0x00c80000, 0x00d80000,
110 0x00280000, 0x00400000, 0x00080000, 0x00680000,
111 0x00500000, 0x00180000, 0x00200000, 0x00100000,
112 0x00700000, 0x00780000, 0x00600000, 0x00380000,
113 0x00300000, 0x00000000, 0x00480000, 0x00580000,
114 0x04280000, 0x04400000, 0x04080000, 0x04680000,
115 0x04500000, 0x04180000, 0x04200000, 0x04100000,
116 0x04700000, 0x04780000, 0x04600000, 0x04380000,
117 0x04300000, 0x04000000, 0x04480000, 0x04580000,
118 0x04a80000, 0x04c00000, 0x04880000, 0x04e80000,
119 0x04d00000, 0x04980000, 0x04a00000, 0x04900000,
120 0x04f00000, 0x04f80000, 0x04e00000, 0x04b80000,
121 0x04b00000, 0x04800000, 0x04c80000, 0x04d80000,
122 0x07a80000, 0x07c00000, 0x07880000, 0x07e80000,
123 0x07d00000, 0x07980000, 0x07a00000, 0x07900000,
124 0x07f00000, 0x07f80000, 0x07e00000, 0x07b80000,
125 0x07b00000, 0x07800000, 0x07c80000, 0x07d80000,
126 0x07280000, 0x07400000, 0x07080000, 0x07680000,
127 0x07500000, 0x07180000, 0x07200000, 0x07100000,
128 0x07700000, 0x07780000, 0x07600000, 0x07380000,
129 0x07300000, 0x07000000, 0x07480000, 0x07580000,
130 0x02280000, 0x02400000, 0x02080000, 0x02680000,
131 0x02500000, 0x02180000, 0x02200000, 0x02100000,
132 0x02700000, 0x02780000, 0x02600000, 0x02380000,
133 0x02300000, 0x02000000, 0x02480000, 0x02580000,
134 0x03280000, 0x03400000, 0x03080000, 0x03680000,
135 0x03500000, 0x03180000, 0x03200000, 0x03100000,
136 0x03700000, 0x03780000, 0x03600000, 0x03380000,
137 0x03300000, 0x03000000, 0x03480000, 0x03580000,
138 0x06280000, 0x06400000, 0x06080000, 0x06680000,
139 0x06500000, 0x06180000, 0x06200000, 0x06100000,
140 0x06700000, 0x06780000, 0x06600000, 0x06380000,
141 0x06300000, 0x06000000, 0x06480000, 0x06580000,
142 0x05a80000, 0x05c00000, 0x05880000, 0x05e80000,
143 0x05d00000, 0x05980000, 0x05a00000, 0x05900000,
144 0x05f00000, 0x05f80000, 0x05e00000, 0x05b80000,
145 0x05b00000, 0x05800000, 0x05c80000, 0x05d80000,
146 0x01280000, 0x01400000, 0x01080000, 0x01680000,
147 0x01500000, 0x01180000, 0x01200000, 0x01100000,
148 0x01700000, 0x01780000, 0x01600000, 0x01380000,
149 0x01300000, 0x01000000, 0x01480000, 0x01580000,
150 0x02a80000, 0x02c00000, 0x02880000, 0x02e80000,
151 0x02d00000, 0x02980000, 0x02a00000, 0x02900000,
152 0x02f00000, 0x02f80000, 0x02e00000, 0x02b80000,
153 0x02b00000, 0x02800000, 0x02c80000, 0x02d80000,
154 0x01a80000, 0x01c00000, 0x01880000, 0x01e80000,
155 0x01d00000, 0x01980000, 0x01a00000, 0x01900000,
156 0x01f00000, 0x01f80000, 0x01e00000, 0x01b80000,
157 0x01b00000, 0x01800000, 0x01c80000, 0x01d80000,
160 0x30000002, 0x60000002, 0x38000002, 0x08000002,
161 0x28000002, 0x78000002, 0x68000002, 0x40000002,
162 0x20000002, 0x50000002, 0x48000002, 0x70000002,
163 0x00000002, 0x18000002, 0x58000002, 0x10000002,
164 0xb0000005, 0xe0000005, 0xb8000005, 0x88000005,
165 0xa8000005, 0xf8000005, 0xe8000005, 0xc0000005,
166 0xa0000005, 0xd0000005, 0xc8000005, 0xf0000005,
167 0x80000005, 0x98000005, 0xd8000005, 0x90000005,
168 0x30000005, 0x60000005, 0x38000005, 0x08000005,
169 0x28000005, 0x78000005, 0x68000005, 0x40000005,
170 0x20000005, 0x50000005, 0x48000005, 0x70000005,
171 0x00000005, 0x18000005, 0x58000005, 0x10000005,
172 0x30000000, 0x60000000, 0x38000000, 0x08000000,
173 0x28000000, 0x78000000, 0x68000000, 0x40000000,
174 0x20000000, 0x50000000, 0x48000000, 0x70000000,
175 0x00000000, 0x18000000, 0x58000000, 0x10000000,
176 0xb0000003, 0xe0000003, 0xb8000003, 0x88000003,
177 0xa8000003, 0xf8000003, 0xe8000003, 0xc0000003,
178 0xa0000003, 0xd0000003, 0xc8000003, 0xf0000003,
179 0x80000003, 0x98000003, 0xd8000003, 0x90000003,
180 0x30000001, 0x60000001, 0x38000001, 0x08000001,
181 0x28000001, 0x78000001, 0x68000001, 0x40000001,
182 0x20000001, 0x50000001, 0x48000001, 0x70000001,
183 0x00000001, 0x18000001, 0x58000001, 0x10000001,
184 0xb0000000, 0xe0000000, 0xb8000000, 0x88000000,
185 0xa8000000, 0xf8000000, 0xe8000000, 0xc0000000,
186 0xa0000000, 0xd0000000, 0xc8000000, 0xf0000000,
187 0x80000000, 0x98000000, 0xd8000000, 0x90000000,
188 0xb0000006, 0xe0000006, 0xb8000006, 0x88000006,
189 0xa8000006, 0xf8000006, 0xe8000006, 0xc0000006,
190 0xa0000006, 0xd0000006, 0xc8000006, 0xf0000006,
191 0x80000006, 0x98000006, 0xd8000006, 0x90000006,
192 0xb0000001, 0xe0000001, 0xb8000001, 0x88000001,
193 0xa8000001, 0xf8000001, 0xe8000001, 0xc0000001,
194 0xa0000001, 0xd0000001, 0xc8000001, 0xf0000001,
195 0x80000001, 0x98000001, 0xd8000001, 0x90000001,
196 0x30000003, 0x60000003, 0x38000003, 0x08000003,
197 0x28000003, 0x78000003, 0x68000003, 0x40000003,
198 0x20000003, 0x50000003, 0x48000003, 0x70000003,
199 0x00000003, 0x18000003, 0x58000003, 0x10000003,
200 0x30000004, 0x60000004, 0x38000004, 0x08000004,
201 0x28000004, 0x78000004, 0x68000004, 0x40000004,
202 0x20000004, 0x50000004, 0x48000004, 0x70000004,
203 0x00000004, 0x18000004, 0x58000004, 0x10000004,
204 0xb0000002, 0xe0000002, 0xb8000002, 0x88000002,
205 0xa8000002, 0xf8000002, 0xe8000002, 0xc0000002,
206 0xa0000002, 0xd0000002, 0xc8000002, 0xf0000002,
207 0x80000002, 0x98000002, 0xd8000002, 0x90000002,
208 0xb0000004, 0xe0000004, 0xb8000004, 0x88000004,
209 0xa8000004, 0xf8000004, 0xe8000004, 0xc0000004,
210 0xa0000004, 0xd0000004, 0xc8000004, 0xf0000004,
211 0x80000004, 0x98000004, 0xd8000004, 0x90000004,
212 0x30000006, 0x60000006, 0x38000006, 0x08000006,
213 0x28000006, 0x78000006, 0x68000006, 0x40000006,
214 0x20000006, 0x50000006, 0x48000006, 0x70000006,
215 0x00000006, 0x18000006, 0x58000006, 0x10000006,
216 0xb0000007, 0xe0000007, 0xb8000007, 0x88000007,
217 0xa8000007, 0xf8000007, 0xe8000007, 0xc0000007,
218 0xa0000007, 0xd0000007, 0xc8000007, 0xf0000007,
219 0x80000007, 0x98000007, 0xd8000007, 0x90000007,
220 0x30000007, 0x60000007, 0x38000007, 0x08000007,
221 0x28000007, 0x78000007, 0x68000007, 0x40000007,
222 0x20000007, 0x50000007, 0x48000007, 0x70000007,
223 0x00000007, 0x18000007, 0x58000007, 0x10000007,
226 0x000000e8, 0x000000d8, 0x000000a0, 0x00000088,
227 0x00000098, 0x000000f8, 0x000000a8, 0x000000c8,
228 0x00000080, 0x000000d0, 0x000000f0, 0x000000b8,
229 0x000000b0, 0x000000c0, 0x00000090, 0x000000e0,
230 0x000007e8, 0x000007d8, 0x000007a0, 0x00000788,
231 0x00000798, 0x000007f8, 0x000007a8, 0x000007c8,
232 0x00000780, 0x000007d0, 0x000007f0, 0x000007b8,
233 0x000007b0, 0x000007c0, 0x00000790, 0x000007e0,
234 0x000006e8, 0x000006d8, 0x000006a0, 0x00000688,
235 0x00000698, 0x000006f8, 0x000006a8, 0x000006c8,
236 0x00000680, 0x000006d0, 0x000006f0, 0x000006b8,
237 0x000006b0, 0x000006c0, 0x00000690, 0x000006e0,
238 0x00000068, 0x00000058, 0x00000020, 0x00000008,
239 0x00000018, 0x00000078, 0x00000028, 0x00000048,
240 0x00000000, 0x00000050, 0x00000070, 0x00000038,
241 0x00000030, 0x00000040, 0x00000010, 0x00000060,
242 0x000002e8, 0x000002d8, 0x000002a0, 0x00000288,
243 0x00000298, 0x000002f8, 0x000002a8, 0x000002c8,
244 0x00000280, 0x000002d0, 0x000002f0, 0x000002b8,
245 0x000002b0, 0x000002c0, 0x00000290, 0x000002e0,
246 0x000003e8, 0x000003d8, 0x000003a0, 0x00000388,
247 0x00000398, 0x000003f8, 0x000003a8, 0x000003c8,
248 0x00000380, 0x000003d0, 0x000003f0, 0x000003b8,
249 0x000003b0, 0x000003c0, 0x00000390, 0x000003e0,
250 0x00000568, 0x00000558, 0x00000520, 0x00000508,
251 0x00000518, 0x00000578, 0x00000528, 0x00000548,
252 0x00000500, 0x00000550, 0x00000570, 0x00000538,
253 0x00000530, 0x00000540, 0x00000510, 0x00000560,
254 0x00000268, 0x00000258, 0x00000220, 0x00000208,
255 0x00000218, 0x00000278, 0x00000228, 0x00000248,
256 0x00000200, 0x00000250, 0x00000270, 0x00000238,
257 0x00000230, 0x00000240, 0x00000210, 0x00000260,
258 0x000004e8, 0x000004d8, 0x000004a0, 0x00000488,
259 0x00000498, 0x000004f8, 0x000004a8, 0x000004c8,
260 0x00000480, 0x000004d0, 0x000004f0, 0x000004b8,
261 0x000004b0, 0x000004c0, 0x00000490, 0x000004e0,
262 0x00000168, 0x00000158, 0x00000120, 0x00000108,
263 0x00000118, 0x00000178, 0x00000128, 0x00000148,
264 0x00000100, 0x00000150, 0x00000170, 0x00000138,
265 0x00000130, 0x00000140, 0x00000110, 0x00000160,
266 0x000001e8, 0x000001d8, 0x000001a0, 0x00000188,
267 0x00000198, 0x000001f8, 0x000001a8, 0x000001c8,
268 0x00000180, 0x000001d0, 0x000001f0, 0x000001b8,
269 0x000001b0, 0x000001c0, 0x00000190, 0x000001e0,
270 0x00000768, 0x00000758, 0x00000720, 0x00000708,
271 0x00000718, 0x00000778, 0x00000728, 0x00000748,
272 0x00000700, 0x00000750, 0x00000770, 0x00000738,
273 0x00000730, 0x00000740, 0x00000710, 0x00000760,
274 0x00000368, 0x00000358, 0x00000320, 0x00000308,
275 0x00000318, 0x00000378, 0x00000328, 0x00000348,
276 0x00000300, 0x00000350, 0x00000370, 0x00000338,
277 0x00000330, 0x00000340, 0x00000310, 0x00000360,
278 0x000005e8, 0x000005d8, 0x000005a0, 0x00000588,
279 0x00000598, 0x000005f8, 0x000005a8, 0x000005c8,
280 0x00000580, 0x000005d0, 0x000005f0, 0x000005b8,
281 0x000005b0, 0x000005c0, 0x00000590, 0x000005e0,
282 0x00000468, 0x00000458, 0x00000420, 0x00000408,
283 0x00000418, 0x00000478, 0x00000428, 0x00000448,
284 0x00000400, 0x00000450, 0x00000470, 0x00000438,
285 0x00000430, 0x00000440, 0x00000410, 0x00000460,
286 0x00000668, 0x00000658, 0x00000620, 0x00000608,
287 0x00000618, 0x00000678, 0x00000628, 0x00000648,
288 0x00000600, 0x00000650, 0x00000670, 0x00000638,
289 0x00000630, 0x00000640, 0x00000610, 0x00000660,
294 #define BOX(i,n,S) (S)[(n)][(i)]
296 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
298 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
300 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
303 #define _round(k1,k2,tbl) \
307 l ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
308 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
309 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
310 BOX (((t >> 24) & 0xff), 3, tbl); \
312 r ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
313 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
314 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
315 BOX (((t >> 24) & 0xff), 3, tbl); \
318 #define R(k,h,s,i,t) \
324 _round (k[0], k[1], t); \
325 _round (k[2], k[3], t); \
326 _round (k[4], k[5], t); \
327 _round (k[6], k[7], t); \
328 _round (k[0], k[1], t); \
329 _round (k[2], k[3], t); \
330 _round (k[4], k[5], t); \
331 _round (k[6], k[7], t); \
332 _round (k[0], k[1], t); \
333 _round (k[2], k[3], t); \
334 _round (k[4], k[5], t); \
335 _round (k[6], k[7], t); \
336 _round (k[7], k[6], t); \
337 _round (k[5], k[4], t); \
338 _round (k[3], k[2], t); \
339 _round (k[1], k[0], t); \
345 w[0] = u[0] ^ v[0]; \
346 w[1] = u[1] ^ v[1]; \
347 w[2] = u[2] ^ v[2]; \
348 w[3] = u[3] ^ v[3]; \
349 w[4] = u[4] ^ v[4]; \
350 w[5] = u[5] ^ v[5]; \
351 w[6] = u[6] ^ v[6]; \
355 k[0] = ((w[0] & 0x000000ff) << 0) \
356 | ((w[2] & 0x000000ff) << 8) \
357 | ((w[4] & 0x000000ff) << 16) \
358 | ((w[6] & 0x000000ff) << 24); \
359 k[1] = ((w[0] & 0x0000ff00) >> 8) \
360 | ((w[2] & 0x0000ff00) >> 0) \
361 | ((w[4] & 0x0000ff00) << 8) \
362 | ((w[6] & 0x0000ff00) << 16); \
363 k[2] = ((w[0] & 0x00ff0000) >> 16) \
364 | ((w[2] & 0x00ff0000) >> 8) \
365 | ((w[4] & 0x00ff0000) << 0) \
366 | ((w[6] & 0x00ff0000) << 8); \
367 k[3] = ((w[0] & 0xff000000) >> 24) \
368 | ((w[2] & 0xff000000) >> 16) \
369 | ((w[4] & 0xff000000) >> 8) \
370 | ((w[6] & 0xff000000) >> 0); \
371 k[4] = ((w[1] & 0x000000ff) << 0) \
372 | ((w[3] & 0x000000ff) << 8) \
373 | ((w[5] & 0x000000ff) << 16) \
374 | ((w[7] & 0x000000ff) << 24); \
375 k[5] = ((w[1] & 0x0000ff00) >> 8) \
376 | ((w[3] & 0x0000ff00) >> 0) \
377 | ((w[5] & 0x0000ff00) << 8) \
378 | ((w[7] & 0x0000ff00) << 16); \
379 k[6] = ((w[1] & 0x00ff0000) >> 16) \
380 | ((w[3] & 0x00ff0000) >> 8) \
381 | ((w[5] & 0x00ff0000) << 0) \
382 | ((w[7] & 0x00ff0000) << 8); \
383 k[7] = ((w[1] & 0xff000000) >> 24) \
384 | ((w[3] & 0xff000000) >> 16) \
385 | ((w[5] & 0xff000000) >> 8) \
386 | ((w[7] & 0xff000000) >> 0);
423 x[0] ^= 0xff00ff00; \
424 x[1] ^= 0xff00ff00; \
425 x[2] ^= 0x00ff00ff; \
426 x[3] ^= 0x00ff00ff; \
427 x[4] ^= 0x00ffff00; \
428 x[5] ^= 0xff0000ff; \
429 x[6] ^= 0x000000ff; \
432 #define SHIFT12(u,m,s) \
433 u[0] = m[0] ^ s[6]; \
434 u[1] = m[1] ^ s[7]; \
435 u[2] = m[2] ^ (s[0] << 16) \
437 ^ (s[0] & 0x0000ffff) \
438 ^ (s[1] & 0x0000ffff) \
443 ^ (s[7] & 0xffff0000) \
445 u[3] = m[3] ^ (s[0] & 0x0000ffff) \
447 ^ (s[1] & 0x0000ffff) \
456 ^ (s[7] & 0x0000ffff) \
459 u[4] = m[4] ^ (s[0] & 0xffff0000) \
462 ^ (s[1] & 0xffff0000) \
471 ^ (s[7] & 0x0000ffff) \
474 u[5] = m[5] ^ (s[0] << 16) \
476 ^ (s[0] & 0xffff0000) \
477 ^ (s[1] & 0x0000ffff) \
487 ^ (s[7] & 0xffff0000) \
503 u[7] = m[7] ^ (s[0] & 0xffff0000) \
505 ^ (s[1] & 0x0000ffff) \
514 ^ (s[7] & 0x0000ffff) \
518 #define SHIFT16(h,v,u) \
519 v[0] = h[0] ^ (u[1] << 16) \
521 v[1] = h[1] ^ (u[2] << 16) \
523 v[2] = h[2] ^ (u[3] << 16) \
525 v[3] = h[3] ^ (u[4] << 16) \
527 v[4] = h[4] ^ (u[5] << 16) \
529 v[5] = h[5] ^ (u[6] << 16) \
531 v[6] = h[6] ^ (u[7] << 16) \
533 v[7] = h[7] ^ (u[0] & 0xffff0000) \
536 ^ (u[1] & 0xffff0000) \
539 ^ (u[7] & 0xffff0000);
541 #define SHIFT61(h,v) \
542 h[0] = (v[0] & 0xffff0000) \
546 ^ (v[1] & 0xffff0000) \
555 ^ (v[7] & 0x0000ffff); \
556 h[1] = (v[0] << 16) \
558 ^ (v[0] & 0xffff0000) \
559 ^ (v[1] & 0x0000ffff) \
567 ^ (v[7] & 0xffff0000) \
569 h[2] = (v[0] & 0x0000ffff) \
573 ^ (v[1] & 0xffff0000) \
581 ^ (v[7] & 0x0000ffff) \
584 h[3] = (v[0] << 16) \
586 ^ (v[0] & 0xffff0000) \
587 ^ (v[1] & 0xffff0000) \
597 ^ (v[7] & 0x0000ffff) \
599 h[4] = (v[0] >> 16) \
613 h[5] = (v[0] << 16) \
614 ^ (v[0] & 0xffff0000) \
617 ^ (v[1] & 0xffff0000) \
631 ^ (v[7] & 0xffff0000); \
663 #define PASS0(h,s,u,v,t) \
674 #define PASS2(h,s,u,v,t) \
686 #define PASS4(h,s,u,v,t) \
697 #define PASS6(h,s,u,v,t) \
706 static void m06900m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 (*s_tables)[256])
712 const u32 gid = get_global_id (0);
713 const u32 lid = get_local_id (0);
719 const u32 w14 = pw_len * 8;
727 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE)
729 const u32x w0r = w0r_create_bft (bfs_buf, il_pos);
731 const u32x w0lr = w0l | w0r;
768 state_m[0] = state[0];
769 state_m[1] = state[1];
770 state_m[2] = state[2];
771 state_m[3] = state[3];
772 state_m[4] = state[4];
773 state_m[5] = state[5];
774 state_m[6] = state[6];
775 state_m[7] = state[7];
790 PASS0 (state, tmp, state_m, data_m, s_tables);
791 PASS2 (state, tmp, state_m, data_m, s_tables);
792 PASS4 (state, tmp, state_m, data_m, s_tables);
793 PASS6 (state, tmp, state_m, data_m, s_tables);
795 SHIFT12 (state_m, data, tmp);
796 SHIFT16 (state, data_m, state_m);
797 SHIFT61 (state, data_m);
811 state_m[0] = state[0];
812 state_m[1] = state[1];
813 state_m[2] = state[2];
814 state_m[3] = state[3];
815 state_m[4] = state[4];
816 state_m[5] = state[5];
817 state_m[6] = state[6];
818 state_m[7] = state[7];
829 PASS0 (state, tmp, state_m, data_m, s_tables);
830 PASS2 (state, tmp, state_m, data_m, s_tables);
831 PASS4 (state, tmp, state_m, data_m, s_tables);
832 PASS6 (state, tmp, state_m, data_m, s_tables);
834 SHIFT12 (state_m, data, tmp);
835 SHIFT16 (state, data_m, state_m);
836 SHIFT61 (state, data_m);
849 state_m[0] = state[0];
850 state_m[1] = state[1];
851 state_m[2] = state[2];
852 state_m[3] = state[3];
853 state_m[4] = state[4];
854 state_m[5] = state[5];
855 state_m[6] = state[6];
856 state_m[7] = state[7];
867 PASS0 (state, tmp, state_m, data_m, s_tables);
868 PASS2 (state, tmp, state_m, data_m, s_tables);
869 PASS4 (state, tmp, state_m, data_m, s_tables);
870 PASS6 (state, tmp, state_m, data_m, s_tables);
872 SHIFT12 (state_m, data, tmp);
873 SHIFT16 (state, data_m, state_m);
874 SHIFT61 (state, data_m);
878 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
882 static void m06900s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 (*s_tables)[256])
888 const u32 gid = get_global_id (0);
889 const u32 lid = get_local_id (0);
895 const u32 w14 = pw_len * 8;
901 const u32 search[4] =
903 digests_buf[digests_offset].digest_buf[DGST_R0],
904 digests_buf[digests_offset].digest_buf[DGST_R1],
905 digests_buf[digests_offset].digest_buf[DGST_R2],
906 digests_buf[digests_offset].digest_buf[DGST_R3]
915 for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos += VECT_SIZE)
917 const u32x w0r = w0r_create_bft (bfs_buf, il_pos);
919 const u32x w0lr = w0l | w0r;
956 state_m[0] = state[0];
957 state_m[1] = state[1];
958 state_m[2] = state[2];
959 state_m[3] = state[3];
960 state_m[4] = state[4];
961 state_m[5] = state[5];
962 state_m[6] = state[6];
963 state_m[7] = state[7];
978 PASS0 (state, tmp, state_m, data_m, s_tables);
979 PASS2 (state, tmp, state_m, data_m, s_tables);
980 PASS4 (state, tmp, state_m, data_m, s_tables);
981 PASS6 (state, tmp, state_m, data_m, s_tables);
983 SHIFT12 (state_m, data, tmp);
984 SHIFT16 (state, data_m, state_m);
985 SHIFT61 (state, data_m);
999 state_m[0] = state[0];
1000 state_m[1] = state[1];
1001 state_m[2] = state[2];
1002 state_m[3] = state[3];
1003 state_m[4] = state[4];
1004 state_m[5] = state[5];
1005 state_m[6] = state[6];
1006 state_m[7] = state[7];
1008 data_m[0] = data[0];
1009 data_m[1] = data[1];
1010 data_m[2] = data[2];
1011 data_m[3] = data[3];
1012 data_m[4] = data[4];
1013 data_m[5] = data[5];
1014 data_m[6] = data[6];
1015 data_m[7] = data[7];
1017 PASS0 (state, tmp, state_m, data_m, s_tables);
1018 PASS2 (state, tmp, state_m, data_m, s_tables);
1019 PASS4 (state, tmp, state_m, data_m, s_tables);
1020 PASS6 (state, tmp, state_m, data_m, s_tables);
1022 SHIFT12 (state_m, data, tmp);
1023 SHIFT16 (state, data_m, state_m);
1024 SHIFT61 (state, data_m);
1028 data[0] = state[ 8];
1029 data[1] = state[ 9];
1030 data[2] = state[10];
1031 data[3] = state[11];
1032 data[4] = state[12];
1033 data[5] = state[13];
1034 data[6] = state[14];
1035 data[7] = state[15];
1037 state_m[0] = state[0];
1038 state_m[1] = state[1];
1039 state_m[2] = state[2];
1040 state_m[3] = state[3];
1041 state_m[4] = state[4];
1042 state_m[5] = state[5];
1043 state_m[6] = state[6];
1044 state_m[7] = state[7];
1046 data_m[0] = data[0];
1047 data_m[1] = data[1];
1048 data_m[2] = data[2];
1049 data_m[3] = data[3];
1050 data_m[4] = data[4];
1051 data_m[5] = data[5];
1052 data_m[6] = data[6];
1053 data_m[7] = data[7];
1055 PASS0 (state, tmp, state_m, data_m, s_tables);
1056 PASS2 (state, tmp, state_m, data_m, s_tables);
1057 PASS4 (state, tmp, state_m, data_m, s_tables);
1058 PASS6 (state, tmp, state_m, data_m, s_tables);
1060 SHIFT12 (state_m, data, tmp);
1061 SHIFT16 (state, data_m, state_m);
1062 SHIFT61 (state, data_m);
1066 COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
1070 __kernel void m06900_m04 (__global pw_t *pws, __global kernel_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1076 const u32 gid = get_global_id (0);
1077 const u32 lid = get_local_id (0);
1078 const u32 lsz = get_local_size (0);
1084 __local u32 s_tables[4][256];
1086 for (u32 i = lid; i < 256; i += lsz)
1088 s_tables[0][i] = c_tables[0][i];
1089 s_tables[1][i] = c_tables[1][i];
1090 s_tables[2][i] = c_tables[2][i];
1091 s_tables[3][i] = c_tables[3][i];
1094 barrier (CLK_LOCAL_MEM_FENCE);
1096 if (gid >= gid_max) return;
1104 w0[0] = pws[gid].i[ 0];
1105 w0[1] = pws[gid].i[ 1];
1106 w0[2] = pws[gid].i[ 2];
1107 w0[3] = pws[gid].i[ 3];
1130 const u32 pw_len = pws[gid].pw_len;
1136 m06900m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, s_tables);
1139 __kernel void m06900_m08 (__global pw_t *pws, __global kernel_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1145 const u32 gid = get_global_id (0);
1146 const u32 lid = get_local_id (0);
1147 const u32 lsz = get_local_size (0);
1153 __local u32 s_tables[4][256];
1155 for (u32 i = lid; i < 256; i += lsz)
1157 s_tables[0][i] = c_tables[0][i];
1158 s_tables[1][i] = c_tables[1][i];
1159 s_tables[2][i] = c_tables[2][i];
1160 s_tables[3][i] = c_tables[3][i];
1163 barrier (CLK_LOCAL_MEM_FENCE);
1165 if (gid >= gid_max) return;
1173 w0[0] = pws[gid].i[ 0];
1174 w0[1] = pws[gid].i[ 1];
1175 w0[2] = pws[gid].i[ 2];
1176 w0[3] = pws[gid].i[ 3];
1180 w1[0] = pws[gid].i[ 4];
1181 w1[1] = pws[gid].i[ 5];
1182 w1[2] = pws[gid].i[ 6];
1183 w1[3] = pws[gid].i[ 7];
1199 const u32 pw_len = pws[gid].pw_len;
1205 m06900m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, s_tables);
1208 __kernel void m06900_m16 (__global pw_t *pws, __global kernel_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1212 __kernel void m06900_s04 (__global pw_t *pws, __global kernel_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1218 const u32 gid = get_global_id (0);
1219 const u32 lid = get_local_id (0);
1220 const u32 lsz = get_local_size (0);
1226 __local u32 s_tables[4][256];
1228 for (u32 i = lid; i < 256; i += lsz)
1230 s_tables[0][i] = c_tables[0][i];
1231 s_tables[1][i] = c_tables[1][i];
1232 s_tables[2][i] = c_tables[2][i];
1233 s_tables[3][i] = c_tables[3][i];
1236 barrier (CLK_LOCAL_MEM_FENCE);
1238 if (gid >= gid_max) return;
1246 w0[0] = pws[gid].i[ 0];
1247 w0[1] = pws[gid].i[ 1];
1248 w0[2] = pws[gid].i[ 2];
1249 w0[3] = pws[gid].i[ 3];
1272 const u32 pw_len = pws[gid].pw_len;
1278 m06900s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, s_tables);
1281 __kernel void m06900_s08 (__global pw_t *pws, __global kernel_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1287 const u32 gid = get_global_id (0);
1288 const u32 lid = get_local_id (0);
1289 const u32 lsz = get_local_size (0);
1295 __local u32 s_tables[4][256];
1297 for (u32 i = lid; i < 256; i += lsz)
1299 s_tables[0][i] = c_tables[0][i];
1300 s_tables[1][i] = c_tables[1][i];
1301 s_tables[2][i] = c_tables[2][i];
1302 s_tables[3][i] = c_tables[3][i];
1305 barrier (CLK_LOCAL_MEM_FENCE);
1307 if (gid >= gid_max) return;
1315 w0[0] = pws[gid].i[ 0];
1316 w0[1] = pws[gid].i[ 1];
1317 w0[2] = pws[gid].i[ 2];
1318 w0[3] = pws[gid].i[ 3];
1322 w1[0] = pws[gid].i[ 4];
1323 w1[1] = pws[gid].i[ 5];
1324 w1[2] = pws[gid].i[ 6];
1325 w1[3] = pws[gid].i[ 7];
1341 const u32 pw_len = pws[gid].pw_len;
1347 m06900s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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, s_tables);
1350 __kernel void m06900_s16 (__global pw_t *pws, __global kernel_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)