GOSTSec
/
ccminer-gostd-lite
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity
Browse Source

~10% speedup

master
pallas1 9 years ago committed by Tanguy Pruvot
parent
be1f64446a
commit
ebf885d482
1 changed files with 138 additions and 118 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 238
      Algo256/decred.cu

238
Algo256/decred.cu
Unescape View File

@ -2,6 +2,8 @@
* Blake-256 Decred 180-Bytes input Cuda Kernel (Tested on SM 5/5.2) * Blake-256 Decred 180-Bytes input Cuda Kernel (Tested on SM 5/5.2)
* *
* Tanguy Pruvot - Feb 2016 * Tanguy Pruvot - Feb 2016
*
* Revised for optimisation by pallas @ bitcointalk - Apr 2016
*/ */
#include <stdint.h> #include <stdint.h>
@ -14,7 +16,9 @@ extern "C" {
} }
/* threads per block */ /* threads per block */
#define TPB 256 #define TPB 512
/* nonces per round */
#define NPR 128
/* hash by cpu with blake 256 */ /* hash by cpu with blake 256 */
extern "C" void decred_hash(void *output, const void *input) extern "C" void decred_hash(void *output, const void *input)
@ -46,6 +50,7 @@ static uint32_t *h_resNonce[MAX_GPUS];
static uint32_t extra_results[NBN] = { UINT32_MAX }; static uint32_t extra_results[NBN] = { UINT32_MAX };
#endif #endif
/* ############################################################################################################################### */ /* ############################################################################################################################### */
#define GSPREC(a,b,c,d,x,y) { \ #define GSPREC(a,b,c,d,x,y) { \
@ -59,107 +64,106 @@ static uint32_t extra_results[NBN] = { UINT32_MAX };
v[b] = SPH_ROTR32(v[b] ^ v[c], 7); \ v[b] = SPH_ROTR32(v[b] ^ v[c], 7); \
} }
__device__ __forceinline__
void blake256_compress_14(uint32_t *h, const uint32_t nonce, const uint32_t T0)
{
uint32_t v[16];
#pragma unroll 8 #define GSPREC4(a0,b0,c0,d0,x0,y0,a1,b1,c1,d1,x1,y1,a2,b2,c2,d2,x2,y2,a3,b3,c3,d3,x3,y3) { \
for(uint32_t i = 0; i < 8; i++) v[a0] += (m[x0] ^ c_u256[y0]) + v[b0]; \
v[i] = h[i]; v[a1] += (m[x1] ^ c_u256[y1]) + v[b1]; \
v[a2] += (m[x2] ^ c_u256[y2]) + v[b2]; \
v[a3] += (m[x3] ^ c_u256[y3]) + v[b3]; \
v[d0] = __byte_perm(v[d0] ^ v[a0], 0, 0x1032); \
v[d1] = __byte_perm(v[d1] ^ v[a1], 0, 0x1032); \
v[d2] = __byte_perm(v[d2] ^ v[a2], 0, 0x1032); \
v[d3] = __byte_perm(v[d3] ^ v[a3], 0, 0x1032); \
v[c0] += v[d0]; \
v[c1] += v[d1]; \
v[c2] += v[d2]; \
v[c3] += v[d3]; \
v[b0] = SPH_ROTR32(v[b0] ^ v[c0], 12); \
v[b1] = SPH_ROTR32(v[b1] ^ v[c1], 12); \
v[b2] = SPH_ROTR32(v[b2] ^ v[c2], 12); \
v[b3] = SPH_ROTR32(v[b3] ^ v[c3], 12); \
v[a0] += (m[y0] ^ c_u256[x0]) + v[b0]; \
v[a1] += (m[y1] ^ c_u256[x1]) + v[b1]; \
v[a2] += (m[y2] ^ c_u256[x2]) + v[b2]; \
v[a3] += (m[y3] ^ c_u256[x3]) + v[b3]; \
v[d0] = __byte_perm(v[d0] ^ v[a0], 0, 0x0321); \
v[d1] = __byte_perm(v[d1] ^ v[a1], 0, 0x0321); \
v[d2] = __byte_perm(v[d2] ^ v[a2], 0, 0x0321); \
v[d3] = __byte_perm(v[d3] ^ v[a3], 0, 0x0321); \
v[c0] += v[d0]; \
v[c1] += v[d1]; \
v[c2] += v[d2]; \
v[c3] += v[d3]; \
v[b0] = SPH_ROTR32(v[b0] ^ v[c0], 7); \
v[b1] = SPH_ROTR32(v[b1] ^ v[c1], 7); \
v[b2] = SPH_ROTR32(v[b2] ^ v[c2], 7); \
v[b3] = SPH_ROTR32(v[b3] ^ v[c3], 7); \
}
const uint32_t c_u256[16] = { static const __constant__ uint32_t c_u256[16] = {
0x243F6A88, 0x85A308D3, 0x13198A2E, 0x03707344, 0x243F6A88, 0x85A308D3, 0x13198A2E, 0x03707344,
0xA4093822, 0x299F31D0, 0x082EFA98, 0xEC4E6C89, 0xA4093822, 0x299F31D0, 0x082EFA98, 0xEC4E6C89,
0x452821E6, 0x38D01377, 0xBE5466CF, 0x34E90C6C, 0x452821E6, 0x38D01377, 0xBE5466CF, 0x34E90C6C,
0xC0AC29B7, 0xC97C50DD, 0x3F84D5B5, 0xB5470917 0xC0AC29B7, 0xC97C50DD, 0x3F84D5B5, 0xB5470917
}; };
v[ 8] = c_u256[0];
v[ 9] = c_u256[1];
v[10] = c_u256[2];
v[11] = c_u256[3];
v[12] = c_u256[4] ^ T0;
v[13] = c_u256[5] ^ T0;
v[14] = c_u256[6];
v[15] = c_u256[7];
uint32_t m[16];
m[0] = d_data[8]; __device__ __forceinline__
m[1] = d_data[9]; uint32_t blake256_compress_14(uint32_t *m, uint32_t *v_init, uint32_t d_data6, uint32_t d_data7)
m[2] = d_data[10]; {
m[3] = nonce; uint32_t v[16];
#pragma unroll #pragma unroll
for (uint32_t i = 4; i < 16; i++) { for (uint32_t i = 0; i < 16; i++) v[i] = v_init[i];
m[i] = d_data[i+8U]; // these two are not modified:
} v[ 9] = 0x85A308D3;
v[13] = 0x299F31D0 ^ (180U*8U);
// round 1 // round 1 with nonce
GSPREC(0, 4, 0x8, 0xC, 0, 1);
GSPREC(1, 5, 0x9, 0xD, 2, 3); GSPREC(1, 5, 0x9, 0xD, 2, 3);
GSPREC(2, 6, 0xA, 0xE, 4, 5);
GSPREC(3, 7, 0xB, 0xF, 6, 7);
GSPREC(0, 5, 0xA, 0xF, 8, 9); GSPREC(0, 5, 0xA, 0xF, 8, 9);
GSPREC(1, 6, 0xB, 0xC, 10, 11); GSPREC(1, 6, 0xB, 0xC, 10, 11);
GSPREC(2, 7, 0x8, 0xD, 12, 13); GSPREC(2, 7, 0x8, 0xD, 12, 13);
GSPREC(3, 4, 0x9, 0xE, 14, 15); GSPREC(3, 4, 0x9, 0xE, 14, 15);
// round 2 // round 2
GSPREC(0, 4, 0x8, 0xC, 14, 10); GSPREC4(0, 4, 0x8, 0xC, 14, 10, 1, 5, 0x9, 0xD, 4, 8, 2, 6, 0xA, 0xE, 9, 15, 3, 7, 0xB, 0xF, 13, 6);
GSPREC(1, 5, 0x9, 0xD, 4, 8); GSPREC4(0, 5, 0xA, 0xF, 1, 12, 1, 6, 0xB, 0xC, 0, 2, 2, 7, 0x8, 0xD, 11, 7, 3, 4, 0x9, 0xE, 5, 3);
GSPREC(2, 6, 0xA, 0xE, 9, 15);
GSPREC(3, 7, 0xB, 0xF, 13, 6);
GSPREC(0, 5, 0xA, 0xF, 1, 12);
GSPREC(1, 6, 0xB, 0xC, 0, 2);
GSPREC(2, 7, 0x8, 0xD, 11, 7);
GSPREC(3, 4, 0x9, 0xE, 5, 3);
// round 3 // round 3
GSPREC(0, 4, 0x8, 0xC, 11, 8); GSPREC4(0, 4, 0x8, 0xC, 11, 8, 1, 5, 0x9, 0xD, 12, 0, 2, 6, 0xA, 0xE, 5, 2, 3, 7, 0xB, 0xF, 15, 13);
GSPREC(1, 5, 0x9, 0xD, 12, 0); GSPREC4(0, 5, 0xA, 0xF, 10, 14, 1, 6, 0xB, 0xC, 3, 6, 2, 7, 0x8, 0xD, 7, 1, 3, 4, 0x9, 0xE, 9, 4);
GSPREC(2, 6, 0xA, 0xE, 5, 2);
GSPREC(3, 7, 0xB, 0xF, 15, 13);
GSPREC(0, 5, 0xA, 0xF, 10, 14);
GSPREC(1, 6, 0xB, 0xC, 3, 6);
GSPREC(2, 7, 0x8, 0xD, 7, 1);
GSPREC(3, 4, 0x9, 0xE, 9, 4);
// round 4 // round 4
GSPREC(0, 4, 0x8, 0xC, 7, 9); GSPREC4(0, 4, 0x8, 0xC, 7, 9, 1, 5, 0x9, 0xD, 3, 1, 2, 6, 0xA, 0xE, 13, 12, 3, 7, 0xB, 0xF, 11, 14);
GSPREC(1, 5, 0x9, 0xD, 3, 1); GSPREC4(0, 5, 0xA, 0xF, 2, 6, 1, 6, 0xB, 0xC, 5, 10, 2, 7, 0x8, 0xD, 4, 0, 3, 4, 0x9, 0xE, 15, 8);
GSPREC(2, 6, 0xA, 0xE, 13, 12);
GSPREC(3, 7, 0xB, 0xF, 11, 14);
GSPREC(0, 5, 0xA, 0xF, 2, 6);
GSPREC(1, 6, 0xB, 0xC, 5, 10);
GSPREC(2, 7, 0x8, 0xD, 4, 0);
GSPREC(3, 4, 0x9, 0xE, 15, 8);
// round 5 // round 5
GSPREC(0, 4, 0x8, 0xC, 9, 0); GSPREC4(0, 4, 0x8, 0xC, 9, 0, 1, 5, 0x9, 0xD, 5, 7, 2, 6, 0xA, 0xE, 2, 4, 3, 7, 0xB, 0xF, 10, 15);
GSPREC(1, 5, 0x9, 0xD, 5, 7); GSPREC4(0, 5, 0xA, 0xF, 14, 1, 1, 6, 0xB, 0xC, 11, 12, 2, 7, 0x8, 0xD, 6, 8, 3, 4, 0x9, 0xE, 3, 13);
GSPREC(2, 6, 0xA, 0xE, 2, 4);
GSPREC(3, 7, 0xB, 0xF, 10, 15);
GSPREC(0, 5, 0xA, 0xF, 14, 1);
GSPREC(1, 6, 0xB, 0xC, 11, 12);
GSPREC(2, 7, 0x8, 0xD, 6, 8);
GSPREC(3, 4, 0x9, 0xE, 3, 13);
// round 6 // round 6
GSPREC(0, 4, 0x8, 0xC, 2, 12); GSPREC4(0, 4, 0x8, 0xC, 2, 12, 1, 5, 0x9, 0xD, 6, 10, 2, 6, 0xA, 0xE, 0, 11, 3, 7, 0xB, 0xF, 8, 3);
GSPREC(1, 5, 0x9, 0xD, 6, 10); GSPREC4(0, 5, 0xA, 0xF, 4, 13, 1, 6, 0xB, 0xC, 7, 5, 2, 7, 0x8, 0xD, 15,14, 3, 4, 0x9, 0xE, 1, 9);
GSPREC(2, 6, 0xA, 0xE, 0, 11);
GSPREC(3, 7, 0xB, 0xF, 8, 3);
GSPREC(0, 5, 0xA, 0xF, 4, 13);
GSPREC(1, 6, 0xB, 0xC, 7, 5);
GSPREC(2, 7, 0x8, 0xD, 15,14);
GSPREC(3, 4, 0x9, 0xE, 1, 9);
// round 7 // round 7
GSPREC(0, 4, 0x8, 0xC, 12, 5); GSPREC4(0, 4, 0x8, 0xC, 12, 5, 1, 5, 0x9, 0xD, 1, 15, 2, 6, 0xA, 0xE, 14,13, 3, 7, 0xB, 0xF, 4, 10);
GSPREC(1, 5, 0x9, 0xD, 1, 15); GSPREC4(0, 5, 0xA, 0xF, 0, 7, 1, 6, 0xB, 0xC, 6, 3, 2, 7, 0x8, 0xD, 9, 2, 3, 4, 0x9, 0xE, 8, 11);
GSPREC(2, 6, 0xA, 0xE, 14,13); /*
GSPREC(3, 7, 0xB, 0xF, 4, 10); // round 8
GSPREC(0, 5, 0xA, 0xF, 0, 7); GSPREC4(0, 4, 0x8, 0xC, 13,11, 1, 5, 0x9, 0xD, 7, 14, 2, 6, 0xA, 0xE, 12, 1, 3, 7, 0xB, 0xF, 3, 9);
GSPREC(1, 6, 0xB, 0xC, 6, 3); GSPREC4(0, 5, 0xA, 0xF, 5, 0, 1, 6, 0xB, 0xC, 15, 4, 2, 7, 0x8, 0xD, 8, 6, 3, 4, 0x9, 0xE, 2, 10);
GSPREC(2, 7, 0x8, 0xD, 9, 2); // round 9
GSPREC(3, 4, 0x9, 0xE, 8, 11); GSPREC4(0, 4, 0x8, 0xC, 6, 15, 1, 5, 0x9, 0xD, 14, 9, 2, 6, 0xA, 0xE, 11, 3, 3, 7, 0xB, 0xF, 0, 8);
GSPREC4(0, 5, 0xA, 0xF, 12, 2, 1, 6, 0xB, 0xC, 13, 7, 2, 7, 0x8, 0xD, 1, 4, 3, 4, 0x9, 0xE, 10, 5);
// round 10
GSPREC4(0, 4, 0x8, 0xC, 10, 2, 1, 5, 0x9, 0xD, 8, 4, 2, 6, 0xA, 0xE, 7, 6, 3, 7, 0xB, 0xF, 1, 5);
GSPREC4(0, 5, 0xA, 0xF, 15,11, 1, 6, 0xB, 0xC, 9, 14, 2, 7, 0x8, 0xD, 3, 12, 3, 4, 0x9, 0xE, 13, 0);
// round 11
GSPREC4(0, 4, 0x8, 0xC, 0, 1, 1, 5, 0x9, 0xD, 2, 3, 2, 6, 0xA, 0xE, 4, 5, 3, 7, 0xB, 0xF, 6, 7);
GSPREC4(0, 5, 0xA, 0xF, 8, 9, 1, 6, 0xB, 0xC, 10,11, 2, 7, 0x8, 0xD, 12,13, 3, 4, 0x9, 0xE, 14,15);
// round 12
GSPREC4(0, 4, 0x8, 0xC, 14,10, 1, 5, 0x9, 0xD, 4, 8, 2, 6, 0xA, 0xE, 9, 15, 3, 7, 0xB, 0xF, 13, 6);
GSPREC4(0, 5, 0xA, 0xF, 1, 12, 1, 6, 0xB, 0xC, 0, 2, 2, 7, 0x8, 0xD, 11, 7, 3, 4, 0x9, 0xE, 5, 3);
// round 13
GSPREC4(0, 4, 0x8, 0xC, 11, 8, 1, 5, 0x9, 0xD, 12, 0, 2, 6, 0xA, 0xE, 5, 2, 3, 7, 0xB, 0xF, 15,13);
GSPREC4(0, 5, 0xA, 0xF, 10,14, 1, 6, 0xB, 0xC, 3, 6, 2, 7, 0x8, 0xD, 7, 1, 3, 4, 0x9, 0xE, 9, 4);
*/
// round 8 // round 8
GSPREC(0, 4, 0x8, 0xC, 13,11); GSPREC(0, 4, 0x8, 0xC, 13,11);
GSPREC(1, 5, 0x9, 0xD, 7, 14); GSPREC(1, 5, 0x9, 0xD, 7, 14);
@ -220,68 +224,81 @@ void blake256_compress_14(uint32_t *h, const uint32_t nonce, const uint32_t T0)
GSPREC(2, 6, 0xA, 0xE, 13,12); GSPREC(2, 6, 0xA, 0xE, 13,12);
GSPREC(3, 7, 0xB, 0xF, 11,14); GSPREC(3, 7, 0xB, 0xF, 11,14);
GSPREC(0, 5, 0xA, 0xF, 2, 6); GSPREC(0, 5, 0xA, 0xF, 2, 6);
GSPREC(1, 6, 0xB, 0xC, 5, 10);
GSPREC(2, 7, 0x8, 0xD, 4, 0); GSPREC(2, 7, 0x8, 0xD, 4, 0);
//GSPREC(3, 4, 0x9, 0xE, 15, 8);
if ((d_data7 ^ v[7] ^ v[15]) == 0) {
v[3] += (m[15] ^ c_u256[8]) + v[4]; GSPREC(1, 6, 0xB, 0xC, 5, 10);
v[14] = __byte_perm(v[14] ^ v[3], 0, 0x1032); GSPREC(3, 4, 0x9, 0xE, 15, 8);
v[9] += v[14]; \ return (d_data6 ^ v[6] ^ v[14]);
v[4] = SPH_ROTR32(v[4] ^ v[9], 12);
v[3] += (m[8] ^ c_u256[15]) + v[4];
v[14] = __byte_perm(v[14] ^ v[3], 0, 0x0321);
// only compute h6 & 7
h[6] ^= v[6] ^ v[14];
h[7] ^= v[7] ^ v[15];
} }
return UINT32_MAX;
}
/* ############################################################################################################################### */ /* ############################################################################################################################### */
// ------ Close: Last 52/64 bytes ------
__global__ __global__
void blake256_gpu_hash_nonce(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce, const uint64_t highTarget) void blake256_gpu_hash_nonce(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce, const uint64_t highTarget)
{ {
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads) // if (thread < threads)
{ {
const uint32_t nonce = startNonce + thread; const uint32_t nonce = startNonce + thread * NPR;
uint32_t h[8]; uint32_t m[16], v[16], temp;
const uint32_t d_data6 = d_data[6], d_data7 = d_data[7];
#pragma unroll #pragma unroll
for(int i=0; i < 8; i++) { for(int i = 0; i < 8; i++) v[i] = d_data[i];
h[i] = d_data[i];
}
// ------ Close: Last 52/64 bytes ------ #pragma unroll
for (uint32_t i = 0; i < 16; i++) m[i] = d_data[i+8U];
v[ 8] = 0x243F6A88;
v[ 9] = 0x85A308D3;
v[10] = 0x13198A2E;
v[11] = 0x03707344;
blake256_compress_14(h, nonce, (180U*8U)); v[12] = 0xA4093822 ^ (180U*8U);
v[13] = 0x299F31D0 ^ (180U*8U);
v[14] = 0x082EFA98;
v[15] = 0xEC4E6C89;
// round 1 without nonce
GSPREC(0, 4, 0x8, 0xC, 0, 1);
GSPREC(2, 6, 0xA, 0xE, 4, 5);
GSPREC(3, 7, 0xB, 0xF, 6, 7);
if (h[7] == 0 && cuda_swab32(h[6]) <= highTarget) { for (m[3] = nonce; m[3] < nonce + NPR; m[3]++) {
temp = blake256_compress_14(m, v, d_data6, d_data7);
if (temp != UINT32_MAX && cuda_swab32(temp) <= highTarget) {
#if NBN == 2 #if NBN == 2
if (resNonce[0] != UINT32_MAX) if (resNonce[0] != UINT32_MAX) resNonce[1] = m[3];
resNonce[1] = nonce; else resNonce[0] = m[3];
else
resNonce[0] = nonce;
#else #else
resNonce[0] = nonce; resNonce[0] = m[3];
#endif #endif
} }
} }
} }
}
__host__ __host__
static uint32_t decred_cpu_hash_nonce(const int thr_id, const uint32_t threads, const uint32_t startNonce, const uint64_t highTarget) static uint32_t decred_cpu_hash_nonce(const int thr_id, const uint32_t threads, const uint32_t startNonce, const uint64_t highTarget)
{ {
uint32_t result = UINT32_MAX; uint32_t result = UINT32_MAX;
const uint32_t real_threads = threads / NPR;
dim3 grid((threads + TPB-1)/TPB); dim3 grid((real_threads + TPB-1)/TPB);
dim3 block(TPB); dim3 block(TPB);
/* Check error on Ctrl+C or kill to prevent segfaults on exit */ /* Check error on Ctrl+C or kill to prevent segfaults on exit */
if (cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t)) != cudaSuccess) if (cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t)) != cudaSuccess)
return result; return result;
blake256_gpu_hash_nonce <<<grid, block>>> (threads, startNonce, d_resNonce[thr_id], highTarget); blake256_gpu_hash_nonce <<<grid, block>>> (real_threads, startNonce, d_resNonce[thr_id], highTarget);
cudaThreadSynchronize(); cudaThreadSynchronize();
if (cudaSuccess == cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) { if (cudaSuccess == cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) {
@ -294,6 +311,7 @@ static uint32_t decred_cpu_hash_nonce(const int thr_id, const uint32_t threads,
return result; return result;
} }
__host__ __host__
static void decred_midstate_128(uint32_t *output, const uint32_t *input) static void decred_midstate_128(uint32_t *output, const uint32_t *input)
{ {
@ -307,6 +325,7 @@ static void decred_midstate_128(uint32_t *output, const uint32_t *input)
memcpy(output, (void*)ctx.H, 32); memcpy(output, (void*)ctx.H, 32);
} }
__host__ __host__
void decred_cpu_setBlock_52(uint32_t *penddata, const uint32_t *midstate, const uint32_t *ptarget) void decred_cpu_setBlock_52(uint32_t *penddata, const uint32_t *midstate, const uint32_t *ptarget)
{ {
@ -321,6 +340,7 @@ void decred_cpu_setBlock_52(uint32_t *penddata, const uint32_t *midstate, const
CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_data, data, 32 + 64, 0, cudaMemcpyHostToDevice)); CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_data, data, 32 + 64, 0, cudaMemcpyHostToDevice));
} }
/* ############################################################################################################################### */ /* ############################################################################################################################### */
static bool init[MAX_GPUS] = { 0 }; static bool init[MAX_GPUS] = { 0 };
@ -425,6 +445,7 @@ extern "C" int scanhash_decred(int thr_id, struct work* work, uint32_t max_nonce
return rc; return rc;
} }
// cleanup // cleanup
extern "C" void free_decred(int thr_id) extern "C" void free_decred(int thr_id)
{ {
@ -440,4 +461,3 @@ extern "C" void free_decred(int thr_id)
cudaDeviceSynchronize(); cudaDeviceSynchronize();
} }

Write Preview
Loading…
Cancel
Save