From a22edd2a7f42ff61a37daf6f334277dc1f2b8348 Mon Sep 17 00:00:00 2001 From: Con Kolivas Date: Sat, 21 Jul 2012 10:25:33 +1000 Subject: [PATCH] Test the target in the actual scrypt kernel itself saving further calculations. --- driver-opencl.c | 3 +++ scrypt120713.cl | 6 ++++-- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/driver-opencl.c b/driver-opencl.c index 3ec60dd3..cd2c9ab3 100644 --- a/driver-opencl.c +++ b/driver-opencl.c @@ -998,8 +998,10 @@ static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u char *midstate = blk->work->midstate; cl_kernel *kernel = &clState->kernel; unsigned int num = 0; + cl_uint le_target; cl_int status = 0; + le_target = ~swab32((uint32_t)blk->work->target[7]); clState->cldata = blk->work->data; status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL,NULL); @@ -1008,6 +1010,7 @@ static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u CL_SET_ARG(clState->padbuffer8); CL_SET_VARG(4, &midstate[0]); CL_SET_VARG(4, &midstate[16]); + CL_SET_ARG(le_target); return status; } diff --git a/scrypt120713.cl b/scrypt120713.cl index f7c1a6ce..95b006e9 100644 --- a/scrypt120713.cl +++ b/scrypt120713.cl @@ -689,7 +689,9 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup) #define NFLAG (0x7F) __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) -__kernel void search(__global const uint4 * restrict input, __global uint*restrict output, __global uint4*restrict padcache, const uint4 midstate0, const uint4 midstate16) +__kernel void search(__global const uint4 * restrict input, +__global uint*restrict output, __global uint4*restrict padcache, +const uint4 midstate0, const uint4 midstate16, const uint target) { uint gid = get_global_id(0); uint4 X[8]; @@ -722,7 +724,7 @@ __kernel void search(__global const uint4 * restrict input, __global uint*restri SHA256_fixed(&tmp0,&tmp1); SHA256(&ostate0,&ostate1, tmp0, tmp1, (uint4)(0x80000000U, 0U, 0U, 0U), (uint4)(0U, 0U, 0U, 0x300U)); - if ((ostate1.w&0xFFFF) == 0) + if (!(ostate1.w&target)) output[FOUND] = output[NFLAG & gid] = gid; }