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,
295 #define BOX(i,n,S) (S)[(n)][(i)]
297 #define _round(k1,k2,tbl) \
301 l ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
302 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
303 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
304 BOX (((t >> 24) & 0xff), 3, tbl); \
306 r ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
307 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
308 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
309 BOX (((t >> 24) & 0xff), 3, tbl); \
312 #define R(k,h,s,i,t) \
318 _round (k[0], k[1], t); \
319 _round (k[2], k[3], t); \
320 _round (k[4], k[5], t); \
321 _round (k[6], k[7], t); \
322 _round (k[0], k[1], t); \
323 _round (k[2], k[3], t); \
324 _round (k[4], k[5], t); \
325 _round (k[6], k[7], 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[7], k[6], t); \
331 _round (k[5], k[4], t); \
332 _round (k[3], k[2], t); \
333 _round (k[1], k[0], t); \
339 w[0] = u[0] ^ v[0]; \
340 w[1] = u[1] ^ v[1]; \
341 w[2] = u[2] ^ v[2]; \
342 w[3] = u[3] ^ v[3]; \
343 w[4] = u[4] ^ v[4]; \
344 w[5] = u[5] ^ v[5]; \
345 w[6] = u[6] ^ v[6]; \
349 k[0] = ((w[0] & 0x000000ff) << 0) \
350 | ((w[2] & 0x000000ff) << 8) \
351 | ((w[4] & 0x000000ff) << 16) \
352 | ((w[6] & 0x000000ff) << 24); \
353 k[1] = ((w[0] & 0x0000ff00) >> 8) \
354 | ((w[2] & 0x0000ff00) >> 0) \
355 | ((w[4] & 0x0000ff00) << 8) \
356 | ((w[6] & 0x0000ff00) << 16); \
357 k[2] = ((w[0] & 0x00ff0000) >> 16) \
358 | ((w[2] & 0x00ff0000) >> 8) \
359 | ((w[4] & 0x00ff0000) << 0) \
360 | ((w[6] & 0x00ff0000) << 8); \
361 k[3] = ((w[0] & 0xff000000) >> 24) \
362 | ((w[2] & 0xff000000) >> 16) \
363 | ((w[4] & 0xff000000) >> 8) \
364 | ((w[6] & 0xff000000) >> 0); \
365 k[4] = ((w[1] & 0x000000ff) << 0) \
366 | ((w[3] & 0x000000ff) << 8) \
367 | ((w[5] & 0x000000ff) << 16) \
368 | ((w[7] & 0x000000ff) << 24); \
369 k[5] = ((w[1] & 0x0000ff00) >> 8) \
370 | ((w[3] & 0x0000ff00) >> 0) \
371 | ((w[5] & 0x0000ff00) << 8) \
372 | ((w[7] & 0x0000ff00) << 16); \
373 k[6] = ((w[1] & 0x00ff0000) >> 16) \
374 | ((w[3] & 0x00ff0000) >> 8) \
375 | ((w[5] & 0x00ff0000) << 0) \
376 | ((w[7] & 0x00ff0000) << 8); \
377 k[7] = ((w[1] & 0xff000000) >> 24) \
378 | ((w[3] & 0xff000000) >> 16) \
379 | ((w[5] & 0xff000000) >> 8) \
380 | ((w[7] & 0xff000000) >> 0);
417 x[0] ^= 0xff00ff00; \
418 x[1] ^= 0xff00ff00; \
419 x[2] ^= 0x00ff00ff; \
420 x[3] ^= 0x00ff00ff; \
421 x[4] ^= 0x00ffff00; \
422 x[5] ^= 0xff0000ff; \
423 x[6] ^= 0x000000ff; \
426 #define SHIFT12(u,m,s) \
427 u[0] = m[0] ^ s[6]; \
428 u[1] = m[1] ^ s[7]; \
429 u[2] = m[2] ^ (s[0] << 16) \
431 ^ (s[0] & 0x0000ffff) \
432 ^ (s[1] & 0x0000ffff) \
437 ^ (s[7] & 0xffff0000) \
439 u[3] = m[3] ^ (s[0] & 0x0000ffff) \
441 ^ (s[1] & 0x0000ffff) \
450 ^ (s[7] & 0x0000ffff) \
453 u[4] = m[4] ^ (s[0] & 0xffff0000) \
456 ^ (s[1] & 0xffff0000) \
465 ^ (s[7] & 0x0000ffff) \
468 u[5] = m[5] ^ (s[0] << 16) \
470 ^ (s[0] & 0xffff0000) \
471 ^ (s[1] & 0x0000ffff) \
481 ^ (s[7] & 0xffff0000) \
497 u[7] = m[7] ^ (s[0] & 0xffff0000) \
499 ^ (s[1] & 0x0000ffff) \
508 ^ (s[7] & 0x0000ffff) \
512 #define SHIFT16(h,v,u) \
513 v[0] = h[0] ^ (u[1] << 16) \
515 v[1] = h[1] ^ (u[2] << 16) \
517 v[2] = h[2] ^ (u[3] << 16) \
519 v[3] = h[3] ^ (u[4] << 16) \
521 v[4] = h[4] ^ (u[5] << 16) \
523 v[5] = h[5] ^ (u[6] << 16) \
525 v[6] = h[6] ^ (u[7] << 16) \
527 v[7] = h[7] ^ (u[0] & 0xffff0000) \
530 ^ (u[1] & 0xffff0000) \
533 ^ (u[7] & 0xffff0000);
535 #define SHIFT61(h,v) \
536 h[0] = (v[0] & 0xffff0000) \
540 ^ (v[1] & 0xffff0000) \
549 ^ (v[7] & 0x0000ffff); \
550 h[1] = (v[0] << 16) \
552 ^ (v[0] & 0xffff0000) \
553 ^ (v[1] & 0x0000ffff) \
561 ^ (v[7] & 0xffff0000) \
563 h[2] = (v[0] & 0x0000ffff) \
567 ^ (v[1] & 0xffff0000) \
575 ^ (v[7] & 0x0000ffff) \
578 h[3] = (v[0] << 16) \
580 ^ (v[0] & 0xffff0000) \
581 ^ (v[1] & 0xffff0000) \
591 ^ (v[7] & 0x0000ffff) \
593 h[4] = (v[0] >> 16) \
607 h[5] = (v[0] << 16) \
608 ^ (v[0] & 0xffff0000) \
611 ^ (v[1] & 0xffff0000) \
625 ^ (v[7] & 0xffff0000); \
657 #define PASS0(h,s,u,v,t) \
668 #define PASS2(h,s,u,v,t) \
680 #define PASS4(h,s,u,v,t) \
691 #define PASS6(h,s,u,v,t) \
700 __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)
706 const u32 gid = get_global_id (0);
707 const u32 lid = get_local_id (0);
708 const u32 lsz = get_local_size (0);
714 __local u32 s_tables[4][256];
716 for (u32 i = lid; i < 256; i += lsz)
718 s_tables[0][i] = c_tables[0][i];
719 s_tables[1][i] = c_tables[1][i];
720 s_tables[2][i] = c_tables[2][i];
721 s_tables[3][i] = c_tables[3][i];
724 barrier (CLK_LOCAL_MEM_FENCE);
726 if (gid >= gid_max) return;
734 pw_buf0[0] = pws[gid].i[ 0];
735 pw_buf0[1] = pws[gid].i[ 1];
736 pw_buf0[2] = pws[gid].i[ 2];
737 pw_buf0[3] = pws[gid].i[ 3];
741 pw_buf1[0] = pws[gid].i[ 4];
742 pw_buf1[1] = pws[gid].i[ 5];
743 pw_buf1[2] = pws[gid].i[ 6];
744 pw_buf1[3] = pws[gid].i[ 7];
746 const u32 pw_len = pws[gid].pw_len;
752 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
782 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
784 u32 w14 = out_len * 8;
821 state_m[0] = state[0];
822 state_m[1] = state[1];
823 state_m[2] = state[2];
824 state_m[3] = state[3];
825 state_m[4] = state[4];
826 state_m[5] = state[5];
827 state_m[6] = state[6];
828 state_m[7] = state[7];
843 PASS0 (state, tmp, state_m, data_m, s_tables);
844 PASS2 (state, tmp, state_m, data_m, s_tables);
845 PASS4 (state, tmp, state_m, data_m, s_tables);
846 PASS6 (state, tmp, state_m, data_m, s_tables);
848 SHIFT12 (state_m, data, tmp);
849 SHIFT16 (state, data_m, state_m);
850 SHIFT61 (state, data_m);
864 state_m[0] = state[0];
865 state_m[1] = state[1];
866 state_m[2] = state[2];
867 state_m[3] = state[3];
868 state_m[4] = state[4];
869 state_m[5] = state[5];
870 state_m[6] = state[6];
871 state_m[7] = state[7];
882 PASS0 (state, tmp, state_m, data_m, s_tables);
883 PASS2 (state, tmp, state_m, data_m, s_tables);
884 PASS4 (state, tmp, state_m, data_m, s_tables);
885 PASS6 (state, tmp, state_m, data_m, s_tables);
887 SHIFT12 (state_m, data, tmp);
888 SHIFT16 (state, data_m, state_m);
889 SHIFT61 (state, data_m);
902 state_m[0] = state[0];
903 state_m[1] = state[1];
904 state_m[2] = state[2];
905 state_m[3] = state[3];
906 state_m[4] = state[4];
907 state_m[5] = state[5];
908 state_m[6] = state[6];
909 state_m[7] = state[7];
920 PASS0 (state, tmp, state_m, data_m, s_tables);
921 PASS2 (state, tmp, state_m, data_m, s_tables);
922 PASS4 (state, tmp, state_m, data_m, s_tables);
923 PASS6 (state, tmp, state_m, data_m, s_tables);
925 SHIFT12 (state_m, data, tmp);
926 SHIFT16 (state, data_m, state_m);
927 SHIFT61 (state, data_m);
931 const u32 r0 = state[0];
932 const u32 r1 = state[1];
933 const u32 r2 = state[2];
934 const u32 r3 = state[3];
940 __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)
944 __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)
948 __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)
954 const u32 gid = get_global_id (0);
955 const u32 lid = get_local_id (0);
956 const u32 lsz = get_local_size (0);
962 __local u32 s_tables[4][256];
964 for (u32 i = lid; i < 256; i += lsz)
966 s_tables[0][i] = c_tables[0][i];
967 s_tables[1][i] = c_tables[1][i];
968 s_tables[2][i] = c_tables[2][i];
969 s_tables[3][i] = c_tables[3][i];
972 barrier (CLK_LOCAL_MEM_FENCE);
974 if (gid >= gid_max) return;
982 pw_buf0[0] = pws[gid].i[ 0];
983 pw_buf0[1] = pws[gid].i[ 1];
984 pw_buf0[2] = pws[gid].i[ 2];
985 pw_buf0[3] = pws[gid].i[ 3];
989 pw_buf1[0] = pws[gid].i[ 4];
990 pw_buf1[1] = pws[gid].i[ 5];
991 pw_buf1[2] = pws[gid].i[ 6];
992 pw_buf1[3] = pws[gid].i[ 7];
994 const u32 pw_len = pws[gid].pw_len;
1000 const u32 search[4] =
1002 digests_buf[digests_offset].digest_buf[DGST_R0],
1003 digests_buf[digests_offset].digest_buf[DGST_R1],
1004 digests_buf[digests_offset].digest_buf[DGST_R2],
1005 digests_buf[digests_offset].digest_buf[DGST_R3]
1012 for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
1042 const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
1044 u32 w14 = out_len * 8;
1067 state[ 8] = data[0];
1068 state[ 9] = data[1];
1069 state[10] = data[2];
1070 state[11] = data[3];
1071 state[12] = data[4];
1072 state[13] = data[5];
1073 state[14] = data[6];
1074 state[15] = data[7];
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];
1103 PASS0 (state, tmp, state_m, data_m, s_tables);
1104 PASS2 (state, tmp, state_m, data_m, s_tables);
1105 PASS4 (state, tmp, state_m, data_m, s_tables);
1106 PASS6 (state, tmp, state_m, data_m, s_tables);
1108 SHIFT12 (state_m, data, tmp);
1109 SHIFT16 (state, data_m, state_m);
1110 SHIFT61 (state, data_m);
1124 state_m[0] = state[0];
1125 state_m[1] = state[1];
1126 state_m[2] = state[2];
1127 state_m[3] = state[3];
1128 state_m[4] = state[4];
1129 state_m[5] = state[5];
1130 state_m[6] = state[6];
1131 state_m[7] = state[7];
1133 data_m[0] = data[0];
1134 data_m[1] = data[1];
1135 data_m[2] = data[2];
1136 data_m[3] = data[3];
1137 data_m[4] = data[4];
1138 data_m[5] = data[5];
1139 data_m[6] = data[6];
1140 data_m[7] = data[7];
1142 PASS0 (state, tmp, state_m, data_m, s_tables);
1143 PASS2 (state, tmp, state_m, data_m, s_tables);
1144 PASS4 (state, tmp, state_m, data_m, s_tables);
1145 PASS6 (state, tmp, state_m, data_m, s_tables);
1147 SHIFT12 (state_m, data, tmp);
1148 SHIFT16 (state, data_m, state_m);
1149 SHIFT61 (state, data_m);
1153 data[0] = state[ 8];
1154 data[1] = state[ 9];
1155 data[2] = state[10];
1156 data[3] = state[11];
1157 data[4] = state[12];
1158 data[5] = state[13];
1159 data[6] = state[14];
1160 data[7] = state[15];
1162 state_m[0] = state[0];
1163 state_m[1] = state[1];
1164 state_m[2] = state[2];
1165 state_m[3] = state[3];
1166 state_m[4] = state[4];
1167 state_m[5] = state[5];
1168 state_m[6] = state[6];
1169 state_m[7] = state[7];
1171 data_m[0] = data[0];
1172 data_m[1] = data[1];
1173 data_m[2] = data[2];
1174 data_m[3] = data[3];
1175 data_m[4] = data[4];
1176 data_m[5] = data[5];
1177 data_m[6] = data[6];
1178 data_m[7] = data[7];
1180 PASS0 (state, tmp, state_m, data_m, s_tables);
1181 PASS2 (state, tmp, state_m, data_m, s_tables);
1182 PASS4 (state, tmp, state_m, data_m, s_tables);
1183 PASS6 (state, tmp, state_m, data_m, s_tables);
1185 SHIFT12 (state_m, data, tmp);
1186 SHIFT16 (state, data_m, state_m);
1187 SHIFT61 (state, data_m);
1191 const u32 r0 = state[0];
1192 const u32 r1 = state[1];
1193 const u32 r2 = state[2];
1194 const u32 r3 = state[3];
1200 __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)
1204 __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)