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])
301 #elif VECT_SIZE == 16
302 #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], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
305 #define _round(k1,k2,tbl) \
309 l ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
310 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
311 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
312 BOX (((t >> 24) & 0xff), 3, tbl); \
314 r ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
315 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
316 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
317 BOX (((t >> 24) & 0xff), 3, tbl); \
320 #define R(k,h,s,i,t) \
326 _round (k[0], k[1], t); \
327 _round (k[2], k[3], t); \
328 _round (k[4], k[5], t); \
329 _round (k[6], k[7], t); \
330 _round (k[0], k[1], t); \
331 _round (k[2], k[3], t); \
332 _round (k[4], k[5], t); \
333 _round (k[6], k[7], t); \
334 _round (k[0], k[1], t); \
335 _round (k[2], k[3], t); \
336 _round (k[4], k[5], t); \
337 _round (k[6], k[7], t); \
338 _round (k[7], k[6], t); \
339 _round (k[5], k[4], t); \
340 _round (k[3], k[2], t); \
341 _round (k[1], k[0], t); \
347 w[0] = u[0] ^ v[0]; \
348 w[1] = u[1] ^ v[1]; \
349 w[2] = u[2] ^ v[2]; \
350 w[3] = u[3] ^ v[3]; \
351 w[4] = u[4] ^ v[4]; \
352 w[5] = u[5] ^ v[5]; \
353 w[6] = u[6] ^ v[6]; \
357 k[0] = ((w[0] & 0x000000ff) << 0) \
358 | ((w[2] & 0x000000ff) << 8) \
359 | ((w[4] & 0x000000ff) << 16) \
360 | ((w[6] & 0x000000ff) << 24); \
361 k[1] = ((w[0] & 0x0000ff00) >> 8) \
362 | ((w[2] & 0x0000ff00) >> 0) \
363 | ((w[4] & 0x0000ff00) << 8) \
364 | ((w[6] & 0x0000ff00) << 16); \
365 k[2] = ((w[0] & 0x00ff0000) >> 16) \
366 | ((w[2] & 0x00ff0000) >> 8) \
367 | ((w[4] & 0x00ff0000) << 0) \
368 | ((w[6] & 0x00ff0000) << 8); \
369 k[3] = ((w[0] & 0xff000000) >> 24) \
370 | ((w[2] & 0xff000000) >> 16) \
371 | ((w[4] & 0xff000000) >> 8) \
372 | ((w[6] & 0xff000000) >> 0); \
373 k[4] = ((w[1] & 0x000000ff) << 0) \
374 | ((w[3] & 0x000000ff) << 8) \
375 | ((w[5] & 0x000000ff) << 16) \
376 | ((w[7] & 0x000000ff) << 24); \
377 k[5] = ((w[1] & 0x0000ff00) >> 8) \
378 | ((w[3] & 0x0000ff00) >> 0) \
379 | ((w[5] & 0x0000ff00) << 8) \
380 | ((w[7] & 0x0000ff00) << 16); \
381 k[6] = ((w[1] & 0x00ff0000) >> 16) \
382 | ((w[3] & 0x00ff0000) >> 8) \
383 | ((w[5] & 0x00ff0000) << 0) \
384 | ((w[7] & 0x00ff0000) << 8); \
385 k[7] = ((w[1] & 0xff000000) >> 24) \
386 | ((w[3] & 0xff000000) >> 16) \
387 | ((w[5] & 0xff000000) >> 8) \
388 | ((w[7] & 0xff000000) >> 0);
425 x[0] ^= 0xff00ff00; \
426 x[1] ^= 0xff00ff00; \
427 x[2] ^= 0x00ff00ff; \
428 x[3] ^= 0x00ff00ff; \
429 x[4] ^= 0x00ffff00; \
430 x[5] ^= 0xff0000ff; \
431 x[6] ^= 0x000000ff; \
434 #define SHIFT12(u,m,s) \
435 u[0] = m[0] ^ s[6]; \
436 u[1] = m[1] ^ s[7]; \
437 u[2] = m[2] ^ (s[0] << 16) \
439 ^ (s[0] & 0x0000ffff) \
440 ^ (s[1] & 0x0000ffff) \
445 ^ (s[7] & 0xffff0000) \
447 u[3] = m[3] ^ (s[0] & 0x0000ffff) \
449 ^ (s[1] & 0x0000ffff) \
458 ^ (s[7] & 0x0000ffff) \
461 u[4] = m[4] ^ (s[0] & 0xffff0000) \
464 ^ (s[1] & 0xffff0000) \
473 ^ (s[7] & 0x0000ffff) \
476 u[5] = m[5] ^ (s[0] << 16) \
478 ^ (s[0] & 0xffff0000) \
479 ^ (s[1] & 0x0000ffff) \
489 ^ (s[7] & 0xffff0000) \
505 u[7] = m[7] ^ (s[0] & 0xffff0000) \
507 ^ (s[1] & 0x0000ffff) \
516 ^ (s[7] & 0x0000ffff) \
520 #define SHIFT16(h,v,u) \
521 v[0] = h[0] ^ (u[1] << 16) \
523 v[1] = h[1] ^ (u[2] << 16) \
525 v[2] = h[2] ^ (u[3] << 16) \
527 v[3] = h[3] ^ (u[4] << 16) \
529 v[4] = h[4] ^ (u[5] << 16) \
531 v[5] = h[5] ^ (u[6] << 16) \
533 v[6] = h[6] ^ (u[7] << 16) \
535 v[7] = h[7] ^ (u[0] & 0xffff0000) \
538 ^ (u[1] & 0xffff0000) \
541 ^ (u[7] & 0xffff0000);
543 #define SHIFT61(h,v) \
544 h[0] = (v[0] & 0xffff0000) \
548 ^ (v[1] & 0xffff0000) \
557 ^ (v[7] & 0x0000ffff); \
558 h[1] = (v[0] << 16) \
560 ^ (v[0] & 0xffff0000) \
561 ^ (v[1] & 0x0000ffff) \
569 ^ (v[7] & 0xffff0000) \
571 h[2] = (v[0] & 0x0000ffff) \
575 ^ (v[1] & 0xffff0000) \
583 ^ (v[7] & 0x0000ffff) \
586 h[3] = (v[0] << 16) \
588 ^ (v[0] & 0xffff0000) \
589 ^ (v[1] & 0xffff0000) \
599 ^ (v[7] & 0x0000ffff) \
601 h[4] = (v[0] >> 16) \
615 h[5] = (v[0] << 16) \
616 ^ (v[0] & 0xffff0000) \
619 ^ (v[1] & 0xffff0000) \
633 ^ (v[7] & 0xffff0000); \
665 #define PASS0(h,s,u,v,t) \
676 #define PASS2(h,s,u,v,t) \
688 #define PASS4(h,s,u,v,t) \
699 #define PASS6(h,s,u,v,t) \
708 __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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
714 const u32 gid = get_global_id (0);
715 const u32 lid = get_local_id (0);
716 const u32 lsz = get_local_size (0);
722 __local u32 s_tables[4][256];
724 for (u32 i = lid; i < 256; i += lsz)
726 s_tables[0][i] = c_tables[0][i];
727 s_tables[1][i] = c_tables[1][i];
728 s_tables[2][i] = c_tables[2][i];
729 s_tables[3][i] = c_tables[3][i];
732 barrier (CLK_LOCAL_MEM_FENCE);
734 if (gid >= gid_max) return;
743 pw_buf0[0] = pws[gid].i[0];
744 pw_buf0[1] = pws[gid].i[1];
745 pw_buf0[2] = pws[gid].i[2];
746 pw_buf0[3] = pws[gid].i[3];
747 pw_buf1[0] = pws[gid].i[4];
748 pw_buf1[1] = pws[gid].i[5];
749 pw_buf1[2] = pws[gid].i[6];
750 pw_buf1[3] = pws[gid].i[7];
752 const u32 pw_l_len = pws[gid].pw_len;
758 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
760 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
762 const u32x pw_len = pw_l_len + pw_r_len;
765 * concat password candidate
768 u32x wordl0[4] = { 0 };
769 u32x wordl1[4] = { 0 };
770 u32x wordl2[4] = { 0 };
771 u32x wordl3[4] = { 0 };
773 wordl0[0] = pw_buf0[0];
774 wordl0[1] = pw_buf0[1];
775 wordl0[2] = pw_buf0[2];
776 wordl0[3] = pw_buf0[3];
777 wordl1[0] = pw_buf1[0];
778 wordl1[1] = pw_buf1[1];
779 wordl1[2] = pw_buf1[2];
780 wordl1[3] = pw_buf1[3];
782 u32x wordr0[4] = { 0 };
783 u32x wordr1[4] = { 0 };
784 u32x wordr2[4] = { 0 };
785 u32x wordr3[4] = { 0 };
787 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
788 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
789 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
790 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
791 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
792 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
793 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
794 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
796 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
798 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
802 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
808 w0[0] = wordl0[0] | wordr0[0];
809 w0[1] = wordl0[1] | wordr0[1];
810 w0[2] = wordl0[2] | wordr0[2];
811 w0[3] = wordl0[3] | wordr0[3];
812 w1[0] = wordl1[0] | wordr1[0];
813 w1[1] = wordl1[1] | wordr1[1];
814 w1[2] = wordl1[2] | wordr1[2];
815 w1[3] = wordl1[3] | wordr1[3];
856 state_m[0] = state[0];
857 state_m[1] = state[1];
858 state_m[2] = state[2];
859 state_m[3] = state[3];
860 state_m[4] = state[4];
861 state_m[5] = state[5];
862 state_m[6] = state[6];
863 state_m[7] = state[7];
876 //if (pw_len > 0) // not really SIMD compatible
878 PASS0 (state, tmp, state_m, data_m, s_tables);
879 PASS2 (state, tmp, state_m, data_m, s_tables);
880 PASS4 (state, tmp, state_m, data_m, s_tables);
881 PASS6 (state, tmp, state_m, data_m, s_tables);
883 SHIFT12 (state_m, data, tmp);
884 SHIFT16 (state, data_m, state_m);
885 SHIFT61 (state, data_m);
888 data[0] = pw_len * 8;
899 state_m[0] = state[0];
900 state_m[1] = state[1];
901 state_m[2] = state[2];
902 state_m[3] = state[3];
903 state_m[4] = state[4];
904 state_m[5] = state[5];
905 state_m[6] = state[6];
906 state_m[7] = state[7];
917 PASS0 (state, tmp, state_m, data_m, s_tables);
918 PASS2 (state, tmp, state_m, data_m, s_tables);
919 PASS4 (state, tmp, state_m, data_m, s_tables);
920 PASS6 (state, tmp, state_m, data_m, s_tables);
922 SHIFT12 (state_m, data, tmp);
923 SHIFT16 (state, data_m, state_m);
924 SHIFT61 (state, data_m);
937 state_m[0] = state[0];
938 state_m[1] = state[1];
939 state_m[2] = state[2];
940 state_m[3] = state[3];
941 state_m[4] = state[4];
942 state_m[5] = state[5];
943 state_m[6] = state[6];
944 state_m[7] = state[7];
955 PASS0 (state, tmp, state_m, data_m, s_tables);
956 PASS2 (state, tmp, state_m, data_m, s_tables);
957 PASS4 (state, tmp, state_m, data_m, s_tables);
958 PASS6 (state, tmp, state_m, data_m, s_tables);
960 SHIFT12 (state_m, data, tmp);
961 SHIFT16 (state, data_m, state_m);
962 SHIFT61 (state, data_m);
966 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
970 __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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
974 __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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
978 __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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
984 const u32 gid = get_global_id (0);
985 const u32 lid = get_local_id (0);
986 const u32 lsz = get_local_size (0);
992 __local u32 s_tables[4][256];
994 for (u32 i = lid; i < 256; i += lsz)
996 s_tables[0][i] = c_tables[0][i];
997 s_tables[1][i] = c_tables[1][i];
998 s_tables[2][i] = c_tables[2][i];
999 s_tables[3][i] = c_tables[3][i];
1002 barrier (CLK_LOCAL_MEM_FENCE);
1004 if (gid >= gid_max) return;
1013 pw_buf0[0] = pws[gid].i[0];
1014 pw_buf0[1] = pws[gid].i[1];
1015 pw_buf0[2] = pws[gid].i[2];
1016 pw_buf0[3] = pws[gid].i[3];
1017 pw_buf1[0] = pws[gid].i[4];
1018 pw_buf1[1] = pws[gid].i[5];
1019 pw_buf1[2] = pws[gid].i[6];
1020 pw_buf1[3] = pws[gid].i[7];
1022 const u32 pw_l_len = pws[gid].pw_len;
1028 const u32 search[4] =
1030 digests_buf[digests_offset].digest_buf[DGST_R0],
1031 digests_buf[digests_offset].digest_buf[DGST_R1],
1032 digests_buf[digests_offset].digest_buf[DGST_R2],
1033 digests_buf[digests_offset].digest_buf[DGST_R3]
1040 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
1042 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
1044 const u32x pw_len = pw_l_len + pw_r_len;
1047 * concat password candidate
1050 u32x wordl0[4] = { 0 };
1051 u32x wordl1[4] = { 0 };
1052 u32x wordl2[4] = { 0 };
1053 u32x wordl3[4] = { 0 };
1055 wordl0[0] = pw_buf0[0];
1056 wordl0[1] = pw_buf0[1];
1057 wordl0[2] = pw_buf0[2];
1058 wordl0[3] = pw_buf0[3];
1059 wordl1[0] = pw_buf1[0];
1060 wordl1[1] = pw_buf1[1];
1061 wordl1[2] = pw_buf1[2];
1062 wordl1[3] = pw_buf1[3];
1064 u32x wordr0[4] = { 0 };
1065 u32x wordr1[4] = { 0 };
1066 u32x wordr2[4] = { 0 };
1067 u32x wordr3[4] = { 0 };
1069 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
1070 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
1071 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
1072 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
1073 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
1074 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
1075 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
1076 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
1078 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
1080 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
1084 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
1090 w0[0] = wordl0[0] | wordr0[0];
1091 w0[1] = wordl0[1] | wordr0[1];
1092 w0[2] = wordl0[2] | wordr0[2];
1093 w0[3] = wordl0[3] | wordr0[3];
1094 w1[0] = wordl1[0] | wordr1[0];
1095 w1[1] = wordl1[1] | wordr1[1];
1096 w1[2] = wordl1[2] | wordr1[2];
1097 w1[3] = wordl1[3] | wordr1[3];
1124 state[ 8] = data[0];
1125 state[ 9] = data[1];
1126 state[10] = data[2];
1127 state[11] = data[3];
1128 state[12] = data[4];
1129 state[13] = data[5];
1130 state[14] = data[6];
1131 state[15] = data[7];
1138 state_m[0] = state[0];
1139 state_m[1] = state[1];
1140 state_m[2] = state[2];
1141 state_m[3] = state[3];
1142 state_m[4] = state[4];
1143 state_m[5] = state[5];
1144 state_m[6] = state[6];
1145 state_m[7] = state[7];
1147 data_m[0] = data[0];
1148 data_m[1] = data[1];
1149 data_m[2] = data[2];
1150 data_m[3] = data[3];
1151 data_m[4] = data[4];
1152 data_m[5] = data[5];
1153 data_m[6] = data[6];
1154 data_m[7] = data[7];
1158 //if (pw_len > 0) // not really SIMD compatible
1160 PASS0 (state, tmp, state_m, data_m, s_tables);
1161 PASS2 (state, tmp, state_m, data_m, s_tables);
1162 PASS4 (state, tmp, state_m, data_m, s_tables);
1163 PASS6 (state, tmp, state_m, data_m, s_tables);
1165 SHIFT12 (state_m, data, tmp);
1166 SHIFT16 (state, data_m, state_m);
1167 SHIFT61 (state, data_m);
1170 data[0] = pw_len * 8;
1181 state_m[0] = state[0];
1182 state_m[1] = state[1];
1183 state_m[2] = state[2];
1184 state_m[3] = state[3];
1185 state_m[4] = state[4];
1186 state_m[5] = state[5];
1187 state_m[6] = state[6];
1188 state_m[7] = state[7];
1190 data_m[0] = data[0];
1191 data_m[1] = data[1];
1192 data_m[2] = data[2];
1193 data_m[3] = data[3];
1194 data_m[4] = data[4];
1195 data_m[5] = data[5];
1196 data_m[6] = data[6];
1197 data_m[7] = data[7];
1199 PASS0 (state, tmp, state_m, data_m, s_tables);
1200 PASS2 (state, tmp, state_m, data_m, s_tables);
1201 PASS4 (state, tmp, state_m, data_m, s_tables);
1202 PASS6 (state, tmp, state_m, data_m, s_tables);
1204 SHIFT12 (state_m, data, tmp);
1205 SHIFT16 (state, data_m, state_m);
1206 SHIFT61 (state, data_m);
1210 data[0] = state[ 8];
1211 data[1] = state[ 9];
1212 data[2] = state[10];
1213 data[3] = state[11];
1214 data[4] = state[12];
1215 data[5] = state[13];
1216 data[6] = state[14];
1217 data[7] = state[15];
1219 state_m[0] = state[0];
1220 state_m[1] = state[1];
1221 state_m[2] = state[2];
1222 state_m[3] = state[3];
1223 state_m[4] = state[4];
1224 state_m[5] = state[5];
1225 state_m[6] = state[6];
1226 state_m[7] = state[7];
1228 data_m[0] = data[0];
1229 data_m[1] = data[1];
1230 data_m[2] = data[2];
1231 data_m[3] = data[3];
1232 data_m[4] = data[4];
1233 data_m[5] = data[5];
1234 data_m[6] = data[6];
1235 data_m[7] = data[7];
1237 PASS0 (state, tmp, state_m, data_m, s_tables);
1238 PASS2 (state, tmp, state_m, data_m, s_tables);
1239 PASS4 (state, tmp, state_m, data_m, s_tables);
1240 PASS6 (state, tmp, state_m, data_m, s_tables);
1242 SHIFT12 (state_m, data, tmp);
1243 SHIFT16 (state, data_m, state_m);
1244 SHIFT61 (state, data_m);
1248 COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
1252 __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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1256 __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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)