diff --git a/Makefile.am b/Makefile.am index d36f478..7b343c1 100644 --- a/Makefile.am +++ b/Makefile.am @@ -41,8 +41,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \ JHA/cuda_jha_compactionTest.cu cuda_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 \ + quark/quarkcoin.cu quark/cuda_quark_compactionTest.cu \ neoscrypt/neoscrypt.cpp neoscrypt/neoscrypt-cpu.c neoscrypt/cuda_neoscrypt.cu \ cuda_nist5.cu pentablake.cu skein.cu cuda_skeincoin.cu skein2.cpp zr5.cu \ sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c \ diff --git a/README.txt b/README.txt index eb0104e..27fffea 100644 --- a/README.txt +++ b/README.txt @@ -64,7 +64,6 @@ This code is based on the pooler cpuminer and inherits its command line interface and options. -a, --algo=ALGO specify the algorithm to use - anime use to mine Animecoin blake use to mine Saffroncoin (Blake 256) blakecoin use to mine Old Blake 256 c11/flax use to mine Chaincoin and Flax @@ -226,6 +225,7 @@ features. August 2015... Add Lyra2REv2 algo (Vertcoin/Zoom) Restore WhirlpoolX algo (VNL) + Drop animecoin support July 06th 2015 v1.6.5-C11 Nvml api power limits diff --git a/ccminer.cpp b/ccminer.cpp index 7816faa..429257c 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -84,7 +84,6 @@ struct workio_cmd { }; enum sha_algos { - ALGO_ANIME, ALGO_BLAKE, ALGO_BLAKECOIN, ALGO_C11, @@ -122,7 +121,6 @@ enum sha_algos { }; static const char *algo_names[] = { - "anime", "blake", "blakecoin", "c11", @@ -280,7 +278,6 @@ static char const usage[] = "\ Usage: " PROGRAM_NAME " [OPTIONS]\n\ Options:\n\ -a, --algo=ALGO specify the hash algorithm to use\n\ - anime Animecoin\n\ blake Blake 256 (SFR)\n\ blakecoin Fast Blake 256 (8 rounds)\n\ c11/flax X11 variant\n\ @@ -604,7 +601,6 @@ static void calc_network_diff(struct work *work) int shfb = 8 * (26 - (shift - 3)); switch (opt_algo) { - case ALGO_ANIME: case ALGO_QUARK: diffone = 0xFFFFFF0000000000ull; break; @@ -630,7 +626,6 @@ static void calc_network_diff(struct work *work) case ALGO_HEAVY: data64 = (uint64_t*)(rtarget + 2); break; - case ALGO_ANIME: case ALGO_QUARK: data64 = (uint64_t*)(rtarget + 3); break; @@ -1885,11 +1880,6 @@ static void *miner_thread(void *userdata) max_nonce, &hashes_done); break; - case ALGO_ANIME: - rc = scanhash_anime(thr_id, work.data, work.target, - max_nonce, &hashes_done); - break; - case ALGO_BLAKECOIN: rc = scanhash_blake256(thr_id, work.data, work.target, max_nonce, &hashes_done, 8); diff --git a/ccminer.vcxproj b/ccminer.vcxproj index 4f6584e..badbe7e 100644 --- a/ccminer.vcxproj +++ b/ccminer.vcxproj @@ -405,8 +405,6 @@ --ptxas-options="-dlcm=cg" %(AdditionalOptions) true - - 128 diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters index fd3cf83..bb4573d 100644 --- a/ccminer.vcxproj.filters +++ b/ccminer.vcxproj.filters @@ -466,9 +466,6 @@ Source Files\CUDA\heavy - - Source Files\CUDA\quark - Source Files\CUDA\quark diff --git a/miner.h b/miner.h index 835eb39..4c5822c 100644 --- a/miner.h +++ b/miner.h @@ -299,10 +299,6 @@ extern int scanhash_quark(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, unsigned long *hashes_done); -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_blake256(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, unsigned long *hashes_done, int8_t blakerounds); @@ -772,7 +768,6 @@ void applog_hash(unsigned char *hash); void applog_compare_hash(unsigned char *hash, unsigned char *hash2); void print_hash_tests(void); -void animehash(void *state, const void *input); void blake256hash(void *output, const void *input, int8_t rounds); void c11hash(void *output, const void *input); void deephash(void *state, const void *input); diff --git a/quark/animecoin.cu b/quark/animecoin.cu deleted file mode 100644 index 203a0c9..0000000 --- a/quark/animecoin.cu +++ /dev/null @@ -1,292 +0,0 @@ -extern "C" -{ -#include "sph/sph_blake.h" -#include "sph/sph_bmw.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 "cuda_helper.h" - -static uint32_t *d_hash[MAX_GPUS]; - -// Speicher zur Generierung der Noncevektoren für die bedingten Hashes -static uint32_t *d_branch1Nonces[MAX_GPUS]; -static uint32_t *d_branch2Nonces[MAX_GPUS]; -static uint32_t *d_branch3Nonces[MAX_GPUS]; - -extern void quark_blake512_cpu_init(int thr_id, uint32_t threads); -extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); - -extern void quark_bmw512_cpu_init(int thr_id, uint32_t threads); -extern void quark_bmw512_cpu_setBlock_80(void *pdata); -extern void quark_bmw512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_outputHash, int order); -extern void quark_bmw512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_outputHash, int order); - -extern void quark_groestl512_cpu_init(int thr_id, uint32_t threads); -extern void quark_groestl512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); -extern void quark_doublegroestl512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); - -extern void quark_skein512_cpu_init(int thr_id, uint32_t threads); -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 void quark_keccak512_cpu_init(int thr_id, uint32_t threads); -extern void quark_keccak512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); - -extern void quark_jh512_cpu_init(int thr_id, uint32_t threads); -extern void quark_jh512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); - -extern void quark_compactTest_cpu_init(int thr_id, uint32_t threads); -extern void quark_compactTest_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable, - uint32_t *d_nonces1, uint32_t *nrm1, - uint32_t *d_nonces2, uint32_t *nrm2, - int order); -extern void quark_compactTest_single_false_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable, - uint32_t *d_nonces1, uint32_t *nrm1, - int order); - -extern 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); - -/* CPU Hash */ -extern "C" void animehash(void *state, const void *input) -{ - sph_blake512_context ctx_blake; - sph_bmw512_context ctx_bmw; - 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_bmw512_init(&ctx_bmw); - // ZBMW; - sph_bmw512 (&ctx_bmw, (const void*) input, 80); - sph_bmw512_close(&ctx_bmw, (void*) hash); - - sph_blake512_init(&ctx_blake); - // ZBLAKE; - sph_blake512 (&ctx_blake, hash, 64); - sph_blake512_close(&ctx_blake, (void*) hash); - - if (hash[0] & 0x8) - { - sph_groestl512_init(&ctx_groestl); - // ZGROESTL; - sph_groestl512 (&ctx_groestl, (const void*) hash, 64); - sph_groestl512_close(&ctx_groestl, (void*) hash); - } - else - { - sph_skein512_init(&ctx_skein); - // ZSKEIN; - sph_skein512 (&ctx_skein, (const void*) hash, 64); - sph_skein512_close(&ctx_skein, (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); - - if (hash[0] & 0x8) - { - sph_blake512_init(&ctx_blake); - // ZBLAKE; - sph_blake512 (&ctx_blake, (const void*) hash, 64); - sph_blake512_close(&ctx_blake, (void*) hash); - } - else - { - sph_bmw512_init(&ctx_bmw); - // ZBMW; - sph_bmw512 (&ctx_bmw, (const void*) hash, 64); - sph_bmw512_close(&ctx_bmw, (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); - // SKEIN; - sph_skein512 (&ctx_skein, (const void*) hash, 64); - sph_skein512_close(&ctx_skein, (void*) hash); - - if (hash[0] & 0x8) - { - sph_keccak512_init(&ctx_keccak); - // ZKECCAK; - sph_keccak512 (&ctx_keccak, (const void*) hash, 64); - sph_keccak512_close(&ctx_keccak, (void*) hash); - } - else - { - sph_jh512_init(&ctx_jh); - // ZJH; - sph_jh512 (&ctx_jh, (const void*) hash, 64); - sph_jh512_close(&ctx_jh, (void*) hash); - } - - memcpy(state, hash, 32); -} - -/* -struct HashPredicate -{ - HashPredicate(uint32_t *hashes, uint32_t startNonce) : - m_hashes(hashes), - m_startNonce(startNonce) - { } - - __device__ - bool operator()(const uint32_t x) - { - uint32_t *hash = &m_hashes[(x - m_startNonce)*16]; - return hash[0] & 0x8; - } - - uint32_t *m_hashes; - uint32_t m_startNonce; -}; -*/ - -static bool init[MAX_GPUS] = { 0 }; - -extern "C" int scanhash_anime(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]; - uint32_t throughput = device_intensity(thr_id, __func__, 1 << 19); // 256*256*8 - throughput = min(throughput, max_nonce - first_nonce); - - if (opt_benchmark) - ((uint32_t*)ptarget)[7] = 0x00000f; - - if (!init[thr_id]) - { - cudaSetDevice(device_map[thr_id]); - - CUDA_SAFE_CALL(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_skein512_cpu_init(thr_id, throughput); - quark_bmw512_cpu_init(thr_id, throughput); - quark_keccak512_cpu_init(thr_id, throughput); - quark_jh512_cpu_init(thr_id, throughput); - cuda_check_cpu_init(thr_id, throughput); - quark_compactTest_cpu_init(thr_id, throughput); - - CUDA_SAFE_CALL(cudaMalloc(&d_branch1Nonces[thr_id], sizeof(uint32_t)*throughput)); - CUDA_SAFE_CALL(cudaMalloc(&d_branch2Nonces[thr_id], sizeof(uint32_t)*throughput)); - CUDA_SAFE_CALL(cudaMalloc(&d_branch3Nonces[thr_id], sizeof(uint32_t)*throughput)); - - init[thr_id] = true; - } - - uint32_t endiandata[20]; - for (int k=0; k < 20; k++) - be32enc(&endiandata[k], pdata[k]); - - quark_bmw512_cpu_setBlock_80((void*)endiandata); - cuda_check_cpu_setTarget(ptarget); - - do { - int order = 0; - uint32_t nrm1=0, nrm2=0, nrm3=0; - - // erstes BMW512 Hash mit CUDA - quark_bmw512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); - - // das ist der unbedingte Branch für Blake512 - quark_blake512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); - - quark_compactTest_single_false_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], NULL, - d_branch3Nonces[thr_id], &nrm3, - order++); - - // nur den Skein Branch weiterverfolgen - quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++); - - // das ist der unbedingte Branch für Groestl512 - quark_groestl512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++); - - // das ist der unbedingte Branch für JH512 - quark_jh512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++); - - // quarkNonces in branch1 und branch2 aufsplitten gemäss if (hash[0] & 0x8) - quark_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id], - d_branch1Nonces[thr_id], &nrm1, - d_branch2Nonces[thr_id], &nrm2, - order++); - - // das ist der bedingte Branch für Blake512 - quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id], order++); - - // das ist der bedingte Branch für Bmw512 - quark_bmw512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id], order++); - - // das ist der unbedingte Branch für Keccak512 - quark_keccak512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++); - - // das ist der unbedingte Branch für Skein512 - quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++); - - // quarkNonces in branch1 und branch2 aufsplitten gemäss if (hash[0] & 0x8) - quark_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id], - d_branch1Nonces[thr_id], &nrm1, - d_branch2Nonces[thr_id], &nrm2, - order++); - - // das ist der bedingte Branch für Keccak512 - quark_keccak512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id], order++); - - // das ist der bedingte Branch für JH512 - quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id], order++); - - // Scan nach Gewinner Hashes auf der GPU - uint32_t foundNonce = cuda_check_hash_branch(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++); - if (foundNonce != UINT32_MAX) - { - const uint32_t Htarg = ptarget[7]; - uint32_t vhash64[8]; - be32enc(&endiandata[19], foundNonce); - animehash(vhash64, endiandata); - - if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) { - int res = 1; - uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1); - *hashes_done = pdata[19] - first_nonce + throughput; - if (secNonce != 0) { - pdata[21] = secNonce; - res++; - } - pdata[19] = foundNonce; - return res; - } else { - 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 > (uint64_t)max_nonce) { - pdata[19] = max_nonce; - break; - } - - pdata[19] += throughput; - - } while (!work_restart[thr_id].restart); - - *hashes_done = pdata[19] - first_nonce + 1; - return 0; -} diff --git a/util.cpp b/util.cpp index 8efec8c..a4ddbf7 100644 --- a/util.cpp +++ b/util.cpp @@ -1806,9 +1806,6 @@ void print_hash_tests(void) printf(CL_WHT "CPU HASH ON EMPTY BUFFER RESULTS:" CL_N "\n"); - animehash(&hash[0], &buf[0]); - printpfx("anime", hash); - blake256hash(&hash[0], &buf[0], 8); printpfx("blakecoin", hash);