GOSTSec
/
sgminer
mirror of https://github.com/GOSTSec/sgminer
1
0
Fork 0
Code Issues Releases Wiki Activity
Browse Source

Added support for MyriadCoin groestl algorithm.

djm34
phm 11 years ago
parent
3370e0846f
commit
8fe1b860c5
12 changed files with 1517 additions and 1 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 1
      Makefile.am
  2. 1
      configure.ac
  3. 6
      driver-opencl.c
  4. 396
      kernel/myriadcoin-groestl.cl
  5. 1
      miner.h
  6. 153
      myriadcoin-groestl.c
  7. 10
      myriadcoin-groestl.h
  8. 5
      ocl.c
  9. 6
      sgminer.c
  10. 2
      sph/Makefile.am
  11. 690
      sph/sha2.c
  12. 247
      sph/sha2big.c

1
Makefile.am
Unescape View File

@ -45,6 +45,7 @@ sgminer_SOURCES += scrypt.c scrypt.h @@ -45,6 +45,7 @@ sgminer_SOURCES += scrypt.c scrypt.h
sgminer_SOURCES += darkcoin.c darkcoin.h
sgminer_SOURCES += qubitcoin.c qubitcoin.h
sgminer_SOURCES += quarkcoin.c quarkcoin.h
sgminer_SOURCES += myriadcoin-groestl.c myriadcoin-groestl.h
sgminer_SOURCES += kernel/*.cl
bin_SCRIPTS = $(top_srcdir)/kernel/*.cl

1
configure.ac
Unescape View File

@ -348,6 +348,7 @@ AC_DEFINE_UNQUOTED([PSW_KERNNAME], ["psw"], [Filename for psw's experimental ker @@ -348,6 +348,7 @@ AC_DEFINE_UNQUOTED([PSW_KERNNAME], ["psw"], [Filename for psw's experimental ker
AC_DEFINE_UNQUOTED([DARKCOIN_KERNNAME], ["darkcoin"], [Filename for DarkCoin optimised kernel])
AC_DEFINE_UNQUOTED([QUBITCOIN_KERNNAME], ["qubitcoin"], [Filename for QubitCoin optimised kernel])
AC_DEFINE_UNQUOTED([QUARKCOIN_KERNNAME], ["quarkcoin"], [Filename for QuarkCoin optimised kernel])
AC_DEFINE_UNQUOTED([MYRIADCOIN_GROESTL_KERNNAME], ["myriadcoin-groestl"], [Filename for MyriadCoin-Groestl optimised kernel])
AC_SUBST(OPENCL_LIBS)
AC_SUBST(OPENCL_FLAGS)

6
driver-opencl.c
Unescape View File

@ -213,6 +213,8 @@ static enum cl_kernels select_kernel(char *arg) @@ -213,6 +213,8 @@ static enum cl_kernels select_kernel(char *arg)
return KL_QUBITCOIN;
if (!strcmp(arg, QUARKCOIN_KERNNAME))
return KL_QUARKCOIN;
if (!strcmp(arg, MYRIADCOIN_GROESTL_KERNNAME))
return KL_MYRIADCOIN_GROESTL;
return KL_NONE;
}
@ -1362,6 +1364,9 @@ static bool opencl_thread_prepare(struct thr_info *thr) @@ -1362,6 +1364,9 @@ static bool opencl_thread_prepare(struct thr_info *thr)
case KL_QUARKCOIN:
cgpu->kname = QUARKCOIN_KERNNAME;
break;
case KL_MYRIADCOIN_GROESTL:
cgpu->kname = MYRIADCOIN_GROESTL_KERNNAME;
break;
default:
break;
}
@ -1400,6 +1405,7 @@ static bool opencl_thread_init(struct thr_info *thr) @@ -1400,6 +1405,7 @@ static bool opencl_thread_init(struct thr_info *thr)
case KL_DARKCOIN:
case KL_QUBITCOIN:
case KL_QUARKCOIN:
case KL_MYRIADCOIN_GROESTL:
thrdata->queue_kernel_parameters = &queue_sph_kernel;
break;
default:

396
kernel/myriadcoin-groestl.cl
Unescape View File

@ -0,0 +1,396 @@ @@ -0,0 +1,396 @@
/*
* MyriadCoin Groestl kernel implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2014 phm
*
* Permission is hereby granted, free of charge, to any person obtaining
* a copy of this software and associated documentation files (the
* "Software"), to deal in the Software without restriction, including
* without limitation the rights to use, copy, modify, merge, publish,
* distribute, sublicense, and/or sell copies of the Software, and to
* permit persons to whom the Software is furnished to do so, subject to
* the following conditions:
*
* The above copyright notice and this permission notice shall be
* included in all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
* IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
* CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
* TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
* SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
*
* ===========================(LICENSE END)=============================
*
* @author phm <phm@inbox.com>
*/
#ifndef MYRIADCOIN_GROESTL_CL
#define MYRIADCOIN_GROESTL_CL
#if __ENDIAN_LITTLE__
#define SPH_LITTLE_ENDIAN 1
#else
#define SPH_BIG_ENDIAN 1
#endif
#define SPH_UPTR sph_u64
typedef unsigned int sph_u32;
typedef int sph_s32;
#ifndef __OPENCL_VERSION__
typedef unsigned long long sph_u64;
typedef long long sph_s64;
#else
typedef unsigned long sph_u64;
typedef long sph_s64;
#endif
#define SPH_64 1
#define SPH_64_TRUE 1
#define SPH_C32(x) ((sph_u32)(x ## U))
#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF))
#define SPH_ROTL32(x, n) SPH_T32(((x) << (n)) | ((x) >> (32 - (n))))
#define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n)))
#define SPH_C64(x) ((sph_u64)(x ## UL))
#define SPH_T64(x) ((x) & SPH_C64(0xFFFFFFFFFFFFFFFF))
#define SPH_ROTL64(x, n) SPH_T64(((x) << (n)) | ((x) >> (64 - (n))))
#define SPH_ROTR64(x, n) SPH_ROTL64(x, (64 - (n)))
#define SPH_ECHO_64 1
#define SPH_SIMD_NOCOPY 0
#define SPH_LUFFA_PARALLEL 0
#define SPH_CUBEHASH_UNROLL 0
#include "groestl.cl"
#define SWAP4(x) as_uint(as_uchar4(x).wzyx)
#define SWAP8(x) as_ulong(as_uchar8(x).s76543210)
#if SPH_BIG_ENDIAN
#define ENC64E(x) SWAP8(x)
#define DEC64E(x) SWAP8(*(const __global sph_u64 *) (x));
#else
#define ENC64E(x) (x)
#define DEC64E(x) (*(const __global sph_u64 *) (x));
#endif
#define ROL32(x, n) rotate(x, (uint) n)
#define SHR(x, n) ((x) >> n)
#define SWAP32(a) (as_uint(as_uchar4(a).wzyx))
#define S0(x) (ROL32(x, 25) ^ ROL32(x, 14) ^ SHR(x, 3))
#define S1(x) (ROL32(x, 15) ^ ROL32(x, 13) ^ SHR(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 RD14 (S1(W12) + W7 + S0(W15) + W14)
#define RD15 (S1(W13) + W8 + S0(W0) + W15)
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void search(__global unsigned char* block, volatile __global uint* output, const ulong target)
{
uint gid = get_global_id(0);
union {
unsigned char h1[64];
uint h4[16];
ulong h8[8];
} hash;
__local sph_u64 T0_L[256], T1_L[256], T2_L[256], T3_L[256], T4_L[256], T5_L[256], T6_L[256], T7_L[256];
int init = get_local_id(0);
int step = get_local_size(0);
for (int i = init; i < 256; i += step)
{
T0_L[i] = T0[i];
T1_L[i] = T1[i];
T2_L[i] = T2[i];
T3_L[i] = T3[i];
T4_L[i] = T4[i];
T5_L[i] = T5[i];
T6_L[i] = T6[i];
T7_L[i] = T7[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
#define T0 T0_L
#define T1 T1_L
#define T2 T2_L
#define T3 T3_L
#define T4 T4_L
#define T5 T5_L
#define T6 T6_L
#define T7 T7_L
// groestl
sph_u64 H[16];
for (unsigned int u = 0; u < 15; u ++)
H[u] = 0;
#if USE_LE
H[15] = ((sph_u64)(512 & 0xFF) << 56) | ((sph_u64)(512 & 0xFF00) << 40);
#else
H[15] = (sph_u64)512;
#endif
sph_u64 g[16], m[16];
m[0] = DEC64E(block + 0 * 8);
m[1] = DEC64E(block + 1 * 8);
m[2] = DEC64E(block + 2 * 8);
m[3] = DEC64E(block + 3 * 8);
m[4] = DEC64E(block + 4 * 8);
m[5] = DEC64E(block + 5 * 8);
m[6] = DEC64E(block + 6 * 8);
m[7] = DEC64E(block + 7 * 8);
m[8] = DEC64E(block + 8 * 8);
m[9] = DEC64E(block + 9 * 8);
m[9] &= 0x00000000FFFFFFFF;
m[9] |= ((sph_u64) gid << 32);
m[10] = 0x80;
m[11] = 0;
m[12] = 0;
m[13] = 0;
m[14] = 0;
m[15] = 0x100000000000000;
for (unsigned int u = 0; u < 16; u ++)
g[u] = m[u] ^ H[u];
PERM_BIG_P(g);
PERM_BIG_Q(m);
for (unsigned int u = 0; u < 16; u ++)
H[u] ^= g[u] ^ m[u];
sph_u64 xH[16];
for (unsigned int u = 0; u < 16; u ++)
xH[u] = H[u];
PERM_BIG_P(xH);
for (unsigned int u = 0; u < 16; u ++)
H[u] ^= xH[u];
for (unsigned int u = 0; u < 8; u ++)
hash.h8[u] = ENC64E(H[u + 8]);
uint temp1;
uint W0 = SWAP32(hash.h4[0x0]);
uint W1 = SWAP32(hash.h4[0x1]);
uint W2 = SWAP32(hash.h4[0x2]);
uint W3 = SWAP32(hash.h4[0x3]);
uint W4 = SWAP32(hash.h4[0x4]);
uint W5 = SWAP32(hash.h4[0x5]);
uint W6 = SWAP32(hash.h4[0x6]);
uint W7 = SWAP32(hash.h4[0x7]);
uint W8 = SWAP32(hash.h4[0x8]);
uint W9 = SWAP32(hash.h4[0x9]);
uint W10 = SWAP32(hash.h4[0xA]);
uint W11 = SWAP32(hash.h4[0xB]);
uint W12 = SWAP32(hash.h4[0xC]);
uint W13 = SWAP32(hash.h4[0xD]);
uint W14 = SWAP32(hash.h4[0xE]);
uint W15 = SWAP32(hash.h4[0xF]);
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;
uint s0 = v0;
v1 += 0xBB67AE85;
uint s1 = v1;
v2 += 0x3C6EF372;
uint s2 = v2;
v3 += 0xA54FF53A;
uint s3 = v3;
v4 += 0x510E527F;
uint s4 = v4;
v5 += 0x9B05688C;
uint s5 = v5;
v6 += 0x1F83D9AB;
uint s6 = v6;
v7 += 0x5BE0CD19;
uint s7 = v7;
P( v0, v1, v2, v3, v4, v5, v6, v7, 0x80000000, 0x428A2F98 );
P( v7, v0, v1, v2, v3, v4, v5, v6, 0, 0x71374491 );
P( v6, v7, v0, v1, v2, v3, v4, v5, 0, 0xB5C0FBCF );
P( v5, v6, v7, v0, v1, v2, v3, v4, 0, 0xE9B5DBA5 );
P( v4, v5, v6, v7, v0, v1, v2, v3, 0, 0x3956C25B );
P( v3, v4, v5, v6, v7, v0, v1, v2, 0, 0x59F111F1 );
P( v2, v3, v4, v5, v6, v7, v0, v1, 0, 0x923F82A4 );
P( v1, v2, v3, v4, v5, v6, v7, v0, 0, 0xAB1C5ED5 );
P( v0, v1, v2, v3, v4, v5, v6, v7, 0, 0xD807AA98 );
P( v7, v0, v1, v2, v3, v4, v5, v6, 0, 0x12835B01 );
P( v6, v7, v0, v1, v2, v3, v4, v5, 0, 0x243185BE );
P( v5, v6, v7, v0, v1, v2, v3, v4, 0, 0x550C7DC3 );
P( v4, v5, v6, v7, v0, v1, v2, v3, 0, 0x72BE5D74 );
P( v3, v4, v5, v6, v7, v0, v1, v2, 0, 0x80DEB1FE );
P( v2, v3, v4, v5, v6, v7, v0, v1, 0, 0x9BDC06A7 );
P( v1, v2, v3, v4, v5, v6, v7, v0, 512, 0xC19BF174 );
P( v0, v1, v2, v3, v4, v5, v6, v7, 0x80000000U, 0xE49B69C1U );
P( v7, v0, v1, v2, v3, v4, v5, v6, 0x01400000U, 0xEFBE4786U );
P( v6, v7, v0, v1, v2, v3, v4, v5, 0x00205000U, 0x0FC19DC6U );
P( v5, v6, v7, v0, v1, v2, v3, v4, 0x00005088U, 0x240CA1CCU );
P( v4, v5, v6, v7, v0, v1, v2, v3, 0x22000800U, 0x2DE92C6FU );
P( v3, v4, v5, v6, v7, v0, v1, v2, 0x22550014U, 0x4A7484AAU );
P( v2, v3, v4, v5, v6, v7, v0, v1, 0x05089742U, 0x5CB0A9DCU );
P( v1, v2, v3, v4, v5, v6, v7, v0, 0xa0000020U, 0x76F988DAU );
P( v0, v1, v2, v3, v4, v5, v6, v7, 0x5a880000U, 0x983E5152U );
P( v7, v0, v1, v2, v3, v4, v5, v6, 0x005c9400U, 0xA831C66DU );
P( v6, v7, v0, v1, v2, v3, v4, v5, 0x0016d49dU, 0xB00327C8U );
P( v5, v6, v7, v0, v1, v2, v3, v4, 0xfa801f00U, 0xBF597FC7U );
P( v4, v5, v6, v7, v0, v1, v2, v3, 0xd33225d0U, 0xC6E00BF3U );
P( v3, v4, v5, v6, v7, v0, v1, v2, 0x11675959U, 0xD5A79147U );
P( v2, v3, v4, v5, v6, v7, v0, v1, 0xf6e6bfdaU, 0x06CA6351U );
P( v1, v2, v3, v4, v5, v6, v7, v0, 0xb30c1549U, 0x14292967U );
P( v0, v1, v2, v3, v4, v5, v6, v7, 0x08b2b050U, 0x27B70A85U );
P( v7, v0, v1, v2, v3, v4, v5, v6, 0x9d7c4c27U, 0x2E1B2138U );
P( v6, v7, v0, v1, v2, v3, v4, v5, 0x0ce2a393U, 0x4D2C6DFCU );
P( v5, v6, v7, v0, v1, v2, v3, v4, 0x88e6e1eaU, 0x53380D13U );
P( v4, v5, v6, v7, v0, v1, v2, v3, 0xa52b4335U, 0x650A7354U );
P( v3, v4, v5, v6, v7, v0, v1, v2, 0x67a16f49U, 0x766A0ABBU );
P( v2, v3, v4, v5, v6, v7, v0, v1, 0xd732016fU, 0x81C2C92EU );
P( v1, v2, v3, v4, v5, v6, v7, v0, 0x4eeb2e91U, 0x92722C85U );
P( v0, v1, v2, v3, v4, v5, v6, v7, 0x5dbf55e5U, 0xA2BFE8A1U );
P( v7, v0, v1, v2, v3, v4, v5, v6, 0x8eee2335U, 0xA81A664BU );
P( v6, v7, v0, v1, v2, v3, v4, v5, 0xe2bc5ec2U, 0xC24B8B70U );
P( v5, v6, v7, v0, v1, v2, v3, v4, 0xa83f4394U, 0xC76C51A3U );
P( v4, v5, v6, v7, v0, v1, v2, v3, 0x45ad78f7U, 0xD192E819U );
P( v3, v4, v5, v6, v7, v0, v1, v2, 0x36f3d0cdU, 0xD6990624U );
P( v2, v3, v4, v5, v6, v7, v0, v1, 0xd99c05e8U, 0xF40E3585U );
P( v1, v2, v3, v4, v5, v6, v7, v0, 0xb0511dc7U, 0x106AA070U );
P( v0, v1, v2, v3, v4, v5, v6, v7, 0x69bc7ac4U, 0x19A4C116U );
P( v7, v0, v1, v2, v3, v4, v5, v6, 0xbd11375bU, 0x1E376C08U );
P( v6, v7, v0, v1, v2, v3, v4, v5, 0xe3ba71e5U, 0x2748774CU );
P( v5, v6, v7, v0, v1, v2, v3, v4, 0x3b209ff2U, 0x34B0BCB5U );
P( v4, v5, v6, v7, v0, v1, v2, v3, 0x18feee17U, 0x391C0CB3U );
P( v3, v4, v5, v6, v7, v0, v1, v2, 0xe25ad9e7U, 0x4ED8AA4AU );
P( v2, v3, v4, v5, v6, v7, v0, v1, 0x13375046U, 0x5B9CCA4FU );
P( v1, v2, v3, v4, v5, v6, v7, v0, 0x0515089dU, 0x682E6FF3U );
P( v0, v1, v2, v3, v4, v5, v6, v7, 0x4f0d0f04U, 0x748F82EEU );
P( v7, v0, v1, v2, v3, v4, v5, v6, 0x2627484eU, 0x78A5636FU );
P( v6, v7, v0, v1, v2, v3, v4, v5, 0x310128d2U, 0x84C87814U );
P( v5, v6, v7, v0, v1, v2, v3, v4, 0xc668b434U, 0x8CC70208U );
PLAST( v4, v5, v6, v7, v0, v1, v2, v3, 0x420841ccU, 0x90BEFFFAU );
hash.h4[0] = SWAP4(v0 + s0);
hash.h4[1] = SWAP4(v1 + s1);
hash.h4[2] = SWAP4(v2 + s2);
hash.h4[3] = SWAP4(v3 + s3);
hash.h4[4] = SWAP4(v4 + s4);
hash.h4[5] = SWAP4(v5 + s5);
hash.h4[6] = SWAP4(v6 + s6);
hash.h4[7] = SWAP4(v7 + s7);
bool result = (hash.h8[3] <= target);
if (result)
output[output[0xFF]++] = SWAP4(gid);
}
#endif // MYRIADCOIN_GROESTL_CL

1
miner.h
Unescape View File

@ -384,6 +384,7 @@ enum cl_kernels { @@ -384,6 +384,7 @@ enum cl_kernels {
KL_QUARKCOIN, // kernels starting from this will have difficulty calculated by using quarkcoin algorithm
KL_QUBITCOIN,
KL_DARKCOIN, // kernels starting from this will have difficulty calculated by using bitcoin algorithm
KL_MYRIADCOIN_GROESTL,
};
enum dev_reason {

153
myriadcoin-groestl.c
Unescape View File

@ -0,0 +1,153 @@ @@ -0,0 +1,153 @@
/*-
* Copyright 2009 Colin Percival, 2014 phm
* All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions
* are met:
* 1. Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
* 2. Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
*
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR AND CONTRIBUTORS ``AS IS'' AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
* ARE DISCLAIMED. IN NO EVENT SHALL THE AUTHOR OR CONTRIBUTORS BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
* DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS
* OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)
* HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY
* OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF
* SUCH DAMAGE.
*
* This file was originally written by Colin Percival as part of the Tarsnap
* online backup system.
*/
#include "config.h"
#include "miner.h"
#include <stdlib.h>
#include <stdint.h>
#include <string.h>
#include "sph/sph_groestl.h"
#include "sph/sph_sha2.h"
/*
* Encode a length len/4 vector of (uint32_t) into a length len vector of
* (unsigned char) in big-endian form. Assumes len is a multiple of 4.
*/
static inline void
be32enc_vect(uint32_t *dst, const uint32_t *src, uint32_t len)
{
uint32_t i;
for (i = 0; i < len; i++)
dst[i] = htobe32(src[i]);
}
inline void mghash(void *state, const void *input)
{
sph_groestl512_context ctx_groestl;
sph_sha256_context ctx_sha2;
uint32_t hash[16];
sph_groestl512_init(&ctx_groestl);
sph_groestl512(&ctx_groestl, input, 80);
sph_groestl512_close(&ctx_groestl, (void*) hash);
sph_sha256_init(&ctx_sha2);
sph_sha256(&ctx_sha2, hash, 64);
sph_sha256_close(&ctx_sha2, (void*) hash);
memcpy(state, hash, 32);
}
static const uint32_t diff1targ = 0x0000ffff;
/* Used externally as confirmation of correct OCL code */
int myriadcoin_groestl_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce)
{
uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]);
uint32_t data[20], ohash[8];
//char *scratchbuf;
be32enc_vect(data, (const uint32_t *)pdata, 19);
data[19] = htobe32(nonce);
//scratchbuf = alloca(SCRATCHBUF_SIZE);
mghash(ohash, data);
tmp_hash7 = be32toh(ohash[7]);
applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx",
(long unsigned int)Htarg,
(long unsigned int)diff1targ,
(long unsigned int)tmp_hash7);
if (tmp_hash7 > diff1targ)
return -1;
if (tmp_hash7 > Htarg)
return 0;
return 1;
}
void myriadcoin_groestl_regenhash(struct work *work)
{
uint32_t data[20];
char *scratchbuf;
uint32_t *nonce = (uint32_t *)(work->data + 76);
uint32_t *ohash = (uint32_t *)(work->hash);
be32enc_vect(data, (const uint32_t *)work->data, 19);
data[19] = htobe32(*nonce);
mghash(ohash, data);
}
bool scanhash_myriadcoin_groestl(struct thr_info *thr, const unsigned char __maybe_unused *pmidstate,
unsigned char *pdata, unsigned char __maybe_unused *phash1,
unsigned char __maybe_unused *phash, const unsigned char *ptarget,
uint32_t max_nonce, uint32_t *last_nonce, uint32_t n)
{
uint32_t *nonce = (uint32_t *)(pdata + 76);
char *scratchbuf;
uint32_t data[20];
uint32_t tmp_hash7;
uint32_t Htarg = le32toh(((const uint32_t *)ptarget)[7]);
bool ret = false;
be32enc_vect(data, (const uint32_t *)pdata, 19);
while(1) {
uint32_t ostate[8];
*nonce = ++n;
data[19] = (n);
mghash(ostate, data);
tmp_hash7 = (ostate[7]);
applog(LOG_INFO, "data7 %08lx",
(long unsigned int)data[7]);
if (unlikely(tmp_hash7 <= Htarg)) {
((uint32_t *)pdata)[19] = htobe32(n);
*last_nonce = n;
ret = true;
break;
}
if (unlikely((n >= max_nonce) || thr->work_restart)) {
*last_nonce = n;
break;
}
}
return ret;
}

