From 11a512f2a78b40d2c1b2b3a19a31020fb1e75055 Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Sat, 5 Aug 2017 04:16:28 +0200 Subject: [PATCH] change defaults to handle cuda 9+, disable heavy and SM 2.x Heavy is the only algo using thrust which is generally broken on new cuda releases mjollnir dropped too... never seen this coin anyway... --- Makefile.am | 7 ++-- README.txt | 4 --- bench.cpp | 3 ++ ccminer.cpp | 19 ++++++----- ccminer.vcxproj | 14 ++++---- heavy/heavy.cu | 89 ++++++++++++++++++++++++++++++------------------- 6 files changed, 78 insertions(+), 58 deletions(-) diff --git a/Makefile.am b/Makefile.am index edaad61..4360cdd 100644 --- a/Makefile.am +++ b/Makefile.am @@ -116,7 +116,6 @@ nvcc_ARCH = -gencode=arch=compute_50,code=\"sm_50,compute_50\" nvcc_ARCH += -gencode=arch=compute_52,code=\"sm_52,compute_52\" #nvcc_ARCH += -gencode=arch=compute_35,code=\"sm_35,compute_35\" #nvcc_ARCH += -gencode=arch=compute_30,code=\"sm_30,compute_30\" -#nvcc_ARCH += -gencode=arch=compute_20,code=\"sm_21,compute_20\" nvcc_FLAGS = $(nvcc_ARCH) @CUDA_INCLUDES@ -I. @CUDA_CFLAGS@ nvcc_FLAGS += $(JANSSON_INCLUDES) --ptxas-options="-v" @@ -176,15 +175,15 @@ JHA/cuda_jha_compactionTest.o: JHA/cuda_jha_compactionTest.cu # This object does not use cuda device code but call the different kernels (autotune) scrypt/salsa_kernel.o: scrypt/salsa_kernel.cu - $(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_20,code=\"sm_21,compute_20\" -o $@ -c $< + $(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o $@ -c $< # These kernels are for older devices (SM) scrypt/test_kernel.o: scrypt/test_kernel.cu - $(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_20,code=\"sm_20,compute_20\" -o $@ -c $< + $(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o $@ -c $< scrypt/fermi_kernel.o: scrypt/fermi_kernel.cu - $(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_20,code=\"sm_21,compute_20\" -o $@ -c $< + $(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o $@ -c $< scrypt/kepler_kernel.o: scrypt/kepler_kernel.cu $(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o $@ -c $< diff --git a/README.txt b/README.txt index 14d27ad..23f027a 100644 --- a/README.txt +++ b/README.txt @@ -89,7 +89,6 @@ its command line interface and options. fresh use to mine Freshcoin fugue256 use to mine Fuguecoin groestl use to mine Groestlcoin - heavy use to mine Heavycoin hsr use to mine Hshare jackpot use to mine Sweepcoin keccak use to mine Maxcoin @@ -99,7 +98,6 @@ its command line interface and options. lyra2 use to mine CryptoCoin lyra2v2 use to mine Vertcoin lyra2z use to mine Zerocoin (XZC) - mjollnir use to mine Mjollnircoin myr-gr use to mine Myriad-Groest neoscrypt use to mine FeatherCoin nist5 use to mine TalkCoin @@ -142,8 +140,6 @@ its command line interface and options. --cuda-schedule Set device threads scheduling mode (default: auto) -f, --diff-factor Divide difficulty by this factor (default 1.0) -m, --diff-multiplier Multiply difficulty by this value (default 1.0) - --vote=VOTE block reward vote (for HeavyCoin) - --trust-pool trust the max block reward vote (maxvote) sent by the pool -o, --url=URL URL of mining server -O, --userpass=U:P username:password pair for mining server -u, --user=USERNAME username for mining server diff --git a/bench.cpp b/bench.cpp index dc670e7..baa999d 100644 --- a/bench.cpp +++ b/bench.cpp @@ -64,7 +64,9 @@ void algo_free_all(int thr_id) free_fresh(thr_id); free_fugue256(thr_id); free_groestlcoin(thr_id); +#ifdef WITH_HEAVY_ALGO free_heavy(thr_id); +#endif free_hmq17(thr_id); free_hsr(thr_id); free_jackpot(thr_id); @@ -125,6 +127,7 @@ bool bench_algo_switch_next(int thr_id) // skip some duplicated algos if (algo == ALGO_C11) algo++; // same as x11 if (algo == ALGO_DMD_GR) algo++; // same as groestl + if (algo == ALGO_HEAVY) algo++; // dead if (algo == ALGO_MJOLLNIR) algo++; // same as heavy if (algo == ALGO_KECCAKC) algo++; // same as keccak if (algo == ALGO_WHIRLCOIN) algo++; // same as whirlpool diff --git a/ccminer.cpp b/ccminer.cpp index e1d81f7..6c6d33f 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -251,9 +251,11 @@ Options:\n\ dmd-gr Diamond-Groestl\n\ fresh Freshcoin (shavite 80)\n\ fugue256 Fuguecoin\n\ - groestl Groestlcoin\n\ - heavy Heavycoin\n\ - hmq1725 Doubloons / Espers\n\ + groestl Groestlcoin\n" +#ifdef WITH_HEAVY_ALGO +" heavy Heavycoin\n" +#endif +" hmq1725 Doubloons / Espers\n\ jackpot JHA v8\n\ keccak Deprecated Keccak-256\n\ keccakc Keccak-256 (CreativeCoin)\n\ @@ -262,7 +264,6 @@ Options:\n\ lyra2 CryptoCoin\n\ lyra2v2 VertCoin\n\ lyra2z ZeroCoin (3rd impl)\n\ - mjollnir Mjollnircoin\n\ myr-gr Myriad-Groestl\n\ neoscrypt FeatherCoin, Phoenix, UFO...\n\ nist5 NIST5 (TalkCoin)\n\ @@ -304,8 +305,6 @@ Options:\n\ --cuda-schedule Set device threads scheduling mode (default: auto)\n\ -f, --diff-factor Divide difficulty by this factor (default 1.0) \n\ -m, --diff-multiplier Multiply difficulty by this value (default 1.0) \n\ - --vote=VOTE vote (for HeavyCoin)\n\ - --trust-pool trust the max block reward vote (maxvote) sent by the pool\n\ -o, --url=URL URL of mining server\n\ -O, --userpass=U:P username:password pair for mining server\n\ -u, --user=USERNAME username for mining server\n\ @@ -1555,10 +1554,12 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work) case ALGO_SIA: // getwork over stratum, no merkle to generate break; +#ifdef WITH_HEAVY_ALGO case ALGO_HEAVY: case ALGO_MJOLLNIR: heavycoin_hash(merkle_root, sctx->job.coinbase, (int)sctx->job.coinbase_size); break; +#endif case ALGO_FUGUE256: case ALGO_GROESTL: case ALGO_KECCAK: @@ -1573,9 +1574,11 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work) for (i = 0; i < sctx->job.merkle_count; i++) { memcpy(merkle_root + 32, sctx->job.merkle[i], 32); +#ifdef WITH_HEAVY_ALGO if (opt_algo == ALGO_HEAVY || opt_algo == ALGO_MJOLLNIR) heavycoin_hash(merkle_root, merkle_root, 64); else +#endif sha256d(merkle_root, merkle_root, 64); } @@ -2368,14 +2371,14 @@ static void *miner_thread(void *userdata) case ALGO_HSR: rc = scanhash_hsr(thr_id, &work, max_nonce, &hashes_done); break; - +#ifdef WITH_HEAVY_ALGO case ALGO_HEAVY: rc = scanhash_heavy(thr_id, &work, max_nonce, &hashes_done, work.maxvote, HEAVYCOIN_BLKHDR_SZ); break; case ALGO_MJOLLNIR: rc = scanhash_heavy(thr_id, &work, max_nonce, &hashes_done, 0, MNR_BLKHDR_SZ); break; - +#endif case ALGO_KECCAK: case ALGO_KECCAKC: rc = scanhash_keccak256(thr_id, &work, max_nonce, &hashes_done); diff --git a/ccminer.vcxproj b/ccminer.vcxproj index 21fa5fa..2a37505 100644 --- a/ccminer.vcxproj +++ b/ccminer.vcxproj @@ -39,7 +39,7 @@ - + @@ -155,7 +155,7 @@ 80 true true - compute_50,sm_50;compute_52,sm_52;compute_30,sm_30;compute_20,sm_21 + compute_50,sm_50;compute_52,sm_52;compute_30,sm_30 --ptxas-options="-O2" --Wno-deprecated-gpu-targets %(AdditionalOptions) O2 @@ -198,7 +198,7 @@ 80 true true - compute_52,sm_52;compute_50,sm_50;compute_35,sm_35;compute_30,sm_30;compute_20,sm_21 + compute_52,sm_52;compute_50,sm_50;compute_35,sm_35;compute_30,sm_30 $(NVTOOLSEXT_PATH)\include O3 64 @@ -410,10 +410,10 @@ - compute_20,sm_21 + compute_30,sm_30 - compute_20,sm_21 + compute_30,sm_30 compute_30,sm_30 @@ -425,7 +425,7 @@ compute_35,sm_35;compute_50,sm_50;compute_52,sm_52 - compute_20,sm_21 + compute_30,sm_30 compute_35,sm_35;compute_50,sm_50 @@ -606,7 +606,7 @@ - + diff --git a/heavy/heavy.cu b/heavy/heavy.cu index 4e52fa4..a5349c2 100644 --- a/heavy/heavy.cu +++ b/heavy/heavy.cu @@ -3,6 +3,21 @@ #include #include +#ifndef WITH_HEAVY_ALGO +#include +#include "miner.h" +// nonce array also used in other algos +uint32_t *heavy_nonceVector[MAX_GPUS]; +int scanhash_heavy(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done, uint32_t maxvote, int blocklen) +{ + applog(LOG_ERR, "heavy algo not included in this build!"); + sleep(3); + return -1; +} +void free_heavy(int thr_id) {} + +#else + // include thrust if possible #if defined(__GNUC__) && __GNUC__ == 5 && __GNUC_MINOR__ >= 2 && CUDA_VERSION < 7000 #warning "Heavy: incompatible GCC version!" @@ -17,16 +32,11 @@ #endif #include "miner.h" - -extern "C" { -#include "sph/sph_keccak.h" -#include "sph/sph_blake.h" -#include "sph/sph_groestl.h" -} -#include "hefty1.h" -#include "heavy/heavy.h" #include "cuda_helper.h" +// nonce array also used in other algos +uint32_t *heavy_nonceVector[MAX_GPUS]; + extern uint32_t *d_hash2output[MAX_GPUS]; extern uint32_t *d_hash3output[MAX_GPUS]; extern uint32_t *d_hash4output[MAX_GPUS]; @@ -35,35 +45,8 @@ extern uint32_t *d_hash5output[MAX_GPUS]; #define HEAVYCOIN_BLKHDR_SZ 84 #define MNR_BLKHDR_SZ 80 -// nonce-array für die threads -uint32_t *heavy_nonceVector[MAX_GPUS]; - extern uint32_t *heavy_heftyHashes[MAX_GPUS]; -/* Combines top 64-bits from each hash into a single hash */ -static void combine_hashes(uint32_t *out, const uint32_t *hash1, const uint32_t *hash2, const uint32_t *hash3, const uint32_t *hash4) -{ - const uint32_t *hash[4] = { hash1, hash2, hash3, hash4 }; - int bits; - unsigned int i; - uint32_t mask; - unsigned int k; - - /* Transpose first 64 bits of each hash into out */ - memset(out, 0, 32); - bits = 0; - for (i = 7; i >= 6; i--) { - for (mask = 0x80000000; mask; mask >>= 1) { - for (k = 0; k < 4; k++) { - out[(255 - bits)/32] <<= 1; - if ((hash[k][i] & mask) != 0) - out[(255 - bits)/32] |= 1; - bits++; - } - } - } -} - #ifdef _MSC_VER #include static uint32_t __inline bitsset( uint32_t x ) @@ -349,6 +332,42 @@ extern "C" void free_heavy(int thr_id) cudaDeviceSynchronize(); } +#endif + +extern "C" { +#include "sph/sph_keccak.h" +#include "sph/sph_blake.h" +#include "sph/sph_groestl.h" +} +#include "hefty1.h" +#include "heavy/heavy.h" + +/* Combines top 64-bits from each hash into a single hash */ +__host__ +static void combine_hashes(uint32_t *out, const uint32_t *hash1, const uint32_t *hash2, const uint32_t *hash3, const uint32_t *hash4) +{ + const uint32_t *hash[4] = { hash1, hash2, hash3, hash4 }; + int bits; + unsigned int i; + uint32_t mask; + unsigned int k; + + /* Transpose first 64 bits of each hash into out */ + memset(out, 0, 32); + bits = 0; + for (i = 7; i >= 6; i--) { + for (mask = 0x80000000; mask; mask >>= 1) { + for (k = 0; k < 4; k++) { + out[(255 - bits) / 32] <<= 1; + if ((hash[k][i] & mask) != 0) + out[(255 - bits) / 32] |= 1; + bits++; + } + } + } +} + +// CPU hash function __host__ void heavycoin_hash(uchar* output, const uchar* input, int len) {