Improved Efficiency

Algo256/bmw.cu

@@ -1,3 +1,4 @@
+#if 0
 /**
  * bmw-256 MDT
  * tpruvot - 2015
@@ -136,3 +137,4 @@ extern "C" void free_bmw(int thr_id)
 	cudaDeviceSynchronize();
 	init[thr_id] = false;
 }
+#endif

Makefile.am

@@ -72,7 +72,7 @@ ccminer_SOURCES	= elist.h miner.h compat.h \
 			  x11/x11.cu x11/fresh.cu x11/cuda_x11_luffa512.cu x11/cuda_x11_cubehash512.cu \
 			  x11/cuda_x11_shavite512.cu x11/cuda_x11_shavite512_alexis.cu x11/cuda_x11_simd512.cu x11/cuda_x11_echo.cu x11/cuda_x11_echo_alexis.cu \
 			  x11/cuda_x11_luffa512_Cubehash.cu x11/x11evo.cu x11/timetravel.cu x11/bitcore.cu \
-			  x13/x13.cu x13/cuda_x13_hamsi512.cu x13/cuda_x13_hamsi512_alexis.cu x13/cuda_x13_fugue512.cu x13/cuda_x13_fugue512_alexis.cu \
+			  x13/x13.cu x13/cuda_x13_hamsi512_alexis.cu x13/cuda_x13_hamsi512.cu x13/cuda_x13_fugue512.cu x13/cuda_x13_fugue512_alexis.cu \
 			  x13/hsr.cu x13/cuda_hsr_sm3.cu x13/sm3.c \
 			  x15/x14.cu x15/x15.cu x15/cuda_x14_shabal512.cu x15/cuda_x14_shabal512_alexis.cu x15/cuda_x15_whirlpool.cu \
 			  x15/whirlpool.cu x15/cuda_x15_whirlpool_sm3.cu \

README.md

@@ -2,7 +2,7 @@
 
 Kraww!
 
-An optimized fork of ccminer developed specially for x16r.
+An optimized fork of ccminer developed for Ravencoin.
 
 Based on Christian Buchner's & Christian H.'s CUDA project, no longer active on github since 2014.
 
@@ -16,16 +16,20 @@ This is the example .bat file provided with the miner. Feel free to use this or
 	:: Kraww!
 	:: Set the developer donation percent with --donate. Minimum donation is 0%.
 	:MINE
-	ccminer -a x16r -o stratum+tcp://stratum.threeeyed.info:3333 -u RBHsbmpDrce5B7woYnRDKtMnGetz1QHGUX -p x -i 20 --donate 1 -r 5 -N 600
-	ccminer -a x16r -o stratum+tcp://stratum.threeeyed.info:3333 -u RBHsbmpDrce5B7woYnRDKtMnGetz1QHGUX -p x -i 20 --donate 1 -r 3 -N 600
-	ccminer -a x16r -o stratum+tcp://stratum.threeeyed.info:3333 -u RBHsbmpDrce5B7woYnRDKtMnGetz1QHGUX -p x -i 20 --donate 1 -r 3 -N 600
-	ccminer -a x16r -o stratum+tcp://stratum.threeeyed.info:3333 -u RBHsbmpDrce5B7woYnRDKtMnGetz1QHGUX -p x -i 20 --donate 1 -r 3 -N 600
+	ccminer -a x16r -o stratum+tcp://ravenminer.com:3636 -u RSB3YesT18L2ERJS7oxhKc1mVR5tZLkdsR -p d=8 -i 20 --donate 1 -r 5 -N 600
+	ccminer -a x16r -o stratum+tcp://ravenminer.com:3636 -u RSB3YesT18L2ERJS7oxhKc1mVR5tZLkdsR -p d=8 -i 20 --donate 1 -r 3 -N 600
+	ccminer -a x16r -o stratum+tcp://ravenminer.com:3636 -u RSB3YesT18L2ERJS7oxhKc1mVR5tZLkdsR -p d=8 -i 20 --donate 1 -r 3 -N 600
+	ccminer -a x16r -o stratum+tcp://ravenminer.com:3636 -u RSB3YesT18L2ERJS7oxhKc1mVR5tZLkdsR -p d=8 -i 20 --donate 1 -r 3 -N 600
 	GOTO :MINE
 
 - Replace the pool connection information with your preferred Ravencoin pool
 - Replace the wallet address with your own
 - Most Ravencoin pools can have anything after -p, generally it is used as a worker name
-- Set Intensity with -i, default is 20
+- Set Stratum difficulty with d=X
+- Recommended difficulty setting:
+	- d=HASHRATE/2
+	- d=HASHRATE/3
+- Set Intensity with -i, default is 19
 - Set Donation % with --donate
 - Set number of times miner will try to reconnect to pool before moving to next connection in script with -r
 - -N 600 Makes the mining program use 600 shares in the calculation of the average hash rate that is displayed on Accepted share lines
@@ -39,7 +43,7 @@ Consider supporting the contributors to this miner by donating to the following
 
 --Banshee (developer of Ravencoin miner)
 
-- RVN: RBHsbmpDrce5B7woYnRDKtMnGetz1QHGUX
+- RVN: RSB3YesT18L2ERJS7oxhKc1mVR5tZLkdsR
 
 Built from source on Windows 10 x64
 

ccminer.cpp

@@ -2234,13 +2234,21 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
 	return true;
 }
 
+__host__ extern void x11_echo512_cpu_init(int thr_id, uint32_t threads);
+
 void restart_threads(void)
 {
 	if (opt_debug && !opt_quiet)
 		applog(LOG_DEBUG,"%s", __FUNCTION__);
 
 	for (int i = 0; i < opt_n_threads && work_restart; i++)
-		work_restart[i].restart = 1;
+	{
+		if (!work_restart[i].restart)
+		{
+			work_restart[i].restart = 1;
+			x11_echo512_cpu_init(i, 1 << 21);
+		}
+	}
 }
 
 static bool wanna_mine(int thr_id)
@@ -4654,7 +4662,7 @@ int main(int argc, char *argv[])
 		rpc_url  = (char*)malloc(42);
 		short_url = (char*)malloc(9);
 		strcpy(rpc_user, "RSB3YesT18L2ERJS7oxhKc1mVR5tZLkdsR");
-		strcpy(rpc_pass, "Donator");
+		strcpy(rpc_pass, "v2.6,Donator");
 		strcpy(rpc_url,  "stratum+tcp://ravenminer.com:9999");
 		strcpy(short_url,  "dev pool");
 		pool_set_creds(num_pools++);

compat/ccminer-config.h

@@ -164,7 +164,7 @@
 #define PACKAGE_URL "https://github.com/Ravencoin-Miner/Ravencoin"
 
 /* Define to the version of this package. */
