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

blake: merge sp improvements, start 1.7.2 dev.. 

to be tested on old arch too...
master
Tanguy Pruvot 9 years ago
parent
a237601747
commit
4a7e239d7c
3 changed files with 350 additions and 242 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. 584
      Algo256/blake256.cu
  2. 2
      configure.ac
  3. 6
      cpuminer-config.h

584
Algo256/blake256.cu
Unescape View File

@ -1,23 +1,20 @@
/** /**
* Blake-256 Cuda Kernel (Tested on SM 5.0) * Blake-256 Cuda Kernel (Tested on SM 5/5.2)
* *
* Tanguy Pruvot - Nov. 2014 * Tanguy Pruvot / SP - Jan 2016
*/ */
#define PRECALC64 1 #include <stdint.h>
#include <memory.h>
#include "miner.h" #include "miner.h"
extern "C" { extern "C" {
#include "sph/sph_blake.h" #include "sph/sph_blake.h"
//extern int blake256_rounds;
} }
#include <stdint.h> /* threads per block */
#include <memory.h> #define TPB 512
/* threads per block and throughput (intensity) */
#define TPB 128
/* hash by cpu with blake 256 */ /* hash by cpu with blake 256 */
extern "C" void blake256hash(void *output, const void *input, int8_t rounds = 14) extern "C" void blake256hash(void *output, const void *input, int8_t rounds = 14)
@ -36,16 +33,7 @@ extern "C" void blake256hash(void *output, const void *input, int8_t rounds = 14
#include "cuda_helper.h" #include "cuda_helper.h"
#if PRECALC64
__constant__ uint32_t _ALIGN(32) d_data[12]; __constant__ uint32_t _ALIGN(32) d_data[12];
#else
__constant__ static uint32_t _ALIGN(32) c_data[20];
/* midstate hash cache, this algo is run on 2 parts */
__device__ static uint32_t cache[8];
__device__ static uint32_t prevsum = 0;
/* crc32.c */
extern "C" uint32_t crc32_u32t(const uint32_t *buf, size_t size);
#endif
/* 8 adapters max */ /* 8 adapters max */
static uint32_t *d_resNonce[MAX_GPUS]; static uint32_t *d_resNonce[MAX_GPUS];
@ -55,73 +43,19 @@ static uint32_t *h_resNonce[MAX_GPUS];
#define NBN 2 #define NBN 2
static uint32_t extra_results[NBN] = { UINT32_MAX }; static uint32_t extra_results[NBN] = { UINT32_MAX };
/* prefer uint32_t to prevent size conversions = speed +5/10 % */ #define GSPREC(a,b,c,d,x,y) { \
__constant__ v[a] += (m[x] ^ c_u256[y]) + v[b]; \
static uint32_t _ALIGN(32) c_sigma[16][16] = { v[d] = __byte_perm(v[d] ^ v[a],0, 0x1032); \
{ 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 },
{ 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 }
};
#if !PRECALC64
__device__ __constant__
static const uint32_t __align__(32) c_IV256[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)
};
#endif
__device__ __constant__
static const uint32_t __align__(32) c_u256[16] = {
SPH_C32(0x243F6A88), SPH_C32(0x85A308D3),
SPH_C32(0x13198A2E), SPH_C32(0x03707344),
SPH_C32(0xA4093822), SPH_C32(0x299F31D0),
SPH_C32(0x082EFA98), SPH_C32(0xEC4E6C89),
SPH_C32(0x452821E6), SPH_C32(0x38D01377),
SPH_C32(0xBE5466CF), SPH_C32(0x34E90C6C),
SPH_C32(0xC0AC29B7), SPH_C32(0xC97C50DD),
SPH_C32(0x3F84D5B5), SPH_C32(0xB5470917)
};
#define GS(a,b,c,d,x) { \
const uint32_t idx1 = c_sigma[r][x]; \
const uint32_t idx2 = c_sigma[r][x+1]; \
v[a] += (m[idx1] ^ c_u256[idx2]) + v[b]; \
v[d] = SPH_ROTL32(v[d] ^ v[a], 16); \
v[c] += v[d]; \ v[c] += v[d]; \
v[b] = SPH_ROTR32(v[b] ^ v[c], 12); \ v[b] = SPH_ROTR32(v[b] ^ v[c], 12); \
\ v[a] += (m[y] ^ c_u256[x]) + v[b]; \
v[a] += (m[idx2] ^ c_u256[idx1]) + v[b]; \ v[d] = __byte_perm(v[d] ^ v[a],0, 0x0321); \
v[d] = SPH_ROTR32(v[d] ^ v[a], 8); \
v[c] += v[d]; \ v[c] += v[d]; \
v[b] = SPH_ROTR32(v[b] ^ v[c], 7); \ v[b] = SPH_ROTR32(v[b] ^ v[c], 7); \
} }
/* Second part (64-80) msg never change, store it */ /* Second part (64-80) msg never change, store it */
__device__ __constant__ __device__ __forceinline__
static const uint32_t __align__(32) c_Padding[16] = {
0, 0, 0, 0,
0x80000000UL, 0, 0, 0,
0, 0, 0, 0,
0, 1, 0, 640,
};
__device__ static
void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, const int rounds) void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, const int rounds)
{ {
uint32_t /*_ALIGN(8)*/ m[16]; uint32_t /*_ALIGN(8)*/ m[16];
@ -132,12 +66,23 @@ void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, co
m[2] = block[2]; m[2] = block[2];
m[3] = block[3]; m[3] = block[3];
const uint32_t c_u256[16] = {
0x243F6A88, 0x85A308D3, 0x13198A2E, 0x03707344,
0xA4093822, 0x299F31D0, 0x082EFA98, 0xEC4E6C89,
0x452821E6, 0x38D01377, 0xBE5466CF, 0x34E90C6C,
0xC0AC29B7, 0xC97C50DD, 0x3F84D5B5, 0xB5470917
};
const uint32_t c_Padding[16] = {
0, 0, 0, 0,
0x80000000UL, 0, 0, 0,
0, 0, 0, 0,
0, 1, 0, 640,
};
#pragma unroll
for (uint32_t i = 4; i < 16; i++) { for (uint32_t i = 4; i < 16; i++) {
#if PRECALC64
m[i] = c_Padding[i]; m[i] = c_Padding[i];
#else
m[i] = (T0 == 0x200) ? block[i] : c_Padding[i];
#endif
} }
//#pragma unroll 8 //#pragma unroll 8
@ -154,153 +99,333 @@ void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0, co
v[14] = c_u256[6]; v[14] = c_u256[6];
v[15] = c_u256[7]; v[15] = c_u256[7];
for (int r = 0; r < rounds; r++) { // { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
/* column step */ GSPREC(0, 4, 0x8, 0xC,0,1);
GS(0, 4, 0x8, 0xC, 0x0); GSPREC(1, 5, 0x9, 0xD,2,3);
GS(1, 5, 0x9, 0xD, 0x2); GSPREC(2, 6, 0xA, 0xE, 4,5);
GS(2, 6, 0xA, 0xE, 0x4); GSPREC(3, 7, 0xB, 0xF, 6,7);
GS(3, 7, 0xB, 0xF, 0x6); GSPREC(0, 5, 0xA, 0xF, 8,9);
/* diagonal step */ GSPREC(1, 6, 0xB, 0xC, 10,11);
GS(0, 5, 0xA, 0xF, 0x8); GSPREC(2, 7, 0x8, 0xD, 12,13);
GS(1, 6, 0xB, 0xC, 0xA); GSPREC(3, 4, 0x9, 0xE, 14,15);
GS(2, 7, 0x8, 0xD, 0xC); // { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
GS(3, 4, 0x9, 0xE, 0xE); GSPREC(0, 4, 0x8, 0xC, 14, 10);
} GSPREC(1, 5, 0x9, 0xD, 4, 8);
#if PRECALC64 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);
// { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
GSPREC(0, 4, 0x8, 0xC, 11, 8);
GSPREC(1, 5, 0x9, 0xD, 12, 0);
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);
// { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
GSPREC(0, 4, 0x8, 0xC, 7, 9);
GSPREC(1, 5, 0x9, 0xD, 3, 1);
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);
// { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
GSPREC(0, 4, 0x8, 0xC, 9, 0);
GSPREC(1, 5, 0x9, 0xD, 5, 7);
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);
// { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
GSPREC(0, 4, 0x8, 0xC, 2, 12);
GSPREC(1, 5, 0x9, 0xD, 6, 10);
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);
// { 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
GSPREC(0, 4, 0x8, 0xC, 12, 5);
GSPREC(1, 5, 0x9, 0xD, 1, 15);
GSPREC(2, 6, 0xA, 0xE, 14, 13);
GSPREC(3, 7, 0xB, 0xF, 4, 10);
GSPREC(0, 5, 0xA, 0xF, 0, 7);
GSPREC(1, 6, 0xB, 0xC, 6, 3);
GSPREC(2, 7, 0x8, 0xD, 9, 2);
GSPREC(3, 4, 0x9, 0xE, 8, 11);
// { 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
GSPREC(0, 4, 0x8, 0xC, 13, 11);
GSPREC(1, 5, 0x9, 0xD, 7, 14);
GSPREC(2, 6, 0xA, 0xE, 12, 1);
GSPREC(3, 7, 0xB, 0xF, 3, 9);
GSPREC(0, 5, 0xA, 0xF, 5, 0);
GSPREC(1, 6, 0xB, 0xC, 15, 4);
GSPREC(2, 7, 0x8, 0xD, 8, 6);
GSPREC(3, 4, 0x9, 0xE, 2, 10);
// { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
GSPREC(0, 4, 0x8, 0xC, 6, 15);
GSPREC(1, 5, 0x9, 0xD, 14, 9);
GSPREC(2, 6, 0xA, 0xE, 11, 3);
GSPREC(3, 7, 0xB, 0xF, 0, 8);
GSPREC(0, 5, 0xA, 0xF, 12, 2);
GSPREC(1, 6, 0xB, 0xC, 13, 7);
GSPREC(2, 7, 0x8, 0xD, 1, 4);
GSPREC(3, 4, 0x9, 0xE, 10, 5);
// { 10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13, 0 },
GSPREC(0, 4, 0x8, 0xC, 10, 2);
GSPREC(1, 5, 0x9, 0xD, 8, 4);
GSPREC(2, 6, 0xA, 0xE, 7, 6);
GSPREC(3, 7, 0xB, 0xF, 1, 5);
GSPREC(0, 5, 0xA, 0xF, 15, 11);
GSPREC(1, 6, 0xB, 0xC, 9, 14);
GSPREC(2, 7, 0x8, 0xD, 3, 12);
GSPREC(3, 4, 0x9, 0xE, 13, 0);
// { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
GSPREC(0, 4, 0x8, 0xC, 0, 1);
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(1, 6, 0xB, 0xC, 10, 11);
GSPREC(2, 7, 0x8, 0xD, 12, 13);
GSPREC(3, 4, 0x9, 0xE, 14, 15);
// { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
GSPREC(0, 4, 0x8, 0xC, 14, 10);
GSPREC(1, 5, 0x9, 0xD, 4, 8);
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);
// { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
GSPREC(0, 4, 0x8, 0xC, 11, 8);
GSPREC(1, 5, 0x9, 0xD, 12, 0);
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);
// { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
GSPREC(0, 4, 0x8, 0xC, 7, 9);
GSPREC(1, 5, 0x9, 0xD, 3, 1);
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);
// only compute h6 & 7 // only compute h6 & 7
h[6U] ^= v[6U] ^ v[14U]; h[6U] ^= v[6U] ^ v[14U];
h[7U] ^= v[7U] ^ v[15U]; h[7U] ^= v[7U] ^ v[15U];
#else
//#pragma unroll 16
for (uint32_t i = 0; i < 16; i++) {
uint32_t j = i % 8U;
h[j] ^= v[i];
}
#endif
} }
#if !PRECALC64 /* original method */ /* ############################################################################################################################### */
__global__ /* Precalculated 1st 64-bytes block (midstate) method */
void blake256_gpu_hash_80(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce,
const uint64_t highTarget, const int crcsum, const int rounds) __global__ __launch_bounds__(1024,1)
void blake256_gpu_hash_16(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce,
const uint64_t highTarget, const int rounds, const bool trace)
{ {
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;
uint32_t h[8]; uint32_t _ALIGN(16) h[8];
#pragma unroll #pragma unroll
for(int i=0; i<8; i++) { for(int i=0; i < 8; i++) {
h[i] = c_IV256[i]; h[i] = d_data[i];
}
if (crcsum != prevsum) {
prevsum = crcsum;
blake256_compress(h, c_data, 512, rounds);
#pragma unroll
for(int i=0; i<8; i++) {
cache[i] = h[i];
}
} else {
#pragma unroll
for(int i=0; i<8; i++) {
h[i] = cache[i];
}
} }
// ------ Close: Bytes 64 to 80 ------ // ------ Close: Bytes 64 to 80 ------
uint32_t ending[4]; uint32_t _ALIGN(16) ending[4];
ending[0] = c_data[16]; ending[0] = d_data[8];
ending[1] = c_data[17]; ending[1] = d_data[9];
ending[2] = c_data[18]; ending[2] = d_data[10];
ending[3] = nonce; /* our tested value */ ending[3] = nonce; /* our tested value */
blake256_compress(h, ending, 640, rounds); blake256_compress(h, ending, 640, rounds);
// not sure why, h[7] is ok if (h[7] == 0 && cuda_swab32(h[6]) <= highTarget) {
h[6] = cuda_swab32(h[6]);
// compare count of leading zeros h[6] + h[7]
uint64_t high64 = ((uint64_t*)h)[3];
if (high64 <= highTarget)
#if NBN == 2 #if NBN == 2
/* keep the smallest nonce, + extra one if found */ /* keep the smallest nonce, + extra one if found */
if (resNonce[0] > nonce) { if (resNonce[0] > nonce) {
// printf("%llx %llx \n", high64, highTarget); resNonce[1] = resNonce[0];
resNonce[1] = resNonce[0]; resNonce[0] = nonce;
resNonce[0] = nonce; }
} else
else resNonce[1] = nonce;
resNonce[1] = nonce;
#else #else
resNonce[0] = nonce; resNonce[0] = nonce;
#endif #endif
#ifdef _DEBUG
if (trace) {
uint64_t high64 = ((uint64_t*)h)[3];
printf("gpu: %16llx\n", high64);
printf("gpu: %08x.%08x\n", h[7], h[6]);
printf("tgt: %16llx\n", highTarget);
}
#endif
}
} }
} }
__host__
uint32_t blake256_cpu_hash_80(const int thr_id, const uint32_t threads, const uint32_t startNonce, const uint64_t highTarget,
const uint32_t crcsum, const int8_t rounds)
{
const uint32_t threadsperblock = TPB;
uint32_t result = UINT32_MAX;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
size_t shared_size = 0;
/* Check error on Ctrl+C or kill to prevent segfaults on exit */
if (cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t)) != cudaSuccess)
return result;
blake256_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNonce, d_resNonce[thr_id], highTarget, crcsum, (int) rounds);
//MyStreamSynchronize(NULL, 0, thr_id);
if (cudaSuccess == cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) {
result = h_resNonce[thr_id][0];
for (int n=0; n < (NBN-1); n++)
extra_results[n] = h_resNonce[thr_id][n+1];
}
return result;
}
__host__
void blake256_cpu_setBlock_80(uint32_t *pdata, const uint32_t *ptarget)
{
uint32_t data[20];
memcpy(data, pdata, 80);
CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_data, data, sizeof(data), 0, cudaMemcpyHostToDevice));
}
#else
/* ############################################################################################################################### */
/* Precalculated 1st 64-bytes block (midstate) method */
__global__ __global__
void blake256_gpu_hash_16(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce, void blake256_gpu_hash_16_8(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonce,
const uint64_t highTarget, const int rounds, const bool trace) const uint64_t highTarget, const int rounds, const bool trace)
{ {
const 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;
uint32_t _ALIGN(16) h[8]; uint32_t _ALIGN(16) h[8];
#pragma unroll #pragma unroll
for(int i=0; i < 8; i++) { for (int i = 0; i < 8; i++) {
h[i] = d_data[i]; h[i] = d_data[i];
} }
// ------ Close: Bytes 64 to 80 ------ // ------ Close: Bytes 64 to 80 ------
uint32_t _ALIGN(16) ending[4]; uint32_t _ALIGN(16) block[4];
ending[0] = d_data[8]; block[0] = d_data[8];
ending[1] = d_data[9]; block[1] = d_data[9];
ending[2] = d_data[10]; block[2] = d_data[10];
ending[3] = nonce; /* our tested value */ block[3] = nonce; /* our tested value */
blake256_compress(h, ending, 640, rounds); // blake256_compress_8(h, block, 640, rounds);
uint32_t /*_ALIGN(8)*/ m[16];
uint32_t v[16];
m[0] = block[0];
m[1] = block[1];
m[2] = block[2];
m[3] = block[3];
const uint32_t c_u256[16] = {
0x243F6A88, 0x85A308D3, 0x13198A2E, 0x03707344,
0xA4093822, 0x299F31D0, 0x082EFA98, 0xEC4E6C89,
0x452821E6, 0x38D01377, 0xBE5466CF, 0x34E90C6C,
0xC0AC29B7, 0xC97C50DD, 0x3F84D5B5, 0xB5470917
};
const uint32_t c_Padding[16] = {
0, 0, 0, 0,
0x80000000UL, 0, 0, 0,
0, 0, 0, 0,
0, 1, 0, 640,
};
#pragma unroll
for (uint32_t i = 4; i < 16; i++) {
m[i] = c_Padding[i];
}
//#pragma unroll 8
for (uint32_t i = 0; i < 8; i++)
v[i] = h[i];
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] ^ 640;
v[13] = c_u256[5] ^ 640;
v[14] = c_u256[6];
v[15] = c_u256[7];
// { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
GSPREC(0, 4, 0x8, 0xC, 0, 1);
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(1, 6, 0xB, 0xC, 10, 11);
GSPREC(2, 7, 0x8, 0xD, 12, 13);
GSPREC(3, 4, 0x9, 0xE, 14, 15);
// { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
GSPREC(0, 4, 0x8, 0xC, 14, 10);
GSPREC(1, 5, 0x9, 0xD, 4, 8);
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);
// { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
GSPREC(0, 4, 0x8, 0xC, 11, 8);
GSPREC(1, 5, 0x9, 0xD, 12, 0);
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);
// { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
GSPREC(0, 4, 0x8, 0xC, 7, 9);
GSPREC(1, 5, 0x9, 0xD, 3, 1);
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);
// { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
GSPREC(0, 4, 0x8, 0xC, 9, 0);
GSPREC(1, 5, 0x9, 0xD, 5, 7);
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);
// { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
GSPREC(0, 4, 0x8, 0xC, 2, 12);
GSPREC(1, 5, 0x9, 0xD, 6, 10);
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);
// { 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
GSPREC(0, 4, 0x8, 0xC, 12, 5);
GSPREC(1, 5, 0x9, 0xD, 1, 15);
GSPREC(2, 6, 0xA, 0xE, 14, 13);
GSPREC(3, 7, 0xB, 0xF, 4, 10);
GSPREC(0, 5, 0xA, 0xF, 0, 7);
GSPREC(1, 6, 0xB, 0xC, 6, 3);
GSPREC(2, 7, 0x8, 0xD, 9, 2);
GSPREC(3, 4, 0x9, 0xE, 8, 11);
// { 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
GSPREC(0, 4, 0x8, 0xC, 13, 11);
GSPREC(1, 5, 0x9, 0xD, 7, 14);
GSPREC(2, 6, 0xA, 0xE, 12, 1);
GSPREC(3, 7, 0xB, 0xF, 3, 9);
GSPREC(0, 5, 0xA, 0xF, 5, 0);
GSPREC(1, 6, 0xB, 0xC, 15, 4);
GSPREC(2, 7, 0x8, 0xD, 8, 6);
GSPREC(3, 4, 0x9, 0xE, 2, 10);
// { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
// only compute h6 & 7
h[6U] ^= v[6U] ^ v[14U];
h[7U] ^= v[7U] ^ v[15U];
if (h[7] == 0 && cuda_swab32(h[6]) <= highTarget) { if (h[7] == 0 && cuda_swab32(h[6]) <= highTarget) {
#if NBN == 2 #if NBN == 2
@ -330,26 +455,27 @@ __host__
static uint32_t blake256_cpu_hash_16(const int thr_id, const uint32_t threads, const uint32_t startNonce, const uint64_t highTarget, static uint32_t blake256_cpu_hash_16(const int thr_id, const uint32_t threads, const uint32_t startNonce, const uint64_t highTarget,
const int8_t rounds) const int8_t rounds)
{ {
const uint32_t threadsperblock = TPB;
uint32_t result = UINT32_MAX; uint32_t result = UINT32_MAX;
dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 grid((threads + TPB-1)/TPB);
dim3 block(threadsperblock); dim3 block(TPB);
cudaGetLastError();
/* 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_16 <<<grid, block>>> (threads, startNonce, d_resNonce[thr_id], highTarget, (int) rounds, opt_tracegpu); if (rounds == 8)
//MyStreamSynchronize(NULL, 0, thr_id); blake256_gpu_hash_16_8 <<<grid, block>>> (threads, startNonce, d_resNonce[thr_id], highTarget, (int)rounds, opt_tracegpu);
else
blake256_gpu_hash_16 <<<grid, block>>> (threads, startNonce, d_resNonce[thr_id], highTarget, (int)rounds, opt_tracegpu);
// cudaDeviceSynchronize();
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)) {
//cudaDeviceSynchronize(); /* seems no more required */
result = h_resNonce[thr_id][0]; result = h_resNonce[thr_id][0];
for (int n=0; n < (NBN-1); n++) for (int n=0; n < (NBN-1); n++)
extra_results[n] = h_resNonce[thr_id][n+1]; extra_results[n] = h_resNonce[thr_id][n+1];
} }
CUDA_LOG_ERROR();
return result; return result;
} }
@ -376,32 +502,29 @@ void blake256_cpu_setBlock_16(uint32_t *penddata, const uint32_t *midstate, cons
data[10]= penddata[2]; data[10]= penddata[2];
CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_data, data, 32 + 12, 0, cudaMemcpyHostToDevice)); CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_data, data, 32 + 12, 0, cudaMemcpyHostToDevice));
} }
#endif
static bool init[MAX_GPUS] = { 0 }; static bool init[MAX_GPUS] = { 0 };
extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done, int8_t blakerounds=14) extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done, int8_t blakerounds=14)
{ {
uint32_t _ALIGN(64) endiandata[20]; uint32_t _ALIGN(64) endiandata[20];
#if PRECALC64
uint32_t _ALIGN(64) midstate[8]; uint32_t _ALIGN(64) midstate[8];
#else
uint32_t crcsum;
#endif
uint32_t *pdata = work->data; uint32_t *pdata = work->data;
uint32_t *ptarget = work->target; uint32_t *ptarget = work->target;
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint64_t targetHigh = ((uint64_t*)ptarget)[3]; uint64_t targetHigh = ((uint64_t*)ptarget)[3];
int intensity = (device_sm[device_map[thr_id]] > 500) ? 22 : 20;
uint32_t intensity = (device_sm[device_map[thr_id]] > 500) ? 31 : 28;
uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity); uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity);
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
int rc = 0; int rc = 0;
if (opt_benchmark) { if (opt_benchmark) {
ptarget[7] = 0; targetHigh = 0x1ULL << 32;
ptarget[6] = swab32(0xff); ptarget[6] = swab32(0x00ff);
targetHigh = 0xffULL << 32;
} }
if (opt_tracegpu) { if (opt_tracegpu) {
@ -411,44 +534,33 @@ extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_non
pdata[k] = swab32(pdata[k]); pdata[k] = swab32(pdata[k]);
} }
if (!init[thr_id]) { if (!init[thr_id])
{
cudaSetDevice(device_map[thr_id]); cudaSetDevice(device_map[thr_id]);
if (opt_cudaschedule == -1 && gpu_threads == 1) { if (opt_cudaschedule == -1 && gpu_threads == 1) {
cudaDeviceReset(); cudaDeviceReset();
// reduce cpu usage (linux) // reduce cpu usage (linux)
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
CUDA_LOG_ERROR(); CUDA_LOG_ERROR();
} }
cudaMallocHost(&h_resNonce[thr_id], NBN * sizeof(uint32_t)); CUDA_CALL_OR_RET_X(cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t)), -1);
cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t)); CUDA_CALL_OR_RET_X(cudaMallocHost(&h_resNonce[thr_id], NBN * sizeof(uint32_t)), -1);
CUDA_LOG_ERROR();
init[thr_id] = true; init[thr_id] = true;
} }
#if PRECALC64
for (int k = 0; k < 16; k++) for (int k = 0; k < 16; k++)
be32enc(&endiandata[k], pdata[k]); be32enc(&endiandata[k], pdata[k]);
blake256mid(midstate, endiandata, blakerounds); blake256mid(midstate, endiandata, blakerounds);
blake256_cpu_setBlock_16(&pdata[16], midstate, ptarget); blake256_cpu_setBlock_16(&pdata[16], midstate, ptarget);
#else
blake256_cpu_setBlock_80(pdata, ptarget);
crcsum = crc32_u32t(pdata, 64);
#endif /* PRECALC64 */
do { do {
*hashes_done = pdata[19] - first_nonce + throughput;
uint32_t foundNonce =
#if PRECALC64
// GPU HASH (second block only, first is midstate) // GPU HASH (second block only, first is midstate)
blake256_cpu_hash_16(thr_id, throughput, pdata[19], targetHigh, blakerounds); uint32_t foundNonce = blake256_cpu_hash_16(thr_id, throughput, pdata[19], targetHigh, blakerounds);
#else
// GPU FULL HASH if (foundNonce != UINT32_MAX)
blake256_cpu_hash_80(thr_id, throughput, pdata[19], targetHigh, crcsum, blakerounds);
#endif
if (foundNonce != UINT32_MAX && bench_algo == -1)
{ {
uint32_t vhashcpu[8]; uint32_t vhashcpu[8];
uint32_t Htarg = (uint32_t)targetHigh; uint32_t Htarg = (uint32_t)targetHigh;
@ -463,16 +575,15 @@ extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_non
{ {
rc = 1; rc = 1;
work_set_target_ratio(work, vhashcpu); work_set_target_ratio(work, vhashcpu);
*hashes_done = pdata[19] - first_nonce + throughput;
pdata[19] = foundNonce; pdata[19] = foundNonce;
#if NBN > 1 #if NBN > 1
if (extra_results[0] != UINT32_MAX) { if (extra_results[0] != UINT32_MAX) {
be32enc(&endiandata[19], extra_results[0]); be32enc(&endiandata[19], extra_results[0]);
blake256hash(vhashcpu, endiandata, blakerounds); blake256hash(vhashcpu, endiandata, blakerounds);
if (vhashcpu[6] <= Htarg && fulltest(vhashcpu, ptarget)) { if (vhashcpu[6] <= Htarg /* && fulltest(vhashcpu, ptarget) */) {
pdata[21] = extra_results[0]; pdata[21] = extra_results[0];
applog(LOG_BLUE, "1:%x 2:%x", foundNonce, extra_results[0]); applog(LOG_BLUE, "1:%x 2:%x", foundNonce, extra_results[0]);
if (bn_hash_target_ratio(vhashcpu, ptarget) > work->shareratio)
work_set_target_ratio(work, vhashcpu);
rc = 2; rc = 2;
} }
extra_results[0] = UINT32_MAX; extra_results[0] = UINT32_MAX;
@ -480,24 +591,20 @@ extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_non
#endif #endif
return rc; return rc;
} }
else if (vhashcpu[7] > ptarget[7] && opt_debug) { else if (opt_debug) {
applog_hash((uchar*)ptarget); applog_hash((uchar*)ptarget);
applog_compare_hash((uchar*)vhashcpu, (uchar*)ptarget); applog_compare_hash((uchar*)vhashcpu, (uchar*)ptarget);
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce); gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
} }
} }
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart && ((uint64_t)max_nonce > (uint64_t)throughput + pdata[19]));
*hashes_done = pdata[19] - first_nonce; *hashes_done = pdata[19] - first_nonce;
MyStreamSynchronize(NULL, 0, device_map[thr_id]);
return rc; return rc;
} }
@ -507,7 +614,7 @@ extern "C" void free_blake256(int thr_id)
if (!init[thr_id]) if (!init[thr_id])
return; return;
cudaThreadSynchronize(); cudaDeviceSynchronize();
cudaFreeHost(h_resNonce[thr_id]); cudaFreeHost(h_resNonce[thr_id]);
cudaFree(d_resNonce[thr_id]); cudaFree(d_resNonce[thr_id]);
@ -516,3 +623,4 @@ extern "C" void free_blake256(int thr_id)
cudaDeviceSynchronize(); cudaDeviceSynchronize();
} }

