2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
10 #include "include/constants.h"
11 #include "include/kernel_vendor.h"
18 #include "include/kernel_functions.c"
19 #include "OpenCL/types_ocl.c"
20 #include "OpenCL/common.c"
21 #include "include/rp_kernel.h"
22 #include "OpenCL/rp.c"
24 #define COMPARE_S "OpenCL/check_single_comp4.c"
25 #define COMPARE_M "OpenCL/check_multi_comp4.c"
27 __constant u32 c_tables[4][256] =
30 0x00072000, 0x00075000, 0x00074800, 0x00071000,
31 0x00076800, 0x00074000, 0x00070000, 0x00077000,
32 0x00073000, 0x00075800, 0x00070800, 0x00076000,
33 0x00073800, 0x00077800, 0x00072800, 0x00071800,
34 0x0005a000, 0x0005d000, 0x0005c800, 0x00059000,
35 0x0005e800, 0x0005c000, 0x00058000, 0x0005f000,
36 0x0005b000, 0x0005d800, 0x00058800, 0x0005e000,
37 0x0005b800, 0x0005f800, 0x0005a800, 0x00059800,
38 0x00022000, 0x00025000, 0x00024800, 0x00021000,
39 0x00026800, 0x00024000, 0x00020000, 0x00027000,
40 0x00023000, 0x00025800, 0x00020800, 0x00026000,
41 0x00023800, 0x00027800, 0x00022800, 0x00021800,
42 0x00062000, 0x00065000, 0x00064800, 0x00061000,
43 0x00066800, 0x00064000, 0x00060000, 0x00067000,
44 0x00063000, 0x00065800, 0x00060800, 0x00066000,
45 0x00063800, 0x00067800, 0x00062800, 0x00061800,
46 0x00032000, 0x00035000, 0x00034800, 0x00031000,
47 0x00036800, 0x00034000, 0x00030000, 0x00037000,
48 0x00033000, 0x00035800, 0x00030800, 0x00036000,
49 0x00033800, 0x00037800, 0x00032800, 0x00031800,
50 0x0006a000, 0x0006d000, 0x0006c800, 0x00069000,
51 0x0006e800, 0x0006c000, 0x00068000, 0x0006f000,
52 0x0006b000, 0x0006d800, 0x00068800, 0x0006e000,
53 0x0006b800, 0x0006f800, 0x0006a800, 0x00069800,
54 0x0007a000, 0x0007d000, 0x0007c800, 0x00079000,
55 0x0007e800, 0x0007c000, 0x00078000, 0x0007f000,
56 0x0007b000, 0x0007d800, 0x00078800, 0x0007e000,
57 0x0007b800, 0x0007f800, 0x0007a800, 0x00079800,
58 0x00052000, 0x00055000, 0x00054800, 0x00051000,
59 0x00056800, 0x00054000, 0x00050000, 0x00057000,
60 0x00053000, 0x00055800, 0x00050800, 0x00056000,
61 0x00053800, 0x00057800, 0x00052800, 0x00051800,
62 0x00012000, 0x00015000, 0x00014800, 0x00011000,
63 0x00016800, 0x00014000, 0x00010000, 0x00017000,
64 0x00013000, 0x00015800, 0x00010800, 0x00016000,
65 0x00013800, 0x00017800, 0x00012800, 0x00011800,
66 0x0001a000, 0x0001d000, 0x0001c800, 0x00019000,
67 0x0001e800, 0x0001c000, 0x00018000, 0x0001f000,
68 0x0001b000, 0x0001d800, 0x00018800, 0x0001e000,
69 0x0001b800, 0x0001f800, 0x0001a800, 0x00019800,
70 0x00042000, 0x00045000, 0x00044800, 0x00041000,
71 0x00046800, 0x00044000, 0x00040000, 0x00047000,
72 0x00043000, 0x00045800, 0x00040800, 0x00046000,
73 0x00043800, 0x00047800, 0x00042800, 0x00041800,
74 0x0000a000, 0x0000d000, 0x0000c800, 0x00009000,
75 0x0000e800, 0x0000c000, 0x00008000, 0x0000f000,
76 0x0000b000, 0x0000d800, 0x00008800, 0x0000e000,
77 0x0000b800, 0x0000f800, 0x0000a800, 0x00009800,
78 0x00002000, 0x00005000, 0x00004800, 0x00001000,
79 0x00006800, 0x00004000, 0x00000000, 0x00007000,
80 0x00003000, 0x00005800, 0x00000800, 0x00006000,
81 0x00003800, 0x00007800, 0x00002800, 0x00001800,
82 0x0003a000, 0x0003d000, 0x0003c800, 0x00039000,
83 0x0003e800, 0x0003c000, 0x00038000, 0x0003f000,
84 0x0003b000, 0x0003d800, 0x00038800, 0x0003e000,
85 0x0003b800, 0x0003f800, 0x0003a800, 0x00039800,
86 0x0002a000, 0x0002d000, 0x0002c800, 0x00029000,
87 0x0002e800, 0x0002c000, 0x00028000, 0x0002f000,
88 0x0002b000, 0x0002d800, 0x00028800, 0x0002e000,
89 0x0002b800, 0x0002f800, 0x0002a800, 0x00029800,
90 0x0004a000, 0x0004d000, 0x0004c800, 0x00049000,
91 0x0004e800, 0x0004c000, 0x00048000, 0x0004f000,
92 0x0004b000, 0x0004d800, 0x00048800, 0x0004e000,
93 0x0004b800, 0x0004f800, 0x0004a800, 0x00049800,
96 0x03a80000, 0x03c00000, 0x03880000, 0x03e80000,
97 0x03d00000, 0x03980000, 0x03a00000, 0x03900000,
98 0x03f00000, 0x03f80000, 0x03e00000, 0x03b80000,
99 0x03b00000, 0x03800000, 0x03c80000, 0x03d80000,
100 0x06a80000, 0x06c00000, 0x06880000, 0x06e80000,
101 0x06d00000, 0x06980000, 0x06a00000, 0x06900000,
102 0x06f00000, 0x06f80000, 0x06e00000, 0x06b80000,
103 0x06b00000, 0x06800000, 0x06c80000, 0x06d80000,
104 0x05280000, 0x05400000, 0x05080000, 0x05680000,
105 0x05500000, 0x05180000, 0x05200000, 0x05100000,
106 0x05700000, 0x05780000, 0x05600000, 0x05380000,
107 0x05300000, 0x05000000, 0x05480000, 0x05580000,
108 0x00a80000, 0x00c00000, 0x00880000, 0x00e80000,
109 0x00d00000, 0x00980000, 0x00a00000, 0x00900000,
110 0x00f00000, 0x00f80000, 0x00e00000, 0x00b80000,
111 0x00b00000, 0x00800000, 0x00c80000, 0x00d80000,
112 0x00280000, 0x00400000, 0x00080000, 0x00680000,
113 0x00500000, 0x00180000, 0x00200000, 0x00100000,
114 0x00700000, 0x00780000, 0x00600000, 0x00380000,
115 0x00300000, 0x00000000, 0x00480000, 0x00580000,
116 0x04280000, 0x04400000, 0x04080000, 0x04680000,
117 0x04500000, 0x04180000, 0x04200000, 0x04100000,
118 0x04700000, 0x04780000, 0x04600000, 0x04380000,
119 0x04300000, 0x04000000, 0x04480000, 0x04580000,
120 0x04a80000, 0x04c00000, 0x04880000, 0x04e80000,
121 0x04d00000, 0x04980000, 0x04a00000, 0x04900000,
122 0x04f00000, 0x04f80000, 0x04e00000, 0x04b80000,
123 0x04b00000, 0x04800000, 0x04c80000, 0x04d80000,
124 0x07a80000, 0x07c00000, 0x07880000, 0x07e80000,
125 0x07d00000, 0x07980000, 0x07a00000, 0x07900000,
126 0x07f00000, 0x07f80000, 0x07e00000, 0x07b80000,
127 0x07b00000, 0x07800000, 0x07c80000, 0x07d80000,
128 0x07280000, 0x07400000, 0x07080000, 0x07680000,
129 0x07500000, 0x07180000, 0x07200000, 0x07100000,
130 0x07700000, 0x07780000, 0x07600000, 0x07380000,
131 0x07300000, 0x07000000, 0x07480000, 0x07580000,
132 0x02280000, 0x02400000, 0x02080000, 0x02680000,
133 0x02500000, 0x02180000, 0x02200000, 0x02100000,
134 0x02700000, 0x02780000, 0x02600000, 0x02380000,
135 0x02300000, 0x02000000, 0x02480000, 0x02580000,
136 0x03280000, 0x03400000, 0x03080000, 0x03680000,
137 0x03500000, 0x03180000, 0x03200000, 0x03100000,
138 0x03700000, 0x03780000, 0x03600000, 0x03380000,
139 0x03300000, 0x03000000, 0x03480000, 0x03580000,
140 0x06280000, 0x06400000, 0x06080000, 0x06680000,
141 0x06500000, 0x06180000, 0x06200000, 0x06100000,
142 0x06700000, 0x06780000, 0x06600000, 0x06380000,
143 0x06300000, 0x06000000, 0x06480000, 0x06580000,
144 0x05a80000, 0x05c00000, 0x05880000, 0x05e80000,
145 0x05d00000, 0x05980000, 0x05a00000, 0x05900000,
146 0x05f00000, 0x05f80000, 0x05e00000, 0x05b80000,
147 0x05b00000, 0x05800000, 0x05c80000, 0x05d80000,
148 0x01280000, 0x01400000, 0x01080000, 0x01680000,
149 0x01500000, 0x01180000, 0x01200000, 0x01100000,
150 0x01700000, 0x01780000, 0x01600000, 0x01380000,
151 0x01300000, 0x01000000, 0x01480000, 0x01580000,
152 0x02a80000, 0x02c00000, 0x02880000, 0x02e80000,
153 0x02d00000, 0x02980000, 0x02a00000, 0x02900000,
154 0x02f00000, 0x02f80000, 0x02e00000, 0x02b80000,
155 0x02b00000, 0x02800000, 0x02c80000, 0x02d80000,
156 0x01a80000, 0x01c00000, 0x01880000, 0x01e80000,
157 0x01d00000, 0x01980000, 0x01a00000, 0x01900000,
158 0x01f00000, 0x01f80000, 0x01e00000, 0x01b80000,
159 0x01b00000, 0x01800000, 0x01c80000, 0x01d80000,
162 0x30000002, 0x60000002, 0x38000002, 0x08000002,
163 0x28000002, 0x78000002, 0x68000002, 0x40000002,
164 0x20000002, 0x50000002, 0x48000002, 0x70000002,
165 0x00000002, 0x18000002, 0x58000002, 0x10000002,
166 0xb0000005, 0xe0000005, 0xb8000005, 0x88000005,
167 0xa8000005, 0xf8000005, 0xe8000005, 0xc0000005,
168 0xa0000005, 0xd0000005, 0xc8000005, 0xf0000005,
169 0x80000005, 0x98000005, 0xd8000005, 0x90000005,
170 0x30000005, 0x60000005, 0x38000005, 0x08000005,
171 0x28000005, 0x78000005, 0x68000005, 0x40000005,
172 0x20000005, 0x50000005, 0x48000005, 0x70000005,
173 0x00000005, 0x18000005, 0x58000005, 0x10000005,
174 0x30000000, 0x60000000, 0x38000000, 0x08000000,
175 0x28000000, 0x78000000, 0x68000000, 0x40000000,
176 0x20000000, 0x50000000, 0x48000000, 0x70000000,
177 0x00000000, 0x18000000, 0x58000000, 0x10000000,
178 0xb0000003, 0xe0000003, 0xb8000003, 0x88000003,
179 0xa8000003, 0xf8000003, 0xe8000003, 0xc0000003,
180 0xa0000003, 0xd0000003, 0xc8000003, 0xf0000003,
181 0x80000003, 0x98000003, 0xd8000003, 0x90000003,
182 0x30000001, 0x60000001, 0x38000001, 0x08000001,
183 0x28000001, 0x78000001, 0x68000001, 0x40000001,
184 0x20000001, 0x50000001, 0x48000001, 0x70000001,
185 0x00000001, 0x18000001, 0x58000001, 0x10000001,
186 0xb0000000, 0xe0000000, 0xb8000000, 0x88000000,
187 0xa8000000, 0xf8000000, 0xe8000000, 0xc0000000,
188 0xa0000000, 0xd0000000, 0xc8000000, 0xf0000000,
189 0x80000000, 0x98000000, 0xd8000000, 0x90000000,
190 0xb0000006, 0xe0000006, 0xb8000006, 0x88000006,
191 0xa8000006, 0xf8000006, 0xe8000006, 0xc0000006,
192 0xa0000006, 0xd0000006, 0xc8000006, 0xf0000006,
193 0x80000006, 0x98000006, 0xd8000006, 0x90000006,
194 0xb0000001, 0xe0000001, 0xb8000001, 0x88000001,
195 0xa8000001, 0xf8000001, 0xe8000001, 0xc0000001,
196 0xa0000001, 0xd0000001, 0xc8000001, 0xf0000001,
197 0x80000001, 0x98000001, 0xd8000001, 0x90000001,
198 0x30000003, 0x60000003, 0x38000003, 0x08000003,
199 0x28000003, 0x78000003, 0x68000003, 0x40000003,
200 0x20000003, 0x50000003, 0x48000003, 0x70000003,
201 0x00000003, 0x18000003, 0x58000003, 0x10000003,
202 0x30000004, 0x60000004, 0x38000004, 0x08000004,
203 0x28000004, 0x78000004, 0x68000004, 0x40000004,
204 0x20000004, 0x50000004, 0x48000004, 0x70000004,
205 0x00000004, 0x18000004, 0x58000004, 0x10000004,
206 0xb0000002, 0xe0000002, 0xb8000002, 0x88000002,
207 0xa8000002, 0xf8000002, 0xe8000002, 0xc0000002,
208 0xa0000002, 0xd0000002, 0xc8000002, 0xf0000002,
209 0x80000002, 0x98000002, 0xd8000002, 0x90000002,
210 0xb0000004, 0xe0000004, 0xb8000004, 0x88000004,
211 0xa8000004, 0xf8000004, 0xe8000004, 0xc0000004,
212 0xa0000004, 0xd0000004, 0xc8000004, 0xf0000004,
213 0x80000004, 0x98000004, 0xd8000004, 0x90000004,
214 0x30000006, 0x60000006, 0x38000006, 0x08000006,
215 0x28000006, 0x78000006, 0x68000006, 0x40000006,
216 0x20000006, 0x50000006, 0x48000006, 0x70000006,
217 0x00000006, 0x18000006, 0x58000006, 0x10000006,
218 0xb0000007, 0xe0000007, 0xb8000007, 0x88000007,
219 0xa8000007, 0xf8000007, 0xe8000007, 0xc0000007,
220 0xa0000007, 0xd0000007, 0xc8000007, 0xf0000007,
221 0x80000007, 0x98000007, 0xd8000007, 0x90000007,
222 0x30000007, 0x60000007, 0x38000007, 0x08000007,
223 0x28000007, 0x78000007, 0x68000007, 0x40000007,
224 0x20000007, 0x50000007, 0x48000007, 0x70000007,
225 0x00000007, 0x18000007, 0x58000007, 0x10000007,
228 0x000000e8, 0x000000d8, 0x000000a0, 0x00000088,
229 0x00000098, 0x000000f8, 0x000000a8, 0x000000c8,
230 0x00000080, 0x000000d0, 0x000000f0, 0x000000b8,
231 0x000000b0, 0x000000c0, 0x00000090, 0x000000e0,
232 0x000007e8, 0x000007d8, 0x000007a0, 0x00000788,
233 0x00000798, 0x000007f8, 0x000007a8, 0x000007c8,
234 0x00000780, 0x000007d0, 0x000007f0, 0x000007b8,
235 0x000007b0, 0x000007c0, 0x00000790, 0x000007e0,
236 0x000006e8, 0x000006d8, 0x000006a0, 0x00000688,
237 0x00000698, 0x000006f8, 0x000006a8, 0x000006c8,
238 0x00000680, 0x000006d0, 0x000006f0, 0x000006b8,
239 0x000006b0, 0x000006c0, 0x00000690, 0x000006e0,
240 0x00000068, 0x00000058, 0x00000020, 0x00000008,
241 0x00000018, 0x00000078, 0x00000028, 0x00000048,
242 0x00000000, 0x00000050, 0x00000070, 0x00000038,
243 0x00000030, 0x00000040, 0x00000010, 0x00000060,
244 0x000002e8, 0x000002d8, 0x000002a0, 0x00000288,
245 0x00000298, 0x000002f8, 0x000002a8, 0x000002c8,
246 0x00000280, 0x000002d0, 0x000002f0, 0x000002b8,
247 0x000002b0, 0x000002c0, 0x00000290, 0x000002e0,
248 0x000003e8, 0x000003d8, 0x000003a0, 0x00000388,
249 0x00000398, 0x000003f8, 0x000003a8, 0x000003c8,
250 0x00000380, 0x000003d0, 0x000003f0, 0x000003b8,
251 0x000003b0, 0x000003c0, 0x00000390, 0x000003e0,
252 0x00000568, 0x00000558, 0x00000520, 0x00000508,
253 0x00000518, 0x00000578, 0x00000528, 0x00000548,
254 0x00000500, 0x00000550, 0x00000570, 0x00000538,
255 0x00000530, 0x00000540, 0x00000510, 0x00000560,
256 0x00000268, 0x00000258, 0x00000220, 0x00000208,
257 0x00000218, 0x00000278, 0x00000228, 0x00000248,
258 0x00000200, 0x00000250, 0x00000270, 0x00000238,
259 0x00000230, 0x00000240, 0x00000210, 0x00000260,
260 0x000004e8, 0x000004d8, 0x000004a0, 0x00000488,
261 0x00000498, 0x000004f8, 0x000004a8, 0x000004c8,
262 0x00000480, 0x000004d0, 0x000004f0, 0x000004b8,
263 0x000004b0, 0x000004c0, 0x00000490, 0x000004e0,
264 0x00000168, 0x00000158, 0x00000120, 0x00000108,
265 0x00000118, 0x00000178, 0x00000128, 0x00000148,
266 0x00000100, 0x00000150, 0x00000170, 0x00000138,
267 0x00000130, 0x00000140, 0x00000110, 0x00000160,
268 0x000001e8, 0x000001d8, 0x000001a0, 0x00000188,
269 0x00000198, 0x000001f8, 0x000001a8, 0x000001c8,
270 0x00000180, 0x000001d0, 0x000001f0, 0x000001b8,
271 0x000001b0, 0x000001c0, 0x00000190, 0x000001e0,
272 0x00000768, 0x00000758, 0x00000720, 0x00000708,
273 0x00000718, 0x00000778, 0x00000728, 0x00000748,
274 0x00000700, 0x00000750, 0x00000770, 0x00000738,
275 0x00000730, 0x00000740, 0x00000710, 0x00000760,
276 0x00000368, 0x00000358, 0x00000320, 0x00000308,
277 0x00000318, 0x00000378, 0x00000328, 0x00000348,
278 0x00000300, 0x00000350, 0x00000370, 0x00000338,
279 0x00000330, 0x00000340, 0x00000310, 0x00000360,
280 0x000005e8, 0x000005d8, 0x000005a0, 0x00000588,
281 0x00000598, 0x000005f8, 0x000005a8, 0x000005c8,
282 0x00000580, 0x000005d0, 0x000005f0, 0x000005b8,
283 0x000005b0, 0x000005c0, 0x00000590, 0x000005e0,
284 0x00000468, 0x00000458, 0x00000420, 0x00000408,
285 0x00000418, 0x00000478, 0x00000428, 0x00000448,
286 0x00000400, 0x00000450, 0x00000470, 0x00000438,
287 0x00000430, 0x00000440, 0x00000410, 0x00000460,
288 0x00000668, 0x00000658, 0x00000620, 0x00000608,
289 0x00000618, 0x00000678, 0x00000628, 0x00000648,
290 0x00000600, 0x00000650, 0x00000670, 0x00000638,
291 0x00000630, 0x00000640, 0x00000610, 0x00000660,
296 #define BOX(i,n,S) (S)[(n)][(i)]
298 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
300 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
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])
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 rules_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;
742 pw_buf0[0] = pws[gid].i[ 0];
743 pw_buf0[1] = pws[gid].i[ 1];
744 pw_buf0[2] = pws[gid].i[ 2];
745 pw_buf0[3] = pws[gid].i[ 3];
749 pw_buf1[0] = pws[gid].i[ 4];
750 pw_buf1[1] = pws[gid].i[ 5];
751 pw_buf1[2] = pws[gid].i[ 6];
752 pw_buf1[3] = pws[gid].i[ 7];
754 const u32 pw_len = pws[gid].pw_len;
760 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
790 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
792 u32 w14 = out_len * 8;
829 state_m[0] = state[0];
830 state_m[1] = state[1];
831 state_m[2] = state[2];
832 state_m[3] = state[3];
833 state_m[4] = state[4];
834 state_m[5] = state[5];
835 state_m[6] = state[6];
836 state_m[7] = state[7];
851 PASS0 (state, tmp, state_m, data_m, s_tables);
852 PASS2 (state, tmp, state_m, data_m, s_tables);
853 PASS4 (state, tmp, state_m, data_m, s_tables);
854 PASS6 (state, tmp, state_m, data_m, s_tables);
856 SHIFT12 (state_m, data, tmp);
857 SHIFT16 (state, data_m, state_m);
858 SHIFT61 (state, data_m);
872 state_m[0] = state[0];
873 state_m[1] = state[1];
874 state_m[2] = state[2];
875 state_m[3] = state[3];
876 state_m[4] = state[4];
877 state_m[5] = state[5];
878 state_m[6] = state[6];
879 state_m[7] = state[7];
890 PASS0 (state, tmp, state_m, data_m, s_tables);
891 PASS2 (state, tmp, state_m, data_m, s_tables);
892 PASS4 (state, tmp, state_m, data_m, s_tables);
893 PASS6 (state, tmp, state_m, data_m, s_tables);
895 SHIFT12 (state_m, data, tmp);
896 SHIFT16 (state, data_m, state_m);
897 SHIFT61 (state, data_m);
910 state_m[0] = state[0];
911 state_m[1] = state[1];
912 state_m[2] = state[2];
913 state_m[3] = state[3];
914 state_m[4] = state[4];
915 state_m[5] = state[5];
916 state_m[6] = state[6];
917 state_m[7] = state[7];
928 PASS0 (state, tmp, state_m, data_m, s_tables);
929 PASS2 (state, tmp, state_m, data_m, s_tables);
930 PASS4 (state, tmp, state_m, data_m, s_tables);
931 PASS6 (state, tmp, state_m, data_m, s_tables);
933 SHIFT12 (state_m, data, tmp);
934 SHIFT16 (state, data_m, state_m);
935 SHIFT61 (state, data_m);
939 const u32 r0 = state[0];
940 const u32 r1 = state[1];
941 const u32 r2 = state[2];
942 const u32 r3 = state[3];
948 __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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
952 __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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
956 __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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
962 const u32 gid = get_global_id (0);
963 const u32 lid = get_local_id (0);
964 const u32 lsz = get_local_size (0);
970 __local u32 s_tables[4][256];
972 for (u32 i = lid; i < 256; i += lsz)
974 s_tables[0][i] = c_tables[0][i];
975 s_tables[1][i] = c_tables[1][i];
976 s_tables[2][i] = c_tables[2][i];
977 s_tables[3][i] = c_tables[3][i];
980 barrier (CLK_LOCAL_MEM_FENCE);
982 if (gid >= gid_max) return;
990 pw_buf0[0] = pws[gid].i[ 0];
991 pw_buf0[1] = pws[gid].i[ 1];
992 pw_buf0[2] = pws[gid].i[ 2];
993 pw_buf0[3] = pws[gid].i[ 3];
997 pw_buf1[0] = pws[gid].i[ 4];
998 pw_buf1[1] = pws[gid].i[ 5];
999 pw_buf1[2] = pws[gid].i[ 6];
1000 pw_buf1[3] = pws[gid].i[ 7];
1002 const u32 pw_len = pws[gid].pw_len;
1008 const u32 search[4] =
1010 digests_buf[digests_offset].digest_buf[DGST_R0],
1011 digests_buf[digests_offset].digest_buf[DGST_R1],
1012 digests_buf[digests_offset].digest_buf[DGST_R2],
1013 digests_buf[digests_offset].digest_buf[DGST_R3]
1020 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
1050 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
1052 u32 w14 = out_len * 8;
1075 state[ 8] = data[0];
1076 state[ 9] = data[1];
1077 state[10] = data[2];
1078 state[11] = data[3];
1079 state[12] = data[4];
1080 state[13] = data[5];
1081 state[14] = data[6];
1082 state[15] = data[7];
1089 state_m[0] = state[0];
1090 state_m[1] = state[1];
1091 state_m[2] = state[2];
1092 state_m[3] = state[3];
1093 state_m[4] = state[4];
1094 state_m[5] = state[5];
1095 state_m[6] = state[6];
1096 state_m[7] = state[7];
1098 data_m[0] = data[0];
1099 data_m[1] = data[1];
1100 data_m[2] = data[2];
1101 data_m[3] = data[3];
1102 data_m[4] = data[4];
1103 data_m[5] = data[5];
1104 data_m[6] = data[6];
1105 data_m[7] = data[7];
1111 PASS0 (state, tmp, state_m, data_m, s_tables);
1112 PASS2 (state, tmp, state_m, data_m, s_tables);
1113 PASS4 (state, tmp, state_m, data_m, s_tables);
1114 PASS6 (state, tmp, state_m, data_m, s_tables);
1116 SHIFT12 (state_m, data, tmp);
1117 SHIFT16 (state, data_m, state_m);
1118 SHIFT61 (state, data_m);
1132 state_m[0] = state[0];
1133 state_m[1] = state[1];
1134 state_m[2] = state[2];
1135 state_m[3] = state[3];
1136 state_m[4] = state[4];
1137 state_m[5] = state[5];
1138 state_m[6] = state[6];
1139 state_m[7] = state[7];
1141 data_m[0] = data[0];
1142 data_m[1] = data[1];
1143 data_m[2] = data[2];
1144 data_m[3] = data[3];
1145 data_m[4] = data[4];
1146 data_m[5] = data[5];
1147 data_m[6] = data[6];
1148 data_m[7] = data[7];
1150 PASS0 (state, tmp, state_m, data_m, s_tables);
1151 PASS2 (state, tmp, state_m, data_m, s_tables);
1152 PASS4 (state, tmp, state_m, data_m, s_tables);
1153 PASS6 (state, tmp, state_m, data_m, s_tables);
1155 SHIFT12 (state_m, data, tmp);
1156 SHIFT16 (state, data_m, state_m);
1157 SHIFT61 (state, data_m);
1161 data[0] = state[ 8];
1162 data[1] = state[ 9];
1163 data[2] = state[10];
1164 data[3] = state[11];
1165 data[4] = state[12];
1166 data[5] = state[13];
1167 data[6] = state[14];
1168 data[7] = state[15];
1170 state_m[0] = state[0];
1171 state_m[1] = state[1];
1172 state_m[2] = state[2];
1173 state_m[3] = state[3];
1174 state_m[4] = state[4];
1175 state_m[5] = state[5];
1176 state_m[6] = state[6];
1177 state_m[7] = state[7];
1179 data_m[0] = data[0];
1180 data_m[1] = data[1];
1181 data_m[2] = data[2];
1182 data_m[3] = data[3];
1183 data_m[4] = data[4];
1184 data_m[5] = data[5];
1185 data_m[6] = data[6];
1186 data_m[7] = data[7];
1188 PASS0 (state, tmp, state_m, data_m, s_tables);
1189 PASS2 (state, tmp, state_m, data_m, s_tables);
1190 PASS4 (state, tmp, state_m, data_m, s_tables);
1191 PASS6 (state, tmp, state_m, data_m, s_tables);
1193 SHIFT12 (state_m, data, tmp);
1194 SHIFT16 (state, data_m, state_m);
1195 SHIFT61 (state, data_m);
1199 const u32 r0 = state[0];
1200 const u32 r1 = state[1];
1201 const u32 r2 = state[2];
1202 const u32 r3 = state[3];
1208 __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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1212 __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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)