?  search:
summary | log | graphiclog1 | graphiclog2 | commit | commitdiff | tree | refs | edit | fork
blame | history | raw | HEAD
hf mf fchk - output style
[RRG-proxmark3.git] / tools / hitag2crack / crack5opencl / ht2crack5opencl_kernel.cl
blob39c5737f27745df7ee8040b00f3825f46bc2fa0d
1 /* ht2crack5opencl_kernel.cl
2 * -------------------------
3 * This code is heavily based on crack5gpu implementation.
4 *
5 * Additional changes done by Gabriele 'matrix' Gristina <gabriele.gristina@gmail.com>
6 * - generic code optimizations
7 * - using local memory for keystream, if enabled by OpenCL host engine
8 * - added the two macros bs_res () and bs_res_lut_1, used during the generation of intermediate results
9 * - split lut3 function to some variants (tentative to reduce registers usage)
10 * - add support for devices without lop3.b32 instruction (for ! NVIDIA platforms/devices, like Intel and Apple CPU/GPU, not tested on ADM)
11 * - add HITAG2 routine to perform key verification, if enabled
12 * - using local memory for uid, aR2, nR1, nR2 (if HITAG2 routine is enabled)
13 */
14
15 #define MAX_BITSLICES 32
16 #define KEYSTREAM_LENGTH 32
17
18 typedef uint bitslice_t __attribute__((aligned(MAX_BITSLICES / 8)));
19
20 #ifndef HAVE_LOP3
21 // (0xf0&(0xcc^0xaa))
22 inline static uint lut3_0x60(uint a, uint b, uint c) {
23 const uint r = b ^ c;
24 return a & r;
25 }
26
27 // (~((0xf0|0xcc)^0xaa))
28 inline static uint lut3_0xa9(uint a, uint b, uint c) {
29 const uint r = a | b;
30 return ~(r ^ c);
31 }
32
33 // (~((0xf0|0xcc|0xaa)))
34 inline static uint lut3_0x01(uint a, uint b, uint c) {
35 const uint r = a | b;
36 return ~(r | c);
37 }
38
39 // (((0xf0|0xcc)&0xaa))
40 inline static uint lut3_0xa8(uint a, uint b, uint c) {
41 const uint r = a | b;
42 return r & c;
43 }
44
45 // (((0xf0|0xcc)&0xaa)^0xcc)
46 inline static uint lut3_0x64(uint a, uint b, uint c) {
47 const uint r = (a | b) & c;
48 return r ^ b;
49 }
50
51 // (0xf0^0xaa^0xcc)
52 inline static uint lut3_0x96(uint a, uint b, uint c) {
53 const uint r = a ^ b;
54 return r ^ c;
55 }
56
57 // (((0xf0^0xcc)|0xaa))
58 inline static uint lut3_0xbe(uint a, uint b, uint c) {
59 const uint r = a ^ b;
60 return r | c;
61 }
62
63 // (~((0xf0^0xcc)&(0xaa^0xcc)))
64 inline static uint lut3_0xdb(uint a, uint b, uint c) {
65 const uint r = a ^ b;
66 const uint r2 = c ^ b;
67 return ~(r & r2);
68 }
69 /*
70 // (0xf0|(0xcc&0xaa))
71 inline static uint lut3_0xf8(uint a, uint b, uint c) {
72 const uint r = b & c;
73 return a | r;
74 }
75 */
76 /*
77 // (0xf0|(0xcc&(0x01)))
78 inline static uint lut3_0xf8_0x1(uint a, uint b) {
79 const uint r = b & 0x1;
80 return a | r;
81 }
82 */
83
84 #ifdef WITH_HITAG2_FULL
85 // (0xf0|(0xcc&(0xC)))
86 inline static uint lut3_0xf8_0xC(uint a, uint b) {
87 const uint r = b & 0xC;
88 return a | r;
89 }
90
91 // (0xf0|0xcc|0xaa)
92 inline static uint lut3_0xfe(uint a, uint b, uint c) {
93 const uint r = a | b;
94 return r | c;
95 }
96 #endif // WITH_HITAG2_FULL
97
98 #else // HAVE_LOP3
99
100 inline static uint lut3_0x01(uint a, uint b, uint c) {
101 uint r;
102 asm("lop3.b32 %0, %1, %2, %3, 0x01;" : "=r"(r): "r"(a), "r"(b), "r"(c));
103 return r;
104 }
105
106 inline static uint lut3_0x60(uint a, uint b, uint c) {
107 uint r;
108 asm("lop3.b32 %0, %1, %2, %3, 0x60;" : "=r"(r): "r"(a), "r"(b), "r"(c));
109 return r;
110 }
111
112 inline static uint lut3_0x64(uint a, uint b, uint c) {
113 uint r;
114 asm("lop3.b32 %0, %1, %2, %3, 0x64;" : "=r"(r): "r"(a), "r"(b), "r"(c));
115 return r;
116 }
117
118 inline static uint lut3_0x96(uint a, uint b, uint c) {
119 uint r;
120 asm("lop3.b32 %0, %1, %2, %3, 0x96;" : "=r"(r): "r"(a), "r"(b), "r"(c));
121 return r;
122 }
123
124 inline static uint lut3_0xa8(uint a, uint b, uint c) {
125 uint r;
126 asm("lop3.b32 %0, %1, %2, %3, 0xa8;" : "=r"(r): "r"(a), "r"(b), "r"(c));
127 return r;
128 }
129
130 inline static uint lut3_0xa9(uint a, uint b, uint c) {
131 uint r;
132 asm("lop3.b32 %0, %1, %2, %3, 0xa9;" : "=r"(r): "r"(a), "r"(b), "r"(c));
133 return r;
134 }
135
136 inline static uint lut3_0xbe(uint a, uint b, uint c) {
137 uint r;
138 asm("lop3.b32 %0, %1, %2, %3, 0xbe;" : "=r"(r): "r"(a), "r"(b), "r"(c));
139 return r;
140 }
141
142 inline static uint lut3_0xdb(uint a, uint b, uint c) {
143 uint r;
144 asm("lop3.b32 %0, %1, %2, %3, 0xdb;" : "=r"(r): "r"(a), "r"(b), "r"(c));
145 return r;
146 }
147
148 /*
149 inline static uint lut3_0xf8(uint a, uint b, uint c) {
150 uint r;
151 asm("lop3.b32 %0, %1, %2, %3, 0xf8;" : "=r"(r): "r"(a), "r"(b), "r"(c));
152 return r;
153 }
154 */
155 /*
156 inline static uint lut3_0xf8_0x1(uint a, uint b) {
157 uint r;
158 asm("lop3.b32 %0, %1, %2, 0x1, 0xf8;" : "=r"(r): "r"(a), "r"(b));
159 return r;
160 }
161 */
162
163 #ifdef WITH_HITAG2_FULL
164 inline static uint lut3_0xf8_0xC(uint a, uint b) {
165 uint r;
166 asm("lop3.b32 %0, %1, %2, 0xC, 0xf8;" : "=r"(r): "r"(a), "r"(b));
167 return r;
168 }
169
170 inline static uint lut3_0xfe(uint a, uint b, uint c) {
171 uint r;
172 asm("lop3.b32 %0, %1, %2, %3, 0xfe;" : "=r"(r): "r"(a), "r"(b), "r"(c));
173 return r;
174 }
175 #endif // WITH_HITAG2_FULL
176
177 #endif // HAVE_LOP3
178
179 #define f_a_bs(a,b,c,d) (lut3_0xa9(a,d,lut3_0x64(a,b,c))) // 2 luts
180 #define f_b_bs(a,b,c,d) (lut3_0xa8(d,c,a^b) ^ lut3_0x01(d,a,b)) // 2 luts, 2 xors
181 #define f_c_bs(a,b,c,d,e) (((lut3_0xdb((lut3_0xbe(c,e,d) & a), b, c)) ^ (lut3_0xbe(d,e,a) & lut3_0xbe(d,b,c))))
182 #define bs_res(a,b,c) (lut3_0x60(a,b,c))
183
184 #define lfsr_bs(i) (lut3_0x96(lut3_0x96(lut3_0x96(state[-2+i+ 0], state[-2+i+ 2], state[-2+i+ 3]), \
185 lut3_0x96(state[-2+i+ 6], state[-2+i+ 7], state[-2+i+ 8]), \
186 lut3_0x96(state[-2+i+16], state[-2+i+22], state[-2+i+23])), \
187 lut3_0x96(state[-2+i+26], state[-2+i+30], state[-2+i+41]), \
188 lut3_0x96(state[-2+i+42], state[-2+i+43], state[-2+i+46])) ^ state[-2+i+47])
189
190 // 46 iterations * 4 ops
191 static inline void bitslice (bitslice_t *restrict b, ulong x) {
192 for (uint i = 0; i < 46; ++i) {
193 b[i] = -(x & 1);
194 x >>= 1;
195 }
196 }
197
198 // don't care about the complexity of this function
199 static inline ulong unbitslice (const bitslice_t *restrict b, const uint s) {
200 const bitslice_t mask = ((bitslice_t) 1) << s;
201 ulong result = 0;
202
203 for (int i = 47; i >= 0; --i) {
204 result <<= 1;
205 result |= (bool)(b[i] & mask);
206 }
207 return result;
208 }
209
210 ///////////////////////////////
211
212 #ifdef WITH_HITAG2_FULL
213
214 /*
215 * Hitag Crypto support macros
216 * These macros reverse the bit order in a byte, or *within* each byte of a
217 * 16 , 32 or 64 bit unsigned integer. (Not across the whole 16 etc bits.)
218 */
219 #define rev8(X) ((((X) >> 7) &1) + (((X) >> 5) &2) + (((X) >> 3) &4) \
220 + (((X) >> 1) &8) + (((X) << 1) &16) + (((X) << 3) &32) \
221 + (((X) << 5) &64) + (((X) << 7) &128) )
222 #define rev16(X) (rev8 (X) + (rev8 (X >> 8) << 8))
223 #define rev32(X) (rev16(X) + (rev16(X >> 16) << 16))
224 #define rev64(X) (rev32(X) + (rev32(X >> 32) << 32))
225
226 // (0xf0|(0xcc&0xaa))
227 #define pickbits2_2_lut(a,b) (lut3_0xf8_0xC(a,b))
228 #define pickbits2_2(S) pickbits2_2_lut( ((S >> 1) & 3) , (S >> 2) )
229
230 // (0xf0|0xcc|0xaa)
231 #define pickbits1_1_2_lut(a,b,c) (lut3_0xfe(a,b,c))
232 #define pickbits1_1_2(S) pickbits1_1_2_lut( ((S >> 7) & 1) , ((S >> 10) & 2) , ((S >> 11) & 0xC) )
233
234 #define pickbits1x4_lut(a,b,c,d) (lut3_0xfe(a,b,c) | d)
235 #define pickbits1x4(S) pickbits1x4_lut( ((S >> 16) & 1) , ((S >> 19) & 2) , ((S >> 20) & 4) , ((S >> 22) & 8) )
236
237 #define pickbits2_1_1_lut(a,b,c) (lut3_0xfe(a,b,c))
238 #define pickbits2_1_1(S) pickbits2_1_1_lut( ((S >> 27) & 3) , ((S >> 28) & 4) , ((S >> 29) & 8) )
239
240 #define pickbits1_2_1_lut(a,b,c) (lut3_0xfe(a,b,c))
241 #define pickbits1_2_1(S) pickbits1_2_1_lut( ((S >> 33) & 1) , ((S >> 41) & 6) , ((S >> 42) & 8) )
242
243 static uint hitag2_crypt (ulong x)
244 {
245 const uint ht2_function4a = 0x2C79; // 0010 1100 0111 1001
246 const uint ht2_function4b = 0x6671; // 0110 0110 0111 0001
247 const uint ht2_function5c = 0x7907287B; // 0111 1001 0000 0111 0010 1000 0111 1011
248
249 uint bitindex;
250 bitindex = (ht2_function4a >> pickbits2_2(x)) & 1;
251 bitindex |= ((ht2_function4b << 1) >> pickbits1_1_2(x)) & 0x02;
252 bitindex |= ((ht2_function4b << 2) >> pickbits1x4(x)) & 0x04;
253 bitindex |= ((ht2_function4b << 3) >> pickbits2_1_1(x)) & 0x08;
254 bitindex |= ((ht2_function4a << 4) >> pickbits1_2_1(x)) & 0x10;
255
256 return (ht2_function5c >> bitindex) & 1;
257 }
258
259 static void hitag2_init2 (ulong *shiftreg, ulong *lfsr, const ulong sharedkey, const uint serialnum, uint initvector)
260 {
261 ulong state = ((sharedkey & 0xFFFF) << 32) | serialnum;
262
263 initvector ^= (uint)(sharedkey >> 16);
264 state |= (ulong) initvector << 48;
265
266 initvector >>= 16;
267 state >>= 1;
268
269 for (uint x = 0; x < 16; x++) state = (state >> 1) ^ (ulong) hitag2_crypt (state) << 46;
270
271 state |= (ulong) initvector << 47;
272
273 for (uint x = 0; x < 15; x++) state = (state >> 1) ^ (ulong) hitag2_crypt (state) << 46;
274
275 state ^= (ulong) hitag2_crypt(state) << 47;
276
277 *shiftreg = state;
278
279 // ulong temp = state ^ (state >> 1);
280 // *lfsr = state ^ (state >> 6) ^ (state >> 16) ^ (state >> 26) ^ (state >> 30) ^ (state >> 41) ^
281 // (temp >> 2) ^ (temp >> 7) ^ (temp >> 22) ^ (temp >> 42) ^ (temp >> 46);
282
283 *lfsr = (ulong) state ^
284 lut3_0x96 ((state >> 2), (state >> 3), (state >> 6)) ^
285 lut3_0x96 ((state >> 7), (state >> 8), (state >> 16)) ^
286 lut3_0x96 ((state >> 22), (state >> 23), (state >> 26)) ^
287 lut3_0x96 ((state >> 30), (state >> 41), (state >> 42)) ^
288 lut3_0x96 ((state >> 43), (state >> 46), (state >> 47));
289 }
290
291 static uint hitag2_nstep2 (ulong state, ulong lfsr)
292 {
293 uint result = 0;
294 for (uint nsteps = 32; nsteps > 0; nsteps--)
295 {
296 if (lfsr & 1)
297 {
298 state = (state >> 1) | 0x800000000000;
299 lfsr = (lfsr >> 1) ^ 0xB38083220073;
300 result = (result << 1) | hitag2_crypt (state);
301 }
302 else
303 {
304 state >>= 1;
305 lfsr >>= 1;
306 result = (result << 1) | hitag2_crypt (state);
307 }
308 }
309 return result;
310 }
311
312 inline static int bitn(ulong x, int bit)
313 {
314 const ulong bitmask = 1UL << bit;
315 return (x & bitmask) ? 1 : 0;
316 }
317
318 static int fnR (ulong x)
319 {
320 return (bitn(x, 1) ^ bitn(x, 2) ^ bitn(x, 5) ^ bitn(x, 6) ^ bitn(x, 7) ^
321 bitn(x, 15) ^ bitn(x, 21) ^ bitn(x, 22) ^ bitn(x, 25) ^ bitn(x, 29) ^ bitn(x, 40) ^
322 bitn(x, 41) ^ bitn(x, 42) ^ bitn(x, 45) ^ bitn(x, 46) ^ bitn(x, 47));
323 }
324
325 inline static int fa(unsigned int i) {
326 return bitn(0x2C79, i);
327 }
328
329 inline static int fb(unsigned int i) {
330 return bitn(0x6671, i);
331 }
332
333 static int fnf (ulong s)
334 {
335 const uint x1 = (bitn(s, 2) << 0) | lut3_0x96( (bitn(s, 3) << 1), (bitn(s, 5) << 2), (bitn(s, 6) << 3));
336 const uint x2 = (bitn(s, 8) << 0) | lut3_0x96( (bitn(s, 12) << 1), (bitn(s, 14) << 2), (bitn(s, 15) << 3));
337 const uint x3 = (bitn(s, 17) << 0) | lut3_0x96( (bitn(s, 21) << 1), (bitn(s, 23) << 2), (bitn(s, 26) << 3));
338 const uint x4 = (bitn(s, 28) << 0) | lut3_0x96( (bitn(s, 29) << 1), (bitn(s, 31) << 2), (bitn(s, 33) << 3));
339 const uint x5 = (bitn(s, 34) << 0) | lut3_0x96( (bitn(s, 43) << 1), (bitn(s, 44) << 2), (bitn(s, 46) << 3));
340 const uint x6 = lut3_0x96( (fa(x1) << 0), (fb(x2) << 1), lut3_0x96( (fb(x3) << 2), (fb(x4) << 3), (fa(x5) << 4)));
341
342 return bitn(0x7907287B, x6);
343 }
344
345 #endif // WITH_HITAG2_FULL
346
347 // format this array with 32 bitsliced vectors of ones and zeroes representing the inverted keystream
348
349 __kernel
350 __attribute__((vec_type_hint(bitslice_t)))
351 void find_state(const uint candidate_index_base,
352 __global const ushort *restrict candidates,
353 // __global const ulong *restrict candidates,
354 __global const bitslice_t *restrict _keystream,
355 __global ulong *restrict matches,
356 #ifndef WITH_HITAG2_FULL
357 __global uint *restrict matches_found)
358 #else
359 __global uint *restrict matches_found,
360 __global const uint *restrict _checks)
361 #endif
362 {
363 const size_t gid[2] = { get_global_id(0), get_global_id(1) };
364
365 // if (gid[0] == 0) printf("work-item 1,%u\n", gid[1]);
366
367 #ifdef HAVE_LOCAL_MEMORY
368 const size_t lid = get_local_id(0);
369 const size_t lsize = get_local_size(0);
370 #endif // HAVE_LOCAL_MEMORY
371
372 const uint index = 3 * (candidate_index_base + gid[0]); // dimension 0 should at least keep the execution units saturated - 8k is fine
373
374 const ulong3 c = { candidates[index], candidates[index + 1], candidates[index + 2] };
375
376 const ulong candidate = ( c.x << 32 | c.y << 16 | c.z );
377
378 #ifdef HAVE_LOCAL_MEMORY
379 // store keystream in local memory
380 __local bitslice_t keystream[32];
381
382 for (size_t i = lid; i < 32; i+= lsize) keystream[i] = _keystream[i];
383
384 #ifdef WITH_HITAG2_FULL
385 // store uid, aR2, nR1, nR2 in local memory
386 __local uint checks[4];
387
388 for (uint i = lid; i < 4; i+= lsize) checks[i] = _checks[i];
389 #endif
390
391 // threads syncronization
392 barrier (CLK_LOCAL_MEM_FENCE);
393 #else
394 #define keystream _keystream
395 #define checks _checks
396 #endif // HAVE_LOCAL_MEMORY
397
398 // we never actually set or use the lowest 2 bits the initial state, so we can save 2 bitslices everywhere
399 bitslice_t state[-2 + 48 + KEYSTREAM_LENGTH] = { 0 };
400
401 // set bits 0+2, 0+3, 0+5, 0+6, 0+8, 0+12, 0+14, 0+15, 0+17, 0+21, 0+23, 0+26, 0+28, 0+29, 0+31, 0+33, 0+34, 0+43, 0+44, 0+46
402 // get the 48-bit cipher states as 3 16-bit words from the host memory queue (to save 25% throughput)
403
404 // using 64bit candidate
405 // const uint index = (candidate_index_base + gid[0]); // dimension 0 should at least keep the execution units saturated - 8k is fine
406 // const ulong candidate = candidates[index];
407 // bitslice (&state[-2 + 2], candidate >> 2);
408
409 // set all 48 state bits except the lowest 2
410 bitslice (&state[-2 + 2], candidate);
411
412 // set bits 3, 6, 8, 12, 15
413 state[-2 + 1 + 3] = 0xaaaaaaaa;
414 state[-2 + 1 + 6] = 0xcccccccc;
415 state[-2 + 1 + 8] = 0xf0f0f0f0;
416 state[-2 + 1 + 12] = 0xff00ff00;
417 state[-2 + 1 + 15] = 0xffff0000;
418
419 const uint i1 = gid[1]; // dimension 1 should be 1024
420 state[-2 + 18] = -((bool)(i1 & 0x1));
421 state[-2 + 22] = -((bool)(i1 & 0x2));
422 state[-2 + 24] = -((bool)(i1 & 0x4));
423 state[-2 + 27] = -((bool)(i1 & 0x8));
424 state[-2 + 30] = -((bool)(i1 & 0x10));
425 state[-2 + 32] = -((bool)(i1 & 0x20));
426 state[-2 + 35] = -((bool)(i1 & 0x40));
427 state[-2 + 45] = -((bool)(i1 & 0x80));
428 state[-2 + 47] = -((bool)(i1 & 0x100));
429 state[-2 + 48] = -((bool)(i1 & 0x200)); // guess lfsr output 0
430
431 // 0xfc07fef3f9fe
432 const bitslice_t filter1_0 = f_a_bs(state[-2 + 3], state[-2 + 4], state[-2 + 6], state[-2 + 7]);
433 const bitslice_t filter1_1 = f_b_bs(state[-2 + 9], state[-2 + 13], state[-2 + 15], state[-2 + 16]);
434 const bitslice_t filter1_2 = f_b_bs(state[-2 + 18], state[-2 + 22], state[-2 + 24], state[-2 + 27]);
435 const bitslice_t filter1_3 = f_b_bs(state[-2 + 29], state[-2 + 30], state[-2 + 32], state[-2 + 34]);
436 const bitslice_t filter1_4 = f_a_bs(state[-2 + 35], state[-2 + 44], state[-2 + 45], state[-2 + 47]);
437 const bitslice_t filter1 = f_c_bs(filter1_0, filter1_1, filter1_2, filter1_3, filter1_4);
438
439 const bitslice_t results1 = filter1 ^ keystream[1];
440 if (!results1) return;
441
442 const bitslice_t filter2_0 = f_a_bs(state[-2 + 4], state[-2 + 5], state[-2 + 7], state[-2 + 8]);
443 const bitslice_t filter2_3 = f_b_bs(state[-2 + 30], state[-2 + 31], state[-2 + 33], state[-2 + 35]);
444 const bitslice_t filter3_0 = f_a_bs(state[-2 + 5], state[-2 + 6], state[-2 + 8], state[-2 + 9]);
445 const bitslice_t filter5_2 = f_b_bs(state[-2 + 22], state[-2 + 26], state[-2 + 28], state[-2 + 31]);
446 const bitslice_t filter6_2 = f_b_bs(state[-2 + 23], state[-2 + 27], state[-2 + 29], state[-2 + 32]);
447 const bitslice_t filter7_2 = f_b_bs(state[-2 + 24], state[-2 + 28], state[-2 + 30], state[-2 + 33]);
448 const bitslice_t filter9_1 = f_b_bs(state[-2 + 17], state[-2 + 21], state[-2 + 23], state[-2 + 24]);
449 const bitslice_t filter9_2 = f_b_bs(state[-2 + 26], state[-2 + 30], state[-2 + 32], state[-2 + 35]);
450 const bitslice_t filter10_0 = f_a_bs(state[-2 + 12], state[-2 + 13], state[-2 + 15], state[-2 + 16]);
451 const bitslice_t filter11_0 = f_a_bs(state[-2 + 13], state[-2 + 14], state[-2 + 16], state[-2 + 17]);
452 const bitslice_t filter12_0 = f_a_bs(state[-2 + 14], state[-2 + 15], state[-2 + 17], state[-2 + 18]);
453 const bitslice_t filter14_1 = f_b_bs(state[-2 + 22], state[-2 + 26], state[-2 + 28], state[-2 + 29]);
454 const bitslice_t filter15_1 = f_b_bs(state[-2 + 23], state[-2 + 27], state[-2 + 29], state[-2 + 30]);
455 const bitslice_t filter15_3 = f_b_bs(state[-2 + 43], state[-2 + 44], state[-2 + 46], state[-2 + 48]);
456 const bitslice_t filter16_1 = f_b_bs(state[-2 + 24], state[-2 + 28], state[-2 + 30], state[-2 + 31]);
457
458 for (uint i2 = 0; i2 < 32; i2++) {
459 state[-2 + 10] = -((bool)(i2 & 0x1));
460 state[-2 + 19] = -((bool)(i2 & 0x2));
461 state[-2 + 25] = -((bool)(i2 & 0x4));
462 state[-2 + 36] = -((bool)(i2 & 0x8));
463 state[-2 + 49] = -((bool)(i2 & 0x10)); // guess lfsr output 1
464
465 // 0xfe07fffbfdff
466 const bitslice_t filter2_1 = f_b_bs(state[-2 + 10], state[-2 + 14], state[-2 + 16], state[-2 + 17]);
467 const bitslice_t filter2_2 = f_b_bs(state[-2 + 19], state[-2 + 23], state[-2 + 25], state[-2 + 28]);
468 const bitslice_t filter2_4 = f_a_bs(state[-2 + 36], state[-2 + 45], state[-2 + 46], state[-2 + 48]);
469 const bitslice_t filter2 = f_c_bs(filter2_0, filter2_1, filter2_2, filter2_3, filter2_4);
470
471 const bitslice_t results2 = bs_res(results1,filter2,keystream[2]);
472 if (!results2) continue;
473
474 state[-2 + 50] = lfsr_bs(2);
475 const bitslice_t filter3_3 = f_b_bs(state[-2 + 31], state[-2 + 32], state[-2 + 34], state[-2 + 36]);
476 const bitslice_t filter4_0 = f_a_bs(state[-2 + 6], state[-2 + 7], state[-2 + 9], state[-2 + 10]);
477 const bitslice_t filter4_1 = f_b_bs(state[-2 + 12], state[-2 + 16], state[-2 + 18], state[-2 + 19]);
478 const bitslice_t filter4_2 = f_b_bs(state[-2 + 21], state[-2 + 25], state[-2 + 27], state[-2 + 30]);
479 const bitslice_t filter7_0 = f_a_bs(state[-2 + 9], state[-2 + 10], state[-2 + 12], state[-2 + 13]);
480 const bitslice_t filter7_1 = f_b_bs(state[-2 + 15], state[-2 + 19], state[-2 + 21], state[-2 + 22]);
481 const bitslice_t filter8_2 = f_b_bs(state[-2 + 25], state[-2 + 29], state[-2 + 31], state[-2 + 34]);
482 const bitslice_t filter10_1 = f_b_bs(state[-2 + 18], state[-2 + 22], state[-2 + 24], state[-2 + 25]);
483 const bitslice_t filter10_2 = f_b_bs(state[-2 + 27], state[-2 + 31], state[-2 + 33], state[-2 + 36]);
484 const bitslice_t filter11_1 = f_b_bs(state[-2 + 19], state[-2 + 23], state[-2 + 25], state[-2 + 26]);
485 const bitslice_t filter13_0 = f_a_bs(state[-2 + 15], state[-2 + 16], state[-2 + 18], state[-2 + 19]);
486 const bitslice_t filter13_1 = f_b_bs(state[-2 + 21], state[-2 + 25], state[-2 + 27], state[-2 + 28]);
487 const bitslice_t filter16_0 = f_a_bs(state[-2 + 18], state[-2 + 19], state[-2 + 21], state[-2 + 22]);
488 const bitslice_t filter16_3 = f_b_bs(state[-2 + 44], state[-2 + 45], state[-2 + 47], state[-2 + 49]);
489 const bitslice_t filter17_1 = f_b_bs(state[-2 + 25], state[-2 + 29], state[-2 + 31], state[-2 + 32]);
490 const bitslice_t filter17_3 = f_b_bs(state[-2 + 45], state[-2 + 46], state[-2 + 48], state[-2 + 50]);
491
492 for (uint i3 = 0; i3 < 8; i3++) {
493 state[-2 + 11] = -((bool)(i3 & 0x1));
494 state[-2 + 20] = -((bool)(i3 & 0x2));
495 state[-2 + 37] = -((bool)(i3 & 0x4));
496
497 // 0xff07ffffffff
498 const bitslice_t filter3_1 = f_b_bs(state[-2 + 11], state[-2 + 15], state[-2 + 17], state[-2 + 18]);
499 const bitslice_t filter3_2 = f_b_bs(state[-2 + 20], state[-2 + 24], state[-2 + 26], state[-2 + 29]);
500 const bitslice_t filter3_4 = f_a_bs(state[-2 + 37], state[-2 + 46], state[-2 + 47], state[-2 + 49]);
501 const bitslice_t filter3 = f_c_bs(filter3_0, filter3_1, filter3_2, filter3_3, filter3_4);
502
503 const bitslice_t results3 = bs_res(results2,filter3,keystream[3]);
504 if (!results3) continue;
505
506 state[-2 + 51] = lfsr_bs(3);
507 state[-2 + 52] = lfsr_bs(4);
508 state[-2 + 53] = lfsr_bs(5);
509 state[-2 + 54] = lfsr_bs(6);
510 state[-2 + 55] = lfsr_bs(7);
511
512 const bitslice_t filter4_3 = f_b_bs(state[-2 + 32], state[-2 + 33], state[-2 + 35], state[-2 + 37]);
513 const bitslice_t filter5_0 = f_a_bs(state[-2 + 7], state[-2 + 8], state[-2 + 10], state[-2 + 11]);
514 const bitslice_t filter5_1 = f_b_bs(state[-2 + 13], state[-2 + 17], state[-2 + 19], state[-2 + 20]);
515 const bitslice_t filter6_0 = f_a_bs(state[-2 + 8], state[-2 + 9], state[-2 + 11], state[-2 + 12]);
516 const bitslice_t filter6_1 = f_b_bs(state[-2 + 14], state[-2 + 18], state[-2 + 20], state[-2 + 21]);
517 const bitslice_t filter8_0 = f_a_bs(state[-2 + 10], state[-2 + 11], state[-2 + 13], state[-2 + 14]);
518 const bitslice_t filter8_1 = f_b_bs(state[-2 + 16], state[-2 + 20], state[-2 + 22], state[-2 + 23]);
519 const bitslice_t filter9_0 = f_a_bs(state[-2 + 11], state[-2 + 12], state[-2 + 14], state[-2 + 15]);
520 const bitslice_t filter9_4 = f_a_bs(state[-2 + 43], state[-2 + 52], state[-2 + 53], state[-2 + 55]);
521 const bitslice_t filter11_2 = f_b_bs(state[-2 + 28], state[-2 + 32], state[-2 + 34], state[-2 + 37]);
522 const bitslice_t filter12_1 = f_b_bs(state[-2 + 20], state[-2 + 24], state[-2 + 26], state[-2 + 27]);
523 const bitslice_t filter14_0 = f_a_bs(state[-2 + 16], state[-2 + 17], state[-2 + 19], state[-2 + 20]);
524 const bitslice_t filter15_0 = f_a_bs(state[-2 + 17], state[-2 + 18], state[-2 + 20], state[-2 + 21]);
525 const bitslice_t filter17_0 = f_a_bs(state[-2 + 19], state[-2 + 20], state[-2 + 22], state[-2 + 23]);
526
527 for (uint i4 = 0; i4 < 2; i4++) {
528 state[-2 + 38] = -i4;
529
530 // 0xff87ffffffff
531 const bitslice_t filter4_4 = f_a_bs(state[-2 + 38], state[-2 + 47], state[-2 + 48], state[-2 + 50]);
532 const bitslice_t filter4 = f_c_bs(filter4_0, filter4_1, filter4_2, filter4_3, filter4_4);
533
534 const bitslice_t results4 = bs_res(results3,filter4,keystream[4]);
535 if (!results4) continue;
536
537 state[-2 + 56] = lfsr_bs(8);
538 const bitslice_t filter5_3 = f_b_bs(state[-2 + 33], state[-2 + 34], state[-2 + 36], state[-2 + 38]);
539 const bitslice_t filter10_4 = f_a_bs(state[-2 + 44], state[-2 + 53], state[-2 + 54], state[-2 + 56]);
540 const bitslice_t filter12_2 = f_b_bs(state[-2 + 29], state[-2 + 33], state[-2 + 35], state[-2 + 38]);
541
542 for (uint i5 = 0; i5 < 2; i5++) {
543 state[-2 + 39] = -i5;
544
545 // 0xffc7ffffffff
546 const bitslice_t filter5_4 = f_a_bs(state[-2 + 39], state[-2 + 48], state[-2 + 49], state[-2 + 51]);
547 const bitslice_t filter5 = f_c_bs(filter5_0, filter5_1, filter5_2, filter5_3, filter5_4);
548
549 const bitslice_t results5 = bs_res(results4,filter5,keystream[5]);
550 if (!results5) continue;
551
552 state[-2 + 57] = lfsr_bs(9);
553 const bitslice_t filter6_3 = f_b_bs(state[-2 + 34], state[-2 + 35], state[-2 + 37], state[-2 + 39]);
554 const bitslice_t filter11_4 = f_a_bs(state[-2 + 45], state[-2 + 54], state[-2 + 55], state[-2 + 57]);
555 const bitslice_t filter13_2 = f_b_bs(state[-2 + 30], state[-2 + 34], state[-2 + 36], state[-2 + 39]);
556
557 for (uint i6 = 0; i6 < 2; i6++) {
558 state[-2 + 40] = -i6;
559
560 // 0xffe7ffffffff
561 const bitslice_t filter6_4 = f_a_bs(state[-2 + 40], state[-2 + 49], state[-2 + 50], state[-2 + 52]);
562 const bitslice_t filter6 = f_c_bs(filter6_0, filter6_1, filter6_2, filter6_3, filter6_4);
563
564 const bitslice_t results6 = bs_res(results5,filter6,keystream[6]);
565 if (!results6) continue;
566
567 state[-2 + 58] = lfsr_bs(10);
568 const bitslice_t filter7_3 = f_b_bs(state[-2 + 35], state[-2 + 36], state[-2 + 38], state[-2 + 40]);
569 const bitslice_t filter12_4 = f_a_bs(state[-2 + 46], state[-2 + 55], state[-2 + 56], state[-2 + 58]);
570 const bitslice_t filter14_2 = f_b_bs(state[-2 + 31], state[-2 + 35], state[-2 + 37], state[-2 + 40]);
571 const bitslice_t filter17_2 = f_b_bs(state[-2 + 34], state[-2 + 38], state[-2 + 40], state[-2 + 43]);
572
573 #pragma unroll
574 for (uint i7 = 0; i7 < 2; i7++) {
575 state[-2 + 41] = -i7;
576
577 // 0xfff7ffffffff
578 const bitslice_t filter7_4 = f_a_bs(state[-2 + 41], state[-2 + 50], state[-2 + 51], state[-2 + 53]);
579 const bitslice_t filter7 = f_c_bs(filter7_0, filter7_1, filter7_2, filter7_3, filter7_4);
580
581 const bitslice_t results7 = bs_res(results6,filter7,keystream[7]);
582 if (!results7) continue;
583
584 state[-2 + 59] = lfsr_bs(11);
585 const bitslice_t filter8_3 = f_b_bs(state[-2 + 36], state[-2 + 37], state[-2 + 39], state[-2 + 41]);
586 const bitslice_t filter10_3 = f_b_bs(state[-2 + 38], state[-2 + 39], state[-2 + 41], state[-2 + 43]);
587 const bitslice_t filter10 = f_c_bs(filter10_0, filter10_1, filter10_2, filter10_3, filter10_4);
588 const bitslice_t filter12_3 = f_b_bs(state[-2 + 40], state[-2 + 41], state[-2 + 43], state[-2 + 45]);
589 const bitslice_t filter12 = f_c_bs(filter12_0, filter12_1, filter12_2, filter12_3, filter12_4);
590 const bitslice_t filter13_4 = f_a_bs(state[-2 + 47], state[-2 + 56], state[-2 + 57], state[-2 + 59]);
591 const bitslice_t filter15_2 = f_b_bs(state[-2 + 32], state[-2 + 36], state[-2 + 38], state[-2 + 41]);
592
593 #pragma unroll
594 for (uint i8 = 0; i8 < 2; i8++) {
595 state[-2 + 42] = -i8;
596
597 // 0xffffffffffff
598 const bitslice_t filter8_4 = f_a_bs(state[-2 + 42], state[-2 + 51], state[-2 + 52], state[-2 + 54]);
599 const bitslice_t filter8 = f_c_bs(filter8_0, filter8_1, filter8_2, filter8_3, filter8_4);
600
601 bitslice_t results8 = bs_res(results7,filter8,keystream[8]);
602 if (!results8) continue;
603
604 const bitslice_t filter9_3 = f_b_bs(state[-2 + 37], state[-2 + 38], state[-2 + 40], state[-2 + 42]);
605 const bitslice_t filter9 = f_c_bs(filter9_0, filter9_1, filter9_2, filter9_3, filter9_4);
606
607 results8 &= (filter9 ^ keystream[9]);
608 if (!results8) continue;
609
610 results8 &= (filter10 ^ keystream[10]);
611 if (!results8) continue;
612
613 const bitslice_t filter11_3 = f_b_bs(state[-2 + 39], state[-2 + 40], state[-2 + 42], state[-2 + 44]);
614 const bitslice_t filter11 = f_c_bs(filter11_0, filter11_1, filter11_2, filter11_3, filter11_4);
615
616 results8 &= (filter11 ^ keystream[11]);
617 if (!results8) continue;
618
619 results8 &= (filter12 ^ keystream[12]);
620 if (!results8) continue;
621
622 const bitslice_t filter13_3 = f_b_bs(state[-2 + 41], state[-2 + 42], state[-2 + 44], state[-2 + 46]);
623 const bitslice_t filter13 = f_c_bs(filter13_0, filter13_1, filter13_2, filter13_3, filter13_4);
624
625 results8 &= (filter13 ^ keystream[13]);
626 if (!results8) continue;
627
628 state[-2 + 60] = lfsr_bs(12);
629 const bitslice_t filter14_3 = f_b_bs(state[-2 + 42], state[-2 + 43], state[-2 + 45], state[-2 + 47]);
630 const bitslice_t filter14_4 = f_a_bs(state[-2 + 48], state[-2 + 57], state[-2 + 58], state[-2 + 60]);
631 const bitslice_t filter14 = f_c_bs(filter14_0, filter14_1, filter14_2, filter14_3, filter14_4);
632
633 results8 &= (filter14 ^ keystream[14]);
634 if (!results8) continue;
635
636 state[-2 + 61] = lfsr_bs(13);
637 const bitslice_t filter15_4 = f_a_bs(state[-2 + 49], state[-2 + 58], state[-2 + 59], state[-2 + 61]);
638 const bitslice_t filter15 = f_c_bs(filter15_0, filter15_1, filter15_2, filter15_3, filter15_4);
639
640 results8 &= (filter15 ^ keystream[15]);
641 if (!results8) continue;
642
643 state[-2 + 62] = lfsr_bs(14);
644 const bitslice_t filter16_2 = f_b_bs(state[-2 + 33], state[-2 + 37], state[-2 + 39], state[-2 + 42]);
645 const bitslice_t filter16_4 = f_a_bs(state[-2 + 50], state[-2 + 59], state[-2 + 60], state[-2 + 62]);
646 const bitslice_t filter16 = f_c_bs(filter16_0, filter16_1, filter16_2, filter16_3, filter16_4);
647
648 results8 &= (filter16 ^ keystream[16]);
649 if (!results8) continue;
650
651 state[-2 + 63] = lfsr_bs(15);
652 const bitslice_t filter17_4 = f_a_bs(state[-2 + 51], state[-2 + 60], state[-2 + 61], state[-2 + 63]);
653 const bitslice_t filter17 = f_c_bs(filter17_0, filter17_1, filter17_2, filter17_3, filter17_4);
654
655 results8 &= (filter17 ^ keystream[17]);
656 if (!results8) continue;
657
658 state[-2 + 64] = lfsr_bs(16);
659 const bitslice_t filter18_0 = f_a_bs(state[-2 + 20], state[-2 + 21], state[-2 + 23], state[-2 + 24]);
660 const bitslice_t filter18_1 = f_b_bs(state[-2 + 26], state[-2 + 30], state[-2 + 32], state[-2 + 33]);
661 const bitslice_t filter18_2 = f_b_bs(state[-2 + 35], state[-2 + 39], state[-2 + 41], state[-2 + 44]);
662 const bitslice_t filter18_3 = f_b_bs(state[-2 + 46], state[-2 + 47], state[-2 + 49], state[-2 + 51]);
663 const bitslice_t filter18_4 = f_a_bs(state[-2 + 52], state[-2 + 61], state[-2 + 62], state[-2 + 64]);
664 const bitslice_t filter18 = f_c_bs(filter18_0, filter18_1, filter18_2, filter18_3, filter18_4);
665
666 results8 &= (filter18 ^ keystream[18]);
667 if (!results8) continue;
668
669 state[-2 + 65] = lfsr_bs(17);
670 const bitslice_t filter19_0 = f_a_bs(state[-2 + 21], state[-2 + 22], state[-2 + 24], state[-2 + 25]);
671 const bitslice_t filter19_1 = f_b_bs(state[-2 + 27], state[-2 + 31], state[-2 + 33], state[-2 + 34]);
672 const bitslice_t filter19_2 = f_b_bs(state[-2 + 36], state[-2 + 40], state[-2 + 42], state[-2 + 45]);
673 const bitslice_t filter19_3 = f_b_bs(state[-2 + 47], state[-2 + 48], state[-2 + 50], state[-2 + 52]);
674 const bitslice_t filter19_4 = f_a_bs(state[-2 + 53], state[-2 + 62], state[-2 + 63], state[-2 + 65]);
675 const bitslice_t filter19 = f_c_bs(filter19_0, filter19_1, filter19_2, filter19_3, filter19_4);
676
677 results8 &= (filter19 ^ keystream[19]);
678 if (!results8) continue;
679
680 state[-2 + 66] = lfsr_bs(18);
681 const bitslice_t filter20_0 = f_a_bs(state[-2 + 22], state[-2 + 23], state[-2 + 25], state[-2 + 26]);
682 const bitslice_t filter20_1 = f_b_bs(state[-2 + 28], state[-2 + 32], state[-2 + 34], state[-2 + 35]);
683 const bitslice_t filter20_2 = f_b_bs(state[-2 + 37], state[-2 + 41], state[-2 + 43], state[-2 + 46]);
684 const bitslice_t filter20_3 = f_b_bs(state[-2 + 48], state[-2 + 49], state[-2 + 51], state[-2 + 53]);
685 const bitslice_t filter20_4 = f_a_bs(state[-2 + 54], state[-2 + 63], state[-2 + 64], state[-2 + 66]);
686 const bitslice_t filter20 = f_c_bs(filter20_0, filter20_1, filter20_2, filter20_3, filter20_4);
687
688 results8 &= (filter20 ^ keystream[20]);
689 if (!results8) continue;
690
691 state[-2 + 67] = lfsr_bs(19);
692 const bitslice_t filter21_0 = f_a_bs(state[-2 + 23], state[-2 + 24], state[-2 + 26], state[-2 + 27]);
693 const bitslice_t filter21_1 = f_b_bs(state[-2 + 29], state[-2 + 33], state[-2 + 35], state[-2 + 36]);
694 const bitslice_t filter21_2 = f_b_bs(state[-2 + 38], state[-2 + 42], state[-2 + 44], state[-2 + 47]);
695 const bitslice_t filter21_3 = f_b_bs(state[-2 + 49], state[-2 + 50], state[-2 + 52], state[-2 + 54]);
696 const bitslice_t filter21_4 = f_a_bs(state[-2 + 55], state[-2 + 64], state[-2 + 65], state[-2 + 67]);
697 const bitslice_t filter21 = f_c_bs(filter21_0, filter21_1, filter21_2, filter21_3, filter21_4);
698
699 results8 &= (filter21 ^ keystream[21]);
700 if (!results8) continue;
701
702 state[-2 + 68] = lfsr_bs(20);
703 const bitslice_t filter22_0 = f_a_bs(state[-2 + 24], state[-2 + 25], state[-2 + 27], state[-2 + 28]);
704 const bitslice_t filter22_1 = f_b_bs(state[-2 + 30], state[-2 + 34], state[-2 + 36], state[-2 + 37]);
705 const bitslice_t filter22_2 = f_b_bs(state[-2 + 39], state[-2 + 43], state[-2 + 45], state[-2 + 48]);
706 const bitslice_t filter22_3 = f_b_bs(state[-2 + 50], state[-2 + 51], state[-2 + 53], state[-2 + 55]);
707 const bitslice_t filter22_4 = f_a_bs(state[-2 + 56], state[-2 + 65], state[-2 + 66], state[-2 + 68]);
708 const bitslice_t filter22 = f_c_bs(filter22_0, filter22_1, filter22_2, filter22_3, filter22_4);
709
710 results8 &= (filter22 ^ keystream[22]);
711 if (!results8) continue;
712
713 state[-2 + 69] = lfsr_bs(21);
714 const bitslice_t filter23_0 = f_a_bs(state[-2 + 25], state[-2 + 26], state[-2 + 28], state[-2 + 29]);
715 const bitslice_t filter23_1 = f_b_bs(state[-2 + 31], state[-2 + 35], state[-2 + 37], state[-2 + 38]);
716 const bitslice_t filter23_2 = f_b_bs(state[-2 + 40], state[-2 + 44], state[-2 + 46], state[-2 + 49]);
717 const bitslice_t filter23_3 = f_b_bs(state[-2 + 51], state[-2 + 52], state[-2 + 54], state[-2 + 56]);
718 const bitslice_t filter23_4 = f_a_bs(state[-2 + 57], state[-2 + 66], state[-2 + 67], state[-2 + 69]);
719 const bitslice_t filter23 = f_c_bs(filter23_0, filter23_1, filter23_2, filter23_3, filter23_4);
720
721 results8 &= (filter23 ^ keystream[23]);
722 if (!results8) continue;
723
724 state[-2 + 70] = lfsr_bs(22);
725 const bitslice_t filter24_0 = f_a_bs(state[-2 + 26], state[-2 + 27], state[-2 + 29], state[-2 + 30]);
726 const bitslice_t filter24_1 = f_b_bs(state[-2 + 32], state[-2 + 36], state[-2 + 38], state[-2 + 39]);
727 const bitslice_t filter24_2 = f_b_bs(state[-2 + 41], state[-2 + 45], state[-2 + 47], state[-2 + 50]);
728 const bitslice_t filter24_3 = f_b_bs(state[-2 + 52], state[-2 + 53], state[-2 + 55], state[-2 + 57]);
729 const bitslice_t filter24_4 = f_a_bs(state[-2 + 58], state[-2 + 67], state[-2 + 68], state[-2 + 70]);
730 const bitslice_t filter24 = f_c_bs(filter24_0, filter24_1, filter24_2, filter24_3, filter24_4);
731
732 results8 &= (filter24 ^ keystream[24]);
733 if (!results8) continue;
734
735 state[-2 + 71] = lfsr_bs(23);
736 const bitslice_t filter25_0 = f_a_bs(state[-2 + 27], state[-2 + 28], state[-2 + 30], state[-2 + 31]);
737 const bitslice_t filter25_1 = f_b_bs(state[-2 + 33], state[-2 + 37], state[-2 + 39], state[-2 + 40]);
738 const bitslice_t filter25_2 = f_b_bs(state[-2 + 42], state[-2 + 46], state[-2 + 48], state[-2 + 51]);
739 const bitslice_t filter25_3 = f_b_bs(state[-2 + 53], state[-2 + 54], state[-2 + 56], state[-2 + 58]);
740 const bitslice_t filter25_4 = f_a_bs(state[-2 + 59], state[-2 + 68], state[-2 + 69], state[-2 + 71]);
741 const bitslice_t filter25 = f_c_bs(filter25_0, filter25_1, filter25_2, filter25_3, filter25_4);
742
743 results8 &= (filter25 ^ keystream[25]);
744 if (!results8) continue;
745
746 state[-2 + 72] = lfsr_bs(24);
747 const bitslice_t filter26_0 = f_a_bs(state[-2 + 28], state[-2 + 29], state[-2 + 31], state[-2 + 32]);
748 const bitslice_t filter26_1 = f_b_bs(state[-2 + 34], state[-2 + 38], state[-2 + 40], state[-2 + 41]);
749 const bitslice_t filter26_2 = f_b_bs(state[-2 + 43], state[-2 + 47], state[-2 + 49], state[-2 + 52]);
750 const bitslice_t filter26_3 = f_b_bs(state[-2 + 54], state[-2 + 55], state[-2 + 57], state[-2 + 59]);
751 const bitslice_t filter26_4 = f_a_bs(state[-2 + 60], state[-2 + 69], state[-2 + 70], state[-2 + 72]);
752 const bitslice_t filter26 = f_c_bs(filter26_0, filter26_1, filter26_2, filter26_3, filter26_4);
753
754 results8 &= (filter26 ^ keystream[26]);
755 if (!results8) continue;
756
757 state[-2 + 73] = lfsr_bs(25);
758 const bitslice_t filter27_0 = f_a_bs(state[-2 + 29], state[-2 + 30], state[-2 + 32], state[-2 + 33]);
759 const bitslice_t filter27_1 = f_b_bs(state[-2 + 35], state[-2 + 39], state[-2 + 41], state[-2 + 42]);
760 const bitslice_t filter27_2 = f_b_bs(state[-2 + 44], state[-2 + 48], state[-2 + 50], state[-2 + 53]);
761 const bitslice_t filter27_3 = f_b_bs(state[-2 + 55], state[-2 + 56], state[-2 + 58], state[-2 + 60]);
762 const bitslice_t filter27_4 = f_a_bs(state[-2 + 61], state[-2 + 70], state[-2 + 71], state[-2 + 73]);
763 const bitslice_t filter27 = f_c_bs(filter27_0, filter27_1, filter27_2, filter27_3, filter27_4);
764
765 results8 &= (filter27 ^ keystream[27]);
766 if (!results8) continue;
767
768 state[-2 + 74] = lfsr_bs(26);
769 const bitslice_t filter28_0 = f_a_bs(state[-2 + 30], state[-2 + 31], state[-2 + 33], state[-2 + 34]);
770 const bitslice_t filter28_1 = f_b_bs(state[-2 + 36], state[-2 + 40], state[-2 + 42], state[-2 + 43]);
771 const bitslice_t filter28_2 = f_b_bs(state[-2 + 45], state[-2 + 49], state[-2 + 51], state[-2 + 54]);
772 const bitslice_t filter28_3 = f_b_bs(state[-2 + 56], state[-2 + 57], state[-2 + 59], state[-2 + 61]);
773 const bitslice_t filter28_4 = f_a_bs(state[-2 + 62], state[-2 + 71], state[-2 + 72], state[-2 + 74]);
774 const bitslice_t filter28 = f_c_bs(filter28_0, filter28_1, filter28_2, filter28_3, filter28_4);
775
776 results8 &= (filter28 ^ keystream[28]);
777 if (!results8) continue;
778
779 state[-2 + 75] = lfsr_bs(27);
780 const bitslice_t filter29_0 = f_a_bs(state[-2 + 31], state[-2 + 32], state[-2 + 34], state[-2 + 35]);
781 const bitslice_t filter29_1 = f_b_bs(state[-2 + 37], state[-2 + 41], state[-2 + 43], state[-2 + 44]);
782 const bitslice_t filter29_2 = f_b_bs(state[-2 + 46], state[-2 + 50], state[-2 + 52], state[-2 + 55]);
783 const bitslice_t filter29_3 = f_b_bs(state[-2 + 57], state[-2 + 58], state[-2 + 60], state[-2 + 62]);
784 const bitslice_t filter29_4 = f_a_bs(state[-2 + 63], state[-2 + 72], state[-2 + 73], state[-2 + 75]);
785 const bitslice_t filter29 = f_c_bs(filter29_0, filter29_1, filter29_2, filter29_3, filter29_4);
786
787 results8 &= (filter29 ^ keystream[29]);
788 if (!results8) continue;
789
790 state[-2 + 76] = lfsr_bs(28);
791 const bitslice_t filter30_0 = f_a_bs(state[-2 + 32], state[-2 + 33], state[-2 + 35], state[-2 + 36]);
792 const bitslice_t filter30_1 = f_b_bs(state[-2 + 38], state[-2 + 42], state[-2 + 44], state[-2 + 45]);
793 const bitslice_t filter30_2 = f_b_bs(state[-2 + 47], state[-2 + 51], state[-2 + 53], state[-2 + 56]);
794 const bitslice_t filter30_3 = f_b_bs(state[-2 + 58], state[-2 + 59], state[-2 + 61], state[-2 + 63]);
795 const bitslice_t filter30_4 = f_a_bs(state[-2 + 64], state[-2 + 73], state[-2 + 74], state[-2 + 76]);
796 const bitslice_t filter30 = f_c_bs(filter30_0, filter30_1, filter30_2, filter30_3, filter30_4);
797
798 results8 &= (filter30 ^ keystream[30]);
799 if (!results8) continue;
800
801 state[-2 + 77] = lfsr_bs(29);
802 const bitslice_t filter31_0 = f_a_bs(state[-2 + 33], state[-2 + 34], state[-2 + 36], state[-2 + 37]);
803 const bitslice_t filter31_1 = f_b_bs(state[-2 + 39], state[-2 + 43], state[-2 + 45], state[-2 + 46]);
804 const bitslice_t filter31_2 = f_b_bs(state[-2 + 48], state[-2 + 52], state[-2 + 54], state[-2 + 57]);
805 const bitslice_t filter31_3 = f_b_bs(state[-2 + 59], state[-2 + 60], state[-2 + 62], state[-2 + 64]);
806 const bitslice_t filter31_4 = f_a_bs(state[-2 + 65], state[-2 + 74], state[-2 + 75], state[-2 + 77]);
807 const bitslice_t filter31 = f_c_bs(filter31_0, filter31_1, filter31_2, filter31_3, filter31_4);
808
809 results8 &= (filter31 ^ keystream[31]);
810 if (!results8) continue;
811
812 for (uint match_index = 0; match_index < MAX_BITSLICES && results8;) {
813 const uint shift = clz(results8) + 1;
814 match_index += shift;
815
816 #ifdef WITH_HITAG2_FULL
817
818 ulong state_check = unbitslice (&state[-2 + 2], MAX_BITSLICES - match_index);
819
820 // 2 rollback
821 state_check = (ulong)(((state_check << 1) & 0xffffffffffff) | (ulong)fnR(state_check));
822 state_check = (ulong)(((state_check << 1) & 0xffffffffffff) | (ulong)fnR(state_check));
823
824 // recover key
825 ulong keyrev = state_check & 0xffff;
826 ulong nR1xk = (state_check >> 16) & 0xffffffff;
827
828 uint b = 0;
829
830 for (uint i = 0; i < 32; i++) {
831 state_check = ((state_check) << 1) | ((checks[0] >> (31 - i)) & 0x1);
832 b = (b << 1) | fnf (state_check);
833 }
834
835 keyrev |= (nR1xk ^ checks[2] ^ b) << 16;
836
837 ulong lfsr = 0;
838
839 // test key
840 hitag2_init2 (&state_check, &lfsr, keyrev, checks[0], checks[3]);
841
842 if ((checks[1] ^ hitag2_nstep2 (state_check, lfsr)) == 0xffffffff)
843 {
844 // there can be only one (Highlander) :P
845 matches[atomic_inc(matches_found)] = rev64 (keyrev);
846 return;
847 }
848
849 #else
850
851 // take the state from layer 2 so we can recover the lowest 2 bits on the host by inverting the LFSR
852 matches[atomic_inc(matches_found)] = unbitslice (&state[-2 + 2], MAX_BITSLICES - match_index);
853
854
855 #endif // WITH_HITAG2_FULL
856
857 results8 <<= shift;
858 } // key check
859 } // 8
860 } // 7
861 } // 6
862 } // 5
863 } // 4
864 } // 3
865 } // 2
866 } // 1
Proxmark3 - the RFID research Swiss army knife. Iceman firmware