GOSTSec
/
ccminer
mirror of https://github.com/GOSTSec/ccminer
1
0
Fork 0
Code Issues Releases Wiki Activity
Browse Source

groestl: explain code and improve perf on SM 2.x 

Signed-off-by: Tanguy Pruvot <tanguy.pruvot@gmail.com>
2upstream
Tanguy Pruvot 8 years ago
parent
feb99d020f
commit
5a77d36635
6 changed files with 455 additions and 410 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. 1
      ccminer.vcxproj
  2. 3
      ccminer.vcxproj.filters
  3. 141
      quark/cuda_quark_groestl512.cu
  4. 267
      quark/cuda_quark_groestl512_sm2.cuh
  5. 342
      quark/cuda_quark_groestl512_sm20.cu
  6. 111
      quark/groestl_simple.cuh

1
ccminer.vcxproj
Unescape View File

@ -353,6 +353,7 @@
<ClInclude Include="miner.h" /> <ClInclude Include="miner.h" />
<ClInclude Include="nvml.h" /> <ClInclude Include="nvml.h" />
<ClInclude Include="quark\cuda_bmw512_sm3.cuh" /> <ClInclude Include="quark\cuda_bmw512_sm3.cuh" />
<ClInclude Include="quark\cuda_quark_groestl512_sm2.cuh" />
<ClInclude Include="quark\cuda_quark_blake512_sp.cuh" /> <ClInclude Include="quark\cuda_quark_blake512_sp.cuh" />
<ClInclude Include="quark\cuda_skein512_sp.cuh" /> <ClInclude Include="quark\cuda_skein512_sp.cuh" />
<ClInclude Include="res\resource.h" /> <ClInclude Include="res\resource.h" />

3
ccminer.vcxproj.filters
Unescape View File

@ -512,6 +512,9 @@
<ClInclude Include="quark\cuda_bmw512_sm3.cuh"> <ClInclude Include="quark\cuda_bmw512_sm3.cuh">
<Filter>Source Files\CUDA\quark</Filter> <Filter>Source Files\CUDA\quark</Filter>
</ClInclude> </ClInclude>
<ClInclude Include="quark\cuda_quark_groestl512_sm2.cuh">
<Filter>Source Files\CUDA\quark</Filter>
</ClInclude>
<ClInclude Include="x11\cuda_x11_simd512_sm2.cuh"> <ClInclude Include="x11\cuda_x11_simd512_sm2.cuh">
<Filter>Source Files\CUDA\x11</Filter> <Filter>Source Files\CUDA\x11</Filter>
</ClInclude> </ClInclude>

141
quark/cuda_quark_groestl512.cu
Unescape View File