10
myriadcoin-groestl.h
Unescape View File

@ -0,0 +1,10 @@ @@ -0,0 +1,10 @@
#ifndef MYRIADCOIN_GROESTL_H
#define MYRIADCOIN_GROESTL_H
#include "miner.h"
extern int myriadcoin_groestl_test(unsigned char *pdata, const unsigned char *ptarget,
uint32_t nonce);
extern void myriadcoin_groestl_regenhash(struct work *work);
#endif /* MYRIADCOIN_GROESTL_H */

5
ocl.c
Unescape View File

@ -469,6 +469,11 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) @@ -469,6 +469,11 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
strcpy(filename, QUARKCOIN_KERNNAME".cl");
strcpy(binaryfilename, QUARKCOIN_KERNNAME);
break;
case KL_MYRIADCOIN_GROESTL:
applog(LOG_WARNING, "Kernel myriadcoin-groestl is experimental.");
strcpy(filename, MYRIADCOIN_GROESTL_KERNNAME".cl");
strcpy(binaryfilename, MYRIADCOIN_GROESTL_KERNNAME);
break;
case KL_NONE: /* Shouldn't happen */
break;
}

6
sgminer.c
Unescape View File

@ -4236,6 +4236,9 @@ void write_config(FILE *fcfg) @@ -4236,6 +4236,9 @@ void write_config(FILE *fcfg)
case KL_QUARKCOIN:
fprintf(fcfg, QUARKCOIN_KERNNAME);
break;
case KL_MYRIADCOIN_GROESTL:
fprintf(fcfg, MYRIADCOIN_GROESTL_KERNNAME);
break;
}
}
@ -6049,6 +6052,9 @@ static void rebuild_nonce(struct work *work, uint32_t nonce) @@ -6049,6 +6052,9 @@ static void rebuild_nonce(struct work *work, uint32_t nonce)
case KL_QUARKCOIN:
quarkcoin_regenhash(work);
break;
case KL_MYRIADCOIN_GROESTL:
myriadcoin_groestl_regenhash(work);
break;
default:
scrypt_regenhash(work);
break;

