diff options
| author | David Phillips <david@sighup.nz> | 2018-04-24 17:31:01 +1200 |
|---|---|---|
| committer | David Phillips <david@sighup.nz> | 2018-04-24 17:31:01 +1200 |
| commit | 579a6ffe7f0ec2e460d25ef6c6b4e732228e707d (patch) | |
| tree | 683bbb2c6a16ce880924f0b928438c1195fa9873 | |
| parent | 133dd90a0de29b7c24217ca87af57ede4da1247b (diff) | |
| download | sand-leek-579a6ffe7f0ec2e460d25ef6c6b4e732228e707d.tar.xz | |
Remove magic constant for thread count
| -rw-r--r-- | sand-leek-cl.c | 12 | ||||
| -rw-r--r-- | trampoline.c | 6 | ||||
| -rw-r--r-- | 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); |