From 009b013d25aa1ee1912792083308c33712cbeb44 Mon Sep 17 00:00:00 2001
From: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date: Tue, 27 Sep 2016 00:53:56 +0200
Subject: [PATCH] nist5: rename and move source file

build tip: autoreconf && make -j
---
 Makefile.am                     |  3 ++-
 ccminer.vcxproj                 | 18 ++++++------------
 ccminer.vcxproj.filters         |  6 +++---
 cuda_nist5.cu => quark/nist5.cu | 12 +++++++++---
 4 files changed, 20 insertions(+), 19 deletions(-)
 rename cuda_nist5.cu => quark/nist5.cu (94%)

diff --git a/Makefile.am b/Makefile.am
index 73702d8..45a1695 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -43,9 +43,10 @@ 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/nist5.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 \
+			  pentablake.cu skein.cu cuda_skeincoin.cu skein2.cpp zr5.cu \
 			  sia.cu sph/blake2b.c \
 			  sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c \
 			  sph/cubehash.c sph/echo.c sph/luffa.c sph/sha2.c sph/shavite.c sph/simd.c \
diff --git a/ccminer.vcxproj b/ccminer.vcxproj
index d4243bd..506dbc3 100644
--- a/ccminer.vcxproj
+++ b/ccminer.vcxproj
@@ -357,8 +357,6 @@
     <CudaCompile Include="cuda.cpp" />
     <CudaCompile Include="cuda_groestlcoin.cu" />
     <CudaCompile Include="cuda_myriadgroestl.cu" />
-    <CudaCompile Include="cuda_nist5.cu">
-    </CudaCompile>
     <CudaCompile Include="scrypt\blake.cu" />
     <CudaCompile Include="scrypt\keccak.cu" />
     <CudaCompile Include="scrypt\sha256.cu" />
@@ -459,16 +457,12 @@
     </CudaCompile>
     <CudaCompile Include="quark\cuda_skein512.cu">
     </CudaCompile>
-    <CudaCompile Include="quark\quarkcoin.cu">
-    </CudaCompile>
-    <CudaCompile Include="qubit\deep.cu">
-    </CudaCompile>
-    <CudaCompile Include="qubit\luffa.cu">
-    </CudaCompile>
-    <CudaCompile Include="qubit\qubit.cu">
-    </CudaCompile>
-    <CudaCompile Include="qubit\qubit_luffa512.cu">
-    </CudaCompile>
+    <CudaCompile Include="quark\quarkcoin.cu" />
+    <CudaCompile Include="quark\nist5.cu" />
+    <CudaCompile Include="qubit\deep.cu" />
+    <CudaCompile Include="qubit\luffa.cu" />
+    <CudaCompile Include="qubit\qubit.cu" />
+    <CudaCompile Include="qubit\qubit_luffa512.cu" />
     <CudaCompile Include="lyra2\lyra2RE.cu" />
     <CudaCompile Include="lyra2\cuda_lyra2.cu" />
     <CudaCompile Include="lyra2\lyra2REv2.cu" />
diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters
index 6861319..ed15dfc 100644
--- a/ccminer.vcxproj.filters
+++ b/ccminer.vcxproj.filters
@@ -535,6 +535,9 @@
     <CudaCompile Include="quark\cuda_quark_keccak512.cu">
       <Filter>Source Files\CUDA\quark</Filter>
     </CudaCompile>
+    <CudaCompile Include="quark\nist5.cu">
+      <Filter>Source Files\CUDA\quark</Filter>
+    </CudaCompile>
     <CudaCompile Include="quark\quarkcoin.cu">
       <Filter>Source Files\CUDA\quark</Filter>
     </CudaCompile>
@@ -550,9 +553,6 @@
     <CudaCompile Include="qubit\qubit_luffa512.cu">
       <Filter>Source Files\CUDA\qubit</Filter>
     </CudaCompile>
-    <CudaCompile Include="cuda_nist5.cu">
-      <Filter>Source Files\CUDA</Filter>
-    </CudaCompile>
     <ClInclude Include="x11\cuda_x11_aes.cuh">
       <Filter>Source Files\CUDA\x11</Filter>
     </ClInclude>
diff --git a/cuda_nist5.cu b/quark/nist5.cu
similarity index 94%
rename from cuda_nist5.cu
rename to quark/nist5.cu
index 57a2a90..ddd4363 100644
--- a/cuda_nist5.cu
+++ b/quark/nist5.cu
@@ -22,13 +22,13 @@ extern "C" void nist5hash(void *state, const void *input)
     sph_jh512_context ctx_jh;
     sph_keccak512_context ctx_keccak;
     sph_skein512_context ctx_skein;
-    
+
     uint8_t hash[64];
 
     sph_blake512_init(&ctx_blake);
     sph_blake512 (&ctx_blake, input, 80);
     sph_blake512_close(&ctx_blake, (void*) hash);
-    
+
     sph_groestl512_init(&ctx_groestl);
     sph_groestl512 (&ctx_groestl, (const void*) hash, 64);
     sph_groestl512_close(&ctx_groestl, (void*) hash);
@@ -66,8 +66,14 @@ extern "C" int scanhash_nist5(int thr_id, struct work *work, uint32_t max_nonce,
 
 	if (!init[thr_id])
 	{
-		cudaDeviceSynchronize();
 		cudaSetDevice(device_map[thr_id]);
+		if (opt_cudaschedule == -1 && gpu_threads == 1) {
+			cudaDeviceReset();
+			// reduce cpu usage
+			cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
+			CUDA_LOG_ERROR();
+		}
+		gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);
 
 		// Constants copy/init (no device alloc in these algos)
 		quark_blake512_cpu_init(thr_id, throughput);