Bitcoin

Algo256/blake256.cu

@@ -339,13 +339,10 @@ static uint32_t blake256_cpu_hash_16(const int thr_id, const uint32_t threads, c
 		return result;
 
 	blake256_gpu_hash_16 <<<grid, block>>> (threads, startNonce, d_resNonce[thr_id], highTarget, (int) rounds, opt_tracegpu);
-	cudaDeviceSynchronize();
-	if (cudaSuccess == cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost)) {
-		//cudaDeviceSynchronize(); /* seems no more required */
-		result = h_resNonce[thr_id][0];
-		for (int n=0; n < (NBN-1); n++)
-			extra_results[n] = h_resNonce[thr_id][n+1];
-	}
+	CUDA_SAFE_CALL(cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost));
+	result = h_resNonce[thr_id][0];
+	for (int n=0; n < (NBN-1); n++)
+		extra_results[n] = h_resNonce[thr_id][n+1];
 	return result;
 }
 

Algo256/cuda_fugue256.cu

@@ -765,5 +765,5 @@ __host__ void fugue256_cpu_hash(int thr_id, uint32_t threads, int startNounce, v
 	fugue256_gpu_hash<<<grid, block>>>(thr_id, threads, startNounce, d_fugue256_hashoutput[thr_id], d_resultNonce[thr_id]);
 
 	//cudaMemcpy(outputHashes, d_fugue256_hashoutput[thr_id], 8 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
-	cudaMemcpy(nounce, d_resultNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost);
+	CUDA_SAFE_CALL(cudaMemcpy(nounce, d_resultNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost));
 }

Algo256/cuda_keccak256.cu

@@ -629,7 +629,7 @@ void keccak256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, u
 
 	keccak256_gpu_hash_80<<<grid, block>>>(threads, startNounce, d_outputHash, d_KNonce[thr_id]);
 	//MyStreamSynchronize(NULL, order, thr_id);
-	cudaMemcpy(h_nounce, d_KNonce[thr_id], 4 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
+	CUDA_SAFE_CALL(cudaMemcpy(h_nounce, d_KNonce[thr_id], 4 * sizeof(uint32_t), cudaMemcpyDeviceToHost));
 }
 
 __global__ __launch_bounds__(256,3)

Algo256/keccak256.cu

@@ -49,7 +49,7 @@ extern "C" int scanhash_keccak256(int thr_id, uint32_t *pdata,
 		((uint32_t*)ptarget)[7] = 0x0005;
 
 	if (!init[thr_id]) {
-		cudaSetDevice(device_map[thr_id]);
+		CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
 		cudaDeviceReset();
 		cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
 		cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);

README.txt

@@ -62,6 +62,7 @@ its command line interface and options.
 
   -a, --algo=ALGO       specify the algorithm to use
                           anime       use to mine Animecoin
+						  bitcoin     use to mine Bitcoin
                           blake       use to mine NEOS (Blake 256)
                           blakecoin   use to mine Old Blake 256
                           deep        use to mine Deepcoin

bitcoin.cu

@@ -112,13 +112,13 @@ void bitcoin_hash(uint32_t *output, const uint32_t *data, uint32_t nonce, const
 		b = a;
 		a = t1 + t2;
 	}
-	be32enc(&output[0], a + hc[0]);
-	be32enc(&output[1], b + hc[1]);
-	be32enc(&output[2], c + hc[2]);
-	be32enc(&output[3], d + hc[3]);
-	be32enc(&output[4], e + hc[4]);
-	be32enc(&output[5], f + hc[5]);
-	be32enc(&output[6], g + hc[6]);
+	be32enc(&output[0], a + hc[0]);
+	be32enc(&output[1], b + hc[1]);
+	be32enc(&output[2], c + hc[2]);
+	be32enc(&output[3], d + hc[3]);
+	be32enc(&output[4], e + hc[4]);
+	be32enc(&output[5], f + hc[5]);
+	be32enc(&output[6], g + hc[6]);
 	be32enc(&output[7], h + hc[7]);
 }
 
@@ -137,7 +137,7 @@ int scanhash_bitcoin(int thr_id, uint32_t *pdata,
 
 	if (!init[thr_id])
 	{
-		cudaSetDevice(device_map[thr_id]);
+		CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
 		cudaDeviceReset();
 		cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
 		cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);

bitslice_transformations_quad.cu

@@ -166,3 +166,35 @@ void from_bitslice_quad(const uint32_t *const __restrict__ input, uint32_t *cons
 	}
 */
 }
