2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
4 * magnum <john.magnum@hushmail.com>
13 #include "inc_hash_constants.h"
14 #include "inc_vendor.cl"
21 #include "inc_hash_functions.cl"
22 #include "inc_types.cl"
23 #include "inc_common.cl"
26 #include "inc_simd.cl"
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;
746 pw_buf0[0] = pws[gid].i[0];
747 pw_buf0[1] = pws[gid].i[1];
748 pw_buf0[2] = pws[gid].i[2];
749 pw_buf0[3] = pws[gid].i[3];
750 pw_buf1[0] = pws[gid].i[4];
751 pw_buf1[1] = pws[gid].i[5];
752 pw_buf1[2] = pws[gid].i[6];
753 pw_buf1[3] = pws[gid].i[7];
755 const u32 pw_len = pws[gid].pw_len;
761 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
768 const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
805 state_m[0] = state[0];
806 state_m[1] = state[1];
807 state_m[2] = state[2];
808 state_m[3] = state[3];
809 state_m[4] = state[4];
810 state_m[5] = state[5];
811 state_m[6] = state[6];
812 state_m[7] = state[7];
825 //if (pw_len > 0) // not really SIMD compatible
827 PASS0 (state, tmp, state_m, data_m, s_tables);
828 PASS2 (state, tmp, state_m, data_m, s_tables);
829 PASS4 (state, tmp, state_m, data_m, s_tables);
830 PASS6 (state, tmp, state_m, data_m, s_tables);
832 SHIFT12 (state_m, data, tmp);
833 SHIFT16 (state, data_m, state_m);
834 SHIFT61 (state, data_m);
837 data[0] = out_len * 8;
848 state_m[0] = state[0];
849 state_m[1] = state[1];
850 state_m[2] = state[2];
851 state_m[3] = state[3];
852 state_m[4] = state[4];
853 state_m[5] = state[5];
854 state_m[6] = state[6];
855 state_m[7] = state[7];
866 PASS0 (state, tmp, state_m, data_m, s_tables);
867 PASS2 (state, tmp, state_m, data_m, s_tables);
868 PASS4 (state, tmp, state_m, data_m, s_tables);
869 PASS6 (state, tmp, state_m, data_m, s_tables);
871 SHIFT12 (state_m, data, tmp);
872 SHIFT16 (state, data_m, state_m);
873 SHIFT61 (state, data_m);
886 state_m[0] = state[0];
887 state_m[1] = state[1];
888 state_m[2] = state[2];
889 state_m[3] = state[3];
890 state_m[4] = state[4];
891 state_m[5] = state[5];
892 state_m[6] = state[6];
893 state_m[7] = state[7];
904 PASS0 (state, tmp, state_m, data_m, s_tables);
905 PASS2 (state, tmp, state_m, data_m, s_tables);
906 PASS4 (state, tmp, state_m, data_m, s_tables);
907 PASS6 (state, tmp, state_m, data_m, s_tables);
909 SHIFT12 (state_m, data, tmp);
910 SHIFT16 (state, data_m, state_m);
911 SHIFT61 (state, data_m);
915 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
919 __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)
923 __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)
927 __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)
933 const u32 gid = get_global_id (0);
934 const u32 lid = get_local_id (0);
935 const u32 lsz = get_local_size (0);
941 __local u32 s_tables[4][256];
943 for (u32 i = lid; i < 256; i += lsz)
945 s_tables[0][i] = c_tables[0][i];
946 s_tables[1][i] = c_tables[1][i];
947 s_tables[2][i] = c_tables[2][i];
948 s_tables[3][i] = c_tables[3][i];
951 barrier (CLK_LOCAL_MEM_FENCE);
953 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];
966 pw_buf1[0] = pws[gid].i[4];
967 pw_buf1[1] = pws[gid].i[5];
968 pw_buf1[2] = pws[gid].i[6];
969 pw_buf1[3] = pws[gid].i[7];
971 const u32 pw_len = pws[gid].pw_len;
977 const u32 search[4] =
979 digests_buf[digests_offset].digest_buf[DGST_R0],
980 digests_buf[digests_offset].digest_buf[DGST_R1],
981 digests_buf[digests_offset].digest_buf[DGST_R2],
982 digests_buf[digests_offset].digest_buf[DGST_R3]
989 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
996 const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
1023 state[ 8] = data[0];
1024 state[ 9] = data[1];
1025 state[10] = data[2];
1026 state[11] = data[3];
1027 state[12] = data[4];
1028 state[13] = data[5];
1029 state[14] = data[6];
1030 state[15] = data[7];
1037 state_m[0] = state[0];
1038 state_m[1] = state[1];
1039 state_m[2] = state[2];
1040 state_m[3] = state[3];
1041 state_m[4] = state[4];
1042 state_m[5] = state[5];
1043 state_m[6] = state[6];
1044 state_m[7] = state[7];
1046 data_m[0] = data[0];
1047 data_m[1] = data[1];
1048 data_m[2] = data[2];
1049 data_m[3] = data[3];
1050 data_m[4] = data[4];
1051 data_m[5] = data[5];
1052 data_m[6] = data[6];
1053 data_m[7] = data[7];
1057 //if (pw_len > 0) // not really SIMD compatible
1059 PASS0 (state, tmp, state_m, data_m, s_tables);
1060 PASS2 (state, tmp, state_m, data_m, s_tables);
1061 PASS4 (state, tmp, state_m, data_m, s_tables);
1062 PASS6 (state, tmp, state_m, data_m, s_tables);
1064 SHIFT12 (state_m, data, tmp);
1065 SHIFT16 (state, data_m, state_m);
1066 SHIFT61 (state, data_m);
1069 data[0] = out_len * 8;
1080 state_m[0] = state[0];
1081 state_m[1] = state[1];
1082 state_m[2] = state[2];
1083 state_m[3] = state[3];
1084 state_m[4] = state[4];
1085 state_m[5] = state[5];
1086 state_m[6] = state[6];
1087 state_m[7] = state[7];
1089 data_m[0] = data[0];
1090 data_m[1] = data[1];
1091 data_m[2] = data[2];
1092 data_m[3] = data[3];
1093 data_m[4] = data[4];
1094 data_m[5] = data[5];
1095 data_m[6] = data[6];
1096 data_m[7] = data[7];
1098 PASS0 (state, tmp, state_m, data_m, s_tables);
1099 PASS2 (state, tmp, state_m, data_m, s_tables);
1100 PASS4 (state, tmp, state_m, data_m, s_tables);
1101 PASS6 (state, tmp, state_m, data_m, s_tables);
1103 SHIFT12 (state_m, data, tmp);
1104 SHIFT16 (state, data_m, state_m);
1105 SHIFT61 (state, data_m);
1109 data[0] = state[ 8];
1110 data[1] = state[ 9];
1111 data[2] = state[10];
1112 data[3] = state[11];
1113 data[4] = state[12];
1114 data[5] = state[13];
1115 data[6] = state[14];
1116 data[7] = state[15];
1118 state_m[0] = state[0];
1119 state_m[1] = state[1];
1120 state_m[2] = state[2];
1121 state_m[3] = state[3];
1122 state_m[4] = state[4];
1123 state_m[5] = state[5];
1124 state_m[6] = state[6];
1125 state_m[7] = state[7];
1127 data_m[0] = data[0];
1128 data_m[1] = data[1];
1129 data_m[2] = data[2];
1130 data_m[3] = data[3];
1131 data_m[4] = data[4];
1132 data_m[5] = data[5];
1133 data_m[6] = data[6];
1134 data_m[7] = data[7];
1136 PASS0 (state, tmp, state_m, data_m, s_tables);
1137 PASS2 (state, tmp, state_m, data_m, s_tables);
1138 PASS4 (state, tmp, state_m, data_m, s_tables);
1139 PASS6 (state, tmp, state_m, data_m, s_tables);
1141 SHIFT12 (state_m, data, tmp);
1142 SHIFT16 (state, data_m, state_m);
1143 SHIFT61 (state, data_m);
1147 COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
1151 __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)
1155 __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)