@ -4,7 +4,7 @@
#include <memory.h> #include <memory.h>
#include <sys/types.h> // off_t #include <sys/types.h> // off_t
#include "cuda_helper.h" #include <cuda_helper.h>
#ifdef __INTELLISENSE__ #ifdef __INTELLISENSE__
#define __CUDA_ARCH__ 500 #define __CUDA_ARCH__ 500
@ -14,98 +14,103 @@
#define THF 4U #define THF 4U
#if __CUDA_ARCH__ >= 300 #if __CUDA_ARCH__ >= 300
#include "quark/groestl_functions_quad.h" #include "groestl_functions_quad.h"
#include "quark/groestl_transf_quad.h" #include "groestl_transf_quad.h"
#endif #endif
#include "quark/cuda_quark_groestl512_sm20.cu" #include "cuda_quark_groestl512_sm2.cuh"
__global__ __launch_bounds__(TPB, THF) __global__ __launch_bounds__(TPB, THF)
void quark_groestl512_gpu_hash_64_quad(const uint32_t threads, const uint32_t startNounce, uint32_t * g_hash, uint32_t * __restrict g_nonceVector) void quark_groestl512_gpu_hash_64_quad(const uint32_t threads, const uint32_t startNounce, uint32_t * g_hash, uint32_t * __restrict g_nonceVector)
{ {
#if __CUDA_ARCH__ >= 300 #if __CUDA_ARCH__ >= 300
// durch 4 dividieren, weil jeweils 4 Threads zusammen ein Hash berechnen
const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x) >> 2; // BEWARE : 4-WAY CODE (one hash need 4 threads)
if (thread < threads) const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x) >> 2;
{
// GROESTL if (thread < threads)
uint32_t message[8]; {
uint32_t state[8]; uint32_t message[8];
uint32_t state[8];
uint32_t nounce = g_nonceVector ? g_nonceVector[thread] : (startNounce + thread);
off_t hashPosition = nounce - startNounce; uint32_t nounce = g_nonceVector ? g_nonceVector[thread] : (startNounce + thread);
uint32_t *pHash = &g_hash[hashPosition << 4]; off_t hashPosition = nounce - startNounce;
uint32_t *pHash = &g_hash[hashPosition << 4];
const uint32_t thr = threadIdx.x & 0x3; // % THF
const uint32_t thr = threadIdx.x & 0x3; // % THF
#pragma unroll
for(int k=0;k<4;k++) message[k] = pHash[thr + (k * THF)]; /*| M0 M1 M2 M3 | M4 M5 M6 M7 | (input)
--|-------------|-------------|
#pragma unroll T0| 0 4 8 12 | 80 |
for(int k=4;k<8;k++) message[k] = 0; T1| 1 5 13 | |
T2| 2 6 14 | |
if (thr == 0) message[4] = 0x80U; T3| 3 7 15 | 01 |
if (thr == 3) message[7] = 0x01000000U; --|-------------|-------------| */
uint32_t msgBitsliced[8]; #pragma unroll
to_bitslice_quad(message, msgBitsliced); for(int k=0;k<4;k++) message[k] = pHash[thr + (k * THF)];
groestl512_progressMessage_quad(state, msgBitsliced); #pragma unroll
for(int k=4;k<8;k++) message[k] = 0;
// Nur der erste von jeweils 4 Threads bekommt das Ergebns-Hash
uint32_t __align__(16) hash[16]; if (thr == 0) message[4] = 0x80U; // end of data tag
from_bitslice_quad(state, hash); if (thr == 3) message[7] = 0x01000000U;
// uint4 = 4x4 uint32_t = 16 bytes uint32_t msgBitsliced[8];
if (thr == 0) { to_bitslice_quad(message, msgBitsliced);
uint4 *phash = (uint4*) hash;
uint4 *outpt = (uint4*) pHash; groestl512_progressMessage_quad(state, msgBitsliced);
outpt[0] = phash[0];
outpt[1] = phash[1]; uint32_t hash[16];
outpt[2] = phash[2]; from_bitslice_quad(state, hash);
outpt[3] = phash[3];
} // uint4 = 4x4 uint32_t = 16 bytes
} if (thr == 0) {
uint4 *phash = (uint4*) hash;
uint4 *outpt = (uint4*) pHash;
outpt[0] = phash[0];
outpt[1] = phash[1];
outpt[2] = phash[2];
outpt[3] = phash[3];
}
}
#endif #endif
} }
__host__ __host__
void quark_groestl512_cpu_init(int thr_id, uint32_t threads) void quark_groestl512_cpu_init(int thr_id, uint32_t threads)
{ {
int dev_id = device_map[thr_id]; int dev_id = device_map[thr_id];
cuda_get_arch(thr_id); cuda_get_arch(thr_id);
if (device_sm[dev_id] < 300 || cuda_arch[dev_id] < 300) if (device_sm[dev_id] < 300 || cuda_arch[dev_id] < 300)
quark_groestl512_sm20_init(thr_id, threads); quark_groestl512_sm20_init(thr_id, threads);
} }
__host__ __host__
void quark_groestl512_cpu_free(int thr_id) void quark_groestl512_cpu_free(int thr_id)
{ {
int dev_id = device_map[thr_id]; int dev_id = device_map[thr_id];
if (device_sm[dev_id] < 300 || cuda_arch[dev_id] < 300) if (device_sm[dev_id] < 300 || cuda_arch[dev_id] < 300)
quark_groestl512_sm20_free(thr_id); quark_groestl512_sm20_free(thr_id);
} }
__host__ __host__
void quark_groestl512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) void quark_groestl512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{ {
uint32_t threadsperblock = TPB; uint32_t threadsperblock = TPB;
// Compute 3.0 benutzt die registeroptimierte Quad Variante mit Warp Shuffle
// mit den Quad Funktionen brauchen wir jetzt 4 threads pro Hash, daher Faktor 4 bei der Blockzahl
const uint32_t factor = THF;
// berechne wie viele Thread Blocks wir brauchen // Compute 3.0 benutzt die registeroptimierte Quad Variante mit Warp Shuffle
dim3 grid(factor*((threads + threadsperblock-1)/threadsperblock)); // mit den Quad Funktionen brauchen wir jetzt 4 threads pro Hash, daher Faktor 4 bei der Blockzahl
dim3 block(threadsperblock); const uint32_t factor = THF;
int dev_id = device_map[thr_id]; dim3 grid(factor*((threads + threadsperblock-1)/threadsperblock));
dim3 block(threadsperblock);
if (device_sm[dev_id] >= 300 && cuda_arch[dev_id] >= 300) int dev_id = device_map[thr_id];
quark_groestl512_gpu_hash_64_quad<<<grid, block>>>(threads, startNounce, d_hash, d_nonceVector);
else
quark_groestl512_sm20_hash_64(thr_id, threads, startNounce, d_nonceVector, d_hash, order);
// Strategisches Sleep Kommando zur Senkung der CPU Last if (device_sm[dev_id] >= 300 && cuda_arch[dev_id] >= 300)
// MyStreamSynchronize(NULL, order, thr_id); quark_groestl512_gpu_hash_64_quad<<<grid, block>>>(threads, startNounce, d_hash, d_nonceVector);
else
quark_groestl512_sm20_hash_64(thr_id, threads, startNounce, d_nonceVector, d_hash, order);
} }

