From 0f782ba6bd231565aef8901d30371364b0979b09 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Wed, 17 Aug 2011 15:47:18 +1000
Subject: [PATCH] Update poclbm kernel to FF sized mask and only check that
 range.

---
 Makefile.am                        |  4 ++--
 findnonce.c                        |  6 +++---
 findnonce.h                        |  1 +
 ocl.c                              |  4 ++--
 poclbm110717.cl => poclbm110816.cl | 14 ++++++--------
 5 files changed, 14 insertions(+), 15 deletions(-)
 rename poclbm110717.cl => poclbm110816.cl (99%)

diff --git a/Makefile.am b/Makefile.am
index 83dcab43..0fed929d 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -15,7 +15,7 @@ INCLUDES	= $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES)
 
 bin_PROGRAMS	= cgminer
 
-bin_SCRIPTS	= phatk110816.cl poclbm110717.cl
+bin_SCRIPTS	= phatk110816.cl poclbm110816.cl
 
 cgminer_SOURCES	= elist.h miner.h compat.h bench_block.h	\
 		  main.c util.c					\
@@ -23,7 +23,7 @@ cgminer_SOURCES	= elist.h miner.h compat.h bench_block.h	\
 		  sha256_generic.c sha256_4way.c sha256_via.c	\
 		  sha256_cryptopp.c sha256_sse2_amd64.c		\
 		  sha256_sse4_amd64.c \
-		  phatk110816.cl poclbm110717.cl
+		  phatk110816.cl poclbm110816.cl
 
 cgminer_LDFLAGS	= $(PTHREAD_FLAGS)
 cgminer_LDADD	= @LIBCURL_LIBS@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @OPENCL_LIBS@ @NCURSES_LIBS@ @PDCURSES_LIBS@ lib/libgnu.a ccan/libccan.a
diff --git a/findnonce.c b/findnonce.c
index 229b3e6d..23c01e62 100644
--- a/findnonce.c
+++ b/findnonce.c
@@ -181,14 +181,14 @@ static void *postcalc_hash(void *userdata)
 
 	pthread_detach(pthread_self());
 cycle:
-	while (entry < MAXBUFFERS) {
+	while (entry < OUTBUFFERS) {
 		if (pcd->res[entry]) {
 			nonce = pcd->res[entry++];
 			break;
 		}
 		entry++;
 	}
-	if (entry == MAXBUFFERS)
+	if (entry == OUTBUFFERS)
 		goto out;
 
 	A = blk->cty_a; B = blk->cty_b;
@@ -231,7 +231,7 @@ cycle:
 		hw_errors++;
 		thr->cgpu->hw_errors++;
 	}
-	if (entry < MAXBUFFERS)
+	if (entry < OUTBUFFERS)
 		goto cycle;
 out:
 	free(pcd);
diff --git a/findnonce.h b/findnonce.h
index 3ca7c8bd..b22afeb6 100644
--- a/findnonce.h
+++ b/findnonce.h
@@ -7,6 +7,7 @@
 /* Maximum worksize 4k to match page size */
 #define MAXBUFFERS (4095)
 #define BUFFERSIZE (sizeof(uint32_t) * (MAXBUFFERS + 1))
+#define OUTBUFFERS (0xFF)
 
 #ifdef HAVE_OPENCL
 extern void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data);
diff --git a/ocl.c b/ocl.c
index 6744eeb5..c6c386c4 100644
--- a/ocl.c
+++ b/ocl.c
@@ -350,8 +350,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 
 	switch (chosen_kernel) {
 		case KL_POCLBM:
-			strcpy(filename, "poclbm110717.cl");
-			strcpy(binaryfilename, "poclbm110717");
+			strcpy(filename, "poclbm110816.cl");
+			strcpy(binaryfilename, "poclbm110816");
 			break;
 		case KL_NONE: /* Shouldn't happen */
 		case KL_PHATK:
diff --git a/poclbm110717.cl b/poclbm110816.cl
similarity index 99%
rename from poclbm110717.cl
rename to poclbm110816.cl
index 5b2c1994..a16e4de8 100644
--- a/poclbm110717.cl
+++ b/poclbm110816.cl
@@ -625,32 +625,30 @@ __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];
 
-	Vals[7]+=0x5be0cd19U;
-
 #define MAXBUFFERS (4095)
-#define NFLAG (0xFFEUL)
+#define NFLAG (0xFF)
 
 #if defined(VECTORS4) || defined(VECTORS2)
-	if (Vals[7].x == 0)
+	if (Vals[7].x == -0x5be0cd19U)
 	{
 		output[MAXBUFFERS] = output[NFLAG & nonce.x] =  nonce.x;
 	}
-	if (Vals[7].y == 0)
+	if (Vals[7].y == -0x5be0cd19U)
 	{
 		output[MAXBUFFERS] = output[NFLAG & nonce.y] =  nonce.y;
 	}
 #ifdef VECTORS4
-	if (Vals[7].z == 0)
+	if (Vals[7].z == -0x5be0cd19U)
 	{
 		output[MAXBUFFERS] = output[NFLAG & nonce.z] =  nonce.z;
 	}
-	if (Vals[7].w == 0)
+	if (Vals[7].w == -0x5be0cd19U)
 	{
 		output[MAXBUFFERS] = output[NFLAG & nonce.w] =  nonce.w;
 	}
 #endif
 #else
-	if (Vals[7] == 0)
+	if (Vals[7] == -0x5be0cd19U)
 	{
 		output[MAXBUFFERS] = output[NFLAG & nonce] =  nonce;
 	}