heavy140330.cl

#define ROL32(x, n) rotate(x, (uint) n)

#define SWAP32(a) (as_uint(as_uchar4(a).wzyx))
#define SWAP64(n)       (as_ulong(as_uchar8(n).s76543210))


#define rol64hackl(n) \
inline ulong rol64_ ## n  (ulong l) \
{ \
    uint2 t = rotate(as_uint2(l), (n)); \
    return as_ulong((uint2)(bitselect(t.s0, t.s1, (uint)(1 << (n)) - 1), bitselect(t.s0, t.s1, (uint)(~((1 << (n)) - 1))))); \
}

#define rol64hackr(n) \
inline ulong rol64_ ## n  (ulong l) \
{ \
    uint2 t = rotate(as_uint2(l), (n - 32)); \
    return as_ulong((uint2)(bitselect(t.s1, t.s0, (uint)(1 << (n - 32)) - 1), bitselect(t.s1, t.s0, (uint)(~((1 << (n - 32)) - 1))))); \
}

#define rol64_32(l) as_ulong(as_uint2(l).yx)

rol64hackl(1)
rol64hackl(2)
rol64hackl(3)
rol64hackl(6)
rol64hackl(8)
rol64hackl(10)
rol64hackl(14)
rol64hackl(15)
rol64hackl(18)
rol64hackl(20)
rol64hackl(21)
rol64hackl(25)
rol64hackl(27)
rol64hackl(28)
rol64hackr(36)
rol64hackr(39)
rol64hackr(41)
rol64hackr(43)
rol64hackr(44)
rol64hackr(45)
rol64hackr(48)
rol64hackr(53)
rol64hackr(55)
rol64hackr(56)
rol64hackr(61)
rol64hackr(62)

__constant uint K[64] = {
    0x428a2f98U, 0x71374491U, 0xb5c0fbcfU, 0xe9b5dba5U,
    0x3956c25bU, 0x59f111f1U, 0x923f82a4U, 0xab1c5ed5U,
    0xd807aa98U, 0x12835b01U, 0x243185beU, 0x550c7dc3U,
    0x72be5d74U, 0x80deb1feU, 0x9bdc06a7U, 0xc19bf174U,
    0xe49b69c1U, 0xefbe4786U, 0x0fc19dc6U, 0x240ca1ccU,
    0x2de92c6fU, 0x4a7484aaU, 0x5cb0a9dcU, 0x76f988daU,
    0x983e5152U, 0xa831c66dU, 0xb00327c8U, 0xbf597fc7U,
    0xc6e00bf3U, 0xd5a79147U, 0x06ca6351U, 0x14292967U,
    0x27b70a85U, 0x2e1b2138U, 0x4d2c6dfcU, 0x53380d13U,
    0x650a7354U, 0x766a0abbU, 0x81c2c92eU, 0x92722c85U,
    0xa2bfe8a1U, 0xa81a664bU, 0xc24b8b70U, 0xc76c51a3U,
    0xd192e819U, 0xd6990624U, 0xf40e3585U, 0x106aa070U,
    0x19a4c116U, 0x1e376c08U, 0x2748774cU, 0x34b0bcb5U,
    0x391c0cb3U, 0x4ed8aa4aU, 0x5b9cca4fU, 0x682e6ff3U,
    0x748f82eeU, 0x78a5636fU, 0x84c87814U, 0x8cc70208U,
    0x90befffaU, 0xa4506cebU, 0xbef9a3f7U, 0xc67178f2U,
};

__constant uint H[] = {0x6a09e667U, 0xbb67ae85U, 0x3c6ef372U, 0xa54ff53aU,
                       0x510e527fU, 0x9b05688cU, 0x1f83d9abU, 0x5be0cd19U};

#define S0(x) (ROL32(x, 25) ^ ROL32(x, 14) ^  (x >> 3))
#define S1(x) (ROL32(x, 15) ^ ROL32(x, 13) ^  (x >> 10))

#define S2(x) (ROL32(x, 30) ^ ROL32(x, 19) ^ ROL32(x, 10))
#define S3(x) (ROL32(x, 26) ^ ROL32(x, 21) ^ ROL32(x, 7))

#define P(a,b,c,d,e,f,g,h,x,K)                  \
{                                               \
    temp1 = h + S3(e) + F1(e,f,g) + (K + x);      \
    d += temp1; h = temp1 + S2(a) + F0(a,b,c);  \
}

#define PLAST(a,b,c,d,e,f,g,h,x,K)                  \
{                                               \
    d += h + S3(e) + F1(e,f,g) + (x + K);              \
}

#define F0(y, x, z) bitselect(z, y, z ^ x)
#define F1(x, y, z) bitselect(z, y, x)

#define R0 (W0 = S1(W14) + W9 + S0(W1) + W0)
#define R1 (W1 = S1(W15) + W10 + S0(W2) + W1)
#define R2 (W2 = S1(W0) + W11 + S0(W3) + W2)
#define R3 (W3 = S1(W1) + W12 + S0(W4) + W3)
#define R4 (W4 = S1(W2) + W13 + S0(W5) + W4)
#define R5 (W5 = S1(W3) + W14 + S0(W6) + W5)
#define R6 (W6 = S1(W4) + W15 + S0(W7) + W6)
#define R7 (W7 = S1(W5) + W0 + S0(W8) + W7)
#define R8 (W8 = S1(W6) + W1 + S0(W9) + W8)
#define R9 (W9 = S1(W7) + W2 + S0(W10) + W9)
#define R10 (W10 = S1(W8) + W3 + S0(W11) + W10)
#define R11 (W11 = S1(W9) + W4 + S0(W12) + W11)
#define R12 (W12 = S1(W10) + W5 + S0(W13) + W12)
#define R13 (W13 = S1(W11) + W6 + S0(W14) + W13)
#define R14 (W14 = S1(W12) + W7 + S0(W15) + W14)
#define R15 (W15 = S1(W13) + W8 + S0(W0) + W15)
#define RD11 (S1(W9) + W4 + S0(W12) + W11)
#define RD12 (S1(W10) + W5 + S0(W13) + W12)
#define RD14 (W14 = S1(W12) + W7 + S0(W15) + W14)
#define RD15 (S1(W13) + W8 + S0(W0) + W15)

inline uint sha256_116_last(const uint *data)
{
    uint temp1;
    uint W0 = data[0];
    uint W1 = data[1];
    uint W2 = data[2];
    uint W3 = data[3];
    uint W4 = data[4];
    uint W5 = data[5];
    uint W6 = data[6];
    uint W7 = data[7];
    uint W8 = data[8];
    uint W9 = data[9];
    uint W10 = data[10];
    uint W11 = data[11];
    uint W12 = data[12];
    uint W13 = data[13];
    uint W14 = data[14];
    uint W15 = data[15];

    uint s7;

    uint v0 = 0x6A09E667;
    uint v1 = 0xBB67AE85;
    uint v2 = 0x3C6EF372;
    uint v3 = 0xA54FF53A;
    uint v4 = 0x510E527F;
    uint v5 = 0x9B05688C;
    uint v6 = 0x1F83D9AB;
    uint v7 = 0x5BE0CD19;

    P( v0, v1, v2, v3, v4, v5, v6, v7, W0, 0x428A2F98 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, W1, 0x71374491 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, W2, 0xB5C0FBCF );
    P( v5, v6, v7, v0, v1, v2, v3, v4, W3, 0xE9B5DBA5 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, W4, 0x3956C25B );
    P( v3, v4, v5, v6, v7, v0, v1, v2, W5, 0x59F111F1 );
    P( v2, v3, v4, v5, v6, v7, v0, v1, W6, 0x923F82A4 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, W7, 0xAB1C5ED5 );
    P( v0, v1, v2, v3, v4, v5, v6, v7, W8, 0xD807AA98 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, W9, 0x12835B01 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, W10, 0x243185BE );
    P( v5, v6, v7, v0, v1, v2, v3, v4, W11, 0x550C7DC3 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, W12, 0x72BE5D74 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, W13, 0x80DEB1FE );
    P( v2, v3, v4, v5, v6, v7, v0, v1, W14, 0x9BDC06A7 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, W15, 0xC19BF174 );

    P( v0, v1, v2, v3, v4, v5, v6, v7, R0, 0xE49B69C1 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R1, 0xEFBE4786 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R2, 0x0FC19DC6 );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R3, 0x240CA1CC );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R4, 0x2DE92C6F );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R5, 0x4A7484AA );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R6, 0x5CB0A9DC );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R7, 0x76F988DA );
    P( v0, v1, v2, v3, v4, v5, v6, v7, R8, 0x983E5152 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R9, 0xA831C66D );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R10, 0xB00327C8 );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R11, 0xBF597FC7 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R12, 0xC6E00BF3 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R13, 0xD5A79147 );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R14, 0x06CA6351 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R15, 0x14292967 );

    P( v0, v1, v2, v3, v4, v5, v6, v7, R0,  0x27B70A85 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R1,  0x2E1B2138 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R2,  0x4D2C6DFC );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R3,  0x53380D13 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R4,  0x650A7354 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R5,  0x766A0ABB );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R6,  0x81C2C92E );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R7,  0x92722C85 );
    P( v0, v1, v2, v3, v4, v5, v6, v7, R8,  0xA2BFE8A1 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R9,  0xA81A664B );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R10, 0xC24B8B70 );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R11, 0xC76C51A3 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R12, 0xD192E819 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R13, 0xD6990624 );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R14, 0xF40E3585 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R15, 0x106AA070 );

    P( v0, v1, v2, v3, v4, v5, v6, v7, R0,  0x19A4C116 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R1,  0x1E376C08 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R2,  0x2748774C );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R3,  0x34B0BCB5 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R4,  0x391C0CB3 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R5,  0x4ED8AA4A );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R6,  0x5B9CCA4F );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R7,  0x682E6FF3 );
    P( v0, v1, v2, v3, v4, v5, v6, v7, R8,  0x748F82EE );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R9,  0x78A5636F );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R10, 0x84C87814 );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R11, 0x8CC70208 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R12, 0x90BEFFFA );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R13, 0xA4506CEB );
    P( v2, v3, v4, v5, v6, v7, v0, v1, RD14, 0xBEF9A3F7 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, RD15, 0xC67178F2 );

    v0 += 0x6A09E667;
    v1 += 0xBB67AE85;
    v2 += 0x3C6EF372;
    v3 += 0xA54FF53A;
    v4 += 0x510E527F;
    v5 += 0x9B05688C;
    v6 += 0x1F83D9AB;
    v7 += 0x5BE0CD19;

    s7 = v7;

    W0 = data[16];
    W1 = data[17];
    W2 = data[18];
    W3 = data[19];
    W4 = data[20];
    W5 = data[21];
    W6 = data[22];
    W7 = data[23];
    W8 = data[24];
    W9 = data[25];
    W10 = data[26];
    W11 = data[27];
    W12 = data[28];
    W13 = 0x80000000;
    W14 = 0;
    W15 = 116*8;

    P( v0, v1, v2, v3, v4, v5, v6, v7, W0, 0x428A2F98 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, W1, 0x71374491 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, W2, 0xB5C0FBCF );
    P( v5, v6, v7, v0, v1, v2, v3, v4, W3, 0xE9B5DBA5 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, W4, 0x3956C25B );
    P( v3, v4, v5, v6, v7, v0, v1, v2, W5, 0x59F111F1 );
    P( v2, v3, v4, v5, v6, v7, v0, v1, W6, 0x923F82A4 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, W7, 0xAB1C5ED5 );
    P( v0, v1, v2, v3, v4, v5, v6, v7, W8, 0xD807AA98 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, W9, 0x12835B01 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, W10, 0x243185BE );
    P( v5, v6, v7, v0, v1, v2, v3, v4, W11, 0x550C7DC3 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, W12, 0x72BE5D74 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, W13, 0x80DEB1FE );
    P( v2, v3, v4, v5, v6, v7, v0, v1, W14, 0x9BDC06A7 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, W15, 0xC19BF174 );

    P( v0, v1, v2, v3, v4, v5, v6, v7, R0, 0xE49B69C1 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R1, 0xEFBE4786 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R2, 0x0FC19DC6 );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R3, 0x240CA1CC );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R4, 0x2DE92C6F );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R5, 0x4A7484AA );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R6, 0x5CB0A9DC );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R7, 0x76F988DA );
    P( v0, v1, v2, v3, v4, v5, v6, v7, R8, 0x983E5152 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R9, 0xA831C66D );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R10, 0xB00327C8 );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R11, 0xBF597FC7 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R12, 0xC6E00BF3 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R13, 0xD5A79147 );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R14, 0x06CA6351 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R15, 0x14292967 );

    P( v0, v1, v2, v3, v4, v5, v6, v7, R0,  0x27B70A85 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R1,  0x2E1B2138 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R2,  0x4D2C6DFC );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R3,  0x53380D13 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R4,  0x650A7354 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R5,  0x766A0ABB );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R6,  0x81C2C92E );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R7,  0x92722C85 );
    P( v0, v1, v2, v3, v4, v5, v6, v7, R8,  0xA2BFE8A1 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R9,  0xA81A664B );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R10, 0xC24B8B70 );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R11, 0xC76C51A3 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R12, 0xD192E819 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R13, 0xD6990624 );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R14, 0xF40E3585 );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R15, 0x106AA070 );

    P( v0, v1, v2, v3, v4, v5, v6, v7, R0,  0x19A4C116 );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R1,  0x1E376C08 );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R2,  0x2748774C );
    P( v5, v6, v7, v0, v1, v2, v3, v4, R3,  0x34B0BCB5 );
    P( v4, v5, v6, v7, v0, v1, v2, v3, R4,  0x391C0CB3 );
    P( v3, v4, v5, v6, v7, v0, v1, v2, R5,  0x4ED8AA4A );
    P( v2, v3, v4, v5, v6, v7, v0, v1, R6,  0x5B9CCA4F );
    P( v1, v2, v3, v4, v5, v6, v7, v0, R7,  0x682E6FF3 );
    P( v0, v1, v2, v3, v4, v5, v6, v7, R8,  0x748F82EE );
    P( v7, v0, v1, v2, v3, v4, v5, v6, R9,  0x78A5636F );
    P( v6, v7, v0, v1, v2, v3, v4, v5, R10, 0x84C87814 );
    P( v5, v6, v7, v0, v1, v2, v3, v4, RD11, 0x8CC70208 );
    PLAST( v4, v5, v6, v7, v0, v1, v2, v3, RD12, 0x90BEFFFA );

    v7 += s7;
    return SWAP32(v7);
}

#define BLAKE_CST0 0x243F6A8885A308D3UL
#define BLAKE_CST1 0x13198A2E03707344UL
#define BLAKE_CST2 0xA4093822299F31D0UL
#define BLAKE_CST3 0x082EFA98EC4E6C89UL
#define BLAKE_CST4 0x452821E638D01377UL
#define BLAKE_CST5 0xBE5466CF34E90C6CUL
#define BLAKE_CST6 0xC0AC29B7C97C50DDUL
#define BLAKE_CST7 0x3F84D5B5B5470917UL
#define BLAKE_CST8 0x9216D5D98979FB1BUL
#define BLAKE_CST9 0xD1310BA698DFB5ACUL
#define BLAKE_CSTA 0x2FFD72DBD01ADFB7UL
#define BLAKE_CSTB 0xB8E1AFED6A267E96UL
#define BLAKE_CSTC 0xBA7C9045F12C7F99UL
#define BLAKE_CSTD 0x24A19947B3916CF7UL
#define BLAKE_CSTE 0x0801F2E2858EFC16UL
#define BLAKE_CSTF 0x636920D871574E69UL

#define blake512_mut00_0(V, MSG) V##0 += (MSG##0 ^ BLAKE_CST1) + V##4; V##1 += (MSG##2 ^ BLAKE_CST3) + V##5; V##2 += (MSG##4 ^ BLAKE_CST5) + V##6; V##3 += (MSG##6 ^ BLAKE_CST7) + V##7;
#define blake512_mut00_1(V, MSG) V##0 += (MSG##E ^ BLAKE_CSTA) + V##4; V##1 += (MSG##4 ^ BLAKE_CST8) + V##5; V##2 += (MSG##9 ^ BLAKE_CSTF) + V##6; V##3 += (MSG##D ^ BLAKE_CST6) + V##7;
#define blake512_mut00_2(V, MSG) V##0 += (MSG##B ^ BLAKE_CST8) + V##4; V##1 += (MSG##C ^ BLAKE_CST0) + V##5; V##2 += (MSG##5 ^ BLAKE_CST2) + V##6; V##3 += (MSG##F ^ BLAKE_CSTD) + V##7;
#define blake512_mut00_3(V, MSG) V##0 += (MSG##7 ^ BLAKE_CST9) + V##4; V##1 += (MSG##3 ^ BLAKE_CST1) + V##5; V##2 += (MSG##D ^ BLAKE_CSTC) + V##6; V##3 += (MSG##B ^ BLAKE_CSTE) + V##7;
#define blake512_mut00_4(V, MSG) V##0 += (MSG##9 ^ BLAKE_CST0) + V##4; V##1 += (MSG##5 ^ BLAKE_CST7) + V##5; V##2 += (MSG##2 ^ BLAKE_CST4) + V##6; V##3 += (MSG##A ^ BLAKE_CSTF) + V##7;
#define blake512_mut00_5(V, MSG) V##0 += (MSG##2 ^ BLAKE_CSTC) + V##4; V##1 += (MSG##6 ^ BLAKE_CSTA) + V##5; V##2 += (MSG##0 ^ BLAKE_CSTB) + V##6; V##3 += (MSG##8 ^ BLAKE_CST3) + V##7;
#define blake512_mut00_6(V, MSG) V##0 += (MSG##C ^ BLAKE_CST5) + V##4; V##1 += (MSG##1 ^ BLAKE_CSTF) + V##5; V##2 += (MSG##E ^ BLAKE_CSTD) + V##6; V##3 += (MSG##4 ^ BLAKE_CSTA) + V##7;
#define blake512_mut00_7(V, MSG) V##0 += (MSG##D ^ BLAKE_CSTB) + V##4; V##1 += (MSG##7 ^ BLAKE_CSTE) + V##5; V##2 += (MSG##C ^ BLAKE_CST1) + V##6; V##3 += (MSG##3 ^ BLAKE_CST9) + V##7;
#define blake512_mut00_8(V, MSG) V##0 += (MSG##6 ^ BLAKE_CSTF) + V##4; V##1 += (MSG##E ^ BLAKE_CST9) + V##5; V##2 += (MSG##B ^ BLAKE_CST3) + V##6; V##3 += (MSG##0 ^ BLAKE_CST8) + V##7;
#define blake512_mut00_9(V, MSG) V##0 += (MSG##A ^ BLAKE_CST2) + V##4; V##1 += (MSG##8 ^ BLAKE_CST4) + V##5; V##2 += (MSG##7 ^ BLAKE_CST6) + V##6; V##3 += (MSG##1 ^ BLAKE_CST5) + V##7;