267
quark/cuda_quark_groestl512_sm2.cuh
Unescape View File

@ -0,0 +1,267 @@
// SM 2.x variant (tpruvot)
#ifdef __INTELLISENSE__
//#define __CUDA_ARCH__ 210
#define __CUDACC__
#include <cuda_helper.h>
#include <cuda_texture_types.h>
#define __byte_perm(a,b,c) (a)
#define tex1Dfetch(t, n) (n)
#endif
#define USE_SHARED 1
static unsigned int *d_textures[MAX_GPUS][8];
#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 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))))
texture<unsigned int, 1, cudaReadModeElementType> t0up1;
texture<unsigned int, 1, cudaReadModeElementType> t0dn1;
texture<unsigned int, 1, cudaReadModeElementType> t1up1;
texture<unsigned int, 1, cudaReadModeElementType> t1dn1;
texture<unsigned int, 1, cudaReadModeElementType> t2up1;
texture<unsigned int, 1, cudaReadModeElementType> t2dn1;
texture<unsigned int, 1, cudaReadModeElementType> t3up1;
texture<unsigned int, 1, cudaReadModeElementType> t3dn1;
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[];
#if __CUDA_ARCH__ < 300 || defined(_DEBUG)
#if (!USE_SHARED)
#include "groestl_simple.cuh"
#endif
__device__ __forceinline__
void quark_groestl512_perm_P(uint32_t *a, char *mixtabs)
{
#pragma unroll 1
for(int r=0; r<14; r++)
{
uint32_t t[32];
#pragma unroll 16
for (int k=0; k<16; k++)
a[(k*2)+0] ^= PC32up(k<< 4, r);
#pragma unroll 16
for(int k=0;k<32;k+=2) {
uint32_t t0_0 = B32_0(a[(k ) & 0x1f]), t9_0 = B32_0(a[(k + 9) & 0x1f]);
uint32_t t2_1 = B32_1(a[(k + 2) & 0x1f]), t11_1 = B32_1(a[(k + 11) & 0x1f]);
uint32_t t4_2 = B32_2(a[(k + 4) & 0x1f]), t13_2 = B32_2(a[(k + 13) & 0x1f]);
uint32_t t6_3 = B32_3(a[(k + 6) & 0x1f]), t23_3 = B32_3(a[(k + 23) & 0x1f]);
t[k + 0] = T0up( t0_0 ) ^ T1up( t2_1 ) ^ T2up( t4_2 ) ^ T3up( t6_3 ) ^
T0dn( t9_0 ) ^ T1dn( t11_1 ) ^ T2dn( t13_2 ) ^ T3dn( t23_3 );
t[k + 1] = T0dn( t0_0 ) ^ T1dn( t2_1 ) ^ T2dn( t4_2 ) ^ T3dn( t6_3 ) ^
T0up( t9_0 ) ^ T1up( t11_1 ) ^ T2up( t13_2 ) ^ T3up( t23_3 );
}
#pragma unroll 32
for(int k=0; k<32; k++)
a[k] = t[k];
}
}
__device__ __forceinline__
void quark_groestl512_perm_Q(uint32_t *a, char *mixtabs)
{
#pragma unroll 1
for(int r=0; r<14; r++)
{
uint32_t t[32];
#pragma unroll 16
for (int k=0; k<16; k++) {
a[(k*2)+0] ^= QC32up(k << 4, r);
a[(k*2)+1] ^= QC32dn(k << 4, r);
}
#pragma unroll 16
for(int k=0;k<32;k+=2)
{
uint32_t t2_0 = B32_0(a[(k + 2) & 0x1f]), t1_0 = B32_0(a[(k + 1) & 0x1f]);
uint32_t t6_1 = B32_1(a[(k + 6) & 0x1f]), t5_1 = B32_1(a[(k + 5) & 0x1f]);
uint32_t t10_2 = B32_2(a[(k + 10) & 0x1f]), t9_2 = B32_2(a[(k + 9) & 0x1f]);
uint32_t t22_3 = B32_3(a[(k + 22) & 0x1f]), t13_3 = B32_3(a[(k + 13) & 0x1f]);
t[k + 0] = T0up( t2_0 ) ^ T1up( t6_1 ) ^ T2up( t10_2 ) ^ T3up( t22_3 ) ^
T0dn( t1_0 ) ^ T1dn( t5_1 ) ^ T2dn( t9_2 ) ^ T3dn( t13_3 );
t[k + 1] = T0dn( t2_0 ) ^ T1dn( t6_1 ) ^ T2dn( t10_2 ) ^ T3dn( t22_3 ) ^
T0up( t1_0 ) ^ T1up( t5_1 ) ^ T2up( t9_2 ) ^ T3up( t13_3 );
}
#pragma unroll 32
for(int k=0; k<32; k++)
a[k] = t[k];
}
}
#endif
__global__
void quark_groestl512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint32_t *g_hash, uint32_t *g_nonceVector)
{
#if __CUDA_ARCH__ < 300 || defined(_DEBUG)
#if USE_SHARED
__shared__ char mixtabs[8 * 1024];
if (threadIdx.x < 256) {
*((uint32_t*)mixtabs + ( threadIdx.x)) = tex1Dfetch(t0up1, threadIdx.x);
*((uint32_t*)mixtabs + ( 256+threadIdx.x)) = tex1Dfetch(t0dn1, threadIdx.x);
*((uint32_t*)mixtabs + ( 512+threadIdx.x)) = tex1Dfetch(t1up1, threadIdx.x);
*((uint32_t*)mixtabs + ( 768+threadIdx.x)) = tex1Dfetch(t1dn1, threadIdx.x);
*((uint32_t*)mixtabs + (1024+threadIdx.x)) = tex1Dfetch(t2up1, threadIdx.x);
*((uint32_t*)mixtabs + (1280+threadIdx.x)) = tex1Dfetch(t2dn1, threadIdx.x);
*((uint32_t*)mixtabs + (1536+threadIdx.x)) = tex1Dfetch(t3up1, threadIdx.x);
*((uint32_t*)mixtabs + (1792+threadIdx.x)) = tex1Dfetch(t3dn1, threadIdx.x);
}
__syncthreads();
#endif
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
// GROESTL
uint32_t message[32];
uint32_t state[32];
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
off_t hashPosition = nounce - startNounce;
uint32_t *pHash = &g_hash[hashPosition * 16];
#pragma unroll 4
for (int i=0; i<16; i += 4)
AS_UINT4(&message[i]) = AS_UINT4(&pHash[i]);
message[16] = 0x80U;
#pragma unroll 14
for(int i=17; i<31; i++) message[i] = 0;
message[31] = 0x01000000U;
#pragma unroll 32
for(int i=0; i<32; i++) state[i] = message[i];
state[31] ^= 0x20000U;
// Perm
#if USE_SHARED
quark_groestl512_perm_P(state, mixtabs);
state[31] ^= 0x20000U;
quark_groestl512_perm_Q(message, mixtabs);
#pragma unroll 32
for(int i=0; i<32; i++) state[i] ^= message[i];
#pragma unroll 16
for(int i=16; i<32; i++) message[i] = state[i];
quark_groestl512_perm_P(state, mixtabs);
#else
tex_groestl512_perm_P(state);
state[31] ^= 0x20000U;
tex_groestl512_perm_Q(message);
#pragma unroll 32
for(int i=0; i<32; i++) state[i] ^= message[i];
#pragma unroll 16
for(int i=16; i<32; i++) message[i] = state[i];
tex_groestl512_perm_P(state);
#endif
#pragma unroll 16
for(int i=16; i<32; i++) state[i] ^= message[i];
uint4 *outpt = (uint4*)(pHash);
uint4 *phash = (uint4*)(&state[16]);
outpt[0] = phash[0];
outpt[1] = phash[1];
outpt[2] = phash[2];
outpt[3] = phash[3];
}
#endif
}
#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 quark_groestl512_sm20_init(int thr_id, uint32_t threads)
{
texDef(0, t0up1, d_T0up, T0up_cpu, sizeof(uint32_t)*256);
texDef(1, t0dn1, d_T0dn, T0dn_cpu, sizeof(uint32_t)*256);
texDef(2, t1up1, d_T1up, T1up_cpu, sizeof(uint32_t)*256);
texDef(3, t1dn1, d_T1dn, T1dn_cpu, sizeof(uint32_t)*256);
texDef(4, t2up1, d_T2up, T2up_cpu, sizeof(uint32_t)*256);
texDef(5, t2dn1, d_T2dn, T2dn_cpu, sizeof(uint32_t)*256);
texDef(6, t3up1, d_T3up, T3up_cpu, sizeof(uint32_t)*256);
texDef(7, t3dn1, d_T3dn, T3dn_cpu, sizeof(uint32_t)*256);
}
__host__
void quark_groestl512_sm20_free(int thr_id)
{
if (!d_textures[thr_id][0]) return;
for (int i=0; i<8; i++)
cudaFree(d_textures[thr_id][i]);
d_textures[thr_id][0] = NULL;
}
__host__
void quark_groestl512_sm20_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{
int threadsperblock = 512;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
quark_groestl512_gpu_hash_64<<<grid, block>>>(threads, startNounce, d_hash, d_nonceVector);
}
__host__
void quark_doublegroestl512_sm20_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{
int threadsperblock = 512;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
quark_groestl512_gpu_hash_64<<<grid, block>>>(threads, startNounce, d_hash, d_nonceVector);
quark_groestl512_gpu_hash_64<<<grid, block>>>(threads, startNounce, d_hash, d_nonceVector);
}

