From cf54f9b850ff70e22b4251ca611e90f390399ed5 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Wed, 17 Aug 2011 16:07:15 +1000
Subject: [PATCH] Move to 256 sized buffers and don't risk overwrite by using
 only 127 mask.

---
 findnonce.c     |  6 +++---
 findnonce.h     |  7 ++++---
 main.c          |  2 +-
 phatk110816.cl  | 20 ++++++++++----------
 poclbm110816.cl | 16 ++++++++--------
 5 files changed, 26 insertions(+), 25 deletions(-)

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
+}