2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
4 * magnum <john.magnum@hushmail.com>
13 #include "include/constants.h"
14 #include "include/kernel_vendor.h"
21 #include "include/kernel_functions.c"
22 #include "OpenCL/types_ocl.c"
23 #include "OpenCL/common.c"
24 #include "include/rp_kernel.h"
25 #include "OpenCL/rp.c"
26 #include "OpenCL/simd.c"
28 __constant u32 c_tables[4][256] =
31 0x00072000, 0x00075000, 0x00074800, 0x00071000,
32 0x00076800, 0x00074000, 0x00070000, 0x00077000,
33 0x00073000, 0x00075800, 0x00070800, 0x00076000,
34 0x00073800, 0x00077800, 0x00072800, 0x00071800,
35 0x0005a000, 0x0005d000, 0x0005c800, 0x00059000,
36 0x0005e800, 0x0005c000, 0x00058000, 0x0005f000,
37 0x0005b000, 0x0005d800, 0x00058800, 0x0005e000,
38 0x0005b800, 0x0005f800, 0x0005a800, 0x00059800,
39 0x00022000, 0x00025000, 0x00024800, 0x00021000,
40 0x00026800, 0x00024000, 0x00020000, 0x00027000,
41 0x00023000, 0x00025800, 0x00020800, 0x00026000,
42 0x00023800, 0x00027800, 0x00022800, 0x00021800,
43 0x00062000, 0x00065000, 0x00064800, 0x00061000,
44 0x00066800, 0x00064000, 0x00060000, 0x00067000,
45 0x00063000, 0x00065800, 0x00060800, 0x00066000,
46 0x00063800, 0x00067800, 0x00062800, 0x00061800,
47 0x00032000, 0x00035000, 0x00034800, 0x00031000,
48 0x00036800, 0x00034000, 0x00030000, 0x00037000,
49 0x00033000, 0x00035800, 0x00030800, 0x00036000,
50 0x00033800, 0x00037800, 0x00032800, 0x00031800,
51 0x0006a000, 0x0006d000, 0x0006c800, 0x00069000,
52 0x0006e800, 0x0006c000, 0x00068000, 0x0006f000,
53 0x0006b000, 0x0006d800, 0x00068800, 0x0006e000,
54 0x0006b800, 0x0006f800, 0x0006a800, 0x00069800,
55 0x0007a000, 0x0007d000, 0x0007c800, 0x00079000,
56 0x0007e800, 0x0007c000, 0x00078000, 0x0007f000,
57 0x0007b000, 0x0007d800, 0x00078800, 0x0007e000,
58 0x0007b800, 0x0007f800, 0x0007a800, 0x00079800,
59 0x00052000, 0x00055000, 0x00054800, 0x00051000,
60 0x00056800, 0x00054000, 0x00050000, 0x00057000,
61 0x00053000, 0x00055800, 0x00050800, 0x00056000,
62 0x00053800, 0x00057800, 0x00052800, 0x00051800,
63 0x00012000, 0x00015000, 0x00014800, 0x00011000,
64 0x00016800, 0x00014000, 0x00010000, 0x00017000,
65 0x00013000, 0x00015800, 0x00010800, 0x00016000,
66 0x00013800, 0x00017800, 0x00012800, 0x00011800,
67 0x0001a000, 0x0001d000, 0x0001c800, 0x00019000,
68 0x0001e800, 0x0001c000, 0x00018000, 0x0001f000,
69 0x0001b000, 0x0001d800, 0x00018800, 0x0001e000,
70 0x0001b800, 0x0001f800, 0x0001a800, 0x00019800,
71 0x00042000, 0x00045000, 0x00044800, 0x00041000,
72 0x00046800, 0x00044000, 0x00040000, 0x00047000,
73 0x00043000, 0x00045800, 0x00040800, 0x00046000,
74 0x00043800, 0x00047800, 0x00042800, 0x00041800,
75 0x0000a000, 0x0000d000, 0x0000c800, 0x00009000,
76 0x0000e800, 0x0000c000, 0x00008000, 0x0000f000,
77 0x0000b000, 0x0000d800, 0x00008800, 0x0000e000,
78 0x0000b800, 0x0000f800, 0x0000a800, 0x00009800,
79 0x00002000, 0x00005000, 0x00004800, 0x00001000,
80 0x00006800, 0x00004000, 0x00000000, 0x00007000,
81 0x00003000, 0x00005800, 0x00000800, 0x00006000,
82 0x00003800, 0x00007800, 0x00002800, 0x00001800,
83 0x0003a000, 0x0003d000, 0x0003c800, 0x00039000,
84 0x0003e800, 0x0003c000, 0x00038000, 0x0003f000,
85 0x0003b000, 0x0003d800, 0x00038800, 0x0003e000,
86 0x0003b800, 0x0003f800, 0x0003a800, 0x00039800,
87 0x0002a000, 0x0002d000, 0x0002c800, 0x00029000,
88 0x0002e800, 0x0002c000, 0x00028000, 0x0002f000,
89 0x0002b000, 0x0002d800, 0x00028800, 0x0002e000,
90 0x0002b800, 0x0002f800, 0x0002a800, 0x00029800,
91 0x0004a000, 0x0004d000, 0x0004c800, 0x00049000,
92 0x0004e800, 0x0004c000, 0x00048000, 0x0004f000,
93 0x0004b000, 0x0004d800, 0x00048800, 0x0004e000,
94 0x0004b800, 0x0004f800, 0x0004a800, 0x00049800,
97 0x03a80000, 0x03c00000, 0x03880000, 0x03e80000,
98 0x03d00000, 0x03980000, 0x03a00000, 0x03900000,
99 0x03f00000, 0x03f80000, 0x03e00000, 0x03b80000,
100 0x03b00000, 0x03800000, 0x03c80000, 0x03d80000,
101 0x06a80000, 0x06c00000, 0x06880000, 0x06e80000,
102 0x06d00000, 0x06980000, 0x06a00000, 0x06900000,
103 0x06f00000, 0x06f80000, 0x06e00000, 0x06b80000,
104 0x06b00000, 0x06800000, 0x06c80000, 0x06d80000,
105 0x05280000, 0x05400000, 0x05080000, 0x05680000,
106 0x05500000, 0x05180000, 0x05200000, 0x05100000,
107 0x05700000, 0x05780000, 0x05600000, 0x05380000,
108 0x05300000, 0x05000000, 0x05480000, 0x05580000,
109 0x00a80000, 0x00c00000, 0x00880000, 0x00e80000,
110 0x00d00000, 0x00980000, 0x00a00000, 0x00900000,
111 0x00f00000, 0x00f80000, 0x00e00000, 0x00b80000,
112 0x00b00000, 0x00800000, 0x00c80000, 0x00d80000,
113 0x00280000, 0x00400000, 0x00080000, 0x00680000,
114 0x00500000, 0x00180000, 0x00200000, 0x00100000,
115 0x00700000, 0x00780000, 0x00600000, 0x00380000,
116 0x00300000, 0x00000000, 0x00480000, 0x00580000,
117 0x04280000, 0x04400000, 0x04080000, 0x04680000,
118 0x04500000, 0x04180000, 0x04200000, 0x04100000,
119 0x04700000, 0x04780000, 0x04600000, 0x04380000,
120 0x04300000, 0x04000000, 0x04480000, 0x04580000,
121 0x04a80000, 0x04c00000, 0x04880000, 0x04e80000,
122 0x04d00000, 0x04980000, 0x04a00000, 0x04900000,
123 0x04f00000, 0x04f80000, 0x04e00000, 0x04b80000,
124 0x04b00000, 0x04800000, 0x04c80000, 0x04d80000,
125 0x07a80000, 0x07c00000, 0x07880000, 0x07e80000,
126 0x07d00000, 0x07980000, 0x07a00000, 0x07900000,
127 0x07f00000, 0x07f80000, 0x07e00000, 0x07b80000,
128 0x07b00000, 0x07800000, 0x07c80000, 0x07d80000,
129 0x07280000, 0x07400000, 0x07080000, 0x07680000,
130 0x07500000, 0x07180000, 0x07200000, 0x07100000,
131 0x07700000, 0x07780000, 0x07600000, 0x07380000,
132 0x07300000, 0x07000000, 0x07480000, 0x07580000,
133 0x02280000, 0x02400000, 0x02080000, 0x02680000,
134 0x02500000, 0x02180000, 0x02200000, 0x02100000,
135 0x02700000, 0x02780000, 0x02600000, 0x02380000,
136 0x02300000, 0x02000000, 0x02480000, 0x02580000,
137 0x03280000, 0x03400000, 0x03080000, 0x03680000,
138 0x03500000, 0x03180000, 0x03200000, 0x03100000,
139 0x03700000, 0x03780000, 0x03600000, 0x03380000,
140 0x03300000, 0x03000000, 0x03480000, 0x03580000,
141 0x06280000, 0x06400000, 0x06080000, 0x06680000,
142 0x06500000, 0x06180000, 0x06200000, 0x06100000,
143 0x06700000, 0x06780000, 0x06600000, 0x06380000,
144 0x06300000, 0x06000000, 0x06480000, 0x06580000,
145 0x05a80000, 0x05c00000, 0x05880000, 0x05e80000,
146 0x05d00000, 0x05980000, 0x05a00000, 0x05900000,
147 0x05f00000, 0x05f80000, 0x05e00000, 0x05b80000,
148 0x05b00000, 0x05800000, 0x05c80000, 0x05d80000,
149 0x01280000, 0x01400000, 0x01080000, 0x01680000,
150 0x01500000, 0x01180000, 0x01200000, 0x01100000,
151 0x01700000, 0x01780000, 0x01600000, 0x01380000,
152 0x01300000, 0x01000000, 0x01480000, 0x01580000,
153 0x02a80000, 0x02c00000, 0x02880000, 0x02e80000,
154 0x02d00000, 0x02980000, 0x02a00000, 0x02900000,
155 0x02f00000, 0x02f80000, 0x02e00000, 0x02b80000,
156 0x02b00000, 0x02800000, 0x02c80000, 0x02d80000,
157 0x01a80000, 0x01c00000, 0x01880000, 0x01e80000,
158 0x01d00000, 0x01980000, 0x01a00000, 0x01900000,
159 0x01f00000, 0x01f80000, 0x01e00000, 0x01b80000,
160 0x01b00000, 0x01800000, 0x01c80000, 0x01d80000,
163 0x30000002, 0x60000002, 0x38000002, 0x08000002,
164 0x28000002, 0x78000002, 0x68000002, 0x40000002,
165 0x20000002, 0x50000002, 0x48000002, 0x70000002,
166 0x00000002, 0x18000002, 0x58000002, 0x10000002,
167 0xb0000005, 0xe0000005, 0xb8000005, 0x88000005,
168 0xa8000005, 0xf8000005, 0xe8000005, 0xc0000005,
169 0xa0000005, 0xd0000005, 0xc8000005, 0xf0000005,
170 0x80000005, 0x98000005, 0xd8000005, 0x90000005,
171 0x30000005, 0x60000005, 0x38000005, 0x08000005,
172 0x28000005, 0x78000005, 0x68000005, 0x40000005,
173 0x20000005, 0x50000005, 0x48000005, 0x70000005,
174 0x00000005, 0x18000005, 0x58000005, 0x10000005,
175 0x30000000, 0x60000000, 0x38000000, 0x08000000,
176 0x28000000, 0x78000000, 0x68000000, 0x40000000,
177 0x20000000, 0x50000000, 0x48000000, 0x70000000,
178 0x00000000, 0x18000000, 0x58000000, 0x10000000,
179 0xb0000003, 0xe0000003, 0xb8000003, 0x88000003,
180 0xa8000003, 0xf8000003, 0xe8000003, 0xc0000003,
181 0xa0000003, 0xd0000003, 0xc8000003, 0xf0000003,
182 0x80000003, 0x98000003, 0xd8000003, 0x90000003,
183 0x30000001, 0x60000001, 0x38000001, 0x08000001,
184 0x28000001, 0x78000001, 0x68000001, 0x40000001,
185 0x20000001, 0x50000001, 0x48000001, 0x70000001,
186 0x00000001, 0x18000001, 0x58000001, 0x10000001,
187 0xb0000000, 0xe0000000, 0xb8000000, 0x88000000,
188 0xa8000000, 0xf8000000, 0xe8000000, 0xc0000000,
189 0xa0000000, 0xd0000000, 0xc8000000, 0xf0000000,
190 0x80000000, 0x98000000, 0xd8000000, 0x90000000,
191 0xb0000006, 0xe0000006, 0xb8000006, 0x88000006,
192 0xa8000006, 0xf8000006, 0xe8000006, 0xc0000006,
193 0xa0000006, 0xd0000006, 0xc8000006, 0xf0000006,
194 0x80000006, 0x98000006, 0xd8000006, 0x90000006,
195 0xb0000001, 0xe0000001, 0xb8000001, 0x88000001,
196 0xa8000001, 0xf8000001, 0xe8000001, 0xc0000001,
197 0xa0000001, 0xd0000001, 0xc8000001, 0xf0000001,
198 0x80000001, 0x98000001, 0xd8000001, 0x90000001,
199 0x30000003, 0x60000003, 0x38000003, 0x08000003,
200 0x28000003, 0x78000003, 0x68000003, 0x40000003,
201 0x20000003, 0x50000003, 0x48000003, 0x70000003,
202 0x00000003, 0x18000003, 0x58000003, 0x10000003,
203 0x30000004, 0x60000004, 0x38000004, 0x08000004,
204 0x28000004, 0x78000004, 0x68000004, 0x40000004,
205 0x20000004, 0x50000004, 0x48000004, 0x70000004,
206 0x00000004, 0x18000004, 0x58000004, 0x10000004,
207 0xb0000002, 0xe0000002, 0xb8000002, 0x88000002,
208 0xa8000002, 0xf8000002, 0xe8000002, 0xc0000002,
209 0xa0000002, 0xd0000002, 0xc8000002, 0xf0000002,
210 0x80000002, 0x98000002, 0xd8000002, 0x90000002,
211 0xb0000004, 0xe0000004, 0xb8000004, 0x88000004,
212 0xa8000004, 0xf8000004, 0xe8000004, 0xc0000004,
213 0xa0000004, 0xd0000004, 0xc8000004, 0xf0000004,
214 0x80000004, 0x98000004, 0xd8000004, 0x90000004,
215 0x30000006, 0x60000006, 0x38000006, 0x08000006,
216 0x28000006, 0x78000006, 0x68000006, 0x40000006,
217 0x20000006, 0x50000006, 0x48000006, 0x70000006,
218 0x00000006, 0x18000006, 0x58000006, 0x10000006,
219 0xb0000007, 0xe0000007, 0xb8000007, 0x88000007,
220 0xa8000007, 0xf8000007, 0xe8000007, 0xc0000007,
221 0xa0000007, 0xd0000007, 0xc8000007, 0xf0000007,
222 0x80000007, 0x98000007, 0xd8000007, 0x90000007,
223 0x30000007, 0x60000007, 0x38000007, 0x08000007,
224 0x28000007, 0x78000007, 0x68000007, 0x40000007,
225 0x20000007, 0x50000007, 0x48000007, 0x70000007,
226 0x00000007, 0x18000007, 0x58000007, 0x10000007,
229 0x000000e8, 0x000000d8, 0x000000a0, 0x00000088,
230 0x00000098, 0x000000f8, 0x000000a8, 0x000000c8,
231 0x00000080, 0x000000d0, 0x000000f0, 0x000000b8,
232 0x000000b0, 0x000000c0, 0x00000090, 0x000000e0,
233 0x000007e8, 0x000007d8, 0x000007a0, 0x00000788,
234 0x00000798, 0x000007f8, 0x000007a8, 0x000007c8,
235 0x00000780, 0x000007d0, 0x000007f0, 0x000007b8,
236 0x000007b0, 0x000007c0, 0x00000790, 0x000007e0,
237 0x000006e8, 0x000006d8, 0x000006a0, 0x00000688,
238 0x00000698, 0x000006f8, 0x000006a8, 0x000006c8,
239 0x00000680, 0x000006d0, 0x000006f0, 0x000006b8,
240 0x000006b0, 0x000006c0, 0x00000690, 0x000006e0,
241 0x00000068, 0x00000058, 0x00000020, 0x00000008,
242 0x00000018, 0x00000078, 0x00000028, 0x00000048,
243 0x00000000, 0x00000050, 0x00000070, 0x00000038,
244 0x00000030, 0x00000040, 0x00000010, 0x00000060,
245 0x000002e8, 0x000002d8, 0x000002a0, 0x00000288,
246 0x00000298, 0x000002f8, 0x000002a8, 0x000002c8,
247 0x00000280, 0x000002d0, 0x000002f0, 0x000002b8,
248 0x000002b0, 0x000002c0, 0x00000290, 0x000002e0,
249 0x000003e8, 0x000003d8, 0x000003a0, 0x00000388,
250 0x00000398, 0x000003f8, 0x000003a8, 0x000003c8,
251 0x00000380, 0x000003d0, 0x000003f0, 0x000003b8,
252 0x000003b0, 0x000003c0, 0x00000390, 0x000003e0,
253 0x00000568, 0x00000558, 0x00000520, 0x00000508,
254 0x00000518, 0x00000578, 0x00000528, 0x00000548,
255 0x00000500, 0x00000550, 0x00000570, 0x00000538,
256 0x00000530, 0x00000540, 0x00000510, 0x00000560,
257 0x00000268, 0x00000258, 0x00000220, 0x00000208,
258 0x00000218, 0x00000278, 0x00000228, 0x00000248,
259 0x00000200, 0x00000250, 0x00000270, 0x00000238,
260 0x00000230, 0x00000240, 0x00000210, 0x00000260,
261 0x000004e8, 0x000004d8, 0x000004a0, 0x00000488,
262 0x00000498, 0x000004f8, 0x000004a8, 0x000004c8,
263 0x00000480, 0x000004d0, 0x000004f0, 0x000004b8,
264 0x000004b0, 0x000004c0, 0x00000490, 0x000004e0,
265 0x00000168, 0x00000158, 0x00000120, 0x00000108,
266 0x00000118, 0x00000178, 0x00000128, 0x00000148,
267 0x00000100, 0x00000150, 0x00000170, 0x00000138,
268 0x00000130, 0x00000140, 0x00000110, 0x00000160,
269 0x000001e8, 0x000001d8, 0x000001a0, 0x00000188,
270 0x00000198, 0x000001f8, 0x000001a8, 0x000001c8,
271 0x00000180, 0x000001d0, 0x000001f0, 0x000001b8,
272 0x000001b0, 0x000001c0, 0x00000190, 0x000001e0,
273 0x00000768, 0x00000758, 0x00000720, 0x00000708,
274 0x00000718, 0x00000778, 0x00000728, 0x00000748,
275 0x00000700, 0x00000750, 0x00000770, 0x00000738,
276 0x00000730, 0x00000740, 0x00000710, 0x00000760,
277 0x00000368, 0x00000358, 0x00000320, 0x00000308,
278 0x00000318, 0x00000378, 0x00000328, 0x00000348,
279 0x00000300, 0x00000350, 0x00000370, 0x00000338,
280 0x00000330, 0x00000340, 0x00000310, 0x00000360,
281 0x000005e8, 0x000005d8, 0x000005a0, 0x00000588,
282 0x00000598, 0x000005f8, 0x000005a8, 0x000005c8,
283 0x00000580, 0x000005d0, 0x000005f0, 0x000005b8,
284 0x000005b0, 0x000005c0, 0x00000590, 0x000005e0,
285 0x00000468, 0x00000458, 0x00000420, 0x00000408,
286 0x00000418, 0x00000478, 0x00000428, 0x00000448,
287 0x00000400, 0x00000450, 0x00000470, 0x00000438,
288 0x00000430, 0x00000440, 0x00000410, 0x00000460,
289 0x00000668, 0x00000658, 0x00000620, 0x00000608,
290 0x00000618, 0x00000678, 0x00000628, 0x00000648,
291 0x00000600, 0x00000650, 0x00000670, 0x00000638,
292 0x00000630, 0x00000640, 0x00000610, 0x00000660,
297 #define BOX(i,n,S) (S)[(n)][(i)]
299 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
301 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
303 #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])
304 #elif VECT_SIZE == 16
305 #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])
308 #define _round(k1,k2,tbl) \
312 l ^= 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); \
317 r ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
318 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
319 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
320 BOX (((t >> 24) & 0xff), 3, tbl); \
323 #define R(k,h,s,i,t) \
329 _round (k[0], k[1], t); \
330 _round (k[2], k[3], t); \
331 _round (k[4], k[5], t); \
332 _round (k[6], k[7], t); \
333 _round (k[0], k[1], t); \
334 _round (k[2], k[3], t); \
335 _round (k[4], k[5], t); \
336 _round (k[6], k[7], t); \
337 _round (k[0], k[1], t); \
338 _round (k[2], k[3], t); \
339 _round (k[4], k[5], t); \
340 _round (k[6], k[7], t); \
341 _round (k[7], k[6], t); \
342 _round (k[5], k[4], t); \
343 _round (k[3], k[2], t); \
344 _round (k[1], k[0], t); \
350 w[0] = u[0] ^ v[0]; \
351 w[1] = u[1] ^ v[1]; \
352 w[2] = u[2] ^ v[2]; \
353 w[3] = u[3] ^ v[3]; \
354 w[4] = u[4] ^ v[4]; \
355 w[5] = u[5] ^ v[5]; \
356 w[6] = u[6] ^ v[6]; \
360 k[0] = ((w[0] & 0x000000ff) << 0) \
361 | ((w[2] & 0x000000ff) << 8) \
362 | ((w[4] & 0x000000ff) << 16) \
363 | ((w[6] & 0x000000ff) << 24); \
364 k[1] = ((w[0] & 0x0000ff00) >> 8) \
365 | ((w[2] & 0x0000ff00) >> 0) \
366 | ((w[4] & 0x0000ff00) << 8) \
367 | ((w[6] & 0x0000ff00) << 16); \
368 k[2] = ((w[0] & 0x00ff0000) >> 16) \
369 | ((w[2] & 0x00ff0000) >> 8) \
370 | ((w[4] & 0x00ff0000) << 0) \
371 | ((w[6] & 0x00ff0000) << 8); \
372 k[3] = ((w[0] & 0xff000000) >> 24) \
373 | ((w[2] & 0xff000000) >> 16) \
374 | ((w[4] & 0xff000000) >> 8) \
375 | ((w[6] & 0xff000000) >> 0); \
376 k[4] = ((w[1] & 0x000000ff) << 0) \
377 | ((w[3] & 0x000000ff) << 8) \
378 | ((w[5] & 0x000000ff) << 16) \
379 | ((w[7] & 0x000000ff) << 24); \
380 k[5] = ((w[1] & 0x0000ff00) >> 8) \
381 | ((w[3] & 0x0000ff00) >> 0) \
382 | ((w[5] & 0x0000ff00) << 8) \
383 | ((w[7] & 0x0000ff00) << 16); \
384 k[6] = ((w[1] & 0x00ff0000) >> 16) \
385 | ((w[3] & 0x00ff0000) >> 8) \
386 | ((w[5] & 0x00ff0000) << 0) \
387 | ((w[7] & 0x00ff0000) << 8); \
388 k[7] = ((w[1] & 0xff000000) >> 24) \
389 | ((w[3] & 0xff000000) >> 16) \
390 | ((w[5] & 0xff000000) >> 8) \
391 | ((w[7] & 0xff000000) >> 0);
428 x[0] ^= 0xff00ff00; \
429 x[1] ^= 0xff00ff00; \
430 x[2] ^= 0x00ff00ff; \
431 x[3] ^= 0x00ff00ff; \
432 x[4] ^= 0x00ffff00; \
433 x[5] ^= 0xff0000ff; \
434 x[6] ^= 0x000000ff; \
437 #define SHIFT12(u,m,s) \
438 u[0] = m[0] ^ s[6]; \
439 u[1] = m[1] ^ s[7]; \
440 u[2] = m[2] ^ (s[0] << 16) \
442 ^ (s[0] & 0x0000ffff) \
443 ^ (s[1] & 0x0000ffff) \
448 ^ (s[7] & 0xffff0000) \
450 u[3] = m[3] ^ (s[0] & 0x0000ffff) \
452 ^ (s[1] & 0x0000ffff) \
461 ^ (s[7] & 0x0000ffff) \
464 u[4] = m[4] ^ (s[0] & 0xffff0000) \
467 ^ (s[1] & 0xffff0000) \
476 ^ (s[7] & 0x0000ffff) \
479 u[5] = m[5] ^ (s[0] << 16) \
481 ^ (s[0] & 0xffff0000) \
482 ^ (s[1] & 0x0000ffff) \
492 ^ (s[7] & 0xffff0000) \
508 u[7] = m[7] ^ (s[0] & 0xffff0000) \
510 ^ (s[1] & 0x0000ffff) \
519 ^ (s[7] & 0x0000ffff) \
523 #define SHIFT16(h,v,u) \
524 v[0] = h[0] ^ (u[1] << 16) \
526 v[1] = h[1] ^ (u[2] << 16) \
528 v[2] = h[2] ^ (u[3] << 16) \
530 v[3] = h[3] ^ (u[4] << 16) \
532 v[4] = h[4] ^ (u[5] << 16) \
534 v[5] = h[5] ^ (u[6] << 16) \
536 v[6] = h[6] ^ (u[7] << 16) \
538 v[7] = h[7] ^ (u[0] & 0xffff0000) \
541 ^ (u[1] & 0xffff0000) \
544 ^ (u[7] & 0xffff0000);
546 #define SHIFT61(h,v) \
547 h[0] = (v[0] & 0xffff0000) \
551 ^ (v[1] & 0xffff0000) \
560 ^ (v[7] & 0x0000ffff); \
561 h[1] = (v[0] << 16) \
563 ^ (v[0] & 0xffff0000) \
564 ^ (v[1] & 0x0000ffff) \
572 ^ (v[7] & 0xffff0000) \
574 h[2] = (v[0] & 0x0000ffff) \
578 ^ (v[1] & 0xffff0000) \
586 ^ (v[7] & 0x0000ffff) \
589 h[3] = (v[0] << 16) \
591 ^ (v[0] & 0xffff0000) \
592 ^ (v[1] & 0xffff0000) \
602 ^ (v[7] & 0x0000ffff) \
604 h[4] = (v[0] >> 16) \
618 h[5] = (v[0] << 16) \
619 ^ (v[0] & 0xffff0000) \
622 ^ (v[1] & 0xffff0000) \
636 ^ (v[7] & 0xffff0000); \
668 #define PASS0(h,s,u,v,t) \
679 #define PASS2(h,s,u,v,t) \
691 #define PASS4(h,s,u,v,t) \
702 #define PASS6(h,s,u,v,t) \
711 __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)
717 const u32 gid = get_global_id (0);
718 const u32 lid = get_local_id (0);
719 const u32 lsz = get_local_size (0);
725 __local u32 s_tables[4][256];
727 for (u32 i = lid; i < 256; i += lsz)
729 s_tables[0][i] = c_tables[0][i];
730 s_tables[1][i] = c_tables[1][i];
731 s_tables[2][i] = c_tables[2][i];
732 s_tables[3][i] = c_tables[3][i];
735 barrier (CLK_LOCAL_MEM_FENCE);
737 if (gid >= gid_max) return;
745 pw_buf0[0] = pws[gid].i[ 0];
746 pw_buf0[1] = pws[gid].i[ 1];
747 pw_buf0[2] = pws[gid].i[ 2];
748 pw_buf0[3] = pws[gid].i[ 3];
752 pw_buf1[0] = pws[gid].i[ 4];
753 pw_buf1[1] = pws[gid].i[ 5];
754 pw_buf1[2] = pws[gid].i[ 6];
755 pw_buf1[3] = pws[gid].i[ 7];
757 const u32 pw_len = pws[gid].pw_len;
763 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
770 const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
772 u32x w14 = out_len * 8;
809 state_m[0] = state[0];
810 state_m[1] = state[1];
811 state_m[2] = state[2];
812 state_m[3] = state[3];
813 state_m[4] = state[4];
814 state_m[5] = state[5];
815 state_m[6] = state[6];
816 state_m[7] = state[7];
831 PASS0 (state, tmp, state_m, data_m, s_tables);
832 PASS2 (state, tmp, state_m, data_m, s_tables);
833 PASS4 (state, tmp, state_m, data_m, s_tables);
834 PASS6 (state, tmp, state_m, data_m, s_tables);
836 SHIFT12 (state_m, data, tmp);
837 SHIFT16 (state, data_m, state_m);
838 SHIFT61 (state, data_m);
852 state_m[0] = state[0];
853 state_m[1] = state[1];
854 state_m[2] = state[2];
855 state_m[3] = state[3];
856 state_m[4] = state[4];
857 state_m[5] = state[5];
858 state_m[6] = state[6];
859 state_m[7] = state[7];
870 PASS0 (state, tmp, state_m, data_m, s_tables);
871 PASS2 (state, tmp, state_m, data_m, s_tables);
872 PASS4 (state, tmp, state_m, data_m, s_tables);
873 PASS6 (state, tmp, state_m, data_m, s_tables);
875 SHIFT12 (state_m, data, tmp);
876 SHIFT16 (state, data_m, state_m);
877 SHIFT61 (state, data_m);
890 state_m[0] = state[0];
891 state_m[1] = state[1];
892 state_m[2] = state[2];
893 state_m[3] = state[3];
894 state_m[4] = state[4];
895 state_m[5] = state[5];
896 state_m[6] = state[6];
897 state_m[7] = state[7];
908 PASS0 (state, tmp, state_m, data_m, s_tables);
909 PASS2 (state, tmp, state_m, data_m, s_tables);
910 PASS4 (state, tmp, state_m, data_m, s_tables);
911 PASS6 (state, tmp, state_m, data_m, s_tables);
913 SHIFT12 (state_m, data, tmp);
914 SHIFT16 (state, data_m, state_m);
915 SHIFT61 (state, data_m);
919 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
923 __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)
927 __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)
931 __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)
937 const u32 gid = get_global_id (0);
938 const u32 lid = get_local_id (0);
939 const u32 lsz = get_local_size (0);
945 __local u32 s_tables[4][256];
947 for (u32 i = lid; i < 256; i += lsz)
949 s_tables[0][i] = c_tables[0][i];
950 s_tables[1][i] = c_tables[1][i];
951 s_tables[2][i] = c_tables[2][i];
952 s_tables[3][i] = c_tables[3][i];
955 barrier (CLK_LOCAL_MEM_FENCE);
957 if (gid >= gid_max) return;
965 pw_buf0[0] = pws[gid].i[ 0];
966 pw_buf0[1] = pws[gid].i[ 1];
967 pw_buf0[2] = pws[gid].i[ 2];
968 pw_buf0[3] = pws[gid].i[ 3];
972 pw_buf1[0] = pws[gid].i[ 4];
973 pw_buf1[1] = pws[gid].i[ 5];
974 pw_buf1[2] = pws[gid].i[ 6];
975 pw_buf1[3] = pws[gid].i[ 7];
977 const u32 pw_len = pws[gid].pw_len;
983 const u32 search[4] =
985 digests_buf[digests_offset].digest_buf[DGST_R0],
986 digests_buf[digests_offset].digest_buf[DGST_R1],
987 digests_buf[digests_offset].digest_buf[DGST_R2],
988 digests_buf[digests_offset].digest_buf[DGST_R3]
995 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
1002 const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
1004 u32x w14 = out_len * 8;
1027 state[ 8] = data[0];
1028 state[ 9] = data[1];
1029 state[10] = data[2];
1030 state[11] = data[3];
1031 state[12] = data[4];
1032 state[13] = data[5];
1033 state[14] = data[6];
1034 state[15] = data[7];
1041 state_m[0] = state[0];
1042 state_m[1] = state[1];
1043 state_m[2] = state[2];
1044 state_m[3] = state[3];
1045 state_m[4] = state[4];
1046 state_m[5] = state[5];
1047 state_m[6] = state[6];
1048 state_m[7] = state[7];
1050 data_m[0] = data[0];
1051 data_m[1] = data[1];
1052 data_m[2] = data[2];
1053 data_m[3] = data[3];
1054 data_m[4] = data[4];
1055 data_m[5] = data[5];
1056 data_m[6] = data[6];
1057 data_m[7] = data[7];
1063 PASS0 (state, tmp, state_m, data_m, s_tables);
1064 PASS2 (state, tmp, state_m, data_m, s_tables);
1065 PASS4 (state, tmp, state_m, data_m, s_tables);
1066 PASS6 (state, tmp, state_m, data_m, s_tables);
1068 SHIFT12 (state_m, data, tmp);
1069 SHIFT16 (state, data_m, state_m);
1070 SHIFT61 (state, data_m);
1084 state_m[0] = state[0];
1085 state_m[1] = state[1];
1086 state_m[2] = state[2];
1087 state_m[3] = state[3];
1088 state_m[4] = state[4];
1089 state_m[5] = state[5];
1090 state_m[6] = state[6];
1091 state_m[7] = state[7];
1093 data_m[0] = data[0];
1094 data_m[1] = data[1];
1095 data_m[2] = data[2];
1096 data_m[3] = data[3];
1097 data_m[4] = data[4];
1098 data_m[5] = data[5];
1099 data_m[6] = data[6];
1100 data_m[7] = data[7];
1102 PASS0 (state, tmp, state_m, data_m, s_tables);
1103 PASS2 (state, tmp, state_m, data_m, s_tables);
1104 PASS4 (state, tmp, state_m, data_m, s_tables);
1105 PASS6 (state, tmp, state_m, data_m, s_tables);
1107 SHIFT12 (state_m, data, tmp);
1108 SHIFT16 (state, data_m, state_m);
1109 SHIFT61 (state, data_m);
1113 data[0] = state[ 8];
1114 data[1] = state[ 9];
1115 data[2] = state[10];
1116 data[3] = state[11];
1117 data[4] = state[12];
1118 data[5] = state[13];
1119 data[6] = state[14];
1120 data[7] = state[15];
1122 state_m[0] = state[0];
1123 state_m[1] = state[1];
1124 state_m[2] = state[2];
1125 state_m[3] = state[3];
1126 state_m[4] = state[4];
1127 state_m[5] = state[5];
1128 state_m[6] = state[6];
1129 state_m[7] = state[7];
1131 data_m[0] = data[0];
1132 data_m[1] = data[1];
1133 data_m[2] = data[2];
1134 data_m[3] = data[3];
1135 data_m[4] = data[4];
1136 data_m[5] = data[5];
1137 data_m[6] = data[6];
1138 data_m[7] = data[7];
1140 PASS0 (state, tmp, state_m, data_m, s_tables);
1141 PASS2 (state, tmp, state_m, data_m, s_tables);
1142 PASS4 (state, tmp, state_m, data_m, s_tables);
1143 PASS6 (state, tmp, state_m, data_m, s_tables);
1145 SHIFT12 (state_m, data, tmp);
1146 SHIFT16 (state, data_m, state_m);
1147 SHIFT61 (state, data_m);
1151 COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
1155 __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)
1159 __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)