Move to 256 sized buffers and don't risk overwrite by using only 127 mask.

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);

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);

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);

phatk110816.cl

@@ -388,41 +388,41 @@ 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
 }
-

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
 }