From 579a6ffe7f0ec2e460d25ef6c6b4e732228e707d Mon Sep 17 00:00:00 2001
From: David Phillips <david@sighup.nz>
Date: Tue, 24 Apr 2018 17:31:01 +1200
Subject: Remove magic constant for thread count

---
 sand-leek-cl.c | 12 ++++++------
 trampoline.c   |  6 +++---
 trampoline.h   |  2 ++
 3 files changed, 11 insertions(+), 9 deletions(-)

diff --git a/sand-leek-cl.c b/sand-leek-cl.c
index de512ec..ce293a9 100644
--- a/sand-leek-cl.c
+++ b/sand-leek-cl.c
@@ -1,9 +1,9 @@
-/* FIXME magic */
-/* 32768 kernels doing 32767 each, except if it's 0xxxxxxx */
-#define HASH_PER_RUN ((32768UL*32767UL) - (1<<24))
+/* SL_WORK_THREADS kernels doing 32767 exponents each, except if it's 00xxxxxx */
+#define HASH_PER_RUN ((SL_WORK_THREADS*32767UL) - (1<<24))
 #define EXPONENT_MAX 0x1FFFFFFFUL
 #define EXPONENT_SIZE_BYTES 4
 #define RSA_KEY_BITS 1024
+
 #include <stdio.h>
 #include <stdlib.h>
 #include <unistd.h>
@@ -197,7 +197,7 @@ int main(int argc, char **argv)
 	}
 	fprintf(stderr, "Done.\n");
 
-	/* FIXME */cl_int *buffer = malloc(sizeof(cl_int)*65536);
+	/* FIXME */cl_int *buffer = malloc(sizeof(cl_int)*SL_WORK_THREADS);
 
 	/* FIXME check for error */
 	bignum_e = BN_new();
@@ -278,7 +278,7 @@ int main(int argc, char **argv)
 			perror("host data buffer malloc");
 			return 1;
 		}
-		if (tramp_copy_data((void*)&buffer, sizeof(cl_int)*65536)) {
+		if (tramp_copy_data((void*)&buffer, sizeof(cl_int)*SL_WORK_THREADS)) {
 			fprintf(stderr, "Failed.\n");
 			return 1;
 		}
@@ -287,7 +287,7 @@ int main(int argc, char **argv)
 		/* FIXME */ int count = 0;
 		/* FIXME BUG: temporarily looping backwards to increase chance of using
 		 * something beginning with bit '1' as our exponent to highligt bug */
-		for (i = 0; i < 65536; i++) {
+		for (i = 0; i < SL_WORK_THREADS; i++) {
 			if (buffer[i] != 0) {
 				count++;
 				/* FIXME */uint16_t smalls = (unsigned int)buffer[i];
diff --git a/trampoline.c b/trampoline.c
index 8dd9040..28ae3fd 100644
--- a/trampoline.c
+++ b/trampoline.c
@@ -2,9 +2,9 @@
 #include <string.h>
 #include <opencl.h>
 
+#include "trampoline.h"
 #include "cl_error.h"
 #include "slurp.h"
-#include "sha1.h"
 
 static cl_platform_id platform;
 static cl_context context;
@@ -395,7 +395,7 @@ int tramp_run_kernel()
 	cl_event event;
 	cl_int ret = 0;
 	size_t workgroup_sizes[2];
-	workgroup_sizes[0] = 32768;
+	workgroup_sizes[0] = SL_WORK_THREADS;
 	workgroup_sizes[1] = 1;
 
 	ret = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, workgroup_sizes, NULL, 0, NULL, &event);
@@ -457,7 +457,7 @@ int tramp_copy_sha(struct sha_data *sha)
  *
  * Returns 0 on success, non-zero otherwise.
  */
-int tramp_copy_search(unsigned int search_raw[10])
+int tramp_copy_search(unsigned char search_raw[10])
 {
 	cl_event event;
 	cl_int ret = 0;
diff --git a/trampoline.h b/trampoline.h
index 5f4a76e..814f0db 100644
--- a/trampoline.h
+++ b/trampoline.h
@@ -1,5 +1,7 @@
 #include "sha1.h"
 
+#define SL_WORK_THREADS 32768
+
 int tramp_init(const char *preferred_platform);
 void tramp_destroy(void);
 int tramp_load_kernel(const char *filename);
-- 
cgit v1.1