 
 		// squeeze this in here
 		/* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */
-		u = s[0]; v = s[1]; 
+		u = s[0]; 
 		s[0] = chi(s[0], s[1], s[2]);
 
 		/* iota: a[0,0] ^= round constant */
 		s[0] ^= vectorize(keccak_round_constants[i]);
 		if (i == 23 && out_size == 1) return;
 
+		v = s[1];
 		// continue chi
 		s[1] = chi(s[1], s[2], s[3]);
 		s[2] = chi(s[2], s[3], s[4]);


		/* theta: d[i] = c[i+4] ^ rotl(c[i+1],1) */
		/* theta: a[0,i], a[1,i], .. a[4,i] ^= d[i] */
		u = t[4] ^ ROL2(t[1], 1);
		s[0] ^= u; s[5] ^= u; s[10] ^= u; s[15] ^= u; s[20] ^= u;
		u = t[0] ^ ROL2(t[2], 1);
		s[1] ^= u; s[6] ^= u; s[11] ^= u; s[16] ^= u; s[21] ^= u;
		u = t[1] ^ ROL2(t[3], 1);
		s[2] ^= u; s[7] ^= u; s[12] ^= u; s[17] ^= u; s[22] ^= u;
		u = t[2] ^ ROL2(t[4], 1);
		s[3] ^= u; s[8] ^= u; s[13] ^= u; s[18] ^= u; s[23] ^= u;
		u = t[3] ^ ROL2(t[0], 1);
		s[4] ^= u; s[9] ^= u; s[14] ^= u; s[19] ^= u; s[24] ^= u;

		/* rho pi: b[..] = rotl(a[..], ..) */
		u = s[1];

__device__ __forceinline__
uint2 xor5(const uint2 a, const uint2 b, const uint2 c, const uint2 d, const uint2 e) {
	return a ^ b ^ c ^ d ^ e;
}
__device__ __forceinline__
uint2 lop3xor(const uint2 a, const uint2 b, const uint2 c) {
	return a ^ b ^ c;
}
#endif