342
quark/cuda_quark_groestl512_sm20.cu
Unescape View File

@ -1,342 +0,0 @@
// SM 2.1 variant
// #include "cuda_helper.h"
#define MAXWELL_OR_FERMI 0
#define USE_SHARED 1
static unsigned int *d_textures[MAX_GPUS][8];
// #define SPH_C32(x) ((uint32_t)(x ## U))
// #define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF))
#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)
// 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(t0dn1, x)
#define T1up(x) tex1Dfetch(t1up1, x)
#define T1dn(x) (*((uint32_t*)mixtabs + (768+(x))))
#define T2up(x) tex1Dfetch(t2up1, x)
#define T2dn(x) (*((uint32_t*)mixtabs + (1280+(x))))
#define T3up(x) (*((uint32_t*)mixtabs + (1536+(x))))
#define T3dn(x) tex1Dfetch(t3dn1, x)
texture<unsigned int, 1, cudaReadModeElementType> t0up1;
texture<unsigned int, 1, cudaReadModeElementType> t0dn1;
texture<unsigned int, 1, cudaReadModeElementType> t1up1;
texture<unsigned int, 1, cudaReadModeElementType> t1dn1;
texture<unsigned int, 1, cudaReadModeElementType> t2up1;
texture<unsigned int, 1, cudaReadModeElementType> t2dn1;
texture<unsigned int, 1, cudaReadModeElementType> t3up1;
texture<unsigned int, 1, cudaReadModeElementType> t3dn1;
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[];
#if __CUDA_ARCH__ < 300 || defined(_DEBUG)
__device__ __forceinline__
void quark_groestl512_perm_P(uint32_t *a, char *mixtabs)
{
uint32_t t[32];
for(int r=0; r<14; r++)
{
switch(r) {
case 0:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 0); break;
case 1:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 1); break;
case 2:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 2); break;
case 3:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 3); break;
case 4:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 4); break;
case 5:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 5); break;
case 6:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 6); break;
case 7:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 7); break;
case 8:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 8); break;
case 9:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 9); break;
case 10:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 10); break;
case 11:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 11); break;
case 12:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 12); break;
case 13:
#pragma unroll 16
for(int k=0;k<16;k++) a[(k*2)+0] ^= PC32up(k<< 4, 13); break;
}
// RBTT
#pragma unroll 16
for(int k=0;k<32;k+=2) {
uint32_t t0_0 = B32_0(a[(k ) & 0x1f]), t9_0 = B32_0(a[(k + 9) & 0x1f]);
uint32_t t2_1 = B32_1(a[(k + 2) & 0x1f]), t11_1 = B32_1(a[(k + 11) & 0x1f]);
uint32_t t4_2 = B32_2(a[(k + 4) & 0x1f]), t13_2 = B32_2(a[(k + 13) & 0x1f]);
uint32_t t6_3 = B32_3(a[(k + 6) & 0x1f]), t23_3 = B32_3(a[(k + 23) & 0x1f]);
t[k + 0] = T0up( t0_0 ) ^ T1up( t2_1 ) ^ T2up( t4_2 ) ^ T3up( t6_3 ) ^
T0dn( t9_0 ) ^ T1dn( t11_1 ) ^ T2dn( t13_2 ) ^ T3dn( t23_3 );
t[k + 1] = T0dn( t0_0 ) ^ T1dn( t2_1 ) ^ T2dn( t4_2 ) ^ T3dn( t6_3 ) ^
T0up( t9_0 ) ^ T1up( t11_1 ) ^ T2up( t13_2 ) ^ T3up( t23_3 );
}
#pragma unroll 32
for(int k=0; k<32; k++) {
a[k] = t[k];
}
}
}
__device__ __forceinline__
void quark_groestl512_perm_Q(uint32_t *a, char *mixtabs)
{
for(int r=0; r<14; r++)
{
uint32_t t[32];
switch(r) {
case 0:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 0); a[(k*2)+1] ^= QC32dn(k<< 4, 0);} break;
case 1:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 1); a[(k*2)+1] ^= QC32dn(k<< 4, 1);} break;
case 2:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 2); a[(k*2)+1] ^= QC32dn(k<< 4, 2);} break;
case 3:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 3); a[(k*2)+1] ^= QC32dn(k<< 4, 3);} break;
case 4:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 4); a[(k*2)+1] ^= QC32dn(k<< 4, 4);} break;
case 5:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 5); a[(k*2)+1] ^= QC32dn(k<< 4, 5);} break;
case 6:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 6); a[(k*2)+1] ^= QC32dn(k<< 4, 6);} break;
case 7:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 7); a[(k*2)+1] ^= QC32dn(k<< 4, 7);} break;
case 8:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 8); a[(k*2)+1] ^= QC32dn(k<< 4, 8);} break;
case 9:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 9); a[(k*2)+1] ^= QC32dn(k<< 4, 9);} break;
case 10:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 10); a[(k*2)+1] ^= QC32dn(k<< 4, 10);} break;
case 11:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 11); a[(k*2)+1] ^= QC32dn(k<< 4, 11);} break;
case 12:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 12); a[(k*2)+1] ^= QC32dn(k<< 4, 12);} break;
case 13:
#pragma unroll 16
for(int k=0;k<16;k++) { a[(k*2)+0] ^= QC32up(k<< 4, 13); a[(k*2)+1] ^= QC32dn(k<< 4, 13);} break;
}
// RBTT
#pragma unroll 16
for(int k=0;k<32;k+=2)
{
uint32_t t2_0 = B32_0(a[(k + 2) & 0x1f]), t1_0 = B32_0(a[(k + 1) & 0x1f]);
uint32_t t6_1 = B32_1(a[(k + 6) & 0x1f]), t5_1 = B32_1(a[(k + 5) & 0x1f]);
uint32_t t10_2 = B32_2(a[(k + 10) & 0x1f]), t9_2 = B32_2(a[(k + 9) & 0x1f]);
uint32_t t22_3 = B32_3(a[(k + 22) & 0x1f]), t13_3 = B32_3(a[(k + 13) & 0x1f]);
t[k + 0] = T0up( t2_0 ) ^ T1up( t6_1 ) ^ T2up( t10_2 ) ^ T3up( t22_3 ) ^
T0dn( t1_0 ) ^ T1dn( t5_1 ) ^ T2dn( t9_2 ) ^ T3dn( t13_3 );
t[k + 1] = T0dn( t2_0 ) ^ T1dn( t6_1 ) ^ T2dn( t10_2 ) ^ T3dn( t22_3 ) ^
T0up( t1_0 ) ^ T1up( t5_1 ) ^ T2up( t9_2 ) ^ T3up( t13_3 );
}
#pragma unroll 32
for(int k=0;k<32;k++)
a[k] = t[k];
}
}
#endif
__global__
void quark_groestl512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint32_t *g_hash, uint32_t *g_nonceVector)
{
#if __CUDA_ARCH__ < 300 || defined(_DEBUG)
extern __shared__ char mixtabs[];
if (threadIdx.x < 256)
{
*((uint32_t*)mixtabs + ( threadIdx.x)) = tex1Dfetch(t0up1, threadIdx.x);
*((uint32_t*)mixtabs + (256+threadIdx.x)) = tex1Dfetch(t0dn1, threadIdx.x);
*((uint32_t*)mixtabs + (512+threadIdx.x)) = tex1Dfetch(t1up1, threadIdx.x);
*((uint32_t*)mixtabs + (768+threadIdx.x)) = tex1Dfetch(t1dn1, threadIdx.x);
*((uint32_t*)mixtabs + (1024+threadIdx.x)) = tex1Dfetch(t2up1, threadIdx.x);
*((uint32_t*)mixtabs + (1280+threadIdx.x)) = tex1Dfetch(t2dn1, threadIdx.x);
*((uint32_t*)mixtabs + (1536+threadIdx.x)) = tex1Dfetch(t3up1, threadIdx.x);
*((uint32_t*)mixtabs + (1792+threadIdx.x)) = tex1Dfetch(t3dn1, threadIdx.x);
}
__syncthreads();
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
// GROESTL
uint32_t message[32];
uint32_t state[32];
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
off_t hashPosition = nounce - startNounce;
uint32_t *inpHash = &g_hash[hashPosition * 16];
#pragma unroll 16
for(int k=0; k<16; k++)
message[k] = inpHash[k];
#pragma unroll 14
for(int k=1; k<15; k++)
message[k+16] = 0;
message[16] = 0x80;
message[31] = 0x01000000;
#pragma unroll 32
for(int u=0; u<32; u++)
state[u] = message[u];
state[31] ^= 0x20000;
// Perm
quark_groestl512_perm_P(state, mixtabs);
state[31] ^= 0x20000;
quark_groestl512_perm_Q(message, mixtabs);
#pragma unroll 32
for(int u=0;u<32;u++) state[u] ^= message[u];
#pragma unroll 32
for(int u=0;u<32;u++) message[u] = state[u];
quark_groestl512_perm_P(message, mixtabs);
#pragma unroll 32
for(int u=0;u<32;u++) state[u] ^= message[u];
// Erzeugten Hash rausschreiben
uint32_t *outpHash = &g_hash[hashPosition * 16];
#pragma unroll 16
for(int k=0;k<16;k++) outpHash[k] = state[k+16];
}
#endif
}
#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 quark_groestl512_sm20_init(int thr_id, uint32_t threads)
{
// Texturen mit obigem Makro initialisieren
texDef(0, t0up1, d_T0up, T0up_cpu, sizeof(uint32_t)*256);
texDef(1, t0dn1, d_T0dn, T0dn_cpu, sizeof(uint32_t)*256);
texDef(2, t1up1, d_T1up, T1up_cpu, sizeof(uint32_t)*256);
texDef(3, t1dn1, d_T1dn, T1dn_cpu, sizeof(uint32_t)*256);
texDef(4, t2up1, d_T2up, T2up_cpu, sizeof(uint32_t)*256);
texDef(5, t2dn1, d_T2dn, T2dn_cpu, sizeof(uint32_t)*256);
texDef(6, t3up1, d_T3up, T3up_cpu, sizeof(uint32_t)*256);
texDef(7, t3dn1, d_T3dn, T3dn_cpu, sizeof(uint32_t)*256);
}
__host__
void quark_groestl512_sm20_free(int thr_id)
{
for (int i=0; i<8; i++)
cudaFree(d_textures[thr_id][i]);
}
__host__
void quark_groestl512_sm20_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{
int threadsperblock = 512;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
size_t shared_size = 8 * 256 * sizeof(uint32_t);
quark_groestl512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector);
// MyStreamSynchronize(NULL, order, thr_id);
}
__host__
void quark_doublegroestl512_sm20_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{
int threadsperblock = 512;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
size_t shared_size = 8 * 256 * sizeof(uint32_t);
quark_groestl512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector);
quark_groestl512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector);
// MyStreamSynchronize(NULL, order, thr_id);
}