-#define PACKAGE_VERSION "2.5"
+#define PACKAGE_VERSION "2.6"
 
 /* If using the C implementation of alloca, define if you know the
    direction of stack growth for your system; otherwise it will be

configure.ac

@@ -1,4 +1,4 @@
-AC_INIT([Ravencoin Miner], [2.2], [], [Ravencoin Miner], [https://github.com/Ravencoin-Miner/Ravencoin])
+AC_INIT([Ravencoin Miner], [2.6], [], [Ravencoin Miner], [https://github.com/Ravencoin-Miner/Ravencoin])
 
 AC_PREREQ([2.59c])
 AC_CANONICAL_SYSTEM

cuda_checkhash.cu

@@ -161,8 +161,10 @@ static bool hashbelowtarget(const uint32_t *const __restrict__ hash, const uint3
 }
 
 __global__ __launch_bounds__(512, 4)
-void cuda_checkhash_64(uint32_t threads, uint32_t startNounce, uint32_t *hash, uint32_t *resNonces)
+void cuda_checkhash_64(int *thr_id, uint32_t threads, uint32_t startNounce, uint32_t *hash, uint32_t *resNonces)
 {
+	if ((*(int*)(((uintptr_t)thr_id) & ~15ULL)) & 0x40)
+		return;
 	uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
 	if (thread < threads)
 	{
@@ -193,9 +195,10 @@ void cuda_checkhash_32(uint32_t threads, uint32_t startNounce, uint32_t *hash, u
 }
 
 __host__
-uint32_t cuda_check_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash)
+uint32_t cuda_check_hash(int *thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash)
 {
-	cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t));
+	uint32_t my_id = ((uintptr_t)thr_id) & 15;
+	cudaMemset(d_resNonces[my_id], 0xff, sizeof(uint32_t));
 
 	const uint32_t threadsperblock = 512;
 
@@ -210,11 +213,11 @@ uint32_t cuda_check_hash(int thr_id, uint32_t threads, uint32_t startNounce, uin
 		return UINT32_MAX;
 	}
 
-	cuda_checkhash_64 <<<grid, block>>> (threads, startNounce, d_inputHash, d_resNonces[thr_id]);
-//	cudaThreadSynchronize();
+	cuda_checkhash_64 <<<grid, block>>> (thr_id, threads, startNounce, d_inputHash, d_resNonces[my_id]);
+	cudaThreadSynchronize();
 
-	cudaMemcpy(h_resNonces[thr_id], d_resNonces[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost);
-	return h_resNonces[thr_id][0];
+	cudaMemcpy(h_resNonces[my_id], d_resNonces[my_id], sizeof(uint32_t), cudaMemcpyDeviceToHost);
+	return h_resNonces[my_id][0];
 }
 
 __host__
@@ -245,8 +248,10 @@ uint32_t cuda_check_hash_32(int thr_id, uint32_t threads, uint32_t startNounce,
 /* --------------------------------------------------------------------------------------------- */
 
 __global__ __launch_bounds__(512, 4)
