benchmark: enhance the mem leak detection

reduce "false" warnings, and ignore unrelated/small ones <= 1 MB

On windows the gpu memory can be allocated by other processes

+ some cleanup in algos... (free/gpulog)

Algo256/blake256.cu

@@ -474,7 +474,7 @@ extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_non
 			else if (vhashcpu[7] > ptarget[7] && opt_debug) {
 				applog_hash((uchar*)ptarget);
 				applog_compare_hash((uchar*)vhashcpu, (uchar*)ptarget);
-				gpulog(LOG_WARNING, thr_id, "result for nonce %08x does not validate on CPU!", foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 

Algo256/bmw.cu

@@ -87,7 +87,7 @@ extern "C" int scanhash_bmw(int thr_id, struct work* work, uint32_t max_nonce, u
 				return 1;
 			}
 			else {
-				gpulog(LOG_WARNING, thr_id, "result for nonce %08x does not validate on CPU!", foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 

Algo256/keccak256.cu

@@ -80,7 +80,7 @@ extern "C" int scanhash_keccak256(int thr_id, struct work* work, uint32_t max_no
 				return 1;
 			}
 			else {
-				gpulog(LOG_WARNING, thr_id, "result for nonce %08x does not validate on CPU!", foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 

JHA/jackpotcoin.cu

@@ -247,7 +247,7 @@ extern "C" int scanhash_jackpot(int thr_id, struct work *work, uint32_t max_nonc
 				pdata[19] = foundNonce;
 				return res;
 			} else {
-				gpulog(LOG_WARNING, thr_id, "result for nonce %08x does not validate on CPU!", foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 

bench.cpp

@@ -4,11 +4,11 @@
  * 2015 - tpruvot@github
  */
 
+#include <unistd.h>
+
 #include "miner.h"
 #include "algos.h"
 
-#include <unistd.h>
-
 int bench_algo = -1;
 
 static double algo_hashrates[MAX_GPUS][ALGO_COUNT] = { 0 };
@@ -120,7 +120,13 @@ bool bench_algo_switch_next(int thr_id)
 	// free current algo memory and track mem usage
 	mused = cuda_available_memory(thr_id);
 	algo_free_all(thr_id);
+
+	// device can take some time to free
 	mfree = cuda_available_memory(thr_id);
+	if (device_mem_free[thr_id] > mfree) {
+		sleep(1);
+		mfree = cuda_available_memory(thr_id);
+	}
 
 	// we need to wait completion on all cards before the switch
 	if (opt_n_threads > 1) {
@@ -132,10 +138,15 @@ bool bench_algo_switch_next(int thr_id)
 	format_hashrate(hashrate, rate);
 	gpulog(LOG_NOTICE, thr_id, "%s hashrate = %s", algo_names[prev_algo], rate);
 
-	// check if there is memory leak
+	// ensure memory leak is still real after the barrier
 	if (device_mem_free[thr_id] > mfree) {
-		gpulog(LOG_WARNING, thr_id, "memory leak detected in %s ! %d MB free",
-			algo_names[prev_algo], mfree);
+		mfree = cuda_available_memory(thr_id);
+	}
+
+	// check if there is memory leak
+	if (device_mem_free[thr_id] - mfree > 1) {
+		gpulog(LOG_WARNING, thr_id, "possible %d MB memory leak in %s! %d MB free",
+			(device_mem_free[thr_id] - mfree), algo_names[prev_algo], mfree);
 		cuda_reset_device(thr_id, NULL); // force to free the leak
 		mfree = cuda_available_memory(thr_id);
 	}

cuda.cpp

@@ -4,10 +4,6 @@
 #include <unistd.h>
 #include <map>
 
-#ifndef _WIN32
-#include <unistd.h>
-#endif
-
 // include thrust
 #ifndef __cplusplus
 #include <thrust/version.h>

cuda_groestlcoin.cu

@@ -99,7 +99,7 @@ void groestlcoin_cpu_init(int thr_id, uint32_t threads)
 	// to check if the binary supports SM3+
 	cuda_get_arch(thr_id);
 
-	cudaMalloc(&d_resultNonce[thr_id], sizeof(uint32_t));
+	CUDA_SAFE_CALL(cudaMalloc(&d_resultNonce[thr_id], sizeof(uint32_t)));
 }
 
 __host__

cuda_nist5.cu

@@ -148,7 +148,7 @@ extern "C" int scanhash_nist5(int thr_id, struct work *work, uint32_t max_nonce,
 				goto out;
 			}
 			else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 
@@ -171,7 +171,7 @@ extern "C" void free_nist5(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 

fuguecoin.cpp

@@ -82,8 +82,7 @@ int scanhash_fugue256(int thr_id, struct work* work, uint32_t max_nonce, unsigne
 				*hashes_done = foundNounce - start_nonce + 1;
 				return 1;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!",
-					device_map[thr_id], foundNounce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce);
 			}
 		}
 
@@ -106,7 +105,7 @@ void free_fugue256(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	fugue256_cpu_free(thr_id);
 

groestlcoin.cpp

@@ -36,14 +36,15 @@ int scanhash_groestlcoin(int thr_id, struct work *work, uint32_t max_nonce, unsi
 	uint32_t throughput = cuda_default_throughput(thr_id, 1 << 19); // 256*256*8
 	if (init[thr_id]) throughput = min(throughput, max_nonce - start_nonce);
 
-	uint32_t *outputHash = (uint32_t*)malloc(throughput * 64);
+	uint32_t *outputHash = (uint32_t*)malloc((size_t) 64* throughput);
 
 	if (opt_benchmark)
-		((uint32_t*)ptarget)[7] = 0x000000ff;
+		ptarget[7] = 0x000ff;
 
 	if (!init[thr_id])
 	{
 		cudaSetDevice(device_map[thr_id]);
+		CUDA_LOG_ERROR();
 		groestlcoin_cpu_init(thr_id, throughput);
 		init[thr_id] = true;
 	}
@@ -73,8 +74,7 @@ int scanhash_groestlcoin(int thr_id, struct work *work, uint32_t max_nonce, unsi
 				free(outputHash);
 				return true;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!",
-					device_map[thr_id], foundNounce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce);
 			}
 		}
 
@@ -97,7 +97,7 @@ void free_groestlcoin(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	groestlcoin_cpu_free(thr_id);
 	init[thr_id] = false;

heavy/heavy.cu

@@ -275,8 +275,8 @@ int scanhash_heavy(int thr_id, struct work *work, uint32_t max_nonce, unsigned l
                     uint32_t vhash[8];
                     pdata[19] += nonce - pdata[19];
                     heavycoin_hash((uchar*)vhash, (uchar*)pdata, blocklen);
-                    if (memcmp(vhash, foundhash, 8*sizeof(uint32_t))) {
-                        applog(LOG_ERR, "hash for nonce %08x does not validate on CPU!\n", nonce);
+                    if (memcmp(vhash, foundhash, 32)) {
+                        gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", nonce);
                     } else {
                         *hashes_done = pdata[19] - first_nonce;
                         work_set_target_ratio(work, vhash);
@@ -306,7 +306,7 @@ extern "C" void free_heavy(int thr_id)
     if (!init[thr_id])
         return;
 
-    cudaSetDevice(device_map[thr_id]);
+    cudaThreadSynchronize();
 
     cudaFree(heavy_nonceVector[thr_id]);
 

myriadgroestl.cpp

@@ -78,8 +78,7 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
 				free(outputHash);
 				return 1;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!",
-					device_map[thr_id], foundNounce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce);
 			}
 		}
 
@@ -102,7 +101,7 @@ void free_myriad(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	myriadgroestl_cpu_free(thr_id);
 	init[thr_id] = false;

neoscrypt/neoscrypt.cpp

@@ -75,7 +75,7 @@ int scanhash_neoscrypt(int thr_id, struct work* work, uint32_t max_nonce, unsign
 				pdata[19] = foundNonce;
 				return 1;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!", dev_id, foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 
@@ -93,7 +93,7 @@ void free_neoscrypt(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	neoscrypt_cpu_free(thr_id);
 	init[thr_id] = false;

pentablake.cu

@@ -430,7 +430,7 @@ extern "C" int scanhash_pentablake(int thr_id, struct work *work, uint32_t max_n
 				pdata[19] = foundNonce;
 				return rc;
 			} else {
-				gpulog(LOG_WARNING, thr_id, "result for nonce %08x does not validate on CPU!", foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 

qubit/deep.cu

@@ -67,8 +67,9 @@ extern "C" int scanhash_deep(int thr_id, struct work* work, uint32_t max_nonce,
 	if (!init[thr_id])
 	{
 		cudaSetDevice(device_map[thr_id]);
+		CUDA_LOG_ERROR();
 
-		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], throughput * 64));
+		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput));
 
 		qubit_luffa512_cpu_init(thr_id, throughput);
 		x11_cubehash512_cpu_init(thr_id, throughput);
@@ -117,7 +118,7 @@ extern "C" int scanhash_deep(int thr_id, struct work* work, uint32_t max_nonce,
 				return res;
 			}
 			else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 
@@ -135,7 +136,7 @@ extern "C" void free_deep(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 

qubit/luffa.cu

@@ -48,7 +48,7 @@ extern "C" int scanhash_luffa(int thr_id, struct work* work, uint32_t max_nonce,
 		CUDA_LOG_ERROR();
 		//if (opt_cudaschedule == -1) // to reduce cpu usage...
 		//	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
-		CUDA_LOG_ERROR();
+		//CUDA_LOG_ERROR();
 
 		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput));
 
@@ -82,7 +82,7 @@ extern "C" int scanhash_luffa(int thr_id, struct work* work, uint32_t max_nonce,
 				pdata[19] = foundNonce;
 				return 1;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 
@@ -111,6 +111,6 @@ extern "C" void free_luffa(int thr_id)
 
 	cuda_check_cpu_free(thr_id);
 
-	cudaDeviceSynchronize();
 	init[thr_id] = false;
+	cudaDeviceSynchronize();
 }

qubit/qubit.cu

@@ -160,7 +160,7 @@ extern "C" void free_qubit(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 
@@ -170,4 +170,4 @@ extern "C" void free_qubit(int thr_id)
 	init[thr_id] = false;
 
 	cudaDeviceSynchronize();
-}
+}

skein.cu

@@ -449,7 +449,7 @@ extern "C" int scanhash_skeincoin(int thr_id, struct work* work, uint32_t max_no
 				return res;
 			}
 			else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 
@@ -473,7 +473,7 @@ extern "C" void free_skeincoin(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	if (sm5)
 		skeincoin_free(thr_id);

skein2.cpp

@@ -106,7 +106,7 @@ int scanhash_skein2(int thr_id, struct work* work, uint32_t max_nonce, unsigned
 				pdata[19] = swab32(foundNonce);
 				return res;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce %08x does not validate on CPU!", dev_id, foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 
@@ -129,7 +129,7 @@ void free_skein2(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 

x11/c11.cu

@@ -235,7 +235,7 @@ extern "C" int scanhash_c11(int thr_id, struct work* work, uint32_t max_nonce, u
 				pdata[19] = foundNonce;
 				return res;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 				pdata[19] = foundNonce + 1;
 			}
 		}
@@ -254,7 +254,7 @@ extern "C" void free_c11(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 	quark_groestl512_cpu_free(thr_id);

x11/fresh.cu

@@ -143,7 +143,7 @@ extern "C" int scanhash_fresh(int thr_id, struct work* work, uint32_t max_nonce,
 				pdata[19] = foundNonce;
 				return res;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 

x11/s3.cu

@@ -126,7 +126,7 @@ extern "C" int scanhash_s3(int thr_id, struct work* work, uint32_t max_nonce, un
 				return res;
 
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for nonce $%08X does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 
@@ -144,7 +144,7 @@ extern "C" void free_s3(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 	x11_simd512_cpu_free(thr_id);
@@ -153,4 +153,4 @@ extern "C" void free_s3(int thr_id)
 	init[thr_id] = false;
 
 	cudaDeviceSynchronize();
-}
+}

x11/x11.cu

@@ -233,7 +233,7 @@ extern "C" int scanhash_x11(int thr_id, struct work* work, uint32_t max_nonce, u
 				pdata[19] = foundNonce;
 				return res;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 				pdata[19] = foundNonce + 1;
 			}
 		}
@@ -252,7 +252,7 @@ extern "C" void free_x11(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 
@@ -263,4 +263,4 @@ extern "C" void free_x11(int thr_id)
 	init[thr_id] = false;
 
 	cudaDeviceSynchronize();
-}
+}

x13/x13.cu

@@ -234,7 +234,7 @@ extern "C" int scanhash_x13(int thr_id, struct work* work, uint32_t max_nonce, u
 				}
 				return res;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 
@@ -272,4 +272,4 @@ extern "C" void free_x13(int thr_id)
 
 	cudaDeviceSynchronize();
 	init[thr_id] = false;
-}
+}

x15/whirlpool.cu

@@ -126,7 +126,7 @@ extern "C" void free_whirl(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 

x15/whirlpoolx.cu

@@ -107,7 +107,7 @@ extern "C" void free_whirlx(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);
 

x15/x14.cu

@@ -248,7 +248,7 @@ extern "C" int scanhash_x14(int thr_id,  struct work* work, uint32_t max_nonce,
 				pdata[19] = foundNonce;
 				return res;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 		pdata[19] += throughput;

x15/x15.cu

@@ -254,7 +254,7 @@ extern "C" int scanhash_x15(int thr_id,  struct work* work, uint32_t max_nonce,
 				pdata[19] = foundNonce;
 				return res;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 

x17/x17.cu

@@ -279,7 +279,7 @@ extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, u
 				pdata[19] = foundNonce;
 				return res;
 			} else {
-				applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce);
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
 			}
 		}
 

zr5.cu

@@ -481,7 +481,7 @@ extern "C" void free_zr5(int thr_id)
 	if (!init[thr_id])
 		return;
 
-	cudaSetDevice(device_map[thr_id]);
+	cudaThreadSynchronize();
 
 	cudaFree(d_hash[thr_id]);