2
sph/Makefile.am
Unescape View File

@ -1,3 +1,3 @@ @@ -1,3 +1,3 @@
noinst_LIBRARIES = libsph.a
libsph_a_SOURCES = bmw.c echo.c jh.c luffa.c simd.c blake.c cubehash.c groestl.c keccak.c shavite.c skein.c
libsph_a_SOURCES = bmw.c echo.c jh.c luffa.c simd.c blake.c cubehash.c groestl.c keccak.c shavite.c skein.c sha2.c sha2big.c

690
sph/sha2.c
Unescape View File

@ -0,0 +1,690 @@ @@ -0,0 +1,690 @@
/* $Id: sha2.c 227 2010-06-16 17:28:38Z tp $ */
/*
* SHA-224 / SHA-256 implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* Permission is hereby granted, free of charge, to any person obtaining
* a copy of this software and associated documentation files (the
* "Software"), to deal in the Software without restriction, including
* without limitation the rights to use, copy, modify, merge, publish,
* distribute, sublicense, and/or sell copies of the Software, and to
* permit persons to whom the Software is furnished to do so, subject to
* the following conditions:
*
* The above copyright notice and this permission notice shall be
* included in all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
* IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
* CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
* TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
* SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
*
* ===========================(LICENSE END)=============================
*
* @author Thomas Pornin <thomas.pornin@cryptolog.com>
*/
#include <stddef.h>
#include <string.h>
#include "sph_sha2.h"
#if SPH_SMALL_FOOTPRINT && !defined SPH_SMALL_FOOTPRINT_SHA2
#define SPH_SMALL_FOOTPRINT_SHA2 1
#endif
#define CH(X, Y, Z) ((((Y) ^ (Z)) & (X)) ^ (Z))
#define MAJ(X, Y, Z) (((Y) & (Z)) | (((Y) | (Z)) & (X)))
#define ROTR SPH_ROTR32
#define BSG2_0(x) (ROTR(x, 2) ^ ROTR(x, 13) ^ ROTR(x, 22))
#define BSG2_1(x) (ROTR(x, 6) ^ ROTR(x, 11) ^ ROTR(x, 25))
#define SSG2_0(x) (ROTR(x, 7) ^ ROTR(x, 18) ^ SPH_T32((x) >> 3))
#define SSG2_1(x) (ROTR(x, 17) ^ ROTR(x, 19) ^ SPH_T32((x) >> 10))
static const sph_u32 H224[8] = {
SPH_C32(0xC1059ED8), SPH_C32(0x367CD507), SPH_C32(0x3070DD17),
SPH_C32(0xF70E5939), SPH_C32(0xFFC00B31), SPH_C32(0x68581511),
SPH_C32(0x64F98FA7), SPH_C32(0xBEFA4FA4)
};
static const sph_u32 H256[8] = {
SPH_C32(0x6A09E667), SPH_C32(0xBB67AE85), SPH_C32(0x3C6EF372),
SPH_C32(0xA54FF53A), SPH_C32(0x510E527F), SPH_C32(0x9B05688C),
SPH_C32(0x1F83D9AB), SPH_C32(0x5BE0CD19)
};
/*
* The SHA2_ROUND_BODY defines the body for a SHA-224 / SHA-256
* compression function implementation. The "in" parameter should
* evaluate, when applied to a numerical input parameter from 0 to 15,
* to an expression which yields the corresponding input block. The "r"
* parameter should evaluate to an array or pointer expression
* designating the array of 8 words which contains the input and output
* of the compression function.
*/
#if SPH_SMALL_FOOTPRINT_SHA2
static const sph_u32 K[64] = {
SPH_C32(0x428A2F98), SPH_C32(0x71374491),
SPH_C32(0xB5C0FBCF), SPH_C32(0xE9B5DBA5),
SPH_C32(0x3956C25B), SPH_C32(0x59F111F1),
SPH_C32(0x923F82A4), SPH_C32(0xAB1C5ED5),
SPH_C32(0xD807AA98), SPH_C32(0x12835B01),
SPH_C32(0x243185BE), SPH_C32(0x550C7DC3),
SPH_C32(0x72BE5D74), SPH_C32(0x80DEB1FE),
SPH_C32(0x9BDC06A7), SPH_C32(0xC19BF174),
SPH_C32(0xE49B69C1), SPH_C32(0xEFBE4786),
SPH_C32(0x0FC19DC6), SPH_C32(0x240CA1CC),
SPH_C32(0x2DE92C6F), SPH_C32(0x4A7484AA),
SPH_C32(0x5CB0A9DC), SPH_C32(0x76F988DA),
SPH_C32(0x983E5152), SPH_C32(0xA831C66D),
SPH_C32(0xB00327C8), SPH_C32(0xBF597FC7),
SPH_C32(0xC6E00BF3), SPH_C32(0xD5A79147),
SPH_C32(0x06CA6351), SPH_C32(0x14292967),
SPH_C32(0x27B70A85), SPH_C32(0x2E1B2138),
SPH_C32(0x4D2C6DFC), SPH_C32(0x53380D13),
SPH_C32(0x650A7354), SPH_C32(0x766A0ABB),
SPH_C32(0x81C2C92E), SPH_C32(0x92722C85),
SPH_C32(0xA2BFE8A1), SPH_C32(0xA81A664B),
SPH_C32(0xC24B8B70), SPH_C32(0xC76C51A3),
SPH_C32(0xD192E819), SPH_C32(0xD6990624),
SPH_C32(0xF40E3585), SPH_C32(0x106AA070),
SPH_C32(0x19A4C116), SPH_C32(0x1E376C08),
SPH_C32(0x2748774C), SPH_C32(0x34B0BCB5),
SPH_C32(0x391C0CB3), SPH_C32(0x4ED8AA4A),
SPH_C32(0x5B9CCA4F), SPH_C32(0x682E6FF3),
SPH_C32(0x748F82EE), SPH_C32(0x78A5636F),
SPH_C32(0x84C87814), SPH_C32(0x8CC70208),
SPH_C32(0x90BEFFFA), SPH_C32(0xA4506CEB),
SPH_C32(0xBEF9A3F7), SPH_C32(0xC67178F2)
};
#define SHA2_MEXP1(in, pc) do { \
W[pc] = in(pc); \
} while (0)
#define SHA2_MEXP2(in, pc) do { \
W[(pc) & 0x0F] = SPH_T32(SSG2_1(W[((pc) - 2) & 0x0F]) \
+ W[((pc) - 7) & 0x0F] \
+ SSG2_0(W[((pc) - 15) & 0x0F]) + W[(pc) & 0x0F]); \
} while (0)
#define SHA2_STEPn(n, a, b, c, d, e, f, g, h, in, pc) do { \
sph_u32 t1, t2; \
SHA2_MEXP ## n(in, pc); \
t1 = SPH_T32(h + BSG2_1(e) + CH(e, f, g) \
+ K[pcount + (pc)] + W[(pc) & 0x0F]); \
t2 = SPH_T32(BSG2_0(a) + MAJ(a, b, c)); \
d = SPH_T32(d + t1); \
h = SPH_T32(t1 + t2); \
} while (0)
#define SHA2_STEP1(a, b, c, d, e, f, g, h, in, pc) \
SHA2_STEPn(1, a, b, c, d, e, f, g, h, in, pc)
#define SHA2_STEP2(a, b, c, d, e, f, g, h, in, pc) \
SHA2_STEPn(2, a, b, c, d, e, f, g, h, in, pc)
#define SHA2_ROUND_BODY(in, r) do { \
sph_u32 A, B, C, D, E, F, G, H; \
sph_u32 W[16]; \
unsigned pcount; \
\
A = (r)[0]; \
B = (r)[1]; \
C = (r)[2]; \
D = (r)[3]; \
E = (r)[4]; \
F = (r)[5]; \
G = (r)[6]; \
H = (r)[7]; \
pcount = 0; \
SHA2_STEP1(A, B, C, D, E, F, G, H, in, 0); \
SHA2_STEP1(H, A, B, C, D, E, F, G, in, 1); \
SHA2_STEP1(G, H, A, B, C, D, E, F, in, 2); \
SHA2_STEP1(F, G, H, A, B, C, D, E, in, 3); \
SHA2_STEP1(E, F, G, H, A, B, C, D, in, 4); \
SHA2_STEP1(D, E, F, G, H, A, B, C, in, 5); \
SHA2_STEP1(C, D, E, F, G, H, A, B, in, 6); \
SHA2_STEP1(B, C, D, E, F, G, H, A, in, 7); \
SHA2_STEP1(A, B, C, D, E, F, G, H, in, 8); \
SHA2_STEP1(H, A, B, C, D, E, F, G, in, 9); \
SHA2_STEP1(G, H, A, B, C, D, E, F, in, 10); \
SHA2_STEP1(F, G, H, A, B, C, D, E, in, 11); \
SHA2_STEP1(E, F, G, H, A, B, C, D, in, 12); \
SHA2_STEP1(D, E, F, G, H, A, B, C, in, 13); \
SHA2_STEP1(C, D, E, F, G, H, A, B, in, 14); \
SHA2_STEP1(B, C, D, E, F, G, H, A, in, 15); \
for (pcount = 16; pcount < 64; pcount += 16) { \
SHA2_STEP2(A, B, C, D, E, F, G, H, in, 0); \
SHA2_STEP2(H, A, B, C, D, E, F, G, in, 1); \
SHA2_STEP2(G, H, A, B, C, D, E, F, in, 2); \
SHA2_STEP2(F, G, H, A, B, C, D, E, in, 3); \
SHA2_STEP2(E, F, G, H, A, B, C, D, in, 4); \
SHA2_STEP2(D, E, F, G, H, A, B, C, in, 5); \
SHA2_STEP2(C, D, E, F, G, H, A, B, in, 6); \
SHA2_STEP2(B, C, D, E, F, G, H, A, in, 7); \
SHA2_STEP2(A, B, C, D, E, F, G, H, in, 8); \
SHA2_STEP2(H, A, B, C, D, E, F, G, in, 9); \
SHA2_STEP2(G, H, A, B, C, D, E, F, in, 10); \
SHA2_STEP2(F, G, H, A, B, C, D, E, in, 11); \
SHA2_STEP2(E, F, G, H, A, B, C, D, in, 12); \
SHA2_STEP2(D, E, F, G, H, A, B, C, in, 13); \
SHA2_STEP2(C, D, E, F, G, H, A, B, in, 14); \
SHA2_STEP2(B, C, D, E, F, G, H, A, in, 15); \
} \
(r)[0] = SPH_T32((r)[0] + A); \
(r)[1] = SPH_T32((r)[1] + B); \
(r)[2] = SPH_T32((r)[2] + C); \
(r)[3] = SPH_T32((r)[3] + D); \
(r)[4] = SPH_T32((r)[4] + E); \
(r)[5] = SPH_T32((r)[5] + F); \
(r)[6] = SPH_T32((r)[6] + G); \
(r)[7] = SPH_T32((r)[7] + H); \
} while (0)
#else
#define SHA2_ROUND_BODY(in, r) do { \
sph_u32 A, B, C, D, E, F, G, H, T1, T2; \
sph_u32 W00, W01, W02, W03, W04, W05, W06, W07; \
sph_u32 W08, W09, W10, W11, W12, W13, W14, W15; \
\
A = (r)[0]; \
B = (r)[1]; \
C = (r)[2]; \
D = (r)[3]; \
E = (r)[4]; \
F = (r)[5]; \
G = (r)[6]; \
H = (r)[7]; \
W00 = in(0); \
T1 = SPH_T32(H + BSG2_1(E) + CH(E, F, G) \
+ SPH_C32(0x428A2F98) + W00); \
T2 = SPH_T32(BSG2_0(A) + MAJ(A, B, C)); \
D = SPH_T32(D + T1); \
H = SPH_T32(T1 + T2); \
W01 = in(1); \
T1 = SPH_T32(G + BSG2_1(D) + CH(D, E, F) \
+ SPH_C32(0x71374491) + W01); \
T2 = SPH_T32(BSG2_0(H) + MAJ(H, A, B)); \
C = SPH_T32(C + T1); \
G = SPH_T32(T1 + T2); \
W02 = in(2); \
T1 = SPH_T32(F + BSG2_1(C) + CH(C, D, E) \
+ SPH_C32(0xB5C0FBCF) + W02); \
T2 = SPH_T32(BSG2_0(G) + MAJ(G, H, A)); \
B = SPH_T32(B + T1); \
F = SPH_T32(T1 + T2); \
W03 = in(3); \
T1 = SPH_T32(E + BSG2_1(B) + CH(B, C, D) \
+ SPH_C32(0xE9B5DBA5) + W03); \
T2 = SPH_T32(BSG2_0(F) + MAJ(F, G, H)); \
A = SPH_T32(A + T1); \
E = SPH_T32(T1 + T2); \
W04 = in(4); \
T1 = SPH_T32(D + BSG2_1(A) + CH(A, B, C) \
+ SPH_C32(0x3956C25B) + W04); \
T2 = SPH_T32(BSG2_0(E) + MAJ(E, F, G)); \
H = SPH_T32(H + T1); \
D = SPH_T32(T1 + T2); \
W05 = in(5); \
T1 = SPH_T32(C + BSG2_1(H) + CH(H, A, B) \
+ SPH_C32(0x59F111F1) + W05); \
T2 = SPH_T32(BSG2_0(D) + MAJ(D, E, F)); \
G = SPH_T32(G + T1); \
C = SPH_T32(T1 + T2); \
W06 = in(6); \
T1 = SPH_T32(B + BSG2_1(G) + CH(G, H, A) \
+ SPH_C32(0x923F82A4) + W06); \
T2 = SPH_T32(BSG2_0(C) + MAJ(C, D, E)); \
F = SPH_T32(F + T1); \
B = SPH_T32(T1 + T2); \
W07 = in(7); \
T1 = SPH_T32(A + BSG2_1(F) + CH(F, G, H) \
+ SPH_C32(0xAB1C5ED5) + W07); \
T2 = SPH_T32(BSG2_0(B) + MAJ(B, C, D)); \
E = SPH_T32(E + T1); \
A = SPH_T32(T1 + T2); \
W08 = in(8); \
T1 = SPH_T32(H + BSG2_1(E) + CH(E, F, G) \
+ SPH_C32(0xD807AA98) + W08); \
T2 = SPH_T32(BSG2_0(A) + MAJ(A, B, C)); \
D = SPH_T32(D + T1); \
H = SPH_T32(T1 + T2); \
W09 = in(9); \
T1 = SPH_T32(G + BSG2_1(D) + CH(D, E, F) \
+ SPH_C32(0x12835B01) + W09); \
T2 = SPH_T32(BSG2_0(H) + MAJ(H, A, B)); \
C = SPH_T32(C + T1); \
G = SPH_T32(T1 + T2); \
W10 = in(10); \
T1 = SPH_T32(F + BSG2_1(C) + CH(C, D, E) \
+ SPH_C32(0x243185BE) + W10); \
T2 = SPH_T32(BSG2_0(G) + MAJ(G, H, A)); \
B = SPH_T32(B + T1); \
F = SPH_T32(T1 + T2); \
W11 = in(11); \
T1 = SPH_T32(E + BSG2_1(B) + CH(B, C, D) \
+ SPH_C32(0x550C7DC3) + W11); \
T2 = SPH_T32(BSG2_0(F) + MAJ(F, G, H)); \
A = SPH_T32(A + T1); \
E = SPH_T32(T1 + T2); \
W12 = in(12); \
T1 = SPH_T32(D + BSG2_1(A) + CH(A, B, C) \
+ SPH_C32(0x72BE5D74) + W12); \
T2 = SPH_T32(BSG2_0(E) + MAJ(E, F, G)); \
H = SPH_T32(H + T1); \
D = SPH_T32(T1 + T2); \
W13 = in(13); \
T1 = SPH_T32(C + BSG2_1(H) + CH(H, A, B) \
+ SPH_C32(0x80DEB1FE) + W13); \
T2 = SPH_T32(BSG2_0(D) + MAJ(D, E, F)); \
G = SPH_T32(G + T1); \
C = SPH_T32(T1 + T2); \
W14 = in(14); \
T1 = SPH_T32(B + BSG2_1(G) + CH(G, H, A) \
+ SPH_C32(0x9BDC06A7) + W14); \
T2 = SPH_T32(BSG2_0(C) + MAJ(C, D, E)); \
F = SPH_T32(F + T1); \
B = SPH_T32(T1 + T2); \
W15 = in(15); \
T1 = SPH_T32(A + BSG2_1(F) + CH(F, G, H) \
+ SPH_C32(0xC19BF174) + W15); \
T2 = SPH_T32(BSG2_0(B) + MAJ(B, C, D)); \
E = SPH_T32(E + T1); \
A = SPH_T32(T1 + T2); \
W00 = SPH_T32(SSG2_1(W14) + W09 + SSG2_0(W01) + W00); \
T1 = SPH_T32(H + BSG2_1(E) + CH(E, F, G) \
+ SPH_C32(0xE49B69C1) + W00); \
T2 = SPH_T32(BSG2_0(A) + MAJ(A, B, C)); \
D = SPH_T32(D + T1); \
H = SPH_T32(T1 + T2); \
W01 = SPH_T32(SSG2_1(W15) + W10 + SSG2_0(W02) + W01); \
T1 = SPH_T32(G + BSG2_1(D) + CH(D, E, F) \
+ SPH_C32(0xEFBE4786) + W01); \
T2 = SPH_T32(BSG2_0(H) + MAJ(H, A, B)); \
C = SPH_T32(C + T1); \
G = SPH_T32(T1 + T2); \
W02 = SPH_T32(SSG2_1(W00) + W11 + SSG2_0(W03) + W02); \
T1 = SPH_T32(F + BSG2_1(C) + CH(C, D, E) \
+ SPH_C32(0x0FC19DC6) + W02); \
T2 = SPH_T32(BSG2_0(G) + MAJ(G, H, A)); \
B = SPH_T32(B + T1); \
F = SPH_T32(T1 + T2); \
W03 = SPH_T32(SSG2_1(W01) + W12 + SSG2_0(W04) + W03); \
T1 = SPH_T32(E + BSG2_1(B) + CH(B, C, D) \
+ SPH_C32(0x240CA1CC) + W03); \
T2 = SPH_T32(BSG2_0(F) + MAJ(F, G, H)); \
A = SPH_T32(A + T1); \
E = SPH_T32(T1 + T2); \
W04 = SPH_T32(SSG2_1(W02) + W13 + SSG2_0(W05) + W04); \
T1 = SPH_T32(D + BSG2_1(A) + CH(A, B, C) \
+ SPH_C32(0x2DE92C6F) + W04); \
T2 = SPH_T32(BSG2_0(E) + MAJ(E, F, G)); \
H = SPH_T32(H + T1); \
D = SPH_T32(T1 + T2); \
W05 = SPH_T32(SSG2_1(W03) + W14 + SSG2_0(W06) + W05); \
T1 = SPH_T32(C + BSG2_1(H) + CH(H, A, B) \
+ SPH_C32(0x4A7484AA) + W05); \
T2 = SPH_T32(BSG2_0(D) + MAJ(D, E, F)); \
G = SPH_T32(G + T1); \
C = SPH_T32(T1 + T2); \
W06 = SPH_T32(SSG2_1(W04) + W15 + SSG2_0(W07) + W06); \
T1 = SPH_T32(B + BSG2_1(G) + CH(G, H, A) \
+ SPH_C32(0x5CB0A9DC) + W06); \
T2 = SPH_T32(BSG2_0(C) + MAJ(C, D, E)); \
F = SPH_T32(F + T1); \
B = SPH_T32(T1 + T2); \
W07 = SPH_T32(SSG2_1(W05) + W00 + SSG2_0(W08) + W07); \
T1 = SPH_T32(A + BSG2_1(F) + CH(F, G, H) \
+ SPH_C32(0x76F988DA) + W07); \
T2 = SPH_T32(BSG2_0(B) + MAJ(B, C, D)); \
E = SPH_T32(E + T1); \
A = SPH_T32(T1 + T2); \
W08 = SPH_T32(SSG2_1(W06) + W01 + SSG2_0(W09) + W08); \
T1 = SPH_T32(H + BSG2_1(E) + CH(E, F, G) \
+ SPH_C32(0x983E5152) + W08); \
T2 = SPH_T32(BSG2_0(A) + MAJ(A, B, C)); \
D = SPH_T32(D + T1); \
H = SPH_T32(T1 + T2); \
W09 = SPH_T32(SSG2_1(W07) + W02 + SSG2_0(W10) + W09); \
T1 = SPH_T32(G + BSG2_1(D) + CH(D, E, F) \
+ SPH_C32(0xA831C66D) + W09); \
T2 = SPH_T32(BSG2_0(H) + MAJ(H, A, B)); \
C = SPH_T32(C + T1); \
G = SPH_T32(T1 + T2); \
W10 = SPH_T32(SSG2_1(W08) + W03 + SSG2_0(W11) + W10); \
T1 = SPH_T32(F + BSG2_1(C) + CH(C, D, E) \
+ SPH_C32(0xB00327C8) + W10); \
T2 = SPH_T32(BSG2_0(G) + MAJ(G, H, A)); \
B = SPH_T32(B + T1); \
F = SPH_T32(T1 + T2); \
W11 = SPH_T32(SSG2_1(W09) + W04 + SSG2_0(W12) + W11); \
T1 = SPH_T32(E + BSG2_1(B) + CH(B, C, D) \
+ SPH_C32(0xBF597FC7) + W11); \
T2 = SPH_T32(BSG2_0(F) + MAJ(F, G, H)); \
A = SPH_T32(A + T1); \
E = SPH_T32(T1 + T2); \
W12 = SPH_T32(SSG2_1(W10) + W05 + SSG2_0(W13) + W12); \
T1 = SPH_T32(D + BSG2_1(A) + CH(A, B, C) \
+ SPH_C32(0xC6E00BF3) + W12); \
T2 = SPH_T32(BSG2_0(E) + MAJ(E, F, G)); \
H = SPH_T32(H + T1); \
D = SPH_T32(T1 + T2); \
W13 = SPH_T32(SSG2_1(W11) + W06 + SSG2_0(W14) + W13); \
T1 = SPH_T32(C + BSG2_1(H) + CH(H, A, B) \
+ SPH_C32(0xD5A79147) + W13); \
T2 = SPH_T32(BSG2_0(D) + MAJ(D, E, F)); \
G = SPH_T32(G + T1); \
C = SPH_T32(T1 + T2); \
W14 = SPH_T32(SSG2_1(W12) + W07 + SSG2_0(W15) + W14); \
T1 = SPH_T32(B + BSG2_1(G) + CH(G, H, A) \
+ SPH_C32(0x06CA6351) + W14); \
T2 = SPH_T32(BSG2_0(C) + MAJ(C, D, E)); \
F = SPH_T32(F + T1); \
B = SPH_T32(T1 + T2); \
W15 = SPH_T32(SSG2_1(W13) + W08 + SSG2_0(W00) + W15); \
T1 = SPH_T32(A + BSG2_1(F) + CH(F, G, H) \
+ SPH_C32(0x14292967) + W15); \
T2 = SPH_T32(BSG2_0(B) + MAJ(B, C, D)); \
E = SPH_T32(E + T1); \
A = SPH_T32(T1 + T2); \
W00 = SPH_T32(SSG2_1(W14) + W09 + SSG2_0(W01) + W00); \
T1 = SPH_T32(H + BSG2_1(E) + CH(E, F, G) \
+ SPH_C32(0x27B70A85) + W00); \
T2 = SPH_T32(BSG2_0(A) + MAJ(A, B, C)); \
D = SPH_T32(D + T1); \
H = SPH_T32(T1 + T2); \
W01 = SPH_T32(SSG2_1(W15) + W10 + SSG2_0(W02) + W01); \
T1 = SPH_T32(G + BSG2_1(D) + CH(D, E, F) \
+ SPH_C32(0x2E1B2138) + W01); \
T2 = SPH_T32(BSG2_0(H) + MAJ(H, A, B)); \
C = SPH_T32(C + T1); \
G = SPH_T32(T1 + T2); \
W02 = SPH_T32(SSG2_1(W00) + W11 + SSG2_0(W03) + W02); \
T1 = SPH_T32(F + BSG2_1(C) + CH(C, D, E) \
+ SPH_C32(0x4D2C6DFC) + W02); \
T2 = SPH_T32(BSG2_0(G) + MAJ(G, H, A)); \
B = SPH_T32(B + T1); \
F = SPH_T32(T1 + T2); \
W03 = SPH_T32(SSG2_1(W01) + W12 + SSG2_0(W04) + W03); \
T1 = SPH_T32(E + BSG2_1(B) + CH(B, C, D) \
+ SPH_C32(0x53380D13) + W03); \
T2 = SPH_T32(BSG2_0(F) + MAJ(F, G, H)); \
A = SPH_T32(A + T1); \
E = SPH_T32(T1 + T2); \
W04 = SPH_T32(SSG2_1(W02) + W13 + SSG2_0(W05) + W04); \
T1 = SPH_T32(D + BSG2_1(A) + CH(A, B, C) \
+ SPH_C32(0x650A7354) + W04); \
T2 = SPH_T32(BSG2_0(E) + MAJ(E, F, G)); \
H = SPH_T32(H + T1); \
D = SPH_T32(T1 + T2); \
W05 = SPH_T32(SSG2_1(W03) + W14 + SSG2_0(W06) + W05); \
T1 = SPH_T32(C + BSG2_1(H) + CH(H, A, B) \
+ SPH_C32(0x766A0ABB) + W05); \
T2 = SPH_T32(BSG2_0(D) + MAJ(D, E, F)); \
G = SPH_T32(G + T1); \
C = SPH_T32(T1 + T2); \
W06 = SPH_T32(SSG2_1(W04) + W15 + SSG2_0(W07) + W06); \
T1 = SPH_T32(B + BSG2_1(G) + CH(G, H, A) \
+ SPH_C32(0x81C2C92E) + W06); \
T2 = SPH_T32(BSG2_0(C) + MAJ(C, D, E)); \
F = SPH_T32(F + T1); \
B = SPH_T32(T1 + T2); \
W07 = SPH_T32(SSG2_1(W05) + W00 + SSG2_0(W08) + W07); \
T1 = SPH_T32(A + BSG2_1(F) + CH(F, G, H) \
+ SPH_C32(0x92722C85) + W07); \
T2 = SPH_T32(BSG2_0(B) + MAJ(B, C, D)); \
E = SPH_T32(E + T1); \
A = SPH_T32(T1 + T2); \
W08 = SPH_T32(SSG2_1(W06) + W01 + SSG2_0(W09) + W08); \
T1 = SPH_T32(H + BSG2_1(E) + CH(E, F, G) \
+ SPH_C32(0xA2BFE8A1) + W08); \
T2 = SPH_T32(BSG2_0(A) + MAJ(A, B, C)); \
D = SPH_T32(D + T1); \
H = SPH_T32(T1 + T2); \
W09 = SPH_T32(SSG2_1(W07) + W02 + SSG2_0(W10) + W09); \
T1 = SPH_T32(G + BSG2_1(D) + CH(D, E, F) \
+ SPH_C32(0xA81A664B) + W09); \
T2 = SPH_T32(BSG2_0(H) + MAJ(H, A, B)); \
C = SPH_T32(C + T1); \
G = SPH_T32(T1 + T2); \
W10 = SPH_T32(SSG2_1(W08) + W03 + SSG2_0(W11) + W10); \
T1 = SPH_T32(F + BSG2_1(C) + CH(C, D, E) \
+ SPH_C32(0xC24B8B70) + W10); \
T2 = SPH_T32(BSG2_0(G) + MAJ(G, H, A)); \
B = SPH_T32(B + T1); \
F = SPH_T32(T1 + T2); \
W11 = SPH_T32(SSG2_1(W09) + W04 + SSG2_0(W12) + W11); \
T1 = SPH_T32(E + BSG2_1(B) + CH(B, C, D) \
+ SPH_C32(0xC76C51A3) + W11); \
T2 = SPH_T32(BSG2_0(F) + MAJ(F, G, H)); \
A = SPH_T32(A + T1); \
E = SPH_T32(T1 + T2); \
W12 = SPH_T32(SSG2_1(W10) + W05 + SSG2_0(W13) + W12); \
T1 = SPH_T32(D + BSG2_1(A) + CH(A, B, C) \
+ SPH_C32(0xD192E819) + W12); \
T2 = SPH_T32(BSG2_0(E) + MAJ(E, F, G)); \
H = SPH_T32(H + T1); \
D = SPH_T32(T1 + T2); \
W13 = SPH_T32(SSG2_1(W11) + W06 + SSG2_0(W14) + W13); \
T1 = SPH_T32(C + BSG2_1(H) + CH(H, A, B) \
+ SPH_C32(0xD6990624) + W13); \
T2 = SPH_T32(BSG2_0(D) + MAJ(D, E, F)); \
G = SPH_T32(G + T1); \
C = SPH_T32(T1 + T2); \
W14 = SPH_T32(SSG2_1(W12) + W07 + SSG2_0(W15) + W14); \
T1 = SPH_T32(B + BSG2_1(G) + CH(G, H, A) \
+ SPH_C32(0xF40E3585) + W14); \
T2 = SPH_T32(BSG2_0(C) + MAJ(C, D, E)); \
F = SPH_T32(F + T1); \
B = SPH_T32(T1 + T2); \
W15 = SPH_T32(SSG2_1(W13) + W08 + SSG2_0(W00) + W15); \
T1 = SPH_T32(A + BSG2_1(F) + CH(F, G, H) \
+ SPH_C32(0x106AA070) + W15); \
T2 = SPH_T32(BSG2_0(B) + MAJ(B, C, D)); \
E = SPH_T32(E + T1); \
A = SPH_T32(T1 + T2); \
W00 = SPH_T32(SSG2_1(W14) + W09 + SSG2_0(W01) + W00); \
T1 = SPH_T32(H + BSG2_1(E) + CH(E, F, G) \
+ SPH_C32(0x19A4C116) + W00); \
T2 = SPH_T32(BSG2_0(A) + MAJ(A, B, C)); \
D = SPH_T32(D + T1); \
H = SPH_T32(T1 + T2); \
W01 = SPH_T32(SSG2_1(W15) + W10 + SSG2_0(W02) + W01); \
T1 = SPH_T32(G + BSG2_1(D) + CH(D, E, F) \
+ SPH_C32(0x1E376C08) + W01); \
T2 = SPH_T32(BSG2_0(H) + MAJ(H, A, B)); \
C = SPH_T32(C + T1); \
G = SPH_T32(T1 + T2); \
W02 = SPH_T32(SSG2_1(W00) + W11 + SSG2_0(W03) + W02); \
T1 = SPH_T32(F + BSG2_1(C) + CH(C, D, E) \
+ SPH_C32(0x2748774C) + W02); \
T2 = SPH_T32(BSG2_0(G) + MAJ(G, H, A)); \
B = SPH_T32(B + T1); \
F = SPH_T32(T1 + T2); \
W03 = SPH_T32(SSG2_1(W01) + W12 + SSG2_0(W04) + W03); \
T1 = SPH_T32(E + BSG2_1(B) + CH(B, C, D) \
+ SPH_C32(0x34B0BCB5) + W03); \
T2 = SPH_T32(BSG2_0(F) + MAJ(F, G, H)); \
A = SPH_T32(A + T1); \
E = SPH_T32(T1 + T2); \
W04 = SPH_T32(SSG2_1(W02) + W13 + SSG2_0(W05) + W04); \
T1 = SPH_T32(D + BSG2_1(A) + CH(A, B, C) \
+ SPH_C32(0x391C0CB3) + W04); \
T2 = SPH_T32(BSG2_0(E) + MAJ(E, F, G)); \
H = SPH_T32(H + T1); \
D = SPH_T32(T1 + T2); \
W05 = SPH_T32(SSG2_1(W03) + W14 + SSG2_0(W06) + W05); \
T1 = SPH_T32(C + BSG2_1(H) + CH(H, A, B) \
+ SPH_C32(0x4ED8AA4A) + W05); \
T2 = SPH_T32(BSG2_0(D) + MAJ(D, E, F)); \
G = SPH_T32(G + T1); \
C = SPH_T32(T1 + T2); \
W06 = SPH_T32(SSG2_1(W04) + W15 + SSG2_0(W07) + W06); \
T1 = SPH_T32(B + BSG2_1(G) + CH(G, H, A) \
+ SPH_C32(0x5B9CCA4F) + W06); \
T2 = SPH_T32(BSG2_0(C) + MAJ(C, D, E)); \
F = SPH_T32(F + T1); \
B = SPH_T32(T1 + T2); \
W07 = SPH_T32(SSG2_1(W05) + W00 + SSG2_0(W08) + W07); \
T1 = SPH_T32(A + BSG2_1(F) + CH(F, G, H) \
+ SPH_C32(0x682E6FF3) + W07); \
T2 = SPH_T32(BSG2_0(B) + MAJ(B, C, D)); \
E = SPH_T32(E + T1); \
A = SPH_T32(T1 + T2); \
W08 = SPH_T32(SSG2_1(W06) + W01 + SSG2_0(W09) + W08); \
T1 = SPH_T32(H + BSG2_1(E) + CH(E, F, G) \
+ SPH_C32(0x748F82EE) + W08); \
T2 = SPH_T32(BSG2_0(A) + MAJ(A, B, C)); \
D = SPH_T32(D + T1); \
H = SPH_T32(T1 + T2); \
W09 = SPH_T32(SSG2_1(W07) + W02 + SSG2_0(W10) + W09); \
T1 = SPH_T32(G + BSG2_1(D) + CH(D, E, F) \
+ SPH_C32(0x78A5636F) + W09); \
T2 = SPH_T32(BSG2_0(H) + MAJ(H, A, B)); \
C = SPH_T32(C + T1); \
G = SPH_T32(T1 + T2); \
W10 = SPH_T32(SSG2_1(W08) + W03 + SSG2_0(W11) + W10); \
T1 = SPH_T32(F + BSG2_1(C) + CH(C, D, E) \
+ SPH_C32(0x84C87814) + W10); \
T2 = SPH_T32(BSG2_0(G) + MAJ(G, H, A)); \
B = SPH_T32(B + T1); \
F = SPH_T32(T1 + T2); \
W11 = SPH_T32(SSG2_1(W09) + W04 + SSG2_0(W12) + W11); \
T1 = SPH_T32(E + BSG2_1(B) + CH(B, C, D) \
+ SPH_C32(0x8CC70208) + W11); \
T2 = SPH_T32(BSG2_0(F) + MAJ(F, G, H)); \
A = SPH_T32(A + T1); \
E = SPH_T32(T1 + T2); \
W12 = SPH_T32(SSG2_1(W10) + W05 + SSG2_0(W13) + W12); \
T1 = SPH_T32(D + BSG2_1(A) + CH(A, B, C) \
+ SPH_C32(0x90BEFFFA) + W12); \
T2 = SPH_T32(BSG2_0(E) + MAJ(E, F, G)); \
H = SPH_T32(H + T1); \
D = SPH_T32(T1 + T2); \
W13 = SPH_T32(SSG2_1(W11) + W06 + SSG2_0(W14) + W13); \
T1 = SPH_T32(C + BSG2_1(H) + CH(H, A, B) \
+ SPH_C32(0xA4506CEB) + W13); \
T2 = SPH_T32(BSG2_0(D) + MAJ(D, E, F)); \
G = SPH_T32(G + T1); \
C = SPH_T32(T1 + T2); \
W14 = SPH_T32(SSG2_1(W12) + W07 + SSG2_0(W15) + W14); \
T1 = SPH_T32(B + BSG2_1(G) + CH(G, H, A) \
+ SPH_C32(0xBEF9A3F7) + W14); \
T2 = SPH_T32(BSG2_0(C) + MAJ(C, D, E)); \
F = SPH_T32(F + T1); \
B = SPH_T32(T1 + T2); \
W15 = SPH_T32(SSG2_1(W13) + W08 + SSG2_0(W00) + W15); \
T1 = SPH_T32(A + BSG2_1(F) + CH(F, G, H) \
+ SPH_C32(0xC67178F2) + W15); \
T2 = SPH_T32(BSG2_0(B) + MAJ(B, C, D)); \
E = SPH_T32(E + T1); \
A = SPH_T32(T1 + T2); \
(r)[0] = SPH_T32((r)[0] + A); \
(r)[1] = SPH_T32((r)[1] + B); \
(r)[2] = SPH_T32((r)[2] + C); \
(r)[3] = SPH_T32((r)[3] + D); \
(r)[4] = SPH_T32((r)[4] + E); \
(r)[5] = SPH_T32((r)[5] + F); \
(r)[6] = SPH_T32((r)[6] + G); \
(r)[7] = SPH_T32((r)[7] + H); \
} while (0)
#endif
/*
* One round of SHA-224 / SHA-256. The data must be aligned for 32-bit access.
*/
static void
sha2_round(const unsigned char *data, sph_u32 r[8])
{
#define SHA2_IN(x) sph_dec32be_aligned(data + (4 * (x)))
SHA2_ROUND_BODY(SHA2_IN, r);
#undef SHA2_IN
}
/* see sph_sha2.h */
void
sph_sha224_init(void *cc)
{
sph_sha224_context *sc;
sc = cc;
memcpy(sc->val, H224, sizeof H224);
#if SPH_64
sc->count = 0;
#else
sc->count_high = sc->count_low = 0;
#endif
}
/* see sph_sha2.h */
void
sph_sha256_init(void *cc)
{
sph_sha256_context *sc;
sc = cc;
memcpy(sc->val, H256, sizeof H256);
#if SPH_64
sc->count = 0;
#else
sc->count_high = sc->count_low = 0;
#endif
}
#define RFUN sha2_round
#define HASH sha224
#define BE32 1
#include "md_helper.c"
/* see sph_sha2.h */
void
sph_sha224_close(void *cc, void *dst)
{
sha224_close(cc, dst, 7);
sph_sha224_init(cc);
}
/* see sph_sha2.h */
void
sph_sha224_addbits_and_close(void *cc, unsigned ub, unsigned n, void *dst)
{
sha224_addbits_and_close(cc, ub, n, dst, 7);
sph_sha224_init(cc);
}
/* see sph_sha2.h */
void
sph_sha256_close(void *cc, void *dst)
{
sha224_close(cc, dst, 8);
sph_sha256_init(cc);
}
/* see sph_sha2.h */
void
sph_sha256_addbits_and_close(void *cc, unsigned ub, unsigned n, void *dst)
{
sha224_addbits_and_close(cc, ub, n, dst, 8);
sph_sha256_init(cc);
}
/* see sph_sha2.h */
void
sph_sha224_comp(const sph_u32 msg[16], sph_u32 val[8])
{
#define SHA2_IN(x) msg[x]
SHA2_ROUND_BODY(SHA2_IN, val);
#undef SHA2_IN
}

