restore the whirlpool algo

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
 

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

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);

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);
 

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;
+}