111
quark/groestl_simple.cuh
Unescape View File

@ -0,0 +1,111 @@
/**
* Simplified groestl512 big perm code
* tpruvot - 2017
*/
#ifdef __INTELLISENSE__
#define __CUDA_ARCH__ 210
#define __CUDACC__
#include <cuda_helper.h>
#include <cuda_texture_types.h>
#define __byte_perm(a,b,c) (a)
#define tex1Dfetch(t, n) (n)
#endif
// todo: merge with cuda_quark_groestl512_sm20.cu (used for groestl512-80)
#if __CUDA_ARCH__ < 300 || defined(_DEBUG)
#ifndef SPH_C32
#define SPH_C32(x) ((uint32_t)(x ## U))
#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF))
#endif
#ifndef PC32up
#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)))
#endif
#define tT0up(x) tex1Dfetch(t0up1, x)
#define tT0dn(x) tex1Dfetch(t0dn1, x)
#define tT1up(x) tex1Dfetch(t1up1, x)
#define tT1dn(x) tex1Dfetch(t1dn1, x)
#define tT2up(x) tex1Dfetch(t2up1, x)
#define tT2dn(x) tex1Dfetch(t2dn1, x)
#define tT3up(x) tex1Dfetch(t3up1, x)
#define tT3dn(x) tex1Dfetch(t3dn1, x)
#undef B32_0
#define B32_0(x) ((x) & 0xFFu)
__device__ __forceinline__
static void tex_groestl512_perm_P(uint32_t *a)
{
#pragma unroll 1
for(int r=0; r<14; r++)
{
uint32_t t[32];
#pragma unroll 16
for (int k=0; k<16; k++)
a[(k*2)+0] ^= PC32up(k<< 4, r);
#pragma unroll 16
for(int k=0; k<32; k+=2)
{
uint32_t t0_0 = B32_0(a[(k ) & 0x1f]), t9_0 = B32_0(a[(k + 9) & 0x1f]);
uint32_t t2_1 = B32_1(a[(k + 2) & 0x1f]), t11_1 = B32_1(a[(k + 11) & 0x1f]);
uint32_t t4_2 = B32_2(a[(k + 4) & 0x1f]), t13_2 = B32_2(a[(k + 13) & 0x1f]);
uint32_t t6_3 = B32_3(a[(k + 6) & 0x1f]), t23_3 = B32_3(a[(k + 23) & 0x1f]);
t[k + 0] = tT0up( t0_0 ) ^ tT1up( t2_1 ) ^ tT2up( t4_2 ) ^ tT3up( t6_3 ) ^
tT0dn( t9_0 ) ^ tT1dn( t11_1 ) ^ tT2dn( t13_2 ) ^ tT3dn( t23_3 );
t[k + 1] = tT0dn( t0_0 ) ^ tT1dn( t2_1 ) ^ tT2dn( t4_2 ) ^ tT3dn( t6_3 ) ^
tT0up( t9_0 ) ^ tT1up( t11_1 ) ^ tT2up( t13_2 ) ^ tT3up( t23_3 );
}
#pragma unroll 32
for(int k=0; k<32; k++)
a[k] = t[k];
}
}
__device__ __forceinline__
static void tex_groestl512_perm_Q(uint32_t *a)
{
#pragma unroll 1
for(int r=0; r<14; r++)
{
uint32_t t[32];
#pragma unroll 16
for (int k=0; k<16; k++) {
a[(k*2)+0] ^= QC32up(k<< 4, r);
a[(k*2)+1] ^= QC32dn(k<< 4, r);
}
#pragma unroll 16
for(int k=0; k<32; k+=2)
{
uint32_t t2_0 = B32_0(a[(k + 2) & 0x1f]), t1_0 = B32_0(a[(k + 1) & 0x1f]);
uint32_t t6_1 = B32_1(a[(k + 6) & 0x1f]), t5_1 = B32_1(a[(k + 5) & 0x1f]);
uint32_t t10_2 = B32_2(a[(k + 10) & 0x1f]), t9_2 = B32_2(a[(k + 9) & 0x1f]);
uint32_t t22_3 = B32_3(a[(k + 22) & 0x1f]), t13_3 = B32_3(a[(k + 13) & 0x1f]);
t[k + 0] = tT0up( t2_0 ) ^ tT1up( t6_1 ) ^ tT2up( t10_2 ) ^ tT3up( t22_3 ) ^
tT0dn( t1_0 ) ^ tT1dn( t5_1 ) ^ tT2dn( t9_2 ) ^ tT3dn( t13_3 );
t[k + 1] = tT0dn( t2_0 ) ^ tT1dn( t6_1 ) ^ tT2dn( t10_2 ) ^ tT3dn( t22_3 ) ^
tT0up( t1_0 ) ^ tT1up( t5_1 ) ^ tT2up( t9_2 ) ^ tT3up( t13_3 );
}
#pragma unroll 32
for(int k=0; k<32; k++)
a[k] = t[k];
}
}
#endif
Write Preview
Loading…
Cancel
Save