#define blake512_mut01_0(V, MSG) V##0 += (MSG##1 ^ BLAKE_CST0) + V##4; V##1 += (MSG##3 ^ BLAKE_CST2) + V##5; V##2 += (MSG##5 ^ BLAKE_CST4) + V##6; V##3 += (MSG##7 ^ BLAKE_CST6) + V##7;
#define blake512_mut01_1(V, MSG) V##0 += (MSG##A ^ BLAKE_CSTE) + V##4; V##1 += (MSG##8 ^ BLAKE_CST4) + V##5; V##2 += (MSG##F ^ BLAKE_CST9) + V##6; V##3 += (MSG##6 ^ BLAKE_CSTD) + V##7;
#define blake512_mut01_2(V, MSG) V##0 += (MSG##8 ^ BLAKE_CSTB) + V##4; V##1 += (MSG##0 ^ BLAKE_CSTC) + V##5; V##2 += (MSG##2 ^ BLAKE_CST5) + V##6; V##3 += (MSG##D ^ BLAKE_CSTF) + V##7;
#define blake512_mut01_3(V, MSG) V##0 += (MSG##9 ^ BLAKE_CST7) + V##4; V##1 += (MSG##1 ^ BLAKE_CST3) + V##5; V##2 += (MSG##C ^ BLAKE_CSTD) + V##6; V##3 += (MSG##E ^ BLAKE_CSTB) + V##7;
#define blake512_mut01_4(V, MSG) V##0 += (MSG##0 ^ BLAKE_CST9) + V##4; V##1 += (MSG##7 ^ BLAKE_CST5) + V##5; V##2 += (MSG##4 ^ BLAKE_CST2) + V##6; V##3 += (MSG##F ^ BLAKE_CSTA) + V##7;
#define blake512_mut01_5(V, MSG) V##0 += (MSG##C ^ BLAKE_CST2) + V##4; V##1 += (MSG##A ^ BLAKE_CST6) + V##5; V##2 += (MSG##B ^ BLAKE_CST0) + V##6; V##3 += (MSG##3 ^ BLAKE_CST8) + V##7;
#define blake512_mut01_6(V, MSG) V##0 += (MSG##5 ^ BLAKE_CSTC) + V##4; V##1 += (MSG##F ^ BLAKE_CST1) + V##5; V##2 += (MSG##D ^ BLAKE_CSTE) + V##6; V##3 += (MSG##A ^ BLAKE_CST4) + V##7;
#define blake512_mut01_7(V, MSG) V##0 += (MSG##B ^ BLAKE_CSTD) + V##4; V##1 += (MSG##E ^ BLAKE_CST7) + V##5; V##2 += (MSG##1 ^ BLAKE_CSTC) + V##6; V##3 += (MSG##9 ^ BLAKE_CST3) + V##7;
#define blake512_mut01_8(V, MSG) V##0 += (MSG##F ^ BLAKE_CST6) + V##4; V##1 += (MSG##9 ^ BLAKE_CSTE) + V##5; V##2 += (MSG##3 ^ BLAKE_CSTB) + V##6; V##3 += (MSG##8 ^ BLAKE_CST0) + V##7;
#define blake512_mut01_9(V, MSG) V##0 += (MSG##2 ^ BLAKE_CSTA) + V##4; V##1 += (MSG##4 ^ BLAKE_CST8) + V##5; V##2 += (MSG##6 ^ BLAKE_CST7) + V##6; V##3 += (MSG##5 ^ BLAKE_CST1) + V##7;

#define blake512_mut10_0(V, MSG) V##0 += (MSG##8 ^ BLAKE_CST9) + V##5; V##1 += (MSG##A ^ BLAKE_CSTB) + V##6; V##2 += (MSG##C ^ BLAKE_CSTD) + V##7; V##3 += (MSG##E ^ BLAKE_CSTF) + V##4;
#define blake512_mut10_1(V, MSG) V##0 += (MSG##1 ^ BLAKE_CSTC) + V##5; V##1 += (MSG##0 ^ BLAKE_CST2) + V##6; V##2 += (MSG##B ^ BLAKE_CST7) + V##7; V##3 += (MSG##5 ^ BLAKE_CST3) + V##4;
#define blake512_mut10_2(V, MSG) V##0 += (MSG##A ^ BLAKE_CSTE) + V##5; V##1 += (MSG##3 ^ BLAKE_CST6) + V##6; V##2 += (MSG##7 ^ BLAKE_CST1) + V##7; V##3 += (MSG##9 ^ BLAKE_CST4) + V##4;
#define blake512_mut10_3(V, MSG) V##0 += (MSG##2 ^ BLAKE_CST6) + V##5; V##1 += (MSG##5 ^ BLAKE_CSTA) + V##6; V##2 += (MSG##4 ^ BLAKE_CST0) + V##7; V##3 += (MSG##F ^ BLAKE_CST8) + V##4;
#define blake512_mut10_4(V, MSG) V##0 += (MSG##E ^ BLAKE_CST1) + V##5; V##1 += (MSG##B ^ BLAKE_CSTC) + V##6; V##2 += (MSG##6 ^ BLAKE_CST8) + V##7; V##3 += (MSG##3 ^ BLAKE_CSTD) + V##4;
#define blake512_mut10_5(V, MSG) V##0 += (MSG##4 ^ BLAKE_CSTD) + V##5; V##1 += (MSG##7 ^ BLAKE_CST5) + V##6; V##2 += (MSG##F ^ BLAKE_CSTE) + V##7; V##3 += (MSG##1 ^ BLAKE_CST9) + V##4;
#define blake512_mut10_6(V, MSG) V##0 += (MSG##0 ^ BLAKE_CST7) + V##5; V##1 += (MSG##6 ^ BLAKE_CST3) + V##6; V##2 += (MSG##9 ^ BLAKE_CST2) + V##7; V##3 += (MSG##8 ^ BLAKE_CSTB) + V##4;
#define blake512_mut10_7(V, MSG) V##0 += (MSG##5 ^ BLAKE_CST0) + V##5; V##1 += (MSG##F ^ BLAKE_CST4) + V##6; V##2 += (MSG##8 ^ BLAKE_CST6) + V##7; V##3 += (MSG##2 ^ BLAKE_CSTA) + V##4;
#define blake512_mut10_8(V, MSG) V##0 += (MSG##C ^ BLAKE_CST2) + V##5; V##1 += (MSG##D ^ BLAKE_CST7) + V##6; V##2 += (MSG##1 ^ BLAKE_CST4) + V##7; V##3 += (MSG##A ^ BLAKE_CST5) + V##4;
#define blake512_mut10_9(V, MSG) V##0 += (MSG##F ^ BLAKE_CSTB) + V##5; V##1 += (MSG##9 ^ BLAKE_CSTE) + V##6; V##2 += (MSG##3 ^ BLAKE_CSTC) + V##7; V##3 += (MSG##D ^ BLAKE_CST0) + V##4;

#define blake512_mut11_0(V, MSG) V##0 += (MSG##9 ^ BLAKE_CST8) + V##5; V##1 += (MSG##B ^ BLAKE_CSTA) + V##6; V##2 += (MSG##D ^ BLAKE_CSTC) + V##7; V##3 += (MSG##F ^ BLAKE_CSTE) + V##4;
#define blake512_mut11_1(V, MSG) V##0 += (MSG##C ^ BLAKE_CST1) + V##5; V##1 += (MSG##2 ^ BLAKE_CST0) + V##6; V##2 += (MSG##7 ^ BLAKE_CSTB) + V##7; V##3 += (MSG##3 ^ BLAKE_CST5) + V##4;
#define blake512_mut11_2(V, MSG) V##0 += (MSG##E ^ BLAKE_CSTA) + V##5; V##1 += (MSG##6 ^ BLAKE_CST3) + V##6; V##2 += (MSG##1 ^ BLAKE_CST7) + V##7; V##3 += (MSG##4 ^ BLAKE_CST9) + V##4;
#define blake512_mut11_3(V, MSG) V##0 += (MSG##6 ^ BLAKE_CST2) + V##5; V##1 += (MSG##A ^ BLAKE_CST5) + V##6; V##2 += (MSG##0 ^ BLAKE_CST4) + V##7; V##3 += (MSG##8 ^ BLAKE_CSTF) + V##4;
#define blake512_mut11_4(V, MSG) V##0 += (MSG##1 ^ BLAKE_CSTE) + V##5; V##1 += (MSG##C ^ BLAKE_CSTB) + V##6; V##2 += (MSG##8 ^ BLAKE_CST6) + V##7; V##3 += (MSG##D ^ BLAKE_CST3) + V##4;
#define blake512_mut11_5(V, MSG) V##0 += (MSG##D ^ BLAKE_CST4) + V##5; V##1 += (MSG##5 ^ BLAKE_CST7) + V##6; V##2 += (MSG##E ^ BLAKE_CSTF) + V##7; V##3 += (MSG##9 ^ BLAKE_CST1) + V##4;
#define blake512_mut11_6(V, MSG) V##0 += (MSG##7 ^ BLAKE_CST0) + V##5; V##1 += (MSG##3 ^ BLAKE_CST6) + V##6; V##2 += (MSG##2 ^ BLAKE_CST9) + V##7; V##3 += (MSG##B ^ BLAKE_CST8) + V##4;
#define blake512_mut11_7(V, MSG) V##0 += (MSG##0 ^ BLAKE_CST5) + V##5; V##1 += (MSG##4 ^ BLAKE_CSTF) + V##6; V##2 += (MSG##6 ^ BLAKE_CST8) + V##7; V##3 += (MSG##A ^ BLAKE_CST2) + V##4;
#define blake512_mut11_8(V, MSG) V##0 += (MSG##2 ^ BLAKE_CSTC) + V##5; V##1 += (MSG##7 ^ BLAKE_CSTD) + V##6; V##2 += (MSG##4 ^ BLAKE_CST1) + V##7; V##3 += (MSG##5 ^ BLAKE_CSTA) + V##4;
#define blake512_mut11_9(V, MSG) V##0 += (MSG##B ^ BLAKE_CSTF) + V##5; V##1 += (MSG##E ^ BLAKE_CST9) + V##6; V##2 += (MSG##C ^ BLAKE_CST3) + V##7; V##3 += (MSG##0 ^ BLAKE_CSTD) + V##4;

#define blake512GV(i, V, MSG)                    \
    blake512_mut00_##i(V, MSG) \
    V##C = rol64_32(V##C ^ V##0);              \
    V##D = rol64_32(V##D ^ V##1);              \
    V##E = rol64_32(V##E ^ V##2);              \
    V##F = rol64_32(V##F ^ V##3);              \
    V##8 += V##C; V##9 += V##D; V##A += V##E; V##B += V##F;                     \
    V##4 ^= V##8; V##5 ^= V##9; V##6 ^= V##A; V##7 ^= V##B;                     \
    V##4 = rol64_39(V##4);              \
    V##5 = rol64_39(V##5);              \
    V##6 = rol64_39(V##6);              \
    V##7 = rol64_39(V##7);              \
    blake512_mut01_##i(V, MSG)   \
    V##C ^= V##0; V##D ^= V##1; V##E ^= V##2; V##F ^= V##3; \
    V##C = rol64_48(V##C);              \
    V##D = rol64_48(V##D);              \
    V##E = rol64_48(V##E);              \
    V##F = rol64_48(V##F);              \
    V##8 += V##C; V##9 += V##D; V##A += V##E; V##B += V##F;                     \
    V##4 ^= V##8; V##5 ^= V##9; V##6 ^= V##A; V##7 ^= V##B; \
    V##4 = rol64_53(V##4); \
    V##5 = rol64_53(V##5); \
    V##6 = rol64_53(V##6); \
    V##7 = rol64_53(V##7); \
    blake512_mut10_##i(V, MSG) \
    V##C = rol64_32(V##C ^ V##1);              \
    V##D = rol64_32(V##D ^ V##2);              \
    V##E = rol64_32(V##E ^ V##3);              \
    V##F = rol64_32(V##F ^ V##0);              \
    V##8 += V##D; V##9 += V##E; V##A += V##F; V##B += V##C;                     \
    V##4 ^= V##9; V##5 ^= V##A; V##6 ^= V##B; V##7 ^= V##8; \
    V##4 = rol64_39(V##4);              \
    V##5 = rol64_39(V##5);              \
    V##6 = rol64_39(V##6);              \
    V##7 = rol64_39(V##7);              \
    blake512_mut11_##i(V, MSG)  \
    V##C ^= V##1; V##D ^= V##2; V##E ^= V##3; V##F ^= V##0; \
    V##C = rol64_48(V##C);              \
    V##D = rol64_48(V##D);              \
    V##E = rol64_48(V##E);              \
    V##F = rol64_48(V##F);              \
    V##8 += V##D; V##9 += V##E; V##A += V##F; V##B += V##C;                     \
    V##5 ^= V##A; V##6 ^= V##B; V##7 ^= V##8; V##4 ^= V##9; \
    V##4 = rol64_53(V##4); \
    V##5 = rol64_53(V##5); \
    V##6 = rol64_53(V##6); \
    V##7 = rol64_53(V##7);

inline uint blake512_116_last(const uint *msg)
{
    ulong MSG0, MSG1, MSG2, MSG3, MSG4, MSG5, MSG6, MSG7, MSG8, MSG9, MSGA, MSGB, MSGC, MSGD, MSGE, MSGF;
    ulong V0, V1, V2, V3, V4, V5, V6, V7, V8, V9, VA, VB, VC, VD, VE, VF;
    ulong H3;

    MSG0 = as_ulong((uint2)(msg[1], msg[0]));
    MSG1 = as_ulong((uint2)(msg[3], msg[2]));
    MSG2 = as_ulong((uint2)(msg[5], msg[4]));
    MSG3 = as_ulong((uint2)(msg[7], msg[6]));
    MSG4 = as_ulong((uint2)(msg[9], msg[8]));
    MSG5 = as_ulong((uint2)(msg[11], msg[10]));
    MSG6 = as_ulong((uint2)(msg[13], msg[12]));
    MSG7 = as_ulong((uint2)(msg[15], msg[14]));
    MSG8 = as_ulong((uint2)(msg[17], msg[16]));
    MSG9 = as_ulong((uint2)(msg[19], msg[18]));
    MSGA = as_ulong((uint2)(msg[21], msg[20]));
    MSGB = as_ulong((uint2)(msg[23], msg[22]));
    MSGC = as_ulong((uint2)(msg[25], msg[24]));
    MSGD = as_ulong((uint2)(msg[27], msg[26]));
    MSGE = as_ulong((uint2)(0x80000000, msg[28]));
    MSGF = 0;

    V0 = 0x6A09E667F3BCC908UL;
    V1 = 0xBB67AE8584CAA73BUL;
    V2 = 0x3C6EF372FE94F82BUL;
    V3 = 0xA54FF53A5F1D36F1UL;
    V4 = 0x510E527FADE682D1UL;
    V5 = 0x9B05688C2B3E6C1FUL;
    V6 = 0x1F83D9ABFB41BD6BUL;
    V7 = 0x5BE0CD19137E2179UL;
    V8 = 0x243F6A8885A308D3UL;
    V9 = 0x13198A2E03707344UL;
    VA = 0xA4093822299F31D0UL;
    VB = 0x082EFA98EC4E6C89UL;
    VC = 0x452821E638D01377UL ^ 0x3A0UL;
    VD = 0xBE5466CF34E90C6CUL ^ 0x3A0UL;
    VE = 0xC0AC29B7C97C50DDUL;
    VF = 0x3F84D5B5B5470917UL;

    blake512GV(0, V, MSG);
    blake512GV(1, V, MSG);
    blake512GV(2, V, MSG);
    blake512GV(3, V, MSG);
    blake512GV(4, V, MSG);
    blake512GV(5, V, MSG);
    blake512GV(6, V, MSG);
    blake512GV(7, V, MSG);
    blake512GV(8, V, MSG);
    blake512GV(9, V, MSG);

    blake512GV(0, V, MSG);
    blake512GV(1, V, MSG);
    blake512GV(2, V, MSG);
    blake512GV(3, V, MSG);
    blake512GV(4, V, MSG);
    blake512GV(5, V, MSG);

    V0 ^= V8 ^ 0x6A09E667F3BCC908UL;
    V1 ^= V9 ^ 0xBB67AE8584CAA73BUL;
    V2 ^= VA ^ 0x3C6EF372FE94F82BUL;
    H3 = V3 ^= VB ^ 0xA54FF53A5F1D36F1UL;
    V4 ^= VC ^ 0x510E527FADE682D1UL;
    V5 ^= VD ^ 0x9B05688C2B3E6C1FUL;
    V6 ^= VE ^ 0x1F83D9ABFB41BD6BUL;
    V7 ^= VF ^ 0x5BE0CD19137E2179UL;

    MSG0 = MSG1 = MSG2 = MSG3 = MSG4 = MSG5 = MSG6 = MSG7 = MSG8 = MSG9 = MSGA = MSGB = MSGC = 0;
    MSGD = 1UL;
    MSGE = 0;
    MSGF = 0x3A0;

    V8 = 0x243F6A8885A308D3UL;
    V9 = 0x13198A2E03707344UL;
    VA = 0xA4093822299F31D0UL;
    VB = 0x082EFA98EC4E6C89UL;
    VC = 0x452821E638D01377UL;
    VD = 0xBE5466CF34E90C6CUL;
    VE = 0xC0AC29B7C97C50DDUL;
    VF = 0x3F84D5B5B5470917UL;

    blake512GV(0, V, MSG);
    blake512GV(1, V, MSG);
    blake512GV(2, V, MSG);
    blake512GV(3, V, MSG);
    blake512GV(4, V, MSG);
    blake512GV(5, V, MSG);
    blake512GV(6, V, MSG);
    blake512GV(7, V, MSG);
    blake512GV(8, V, MSG);
    blake512GV(9, V, MSG);

    blake512GV(0, V, MSG);
    blake512GV(1, V, MSG);
    blake512GV(2, V, MSG);
    blake512GV(3, V, MSG);
    blake512GV(4, V, MSG);
    blake512GV(5, V, MSG);

    H3 ^= V3 ^ VB;
    return as_uint2(SWAP64(H3)).y;
}

