2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
4 * magnum <john.magnum@hushmail.com>
13 #include "inc_vendor.cl"
14 #include "inc_hash_constants.h"
15 #include "inc_hash_functions.cl"
16 #include "inc_types.cl"
17 #include "inc_common.cl"
20 #include "inc_simd.cl"
22 __constant u32 c_tables[4][256] =
25 0x00072000, 0x00075000, 0x00074800, 0x00071000,
26 0x00076800, 0x00074000, 0x00070000, 0x00077000,
27 0x00073000, 0x00075800, 0x00070800, 0x00076000,
28 0x00073800, 0x00077800, 0x00072800, 0x00071800,
29 0x0005a000, 0x0005d000, 0x0005c800, 0x00059000,
30 0x0005e800, 0x0005c000, 0x00058000, 0x0005f000,
31 0x0005b000, 0x0005d800, 0x00058800, 0x0005e000,
32 0x0005b800, 0x0005f800, 0x0005a800, 0x00059800,
33 0x00022000, 0x00025000, 0x00024800, 0x00021000,
34 0x00026800, 0x00024000, 0x00020000, 0x00027000,
35 0x00023000, 0x00025800, 0x00020800, 0x00026000,
36 0x00023800, 0x00027800, 0x00022800, 0x00021800,
37 0x00062000, 0x00065000, 0x00064800, 0x00061000,
38 0x00066800, 0x00064000, 0x00060000, 0x00067000,
39 0x00063000, 0x00065800, 0x00060800, 0x00066000,
40 0x00063800, 0x00067800, 0x00062800, 0x00061800,
41 0x00032000, 0x00035000, 0x00034800, 0x00031000,
42 0x00036800, 0x00034000, 0x00030000, 0x00037000,
43 0x00033000, 0x00035800, 0x00030800, 0x00036000,
44 0x00033800, 0x00037800, 0x00032800, 0x00031800,
45 0x0006a000, 0x0006d000, 0x0006c800, 0x00069000,
46 0x0006e800, 0x0006c000, 0x00068000, 0x0006f000,
47 0x0006b000, 0x0006d800, 0x00068800, 0x0006e000,
48 0x0006b800, 0x0006f800, 0x0006a800, 0x00069800,
49 0x0007a000, 0x0007d000, 0x0007c800, 0x00079000,
50 0x0007e800, 0x0007c000, 0x00078000, 0x0007f000,
51 0x0007b000, 0x0007d800, 0x00078800, 0x0007e000,
52 0x0007b800, 0x0007f800, 0x0007a800, 0x00079800,
53 0x00052000, 0x00055000, 0x00054800, 0x00051000,
54 0x00056800, 0x00054000, 0x00050000, 0x00057000,
55 0x00053000, 0x00055800, 0x00050800, 0x00056000,
56 0x00053800, 0x00057800, 0x00052800, 0x00051800,
57 0x00012000, 0x00015000, 0x00014800, 0x00011000,
58 0x00016800, 0x00014000, 0x00010000, 0x00017000,
59 0x00013000, 0x00015800, 0x00010800, 0x00016000,
60 0x00013800, 0x00017800, 0x00012800, 0x00011800,
61 0x0001a000, 0x0001d000, 0x0001c800, 0x00019000,
62 0x0001e800, 0x0001c000, 0x00018000, 0x0001f000,
63 0x0001b000, 0x0001d800, 0x00018800, 0x0001e000,
64 0x0001b800, 0x0001f800, 0x0001a800, 0x00019800,
65 0x00042000, 0x00045000, 0x00044800, 0x00041000,
66 0x00046800, 0x00044000, 0x00040000, 0x00047000,
67 0x00043000, 0x00045800, 0x00040800, 0x00046000,
68 0x00043800, 0x00047800, 0x00042800, 0x00041800,
69 0x0000a000, 0x0000d000, 0x0000c800, 0x00009000,
70 0x0000e800, 0x0000c000, 0x00008000, 0x0000f000,
71 0x0000b000, 0x0000d800, 0x00008800, 0x0000e000,
72 0x0000b800, 0x0000f800, 0x0000a800, 0x00009800,
73 0x00002000, 0x00005000, 0x00004800, 0x00001000,
74 0x00006800, 0x00004000, 0x00000000, 0x00007000,
75 0x00003000, 0x00005800, 0x00000800, 0x00006000,
76 0x00003800, 0x00007800, 0x00002800, 0x00001800,
77 0x0003a000, 0x0003d000, 0x0003c800, 0x00039000,
78 0x0003e800, 0x0003c000, 0x00038000, 0x0003f000,
79 0x0003b000, 0x0003d800, 0x00038800, 0x0003e000,
80 0x0003b800, 0x0003f800, 0x0003a800, 0x00039800,
81 0x0002a000, 0x0002d000, 0x0002c800, 0x00029000,
82 0x0002e800, 0x0002c000, 0x00028000, 0x0002f000,
83 0x0002b000, 0x0002d800, 0x00028800, 0x0002e000,
84 0x0002b800, 0x0002f800, 0x0002a800, 0x00029800,
85 0x0004a000, 0x0004d000, 0x0004c800, 0x00049000,
86 0x0004e800, 0x0004c000, 0x00048000, 0x0004f000,
87 0x0004b000, 0x0004d800, 0x00048800, 0x0004e000,
88 0x0004b800, 0x0004f800, 0x0004a800, 0x00049800,
91 0x03a80000, 0x03c00000, 0x03880000, 0x03e80000,
92 0x03d00000, 0x03980000, 0x03a00000, 0x03900000,
93 0x03f00000, 0x03f80000, 0x03e00000, 0x03b80000,
94 0x03b00000, 0x03800000, 0x03c80000, 0x03d80000,
95 0x06a80000, 0x06c00000, 0x06880000, 0x06e80000,
96 0x06d00000, 0x06980000, 0x06a00000, 0x06900000,
97 0x06f00000, 0x06f80000, 0x06e00000, 0x06b80000,
98 0x06b00000, 0x06800000, 0x06c80000, 0x06d80000,
99 0x05280000, 0x05400000, 0x05080000, 0x05680000,
100 0x05500000, 0x05180000, 0x05200000, 0x05100000,
101 0x05700000, 0x05780000, 0x05600000, 0x05380000,
102 0x05300000, 0x05000000, 0x05480000, 0x05580000,
103 0x00a80000, 0x00c00000, 0x00880000, 0x00e80000,
104 0x00d00000, 0x00980000, 0x00a00000, 0x00900000,
105 0x00f00000, 0x00f80000, 0x00e00000, 0x00b80000,
106 0x00b00000, 0x00800000, 0x00c80000, 0x00d80000,
107 0x00280000, 0x00400000, 0x00080000, 0x00680000,
108 0x00500000, 0x00180000, 0x00200000, 0x00100000,
109 0x00700000, 0x00780000, 0x00600000, 0x00380000,
110 0x00300000, 0x00000000, 0x00480000, 0x00580000,
111 0x04280000, 0x04400000, 0x04080000, 0x04680000,
112 0x04500000, 0x04180000, 0x04200000, 0x04100000,
113 0x04700000, 0x04780000, 0x04600000, 0x04380000,
114 0x04300000, 0x04000000, 0x04480000, 0x04580000,
115 0x04a80000, 0x04c00000, 0x04880000, 0x04e80000,
116 0x04d00000, 0x04980000, 0x04a00000, 0x04900000,
117 0x04f00000, 0x04f80000, 0x04e00000, 0x04b80000,
118 0x04b00000, 0x04800000, 0x04c80000, 0x04d80000,
119 0x07a80000, 0x07c00000, 0x07880000, 0x07e80000,
120 0x07d00000, 0x07980000, 0x07a00000, 0x07900000,
121 0x07f00000, 0x07f80000, 0x07e00000, 0x07b80000,
122 0x07b00000, 0x07800000, 0x07c80000, 0x07d80000,
123 0x07280000, 0x07400000, 0x07080000, 0x07680000,
124 0x07500000, 0x07180000, 0x07200000, 0x07100000,
125 0x07700000, 0x07780000, 0x07600000, 0x07380000,
126 0x07300000, 0x07000000, 0x07480000, 0x07580000,
127 0x02280000, 0x02400000, 0x02080000, 0x02680000,
128 0x02500000, 0x02180000, 0x02200000, 0x02100000,
129 0x02700000, 0x02780000, 0x02600000, 0x02380000,
130 0x02300000, 0x02000000, 0x02480000, 0x02580000,
131 0x03280000, 0x03400000, 0x03080000, 0x03680000,
132 0x03500000, 0x03180000, 0x03200000, 0x03100000,
133 0x03700000, 0x03780000, 0x03600000, 0x03380000,
134 0x03300000, 0x03000000, 0x03480000, 0x03580000,
135 0x06280000, 0x06400000, 0x06080000, 0x06680000,
136 0x06500000, 0x06180000, 0x06200000, 0x06100000,
137 0x06700000, 0x06780000, 0x06600000, 0x06380000,
138 0x06300000, 0x06000000, 0x06480000, 0x06580000,
139 0x05a80000, 0x05c00000, 0x05880000, 0x05e80000,
140 0x05d00000, 0x05980000, 0x05a00000, 0x05900000,
141 0x05f00000, 0x05f80000, 0x05e00000, 0x05b80000,
142 0x05b00000, 0x05800000, 0x05c80000, 0x05d80000,
143 0x01280000, 0x01400000, 0x01080000, 0x01680000,
144 0x01500000, 0x01180000, 0x01200000, 0x01100000,
145 0x01700000, 0x01780000, 0x01600000, 0x01380000,
146 0x01300000, 0x01000000, 0x01480000, 0x01580000,
147 0x02a80000, 0x02c00000, 0x02880000, 0x02e80000,
148 0x02d00000, 0x02980000, 0x02a00000, 0x02900000,
149 0x02f00000, 0x02f80000, 0x02e00000, 0x02b80000,
150 0x02b00000, 0x02800000, 0x02c80000, 0x02d80000,
151 0x01a80000, 0x01c00000, 0x01880000, 0x01e80000,
152 0x01d00000, 0x01980000, 0x01a00000, 0x01900000,
153 0x01f00000, 0x01f80000, 0x01e00000, 0x01b80000,
154 0x01b00000, 0x01800000, 0x01c80000, 0x01d80000,
157 0x30000002, 0x60000002, 0x38000002, 0x08000002,
158 0x28000002, 0x78000002, 0x68000002, 0x40000002,
159 0x20000002, 0x50000002, 0x48000002, 0x70000002,
160 0x00000002, 0x18000002, 0x58000002, 0x10000002,
161 0xb0000005, 0xe0000005, 0xb8000005, 0x88000005,
162 0xa8000005, 0xf8000005, 0xe8000005, 0xc0000005,
163 0xa0000005, 0xd0000005, 0xc8000005, 0xf0000005,
164 0x80000005, 0x98000005, 0xd8000005, 0x90000005,
165 0x30000005, 0x60000005, 0x38000005, 0x08000005,
166 0x28000005, 0x78000005, 0x68000005, 0x40000005,
167 0x20000005, 0x50000005, 0x48000005, 0x70000005,
168 0x00000005, 0x18000005, 0x58000005, 0x10000005,
169 0x30000000, 0x60000000, 0x38000000, 0x08000000,
170 0x28000000, 0x78000000, 0x68000000, 0x40000000,
171 0x20000000, 0x50000000, 0x48000000, 0x70000000,
172 0x00000000, 0x18000000, 0x58000000, 0x10000000,
173 0xb0000003, 0xe0000003, 0xb8000003, 0x88000003,
174 0xa8000003, 0xf8000003, 0xe8000003, 0xc0000003,
175 0xa0000003, 0xd0000003, 0xc8000003, 0xf0000003,
176 0x80000003, 0x98000003, 0xd8000003, 0x90000003,
177 0x30000001, 0x60000001, 0x38000001, 0x08000001,
178 0x28000001, 0x78000001, 0x68000001, 0x40000001,
179 0x20000001, 0x50000001, 0x48000001, 0x70000001,
180 0x00000001, 0x18000001, 0x58000001, 0x10000001,
181 0xb0000000, 0xe0000000, 0xb8000000, 0x88000000,
182 0xa8000000, 0xf8000000, 0xe8000000, 0xc0000000,
183 0xa0000000, 0xd0000000, 0xc8000000, 0xf0000000,
184 0x80000000, 0x98000000, 0xd8000000, 0x90000000,
185 0xb0000006, 0xe0000006, 0xb8000006, 0x88000006,
186 0xa8000006, 0xf8000006, 0xe8000006, 0xc0000006,
187 0xa0000006, 0xd0000006, 0xc8000006, 0xf0000006,
188 0x80000006, 0x98000006, 0xd8000006, 0x90000006,
189 0xb0000001, 0xe0000001, 0xb8000001, 0x88000001,
190 0xa8000001, 0xf8000001, 0xe8000001, 0xc0000001,
191 0xa0000001, 0xd0000001, 0xc8000001, 0xf0000001,
192 0x80000001, 0x98000001, 0xd8000001, 0x90000001,
193 0x30000003, 0x60000003, 0x38000003, 0x08000003,
194 0x28000003, 0x78000003, 0x68000003, 0x40000003,
195 0x20000003, 0x50000003, 0x48000003, 0x70000003,
196 0x00000003, 0x18000003, 0x58000003, 0x10000003,
197 0x30000004, 0x60000004, 0x38000004, 0x08000004,
198 0x28000004, 0x78000004, 0x68000004, 0x40000004,
199 0x20000004, 0x50000004, 0x48000004, 0x70000004,
200 0x00000004, 0x18000004, 0x58000004, 0x10000004,
201 0xb0000002, 0xe0000002, 0xb8000002, 0x88000002,
202 0xa8000002, 0xf8000002, 0xe8000002, 0xc0000002,
203 0xa0000002, 0xd0000002, 0xc8000002, 0xf0000002,
204 0x80000002, 0x98000002, 0xd8000002, 0x90000002,
205 0xb0000004, 0xe0000004, 0xb8000004, 0x88000004,
206 0xa8000004, 0xf8000004, 0xe8000004, 0xc0000004,
207 0xa0000004, 0xd0000004, 0xc8000004, 0xf0000004,
208 0x80000004, 0x98000004, 0xd8000004, 0x90000004,
209 0x30000006, 0x60000006, 0x38000006, 0x08000006,
210 0x28000006, 0x78000006, 0x68000006, 0x40000006,
211 0x20000006, 0x50000006, 0x48000006, 0x70000006,
212 0x00000006, 0x18000006, 0x58000006, 0x10000006,
213 0xb0000007, 0xe0000007, 0xb8000007, 0x88000007,
214 0xa8000007, 0xf8000007, 0xe8000007, 0xc0000007,
215 0xa0000007, 0xd0000007, 0xc8000007, 0xf0000007,
216 0x80000007, 0x98000007, 0xd8000007, 0x90000007,
217 0x30000007, 0x60000007, 0x38000007, 0x08000007,
218 0x28000007, 0x78000007, 0x68000007, 0x40000007,
219 0x20000007, 0x50000007, 0x48000007, 0x70000007,
220 0x00000007, 0x18000007, 0x58000007, 0x10000007,
223 0x000000e8, 0x000000d8, 0x000000a0, 0x00000088,
224 0x00000098, 0x000000f8, 0x000000a8, 0x000000c8,
225 0x00000080, 0x000000d0, 0x000000f0, 0x000000b8,
226 0x000000b0, 0x000000c0, 0x00000090, 0x000000e0,
227 0x000007e8, 0x000007d8, 0x000007a0, 0x00000788,
228 0x00000798, 0x000007f8, 0x000007a8, 0x000007c8,
229 0x00000780, 0x000007d0, 0x000007f0, 0x000007b8,
230 0x000007b0, 0x000007c0, 0x00000790, 0x000007e0,
231 0x000006e8, 0x000006d8, 0x000006a0, 0x00000688,
232 0x00000698, 0x000006f8, 0x000006a8, 0x000006c8,
233 0x00000680, 0x000006d0, 0x000006f0, 0x000006b8,
234 0x000006b0, 0x000006c0, 0x00000690, 0x000006e0,
235 0x00000068, 0x00000058, 0x00000020, 0x00000008,
236 0x00000018, 0x00000078, 0x00000028, 0x00000048,
237 0x00000000, 0x00000050, 0x00000070, 0x00000038,
238 0x00000030, 0x00000040, 0x00000010, 0x00000060,
239 0x000002e8, 0x000002d8, 0x000002a0, 0x00000288,
240 0x00000298, 0x000002f8, 0x000002a8, 0x000002c8,
241 0x00000280, 0x000002d0, 0x000002f0, 0x000002b8,
242 0x000002b0, 0x000002c0, 0x00000290, 0x000002e0,
243 0x000003e8, 0x000003d8, 0x000003a0, 0x00000388,
244 0x00000398, 0x000003f8, 0x000003a8, 0x000003c8,
245 0x00000380, 0x000003d0, 0x000003f0, 0x000003b8,
246 0x000003b0, 0x000003c0, 0x00000390, 0x000003e0,
247 0x00000568, 0x00000558, 0x00000520, 0x00000508,
248 0x00000518, 0x00000578, 0x00000528, 0x00000548,
249 0x00000500, 0x00000550, 0x00000570, 0x00000538,
250 0x00000530, 0x00000540, 0x00000510, 0x00000560,
251 0x00000268, 0x00000258, 0x00000220, 0x00000208,
252 0x00000218, 0x00000278, 0x00000228, 0x00000248,
253 0x00000200, 0x00000250, 0x00000270, 0x00000238,
254 0x00000230, 0x00000240, 0x00000210, 0x00000260,
255 0x000004e8, 0x000004d8, 0x000004a0, 0x00000488,
256 0x00000498, 0x000004f8, 0x000004a8, 0x000004c8,
257 0x00000480, 0x000004d0, 0x000004f0, 0x000004b8,
258 0x000004b0, 0x000004c0, 0x00000490, 0x000004e0,
259 0x00000168, 0x00000158, 0x00000120, 0x00000108,
260 0x00000118, 0x00000178, 0x00000128, 0x00000148,
261 0x00000100, 0x00000150, 0x00000170, 0x00000138,
262 0x00000130, 0x00000140, 0x00000110, 0x00000160,
263 0x000001e8, 0x000001d8, 0x000001a0, 0x00000188,
264 0x00000198, 0x000001f8, 0x000001a8, 0x000001c8,
265 0x00000180, 0x000001d0, 0x000001f0, 0x000001b8,
266 0x000001b0, 0x000001c0, 0x00000190, 0x000001e0,
267 0x00000768, 0x00000758, 0x00000720, 0x00000708,
268 0x00000718, 0x00000778, 0x00000728, 0x00000748,
269 0x00000700, 0x00000750, 0x00000770, 0x00000738,
270 0x00000730, 0x00000740, 0x00000710, 0x00000760,
271 0x00000368, 0x00000358, 0x00000320, 0x00000308,
272 0x00000318, 0x00000378, 0x00000328, 0x00000348,
273 0x00000300, 0x00000350, 0x00000370, 0x00000338,
274 0x00000330, 0x00000340, 0x00000310, 0x00000360,
275 0x000005e8, 0x000005d8, 0x000005a0, 0x00000588,
276 0x00000598, 0x000005f8, 0x000005a8, 0x000005c8,
277 0x00000580, 0x000005d0, 0x000005f0, 0x000005b8,
278 0x000005b0, 0x000005c0, 0x00000590, 0x000005e0,
279 0x00000468, 0x00000458, 0x00000420, 0x00000408,
280 0x00000418, 0x00000478, 0x00000428, 0x00000448,
281 0x00000400, 0x00000450, 0x00000470, 0x00000438,
282 0x00000430, 0x00000440, 0x00000410, 0x00000460,
283 0x00000668, 0x00000658, 0x00000620, 0x00000608,
284 0x00000618, 0x00000678, 0x00000628, 0x00000648,
285 0x00000600, 0x00000650, 0x00000670, 0x00000638,
286 0x00000630, 0x00000640, 0x00000610, 0x00000660,
291 #define BOX(i,n,S) (S)[(n)][(i)]
293 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
295 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
297 #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])
298 #elif VECT_SIZE == 16
299 #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])
302 #define _round(k1,k2,tbl) \
306 l ^= 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); \
311 r ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
312 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
313 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
314 BOX (((t >> 24) & 0xff), 3, tbl); \
317 #define R(k,h,s,i,t) \
323 _round (k[0], k[1], t); \
324 _round (k[2], k[3], t); \
325 _round (k[4], k[5], t); \
326 _round (k[6], k[7], t); \
327 _round (k[0], k[1], t); \
328 _round (k[2], k[3], t); \
329 _round (k[4], k[5], t); \
330 _round (k[6], k[7], t); \
331 _round (k[0], k[1], t); \
332 _round (k[2], k[3], t); \
333 _round (k[4], k[5], t); \
334 _round (k[6], k[7], t); \
335 _round (k[7], k[6], t); \
336 _round (k[5], k[4], t); \
337 _round (k[3], k[2], t); \
338 _round (k[1], k[0], t); \
344 w[0] = u[0] ^ v[0]; \
345 w[1] = u[1] ^ v[1]; \
346 w[2] = u[2] ^ v[2]; \
347 w[3] = u[3] ^ v[3]; \
348 w[4] = u[4] ^ v[4]; \
349 w[5] = u[5] ^ v[5]; \
350 w[6] = u[6] ^ v[6]; \
354 k[0] = ((w[0] & 0x000000ff) << 0) \
355 | ((w[2] & 0x000000ff) << 8) \
356 | ((w[4] & 0x000000ff) << 16) \
357 | ((w[6] & 0x000000ff) << 24); \
358 k[1] = ((w[0] & 0x0000ff00) >> 8) \
359 | ((w[2] & 0x0000ff00) >> 0) \
360 | ((w[4] & 0x0000ff00) << 8) \
361 | ((w[6] & 0x0000ff00) << 16); \
362 k[2] = ((w[0] & 0x00ff0000) >> 16) \
363 | ((w[2] & 0x00ff0000) >> 8) \
364 | ((w[4] & 0x00ff0000) << 0) \
365 | ((w[6] & 0x00ff0000) << 8); \
366 k[3] = ((w[0] & 0xff000000) >> 24) \
367 | ((w[2] & 0xff000000) >> 16) \
368 | ((w[4] & 0xff000000) >> 8) \
369 | ((w[6] & 0xff000000) >> 0); \
370 k[4] = ((w[1] & 0x000000ff) << 0) \
371 | ((w[3] & 0x000000ff) << 8) \
372 | ((w[5] & 0x000000ff) << 16) \
373 | ((w[7] & 0x000000ff) << 24); \
374 k[5] = ((w[1] & 0x0000ff00) >> 8) \
375 | ((w[3] & 0x0000ff00) >> 0) \
376 | ((w[5] & 0x0000ff00) << 8) \
377 | ((w[7] & 0x0000ff00) << 16); \
378 k[6] = ((w[1] & 0x00ff0000) >> 16) \
379 | ((w[3] & 0x00ff0000) >> 8) \
380 | ((w[5] & 0x00ff0000) << 0) \
381 | ((w[7] & 0x00ff0000) << 8); \
382 k[7] = ((w[1] & 0xff000000) >> 24) \
383 | ((w[3] & 0xff000000) >> 16) \
384 | ((w[5] & 0xff000000) >> 8) \
385 | ((w[7] & 0xff000000) >> 0);
422 x[0] ^= 0xff00ff00; \
423 x[1] ^= 0xff00ff00; \
424 x[2] ^= 0x00ff00ff; \
425 x[3] ^= 0x00ff00ff; \
426 x[4] ^= 0x00ffff00; \
427 x[5] ^= 0xff0000ff; \
428 x[6] ^= 0x000000ff; \
431 #define SHIFT12(u,m,s) \
432 u[0] = m[0] ^ s[6]; \
433 u[1] = m[1] ^ s[7]; \
434 u[2] = m[2] ^ (s[0] << 16) \
436 ^ (s[0] & 0x0000ffff) \
437 ^ (s[1] & 0x0000ffff) \
442 ^ (s[7] & 0xffff0000) \
444 u[3] = m[3] ^ (s[0] & 0x0000ffff) \
446 ^ (s[1] & 0x0000ffff) \
455 ^ (s[7] & 0x0000ffff) \
458 u[4] = m[4] ^ (s[0] & 0xffff0000) \
461 ^ (s[1] & 0xffff0000) \
470 ^ (s[7] & 0x0000ffff) \
473 u[5] = m[5] ^ (s[0] << 16) \
475 ^ (s[0] & 0xffff0000) \
476 ^ (s[1] & 0x0000ffff) \
486 ^ (s[7] & 0xffff0000) \
502 u[7] = m[7] ^ (s[0] & 0xffff0000) \
504 ^ (s[1] & 0x0000ffff) \
513 ^ (s[7] & 0x0000ffff) \
517 #define SHIFT16(h,v,u) \
518 v[0] = h[0] ^ (u[1] << 16) \
520 v[1] = h[1] ^ (u[2] << 16) \
522 v[2] = h[2] ^ (u[3] << 16) \
524 v[3] = h[3] ^ (u[4] << 16) \
526 v[4] = h[4] ^ (u[5] << 16) \
528 v[5] = h[5] ^ (u[6] << 16) \
530 v[6] = h[6] ^ (u[7] << 16) \
532 v[7] = h[7] ^ (u[0] & 0xffff0000) \
535 ^ (u[1] & 0xffff0000) \
538 ^ (u[7] & 0xffff0000);
540 #define SHIFT61(h,v) \
541 h[0] = (v[0] & 0xffff0000) \
545 ^ (v[1] & 0xffff0000) \
554 ^ (v[7] & 0x0000ffff); \
555 h[1] = (v[0] << 16) \
557 ^ (v[0] & 0xffff0000) \
558 ^ (v[1] & 0x0000ffff) \
566 ^ (v[7] & 0xffff0000) \
568 h[2] = (v[0] & 0x0000ffff) \
572 ^ (v[1] & 0xffff0000) \
580 ^ (v[7] & 0x0000ffff) \
583 h[3] = (v[0] << 16) \
585 ^ (v[0] & 0xffff0000) \
586 ^ (v[1] & 0xffff0000) \
596 ^ (v[7] & 0x0000ffff) \
598 h[4] = (v[0] >> 16) \
612 h[5] = (v[0] << 16) \
613 ^ (v[0] & 0xffff0000) \
616 ^ (v[1] & 0xffff0000) \
630 ^ (v[7] & 0xffff0000); \
662 #define PASS0(h,s,u,v,t) \
673 #define PASS2(h,s,u,v,t) \
685 #define PASS4(h,s,u,v,t) \
696 #define PASS6(h,s,u,v,t) \
705 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
711 const u32 gid = get_global_id (0);
712 const u32 lid = get_local_id (0);
713 const u32 lsz = get_local_size (0);
719 __local u32 s_tables[4][256];
721 for (u32 i = lid; i < 256; i += lsz)
723 s_tables[0][i] = c_tables[0][i];
724 s_tables[1][i] = c_tables[1][i];
725 s_tables[2][i] = c_tables[2][i];
726 s_tables[3][i] = c_tables[3][i];
729 barrier (CLK_LOCAL_MEM_FENCE);
731 if (gid >= gid_max) return;
740 pw_buf0[0] = pws[gid].i[0];
741 pw_buf0[1] = pws[gid].i[1];
742 pw_buf0[2] = pws[gid].i[2];
743 pw_buf0[3] = pws[gid].i[3];
744 pw_buf1[0] = pws[gid].i[4];
745 pw_buf1[1] = pws[gid].i[5];
746 pw_buf1[2] = pws[gid].i[6];
747 pw_buf1[3] = pws[gid].i[7];
749 const u32 pw_len = pws[gid].pw_len;
755 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
762 const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
799 state_m[0] = state[0];
800 state_m[1] = state[1];
801 state_m[2] = state[2];
802 state_m[3] = state[3];
803 state_m[4] = state[4];
804 state_m[5] = state[5];
805 state_m[6] = state[6];
806 state_m[7] = state[7];
819 //if (pw_len > 0) // not really SIMD compatible
821 PASS0 (state, tmp, state_m, data_m, s_tables);
822 PASS2 (state, tmp, state_m, data_m, s_tables);
823 PASS4 (state, tmp, state_m, data_m, s_tables);
824 PASS6 (state, tmp, state_m, data_m, s_tables);
826 SHIFT12 (state_m, data, tmp);
827 SHIFT16 (state, data_m, state_m);
828 SHIFT61 (state, data_m);
831 data[0] = out_len * 8;
842 state_m[0] = state[0];
843 state_m[1] = state[1];
844 state_m[2] = state[2];
845 state_m[3] = state[3];
846 state_m[4] = state[4];
847 state_m[5] = state[5];
848 state_m[6] = state[6];
849 state_m[7] = state[7];
860 PASS0 (state, tmp, state_m, data_m, s_tables);
861 PASS2 (state, tmp, state_m, data_m, s_tables);
862 PASS4 (state, tmp, state_m, data_m, s_tables);
863 PASS6 (state, tmp, state_m, data_m, s_tables);
865 SHIFT12 (state_m, data, tmp);
866 SHIFT16 (state, data_m, state_m);
867 SHIFT61 (state, data_m);
880 state_m[0] = state[0];
881 state_m[1] = state[1];
882 state_m[2] = state[2];
883 state_m[3] = state[3];
884 state_m[4] = state[4];
885 state_m[5] = state[5];
886 state_m[6] = state[6];
887 state_m[7] = state[7];
898 PASS0 (state, tmp, state_m, data_m, s_tables);
899 PASS2 (state, tmp, state_m, data_m, s_tables);
900 PASS4 (state, tmp, state_m, data_m, s_tables);
901 PASS6 (state, tmp, state_m, data_m, s_tables);
903 SHIFT12 (state_m, data, tmp);
904 SHIFT16 (state, data_m, state_m);
905 SHIFT61 (state, data_m);
909 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
913 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
917 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
921 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 const u32 gid = get_global_id (0);
928 const u32 lid = get_local_id (0);
929 const u32 lsz = get_local_size (0);
935 __local u32 s_tables[4][256];
937 for (u32 i = lid; i < 256; i += lsz)
939 s_tables[0][i] = c_tables[0][i];
940 s_tables[1][i] = c_tables[1][i];
941 s_tables[2][i] = c_tables[2][i];
942 s_tables[3][i] = c_tables[3][i];
945 barrier (CLK_LOCAL_MEM_FENCE);
947 if (gid >= gid_max) return;
956 pw_buf0[0] = pws[gid].i[0];
957 pw_buf0[1] = pws[gid].i[1];
958 pw_buf0[2] = pws[gid].i[2];
959 pw_buf0[3] = pws[gid].i[3];
960 pw_buf1[0] = pws[gid].i[4];
961 pw_buf1[1] = pws[gid].i[5];
962 pw_buf1[2] = pws[gid].i[6];
963 pw_buf1[3] = pws[gid].i[7];
965 const u32 pw_len = pws[gid].pw_len;
971 const u32 search[4] =
973 digests_buf[digests_offset].digest_buf[DGST_R0],
974 digests_buf[digests_offset].digest_buf[DGST_R1],
975 digests_buf[digests_offset].digest_buf[DGST_R2],
976 digests_buf[digests_offset].digest_buf[DGST_R3]
983 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
990 const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
1017 state[ 8] = data[0];
1018 state[ 9] = data[1];
1019 state[10] = data[2];
1020 state[11] = data[3];
1021 state[12] = data[4];
1022 state[13] = data[5];
1023 state[14] = data[6];
1024 state[15] = data[7];
1031 state_m[0] = state[0];
1032 state_m[1] = state[1];
1033 state_m[2] = state[2];
1034 state_m[3] = state[3];
1035 state_m[4] = state[4];
1036 state_m[5] = state[5];
1037 state_m[6] = state[6];
1038 state_m[7] = state[7];
1040 data_m[0] = data[0];
1041 data_m[1] = data[1];
1042 data_m[2] = data[2];
1043 data_m[3] = data[3];
1044 data_m[4] = data[4];
1045 data_m[5] = data[5];
1046 data_m[6] = data[6];
1047 data_m[7] = data[7];
1051 //if (pw_len > 0) // not really SIMD compatible
1053 PASS0 (state, tmp, state_m, data_m, s_tables);
1054 PASS2 (state, tmp, state_m, data_m, s_tables);
1055 PASS4 (state, tmp, state_m, data_m, s_tables);
1056 PASS6 (state, tmp, state_m, data_m, s_tables);
1058 SHIFT12 (state_m, data, tmp);
1059 SHIFT16 (state, data_m, state_m);
1060 SHIFT61 (state, data_m);
1063 data[0] = out_len * 8;
1074 state_m[0] = state[0];
1075 state_m[1] = state[1];
1076 state_m[2] = state[2];
1077 state_m[3] = state[3];
1078 state_m[4] = state[4];
1079 state_m[5] = state[5];
1080 state_m[6] = state[6];
1081 state_m[7] = state[7];
1083 data_m[0] = data[0];
1084 data_m[1] = data[1];
1085 data_m[2] = data[2];
1086 data_m[3] = data[3];
1087 data_m[4] = data[4];
1088 data_m[5] = data[5];
1089 data_m[6] = data[6];
1090 data_m[7] = data[7];
1092 PASS0 (state, tmp, state_m, data_m, s_tables);
1093 PASS2 (state, tmp, state_m, data_m, s_tables);
1094 PASS4 (state, tmp, state_m, data_m, s_tables);
1095 PASS6 (state, tmp, state_m, data_m, s_tables);
1097 SHIFT12 (state_m, data, tmp);
1098 SHIFT16 (state, data_m, state_m);
1099 SHIFT61 (state, data_m);
1103 data[0] = state[ 8];
1104 data[1] = state[ 9];
1105 data[2] = state[10];
1106 data[3] = state[11];
1107 data[4] = state[12];
1108 data[5] = state[13];
1109 data[6] = state[14];
1110 data[7] = state[15];
1112 state_m[0] = state[0];
1113 state_m[1] = state[1];
1114 state_m[2] = state[2];
1115 state_m[3] = state[3];
1116 state_m[4] = state[4];
1117 state_m[5] = state[5];
1118 state_m[6] = state[6];
1119 state_m[7] = state[7];
1121 data_m[0] = data[0];
1122 data_m[1] = data[1];
1123 data_m[2] = data[2];
1124 data_m[3] = data[3];
1125 data_m[4] = data[4];
1126 data_m[5] = data[5];
1127 data_m[6] = data[6];
1128 data_m[7] = data[7];
1130 PASS0 (state, tmp, state_m, data_m, s_tables);
1131 PASS2 (state, tmp, state_m, data_m, s_tables);
1132 PASS4 (state, tmp, state_m, data_m, s_tables);
1133 PASS6 (state, tmp, state_m, data_m, s_tables);
1135 SHIFT12 (state_m, data, tmp);
1136 SHIFT16 (state, data_m, state_m);
1137 SHIFT61 (state, data_m);
1141 COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
1145 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)
1149 __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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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)