GOSTSec
/
ccminer-gostd-lite
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity
GOSTCoin CUDA miner project, compatible with most nvidia cards, containing only gostd algo
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.
333 Commits
1 Branch
55 Tags
65 MiB
 
 
 
 
 
 
Tree: dfa9f63886
Branches Tags
${ item.name }
Create tag ${ searchTerm }
Create branch ${ searchTerm }
from 'dfa9f63886'
${ noResults }
ccminer-gostd-lite/Algo256/cuda_groestl256.cu

309 lines
9.5 KiB
Raw Blame History

#include <memory.h>
#include "cuda_helper.h"
uint32_t *d_gnounce[MAX_GPUS];
uint32_t *d_GNonce[MAX_GPUS];
__constant__ uint32_t pTarget[8];
#define SPH_C32(x) ((uint32_t)(x ## U))
#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF))
#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);
}