From e2bfb4b14f061d1f0b0cd38663b291ebfa02d6b3 Mon Sep 17 00:00:00 2001 From: Christian Buchner Date: Thu, 8 May 2014 09:25:37 +0200 Subject: [PATCH] committing a quick attempt at NIST5 (TalkCoin) --- Makefile.am | 3 +- ccminer.vcxproj | 6 ++ ccminer.vcxproj.filters | 3 + cpu-miner.c | 28 ++++--- cuda_nist5.cu | 161 ++++++++++++++++++++++++++++++++++++++++ miner.h | 4 + 6 files changed, 194 insertions(+), 11 deletions(-) create mode 100644 cuda_nist5.cu diff --git a/Makefile.am b/Makefile.am index e8861988a..8cfdbd14a 100644 --- a/Makefile.am +++ b/Makefile.am @@ -27,12 +27,13 @@ ccminer_SOURCES = elist.h miner.h compat.h \ heavy/cuda_sha256.cu heavy/cuda_sha256.h \ fuguecoin.cpp cuda_fugue256.cu sph/fugue.c sph/sph_fugue.h uint256.h \ groestlcoin.cpp cuda_groestlcoin.cu cuda_groestlcoin.h \ + myriadgroestl.cpp cuda_myriadgroestl.cu \ JHA/jackpotcoin.cu JHA/cuda_jha_keccak512.cu \ JHA/cuda_jha_compactionTest.cu quark/cuda_quark_checkhash.cu \ quark/cuda_jh512.cu quark/cuda_quark_blake512.cu quark/cuda_quark_groestl512.cu quark/cuda_skein512.cu \ quark/cuda_bmw512.cu quark/cuda_quark_keccak512.cu quark/quarkcoin.cu quark/animecoin.cu \ quark/cuda_quark_compactionTest.cu \ - myriadgroestl.cpp cuda_myriadgroestl.cu + cuda_nist5.cu ccminer_LDFLAGS = $(PTHREAD_FLAGS) @CUDA_LDFLAGS@ ccminer_LDADD = @LIBCURL@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @WS2_LIBS@ @CUDA_LIBS@ @OPENMP_CFLAGS@ @LIBS@ diff --git a/ccminer.vcxproj b/ccminer.vcxproj index 0036b8098..7f212611c 100644 --- a/ccminer.vcxproj +++ b/ccminer.vcxproj @@ -305,6 +305,12 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)" -Xptxas "-abi=no -v" %(AdditionalOptions) -Xptxas "-abi=no -v" %(AdditionalOptions) + + -Xptxas "-abi=no -v" %(AdditionalOptions) + -Xptxas "-abi=no -v" %(AdditionalOptions) + -Xptxas "-abi=no -v" %(AdditionalOptions) + -Xptxas "-abi=no -v" %(AdditionalOptions) + -Xptxas "-abi=no -v" %(AdditionalOptions) -Xptxas "-abi=no -v" %(AdditionalOptions) diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters index dba16228d..d86d89d25 100644 --- a/ccminer.vcxproj.filters +++ b/ccminer.vcxproj.filters @@ -310,5 +310,8 @@ Source Files\CUDA\quark + + Source Files\CUDA + \ No newline at end of file diff --git a/cpu-miner.c b/cpu-miner.c index b0ca293be..76c138ebf 100644 --- a/cpu-miner.c +++ b/cpu-miner.c @@ -128,7 +128,8 @@ typedef enum { ALGO_MYR_GR, ALGO_JACKPOT, ALGO_QUARK, - ALGO_ANIME + ALGO_ANIME, + ALGO_NIST5 } sha256_algos; static const char *algo_names[] = { @@ -139,7 +140,8 @@ static const char *algo_names[] = { "myr-gr", "jackpot", "quark", - "anime" + "anime", + "nist5" }; bool opt_debug = false; @@ -209,6 +211,7 @@ Options:\n\ jackpot Jackpot hash\n\ quark Quark hash\n\ anime Animecoin hash\n\ + nist5 NIST5 (TalkCoin) hash\n\ -d, --devices takes a comma separated list of CUDA devices to use.\n\ Device IDs start counting from 0! Alternatively takes\n\ string names of your cards like gtx780ti or gt640#2\n\ @@ -883,15 +886,20 @@ static void *miner_thread(void *userdata) max_nonce, &hashes_done); break; - case ALGO_QUARK: - rc = scanhash_quark(thr_id, work.data, work.target, - max_nonce, &hashes_done); - break; + case ALGO_QUARK: + rc = scanhash_quark(thr_id, work.data, work.target, + max_nonce, &hashes_done); + break; - case ALGO_ANIME: - rc = scanhash_anime(thr_id, work.data, work.target, - max_nonce, &hashes_done); - break; + case ALGO_ANIME: + rc = scanhash_anime(thr_id, work.data, work.target, + max_nonce, &hashes_done); + break; + + case ALGO_NIST5: + rc = scanhash_nist5(thr_id, work.data, work.target, + max_nonce, &hashes_done); + break; default: /* should never happen */ diff --git a/cuda_nist5.cu b/cuda_nist5.cu new file mode 100644 index 000000000..0296b4e99 --- /dev/null +++ b/cuda_nist5.cu @@ -0,0 +1,161 @@ + +extern "C" +{ +#include "sph/sph_blake.h" +#include "sph/sph_groestl.h" +#include "sph/sph_skein.h" +#include "sph/sph_jh.h" +#include "sph/sph_keccak.h" +#include "miner.h" +} + +#include + +// aus cpu-miner.c +extern int device_map[8]; + +// Speicher für Input/Output der verketteten Hashfunktionen +static uint32_t *d_hash[8]; + +extern void quark_blake512_cpu_init(int thr_id, int threads); +extern void quark_blake512_cpu_setBlock_80(void *pdata); +extern void quark_blake512_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_hash, int order); + +extern void quark_groestl512_cpu_init(int thr_id, int threads); +extern void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); + +extern void quark_jh512_cpu_init(int thr_id, int threads); +extern void quark_jh512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); + +extern void quark_keccak512_cpu_init(int thr_id, int threads); +extern void quark_keccak512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); + +extern void quark_skein512_cpu_init(int thr_id, int threads); +extern void quark_skein512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); + +extern void quark_check_cpu_init(int thr_id, int threads); +extern void quark_check_cpu_setTarget(const void *ptarget); +extern uint32_t quark_check_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order); + +// Original nist5hash Funktion aus einem miner Quelltext +inline void nist5hash(void *state, const void *input) +{ + sph_blake512_context ctx_blake; + sph_groestl512_context ctx_groestl; + sph_jh512_context ctx_jh; + sph_keccak512_context ctx_keccak; + sph_skein512_context ctx_skein; + + unsigned char hash[64]; + + sph_blake512_init(&ctx_blake); + // ZBLAKE; + sph_blake512 (&ctx_blake, input, 80); + sph_blake512_close(&ctx_blake, (void*) hash); + + sph_groestl512_init(&ctx_groestl); + // ZGROESTL; + sph_groestl512 (&ctx_groestl, (const void*) hash, 64); + sph_groestl512_close(&ctx_groestl, (void*) hash); + + sph_jh512_init(&ctx_jh); + // ZJH; + sph_jh512 (&ctx_jh, (const void*) hash, 64); + sph_jh512_close(&ctx_jh, (void*) hash); + + sph_keccak512_init(&ctx_keccak); + // ZKECCAK; + sph_keccak512 (&ctx_keccak, (const void*) hash, 64); + sph_keccak512_close(&ctx_keccak, (void*) hash); + + sph_skein512_init(&ctx_skein); + // ZSKEIN; + sph_skein512 (&ctx_skein, (const void*) hash, 64); + sph_skein512_close(&ctx_skein, (void*) hash); + + memcpy(state, hash, 32); +} + + +extern bool opt_benchmark; + +extern "C" int scanhash_nist5(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]; + + // TODO: entfernen für eine Release! Ist nur zum Testen! + if (opt_benchmark) + ((uint32_t*)ptarget)[7] = 0x0000ff; + + const uint32_t Htarg = ptarget[7]; + + const int throughput = 256*4096; // 100; + + static bool init[8] = {0,0,0,0,0,0,0,0}; + if (!init[thr_id]) + { + cudaSetDevice(device_map[thr_id]); + + // Konstanten kopieren, Speicher belegen + cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput); + quark_blake512_cpu_init(thr_id, throughput); + quark_groestl512_cpu_init(thr_id, throughput); + quark_jh512_cpu_init(thr_id, throughput); + quark_keccak512_cpu_init(thr_id, throughput); + quark_skein512_cpu_init(thr_id, throughput); + quark_check_cpu_init(thr_id, throughput); + init[thr_id] = true; + } + + uint32_t endiandata[20]; + for (int k=0; k < 20; k++) + be32enc(&endiandata[k], ((uint32_t*)pdata)[k]); + + quark_blake512_cpu_setBlock_80((void*)endiandata); + quark_check_cpu_setTarget(ptarget); + + do { + int order = 0; + + // erstes Blake512 Hash mit CUDA + quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); + + // das ist der unbedingte Branch für Groestl512 + quark_groestl512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + + // das ist der unbedingte Branch für JH512 + quark_jh512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + + // das ist der unbedingte Branch für Keccak512 + quark_keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + + // das ist der unbedingte Branch für Skein512 + quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + + // Scan nach Gewinner Hashes auf der GPU + uint32_t foundNonce = quark_check_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + if (foundNonce != 0xffffffff) + { + uint32_t vhash64[8]; + be32enc(&endiandata[19], foundNonce); + nist5hash(vhash64, endiandata); + + if ((vhash64[7]<=Htarg) && fulltest(vhash64, ptarget)) { + + pdata[19] = foundNonce; + *hashes_done = foundNonce - first_nonce + 1; + return 1; + } else { + applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU!", thr_id, foundNonce); + } + } + + pdata[19] += throughput; + + } while (pdata[19] < max_nonce && !work_restart[thr_id].restart); + + *hashes_done = pdata[19] - first_nonce + 1; + return 0; +} diff --git a/miner.h b/miner.h index e3f4b8f59..f565d147e 100644 --- a/miner.h +++ b/miner.h @@ -231,6 +231,10 @@ extern int scanhash_anime(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, unsigned long *hashes_done); +extern int scanhash_nist5(int thr_id, uint32_t *pdata, + const uint32_t *ptarget, uint32_t max_nonce, + unsigned long *hashes_done); + extern void fugue256_hash(unsigned char* output, const unsigned char* input, int len); extern void heavycoin_hash(unsigned char* output, const unsigned char* input, int len); extern void groestlcoin_hash(unsigned char* output, const unsigned char* input, int len);