From afe57f83412cc133d8ae2312c1cb11fb4a943475 Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Tue, 12 May 2015 04:37:27 +0200 Subject: [PATCH] rename skein2 to c++, no cuda kernel code and some other changes... --- Makefile.am | 2 +- ccminer.vcxproj | 6 ++---- ccminer.vcxproj.filters | 8 ++++---- cuda_checkhash.cu | 16 ++++++++-------- groestlcoin.cpp | 2 +- myriadgroestl.cpp | 4 ++-- skein.cu | 15 +++++++++++---- skein2.cu => skein2.cpp | 11 +++++++---- 8 files changed, 36 insertions(+), 28 deletions(-) rename skein2.cu => skein2.cpp (90%) diff --git a/Makefile.am b/Makefile.am index a8d19b5..8878096 100644 --- a/Makefile.am +++ b/Makefile.am @@ -42,7 +42,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \ quark/quarkcoin.cu quark/animecoin.cu \ quark/cuda_quark_compactionTest.cu \ neoscrypt/neoscrypt.cpp neoscrypt/neoscrypt-cpu.c neoscrypt/cuda_neoscrypt.cu \ - cuda_nist5.cu pentablake.cu skein.cu skein2.cu zr5.cu \ + cuda_nist5.cu pentablake.cu skein.cu skein2.cpp zr5.cu \ sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c \ sph/cubehash.c sph/echo.c sph/luffa.c sph/sha2.c sph/shavite.c sph/simd.c \ sph/hamsi.c sph/hamsi_helper.c sph/sph_hamsi.h \ diff --git a/ccminer.vcxproj b/ccminer.vcxproj index 485da3d..0a8d629 100644 --- a/ccminer.vcxproj +++ b/ccminer.vcxproj @@ -270,6 +270,7 @@ + @@ -464,9 +465,6 @@ 64 - - 64 - true @@ -538,4 +536,4 @@ - \ No newline at end of file + diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters index 9cc242c..c7042db 100644 --- a/ccminer.vcxproj.filters +++ b/ccminer.vcxproj.filters @@ -240,6 +240,9 @@ Source Files\neoscrypt + + Source Files\CUDA + @@ -601,9 +604,6 @@ Source Files\CUDA - - Source Files\CUDA - Source Files\CUDA\scrypt @@ -653,4 +653,4 @@ Ressources - \ No newline at end of file + diff --git a/cuda_checkhash.cu b/cuda_checkhash.cu index 4bbef59..a40f75e 100644 --- a/cuda_checkhash.cu +++ b/cuda_checkhash.cu @@ -17,15 +17,15 @@ static uint32_t* d_resNonces[MAX_GPUS]; __host__ void cuda_check_cpu_init(int thr_id, uint32_t threads) { - CUDA_CALL_OR_RET(cudaMallocHost(&h_resNonces[thr_id], 8*sizeof(uint32_t))); - CUDA_CALL_OR_RET(cudaMalloc(&d_resNonces[thr_id], 8*sizeof(uint32_t))); + CUDA_CALL_OR_RET(cudaMallocHost(&h_resNonces[thr_id], 32)); + CUDA_CALL_OR_RET(cudaMalloc(&d_resNonces[thr_id], 32)); } // Target Difficulty __host__ void cuda_check_cpu_setTarget(const void *ptarget) { - CUDA_SAFE_CALL(cudaMemcpyToSymbol(pTarget, ptarget, 8*sizeof(uint32_t), 0, cudaMemcpyHostToDevice)); + CUDA_SAFE_CALL(cudaMemcpyToSymbol(pTarget, ptarget, 32, 0, cudaMemcpyHostToDevice)); } /* --------------------------------------------------------------------------------------------- */ @@ -177,11 +177,11 @@ void cuda_check_hash_branch_64(uint32_t threads, uint32_t startNounce, uint32_t __host__ uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order) { - uint32_t result = 0xffffffff; - cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t)); - const uint32_t threadsperblock = 256; + uint32_t result = UINT32_MAX; + cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t)); + dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 block(threadsperblock); @@ -199,9 +199,9 @@ uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNoun /* Function to get the compiled Shader Model version */ int cuda_arch[MAX_GPUS] = { 0 }; -__global__ -void nvcc_get_arch(int *d_version) +__global__ void nvcc_get_arch(int *d_version) { + *d_version = 0; #ifdef __CUDA_ARCH__ *d_version = __CUDA_ARCH__; #endif diff --git a/groestlcoin.cpp b/groestlcoin.cpp index 400b8ef..6956102 100644 --- a/groestlcoin.cpp +++ b/groestlcoin.cpp @@ -78,8 +78,8 @@ int scanhash_groestlcoin(int thr_id, uint32_t *pdata, const uint32_t *ptarget, } if ((uint64_t) pdata[19] + throughput > max_nonce) { + *hashes_done = pdata[19] - start_nonce + 1; pdata[19] = max_nonce; - *hashes_done = max_nonce - start_nonce + 1; break; } pdata[19] += throughput; diff --git a/myriadgroestl.cpp b/myriadgroestl.cpp index 5464819..4186334 100644 --- a/myriadgroestl.cpp +++ b/myriadgroestl.cpp @@ -35,7 +35,7 @@ int scanhash_myriad(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, unsigned long *hashes_done) { uint32_t _ALIGN(64) endiandata[32]; - uint32_t start_nonce = pdata[19]++; + uint32_t start_nonce = pdata[19]; uint32_t throughput = device_intensity(thr_id, __func__, 1 << 17); throughput = min(throughput, max_nonce - start_nonce); @@ -82,8 +82,8 @@ int scanhash_myriad(int thr_id, uint32_t *pdata, const uint32_t *ptarget, } if ((uint64_t) pdata[19] + throughput > max_nonce) { + *hashes_done = pdata[19] - start_nonce; pdata[19] = max_nonce; - *hashes_done = max_nonce - start_nonce + 1; break; } pdata[19] += throughput; diff --git a/skein.cu b/skein.cu index 131812b..38b4bcc 100644 --- a/skein.cu +++ b/skein.cu @@ -396,12 +396,13 @@ extern "C" int scanhash_skeincoin(int thr_id, uint32_t *pdata, const uint32_t *p int res = 1; uint8_t num = res; uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], num); - while (secNonce != 0 && res < 6) + while (secNonce != 0 && res < 2) /* todo: up to 6 */ { endiandata[19] = swab32_if(secNonce, swap); skeincoinhash(vhash64, endiandata); if (vhash64[7] <= ptarget[7] && fulltest(vhash64, ptarget)) { - pdata[19+res] = swab32_if(secNonce, !swap); + // todo: use 19 20 21... zr5 pok to adapt... + pdata[19+res*2] = swab32_if(secNonce, !swap); res++; } num++; @@ -416,10 +417,16 @@ extern "C" int scanhash_skeincoin(int thr_id, uint32_t *pdata, const uint32_t *p applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!", device_map[thr_id], foundNonce); } } + + if ((uint64_t) pdata[19] + throughput > max_nonce) { + *hashes_done = pdata[19] - first_nonce; + pdata[19] = max_nonce; + break; + } + pdata[19] += throughput; - } while (pdata[19] < max_nonce && !work_restart[thr_id].restart); + } while (!work_restart[thr_id].restart); - *hashes_done = pdata[19] - first_nonce + 1; return 0; } diff --git a/skein2.cu b/skein2.cpp similarity index 90% rename from skein2.cu rename to skein2.cpp index b264597..bdb170e 100644 --- a/skein2.cu +++ b/skein2.cpp @@ -2,6 +2,7 @@ * SKEIN512 80 + SKEIN512 64 (Woodcoin) * by tpruvot@github - 2015 */ +#include #include "sph/sph_skein.h" @@ -17,7 +18,7 @@ extern void skein512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNou extern void quark_skein512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); -extern "C" void skein2hash(void *output, const void *input) +void skein2hash(void *output, const void *input) { uint32_t _ALIGN(64) hash[16]; sph_skein512_context ctx_skein; @@ -30,12 +31,12 @@ extern "C" void skein2hash(void *output, const void *input) sph_skein512(&ctx_skein, hash, 64); sph_skein512_close(&ctx_skein, hash); - memcpy(output, hash, 32); + memcpy(output, (void*) hash, 32); } static bool init[MAX_GPUS] = { 0 }; -extern "C" int scanhash_skein2(int thr_id, uint32_t *pdata, const uint32_t *ptarget, +int scanhash_skein2(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, unsigned long *hashes_done) { const uint32_t first_nonce = pdata[19]; @@ -100,7 +101,9 @@ extern "C" int scanhash_skein2(int thr_id, uint32_t *pdata, const uint32_t *ptar } } - if (((uint64_t) throughput + pdata[19]) > max_nonce) { + if ((uint64_t) pdata[19] + throughput > max_nonce) { + *hashes_done = pdata[19] - first_nonce; + pdata[19] = max_nonce; break; }