diff --git a/Makefile.am b/Makefile.am index d086f91..f7c801c 100644 --- a/Makefile.am +++ b/Makefile.am @@ -55,7 +55,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \ x11/cuda_x11_luffa512_Cubehash.cu \ x13/x13.cu x13/cuda_x13_hamsi512.cu x13/cuda_x13_fugue512.cu \ x15/x14.cu x15/x15.cu x15/cuda_x14_shabal512.cu x15/cuda_x15_whirlpool.cu \ - x15/whirlpoolx.cu x15/cuda_whirlpoolx.cu \ + x15/whirlpool.cu x15/whirlpoolx.cu x15/cuda_whirlpoolx.cu \ x17/x17.cu x17/cuda_x17_haval512.cu x17/cuda_x17_sha512.cu \ x11/c11.cu x11/s3.cu diff --git a/ccminer.cpp b/ccminer.cpp index 42fb5a4..9736ed1 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -116,6 +116,7 @@ enum sha_algos { ALGO_X14, ALGO_X15, ALGO_X17, + ALGO_WHIRLCOIN, ALGO_WHIRLPOOLX, ALGO_ZR5, ALGO_COUNT @@ -154,6 +155,7 @@ static const char *algo_names[] = { "x14", "x15", "x17", + "whirlpool", "whirlpoolx", "zr5", "" @@ -315,6 +317,7 @@ Options:\n\ x14 X14\n\ x15 X15\n\ x17 X17\n\ + whirlpool Old Whirlcoin algo\n\ whirlpoolx WhirlpoolX (VNL)\n\ zr5 ZR5 (ZiftrCoin)\n\ -d, --devices Comma separated list of CUDA devices to use.\n\ @@ -1372,6 +1375,7 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work) case ALGO_GROESTL: case ALGO_KECCAK: case ALGO_BLAKECOIN: + case ALGO_WHIRLCOIN: SHA256((uchar*)sctx->job.coinbase, sctx->job.coinbase_size, (uchar*)merkle_root); break; default: @@ -1807,6 +1811,7 @@ static void *miner_thread(void *userdata) case ALGO_S3: case ALGO_X11: case ALGO_X13: + case ALGO_WHIRLCOIN: minmax = 0x400000; break; case ALGO_LYRA2: @@ -1938,6 +1943,9 @@ static void *miner_thread(void *userdata) case ALGO_S3: rc = scanhash_s3(thr_id, &work, max_nonce, &hashes_done); break; + case ALGO_WHIRLCOIN: + rc = scanhash_whirl(thr_id, &work, max_nonce, &hashes_done); + break; case ALGO_WHIRLPOOLX: rc = scanhash_whirlx(thr_id, &work, max_nonce, &hashes_done); break; @@ -2444,6 +2452,8 @@ void parse_arg(int key, char *arg) i = opt_algo = ALGO_LYRA2; else if (!strcasecmp("lyra2rev2", arg)) i = opt_algo = ALGO_LYRA2v2; + else if (!strcasecmp("whirl", arg)) + i = opt_algo = ALGO_WHIRLCOIN; else if (!strcasecmp("ziftr", arg)) i = opt_algo = ALGO_ZR5; else diff --git a/miner.h b/miner.h index db868fa..cd0f58f 100644 --- a/miner.h +++ b/miner.h @@ -286,6 +286,7 @@ extern int scanhash_qubit(int thr_id, struct work* work, uint32_t max_nonce, uns 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); +extern int scanhash_whirl(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_whirlx(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_x11(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_x13(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); @@ -325,6 +326,7 @@ extern void free_qubit(int thr_id); extern void free_skeincoin(int thr_id); extern void free_skein2(int thr_id); extern void free_s3(int thr_id); +extern void free_whirl(int thr_id); extern void free_whirlx(int thr_id); extern void free_x11(int thr_id); extern void free_x13(int thr_id); @@ -757,6 +759,7 @@ void scryptjane_hash(void* output, const void* input); void skeincoinhash(void *output, const void *input); void skein2hash(void *output, const void *input); void s3hash(void *output, const void *input); +void wcoinhash(void *state, const void *input); void whirlxHash(void *state, const void *input); void x11hash(void *output, const void *input); void x13hash(void *output, const void *input); diff --git a/util.cpp b/util.cpp index 38fc8a8..cd26e65 100644 --- a/util.cpp +++ b/util.cpp @@ -1948,6 +1948,9 @@ void print_hash_tests(void) s3hash(&hash[0], &buf[0]); printpfx("S3", hash); + wcoinhash(&hash[0], &buf[0]); + printpfx("whirlpool", hash); + whirlxHash(&hash[0], &buf[0]); printpfx("whirlpoolx", hash); diff --git a/x15/whirlpool.cu b/x15/whirlpool.cu new file mode 100644 index 0000000..cf4ded4 --- /dev/null +++ b/x15/whirlpool.cu @@ -0,0 +1,120 @@ +/* + * whirlpool routine (djm) + */ +extern "C" +{ +#include "sph/sph_whirlpool.h" +#include "miner.h" +} + +#include "cuda_helper.h" + +static uint32_t *d_hash[MAX_GPUS]; + +extern void x15_whirlpool_cpu_init(int thr_id, uint32_t threads, int mode); +extern void x15_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); + +extern void whirlpool512_setBlock_80(void *pdata, const void *ptarget); +extern void whirlpool512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash, int order); +extern uint32_t whirlpool512_cpu_finalhash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); + + +// CPU Hash function +extern "C" void wcoinhash(void *state, const void *input) +{ + sph_whirlpool_context ctx_whirlpool; + + unsigned char hash[128]; // uint32_t hashA[16], hashB[16]; + #define hashB hash+64 + + memset(hash, 0, sizeof hash); + + sph_whirlpool1_init(&ctx_whirlpool); + sph_whirlpool1(&ctx_whirlpool, input, 80); + sph_whirlpool1_close(&ctx_whirlpool, hash); + + sph_whirlpool1_init(&ctx_whirlpool); + sph_whirlpool1(&ctx_whirlpool, hash, 64); + sph_whirlpool1_close(&ctx_whirlpool, hashB); + + sph_whirlpool1_init(&ctx_whirlpool); + sph_whirlpool1(&ctx_whirlpool, hashB, 64); + sph_whirlpool1_close(&ctx_whirlpool, hash); + + sph_whirlpool1_init(&ctx_whirlpool); + sph_whirlpool1(&ctx_whirlpool, hash, 64); + sph_whirlpool1_close(&ctx_whirlpool, hash); + + memcpy(state, hash, 32); +} + +static bool init[MAX_GPUS] = { 0 }; + +extern "C" int scanhash_whirl(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done) +{ + uint32_t _ALIGN(128) endiandata[20]; + uint32_t* pdata = work->data; + uint32_t* ptarget = work->target; + const uint32_t first_nonce = pdata[19]; + + uint32_t throughput = device_intensity(thr_id, __func__, 1U << 19); // 19=256*256*8; + throughput = min(throughput, max_nonce - first_nonce); + + if (opt_benchmark) + ((uint32_t*)ptarget)[7] = 0x0000ff; + + if (!init[thr_id]) { + cudaSetDevice(device_map[thr_id]); + + cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput); + x15_whirlpool_cpu_init(thr_id, throughput, 1 /* old whirlpool */); + + init[thr_id] = true; + } + + for (int k=0; k < 20; k++) { + be32enc(&endiandata[k], pdata[k]); + } + + whirlpool512_setBlock_80((void*)endiandata, ptarget); + + do { + uint32_t foundNonce; + int order = 0; + + whirlpool512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); + x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + + foundNonce = whirlpool512_cpu_finalhash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + if (foundNonce != UINT32_MAX) + { + const uint32_t Htarg = ptarget[7]; + uint32_t vhash[8]; + be32enc(&endiandata[19], foundNonce); + wcoinhash(vhash, endiandata); + + if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) { + int res = 1; + *hashes_done = pdata[19] - first_nonce + throughput; + bn_store_hash_target_ratio(vhash, ptarget, work); + #if 0 + uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1); + if (secNonce != 0) { + pdata[21] = secNonce; + res++; + } + #endif + pdata[19] = foundNonce; + return res; + } else { + applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce); + } + } + pdata[19] += throughput; + + } while (pdata[19] < max_nonce && !work_restart[thr_id].restart); + + *hashes_done = pdata[19] - first_nonce; + return 0; +}