mirror of https://github.com/GOSTSec/ccminer
Christian Buchner
10 years ago
15 changed files with 2078 additions and 88 deletions
@ -0,0 +1,255 @@
@@ -0,0 +1,255 @@
|
||||
#include <cuda.h> |
||||
#include "cuda_runtime.h" |
||||
#include "device_launch_parameters.h" |
||||
#include "sm_30_intrinsics.h" |
||||
|
||||
#include <stdio.h> |
||||
#include <memory.h> |
||||
#include <stdint.h> |
||||
|
||||
// aus cpu-miner.c |
||||
extern int device_map[8]; |
||||
|
||||
// diese Struktur wird in der Init Funktion angefordert |
||||
static cudaDeviceProp props[8]; |
||||
|
||||
static uint32_t *d_tempBranch1Nonces[8]; |
||||
static uint32_t *d_tempBranch2Nonces[8]; |
||||
static size_t *d_numValid[8]; |
||||
static size_t *h_numValid[8]; |
||||
|
||||
static uint32_t *d_partSum1[8], *d_partSum2[8]; // 2x partielle summen |
||||
static uint32_t *d_validTemp1[8], *d_validTemp2[8]; |
||||
|
||||
// Zwischenspeicher |
||||
static uint32_t *d_tempBranchAllNonces[8]; |
||||
|
||||
// aus heavy.cu |
||||
extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); |
||||
|
||||
|
||||
|
||||
// Setup-Funktionen |
||||
__host__ void jackpot_compactTest_cpu_init(int thr_id, int threads) |
||||
{ |
||||
cudaGetDeviceProperties(&props[thr_id], device_map[thr_id]); |
||||
|
||||
// wir brauchen auch Speicherplatz auf dem Device |
||||
cudaMalloc(&d_tempBranchAllNonces[thr_id], sizeof(uint32_t) * threads); |
||||
cudaMalloc(&d_tempBranch1Nonces[thr_id], sizeof(uint32_t) * threads); |
||||
cudaMalloc(&d_tempBranch2Nonces[thr_id], sizeof(uint32_t) * threads); |
||||
cudaMalloc(&d_numValid[thr_id], 2*sizeof(size_t)); |
||||
cudaMallocHost(&h_numValid[thr_id], 2*sizeof(size_t)); |
||||
|
||||
uint32_t s1; |
||||
s1 = threads / 256; |
||||
|
||||
cudaMalloc(&d_partSum1[thr_id], sizeof(uint32_t) * s1); // BLOCKSIZE (Threads/Block) |
||||
cudaMalloc(&d_partSum2[thr_id], sizeof(uint32_t) * s1); // BLOCKSIZE (Threads/Block) |
||||
|
||||
cudaMalloc(&d_validTemp1[thr_id], sizeof(uint32_t) * threads); // BLOCKSIZE (Threads/Block) |
||||
cudaMalloc(&d_validTemp2[thr_id], sizeof(uint32_t) * threads); // BLOCKSIZE (Threads/Block) |
||||
} |
||||
|
||||
// Die Testfunktion (zum Erstellen der TestMap) |
||||
__global__ void jackpot_compactTest_gpu_TEST_64(int threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_noncesFull, |
||||
uint32_t *d_nonces1, uint32_t *d_nonces2, |
||||
uint32_t *d_validT1, uint32_t *d_validT2) |
||||
{ |
||||
int thread = (blockDim.x * blockIdx.x + threadIdx.x); |
||||
if (thread < threads) |
||||
{ |
||||
// bestimme den aktuellen Zähler |
||||
uint32_t nounce = startNounce + thread; |
||||
uint32_t *inpHash = &inpHashes[16 * thread]; |
||||
|
||||
uint32_t tmp = inpHash[0] & 0x01; |
||||
uint32_t val1 = (tmp == 1); |
||||
uint32_t val2 = (tmp == 0); |
||||
|
||||
d_nonces1[thread] = val1; |
||||
d_validT1[thread] = val1; |
||||
d_nonces2[thread] = val2; |
||||
d_validT2[thread] = val2; |
||||
d_noncesFull[thread] = nounce; |
||||
} |
||||
} |
||||
|
||||
// Die Summenfunktion (vom NVIDIA SDK) |
||||
__global__ void jackpot_compactTest_gpu_SCAN(uint32_t *data, int width, uint32_t *partial_sums=NULL) |
||||
{ |
||||
extern __shared__ uint32_t sums[]; |
||||
int id = ((blockIdx.x * blockDim.x) + threadIdx.x); |
||||
//int lane_id = id % warpSize; |
||||
int lane_id = id % width; |
||||
// determine a warp_id within a block |
||||
//int warp_id = threadIdx.x / warpSize; |
||||
int warp_id = threadIdx.x / width; |
||||
|
||||
// Below is the basic structure of using a shfl instruction |
||||
// for a scan. |
||||
// Record "value" as a variable - we accumulate it along the way |
||||
uint32_t value = data[id]; |
||||
|
||||
// Now accumulate in log steps up the chain |
||||
// compute sums, with another thread's value who is |
||||
// distance delta away (i). Note |
||||
// those threads where the thread 'i' away would have |
||||
// been out of bounds of the warp are unaffected. This |
||||
// creates the scan sum. |
||||
#pragma unroll |
||||
|
||||
for (int i=1; i<=width; i*=2) |
||||
{ |
||||
uint32_t n = __shfl_up((int)value, i, width); |
||||
|
||||
if (lane_id >= i) value += n; |
||||
} |
||||
|
||||
// value now holds the scan value for the individual thread |
||||
// next sum the largest values for each warp |
||||
|
||||
// write the sum of the warp to smem |
||||
//if (threadIdx.x % warpSize == warpSize-1) |
||||
if (threadIdx.x % width == width-1) |
||||
{ |
||||
sums[warp_id] = value; |
||||
} |
||||
|
||||
__syncthreads(); |
||||
|
||||
// |
||||
// scan sum the warp sums |
||||
// the same shfl scan operation, but performed on warp sums |
||||
// |
||||
if (warp_id == 0) |
||||
{ |
||||
uint32_t warp_sum = sums[lane_id]; |
||||
|
||||
for (int i=1; i<=width; i*=2) |
||||
{ |
||||
uint32_t n = __shfl_up((int)warp_sum, i, width); |
||||
|
||||
if (lane_id >= i) warp_sum += n; |
||||
} |
||||
|
||||
sums[lane_id] = warp_sum; |
||||
} |
||||
|
||||
__syncthreads(); |
||||
|
||||
// perform a uniform add across warps in the block |
||||
// read neighbouring warp's sum and add it to threads value |
||||
uint32_t blockSum = 0; |
||||
|
||||
if (warp_id > 0) |
||||
{ |
||||
blockSum = sums[warp_id-1]; |
||||
} |
||||
|
||||
value += blockSum; |
||||
|
||||
// Now write out our result |
||||
data[id] = value; |
||||
|
||||
// last thread has sum, write write out the block's sum |
||||
if (partial_sums != NULL && threadIdx.x == blockDim.x-1) |
||||
{ |
||||
partial_sums[blockIdx.x] = value; |
||||
} |
||||
} |
||||
|
||||
// Uniform add: add partial sums array |
||||
__global__ void jackpot_compactTest_gpu_ADD(uint32_t *data, uint32_t *partial_sums, int len) |
||||
{ |
||||
__shared__ uint32_t buf; |
||||
int id = ((blockIdx.x * blockDim.x) + threadIdx.x); |
||||
|
||||
if (id > len) return; |
||||
|
||||
if (threadIdx.x == 0) |
||||
{ |
||||
buf = partial_sums[blockIdx.x]; |
||||
} |
||||
|
||||
__syncthreads(); |
||||
data[id] += buf; |
||||
} |
||||
|
||||
// Der Scatter |
||||
__global__ void jackpot_compactTest_gpu_SCATTER(uint32_t *data, uint32_t *valid, uint32_t *sum, uint32_t *outp) |
||||
{ |
||||
int id = ((blockIdx.x * blockDim.x) + threadIdx.x); |
||||
if( valid[id] ) |
||||
{ |
||||
int idx = sum[id]; |
||||
if(idx > 0) |
||||
outp[idx-1] = data[id]; |
||||
} |
||||
} |
||||
|
||||
////// ACHTUNG: Diese funktion geht aktuell nur mit threads > 65536 (Am besten 256 * 1024 oder 256*2048) |
||||
__host__ void jackpot_compactTest_cpu_dualCompaction(int thr_id, int threads, size_t *nrm, |
||||
uint32_t *d_nonces1, uint32_t *d_nonces2) |
||||
{ |
||||
// threadsPerBlock ausrechnen |
||||
int blockSize = 256; |
||||
int thr1 = threads / blockSize; |
||||
int thr2 = threads / (blockSize*blockSize); |
||||
|
||||
// 1 |
||||
jackpot_compactTest_gpu_SCAN<<<thr1,blockSize, 8*sizeof(uint32_t)>>>(d_tempBranch1Nonces[thr_id], 32, d_partSum1[thr_id]); |
||||
jackpot_compactTest_gpu_SCAN<<<thr2,blockSize, 8*sizeof(uint32_t)>>>(d_partSum1[thr_id], 32, d_partSum2[thr_id]); |
||||
jackpot_compactTest_gpu_SCAN<<<1, thr2, 8*sizeof(uint32_t)>>>(d_partSum2[thr_id], (thr2>32) ? 32 : thr2); |
||||
cudaStreamSynchronize(NULL); |
||||
cudaMemcpy(&nrm[0], &(d_partSum2[thr_id])[thr2-1], sizeof(uint32_t), cudaMemcpyDeviceToHost); |
||||
jackpot_compactTest_gpu_ADD<<<thr2-1, blockSize>>>(d_partSum1[thr_id]+blockSize, d_partSum2[thr_id], blockSize*thr2); |
||||
jackpot_compactTest_gpu_ADD<<<thr1-1, blockSize>>>(d_tempBranch1Nonces[thr_id]+blockSize, d_partSum1[thr_id], threads); |
||||
|
||||
// 2 |
||||
jackpot_compactTest_gpu_SCAN<<<thr1,blockSize, 8*sizeof(uint32_t)>>>(d_tempBranch2Nonces[thr_id], 32, d_partSum1[thr_id]); |
||||
jackpot_compactTest_gpu_SCAN<<<thr2,blockSize, 8*sizeof(uint32_t)>>>(d_partSum1[thr_id], 32, d_partSum2[thr_id]); |
||||
jackpot_compactTest_gpu_SCAN<<<1, thr2, 8*sizeof(uint32_t)>>>(d_partSum2[thr_id], (thr2>32) ? 32 : thr2); |
||||
cudaStreamSynchronize(NULL); |
||||
cudaMemcpy(&nrm[1], &(d_partSum2[thr_id])[thr2-1], sizeof(uint32_t), cudaMemcpyDeviceToHost); |
||||
jackpot_compactTest_gpu_ADD<<<thr2-1, blockSize>>>(d_partSum1[thr_id]+blockSize, d_partSum2[thr_id], blockSize*thr2); |
||||
jackpot_compactTest_gpu_ADD<<<thr1-1, blockSize>>>(d_tempBranch2Nonces[thr_id]+blockSize, d_partSum1[thr_id], threads); |
||||
|
||||
// Hier ist noch eine Besonderheit: in d_tempBranch1Nonces sind die element von 1...nrm1 die Interessanten |
||||
// Schritt 3: Scatter |
||||
jackpot_compactTest_gpu_SCATTER<<<thr1,blockSize,0>>>(d_tempBranchAllNonces[thr_id], d_validTemp1[thr_id], d_tempBranch1Nonces[thr_id], d_nonces1); |
||||
jackpot_compactTest_gpu_SCATTER<<<thr1,blockSize,0>>>(d_tempBranchAllNonces[thr_id], d_validTemp2[thr_id], d_tempBranch2Nonces[thr_id], d_nonces2); |
||||
cudaStreamSynchronize(NULL); |
||||
} |
||||
|
||||
__host__ void jackpot_compactTest_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *inpHashes, |
||||
uint32_t *d_nonces1, size_t *nrm1, |
||||
uint32_t *d_nonces2, size_t *nrm2, |
||||
int order) |
||||
{ |
||||
// Compute 3.x und 5.x Geräte am besten mit 768 Threads ansteuern, |
||||
// alle anderen mit 512 Threads. |
||||
//int threadsperblock = (props[thr_id].major >= 3) ? 768 : 512; |
||||
int threadsperblock = 256; |
||||
|
||||
// berechne wie viele Thread Blocks wir brauchen |
||||
dim3 grid((threads + threadsperblock-1)/threadsperblock); |
||||
dim3 block(threadsperblock); |
||||
|
||||
size_t shared_size = 0; |
||||
|
||||
// fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size); |
||||
|
||||
// Schritt 1: Prüfen der Bedingung und Speicherung in d_tempBranch1/2Nonces |
||||
jackpot_compactTest_gpu_TEST_64<<<grid, block, shared_size>>>(threads, startNounce, inpHashes, d_tempBranchAllNonces[thr_id], |
||||
d_tempBranch1Nonces[thr_id], d_tempBranch2Nonces[thr_id], |
||||
d_validTemp1[thr_id], d_validTemp2[thr_id]); |
||||
|
||||
// Strategisches Sleep Kommando zur Senkung der CPU Last |
||||
jackpot_compactTest_cpu_dualCompaction(thr_id, threads, |
||||
h_numValid[thr_id], d_nonces1, d_nonces2); |
||||
|
||||
cudaStreamSynchronize(NULL); // Das original braucht zwar etwas CPU-Last, ist an dieser Stelle aber evtl besser |
||||
*nrm1 = h_numValid[thr_id][0]; |
||||
*nrm2 = h_numValid[thr_id][1]; |
||||
} |
||||
@ -0,0 +1,356 @@
@@ -0,0 +1,356 @@
|
||||
#include <stdint.h> |
||||
|
||||
// aus heavy.cu |
||||
extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); |
||||
|
||||
typedef struct { |
||||
uint32_t x[8][4]; /*the 1024-bit state, ( x[i][0] || x[i][1] || x[i][2] || x[i][3] ) is the ith row of the state in the pseudocode*/ |
||||
uint32_t buffer[16]; /*the 512-bit message block to be hashed;*/ |
||||
} hashState; |
||||
|
||||
/*42 round constants, each round constant is 32-byte (256-bit)*/ |
||||
__constant__ uint32_t c_INIT_bitslice[8][4]; |
||||
__constant__ unsigned char c_E8_bitslice_roundconstant[42][32]; |
||||
|
||||
const uint32_t h_INIT_bitslice[8][4] = { |
||||
{ 0x964bd16f, 0x17aa003e, 0x052e6a63, 0x43d5157a}, |
||||
{ 0x8d5e228a, 0x0bef970c, 0x591234e9, 0x61c3b3f2}, |
||||
{ 0xc1a01d89, 0x1e806f53, 0x6b05a92a, 0x806d2bea}, |
||||
{ 0xdbcc8e58, 0xa6ba7520, 0x763a0fa9, 0xf73bf8ba}, |
||||
{ 0x05e66901, 0x694ae341, 0x8e8ab546, 0x5ae66f2e}, |
||||
{ 0xd0a74710, 0x243c84c1, 0xb1716e3b, 0x99c15a2d}, |
||||
{ 0xecf657cf, 0x56f8b19d, 0x7c8806a7, 0x56b11657}, |
||||
{ 0xdffcc2e3, 0xfb1785e6, 0x78465a54, 0x4bdd8ccc} }; |
||||
|
||||
const unsigned char h_E8_bitslice_roundconstant[42][32]={ |
||||
{0x72,0xd5,0xde,0xa2,0xdf,0x15,0xf8,0x67,0x7b,0x84,0x15,0xa,0xb7,0x23,0x15,0x57,0x81,0xab,0xd6,0x90,0x4d,0x5a,0x87,0xf6,0x4e,0x9f,0x4f,0xc5,0xc3,0xd1,0x2b,0x40}, |
||||
{0xea,0x98,0x3a,0xe0,0x5c,0x45,0xfa,0x9c,0x3,0xc5,0xd2,0x99,0x66,0xb2,0x99,0x9a,0x66,0x2,0x96,0xb4,0xf2,0xbb,0x53,0x8a,0xb5,0x56,0x14,0x1a,0x88,0xdb,0xa2,0x31}, |
||||
{0x3,0xa3,0x5a,0x5c,0x9a,0x19,0xe,0xdb,0x40,0x3f,0xb2,0xa,0x87,0xc1,0x44,0x10,0x1c,0x5,0x19,0x80,0x84,0x9e,0x95,0x1d,0x6f,0x33,0xeb,0xad,0x5e,0xe7,0xcd,0xdc}, |
||||
{0x10,0xba,0x13,0x92,0x2,0xbf,0x6b,0x41,0xdc,0x78,0x65,0x15,0xf7,0xbb,0x27,0xd0,0xa,0x2c,0x81,0x39,0x37,0xaa,0x78,0x50,0x3f,0x1a,0xbf,0xd2,0x41,0x0,0x91,0xd3}, |
||||
{0x42,0x2d,0x5a,0xd,0xf6,0xcc,0x7e,0x90,0xdd,0x62,0x9f,0x9c,0x92,0xc0,0x97,0xce,0x18,0x5c,0xa7,0xb,0xc7,0x2b,0x44,0xac,0xd1,0xdf,0x65,0xd6,0x63,0xc6,0xfc,0x23}, |
||||
{0x97,0x6e,0x6c,0x3,0x9e,0xe0,0xb8,0x1a,0x21,0x5,0x45,0x7e,0x44,0x6c,0xec,0xa8,0xee,0xf1,0x3,0xbb,0x5d,0x8e,0x61,0xfa,0xfd,0x96,0x97,0xb2,0x94,0x83,0x81,0x97}, |
||||
{0x4a,0x8e,0x85,0x37,0xdb,0x3,0x30,0x2f,0x2a,0x67,0x8d,0x2d,0xfb,0x9f,0x6a,0x95,0x8a,0xfe,0x73,0x81,0xf8,0xb8,0x69,0x6c,0x8a,0xc7,0x72,0x46,0xc0,0x7f,0x42,0x14}, |
||||
{0xc5,0xf4,0x15,0x8f,0xbd,0xc7,0x5e,0xc4,0x75,0x44,0x6f,0xa7,0x8f,0x11,0xbb,0x80,0x52,0xde,0x75,0xb7,0xae,0xe4,0x88,0xbc,0x82,0xb8,0x0,0x1e,0x98,0xa6,0xa3,0xf4}, |
||||
{0x8e,0xf4,0x8f,0x33,0xa9,0xa3,0x63,0x15,0xaa,0x5f,0x56,0x24,0xd5,0xb7,0xf9,0x89,0xb6,0xf1,0xed,0x20,0x7c,0x5a,0xe0,0xfd,0x36,0xca,0xe9,0x5a,0x6,0x42,0x2c,0x36}, |
||||
{0xce,0x29,0x35,0x43,0x4e,0xfe,0x98,0x3d,0x53,0x3a,0xf9,0x74,0x73,0x9a,0x4b,0xa7,0xd0,0xf5,0x1f,0x59,0x6f,0x4e,0x81,0x86,0xe,0x9d,0xad,0x81,0xaf,0xd8,0x5a,0x9f}, |
||||
{0xa7,0x5,0x6,0x67,0xee,0x34,0x62,0x6a,0x8b,0xb,0x28,0xbe,0x6e,0xb9,0x17,0x27,0x47,0x74,0x7,0x26,0xc6,0x80,0x10,0x3f,0xe0,0xa0,0x7e,0x6f,0xc6,0x7e,0x48,0x7b}, |
||||
{0xd,0x55,0xa,0xa5,0x4a,0xf8,0xa4,0xc0,0x91,0xe3,0xe7,0x9f,0x97,0x8e,0xf1,0x9e,0x86,0x76,0x72,0x81,0x50,0x60,0x8d,0xd4,0x7e,0x9e,0x5a,0x41,0xf3,0xe5,0xb0,0x62}, |
||||
{0xfc,0x9f,0x1f,0xec,0x40,0x54,0x20,0x7a,0xe3,0xe4,0x1a,0x0,0xce,0xf4,0xc9,0x84,0x4f,0xd7,0x94,0xf5,0x9d,0xfa,0x95,0xd8,0x55,0x2e,0x7e,0x11,0x24,0xc3,0x54,0xa5}, |
||||
{0x5b,0xdf,0x72,0x28,0xbd,0xfe,0x6e,0x28,0x78,0xf5,0x7f,0xe2,0xf,0xa5,0xc4,0xb2,0x5,0x89,0x7c,0xef,0xee,0x49,0xd3,0x2e,0x44,0x7e,0x93,0x85,0xeb,0x28,0x59,0x7f}, |
||||
{0x70,0x5f,0x69,0x37,0xb3,0x24,0x31,0x4a,0x5e,0x86,0x28,0xf1,0x1d,0xd6,0xe4,0x65,0xc7,0x1b,0x77,0x4,0x51,0xb9,0x20,0xe7,0x74,0xfe,0x43,0xe8,0x23,0xd4,0x87,0x8a}, |
||||
{0x7d,0x29,0xe8,0xa3,0x92,0x76,0x94,0xf2,0xdd,0xcb,0x7a,0x9,0x9b,0x30,0xd9,0xc1,0x1d,0x1b,0x30,0xfb,0x5b,0xdc,0x1b,0xe0,0xda,0x24,0x49,0x4f,0xf2,0x9c,0x82,0xbf}, |
||||
{0xa4,0xe7,0xba,0x31,0xb4,0x70,0xbf,0xff,0xd,0x32,0x44,0x5,0xde,0xf8,0xbc,0x48,0x3b,0xae,0xfc,0x32,0x53,0xbb,0xd3,0x39,0x45,0x9f,0xc3,0xc1,0xe0,0x29,0x8b,0xa0}, |
||||
{0xe5,0xc9,0x5,0xfd,0xf7,0xae,0x9,0xf,0x94,0x70,0x34,0x12,0x42,0x90,0xf1,0x34,0xa2,0x71,0xb7,0x1,0xe3,0x44,0xed,0x95,0xe9,0x3b,0x8e,0x36,0x4f,0x2f,0x98,0x4a}, |
||||
{0x88,0x40,0x1d,0x63,0xa0,0x6c,0xf6,0x15,0x47,0xc1,0x44,0x4b,0x87,0x52,0xaf,0xff,0x7e,0xbb,0x4a,0xf1,0xe2,0xa,0xc6,0x30,0x46,0x70,0xb6,0xc5,0xcc,0x6e,0x8c,0xe6}, |
||||
{0xa4,0xd5,0xa4,0x56,0xbd,0x4f,0xca,0x0,0xda,0x9d,0x84,0x4b,0xc8,0x3e,0x18,0xae,0x73,0x57,0xce,0x45,0x30,0x64,0xd1,0xad,0xe8,0xa6,0xce,0x68,0x14,0x5c,0x25,0x67}, |
||||
{0xa3,0xda,0x8c,0xf2,0xcb,0xe,0xe1,0x16,0x33,0xe9,0x6,0x58,0x9a,0x94,0x99,0x9a,0x1f,0x60,0xb2,0x20,0xc2,0x6f,0x84,0x7b,0xd1,0xce,0xac,0x7f,0xa0,0xd1,0x85,0x18}, |
||||
{0x32,0x59,0x5b,0xa1,0x8d,0xdd,0x19,0xd3,0x50,0x9a,0x1c,0xc0,0xaa,0xa5,0xb4,0x46,0x9f,0x3d,0x63,0x67,0xe4,0x4,0x6b,0xba,0xf6,0xca,0x19,0xab,0xb,0x56,0xee,0x7e}, |
||||
{0x1f,0xb1,0x79,0xea,0xa9,0x28,0x21,0x74,0xe9,0xbd,0xf7,0x35,0x3b,0x36,0x51,0xee,0x1d,0x57,0xac,0x5a,0x75,0x50,0xd3,0x76,0x3a,0x46,0xc2,0xfe,0xa3,0x7d,0x70,0x1}, |
||||
{0xf7,0x35,0xc1,0xaf,0x98,0xa4,0xd8,0x42,0x78,0xed,0xec,0x20,0x9e,0x6b,0x67,0x79,0x41,0x83,0x63,0x15,0xea,0x3a,0xdb,0xa8,0xfa,0xc3,0x3b,0x4d,0x32,0x83,0x2c,0x83}, |
||||
{0xa7,0x40,0x3b,0x1f,0x1c,0x27,0x47,0xf3,0x59,0x40,0xf0,0x34,0xb7,0x2d,0x76,0x9a,0xe7,0x3e,0x4e,0x6c,0xd2,0x21,0x4f,0xfd,0xb8,0xfd,0x8d,0x39,0xdc,0x57,0x59,0xef}, |
||||
{0x8d,0x9b,0xc,0x49,0x2b,0x49,0xeb,0xda,0x5b,0xa2,0xd7,0x49,0x68,0xf3,0x70,0xd,0x7d,0x3b,0xae,0xd0,0x7a,0x8d,0x55,0x84,0xf5,0xa5,0xe9,0xf0,0xe4,0xf8,0x8e,0x65}, |
||||
{0xa0,0xb8,0xa2,0xf4,0x36,0x10,0x3b,0x53,0xc,0xa8,0x7,0x9e,0x75,0x3e,0xec,0x5a,0x91,0x68,0x94,0x92,0x56,0xe8,0x88,0x4f,0x5b,0xb0,0x5c,0x55,0xf8,0xba,0xbc,0x4c}, |
||||
{0xe3,0xbb,0x3b,0x99,0xf3,0x87,0x94,0x7b,0x75,0xda,0xf4,0xd6,0x72,0x6b,0x1c,0x5d,0x64,0xae,0xac,0x28,0xdc,0x34,0xb3,0x6d,0x6c,0x34,0xa5,0x50,0xb8,0x28,0xdb,0x71}, |
||||
{0xf8,0x61,0xe2,0xf2,0x10,0x8d,0x51,0x2a,0xe3,0xdb,0x64,0x33,0x59,0xdd,0x75,0xfc,0x1c,0xac,0xbc,0xf1,0x43,0xce,0x3f,0xa2,0x67,0xbb,0xd1,0x3c,0x2,0xe8,0x43,0xb0}, |
||||
{0x33,0xa,0x5b,0xca,0x88,0x29,0xa1,0x75,0x7f,0x34,0x19,0x4d,0xb4,0x16,0x53,0x5c,0x92,0x3b,0x94,0xc3,0xe,0x79,0x4d,0x1e,0x79,0x74,0x75,0xd7,0xb6,0xee,0xaf,0x3f}, |
||||
{0xea,0xa8,0xd4,0xf7,0xbe,0x1a,0x39,0x21,0x5c,0xf4,0x7e,0x9,0x4c,0x23,0x27,0x51,0x26,0xa3,0x24,0x53,0xba,0x32,0x3c,0xd2,0x44,0xa3,0x17,0x4a,0x6d,0xa6,0xd5,0xad}, |
||||
{0xb5,0x1d,0x3e,0xa6,0xaf,0xf2,0xc9,0x8,0x83,0x59,0x3d,0x98,0x91,0x6b,0x3c,0x56,0x4c,0xf8,0x7c,0xa1,0x72,0x86,0x60,0x4d,0x46,0xe2,0x3e,0xcc,0x8,0x6e,0xc7,0xf6}, |
||||
{0x2f,0x98,0x33,0xb3,0xb1,0xbc,0x76,0x5e,0x2b,0xd6,0x66,0xa5,0xef,0xc4,0xe6,0x2a,0x6,0xf4,0xb6,0xe8,0xbe,0xc1,0xd4,0x36,0x74,0xee,0x82,0x15,0xbc,0xef,0x21,0x63}, |
||||
{0xfd,0xc1,0x4e,0xd,0xf4,0x53,0xc9,0x69,0xa7,0x7d,0x5a,0xc4,0x6,0x58,0x58,0x26,0x7e,0xc1,0x14,0x16,0x6,0xe0,0xfa,0x16,0x7e,0x90,0xaf,0x3d,0x28,0x63,0x9d,0x3f}, |
||||
{0xd2,0xc9,0xf2,0xe3,0x0,0x9b,0xd2,0xc,0x5f,0xaa,0xce,0x30,0xb7,0xd4,0xc,0x30,0x74,0x2a,0x51,0x16,0xf2,0xe0,0x32,0x98,0xd,0xeb,0x30,0xd8,0xe3,0xce,0xf8,0x9a}, |
||||
{0x4b,0xc5,0x9e,0x7b,0xb5,0xf1,0x79,0x92,0xff,0x51,0xe6,0x6e,0x4,0x86,0x68,0xd3,0x9b,0x23,0x4d,0x57,0xe6,0x96,0x67,0x31,0xcc,0xe6,0xa6,0xf3,0x17,0xa,0x75,0x5}, |
||||
{0xb1,0x76,0x81,0xd9,0x13,0x32,0x6c,0xce,0x3c,0x17,0x52,0x84,0xf8,0x5,0xa2,0x62,0xf4,0x2b,0xcb,0xb3,0x78,0x47,0x15,0x47,0xff,0x46,0x54,0x82,0x23,0x93,0x6a,0x48}, |
||||
{0x38,0xdf,0x58,0x7,0x4e,0x5e,0x65,0x65,0xf2,0xfc,0x7c,0x89,0xfc,0x86,0x50,0x8e,0x31,0x70,0x2e,0x44,0xd0,0xb,0xca,0x86,0xf0,0x40,0x9,0xa2,0x30,0x78,0x47,0x4e}, |
||||
{0x65,0xa0,0xee,0x39,0xd1,0xf7,0x38,0x83,0xf7,0x5e,0xe9,0x37,0xe4,0x2c,0x3a,0xbd,0x21,0x97,0xb2,0x26,0x1,0x13,0xf8,0x6f,0xa3,0x44,0xed,0xd1,0xef,0x9f,0xde,0xe7}, |
||||
{0x8b,0xa0,0xdf,0x15,0x76,0x25,0x92,0xd9,0x3c,0x85,0xf7,0xf6,0x12,0xdc,0x42,0xbe,0xd8,0xa7,0xec,0x7c,0xab,0x27,0xb0,0x7e,0x53,0x8d,0x7d,0xda,0xaa,0x3e,0xa8,0xde}, |
||||
{0xaa,0x25,0xce,0x93,0xbd,0x2,0x69,0xd8,0x5a,0xf6,0x43,0xfd,0x1a,0x73,0x8,0xf9,0xc0,0x5f,0xef,0xda,0x17,0x4a,0x19,0xa5,0x97,0x4d,0x66,0x33,0x4c,0xfd,0x21,0x6a}, |
||||
{0x35,0xb4,0x98,0x31,0xdb,0x41,0x15,0x70,0xea,0x1e,0xf,0xbb,0xed,0xcd,0x54,0x9b,0x9a,0xd0,0x63,0xa1,0x51,0x97,0x40,0x72,0xf6,0x75,0x9d,0xbf,0x91,0x47,0x6f,0xe2}}; |
||||
|
||||
/*swapping bit 2i with bit 2i+1 of 32-bit x*/ |
||||
#define SWAP1(x) (x) = ((((x) & 0x55555555UL) << 1) | (((x) & 0xaaaaaaaaUL) >> 1)); |
||||
/*swapping bits 4i||4i+1 with bits 4i+2||4i+3 of 32-bit x*/ |
||||
#define SWAP2(x) (x) = ((((x) & 0x33333333UL) << 2) | (((x) & 0xccccccccUL) >> 2)); |
||||
/*swapping bits 8i||8i+1||8i+2||8i+3 with bits 8i+4||8i+5||8i+6||8i+7 of 32-bit x*/ |
||||
#define SWAP4(x) (x) = ((((x) & 0x0f0f0f0fUL) << 4) | (((x) & 0xf0f0f0f0UL) >> 4)); |
||||
/*swapping bits 16i||16i+1||......||16i+7 with bits 16i+8||16i+9||......||16i+15 of 32-bit x*/ |
||||
#define SWAP8(x) (x) = ((((x) & 0x00ff00ffUL) << 8) | (((x) & 0xff00ff00UL) >> 8)); |
||||
/*swapping bits 32i||32i+1||......||32i+15 with bits 32i+16||32i+17||......||32i+31 of 32-bit x*/ |
||||
#define SWAP16(x) (x) = ((((x) & 0x0000ffffUL) << 16) | (((x) & 0xffff0000UL) >> 16)); |
||||
|
||||
/*The MDS transform*/ |
||||
#define L(m0,m1,m2,m3,m4,m5,m6,m7) \ |
||||
(m4) ^= (m1); \ |
||||
(m5) ^= (m2); \ |
||||
(m6) ^= (m0) ^ (m3); \ |
||||
(m7) ^= (m0); \ |
||||
(m0) ^= (m5); \ |
||||
(m1) ^= (m6); \ |
||||
(m2) ^= (m4) ^ (m7); \ |
||||
(m3) ^= (m4); |
||||
|
||||
/*The Sbox*/ |
||||
#define Sbox(m0,m1,m2,m3,cc) \ |
||||
m3 = ~(m3); \ |
||||
m0 ^= ((~(m2)) & (cc)); \ |
||||
temp0 = (cc) ^ ((m0) & (m1));\ |
||||
m0 ^= ((m2) & (m3)); \ |
||||
m3 ^= ((~(m1)) & (m2)); \ |
||||
m1 ^= ((m0) & (m2)); \ |
||||
m2 ^= ((m0) & (~(m3))); \ |
||||
m0 ^= ((m1) | (m3)); \ |
||||
m3 ^= ((m1) & (m2)); \ |
||||
m1 ^= (temp0 & (m0)); \ |
||||
m2 ^= temp0; |
||||
|
||||
__device__ __forceinline__ void Sbox_and_MDS_layer(hashState* state, uint32_t roundnumber) |
||||
{ |
||||
uint32_t temp0; |
||||
uint32_t cc0, cc1; |
||||
//Sbox and MDS layer |
||||
#pragma unroll 4 |
||||
for (int i = 0; i < 4; i++) { |
||||
cc0 = ((uint32_t*)c_E8_bitslice_roundconstant[roundnumber])[i]; |
||||
cc1 = ((uint32_t*)c_E8_bitslice_roundconstant[roundnumber])[i+4]; |
||||
Sbox(state->x[0][i],state->x[2][i], state->x[4][i], state->x[6][i], cc0); |
||||
Sbox(state->x[1][i],state->x[3][i], state->x[5][i], state->x[7][i], cc1); |
||||
L(state->x[0][i],state->x[2][i],state->x[4][i],state->x[6][i],state->x[1][i],state->x[3][i],state->x[5][i],state->x[7][i]); |
||||
} |
||||
} |
||||
|
||||
__device__ __forceinline__ void RoundFunction0(hashState* state, uint32_t roundnumber) |
||||
{ |
||||
Sbox_and_MDS_layer(state, roundnumber); |
||||
|
||||
#pragma unroll 4 |
||||
for (int j = 1; j < 8; j = j+2) |
||||
{ |
||||
#pragma unroll 4 |
||||
for (int i = 0; i < 4; i++) SWAP1(state->x[j][i]); |
||||
} |
||||
} |
||||
|
||||
__device__ __forceinline__ void RoundFunction1(hashState* state, uint32_t roundnumber) |
||||
{ |
||||
Sbox_and_MDS_layer(state, roundnumber); |
||||
|
||||
#pragma unroll 4 |
||||
for (int j = 1; j < 8; j = j+2) |
||||
{ |
||||
#pragma unroll 4 |
||||
for (int i = 0; i < 4; i++) SWAP2(state->x[j][i]); |
||||
} |
||||
} |
||||
|
||||
__device__ __forceinline__ void RoundFunction2(hashState* state, uint32_t roundnumber) |
||||
{ |
||||
Sbox_and_MDS_layer(state, roundnumber); |
||||
|
||||
#pragma unroll 4 |
||||
for (int j = 1; j < 8; j = j+2) |
||||
{ |
||||
#pragma unroll 4 |
||||
for (int i = 0; i < 4; i++) SWAP4(state->x[j][i]); |
||||
} |
||||
} |
||||
|
||||
__device__ __forceinline__ void RoundFunction3(hashState* state, uint32_t roundnumber) |
||||
{ |
||||
Sbox_and_MDS_layer(state, roundnumber); |
||||
|
||||
#pragma unroll 4 |
||||
for (int j = 1; j < 8; j = j+2) |
||||
{ |
||||
#pragma unroll 4 |
||||
for (int i = 0; i < 4; i++) SWAP8(state->x[j][i]); |
||||
} |
||||
} |
||||
|
||||
__device__ __forceinline__ void RoundFunction4(hashState* state, uint32_t roundnumber) |
||||
{ |
||||
Sbox_and_MDS_layer(state, roundnumber); |
||||
|
||||
#pragma unroll 4 |
||||
for (int j = 1; j < 8; j = j+2) |
||||
{ |
||||
#pragma unroll 4 |
||||
for (int i = 0; i < 4; i++) SWAP16(state->x[j][i]); |
||||
} |
||||
} |
||||
|
||||
__device__ __forceinline__ void RoundFunction5(hashState* state, uint32_t roundnumber) |
||||
{ |
||||
uint32_t temp0; |
||||
|
||||
Sbox_and_MDS_layer(state, roundnumber); |
||||
|
||||
#pragma unroll 4 |
||||
for (int j = 1; j < 8; j = j+2) |
||||
{ |
||||
#pragma unroll 2 |
||||
for (int i = 0; i < 4; i = i+2) { |
||||
temp0 = state->x[j][i]; state->x[j][i] = state->x[j][i+1]; state->x[j][i+1] = temp0; |
||||
} |
||||
} |
||||
} |
||||
|
||||
__device__ __forceinline__ void RoundFunction6(hashState* state, uint32_t roundnumber) |
||||
{ |
||||
uint32_t temp0; |
||||
|
||||
Sbox_and_MDS_layer(state, roundnumber); |
||||
|
||||
#pragma unroll 4 |
||||
for (int j = 1; j < 8; j = j+2) |
||||
{ |
||||
#pragma unroll 2 |
||||
for (int i = 0; i < 2; i++) { |
||||
temp0 = state->x[j][i]; state->x[j][i] = state->x[j][i+2]; state->x[j][i+2] = temp0; |
||||
} |
||||
} |
||||
} |
||||
|
||||
/*The bijective function E8, in bitslice form */ |
||||
__device__ __forceinline__ void E8(hashState *state) |
||||
{ |
||||
/*perform 6 rounds*/ |
||||
//#pragma unroll 6 |
||||
for (int i = 0; i < 42; i+=7) |
||||
{ |
||||
RoundFunction0(state, i); |
||||
RoundFunction1(state, i+1); |
||||
RoundFunction2(state, i+2); |
||||
RoundFunction3(state, i+3); |
||||
RoundFunction4(state, i+4); |
||||
RoundFunction5(state, i+5); |
||||
RoundFunction6(state, i+6); |
||||
} |
||||
} |
||||
|
||||
/*The compression function F8 */ |
||||
__device__ __forceinline__ void F8(hashState *state) |
||||
{ |
||||
/*xor the 512-bit message with the fist half of the 1024-bit hash state*/ |
||||
#pragma unroll 16 |
||||
for (int i = 0; i < 16; i++) state->x[i >> 2][i & 3] ^= ((uint32_t*)state->buffer)[i]; |
||||
|
||||
/*the bijective function E8 */ |
||||
E8(state); |
||||
|
||||
/*xor the 512-bit message with the second half of the 1024-bit hash state*/ |
||||
#pragma unroll 16 |
||||
for (int i = 0; i < 16; i++) state->x[(16+i) >> 2][(16+i) & 3] ^= ((uint32_t*)state->buffer)[i]; |
||||
} |
||||
|
||||
|
||||
__device__ __forceinline__ void JHHash(const uint32_t *data, uint32_t *hashval) |
||||
{ |
||||
hashState state; |
||||
|
||||
/*load the intital hash value H0 into state*/ |
||||
/* |
||||
#define INIT(a,b,c,d) ((a) | ((b)<<8) | ((c)<<16) | ((d)<<24)) |
||||
state.x[0][0] = INIT(0x6f,0xd1,0x4b,0x96); |
||||
state.x[0][1] = INIT(0x3e,0x00,0xaa,0x17); |
||||
state.x[0][2] = INIT(0x63,0x6a,0x2e,0x05); |
||||
state.x[0][3] = INIT(0x7a,0x15,0xd5,0x43); |
||||
state.x[1][0] = INIT(0x8a,0x22,0x5e,0x8d); |
||||
state.x[1][1] = INIT(0x0c,0x97,0xef,0x0b); |
||||
state.x[1][2] = INIT(0xe9,0x34,0x12,0x59); |
||||
state.x[1][3] = INIT(0xf2,0xb3,0xc3,0x61); |
||||
state.x[2][0] = INIT(0x89,0x1d,0xa0,0xc1); |
||||
state.x[2][1] = INIT(0x53,0x6f,0x80,0x1e); |
||||
state.x[2][2] = INIT(0x2a,0xa9,0x05,0x6b); |
||||
state.x[2][3] = INIT(0xea,0x2b,0x6d,0x80); |
||||
state.x[3][0] = INIT(0x58,0x8e,0xcc,0xdb); |
||||
state.x[3][1] = INIT(0x20,0x75,0xba,0xa6); |
||||
state.x[3][2] = INIT(0xa9,0x0f,0x3a,0x76); |
||||
state.x[3][3] = INIT(0xba,0xf8,0x3b,0xf7); |
||||
state.x[4][0] = INIT(0x01,0x69,0xe6,0x05); |
||||
state.x[4][1] = INIT(0x41,0xe3,0x4a,0x69); |
||||
state.x[4][2] = INIT(0x46,0xb5,0x8a,0x8e); |
||||
state.x[4][3] = INIT(0x2e,0x6f,0xe6,0x5a); |
||||
state.x[5][0] = INIT(0x10,0x47,0xa7,0xd0); |
||||
state.x[5][1] = INIT(0xc1,0x84,0x3c,0x24); |
||||
state.x[5][2] = INIT(0x3b,0x6e,0x71,0xb1); |
||||
state.x[5][3] = INIT(0x2d,0x5a,0xc1,0x99); |
||||
state.x[6][0] = INIT(0xcf,0x57,0xf6,0xec); |
||||
state.x[6][1] = INIT(0x9d,0xb1,0xf8,0x56); |
||||
state.x[6][2] = INIT(0xa7,0x06,0x88,0x7c); |
||||
state.x[6][3] = INIT(0x57,0x16,0xb1,0x56); |
||||
state.x[7][0] = INIT(0xe3,0xc2,0xfc,0xdf); |
||||
state.x[7][1] = INIT(0xe6,0x85,0x17,0xfb); |
||||
state.x[7][2] = INIT(0x54,0x5a,0x46,0x78); |
||||
state.x[7][3] = INIT(0xcc,0x8c,0xdd,0x4b); |
||||
*/ |
||||
#pragma unroll 8 |
||||
for(int j=0;j<8;j++) |
||||
{ |
||||
#pragma unroll 4 |
||||
for(int i=0;i<4;i++) |
||||
state.x[j][i] = c_INIT_bitslice[j][i]; |
||||
} |
||||
|
||||
#pragma unroll 16 |
||||
for (int i=0; i < 16; ++i) state.buffer[i] = data[i]; |
||||
F8(&state); |
||||
|
||||
/*pad the message when databitlen is multiple of 512 bits, then process the padded block*/ |
||||
state.buffer[0] = 0x80; |
||||
#pragma unroll 14 |
||||
for (int i=1; i < 15; i++) state.buffer[i] = 0; |
||||
state.buffer[15] = 0x00020000; |
||||
F8(&state); |
||||
|
||||
/*truncating the final hash value to generate the message digest*/ |
||||
#pragma unroll 16 |
||||
for (int i=0; i < 16; ++i) hashval[i] = state.x[4][i]; |
||||
} |
||||
|
||||
// Die Hash-Funktion |
||||
__global__ void quark_jh512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector) |
||||
{ |
||||
int thread = (blockDim.x * blockIdx.x + threadIdx.x); |
||||
if (thread < threads) |
||||
{ |
||||
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread); |
||||
|
||||
int hashPosition = nounce - startNounce; |
||||
uint32_t *Hash = (uint32_t*)&g_hash[8 * hashPosition]; |
||||
|
||||
JHHash(Hash, Hash); |
||||
} |
||||
} |
||||
|
||||
|
||||
// Setup-Funktionen |
||||
__host__ void quark_jh512_cpu_init(int thr_id, int threads) |
||||
{ |
||||
|
||||
cudaMemcpyToSymbol( c_E8_bitslice_roundconstant, |
||||
h_E8_bitslice_roundconstant, |
||||
sizeof(h_E8_bitslice_roundconstant), |
||||
0, cudaMemcpyHostToDevice); |
||||
|
||||
cudaMemcpyToSymbol( c_INIT_bitslice, |
||||
h_INIT_bitslice, |
||||
sizeof(h_INIT_bitslice), |
||||
0, cudaMemcpyHostToDevice); |
||||
} |
||||
|
||||
__host__ void quark_jh512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) |
||||
{ |
||||
const int threadsperblock = 256; |
||||
|
||||
// berechne wie viele Thread Blocks wir brauchen |
||||
dim3 grid((threads + threadsperblock-1)/threadsperblock); |
||||
dim3 block(threadsperblock); |
||||
|
||||
// Größe des dynamischen Shared Memory Bereichs |
||||
size_t shared_size = 0; |
||||
|
||||
// fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size); |
||||
|
||||
quark_jh512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector); |
||||
MyStreamSynchronize(NULL, order, thr_id); |
||||
} |
||||
|
||||
@ -0,0 +1,482 @@
@@ -0,0 +1,482 @@
|
||||
#include <cuda.h> |
||||
#include "cuda_runtime.h" |
||||
#include "device_launch_parameters.h" |
||||
|
||||
#include <stdio.h> |
||||
#include <memory.h> |
||||
|
||||
#define USE_SHUFFLE 0 |
||||
|
||||
// Folgende Definitionen später durch header ersetzen |
||||
typedef unsigned char uint8_t; |
||||
typedef unsigned int uint32_t; |
||||
typedef unsigned long long uint64_t; |
||||
|
||||
// aus heavy.cu |
||||
extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); |
||||
|
||||
// die Message it Padding zur Berechnung auf der GPU |
||||
__constant__ uint64_t c_PaddedMessage80[16]; // padded message (80 bytes + padding) |
||||
|
||||
// ---------------------------- BEGIN CUDA quark_blake512 functions ------------------------------------ |
||||
|
||||
__constant__ uint8_t c_sigma[16][16]; |
||||
|
||||
const uint8_t host_sigma[16][16] = |
||||
{ |
||||
{ 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 } |
||||
}; |
||||
|
||||
// das Hi Word aus einem 64 Bit Typen extrahieren |
||||
static __device__ uint32_t HIWORD(const uint64_t &x) { |
||||
#if __CUDA_ARCH__ >= 130 |
||||
return (uint32_t)__double2hiint(__longlong_as_double(x)); |
||||
#else |
||||
return (uint32_t)(x >> 32); |
||||
#endif |
||||
} |
||||
|
||||
// das Hi Word in einem 64 Bit Typen ersetzen |
||||
static __device__ uint64_t REPLACE_HIWORD(const uint64_t &x, const uint32_t &y) { |
||||
return (x & 0xFFFFFFFFULL) | (((uint64_t)y) << 32ULL); |
||||
} |
||||
|
||||
// das Lo Word aus einem 64 Bit Typen extrahieren |
||||
static __device__ uint32_t LOWORD(const uint64_t &x) { |
||||
#if __CUDA_ARCH__ >= 130 |
||||
return (uint32_t)__double2loint(__longlong_as_double(x)); |
||||
#else |
||||
return (uint32_t)(x & 0xFFFFFFFFULL); |
||||
#endif |
||||
} |
||||
|
||||
// das Lo Word in einem 64 Bit Typen ersetzen |
||||
static __device__ uint64_t REPLACE_LOWORD(const uint64_t &x, const uint32_t &y) { |
||||
return (x & 0xFFFFFFFF00000000ULL) | ((uint64_t)y); |
||||
} |
||||
|
||||
/* |
||||
#define SWAP32(x) \ |
||||
((((x) << 24) & 0xff000000u) | (((x) << 8) & 0x00ff0000u) | \ |
||||
(((x) >> 8) & 0x0000ff00u) | (((x) >> 24) & 0x000000ffu)) |
||||
|
||||
#define SWAP64(x) \ |
||||
((uint64_t)((((uint64_t)(x) & 0xff00000000000000ULL) >> 56) | \ |
||||
(((uint64_t)(x) & 0x00ff000000000000ULL) >> 40) | \ |
||||
(((uint64_t)(x) & 0x0000ff0000000000ULL) >> 24) | \ |
||||
(((uint64_t)(x) & 0x000000ff00000000ULL) >> 8) | \ |
||||
(((uint64_t)(x) & 0x00000000ff000000ULL) << 8) | \ |
||||
(((uint64_t)(x) & 0x0000000000ff0000ULL) << 24) | \ |
||||
(((uint64_t)(x) & 0x000000000000ff00ULL) << 40) | \ |
||||
(((uint64_t)(x) & 0x00000000000000ffULL) << 56))) |
||||
*/ |
||||
|
||||
/* |
||||
__device__ __forceinline__ void SWAP32(uint32_t *x) |
||||
{ |
||||
// Input: 33221100 |
||||
// Output: 00112233 |
||||
x[0] = __byte_perm(x[0], 0, 0x0123); |
||||
} |
||||
*/ |
||||
__device__ __forceinline__ uint64_t SWAP64(uint64_t x) |
||||
{ |
||||
// Input: 77665544 33221100 |
||||
// Output: 00112233 44556677 |
||||
uint64_t temp[2]; |
||||
temp[0] = __byte_perm(HIWORD(x), 0, 0x0123); |
||||
temp[1] = __byte_perm(LOWORD(x), 0, 0x0123); |
||||
|
||||
return temp[0] | (temp[1]<<32); |
||||
} |
||||
|
||||
__constant__ uint64_t c_u512[16]; |
||||
|
||||
const uint64_t host_u512[16] = |
||||
{ |
||||
0x243f6a8885a308d3ULL, 0x13198a2e03707344ULL, |
||||
0xa4093822299f31d0ULL, 0x082efa98ec4e6c89ULL, |
||||
0x452821e638d01377ULL, 0xbe5466cf34e90c6cULL, |
||||
0xc0ac29b7c97c50ddULL, 0x3f84d5b5b5470917ULL, |
||||
0x9216d5d98979fb1bULL, 0xd1310ba698dfb5acULL, |
||||
0x2ffd72dbd01adfb7ULL, 0xb8e1afed6a267e96ULL, |
||||
0xba7c9045f12c7f99ULL, 0x24a19947b3916cf7ULL, |
||||
0x0801f2e2858efc16ULL, 0x636920d871574e69ULL |
||||
}; |
||||
|
||||
|
||||
// diese 64 Bit Rotates werden unter Compute 3.5 (und besser) mit dem Funnel Shifter beschleunigt |
||||
#if __CUDA_ARCH__ >= 350 |
||||
__forceinline__ __device__ uint64_t ROTR(const uint64_t value, const int offset) { |
||||
uint2 result; |
||||
if(offset < 32) { |
||||
asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset)); |
||||
asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset)); |
||||
} else { |
||||
asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset)); |
||||
asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset)); |
||||
} |
||||
return __double_as_longlong(__hiloint2double(result.y, result.x)); |
||||
} |
||||
#else |
||||
#define ROTR(x, n) (((x) >> (n)) | ((x) << (64 - (n)))) |
||||
#endif |
||||
|
||||
#define G(a,b,c,d,e) \ |
||||
v[a] += (m[sigma[i][e]] ^ u512[sigma[i][e+1]]) + v[b];\ |
||||
v[d] = ROTR( v[d] ^ v[a],32); \ |
||||
v[c] += v[d]; \ |
||||
v[b] = ROTR( v[b] ^ v[c],25); \ |
||||
v[a] += (m[sigma[i][e+1]] ^ u512[sigma[i][e]])+v[b]; \ |
||||
v[d] = ROTR( v[d] ^ v[a],16); \ |
||||
v[c] += v[d]; \ |
||||
v[b] = ROTR( v[b] ^ v[c],11); |
||||
|
||||
|
||||
__device__ void quark_blake512_compress( uint64_t *h, const uint64_t *block, const uint8_t ((*sigma)[16]), const uint64_t *u512, const int bits ) |
||||
{ |
||||
uint64_t v[16], m[16], i; |
||||
|
||||
#pragma unroll 16 |
||||
for( i = 0; i < 16; ++i ) |
||||
{ |
||||
m[i] = SWAP64(block[i]); |
||||
} |
||||
|
||||
#pragma unroll 8 |
||||
for( i = 0; i < 8; ++i ) v[i] = h[i]; |
||||
|
||||
v[ 8] = u512[0]; |
||||
v[ 9] = u512[1]; |
||||
v[10] = u512[2]; |
||||
v[11] = u512[3]; |
||||
v[12] = u512[4]; |
||||
v[13] = u512[5]; |
||||
v[14] = u512[6]; |
||||
v[15] = u512[7]; |
||||
|
||||
v[12] ^= bits; |
||||
v[13] ^= bits; |
||||
|
||||
//#pragma unroll 16 |
||||
for( i = 0; i < 16; ++i ) |
||||
{ |
||||
/* column step */ |
||||
G( 0, 4, 8, 12, 0 ); |
||||
G( 1, 5, 9, 13, 2 ); |
||||
G( 2, 6, 10, 14, 4 ); |
||||
G( 3, 7, 11, 15, 6 ); |
||||
/* diagonal step */ |
||||
G( 0, 5, 10, 15, 8 ); |
||||
G( 1, 6, 11, 12, 10 ); |
||||
G( 2, 7, 8, 13, 12 ); |
||||
G( 3, 4, 9, 14, 14 ); |
||||
} |
||||
|
||||
#pragma unroll 16 |
||||
for( i = 0; i < 16; ++i ) h[i % 8] ^= v[i]; |
||||
} |
||||
|
||||
// Endian Drehung für 32 Bit Typen |
||||
|
||||
static __device__ uint32_t cuda_swab32(uint32_t x) |
||||
{ |
||||
return __byte_perm(x, 0, 0x0123); |
||||
/* |
||||
return (((x << 24) & 0xff000000u) | ((x << 8) & 0x00ff0000u) |
||||
| ((x >> 8) & 0x0000ff00u) | ((x >> 24) & 0x000000ffu)); |
||||
*/ |
||||
} |
||||
/* |
||||
// Endian Drehung für 64 Bit Typen |
||||
static __device__ uint64_t cuda_swab64(uint64_t x) { |
||||
uint32_t h = (x >> 32); |
||||
uint32_t l = (x & 0xFFFFFFFFULL); |
||||
return (((uint64_t)cuda_swab32(l)) << 32) | ((uint64_t)cuda_swab32(h)); |
||||
} |
||||
*/ |
||||
|
||||
static __constant__ uint64_t d_constMem[8]; |
||||
static const uint64_t h_constMem[8] = { |
||||
0x6a09e667f3bcc908ULL, |
||||
0xbb67ae8584caa73bULL, |
||||
0x3c6ef372fe94f82bULL, |
||||
0xa54ff53a5f1d36f1ULL, |
||||
0x510e527fade682d1ULL, |
||||
0x9b05688c2b3e6c1fULL, |
||||
0x1f83d9abfb41bd6bULL, |
||||
0x5be0cd19137e2179ULL }; |
||||
|
||||
// Hash-Padding |
||||
static __constant__ uint64_t d_constHashPadding[8]; |
||||
static const uint64_t h_constHashPadding[8] = { |
||||
0x0000000000000080ull, |
||||
0, |
||||
0, |
||||
0, |
||||
0, |
||||
0x0100000000000000ull, |
||||
0, |
||||
0x0002000000000000ull }; |
||||
|
||||
__global__ void quark_blake512_gpu_hash_64(int threads, uint32_t startNounce, uint32_t *g_nonceVector, uint64_t *g_hash) |
||||
{ |
||||
int thread = (blockDim.x * blockIdx.x + threadIdx.x); |
||||
|
||||
#if USE_SHUFFLE |
||||
const int warpID = threadIdx.x & 0x0F; // 16 warps |
||||
const int warpBlockID = (thread + 15)>>4; // aufrunden auf volle Warp-Blöcke |
||||
const int maxHashPosition = thread<<3; |
||||
#endif |
||||
|
||||
#if USE_SHUFFLE |
||||
if (warpBlockID < ( (threads+15)>>4 )) |
||||
#else |
||||
if (thread < threads) |
||||
#endif |
||||
{ |
||||
// bestimme den aktuellen Zähler |
||||
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread); |
||||
|
||||
int hashPosition = nounce - startNounce; |
||||
//uint64_t *inpHash = &g_hash[8 * hashPosition]; |
||||
uint64_t *inpHash = &g_hash[hashPosition<<3]; |
||||
|
||||
// State vorbereiten |
||||
uint64_t h[8]; |
||||
/* |
||||
h[0] = 0x6a09e667f3bcc908ULL; |
||||
h[1] = 0xbb67ae8584caa73bULL; |
||||
h[2] = 0x3c6ef372fe94f82bULL; |
||||
h[3] = 0xa54ff53a5f1d36f1ULL; |
||||
h[4] = 0x510e527fade682d1ULL; |
||||
h[5] = 0x9b05688c2b3e6c1fULL; |
||||
h[6] = 0x1f83d9abfb41bd6bULL; |
||||
h[7] = 0x5be0cd19137e2179ULL; |
||||
*/ |
||||
#pragma unroll 8 |
||||
for(int i=0;i<8;i++) |
||||
h[i] = d_constMem[i]; |
||||
|
||||
// 128 Byte für die Message |
||||
uint64_t buf[16]; |
||||
|
||||
// Message für die erste Runde in Register holen |
||||
#pragma unroll 8 |
||||
for (int i=0; i < 8; ++i) buf[i] = inpHash[i]; |
||||
|
||||
/* |
||||
buf[ 8] = 0x0000000000000080ull; |
||||
buf[ 9] = 0; |
||||
buf[10] = 0; |
||||
buf[11] = 0; |
||||
buf[12] = 0; |
||||
buf[13] = 0x0100000000000000ull; |
||||
buf[14] = 0; |
||||
buf[15] = 0x0002000000000000ull; |
||||
*/ |
||||
#pragma unroll 8 |
||||
for(int i=0;i<8;i++) |
||||
buf[i+8] = d_constHashPadding[i]; |
||||
|
||||
// die einzige Hashing-Runde |
||||
quark_blake512_compress( h, buf, c_sigma, c_u512, 512 ); |
||||
|
||||
// Hash rauslassen |
||||
#if __CUDA_ARCH__ >= 130 |
||||
// ausschliesslich 32 bit Operationen sofern die SM1.3 double intrinsics verfügbar sind |
||||
uint32_t *outHash = (uint32_t*)&g_hash[8 * hashPosition]; |
||||
#pragma unroll 8 |
||||
for (int i=0; i < 8; ++i) { |
||||
outHash[2*i+0] = cuda_swab32( HIWORD(h[i]) ); |
||||
outHash[2*i+1] = cuda_swab32( LOWORD(h[i]) ); |
||||
} |
||||
#else |
||||
// in dieser Version passieren auch ein paar 64 Bit Shifts |
||||
uint64_t *outHash = &g_hash[8 * hashPosition]; |
||||
#pragma unroll 8 |
||||
for (int i=0; i < 8; ++i) |
||||
{ |
||||
//outHash[i] = cuda_swab64( h[i] ); |
||||
outHash[i] = SWAP64(h[i]); |
||||
} |
||||
#endif |
||||
} |
||||
} |
||||
|
||||
__global__ void quark_blake512_gpu_hash_80(int threads, uint32_t startNounce, void *outputHash) |
||||
{ |
||||
int thread = (blockDim.x * blockIdx.x + threadIdx.x); |
||||
if (thread < threads) |
||||
{ |
||||
// bestimme den aktuellen Zähler |
||||
uint32_t nounce = startNounce + thread; |
||||
|
||||
// State vorbereiten |
||||
uint64_t h[8]; |
||||
/* |
||||
h[0] = 0x6a09e667f3bcc908ULL; |
||||
h[1] = 0xbb67ae8584caa73bULL; |
||||
h[2] = 0x3c6ef372fe94f82bULL; |
||||
h[3] = 0xa54ff53a5f1d36f1ULL; |
||||
h[4] = 0x510e527fade682d1ULL; |
||||
h[5] = 0x9b05688c2b3e6c1fULL; |
||||
h[6] = 0x1f83d9abfb41bd6bULL; |
||||
h[7] = 0x5be0cd19137e2179ULL; |
||||
*/ |
||||
#pragma unroll 8 |
||||
for(int i=0;i<8;i++) |
||||
h[i] = d_constMem[i]; |
||||
// 128 Byte für die Message |
||||
uint64_t buf[16]; |
||||
|
||||
// Message für die erste Runde in Register holen |
||||
#pragma unroll 16 |
||||
for (int i=0; i < 16; ++i) buf[i] = c_PaddedMessage80[i]; |
||||
|
||||
// die Nounce durch die thread-spezifische ersetzen |
||||
buf[9] = REPLACE_HIWORD(buf[9], cuda_swab32(nounce)); |
||||
|
||||
// die einzige Hashing-Runde |
||||
quark_blake512_compress( h, buf, c_sigma, c_u512, 640 ); |
||||
|
||||
// Hash rauslassen |
||||
#if __CUDA_ARCH__ >= 130 |
||||
// ausschliesslich 32 bit Operationen sofern die SM1.3 double intrinsics verfügbar sind |
||||
uint32_t *outHash = (uint32_t *)outputHash + 16 * thread; |
||||
#pragma unroll 8 |
||||
for (int i=0; i < 8; ++i) { |
||||
outHash[2*i+0] = cuda_swab32( HIWORD(h[i]) ); |
||||
outHash[2*i+1] = cuda_swab32( LOWORD(h[i]) ); |
||||
} |
||||
#else |
||||
// in dieser Version passieren auch ein paar 64 Bit Shifts |
||||
uint64_t *outHash = (uint64_t *)outputHash + 8 * thread; |
||||
#pragma unroll 8 |
||||
for (int i=0; i < 8; ++i) |
||||
{ |
||||
//outHash[i] = cuda_swab64( h[i] ); |
||||
outHash[i] = SWAP64(h[i]); |
||||
} |
||||
#endif |
||||
} |
||||
} |
||||
|
||||
|
||||
// ---------------------------- END CUDA quark_blake512 functions ------------------------------------ |
||||
|
||||
// Setup-Funktionen |
||||
__host__ void quark_blake512_cpu_init(int thr_id, int threads) |
||||
{ |
||||
// Kopiere die Hash-Tabellen in den GPU-Speicher |
||||
cudaMemcpyToSymbol( c_sigma, |
||||
host_sigma, |
||||
sizeof(host_sigma), |
||||
0, cudaMemcpyHostToDevice); |
||||
|
||||
cudaMemcpyToSymbol( c_u512, |
||||
host_u512, |
||||
sizeof(host_u512), |
||||
0, cudaMemcpyHostToDevice); |
||||
|
||||
cudaMemcpyToSymbol( d_constMem, |
||||
h_constMem, |
||||
sizeof(h_constMem), |
||||
0, cudaMemcpyHostToDevice); |
||||
|
||||
cudaMemcpyToSymbol( d_constHashPadding, |
||||
h_constHashPadding, |
||||
sizeof(h_constHashPadding), |
||||
0, cudaMemcpyHostToDevice); |
||||
} |
||||
|
||||
// Blake512 für 80 Byte grosse Eingangsdaten |
||||
__host__ void quark_blake512_cpu_setBlock_80(void *pdata) |
||||
{ |
||||
// Message mit Padding bereitstellen |
||||
// lediglich die korrekte Nonce ist noch ab Byte 76 einzusetzen. |
||||
unsigned char PaddedMessage[128]; |
||||
memcpy(PaddedMessage, pdata, 80); |
||||
memset(PaddedMessage+80, 0, 48); |
||||
PaddedMessage[80] = 0x80; |
||||
PaddedMessage[111] = 1; |
||||
PaddedMessage[126] = 0x02; |
||||
PaddedMessage[127] = 0x80; |
||||
|
||||
// die Message zur Berechnung auf der GPU |
||||
cudaMemcpyToSymbol( c_PaddedMessage80, PaddedMessage, 16*sizeof(uint64_t), 0, cudaMemcpyHostToDevice); |
||||
} |
||||
|
||||
#if 0 |
||||
// Blake512 für 64 Byte grosse Eingangsdaten |
||||
// evtl. macht es gar keinen Sinn, das alles ins Constant Memory to schicken. Es sind hier sowieso |
||||
// nur die letzten 64 Bytes des Blocks konstant, und die meisten Bytes davon sind 0. Das kann mnan |
||||
// auch im Kernel initialisieren. |
||||
__host__ void quark_blake512_cpu_setBlock_64(void *pdata) |
||||
{ |
||||
// Message mit Padding bereitstellen |
||||
unsigned char PaddedMessage[128]; |
||||
memcpy(PaddedMessage, pdata, 64); // Hinweis: diese 64 Bytes sind nonce-spezifisch und ändern sich KOMPLETT für jede Nonce! |
||||
memset(PaddedMessage+64, 0, 64); |
||||
PaddedMessage[64] = 0x80; |
||||
PaddedMessage[111] = 1; |
||||
PaddedMessage[126] = 0x02; |
||||
PaddedMessage[127] = 0x00; |
||||
|
||||
// die Message zur Berechnung auf der GPU |
||||
cudaMemcpyToSymbol( c_PaddedMessage80, PaddedMessage, 16*sizeof(uint64_t), 0, cudaMemcpyHostToDevice); |
||||
} |
||||
#endif |
||||
|
||||
__host__ void quark_blake512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_outputHash, int order) |
||||
{ |
||||
const int threadsperblock = 256; |
||||
|
||||
// berechne wie viele Thread Blocks wir brauchen |
||||
dim3 grid((threads + threadsperblock-1)/threadsperblock); |
||||
dim3 block(threadsperblock); |
||||
|
||||
// Größe des dynamischen Shared Memory Bereichs |
||||
size_t shared_size = 0; |
||||
|
||||
// fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size); |
||||
|
||||
quark_blake512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_nonceVector, (uint64_t*)d_outputHash); |
||||
|
||||
// Strategisches Sleep Kommando zur Senkung der CPU Last |
||||
MyStreamSynchronize(NULL, order, thr_id); |
||||
} |
||||
|
||||
__host__ void quark_blake512_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_outputHash, int order) |
||||
{ |
||||
const int threadsperblock = 256; |
||||
|
||||
// berechne wie viele Thread Blocks wir brauchen |
||||
dim3 grid((threads + threadsperblock-1)/threadsperblock); |
||||
dim3 block(threadsperblock); |
||||
|
||||
// Größe des dynamischen Shared Memory Bereichs |
||||
size_t shared_size = 0; |
||||
|
||||
// fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size); |
||||
|
||||
quark_blake512_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNounce, d_outputHash); |
||||
|
||||
// Strategisches Sleep Kommando zur Senkung der CPU Last |
||||
MyStreamSynchronize(NULL, order, thr_id); |
||||
} |
||||
@ -0,0 +1,397 @@
@@ -0,0 +1,397 @@
|
||||
// Auf QuarkCoin spezialisierte Version von Groestl |
||||
|
||||
#include <cuda.h> |
||||
#include "cuda_runtime.h" |
||||
#include "device_launch_parameters.h" |
||||
|
||||
#include <stdio.h> |
||||
#include <memory.h> |
||||
|
||||
// it's unfortunate that this is a compile time constant. |
||||
#define MAXWELL_OR_FERMI 0 |
||||
|
||||
// aus cpu-miner.c |
||||
extern int device_map[8]; |
||||
|
||||
// aus heavy.cu |
||||
extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); |
||||
|
||||
// Folgende Definitionen später durch header ersetzen |
||||
typedef unsigned char uint8_t; |
||||
typedef unsigned int uint32_t; |
||||
typedef unsigned long long uint64_t; |
||||
|
||||
// diese Struktur wird in der Init Funktion angefordert |
||||
static cudaDeviceProp props[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) |
||||
|
||||
#if MAXWELL_OR_FEMRI |
||||
#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(t0up1, x) |
||||
#define T0dn(x) tex1Dfetch(t0dn1, x) |
||||
#define T1up(x) tex1Dfetch(t1up1, x) |
||||
#define T1dn(x) tex1Dfetch(t1dn1, x) |
||||
#define T2up(x) tex1Dfetch(t2up1, x) |
||||
#define T2dn(x) tex1Dfetch(t2dn1, x) |
||||
#define T3up(x) tex1Dfetch(t3up1, x) |
||||
#define T3dn(x) tex1Dfetch(t3dn1, 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(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) |
||||
#endif |
||||
|
||||
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[]; |
||||
|
||||
__device__ __forceinline__ void quark_groestl512_perm_P(uint32_t *a, char *mixtabs) |
||||
{ |
||||
uint32_t t[32]; |
||||
|
||||
//#pragma unroll 14 |
||||
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) |
||||
{ |
||||
//#pragma unroll 14 |
||||
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]; |
||||
} |
||||
} |
||||
__global__ void quark_groestl512_gpu_hash_64(int threads, uint32_t startNounce, uint32_t *g_hash, uint32_t *g_nonceVector) |
||||
{ |
||||
#if USE_SHARED |
||||
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(); |
||||
#endif |
||||
|
||||
int 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); |
||||
|
||||
int hashPosition = nounce - startNounce; |
||||
uint32_t *inpHash = &g_hash[16 * hashPosition]; |
||||
|
||||
#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 |
||||
#if USE_SHARED |
||||
quark_groestl512_perm_P(state, mixtabs); |
||||
state[31] ^= 0x20000; |
||||
quark_groestl512_perm_Q(message, mixtabs); |
||||
#else |
||||
quark_groestl512_perm_P(state, NULL); |
||||
state[31] ^= 0x20000; |
||||
quark_groestl512_perm_Q(message, NULL); |
||||
#endif |
||||
#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]; |
||||
|
||||
#if USE_SHARED |
||||
quark_groestl512_perm_P(message, mixtabs); |
||||
#else |
||||
quark_groestl512_perm_P(message, NULL); |
||||
#endif |
||||
|
||||
#pragma unroll 32 |
||||
for(int u=0;u<32;u++) state[u] ^= message[u]; |
||||
// Erzeugten Hash rausschreiben |
||||
|
||||
uint32_t *outpHash = &g_hash[16 * hashPosition]; |
||||
|
||||
#pragma unroll 16 |
||||
for(int k=0;k<16;k++) outpHash[k] = state[k+16]; |
||||
} |
||||
} |
||||
|
||||
#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 ); } \ |
||||
|
||||
// Setup-Funktionen |
||||
__host__ void quark_groestl512_cpu_init(int thr_id, int threads) |
||||
{ |
||||
cudaGetDeviceProperties(&props[thr_id], device_map[thr_id]); |
||||
|
||||
// Texturen mit obigem Makro initialisieren |
||||
texDef(t0up1, d_T0up, T0up_cpu, sizeof(uint32_t)*256); |
||||
texDef(t0dn1, d_T0dn, T0dn_cpu, sizeof(uint32_t)*256); |
||||
texDef(t1up1, d_T1up, T1up_cpu, sizeof(uint32_t)*256); |
||||
texDef(t1dn1, d_T1dn, T1dn_cpu, sizeof(uint32_t)*256); |
||||
texDef(t2up1, d_T2up, T2up_cpu, sizeof(uint32_t)*256); |
||||
texDef(t2dn1, d_T2dn, T2dn_cpu, sizeof(uint32_t)*256); |
||||
texDef(t3up1, d_T3up, T3up_cpu, sizeof(uint32_t)*256); |
||||
texDef(t3dn1, d_T3dn, T3dn_cpu, sizeof(uint32_t)*256); |
||||
} |
||||
|
||||
__host__ void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) |
||||
{ |
||||
// Compute 3.5 und 5.x Geräte am besten mit 768 Threads ansteuern, |
||||
// alle anderen mit 512 Threads. |
||||
int threadsperblock = ((props[thr_id].major == 3 && props[thr_id].minor == 5) || props[thr_id].major > 3) ? 768 : 512; |
||||
|
||||
// berechne wie viele Thread Blocks wir brauchen |
||||
dim3 grid((threads + threadsperblock-1)/threadsperblock); |
||||
dim3 block(threadsperblock); |
||||
|
||||
// Größe des dynamischen Shared Memory Bereichs |
||||
#if USE_SHARED |
||||
size_t shared_size = 8 * 256 * sizeof(uint32_t); |
||||
#else |
||||
size_t shared_size = 0; |
||||
#endif |
||||
|
||||
// fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size); |
||||
//fprintf(stderr, "ThrID: %d\n", thr_id); |
||||
quark_groestl512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_hash, d_nonceVector); |
||||
|
||||
// Strategisches Sleep Kommando zur Senkung der CPU Last |
||||
MyStreamSynchronize(NULL, order, thr_id); |
||||
} |
||||
|
||||
__host__ void quark_doublegroestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) |
||||
{ |
||||
// Compute 3.5 und 5.x Geräte am besten mit 768 Threads ansteuern, |
||||
// alle anderen mit 512 Threads. |
||||
int threadsperblock = ((props[thr_id].major == 3 && props[thr_id].minor == 5) || props[thr_id].major > 3) ? 768 : 512; |
||||
|
||||
// berechne wie viele Thread Blocks wir brauchen |
||||
dim3 grid((threads + threadsperblock-1)/threadsperblock); |
||||
dim3 block(threadsperblock); |
||||
|
||||
// Größe des dynamischen Shared Memory Bereichs |
||||
#if USE_SHARED |
||||
size_t shared_size = 8 * 256 * sizeof(uint32_t); |
||||
#else |
||||
size_t shared_size = 0; |
||||
#endif |
||||
|
||||
// fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size); |
||||
//fprintf(stderr, "ThrID: %d\n", thr_id); |
||||
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); |
||||
|
||||
// Strategisches Sleep Kommando zur Senkung der CPU Last |
||||
MyStreamSynchronize(NULL, order, thr_id); |
||||
} |
||||
@ -0,0 +1,450 @@
@@ -0,0 +1,450 @@
|
||||
#include <cuda.h> |
||||
#include "cuda_runtime.h" |
||||
#include "device_launch_parameters.h" |
||||
|
||||
#include <stdio.h> |
||||
#include <memory.h> |
||||
|
||||
// Folgende Definitionen später durch header ersetzen |
||||
typedef unsigned char uint8_t; |
||||
typedef unsigned int uint32_t; |
||||
typedef unsigned long long uint64_t; |
||||
|
||||
#define SPH_C64(x) ((uint64_t)(x ## ULL)) |
||||
|
||||
// aus cpu-miner.c |
||||
extern "C" extern int device_map[8]; |
||||
// aus heavy.cu |
||||
extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); |
||||
|
||||
// Take a look at: https://www.schneier.com/skein1.3.pdf |
||||
|
||||
#if __CUDA_ARCH__ >= 350 |
||||
__forceinline__ __device__ uint64_t ROTL64(const uint64_t value, const int offset) { |
||||
uint2 result; |
||||
if(offset >= 32) { |
||||
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset)); |
||||
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset)); |
||||
} else { |
||||
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset)); |
||||
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset)); |
||||
} |
||||
return __double_as_longlong(__hiloint2double(result.y, result.x)); |
||||
} |
||||
#else |
||||
#define ROTL64(x, n) (((x) << (n)) | ((x) >> (64 - (n)))) |
||||
#endif |
||||
#define SHL(x, n) ((x) << (n)) |
||||
#define SHR(x, n) ((x) >> (n)) |
||||
|
||||
// Zum testen Hostcode... |
||||
/* Hier erstmal die Tabelle mit den Konstanten für die Mix-Funktion. Kann später vll. |
||||
mal direkt in den Code eingesetzt werden |
||||
*/ |
||||
|
||||
/* |
||||
* M9_ ## s ## _ ## i evaluates to s+i mod 9 (0 <= s <= 18, 0 <= i <= 7). |
||||
*/ |
||||
|
||||
#define M9_0_0 0 |
||||
#define M9_0_1 1 |
||||
#define M9_0_2 2 |
||||
#define M9_0_3 3 |
||||
#define M9_0_4 4 |
||||
#define M9_0_5 5 |
||||
#define M9_0_6 6 |
||||
#define M9_0_7 7 |
||||
|
||||
#define M9_1_0 1 |
||||
#define M9_1_1 2 |
||||
#define M9_1_2 3 |
||||
#define M9_1_3 4 |
||||
#define M9_1_4 5 |
||||
#define M9_1_5 6 |
||||
#define M9_1_6 7 |
||||
#define M9_1_7 8 |
||||
|
||||
#define M9_2_0 2 |
||||
#define M9_2_1 3 |
||||
#define M9_2_2 4 |
||||
#define M9_2_3 5 |
||||
#define M9_2_4 6 |
||||
#define M9_2_5 7 |
||||
#define M9_2_6 8 |
||||
#define M9_2_7 0 |
||||
|
||||
#define M9_3_0 3 |
||||
#define M9_3_1 4 |
||||
#define M9_3_2 5 |
||||
#define M9_3_3 6 |
||||
#define M9_3_4 7 |
||||
#define M9_3_5 8 |
||||
#define M9_3_6 0 |
||||
#define M9_3_7 1 |
||||
|
||||
#define M9_4_0 4 |
||||
#define M9_4_1 5 |
||||
#define M9_4_2 6 |
||||
#define M9_4_3 7 |
||||
#define M9_4_4 8 |
||||
#define M9_4_5 0 |
||||
#define M9_4_6 1 |
||||
#define M9_4_7 2 |
||||
|
||||
#define M9_5_0 5 |
||||
#define M9_5_1 6 |
||||
#define M9_5_2 7 |
||||
#define M9_5_3 8 |
||||
#define M9_5_4 0 |
||||
#define M9_5_5 1 |
||||
#define M9_5_6 2 |
||||
#define M9_5_7 3 |
||||
|
||||
#define M9_6_0 6 |
||||
#define M9_6_1 7 |
||||
#define M9_6_2 8 |
||||
#define M9_6_3 0 |
||||
#define M9_6_4 1 |
||||
#define M9_6_5 2 |
||||
#define M9_6_6 3 |
||||
#define M9_6_7 4 |
||||
|
||||
#define M9_7_0 7 |
||||
#define M9_7_1 8 |
||||
#define M9_7_2 0 |
||||
#define M9_7_3 1 |
||||
#define M9_7_4 2 |
||||
#define M9_7_5 3 |
||||
#define M9_7_6 4 |
||||
#define M9_7_7 5 |
||||
|
||||
#define M9_8_0 8 |
||||
#define M9_8_1 0 |
||||
#define M9_8_2 1 |
||||
#define M9_8_3 2 |
||||
#define M9_8_4 3 |
||||
#define M9_8_5 4 |
||||
#define M9_8_6 5 |
||||
#define M9_8_7 6 |
||||
|
||||
#define M9_9_0 0 |
||||
#define M9_9_1 1 |
||||
#define M9_9_2 2 |
||||
#define M9_9_3 3 |
||||
#define M9_9_4 4 |
||||
#define M9_9_5 5 |
||||
#define M9_9_6 6 |
||||
#define M9_9_7 7 |
||||
|
||||
#define M9_10_0 1 |
||||
#define M9_10_1 2 |
||||
#define M9_10_2 3 |
||||
#define M9_10_3 4 |
||||
#define M9_10_4 5 |
||||
#define M9_10_5 6 |
||||
#define M9_10_6 7 |
||||
#define M9_10_7 8 |
||||
|
||||
#define M9_11_0 2 |
||||
#define M9_11_1 3 |
||||
#define M9_11_2 4 |
||||
#define M9_11_3 5 |
||||
#define M9_11_4 6 |
||||
#define M9_11_5 7 |
||||
#define M9_11_6 8 |
||||
#define M9_11_7 0 |
||||
|
||||
#define M9_12_0 3 |
||||
#define M9_12_1 4 |
||||
#define M9_12_2 5 |
||||
#define M9_12_3 6 |
||||
#define M9_12_4 7 |
||||
#define M9_12_5 8 |
||||
#define M9_12_6 0 |
||||
#define M9_12_7 1 |
||||
|
||||
#define M9_13_0 4 |
||||
#define M9_13_1 5 |
||||
#define M9_13_2 6 |
||||
#define M9_13_3 7 |
||||
#define M9_13_4 8 |
||||
#define M9_13_5 0 |
||||
#define M9_13_6 1 |
||||
#define M9_13_7 2 |
||||
|
||||
#define M9_14_0 5 |
||||
#define M9_14_1 6 |
||||
#define M9_14_2 7 |
||||
#define M9_14_3 8 |
||||
#define M9_14_4 0 |
||||
#define M9_14_5 1 |
||||
#define M9_14_6 2 |
||||
#define M9_14_7 3 |
||||
|
||||
#define M9_15_0 6 |
||||
#define M9_15_1 7 |
||||
#define M9_15_2 8 |
||||
#define M9_15_3 0 |
||||
#define M9_15_4 1 |
||||
#define M9_15_5 2 |
||||
#define M9_15_6 3 |
||||
#define M9_15_7 4 |
||||
|
||||
#define M9_16_0 7 |
||||
#define M9_16_1 8 |
||||
#define M9_16_2 0 |
||||
#define M9_16_3 1 |
||||
#define M9_16_4 2 |
||||
#define M9_16_5 3 |
||||
#define M9_16_6 4 |
||||
#define M9_16_7 5 |
||||
|
||||
#define M9_17_0 8 |
||||
#define M9_17_1 0 |
||||
#define M9_17_2 1 |
||||
#define M9_17_3 2 |
||||
#define M9_17_4 3 |
||||
#define M9_17_5 4 |
||||
#define M9_17_6 5 |
||||
#define M9_17_7 6 |
||||
|
||||
#define M9_18_0 0 |
||||
#define M9_18_1 1 |
||||
#define M9_18_2 2 |
||||
#define M9_18_3 3 |
||||
#define M9_18_4 4 |
||||
#define M9_18_5 5 |
||||
#define M9_18_6 6 |
||||
#define M9_18_7 7 |
||||
|
||||
/* |
||||
* M3_ ## s ## _ ## i evaluates to s+i mod 3 (0 <= s <= 18, 0 <= i <= 1). |
||||
*/ |
||||
|
||||
#define M3_0_0 0 |
||||
#define M3_0_1 1 |
||||
#define M3_1_0 1 |
||||
#define M3_1_1 2 |
||||
#define M3_2_0 2 |
||||
#define M3_2_1 0 |
||||
#define M3_3_0 0 |
||||
#define M3_3_1 1 |
||||
#define M3_4_0 1 |
||||
#define M3_4_1 2 |
||||
#define M3_5_0 2 |
||||
#define M3_5_1 0 |
||||
#define M3_6_0 0 |
||||
#define M3_6_1 1 |
||||
#define M3_7_0 1 |
||||
#define M3_7_1 2 |
||||
#define M3_8_0 2 |
||||
#define M3_8_1 0 |
||||
#define M3_9_0 0 |
||||
#define M3_9_1 1 |
||||
#define M3_10_0 1 |
||||
#define M3_10_1 2 |
||||
#define M3_11_0 2 |
||||
#define M3_11_1 0 |
||||
#define M3_12_0 0 |
||||
#define M3_12_1 1 |
||||
#define M3_13_0 1 |
||||
#define M3_13_1 2 |
||||
#define M3_14_0 2 |
||||
#define M3_14_1 0 |
||||
#define M3_15_0 0 |
||||
#define M3_15_1 1 |
||||
#define M3_16_0 1 |
||||
#define M3_16_1 2 |
||||
#define M3_17_0 2 |
||||
#define M3_17_1 0 |
||||
#define M3_18_0 0 |
||||
#define M3_18_1 1 |
||||
|
||||
#define XCAT(x, y) XCAT_(x, y) |
||||
#define XCAT_(x, y) x ## y |
||||
|
||||
#define SKBI(k, s, i) XCAT(k, XCAT(XCAT(XCAT(M9_, s), _), i)) |
||||
#define SKBT(t, s, v) XCAT(t, XCAT(XCAT(XCAT(M3_, s), _), v)) |
||||
|
||||
#define TFBIG_KINIT(k0, k1, k2, k3, k4, k5, k6, k7, k8, t0, t1, t2) { \ |
||||
k8 = ((k0 ^ k1) ^ (k2 ^ k3)) ^ ((k4 ^ k5) ^ (k6 ^ k7)) \ |
||||
^ SPH_C64(0x1BD11BDAA9FC1A22); \ |
||||
t2 = t0 ^ t1; \ |
||||
} |
||||
|
||||
#define TFBIG_ADDKEY(w0, w1, w2, w3, w4, w5, w6, w7, k, t, s) { \ |
||||
w0 = (w0 + SKBI(k, s, 0)); \ |
||||
w1 = (w1 + SKBI(k, s, 1)); \ |
||||
w2 = (w2 + SKBI(k, s, 2)); \ |
||||
w3 = (w3 + SKBI(k, s, 3)); \ |
||||
w4 = (w4 + SKBI(k, s, 4)); \ |
||||
w5 = (w5 + SKBI(k, s, 5) + SKBT(t, s, 0)); \ |
||||
w6 = (w6 + SKBI(k, s, 6) + SKBT(t, s, 1)); \ |
||||
w7 = (w7 + SKBI(k, s, 7) + (uint64_t)s); \ |
||||
} |
||||
|
||||
#define TFBIG_MIX(x0, x1, rc) { \ |
||||
x0 = x0 + x1; \ |
||||
x1 = ROTL64(x1, rc) ^ x0; \ |
||||
} |
||||
|
||||
#define TFBIG_MIX8(w0, w1, w2, w3, w4, w5, w6, w7, rc0, rc1, rc2, rc3) { \ |
||||
TFBIG_MIX(w0, w1, rc0); \ |
||||
TFBIG_MIX(w2, w3, rc1); \ |
||||
TFBIG_MIX(w4, w5, rc2); \ |
||||
TFBIG_MIX(w6, w7, rc3); \ |
||||
} |
||||
|
||||
#define TFBIG_4e(s) { \ |
||||
TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \ |
||||
TFBIG_MIX8(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 46, 36, 19, 37); \ |
||||
TFBIG_MIX8(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 33, 27, 14, 42); \ |
||||
TFBIG_MIX8(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 17, 49, 36, 39); \ |
||||
TFBIG_MIX8(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 44, 9, 54, 56); \ |
||||
} |
||||
|
||||
#define TFBIG_4o(s) { \ |
||||
TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \ |
||||
TFBIG_MIX8(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 39, 30, 34, 24); \ |
||||
TFBIG_MIX8(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 13, 50, 10, 17); \ |
||||
TFBIG_MIX8(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 25, 29, 39, 43); \ |
||||
TFBIG_MIX8(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 8, 35, 56, 22); \ |
||||
} |
||||
|
||||
static __constant__ uint64_t d_constMem[8]; |
||||
static uint64_t h_constMem[8] = { |
||||
SPH_C64(0x4903ADFF749C51CE), |
||||
SPH_C64(0x0D95DE399746DF03), |
||||
SPH_C64(0x8FD1934127C79BCE), |
||||
SPH_C64(0x9A255629FF352CB1), |
||||
SPH_C64(0x5DB62599DF6CA7B0), |
||||
SPH_C64(0xEABE394CA9D5C3F4), |
||||
SPH_C64(0x991112C71A75B523), |
||||
SPH_C64(0xAE18A40B660FCC33) }; |
||||
|
||||
__global__ void quark_skein512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector) |
||||
{ |
||||
int thread = (blockDim.x * blockIdx.x + threadIdx.x); |
||||
if (thread < threads) |
||||
{ |
||||
// Skein |
||||
uint64_t p[8]; |
||||
uint64_t h0, h1, h2, h3, h4, h5, h6, h7, h8; |
||||
uint64_t t0, t1, t2; |
||||
|
||||
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread); |
||||
|
||||
int hashPosition = nounce - startNounce; |
||||
uint64_t *inpHash = &g_hash[8 * hashPosition]; |
||||
|
||||
// Initialisierung |
||||
h0 = d_constMem[0]; |
||||
h1 = d_constMem[1]; |
||||
h2 = d_constMem[2]; |
||||
h3 = d_constMem[3]; |
||||
h4 = d_constMem[4]; |
||||
h5 = d_constMem[5]; |
||||
h6 = d_constMem[6]; |
||||
h7 = d_constMem[7]; |
||||
|
||||
// 1. Runde -> etype = 480, ptr = 64, bcount = 0, data = msg |
||||
#pragma unroll 8 |
||||
for(int i=0;i<8;i++) |
||||
p[i] = inpHash[i]; |
||||
|
||||
t0 = 64; // ptr |
||||
t1 = 480ull << 55; // etype |
||||
TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2); |
||||
TFBIG_4e(0); |
||||
TFBIG_4o(1); |
||||
TFBIG_4e(2); |
||||
TFBIG_4o(3); |
||||
TFBIG_4e(4); |
||||
TFBIG_4o(5); |
||||
TFBIG_4e(6); |
||||
TFBIG_4o(7); |
||||
TFBIG_4e(8); |
||||
TFBIG_4o(9); |
||||
TFBIG_4e(10); |
||||
TFBIG_4o(11); |
||||
TFBIG_4e(12); |
||||
TFBIG_4o(13); |
||||
TFBIG_4e(14); |
||||
TFBIG_4o(15); |
||||
TFBIG_4e(16); |
||||
TFBIG_4o(17); |
||||
TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18); |
||||
|
||||
h0 = inpHash[0] ^ p[0]; |
||||
h1 = inpHash[1] ^ p[1]; |
||||
h2 = inpHash[2] ^ p[2]; |
||||
h3 = inpHash[3] ^ p[3]; |
||||
h4 = inpHash[4] ^ p[4]; |
||||
h5 = inpHash[5] ^ p[5]; |
||||
h6 = inpHash[6] ^ p[6]; |
||||
h7 = inpHash[7] ^ p[7]; |
||||
|
||||
// 2. Runde -> etype = 510, ptr = 8, bcount = 0, data = 0 |
||||
#pragma unroll 8 |
||||
for(int i=0;i<8;i++) |
||||
p[i] = 0; |
||||
|
||||
t0 = 8; // ptr |
||||
t1 = 510ull << 55; // etype |
||||
TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2); |
||||
TFBIG_4e(0); |
||||
TFBIG_4o(1); |
||||
TFBIG_4e(2); |
||||
TFBIG_4o(3); |
||||
TFBIG_4e(4); |
||||
TFBIG_4o(5); |
||||
TFBIG_4e(6); |
||||
TFBIG_4o(7); |
||||
TFBIG_4e(8); |
||||
TFBIG_4o(9); |
||||
TFBIG_4e(10); |
||||
TFBIG_4o(11); |
||||
TFBIG_4e(12); |
||||
TFBIG_4o(13); |
||||
TFBIG_4e(14); |
||||
TFBIG_4o(15); |
||||
TFBIG_4e(16); |
||||
TFBIG_4o(17); |
||||
TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18); |
||||
|
||||
// fertig |
||||
uint64_t *outpHash = &g_hash[8 * hashPosition]; |
||||
|
||||
#pragma unroll 8 |
||||
for(int i=0;i<8;i++) |
||||
outpHash[i] = p[i]; |
||||
} |
||||
} |
||||
|
||||
// Setup-Funktionen |
||||
__host__ void quark_skein512_cpu_init(int thr_id, int threads) |
||||
{ |
||||
// nix zu tun ;-) |
||||
cudaMemcpyToSymbol( d_constMem, |
||||
h_constMem, |
||||
sizeof(h_constMem), |
||||
0, cudaMemcpyHostToDevice); |
||||
} |
||||
|
||||
__host__ void quark_skein512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) |
||||
{ |
||||
const int threadsperblock = 256; |
||||
|
||||
// berechne wie viele Thread Blocks wir brauchen |
||||
dim3 grid((threads + threadsperblock-1)/threadsperblock); |
||||
dim3 block(threadsperblock); |
||||
|
||||
// Größe des dynamischen Shared Memory Bereichs |
||||
size_t shared_size = 0; |
||||
|
||||
// fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size); |
||||
quark_skein512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector); |
||||
|
||||
// Strategisches Sleep Kommando zur Senkung der CPU Last |
||||
MyStreamSynchronize(NULL, order, thr_id); |
||||
} |
||||
Loading…
Reference in new issue