GOSTSec/ccminer-gostd-lite
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

x17: restore nonceVector parameters

Browse Source 
could be used by new algos
This commit is contained in:
Tanguy Pruvot 2016-05-09 21:02:45 +02:00
parent 8d67f0f774
commit 81051dd75f
3 changed files with 27 additions and 24 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
x17/cuda_x17_haval256.cu
View File

@ -259,12 +259,14 @@
} }
__global__ /* __launch_bounds__(256, 6) */ __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); uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads) 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; uint32_t s0, s1, s2, s3, s4, s5, s6, s7;
const uint32_t u0 = s0 = 0x243F6A88; const uint32_t u0 = s0 = 0x243F6A88;
@ -338,14 +340,14 @@ void x17_haval256_cpu_init(int thr_id, uint32_t threads)
} }
__host__ __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; const uint32_t threadsperblock = 256;
dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock); dim3 block(threadsperblock);
x17_haval256_gpu_hash_64 <<<grid, block>>> (threads, (uint64_t*)d_hash); x17_haval256_gpu_hash_64 <<<grid, block>>> (threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
//MyStreamSynchronize(NULL, order, thr_id); //MyStreamSynchronize(NULL, order, thr_id);
} }

13
x17/cuda_x17_sha512.cu
View File

@ -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; \ r[(a+7) & 7] = T1 + T2; \
} }
__global__ /*__launch_bounds__(256, 4)*/ __global__
void x17_sha512_gpu_hash_64(const uint32_t threads, uint64_t *g_hash) /*__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); const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads) 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]; uint64_t W[80];
#pragma unroll #pragma unroll
@ -158,14 +161,14 @@ void x17_sha512_cpu_init(int thr_id, uint32_t threads)
} }
__host__ __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; const uint32_t threadsperblock = 256;
dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock); dim3 block(threadsperblock);
x17_sha512_gpu_hash_64 <<<grid, block>>> (threads, (uint64_t*)d_hash); x17_sha512_gpu_hash_64 <<<grid, block>>> (threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
//MyStreamSynchronize(NULL, order, thr_id); //MyStreamSynchronize(NULL, order, thr_id);
} }

28
x17/x17.cu
View File

@ -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_blake.h"
#include "sph/sph_bmw.h" #include "sph/sph_bmw.h"
#include "sph/sph_groestl.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 x15_whirlpool_cpu_free(int thr_id);
extern void x17_sha512_cpu_init(int thr_id, uint32_t threads); 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_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) 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_blake512_context ctx_blake;
sph_bmw512_context ctx_bmw; sph_bmw512_context ctx_bmw;
@ -78,8 +78,6 @@ extern "C" void x17hash(void *output, const void *input)
sph_sha512_context ctx_sha512; sph_sha512_context ctx_sha512;
sph_haval256_5_context ctx_haval; 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_init(&ctx_blake);
sph_blake512(&ctx_blake, input, 80); 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; init[thr_id] = true;
} }
uint32_t endiandata[20]; uint32_t _ALIGN(64) endiandata[20];
for (int k=0; k < 20; k++) for (int k=0; k < 20; k++)
be32enc(&endiandata[k], pdata[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++); 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++); 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++); 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_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], 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; *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) if (foundNonce != UINT32_MAX)
{ {
const uint32_t Htarg = ptarget[7]; const uint32_t Htarg = ptarget[7];
uint32_t vhash64[8]; uint32_t _ALIGN(64) vhash64[8];
be32enc(&endiandata[19], foundNonce); be32enc(&endiandata[19], foundNonce);
x17hash(vhash64, endiandata); x17hash(vhash64, endiandata);