247
sph/sha2big.c
Unescape View File

@ -0,0 +1,247 @@ @@ -0,0 +1,247 @@
/* $Id: sha2big.c 216 2010-06-08 09:46:57Z tp $ */
/*
* SHA-384 / SHA-512 implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* Permission is hereby granted, free of charge, to any person obtaining
* a copy of this software and associated documentation files (the
* "Software"), to deal in the Software without restriction, including
* without limitation the rights to use, copy, modify, merge, publish,
* distribute, sublicense, and/or sell copies of the Software, and to
* permit persons to whom the Software is furnished to do so, subject to
* the following conditions:
*
* The above copyright notice and this permission notice shall be
* included in all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
* IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
* CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
* TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
* SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
*
* ===========================(LICENSE END)=============================
*
* @author Thomas Pornin <thomas.pornin@cryptolog.com>
*/
#include <stddef.h>
#include <string.h>
#include "sph_sha2.h"
#if SPH_64
#define CH(X, Y, Z) ((((Y) ^ (Z)) & (X)) ^ (Z))
#define MAJ(X, Y, Z) (((X) & (Y)) | (((X) | (Y)) & (Z)))
#define ROTR64 SPH_ROTR64
#define BSG5_0(x) (ROTR64(x, 28) ^ ROTR64(x, 34) ^ ROTR64(x, 39))
#define BSG5_1(x) (ROTR64(x, 14) ^ ROTR64(x, 18) ^ ROTR64(x, 41))
#define SSG5_0(x) (ROTR64(x, 1) ^ ROTR64(x, 8) ^ SPH_T64((x) >> 7))
#define SSG5_1(x) (ROTR64(x, 19) ^ ROTR64(x, 61) ^ SPH_T64((x) >> 6))
static const sph_u64 K512[80] = {
SPH_C64(0x428A2F98D728AE22), SPH_C64(0x7137449123EF65CD),
SPH_C64(0xB5C0FBCFEC4D3B2F), SPH_C64(0xE9B5DBA58189DBBC),
SPH_C64(0x3956C25BF348B538), SPH_C64(0x59F111F1B605D019),
SPH_C64(0x923F82A4AF194F9B), SPH_C64(0xAB1C5ED5DA6D8118),
SPH_C64(0xD807AA98A3030242), SPH_C64(0x12835B0145706FBE),
SPH_C64(0x243185BE4EE4B28C), SPH_C64(0x550C7DC3D5FFB4E2),
SPH_C64(0x72BE5D74F27B896F), SPH_C64(0x80DEB1FE3B1696B1),
SPH_C64(0x9BDC06A725C71235), SPH_C64(0xC19BF174CF692694),
SPH_C64(0xE49B69C19EF14AD2), SPH_C64(0xEFBE4786384F25E3),
SPH_C64(0x0FC19DC68B8CD5B5), SPH_C64(0x240CA1CC77AC9C65),
SPH_C64(0x2DE92C6F592B0275), SPH_C64(0x4A7484AA6EA6E483),
SPH_C64(0x5CB0A9DCBD41FBD4), SPH_C64(0x76F988DA831153B5),
SPH_C64(0x983E5152EE66DFAB), SPH_C64(0xA831C66D2DB43210),
SPH_C64(0xB00327C898FB213F), SPH_C64(0xBF597FC7BEEF0EE4),
SPH_C64(0xC6E00BF33DA88FC2), SPH_C64(0xD5A79147930AA725),
SPH_C64(0x06CA6351E003826F), SPH_C64(0x142929670A0E6E70),
SPH_C64(0x27B70A8546D22FFC), SPH_C64(0x2E1B21385C26C926),
SPH_C64(0x4D2C6DFC5AC42AED), SPH_C64(0x53380D139D95B3DF),
SPH_C64(0x650A73548BAF63DE), SPH_C64(0x766A0ABB3C77B2A8),
SPH_C64(0x81C2C92E47EDAEE6), SPH_C64(0x92722C851482353B),
SPH_C64(0xA2BFE8A14CF10364), SPH_C64(0xA81A664BBC423001),
SPH_C64(0xC24B8B70D0F89791), SPH_C64(0xC76C51A30654BE30),
SPH_C64(0xD192E819D6EF5218), SPH_C64(0xD69906245565A910),
SPH_C64(0xF40E35855771202A), SPH_C64(0x106AA07032BBD1B8),
SPH_C64(0x19A4C116B8D2D0C8), SPH_C64(0x1E376C085141AB53),
SPH_C64(0x2748774CDF8EEB99), SPH_C64(0x34B0BCB5E19B48A8),
SPH_C64(0x391C0CB3C5C95A63), SPH_C64(0x4ED8AA4AE3418ACB),
SPH_C64(0x5B9CCA4F7763E373), SPH_C64(0x682E6FF3D6B2B8A3),
SPH_C64(0x748F82EE5DEFB2FC), SPH_C64(0x78A5636F43172F60),
SPH_C64(0x84C87814A1F0AB72), SPH_C64(0x8CC702081A6439EC),
SPH_C64(0x90BEFFFA23631E28), SPH_C64(0xA4506CEBDE82BDE9),
SPH_C64(0xBEF9A3F7B2C67915), SPH_C64(0xC67178F2E372532B),
SPH_C64(0xCA273ECEEA26619C), SPH_C64(0xD186B8C721C0C207),
SPH_C64(0xEADA7DD6CDE0EB1E), SPH_C64(0xF57D4F7FEE6ED178),
SPH_C64(0x06F067AA72176FBA), SPH_C64(0x0A637DC5A2C898A6),
SPH_C64(0x113F9804BEF90DAE), SPH_C64(0x1B710B35131C471B),
SPH_C64(0x28DB77F523047D84), SPH_C64(0x32CAAB7B40C72493),
SPH_C64(0x3C9EBE0A15C9BEBC), SPH_C64(0x431D67C49C100D4C),
SPH_C64(0x4CC5D4BECB3E42B6), SPH_C64(0x597F299CFC657E2A),
SPH_C64(0x5FCB6FAB3AD6FAEC), SPH_C64(0x6C44198C4A475817)
};
static const sph_u64 H384[8] = {
SPH_C64(0xCBBB9D5DC1059ED8), SPH_C64(0x629A292A367CD507),
SPH_C64(0x9159015A3070DD17), SPH_C64(0x152FECD8F70E5939),
SPH_C64(0x67332667FFC00B31), SPH_C64(0x8EB44A8768581511),
SPH_C64(0xDB0C2E0D64F98FA7), SPH_C64(0x47B5481DBEFA4FA4)
};
static const sph_u64 H512[8] = {
SPH_C64(0x6A09E667F3BCC908), SPH_C64(0xBB67AE8584CAA73B),
SPH_C64(0x3C6EF372FE94F82B), SPH_C64(0xA54FF53A5F1D36F1),
SPH_C64(0x510E527FADE682D1), SPH_C64(0x9B05688C2B3E6C1F),
SPH_C64(0x1F83D9ABFB41BD6B), SPH_C64(0x5BE0CD19137E2179)
};
/*
* This macro defines the body for a SHA-384 / SHA-512 compression function
* implementation. The "in" parameter should evaluate, when applied to a
* numerical input parameter from 0 to 15, to an expression which yields
* the corresponding input block. The "r" parameter should evaluate to
* an array or pointer expression designating the array of 8 words which
* contains the input and output of the compression function.
*
* SHA-512 is hard for the compiler. If the loop is completely unrolled,
* then the code will be quite huge (possibly more than 100 kB), and the
* performance will be degraded due to cache misses on the code. We
* unroll only eight steps, which avoids all needless copies when
* 64-bit registers are swapped.
*/
#define SHA3_STEP(A, B, C, D, E, F, G, H, i) do { \
sph_u64 T1, T2; \
T1 = SPH_T64(H + BSG5_1(E) + CH(E, F, G) + K512[i] + W[i]); \
T2 = SPH_T64(BSG5_0(A) + MAJ(A, B, C)); \
D = SPH_T64(D + T1); \
H = SPH_T64(T1 + T2); \
} while (0)
#define SHA3_ROUND_BODY(in, r) do { \
int i; \
sph_u64 A, B, C, D, E, F, G, H; \
sph_u64 W[80]; \
\
for (i = 0; i < 16; i ++) \
W[i] = in(i); \
for (i = 16; i < 80; i ++) \
W[i] = SPH_T64(SSG5_1(W[i - 2]) + W[i - 7] \
+ SSG5_0(W[i - 15]) + W[i - 16]); \
A = (r)[0]; \
B = (r)[1]; \
C = (r)[2]; \
D = (r)[3]; \
E = (r)[4]; \
F = (r)[5]; \
G = (r)[6]; \
H = (r)[7]; \
for (i = 0; i < 80; i += 8) { \
SHA3_STEP(A, B, C, D, E, F, G, H, i + 0); \
SHA3_STEP(H, A, B, C, D, E, F, G, i + 1); \
SHA3_STEP(G, H, A, B, C, D, E, F, i + 2); \
SHA3_STEP(F, G, H, A, B, C, D, E, i + 3); \
SHA3_STEP(E, F, G, H, A, B, C, D, i + 4); \
SHA3_STEP(D, E, F, G, H, A, B, C, i + 5); \
SHA3_STEP(C, D, E, F, G, H, A, B, i + 6); \
SHA3_STEP(B, C, D, E, F, G, H, A, i + 7); \
} \
(r)[0] = SPH_T64((r)[0] + A); \
(r)[1] = SPH_T64((r)[1] + B); \
(r)[2] = SPH_T64((r)[2] + C); \
(r)[3] = SPH_T64((r)[3] + D); \
(r)[4] = SPH_T64((r)[4] + E); \
(r)[5] = SPH_T64((r)[5] + F); \
(r)[6] = SPH_T64((r)[6] + G); \
(r)[7] = SPH_T64((r)[7] + H); \
} while (0)
/*
* One round of SHA-384 / SHA-512. The data must be aligned for 64-bit access.
*/
static void
sha3_round(const unsigned char *data, sph_u64 r[8])
{
#define SHA3_IN(x) sph_dec64be_aligned(data + (8 * (x)))
SHA3_ROUND_BODY(SHA3_IN, r);
#undef SHA3_IN
}
/* see sph_sha3.h */
void
sph_sha384_init(void *cc)
{
sph_sha384_context *sc;
sc = cc;
memcpy(sc->val, H384, sizeof H384);
sc->count = 0;
}
/* see sph_sha3.h */
void
sph_sha512_init(void *cc)
{
sph_sha512_context *sc;
sc = cc;
memcpy(sc->val, H512, sizeof H512);
sc->count = 0;
}
#define RFUN sha3_round
#define HASH sha384
#define BE64 1
#include "md_helper.c"
/* see sph_sha3.h */
void
sph_sha384_close(void *cc, void *dst)
{
sha384_close(cc, dst, 6);
sph_sha384_init(cc);
}
/* see sph_sha3.h */
void
sph_sha384_addbits_and_close(void *cc, unsigned ub, unsigned n, void *dst)
{
sha384_addbits_and_close(cc, ub, n, dst, 6);
sph_sha384_init(cc);
}
/* see sph_sha3.h */
void
sph_sha512_close(void *cc, void *dst)
{
sha384_close(cc, dst, 8);
sph_sha512_init(cc);
}
/* see sph_sha3.h */
void
sph_sha512_addbits_and_close(void *cc, unsigned ub, unsigned n, void *dst)
{
sha384_addbits_and_close(cc, ub, n, dst, 8);
sph_sha512_init(cc);
}
/* see sph_sha3.h */
void
sph_sha384_comp(const sph_u64 msg[16], sph_u64 val[8])
{
#define SHA3_IN(x) msg[x]
SHA3_ROUND_BODY(SHA3_IN, val);
#undef SHA3_IN
}
#endif
Write Preview
Loading…
Cancel
Save