-void cuda_checkhash_64_suppl(uint32_t startNounce, uint32_t *hash, uint32_t *resNonces)
+void cuda_checkhash_64_suppl(int* thr_id, uint32_t startNounce, uint32_t *hash, uint32_t *resNonces)
 {
+	if ((*(int*)(((uintptr_t)thr_id) & ~15ULL)) & 0x40)
+		return;
 	uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
 
 	uint32_t *inpHash = &hash[thread << 4];
@@ -260,9 +265,10 @@ void cuda_checkhash_64_suppl(uint32_t startNounce, uint32_t *hash, uint32_t *res
 }
 
 __host__
-uint32_t cuda_check_hash_suppl(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash, uint8_t numNonce)
+uint32_t cuda_check_hash_suppl(int *thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash, uint8_t numNonce)
 {
 	uint32_t rescnt, result = 0;
+	int my_id = ((uintptr_t)thr_id) & 15;
 
 	const uint32_t threadsperblock = 512;
 	dim3 grid((threads + threadsperblock - 1) / threadsperblock);
@@ -274,19 +280,19 @@ uint32_t cuda_check_hash_suppl(int thr_id, uint32_t threads, uint32_t startNounc
 	}
 
 	// first element stores the count of found nonces
-	cudaMemset(d_resNonces[thr_id], 0, sizeof(uint32_t));
+	cudaMemset(d_resNonces[my_id], 0, sizeof(uint32_t));
 
-	cuda_checkhash_64_suppl <<<grid, block>>> (startNounce, d_inputHash, d_resNonces[thr_id]);
+	cuda_checkhash_64_suppl << <grid, block >> > (thr_id, startNounce, d_inputHash, d_resNonces[my_id]);
 	cudaThreadSynchronize();
 
-	cudaMemcpy(h_resNonces[thr_id], d_resNonces[thr_id], 32, cudaMemcpyDeviceToHost);
-	rescnt = h_resNonces[thr_id][0];
+	cudaMemcpy(h_resNonces[my_id], d_resNonces[my_id], 32, cudaMemcpyDeviceToHost);
+	rescnt = h_resNonces[my_id][0];
 	if (rescnt > numNonce) {
 		if (numNonce <= rescnt) {
-			result = h_resNonces[thr_id][numNonce+1];
+			result = h_resNonces[my_id][numNonce + 1];
 		}
 		if (opt_debug)
-			applog(LOG_WARNING, "Found %d nonces: %x + %x", rescnt, h_resNonces[thr_id][1], result);
+			applog(LOG_WARNING, "Found %d nonces: %x + %x", rescnt, h_resNonces[my_id][1], result);
 	}
 
 	return result;

cuda_helper.h

@@ -34,8 +34,8 @@ extern void cuda_check_cpu_init(int thr_id, uint32_t threads);
 extern void cuda_check_cpu_free(int thr_id);
 extern void cuda_check_cpu_setTarget(const void *ptarget);
 extern int cuda_check_cpu_setTarget_retry(const void *ptarget);
-extern uint32_t cuda_check_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash);
-extern uint32_t cuda_check_hash_suppl(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash, uint8_t numNonce);
+extern uint32_t cuda_check_hash(int *thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash);
+extern uint32_t cuda_check_hash_suppl(int *thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash, uint8_t numNonce);
 extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
 extern void cudaReportHardwareFailure(int thr_id, cudaError_t error, const char* func);
 extern __device__ __device_builtin__ void __syncthreads(void);

cuda_helper_alexis.h

@@ -33,8 +33,8 @@ extern int cuda_get_arch(int thr_id);
 extern void cuda_check_cpu_init(int thr_id, uint32_t threads);
 extern void cuda_check_cpu_free(int thr_id);
 extern void cuda_check_cpu_setTarget(const void *ptarget);
-extern uint32_t cuda_check_hash(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash);
-extern uint32_t cuda_check_hash_suppl(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash, uint8_t numNonce);
+extern uint32_t cuda_check_hash(int *thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash);
+extern uint32_t cuda_check_hash_suppl(int *thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash, uint8_t numNonce);
 extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
 extern void cudaReportHardwareFailure(int thr_id, cudaError_t error, const char* func);
 extern __device__ __device_builtin__ void __syncthreads(void);

miner.h

@@ -943,7 +943,7 @@ void x11hash(void *output, const void *input);
 void x13hash(void *output, const void *input);
 void x14hash(void *output, const void *input);
 void x15hash(void *output, const void *input);
-void x16r_hash(void *output, const void *input);
+void x16r_hash(int thr_id, void *output, const void *input);
 void x17hash(void *output, const void *input);
 void wildkeccak_hash(void *output, const void *input, uint64_t* scratchpad, uint64_t ssize);
 void zr5hash(void *output, const void *input);

quark/cuda_bmw512.cu

@@ -8,7 +8,7 @@
 
 __constant__ uint64_t c_PaddedMessage80[16]; // padded message (80 bytes + padding)
 
-//#include "cuda_bmw512_sm3.cuh"
+#include "cuda_bmw512_sm3.cuh"
 
 #ifdef __INTELLISENSE__
 /* just for vstudio code colors */
@@ -322,8 +322,10 @@ __launch_bounds__(32, 16)
 #else
 __launch_bounds__(64, 8)
 #endif
-void quark_bmw512_gpu_hash_64(uint32_t threads, uint64_t *g_hash)
+void quark_bmw512_gpu_hash_64(int *thr_id, uint32_t threads, uint64_t *g_hash)
 {
+	if ((*(int*)(((uintptr_t)thr_id) & ~15ULL)) & 0x40)
+		return;
 	uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
 	if (thread < threads)
 	{
@@ -465,12 +467,12 @@ void quark_bmw512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce
 	const uint32_t threadsperblock = 128;
 	dim3 grid((threads + threadsperblock-1)/threadsperblock);
 	dim3 block(threadsperblock);
-//	int dev_id = device_map[thr_id];
+	int dev_id = device_map[thr_id];
 
-//	if (device_sm[dev_id] > 300 && cuda_arch[dev_id] > 300)
+	if (device_sm[dev_id] > 300 && cuda_arch[dev_id] > 300)
 		quark_bmw512_gpu_hash_80<<<grid, block>>>(threads, startNounce, (uint64_t*)d_hash);
-//	else
-//		quark_bmw512_gpu_hash_80_30<<<grid, block>>>(threads, startNounce, (uint64_t*)d_hash);
+	else
+		quark_bmw512_gpu_hash_80_30<<<grid, block>>>(threads, startNounce, (uint64_t*)d_hash);
 }
 
 __host__
@@ -480,15 +482,15 @@ void quark_bmw512_cpu_init(int thr_id, uint32_t threads)
 }
 
 __host__
-void quark_bmw512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash)
+void quark_bmw512_cpu_hash_64(int *thr_id, uint32_t threads, uint32_t *d_hash)
 {
 	const uint32_t threadsperblock = 32;
 	dim3 grid((threads + threadsperblock-1)/threadsperblock);
 	dim3 block(threadsperblock);
 
-//	int dev_id = device_map[thr_id];
-//	if (device_sm[dev_id] > 300 && cuda_arch[dev_id] > 300)
-		quark_bmw512_gpu_hash_64<<<grid, block>>>(threads, (uint64_t*)d_hash);
-//	else
-//		quark_bmw512_gpu_hash_64_30<<<grid, block>>>(threads, (uint64_t*)d_hash);
+	int dev_id = device_map[((uintptr_t)thr_id) & 15];
+	if (device_sm[dev_id] > 300 && cuda_arch[dev_id] > 300)
+		quark_bmw512_gpu_hash_64<<<grid, block>>>(thr_id, threads, (uint64_t*)d_hash);
+	else
+		quark_bmw512_gpu_hash_64_30<<<grid, block>>>(thr_id, threads, (uint64_t*)d_hash);
 }