2
configure.ac
Unescape View File

@ -1,4 +1,4 @@
AC_INIT([ccminer], [1.7.1], [], [ccminer], [http://github.com/tpruvot/ccminer]) AC_INIT([ccminer], [1.7.2], [], [ccminer], [http://github.com/tpruvot/ccminer])
AC_PREREQ([2.59c]) AC_PREREQ([2.59c])
AC_CANONICAL_SYSTEM AC_CANONICAL_SYSTEM

6
cpuminer-config.h
Unescape View File

@ -162,7 +162,7 @@
#define PACKAGE_NAME "ccminer" #define PACKAGE_NAME "ccminer"
/* Define to the full name and version of this package. */ /* Define to the full name and version of this package. */
#define PACKAGE_STRING "ccminer 1.7.1" #define PACKAGE_STRING "ccminer 1.7.2"
/* Define to the one symbol short name of this package. */ /* Define to the one symbol short name of this package. */
#define PACKAGE_TARNAME "ccminer" #define PACKAGE_TARNAME "ccminer"
@ -171,7 +171,7 @@
#define PACKAGE_URL "http://github.com/tpruvot/ccminer" #define PACKAGE_URL "http://github.com/tpruvot/ccminer"
/* Define to the version of this package. */ /* Define to the version of this package. */
#define PACKAGE_VERSION "1.7.1" #define PACKAGE_VERSION "1.7.2"
/* If using the C implementation of alloca, define if you know the /* If using the C implementation of alloca, define if you know the
direction of stack growth for your system; otherwise it will be direction of stack growth for your system; otherwise it will be
@ -185,7 +185,7 @@
#define STDC_HEADERS 1 #define STDC_HEADERS 1
/* Version number of package */ /* Version number of package */
#define VERSION "1.7.1" #define VERSION "1.7.2"
/* Define curl_free() as free() if our version of curl lacks curl_free. */ /* Define curl_free() as free() if our version of curl lacks curl_free. */
/* #undef curl_free */ /* #undef curl_free */

Write Preview
Loading…
Cancel
Save