mirror of https://github.com/GOSTSec/ccminer
You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
335 lines
10 KiB
335 lines
10 KiB
#include <memory.h> |
|
|
|
#define SPH_C32(x) ((uint32_t)(x ## U)) |
|
#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF)) |
|
|
|
#include "cuda_helper.h" |
|
|
|
static uint32_t *h_GNonces[MAX_GPUS]; |
|
static uint32_t *d_GNonces[MAX_GPUS]; |
|
static unsigned int* d_textures[MAX_GPUS][8]; |
|
|
|
__constant__ uint32_t pTarget[8]; |
|
|
|
#define C32e(x) \ |
|
((SPH_C32(x) >> 24) \ |
|
| ((SPH_C32(x) >> 8) & SPH_C32(0x0000FF00)) \ |
|
| ((SPH_C32(x) << 8) & SPH_C32(0x00FF0000)) \ |
|
| ((SPH_C32(x) << 24) & SPH_C32(0xFF000000))) |
|
|
|
#define PC32up(j, r) ((uint32_t)((j) + (r))) |
|
#define PC32dn(j, r) 0 |
|
#define QC32up(j, r) 0xFFFFFFFF |
|
#define QC32dn(j, r) (((uint32_t)(r) << 24) ^ SPH_T32(~((uint32_t)(j) << 24))) |
|
|
|
#define B32_0(x) __byte_perm(x, 0, 0x4440) |
|
//((x) & 0xFF) |
|
#define B32_1(x) __byte_perm(x, 0, 0x4441) |
|
//(((x) >> 8) & 0xFF) |
|
#define B32_2(x) __byte_perm(x, 0, 0x4442) |
|
//(((x) >> 16) & 0xFF) |
|
#define B32_3(x) __byte_perm(x, 0, 0x4443) |
|
//((x) >> 24) |
|
|
|
#define MAXWELL_OR_FERMI 1 |
|
#if MAXWELL_OR_FERMI |
|
#define USE_SHARED 1 |
|
// Maxwell and Fermi cards get the best speed with SHARED access it seems. |
|
#if USE_SHARED |
|
#define T0up(x) (*((uint32_t*)mixtabs + ( (x)))) |
|
#define T0dn(x) (*((uint32_t*)mixtabs + (256+(x)))) |
|
#define T1up(x) (*((uint32_t*)mixtabs + (512+(x)))) |
|
#define T1dn(x) (*((uint32_t*)mixtabs + (768+(x)))) |
|
#define T2up(x) (*((uint32_t*)mixtabs + (1024+(x)))) |
|
#define T2dn(x) (*((uint32_t*)mixtabs + (1280+(x)))) |
|
#define T3up(x) (*((uint32_t*)mixtabs + (1536+(x)))) |
|
#define T3dn(x) (*((uint32_t*)mixtabs + (1792+(x)))) |
|
#else |
|
#define T0up(x) tex1Dfetch(t0up2, x) |
|
#define T0dn(x) tex1Dfetch(t0dn2, x) |
|
#define T1up(x) tex1Dfetch(t1up2, x) |
|
#define T1dn(x) tex1Dfetch(t1dn2, x) |
|
#define T2up(x) tex1Dfetch(t2up2, x) |
|
#define T2dn(x) tex1Dfetch(t2dn2, x) |
|
#define T3up(x) tex1Dfetch(t3up2, x) |
|
#define T3dn(x) tex1Dfetch(t3dn2, x) |
|
#endif |
|
#else |
|
#define USE_SHARED 1 |
|
// a healthy mix between shared and textured access provides the highest speed on Compute 3.0 and 3.5! |
|
#define T0up(x) (*((uint32_t*)mixtabs + ( (x)))) |
|
#define T0dn(x) tex1Dfetch(t0dn2, x) |
|
#define T1up(x) tex1Dfetch(t1up2, x) |
|
#define T1dn(x) (*((uint32_t*)mixtabs + (768+(x)))) |
|
#define T2up(x) tex1Dfetch(t2up2, x) |
|
#define T2dn(x) (*((uint32_t*)mixtabs + (1280+(x)))) |
|
#define T3up(x) (*((uint32_t*)mixtabs + (1536+(x)))) |
|
#define T3dn(x) tex1Dfetch(t3dn2, x) |
|
#endif |
|
|
|
static texture<unsigned int, 1, cudaReadModeElementType> t0up2; |
|
static texture<unsigned int, 1, cudaReadModeElementType> t0dn2; |
|
static texture<unsigned int, 1, cudaReadModeElementType> t1up2; |
|
static texture<unsigned int, 1, cudaReadModeElementType> t1dn2; |
|
static texture<unsigned int, 1, cudaReadModeElementType> t2up2; |
|
static texture<unsigned int, 1, cudaReadModeElementType> t2dn2; |
|
static texture<unsigned int, 1, cudaReadModeElementType> t3up2; |
|
static texture<unsigned int, 1, cudaReadModeElementType> t3dn2; |
|
|
|
#define RSTT(d0, d1, a, b0, b1, b2, b3, b4, b5, b6, b7) do { \ |
|
t[d0] = T0up(B32_0(a[b0])) \ |
|
^ T1up(B32_1(a[b1])) \ |
|
^ T2up(B32_2(a[b2])) \ |
|
^ T3up(B32_3(a[b3])) \ |
|
^ T0dn(B32_0(a[b4])) \ |
|
^ T1dn(B32_1(a[b5])) \ |
|
^ T2dn(B32_2(a[b6])) \ |
|
^ T3dn(B32_3(a[b7])); \ |
|
t[d1] = T0dn(B32_0(a[b0])) \ |
|
^ T1dn(B32_1(a[b1])) \ |
|
^ T2dn(B32_2(a[b2])) \ |
|
^ T3dn(B32_3(a[b3])) \ |
|
^ T0up(B32_0(a[b4])) \ |
|
^ T1up(B32_1(a[b5])) \ |
|
^ T2up(B32_2(a[b6])) \ |
|
^ T3up(B32_3(a[b7])); \ |
|
} while (0) |
|
|
|
|
|
extern uint32_t T0up_cpu[]; |
|
extern uint32_t T0dn_cpu[]; |
|
extern uint32_t T1up_cpu[]; |
|
extern uint32_t T1dn_cpu[]; |
|
extern uint32_t T2up_cpu[]; |
|
extern uint32_t T2dn_cpu[]; |
|
extern uint32_t T3up_cpu[]; |
|
extern uint32_t T3dn_cpu[]; |
|
|
|
__device__ __forceinline__ |
|
void groestl256_perm_P(uint32_t thread,uint32_t *a, char *mixtabs) |
|
{ |
|
#pragma unroll 10 |
|
for (int r = 0; r<10; r++) |
|
{ |
|
uint32_t t[16]; |
|
|
|
a[0x0] ^= PC32up(0x00, r); |
|
a[0x2] ^= PC32up(0x10, r); |
|
a[0x4] ^= PC32up(0x20, r); |
|
a[0x6] ^= PC32up(0x30, r); |
|
a[0x8] ^= PC32up(0x40, r); |
|
a[0xA] ^= PC32up(0x50, r); |
|
a[0xC] ^= PC32up(0x60, r); |
|
a[0xE] ^= PC32up(0x70, r); |
|
RSTT(0x0, 0x1, a, 0x0, 0x2, 0x4, 0x6, 0x9, 0xB, 0xD, 0xF); |
|
RSTT(0x2, 0x3, a, 0x2, 0x4, 0x6, 0x8, 0xB, 0xD, 0xF, 0x1); |
|
RSTT(0x4, 0x5, a, 0x4, 0x6, 0x8, 0xA, 0xD, 0xF, 0x1, 0x3); |
|
RSTT(0x6, 0x7, a, 0x6, 0x8, 0xA, 0xC, 0xF, 0x1, 0x3, 0x5); |
|
RSTT(0x8, 0x9, a, 0x8, 0xA, 0xC, 0xE, 0x1, 0x3, 0x5, 0x7); |
|
RSTT(0xA, 0xB, a, 0xA, 0xC, 0xE, 0x0, 0x3, 0x5, 0x7, 0x9); |
|
RSTT(0xC, 0xD, a, 0xC, 0xE, 0x0, 0x2, 0x5, 0x7, 0x9, 0xB); |
|
RSTT(0xE, 0xF, a, 0xE, 0x0, 0x2, 0x4, 0x7, 0x9, 0xB, 0xD); |
|
|
|
#pragma unroll 16 |
|
for (int k = 0; k<16; k++) |
|
a[k] = t[k]; |
|
} |
|
} |
|
|
|
__device__ __forceinline__ |
|
void groestl256_perm_Q(uint32_t thread, uint32_t *a, char *mixtabs) |
|
{ |
|
#pragma unroll |
|
for (int r = 0; r<10; r++) |
|
{ |
|
uint32_t t[16]; |
|
|
|
a[0x0] ^= QC32up(0x00, r); |
|
a[0x1] ^= QC32dn(0x00, r); |
|
a[0x2] ^= QC32up(0x10, r); |
|
a[0x3] ^= QC32dn(0x10, r); |
|
a[0x4] ^= QC32up(0x20, r); |
|
a[0x5] ^= QC32dn(0x20, r); |
|
a[0x6] ^= QC32up(0x30, r); |
|
a[0x7] ^= QC32dn(0x30, r); |
|
a[0x8] ^= QC32up(0x40, r); |
|
a[0x9] ^= QC32dn(0x40, r); |
|
a[0xA] ^= QC32up(0x50, r); |
|
a[0xB] ^= QC32dn(0x50, r); |
|
a[0xC] ^= QC32up(0x60, r); |
|
a[0xD] ^= QC32dn(0x60, r); |
|
a[0xE] ^= QC32up(0x70, r); |
|
a[0xF] ^= QC32dn(0x70, r); |
|
RSTT(0x0, 0x1, a, 0x2, 0x6, 0xA, 0xE, 0x1, 0x5, 0x9, 0xD); |
|
RSTT(0x2, 0x3, a, 0x4, 0x8, 0xC, 0x0, 0x3, 0x7, 0xB, 0xF); |
|
RSTT(0x4, 0x5, a, 0x6, 0xA, 0xE, 0x2, 0x5, 0x9, 0xD, 0x1); |
|
RSTT(0x6, 0x7, a, 0x8, 0xC, 0x0, 0x4, 0x7, 0xB, 0xF, 0x3); |
|
RSTT(0x8, 0x9, a, 0xA, 0xE, 0x2, 0x6, 0x9, 0xD, 0x1, 0x5); |
|
RSTT(0xA, 0xB, a, 0xC, 0x0, 0x4, 0x8, 0xB, 0xF, 0x3, 0x7); |
|
RSTT(0xC, 0xD, a, 0xE, 0x2, 0x6, 0xA, 0xD, 0x1, 0x5, 0x9); |
|
RSTT(0xE, 0xF, a, 0x0, 0x4, 0x8, 0xC, 0xF, 0x3, 0x7, 0xB); |
|
|
|
#pragma unroll |
|
for (int k = 0; k<16; k++) |
|
a[k] = t[k]; |
|
} |
|
} |
|
|
|
__global__ __launch_bounds__(256,1) |
|
void groestl256_gpu_hash_32(uint32_t threads, uint32_t startNounce, uint64_t *outputHash, uint32_t *resNonces) |
|
{ |
|
#if USE_SHARED |
|
extern __shared__ char mixtabs[]; |
|
|
|
if (threadIdx.x < 256) { |
|
*((uint32_t*)mixtabs + (threadIdx.x)) = tex1Dfetch(t0up2, threadIdx.x); |
|
*((uint32_t*)mixtabs + (256 + threadIdx.x)) = tex1Dfetch(t0dn2, threadIdx.x); |
|
*((uint32_t*)mixtabs + (512 + threadIdx.x)) = tex1Dfetch(t1up2, threadIdx.x); |
|
*((uint32_t*)mixtabs + (768 + threadIdx.x)) = tex1Dfetch(t1dn2, threadIdx.x); |
|
*((uint32_t*)mixtabs + (1024 + threadIdx.x)) = tex1Dfetch(t2up2, threadIdx.x); |
|
*((uint32_t*)mixtabs + (1280 + threadIdx.x)) = tex1Dfetch(t2dn2, threadIdx.x); |
|
*((uint32_t*)mixtabs + (1536 + threadIdx.x)) = tex1Dfetch(t3up2, threadIdx.x); |
|
*((uint32_t*)mixtabs + (1792 + threadIdx.x)) = tex1Dfetch(t3dn2, threadIdx.x); |
|
} |
|
|
|
__syncthreads(); |
|
#endif |
|
|
|
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); |
|
if (thread < threads) |
|
{ |
|
// GROESTL |
|
uint32_t message[16]; |
|
uint32_t state[16]; |
|
|
|
#pragma unroll |
|
for (int k = 0; k<4; k++) |
|
LOHI(message[2*k], message[2*k+1], outputHash[k*threads+thread]); |
|
|
|
#pragma unroll |
|
for (int k = 9; k<15; k++) |
|
message[k] = 0; |
|
|
|
message[8] = 0x80; |
|
message[15] = 0x01000000; |
|
|
|
#pragma unroll 16 |
|
for (int u = 0; u<16; u++) |
|
state[u] = message[u]; |
|
|
|
state[15] ^= 0x10000; |
|
|
|
// Perm |
|
|
|
#if USE_SHARED |
|
groestl256_perm_P(thread, state, mixtabs); |
|
state[15] ^= 0x10000; |
|
groestl256_perm_Q(thread, message, mixtabs); |
|
#else |
|
groestl256_perm_P(thread, state, NULL); |
|
state[15] ^= 0x10000; |
|
groestl256_perm_P(thread, message, NULL); |
|
#endif |
|
#pragma unroll 16 |
|
for (int u = 0; u<16; u++) state[u] ^= message[u]; |
|
#pragma unroll 16 |
|
for (int u = 0; u<16; u++) message[u] = state[u]; |
|
#if USE_SHARED |
|
groestl256_perm_P(thread, message, mixtabs); |
|
#else |
|
groestl256_perm_P(thread, message, NULL); |
|
#endif |
|
state[14] ^= message[14]; |
|
state[15] ^= message[15]; |
|
|
|
uint32_t nonce = startNounce + thread; |
|
if (state[15] <= pTarget[7]) { |
|
atomicMin(&resNonces[1], resNonces[0]); |
|
atomicMin(&resNonces[0], nonce); |
|
} |
|
} |
|
} |
|
|
|
#define texDef(id, texname, texmem, texsource, texsize) { \ |
|
unsigned int *texmem; \ |
|
cudaMalloc(&texmem, texsize); \ |
|
d_textures[thr_id][id] = texmem; \ |
|
cudaMemcpy(texmem, texsource, texsize, cudaMemcpyHostToDevice); \ |
|
texname.normalized = 0; \ |
|
texname.filterMode = cudaFilterModePoint; \ |
|
texname.addressMode[0] = cudaAddressModeClamp; \ |
|
{ cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned int>(); \ |
|
cudaBindTexture(NULL, &texname, texmem, &channelDesc, texsize ); \ |
|
} \ |
|
} |
|
|
|
__host__ |
|
void groestl256_cpu_init(int thr_id, uint32_t threads) |
|
{ |
|
// Texturen mit obigem Makro initialisieren |
|
texDef(0, t0up2, d_T0up, T0up_cpu, sizeof(uint32_t) * 256); |
|
texDef(1, t0dn2, d_T0dn, T0dn_cpu, sizeof(uint32_t) * 256); |
|
texDef(2, t1up2, d_T1up, T1up_cpu, sizeof(uint32_t) * 256); |
|
texDef(3, t1dn2, d_T1dn, T1dn_cpu, sizeof(uint32_t) * 256); |
|
texDef(4, t2up2, d_T2up, T2up_cpu, sizeof(uint32_t) * 256); |
|
texDef(5, t2dn2, d_T2dn, T2dn_cpu, sizeof(uint32_t) * 256); |
|
texDef(6, t3up2, d_T3up, T3up_cpu, sizeof(uint32_t) * 256); |
|
texDef(7, t3dn2, d_T3dn, T3dn_cpu, sizeof(uint32_t) * 256); |
|
|
|
cudaMalloc(&d_GNonces[thr_id], 2*sizeof(uint32_t)); |
|
cudaMallocHost(&h_GNonces[thr_id], 2*sizeof(uint32_t)); |
|
} |
|
|
|
__host__ |
|
void groestl256_cpu_free(int thr_id) |
|
{ |
|
for (int i=0; i<8; i++) |
|
cudaFree(d_textures[thr_id][i]); |
|
|
|
cudaFree(d_GNonces[thr_id]); |
|
cudaFreeHost(h_GNonces[thr_id]); |
|
} |
|
|
|
__host__ |
|
uint32_t groestl256_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uint64_t *d_outputHash, int order) |
|
{ |
|
uint32_t result = UINT32_MAX; |
|
cudaMemset(d_GNonces[thr_id], 0xff, 2*sizeof(uint32_t)); |
|
const uint32_t threadsperblock = 256; |
|
|
|
// berechne wie viele Thread Blocks wir brauchen |
|
dim3 grid((threads + threadsperblock-1)/threadsperblock); |
|
dim3 block(threadsperblock); |
|
|
|
#if USE_SHARED |
|
size_t shared_size = 8 * 256 * sizeof(uint32_t); |
|
#else |
|
size_t shared_size = 0; |
|
#endif |
|
groestl256_gpu_hash_32<<<grid, block, shared_size>>>(threads, startNounce, d_outputHash, d_GNonces[thr_id]); |
|
|
|
MyStreamSynchronize(NULL, order, thr_id); |
|
|
|
// get first found nonce |
|
cudaMemcpy(h_GNonces[thr_id], d_GNonces[thr_id], 1*sizeof(uint32_t), cudaMemcpyDeviceToHost); |
|
result = *h_GNonces[thr_id]; |
|
|
|
return result; |
|
} |
|
|
|
__host__ |
|
uint32_t groestl256_getSecNonce(int thr_id, int num) |
|
{ |
|
uint32_t results[2]; |
|
memset(results, 0xFF, sizeof(results)); |
|
cudaMemcpy(results, d_GNonces[thr_id], sizeof(results), cudaMemcpyDeviceToHost); |
|
if (results[1] == results[0]) |
|
return UINT32_MAX; |
|
return results[num]; |
|
} |
|
|
|
__host__ |
|
void groestl256_setTarget(const void *pTargetIn) |
|
{ |
|
cudaMemcpyToSymbol(pTarget, pTargetIn, 32, 0, cudaMemcpyHostToDevice); |
|
}
|
|
|