+__device__ __forceinline__
+void from_bitslice_quad_final(const uint32_t *const __restrict__ input, uint32_t *const __restrict__ output)
+{
+
+	uint32_t t;
+
+	output[0] = __byte_perm(input[0], input[4], 0x7531);
+	output[2] = __byte_perm(input[1], input[5], 0x7531);
+	output[8] = __byte_perm(input[2], input[6], 0x7531);
+	output[10] = __byte_perm(input[3], input[7], 0x7531);
+
+	SWAP1(output[0], output[2]);
+	SWAP1(output[8], output[10]);
+
+	SWAP2(output[2], output[10]);
+
+	output[6] = __byte_perm(output[2], output[10], 0x5410);
+	output[10] = __byte_perm(output[2], output[10], 0x7632);
+	output[2] = output[6];
+
+	SWAP4(output[2], output[10]);
+
+	output[6] = output[2];
+
+	if (threadIdx.x & 1)
+	{
+		output[6] = __byte_perm(output[2], 0, 0x3232);
+	}
+
+	output[6] = __byte_perm(output[6], __shfl((int)output[6], (threadIdx.x + 1) & 3, 4), 0x7632);
+	output[6 + 1] = __shfl((int)output[6], (threadIdx.x + 2) & 3, 4);
+}

ccminer.cpp

@@ -222,7 +222,7 @@ Usage: " PROGRAM_NAME " [OPTIONS]\n\
 Options:\n\
   -a, --algo=ALGO       specify the hash algorithm to use\n\
 			anime       Animecoin\n\
-			bitcoin     use to mine Bitcoin\n\
+			bitcoin     Bitcoin\n\
 			blake       Blake 256 (SFR/NEOS)\n\
 			blakecoin   Fast Blake 256 (8 rounds)\n\
 			deep        Deepcoin\n\

cuda.cpp

@@ -20,6 +20,8 @@
 
 #include "cuda_runtime.h"
 
+cudaDeviceProp device_props[MAX_GPUS];
+
 // CUDA Devices on the System
 int cuda_num_devices()
 {
@@ -32,7 +34,7 @@ int cuda_num_devices()
 	}
 
 	int maj = version / 1000, min = version % 100; // same as in deviceQuery sample
-	if (maj < 5 || (maj == 5 && min < 5))
+	if (maj < 5 || (maj == 6 && min < 5))
 	{
 		applog(LOG_ERR, "Driver does not support CUDA %d.%d API! Update your nVidia driver!", 5, 5);
 		exit(1);
@@ -64,6 +66,7 @@ void cuda_devicenames()
 		cudaDeviceProp props;
 		cudaGetDeviceProperties(&props, device_map[i]);
 
+		device_props[i] = props;
 		device_name[i] = strdup(props.name);
 		device_sm[i] = (props.major * 100 + props.minor * 10);
 	}

cuda_checkhash.cu

@@ -90,7 +90,7 @@ void cuda_checkhash_64(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)
 {
-	cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t));
+	CUDA_SAFE_CALL(cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t)));
 
 	const uint32_t threadsperblock = 512;
 
