diff --git a/findnonce.c b/findnonce.c index 23c01e62..45b5910e 100644 --- a/findnonce.c +++ b/findnonce.c @@ -181,14 +181,14 @@ static void *postcalc_hash(void *userdata) pthread_detach(pthread_self()); cycle: - while (entry < OUTBUFFERS) { + while (entry < FOUND) { if (pcd->res[entry]) { nonce = pcd->res[entry++]; break; } entry++; } - if (entry == OUTBUFFERS) + if (entry == FOUND) goto out; A = blk->cty_a; B = blk->cty_b; @@ -231,7 +231,7 @@ cycle: hw_errors++; thr->cgpu->hw_errors++; } - if (entry < OUTBUFFERS) + if (entry < FOUND) goto cycle; out: free(pcd); diff --git a/findnonce.h b/findnonce.h index b22afeb6..93cd1fe6 100644 --- a/findnonce.h +++ b/findnonce.h @@ -5,9 +5,10 @@ #define MAXTHREADS (0xFFFFFFFEULL) /* Maximum worksize 4k to match page size */ -#define MAXBUFFERS (4095) -#define BUFFERSIZE (sizeof(uint32_t) * (MAXBUFFERS + 1)) -#define OUTBUFFERS (0xFF) +#define MAXBUFFERS (0xFF) +#define BUFFERSIZE (sizeof(uint32_t) * MAXBUFFERS) +#define FOUND (0x80) +/* #define NFLAG (0x7F) Just for reference */ #ifdef HAVE_OPENCL extern void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data); diff --git a/main.c b/main.c index e31718b8..03fc3ca3 100644 --- a/main.c +++ b/main.c @@ -3610,7 +3610,7 @@ static void *gpuminer_thread(void *userdata) { applog(LOG_ERR, "Error: clSetKernelArg of all params failed."); goto out; } /* MAXBUFFERS entry is used as a flag to say nonces exist */ - if (res[MAXBUFFERS]) { + if (res[FOUND]) { /* Clear the buffer again */ status = clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, CL_FALSE, 0, BUFFERSIZE, blank_res, 0, NULL, NULL); diff --git a/phatk110816.cl b/phatk110816.cl index 9b0b777b..ebcf0bb0 100644 --- a/phatk110816.cl +++ b/phatk110816.cl @@ -387,42 +387,42 @@ void search( const uint state0, const uint state1, const uint state2, const uint u v = W[117] + W[108] + Vals[3] + Vals[7] + P2(124) + P1(124) + Ch((Vals[0] + Vals[4]) + (K[59] + W(59+64)) + s1(64+59)+ ch(59+64),Vals[1],Vals[2]); u g = -(K[60] + H[7]) - S1((Vals[0] + Vals[4]) + (K[59] + W(59+64)) + s1(64+59)+ ch(59+64)); - -#define NFLAG (0xFF) + +#define FOUND (0x80) +#define NFLAG (0x7F) #ifdef VECTORS4 if (v.x == g.x) { - output[MAXBUFFERS] = output[NFLAG & W[3].x] = W[3].x; + output[FOUND] = output[NFLAG & W[3].x] = W[3].x; } if (v.y == g.y) { - output[MAXBUFFERS] = output[NFLAG & W[3].y] = W[3].y; + output[FOUND] = output[NFLAG & W[3].y] = W[3].y; } if (v.z == g.z) { - output[MAXBUFFERS] = output[NFLAG & W[3].z] = W[3].z; + output[FOUND] = output[NFLAG & W[3].z] = W[3].z; } if (v.w == g.w) { - output[MAXBUFFERS] = output[NFLAG & W[3].w] = W[3].w; + output[FOUND] = output[NFLAG & W[3].w] = W[3].w; } #else #ifdef VECTORS2 if (v.x == g.x) { - output[MAXBUFFERS] = output[NFLAG & W[3].x] = W[3].x; + output[FOUND] = output[NFLAG & W[3].x] = W[3].x; } if (v.y == g.y) { - output[MAXBUFFERS] = output[NFLAG & W[3].y] = W[3].y; + output[FOUND] = output[NFLAG & W[3].y] = W[3].y; } #else if (v == g) { - output[MAXBUFFERS] = output[NFLAG & W[3]] = W[3]; + output[FOUND] = output[NFLAG & W[3]] = W[3]; } #endif #endif } - diff --git a/poclbm110816.cl b/poclbm110816.cl index a16e4de8..a222c6e6 100644 --- a/poclbm110816.cl +++ b/poclbm110816.cl @@ -625,32 +625,32 @@ __kernel void search( const uint state0, const uint state1, const uint state2, c Vals[7] = Vals[7] + Vals[3] + (rotr(Vals[0], 6) ^ rotr(Vals[0], 11) ^ rotr(Vals[0], 25)) + ch(Vals[0], Vals[1], Vals[2]) + K[60] + W[12]; -#define MAXBUFFERS (4095) -#define NFLAG (0xFF) +#define FOUND (0x80) +#define NFLAG (0x7F) #if defined(VECTORS4) || defined(VECTORS2) if (Vals[7].x == -0x5be0cd19U) { - output[MAXBUFFERS] = output[NFLAG & nonce.x] = nonce.x; + output[FOUND] = output[NFLAG & nonce.x] = nonce.x; } if (Vals[7].y == -0x5be0cd19U) { - output[MAXBUFFERS] = output[NFLAG & nonce.y] = nonce.y; + output[FOUND] = output[NFLAG & nonce.y] = nonce.y; } #ifdef VECTORS4 if (Vals[7].z == -0x5be0cd19U) { - output[MAXBUFFERS] = output[NFLAG & nonce.z] = nonce.z; + output[FOUND] = output[NFLAG & nonce.z] = nonce.z; } if (Vals[7].w == -0x5be0cd19U) { - output[MAXBUFFERS] = output[NFLAG & nonce.w] = nonce.w; + output[FOUND] = output[NFLAG & nonce.w] = nonce.w; } #endif #else if (Vals[7] == -0x5be0cd19U) { - output[MAXBUFFERS] = output[NFLAG & nonce] = nonce; + output[FOUND] = output[NFLAG & nonce] = nonce; } #endif -} \ No newline at end of file +}