2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
12 #include "inc_vendor.cl"
13 #include "inc_hash_constants.h"
14 #include "inc_hash_functions.cl"
15 #include "inc_types.cl"
16 #include "inc_common.cl"
17 #include "inc_simd.cl"
19 __constant u32 c_tables[4][256] =
22 0x00072000, 0x00075000, 0x00074800, 0x00071000,
23 0x00076800, 0x00074000, 0x00070000, 0x00077000,
24 0x00073000, 0x00075800, 0x00070800, 0x00076000,
25 0x00073800, 0x00077800, 0x00072800, 0x00071800,
26 0x0005a000, 0x0005d000, 0x0005c800, 0x00059000,
27 0x0005e800, 0x0005c000, 0x00058000, 0x0005f000,
28 0x0005b000, 0x0005d800, 0x00058800, 0x0005e000,
29 0x0005b800, 0x0005f800, 0x0005a800, 0x00059800,
30 0x00022000, 0x00025000, 0x00024800, 0x00021000,
31 0x00026800, 0x00024000, 0x00020000, 0x00027000,
32 0x00023000, 0x00025800, 0x00020800, 0x00026000,
33 0x00023800, 0x00027800, 0x00022800, 0x00021800,
34 0x00062000, 0x00065000, 0x00064800, 0x00061000,
35 0x00066800, 0x00064000, 0x00060000, 0x00067000,
36 0x00063000, 0x00065800, 0x00060800, 0x00066000,
37 0x00063800, 0x00067800, 0x00062800, 0x00061800,
38 0x00032000, 0x00035000, 0x00034800, 0x00031000,
39 0x00036800, 0x00034000, 0x00030000, 0x00037000,
40 0x00033000, 0x00035800, 0x00030800, 0x00036000,
41 0x00033800, 0x00037800, 0x00032800, 0x00031800,
42 0x0006a000, 0x0006d000, 0x0006c800, 0x00069000,
43 0x0006e800, 0x0006c000, 0x00068000, 0x0006f000,
44 0x0006b000, 0x0006d800, 0x00068800, 0x0006e000,
45 0x0006b800, 0x0006f800, 0x0006a800, 0x00069800,
46 0x0007a000, 0x0007d000, 0x0007c800, 0x00079000,
47 0x0007e800, 0x0007c000, 0x00078000, 0x0007f000,
48 0x0007b000, 0x0007d800, 0x00078800, 0x0007e000,
49 0x0007b800, 0x0007f800, 0x0007a800, 0x00079800,
50 0x00052000, 0x00055000, 0x00054800, 0x00051000,
51 0x00056800, 0x00054000, 0x00050000, 0x00057000,
52 0x00053000, 0x00055800, 0x00050800, 0x00056000,
53 0x00053800, 0x00057800, 0x00052800, 0x00051800,
54 0x00012000, 0x00015000, 0x00014800, 0x00011000,
55 0x00016800, 0x00014000, 0x00010000, 0x00017000,
56 0x00013000, 0x00015800, 0x00010800, 0x00016000,
57 0x00013800, 0x00017800, 0x00012800, 0x00011800,
58 0x0001a000, 0x0001d000, 0x0001c800, 0x00019000,
59 0x0001e800, 0x0001c000, 0x00018000, 0x0001f000,
60 0x0001b000, 0x0001d800, 0x00018800, 0x0001e000,
61 0x0001b800, 0x0001f800, 0x0001a800, 0x00019800,
62 0x00042000, 0x00045000, 0x00044800, 0x00041000,
63 0x00046800, 0x00044000, 0x00040000, 0x00047000,
64 0x00043000, 0x00045800, 0x00040800, 0x00046000,
65 0x00043800, 0x00047800, 0x00042800, 0x00041800,
66 0x0000a000, 0x0000d000, 0x0000c800, 0x00009000,
67 0x0000e800, 0x0000c000, 0x00008000, 0x0000f000,
68 0x0000b000, 0x0000d800, 0x00008800, 0x0000e000,
69 0x0000b800, 0x0000f800, 0x0000a800, 0x00009800,
70 0x00002000, 0x00005000, 0x00004800, 0x00001000,
71 0x00006800, 0x00004000, 0x00000000, 0x00007000,
72 0x00003000, 0x00005800, 0x00000800, 0x00006000,
73 0x00003800, 0x00007800, 0x00002800, 0x00001800,
74 0x0003a000, 0x0003d000, 0x0003c800, 0x00039000,
75 0x0003e800, 0x0003c000, 0x00038000, 0x0003f000,
76 0x0003b000, 0x0003d800, 0x00038800, 0x0003e000,
77 0x0003b800, 0x0003f800, 0x0003a800, 0x00039800,
78 0x0002a000, 0x0002d000, 0x0002c800, 0x00029000,
79 0x0002e800, 0x0002c000, 0x00028000, 0x0002f000,
80 0x0002b000, 0x0002d800, 0x00028800, 0x0002e000,
81 0x0002b800, 0x0002f800, 0x0002a800, 0x00029800,
82 0x0004a000, 0x0004d000, 0x0004c800, 0x00049000,
83 0x0004e800, 0x0004c000, 0x00048000, 0x0004f000,
84 0x0004b000, 0x0004d800, 0x00048800, 0x0004e000,
85 0x0004b800, 0x0004f800, 0x0004a800, 0x00049800,
88 0x03a80000, 0x03c00000, 0x03880000, 0x03e80000,
89 0x03d00000, 0x03980000, 0x03a00000, 0x03900000,
90 0x03f00000, 0x03f80000, 0x03e00000, 0x03b80000,
91 0x03b00000, 0x03800000, 0x03c80000, 0x03d80000,
92 0x06a80000, 0x06c00000, 0x06880000, 0x06e80000,
93 0x06d00000, 0x06980000, 0x06a00000, 0x06900000,
94 0x06f00000, 0x06f80000, 0x06e00000, 0x06b80000,
95 0x06b00000, 0x06800000, 0x06c80000, 0x06d80000,
96 0x05280000, 0x05400000, 0x05080000, 0x05680000,
97 0x05500000, 0x05180000, 0x05200000, 0x05100000,
98 0x05700000, 0x05780000, 0x05600000, 0x05380000,
99 0x05300000, 0x05000000, 0x05480000, 0x05580000,
100 0x00a80000, 0x00c00000, 0x00880000, 0x00e80000,
101 0x00d00000, 0x00980000, 0x00a00000, 0x00900000,
102 0x00f00000, 0x00f80000, 0x00e00000, 0x00b80000,
103 0x00b00000, 0x00800000, 0x00c80000, 0x00d80000,
104 0x00280000, 0x00400000, 0x00080000, 0x00680000,
105 0x00500000, 0x00180000, 0x00200000, 0x00100000,
106 0x00700000, 0x00780000, 0x00600000, 0x00380000,
107 0x00300000, 0x00000000, 0x00480000, 0x00580000,
108 0x04280000, 0x04400000, 0x04080000, 0x04680000,
109 0x04500000, 0x04180000, 0x04200000, 0x04100000,
110 0x04700000, 0x04780000, 0x04600000, 0x04380000,
111 0x04300000, 0x04000000, 0x04480000, 0x04580000,
112 0x04a80000, 0x04c00000, 0x04880000, 0x04e80000,
113 0x04d00000, 0x04980000, 0x04a00000, 0x04900000,
114 0x04f00000, 0x04f80000, 0x04e00000, 0x04b80000,
115 0x04b00000, 0x04800000, 0x04c80000, 0x04d80000,
116 0x07a80000, 0x07c00000, 0x07880000, 0x07e80000,
117 0x07d00000, 0x07980000, 0x07a00000, 0x07900000,
118 0x07f00000, 0x07f80000, 0x07e00000, 0x07b80000,
119 0x07b00000, 0x07800000, 0x07c80000, 0x07d80000,
120 0x07280000, 0x07400000, 0x07080000, 0x07680000,
121 0x07500000, 0x07180000, 0x07200000, 0x07100000,
122 0x07700000, 0x07780000, 0x07600000, 0x07380000,
123 0x07300000, 0x07000000, 0x07480000, 0x07580000,
124 0x02280000, 0x02400000, 0x02080000, 0x02680000,
125 0x02500000, 0x02180000, 0x02200000, 0x02100000,
126 0x02700000, 0x02780000, 0x02600000, 0x02380000,
127 0x02300000, 0x02000000, 0x02480000, 0x02580000,
128 0x03280000, 0x03400000, 0x03080000, 0x03680000,
129 0x03500000, 0x03180000, 0x03200000, 0x03100000,
130 0x03700000, 0x03780000, 0x03600000, 0x03380000,
131 0x03300000, 0x03000000, 0x03480000, 0x03580000,
132 0x06280000, 0x06400000, 0x06080000, 0x06680000,
133 0x06500000, 0x06180000, 0x06200000, 0x06100000,
134 0x06700000, 0x06780000, 0x06600000, 0x06380000,
135 0x06300000, 0x06000000, 0x06480000, 0x06580000,
136 0x05a80000, 0x05c00000, 0x05880000, 0x05e80000,
137 0x05d00000, 0x05980000, 0x05a00000, 0x05900000,
138 0x05f00000, 0x05f80000, 0x05e00000, 0x05b80000,
139 0x05b00000, 0x05800000, 0x05c80000, 0x05d80000,
140 0x01280000, 0x01400000, 0x01080000, 0x01680000,
141 0x01500000, 0x01180000, 0x01200000, 0x01100000,
142 0x01700000, 0x01780000, 0x01600000, 0x01380000,
143 0x01300000, 0x01000000, 0x01480000, 0x01580000,
144 0x02a80000, 0x02c00000, 0x02880000, 0x02e80000,
145 0x02d00000, 0x02980000, 0x02a00000, 0x02900000,
146 0x02f00000, 0x02f80000, 0x02e00000, 0x02b80000,
147 0x02b00000, 0x02800000, 0x02c80000, 0x02d80000,
148 0x01a80000, 0x01c00000, 0x01880000, 0x01e80000,
149 0x01d00000, 0x01980000, 0x01a00000, 0x01900000,
150 0x01f00000, 0x01f80000, 0x01e00000, 0x01b80000,
151 0x01b00000, 0x01800000, 0x01c80000, 0x01d80000,
154 0x30000002, 0x60000002, 0x38000002, 0x08000002,
155 0x28000002, 0x78000002, 0x68000002, 0x40000002,
156 0x20000002, 0x50000002, 0x48000002, 0x70000002,
157 0x00000002, 0x18000002, 0x58000002, 0x10000002,
158 0xb0000005, 0xe0000005, 0xb8000005, 0x88000005,
159 0xa8000005, 0xf8000005, 0xe8000005, 0xc0000005,
160 0xa0000005, 0xd0000005, 0xc8000005, 0xf0000005,
161 0x80000005, 0x98000005, 0xd8000005, 0x90000005,
162 0x30000005, 0x60000005, 0x38000005, 0x08000005,
163 0x28000005, 0x78000005, 0x68000005, 0x40000005,
164 0x20000005, 0x50000005, 0x48000005, 0x70000005,
165 0x00000005, 0x18000005, 0x58000005, 0x10000005,
166 0x30000000, 0x60000000, 0x38000000, 0x08000000,
167 0x28000000, 0x78000000, 0x68000000, 0x40000000,
168 0x20000000, 0x50000000, 0x48000000, 0x70000000,
169 0x00000000, 0x18000000, 0x58000000, 0x10000000,
170 0xb0000003, 0xe0000003, 0xb8000003, 0x88000003,
171 0xa8000003, 0xf8000003, 0xe8000003, 0xc0000003,
172 0xa0000003, 0xd0000003, 0xc8000003, 0xf0000003,
173 0x80000003, 0x98000003, 0xd8000003, 0x90000003,
174 0x30000001, 0x60000001, 0x38000001, 0x08000001,
175 0x28000001, 0x78000001, 0x68000001, 0x40000001,
176 0x20000001, 0x50000001, 0x48000001, 0x70000001,
177 0x00000001, 0x18000001, 0x58000001, 0x10000001,
178 0xb0000000, 0xe0000000, 0xb8000000, 0x88000000,
179 0xa8000000, 0xf8000000, 0xe8000000, 0xc0000000,
180 0xa0000000, 0xd0000000, 0xc8000000, 0xf0000000,
181 0x80000000, 0x98000000, 0xd8000000, 0x90000000,
182 0xb0000006, 0xe0000006, 0xb8000006, 0x88000006,
183 0xa8000006, 0xf8000006, 0xe8000006, 0xc0000006,
184 0xa0000006, 0xd0000006, 0xc8000006, 0xf0000006,
185 0x80000006, 0x98000006, 0xd8000006, 0x90000006,
186 0xb0000001, 0xe0000001, 0xb8000001, 0x88000001,
187 0xa8000001, 0xf8000001, 0xe8000001, 0xc0000001,
188 0xa0000001, 0xd0000001, 0xc8000001, 0xf0000001,
189 0x80000001, 0x98000001, 0xd8000001, 0x90000001,
190 0x30000003, 0x60000003, 0x38000003, 0x08000003,
191 0x28000003, 0x78000003, 0x68000003, 0x40000003,
192 0x20000003, 0x50000003, 0x48000003, 0x70000003,
193 0x00000003, 0x18000003, 0x58000003, 0x10000003,
194 0x30000004, 0x60000004, 0x38000004, 0x08000004,
195 0x28000004, 0x78000004, 0x68000004, 0x40000004,
196 0x20000004, 0x50000004, 0x48000004, 0x70000004,
197 0x00000004, 0x18000004, 0x58000004, 0x10000004,
198 0xb0000002, 0xe0000002, 0xb8000002, 0x88000002,
199 0xa8000002, 0xf8000002, 0xe8000002, 0xc0000002,
200 0xa0000002, 0xd0000002, 0xc8000002, 0xf0000002,
201 0x80000002, 0x98000002, 0xd8000002, 0x90000002,
202 0xb0000004, 0xe0000004, 0xb8000004, 0x88000004,
203 0xa8000004, 0xf8000004, 0xe8000004, 0xc0000004,
204 0xa0000004, 0xd0000004, 0xc8000004, 0xf0000004,
205 0x80000004, 0x98000004, 0xd8000004, 0x90000004,
206 0x30000006, 0x60000006, 0x38000006, 0x08000006,
207 0x28000006, 0x78000006, 0x68000006, 0x40000006,
208 0x20000006, 0x50000006, 0x48000006, 0x70000006,
209 0x00000006, 0x18000006, 0x58000006, 0x10000006,
210 0xb0000007, 0xe0000007, 0xb8000007, 0x88000007,
211 0xa8000007, 0xf8000007, 0xe8000007, 0xc0000007,
212 0xa0000007, 0xd0000007, 0xc8000007, 0xf0000007,
213 0x80000007, 0x98000007, 0xd8000007, 0x90000007,
214 0x30000007, 0x60000007, 0x38000007, 0x08000007,
215 0x28000007, 0x78000007, 0x68000007, 0x40000007,
216 0x20000007, 0x50000007, 0x48000007, 0x70000007,
217 0x00000007, 0x18000007, 0x58000007, 0x10000007,
220 0x000000e8, 0x000000d8, 0x000000a0, 0x00000088,
221 0x00000098, 0x000000f8, 0x000000a8, 0x000000c8,
222 0x00000080, 0x000000d0, 0x000000f0, 0x000000b8,
223 0x000000b0, 0x000000c0, 0x00000090, 0x000000e0,
224 0x000007e8, 0x000007d8, 0x000007a0, 0x00000788,
225 0x00000798, 0x000007f8, 0x000007a8, 0x000007c8,
226 0x00000780, 0x000007d0, 0x000007f0, 0x000007b8,
227 0x000007b0, 0x000007c0, 0x00000790, 0x000007e0,
228 0x000006e8, 0x000006d8, 0x000006a0, 0x00000688,
229 0x00000698, 0x000006f8, 0x000006a8, 0x000006c8,
230 0x00000680, 0x000006d0, 0x000006f0, 0x000006b8,
231 0x000006b0, 0x000006c0, 0x00000690, 0x000006e0,
232 0x00000068, 0x00000058, 0x00000020, 0x00000008,
233 0x00000018, 0x00000078, 0x00000028, 0x00000048,
234 0x00000000, 0x00000050, 0x00000070, 0x00000038,
235 0x00000030, 0x00000040, 0x00000010, 0x00000060,
236 0x000002e8, 0x000002d8, 0x000002a0, 0x00000288,
237 0x00000298, 0x000002f8, 0x000002a8, 0x000002c8,
238 0x00000280, 0x000002d0, 0x000002f0, 0x000002b8,
239 0x000002b0, 0x000002c0, 0x00000290, 0x000002e0,
240 0x000003e8, 0x000003d8, 0x000003a0, 0x00000388,
241 0x00000398, 0x000003f8, 0x000003a8, 0x000003c8,
242 0x00000380, 0x000003d0, 0x000003f0, 0x000003b8,
243 0x000003b0, 0x000003c0, 0x00000390, 0x000003e0,
244 0x00000568, 0x00000558, 0x00000520, 0x00000508,
245 0x00000518, 0x00000578, 0x00000528, 0x00000548,
246 0x00000500, 0x00000550, 0x00000570, 0x00000538,
247 0x00000530, 0x00000540, 0x00000510, 0x00000560,
248 0x00000268, 0x00000258, 0x00000220, 0x00000208,
249 0x00000218, 0x00000278, 0x00000228, 0x00000248,
250 0x00000200, 0x00000250, 0x00000270, 0x00000238,
251 0x00000230, 0x00000240, 0x00000210, 0x00000260,
252 0x000004e8, 0x000004d8, 0x000004a0, 0x00000488,
253 0x00000498, 0x000004f8, 0x000004a8, 0x000004c8,
254 0x00000480, 0x000004d0, 0x000004f0, 0x000004b8,
255 0x000004b0, 0x000004c0, 0x00000490, 0x000004e0,
256 0x00000168, 0x00000158, 0x00000120, 0x00000108,
257 0x00000118, 0x00000178, 0x00000128, 0x00000148,
258 0x00000100, 0x00000150, 0x00000170, 0x00000138,
259 0x00000130, 0x00000140, 0x00000110, 0x00000160,
260 0x000001e8, 0x000001d8, 0x000001a0, 0x00000188,
261 0x00000198, 0x000001f8, 0x000001a8, 0x000001c8,
262 0x00000180, 0x000001d0, 0x000001f0, 0x000001b8,
263 0x000001b0, 0x000001c0, 0x00000190, 0x000001e0,
264 0x00000768, 0x00000758, 0x00000720, 0x00000708,
265 0x00000718, 0x00000778, 0x00000728, 0x00000748,
266 0x00000700, 0x00000750, 0x00000770, 0x00000738,
267 0x00000730, 0x00000740, 0x00000710, 0x00000760,
268 0x00000368, 0x00000358, 0x00000320, 0x00000308,
269 0x00000318, 0x00000378, 0x00000328, 0x00000348,
270 0x00000300, 0x00000350, 0x00000370, 0x00000338,
271 0x00000330, 0x00000340, 0x00000310, 0x00000360,
272 0x000005e8, 0x000005d8, 0x000005a0, 0x00000588,
273 0x00000598, 0x000005f8, 0x000005a8, 0x000005c8,
274 0x00000580, 0x000005d0, 0x000005f0, 0x000005b8,
275 0x000005b0, 0x000005c0, 0x00000590, 0x000005e0,
276 0x00000468, 0x00000458, 0x00000420, 0x00000408,
277 0x00000418, 0x00000478, 0x00000428, 0x00000448,
278 0x00000400, 0x00000450, 0x00000470, 0x00000438,
279 0x00000430, 0x00000440, 0x00000410, 0x00000460,
280 0x00000668, 0x00000658, 0x00000620, 0x00000608,
281 0x00000618, 0x00000678, 0x00000628, 0x00000648,
282 0x00000600, 0x00000650, 0x00000670, 0x00000638,
283 0x00000630, 0x00000640, 0x00000610, 0x00000660,
288 #define BOX(i,n,S) (S)[(n)][(i)]
290 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
292 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
294 #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])
295 #elif VECT_SIZE == 16
296 #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])
299 #define _round(k1,k2,tbl) \
303 l ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
304 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
305 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
306 BOX (((t >> 24) & 0xff), 3, tbl); \
308 r ^= BOX (((t >> 0) & 0xff), 0, tbl) ^ \
309 BOX (((t >> 8) & 0xff), 1, tbl) ^ \
310 BOX (((t >> 16) & 0xff), 2, tbl) ^ \
311 BOX (((t >> 24) & 0xff), 3, tbl); \
314 #define R(k,h,s,i,t) \
320 _round (k[0], k[1], t); \
321 _round (k[2], k[3], t); \
322 _round (k[4], k[5], t); \
323 _round (k[6], k[7], t); \
324 _round (k[0], k[1], t); \
325 _round (k[2], k[3], t); \
326 _round (k[4], k[5], t); \
327 _round (k[6], k[7], t); \
328 _round (k[0], k[1], t); \
329 _round (k[2], k[3], t); \
330 _round (k[4], k[5], t); \
331 _round (k[6], k[7], t); \
332 _round (k[7], k[6], t); \
333 _round (k[5], k[4], t); \
334 _round (k[3], k[2], t); \
335 _round (k[1], k[0], t); \
341 w[0] = u[0] ^ v[0]; \
342 w[1] = u[1] ^ v[1]; \
343 w[2] = u[2] ^ v[2]; \
344 w[3] = u[3] ^ v[3]; \
345 w[4] = u[4] ^ v[4]; \
346 w[5] = u[5] ^ v[5]; \
347 w[6] = u[6] ^ v[6]; \
351 k[0] = ((w[0] & 0x000000ff) << 0) \
352 | ((w[2] & 0x000000ff) << 8) \
353 | ((w[4] & 0x000000ff) << 16) \
354 | ((w[6] & 0x000000ff) << 24); \
355 k[1] = ((w[0] & 0x0000ff00) >> 8) \
356 | ((w[2] & 0x0000ff00) >> 0) \
357 | ((w[4] & 0x0000ff00) << 8) \
358 | ((w[6] & 0x0000ff00) << 16); \
359 k[2] = ((w[0] & 0x00ff0000) >> 16) \
360 | ((w[2] & 0x00ff0000) >> 8) \
361 | ((w[4] & 0x00ff0000) << 0) \
362 | ((w[6] & 0x00ff0000) << 8); \
363 k[3] = ((w[0] & 0xff000000) >> 24) \
364 | ((w[2] & 0xff000000) >> 16) \
365 | ((w[4] & 0xff000000) >> 8) \
366 | ((w[6] & 0xff000000) >> 0); \
367 k[4] = ((w[1] & 0x000000ff) << 0) \
368 | ((w[3] & 0x000000ff) << 8) \
369 | ((w[5] & 0x000000ff) << 16) \
370 | ((w[7] & 0x000000ff) << 24); \
371 k[5] = ((w[1] & 0x0000ff00) >> 8) \
372 | ((w[3] & 0x0000ff00) >> 0) \
373 | ((w[5] & 0x0000ff00) << 8) \
374 | ((w[7] & 0x0000ff00) << 16); \
375 k[6] = ((w[1] & 0x00ff0000) >> 16) \
376 | ((w[3] & 0x00ff0000) >> 8) \
377 | ((w[5] & 0x00ff0000) << 0) \
378 | ((w[7] & 0x00ff0000) << 8); \
379 k[7] = ((w[1] & 0xff000000) >> 24) \
380 | ((w[3] & 0xff000000) >> 16) \
381 | ((w[5] & 0xff000000) >> 8) \
382 | ((w[7] & 0xff000000) >> 0);
419 x[0] ^= 0xff00ff00; \
420 x[1] ^= 0xff00ff00; \
421 x[2] ^= 0x00ff00ff; \
422 x[3] ^= 0x00ff00ff; \
423 x[4] ^= 0x00ffff00; \
424 x[5] ^= 0xff0000ff; \
425 x[6] ^= 0x000000ff; \
428 #define SHIFT12(u,m,s) \
429 u[0] = m[0] ^ s[6]; \
430 u[1] = m[1] ^ s[7]; \
431 u[2] = m[2] ^ (s[0] << 16) \
433 ^ (s[0] & 0x0000ffff) \
434 ^ (s[1] & 0x0000ffff) \
439 ^ (s[7] & 0xffff0000) \
441 u[3] = m[3] ^ (s[0] & 0x0000ffff) \
443 ^ (s[1] & 0x0000ffff) \
452 ^ (s[7] & 0x0000ffff) \
455 u[4] = m[4] ^ (s[0] & 0xffff0000) \
458 ^ (s[1] & 0xffff0000) \
467 ^ (s[7] & 0x0000ffff) \
470 u[5] = m[5] ^ (s[0] << 16) \
472 ^ (s[0] & 0xffff0000) \
473 ^ (s[1] & 0x0000ffff) \
483 ^ (s[7] & 0xffff0000) \
499 u[7] = m[7] ^ (s[0] & 0xffff0000) \
501 ^ (s[1] & 0x0000ffff) \
510 ^ (s[7] & 0x0000ffff) \
514 #define SHIFT16(h,v,u) \
515 v[0] = h[0] ^ (u[1] << 16) \
517 v[1] = h[1] ^ (u[2] << 16) \
519 v[2] = h[2] ^ (u[3] << 16) \
521 v[3] = h[3] ^ (u[4] << 16) \
523 v[4] = h[4] ^ (u[5] << 16) \
525 v[5] = h[5] ^ (u[6] << 16) \
527 v[6] = h[6] ^ (u[7] << 16) \
529 v[7] = h[7] ^ (u[0] & 0xffff0000) \
532 ^ (u[1] & 0xffff0000) \
535 ^ (u[7] & 0xffff0000);
537 #define SHIFT61(h,v) \
538 h[0] = (v[0] & 0xffff0000) \
542 ^ (v[1] & 0xffff0000) \
551 ^ (v[7] & 0x0000ffff); \
552 h[1] = (v[0] << 16) \
554 ^ (v[0] & 0xffff0000) \
555 ^ (v[1] & 0x0000ffff) \
563 ^ (v[7] & 0xffff0000) \
565 h[2] = (v[0] & 0x0000ffff) \
569 ^ (v[1] & 0xffff0000) \
577 ^ (v[7] & 0x0000ffff) \
580 h[3] = (v[0] << 16) \
582 ^ (v[0] & 0xffff0000) \
583 ^ (v[1] & 0xffff0000) \
593 ^ (v[7] & 0x0000ffff) \
595 h[4] = (v[0] >> 16) \
609 h[5] = (v[0] << 16) \
610 ^ (v[0] & 0xffff0000) \
613 ^ (v[1] & 0xffff0000) \
627 ^ (v[7] & 0xffff0000); \
659 #define PASS0(h,s,u,v,t) \
670 #define PASS2(h,s,u,v,t) \
682 #define PASS4(h,s,u,v,t) \
693 #define PASS6(h,s,u,v,t) \
702 __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)
708 const u32 gid = get_global_id (0);
709 const u32 lid = get_local_id (0);
710 const u32 lsz = get_local_size (0);
716 __local u32 s_tables[4][256];
718 for (u32 i = lid; i < 256; i += lsz)
720 s_tables[0][i] = c_tables[0][i];
721 s_tables[1][i] = c_tables[1][i];
722 s_tables[2][i] = c_tables[2][i];
723 s_tables[3][i] = c_tables[3][i];
726 barrier (CLK_LOCAL_MEM_FENCE);
728 if (gid >= gid_max) return;
737 pw_buf0[0] = pws[gid].i[0];
738 pw_buf0[1] = pws[gid].i[1];
739 pw_buf0[2] = pws[gid].i[2];
740 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_l_len = pws[gid].pw_len;
752 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
754 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
756 const u32x pw_len = pw_l_len + pw_r_len;
759 * concat password candidate
762 u32x wordl0[4] = { 0 };
763 u32x wordl1[4] = { 0 };
764 u32x wordl2[4] = { 0 };
765 u32x wordl3[4] = { 0 };
767 wordl0[0] = pw_buf0[0];
768 wordl0[1] = pw_buf0[1];
769 wordl0[2] = pw_buf0[2];
770 wordl0[3] = pw_buf0[3];
771 wordl1[0] = pw_buf1[0];
772 wordl1[1] = pw_buf1[1];
773 wordl1[2] = pw_buf1[2];
774 wordl1[3] = pw_buf1[3];
776 u32x wordr0[4] = { 0 };
777 u32x wordr1[4] = { 0 };
778 u32x wordr2[4] = { 0 };
779 u32x wordr3[4] = { 0 };
781 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
782 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
783 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
784 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
785 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
786 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
787 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
788 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
790 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
792 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
796 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
802 w0[0] = wordl0[0] | wordr0[0];
803 w0[1] = wordl0[1] | wordr0[1];
804 w0[2] = wordl0[2] | wordr0[2];
805 w0[3] = wordl0[3] | wordr0[3];
806 w1[0] = wordl1[0] | wordr1[0];
807 w1[1] = wordl1[1] | wordr1[1];
808 w1[2] = wordl1[2] | wordr1[2];
809 w1[3] = wordl1[3] | wordr1[3];
850 state_m[0] = state[0];
851 state_m[1] = state[1];
852 state_m[2] = state[2];
853 state_m[3] = state[3];
854 state_m[4] = state[4];
855 state_m[5] = state[5];
856 state_m[6] = state[6];
857 state_m[7] = state[7];
870 //if (pw_len > 0) // not really SIMD compatible
872 PASS0 (state, tmp, state_m, data_m, s_tables);
873 PASS2 (state, tmp, state_m, data_m, s_tables);
874 PASS4 (state, tmp, state_m, data_m, s_tables);
875 PASS6 (state, tmp, state_m, data_m, s_tables);
877 SHIFT12 (state_m, data, tmp);
878 SHIFT16 (state, data_m, state_m);
879 SHIFT61 (state, data_m);
882 data[0] = pw_len * 8;
893 state_m[0] = state[0];
894 state_m[1] = state[1];
895 state_m[2] = state[2];
896 state_m[3] = state[3];
897 state_m[4] = state[4];
898 state_m[5] = state[5];
899 state_m[6] = state[6];
900 state_m[7] = state[7];
911 PASS0 (state, tmp, state_m, data_m, s_tables);
912 PASS2 (state, tmp, state_m, data_m, s_tables);
913 PASS4 (state, tmp, state_m, data_m, s_tables);
914 PASS6 (state, tmp, state_m, data_m, s_tables);
916 SHIFT12 (state_m, data, tmp);
917 SHIFT16 (state, data_m, state_m);
918 SHIFT61 (state, data_m);
931 state_m[0] = state[0];
932 state_m[1] = state[1];
933 state_m[2] = state[2];
934 state_m[3] = state[3];
935 state_m[4] = state[4];
936 state_m[5] = state[5];
937 state_m[6] = state[6];
938 state_m[7] = state[7];
949 PASS0 (state, tmp, state_m, data_m, s_tables);
950 PASS2 (state, tmp, state_m, data_m, s_tables);
951 PASS4 (state, tmp, state_m, data_m, s_tables);
952 PASS6 (state, tmp, state_m, data_m, s_tables);
954 SHIFT12 (state_m, data, tmp);
955 SHIFT16 (state, data_m, state_m);
956 SHIFT61 (state, data_m);
960 COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
964 __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)
968 __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)
972 __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)
978 const u32 gid = get_global_id (0);
979 const u32 lid = get_local_id (0);
980 const u32 lsz = get_local_size (0);
986 __local u32 s_tables[4][256];
988 for (u32 i = lid; i < 256; i += lsz)
990 s_tables[0][i] = c_tables[0][i];
991 s_tables[1][i] = c_tables[1][i];
992 s_tables[2][i] = c_tables[2][i];
993 s_tables[3][i] = c_tables[3][i];
996 barrier (CLK_LOCAL_MEM_FENCE);
998 if (gid >= gid_max) return;
1007 pw_buf0[0] = pws[gid].i[0];
1008 pw_buf0[1] = pws[gid].i[1];
1009 pw_buf0[2] = pws[gid].i[2];
1010 pw_buf0[3] = pws[gid].i[3];
1011 pw_buf1[0] = pws[gid].i[4];
1012 pw_buf1[1] = pws[gid].i[5];
1013 pw_buf1[2] = pws[gid].i[6];
1014 pw_buf1[3] = pws[gid].i[7];
1016 const u32 pw_l_len = pws[gid].pw_len;
1022 const u32 search[4] =
1024 digests_buf[digests_offset].digest_buf[DGST_R0],
1025 digests_buf[digests_offset].digest_buf[DGST_R1],
1026 digests_buf[digests_offset].digest_buf[DGST_R2],
1027 digests_buf[digests_offset].digest_buf[DGST_R3]
1034 for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
1036 const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
1038 const u32x pw_len = pw_l_len + pw_r_len;
1041 * concat password candidate
1044 u32x wordl0[4] = { 0 };
1045 u32x wordl1[4] = { 0 };
1046 u32x wordl2[4] = { 0 };
1047 u32x wordl3[4] = { 0 };
1049 wordl0[0] = pw_buf0[0];
1050 wordl0[1] = pw_buf0[1];
1051 wordl0[2] = pw_buf0[2];
1052 wordl0[3] = pw_buf0[3];
1053 wordl1[0] = pw_buf1[0];
1054 wordl1[1] = pw_buf1[1];
1055 wordl1[2] = pw_buf1[2];
1056 wordl1[3] = pw_buf1[3];
1058 u32x wordr0[4] = { 0 };
1059 u32x wordr1[4] = { 0 };
1060 u32x wordr2[4] = { 0 };
1061 u32x wordr3[4] = { 0 };
1063 wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
1064 wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
1065 wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
1066 wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
1067 wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
1068 wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
1069 wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
1070 wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
1072 if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
1074 switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
1078 switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
1084 w0[0] = wordl0[0] | wordr0[0];
1085 w0[1] = wordl0[1] | wordr0[1];
1086 w0[2] = wordl0[2] | wordr0[2];
1087 w0[3] = wordl0[3] | wordr0[3];
1088 w1[0] = wordl1[0] | wordr1[0];
1089 w1[1] = wordl1[1] | wordr1[1];
1090 w1[2] = wordl1[2] | wordr1[2];
1091 w1[3] = wordl1[3] | wordr1[3];
1118 state[ 8] = data[0];
1119 state[ 9] = data[1];
1120 state[10] = data[2];
1121 state[11] = data[3];
1122 state[12] = data[4];
1123 state[13] = data[5];
1124 state[14] = data[6];
1125 state[15] = data[7];
1132 state_m[0] = state[0];
1133 state_m[1] = state[1];
1134 state_m[2] = state[2];
1135 state_m[3] = state[3];
1136 state_m[4] = state[4];
1137 state_m[5] = state[5];
1138 state_m[6] = state[6];
1139 state_m[7] = state[7];
1141 data_m[0] = data[0];
1142 data_m[1] = data[1];
1143 data_m[2] = data[2];
1144 data_m[3] = data[3];
1145 data_m[4] = data[4];
1146 data_m[5] = data[5];
1147 data_m[6] = data[6];
1148 data_m[7] = data[7];
1152 //if (pw_len > 0) // not really SIMD compatible
1154 PASS0 (state, tmp, state_m, data_m, s_tables);
1155 PASS2 (state, tmp, state_m, data_m, s_tables);
1156 PASS4 (state, tmp, state_m, data_m, s_tables);
1157 PASS6 (state, tmp, state_m, data_m, s_tables);
1159 SHIFT12 (state_m, data, tmp);
1160 SHIFT16 (state, data_m, state_m);
1161 SHIFT61 (state, data_m);
1164 data[0] = pw_len * 8;
1175 state_m[0] = state[0];
1176 state_m[1] = state[1];
1177 state_m[2] = state[2];
1178 state_m[3] = state[3];
1179 state_m[4] = state[4];
1180 state_m[5] = state[5];
1181 state_m[6] = state[6];
1182 state_m[7] = state[7];
1184 data_m[0] = data[0];
1185 data_m[1] = data[1];
1186 data_m[2] = data[2];
1187 data_m[3] = data[3];
1188 data_m[4] = data[4];
1189 data_m[5] = data[5];
1190 data_m[6] = data[6];
1191 data_m[7] = data[7];
1193 PASS0 (state, tmp, state_m, data_m, s_tables);
1194 PASS2 (state, tmp, state_m, data_m, s_tables);
1195 PASS4 (state, tmp, state_m, data_m, s_tables);
1196 PASS6 (state, tmp, state_m, data_m, s_tables);
1198 SHIFT12 (state_m, data, tmp);
1199 SHIFT16 (state, data_m, state_m);
1200 SHIFT61 (state, data_m);
1204 data[0] = state[ 8];
1205 data[1] = state[ 9];
1206 data[2] = state[10];
1207 data[3] = state[11];
1208 data[4] = state[12];
1209 data[5] = state[13];
1210 data[6] = state[14];
1211 data[7] = state[15];
1213 state_m[0] = state[0];
1214 state_m[1] = state[1];
1215 state_m[2] = state[2];
1216 state_m[3] = state[3];
1217 state_m[4] = state[4];
1218 state_m[5] = state[5];
1219 state_m[6] = state[6];
1220 state_m[7] = state[7];
1222 data_m[0] = data[0];
1223 data_m[1] = data[1];
1224 data_m[2] = data[2];
1225 data_m[3] = data[3];
1226 data_m[4] = data[4];
1227 data_m[5] = data[5];
1228 data_m[6] = data[6];
1229 data_m[7] = data[7];
1231 PASS0 (state, tmp, state_m, data_m, s_tables);
1232 PASS2 (state, tmp, state_m, data_m, s_tables);
1233 PASS4 (state, tmp, state_m, data_m, s_tables);
1234 PASS6 (state, tmp, state_m, data_m, s_tables);
1236 SHIFT12 (state_m, data, tmp);
1237 SHIFT16 (state, data_m, state_m);
1238 SHIFT61 (state, data_m);
1242 COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
1246 __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)
1250 __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)