__constant ulong GROESTL_T0[256] = {
    0xC6A597F4A5F432C6UL,0xF884EB9784976FF8UL,0xEE99C7B099B05EEEUL,0xF68DF78C8D8C7AF6UL,
    0xFF0DE5170D17E8FFUL,0xD6BDB7DCBDDC0AD6UL,0xDEB1A7C8B1C816DEUL,0x915439FC54FC6D91UL,
    0x6050C0F050F09060UL,0x0203040503050702UL,0xCEA987E0A9E02ECEUL,0x567DAC877D87D156UL,
    0xE719D52B192BCCE7UL,0xB56271A662A613B5UL,0x4DE69A31E6317C4DUL,0xEC9AC3B59AB559ECUL,
    0x8F4505CF45CF408FUL,0x1F9D3EBC9DBCA31FUL,0x894009C040C04989UL,0xFA87EF92879268FAUL,
    0xEF15C53F153FD0EFUL,0xB2EB7F26EB2694B2UL,0x8EC90740C940CE8EUL,0xFB0BED1D0B1DE6FBUL,
    0x41EC822FEC2F6E41UL,0xB3677DA967A91AB3UL,0x5FFDBE1CFD1C435FUL,0x45EA8A25EA256045UL,
    0x23BF46DABFDAF923UL,0x53F7A602F7025153UL,0xE496D3A196A145E4UL,0x9B5B2DED5BED769BUL,
    0x75C2EA5DC25D2875UL,0xE11CD9241C24C5E1UL,0x3DAE7AE9AEE9D43DUL,0x4C6A98BE6ABEF24CUL,
    0x6C5AD8EE5AEE826CUL,0x7E41FCC341C3BD7EUL,0xF502F1060206F3F5UL,0x834F1DD14FD15283UL,
    0x685CD0E45CE48C68UL,0x51F4A207F4075651UL,0xD134B95C345C8DD1UL,0xF908E9180818E1F9UL,
    0xE293DFAE93AE4CE2UL,0xAB734D9573953EABUL,0x6253C4F553F59762UL,0x2A3F54413F416B2AUL,
    0x080C10140C141C08UL,0x955231F652F66395UL,0x46658CAF65AFE946UL,0x9D5E21E25EE27F9DUL,
    0x3028607828784830UL,0x37A16EF8A1F8CF37UL,0x0A0F14110F111B0AUL,0x2FB55EC4B5C4EB2FUL,
    0x0E091C1B091B150EUL,0x2436485A365A7E24UL,0x1B9B36B69BB6AD1BUL,0xDF3DA5473D4798DFUL,
    0xCD26816A266AA7CDUL,0x4E699CBB69BBF54EUL,0x7FCDFE4CCD4C337FUL,0xEA9FCFBA9FBA50EAUL,
    0x121B242D1B2D3F12UL,0x1D9E3AB99EB9A41DUL,0x5874B09C749CC458UL,0x342E68722E724634UL,
    0x362D6C772D774136UL,0xDCB2A3CDB2CD11DCUL,0xB4EE7329EE299DB4UL,0x5BFBB616FB164D5BUL,
    0xA4F65301F601A5A4UL,0x764DECD74DD7A176UL,0xB76175A361A314B7UL,0x7DCEFA49CE49347DUL,
    0x527BA48D7B8DDF52UL,0xDD3EA1423E429FDDUL,0x5E71BC937193CD5EUL,0x139726A297A2B113UL,
    0xA6F55704F504A2A6UL,0xB96869B868B801B9UL,0x0000000000000000UL,0xC12C99742C74B5C1UL,
    0x406080A060A0E040UL,0xE31FDD211F21C2E3UL,0x79C8F243C8433A79UL,0xB6ED772CED2C9AB6UL,
    0xD4BEB3D9BED90DD4UL,0x8D4601CA46CA478DUL,0x67D9CE70D9701767UL,0x724BE4DD4BDDAF72UL,
    0x94DE3379DE79ED94UL,0x98D42B67D467FF98UL,0xB0E87B23E82393B0UL,0x854A11DE4ADE5B85UL,
    0xBB6B6DBD6BBD06BBUL,0xC52A917E2A7EBBC5UL,0x4FE59E34E5347B4FUL,0xED16C13A163AD7EDUL,
    0x86C51754C554D286UL,0x9AD72F62D762F89AUL,0x6655CCFF55FF9966UL,0x119422A794A7B611UL,
    0x8ACF0F4ACF4AC08AUL,0xE910C9301030D9E9UL,0x0406080A060A0E04UL,0xFE81E798819866FEUL,
    0xA0F05B0BF00BABA0UL,0x7844F0CC44CCB478UL,0x25BA4AD5BAD5F025UL,0x4BE3963EE33E754BUL,
    0xA2F35F0EF30EACA2UL,0x5DFEBA19FE19445DUL,0x80C01B5BC05BDB80UL,0x058A0A858A858005UL,
    0x3FAD7EECADECD33FUL,0x21BC42DFBCDFFE21UL,0x7048E0D848D8A870UL,0xF104F90C040CFDF1UL,
    0x63DFC67ADF7A1963UL,0x77C1EE58C1582F77UL,0xAF75459F759F30AFUL,0x426384A563A5E742UL,
    0x2030405030507020UL,0xE51AD12E1A2ECBE5UL,0xFD0EE1120E12EFFDUL,0xBF6D65B76DB708BFUL,
    0x814C19D44CD45581UL,0x1814303C143C2418UL,0x26354C5F355F7926UL,0xC32F9D712F71B2C3UL,
    0xBEE16738E13886BEUL,0x35A26AFDA2FDC835UL,0x88CC0B4FCC4FC788UL,0x2E395C4B394B652EUL,
    0x93573DF957F96A93UL,0x55F2AA0DF20D5855UL,0xFC82E39D829D61FCUL,0x7A47F4C947C9B37AUL,
    0xC8AC8BEFACEF27C8UL,0xBAE76F32E73288BAUL,0x322B647D2B7D4F32UL,0xE695D7A495A442E6UL,
    0xC0A09BFBA0FB3BC0UL,0x199832B398B3AA19UL,0x9ED12768D168F69EUL,0xA37F5D817F8122A3UL,
    0x446688AA66AAEE44UL,0x547EA8827E82D654UL,0x3BAB76E6ABE6DD3BUL,0x0B83169E839E950BUL,
    0x8CCA0345CA45C98CUL,0xC729957B297BBCC7UL,0x6BD3D66ED36E056BUL,0x283C50443C446C28UL,
    0xA779558B798B2CA7UL,0xBCE2633DE23D81BCUL,0x161D2C271D273116UL,0xAD76419A769A37ADUL,
    0xDB3BAD4D3B4D96DBUL,0x6456C8FA56FA9E64UL,0x744EE8D24ED2A674UL,0x141E28221E223614UL,
    0x92DB3F76DB76E492UL,0x0C0A181E0A1E120CUL,0x486C90B46CB4FC48UL,0xB8E46B37E4378FB8UL,
    0x9F5D25E75DE7789FUL,0xBD6E61B26EB20FBDUL,0x43EF862AEF2A6943UL,0xC4A693F1A6F135C4UL,
    0x39A872E3A8E3DA39UL,0x31A462F7A4F7C631UL,0xD337BD5937598AD3UL,0xF28BFF868B8674F2UL,
    0xD532B156325683D5UL,0x8B430DC543C54E8BUL,0x6E59DCEB59EB856EUL,0xDAB7AFC2B7C218DAUL,
    0x018C028F8C8F8E01UL,0xB16479AC64AC1DB1UL,0x9CD2236DD26DF19CUL,0x49E0923BE03B7249UL,
    0xD8B4ABC7B4C71FD8UL,0xACFA4315FA15B9ACUL,0xF307FD090709FAF3UL,0xCF25856F256FA0CFUL,
    0xCAAF8FEAAFEA20CAUL,0xF48EF3898E897DF4UL,0x47E98E20E9206747UL,0x1018202818283810UL,
    0x6FD5DE64D5640B6FUL,0xF088FB83888373F0UL,0x4A6F94B16FB1FB4AUL,0x5C72B8967296CA5CUL,
    0x3824706C246C5438UL,0x57F1AE08F1085F57UL,0x73C7E652C7522173UL,0x975135F351F36497UL,
    0xCB238D652365AECBUL,0xA17C59847C8425A1UL,0xE89CCBBF9CBF57E8UL,0x3E217C6321635D3EUL,
    0x96DD377CDD7CEA96UL,0x61DCC27FDC7F1E61UL,0x0D861A9186919C0DUL,0x0F851E9485949B0FUL,
    0xE090DBAB90AB4BE0UL,0x7C42F8C642C6BA7CUL,0x71C4E257C4572671UL,0xCCAA83E5AAE529CCUL,
    0x90D83B73D873E390UL,0x06050C0F050F0906UL,0xF701F5030103F4F7UL,0x1C12383612362A1CUL,
    0xC2A39FFEA3FE3CC2UL,0x6A5FD4E15FE18B6AUL,0xAEF94710F910BEAEUL,0x69D0D26BD06B0269UL,
    0x17912EA891A8BF17UL,0x995829E858E87199UL,0x3A2774692769533AUL,0x27B94ED0B9D0F727UL,
    0xD938A948384891D9UL,0xEB13CD351335DEEBUL,0x2BB356CEB3CEE52BUL,0x2233445533557722UL,
    0xD2BBBFD6BBD604D2UL,0xA9704990709039A9UL,0x07890E8089808707UL,0x33A766F2A7F2C133UL,
    0x2DB65AC1B6C1EC2DUL,0x3C22786622665A3CUL,0x15922AAD92ADB815UL,0xC92089602060A9C9UL,
    0x874915DB49DB5C87UL,0xAAFF4F1AFF1AB0AAUL,0x5078A0887888D850UL,0xA57A518E7A8E2BA5UL,
    0x038F068A8F8A8903UL,0x59F8B213F8134A59UL,0x0980129B809B9209UL,0x1A1734391739231AUL,
    0x65DACA75DA751065UL,0xD731B553315384D7UL,0x84C61351C651D584UL,0xD0B8BBD3B8D303D0UL,
    0x82C31F5EC35EDC82UL,0x29B052CBB0CBE229UL,0x5A77B4997799C35AUL,0x1E113C3311332D1EUL,
    0x7BCBF646CB463D7BUL,0xA8FC4B1FFC1FB7A8UL,0x6DD6DA61D6610C6DUL,0x2C3A584E3A4E622CUL,
};

__constant ulong GROESTL_T1[256] = {
    0xA597F4A5F432C6C6UL,0x84EB9784976FF8F8UL,0x99C7B099B05EEEEEUL,0x8DF78C8D8C7AF6F6UL,
    0x0DE5170D17E8FFFFUL,0xBDB7DCBDDC0AD6D6UL,0xB1A7C8B1C816DEDEUL,0x5439FC54FC6D9191UL,
    0x50C0F050F0906060UL,0x0304050305070202UL,0xA987E0A9E02ECECEUL,0x7DAC877D87D15656UL,
    0x19D52B192BCCE7E7UL,0x6271A662A613B5B5UL,0xE69A31E6317C4D4DUL,0x9AC3B59AB559ECECUL,
    0x4505CF45CF408F8FUL,0x9D3EBC9DBCA31F1FUL,0x4009C040C0498989UL,0x87EF92879268FAFAUL,
    0x15C53F153FD0EFEFUL,0xEB7F26EB2694B2B2UL,0xC90740C940CE8E8EUL,0x0BED1D0B1DE6FBFBUL,
    0xEC822FEC2F6E4141UL,0x677DA967A91AB3B3UL,0xFDBE1CFD1C435F5FUL,0xEA8A25EA25604545UL,
    0xBF46DABFDAF92323UL,0xF7A602F702515353UL,0x96D3A196A145E4E4UL,0x5B2DED5BED769B9BUL,
    0xC2EA5DC25D287575UL,0x1CD9241C24C5E1E1UL,0xAE7AE9AEE9D43D3DUL,0x6A98BE6ABEF24C4CUL,
    0x5AD8EE5AEE826C6CUL,0x41FCC341C3BD7E7EUL,0x02F1060206F3F5F5UL,0x4F1DD14FD1528383UL,
    0x5CD0E45CE48C6868UL,0xF4A207F407565151UL,0x34B95C345C8DD1D1UL,0x08E9180818E1F9F9UL,
    0x93DFAE93AE4CE2E2UL,0x734D9573953EABABUL,0x53C4F553F5976262UL,0x3F54413F416B2A2AUL,
    0x0C10140C141C0808UL,0x5231F652F6639595UL,0x658CAF65AFE94646UL,0x5E21E25EE27F9D9DUL,
    0x2860782878483030UL,0xA16EF8A1F8CF3737UL,0x0F14110F111B0A0AUL,0xB55EC4B5C4EB2F2FUL,
    0x091C1B091B150E0EUL,0x36485A365A7E2424UL,0x9B36B69BB6AD1B1BUL,0x3DA5473D4798DFDFUL,
    0x26816A266AA7CDCDUL,0x699CBB69BBF54E4EUL,0xCDFE4CCD4C337F7FUL,0x9FCFBA9FBA50EAEAUL,
    0x1B242D1B2D3F1212UL,0x9E3AB99EB9A41D1DUL,0x74B09C749CC45858UL,0x2E68722E72463434UL,
    0x2D6C772D77413636UL,0xB2A3CDB2CD11DCDCUL,0xEE7329EE299DB4B4UL,0xFBB616FB164D5B5BUL,
    0xF65301F601A5A4A4UL,0x4DECD74DD7A17676UL,0x6175A361A314B7B7UL,0xCEFA49CE49347D7DUL,
    0x7BA48D7B8DDF5252UL,0x3EA1423E429FDDDDUL,0x71BC937193CD5E5EUL,0x9726A297A2B11313UL,
    0xF55704F504A2A6A6UL,0x6869B868B801B9B9UL,0x0000000000000000UL,0x2C99742C74B5C1C1UL,
    0x6080A060A0E04040UL,0x1FDD211F21C2E3E3UL,0xC8F243C8433A7979UL,0xED772CED2C9AB6B6UL,
    0xBEB3D9BED90DD4D4UL,0x4601CA46CA478D8DUL,0xD9CE70D970176767UL,0x4BE4DD4BDDAF7272UL,
    0xDE3379DE79ED9494UL,0xD42B67D467FF9898UL,0xE87B23E82393B0B0UL,0x4A11DE4ADE5B8585UL,
    0x6B6DBD6BBD06BBBBUL,0x2A917E2A7EBBC5C5UL,0xE59E34E5347B4F4FUL,0x16C13A163AD7EDEDUL,
    0xC51754C554D28686UL,0xD72F62D762F89A9AUL,0x55CCFF55FF996666UL,0x9422A794A7B61111UL,
    0xCF0F4ACF4AC08A8AUL,0x10C9301030D9E9E9UL,0x06080A060A0E0404UL,0x81E798819866FEFEUL,
    0xF05B0BF00BABA0A0UL,0x44F0CC44CCB47878UL,0xBA4AD5BAD5F02525UL,0xE3963EE33E754B4BUL,
    0xF35F0EF30EACA2A2UL,0xFEBA19FE19445D5DUL,0xC01B5BC05BDB8080UL,0x8A0A858A85800505UL,
    0xAD7EECADECD33F3FUL,0xBC42DFBCDFFE2121UL,0x48E0D848D8A87070UL,0x04F90C040CFDF1F1UL,
    0xDFC67ADF7A196363UL,0xC1EE58C1582F7777UL,0x75459F759F30AFAFUL,0x6384A563A5E74242UL,
    0x3040503050702020UL,0x1AD12E1A2ECBE5E5UL,0x0EE1120E12EFFDFDUL,0x6D65B76DB708BFBFUL,
    0x4C19D44CD4558181UL,0x14303C143C241818UL,0x354C5F355F792626UL,0x2F9D712F71B2C3C3UL,
    0xE16738E13886BEBEUL,0xA26AFDA2FDC83535UL,0xCC0B4FCC4FC78888UL,0x395C4B394B652E2EUL,
    0x573DF957F96A9393UL,0xF2AA0DF20D585555UL,0x82E39D829D61FCFCUL,0x47F4C947C9B37A7AUL,
    0xAC8BEFACEF27C8C8UL,0xE76F32E73288BABAUL,0x2B647D2B7D4F3232UL,0x95D7A495A442E6E6UL,
    0xA09BFBA0FB3BC0C0UL,0x9832B398B3AA1919UL,0xD12768D168F69E9EUL,0x7F5D817F8122A3A3UL,
    0x6688AA66AAEE4444UL,0x7EA8827E82D65454UL,0xAB76E6ABE6DD3B3BUL,0x83169E839E950B0BUL,
    0xCA0345CA45C98C8CUL,0x29957B297BBCC7C7UL,0xD3D66ED36E056B6BUL,0x3C50443C446C2828UL,
    0x79558B798B2CA7A7UL,0xE2633DE23D81BCBCUL,0x1D2C271D27311616UL,0x76419A769A37ADADUL,
    0x3BAD4D3B4D96DBDBUL,0x56C8FA56FA9E6464UL,0x4EE8D24ED2A67474UL,0x1E28221E22361414UL,
    0xDB3F76DB76E49292UL,0x0A181E0A1E120C0CUL,0x6C90B46CB4FC4848UL,0xE46B37E4378FB8B8UL,
    0x5D25E75DE7789F9FUL,0x6E61B26EB20FBDBDUL,0xEF862AEF2A694343UL,0xA693F1A6F135C4C4UL,
    0xA872E3A8E3DA3939UL,0xA462F7A4F7C63131UL,0x37BD5937598AD3D3UL,0x8BFF868B8674F2F2UL,
    0x32B156325683D5D5UL,0x430DC543C54E8B8BUL,0x59DCEB59EB856E6EUL,0xB7AFC2B7C218DADAUL,
    0x8C028F8C8F8E0101UL,0x6479AC64AC1DB1B1UL,0xD2236DD26DF19C9CUL,0xE0923BE03B724949UL,
    0xB4ABC7B4C71FD8D8UL,0xFA4315FA15B9ACACUL,0x07FD090709FAF3F3UL,0x25856F256FA0CFCFUL,
    0xAF8FEAAFEA20CACAUL,0x8EF3898E897DF4F4UL,0xE98E20E920674747UL,0x1820281828381010UL,
    0xD5DE64D5640B6F6FUL,0x88FB83888373F0F0UL,0x6F94B16FB1FB4A4AUL,0x72B8967296CA5C5CUL,
    0x24706C246C543838UL,0xF1AE08F1085F5757UL,0xC7E652C752217373UL,0x5135F351F3649797UL,
    0x238D652365AECBCBUL,0x7C59847C8425A1A1UL,0x9CCBBF9CBF57E8E8UL,0x217C6321635D3E3EUL,
    0xDD377CDD7CEA9696UL,0xDCC27FDC7F1E6161UL,0x861A9186919C0D0DUL,0x851E9485949B0F0FUL,
    0x90DBAB90AB4BE0E0UL,0x42F8C642C6BA7C7CUL,0xC4E257C457267171UL,0xAA83E5AAE529CCCCUL,
    0xD83B73D873E39090UL,0x050C0F050F090606UL,0x01F5030103F4F7F7UL,0x12383612362A1C1CUL,
    0xA39FFEA3FE3CC2C2UL,0x5FD4E15FE18B6A6AUL,0xF94710F910BEAEAEUL,0xD0D26BD06B026969UL,
    0x912EA891A8BF1717UL,0x5829E858E8719999UL,0x2774692769533A3AUL,0xB94ED0B9D0F72727UL,
    0x38A948384891D9D9UL,0x13CD351335DEEBEBUL,0xB356CEB3CEE52B2BUL,0x3344553355772222UL,
    0xBBBFD6BBD604D2D2UL,0x704990709039A9A9UL,0x890E808980870707UL,0xA766F2A7F2C13333UL,
    0xB65AC1B6C1EC2D2DUL,0x22786622665A3C3CUL,0x922AAD92ADB81515UL,0x2089602060A9C9C9UL,
    0x4915DB49DB5C8787UL,0xFF4F1AFF1AB0AAAAUL,0x78A0887888D85050UL,0x7A518E7A8E2BA5A5UL,
    0x8F068A8F8A890303UL,0xF8B213F8134A5959UL,0x80129B809B920909UL,0x1734391739231A1AUL,
    0xDACA75DA75106565UL,0x31B553315384D7D7UL,0xC61351C651D58484UL,0xB8BBD3B8D303D0D0UL,
    0xC31F5EC35EDC8282UL,0xB052CBB0CBE22929UL,0x77B4997799C35A5AUL,0x113C3311332D1E1EUL,
    0xCBF646CB463D7B7BUL,0xFC4B1FFC1FB7A8A8UL,0xD6DA61D6610C6D6DUL,0x3A584E3A4E622C2CUL,
};