@@ -212,7 +212,7 @@ uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNoun
 __host__
 void cuda_check_quarkcoin(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order, uint32_t *resNonces)
 {
-	cudaMemset(d_resNonces[thr_id], 0xff, 2*sizeof(uint32_t));
+	CUDA_SAFE_CALL(cudaMemset(d_resNonces[thr_id], 0xff, 2*sizeof(uint32_t)));
 
 	const uint32_t threadsperblock = 256;
 

cuda_groestlcoin.cu

@@ -22,7 +22,7 @@ __global__ __launch_bounds__(256, 4)
 void groestlcoin_gpu_hash_quad(uint32_t threads, uint32_t startNounce, uint32_t *resNounce)
 {
     // durch 4 dividieren, weil jeweils 4 Threads zusammen ein Hash berechnen
-    uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x) / 4;
+    uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x) >> 2;
     if (thread < threads)
     {
         // GROESTL
@@ -38,52 +38,31 @@ void groestlcoin_gpu_hash_quad(uint32_t threads, uint32_t startNounce, uint32_t
         to_bitslice_quad(paddedInput, msgBitsliced);
 
         uint32_t state[8];
-        for (int round=0; round<2; round++)
-        {
-            groestl512_progressMessage_quad(state, msgBitsliced);
-
-            if (round < 1)
-            {
-                // Verkettung zweier Runden inclusive Padding.
-                msgBitsliced[ 0] = __byte_perm(state[ 0], 0x00800100, 0x4341 + ((threadIdx.x & 3)==3)*0x2000);
-                msgBitsliced[ 1] = __byte_perm(state[ 1], 0x00800100, 0x4341);
-                msgBitsliced[ 2] = __byte_perm(state[ 2], 0x00800100, 0x4341);
-                msgBitsliced[ 3] = __byte_perm(state[ 3], 0x00800100, 0x4341);
-                msgBitsliced[ 4] = __byte_perm(state[ 4], 0x00800100, 0x4341);
-                msgBitsliced[ 5] = __byte_perm(state[ 5], 0x00800100, 0x4341);
-                msgBitsliced[ 6] = __byte_perm(state[ 6], 0x00800100, 0x4341);
-				msgBitsliced[7] = __byte_perm(state[7], 0x00800100, 0x4341 + ((threadIdx.x & 3) == 0) * 0x0010);
-            }
-        }
-
-        // Nur der erste von jeweils 4 Threads bekommt das Ergebns-Hash
-        uint32_t out_state[16];
-        from_bitslice_quad(state, out_state);
+		groestl512_progressMessage_quad(state, msgBitsliced);
+
+		msgBitsliced[ 0] = __byte_perm(state[ 0], 0x00800100, 0x4341 + ((threadIdx.x & 3)==3)*0x2000);
+		msgBitsliced[ 1] = __byte_perm(state[ 1], 0x00800100, 0x4341);
+		msgBitsliced[ 2] = __byte_perm(state[ 2], 0x00800100, 0x4341);
+		msgBitsliced[ 3] = __byte_perm(state[ 3], 0x00800100, 0x4341);
+		msgBitsliced[ 4] = __byte_perm(state[ 4], 0x00800100, 0x4341);
+		msgBitsliced[ 5] = __byte_perm(state[ 5], 0x00800100, 0x4341);
+		msgBitsliced[ 6] = __byte_perm(state[ 6], 0x00800100, 0x4341);
+		msgBitsliced[7] = __byte_perm(state[7], 0x00800100, 0x4341 + ((threadIdx.x & 3) == 0) * 0x0010);
+
+		groestl512_progressMessage_quad(state, msgBitsliced);
+
+		uint32_t out_state[16];
+        from_bitslice_quad_final(state, out_state);
 
 		if ((threadIdx.x & 3) == 0)
         {
-            int i, position = -1;
-            bool rc = true;
-
-    #pragma unroll 8
-            for (i = 7; i >= 0; i--) {
-                if (out_state[i] > pTarget[i]) {
-                    if(position < i) {
-                        position = i;
-                        rc = false;
-                    }
-                 }
-                 if (out_state[i] < pTarget[i]) {
-                    if(position < i) {
-                        position = i;
-                        rc = true;
-                    }
-                 }
-            }
-
-            if(rc == true)
-                if(resNounce[0] > nounce)
-                    resNounce[0] = nounce;
+			if (out_state[7] <= pTarget[7])
+			{
+				uint32_t tmp = atomicExch(resNounce, nounce);
+				if (tmp != 0xffffffff)
+					resNounce[1] = tmp;
+			}
+
         }
     }
 }
@@ -92,9 +71,12 @@ void groestlcoin_gpu_hash_quad(uint32_t threads, uint32_t startNounce, uint32_t
 __host__ void groestlcoin_cpu_init(int thr_id, uint32_t threads)
 {
     CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
+	cudaDeviceReset();
+	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
+	cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
 
     // Speicher für Gewinner-Nonce belegen
-    cudaMalloc(&d_resultNonce[thr_id], sizeof(uint32_t)); 
+    cudaMalloc(&d_resultNonce[thr_id], 2 * sizeof(uint32_t)); 
 }
 
 __host__ void groestlcoin_cpu_setBlock(int thr_id, void *data, void *pTargetIn)
@@ -136,8 +118,8 @@ __host__ void groestlcoin_cpu_hash(int thr_id, uint32_t threads, uint32_t startN
     dim3 grid(factor*((threads + threadsperblock-1)/threadsperblock));
     dim3 block(threadsperblock);
 
-    cudaMemset(d_resultNonce[thr_id], 0xFF, sizeof(uint32_t));
+    cudaMemset(d_resultNonce[thr_id], 0xFF, 2 * sizeof(uint32_t));
     groestlcoin_gpu_hash_quad<<<grid, block>>>(threads, startNounce, d_resultNonce[thr_id]);
 
-    cudaMemcpy(nounce, d_resultNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost);
+    CUDA_SAFE_CALL(cudaMemcpy(nounce, d_resultNonce[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost));
 }

cuda_helper.h

@@ -518,8 +518,8 @@ __device__ __inline__ uint2 ROR2(const uint2 v, const int n)
 	}
 	else 
 	{
-		result.y = ((v.x >> (n - 32)) | (v.y << (32 - n)));
-		result.x = ((v.y >> (n - 32)) | (v.x << (32 - n)));
+		result.y = ((v.x >> (n - 32)) | (v.y << (64 - n)));
+		result.x = ((v.y >> (n - 32)) | (v.x << (64 - n)));
 	}
 	return result;
 }
