From 81051dd75fe53b766012f619f2f5e7bea8f91005 Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Mon, 9 May 2016 21:02:45 +0200 Subject: [PATCH] x17: restore nonceVector parameters could be used by new algos --- x17/cuda_x17_haval256.cu | 10 ++++++---- x17/cuda_x17_sha512.cu | 13 ++++++++----- x17/x17.cu | 28 +++++++++++++--------------- 3 files changed, 27 insertions(+), 24 deletions(-) diff --git a/x17/cuda_x17_haval256.cu b/x17/cuda_x17_haval256.cu index 8573e64..a7b9463 100644 --- a/x17/cuda_x17_haval256.cu +++ b/x17/cuda_x17_haval256.cu @@ -259,12 +259,14 @@ } __global__ /* __launch_bounds__(256, 6) */ -void x17_haval256_gpu_hash_64(uint32_t threads, uint64_t *g_hash) +void x17_haval256_gpu_hash_64(const uint32_t threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector) { uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); if (thread < threads) { - uint64_t *pHash = &g_hash[thread*8U]; + uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread); + uint64_t hashPosition = nounce - startNounce; + uint64_t *pHash = &g_hash[hashPosition*8U]; uint32_t s0, s1, s2, s3, s4, s5, s6, s7; const uint32_t u0 = s0 = 0x243F6A88; @@ -338,14 +340,14 @@ void x17_haval256_cpu_init(int thr_id, uint32_t threads) } __host__ -void x17_haval256_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash, int order) +void x17_haval256_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) { const uint32_t threadsperblock = 256; dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 block(threadsperblock); - x17_haval256_gpu_hash_64 <<>> (threads, (uint64_t*)d_hash); + x17_haval256_gpu_hash_64 <<>> (threads, startNounce, (uint64_t*)d_hash, d_nonceVector); //MyStreamSynchronize(NULL, order, thr_id); } diff --git a/x17/cuda_x17_sha512.cu b/x17/cuda_x17_sha512.cu index 088385e..6c97405 100644 --- a/x17/cuda_x17_sha512.cu +++ b/x17/cuda_x17_sha512.cu @@ -87,13 +87,16 @@ uint64_t Tone(uint64_t* K, uint64_t* r, uint64_t* W, const int a, const int i) r[(a+7) & 7] = T1 + T2; \ } -__global__ /*__launch_bounds__(256, 4)*/ -void x17_sha512_gpu_hash_64(const uint32_t threads, uint64_t *g_hash) +__global__ +/*__launch_bounds__(256, 4)*/ +void x17_sha512_gpu_hash_64(const uint32_t threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector) { const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); if (thread < threads) { - uint64_t *pHash = &g_hash[thread*8U]; + uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread); + uint64_t hashPosition = nounce - startNounce; + uint64_t *pHash = &g_hash[hashPosition*8U]; uint64_t W[80]; #pragma unroll @@ -158,14 +161,14 @@ void x17_sha512_cpu_init(int thr_id, uint32_t threads) } __host__ -void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash, int order) +void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) { const uint32_t threadsperblock = 256; dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 block(threadsperblock); - x17_sha512_gpu_hash_64 <<>> (threads, (uint64_t*)d_hash); + x17_sha512_gpu_hash_64 <<>> (threads, startNounce, (uint64_t*)d_hash, d_nonceVector); //MyStreamSynchronize(NULL, order, thr_id); } diff --git a/x17/x17.cu b/x17/x17.cu index 92d5de9..b20e574 100644 --- a/x17/x17.cu +++ b/x17/x17.cu @@ -1,10 +1,8 @@ -/* - * X17 algorithm built on cbuchner1's original X11 - * +/** + * X17 algorithm (X15 + sha512 + haval256) */ -extern "C" -{ +extern "C" { #include "sph/sph_blake.h" #include "sph/sph_bmw.h" #include "sph/sph_groestl.h" @@ -49,16 +47,18 @@ extern void x15_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t sta extern void x15_whirlpool_cpu_free(int thr_id); extern void x17_sha512_cpu_init(int thr_id, uint32_t threads); -extern void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash, int order); +extern void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); extern void x17_haval256_cpu_init(int thr_id, uint32_t threads); -extern void x17_haval256_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash, int order); +extern void x17_haval256_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); -// X17 Hashfunktion +// X17 CPU Hash (Validation) extern "C" void x17hash(void *output, const void *input) { - // blake1-bmw2-grs3-skein4-jh5-keccak6-luffa7-cubehash8-shavite9-simd10-echo11-hamsi12-fugue13-shabal14-whirlpool15-sha512-haval17 + unsigned char _ALIGN(64) hash[128]; + + // x11 + hamsi12-fugue13-shabal14-whirlpool15-sha512-haval256 sph_blake512_context ctx_blake; sph_bmw512_context ctx_bmw; @@ -78,8 +78,6 @@ extern "C" void x17hash(void *output, const void *input) sph_sha512_context ctx_sha512; sph_haval256_5_context ctx_haval; - unsigned char hash[128]; // uint32_t hashA[16], hashB[16]; - #define hashB hash+64 sph_blake512_init(&ctx_blake); sph_blake512(&ctx_blake, input, 80); @@ -199,7 +197,7 @@ extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, u init[thr_id] = true; } - uint32_t endiandata[20]; + uint32_t _ALIGN(64) endiandata[20]; for (int k=0; k < 20; k++) be32enc(&endiandata[k], pdata[k]); @@ -226,8 +224,8 @@ extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, u x13_fugue512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); x14_shabal512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); - x17_sha512_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], order++); - x17_haval256_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], order++); + x17_sha512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + x17_haval256_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); *hashes_done = pdata[19] - first_nonce + throughput; @@ -235,7 +233,7 @@ extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, u if (foundNonce != UINT32_MAX) { const uint32_t Htarg = ptarget[7]; - uint32_t vhash64[8]; + uint32_t _ALIGN(64) vhash64[8]; be32enc(&endiandata[19], foundNonce); x17hash(vhash64, endiandata);