GOSTcoin support for ccminer CUDA miner project, compatible with most nvidia cards
 
 
 
 
 

565 lines
19 KiB

/**
* Based on the SPH implementation of blake2s
* Provos Alexis - 2016
*/
#include "miner.h"
#include <string.h>
#include <stdint.h>
#include "sph/blake2s.h"
#include "sph/sph_types.h"
#ifdef __INTELLISENSE__
#define __byte_perm(x, y, b) x
#endif
#include "cuda_helper.h"
#ifdef __CUDA_ARCH__
__device__ __forceinline__
uint32_t ROR8(const uint32_t a) {
return __byte_perm(a, 0, 0x0321);
}
__device__ __forceinline__
uint32_t ROL16(const uint32_t a) {
return __byte_perm(a, 0, 0x1032);
}
#else
#define ROR8(u) (u >> 8)
#define ROL16(u) (u << 16)
#endif
__device__ __forceinline__
uint32_t xor3x(uint32_t a, uint32_t b, uint32_t c)
{
uint32_t result;
#if __CUDA_ARCH__ >= 500 && CUDA_VERSION >= 7050
asm ("lop3.b32 %0, %1, %2, %3, 0x96;" : "=r"(result) : "r"(a), "r"(b),"r"(c)); //0x96 = 0xF0 ^ 0xCC ^ 0xAA
#else
result = a^b^c;
#endif
return result;
}
static const uint32_t blake2s_IV[8] = {
0x6A09E667UL, 0xBB67AE85UL, 0x3C6EF372UL, 0xA54FF53AUL,
0x510E527FUL, 0x9B05688CUL, 0x1F83D9ABUL, 0x5BE0CD19UL
};
static const uint8_t blake2s_sigma[10][16] = {
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
{ 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
{ 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
{ 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
{ 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
{ 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
{ 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
{ 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
{ 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
{ 10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13 , 0 },
};
#define G(r,i,a,b,c,d) \
do { \
a = a + b + m[blake2s_sigma[r][2*i+0]]; \
d = SPH_ROTR32(d ^ a, 16); \
c = c + d; \
b = SPH_ROTR32(b ^ c, 12); \
a = a + b + m[blake2s_sigma[r][2*i+1]]; \
d = SPH_ROTR32(d ^ a, 8); \
c = c + d; \
b = SPH_ROTR32(b ^ c, 7); \
} while(0)
#define ROUND(r) \
do { \
G(r,0,v[0],v[4],v[ 8],v[12]); \
G(r,1,v[1],v[5],v[ 9],v[13]); \
G(r,2,v[2],v[6],v[10],v[14]); \
G(r,3,v[3],v[7],v[11],v[15]); \
G(r,4,v[0],v[5],v[10],v[15]); \
G(r,5,v[1],v[6],v[11],v[12]); \
G(r,6,v[2],v[7],v[ 8],v[13]); \
G(r,7,v[3],v[4],v[ 9],v[14]); \
} while(0)
extern "C" void blake2s_hash(void *output, const void *input)
{
uint32_t m[16];
uint32_t v[16];
uint32_t h[8];
uint32_t *in = (uint32_t*)input;
// COMPRESS
for(int i = 0; i < 16; ++i )
m[i] = in[i];
h[0] = 0x01010020 ^ blake2s_IV[0];
h[1] = blake2s_IV[1];
h[2] = blake2s_IV[2];
h[3] = blake2s_IV[3];
h[4] = blake2s_IV[4];
h[5] = blake2s_IV[5];
h[6] = blake2s_IV[6];
h[7] = blake2s_IV[7];
for(int i = 0; i < 8; ++i )
v[i] = h[i];
v[ 8] = blake2s_IV[0]; v[ 9] = blake2s_IV[1];
v[10] = blake2s_IV[2]; v[11] = blake2s_IV[3];
v[12] = 64 ^ blake2s_IV[4]; v[13] = blake2s_IV[5];
v[14] = blake2s_IV[6]; v[15] = blake2s_IV[7];
ROUND( 0 ); ROUND( 1 );
ROUND( 2 ); ROUND( 3 );
ROUND( 4 ); ROUND( 5 );
ROUND( 6 ); ROUND( 7 );
ROUND( 8 ); ROUND( 9 );
for(size_t i = 0; i < 8; ++i)
h[i] ^= v[i] ^ v[i + 8];
// COMPRESS
m[0] = in[16]; m[1] = in[17];
m[2] = in[18]; m[3] = in[19];
for(size_t i = 4; i < 16; ++i)
m[i] = 0;
for(size_t i = 0; i < 8; ++i)
v[i] = h[i];
v[ 8] = blake2s_IV[0]; v[ 9] = blake2s_IV[1];
v[10] = blake2s_IV[2]; v[11] = blake2s_IV[3];
v[12] = 0x50 ^ blake2s_IV[4]; v[13] = blake2s_IV[5];
v[14] = ~blake2s_IV[6]; v[15] = blake2s_IV[7];
ROUND( 0 ); ROUND( 1 );
ROUND( 2 ); ROUND( 3 );
ROUND( 4 ); ROUND( 5 );
ROUND( 6 ); ROUND( 7 );
ROUND( 8 ); ROUND( 9 );
for(size_t i = 0; i < 8; ++i)
h[i] ^= v[i] ^ v[i + 8];
memcpy(output, h, 32);
}
#define TPB 1024
#define NPT 256
#define maxResults 16
#define NBN 1
__constant__ uint32_t _ALIGN(32) midstate[20];
static uint32_t *d_resNonce[MAX_GPUS];
static uint32_t *h_resNonce[MAX_GPUS];
#define GS4(a,b,c,d,e,f,a1,b1,c1,d1,e1,f1,a2,b2,c2,d2,e2,f2,a3,b3,c3,d3,e3,f3){ \
a += b + e; a1+= b1 + e1; a2+= b2 + e2; a3+= b3 + e3; \
d = ROL16( d ^ a); d1 = ROL16(d1 ^ a1); d2 = ROL16(d2 ^ a2); d3 = ROL16(d3 ^ a3); \
c +=d; c1+=d1; c2+=d2; c3+=d3;\
b = ROTR32(b ^ c, 12); b1 = ROTR32(b1^c1, 12); b2 = ROTR32(b2^c2, 12); b3 = ROTR32(b3^c3, 12); \
a += b + f; a1+= b1 + f1; a2+= b2 + f2; a3+= b3 + f3; \
d = ROR8(d ^ a); d1 = ROR8(d1^a1); d2 = ROR8(d2^a2); d3 = ROR8(d3^a3); \
c += d; c1 += d1; c2 += d2; c3 += d3;\
b = ROTR32(b ^ c, 7); b1 = ROTR32(b1^c1, 7); b2 = ROTR32(b2^c2, 7); b3 = ROTR32(b3^c3, 7); \
}
__global__ __launch_bounds__(TPB,1)
void blake2s_gpu_hash_nonce(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce, const uint32_t ptarget7)
{
const uint32_t step = gridDim.x * blockDim.x;
uint32_t m[ 3];
uint32_t v[16];
m[0] = midstate[16];
m[1] = midstate[17];
m[2] = midstate[18];
const uint32_t h7 = midstate[19];
for(uint32_t thread = blockDim.x * blockIdx.x + threadIdx.x ; thread <threads; thread+=step){
#pragma unroll
for(int i=0;i<16;i++){
v[ i] = midstate[ i];
}
uint32_t nonce = cuda_swab32(startNonce + thread);
// Round( 0 );
v[ 1] += nonce;
v[13] = ROR8(v[13] ^ v[ 1]);
v[ 9] += v[13];
v[ 5] = ROTR32(v[ 5] ^ v[ 9], 7);
v[ 1]+= v[ 6];
v[ 0]+= v[ 5];
v[12] = ROL16(v[12] ^ v[ 1]);
v[13] = ROL16(v[13] ^ v[ 2]);
v[15] = ROL16(v[15] ^ v[ 0]);
v[11]+= v[12]; v[ 8]+= v[13]; v[ 9]+= v[14]; v[10]+= v[15];
v[ 6] = ROTR32(v[ 6] ^ v[11], 12); v[ 7] = ROTR32(v[ 7] ^ v[ 8], 12); v[ 4] = ROTR32(v[ 4] ^ v[ 9], 12); v[ 5] = ROTR32(v[ 5] ^ v[10], 12);
v[ 1]+= v[ 6]; v[ 2]+= v[ 7]; v[ 3]+= v[ 4]; v[ 0]+= v[ 5];
v[12] = ROR8(v[12] ^ v[ 1]); v[13] = ROR8(v[13] ^ v[ 2]); v[14] = ROR8(v[14] ^ v[ 3]); v[15] = ROR8(v[15] ^ v[ 0]);
v[11]+= v[12]; v[ 8]+= v[13]; v[ 9]+= v[14]; v[10]+= v[15];
v[ 6] = ROTR32(v[ 6] ^ v[11], 7); v[ 7] = ROTR32(v[ 7] ^ v[ 8], 7); v[ 4] = ROTR32(v[ 4] ^ v[ 9], 7); v[ 5] = ROTR32(v[ 5] ^ v[10], 7);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],0,0, v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],m[ 1],0, v[ 1],v[ 6],v[11],v[12],m[ 0],m[ 2], v[ 2],v[ 7],v[ 8],v[13],0,0, v[ 3],v[ 4],v[ 9],v[14],0,nonce);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],0,m[ 0], v[ 2],v[ 6],v[10],v[14],0,m[ 2], v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],0,0, v[ 1],v[ 6],v[11],v[12],nonce,0, v[ 2],v[ 7],v[ 8],v[13],0,m[ 1], v[ 3],v[ 4],v[ 9],v[14],0,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],nonce,m[ 1], v[ 2],v[ 6],v[10],v[14],0,0, v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],m[ 2],0, v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],0,m[ 0], v[ 3],v[ 4],v[ 9],v[14],0,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,m[ 0], v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],m[ 2],0, v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],0,m[ 1], v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],0,0, v[ 3],v[ 4],v[ 9],v[14],nonce,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],m[ 2],0, v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],m[ 0],0, v[ 3],v[ 7],v[11],v[15],0,nonce);
GS4(v[ 0],v[ 5],v[10],v[15],0,0, v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],0,0, v[ 3],v[ 4],v[ 9],v[14],m[ 1],0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],m[ 1],0, v[ 2],v[ 6],v[10],v[14],0,0, v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],m[ 0],0, v[ 1],v[ 6],v[11],v[12],0,nonce, v[ 2],v[ 7],v[ 8],v[13],0,m[ 2], v[ 3],v[ 4],v[ 9],v[14],0,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],0,m[ 1], v[ 3],v[ 7],v[11],v[15],nonce,0);
GS4(v[ 0],v[ 5],v[10],v[15],0,m[ 0], v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],0,0, v[ 3],v[ 4],v[ 9],v[14],m[ 2],0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],0,nonce, v[ 3],v[ 7],v[11],v[15],m[ 0],0);
GS4(v[ 0],v[ 5],v[10],v[15],0,m[ 2], v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],m[ 1],0, v[ 3],v[ 4],v[ 9],v[14],0,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,m[ 2], v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],0,0, v[ 3],v[ 7],v[11],v[15],m[ 1],0);
// GS(9,4,v[ 0],v[ 5],v[10],v[15]);
v[ 0] += v[ 5];
v[ 2] += v[ 7] + nonce;
v[15] = ROL16(v[15] ^ v[ 0]);
v[13] = ROL16(v[13] ^ v[ 2]);
v[10] += v[15];
v[ 8] += v[13];
v[ 5] = ROTR32(v[ 5] ^ v[10], 12);
v[ 7] = ROTR32(v[ 7] ^ v[ 8], 12);
v[ 0] += v[ 5];
v[ 2] += v[ 7];
v[15] = ROR8(v[15] ^ v[ 0]);
v[13] = ROR8(v[13] ^ v[ 2]);
v[ 8] += v[13];
v[ 7] = ROTR32(v[ 7] ^ v[ 8], 7);
if (xor3x(h7,v[7],v[15]) <= ptarget7){
uint32_t pos = atomicInc(&resNonce[0],0xffffffff)+1;
if(pos < maxResults)
resNonce[pos] = nonce;
return;
}
}
}
__global__ __launch_bounds__(TPB,1)
void blake2s_gpu_hash_nonce(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce)
{
const uint32_t step = gridDim.x * blockDim.x;
uint32_t m[ 3];
uint32_t v[16];
m[0] = midstate[16];
m[1] = midstate[17];
m[2] = midstate[18];
const uint32_t h7 = midstate[19];
for(uint32_t thread = blockDim.x * blockIdx.x + threadIdx.x ; thread <threads; thread+=step)
{
#pragma unroll
for(int i=0;i<16;i++){
v[ i] = midstate[ i];
}
uint32_t nonce = cuda_swab32(startNonce+thread);
// Round( 0 );
v[ 1] += nonce;
v[13] = ROR8(v[13] ^ v[ 1]);
v[ 9] += v[13];
v[ 5] = ROTR32(v[ 5] ^ v[ 9], 7);
v[ 1]+= v[ 6];
v[ 0]+= v[ 5];
v[13] = ROL16(v[13] ^ v[ 2]); v[12] = ROL16(v[12] ^ v[ 1]); v[15] = ROL16(v[15] ^ v[ 0]);
v[ 8]+= v[13]; v[11]+= v[12]; v[ 9]+= v[14]; v[10]+= v[15];
v[ 7] = ROTR32(v[ 7] ^ v[ 8], 12); v[ 6] = ROTR32(v[ 6] ^ v[11], 12); v[ 4] = ROTR32(v[ 4] ^ v[ 9], 12); v[ 5] = ROTR32(v[ 5] ^ v[10], 12);
v[ 2]+= v[ 7]; v[ 1]+= v[ 6]; v[ 3]+= v[ 4]; v[ 0]+= v[ 5];
v[13] = ROR8(v[13] ^ v[ 2]); v[12] = ROR8(v[12] ^ v[ 1]); v[14] = ROR8(v[14] ^ v[ 3]); v[15] = ROR8(v[15] ^ v[ 0]);
v[ 8]+= v[13]; v[11]+= v[12]; v[ 9]+= v[14]; v[10]+= v[15];
v[ 6] = ROTR32(v[ 6] ^ v[11], 7); v[ 7] = ROTR32(v[ 7] ^ v[8], 7); v[ 4] = ROTR32(v[ 4] ^ v[ 9], 7); v[ 5] = ROTR32(v[ 5] ^ v[10], 7);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],0,0, v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],m[ 1],0, v[ 1],v[ 6],v[11],v[12],m[ 0],m[ 2], v[ 2],v[ 7],v[ 8],v[13],0,0, v[ 3],v[ 4],v[ 9],v[14],0,nonce);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],0,m[ 0], v[ 2],v[ 6],v[10],v[14],0,m[ 2], v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],0,0, v[ 1],v[ 6],v[11],v[12],nonce,0, v[ 2],v[ 7],v[ 8],v[13],0,m[ 1], v[ 3],v[ 4],v[ 9],v[14],0,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],nonce,m[ 1], v[ 2],v[ 6],v[10],v[14],0,0, v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],m[ 2],0, v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],0,m[ 0], v[ 3],v[ 4],v[ 9],v[14],0,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,m[ 0], v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],m[ 2],0, v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],0,m[ 1], v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],0,0, v[ 3],v[ 4],v[ 9],v[14],nonce,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],m[ 2],0, v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],m[ 0],0, v[ 3],v[ 7],v[11],v[15],0,nonce);
GS4(v[ 0],v[ 5],v[10],v[15],0,0, v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],0,0, v[ 3],v[ 4],v[ 9],v[14],m[ 1],0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],m[ 1],0, v[ 2],v[ 6],v[10],v[14],0,0, v[ 3],v[ 7],v[11],v[15],0,0);
GS4(v[ 0],v[ 5],v[10],v[15],m[ 0],0, v[ 1],v[ 6],v[11],v[12],0,nonce, v[ 2],v[ 7],v[ 8],v[13],0,m[ 2], v[ 3],v[ 4],v[ 9],v[14],0,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],0,m[ 1], v[ 3],v[ 7],v[11],v[15],nonce,0);
GS4(v[ 0],v[ 5],v[10],v[15],0,m[ 0], v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],0,0, v[ 3],v[ 4],v[ 9],v[14],m[ 2],0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,0, v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],0,nonce, v[ 3],v[ 7],v[11],v[15],m[ 0],0);
GS4(v[ 0],v[ 5],v[10],v[15],0,m[ 2], v[ 1],v[ 6],v[11],v[12],0,0, v[ 2],v[ 7],v[ 8],v[13],m[ 1],0, v[ 3],v[ 4],v[ 9],v[14],0,0);
GS4(v[ 0],v[ 4],v[ 8],v[12],0,m[ 2], v[ 1],v[ 5],v[ 9],v[13],0,0, v[ 2],v[ 6],v[10],v[14],0,0, v[ 3],v[ 7],v[11],v[15],m[ 1],0);
v[ 0] += v[ 5];
v[ 2] += v[ 7] + nonce;
v[15] = ROL16(v[15] ^ v[ 0]);
v[13] = ROL16(v[13] ^ v[ 2]);
v[10] += v[15];
v[ 8] += v[13];
v[ 5] = ROTR32(v[ 5] ^ v[10], 12);
v[ 7] = ROTR32(v[ 7] ^ v[ 8], 12);
v[ 0] += v[ 5];
v[ 2] += v[ 7];
v[15] = ROTR32(v[15] ^ v[ 0],1);
v[13] = ROR8(v[13] ^ v[ 2]);
v[ 8] += v[13];
if(xor3x(v[ 7],h7,v[ 8])==v[15]){
uint32_t pos = atomicInc(&resNonce[0],0xffffffff)+1;
if(pos < maxResults)
resNonce[pos]=nonce;
return;
}
}
}
static void blake2s_setBlock(const uint32_t* input,const uint32_t ptarget7)
{
uint32_t _ALIGN(64) m[16];
uint32_t _ALIGN(64) v[16];
uint32_t _ALIGN(64) h[21];
// COMPRESS
for(int i = 0; i < 16; ++i )
m[i] = input[i];
h[0] = 0x01010020 ^ blake2s_IV[0];
h[1] = blake2s_IV[1];
h[2] = blake2s_IV[2]; h[3] = blake2s_IV[3];
h[4] = blake2s_IV[4]; h[5] = blake2s_IV[5];
h[6] = blake2s_IV[6]; h[7] = blake2s_IV[7];
for(int i = 0; i < 8; ++i )
v[i] = h[i];
v[ 8] = blake2s_IV[0]; v[ 9] = blake2s_IV[1];
v[10] = blake2s_IV[2]; v[11] = blake2s_IV[3];
v[12] = 64 ^ blake2s_IV[4]; v[13] = blake2s_IV[5];
v[14] = blake2s_IV[6]; v[15] = blake2s_IV[7];
ROUND( 0 ); ROUND( 1 );
ROUND( 2 ); ROUND( 3 );
ROUND( 4 ); ROUND( 5 );
ROUND( 6 ); ROUND( 7 );
ROUND( 8 ); ROUND( 9 );
for(int i = 0; i < 8; ++i )
h[i] ^= v[i] ^ v[i + 8];
h[16] = input[16];
h[17] = input[17];
h[18] = input[18];
h[ 8] = 0x6A09E667; h[ 9] = 0xBB67AE85;
h[10] = 0x3C6EF372; h[11] = 0xA54FF53A;
h[12] = 0x510E522F; h[13] = 0x9B05688C;
h[14] =~0x1F83D9AB; h[15] = 0x5BE0CD19;
h[ 0]+= h[ 4] + h[16];
h[12] = SPH_ROTR32(h[12] ^ h[ 0],16);
h[ 8]+= h[12];
h[ 4] = SPH_ROTR32(h[ 4] ^ h[ 8],12);
h[ 0]+= h[ 4] + h[17];
h[12] = SPH_ROTR32(h[12] ^ h[ 0],8);
h[ 8]+= h[12];
h[ 4] = SPH_ROTR32(h[ 4] ^ h[ 8],7);
h[ 1]+= h[ 5] + h[18];
h[13] = SPH_ROTR32(h[13] ^ h[ 1], 16);
h[ 9]+= h[13];
h[ 5] = ROTR32(h[ 5] ^ h[ 9], 12);
h[ 2]+= h[ 6];
h[14] = SPH_ROTR32(h[14] ^ h[ 2],16);
h[10]+= h[14];
h[ 6] = SPH_ROTR32(h[ 6] ^ h[10], 12);
h[ 2]+= h[ 6];
h[14] = SPH_ROTR32(h[14] ^ h[ 2],8);
h[10]+= h[14];
h[ 6] = SPH_ROTR32(h[ 6] ^ h[10], 7);
h[19] = h[7]; //constant h[7] for nonce check
h[ 3]+= h[ 7];
h[15] = SPH_ROTR32(h[15] ^ h[ 3],16);
h[11]+= h[15];
h[ 7] = SPH_ROTR32(h[ 7] ^ h[11], 12);
h[ 3]+= h[ 7];
h[15] = SPH_ROTR32(h[15] ^ h[ 3],8);
h[11]+= h[15];
h[ 7] = SPH_ROTR32(h[ 7] ^ h[11], 7);
h[ 1]+= h[ 5];
h[ 3]+= h[ 4];
h[14] = SPH_ROTR32(h[14] ^ h[ 3],16);
h[ 2]+= h[ 7];
if(ptarget7==0){
h[19] = SPH_ROTL32(h[19],7); //align the rotation with v[7] v[15];
}
cudaMemcpyToSymbol(midstate, h, 20*sizeof(uint32_t), 0, cudaMemcpyHostToDevice);
}
static bool init[MAX_GPUS] = { 0 };
extern "C" int scanhash_blake2s(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done)
{
uint32_t _ALIGN(64) endiandata[20];
uint32_t *pdata = work->data;
uint32_t *ptarget = work->target;
uint32_t *resNonces;
const uint32_t first_nonce = pdata[19];
const int dev_id = device_map[thr_id];
int rc = 0;
int intensity = is_windows() ? 25 : 28;
uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity);
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
const dim3 grid((throughput + (NPT*TPB)-1)/(NPT*TPB));
const dim3 block(TPB);
if (!init[thr_id])
{
cudaSetDevice(dev_id);
if (opt_cudaschedule == -1 && gpu_threads == 1) {
cudaDeviceReset();
// reduce cpu usage (linux)
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
CUDA_LOG_ERROR();
}
gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);
cuda_get_arch(thr_id);
CUDA_CALL_OR_RET_X(cudaMalloc(&d_resNonce[thr_id], maxResults * sizeof(uint32_t)), -1);
CUDA_CALL_OR_RET_X(cudaMallocHost(&h_resNonce[thr_id], maxResults * sizeof(uint32_t)), -1);
init[thr_id] = true;
}
resNonces = h_resNonce[thr_id];
for (int i=0; i < 19; i++) {
be32enc(&endiandata[i], pdata[i]);
}
blake2s_setBlock(endiandata,ptarget[7]);
cudaMemset(d_resNonce[thr_id], 0x00, maxResults*sizeof(uint32_t));
do {
if(ptarget[7]) {
blake2s_gpu_hash_nonce<<<grid, block>>>(throughput,pdata[19],d_resNonce[thr_id],ptarget[7]);
} else {
blake2s_gpu_hash_nonce<<<grid, block>>>(throughput,pdata[19],d_resNonce[thr_id]);
}
cudaMemcpy(resNonces, d_resNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost);
if(resNonces[0])
{
cudaMemcpy(resNonces, d_resNonce[thr_id], maxResults*sizeof(uint32_t), cudaMemcpyDeviceToHost);
cudaMemset(d_resNonce[thr_id], 0x00, sizeof(uint32_t));
if(resNonces[0] >= maxResults) {
gpulog(LOG_WARNING, thr_id, "candidates flood: %u", resNonces[0]);
resNonces[0] = maxResults-1;
}
uint32_t vhashcpu[8];
uint32_t nonce = sph_bswap32(resNonces[1]);
be32enc(&endiandata[19], nonce);
blake2s_hash(vhashcpu, endiandata);
*hashes_done = pdata[19] - first_nonce + throughput;
if(vhashcpu[6] <= ptarget[6] && fulltest(vhashcpu, ptarget))
{
work_set_target_ratio(work, vhashcpu);
work->nonces[0] = nonce;
rc = work->valid_nonces = 1;
// search for 2nd best nonce
for(uint32_t j=2; j <= resNonces[0]; j++)
{
nonce = sph_bswap32(resNonces[j]);
be32enc(&endiandata[19], nonce);
blake2s_hash(vhashcpu, endiandata);
if(vhashcpu[6] <= ptarget[6] && fulltest(vhashcpu, ptarget))
{
gpulog(LOG_DEBUG, thr_id, "Multiple nonces: 1/%08x - %u/%08x", work->nonces[0], j, nonce);
work->nonces[1] = nonce;
if (bn_hash_target_ratio(vhashcpu, ptarget) > work->shareratio[0]) {
work->shareratio[1] = work->shareratio[0];
work->sharediff[1] = work->sharediff[0];
xchg(work->nonces[1], work->nonces[0]);
work_set_target_ratio(work, vhashcpu);
} else if (work->valid_nonces == 1) {
bn_set_target_ratio(work, vhashcpu, 1);
}
work->valid_nonces++;
rc = 2;
break;
}
}
pdata[19] = max(work->nonces[0], work->nonces[1]); // next scan start
return rc;
} else if (vhashcpu[7] > ptarget[7]) {
gpu_increment_reject(thr_id);
if (!opt_quiet)
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", resNonces[0]);
}
}
pdata[19] += throughput;
} while (!work_restart[thr_id].restart && (uint64_t)max_nonce > (uint64_t)throughput + pdata[19]);
*hashes_done = pdata[19] - first_nonce;
return rc;
}
// cleanup
extern "C" void free_blake2s(int thr_id)
{
if (!init[thr_id])
return;
cudaDeviceSynchronize();
cudaFreeHost(h_resNonce[thr_id]);
cudaFree(d_resNonce[thr_id]);
init[thr_id] = false;
cudaDeviceSynchronize();
}