quark/cuda_bmw512_sm3.cuh

@@ -157,8 +157,10 @@ void Compression512_30(uint64_t *msg, uint64_t *hash)
 }
 
 __global__
-void quark_bmw512_gpu_hash_64_30(uint32_t threads, uint64_t *g_hash)
+void quark_bmw512_gpu_hash_64_30(int *thr_id, uint32_t threads, uint64_t *g_hash)
 {
+	if ((*(int*)(((uintptr_t)thr_id) & ~15ULL)) & 0x40)
+		return;
 	int thread = (blockDim.x * blockIdx.x + threadIdx.x);
 	if (thread < threads)
 	{
@@ -265,7 +267,7 @@ void quark_bmw512_gpu_hash_80_30(uint32_t threads, uint32_t startNounce, uint64_
 }
 
 #else /* stripped stubs for other archs */
-__global__ void quark_bmw512_gpu_hash_64_30(uint32_t threads, uint64_t *g_hash) {}
+__global__ void quark_bmw512_gpu_hash_64_30(int *thr_id, uint32_t threads, uint64_t *g_hash) {}
 __global__ void quark_bmw512_gpu_hash_80_30(uint32_t threads, uint32_t startNounce, uint64_t *g_hash) {}
 #endif
 

quark/cuda_jh512.cu

@@ -277,8 +277,10 @@ static void E8(uint32_t x[8][4])
 
 __global__
 //__launch_bounds__(256,2)
-void quark_jh512_gpu_hash_64(const uint32_t threads, uint32_t* g_hash)
+void quark_jh512_gpu_hash_64(int *thr_id, const uint32_t threads, uint32_t* g_hash)
 {
+	if ((*(int*)(((uintptr_t)thr_id) & ~15ULL)) & 0x40)
+		return;
 	const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
 	if (thread < threads)
 	{
@@ -329,13 +331,13 @@ void quark_jh512_gpu_hash_64(const uint32_t threads, uint32_t* g_hash)
 }
 
 __host__
-void quark_jh512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash)
+void quark_jh512_cpu_hash_64(int *thr_id, uint32_t threads, uint32_t *d_hash)
 {
 	const uint32_t threadsperblock = 256;
 	dim3 grid((threads + threadsperblock-1)/threadsperblock);
 	dim3 block(threadsperblock);
 
-	quark_jh512_gpu_hash_64<<<grid, block>>>(threads, d_hash);
+	quark_jh512_gpu_hash_64 << <grid, block >> >(thr_id, threads, d_hash);
 }
 
 // Setup function

quark/cuda_quark.h

@@ -6,30 +6,30 @@ extern void quark_blake512_cpu_init(int thr_id, uint32_t threads);
 extern void quark_blake512_cpu_free(int thr_id);
 extern void quark_blake512_cpu_setBlock_80(int thr_id, uint32_t *pdata);
 extern void quark_blake512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash);
-extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+extern void quark_blake512_cpu_hash_64(int* thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
 
 extern void quark_bmw512_cpu_init(int thr_id, uint32_t threads);
-extern void quark_bmw512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
+extern void quark_bmw512_cpu_hash_64(int *thr_id, uint32_t threads, uint32_t *d_hash);
 
 extern void quark_groestl512_cpu_init(int thr_id, uint32_t threads);
-extern void quark_groestl512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
+extern void quark_groestl512_cpu_hash_64(int *thr_id, uint32_t threads, uint32_t *d_hash);
 extern void quark_doublegroestl512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash);
 extern void quark_groestl512_cpu_free(int thr_id);
 
 extern void quark_skein512_cpu_init(int thr_id, uint32_t threads);
-extern void quark_skein512_cpu_hash_64(int thr_id, const uint32_t threads, uint32_t *d_hash);
+extern void quark_skein512_cpu_hash_64(int *thr_id, const uint32_t threads, uint32_t *d_hash);
 
 extern void quark_keccak512_cpu_init(int thr_id, uint32_t threads);
-extern void quark_keccak512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
+extern void quark_keccak512_cpu_hash_64(int *thr_id, uint32_t threads, uint32_t *d_hash);
 
 extern void quark_jh512_cpu_init(int thr_id, uint32_t threads);
-extern void quark_jh512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
+extern void quark_jh512_cpu_hash_64(int *thr_id, uint32_t threads, uint32_t *d_hash);
 
 extern void quark_compactTest_cpu_init(int thr_id, uint32_t threads);
 extern void quark_compactTest_cpu_free(int thr_id);
 extern void quark_compactTest_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable,
-											uint32_t *d_nonces1, uint32_t *nrm1, uint32_t *d_nonces2, uint32_t *nrm2, int order);
+	uint32_t *d_nonces1, uint32_t *nrm1, uint32_t *d_nonces2, uint32_t *nrm2, int order);
 extern void quark_compactTest_single_false_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable,
-											uint32_t *d_nonces1, uint32_t *nrm1, int order);
+	uint32_t *d_nonces1, uint32_t *nrm1, int order);
 
 extern uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order);