@@ -627,10 +627,10 @@ uint2 SWAPDWORDS2(uint2 value)
 	return make_uint2(value.y, value.x);
 }
 
-static __forceinline__ __device__ uint2 SHL2(uint2 a, int offset)
+static __forceinline__ __device__ uint2 SHL2(const uint2 a, int offset)
 {
-#if __CUDA_ARCH__ > 300
 	uint2 result;
+#if __CUDA_ARCH__ > 300
 	if (offset<32) 
 	{
 		asm("{\n\t"
@@ -646,25 +646,24 @@ static __forceinline__ __device__ uint2 SHL2(uint2 a, int offset)
 			"}\n\t"
 			: "=r"(result.x), "=r"(result.y) : "r"(a.y), "r"(a.x), "r"(offset));
 	}
-	return result;
 #else
 	if (offset<=32) 
 	{
-		a.y = (a.y << offset) | (a.x >> (32 - offset));
-		a.x = (a.x << offset);
+		result.y = (a.y << offset) | (a.x >> (32 - offset));
+		result.x = (a.x << offset);
 	}
 	else
 	{
-		a.y = (a.x << (offset-32));
-		a.x = 0;
+		result.y = (a.x << (offset - 32));
+		result.x = 0;
 	}
-	return a;
 #endif
+	return result;
 }
-static __forceinline__ __device__ uint2 SHR2(uint2 a, int offset)
+static __forceinline__ __device__ uint2 SHR2(const uint2 a, int offset)
 {
-	#if __CUDA_ARCH__ > 300
 	uint2 result;
+#if __CUDA_ARCH__ > 300
 	if (offset<32) {
 		asm("{\n\t"
 			"shf.r.clamp.b32 %0,%2,%3,%4; \n\t"
@@ -679,20 +678,19 @@ static __forceinline__ __device__ uint2 SHR2(uint2 a, int offset)
 			"}\n\t"
 			: "=r"(result.x), "=r"(result.y) : "r"(a.y), "r"(a.x), "r"(offset));
 	}
-	return result;
 	#else
 	if (offset<=32) 
 	{
-		a.x = (a.x >> offset) | (a.y << (32 - offset));
-		a.y = (a.y >> offset);
+		result.x = (a.x >> offset) | (a.y << (32 - offset));
+		result.y = (a.y >> offset);
 	}
 	else
 	{
-		a.x = (a.y >> (offset - 32));
-		a.y = 0;
+		result.x = (a.y >> (offset - 32));
+		result.y = 0;
 	}
-	return a;
-	#endif
+#endif
+	return result;
 }
 
 static __device__ __forceinline__ uint64_t devectorizeswap(uint2 v) { return MAKE_ULONGLONG(cuda_swab32(v.y), cuda_swab32(v.x)); }

cuda_myriadgroestl.cu

@@ -269,7 +269,7 @@ __global__ void __launch_bounds__(256, 3)
 // Setup-Funktionen
 __host__ void myriadgroestl_cpu_init(int thr_id, uint32_t threads)
 {
-	cudaSetDevice(device_map[thr_id]);
+	CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
 	cudaDeviceReset();
 	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
 	cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
@@ -326,5 +326,5 @@ __host__ void myriadgroestl_cpu_hash(int thr_id, uint32_t threads, uint32_t star
     myriadgroestl_gpu_hash_quad2<<<grid2, block>>>(threads, startNounce, d_resultNonce[thr_id], d_outputHashes[thr_id]);
 
 
-    cudaMemcpy(nounce, d_resultNonce[thr_id], 4*sizeof(uint32_t), cudaMemcpyDeviceToHost);
+    CUDA_SAFE_CALL(cudaMemcpy(nounce, d_resultNonce[thr_id], 4*sizeof(uint32_t), cudaMemcpyDeviceToHost));
 }

cuda_nist5.cu

@@ -80,7 +80,7 @@ extern "C" int scanhash_nist5(int thr_id, uint32_t *pdata,
 
 	if (!init[thr_id])
 	{
-		cudaSetDevice(device_map[thr_id]);
+		CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
 		cudaDeviceReset();
 		cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
 		cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);