From 5f12943de5f76a2e5e8ad0b22f43d2b27392fb85 Mon Sep 17 00:00:00 2001
From: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date: Tue, 6 Oct 2015 23:49:40 +0200
Subject: [PATCH] whirlpool: add algo free function + vstudio

---
 Algo256/blake256.cu     |  3 ++-
 ccminer.cpp             |  1 +
 ccminer.vcxproj         |  1 +
 ccminer.vcxproj.filters |  3 +++
 x15/whirlpool.cu        | 18 ++++++++++++++++++
 5 files changed, 25 insertions(+), 1 deletion(-)

diff --git a/Algo256/blake256.cu b/Algo256/blake256.cu
index fd20be1..f335440 100644
--- a/Algo256/blake256.cu
+++ b/Algo256/blake256.cu
@@ -461,7 +461,8 @@ extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_non
 					if (vhashcpu[6] <= Htarg && fulltest(vhashcpu, ptarget)) {
 						pdata[21] = extra_results[0];
 						applog(LOG_BLUE, "1:%x 2:%x", foundNonce, extra_results[0]);
-						work_set_target_ratio(work, vhashcpu);
+						if (bn_hash_target_ratio(vhashcpu, ptarget) > work->shareratio)
+							work_set_target_ratio(work, vhashcpu);
 						rc = 2;
 					}
 					extra_results[0] = UINT32_MAX;
diff --git a/ccminer.cpp b/ccminer.cpp
index 4d2f23a..68e60c1 100644
--- a/ccminer.cpp
+++ b/ccminer.cpp
@@ -1543,6 +1543,7 @@ void miner_free_device(int thr_id)
 	free_skeincoin(thr_id);
 	free_skein2(thr_id);
 	free_s3(thr_id);
+	free_whirl(thr_id);
 	free_whirlx(thr_id);
 	free_x11(thr_id);
 	free_x13(thr_id);
diff --git a/ccminer.vcxproj b/ccminer.vcxproj
index 809a631..42dbab5 100644
--- a/ccminer.vcxproj
+++ b/ccminer.vcxproj
@@ -493,6 +493,7 @@
     <CudaCompile Include="x15\x15.cu" />
     <CudaCompile Include="x15\cuda_whirlpoolx.cu" />
     <CudaCompile Include="x15\whirlpoolx.cu" />
+    <CudaCompile Include="x15\whirlpool.cu" />
     <CudaCompile Include="x17\x17.cu">
     </CudaCompile>
     <CudaCompile Include="x17\cuda_x17_haval512.cu">
diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters
index 8e77e8d..93679fd 100644
--- a/ccminer.vcxproj.filters
+++ b/ccminer.vcxproj.filters
@@ -556,6 +556,9 @@
     <CudaCompile Include="x15\whirlpoolx.cu">
       <Filter>Source Files\CUDA\x15</Filter>
     </CudaCompile>
+    <CudaCompile Include="x15\whirlpool.cu">
+      <Filter>Source Files\CUDA\x15</Filter>
+    </CudaCompile>
     <CudaCompile Include="x17\cuda_x17_haval512.cu">
       <Filter>Source Files\CUDA\x17</Filter>
     </CudaCompile>
diff --git a/x15/whirlpool.cu b/x15/whirlpool.cu
index ac2eb1b..ad49f7b 100644
--- a/x15/whirlpool.cu
+++ b/x15/whirlpool.cu
@@ -13,6 +13,7 @@ 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 x15_whirlpool_cpu_free(int thr_id);
 
 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);
@@ -118,3 +119,20 @@ extern "C" int scanhash_whirl(int thr_id, struct work* work, uint32_t max_nonce,
 	*hashes_done = pdata[19] - first_nonce;
 	return 0;
 }
+
+// cleanup
+extern "C" void free_whirl(int thr_id)
+{
+	if (!init[thr_id])
+		return;
+
+	cudaSetDevice(device_map[thr_id]);
+
+	cudaFree(d_hash[thr_id]);
+
+	x15_whirlpool_cpu_free(thr_id);
+	init[thr_id] = false;
+
+	cudaDeviceSynchronize();
+}
+