__constant ulong GROESTL_T2[256] = {
    0x97F4A5F432C6C6A5UL,0xEB9784976FF8F884UL,0xC7B099B05EEEEE99UL,0xF78C8D8C7AF6F68DUL,
    0xE5170D17E8FFFF0DUL,0xB7DCBDDC0AD6D6BDUL,0xA7C8B1C816DEDEB1UL,0x39FC54FC6D919154UL,
    0xC0F050F090606050UL,0x0405030507020203UL,0x87E0A9E02ECECEA9UL,0xAC877D87D156567DUL,
    0xD52B192BCCE7E719UL,0x71A662A613B5B562UL,0x9A31E6317C4D4DE6UL,0xC3B59AB559ECEC9AUL,
    0x05CF45CF408F8F45UL,0x3EBC9DBCA31F1F9DUL,0x09C040C049898940UL,0xEF92879268FAFA87UL,
    0xC53F153FD0EFEF15UL,0x7F26EB2694B2B2EBUL,0x0740C940CE8E8EC9UL,0xED1D0B1DE6FBFB0BUL,
    0x822FEC2F6E4141ECUL,0x7DA967A91AB3B367UL,0xBE1CFD1C435F5FFDUL,0x8A25EA25604545EAUL,
    0x46DABFDAF92323BFUL,0xA602F702515353F7UL,0xD3A196A145E4E496UL,0x2DED5BED769B9B5BUL,
    0xEA5DC25D287575C2UL,0xD9241C24C5E1E11CUL,0x7AE9AEE9D43D3DAEUL,0x98BE6ABEF24C4C6AUL,
    0xD8EE5AEE826C6C5AUL,0xFCC341C3BD7E7E41UL,0xF1060206F3F5F502UL,0x1DD14FD15283834FUL,
    0xD0E45CE48C68685CUL,0xA207F407565151F4UL,0xB95C345C8DD1D134UL,0xE9180818E1F9F908UL,
    0xDFAE93AE4CE2E293UL,0x4D9573953EABAB73UL,0xC4F553F597626253UL,0x54413F416B2A2A3FUL,
    0x10140C141C08080CUL,0x31F652F663959552UL,0x8CAF65AFE9464665UL,0x21E25EE27F9D9D5EUL,
    0x6078287848303028UL,0x6EF8A1F8CF3737A1UL,0x14110F111B0A0A0FUL,0x5EC4B5C4EB2F2FB5UL,
    0x1C1B091B150E0E09UL,0x485A365A7E242436UL,0x36B69BB6AD1B1B9BUL,0xA5473D4798DFDF3DUL,
    0x816A266AA7CDCD26UL,0x9CBB69BBF54E4E69UL,0xFE4CCD4C337F7FCDUL,0xCFBA9FBA50EAEA9FUL,
    0x242D1B2D3F12121BUL,0x3AB99EB9A41D1D9EUL,0xB09C749CC4585874UL,0x68722E724634342EUL,
    0x6C772D774136362DUL,0xA3CDB2CD11DCDCB2UL,0x7329EE299DB4B4EEUL,0xB616FB164D5B5BFBUL,
    0x5301F601A5A4A4F6UL,0xECD74DD7A176764DUL,0x75A361A314B7B761UL,0xFA49CE49347D7DCEUL,
    0xA48D7B8DDF52527BUL,0xA1423E429FDDDD3EUL,0xBC937193CD5E5E71UL,0x26A297A2B1131397UL,
    0x5704F504A2A6A6F5UL,0x69B868B801B9B968UL,0x0000000000000000UL,0x99742C74B5C1C12CUL,
    0x80A060A0E0404060UL,0xDD211F21C2E3E31FUL,0xF243C8433A7979C8UL,0x772CED2C9AB6B6EDUL,
    0xB3D9BED90DD4D4BEUL,0x01CA46CA478D8D46UL,0xCE70D970176767D9UL,0xE4DD4BDDAF72724BUL,
    0x3379DE79ED9494DEUL,0x2B67D467FF9898D4UL,0x7B23E82393B0B0E8UL,0x11DE4ADE5B85854AUL,
    0x6DBD6BBD06BBBB6BUL,0x917E2A7EBBC5C52AUL,0x9E34E5347B4F4FE5UL,0xC13A163AD7EDED16UL,
    0x1754C554D28686C5UL,0x2F62D762F89A9AD7UL,0xCCFF55FF99666655UL,0x22A794A7B6111194UL,
    0x0F4ACF4AC08A8ACFUL,0xC9301030D9E9E910UL,0x080A060A0E040406UL,0xE798819866FEFE81UL,
    0x5B0BF00BABA0A0F0UL,0xF0CC44CCB4787844UL,0x4AD5BAD5F02525BAUL,0x963EE33E754B4BE3UL,
    0x5F0EF30EACA2A2F3UL,0xBA19FE19445D5DFEUL,0x1B5BC05BDB8080C0UL,0x0A858A858005058AUL,
    0x7EECADECD33F3FADUL,0x42DFBCDFFE2121BCUL,0xE0D848D8A8707048UL,0xF90C040CFDF1F104UL,
    0xC67ADF7A196363DFUL,0xEE58C1582F7777C1UL,0x459F759F30AFAF75UL,0x84A563A5E7424263UL,
    0x4050305070202030UL,0xD12E1A2ECBE5E51AUL,0xE1120E12EFFDFD0EUL,0x65B76DB708BFBF6DUL,
    0x19D44CD45581814CUL,0x303C143C24181814UL,0x4C5F355F79262635UL,0x9D712F71B2C3C32FUL,
    0x6738E13886BEBEE1UL,0x6AFDA2FDC83535A2UL,0x0B4FCC4FC78888CCUL,0x5C4B394B652E2E39UL,
    0x3DF957F96A939357UL,0xAA0DF20D585555F2UL,0xE39D829D61FCFC82UL,0xF4C947C9B37A7A47UL,
    0x8BEFACEF27C8C8ACUL,0x6F32E73288BABAE7UL,0x647D2B7D4F32322BUL,0xD7A495A442E6E695UL,
    0x9BFBA0FB3BC0C0A0UL,0x32B398B3AA191998UL,0x2768D168F69E9ED1UL,0x5D817F8122A3A37FUL,
    0x88AA66AAEE444466UL,0xA8827E82D654547EUL,0x76E6ABE6DD3B3BABUL,0x169E839E950B0B83UL,
    0x0345CA45C98C8CCAUL,0x957B297BBCC7C729UL,0xD66ED36E056B6BD3UL,0x50443C446C28283CUL,
    0x558B798B2CA7A779UL,0x633DE23D81BCBCE2UL,0x2C271D273116161DUL,0x419A769A37ADAD76UL,
    0xAD4D3B4D96DBDB3BUL,0xC8FA56FA9E646456UL,0xE8D24ED2A674744EUL,0x28221E223614141EUL,
    0x3F76DB76E49292DBUL,0x181E0A1E120C0C0AUL,0x90B46CB4FC48486CUL,0x6B37E4378FB8B8E4UL,
    0x25E75DE7789F9F5DUL,0x61B26EB20FBDBD6EUL,0x862AEF2A694343EFUL,0x93F1A6F135C4C4A6UL,
    0x72E3A8E3DA3939A8UL,0x62F7A4F7C63131A4UL,0xBD5937598AD3D337UL,0xFF868B8674F2F28BUL,
    0xB156325683D5D532UL,0x0DC543C54E8B8B43UL,0xDCEB59EB856E6E59UL,0xAFC2B7C218DADAB7UL,
    0x028F8C8F8E01018CUL,0x79AC64AC1DB1B164UL,0x236DD26DF19C9CD2UL,0x923BE03B724949E0UL,
    0xABC7B4C71FD8D8B4UL,0x4315FA15B9ACACFAUL,0xFD090709FAF3F307UL,0x856F256FA0CFCF25UL,
    0x8FEAAFEA20CACAAFUL,0xF3898E897DF4F48EUL,0x8E20E920674747E9UL,0x2028182838101018UL,
    0xDE64D5640B6F6FD5UL,0xFB83888373F0F088UL,0x94B16FB1FB4A4A6FUL,0xB8967296CA5C5C72UL,
    0x706C246C54383824UL,0xAE08F1085F5757F1UL,0xE652C752217373C7UL,0x35F351F364979751UL,
    0x8D652365AECBCB23UL,0x59847C8425A1A17CUL,0xCBBF9CBF57E8E89CUL,0x7C6321635D3E3E21UL,
    0x377CDD7CEA9696DDUL,0xC27FDC7F1E6161DCUL,0x1A9186919C0D0D86UL,0x1E9485949B0F0F85UL,
    0xDBAB90AB4BE0E090UL,0xF8C642C6BA7C7C42UL,0xE257C457267171C4UL,0x83E5AAE529CCCCAAUL,
    0x3B73D873E39090D8UL,0x0C0F050F09060605UL,0xF5030103F4F7F701UL,0x383612362A1C1C12UL,
    0x9FFEA3FE3CC2C2A3UL,0xD4E15FE18B6A6A5FUL,0x4710F910BEAEAEF9UL,0xD26BD06B026969D0UL,
    0x2EA891A8BF171791UL,0x29E858E871999958UL,0x74692769533A3A27UL,0x4ED0B9D0F72727B9UL,
    0xA948384891D9D938UL,0xCD351335DEEBEB13UL,0x56CEB3CEE52B2BB3UL,0x4455335577222233UL,
    0xBFD6BBD604D2D2BBUL,0x4990709039A9A970UL,0x0E80898087070789UL,0x66F2A7F2C13333A7UL,
    0x5AC1B6C1EC2D2DB6UL,0x786622665A3C3C22UL,0x2AAD92ADB8151592UL,0x89602060A9C9C920UL,
    0x15DB49DB5C878749UL,0x4F1AFF1AB0AAAAFFUL,0xA0887888D8505078UL,0x518E7A8E2BA5A57AUL,
    0x068A8F8A8903038FUL,0xB213F8134A5959F8UL,0x129B809B92090980UL,0x34391739231A1A17UL,
    0xCA75DA75106565DAUL,0xB553315384D7D731UL,0x1351C651D58484C6UL,0xBBD3B8D303D0D0B8UL,
    0x1F5EC35EDC8282C3UL,0x52CBB0CBE22929B0UL,0xB4997799C35A5A77UL,0x3C3311332D1E1E11UL,
    0xF646CB463D7B7BCBUL,0x4B1FFC1FB7A8A8FCUL,0xDA61D6610C6D6DD6UL,0x584E3A4E622C2C3AUL,
};

__constant ulong GROESTL_T3[256] = {
    0xF4A5F432C6C6A597UL,0x9784976FF8F884EBUL,0xB099B05EEEEE99C7UL,0x8C8D8C7AF6F68DF7UL,
    0x170D17E8FFFF0DE5UL,0xDCBDDC0AD6D6BDB7UL,0xC8B1C816DEDEB1A7UL,0xFC54FC6D91915439UL,
    0xF050F090606050C0UL,0x0503050702020304UL,0xE0A9E02ECECEA987UL,0x877D87D156567DACUL,
    0x2B192BCCE7E719D5UL,0xA662A613B5B56271UL,0x31E6317C4D4DE69AUL,0xB59AB559ECEC9AC3UL,
    0xCF45CF408F8F4505UL,0xBC9DBCA31F1F9D3EUL,0xC040C04989894009UL,0x92879268FAFA87EFUL,
    0x3F153FD0EFEF15C5UL,0x26EB2694B2B2EB7FUL,0x40C940CE8E8EC907UL,0x1D0B1DE6FBFB0BEDUL,
    0x2FEC2F6E4141EC82UL,0xA967A91AB3B3677DUL,0x1CFD1C435F5FFDBEUL,0x25EA25604545EA8AUL,
    0xDABFDAF92323BF46UL,0x02F702515353F7A6UL,0xA196A145E4E496D3UL,0xED5BED769B9B5B2DUL,
    0x5DC25D287575C2EAUL,0x241C24C5E1E11CD9UL,0xE9AEE9D43D3DAE7AUL,0xBE6ABEF24C4C6A98UL,
    0xEE5AEE826C6C5AD8UL,0xC341C3BD7E7E41FCUL,0x060206F3F5F502F1UL,0xD14FD15283834F1DUL,
    0xE45CE48C68685CD0UL,0x07F407565151F4A2UL,0x5C345C8DD1D134B9UL,0x180818E1F9F908E9UL,
    0xAE93AE4CE2E293DFUL,0x9573953EABAB734DUL,0xF553F597626253C4UL,0x413F416B2A2A3F54UL,
    0x140C141C08080C10UL,0xF652F66395955231UL,0xAF65AFE94646658CUL,0xE25EE27F9D9D5E21UL,
    0x7828784830302860UL,0xF8A1F8CF3737A16EUL,0x110F111B0A0A0F14UL,0xC4B5C4EB2F2FB55EUL,
    0x1B091B150E0E091CUL,0x5A365A7E24243648UL,0xB69BB6AD1B1B9B36UL,0x473D4798DFDF3DA5UL,
    0x6A266AA7CDCD2681UL,0xBB69BBF54E4E699CUL,0x4CCD4C337F7FCDFEUL,0xBA9FBA50EAEA9FCFUL,
    0x2D1B2D3F12121B24UL,0xB99EB9A41D1D9E3AUL,0x9C749CC4585874B0UL,0x722E724634342E68UL,
    0x772D774136362D6CUL,0xCDB2CD11DCDCB2A3UL,0x29EE299DB4B4EE73UL,0x16FB164D5B5BFBB6UL,
    0x01F601A5A4A4F653UL,0xD74DD7A176764DECUL,0xA361A314B7B76175UL,0x49CE49347D7DCEFAUL,
    0x8D7B8DDF52527BA4UL,0x423E429FDDDD3EA1UL,0x937193CD5E5E71BCUL,0xA297A2B113139726UL,
    0x04F504A2A6A6F557UL,0xB868B801B9B96869UL,0x0000000000000000UL,0x742C74B5C1C12C99UL,
    0xA060A0E040406080UL,0x211F21C2E3E31FDDUL,0x43C8433A7979C8F2UL,0x2CED2C9AB6B6ED77UL,
    0xD9BED90DD4D4BEB3UL,0xCA46CA478D8D4601UL,0x70D970176767D9CEUL,0xDD4BDDAF72724BE4UL,
    0x79DE79ED9494DE33UL,0x67D467FF9898D42BUL,0x23E82393B0B0E87BUL,0xDE4ADE5B85854A11UL,
    0xBD6BBD06BBBB6B6DUL,0x7E2A7EBBC5C52A91UL,0x34E5347B4F4FE59EUL,0x3A163AD7EDED16C1UL,
    0x54C554D28686C517UL,0x62D762F89A9AD72FUL,0xFF55FF99666655CCUL,0xA794A7B611119422UL,
    0x4ACF4AC08A8ACF0FUL,0x301030D9E9E910C9UL,0x0A060A0E04040608UL,0x98819866FEFE81E7UL,
    0x0BF00BABA0A0F05BUL,0xCC44CCB4787844F0UL,0xD5BAD5F02525BA4AUL,0x3EE33E754B4BE396UL,
    0x0EF30EACA2A2F35FUL,0x19FE19445D5DFEBAUL,0x5BC05BDB8080C01BUL,0x858A858005058A0AUL,
    0xECADECD33F3FAD7EUL,0xDFBCDFFE2121BC42UL,0xD848D8A8707048E0UL,0x0C040CFDF1F104F9UL,
    0x7ADF7A196363DFC6UL,0x58C1582F7777C1EEUL,0x9F759F30AFAF7545UL,0xA563A5E742426384UL,
    0x5030507020203040UL,0x2E1A2ECBE5E51AD1UL,0x120E12EFFDFD0EE1UL,0xB76DB708BFBF6D65UL,
    0xD44CD45581814C19UL,0x3C143C2418181430UL,0x5F355F792626354CUL,0x712F71B2C3C32F9DUL,
    0x38E13886BEBEE167UL,0xFDA2FDC83535A26AUL,0x4FCC4FC78888CC0BUL,0x4B394B652E2E395CUL,
    0xF957F96A9393573DUL,0x0DF20D585555F2AAUL,0x9D829D61FCFC82E3UL,0xC947C9B37A7A47F4UL,
    0xEFACEF27C8C8AC8BUL,0x32E73288BABAE76FUL,0x7D2B7D4F32322B64UL,0xA495A442E6E695D7UL,
    0xFBA0FB3BC0C0A09BUL,0xB398B3AA19199832UL,0x68D168F69E9ED127UL,0x817F8122A3A37F5DUL,
    0xAA66AAEE44446688UL,0x827E82D654547EA8UL,0xE6ABE6DD3B3BAB76UL,0x9E839E950B0B8316UL,
    0x45CA45C98C8CCA03UL,0x7B297BBCC7C72995UL,0x6ED36E056B6BD3D6UL,0x443C446C28283C50UL,
    0x8B798B2CA7A77955UL,0x3DE23D81BCBCE263UL,0x271D273116161D2CUL,0x9A769A37ADAD7641UL,
    0x4D3B4D96DBDB3BADUL,0xFA56FA9E646456C8UL,0xD24ED2A674744EE8UL,0x221E223614141E28UL,
    0x76DB76E49292DB3FUL,0x1E0A1E120C0C0A18UL,0xB46CB4FC48486C90UL,0x37E4378FB8B8E46BUL,
    0xE75DE7789F9F5D25UL,0xB26EB20FBDBD6E61UL,0x2AEF2A694343EF86UL,0xF1A6F135C4C4A693UL,
    0xE3A8E3DA3939A872UL,0xF7A4F7C63131A462UL,0x5937598AD3D337BDUL,0x868B8674F2F28BFFUL,
    0x56325683D5D532B1UL,0xC543C54E8B8B430DUL,0xEB59EB856E6E59DCUL,0xC2B7C218DADAB7AFUL,
    0x8F8C8F8E01018C02UL,0xAC64AC1DB1B16479UL,0x6DD26DF19C9CD223UL,0x3BE03B724949E092UL,
    0xC7B4C71FD8D8B4ABUL,0x15FA15B9ACACFA43UL,0x090709FAF3F307FDUL,0x6F256FA0CFCF2585UL,
    0xEAAFEA20CACAAF8FUL,0x898E897DF4F48EF3UL,0x20E920674747E98EUL,0x2818283810101820UL,
    0x64D5640B6F6FD5DEUL,0x83888373F0F088FBUL,0xB16FB1FB4A4A6F94UL,0x967296CA5C5C72B8UL,
    0x6C246C5438382470UL,0x08F1085F5757F1AEUL,0x52C752217373C7E6UL,0xF351F36497975135UL,
    0x652365AECBCB238DUL,0x847C8425A1A17C59UL,0xBF9CBF57E8E89CCBUL,0x6321635D3E3E217CUL,
    0x7CDD7CEA9696DD37UL,0x7FDC7F1E6161DCC2UL,0x9186919C0D0D861AUL,0x9485949B0F0F851EUL,
    0xAB90AB4BE0E090DBUL,0xC642C6BA7C7C42F8UL,0x57C457267171C4E2UL,0xE5AAE529CCCCAA83UL,
    0x73D873E39090D83BUL,0x0F050F090606050CUL,0x030103F4F7F701F5UL,0x3612362A1C1C1238UL,
    0xFEA3FE3CC2C2A39FUL,0xE15FE18B6A6A5FD4UL,0x10F910BEAEAEF947UL,0x6BD06B026969D0D2UL,
    0xA891A8BF1717912EUL,0xE858E87199995829UL,0x692769533A3A2774UL,0xD0B9D0F72727B94EUL,
    0x48384891D9D938A9UL,0x351335DEEBEB13CDUL,0xCEB3CEE52B2BB356UL,0x5533557722223344UL,
    0xD6BBD604D2D2BBBFUL,0x90709039A9A97049UL,0x808980870707890EUL,0xF2A7F2C13333A766UL,
    0xC1B6C1EC2D2DB65AUL,0x6622665A3C3C2278UL,0xAD92ADB81515922AUL,0x602060A9C9C92089UL,
    0xDB49DB5C87874915UL,0x1AFF1AB0AAAAFF4FUL,0x887888D8505078A0UL,0x8E7A8E2BA5A57A51UL,
    0x8A8F8A8903038F06UL,0x13F8134A5959F8B2UL,0x9B809B9209098012UL,0x391739231A1A1734UL,
    0x75DA75106565DACAUL,0x53315384D7D731B5UL,0x51C651D58484C613UL,0xD3B8D303D0D0B8BBUL,
    0x5EC35EDC8282C31FUL,0xCBB0CBE22929B052UL,0x997799C35A5A77B4UL,0x3311332D1E1E113CUL,
    0x46CB463D7B7BCBF6UL,0x1FFC1FB7A8A8FC4BUL,0x61D6610C6D6DD6DAUL,0x4E3A4E622C2C3A58UL,
};

