GOSTcoin support for ccminer CUDA miner project, compatible with most nvidia cards
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.

310 lines
9.5 KiB

#include <memory.h>
#define SPH_C32(x) ((uint32_t)(x ## U))
#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF))
#include "cuda_helper.h"
uint32_t *d_gnounce[MAX_GPUS];
uint32_t *d_GNonce[MAX_GPUS];
__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
texture<unsigned int, 1, cudaReadModeElementType> t0up2;
texture<unsigned int, 1, cudaReadModeElementType> t0dn2;
texture<unsigned int, 1, cudaReadModeElementType> t1up2;
texture<unsigned int, 1, cudaReadModeElementType> t1dn2;
texture<unsigned int, 1, cudaReadModeElementType> t2up2;
texture<unsigned int, 1, cudaReadModeElementType> t2dn2;
texture<unsigned int, 1, cudaReadModeElementType> t3up2;
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(int 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(int 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_hash32(int threads, uint32_t startNounce, uint64_t *outputHash, uint32_t *nonceVector)
{
#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
int 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]) {
nonceVector[0] = nonce;
}
}
}
#define texDef(texname, texmem, texsource, texsize) \
unsigned int *texmem; \
cudaMalloc(&texmem, texsize); \
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, int threads)
{
// Texturen mit obigem Makro initialisieren
texDef(t0up2, d_T0up, T0up_cpu, sizeof(uint32_t) * 256);
texDef(t0dn2, d_T0dn, T0dn_cpu, sizeof(uint32_t) * 256);
texDef(t1up2, d_T1up, T1up_cpu, sizeof(uint32_t) * 256);
texDef(t1dn2, d_T1dn, T1dn_cpu, sizeof(uint32_t) * 256);
texDef(t2up2, d_T2up, T2up_cpu, sizeof(uint32_t) * 256);
texDef(t2dn2, d_T2dn, T2dn_cpu, sizeof(uint32_t) * 256);
texDef(t3up2, d_T3up, T3up_cpu, sizeof(uint32_t) * 256);
texDef(t3dn2, d_T3dn, T3dn_cpu, sizeof(uint32_t) * 256);
cudaMalloc(&d_GNonce[thr_id], sizeof(uint32_t));
cudaMallocHost(&d_gnounce[thr_id], 1*sizeof(uint32_t));
}
__host__
uint32_t groestl256_cpu_hash_32(int thr_id, int threads, uint32_t startNounce, uint64_t *d_outputHash, int order)
{
uint32_t result = 0xffffffff;
cudaMemset(d_GNonce[thr_id], 0xff, sizeof(uint32_t));
const int 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_hash32<<<grid, block, shared_size>>>(threads, startNounce, d_outputHash, d_GNonce[thr_id]);
MyStreamSynchronize(NULL, order, thr_id);
cudaMemcpy(d_gnounce[thr_id], d_GNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost);
cudaThreadSynchronize();
result = *d_gnounce[thr_id];
return result;
}
__host__
void groestl256_setTarget(const void *pTargetIn)
{
cudaMemcpyToSymbol(pTarget, pTargetIn, 8 * sizeof(uint32_t), 0, cudaMemcpyHostToDevice);
}