From ef817df79a91a2ae4aa7f44474e158392803592b Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Fri, 23 Oct 2015 08:23:09 +0200 Subject: [PATCH] import sp skein512 unrolled 64-bytes kernel (+0,6% x11) Quark and S3 are now a bit faster (+1 %) x11 get +0.6 % (+20kH/s on a 750ti, +30kH on a 960) 80 bytes implementation to do/test ... (skein/skein2) but keep my previous version for older devices... --- cuda_helper.h | 6 +- quark/cuda_skein512.cu | 16 + quark/cuda_skein512_sp.cuh | 2647 +++++++++++++++++++++++++++++++++++ quark/cuda_vector_uint2x4.h | 547 ++++++++ x11/x11.cu | 5 +- 5 files changed, 3215 insertions(+), 6 deletions(-) create mode 100644 quark/cuda_skein512_sp.cuh create mode 100644 quark/cuda_vector_uint2x4.h diff --git a/cuda_helper.h b/cuda_helper.h index 0e63ab7..ea795b0 100644 --- a/cuda_helper.h +++ b/cuda_helper.h @@ -633,9 +633,9 @@ __device__ __inline__ uint2 ROR24(const uint2 a) return result; } #else -#define ROL8(u) ((u) << 8) -#define ROR16(u) ((u) >> 16) -#define ROR24(u) ((u) >> 24) +#define ROL8(u) SHL2(u, 8) +#define ROR16(u) SHR2(u,16) +#define ROR24(u) SHR2(u,24) #endif /* uint2 for bmw512 - to double check later */ diff --git a/quark/cuda_skein512.cu b/quark/cuda_skein512.cu index 011f5d1..ab3d951 100644 --- a/quark/cuda_skein512.cu +++ b/quark/cuda_skein512.cu @@ -1,9 +1,20 @@ +#define SP_KERNEL + +#ifdef SP_KERNEL +#include "cuda_skein512_sp.cuh" +#undef TFBIG_KINIT +#undef TFBIG_ADDKEY +#undef TFBIG_MIX +#else + #include #include #include #include "cuda_helper.h" +#endif + static __constant__ uint64_t c_PaddedMessage80[20]; // padded message (80 bytes + 72 bytes midstate + align) // Take a look at: https://www.schneier.com/skein1.3.pdf @@ -890,6 +901,11 @@ void quark_skein512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNoun int dev_id = device_map[thr_id]; // uint2 uint64 variants for SM 3.2+ +#ifdef SP_KERNEL + if (device_sm[dev_id] >= 500 && cuda_arch[dev_id] >= 500) + quark_skein512_cpu_hash_64(threads, startNounce, d_nonceVector, d_hash); /* sp.cuh */ + else +#endif if (device_sm[dev_id] > 300 && cuda_arch[dev_id] > 300) quark_skein512_gpu_hash_64 <<>> (threads, startNounce, (uint64_t*)d_hash, d_nonceVector); else diff --git a/quark/cuda_skein512_sp.cuh b/quark/cuda_skein512_sp.cuh new file mode 100644 index 0000000..56ca214 --- /dev/null +++ b/quark/cuda_skein512_sp.cuh @@ -0,0 +1,2647 @@ +/* sp unrolled implentation of skein, used only for SM 5+ and 64 bytes input */ + +//#define WANT_SKEIN_80 + +#include +#include +#include + +#include "cuda_helper.h" + +#include "cuda_vector_uint2x4.h" + +/* ******* SP to TP ******* */ +#define _LOWORD(x) _LODWORD(x) +#define _HIWORD(x) _HIDWORD(x) +// simplified, inline func not faster +#define vectorizelow(/* uint32_t*/ v) make_uint2(v,0) +#define vectorizehigh(/*uint32_t*/ v) make_uint2(0,v) + +__device__ __inline__ uint2 ROL24(const uint2 a) +{ + uint2 result; + result.x = __byte_perm(a.x, a.y, 0x0765); + result.y = __byte_perm(a.y, a.x, 0x0765); + return result; +} +__device__ __inline__ uint2 ROR8(const uint2 a) { + uint2 result; + result.x = __byte_perm(a.y, a.x, 0x0765); + result.y = __byte_perm(a.x, a.y, 0x0765); + return result; +} +/* ************************ */ + +#ifdef WANT_SKEIN_80 +static __constant__ uint64_t c_PaddedMessage16[2]; +__constant__ uint2 precalcvalues[9]; +__constant__ uint32_t sha256_endingTable[64]; +static uint32_t *d_found[MAX_GPUS]; + +// Take a look at: https://www.schneier.com/skein1.3.pdf + +#define SHL(x, n) ((x) << (n)) +#define SHR(x, n) ((x) >> (n)) + +static uint32_t *d_nonce[MAX_GPUS]; +#endif + +/* + * M9_ ## s ## _ ## i evaluates to s+i mod 9 (0 <= s <= 18, 0 <= i <= 7). + */ + +#define M9_0_0 0 +#define M9_0_1 1 +#define M9_0_2 2 +#define M9_0_3 3 +#define M9_0_4 4 +#define M9_0_5 5 +#define M9_0_6 6 +#define M9_0_7 7 + +#define M9_1_0 1 +#define M9_1_1 2 +#define M9_1_2 3 +#define M9_1_3 4 +#define M9_1_4 5 +#define M9_1_5 6 +#define M9_1_6 7 +#define M9_1_7 8 + +#define M9_2_0 2 +#define M9_2_1 3 +#define M9_2_2 4 +#define M9_2_3 5 +#define M9_2_4 6 +#define M9_2_5 7 +#define M9_2_6 8 +#define M9_2_7 0 + +#define M9_3_0 3 +#define M9_3_1 4 +#define M9_3_2 5 +#define M9_3_3 6 +#define M9_3_4 7 +#define M9_3_5 8 +#define M9_3_6 0 +#define M9_3_7 1 + +#define M9_4_0 4 +#define M9_4_1 5 +#define M9_4_2 6 +#define M9_4_3 7 +#define M9_4_4 8 +#define M9_4_5 0 +#define M9_4_6 1 +#define M9_4_7 2 + +#define M9_5_0 5 +#define M9_5_1 6 +#define M9_5_2 7 +#define M9_5_3 8 +#define M9_5_4 0 +#define M9_5_5 1 +#define M9_5_6 2 +#define M9_5_7 3 + +#define M9_6_0 6 +#define M9_6_1 7 +#define M9_6_2 8 +#define M9_6_3 0 +#define M9_6_4 1 +#define M9_6_5 2 +#define M9_6_6 3 +#define M9_6_7 4 + +#define M9_7_0 7 +#define M9_7_1 8 +#define M9_7_2 0 +#define M9_7_3 1 +#define M9_7_4 2 +#define M9_7_5 3 +#define M9_7_6 4 +#define M9_7_7 5 + +#define M9_8_0 8 +#define M9_8_1 0 +#define M9_8_2 1 +#define M9_8_3 2 +#define M9_8_4 3 +#define M9_8_5 4 +#define M9_8_6 5 +#define M9_8_7 6 + +#define M9_9_0 0 +#define M9_9_1 1 +#define M9_9_2 2 +#define M9_9_3 3 +#define M9_9_4 4 +#define M9_9_5 5 +#define M9_9_6 6 +#define M9_9_7 7 + +#define M9_10_0 1 +#define M9_10_1 2 +#define M9_10_2 3 +#define M9_10_3 4 +#define M9_10_4 5 +#define M9_10_5 6 +#define M9_10_6 7 +#define M9_10_7 8 + +#define M9_11_0 2 +#define M9_11_1 3 +#define M9_11_2 4 +#define M9_11_3 5 +#define M9_11_4 6 +#define M9_11_5 7 +#define M9_11_6 8 +#define M9_11_7 0 + +#define M9_12_0 3 +#define M9_12_1 4 +#define M9_12_2 5 +#define M9_12_3 6 +#define M9_12_4 7 +#define M9_12_5 8 +#define M9_12_6 0 +#define M9_12_7 1 + +#define M9_13_0 4 +#define M9_13_1 5 +#define M9_13_2 6 +#define M9_13_3 7 +#define M9_13_4 8 +#define M9_13_5 0 +#define M9_13_6 1 +#define M9_13_7 2 + +#define M9_14_0 5 +#define M9_14_1 6 +#define M9_14_2 7 +#define M9_14_3 8 +#define M9_14_4 0 +#define M9_14_5 1 +#define M9_14_6 2 +#define M9_14_7 3 + +#define M9_15_0 6 +#define M9_15_1 7 +#define M9_15_2 8 +#define M9_15_3 0 +#define M9_15_4 1 +#define M9_15_5 2 +#define M9_15_6 3 +#define M9_15_7 4 + +#define M9_16_0 7 +#define M9_16_1 8 +#define M9_16_2 0 +#define M9_16_3 1 +#define M9_16_4 2 +#define M9_16_5 3 +#define M9_16_6 4 +#define M9_16_7 5 + +#define M9_17_0 8 +#define M9_17_1 0 +#define M9_17_2 1 +#define M9_17_3 2 +#define M9_17_4 3 +#define M9_17_5 4 +#define M9_17_6 5 +#define M9_17_7 6 + +#define M9_18_0 0 +#define M9_18_1 1 +#define M9_18_2 2 +#define M9_18_3 3 +#define M9_18_4 4 +#define M9_18_5 5 +#define M9_18_6 6 +#define M9_18_7 7 + +/* + * M3_ ## s ## _ ## i evaluates to s+i mod 3 (0 <= s <= 18, 0 <= i <= 1). + */ + +#define M3_0_0 0 +#define M3_0_1 1 +#define M3_1_0 1 +#define M3_1_1 2 +#define M3_2_0 2 +#define M3_2_1 0 +#define M3_3_0 0 +#define M3_3_1 1 +#define M3_4_0 1 +#define M3_4_1 2 +#define M3_5_0 2 +#define M3_5_1 0 +#define M3_6_0 0 +#define M3_6_1 1 +#define M3_7_0 1 +#define M3_7_1 2 +#define M3_8_0 2 +#define M3_8_1 0 +#define M3_9_0 0 +#define M3_9_1 1 +#define M3_10_0 1 +#define M3_10_1 2 +#define M3_11_0 2 +#define M3_11_1 0 +#define M3_12_0 0 +#define M3_12_1 1 +#define M3_13_0 1 +#define M3_13_1 2 +#define M3_14_0 2 +#define M3_14_1 0 +#define M3_15_0 0 +#define M3_15_1 1 +#define M3_16_0 1 +#define M3_16_1 2 +#define M3_17_0 2 +#define M3_17_1 0 +#define M3_18_0 0 +#define M3_18_1 1 + +#define XCAT(x, y) XCAT_(x, y) +#define XCAT_(x, y) x ## y + +#define SKBI(k, s, i) XCAT(k, XCAT(XCAT(XCAT(M9_, s), _), i)) +#define SKBT(t, s, v) XCAT(t, XCAT(XCAT(XCAT(M3_, s), _), v)) + +#define TFBIG_KINIT(k0, k1, k2, k3, k4, k5, k6, k7, k8, t0, t1, t2) { \ + k8 = ((k0 ^ k1) ^ (k2 ^ k3)) ^ ((k4 ^ k5) ^ (k6 ^ k7)) \ + ^ make_uint2( 0xA9FC1A22UL,0x1BD11BDA); \ + t2 = t0 ^ t1; \ + } +//vectorize(0x1BD11BDAA9FC1A22ULL); +#define TFBIG_ADDKEY(w0, w1, w2, w3, w4, w5, w6, w7, k, t, s) { \ + w0 = (w0 + SKBI(k, s, 0)); \ + w1 = (w1 + SKBI(k, s, 1)); \ + w2 = (w2 + SKBI(k, s, 2)); \ + w3 = (w3 + SKBI(k, s, 3)); \ + w4 = (w4 + SKBI(k, s, 4)); \ + w5 = (w5 + SKBI(k, s, 5) + SKBT(t, s, 0)); \ + w6 = (w6 + SKBI(k, s, 6) + SKBT(t, s, 1)); \ + w7 = (w7 + SKBI(k, s, 7) + vectorizelow(s)); \ + } + +#define TFBIG_MIX(x0, x1, rc) { \ + x0 = x0 + x1; \ + x1 = ROL2(x1, rc) ^ x0; \ + } + +#define TFBIG_MIX8(w0, w1, w2, w3, w4, w5, w6, w7, rc0, rc1, rc2, rc3) { \ + TFBIG_MIX(w0, w1, rc0); \ + TFBIG_MIX(w2, w3, rc1); \ + TFBIG_MIX(w4, w5, rc2); \ + TFBIG_MIX(w6, w7, rc3); \ + } + +#define TFBIG_4e(s) { \ + TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \ + TFBIG_MIX8(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 46, 36, 19, 37); \ + TFBIG_MIX8(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 33, 27, 14, 42); \ + TFBIG_MIX8(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 17, 49, 36, 39); \ + TFBIG_MIX8(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 44, 9, 54, 56); \ + } + +#define TFBIG_4o(s) { \ + TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \ + TFBIG_MIX8(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 39, 30, 34, 24); \ + TFBIG_MIX8(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 13, 50, 10, 17); \ + TFBIG_MIX8(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 25, 29, 39, 43); \ + TFBIG_MIX8(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 8, 35, 56, 22); \ + } + +/* uint2 variant for SM3.2+ */ + +#define TFBIG_KINIT_UI2(k0, k1, k2, k3, k4, k5, k6, k7, k8, t0, t1, t2) { \ + k8 = ((k0 ^ k1) ^ (k2 ^ k3)) ^ ((k4 ^ k5) ^ (k6 ^ k7)) \ + ^ vectorize(SPH_C64(0x1BD11BDAA9FC1A22)); \ + t2 = t0 ^ t1; \ + } + +#define TFBIG_ADDKEY_UI2(w0, w1, w2, w3, w4, w5, w6, w7, k, t, s) { \ + w0 = (w0 + SKBI(k, s, 0)); \ + w1 = (w1 + SKBI(k, s, 1)); \ + w2 = (w2 + SKBI(k, s, 2)); \ + w3 = (w3 + SKBI(k, s, 3)); \ + w4 = (w4 + SKBI(k, s, 4)); \ + w5 = (w5 + SKBI(k, s, 5) + SKBT(t, s, 0)); \ + w6 = (w6 + SKBI(k, s, 6) + SKBT(t, s, 1)); \ + w7 = (w7 + SKBI(k, s, 7) + vectorize(s)); \ + } + +#define TFBIG_ADDKEY_PRE(w0, w1, w2, w3, w4, w5, w6, w7, k, t, s) { \ + w0 = (w0 + SKBI(k, s, 0)); \ + w1 = (w1 + SKBI(k, s, 1)); \ + w2 = (w2 + SKBI(k, s, 2)); \ + w3 = (w3 + SKBI(k, s, 3)); \ + w4 = (w4 + SKBI(k, s, 4)); \ + w5 = (w5 + SKBI(k, s, 5) + SKBT(t, s, 0)); \ + w6 = (w6 + SKBI(k, s, 6) + SKBT(t, s, 1)); \ + w7 = (w7 + SKBI(k, s, 7) + (s)); \ + } + +#define TFBIG_MIX_UI2(x0, x1, rc) { \ + x0 = x0 + x1; \ + x1 = ROL2(x1, rc) ^ x0; \ + } + +#define TFBIG_MIX_PRE(x0, x1, rc) { \ + x0 = x0 + x1; \ + x1 = ROTL64(x1, rc) ^ x0; \ + } + +#define TFBIG_MIX8_UI2(w0, w1, w2, w3, w4, w5, w6, w7, rc0, rc1, rc2, rc3) { \ + TFBIG_MIX_UI2(w0, w1, rc0); \ + TFBIG_MIX_UI2(w2, w3, rc1); \ + TFBIG_MIX_UI2(w4, w5, rc2); \ + TFBIG_MIX_UI2(w6, w7, rc3); \ + } + +#define TFBIG_MIX8_PRE(w0, w1, w2, w3, w4, w5, w6, w7, rc0, rc1, rc2, rc3) { \ + TFBIG_MIX_PRE(w0, w1, rc0); \ + TFBIG_MIX_PRE(w2, w3, rc1); \ + TFBIG_MIX_PRE(w4, w5, rc2); \ + TFBIG_MIX_PRE(w6, w7, rc3); \ + } + +#define TFBIG_4e_UI2(s) { \ + TFBIG_ADDKEY_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \ + TFBIG_MIX8_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 46, 36, 19, 37); \ + TFBIG_MIX8_UI2(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 33, 27, 14, 42); \ + TFBIG_MIX8_UI2(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 17, 49, 36, 39); \ + TFBIG_MIX8_UI2(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 44, 9, 54, 56); \ + } + +#define TFBIG_4e_PRE(s) { \ + TFBIG_ADDKEY_PRE(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \ + TFBIG_MIX8_PRE(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 46, 36, 19, 37); \ + TFBIG_MIX8_PRE(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 33, 27, 14, 42); \ + TFBIG_MIX8_PRE(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 17, 49, 36, 39); \ + TFBIG_MIX8_PRE(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 44, 9, 54, 56); \ + } + +#define TFBIG_4o_UI2(s) { \ + TFBIG_ADDKEY_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \ + TFBIG_MIX8_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 39, 30, 34, 24); \ + TFBIG_MIX8_UI2(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 13, 50, 10, 17); \ + TFBIG_MIX8_UI2(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 25, 29, 39, 43); \ + TFBIG_MIX8_UI2(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 8, 35, 56, 22); \ + } + +#define TFBIG_4o_PRE(s) { \ + TFBIG_ADDKEY_PRE(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \ + TFBIG_MIX8_PRE(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 39, 30, 34, 24); \ + TFBIG_MIX8_PRE(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 13, 50, 10, 17); \ + TFBIG_MIX8_PRE(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 25, 29, 39, 43); \ + TFBIG_MIX8_PRE(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 8, 35, 56, 22); \ + } +__global__ +#if __CUDA_ARCH__ > 500 +__launch_bounds__(480, 3) +#else +__launch_bounds__(240, 6) +#endif +void quark_skein512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint64_t * const __restrict__ g_hash, const uint32_t *const __restrict__ g_nonceVector) +{ + const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); + if (thread < threads) + { + // Skein + uint2 skein_p[8], h[9]; + + const uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread); + + const int hashPosition = nounce - startNounce; + uint64_t *Hash = &g_hash[8 * hashPosition]; + + + uint2 msg[8]; + + uint28 *phash = (uint28*)Hash; + uint28 *outpt = (uint28*)msg; + outpt[0] = phash[0]; + outpt[1] = phash[1]; + + h[0] = skein_p[0] = (msg[0]); + h[1] = skein_p[1] = (msg[1]); + h[2] = skein_p[2] = (msg[2]); + h[3] = skein_p[3] = (msg[3]); + h[4] = skein_p[4] = (msg[4]); + h[5] = skein_p[5] = (msg[5]); + h[6] = skein_p[6] = (msg[6]); + h[7] = skein_p[7] = (msg[7]); + + skein_p[0] += vectorize(0x4903ADFF749C51CEULL); + skein_p[1] += vectorize(0x0D95DE399746DF03ULL); + skein_p[2] += vectorize(0x8FD1934127C79BCEULL); + skein_p[3] += vectorize(0x9A255629FF352CB1ULL); + skein_p[4] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[5] += vectorize(0xEABE394CA9D5C434ULL); + skein_p[6] += vectorize(0x891112C71A75B523ULL); + skein_p[7] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0x0D95DE399746DF03ULL); + skein_p[1] += vectorize(0x8FD1934127C79BCEULL); + skein_p[2] += vectorize(0x9A255629FF352CB1ULL); + skein_p[3] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[4] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[5] += vectorize(0x891112C71A75B523ULL); + skein_p[6] += vectorize(0x9E18A40B660FCC73ULL); + skein_p[7] += vectorize(0xcab2076d98173ec4ULL+1); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0x8FD1934127C79BCEULL); + skein_p[1] += vectorize(0x9A255629FF352CB1ULL); + skein_p[2] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[3] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[4] += vectorize(0x991112C71A75B523ULL); + skein_p[5] += vectorize(0x9E18A40B660FCC73ULL); + skein_p[6] += vectorize(0xCAB2076D98173F04ULL); + skein_p[7] += vectorize(0x4903ADFF749C51D0ULL); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0x9A255629FF352CB1ULL); + skein_p[1] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[2] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[3] += vectorize(0x991112C71A75B523ULL); + skein_p[4] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[5] += vectorize(0xcab2076d98173f04ULL); + skein_p[6] += vectorize(0x3903ADFF749C51CEULL); + skein_p[7] += vectorize(0x0D95DE399746DF03ULL+3); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[1] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[2] += vectorize(0x991112C71A75B523ULL); + skein_p[3] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[4] += vectorize(0xcab2076d98173ec4ULL); + skein_p[5] += vectorize(0x3903ADFF749C51CEULL); + skein_p[6] += vectorize(0xFD95DE399746DF43ULL); + skein_p[7] += vectorize(0x8FD1934127C79BD2ULL); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[1] += vectorize(0x991112C71A75B523ULL); + skein_p[2] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[3] += vectorize(0xcab2076d98173ec4ULL); + skein_p[4] += vectorize(0x4903ADFF749C51CEULL); + skein_p[5] += vectorize(0x0D95DE399746DF03ULL + 0xf000000000000040ULL); + skein_p[6] += vectorize(0x8FD1934127C79BCEULL + 0x0000000000000040ULL); + skein_p[7] += vectorize(0x9A255629FF352CB1ULL + 5); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0x991112C71A75B523ULL); + skein_p[1] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[2] += vectorize(0xcab2076d98173ec4ULL); + skein_p[3] += vectorize(0x4903ADFF749C51CEULL); + skein_p[4] += vectorize(0x0D95DE399746DF03ULL); + skein_p[5] += vectorize(0x8FD1934127C79BCEULL + 0x0000000000000040ULL); + skein_p[6] += vectorize(0x8A255629FF352CB1ULL); + skein_p[7] += vectorize(0x5DB62599DF6CA7B0ULL + 6); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[1] += vectorize(0xcab2076d98173ec4ULL); + skein_p[2] += vectorize(0x4903ADFF749C51CEULL); + skein_p[3] += vectorize(0x0D95DE399746DF03ULL); + skein_p[4] += vectorize(0x8FD1934127C79BCEULL); + skein_p[5] += vectorize(0x8A255629FF352CB1ULL); + skein_p[6] += vectorize(0x4DB62599DF6CA7F0ULL); + skein_p[7] += vectorize(0xEABE394CA9D5C3F4ULL + 7); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0xcab2076d98173ec4ULL); + skein_p[1] += vectorize(0x4903ADFF749C51CEULL); + skein_p[2] += vectorize(0x0D95DE399746DF03ULL); + skein_p[3] += vectorize(0x8FD1934127C79BCEULL); + skein_p[4] += vectorize(0x9A255629FF352CB1ULL); + skein_p[5] += vectorize(0x4DB62599DF6CA7F0ULL); + skein_p[6] += vectorize(0xEABE394CA9D5C434ULL); + skein_p[7] += vectorize(0x991112C71A75B52BULL); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0x4903ADFF749C51CEULL); + skein_p[1] += vectorize(0x0D95DE399746DF03ULL); + skein_p[2] += vectorize(0x8FD1934127C79BCEULL); + skein_p[3] += vectorize(0x9A255629FF352CB1ULL); + skein_p[4] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[5] += vectorize(0xEABE394CA9D5C434ULL); + skein_p[6] += vectorize(0x891112C71A75B523ULL); + skein_p[7] += vectorize(0xAE18A40B660FCC33ULL + 9); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0x0D95DE399746DF03ULL); + skein_p[1] += vectorize(0x8FD1934127C79BCEULL); + skein_p[2] += vectorize(0x9A255629FF352CB1ULL); + skein_p[3] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[4] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[5] += vectorize(0x891112C71A75B523ULL); + skein_p[6] += vectorize(0x9E18A40B660FCC73ULL); + skein_p[7] += vectorize(0xcab2076d98173eceULL); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0x8FD1934127C79BCEULL); + skein_p[1] += vectorize(0x9A255629FF352CB1ULL); + skein_p[2] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[3] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[4] += vectorize(0x991112C71A75B523ULL); + skein_p[5] += vectorize(0x9E18A40B660FCC73ULL); + skein_p[6] += vectorize(0xcab2076d98173ec4ULL + 0x0000000000000040ULL); + skein_p[7] += vectorize(0x4903ADFF749C51CEULL + 11); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0x9A255629FF352CB1ULL); + skein_p[1] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[2] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[3] += vectorize(0x991112C71A75B523ULL); + skein_p[4] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[5] += vectorize(0xcab2076d98173ec4ULL + 0x0000000000000040ULL); + skein_p[6] += vectorize(0x3903ADFF749C51CEULL); + skein_p[7] += vectorize(0x0D95DE399746DF03ULL + 12); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[1] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[2] += vectorize(0x991112C71A75B523ULL); + skein_p[3] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[4] += vectorize(0xcab2076d98173ec4ULL); + skein_p[5] += vectorize(0x3903ADFF749C51CEULL); + skein_p[6] += vectorize(0x0D95DE399746DF03ULL + 0xf000000000000040ULL); + skein_p[7] += vectorize(0x8FD1934127C79BCEULL + 13); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0xEABE394CA9D5C3F4ULL); + skein_p[1] += vectorize(0x991112C71A75B523ULL); + skein_p[2] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[3] += vectorize(0xcab2076d98173ec4ULL); + skein_p[4] += vectorize(0x4903ADFF749C51CEULL); + skein_p[5] += vectorize(0x0D95DE399746DF03ULL + 0xf000000000000040ULL); + skein_p[6] += vectorize(0x8FD1934127C79BCEULL + 0x0000000000000040ULL); + skein_p[7] += vectorize(0x9A255629FF352CB1ULL + 14); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0x991112C71A75B523ULL); + skein_p[1] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[2] += vectorize(0xcab2076d98173ec4ULL); + skein_p[3] += vectorize(0x4903ADFF749C51CEULL); + skein_p[4] += vectorize(0x0D95DE399746DF03ULL); + skein_p[5] += vectorize(0x8FD1934127C79BCEULL + 0x0000000000000040ULL); + skein_p[6] += vectorize(0x8A255629FF352CB1ULL); + skein_p[7] += vectorize(0x5DB62599DF6CA7B0ULL + 15); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0xAE18A40B660FCC33ULL); + skein_p[1] += vectorize(0xcab2076d98173ec4ULL); + skein_p[2] += vectorize(0x4903ADFF749C51CEULL); + skein_p[3] += vectorize(0x0D95DE399746DF03ULL); + skein_p[4] += vectorize(0x8FD1934127C79BCEULL); + skein_p[5] += vectorize(0x8A255629FF352CB1ULL); + skein_p[6] += vectorize(0x4DB62599DF6CA7F0ULL); + skein_p[7] += vectorize(0xEABE394CA9D5C3F4ULL +16ULL); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 46) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 36) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 19) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 37) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 33) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 27) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 14) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 42) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 17) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 49) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 36) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 39) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 44) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 9) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 54) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROR8(skein_p[3]) ^ skein_p[4]; + skein_p[0] += vectorize(0xcab2076d98173ec4ULL); + skein_p[1] += vectorize(0x4903ADFF749C51CEULL); + skein_p[2] += vectorize(0x0D95DE399746DF03ULL); + skein_p[3] += vectorize(0x8FD1934127C79BCEULL); + skein_p[4] += vectorize(0x9A255629FF352CB1ULL); + skein_p[5] += vectorize(0x4DB62599DF6CA7F0ULL); + skein_p[6] += vectorize(0xEABE394CA9D5C3F4ULL + 0x0000000000000040ULL); + skein_p[7] += vectorize(0x991112C71A75B523ULL + 17); + skein_p[0] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 39) ^ skein_p[0]; + skein_p[2] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 30) ^ skein_p[2]; + skein_p[4] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 34) ^ skein_p[4]; + skein_p[6] += skein_p[7]; + skein_p[7] = ROL24(skein_p[7]) ^ skein_p[6]; + skein_p[2] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 13) ^ skein_p[2]; + skein_p[4] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 50) ^ skein_p[4]; + skein_p[6] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 10) ^ skein_p[6]; + skein_p[0] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 17) ^ skein_p[0]; + skein_p[4] += skein_p[1]; + skein_p[1] = ROL2(skein_p[1], 25) ^ skein_p[4]; + skein_p[6] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 29) ^ skein_p[6]; + skein_p[0] += skein_p[5]; + skein_p[5] = ROL2(skein_p[5], 39) ^ skein_p[0]; + skein_p[2] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 43) ^ skein_p[2]; + skein_p[6] += skein_p[1]; + skein_p[1] = ROL8(skein_p[1]) ^ skein_p[6]; + skein_p[0] += skein_p[7]; + skein_p[7] = ROL2(skein_p[7], 35) ^ skein_p[0]; + skein_p[2] += skein_p[5]; + skein_p[5] = ROR8(skein_p[5]) ^ skein_p[2]; + skein_p[4] += skein_p[3]; + skein_p[3] = ROL2(skein_p[3], 22) ^ skein_p[4]; + skein_p[0] += vectorize(0x4903ADFF749C51CEULL); + skein_p[1] += vectorize(0x0D95DE399746DF03ULL); + skein_p[2] += vectorize(0x8FD1934127C79BCEULL); + skein_p[3] += vectorize(0x9A255629FF352CB1ULL); + skein_p[4] += vectorize(0x5DB62599DF6CA7B0ULL); + skein_p[5] += vectorize(0xEABE394CA9D5C3F4ULL + 0x0000000000000040ULL); + skein_p[6] += vectorize(0x891112C71A75B523ULL); + skein_p[7] += vectorize(0xAE18A40B660FCC33ULL + 18); + +#define h0 skein_p[0] +#define h1 skein_p[1] +#define h2 skein_p[2] +#define h3 skein_p[3] +#define h4 skein_p[4] +#define h5 skein_p[5] +#define h6 skein_p[6] +#define h7 skein_p[7] + h0 ^= h[0]; + h1 ^= h[1]; + h2 ^= h[2]; + h3 ^= h[3]; + h4 ^= h[4]; + h5 ^= h[5]; + h6 ^= h[6]; + h7 ^= h[7]; + + uint2 skein_h8 = h0 ^ h1 ^ h2 ^ h3 ^ h4 ^ h5 ^ h6 ^ h7 ^ vectorize(0x1BD11BDAA9FC1A22ULL); + + uint2 hash64[8]; + + hash64[0] = (h0); +// hash64[1] = (h1); + hash64[2] = (h2); +// hash64[3] = (h3); + hash64[4] = (h4); + hash64[5] = (h5 + vectorizelow(8ULL)); + hash64[6] = (h6 + vectorizehigh(0xff000000UL)); +// hash64[7] = (h7); + + hash64[0] += h1; + hash64[1] = ROL2(h1, 46) ^ hash64[0]; + hash64[2] += h3; + hash64[3] = ROL2(h3, 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += h7; + hash64[7] = ROL2(h7, 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + h1); + hash64[1] = (hash64[1] + h2); + hash64[2] = (hash64[2] + h3); + hash64[3] = (hash64[3] + h4); + hash64[4] = (hash64[4] + h5); + hash64[5] = (hash64[5] + h6 + vectorizehigh(0xff000000UL)); + hash64[6] = (hash64[6] + h7 + vectorize(0xff00000000000008ULL)); + hash64[7] = (hash64[7] + skein_h8 + vectorizelow(1)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + hash64[0] = (hash64[0] + h2); + hash64[1] = (hash64[1] + h3); + hash64[2] = (hash64[2] + h4); + hash64[3] = (hash64[3] + h5); + hash64[4] = (hash64[4] + h6); + hash64[5] = (hash64[5] + h7 + vectorize(0xff00000000000008ULL)); + hash64[6] = (hash64[6] + skein_h8 + vectorizelow(8ULL)); + hash64[7] = (hash64[7] + h0 + vectorize(2)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 46) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL2(hash64[7], 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + h3); + hash64[1] = (hash64[1] + h4); + hash64[2] = (hash64[2] + h5); + hash64[3] = (hash64[3] + h6); + hash64[4] = (hash64[4] + h7); + hash64[5] = (hash64[5] + skein_h8 + vectorizelow(8)); + hash64[6] = (hash64[6] + h0 + vectorizehigh(0xff000000UL)); + hash64[7] = (hash64[7] + h1 + vectorizelow(3)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + hash64[0] = (hash64[0] + h4); + hash64[1] = (hash64[1] + h5); + hash64[2] = (hash64[2] + h6); + hash64[3] = (hash64[3] + h7); + hash64[4] = (hash64[4] + skein_h8); + hash64[5] = (hash64[5] + h0 + vectorizehigh(0xff000000UL)); + hash64[6] = (hash64[6] + h1 + vectorize(0xff00000000000008ULL)); + hash64[7] = (hash64[7] + h2 + vectorizelow(4)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 46) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL2(hash64[7], 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + h5); + hash64[1] = (hash64[1] + h6); + hash64[2] = (hash64[2] + h7); + hash64[3] = (hash64[3] + skein_h8); + hash64[4] = (hash64[4] + h0); + hash64[5] = (hash64[5] + h1 + vectorize(0xff00000000000008ULL)); + hash64[6] = (hash64[6] + h2 + vectorizelow(8ULL)); + hash64[7] = (hash64[7] + h3 + vectorizelow(5)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + hash64[0] = (hash64[0] + h6); + hash64[1] = (hash64[1] + h7); + hash64[2] = (hash64[2] + skein_h8); + hash64[3] = (hash64[3] + h0); + hash64[4] = (hash64[4] + h1); + hash64[5] = (hash64[5] + h2 + vectorizelow(8ULL)); + hash64[6] = (hash64[6] + h3 + vectorizehigh(0xff000000UL)); + hash64[7] = (hash64[7] + h4 + vectorizelow(6)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 46) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL2(hash64[7], 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + h7); + hash64[1] = (hash64[1] + skein_h8); + hash64[2] = (hash64[2] + h0); + hash64[3] = (hash64[3] + h1); + hash64[4] = (hash64[4] + h2); + hash64[5] = (hash64[5] + h3 + vectorizehigh(0xff000000UL)); + hash64[6] = (hash64[6] + h4 + vectorize(0xff00000000000008ULL)); + hash64[7] = (hash64[7] + h5 + vectorizelow(7)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + hash64[0] = (hash64[0] + skein_h8); + hash64[1] = (hash64[1] + h0); + hash64[2] = (hash64[2] + h1); + hash64[3] = (hash64[3] + h2); + hash64[4] = (hash64[4] + h3); + hash64[5] = (hash64[5] + h4 + vectorize(0xff00000000000008ULL)); + hash64[6] = (hash64[6] + h5 + vectorizelow(8)); + hash64[7] = (hash64[7] + h6 + vectorizelow(8)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 46) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL2(hash64[7], 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + h0); + hash64[1] = (hash64[1] + h1); + hash64[2] = (hash64[2] + h2); + hash64[3] = (hash64[3] + h3); + hash64[4] = (hash64[4] + h4); + hash64[5] = (hash64[5] + h5 + vectorizelow(8)); + hash64[6] = (hash64[6] + h6 + vectorizehigh(0xff000000UL)); + hash64[7] = (hash64[7] + h7 + vectorizelow(9)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + + hash64[0] = (hash64[0] + h1); + hash64[1] = (hash64[1] + h2); + hash64[2] = (hash64[2] + h3); + hash64[3] = (hash64[3] + h4); + hash64[4] = (hash64[4] + h5); + hash64[5] = (hash64[5] + h6 + vectorizehigh(0xff000000UL)); + hash64[6] = (hash64[6] + h7 + vectorize(0xff00000000000008ULL)); + hash64[7] = (hash64[7] + skein_h8 + (vectorizelow(10))); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 46) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL2(hash64[7], 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + h2); + hash64[1] = (hash64[1] + h3); + hash64[2] = (hash64[2] + h4); + hash64[3] = (hash64[3] + h5); + hash64[4] = (hash64[4] + h6); + hash64[5] = (hash64[5] + h7 + vectorize(0xff00000000000008ULL)); + hash64[6] = (hash64[6] + skein_h8 + vectorizelow(8ULL)); + hash64[7] = (hash64[7] + h0 + vectorizelow(11)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + hash64[0] = (hash64[0] + h3); + hash64[1] = (hash64[1] + h4); + hash64[2] = (hash64[2] + h5); + hash64[3] = (hash64[3] + h6); + hash64[4] = (hash64[4] + h7); + hash64[5] = (hash64[5] + skein_h8 + vectorizelow(8)); + hash64[6] = (hash64[6] + h0 + vectorizehigh(0xff000000UL)); + hash64[7] = (hash64[7] + h1 + vectorizelow(12)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 46) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL2(hash64[7], 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + h4); + hash64[1] = (hash64[1] + h5); + hash64[2] = (hash64[2] + h6); + hash64[3] = (hash64[3] + h7); + hash64[4] = (hash64[4] + skein_h8); + hash64[5] = (hash64[5] + h0 + vectorizehigh(0xff000000UL)); + hash64[6] = (hash64[6] + h1 + vectorize(0xff00000000000008ULL)); + hash64[7] = (hash64[7] + h2 + vectorizelow(13)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + hash64[0] = (hash64[0] + h5); + hash64[1] = (hash64[1] + h6); + hash64[2] = (hash64[2] + h7); + hash64[3] = (hash64[3] + skein_h8); + hash64[4] = (hash64[4] + h0); + hash64[5] = (hash64[5] + h1 + vectorize(0xff00000000000008ULL)); + hash64[6] = (hash64[6] + h2 + vectorizelow(8ULL)); + hash64[7] = (hash64[7] + h3 + vectorizelow(14)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 46) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL2(hash64[7], 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + h6); + hash64[1] = (hash64[1] + h7); + hash64[2] = (hash64[2] + skein_h8); + hash64[3] = (hash64[3] + h0); + hash64[4] = (hash64[4] + h1); + hash64[5] = (hash64[5] + h2 + vectorizelow(8ULL)); + hash64[6] = (hash64[6] + h3 + vectorizehigh(0xff000000UL)); + hash64[7] = (hash64[7] + h4 + vectorizelow(15)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + hash64[0] = (hash64[0] + h7); + hash64[1] = (hash64[1] + skein_h8); + hash64[2] = (hash64[2] + h0); + hash64[3] = (hash64[3] + h1); + hash64[4] = (hash64[4] + h2); + hash64[5] = (hash64[5] + h3 + vectorizehigh(0xff000000UL)); + hash64[6] = (hash64[6] + h4 + vectorize(0xff00000000000008ULL)); + hash64[7] = (hash64[7] + h5 + vectorizelow(16)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 46) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 36) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 19) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL2(hash64[7], 37) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 33) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 27) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 14) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 42) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 17) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 49) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 36) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 39) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL2(hash64[1], 44) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 9) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROL2(hash64[5], 54) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROR8(hash64[3]) ^ hash64[4]; + hash64[0] = (hash64[0] + skein_h8); + hash64[1] = (hash64[1] + h0); + hash64[2] = (hash64[2] + h1); + hash64[3] = (hash64[3] + h2); + hash64[4] = (hash64[4] + h3); + hash64[5] = (hash64[5] + h4 + vectorize(0xff00000000000008ULL)); + hash64[6] = (hash64[6] + h5 + vectorizelow(8ULL)); + hash64[7] = (hash64[7] + h6 + vectorizelow(17)); + hash64[0] += hash64[1]; + hash64[1] = ROL2(hash64[1], 39) ^ hash64[0]; + hash64[2] += hash64[3]; + hash64[3] = ROL2(hash64[3], 30) ^ hash64[2]; + hash64[4] += hash64[5]; + hash64[5] = ROL2(hash64[5], 34) ^ hash64[4]; + hash64[6] += hash64[7]; + hash64[7] = ROL24(hash64[7]) ^ hash64[6]; + hash64[2] += hash64[1]; + hash64[1] = ROL2(hash64[1], 13) ^ hash64[2]; + hash64[4] += hash64[7]; + hash64[7] = ROL2(hash64[7], 50) ^ hash64[4]; + hash64[6] += hash64[5]; + hash64[5] = ROL2(hash64[5], 10) ^ hash64[6]; + hash64[0] += hash64[3]; + hash64[3] = ROL2(hash64[3], 17) ^ hash64[0]; + hash64[4] += hash64[1]; + hash64[1] = ROL2(hash64[1], 25) ^ hash64[4]; + hash64[6] += hash64[3]; + hash64[3] = ROL2(hash64[3], 29) ^ hash64[6]; + hash64[0] += hash64[5]; + hash64[5] = ROL2(hash64[5], 39) ^ hash64[0]; + hash64[2] += hash64[7]; + hash64[7] = ROL2(hash64[7], 43) ^ hash64[2]; + hash64[6] += hash64[1]; + hash64[1] = ROL8(hash64[1]) ^ hash64[6]; + hash64[0] += hash64[7]; + hash64[7] = ROL2(hash64[7], 35) ^ hash64[0]; + hash64[2] += hash64[5]; + hash64[5] = ROR8(hash64[5]) ^ hash64[2]; + hash64[4] += hash64[3]; + hash64[3] = ROL2(hash64[3], 22) ^ hash64[4]; + + Hash[0] = devectorize(hash64[0] + h0); + Hash[1] = devectorize(hash64[1] + h1); + Hash[2] = devectorize(hash64[2] + h2); + Hash[3] = devectorize(hash64[3] + h3); + Hash[4] = devectorize(hash64[4] + h4); + Hash[5] = devectorize(hash64[5] + h5)+ 8; + Hash[6] = devectorize(hash64[6] + h6)+ 0xff00000000000000ULL; + Hash[7] = devectorize(hash64[7] + h7)+ 18; + +#undef h0 +#undef h1 +#undef h2 +#undef h3 +#undef h4 +#undef h5 +#undef h6 +#undef h7 + } +} + +#if __CUDA_ARCH__ > 500 +#define tp 448 +#else +#define tp 128 +#endif + +__host__ +void quark_skein512_cpu_hash_64(uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash) +{ + dim3 grid((threads + tp - 1) / tp); + dim3 block(tp); + quark_skein512_gpu_hash_64 << > >(threads, startNounce, (uint64_t*)d_hash, d_nonceVector); + +} + +#ifdef WANT_SKEIN_80 + +__host__ void quark_skein512_cpu_init(int thr_id) +{ + cudaMalloc(&d_nonce[thr_id], 2*sizeof(uint32_t)); +} + +__host__ void quark_skein512_setTarget(const void *ptarget) +{ +} +__host__ void quark_skein512_cpu_free(int32_t thr_id) +{ + cudaFreeHost(&d_nonce[thr_id]); +} + + +static __device__ __constant__ uint32_t sha256_hashTable[] = { + 0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19 +}; + + +/* Elementary functions used by SHA256 */ +#define SWAB32(x) cuda_swab32(x) +//#define ROTR32(x, n) (((x) >> (n)) | ((x) << (32 - (n)))) + +#define R(x, n) ((x) >> (n)) +#define Ch(x, y, z) ((x & (y ^ z)) ^ z) +#define Maj(x, y, z) ((x & (y | z)) | (y & z)) +#define S0(x) (ROTR32(x, 2) ^ ROTR32(x, 13) ^ ROTR32(x, 22)) +#define S1(x) (ROTR32(x, 6) ^ ROTR32(x, 11) ^ ROTR32(x, 25)) +#define s0(x) (ROTR32(x, 7) ^ ROTR32(x, 18) ^ R(x, 3)) +#define s1(x) (ROTR32(x, 17) ^ ROTR32(x, 19) ^ R(x, 10)) + + +__constant__ uint32_t sha256_constantTable[64] = { + 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5, + 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174, + 0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da, + 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967, + 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85, + 0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070, + 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3, + 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2 +}; + +__global__ __launch_bounds__(1024) +void skein512_gpu_hash_80_52(uint32_t threads, uint32_t startNounce, uint32_t *const __restrict__ d_found, uint64_t target) +{ + const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); + if (thread < threads) + { + uint2 h0, h1, h2, h3, h4, h5, h6, h7, h8; + uint2 t0, t1, t2; + uint2 p[8]; + + h0 = precalcvalues[0]; + h1 = precalcvalues[1]; + h2 = precalcvalues[2]; + h3 = precalcvalues[3]; + h4 = precalcvalues[4]; + h5 = precalcvalues[5]; + h6 = precalcvalues[6]; + h7 = precalcvalues[7]; + t2 = precalcvalues[8]; + + const uint2 nounce2 = make_uint2(_LOWORD(c_PaddedMessage16[1]), cuda_swab32(startNounce + thread)); + + // skein_big_close -> etype = 0x160, ptr = 16, bcount = 1, extra = 16 + p[0] = vectorize(c_PaddedMessage16[0]); + p[1] = nounce2; + + #pragma unroll + for (int i = 2; i < 8; i++) + p[i] = make_uint2(0,0); + + t0 = vectorizelow(0x50ull); // SPH_T64(bcount << 6) + (sph_u64)(extra); + t1 = vectorizehigh(0xB0000000ul); // (bcount >> 58) + ((sph_u64)(etype) << 55); + TFBIG_KINIT_UI2(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2); + TFBIG_4e_UI2(0); + TFBIG_4o_UI2(1); + TFBIG_4e_UI2(2); + TFBIG_4o_UI2(3); + TFBIG_4e_UI2(4); + TFBIG_4o_UI2(5); + TFBIG_4e_UI2(6); + TFBIG_4o_UI2(7); + TFBIG_4e_UI2(8); + TFBIG_4o_UI2(9); + TFBIG_4e_UI2(10); + TFBIG_4o_UI2(11); + TFBIG_4e_UI2(12); + TFBIG_4o_UI2(13); + TFBIG_4e_UI2(14); + TFBIG_4o_UI2(15); + TFBIG_4e_UI2(16); + TFBIG_4o_UI2(17); + TFBIG_ADDKEY_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18); + + + t0 = vectorizelow(8); // extra + t1 = vectorizehigh(0xFF000000ul); // etype + + h0 = vectorize(c_PaddedMessage16[0]) ^ p[0]; + h1 = nounce2 ^ p[1]; + h2 = p[2]; + h3 = p[3]; + h4 = p[4]; + h5 = p[5]; + h6 = p[6]; + h7 = p[7]; + + h8 = h0 ^ h1 ^ p[2] ^ p[3] ^ p[4] ^ p[5] ^ p[6] ^ p[7] ^ vectorize(0x1BD11BDAA9FC1A22); + t2 = vectorize(0xFF00000000000008ull); + + // p[8] = { 0 }; + #pragma unroll 8 + for (int i = 0; i<8; i++) + p[i] = make_uint2(0, 0); + + TFBIG_4e_UI2(0); + TFBIG_4o_UI2(1); + TFBIG_4e_UI2(2); + TFBIG_4o_UI2(3); + TFBIG_4e_UI2(4); + TFBIG_4o_UI2(5); + TFBIG_4e_UI2(6); + TFBIG_4o_UI2(7); + TFBIG_4e_UI2(8); + TFBIG_4o_UI2(9); + TFBIG_4e_UI2(10); + TFBIG_4o_UI2(11); + TFBIG_4e_UI2(12); + TFBIG_4o_UI2(13); + TFBIG_4e_UI2(14); + TFBIG_4o_UI2(15); + TFBIG_4e_UI2(16); + TFBIG_4o_UI2(17); + TFBIG_ADDKEY_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18); + + uint32_t *message = (uint32_t *)p; + + uint32_t W1[16]; + uint32_t W2[16]; + + uint32_t regs[8]; + uint32_t hash[8]; + + // Init with Hash-Table +#pragma unroll 8 + for (int k = 0; k < 8; k++) + { + hash[k] = regs[k] = sha256_hashTable[k]; + } + +#pragma unroll 16 + for (int k = 0; k<16; k++) + W1[k] = SWAB32(message[k]); + + // Progress W1 +#pragma unroll 16 + for (int j = 0; j<16; j++) + { + uint32_t T1, T2; + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_constantTable[j] + W1[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int k = 6; k >= 0; k--) regs[k + 1] = regs[k]; + regs[0] = T1 + T2; + regs[4] += T1; + } + + // Progress W2...W3 + + ////// PART 1 +#pragma unroll 2 + for (int j = 0; j<2; j++) + W2[j] = s1(W1[14 + j]) + W1[9 + j] + s0(W1[1 + j]) + W1[j]; +#pragma unroll 5 + for (int j = 2; j<7; j++) + W2[j] = s1(W2[j - 2]) + W1[9 + j] + s0(W1[1 + j]) + W1[j]; + +#pragma unroll 8 + for (int j = 7; j<15; j++) + W2[j] = s1(W2[j - 2]) + W2[j - 7] + s0(W1[1 + j]) + W1[j]; + + W2[15] = s1(W2[13]) + W2[8] + s0(W2[0]) + W1[15]; + + // Round function +#pragma unroll 16 + for (int j = 0; j<16; j++) + { + uint32_t T1, T2; + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_constantTable[j + 16] + W2[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int l = 6; l >= 0; l--) regs[l + 1] = regs[l]; + regs[0] = T1 + T2; + regs[4] += T1; + } + + ////// PART 2 +#pragma unroll 2 + for (int j = 0; j<2; j++) + W1[j] = s1(W2[14 + j]) + W2[9 + j] + s0(W2[1 + j]) + W2[j]; + +#pragma unroll 5 + for (int j = 2; j<7; j++) + W1[j] = s1(W1[j - 2]) + W2[9 + j] + s0(W2[1 + j]) + W2[j]; + +#pragma unroll 8 + for (int j = 7; j<15; j++) + W1[j] = s1(W1[j - 2]) + W1[j - 7] + s0(W2[1 + j]) + W2[j]; + + W1[15] = s1(W1[13]) + W1[8] + s0(W1[0]) + W2[15]; + + // Round function +#pragma unroll 16 + for (int j = 0; j<16; j++) + { + uint32_t T1, T2; + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_constantTable[j + 32] + W1[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int l = 6; l >= 0; l--) regs[l + 1] = regs[l]; + regs[0] = T1 + T2; + regs[4] += T1; + } + + ////// PART 3 +#pragma unroll 2 + for (int j = 0; j<2; j++) + W2[j] = s1(W1[14 + j]) + W1[9 + j] + s0(W1[1 + j]) + W1[j]; + +#pragma unroll 5 + for (int j = 2; j<7; j++) + W2[j] = s1(W2[j - 2]) + W1[9 + j] + s0(W1[1 + j]) + W1[j]; + +#pragma unroll 8 + for (int j = 7; j<15; j++) + W2[j] = s1(W2[j - 2]) + W2[j - 7] + s0(W1[1 + j]) + W1[j]; + + W2[15] = s1(W2[13]) + W2[8] + s0(W2[0]) + W1[15]; + + // Round function +#pragma unroll 16 + for (int j = 0; j<16; j++) + { + uint32_t T1, T2; + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_constantTable[j + 48] + W2[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int l = 6; l >= 0; l--) regs[l + 1] = regs[l]; + regs[0] = T1 + T2; + regs[4] += T1; + } + +#pragma unroll 8 + for (int k = 0; k<8; k++) + hash[k] += regs[k]; + + ///// + ///// Second Pass (ending) + ///// +#pragma unroll 8 + for (int k = 0; k<8; k++) + regs[k] = hash[k]; + + // Progress W1 + uint32_t T1, T2; +#pragma unroll + for (int j = 0; j<56; j++) + { + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_endingTable[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int k = 6; k >= 0; k--) + regs[k + 1] = regs[k]; + regs[0] = T1 + T2; + regs[4] += T1; + } + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_endingTable[56]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + regs[7] = T1 + T2; + regs[3] += T1; + + T1 = regs[6] + S1(regs[3]) + Ch(regs[3], regs[4], regs[5]) + sha256_endingTable[57]; + T2 = S0(regs[7]) + Maj(regs[7], regs[0], regs[1]); + regs[6] = T1 + T2; + regs[2] += T1; + //************ + regs[1] += regs[5] + S1(regs[2]) + Ch(regs[2], regs[3], regs[4]) + sha256_endingTable[58]; + regs[0] += regs[4] + S1(regs[1]) + Ch(regs[1], regs[2], regs[3]) + sha256_endingTable[59]; + regs[7] += regs[3] + S1(regs[0]) + Ch(regs[0], regs[1], regs[2]) + sha256_endingTable[60]; + regs[6] += regs[2] + S1(regs[7]) + Ch(regs[7], regs[0], regs[1]) + sha256_endingTable[61]; + + uint64_t test = SWAB32(hash[7] + regs[7]); + test <<= 32; + test |= SWAB32(hash[6] + regs[6]); + if (test <= target) + { + uint32_t tmp = atomicExch(&(d_found[0]), startNounce + thread); + if (tmp != 0xffffffff) + d_found[1] = startNounce + thread; + } + } +} +__global__ +void skein512_gpu_hash_80_50(uint32_t threads, uint32_t startNounce, uint32_t *const __restrict__ d_found, uint64_t target) +{ + const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); + if (thread < threads) + { + uint2 h0, h1, h2, h3, h4, h5, h6, h7, h8; + uint2 t0, t1, t2; + uint2 p[8]; + + h0 = precalcvalues[0]; + h1 = precalcvalues[1]; + h2 = precalcvalues[2]; + h3 = precalcvalues[3]; + h4 = precalcvalues[4]; + h5 = precalcvalues[5]; + h6 = precalcvalues[6]; + h7 = precalcvalues[7]; + t2 = precalcvalues[8]; + + const uint2 nounce2 = make_uint2(_LOWORD(c_PaddedMessage16[1]), cuda_swab32(startNounce + thread)); + + // skein_big_close -> etype = 0x160, ptr = 16, bcount = 1, extra = 16 + p[0] = vectorize(c_PaddedMessage16[0]); + p[1] = nounce2; + +#pragma unroll + for (int i = 2; i < 8; i++) + p[i] = make_uint2(0, 0); + + t0 = vectorizelow(0x50ull); // SPH_T64(bcount << 6) + (sph_u64)(extra); + t1 = vectorizehigh(0xB0000000ul); // (bcount >> 58) + ((sph_u64)(etype) << 55); + TFBIG_KINIT_UI2(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2); + TFBIG_4e_UI2(0); + TFBIG_4o_UI2(1); + TFBIG_4e_UI2(2); + TFBIG_4o_UI2(3); + TFBIG_4e_UI2(4); + TFBIG_4o_UI2(5); + TFBIG_4e_UI2(6); + TFBIG_4o_UI2(7); + TFBIG_4e_UI2(8); + TFBIG_4o_UI2(9); + TFBIG_4e_UI2(10); + TFBIG_4o_UI2(11); + TFBIG_4e_UI2(12); + TFBIG_4o_UI2(13); + TFBIG_4e_UI2(14); + TFBIG_4o_UI2(15); + TFBIG_4e_UI2(16); + TFBIG_4o_UI2(17); + TFBIG_ADDKEY_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18); + + + t0 = vectorizelow(8); // extra + t1 = vectorizehigh(0xFF000000ul); // etype + + h0 = vectorize(c_PaddedMessage16[0]) ^ p[0]; + h1 = nounce2 ^ p[1]; + h2 = p[2]; + h3 = p[3]; + h4 = p[4]; + h5 = p[5]; + h6 = p[6]; + h7 = p[7]; + + h8 = h0 ^ h1 ^ p[2] ^ p[3] ^ p[4] ^ p[5] ^ p[6] ^ p[7] ^ vectorize(0x1BD11BDAA9FC1A22); + t2 = vectorize(0xFF00000000000008ull); + + // p[8] = { 0 }; +#pragma unroll 8 + for (int i = 0; i<8; i++) + p[i] = make_uint2(0, 0); + + TFBIG_4e_UI2(0); + TFBIG_4o_UI2(1); + TFBIG_4e_UI2(2); + TFBIG_4o_UI2(3); + TFBIG_4e_UI2(4); + TFBIG_4o_UI2(5); + TFBIG_4e_UI2(6); + TFBIG_4o_UI2(7); + TFBIG_4e_UI2(8); + TFBIG_4o_UI2(9); + TFBIG_4e_UI2(10); + TFBIG_4o_UI2(11); + TFBIG_4e_UI2(12); + TFBIG_4o_UI2(13); + TFBIG_4e_UI2(14); + TFBIG_4o_UI2(15); + TFBIG_4e_UI2(16); + TFBIG_4o_UI2(17); + TFBIG_ADDKEY_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18); + + uint32_t *message = (uint32_t *)p; + + uint32_t W1[16]; + uint32_t W2[16]; + + uint32_t regs[8]; + uint32_t hash[8]; + + // Init with Hash-Table +#pragma unroll 8 + for (int k = 0; k < 8; k++) + { + hash[k] = regs[k] = sha256_hashTable[k]; + } + +#pragma unroll 16 + for (int k = 0; k<16; k++) + W1[k] = SWAB32(message[k]); + + // Progress W1 +#pragma unroll 16 + for (int j = 0; j<16; j++) + { + uint32_t T1, T2; + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_constantTable[j] + W1[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int k = 6; k >= 0; k--) regs[k + 1] = regs[k]; + regs[0] = T1 + T2; + regs[4] += T1; + } + + // Progress W2...W3 + + ////// PART 1 +#pragma unroll 2 + for (int j = 0; j<2; j++) + W2[j] = s1(W1[14 + j]) + W1[9 + j] + s0(W1[1 + j]) + W1[j]; +#pragma unroll 5 + for (int j = 2; j<7; j++) + W2[j] = s1(W2[j - 2]) + W1[9 + j] + s0(W1[1 + j]) + W1[j]; + +#pragma unroll 8 + for (int j = 7; j<15; j++) + W2[j] = s1(W2[j - 2]) + W2[j - 7] + s0(W1[1 + j]) + W1[j]; + + W2[15] = s1(W2[13]) + W2[8] + s0(W2[0]) + W1[15]; + + // Round function +#pragma unroll 16 + for (int j = 0; j<16; j++) + { + uint32_t T1, T2; + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_constantTable[j + 16] + W2[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int l = 6; l >= 0; l--) regs[l + 1] = regs[l]; + regs[0] = T1 + T2; + regs[4] += T1; + } + + ////// PART 2 +#pragma unroll 2 + for (int j = 0; j<2; j++) + W1[j] = s1(W2[14 + j]) + W2[9 + j] + s0(W2[1 + j]) + W2[j]; + +#pragma unroll 5 + for (int j = 2; j<7; j++) + W1[j] = s1(W1[j - 2]) + W2[9 + j] + s0(W2[1 + j]) + W2[j]; + +#pragma unroll 8 + for (int j = 7; j<15; j++) + W1[j] = s1(W1[j - 2]) + W1[j - 7] + s0(W2[1 + j]) + W2[j]; + + W1[15] = s1(W1[13]) + W1[8] + s0(W1[0]) + W2[15]; + + // Round function +#pragma unroll 16 + for (int j = 0; j<16; j++) + { + uint32_t T1, T2; + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_constantTable[j + 32] + W1[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int l = 6; l >= 0; l--) regs[l + 1] = regs[l]; + regs[0] = T1 + T2; + regs[4] += T1; + } + + ////// PART 3 +#pragma unroll 2 + for (int j = 0; j<2; j++) + W2[j] = s1(W1[14 + j]) + W1[9 + j] + s0(W1[1 + j]) + W1[j]; + +#pragma unroll 5 + for (int j = 2; j<7; j++) + W2[j] = s1(W2[j - 2]) + W1[9 + j] + s0(W1[1 + j]) + W1[j]; + +#pragma unroll 8 + for (int j = 7; j<15; j++) + W2[j] = s1(W2[j - 2]) + W2[j - 7] + s0(W1[1 + j]) + W1[j]; + + W2[15] = s1(W2[13]) + W2[8] + s0(W2[0]) + W1[15]; + + // Round function +#pragma unroll 16 + for (int j = 0; j<16; j++) + { + uint32_t T1, T2; + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_constantTable[j + 48] + W2[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int l = 6; l >= 0; l--) regs[l + 1] = regs[l]; + regs[0] = T1 + T2; + regs[4] += T1; + } + +#pragma unroll 8 + for (int k = 0; k<8; k++) + hash[k] += regs[k]; + + ///// + ///// Second Pass (ending) + ///// +#pragma unroll 8 + for (int k = 0; k<8; k++) + regs[k] = hash[k]; + + // Progress W1 + uint32_t T1, T2; +#pragma unroll 1 + for (int j = 0; j<56; j++)//62 + { + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6]) + sha256_endingTable[j]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + +#pragma unroll 7 + for (int k = 6; k >= 0; k--) + regs[k + 1] = regs[k]; + regs[0] = T1 + T2; + regs[4] += T1; + } + T1 = regs[7] + S1(regs[4]) + Ch(regs[4], regs[5], regs[6])+sha256_endingTable[56]; + T2 = S0(regs[0]) + Maj(regs[0], regs[1], regs[2]); + regs[7] = T1 + T2; + regs[3] += T1; + + T1 = regs[6] + S1(regs[3]) + Ch(regs[3], regs[4], regs[5]) + sha256_endingTable[57]; + T2 = S0(regs[7]) + Maj(regs[7], regs[0], regs[1]); + regs[6] = T1 + T2; + regs[2] += T1; + //************ + regs[1] += regs[5] + S1(regs[2]) + Ch(regs[2], regs[3], regs[4]) + sha256_endingTable[58]; + regs[0] += regs[4] + S1(regs[1]) + Ch(regs[1], regs[2], regs[3]) + sha256_endingTable[59]; + regs[7] += regs[3] + S1(regs[0]) + Ch(regs[0], regs[1], regs[2]) + sha256_endingTable[60]; + regs[6] += regs[2] + S1(regs[7]) + Ch(regs[7], regs[0], regs[1]) + sha256_endingTable[61]; + + uint64_t test = SWAB32(hash[7] + regs[7]); + test <<= 32; + test|= SWAB32(hash[6] + regs[6]); + if (test <= target) + { + uint32_t tmp = atomicCAS(d_found, 0xffffffff, startNounce + thread); + if (tmp != 0xffffffff) + d_found[1] = startNounce + thread; + } + } +} + +static uint64_t PaddedMessage[16]; + +__host__ +static void precalc() +{ + uint64_t h0, h1, h2, h3, h4, h5, h6, h7, h8; + uint64_t t0, t1, t2; + + h0 = 0x4903ADFF749C51CEull; + h1 = 0x0D95DE399746DF03ull; + h2 = 0x8FD1934127C79BCEull; + h3 = 0x9A255629FF352CB1ull; + h4 = 0x5DB62599DF6CA7B0ull; + h5 = 0xEABE394CA9D5C3F4ull; + h6 = 0x991112C71A75B523ull; + h7 = 0xAE18A40B660FCC33ull; + h8 = h0 ^ h1 ^ h2 ^ h3 ^ h4 ^ h5 ^ h6 ^ h7 ^ SPH_C64(0x1BD11BDAA9FC1A22); + + t0 = 64; // ptr + t1 = 0x7000000000000000ull; + t2 = 0x7000000000000040ull; + + uint64_t p[8]; + for (int i = 0; i<8; i++) + p[i] = PaddedMessage[i]; + + TFBIG_4e_PRE(0); + TFBIG_4o_PRE(1); + TFBIG_4e_PRE(2); + TFBIG_4o_PRE(3); + TFBIG_4e_PRE(4); + TFBIG_4o_PRE(5); + TFBIG_4e_PRE(6); + TFBIG_4o_PRE(7); + TFBIG_4e_PRE(8); + TFBIG_4o_PRE(9); + TFBIG_4e_PRE(10); + TFBIG_4o_PRE(11); + TFBIG_4e_PRE(12); + TFBIG_4o_PRE(13); + TFBIG_4e_PRE(14); + TFBIG_4o_PRE(15); + TFBIG_4e_PRE(16); + TFBIG_4o_PRE(17); + TFBIG_ADDKEY_PRE(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18); + + uint64_t buffer[9]; + + buffer[0] = PaddedMessage[0] ^ p[0]; + buffer[1] = PaddedMessage[1] ^ p[1]; + buffer[2] = PaddedMessage[2] ^ p[2]; + buffer[3] = PaddedMessage[3] ^ p[3]; + buffer[4] = PaddedMessage[4] ^ p[4]; + buffer[5] = PaddedMessage[5] ^ p[5]; + buffer[6] = PaddedMessage[6] ^ p[6]; + buffer[7] = PaddedMessage[7] ^ p[7]; + buffer[8] = t2; + CUDA_SAFE_CALL(cudaMemcpyToSymbol(precalcvalues, buffer, sizeof(buffer), 0, cudaMemcpyHostToDevice)); + + uint32_t endingTable[] = { + 0x80000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, + 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000200, + 0x80000000, 0x01400000, 0x00205000, 0x00005088, 0x22000800, 0x22550014, 0x05089742, 0xa0000020, + 0x5a880000, 0x005c9400, 0x0016d49d, 0xfa801f00, 0xd33225d0, 0x11675959, 0xf6e6bfda, 0xb30c1549, + 0x08b2b050, 0x9d7c4c27, 0x0ce2a393, 0x88e6e1ea, 0xa52b4335, 0x67a16f49, 0xd732016f, 0x4eeb2e91, + 0x5dbf55e5, 0x8eee2335, 0xe2bc5ec2, 0xa83f4394, 0x45ad78f7, 0x36f3d0cd, 0xd99c05e8, 0xb0511dc7, + 0x69bc7ac4, 0xbd11375b, 0xe3ba71e5, 0x3b209ff2, 0x18feee17, 0xe25ad9e7, 0x13375046, 0x0515089d, + 0x4f0d0f04, 0x2627484e, 0x310128d2, 0xc668b434, 0x420841cc, 0x62d311b8, 0xe59ba771, 0x85a7a484 + }; + + uint32_t constantTable[64] = { + 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5, + 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174, + 0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da, + 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967, + 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85, + 0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070, + 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3, + 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2 + }; + for (int i = 0; i < 64; i++) + { + endingTable[i] = constantTable[i] + endingTable[i]; + } + CUDA_SAFE_CALL(cudaMemcpyToSymbol(sha256_endingTable, endingTable, sizeof(uint32_t) * 64, 0, cudaMemcpyHostToDevice)); + +} + + + +__host__ +void skein512_cpu_setBlock_80(uint32_t thr_id, void *pdata) +{ + memcpy(&PaddedMessage[0], pdata, 80); + + CUDA_SAFE_CALL( + cudaMemcpyToSymbol(c_PaddedMessage16, &PaddedMessage[8], 16, 0, cudaMemcpyHostToDevice) + ); + CUDA_SAFE_CALL(cudaMalloc(&(d_found[thr_id]), 3 * sizeof(uint32_t))); + + precalc(); +} + +__host__ +void skein512_cpu_hash_80_52(int thr_id, uint32_t threads, uint32_t startNounce, int swapu,uint64_t target, uint32_t *h_found) +{ + dim3 grid((threads + 1024 - 1) / 1024); + dim3 block(1024); + cudaMemset(d_found[thr_id], 0xffffffff, 2 * sizeof(uint32_t)); + skein512_gpu_hash_80_52 << < grid, block >> > (threads, startNounce, d_found[thr_id], target); + cudaMemcpy(h_found, d_found[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost); +} +__host__ +void skein512_cpu_hash_80_50(int thr_id, uint32_t threads, uint32_t startNounce, int swapu, uint64_t target, uint32_t *h_found) +{ + dim3 grid((threads + 256 - 1) / 256); + dim3 block(256); + cudaMemset(d_found[thr_id], 0xffffffff, 2 * sizeof(uint32_t)); + skein512_gpu_hash_80_50 << < grid, block >> > (threads, startNounce, d_found[thr_id], target); + cudaMemcpy(h_found, d_found[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost); +} + +#endif diff --git a/quark/cuda_vector_uint2x4.h b/quark/cuda_vector_uint2x4.h new file mode 100644 index 0000000..2245d3d --- /dev/null +++ b/quark/cuda_vector_uint2x4.h @@ -0,0 +1,547 @@ +/* todo: strip to keep only uint2x4 */ + +#ifndef CUDA_VECTOR_UINT2x4_H +#define CUDA_VECTOR_UINT2x4_H + +/////////////////////////////////////////////////////////////////////////////////// +#if (defined(_MSC_VER) && defined(_WIN64)) || defined(__LP64__) +#define __LDG_PTR "l" +#else +#define __LDG_PTR "r" +#endif + +#include "cuda_helper.h" + +#if __CUDA_ARCH__ < 320 && !defined(__ldg4) +#define __ldg4(x) (*(x)) +#endif + +typedef struct __align__(32) uint8 { + unsigned int s0, s1, s2, s3, s4, s5, s6, s7; +} uint8; + +typedef struct __align__(64) uint2_8 { + uint2 s0, s1, s2, s3, s4, s5, s6, s7; +} uint2_8; + +typedef struct __align__(64) ulonglong2to8 { + ulonglong2 l0,l1,l2,l3; +} ulonglong2to8; + +typedef struct __align__(128) ulonglong8to16 { + ulonglong2to8 lo, hi; +} ulonglong8to16; + +typedef struct __align__(256) ulonglong16to32 { + ulonglong8to16 lo, hi; +} ulonglong16to32; + +typedef struct __align__(512) ulonglong32to64 { + ulonglong16to32 lo, hi; +} ulonglong32to64; + +typedef struct __align__(128) ulonglonglong { + ulonglong2 s0,s1,s2,s3,s4,s5,s6,s7; +} ulonglonglong; + +typedef struct __align__(64) uint16 { + union { + struct {unsigned int s0, s1, s2, s3, s4, s5, s6, s7;}; + uint8 lo; + }; + union { + struct {unsigned int s8, s9, sa, sb, sc, sd, se, sf;}; + uint8 hi; + }; +} uint16; + +typedef struct __align__(128) uint2_16 { + union { + struct { uint2 s0, s1, s2, s3, s4, s5, s6, s7; }; + uint2_8 lo; + }; + union { + struct { uint2 s8, s9, sa, sb, sc, sd, se, sf; }; + uint2_8 hi; + }; +} uint2_16; + +typedef struct __align__(128) uint32 { + uint16 lo,hi; +} uint32; + +struct __align__(128) ulong8 { + ulonglong4 s0, s1, s2, s3; +}; +typedef __device_builtin__ struct ulong8 ulong8; + +typedef struct __align__(256) ulonglong16 { + ulonglong4 s0, s1, s2, s3, s4, s5, s6, s7; +} ulonglong16; + +typedef struct __align__(16) uint28 { + uint2 x, y, z, w; +} uint28; +typedef uint28 uint2x4; // proper name + +typedef struct __builtin_align__(32) uint48 { + uint4 s0,s1; +} uint48; +typedef uint48 uint4x2; // proper name + +typedef struct __align__(256) uint4x16 { + uint4 s0, s1, s2, s3, s4, s5, s6, s7, s8, s9, s10, s11, s12, s13, s14, s15; +} uint4x16; + +static __inline__ __device__ ulonglong2to8 make_ulonglong2to8(ulonglong2 s0, ulonglong2 s1, ulonglong2 s2, ulonglong2 s3) +{ + ulonglong2to8 t; t.l0=s0; t.l1=s1; t.l2=s2; t.l3=s3; + return t; +} + +static __inline__ __device__ ulonglong8to16 make_ulonglong8to16(const ulonglong2to8 &s0, const ulonglong2to8 &s1) +{ + ulonglong8to16 t; t.lo = s0; t.hi = s1; + return t; +} + +static __inline__ __device__ ulonglong16to32 make_ulonglong16to32(const ulonglong8to16 &s0, const ulonglong8to16 &s1) +{ + ulonglong16to32 t; t.lo = s0; t.hi = s1; + return t; +} + +static __inline__ __device__ ulonglong32to64 make_ulonglong32to64(const ulonglong16to32 &s0, const ulonglong16to32 &s1) +{ + ulonglong32to64 t; t.lo = s0; t.hi = s1; + return t; +} + +static __inline__ __host__ __device__ ulonglonglong make_ulonglonglong( + const ulonglong2 &s0, const ulonglong2 &s1, const ulonglong2 &s2, const ulonglong2 &s3, + const ulonglong2 &s4, const ulonglong2 &s5) +{ + ulonglonglong t; t.s0 = s0; t.s1 = s1; t.s2 = s2; t.s3 = s3; t.s4 = s4; t.s5 = s5; + return t; +} + +static __inline__ __device__ uint48 make_uint48(uint4 s0, uint4 s1) +{ + uint48 t; t.s0 = s0; t.s1 = s1; + return t; +} + +static __inline__ __device__ uint28 make_uint28(uint2 s0, uint2 s1, uint2 s2, uint2 s3) +{ + uint28 t; t.x = s0; t.y = s1; t.z = s2; t.w = s3; + return t; +} + +static __inline__ __host__ __device__ uint4x16 make_uint4x16( + uint4 s0, uint4 s1, uint4 s2, uint4 s3, uint4 s4, uint4 s5, uint4 s6, uint4 s7, + uint4 s8, uint4 s9, uint4 sa, uint4 sb, uint4 sc, uint4 sd, uint4 se, uint4 sf) +{ + uint4x16 t; t.s0 = s0; t.s1 = s1; t.s2 = s2; t.s3 = s3; t.s4 = s4; t.s5 = s5; t.s6 = s6; t.s7 = s7; + t.s8 = s8; t.s9 = s9; t.s10 = sa; t.s11 = sb; t.s12 = sc; t.s13 = sd; t.s14 = se; t.s15 = sf; + return t; +} + +static __inline__ __device__ uint2_16 make_uint2_16( + uint2 s0, uint2 s1, uint2 s2, uint2 s3, uint2 s4, uint2 s5, uint2 s6, uint2 s7, + uint2 s8, uint2 s9, uint2 sa, uint2 sb, uint2 sc, uint2 sd, uint2 se, uint2 sf) +{ + uint2_16 t; t.s0 = s0; t.s1 = s1; t.s2 = s2; t.s3 = s3; t.s4 = s4; t.s5 = s5; t.s6 = s6; t.s7 = s7; + t.s8 = s8; t.s9 = s9; t.sa = sa; t.sb = sb; t.sc = sc; t.sd = sd; t.se = se; t.sf = sf; + return t; +} + +static __inline__ __host__ __device__ uint16 make_uint16( + unsigned int s0, unsigned int s1, unsigned int s2, unsigned int s3, unsigned int s4, unsigned int s5, unsigned int s6, unsigned int s7, + unsigned int s8, unsigned int s9, unsigned int sa, unsigned int sb, unsigned int sc, unsigned int sd, unsigned int se, unsigned int sf) +{ + uint16 t; t.s0 = s0; t.s1 = s1; t.s2 = s2; t.s3 = s3; t.s4 = s4; t.s5 = s5; t.s6 = s6; t.s7 = s7; + t.s8 = s8; t.s9 = s9; t.sa = sa; t.sb = sb; t.sc = sc; t.sd = sd; t.se = se; t.sf = sf; + return t; +} + +static __inline__ __host__ __device__ uint16 make_uint16(const uint8 &a, const uint8 &b) +{ + uint16 t; t.lo=a; t.hi=b; return t; +} + +static __inline__ __host__ __device__ uint32 make_uint32(const uint16 &a, const uint16 &b) +{ + uint32 t; t.lo = a; t.hi = b; return t; +} + + +static __inline__ __host__ __device__ uint8 make_uint8( + unsigned int s0, unsigned int s1, unsigned int s2, unsigned int s3, unsigned int s4, unsigned int s5, unsigned int s6, unsigned int s7) +{ + uint8 t; t.s0 = s0; t.s1 = s1; t.s2 = s2; t.s3 = s3; t.s4 = s4; t.s5 = s5; t.s6 = s6; t.s7 = s7; + return t; +} + +static __inline__ __host__ __device__ uint2_8 make_uint2_8( + uint2 s0, uint2 s1, uint2 s2, uint2 s3, uint2 s4, uint2 s5, uint2 s6, uint2 s7) +{ + uint2_8 t; t.s0 = s0; t.s1 = s1; t.s2 = s2; t.s3 = s3; t.s4 = s4; t.s5 = s5; t.s6 = s6; t.s7 = s7; + return t; +} + +static __inline__ __host__ __device__ ulonglong16 make_ulonglong16(const ulonglong4 &s0, const ulonglong4 &s1, + const ulonglong4 &s2, const ulonglong4 &s3, const ulonglong4 &s4, const ulonglong4 &s5, const ulonglong4 &s6, const ulonglong4 &s7) +{ + ulonglong16 t; t.s0 = s0; t.s1 = s1; t.s2 = s2; t.s3 = s3; t.s4 = s4; t.s5 = s5; t.s6 = s6; t.s7 = s7; + return t; +} + +static __inline__ __host__ __device__ ulong8 make_ulong8( + ulonglong4 s0, ulonglong4 s1, ulonglong4 s2, ulonglong4 s3) +{ + ulong8 t; t.s0 = s0; t.s1 = s1; t.s2 = s2; t.s3 = s3;// t.s4 = s4; t.s5 = s5; t.s6 = s6; t.s7 = s7; + return t; +} + + +static __forceinline__ __device__ uchar4 operator^ (uchar4 a, uchar4 b) { return make_uchar4(a.x ^ b.x, a.y ^ b.y, a.z ^ b.z, a.w ^ b.w); } +static __forceinline__ __device__ uchar4 operator+ (uchar4 a, uchar4 b) { return make_uchar4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); } + +static __forceinline__ __device__ uint4 operator+ (uint4 a, uint4 b) { return make_uint4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); } + +static __forceinline__ __device__ ulonglong4 operator^ (ulonglong4 a, ulonglong4 b) { return make_ulonglong4(a.x ^ b.x, a.y ^ b.y, a.z ^ b.z, a.w ^ b.w); } +static __forceinline__ __device__ ulonglong4 operator+ (ulonglong4 a, ulonglong4 b) { return make_ulonglong4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); } +static __forceinline__ __device__ ulonglong2 operator^ (ulonglong2 a, ulonglong2 b) { return make_ulonglong2(a.x ^ b.x, a.y ^ b.y); } +static __forceinline__ __device__ ulonglong2 operator+ (ulonglong2 a, ulonglong2 b) { return make_ulonglong2(a.x + b.x, a.y + b.y); } + +static __forceinline__ __device__ ulong8 operator^ (const ulong8 &a, const ulong8 &b) { + return make_ulong8(a.s0 ^ b.s0, a.s1 ^ b.s1, a.s2 ^ b.s2, a.s3 ^ b.s3); +} + +static __forceinline__ __device__ ulong8 operator+ (const ulong8 &a, const ulong8 &b) { + return make_ulong8(a.s0 + b.s0, a.s1 + b.s1, a.s2 + b.s2, a.s3 + b.s3); +} + +static __forceinline__ __device__ __host__ uint8 operator^ (const uint8 &a, const uint8 &b) { return make_uint8(a.s0 ^ b.s0, a.s1 ^ b.s1, a.s2 ^ b.s2, a.s3 ^ b.s3, a.s4 ^ b.s4, a.s5 ^ b.s5, a.s6 ^ b.s6, a.s7 ^ b.s7); } + +static __forceinline__ __device__ __host__ uint8 operator+ (const uint8 &a, const uint8 &b) { return make_uint8(a.s0 + b.s0, a.s1 + b.s1, a.s2 + b.s2, a.s3 + b.s3, a.s4 + b.s4, a.s5 + b.s5, a.s6 + b.s6, a.s7 + b.s7); } + +static __forceinline__ __device__ uint2_8 operator^ (const uint2_8 &a, const uint2_8 &b) { return make_uint2_8(a.s0 ^ b.s0, a.s1 ^ b.s1, a.s2 ^ b.s2, a.s3 ^ b.s3, a.s4 ^ b.s4, a.s5 ^ b.s5, a.s6 ^ b.s6, a.s7 ^ b.s7); } + +static __forceinline__ __device__ uint2_8 operator+ (const uint2_8 &a, const uint2_8 &b) { return make_uint2_8(a.s0 + b.s0, a.s1 + b.s1, a.s2 + b.s2, a.s3 + b.s3, a.s4 + b.s4, a.s5 + b.s5, a.s6 + b.s6, a.s7 + b.s7); } + + +////////////// mess++ ////// + +static __forceinline__ __device__ uint28 operator^ (const uint28 &a, const uint28 &b) { + return make_uint28(a.x ^ b.x, a.y ^ b.y, a.z ^ b.z, a.w ^ b.w); +} + +static __forceinline__ __device__ uint28 operator+ (const uint28 &a, const uint28 &b) { + return make_uint28(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); +} + +static __forceinline__ __device__ uint48 operator+ (const uint48 &a, const uint48 &b) { + return make_uint48(a.s0 + b.s0, a.s1 + b.s1); +} + +///////////////////////// + +static __forceinline__ __device__ __host__ uint16 operator^ (const uint16 &a, const uint16 &b) { + return make_uint16(a.s0 ^ b.s0, a.s1 ^ b.s1, a.s2 ^ b.s2, a.s3 ^ b.s3, a.s4 ^ b.s4, a.s5 ^ b.s5, a.s6 ^ b.s6, a.s7 ^ b.s7, + a.s8 ^ b.s8, a.s9 ^ b.s9, a.sa ^ b.sa, a.sb ^ b.sb, a.sc ^ b.sc, a.sd ^ b.sd, a.se ^ b.se, a.sf ^ b.sf); +} + +static __forceinline__ __device__ __host__ uint16 operator+ (const uint16 &a, const uint16 &b) { + return make_uint16(a.s0 + b.s0, a.s1 + b.s1, a.s2 + b.s2, a.s3 + b.s3, a.s4 + b.s4, a.s5 + b.s5, a.s6 + b.s6, a.s7 + b.s7, + a.s8 + b.s8, a.s9 + b.s9, a.sa + b.sa, a.sb + b.sb, a.sc + b.sc, a.sd + b.sd, a.se + b.se, a.sf + b.sf); +} + +static __forceinline__ __device__ uint2_16 operator^ (const uint2_16 &a, const uint2_16 &b) { + return make_uint2_16(a.s0 ^ b.s0, a.s1 ^ b.s1, a.s2 ^ b.s2, a.s3 ^ b.s3, a.s4 ^ b.s4, a.s5 ^ b.s5, a.s6 ^ b.s6, a.s7 ^ b.s7, + a.s8 ^ b.s8, a.s9 ^ b.s9, a.sa ^ b.sa, a.sb ^ b.sb, a.sc ^ b.sc, a.sd ^ b.sd, a.se ^ b.se, a.sf ^ b.sf); +} + +static __forceinline__ __device__ uint2_16 operator+ (const uint2_16 &a, const uint2_16 &b) { + return make_uint2_16(a.s0 + b.s0, a.s1 + b.s1, a.s2 + b.s2, a.s3 + b.s3, a.s4 + b.s4, a.s5 + b.s5, a.s6 + b.s6, a.s7 + b.s7, + a.s8 + b.s8, a.s9 + b.s9, a.sa + b.sa, a.sb + b.sb, a.sc + b.sc, a.sd + b.sd, a.se + b.se, a.sf + b.sf); +} + +static __forceinline__ __device__ uint32 operator^ (const uint32 &a, const uint32 &b) { + return make_uint32(a.lo ^ b.lo, a.hi ^ b.hi); +} + +static __forceinline__ __device__ uint32 operator+ (const uint32 &a, const uint32 &b) { + return make_uint32(a.lo + b.lo, a.hi + b.hi); +} + +static __forceinline__ __device__ ulonglong16 operator^ (const ulonglong16 &a, const ulonglong16 &b) { + return make_ulonglong16(a.s0 ^ b.s0, a.s1 ^ b.s1, a.s2 ^ b.s2, a.s3 ^ b.s3, a.s4 ^ b.s4, a.s5 ^ b.s5, a.s6 ^ b.s6, a.s7 ^ b.s7); +} + +static __forceinline__ __device__ ulonglong16 operator+ (const ulonglong16 &a, const ulonglong16 &b) { + return make_ulonglong16(a.s0 + b.s0, a.s1 + b.s1, a.s2 + b.s2, a.s3 + b.s3, a.s4 + b.s4, a.s5 + b.s5, a.s6 + b.s6, a.s7 + b.s7); +} + +static __forceinline__ __device__ void operator^= (ulong8 &a, const ulong8 &b) { a = a ^ b; } + +static __forceinline__ __device__ void operator^= (uint28 &a, const uint28 &b) { a = a ^ b; } +static __forceinline__ __device__ void operator+= (uint28 &a, const uint28 &b) { a = a + b; } + +static __forceinline__ __device__ void operator^= (uint2_8 &a, const uint2_8 &b) { a = a ^ b; } +static __forceinline__ __device__ void operator+= (uint2_8 &a, const uint2_8 &b) { a = a + b; } + +static __forceinline__ __device__ void operator^= (uint32 &a, const uint32 &b) { a = a ^ b; } +static __forceinline__ __device__ void operator+= (uint32 &a, const uint32 &b) { a = a + b; } + +static __forceinline__ __device__ void operator^= (uchar4 &a, uchar4 b) { a = a ^ b; } + +static __forceinline__ __device__ __host__ void operator^= (uint8 &a, const uint8 &b) { a = a ^ b; } +static __forceinline__ __device__ __host__ void operator^= (uint16 &a, const uint16 &b) { a = a ^ b; } + +static __forceinline__ __device__ void operator^= (ulonglong16 &a, const ulonglong16 &b) { a = a ^ b; } +static __forceinline__ __device__ void operator^= (ulonglong4 &a, const ulonglong4 &b) { a = a ^ b; } +static __forceinline__ __device__ void operator+= (ulonglong4 &a, const ulonglong4 &b) { a = a + b; } + +static __forceinline__ __device__ void operator^= (ulonglong2 &a, const ulonglong2 &b) { a = a ^ b; } +static __forceinline__ __device__ void operator+= (ulonglong2 &a, const ulonglong2 &b) { a = a + b; } + +static __forceinline__ __device__ +ulonglong2to8 operator^ (const ulonglong2to8 &a, const ulonglong2to8 &b) +{ + return make_ulonglong2to8(a.l0 ^ b.l0, a.l1 ^ b.l1, a.l2 ^ b.l2, a.l3 ^ b.l3); +} +static __forceinline__ __device__ +ulonglong2to8 operator+ (const ulonglong2to8 &a, const ulonglong2to8 &b) +{ + return make_ulonglong2to8(a.l0 + b.l0, a.l1 + b.l1, a.l2 + b.l2, a.l3 + b.l3); +} + +static __forceinline__ __device__ +ulonglong8to16 operator^ (const ulonglong8to16 &a, const ulonglong8to16 &b) +{ + return make_ulonglong8to16(a.lo ^ b.lo, a.hi ^ b.hi); +} + +static __forceinline__ __device__ +ulonglong8to16 operator+ (const ulonglong8to16 &a, const ulonglong8to16 &b) +{ + return make_ulonglong8to16(a.lo + b.lo, a.hi + b.hi); +} + +static __forceinline__ __device__ +ulonglong16to32 operator^ (const ulonglong16to32 &a, const ulonglong16to32 &b) +{ + return make_ulonglong16to32(a.lo ^ b.lo, a.hi ^ b.hi); +} + +static __forceinline__ __device__ +ulonglong16to32 operator+ (const ulonglong16to32 &a, const ulonglong16to32 &b) +{ + return make_ulonglong16to32(a.lo + b.lo, a.hi + b.hi); +} + +static __forceinline__ __device__ +ulonglong32to64 operator^ (const ulonglong32to64 &a, const ulonglong32to64 &b) +{ + return make_ulonglong32to64(a.lo ^ b.lo, a.hi ^ b.hi); +} + +static __forceinline__ __device__ +ulonglong32to64 operator+ (const ulonglong32to64 &a, const ulonglong32to64 &b) +{ + return make_ulonglong32to64(a.lo + b.lo, a.hi + b.hi); +} + +static __forceinline__ __device__ ulonglonglong operator^ (const ulonglonglong &a, const ulonglonglong &b) { + return make_ulonglonglong(a.s0 ^ b.s0, a.s1 ^ b.s1, a.s2 ^ b.s2, a.s3 ^ b.s3, a.s4 ^ b.s4, a.s5 ^ b.s5); +} + +static __forceinline__ __device__ ulonglonglong operator+ (const ulonglonglong &a, const ulonglonglong &b) { + return make_ulonglonglong(a.s0 + b.s0, a.s1 + b.s1, a.s2 + b.s2, a.s3 + b.s3, a.s4 + b.s4, a.s5 + b.s5); +} + +static __forceinline__ __device__ void operator^= (ulonglong2to8 &a, const ulonglong2to8 &b) { a = a ^ b; } + +static __forceinline__ __device__ void operator+= (uint4 &a, uint4 b) { a = a + b; } +static __forceinline__ __device__ void operator+= (uchar4 &a, uchar4 b) { a = a + b; } +static __forceinline__ __device__ __host__ void operator+= (uint8 &a, const uint8 &b) { a = a + b; } +static __forceinline__ __device__ __host__ void operator+= (uint16 &a, const uint16 &b) { a = a + b; } +static __forceinline__ __device__ void operator+= (uint2_16 &a, const uint2_16 &b) { a = a + b; } +static __forceinline__ __device__ void operator^= (uint2_16 &a, const uint2_16 &b) { a = a + b; } + +static __forceinline__ __device__ void operator+= (ulong8 &a, const ulong8 &b) { a = a + b; } +static __forceinline__ __device__ void operator+= (ulonglong16 &a, const ulonglong16 &b) { a = a + b; } +static __forceinline__ __device__ void operator+= (ulonglong8to16 &a, const ulonglong8to16 &b) { a = a + b; } +static __forceinline__ __device__ void operator^= (ulonglong8to16 &a, const ulonglong8to16 &b) { a = a ^ b; } + +static __forceinline__ __device__ void operator+= (ulonglong16to32 &a, const ulonglong16to32 &b) { a = a + b; } +static __forceinline__ __device__ void operator^= (ulonglong16to32 &a, const ulonglong16to32 &b) { a = a ^ b; } + +static __forceinline__ __device__ void operator+= (ulonglong32to64 &a, const ulonglong32to64 &b) { a = a + b; } +static __forceinline__ __device__ void operator^= (ulonglong32to64 &a, const ulonglong32to64 &b) { a = a ^ b; } + +static __forceinline__ __device__ void operator+= (ulonglonglong &a, const ulonglonglong &b) { a = a + b; } +static __forceinline__ __device__ void operator^= (ulonglonglong &a, const ulonglonglong &b) { a = a ^ b; } + +#if __CUDA_ARCH__ < 320 + +#define rotate ROTL32 +#define rotateR ROTR32 + +#else + +static __forceinline__ __device__ uint4 rotate4(uint4 vec4, uint32_t shift) +{ + uint4 ret; + asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(ret.x) : "r"(vec4.x), "r"(vec4.x), "r"(shift)); + asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(ret.y) : "r"(vec4.y), "r"(vec4.y), "r"(shift)); + asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(ret.z) : "r"(vec4.z), "r"(vec4.z), "r"(shift)); + asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(ret.w) : "r"(vec4.w), "r"(vec4.w), "r"(shift)); + return ret; +} + +static __forceinline__ __device__ uint4 rotate4R(uint4 vec4, uint32_t shift) +{ + uint4 ret; + asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(ret.x) : "r"(vec4.x), "r"(vec4.x), "r"(shift)); + asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(ret.y) : "r"(vec4.y), "r"(vec4.y), "r"(shift)); + asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(ret.z) : "r"(vec4.z), "r"(vec4.z), "r"(shift)); + asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(ret.w) : "r"(vec4.w), "r"(vec4.w), "r"(shift)); + return ret; +} + +static __forceinline__ __device__ uint32_t rotate(uint32_t vec4, uint32_t shift) +{ + uint32_t ret; + asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(ret) : "r"(vec4), "r"(vec4), "r"(shift)); + return ret; +} + +static __forceinline__ __device__ uint32_t rotateR(uint32_t vec4, uint32_t shift) +{ + uint32_t ret; + asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(ret) : "r"(vec4), "r"(vec4), "r"(shift)); + return ret; +} + +static __device__ __inline__ ulonglong4 __ldg4(const ulonglong4 *ptr) +{ + ulonglong4 ret; + asm("ld.global.nc.v2.u64 {%0,%1}, [%2];" : "=l"(ret.x), "=l"(ret.y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v2.u64 {%0,%1}, [%2+16];" : "=l"(ret.z), "=l"(ret.w) : __LDG_PTR(ptr)); + return ret; +} + +static __device__ __inline__ void ldg4(const ulonglong4 *ptr,ulonglong4 *ret) +{ + asm("ld.global.nc.v2.u64 {%0,%1}, [%2];" : "=l"(ret[0].x), "=l"(ret[0].y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v2.u64 {%0,%1}, [%2+16];" : "=l"(ret[0].z), "=l"(ret[0].w) : __LDG_PTR(ptr)); + asm("ld.global.nc.v2.u64 {%0,%1}, [%2+32];" : "=l"(ret[1].x), "=l"(ret[1].y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v2.u64 {%0,%1}, [%2+48];" : "=l"(ret[1].z), "=l"(ret[1].w) : __LDG_PTR(ptr)); + asm("ld.global.nc.v2.u64 {%0,%1}, [%2+64];" : "=l"(ret[2].x), "=l"(ret[2].y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v2.u64 {%0,%1}, [%2+80];" : "=l"(ret[2].z), "=l"(ret[2].w) : __LDG_PTR(ptr)); +} + +static __device__ __inline__ uint28 __ldg4(const uint28 *ptr) +{ + uint28 ret; + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4];" : "=r"(ret.x.x), "=r"(ret.x.y), "=r"(ret.y.x), "=r"(ret.y.y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4+16];" : "=r"(ret.z.x), "=r"(ret.z.y), "=r"(ret.w.x), "=r"(ret.w.y) : __LDG_PTR(ptr)); + return ret; +} + +static __device__ __inline__ uint48 __ldg4(const uint48 *ptr) +{ + uint48 ret; + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4];" : "=r"(ret.s0.x), "=r"(ret.s0.y), "=r"(ret.s0.z), "=r"(ret.s0.w) : __LDG_PTR(ptr)); + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4+16];" : "=r"(ret.s1.x), "=r"(ret.s1.y), "=r"(ret.s1.z), "=r"(ret.s1.w) : __LDG_PTR(ptr)); + return ret; +} + +static __device__ __inline__ void ldg4(const uint28 *ptr, uint28 *ret) +{ + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4];" : "=r"(ret[0].x.x), "=r"(ret[0].x.y), "=r"(ret[0].y.x), "=r"(ret[0].y.y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4+16];" : "=r"(ret[0].z.x), "=r"(ret[0].z.y), "=r"(ret[0].w.x), "=r"(ret[0].w.y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4+32];" : "=r"(ret[1].x.x), "=r"(ret[1].x.y), "=r"(ret[1].y.x), "=r"(ret[1].y.y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4+48];" : "=r"(ret[1].z.x), "=r"(ret[1].z.y), "=r"(ret[1].w.x), "=r"(ret[1].w.y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4+64];" : "=r"(ret[2].x.x), "=r"(ret[2].x.y), "=r"(ret[2].y.x), "=r"(ret[2].y.y) : __LDG_PTR(ptr)); + asm("ld.global.nc.v4.u32 {%0,%1,%2,%3}, [%4+80];" : "=r"(ret[2].z.x), "=r"(ret[2].z.y), "=r"(ret[2].w.x), "=r"(ret[2].w.y) : __LDG_PTR(ptr)); +} + +#endif /* __CUDA_ARCH__ < 320 */ + + +static __forceinline__ __device__ uint8 swapvec(const uint8 &buf) +{ + uint8 vec; + vec.s0 = cuda_swab32(buf.s0); + vec.s1 = cuda_swab32(buf.s1); + vec.s2 = cuda_swab32(buf.s2); + vec.s3 = cuda_swab32(buf.s3); + vec.s4 = cuda_swab32(buf.s4); + vec.s5 = cuda_swab32(buf.s5); + vec.s6 = cuda_swab32(buf.s6); + vec.s7 = cuda_swab32(buf.s7); + return vec; +} + +static __forceinline__ __device__ uint8 swapvec(const uint8 *buf) +{ + uint8 vec; + vec.s0 = cuda_swab32(buf[0].s0); + vec.s1 = cuda_swab32(buf[0].s1); + vec.s2 = cuda_swab32(buf[0].s2); + vec.s3 = cuda_swab32(buf[0].s3); + vec.s4 = cuda_swab32(buf[0].s4); + vec.s5 = cuda_swab32(buf[0].s5); + vec.s6 = cuda_swab32(buf[0].s6); + vec.s7 = cuda_swab32(buf[0].s7); + return vec; +} + +static __forceinline__ __device__ uint16 swapvec(const uint16 *buf) +{ + uint16 vec; + vec.s0 = cuda_swab32(buf[0].s0); + vec.s1 = cuda_swab32(buf[0].s1); + vec.s2 = cuda_swab32(buf[0].s2); + vec.s3 = cuda_swab32(buf[0].s3); + vec.s4 = cuda_swab32(buf[0].s4); + vec.s5 = cuda_swab32(buf[0].s5); + vec.s6 = cuda_swab32(buf[0].s6); + vec.s7 = cuda_swab32(buf[0].s7); + vec.s8 = cuda_swab32(buf[0].s8); + vec.s9 = cuda_swab32(buf[0].s9); + vec.sa = cuda_swab32(buf[0].sa); + vec.sb = cuda_swab32(buf[0].sb); + vec.sc = cuda_swab32(buf[0].sc); + vec.sd = cuda_swab32(buf[0].sd); + vec.se = cuda_swab32(buf[0].se); + vec.sf = cuda_swab32(buf[0].sf); + return vec; +} + +static __forceinline__ __device__ uint16 swapvec(const uint16 &buf) +{ + uint16 vec; + vec.s0 = cuda_swab32(buf.s0); + vec.s1 = cuda_swab32(buf.s1); + vec.s2 = cuda_swab32(buf.s2); + vec.s3 = cuda_swab32(buf.s3); + vec.s4 = cuda_swab32(buf.s4); + vec.s5 = cuda_swab32(buf.s5); + vec.s6 = cuda_swab32(buf.s6); + vec.s7 = cuda_swab32(buf.s7); + vec.s8 = cuda_swab32(buf.s8); + vec.s9 = cuda_swab32(buf.s9); + vec.sa = cuda_swab32(buf.sa); + vec.sb = cuda_swab32(buf.sb); + vec.sc = cuda_swab32(buf.sc); + vec.sd = cuda_swab32(buf.sd); + vec.se = cuda_swab32(buf.se); + vec.sf = cuda_swab32(buf.sf); + return vec; +} + +#endif // H diff --git a/x11/x11.cu b/x11/x11.cu index ee33cc1..9040a4b 100644 --- a/x11/x11.cu +++ b/x11/x11.cu @@ -59,6 +59,8 @@ extern void x11_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t start // X11 CPU Hash extern "C" void x11hash(void *output, const void *input) { + unsigned char _ALIGN(128) hash[128] = { 0 }; + // blake1-bmw2-grs3-skein4-jh5-keccak6-luffa7-cubehash8-shavite9-simd10-echo11 sph_blake512_context ctx_blake; @@ -73,9 +75,6 @@ extern "C" void x11hash(void *output, const void *input) sph_simd512_context ctx_simd; sph_echo512_context ctx_echo; - unsigned char hash[128]; - memset(hash, 0, sizeof hash); - sph_blake512_init(&ctx_blake); sph_blake512 (&ctx_blake, input, 80); sph_blake512_close(&ctx_blake, (void*) hash);