Initial commit

Algo256/bmw.cu

@@ -0,0 +1,138 @@
+/**
+ * bmw-256 MDT
+ * tpruvot - 2015
+ */
+extern "C" {
+#include "sph/sph_bmw.h"
+}
+
+#include <miner.h>
+#include <cuda_helper.h>
+
+static uint32_t *d_hash[MAX_GPUS];
+
+extern void bmw256_midstate_init(int thr_id, uint32_t threads);
+extern void bmw256_midstate_free(int thr_id);
+extern void bmw256_setBlock_80(int thr_id, void *pdata);
+extern void bmw256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_outputHash, int swap);
+
+extern uint32_t cuda_check_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_inputHash);
+
+// CPU Hash
+extern "C" void bmw_hash(void *state, const void *input)
+{
+	uint32_t _ALIGN(64) hash[16];
+	sph_bmw256_context ctx;
+
+	sph_bmw256_init(&ctx);
+	sph_bmw256(&ctx, input, 80);
+	sph_bmw256_close(&ctx, (void*) hash);
+
+	memcpy(state, hash, 32);
+}
+
+static bool init[MAX_GPUS] = { 0 };
+
+extern "C" int scanhash_bmw(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done)
+{
+	uint32_t _ALIGN(64) endiandata[20];
+	uint32_t *pdata = work->data;
+	uint32_t *ptarget = work->target;
+	const uint32_t first_nonce = pdata[19];
+	uint32_t throughput = cuda_default_throughput(thr_id, 1U << 21);
+	if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
+
+	if (opt_benchmark)
+		ptarget[7] = 0x0005;
+
+	if (!init[thr_id]) {
+		cudaSetDevice(device_map[thr_id]);
+		if (opt_cudaschedule == -1 && gpu_threads == 1) {
+			cudaDeviceReset();
+                        // reduce cpu usage
+			cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
+		}
+		gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);
+
+		cuda_check_cpu_init(thr_id, throughput);
+		bmw256_midstate_init(thr_id, throughput);
+
+		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t)32 * throughput));
+
+		init[thr_id] = true;
+	}
+
+	for (int k=0; k < 20; k++) {
+		be32enc(&endiandata[k], ((uint32_t*)pdata)[k]);
+	}
+
+	cudaGetLastError();
+	bmw256_setBlock_80(thr_id, (void*)endiandata);
+
+	cuda_check_cpu_setTarget(ptarget);
+
+	do {
+		bmw256_cpu_hash_80(thr_id, (int) throughput, pdata[19], d_hash[thr_id], 1);
+
+		*hashes_done = pdata[19] - first_nonce + throughput;
+
+		work->nonces[0] = cuda_check_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id]);
+		if (work->nonces[0] != UINT32_MAX)
+		{
+			const uint32_t Htarg = ptarget[7];
+			uint32_t _ALIGN(64) vhash[8];
+			be32enc(&endiandata[19], work->nonces[0]);
+			bmw_hash(vhash, endiandata);
+
+			if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) {
+				work->valid_nonces = 1;
+				work_set_target_ratio(work, vhash);
+				work->nonces[1] = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1);
+				if (work->nonces[1] != 0) {
+					be32enc(&endiandata[19], work->nonces[1]);
+					bmw_hash(vhash, endiandata);
+					bn_set_target_ratio(work, vhash, 1);
+					work->valid_nonces++;
+					pdata[19] = max(work->nonces[0], work->nonces[1]) + 1;
+				} else {
+					pdata[19] = work->nonces[0] + 1; // cursor
+				}
+				return work->valid_nonces;
+			}
+			else if (vhash[7] > Htarg) {
+				gpu_increment_reject(thr_id);
+				if (!opt_quiet)
+				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", work->nonces[0]);
+				pdata[19] = work->nonces[0] + 1;
+				continue;
+			}
+		}
+
+		if ((uint64_t) throughput + pdata[19] >= max_nonce) {
+			pdata[19] = max_nonce;
+			break;
+		}
+
+		pdata[19] += throughput;
+
+	} while (!work_restart[thr_id].restart);
+
+	*hashes_done = pdata[19] - first_nonce;
+	return 0;
+}
+
+// cleanup
+extern "C" void free_bmw(int thr_id)
+{
+	if (!init[thr_id])
+		return;
+
+	cudaThreadSynchronize();
+
+	cudaFree(d_hash[thr_id]);
+	bmw256_midstate_free(thr_id);
+	cuda_check_cpu_free(thr_id);
+
+	cudaDeviceSynchronize();
+	init[thr_id] = false;
+}