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 "include/rp_kernel.h"
24 #include "OpenCL/rp.c"
25 #include "OpenCL/simd.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 += VECT_SIZE)
767 const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
769 u32 w14 = out_len * 8;
806 state_m[0] = state[0];
807 state_m[1] = state[1];
808 state_m[2] = state[2];
809 state_m[3] = state[3];
810 state_m[4] = state[4];
811 state_m[5] = state[5];
812 state_m[6] = state[6];
813 state_m[7] = state[7];
828 PASS0 (state, tmp, state_m, data_m, s_tables);
829 PASS2 (state, tmp, state_m, data_m, s_tables);
830 PASS4 (state, tmp, state_m, data_m, s_tables);
831 PASS6 (state, tmp, state_m, data_m, s_tables);
833 SHIFT12 (state_m, data, tmp);
834 SHIFT16 (state, data_m, state_m);
835 SHIFT61 (state, data_m);
849 state_m[0] = state[0];
850 state_m[1] = state[1];
851 state_m[2] = state[2];
852 state_m[3] = state[3];
853 state_m[4] = state[4];
854 state_m[5] = state[5];
855 state_m[6] = state[6];
856 state_m[7] = state[7];
867 PASS0 (state, tmp, state_m, data_m, s_tables);
868 PASS2 (state, tmp, state_m, data_m, s_tables);
869 PASS4 (state, tmp, state_m, data_m, s_tables);
870 PASS6 (state, tmp, state_m, data_m, s_tables);
872 SHIFT12 (state_m, data, tmp);
873 SHIFT16 (state, data_m, state_m);
874 SHIFT61 (state, data_m);
887 state_m[0] = state[0];
888 state_m[1] = state[1];
889 state_m[2] = state[2];
890 state_m[3] = state[3];
891 state_m[4] = state[4];
892 state_m[5] = state[5];
893 state_m[6] = state[6];
894 state_m[7] = state[7];
905 PASS0 (state, tmp, state_m, data_m, s_tables);
906 PASS2 (state, tmp, state_m, data_m, s_tables);
907 PASS4 (state, tmp, state_m, data_m, s_tables);
908 PASS6 (state, tmp, state_m, data_m, s_tables);
910 SHIFT12 (state_m, data, tmp);
911 SHIFT16 (state, data_m, state_m);
912 SHIFT61 (state, data_m);
916 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
920 __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)
924 __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)
928 __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)
934 const u32 gid = get_global_id (0);
935 const u32 lid = get_local_id (0);
936 const u32 lsz = get_local_size (0);
942 __local u32 s_tables[4][256];
944 for (u32 i = lid; i < 256; i += lsz)
946 s_tables[0][i] = c_tables[0][i];
947 s_tables[1][i] = c_tables[1][i];
948 s_tables[2][i] = c_tables[2][i];
949 s_tables[3][i] = c_tables[3][i];
952 barrier (CLK_LOCAL_MEM_FENCE);
954 if (gid >= gid_max) return;
962 pw_buf0[0] = pws[gid].i[ 0];
963 pw_buf0[1] = pws[gid].i[ 1];
964 pw_buf0[2] = pws[gid].i[ 2];
965 pw_buf0[3] = pws[gid].i[ 3];
969 pw_buf1[0] = pws[gid].i[ 4];
970 pw_buf1[1] = pws[gid].i[ 5];
971 pw_buf1[2] = pws[gid].i[ 6];
972 pw_buf1[3] = pws[gid].i[ 7];
974 const u32 pw_len = pws[gid].pw_len;
980 const u32 search[4] =
982 digests_buf[digests_offset].digest_buf[DGST_R0],
983 digests_buf[digests_offset].digest_buf[DGST_R1],
984 digests_buf[digests_offset].digest_buf[DGST_R2],
985 digests_buf[digests_offset].digest_buf[DGST_R3]
992 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos += VECT_SIZE)
999 const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
1001 u32 w14 = out_len * 8;
1024 state[ 8] = data[0];
1025 state[ 9] = data[1];
1026 state[10] = data[2];
1027 state[11] = data[3];
1028 state[12] = data[4];
1029 state[13] = data[5];
1030 state[14] = data[6];
1031 state[15] = data[7];
1038 state_m[0] = state[0];
1039 state_m[1] = state[1];
1040 state_m[2] = state[2];
1041 state_m[3] = state[3];
1042 state_m[4] = state[4];
1043 state_m[5] = state[5];
1044 state_m[6] = state[6];
1045 state_m[7] = state[7];
1047 data_m[0] = data[0];
1048 data_m[1] = data[1];
1049 data_m[2] = data[2];
1050 data_m[3] = data[3];
1051 data_m[4] = data[4];
1052 data_m[5] = data[5];
1053 data_m[6] = data[6];
1054 data_m[7] = data[7];
1060 PASS0 (state, tmp, state_m, data_m, s_tables);
1061 PASS2 (state, tmp, state_m, data_m, s_tables);
1062 PASS4 (state, tmp, state_m, data_m, s_tables);
1063 PASS6 (state, tmp, state_m, data_m, s_tables);
1065 SHIFT12 (state_m, data, tmp);
1066 SHIFT16 (state, data_m, state_m);
1067 SHIFT61 (state, data_m);
1081 state_m[0] = state[0];
1082 state_m[1] = state[1];
1083 state_m[2] = state[2];
1084 state_m[3] = state[3];
1085 state_m[4] = state[4];
1086 state_m[5] = state[5];
1087 state_m[6] = state[6];
1088 state_m[7] = state[7];
1090 data_m[0] = data[0];
1091 data_m[1] = data[1];
1092 data_m[2] = data[2];
1093 data_m[3] = data[3];
1094 data_m[4] = data[4];
1095 data_m[5] = data[5];
1096 data_m[6] = data[6];
1097 data_m[7] = data[7];
1099 PASS0 (state, tmp, state_m, data_m, s_tables);
1100 PASS2 (state, tmp, state_m, data_m, s_tables);
1101 PASS4 (state, tmp, state_m, data_m, s_tables);
1102 PASS6 (state, tmp, state_m, data_m, s_tables);
1104 SHIFT12 (state_m, data, tmp);
1105 SHIFT16 (state, data_m, state_m);
1106 SHIFT61 (state, data_m);
1110 data[0] = state[ 8];
1111 data[1] = state[ 9];
1112 data[2] = state[10];
1113 data[3] = state[11];
1114 data[4] = state[12];
1115 data[5] = state[13];
1116 data[6] = state[14];
1117 data[7] = state[15];
1119 state_m[0] = state[0];
1120 state_m[1] = state[1];
1121 state_m[2] = state[2];
1122 state_m[3] = state[3];
1123 state_m[4] = state[4];
1124 state_m[5] = state[5];
1125 state_m[6] = state[6];
1126 state_m[7] = state[7];
1128 data_m[0] = data[0];
1129 data_m[1] = data[1];
1130 data_m[2] = data[2];
1131 data_m[3] = data[3];
1132 data_m[4] = data[4];
1133 data_m[5] = data[5];
1134 data_m[6] = data[6];
1135 data_m[7] = data[7];
1137 PASS0 (state, tmp, state_m, data_m, s_tables);
1138 PASS2 (state, tmp, state_m, data_m, s_tables);
1139 PASS4 (state, tmp, state_m, data_m, s_tables);
1140 PASS6 (state, tmp, state_m, data_m, s_tables);
1142 SHIFT12 (state_m, data, tmp);
1143 SHIFT16 (state, data_m, state_m);
1144 SHIFT61 (state, data_m);
1148 COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
1152 __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)
1156 __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)