__constant ulong GROESTL_T4[256] = {
    0xA5F432C6C6A597F4UL,0x84976FF8F884EB97UL,0x99B05EEEEE99C7B0UL,0x8D8C7AF6F68DF78CUL,
    0x0D17E8FFFF0DE517UL,0xBDDC0AD6D6BDB7DCUL,0xB1C816DEDEB1A7C8UL,0x54FC6D91915439FCUL,
    0x50F090606050C0F0UL,0x0305070202030405UL,0xA9E02ECECEA987E0UL,0x7D87D156567DAC87UL,
    0x192BCCE7E719D52BUL,0x62A613B5B56271A6UL,0xE6317C4D4DE69A31UL,0x9AB559ECEC9AC3B5UL,
    0x45CF408F8F4505CFUL,0x9DBCA31F1F9D3EBCUL,0x40C04989894009C0UL,0x879268FAFA87EF92UL,
    0x153FD0EFEF15C53FUL,0xEB2694B2B2EB7F26UL,0xC940CE8E8EC90740UL,0x0B1DE6FBFB0BED1DUL,
    0xEC2F6E4141EC822FUL,0x67A91AB3B3677DA9UL,0xFD1C435F5FFDBE1CUL,0xEA25604545EA8A25UL,
    0xBFDAF92323BF46DAUL,0xF702515353F7A602UL,0x96A145E4E496D3A1UL,0x5BED769B9B5B2DEDUL,
    0xC25D287575C2EA5DUL,0x1C24C5E1E11CD924UL,0xAEE9D43D3DAE7AE9UL,0x6ABEF24C4C6A98BEUL,
    0x5AEE826C6C5AD8EEUL,0x41C3BD7E7E41FCC3UL,0x0206F3F5F502F106UL,0x4FD15283834F1DD1UL,
    0x5CE48C68685CD0E4UL,0xF407565151F4A207UL,0x345C8DD1D134B95CUL,0x0818E1F9F908E918UL,
    0x93AE4CE2E293DFAEUL,0x73953EABAB734D95UL,0x53F597626253C4F5UL,0x3F416B2A2A3F5441UL,
    0x0C141C08080C1014UL,0x52F66395955231F6UL,0x65AFE94646658CAFUL,0x5EE27F9D9D5E21E2UL,
    0x2878483030286078UL,0xA1F8CF3737A16EF8UL,0x0F111B0A0A0F1411UL,0xB5C4EB2F2FB55EC4UL,
    0x091B150E0E091C1BUL,0x365A7E242436485AUL,0x9BB6AD1B1B9B36B6UL,0x3D4798DFDF3DA547UL,
    0x266AA7CDCD26816AUL,0x69BBF54E4E699CBBUL,0xCD4C337F7FCDFE4CUL,0x9FBA50EAEA9FCFBAUL,
    0x1B2D3F12121B242DUL,0x9EB9A41D1D9E3AB9UL,0x749CC4585874B09CUL,0x2E724634342E6872UL,
    0x2D774136362D6C77UL,0xB2CD11DCDCB2A3CDUL,0xEE299DB4B4EE7329UL,0xFB164D5B5BFBB616UL,
    0xF601A5A4A4F65301UL,0x4DD7A176764DECD7UL,0x61A314B7B76175A3UL,0xCE49347D7DCEFA49UL,
    0x7B8DDF52527BA48DUL,0x3E429FDDDD3EA142UL,0x7193CD5E5E71BC93UL,0x97A2B113139726A2UL,
    0xF504A2A6A6F55704UL,0x68B801B9B96869B8UL,0x0000000000000000UL,0x2C74B5C1C12C9974UL,
    0x60A0E040406080A0UL,0x1F21C2E3E31FDD21UL,0xC8433A7979C8F243UL,0xED2C9AB6B6ED772CUL,
    0xBED90DD4D4BEB3D9UL,0x46CA478D8D4601CAUL,0xD970176767D9CE70UL,0x4BDDAF72724BE4DDUL,
    0xDE79ED9494DE3379UL,0xD467FF9898D42B67UL,0xE82393B0B0E87B23UL,0x4ADE5B85854A11DEUL,
    0x6BBD06BBBB6B6DBDUL,0x2A7EBBC5C52A917EUL,0xE5347B4F4FE59E34UL,0x163AD7EDED16C13AUL,
    0xC554D28686C51754UL,0xD762F89A9AD72F62UL,0x55FF99666655CCFFUL,0x94A7B611119422A7UL,
    0xCF4AC08A8ACF0F4AUL,0x1030D9E9E910C930UL,0x060A0E040406080AUL,0x819866FEFE81E798UL,
    0xF00BABA0A0F05B0BUL,0x44CCB4787844F0CCUL,0xBAD5F02525BA4AD5UL,0xE33E754B4BE3963EUL,
    0xF30EACA2A2F35F0EUL,0xFE19445D5DFEBA19UL,0xC05BDB8080C01B5BUL,0x8A858005058A0A85UL,
    0xADECD33F3FAD7EECUL,0xBCDFFE2121BC42DFUL,0x48D8A8707048E0D8UL,0x040CFDF1F104F90CUL,
    0xDF7A196363DFC67AUL,0xC1582F7777C1EE58UL,0x759F30AFAF75459FUL,0x63A5E742426384A5UL,
    0x3050702020304050UL,0x1A2ECBE5E51AD12EUL,0x0E12EFFDFD0EE112UL,0x6DB708BFBF6D65B7UL,
    0x4CD45581814C19D4UL,0x143C24181814303CUL,0x355F792626354C5FUL,0x2F71B2C3C32F9D71UL,
    0xE13886BEBEE16738UL,0xA2FDC83535A26AFDUL,0xCC4FC78888CC0B4FUL,0x394B652E2E395C4BUL,
    0x57F96A9393573DF9UL,0xF20D585555F2AA0DUL,0x829D61FCFC82E39DUL,0x47C9B37A7A47F4C9UL,
    0xACEF27C8C8AC8BEFUL,0xE73288BABAE76F32UL,0x2B7D4F32322B647DUL,0x95A442E6E695D7A4UL,
    0xA0FB3BC0C0A09BFBUL,0x98B3AA19199832B3UL,0xD168F69E9ED12768UL,0x7F8122A3A37F5D81UL,
    0x66AAEE44446688AAUL,0x7E82D654547EA882UL,0xABE6DD3B3BAB76E6UL,0x839E950B0B83169EUL,
    0xCA45C98C8CCA0345UL,0x297BBCC7C729957BUL,0xD36E056B6BD3D66EUL,0x3C446C28283C5044UL,
    0x798B2CA7A779558BUL,0xE23D81BCBCE2633DUL,0x1D273116161D2C27UL,0x769A37ADAD76419AUL,
    0x3B4D96DBDB3BAD4DUL,0x56FA9E646456C8FAUL,0x4ED2A674744EE8D2UL,0x1E223614141E2822UL,
    0xDB76E49292DB3F76UL,0x0A1E120C0C0A181EUL,0x6CB4FC48486C90B4UL,0xE4378FB8B8E46B37UL,
    0x5DE7789F9F5D25E7UL,0x6EB20FBDBD6E61B2UL,0xEF2A694343EF862AUL,0xA6F135C4C4A693F1UL,
    0xA8E3DA3939A872E3UL,0xA4F7C63131A462F7UL,0x37598AD3D337BD59UL,0x8B8674F2F28BFF86UL,
    0x325683D5D532B156UL,0x43C54E8B8B430DC5UL,0x59EB856E6E59DCEBUL,0xB7C218DADAB7AFC2UL,
    0x8C8F8E01018C028FUL,0x64AC1DB1B16479ACUL,0xD26DF19C9CD2236DUL,0xE03B724949E0923BUL,
    0xB4C71FD8D8B4ABC7UL,0xFA15B9ACACFA4315UL,0x0709FAF3F307FD09UL,0x256FA0CFCF25856FUL,
    0xAFEA20CACAAF8FEAUL,0x8E897DF4F48EF389UL,0xE920674747E98E20UL,0x1828381010182028UL,
    0xD5640B6F6FD5DE64UL,0x888373F0F088FB83UL,0x6FB1FB4A4A6F94B1UL,0x7296CA5C5C72B896UL,
    0x246C54383824706CUL,0xF1085F5757F1AE08UL,0xC752217373C7E652UL,0x51F36497975135F3UL,
    0x2365AECBCB238D65UL,0x7C8425A1A17C5984UL,0x9CBF57E8E89CCBBFUL,0x21635D3E3E217C63UL,
    0xDD7CEA9696DD377CUL,0xDC7F1E6161DCC27FUL,0x86919C0D0D861A91UL,0x85949B0F0F851E94UL,
    0x90AB4BE0E090DBABUL,0x42C6BA7C7C42F8C6UL,0xC457267171C4E257UL,0xAAE529CCCCAA83E5UL,
    0xD873E39090D83B73UL,0x050F090606050C0FUL,0x0103F4F7F701F503UL,0x12362A1C1C123836UL,
    0xA3FE3CC2C2A39FFEUL,0x5FE18B6A6A5FD4E1UL,0xF910BEAEAEF94710UL,0xD06B026969D0D26BUL,
    0x91A8BF1717912EA8UL,0x58E87199995829E8UL,0x2769533A3A277469UL,0xB9D0F72727B94ED0UL,
    0x384891D9D938A948UL,0x1335DEEBEB13CD35UL,0xB3CEE52B2BB356CEUL,0x3355772222334455UL,
    0xBBD604D2D2BBBFD6UL,0x709039A9A9704990UL,0x8980870707890E80UL,0xA7F2C13333A766F2UL,
    0xB6C1EC2D2DB65AC1UL,0x22665A3C3C227866UL,0x92ADB81515922AADUL,0x2060A9C9C9208960UL,
    0x49DB5C87874915DBUL,0xFF1AB0AAAAFF4F1AUL,0x7888D8505078A088UL,0x7A8E2BA5A57A518EUL,
    0x8F8A8903038F068AUL,0xF8134A5959F8B213UL,0x809B92090980129BUL,0x1739231A1A173439UL,
    0xDA75106565DACA75UL,0x315384D7D731B553UL,0xC651D58484C61351UL,0xB8D303D0D0B8BBD3UL,
    0xC35EDC8282C31F5EUL,0xB0CBE22929B052CBUL,0x7799C35A5A77B499UL,0x11332D1E1E113C33UL,
    0xCB463D7B7BCBF646UL,0xFC1FB7A8A8FC4B1FUL,0xD6610C6D6DD6DA61UL,0x3A4E622C2C3A584EUL,
};

__constant ulong GROESTL_T5[256] = {
    0xF432C6C6A597F4A5UL,0x976FF8F884EB9784UL,0xB05EEEEE99C7B099UL,0x8C7AF6F68DF78C8DUL,
    0x17E8FFFF0DE5170DUL,0xDC0AD6D6BDB7DCBDUL,0xC816DEDEB1A7C8B1UL,0xFC6D91915439FC54UL,
    0xF090606050C0F050UL,0x0507020203040503UL,0xE02ECECEA987E0A9UL,0x87D156567DAC877DUL,
    0x2BCCE7E719D52B19UL,0xA613B5B56271A662UL,0x317C4D4DE69A31E6UL,0xB559ECEC9AC3B59AUL,
    0xCF408F8F4505CF45UL,0xBCA31F1F9D3EBC9DUL,0xC04989894009C040UL,0x9268FAFA87EF9287UL,
    0x3FD0EFEF15C53F15UL,0x2694B2B2EB7F26EBUL,0x40CE8E8EC90740C9UL,0x1DE6FBFB0BED1D0BUL,
    0x2F6E4141EC822FECUL,0xA91AB3B3677DA967UL,0x1C435F5FFDBE1CFDUL,0x25604545EA8A25EAUL,
    0xDAF92323BF46DABFUL,0x02515353F7A602F7UL,0xA145E4E496D3A196UL,0xED769B9B5B2DED5BUL,
    0x5D287575C2EA5DC2UL,0x24C5E1E11CD9241CUL,0xE9D43D3DAE7AE9AEUL,0xBEF24C4C6A98BE6AUL,
    0xEE826C6C5AD8EE5AUL,0xC3BD7E7E41FCC341UL,0x06F3F5F502F10602UL,0xD15283834F1DD14FUL,
    0xE48C68685CD0E45CUL,0x07565151F4A207F4UL,0x5C8DD1D134B95C34UL,0x18E1F9F908E91808UL,
    0xAE4CE2E293DFAE93UL,0x953EABAB734D9573UL,0xF597626253C4F553UL,0x416B2A2A3F54413FUL,
    0x141C08080C10140CUL,0xF66395955231F652UL,0xAFE94646658CAF65UL,0xE27F9D9D5E21E25EUL,
    0x7848303028607828UL,0xF8CF3737A16EF8A1UL,0x111B0A0A0F14110FUL,0xC4EB2F2FB55EC4B5UL,
    0x1B150E0E091C1B09UL,0x5A7E242436485A36UL,0xB6AD1B1B9B36B69BUL,0x4798DFDF3DA5473DUL,
    0x6AA7CDCD26816A26UL,0xBBF54E4E699CBB69UL,0x4C337F7FCDFE4CCDUL,0xBA50EAEA9FCFBA9FUL,
    0x2D3F12121B242D1BUL,0xB9A41D1D9E3AB99EUL,0x9CC4585874B09C74UL,0x724634342E68722EUL,
    0x774136362D6C772DUL,0xCD11DCDCB2A3CDB2UL,0x299DB4B4EE7329EEUL,0x164D5B5BFBB616FBUL,
    0x01A5A4A4F65301F6UL,0xD7A176764DECD74DUL,0xA314B7B76175A361UL,0x49347D7DCEFA49CEUL,
    0x8DDF52527BA48D7BUL,0x429FDDDD3EA1423EUL,0x93CD5E5E71BC9371UL,0xA2B113139726A297UL,
    0x04A2A6A6F55704F5UL,0xB801B9B96869B868UL,0x0000000000000000UL,0x74B5C1C12C99742CUL,
    0xA0E040406080A060UL,0x21C2E3E31FDD211FUL,0x433A7979C8F243C8UL,0x2C9AB6B6ED772CEDUL,
    0xD90DD4D4BEB3D9BEUL,0xCA478D8D4601CA46UL,0x70176767D9CE70D9UL,0xDDAF72724BE4DD4BUL,
    0x79ED9494DE3379DEUL,0x67FF9898D42B67D4UL,0x2393B0B0E87B23E8UL,0xDE5B85854A11DE4AUL,
    0xBD06BBBB6B6DBD6BUL,0x7EBBC5C52A917E2AUL,0x347B4F4FE59E34E5UL,0x3AD7EDED16C13A16UL,
    0x54D28686C51754C5UL,0x62F89A9AD72F62D7UL,0xFF99666655CCFF55UL,0xA7B611119422A794UL,
    0x4AC08A8ACF0F4ACFUL,0x30D9E9E910C93010UL,0x0A0E040406080A06UL,0x9866FEFE81E79881UL,
    0x0BABA0A0F05B0BF0UL,0xCCB4787844F0CC44UL,0xD5F02525BA4AD5BAUL,0x3E754B4BE3963EE3UL,
    0x0EACA2A2F35F0EF3UL,0x19445D5DFEBA19FEUL,0x5BDB8080C01B5BC0UL,0x858005058A0A858AUL,
    0xECD33F3FAD7EECADUL,0xDFFE2121BC42DFBCUL,0xD8A8707048E0D848UL,0x0CFDF1F104F90C04UL,
    0x7A196363DFC67ADFUL,0x582F7777C1EE58C1UL,0x9F30AFAF75459F75UL,0xA5E742426384A563UL,
    0x5070202030405030UL,0x2ECBE5E51AD12E1AUL,0x12EFFDFD0EE1120EUL,0xB708BFBF6D65B76DUL,
    0xD45581814C19D44CUL,0x3C24181814303C14UL,0x5F792626354C5F35UL,0x71B2C3C32F9D712FUL,
    0x3886BEBEE16738E1UL,0xFDC83535A26AFDA2UL,0x4FC78888CC0B4FCCUL,0x4B652E2E395C4B39UL,
    0xF96A9393573DF957UL,0x0D585555F2AA0DF2UL,0x9D61FCFC82E39D82UL,0xC9B37A7A47F4C947UL,
    0xEF27C8C8AC8BEFACUL,0x3288BABAE76F32E7UL,0x7D4F32322B647D2BUL,0xA442E6E695D7A495UL,
    0xFB3BC0C0A09BFBA0UL,0xB3AA19199832B398UL,0x68F69E9ED12768D1UL,0x8122A3A37F5D817FUL,
    0xAAEE44446688AA66UL,0x82D654547EA8827EUL,0xE6DD3B3BAB76E6ABUL,0x9E950B0B83169E83UL,
    0x45C98C8CCA0345CAUL,0x7BBCC7C729957B29UL,0x6E056B6BD3D66ED3UL,0x446C28283C50443CUL,
    0x8B2CA7A779558B79UL,0x3D81BCBCE2633DE2UL,0x273116161D2C271DUL,0x9A37ADAD76419A76UL,
    0x4D96DBDB3BAD4D3BUL,0xFA9E646456C8FA56UL,0xD2A674744EE8D24EUL,0x223614141E28221EUL,
    0x76E49292DB3F76DBUL,0x1E120C0C0A181E0AUL,0xB4FC48486C90B46CUL,0x378FB8B8E46B37E4UL,
    0xE7789F9F5D25E75DUL,0xB20FBDBD6E61B26EUL,0x2A694343EF862AEFUL,0xF135C4C4A693F1A6UL,
    0xE3DA3939A872E3A8UL,0xF7C63131A462F7A4UL,0x598AD3D337BD5937UL,0x8674F2F28BFF868BUL,
    0x5683D5D532B15632UL,0xC54E8B8B430DC543UL,0xEB856E6E59DCEB59UL,0xC218DADAB7AFC2B7UL,
    0x8F8E01018C028F8CUL,0xAC1DB1B16479AC64UL,0x6DF19C9CD2236DD2UL,0x3B724949E0923BE0UL,
    0xC71FD8D8B4ABC7B4UL,0x15B9ACACFA4315FAUL,0x09FAF3F307FD0907UL,0x6FA0CFCF25856F25UL,
    0xEA20CACAAF8FEAAFUL,0x897DF4F48EF3898EUL,0x20674747E98E20E9UL,0x2838101018202818UL,
    0x640B6F6FD5DE64D5UL,0x8373F0F088FB8388UL,0xB1FB4A4A6F94B16FUL,0x96CA5C5C72B89672UL,
    0x6C54383824706C24UL,0x085F5757F1AE08F1UL,0x52217373C7E652C7UL,0xF36497975135F351UL,
    0x65AECBCB238D6523UL,0x8425A1A17C59847CUL,0xBF57E8E89CCBBF9CUL,0x635D3E3E217C6321UL,
    0x7CEA9696DD377CDDUL,0x7F1E6161DCC27FDCUL,0x919C0D0D861A9186UL,0x949B0F0F851E9485UL,
    0xAB4BE0E090DBAB90UL,0xC6BA7C7C42F8C642UL,0x57267171C4E257C4UL,0xE529CCCCAA83E5AAUL,
    0x73E39090D83B73D8UL,0x0F090606050C0F05UL,0x03F4F7F701F50301UL,0x362A1C1C12383612UL,
    0xFE3CC2C2A39FFEA3UL,0xE18B6A6A5FD4E15FUL,0x10BEAEAEF94710F9UL,0x6B026969D0D26BD0UL,
    0xA8BF1717912EA891UL,0xE87199995829E858UL,0x69533A3A27746927UL,0xD0F72727B94ED0B9UL,
    0x4891D9D938A94838UL,0x35DEEBEB13CD3513UL,0xCEE52B2BB356CEB3UL,0x5577222233445533UL,
    0xD604D2D2BBBFD6BBUL,0x9039A9A970499070UL,0x80870707890E8089UL,0xF2C13333A766F2A7UL,
    0xC1EC2D2DB65AC1B6UL,0x665A3C3C22786622UL,0xADB81515922AAD92UL,0x60A9C9C920896020UL,
    0xDB5C87874915DB49UL,0x1AB0AAAAFF4F1AFFUL,0x88D8505078A08878UL,0x8E2BA5A57A518E7AUL,
    0x8A8903038F068A8FUL,0x134A5959F8B213F8UL,0x9B92090980129B80UL,0x39231A1A17343917UL,
    0x75106565DACA75DAUL,0x5384D7D731B55331UL,0x51D58484C61351C6UL,0xD303D0D0B8BBD3B8UL,
    0x5EDC8282C31F5EC3UL,0xCBE22929B052CBB0UL,0x99C35A5A77B49977UL,0x332D1E1E113C3311UL,
    0x463D7B7BCBF646CBUL,0x1FB7A8A8FC4B1FFCUL,0x610C6D6DD6DA61D6UL,0x4E622C2C3A584E3AUL,
};

