Update poclbm kernel to FF sized mask and only check that range.

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

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

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

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:

poclbm110717.cl → 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;
 	}