From 62fd64df9373f63daedd4a4e382b73bffcd5466a Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Thu, 24 Mar 2016 10:19:31 +0100 Subject: [PATCH] blake kernel optimisation for nvidia cards part of my tasks in the RFP-5 Signed-off-by: Tanguy Pruvot --- blake256.cl | 1768 +++------------------------------------------------ 1 file changed, 91 insertions(+), 1677 deletions(-) diff --git a/blake256.cl b/blake256.cl index f871618ea5..0cdf30b247 100644 --- a/blake256.cl +++ b/blake256.cl @@ -1,28 +1,30 @@ -/* BLAKE256 14-round kernel */ -/* Copyright 2015 Company Zero */ -/* A complete kernel re-write */ -/* with inspiration from the */ -/* Golang BLAKE256 repo over at */ -/* github.com/dchest/blake256 */ - -#define SPH_ROTR32(v,n) rotate((uint)(v),(uint)(32-(n))) - -__constant uint cst0 = 0x243F6A88UL; -__constant uint cst1 = 0x85A308D3UL; -__constant uint cst2 = 0x13198A2EUL; -__constant uint cst3 = 0x03707344UL; -__constant uint cst4 = 0xA4093822UL; -__constant uint cst5 = 0x299F31D0UL; -__constant uint cst6 = 0x082EFA98UL; -__constant uint cst7 = 0xEC4E6C89UL; -__constant uint cst8 = 0x452821E6UL; -__constant uint cst9 = 0x38D01377UL; -__constant uint cstA = 0xBE5466CFUL; -__constant uint cstB = 0x34E90C6CUL; -__constant uint cstC = 0xC0AC29B7UL; -__constant uint cstD = 0xC97C50DDUL; -__constant uint cstE = 0x3F84D5B5UL; -__constant uint cstF = 0xB5470917UL; +/** + * BLAKE256 14-round kernel + * + * Copyright 2015 Company Zero + * A complete kernel re-write + * with inspiration from the Golang BLAKE256 repo (github.com/dchest/blake256) + */ + +/** + * optimized by tpruvot 02/2016 : + * + * GTX 960 | (5s):735.3M (avg):789.3Mh/s + * GTX 750 | (5s):443.3M (avg):476.8Mh/s + * to + * GTX 960 | (5s):875.0M (avg):899.2Mh/s + * GTX 750 | (5s):523.1M (avg):536.8Mh/s + */ +#define ROTR(v,n) rotate(v,(uint)(32U-n)) +#define ROTL(v,n) rotate(v, n) + +#ifdef _AMD_OPENCL +#define SWAP(v) rotate(v, 16U) +#define ROTR8(v) rotate(v, 24U) +#else +#define SWAP(v) as_uint(as_uchar4(v).zwxy) +#define ROTR8(v) as_uint(as_uchar4(v).yzwx) +#endif __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __kernel void search( @@ -37,30 +39,45 @@ __kernel void search( const uint h6, const uint h7, - // last 52 bytes of original message - const uint in32, // M[0] - const uint in33, // M[1] - const uint in34, // M[2] - // const uint in35, = nonce M[3] - - const uint in36, // M[4] - const uint in37, // M[5] - const uint in38, // M[6] - const uint in39, // M[7] - - const uint in40, // M[8] - const uint in41, // M[9] - const uint in42, // M[10] - const uint in43, // M[11] - - const uint in44 // M[12] - // in45 = padding M[13] - // in46 = padding M[14] - // in47 = padding M[15] + // last 52 bytes of data + const uint M0, + const uint M1, + const uint M2, + // const uint M3 : nonce + const uint M4, + const uint M5, + const uint M6, + const uint M7, + const uint M8, + const uint M9, + const uint MA, + const uint MB, + const uint MC ) { - uint M0, M1, M2, M3, M4, M5, M6, M7; - uint M8, M9, MA, MB, MC, MD, ME, MF; + /* Load the block header and padding */ + const uint M3 = get_global_id(0); + const uint MD = 0x80000001UL; + const uint ME = 0x00000000UL; + const uint MF = 0x000005a0UL; + + const uint cst0 = 0x243F6A88UL; + const uint cst1 = 0x85A308D3UL; + const uint cst2 = 0x13198A2EUL; + const uint cst3 = 0x03707344UL; + const uint cst4 = 0xA4093822UL; + const uint cst5 = 0x299F31D0UL; + const uint cst6 = 0x082EFA98UL; + const uint cst7 = 0xEC4E6C89UL; + const uint cst8 = 0x452821E6UL; + const uint cst9 = 0x38D01377UL; + const uint cstA = 0xBE5466CFUL; + const uint cstB = 0x34E90C6CUL; + const uint cstC = 0xC0AC29B7UL; + const uint cstD = 0xC97C50DDUL; + const uint cstE = 0x3F84D5B5UL; + const uint cstF = 0xB5470917UL; + uint V0, V1, V2, V3, V4, V5, V6, V7; uint V8, V9, VA, VB, VC, VD, VE, VF; uint pre7; @@ -74,6 +91,7 @@ __kernel void search( V5 = h5; V6 = h6; pre7 = V7 = h7; + V8 = cst0; V9 = cst1; VA = cst2; @@ -83,1626 +101,22 @@ __kernel void search( VE = cst6; VF = cst7; - uint nonce = get_global_id(0); - - /* Load the block header and padding */ - M0 = in32; - M1 = in33; - M2 = in34; - M3 = nonce; - M4 = in36; - M5 = in37; - M6 = in38; - M7 = in39; - M8 = in40; - M9 = in41; - MA = in42; - MB = in43; - MC = in44; - MD = 0x80000001UL; - ME = 0x00000000UL; - MF = 0x000005a0UL; - - /* Begin the doing the 64-byte block. - * This can probably be optimized to - * get another 10-15% performance out. - */ - - /* Round 1. */ - V0 = V0 + (M0 ^ cst1); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M2 ^ cst3); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M4 ^ cst5); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (M6 ^ cst7); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (M5 ^ cst4); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (M7 ^ cst6); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M3 ^ cst2); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M1 ^ cst0); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M8 ^ cst9); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (MA ^ cstB); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (MC ^ cstD); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (ME ^ cstF); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (MD ^ cstC); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (MF ^ cstE); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (MB ^ cstA); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (M9 ^ cst8); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 2. */ - V0 = V0 + (ME ^ cstA); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M4 ^ cst8); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M9 ^ cstF); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (MD ^ cst6); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (MF ^ cst9); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (M6 ^ cstD); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M8 ^ cst4); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (MA ^ cstE); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M1 ^ cstC); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M0 ^ cst2); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (MB ^ cst7); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (M5 ^ cst3); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M7 ^ cstB); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M3 ^ cst5); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (M2 ^ cst0); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (MC ^ cst1); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 3. */ - V0 = V0 + (MB ^ cst8); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (MC ^ cst0); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M5 ^ cst2); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (MF ^ cstD); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (M2 ^ cst5); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (MD ^ cstF); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M0 ^ cstC); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M8 ^ cstB); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (MA ^ cstE); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M3 ^ cst6); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M7 ^ cst1); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (M9 ^ cst4); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M1 ^ cst7); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M4 ^ cst9); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (M6 ^ cst3); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (ME ^ cstA); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 4. */ - V0 = V0 + (M7 ^ cst9); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M3 ^ cst1); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (MD ^ cstC); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (MB ^ cstE); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (MC ^ cstD); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (ME ^ cstB); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M1 ^ cst3); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M9 ^ cst7); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M2 ^ cst6); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M5 ^ cstA); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M4 ^ cst0); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (MF ^ cst8); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M0 ^ cst4); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M8 ^ cstF); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (MA ^ cst5); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (M6 ^ cst2); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 5. */ - V0 = V0 + (M9 ^ cst0); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M5 ^ cst7); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M2 ^ cst4); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (MA ^ cstF); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (M4 ^ cst2); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (MF ^ cstA); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M7 ^ cst5); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M0 ^ cst9); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (ME ^ cst1); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (MB ^ cstC); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M6 ^ cst8); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (M3 ^ cstD); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M8 ^ cst6); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (MD ^ cst3); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (MC ^ cstB); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (M1 ^ cstE); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 6. */ - V0 = V0 + (M2 ^ cstC); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M6 ^ cstA); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M0 ^ cstB); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (M8 ^ cst3); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (MB ^ cst0); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (M3 ^ cst8); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (MA ^ cst6); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (MC ^ cst2); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M4 ^ cstD); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M7 ^ cst5); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (MF ^ cstE); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (M1 ^ cst9); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (ME ^ cstF); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M9 ^ cst1); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (M5 ^ cst7); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (MD ^ cst4); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 7. */ - V0 = V0 + (MC ^ cst5); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M1 ^ cstF); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (ME ^ cstD); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (M4 ^ cstA); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (MD ^ cstE); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (MA ^ cst4); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (MF ^ cst1); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M5 ^ cstC); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M0 ^ cst7); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M6 ^ cst3); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M9 ^ cst2); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (M8 ^ cstB); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M2 ^ cst9); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (MB ^ cst8); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (M3 ^ cst6); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (M7 ^ cst0); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 8. */ - V0 = V0 + (MD ^ cstB); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M7 ^ cstE); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (MC ^ cst1); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (M3 ^ cst9); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (M1 ^ cstC); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (M9 ^ cst3); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (ME ^ cst7); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (MB ^ cstD); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M5 ^ cst0); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (MF ^ cst4); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M8 ^ cst6); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (M2 ^ cstA); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M6 ^ cst8); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (MA ^ cst2); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (M4 ^ cstF); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (M0 ^ cst5); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 9. */ - V0 = V0 + (M6 ^ cstF); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (ME ^ cst9); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (MB ^ cst3); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (M0 ^ cst8); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (M3 ^ cstB); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (M8 ^ cst0); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M9 ^ cstE); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (MF ^ cst6); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (MC ^ cst2); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (MD ^ cst7); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M1 ^ cst4); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (MA ^ cst5); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M4 ^ cst1); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M5 ^ cstA); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (M7 ^ cstD); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (M2 ^ cstC); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 10. */ - V0 = V0 + (MA ^ cst2); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M8 ^ cst4); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M7 ^ cst6); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (M1 ^ cst5); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (M6 ^ cst7); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (M5 ^ cst1); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M4 ^ cst8); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M2 ^ cstA); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (MF ^ cstB); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M9 ^ cstE); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M3 ^ cstC); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (MD ^ cst0); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (MC ^ cst3); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M0 ^ cstD); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (ME ^ cst9); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (MB ^ cstF); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 11. */ - V0 = V0 + (M0 ^ cst1); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M2 ^ cst3); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M4 ^ cst5); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (M6 ^ cst7); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (M5 ^ cst4); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (M7 ^ cst6); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M3 ^ cst2); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M1 ^ cst0); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M8 ^ cst9); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (MA ^ cstB); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (MC ^ cstD); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (ME ^ cstF); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (MD ^ cstC); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (MF ^ cstE); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (MB ^ cstA); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (M9 ^ cst8); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 12. */ - V0 = V0 + (ME ^ cstA); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M4 ^ cst8); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M9 ^ cstF); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (MD ^ cst6); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (MF ^ cst9); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (M6 ^ cstD); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M8 ^ cst4); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (MA ^ cstE); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M1 ^ cstC); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M0 ^ cst2); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (MB ^ cst7); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (M5 ^ cst3); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M7 ^ cstB); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M3 ^ cst5); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (M2 ^ cst0); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (MC ^ cst1); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 13. */ - V0 = V0 + (MB ^ cst8); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (MC ^ cst0); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (M5 ^ cst2); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (MF ^ cstD); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (M2 ^ cst5); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (MD ^ cstF); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M0 ^ cstC); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M8 ^ cstB); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (MA ^ cstE); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M3 ^ cst6); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M7 ^ cst1); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (M9 ^ cst4); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M1 ^ cst7); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M4 ^ cst9); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (M6 ^ cst3); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (ME ^ cstA); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); - - /* Round 14. */ - V0 = V0 + (M7 ^ cst9); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 16); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 12); - V1 = V1 + (M3 ^ cst1); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 16); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 12); - V2 = V2 + (MD ^ cstC); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 16); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 12); - V3 = V3 + (MB ^ cstE); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 16); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 12); - V2 = V2 + (MC ^ cstD); - V2 = V2 + V6; - VE = VE ^ V2; - VE = SPH_ROTR32(VE, 8); - VA = VA + VE; - V6 = V6 ^ VA; - V6 = SPH_ROTR32(V6, 7); - V3 = V3 + (ME ^ cstB); - V3 = V3 + V7; - VF = VF ^ V3; - VF = SPH_ROTR32(VF, 8); - VB = VB + VF; - V7 = V7 ^ VB; - V7 = SPH_ROTR32(V7, 7); - V1 = V1 + (M1 ^ cst3); - V1 = V1 + V5; - VD = VD ^ V1; - VD = SPH_ROTR32(VD, 8); - V9 = V9 + VD; - V5 = V5 ^ V9; - V5 = SPH_ROTR32(V5, 7); - V0 = V0 + (M9 ^ cst7); - V0 = V0 + V4; - VC = VC ^ V0; - VC = SPH_ROTR32(VC, 8); - V8 = V8 + VC; - V4 = V4 ^ V8; - V4 = SPH_ROTR32(V4, 7); - V0 = V0 + (M2 ^ cst6); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 16); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 12); - V1 = V1 + (M5 ^ cstA); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 16); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 12); - V2 = V2 + (M4 ^ cst0); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 16); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 12); - V3 = V3 + (MF ^ cst8); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 16); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 12); - V2 = V2 + (M0 ^ cst4); - V2 = V2 + V7; - VD = VD ^ V2; - VD = SPH_ROTR32(VD, 8); - V8 = V8 + VD; - V7 = V7 ^ V8; - V7 = SPH_ROTR32(V7, 7); - V3 = V3 + (M8 ^ cstF); - V3 = V3 + V4; - VE = VE ^ V3; - VE = SPH_ROTR32(VE, 8); - V9 = V9 + VE; - V4 = V4 ^ V9; - V4 = SPH_ROTR32(V4, 7); - V1 = V1 + (MA ^ cst5); - V1 = V1 + V6; - VC = VC ^ V1; - VC = SPH_ROTR32(VC, 8); - VB = VB + VC; - V6 = V6 ^ VB; - V6 = SPH_ROTR32(V6, 7); - V0 = V0 + (M6 ^ cst2); - V0 = V0 + V5; - VF = VF ^ V0; - VF = SPH_ROTR32(VF, 8); - VA = VA + VF; - V5 = V5 ^ VA; - V5 = SPH_ROTR32(V5, 7); + /* 14 rounds */ + + V0 = V0 + (M0 ^ cst1); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M2 ^ cst3); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M4 ^ cst5); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (M6 ^ cst7); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (M5 ^ cst4); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (M7 ^ cst6); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M3 ^ cst2); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M1 ^ cst0); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M8 ^ cst9); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (MA ^ cstB); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (MC ^ cstD); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (ME ^ cstF); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (MD ^ cstC); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (MF ^ cstE); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (MB ^ cstA); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (M9 ^ cst8); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (ME ^ cstA); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M4 ^ cst8); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M9 ^ cstF); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (MD ^ cst6); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (MF ^ cst9); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (M6 ^ cstD); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M8 ^ cst4); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (MA ^ cstE); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M1 ^ cstC); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M0 ^ cst2); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (MB ^ cst7); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (M5 ^ cst3); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M7 ^ cstB); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M3 ^ cst5); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (M2 ^ cst0); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (MC ^ cst1); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (MB ^ cst8); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (MC ^ cst0); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M5 ^ cst2); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (MF ^ cstD); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (M2 ^ cst5); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (MD ^ cstF); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M0 ^ cstC); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M8 ^ cstB); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (MA ^ cstE); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M3 ^ cst6); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M7 ^ cst1); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (M9 ^ cst4); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M1 ^ cst7); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M4 ^ cst9); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (M6 ^ cst3); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (ME ^ cstA); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (M7 ^ cst9); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M3 ^ cst1); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (MD ^ cstC); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (MB ^ cstE); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (MC ^ cstD); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (ME ^ cstB); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M1 ^ cst3); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M9 ^ cst7); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M2 ^ cst6); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M5 ^ cstA); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M4 ^ cst0); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (MF ^ cst8); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M0 ^ cst4); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M8 ^ cstF); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (MA ^ cst5); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (M6 ^ cst2); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (M9 ^ cst0); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M5 ^ cst7); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M2 ^ cst4); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (MA ^ cstF); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (M4 ^ cst2); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (MF ^ cstA); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M7 ^ cst5); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M0 ^ cst9); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (ME ^ cst1); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (MB ^ cstC); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M6 ^ cst8); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (M3 ^ cstD); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M8 ^ cst6); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (MD ^ cst3); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (MC ^ cstB); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (M1 ^ cstE); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (M2 ^ cstC); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M6 ^ cstA); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M0 ^ cstB); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (M8 ^ cst3); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (MB ^ cst0); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (M3 ^ cst8); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (MA ^ cst6); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (MC ^ cst2); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M4 ^ cstD); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M7 ^ cst5); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (MF ^ cstE); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (M1 ^ cst9); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (ME ^ cstF); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M9 ^ cst1); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (M5 ^ cst7); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (MD ^ cst4); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (MC ^ cst5); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M1 ^ cstF); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (ME ^ cstD); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (M4 ^ cstA); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (MD ^ cstE); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (MA ^ cst4); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (MF ^ cst1); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M5 ^ cstC); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M0 ^ cst7); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M6 ^ cst3); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M9 ^ cst2); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (M8 ^ cstB); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M2 ^ cst9); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (MB ^ cst8); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (M3 ^ cst6); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (M7 ^ cst0); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (MD ^ cstB); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M7 ^ cstE); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (MC ^ cst1); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (M3 ^ cst9); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (M1 ^ cstC); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (M9 ^ cst3); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (ME ^ cst7); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (MB ^ cstD); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M5 ^ cst0); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (MF ^ cst4); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M8 ^ cst6); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (M2 ^ cstA); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M6 ^ cst8); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (MA ^ cst2); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (M4 ^ cstF); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (M0 ^ cst5); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (M6 ^ cstF); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (ME ^ cst9); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (MB ^ cst3); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (M0 ^ cst8); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (M3 ^ cstB); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (M8 ^ cst0); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M9 ^ cstE); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (MF ^ cst6); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (MC ^ cst2); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (MD ^ cst7); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M1 ^ cst4); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (MA ^ cst5); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M4 ^ cst1); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M5 ^ cstA); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (M7 ^ cstD); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (M2 ^ cstC); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (MA ^ cst2); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M8 ^ cst4); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M7 ^ cst6); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (M1 ^ cst5); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (M6 ^ cst7); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (M5 ^ cst1); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M4 ^ cst8); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M2 ^ cstA); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (MF ^ cstB); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M9 ^ cstE); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M3 ^ cstC); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (MD ^ cst0); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (MC ^ cst3); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M0 ^ cstD); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (ME ^ cst9); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (MB ^ cstF); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (M0 ^ cst1); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M2 ^ cst3); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M4 ^ cst5); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (M6 ^ cst7); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (M5 ^ cst4); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (M7 ^ cst6); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M3 ^ cst2); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M1 ^ cst0); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M8 ^ cst9); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (MA ^ cstB); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (MC ^ cstD); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (ME ^ cstF); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (MD ^ cstC); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (MF ^ cstE); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (MB ^ cstA); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (M9 ^ cst8); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (ME ^ cstA); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M4 ^ cst8); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M9 ^ cstF); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (MD ^ cst6); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (MF ^ cst9); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (M6 ^ cstD); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M8 ^ cst4); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (MA ^ cstE); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M1 ^ cstC); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M0 ^ cst2); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (MB ^ cst7); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (M5 ^ cst3); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M7 ^ cstB); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M3 ^ cst5); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (M2 ^ cst0); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (MC ^ cst1); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (MB ^ cst8); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (MC ^ cst0); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (M5 ^ cst2); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (MF ^ cstD); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (M2 ^ cst5); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (MD ^ cstF); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M0 ^ cstC); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M8 ^ cstB); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (MA ^ cstE); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M3 ^ cst6); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M7 ^ cst1); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (M9 ^ cst4); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M1 ^ cst7); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M4 ^ cst9); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (M6 ^ cst3); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (ME ^ cstA); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U); + V0 = V0 + (M7 ^ cst9); V0 = V0 + V4; VC = VC ^ V0; VC = SWAP(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 12U); V1 = V1 + (M3 ^ cst1); V1 = V1 + V5; VD = VD ^ V1; VD = SWAP(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 12U); V2 = V2 + (MD ^ cstC); V2 = V2 + V6; VE = VE ^ V2; VE = SWAP(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 12U); V3 = V3 + (MB ^ cstE); V3 = V3 + V7; VF = VF ^ V3; VF = SWAP(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 12U); V2 = V2 + (MC ^ cstD); V2 = V2 + V6; VE = VE ^ V2; VE = ROTR8(VE); VA = VA + VE; V6 = V6 ^ VA; V6 = ROTR(V6, 7U); V3 = V3 + (ME ^ cstB); V3 = V3 + V7; VF = VF ^ V3; VF = ROTR8(VF); VB = VB + VF; V7 = V7 ^ VB; V7 = ROTR(V7, 7U); V1 = V1 + (M1 ^ cst3); V1 = V1 + V5; VD = VD ^ V1; VD = ROTR8(VD); V9 = V9 + VD; V5 = V5 ^ V9; V5 = ROTR(V5, 7U); V0 = V0 + (M9 ^ cst7); V0 = V0 + V4; VC = VC ^ V0; VC = ROTR8(VC); V8 = V8 + VC; V4 = V4 ^ V8; V4 = ROTR(V4, 7U); V0 = V0 + (M2 ^ cst6); V0 = V0 + V5; VF = VF ^ V0; VF = SWAP(VF); VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 12U); V1 = V1 + (M5 ^ cstA); V1 = V1 + V6; VC = VC ^ V1; VC = SWAP(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 12U); V2 = V2 + (M4 ^ cst0); V2 = V2 + V7; VD = VD ^ V2; VD = SWAP(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 12U); V3 = V3 + (MF ^ cst8); V3 = V3 + V4; VE = VE ^ V3; VE = SWAP(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 12U); V2 = V2 + (M0 ^ cst4); V2 = V2 + V7; VD = VD ^ V2; VD = ROTR8(VD); V8 = V8 + VD; V7 = V7 ^ V8; V7 = ROTR(V7, 7U); V3 = V3 + (M8 ^ cstF); V3 = V3 + V4; VE = VE ^ V3; VE = ROTR8(VE); V9 = V9 + VE; V4 = V4 ^ V9; V4 = ROTR(V4, 7U); V1 = V1 + (MA ^ cst5); V1 = V1 + V6; VC = VC ^ V1; VC = ROTR8(VC); VB = VB + VC; V6 = V6 ^ VB; V6 = ROTR(V6, 7U); V0 = V0 + (M6 ^ cst2); V0 = V0 + V5; VF = VF ^ V0; VF = ROTR8(VF);/*VA = VA + VF; V5 = V5 ^ VA; V5 = ROTR(V5, 7U);*/ /* The final chunks of the hash * are calculated as: @@ -1714,16 +128,16 @@ __kernel void search( * h5 = h5 ^ V5 ^ VD; * h6 = h6 ^ V6 ^ VE; * h7 = h7 ^ V7 ^ VF; - * - * We just check if the last byte - * is zeroed and if it is, we tell - * cgminer that we've found a - * and to check it against the + * + * We just check if the last byte + * is zeroed and if it is, we tell + * cgminer that we've found a + * and to check it against the * target. */ - - /* Debug code to help you assess the correctness - * of your hashing function in case someone decides + + /* Debug code to help you assess the correctness + * of your hashing function in case someone decides * to try to optimize. if (!((pre7 ^ V7 ^ VF) & 0xFFFF0000)) { printf("hash on gpu %x %x %x %x %x %x %x %x\n", @@ -1739,9 +153,9 @@ __kernel void search( nonce); } */ - + if (pre7 ^ V7 ^ VF) return; /* Push this share */ - output[output[0xFF]++] = nonce; + output[output[0xFF]++] = M3; }