__constant ulong GROESTL_T6[256] = {
    0x32C6C6A597F4A5F4UL,0x6FF8F884EB978497UL,0x5EEEEE99C7B099B0UL,0x7AF6F68DF78C8D8CUL,
    0xE8FFFF0DE5170D17UL,0x0AD6D6BDB7DCBDDCUL,0x16DEDEB1A7C8B1C8UL,0x6D91915439FC54FCUL,
    0x90606050C0F050F0UL,0x0702020304050305UL,0x2ECECEA987E0A9E0UL,0xD156567DAC877D87UL,
    0xCCE7E719D52B192BUL,0x13B5B56271A662A6UL,0x7C4D4DE69A31E631UL,0x59ECEC9AC3B59AB5UL,
    0x408F8F4505CF45CFUL,0xA31F1F9D3EBC9DBCUL,0x4989894009C040C0UL,0x68FAFA87EF928792UL,
    0xD0EFEF15C53F153FUL,0x94B2B2EB7F26EB26UL,0xCE8E8EC90740C940UL,0xE6FBFB0BED1D0B1DUL,
    0x6E4141EC822FEC2FUL,0x1AB3B3677DA967A9UL,0x435F5FFDBE1CFD1CUL,0x604545EA8A25EA25UL,
    0xF92323BF46DABFDAUL,0x515353F7A602F702UL,0x45E4E496D3A196A1UL,0x769B9B5B2DED5BEDUL,
    0x287575C2EA5DC25DUL,0xC5E1E11CD9241C24UL,0xD43D3DAE7AE9AEE9UL,0xF24C4C6A98BE6ABEUL,
    0x826C6C5AD8EE5AEEUL,0xBD7E7E41FCC341C3UL,0xF3F5F502F1060206UL,0x5283834F1DD14FD1UL,
    0x8C68685CD0E45CE4UL,0x565151F4A207F407UL,0x8DD1D134B95C345CUL,0xE1F9F908E9180818UL,
    0x4CE2E293DFAE93AEUL,0x3EABAB734D957395UL,0x97626253C4F553F5UL,0x6B2A2A3F54413F41UL,
    0x1C08080C10140C14UL,0x6395955231F652F6UL,0xE94646658CAF65AFUL,0x7F9D9D5E21E25EE2UL,
    0x4830302860782878UL,0xCF3737A16EF8A1F8UL,0x1B0A0A0F14110F11UL,0xEB2F2FB55EC4B5C4UL,
    0x150E0E091C1B091BUL,0x7E242436485A365AUL,0xAD1B1B9B36B69BB6UL,0x98DFDF3DA5473D47UL,
    0xA7CDCD26816A266AUL,0xF54E4E699CBB69BBUL,0x337F7FCDFE4CCD4CUL,0x50EAEA9FCFBA9FBAUL,
    0x3F12121B242D1B2DUL,0xA41D1D9E3AB99EB9UL,0xC4585874B09C749CUL,0x4634342E68722E72UL,
    0x4136362D6C772D77UL,0x11DCDCB2A3CDB2CDUL,0x9DB4B4EE7329EE29UL,0x4D5B5BFBB616FB16UL,
    0xA5A4A4F65301F601UL,0xA176764DECD74DD7UL,0x14B7B76175A361A3UL,0x347D7DCEFA49CE49UL,
    0xDF52527BA48D7B8DUL,0x9FDDDD3EA1423E42UL,0xCD5E5E71BC937193UL,0xB113139726A297A2UL,
    0xA2A6A6F55704F504UL,0x01B9B96869B868B8UL,0x0000000000000000UL,0xB5C1C12C99742C74UL,
    0xE040406080A060A0UL,0xC2E3E31FDD211F21UL,0x3A7979C8F243C843UL,0x9AB6B6ED772CED2CUL,
    0x0DD4D4BEB3D9BED9UL,0x478D8D4601CA46CAUL,0x176767D9CE70D970UL,0xAF72724BE4DD4BDDUL,
    0xED9494DE3379DE79UL,0xFF9898D42B67D467UL,0x93B0B0E87B23E823UL,0x5B85854A11DE4ADEUL,
    0x06BBBB6B6DBD6BBDUL,0xBBC5C52A917E2A7EUL,0x7B4F4FE59E34E534UL,0xD7EDED16C13A163AUL,
    0xD28686C51754C554UL,0xF89A9AD72F62D762UL,0x99666655CCFF55FFUL,0xB611119422A794A7UL,
    0xC08A8ACF0F4ACF4AUL,0xD9E9E910C9301030UL,0x0E040406080A060AUL,0x66FEFE81E7988198UL,
    0xABA0A0F05B0BF00BUL,0xB4787844F0CC44CCUL,0xF02525BA4AD5BAD5UL,0x754B4BE3963EE33EUL,
    0xACA2A2F35F0EF30EUL,0x445D5DFEBA19FE19UL,0xDB8080C01B5BC05BUL,0x8005058A0A858A85UL,
    0xD33F3FAD7EECADECUL,0xFE2121BC42DFBCDFUL,0xA8707048E0D848D8UL,0xFDF1F104F90C040CUL,
    0x196363DFC67ADF7AUL,0x2F7777C1EE58C158UL,0x30AFAF75459F759FUL,0xE742426384A563A5UL,
    0x7020203040503050UL,0xCBE5E51AD12E1A2EUL,0xEFFDFD0EE1120E12UL,0x08BFBF6D65B76DB7UL,
    0x5581814C19D44CD4UL,0x24181814303C143CUL,0x792626354C5F355FUL,0xB2C3C32F9D712F71UL,
    0x86BEBEE16738E138UL,0xC83535A26AFDA2FDUL,0xC78888CC0B4FCC4FUL,0x652E2E395C4B394BUL,
    0x6A9393573DF957F9UL,0x585555F2AA0DF20DUL,0x61FCFC82E39D829DUL,0xB37A7A47F4C947C9UL,
    0x27C8C8AC8BEFACEFUL,0x88BABAE76F32E732UL,0x4F32322B647D2B7DUL,0x42E6E695D7A495A4UL,
    0x3BC0C0A09BFBA0FBUL,0xAA19199832B398B3UL,0xF69E9ED12768D168UL,0x22A3A37F5D817F81UL,
    0xEE44446688AA66AAUL,0xD654547EA8827E82UL,0xDD3B3BAB76E6ABE6UL,0x950B0B83169E839EUL,
    0xC98C8CCA0345CA45UL,0xBCC7C729957B297BUL,0x056B6BD3D66ED36EUL,0x6C28283C50443C44UL,
    0x2CA7A779558B798BUL,0x81BCBCE2633DE23DUL,0x3116161D2C271D27UL,0x37ADAD76419A769AUL,
    0x96DBDB3BAD4D3B4DUL,0x9E646456C8FA56FAUL,0xA674744EE8D24ED2UL,0x3614141E28221E22UL,
    0xE49292DB3F76DB76UL,0x120C0C0A181E0A1EUL,0xFC48486C90B46CB4UL,0x8FB8B8E46B37E437UL,
    0x789F9F5D25E75DE7UL,0x0FBDBD6E61B26EB2UL,0x694343EF862AEF2AUL,0x35C4C4A693F1A6F1UL,
    0xDA3939A872E3A8E3UL,0xC63131A462F7A4F7UL,0x8AD3D337BD593759UL,0x74F2F28BFF868B86UL,
    0x83D5D532B1563256UL,0x4E8B8B430DC543C5UL,0x856E6E59DCEB59EBUL,0x18DADAB7AFC2B7C2UL,
    0x8E01018C028F8C8FUL,0x1DB1B16479AC64ACUL,0xF19C9CD2236DD26DUL,0x724949E0923BE03BUL,
    0x1FD8D8B4ABC7B4C7UL,0xB9ACACFA4315FA15UL,0xFAF3F307FD090709UL,0xA0CFCF25856F256FUL,
    0x20CACAAF8FEAAFEAUL,0x7DF4F48EF3898E89UL,0x674747E98E20E920UL,0x3810101820281828UL,
    0x0B6F6FD5DE64D564UL,0x73F0F088FB838883UL,0xFB4A4A6F94B16FB1UL,0xCA5C5C72B8967296UL,
    0x54383824706C246CUL,0x5F5757F1AE08F108UL,0x217373C7E652C752UL,0x6497975135F351F3UL,
    0xAECBCB238D652365UL,0x25A1A17C59847C84UL,0x57E8E89CCBBF9CBFUL,0x5D3E3E217C632163UL,
    0xEA9696DD377CDD7CUL,0x1E6161DCC27FDC7FUL,0x9C0D0D861A918691UL,0x9B0F0F851E948594UL,
    0x4BE0E090DBAB90ABUL,0xBA7C7C42F8C642C6UL,0x267171C4E257C457UL,0x29CCCCAA83E5AAE5UL,
    0xE39090D83B73D873UL,0x090606050C0F050FUL,0xF4F7F701F5030103UL,0x2A1C1C1238361236UL,
    0x3CC2C2A39FFEA3FEUL,0x8B6A6A5FD4E15FE1UL,0xBEAEAEF94710F910UL,0x026969D0D26BD06BUL,
    0xBF1717912EA891A8UL,0x7199995829E858E8UL,0x533A3A2774692769UL,0xF72727B94ED0B9D0UL,
    0x91D9D938A9483848UL,0xDEEBEB13CD351335UL,0xE52B2BB356CEB3CEUL,0x7722223344553355UL,
    0x04D2D2BBBFD6BBD6UL,0x39A9A97049907090UL,0x870707890E808980UL,0xC13333A766F2A7F2UL,
    0xEC2D2DB65AC1B6C1UL,0x5A3C3C2278662266UL,0xB81515922AAD92ADUL,0xA9C9C92089602060UL,
    0x5C87874915DB49DBUL,0xB0AAAAFF4F1AFF1AUL,0xD8505078A0887888UL,0x2BA5A57A518E7A8EUL,
    0x8903038F068A8F8AUL,0x4A5959F8B213F813UL,0x92090980129B809BUL,0x231A1A1734391739UL,
    0x106565DACA75DA75UL,0x84D7D731B5533153UL,0xD58484C61351C651UL,0x03D0D0B8BBD3B8D3UL,
    0xDC8282C31F5EC35EUL,0xE22929B052CBB0CBUL,0xC35A5A77B4997799UL,0x2D1E1E113C331133UL,
    0x3D7B7BCBF646CB46UL,0xB7A8A8FC4B1FFC1FUL,0x0C6D6DD6DA61D661UL,0x622C2C3A584E3A4EUL,
};

__constant ulong GROESTL_T7[256] = {
    0xC6C6A597F4A5F432UL,0xF8F884EB9784976FUL,0xEEEE99C7B099B05EUL,0xF6F68DF78C8D8C7AUL,
    0xFFFF0DE5170D17E8UL,0xD6D6BDB7DCBDDC0AUL,0xDEDEB1A7C8B1C816UL,0x91915439FC54FC6DUL,
    0x606050C0F050F090UL,0x0202030405030507UL,0xCECEA987E0A9E02EUL,0x56567DAC877D87D1UL,
    0xE7E719D52B192BCCUL,0xB5B56271A662A613UL,0x4D4DE69A31E6317CUL,0xECEC9AC3B59AB559UL,
    0x8F8F4505CF45CF40UL,0x1F1F9D3EBC9DBCA3UL,0x89894009C040C049UL,0xFAFA87EF92879268UL,
    0xEFEF15C53F153FD0UL,0xB2B2EB7F26EB2694UL,0x8E8EC90740C940CEUL,0xFBFB0BED1D0B1DE6UL,
    0x4141EC822FEC2F6EUL,0xB3B3677DA967A91AUL,0x5F5FFDBE1CFD1C43UL,0x4545EA8A25EA2560UL,
    0x2323BF46DABFDAF9UL,0x5353F7A602F70251UL,0xE4E496D3A196A145UL,0x9B9B5B2DED5BED76UL,
    0x7575C2EA5DC25D28UL,0xE1E11CD9241C24C5UL,0x3D3DAE7AE9AEE9D4UL,0x4C4C6A98BE6ABEF2UL,
    0x6C6C5AD8EE5AEE82UL,0x7E7E41FCC341C3BDUL,0xF5F502F1060206F3UL,0x83834F1DD14FD152UL,
    0x68685CD0E45CE48CUL,0x5151F4A207F40756UL,0xD1D134B95C345C8DUL,0xF9F908E9180818E1UL,
    0xE2E293DFAE93AE4CUL,0xABAB734D9573953EUL,0x626253C4F553F597UL,0x2A2A3F54413F416BUL,
    0x08080C10140C141CUL,0x95955231F652F663UL,0x4646658CAF65AFE9UL,0x9D9D5E21E25EE27FUL,
    0x3030286078287848UL,0x3737A16EF8A1F8CFUL,0x0A0A0F14110F111BUL,0x2F2FB55EC4B5C4EBUL,
    0x0E0E091C1B091B15UL,0x242436485A365A7EUL,0x1B1B9B36B69BB6ADUL,0xDFDF3DA5473D4798UL,
    0xCDCD26816A266AA7UL,0x4E4E699CBB69BBF5UL,0x7F7FCDFE4CCD4C33UL,0xEAEA9FCFBA9FBA50UL,
    0x12121B242D1B2D3FUL,0x1D1D9E3AB99EB9A4UL,0x585874B09C749CC4UL,0x34342E68722E7246UL,
    0x36362D6C772D7741UL,0xDCDCB2A3CDB2CD11UL,0xB4B4EE7329EE299DUL,0x5B5BFBB616FB164DUL,
    0xA4A4F65301F601A5UL,0x76764DECD74DD7A1UL,0xB7B76175A361A314UL,0x7D7DCEFA49CE4934UL,
    0x52527BA48D7B8DDFUL,0xDDDD3EA1423E429FUL,0x5E5E71BC937193CDUL,0x13139726A297A2B1UL,
    0xA6A6F55704F504A2UL,0xB9B96869B868B801UL,0x0000000000000000UL,0xC1C12C99742C74B5UL,
    0x40406080A060A0E0UL,0xE3E31FDD211F21C2UL,0x7979C8F243C8433AUL,0xB6B6ED772CED2C9AUL,
    0xD4D4BEB3D9BED90DUL,0x8D8D4601CA46CA47UL,0x6767D9CE70D97017UL,0x72724BE4DD4BDDAFUL,
    0x9494DE3379DE79EDUL,0x9898D42B67D467FFUL,0xB0B0E87B23E82393UL,0x85854A11DE4ADE5BUL,
    0xBBBB6B6DBD6BBD06UL,0xC5C52A917E2A7EBBUL,0x4F4FE59E34E5347BUL,0xEDED16C13A163AD7UL,
    0x8686C51754C554D2UL,0x9A9AD72F62D762F8UL,0x666655CCFF55FF99UL,0x11119422A794A7B6UL,
    0x8A8ACF0F4ACF4AC0UL,0xE9E910C9301030D9UL,0x040406080A060A0EUL,0xFEFE81E798819866UL,
    0xA0A0F05B0BF00BABUL,0x787844F0CC44CCB4UL,0x2525BA4AD5BAD5F0UL,0x4B4BE3963EE33E75UL,
    0xA2A2F35F0EF30EACUL,0x5D5DFEBA19FE1944UL,0x8080C01B5BC05BDBUL,0x05058A0A858A8580UL,
    0x3F3FAD7EECADECD3UL,0x2121BC42DFBCDFFEUL,0x707048E0D848D8A8UL,0xF1F104F90C040CFDUL,
    0x6363DFC67ADF7A19UL,0x7777C1EE58C1582FUL,0xAFAF75459F759F30UL,0x42426384A563A5E7UL,
    0x2020304050305070UL,0xE5E51AD12E1A2ECBUL,0xFDFD0EE1120E12EFUL,0xBFBF6D65B76DB708UL,
    0x81814C19D44CD455UL,0x181814303C143C24UL,0x2626354C5F355F79UL,0xC3C32F9D712F71B2UL,
    0xBEBEE16738E13886UL,0x3535A26AFDA2FDC8UL,0x8888CC0B4FCC4FC7UL,0x2E2E395C4B394B65UL,
    0x9393573DF957F96AUL,0x5555F2AA0DF20D58UL,0xFCFC82E39D829D61UL,0x7A7A47F4C947C9B3UL,
    0xC8C8AC8BEFACEF27UL,0xBABAE76F32E73288UL,0x32322B647D2B7D4FUL,0xE6E695D7A495A442UL,
    0xC0C0A09BFBA0FB3BUL,0x19199832B398B3AAUL,0x9E9ED12768D168F6UL,0xA3A37F5D817F8122UL,
    0x44446688AA66AAEEUL,0x54547EA8827E82D6UL,0x3B3BAB76E6ABE6DDUL,0x0B0B83169E839E95UL,
    0x8C8CCA0345CA45C9UL,0xC7C729957B297BBCUL,0x6B6BD3D66ED36E05UL,0x28283C50443C446CUL,
    0xA7A779558B798B2CUL,0xBCBCE2633DE23D81UL,0x16161D2C271D2731UL,0xADAD76419A769A37UL,
    0xDBDB3BAD4D3B4D96UL,0x646456C8FA56FA9EUL,0x74744EE8D24ED2A6UL,0x14141E28221E2236UL,
    0x9292DB3F76DB76E4UL,0x0C0C0A181E0A1E12UL,0x48486C90B46CB4FCUL,0xB8B8E46B37E4378FUL,
    0x9F9F5D25E75DE778UL,0xBDBD6E61B26EB20FUL,0x4343EF862AEF2A69UL,0xC4C4A693F1A6F135UL,
    0x3939A872E3A8E3DAUL,0x3131A462F7A4F7C6UL,0xD3D337BD5937598AUL,0xF2F28BFF868B8674UL,
    0xD5D532B156325683UL,0x8B8B430DC543C54EUL,0x6E6E59DCEB59EB85UL,0xDADAB7AFC2B7C218UL,
    0x01018C028F8C8F8EUL,0xB1B16479AC64AC1DUL,0x9C9CD2236DD26DF1UL,0x4949E0923BE03B72UL,
    0xD8D8B4ABC7B4C71FUL,0xACACFA4315FA15B9UL,0xF3F307FD090709FAUL,0xCFCF25856F256FA0UL,
    0xCACAAF8FEAAFEA20UL,0xF4F48EF3898E897DUL,0x4747E98E20E92067UL,0x1010182028182838UL,
    0x6F6FD5DE64D5640BUL,0xF0F088FB83888373UL,0x4A4A6F94B16FB1FBUL,0x5C5C72B8967296CAUL,
    0x383824706C246C54UL,0x5757F1AE08F1085FUL,0x7373C7E652C75221UL,0x97975135F351F364UL,
    0xCBCB238D652365AEUL,0xA1A17C59847C8425UL,0xE8E89CCBBF9CBF57UL,0x3E3E217C6321635DUL,
    0x9696DD377CDD7CEAUL,0x6161DCC27FDC7F1EUL,0x0D0D861A9186919CUL,0x0F0F851E9485949BUL,
    0xE0E090DBAB90AB4BUL,0x7C7C42F8C642C6BAUL,0x7171C4E257C45726UL,0xCCCCAA83E5AAE529UL,
    0x9090D83B73D873E3UL,0x0606050C0F050F09UL,0xF7F701F5030103F4UL,0x1C1C12383612362AUL,
    0xC2C2A39FFEA3FE3CUL,0x6A6A5FD4E15FE18BUL,0xAEAEF94710F910BEUL,0x6969D0D26BD06B02UL,
    0x1717912EA891A8BFUL,0x99995829E858E871UL,0x3A3A277469276953UL,0x2727B94ED0B9D0F7UL,
    0xD9D938A948384891UL,0xEBEB13CD351335DEUL,0x2B2BB356CEB3CEE5UL,0x2222334455335577UL,
    0xD2D2BBBFD6BBD604UL,0xA9A9704990709039UL,0x0707890E80898087UL,0x3333A766F2A7F2C1UL,
    0x2D2DB65AC1B6C1ECUL,0x3C3C22786622665AUL,0x1515922AAD92ADB8UL,0xC9C92089602060A9UL,
    0x87874915DB49DB5CUL,0xAAAAFF4F1AFF1AB0UL,0x505078A0887888D8UL,0xA5A57A518E7A8E2BUL,
    0x03038F068A8F8A89UL,0x5959F8B213F8134AUL,0x090980129B809B92UL,0x1A1A173439173923UL,
    0x6565DACA75DA7510UL,0xD7D731B553315384UL,0x8484C61351C651D5UL,0xD0D0B8BBD3B8D303UL,
    0x8282C31F5EC35EDCUL,0x2929B052CBB0CBE2UL,0x5A5A77B4997799C3UL,0x1E1E113C3311332DUL,
    0x7B7BCBF646CB463DUL,0xA8A8FC4B1FFC1FB7UL,0x6D6DD6DA61D6610CUL,0x2C2C3A584E3A4E62UL
};

