From aa4dcdf58d6138086c920e22b452bbcb7627585d Mon Sep 17 00:00:00 2001 From: orignal Date: Mon, 17 Apr 2017 13:18:32 -0400 Subject: [PATCH] invoke gost --- algos.h | 2 ++ ccminer.cpp | 4 +++ gost/gost.cu | 95 ++++++++++++++++++++++++++++++++++++++++++++++++++++ miner.h | 1 + 4 files changed, 102 insertions(+) diff --git a/algos.h b/algos.h index ea27f7c..6c15fc6 100644 --- a/algos.h +++ b/algos.h @@ -18,6 +18,7 @@ enum sha_algos { ALGO_DMD_GR, ALGO_FRESH, ALGO_FUGUE256, /* Fugue256 */ + ALGO_GOST, ALGO_GROESTL, ALGO_HEAVY, /* Heavycoin hash */ ALGO_HMQ1725, @@ -78,6 +79,7 @@ static const char *algo_names[] = { "dmd-gr", "fresh", "fugue256", + "gost", "groestl", "heavy", "hmq1725", diff --git a/ccminer.cpp b/ccminer.cpp index 42013b2..580cc38 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -236,6 +236,7 @@ Options:\n\ dmd-gr Diamond-Groestl\n\ fresh Freshcoin (shavite 80)\n\ fugue256 Fuguecoin\n\ + gost GOST\n\ groestl Groestlcoin\n\ heavy Heavycoin\n\ hmq1725 Doubloons / Espers\n\ @@ -2323,6 +2324,9 @@ static void *miner_thread(void *userdata) case ALGO_SIB: rc = scanhash_sib(thr_id, &work, max_nonce, &hashes_done); break; + case ALGO_GOST: + rc = scanhash_gost(thr_id, &work, max_nonce, &hashes_done); + break; case ALGO_S3: rc = scanhash_s3(thr_id, &work, max_nonce, &hashes_done); break; diff --git a/gost/gost.cu b/gost/gost.cu index 65bea9f..995897d 100644 --- a/gost/gost.cu +++ b/gost/gost.cu @@ -32,3 +32,98 @@ extern "C" void gosthash(void *output, const void *input) memcpy(output, hash, 32); } +//#define _DEBUG +#define _DEBUG_PREFIX "sib" +#include "cuda_debug.cuh" + +static bool init[MAX_GPUS] = { 0 }; + +extern "C" int scanhash_gost(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done) +{ + uint32_t *pdata = work->data; + uint32_t *ptarget = work->target; + const uint32_t first_nonce = pdata[19]; + const int dev_id = device_map[thr_id]; + int intensity = (device_sm[dev_id] >= 500 && !is_windows()) ? 19 : 18; // 2^18 = 262144 cuda threads + if (device_sm[dev_id] >= 600) intensity = 20; + uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity); + if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); + + if (opt_benchmark) + ptarget[7] = 0xf; + + if (!init[thr_id]) + { + cudaSetDevice(device_map[thr_id]); + if (opt_cudaschedule == -1 && gpu_threads == 1) { + cudaDeviceReset(); + // reduce cpu usage + cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); + CUDA_LOG_ERROR(); + } + gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput); + + CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput), -1); + + cuda_check_cpu_init(thr_id, throughput); + + init[thr_id] = true; + } + + uint32_t endiandata[20]; + for (int k=0; k < 20; k++) + be32enc(&endiandata[k], pdata[k]); + + do { + int order = 0; + + // Hash with CUDA + //streebog_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); + //TRACE("gost :"); + + work->nonces[0] = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]); + if (work->nonces[0] != UINT32_MAX) + { + const uint32_t Htarg = ptarget[7]; + uint32_t _ALIGN(64) vhash[8]; + be32enc(&endiandata[19], work->nonces[0]); + sibhash(vhash, endiandata); + + if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) { + work->valid_nonces = 1; + work_set_target_ratio(work, vhash); + work->nonces[1] =cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1); + *hashes_done = pdata[19] - first_nonce + throughput; + if (work->nonces[1] != 0) { + be32enc(&endiandata[19], work->nonces[1]); + sibhash(vhash, endiandata); + bn_set_target_ratio(work, vhash, 1); + work->valid_nonces++; + pdata[19] = max(work->nonces[0], work->nonces[1]) + 1; + } else { + pdata[19] = work->nonces[0] + 1; // cursor + } + return work->valid_nonces; + } + else if (vhash[7] > Htarg) { + gpu_increment_reject(thr_id); + if (!opt_quiet) + gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", work->nonces[0]); + pdata[19] = work->nonces[0] + 1; + continue; + } + } + + if ((uint64_t) throughput + pdata[19] >= max_nonce) { + pdata[19] = max_nonce; + break; + } + pdata[19] += throughput; + + } while (!work_restart[thr_id].restart); + + *hashes_done = pdata[19] - first_nonce; + + return 0; +} + diff --git a/miner.h b/miner.h index 2f5c812..cea5dc4 100644 --- a/miner.h +++ b/miner.h @@ -304,6 +304,7 @@ extern int scanhash_sha256d(int thr_id, struct work *work, uint32_t max_nonce, u extern int scanhash_sha256t(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_sia(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_sib(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); +extern int scanhash_gost(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_skeincoin(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_skein2(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_s3(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);