#define GROESTL_EXT_BYTE(var,n) as_uchar8(var).s##n

#define GROESTL_COLUMN(gx, gy, gi, c0, c1, c2, c3, c4, c5, c6, c7)           \
  gy##gi = groestl_lt0[GROESTL_EXT_BYTE(gx##c0,0)]             \
    ^ groestl_lt1[GROESTL_EXT_BYTE(gx##c1,1)]                \
    ^ groestl_lt2[GROESTL_EXT_BYTE(gx##c2,2)]                \
    ^ groestl_lt3[GROESTL_EXT_BYTE(gx##c3,3)]                \
    ^ groestl_lt4[GROESTL_EXT_BYTE(gx##c4,4)]                \
    ^ groestl_lt5[GROESTL_EXT_BYTE(gx##c5,5)]                \
    ^ groestl_lt6[GROESTL_EXT_BYTE(gx##c6,6)]                \
    ^ groestl_lt7[GROESTL_EXT_BYTE(gx##c7,7)]

#define GROESTL_RND1024Q(X, Y, R)                     \
    X##0 ^= 0xFFFFFFFFFFFFFFFFUL ^ R; \
    X##1 ^= 0xEFFFFFFFFFFFFFFFUL ^ R; \
    X##2 ^= 0xDFFFFFFFFFFFFFFFUL ^ R; \
    X##3 ^= 0xCFFFFFFFFFFFFFFFUL ^ R; \
    X##4 ^= 0xBFFFFFFFFFFFFFFFUL ^ R; \
    X##5 ^= 0xAFFFFFFFFFFFFFFFUL ^ R; \
    X##6 ^= 0x9FFFFFFFFFFFFFFFUL ^ R; \
    X##7 ^= 0x8FFFFFFFFFFFFFFFUL ^ R; \
    X##8 ^= 0x7FFFFFFFFFFFFFFFUL ^ R; \
    X##9 ^= 0x6FFFFFFFFFFFFFFFUL ^ R; \
    X##A ^= 0x5FFFFFFFFFFFFFFFUL ^ R; \
    X##B ^= 0x4FFFFFFFFFFFFFFFUL ^ R; \
    X##C ^= 0x3FFFFFFFFFFFFFFFUL ^ R; \
    X##D ^= 0x2FFFFFFFFFFFFFFFUL ^ R; \
    X##E ^= 0x1FFFFFFFFFFFFFFFUL ^ R; \
    X##F ^= 0x0FFFFFFFFFFFFFFFUL ^ R; \
    GROESTL_COLUMN(X, Y, F, 0, 2, 4, A, F, 1, 3, 5);         \
    GROESTL_COLUMN(X, Y, E, F, 1, 3, 9, E, 0, 2, 4);         \
    GROESTL_COLUMN(X, Y, D, E, 0, 2, 8, D, F, 1, 3);         \
    GROESTL_COLUMN(X, Y, C, D, F, 1, 7, C, E, 0, 2);         \
    GROESTL_COLUMN(X, Y, B, C, E, 0, 6, B, D, F, 1);         \
    GROESTL_COLUMN(X, Y, A, B, D, F, 5, A, C, E, 0);         \
    GROESTL_COLUMN(X, Y, 9, A, C, E, 4, 9, B, D, F);         \
    GROESTL_COLUMN(X, Y, 8, 9, B, D, 3, 8, A, C, E);         \
    GROESTL_COLUMN(X, Y, 7, 8, A, C, 2, 7, 9, B, D);         \
    GROESTL_COLUMN(X, Y, 6, 7, 9, B, 1, 6, 8, A, C);         \
    GROESTL_COLUMN(X, Y, 5, 6, 8, A, 0, 5, 7, 9, B);         \
    GROESTL_COLUMN(X, Y, 4, 5, 7, 9, F, 4, 6, 8, A);         \
    GROESTL_COLUMN(X, Y, 3, 4, 6, 8, E, 3, 5, 7, 9);         \
    GROESTL_COLUMN(X, Y, 2, 3, 5, 7, D, 2, 4, 6, 8);         \
    GROESTL_COLUMN(X, Y, 1, 2, 4, 6, C, 1, 3, 5, 7);         \
    GROESTL_COLUMN(X, Y, 0, 1, 3, 5, B, 0, 2, 4, 6);

#define GROESTL_RND1024P(X, Y, R)             \
    X##0 ^= R; \
    X##1 ^= 0x10UL ^ R; \
    X##2 ^= 0x20UL ^ R; \
    X##3 ^= 0x30UL ^ R; \
    X##4 ^= 0x40UL ^ R; \
    X##5 ^= 0x50UL ^ R; \
    X##6 ^= 0x60UL ^ R; \
    X##7 ^= 0x70UL ^ R; \
    X##8 ^= 0x80UL ^ R; \
    X##9 ^= 0x90UL ^ R; \
    X##A ^= 0xA0UL ^ R; \
    X##B ^= 0xB0UL ^ R; \
    X##C ^= 0xC0UL ^ R; \
    X##D ^= 0xD0UL ^ R; \
    X##E ^= 0xE0UL ^ R; \
    X##F ^= 0xF0UL ^ R; \
    GROESTL_COLUMN(X, Y, F, F, 0, 1, 2, 3, 4, 5, A); \
    GROESTL_COLUMN(X, Y, E, E, F, 0, 1, 2, 3, 4, 9); \
    GROESTL_COLUMN(X, Y, D, D, E, F, 0, 1, 2, 3, 8); \
    GROESTL_COLUMN(X, Y, C, C, D, E, F, 0, 1, 2, 7); \
    GROESTL_COLUMN(X, Y, B, B, C, D, E, F, 0, 1, 6); \
    GROESTL_COLUMN(X, Y, A, A, B, C, D, E, F, 0, 5); \
    GROESTL_COLUMN(X, Y, 9, 9, A, B, C, D, E, F, 4); \
    GROESTL_COLUMN(X, Y, 8, 8, 9, A, B, C, D, E, 3); \
    GROESTL_COLUMN(X, Y, 7, 7, 8, 9, A, B, C, D, 2); \
    GROESTL_COLUMN(X, Y, 6, 6, 7, 8, 9, A, B, C, 1); \
    GROESTL_COLUMN(X, Y, 5, 5, 6, 7, 8, 9, A, B, 0); \
    GROESTL_COLUMN(X, Y, 4, 4, 5, 6, 7, 8, 9, A, F); \
    GROESTL_COLUMN(X, Y, 3, 3, 4, 5, 6, 7, 8, 9, E); \
    GROESTL_COLUMN(X, Y, 2, 2, 3, 4, 5, 6, 7, 8, D); \
    GROESTL_COLUMN(X, Y, 1, 1, 2, 3, 4, 5, 6, 7, C); \
    GROESTL_COLUMN(X, Y, 0, 0, 1, 2, 3, 4, 5, 6, B);

inline void groestl_load_tables(__local ulong* tables)
{
    for (int i = 0; i < 16; i++)
        vstore16(vload16(i, GROESTL_T0), i, tables);
    tables += 256;
    for (int i = 0; i < 16; i++)
        vstore16(vload16(i, GROESTL_T1), i, tables);
    tables += 256;
    for (int i = 0; i < 16; i++)
        vstore16(vload16(i, GROESTL_T2), i, tables);
    tables += 256;
    for (int i = 0; i < 16; i++)
        vstore16(vload16(i, GROESTL_T3), i, tables);
    tables += 256;
    for (int i = 0; i < 16; i++)
        vstore16(vload16(i, GROESTL_T4), i, tables);
    tables += 256;
    for (int i = 0; i < 16; i++)
        vstore16(vload16(i, GROESTL_T5), i, tables);
    tables += 256;
    for (int i = 0; i < 16; i++)
        vstore16(vload16(i, GROESTL_T6), i, tables);
    tables += 256;
    for (int i = 0; i < 16; i++)
        vstore16(vload16(i, GROESTL_T7), i, tables);
    barrier(CLK_LOCAL_MEM_FENCE);
}


inline uint groestl512_116_last(const uint *msg, __local ulong *tables)
{
#ifdef BROKEN_LOCAL
    __constant ulong* groestl_lt0 = GROESTL_T0;
    __constant ulong* groestl_lt1 = GROESTL_T1;
    __constant ulong* groestl_lt2 = GROESTL_T2;
    __constant ulong* groestl_lt3 = GROESTL_T3;
    __constant ulong* groestl_lt4 = GROESTL_T4;
    __constant ulong* groestl_lt5 = GROESTL_T5;
    __constant ulong* groestl_lt6 = GROESTL_T6;
    __constant ulong* groestl_lt7 = GROESTL_T7;
#else
    groestl_load_tables(tables);
    __local ulong* groestl_lt0 = tables;
    __local ulong* groestl_lt1 = groestl_lt0 + 256;
    __local ulong* groestl_lt2 = groestl_lt1 + 256;
    __local ulong* groestl_lt3 = groestl_lt2 + 256;
    __local ulong* groestl_lt4 = groestl_lt3 + 256;
    __local ulong* groestl_lt5 = groestl_lt4 + 256;
    __local ulong* groestl_lt6 = groestl_lt5 + 256;
    __local ulong* groestl_lt7 = groestl_lt6 + 256;
#endif

    ulong Z0, Z1, Z2, Z3, Z4, Z5, Z6, Z7, Z8, Z9, ZA, ZB, ZC, ZD, ZE, ZF;
    ulong H0, H1, H2, H3, H4, H5, H6, H7, H8, H9, HA, HB, HC, HD, HE, HF;
    ulong Y0, Y1, Y2, Y3, Y4, Y5, Y6, Y7, Y8, Y9, YA, YB, YC, YD, YE, YF;
    ulong T0, T1, T2, T3, T4, T5, T6, T7, T8, T9, TA, TB, TC, TD, TE, TF;

    Z0 = T0 = as_ulong((uint2)(SWAP32(msg[0]), SWAP32(msg[1])));
    Z1 = T1 = as_ulong((uint2)(SWAP32(msg[2]), SWAP32(msg[3])));
    Z2 = T2 = as_ulong((uint2)(SWAP32(msg[4]), SWAP32(msg[5])));
    Z3 = T3 = as_ulong((uint2)(SWAP32(msg[6]), SWAP32(msg[7])));
    Z4 = T4 = as_ulong((uint2)(SWAP32(msg[8]), SWAP32(msg[9])));
    Z5 = T5 = as_ulong((uint2)(SWAP32(msg[10]), SWAP32(msg[11])));
    Z6 = T6 = as_ulong((uint2)(SWAP32(msg[12]), SWAP32(msg[13])));
    Z7 = T7 = as_ulong((uint2)(SWAP32(msg[14]), SWAP32(msg[15])));
    Z8 = T8 = as_ulong((uint2)(SWAP32(msg[16]), SWAP32(msg[17])));
    Z9 = T9 = as_ulong((uint2)(SWAP32(msg[18]), SWAP32(msg[19])));
    ZA = TA = as_ulong((uint2)(SWAP32(msg[20]), SWAP32(msg[21])));
    ZB = TB = as_ulong((uint2)(SWAP32(msg[22]), SWAP32(msg[23])));
    ZC = TC = as_ulong((uint2)(SWAP32(msg[24]), SWAP32(msg[25])));
    ZD = TD = as_ulong((uint2)(SWAP32(msg[26]), SWAP32(msg[27])));
    ZE = TE = as_ulong((uint2)(SWAP32(msg[28]), 0x80));
    ZF = 0x0100000000000000UL;
    TF = 0x0100000000000000UL ^ 0x0002000000000000UL;

    GROESTL_RND1024Q(Z, Y, 0);
    GROESTL_RND1024Q(Y, Z, 0x0100000000000000UL);
    GROESTL_RND1024Q(Z, Y, 0x0200000000000000UL);
    GROESTL_RND1024Q(Y, Z, 0x0300000000000000UL);
    GROESTL_RND1024Q(Z, Y, 0x0400000000000000UL);
    GROESTL_RND1024Q(Y, Z, 0x0500000000000000UL);
    GROESTL_RND1024Q(Z, Y, 0x0600000000000000UL);
    GROESTL_RND1024Q(Y, Z, 0x0700000000000000UL);
    GROESTL_RND1024Q(Z, Y, 0x0800000000000000UL);
    GROESTL_RND1024Q(Y, Z, 0x0900000000000000UL);
    GROESTL_RND1024Q(Z, Y, 0x0a00000000000000UL);
    GROESTL_RND1024Q(Y, Z, 0x0b00000000000000UL);
    GROESTL_RND1024Q(Z, Y, 0x0c00000000000000UL);

    GROESTL_RND1024Q(Y, H, 0x0d00000000000000UL);

    GROESTL_RND1024P(T, Z, 0);
    GROESTL_RND1024P(Z, Y, 1UL);
    GROESTL_RND1024P(Y, Z, 2UL);
    GROESTL_RND1024P(Z, Y, 3UL);
    GROESTL_RND1024P(Y, Z, 4UL);
    GROESTL_RND1024P(Z, Y, 5UL);
    GROESTL_RND1024P(Y, Z, 6UL);
    GROESTL_RND1024P(Z, Y, 7UL);
    GROESTL_RND1024P(Y, Z, 8UL);
    GROESTL_RND1024P(Z, Y, 9UL);
    GROESTL_RND1024P(Y, Z, 10UL);
    GROESTL_RND1024P(Z, Y, 11UL);
    GROESTL_RND1024P(Y, Z, 12UL);
    GROESTL_RND1024P(Z, Y, 13UL);

    H0 ^= Y0; H1 ^= Y1; H2 ^= Y2; H3 ^= Y3;
    H4 ^= Y4; H5 ^= Y5; H6 ^= Y6; H7 ^= Y7;
    H8 ^= Y8; H9 ^= Y9; HA ^= YA; HB ^= YB;
    HC ^= YC; HD ^= YD; HE ^= YE; HF ^= YF ^ 0x0002000000000000UL;

    TB = HB;

    GROESTL_RND1024P(H, Y, 0);
    GROESTL_RND1024P(Y, Z, 1UL);
    GROESTL_RND1024P(Z, Y, 2UL);
    GROESTL_RND1024P(Y, Z, 3UL);
    GROESTL_RND1024P(Z, Y, 4UL);
    GROESTL_RND1024P(Y, Z, 5UL);
    GROESTL_RND1024P(Z, Y, 6UL);
    GROESTL_RND1024P(Y, Z, 7UL);
    GROESTL_RND1024P(Z, Y, 8UL);
    GROESTL_RND1024P(Y, Z, 9UL);
    GROESTL_RND1024P(Z, Y, 10UL);
    GROESTL_RND1024P(Y, Z, 11UL);
    GROESTL_RND1024P(Z, Y, 12UL);
    GROESTL_RND1024P(Y, T, 13UL);
    return as_uint2(HB ^ TB).y;
}

inline uint2 rolv2_1(const uint2 x, const uint y)
{
    return (uint2)((x.x << y) ^ (x.y >> (32 - y)), (x.y << y) ^ (x.x >> (32 - y)));
}
inline uint2 rolv2_2(const uint2 x, const uint y)
{
    return (uint2)((x.y << y) ^ (x.x >> (32 - y)), (x.x << y) ^ (x.y >> (32 - y)));
}

#define keccak_round(const0, const1) \
{ \
    BCa = Aba^Aga^Aka^Ama^Asa; BCe = Abe^Age^Ake^Ame^Ase; \
    BCi = Abi^Agi^Aki^Ami^Asi; BCo = Abo^Ago^Ako^Amo^Aso; BCu = Abu^Agu^Aku^Amu^Asu; \
    Da = BCu ^ rolv2_1(BCe, 1); De = BCa ^ rolv2_1(BCi, 1); \
    Di = BCe ^ rolv2_1(BCo, 1); Do = BCi ^ rolv2_1(BCu, 1); \
    Du = BCo ^ rolv2_1(BCa, 1); \
    Aba ^= Da; BCa = Aba; Age ^= De; BCe = rolv2_2(Age, 44); \
    Aki ^= Di; BCi = rolv2_2(Aki, 43); \
    Amo ^= Do; BCo = rolv2_1(Amo, 21); \
    Asu ^= Du; BCu = rolv2_1(Asu, 14); \
    Eba = bitselect(BCa ^ BCi, BCa, BCe) ^ const0; \
    Ebe = bitselect(BCe ^ BCo, BCe, BCi); \
    Ebi = bitselect(BCi ^ BCu, BCi, BCo); \
    Ebo = bitselect(BCo ^ BCa, BCo, BCu); \
    Ebu = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = rolv2_1(Abo ^ Do, 28); BCe = rolv2_1(Agu ^ Du, 20); BCi = rolv2_1(Aka ^ Da, 3); \
    BCo = rolv2_2(Ame ^ De, 45); BCu = rolv2_2(Asi ^ Di, 61); \
    Ega = bitselect(BCa ^ BCi, BCa, BCe); \
    Ege = bitselect(BCe ^ BCo, BCe, BCi); \
    Egi = bitselect(BCi ^ BCu, BCi, BCo); \
    Ego = bitselect(BCo ^ BCa, BCo, BCu); \
    Egu = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = rolv2_1(Abe ^ De, 1); \
    BCe = rolv2_1(Agi ^ Di, 6); \
    BCi = rolv2_1(Ako ^ Do, 25); \
    BCo = rolv2_1(Amu ^ Du, 8); \
    BCu = rolv2_1(Asa ^ Da, 18); \
    Eka = bitselect(BCa ^ BCi, BCa, BCe); \
    Eke = bitselect(BCe ^ BCo, BCe, BCi); \
    Eki = bitselect(BCi ^ BCu, BCi, BCo); \
    Eko = bitselect(BCo ^ BCa, BCo, BCu); \
    Eku = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = rolv2_1(Abu ^ Du, 27); \
    BCe = rolv2_2(Aga ^ Da, 36); \
    BCi = rolv2_1(Ake ^ De, 10); \
    BCo = rolv2_1(Ami ^ Di, 15); \
    BCu = rolv2_2(Aso ^ Do, 56); \
    Ema = bitselect(BCa ^ BCi, BCa, BCe); \
    Eme = bitselect(BCe ^ BCo, BCe, BCi); \
    Emi = bitselect(BCi ^ BCu, BCi, BCo); \
    Emo = bitselect(BCo ^ BCa, BCo, BCu); \
    Emu = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = rolv2_2(Abi ^ Di, 62); \
    BCe = rolv2_2(Ago ^ Do, 55); \
    BCi = rolv2_2(Aku ^ Du, 39); \
    BCo = rolv2_2(Ama ^ Da, 41); \
    BCu = rolv2_1(Ase ^ De, 2); \
    Esa = bitselect(BCa ^ BCi, BCa, BCe); \
    Ese = bitselect(BCe ^ BCo, BCe, BCi); \
    Esi = bitselect(BCi ^ BCu, BCi, BCo); \
    Eso = bitselect(BCo ^ BCa, BCo, BCu); \
    Esu = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = Eba^Ega^Eka^Ema^Esa; \
    BCe = Ebe^Ege^Eke^Eme^Ese; \
    BCi = Ebi^Egi^Eki^Emi^Esi; \
    BCo = Ebo^Ego^Eko^Emo^Eso; \
    BCu = Ebu^Egu^Eku^Emu^Esu; \
    Da = BCu ^ rolv2_1(BCe, 1); \
    De = BCa ^ rolv2_1(BCi, 1); \
    Di = BCe ^ rolv2_1(BCo, 1); \
    Do = BCi ^ rolv2_1(BCu, 1); \
    Du = BCo ^ rolv2_1(BCa, 1); \
    BCa = Eba ^ Da; \
    BCe = rolv2_2(Ege ^ De, 44); \
    BCi = rolv2_2(Eki ^ Di, 43); \
    BCo = rolv2_1(Emo ^ Do, 21); \
    BCu = rolv2_1(Esu ^ Du, 14); \
    Aba = bitselect(BCa ^ BCi, BCa, BCe) ^ const1; \
    Abe = bitselect(BCe ^ BCo, BCe, BCi); \
    Abi = bitselect(BCi ^ BCu, BCi, BCo); \
    Abo = bitselect(BCo ^ BCa, BCo, BCu); \
    Abu = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = rolv2_1(Ebo ^ Do, 28); \
    BCe = rolv2_1(Egu ^ Du, 20); \
    BCi = rolv2_1(Eka ^ Da, 3); \
    BCo = rolv2_2(Eme ^ De, 45); \
    BCu = rolv2_2(Esi ^ Di, 61); \
    Aga = bitselect(BCa ^ BCi, BCa, BCe); \
    Age = bitselect(BCe ^ BCo, BCe, BCi); \
    Agi = bitselect(BCi ^ BCu, BCi, BCo); \
    Ago = bitselect(BCo ^ BCa, BCo, BCu); \
    Agu = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = rolv2_1(Ebe ^ De, 1); \
    BCe = rolv2_1(Egi ^ Di, 6); \
    BCi = rolv2_1(Eko ^ Do, 25); \
    BCo = rolv2_1(Emu ^ Du, 8); \
    BCu = rolv2_1(Esa ^ Da, 18); \
    Aka = bitselect(BCa ^ BCi, BCa, BCe); \
    Ake = bitselect(BCe ^ BCo, BCe, BCi); \
    Aki = bitselect(BCi ^ BCu, BCi, BCo); \
    Ako = bitselect(BCo ^ BCa, BCo, BCu); \
    Aku = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = rolv2_1(Ebu ^ Du, 27); \
    BCe = rolv2_2(Ega ^ Da, 36); \
    BCi = rolv2_1(Eke ^ De, 10); \
    BCo = rolv2_1(Emi ^ Di, 15); \
    BCu = rolv2_2(Eso ^ Do, 56); \
    Ama = bitselect(BCa ^ BCi, BCa, BCe); \
    Ame = bitselect(BCe ^ BCo, BCe, BCi); \
    Ami = bitselect(BCi ^ BCu, BCi, BCo); \
    Amo = bitselect(BCo ^ BCa, BCo, BCu); \
    Amu = bitselect(BCu ^ BCe, BCu, BCa); \
    BCa = rolv2_2(Ebi ^ Di, 62); \
    BCe = rolv2_2(Ego ^ Do, 55); \
    BCi = rolv2_2(Eku ^ Du, 39); \
    BCo = rolv2_2(Ema ^ Da, 41); \
    BCu = rolv2_1(Ese ^ De, 2); \
    Asa = bitselect(BCa ^ BCi, BCa, BCe); \
    Ase = bitselect(BCe ^ BCo, BCe, BCi); \
    Asi = bitselect(BCi ^ BCu, BCi, BCo); \
    Aso = bitselect(BCo ^ BCa, BCo, BCu); \
    Asu = bitselect(BCu ^ BCe, BCu, BCa); \
}

inline uint keccak512_116_last(uint * restrict msg)
{
    uint2 Aba, Abe, Abi, Abo, Abu;
    uint2 Aga, Age, Agi, Ago, Agu;
    uint2 Aka, Ake, Aki, Ako, Aku;
    uint2 Ama, Ame, Ami, Amo, Amu;
    uint2 Asa, Ase, Asi, Aso, Asu;
    uint2 BCa, BCe, BCi, BCo, BCu;
    uint2 Da, De, Di, Do, Du;
    uint2 Eba, Ebe, Ebi, Ebo, Ebu;
    uint2 Ega, Ege, Egi, Ego, Egu;
    uint2 Eka, Eke, Eki, Eko, Eku;
    uint2 Ema, Eme, Emi, Emo, Emu;
    uint2 Esa, Ese, Esi, Eso, Esu;

    Aba = (uint2)(SWAP32(msg[0]), SWAP32(msg[1]));
    Abe = (uint2)(SWAP32(msg[2]), SWAP32(msg[3]));
    Abi = (uint2)(SWAP32(msg[4]), SWAP32(msg[5]));
    Abo = (uint2)(SWAP32(msg[6]), SWAP32(msg[7]));
    Abu = (uint2)(SWAP32(msg[8]), SWAP32(msg[9]));
    Aga = (uint2)(SWAP32(msg[10]), SWAP32(msg[11]));
    Age = (uint2)(SWAP32(msg[12]), SWAP32(msg[13]));
    Agi = (uint2)(SWAP32(msg[14]), SWAP32(msg[15]));
    Ago = (uint2)(SWAP32(msg[16]), SWAP32(msg[17]));
    Agu = 0; Aka = 0; Ake = 0; Aki = 0; Ako = 0; Aku = 0; Ama = 0; Ame = 0;
    Ami = 0; Amo = 0; Amu = 0; Asa = 0; Ase = 0; Asi = 0; Aso = 0; Asu = 0;

    keccak_round((uint2)(1, 0), (uint2)(0x00008082, 0))
    keccak_round((uint2)(0x0000808AU, 0x80000000U), (uint2)(0x80008000U, 0x80000000U))
    keccak_round((uint2)(0x0000808BU, 0), (uint2)(0x80000001U, 0))
    keccak_round((uint2)(0x80008081U, 0x80000000U), (uint2)(0x00008009U, 0x80000000U))
    keccak_round((uint2)(0x0000008AU, 0), (uint2)(0x00000088U, 0))
    keccak_round((uint2)(0x80008009U, 0), (uint2)(0x8000000AU, 0))
    keccak_round((uint2)(0x8000808BU, 0), (uint2)(0x0000008BU, 0x80000000U))
    keccak_round((uint2)(0x00008089U, 0x80000000U), (uint2)(0x00008003U, 0x80000000U))
    keccak_round((uint2)(0x00008002U, 0x80000000U), (uint2)(0x00000080U, 0x80000000U))
    keccak_round((uint2)(0x0000800AU, 0), (uint2)(0x8000000AU, 0x80000000U))
    keccak_round((uint2)(0x80008081U, 0x80000000U), (uint2)(0x00008080U, 0x80000000U))
    keccak_round((uint2)(0x80000001U, 0), (uint2)(0x80008008U, 0x80000000U))

    Aba ^= (uint2)(SWAP32(msg[18]), SWAP32(msg[19]));
    Abe ^= (uint2)(SWAP32(msg[20]), SWAP32(msg[21]));
    Abi ^= (uint2)(SWAP32(msg[22]), SWAP32(msg[23]));
    Abo ^= (uint2)(SWAP32(msg[24]), SWAP32(msg[25]));
    Abu ^= (uint2)(SWAP32(msg[26]), SWAP32(msg[27]));
    Aga ^= (uint2)(SWAP32(msg[28]), 1);

    Ago ^= (uint2)(0, 0x80000000U);

    keccak_round((uint2)(1, 0), (uint2)(0x00008082, 0))
    keccak_round((uint2)(0x0000808AU, 0x80000000U), (uint2)(0x80008000U, 0x80000000U))
    keccak_round((uint2)(0x0000808BU, 0), (uint2)(0x80000001U, 0))
    keccak_round((uint2)(0x80008081U, 0x80000000U), (uint2)(0x00008009U, 0x80000000U))
    keccak_round((uint2)(0x0000008AU, 0), (uint2)(0x00000088U, 0))
    keccak_round((uint2)(0x80008009U, 0), (uint2)(0x8000000AU, 0))
    keccak_round((uint2)(0x8000808BU, 0), (uint2)(0x0000008BU, 0x80000000U))
    keccak_round((uint2)(0x00008089U, 0x80000000U), (uint2)(0x00008003U, 0x80000000U))
    keccak_round((uint2)(0x00008002U, 0x80000000U), (uint2)(0x00000080U, 0x80000000U))
    keccak_round((uint2)(0x0000800AU, 0), (uint2)(0x8000000AU, 0x80000000U))
    keccak_round((uint2)(0x80008081U, 0x80000000U), (uint2)(0x00008080U, 0x80000000U))
    keccak_round((uint2)(0x80000001U, 0), (uint2)(0x80008008U, 0x80000000U))
    return Abo.y;
}

inline uint Ma(uint x, uint y, uint z) { return bitselect(x, y, z ^ x); }
inline uint Sigma1(uint E) { return S3(E); }
inline uint sigma1(uint X) { return rotate(X, 15U) ^ rotate(X, 13U) ^ (X >> 10); }
inline uint Sigma0(uint A) { return rotate(A, 30U) ^ rotate(A, 19U) ^ rotate(A, 10U); }
inline uint sigma0(uint X) { return rotate(X, 25U) ^ rotate(X, 14U) ^ (X >> 3); }

inline uint Smoosh2(uint X) {
    X ^= X >> 16;
    X ^= X >> 8;
    X ^= X >> 4;
    return (X ^ X >> 2) & 3;
}

#ifdef BITALIGN
#pragma OPENCL EXTENSION cl_amd_media_ops2 : enable
#endif

inline void Mangle(uint S[4])
{
    uint xr = (S[0] ^ (S[0] >> 4)) & 0x0f0f0f0f;

    S[1] ^= rotate(S[0], 32 - (xr >> 24));

    xr += 0x01010101;
    
    uint tmp = Smoosh2(S[1]);
#ifdef BITALIGN
    uint RR = ~rotate(S[0], 32 - amd_bfe(xr, 24 - (tmp << 3), 8));
#else
    uint RR = ~rotate(S[0], 32 - (xr >> (24 - (tmp << 3))));
#endif
    int mask = -((tmp ^ tmp >> 1) & 1);
    S[2] = select(bitselect(S[2] ^ ~RR, S[2] & RR, (uint)mask), S[2] + RR, tmp == 1);
    uint t = S[1] ^ S[2];

    xr += 0x01010101;
    tmp = Smoosh2(t);
#ifdef BITALIGN
    RR = ~rotate(S[0], 32 - amd_bfe(xr, 24 - (tmp << 3), 8));
#else
    RR = ~rotate(S[0], 32 - (xr >> (24 - (tmp << 3))));
#endif
    mask = -((tmp ^ tmp >> 1) & 1);
    S[3] = select(bitselect(S[3] ^ ~RR, S[3] & RR, (uint)mask), S[3] + RR, tmp == 1);
    
    S[0] ^= (t + S[3]);
}

#define Absorb(S, X) S[0] ^= (X); Mangle(S);

inline uint Squeeze(uint S[4])
{
    uint Y = S[0];
    Mangle(S);
    return Y;
}

inline uint Br(uint sponge[4], uint X)
{
    uint R = Squeeze(sponge);
    uint r = R & 3;
#ifdef BITALIGN
    uint r0 = amd_bfe(R, 8, 3);
    uint Xf = amd_bfe(X, r0, 1);
#else
    uint r0 = (R >> 8) & 31;
    uint Xf = (X >> r0) & 1;
#endif
    uint mask = 8 | Xf << 1 | (Xf ^ 1) << 2;
#ifdef BITALIGN
    return X ^ (amd_bfe(mask, r, 1) << r0);
#else
    return X ^ (((mask >> r) & 1) << r0);
#endif
}

inline void RoundFunc(uint sponge[4], uint v[8], const uint W, const uint K) {
        uint brG, tmp1, tmp2, brC, brB, tmp3, tmp4;
        brG = Br(sponge, v[6]);
        tmp1 = F1(v[4], Br(sponge, v[5]), brG) + v[7] + W + K;
        tmp2 = tmp1 + Sigma1(Br(sponge, v[4]));
        brC = Br(sponge, v[2]);
        brB = Br(sponge, v[1]);
        tmp3 = Ma(Br(sponge, v[0]), brB, brC);
        tmp4 = tmp3 + Sigma0(Br(sponge, v[0]));
        for (int i = 7; i > 0; i--)
            v[i] = v[i - 1];
        v[4] += Br(sponge, tmp2);
        v[0] = tmp2 + tmp4;
}

#define HR(i) (W[i & 0xf] = sigma1(W[(i + 14) & 0xf]) + W[(i + 9) & 0xf] + sigma0(W[(i + 1) & 0xf]) + W[i & 0xf])

#define HHALF1(i, k) Absorb(sponge, v[3] ^ v[7]); RoundFunc(sponge, v, W[i], k);
#define HHALF2(i, w, k) Absorb(sponge, v[3] + v[7]); RoundFunc(sponge, v, w, k);

inline void hefty84(uint* restrict hashx, uint sponge[4], __global const uint *restrict h)
{
    uint W[16];
    W[0] = hashx[16];
    W[1] = hashx[17];
    W[2] = hashx[18];
    W[3] = hashx[19];
    W[4] = hashx[20];;
    W[5] = 0x80000000U;
    W[6] = W[7] = W[8] = W[9] = W[10] = W[11] = W[12] = W[13] = W[14] = 0;
    W[15] = 0x2a0;
    uint v[8];
    for (int i = 0; i < 8; i++)
        v[i] = h[i];

    for (int i = 3; i < 16; i++) {
        Absorb(sponge, W[i] ^ K[i]);
    }
    for (int i = 0; i < 16; i++) {
        HHALF1(i, K[i]);
    }
    for (int i = 16; i < 64; i++) {
        HHALF2(i, HR(i), K[i]);
    }
    for (int i = 0; i < 8; i++)
        hashx[21 + i] = h[i] + v[i];
}

__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void search(__global const uint*restrict data,
                                    const uint sha_mask, const uint keccak_mask,
                                    const uint groestl_mask, const uint blake_mask,
                                    __global volatile uint*restrict output)
{
    __local ulong groestl_tables[2048];
    uint nonce = get_global_id(0);
    uint hashx[21 + 8];
    uint sponge[4];
    for (int i = 0; i < 21; i++)
        hashx[i] = data[i];
    for (int i = 21; i < 25; i++)
        sponge[i - 21] = data[i];
#ifndef NOLOCAL
    hashx[19] = SWAP32(nonce);
#endif

    hefty84(hashx, sponge, data + 25);

    if ((sha256_116_last(hashx) & sha_mask) != 0)
        return;
    if ((keccak512_116_last(hashx) & keccak_mask) != 0)
        return;
    if ((blake512_116_last(hashx) & blake_mask) != 0)
        return;
    if ((groestl512_116_last(hashx, groestl_tables) & groestl_mask) != 0)
        return;
#define FOUND (0xF)
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
    SETFOUND(nonce);
}