From c40f51c7c1160bee91eb16cdbfb5f8f441872095 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Wed, 17 Aug 2011 15:06:59 +1000
Subject: [PATCH] Move to cgminer style buffer return and file naming
 convention and fix a compiler warning.

---
 Makefile.am                   |  4 ++--
 ocl.c                         |  6 +++---
 phatk2_2.cl => phatk110816.cl | 23 +++++++++--------------
 3 files changed, 14 insertions(+), 19 deletions(-)
 rename phatk2_2.cl => phatk110816.cl (93%)

diff --git a/Makefile.am b/Makefile.am
index e453ad54..83dcab43 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -15,7 +15,7 @@ INCLUDES	= $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES)
 
 bin_PROGRAMS	= cgminer
 
-bin_SCRIPTS	= phatk2_2.cl poclbm110717.cl
+bin_SCRIPTS	= phatk110816.cl poclbm110717.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 \
-		  phatk2_2.cl poclbm110717.cl
+		  phatk110816.cl poclbm110717.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/ocl.c b/ocl.c
index 15c2517b..6744eeb5 100644
--- a/ocl.c
+++ b/ocl.c
@@ -355,8 +355,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 			break;
 		case KL_NONE: /* Shouldn't happen */
 		case KL_PHATK:
-			strcpy(filename, "phatk2_2.cl");
-			strcpy(binaryfilename, "phatk2_2");
+			strcpy(filename, "phatk110816.cl");
+			strcpy(binaryfilename, "phatk110816");
 			break;
 	}
 
@@ -502,7 +502,7 @@ build:
 
 	/* create a cl program executable for all the devices specified */
 	char CompilerOptions[256];
-	sprintf(CompilerOptions, "%s%i", "-DWORKSIZE=", clState->work_size);
+	sprintf(CompilerOptions, "%s%i", "-DWORKSIZE=", (int)clState->work_size);
 	//int n = 1000;
 	//while(n--)
 	//	printf("%s", CompilerOptions);
diff --git a/phatk2_2.cl b/phatk110816.cl
similarity index 93%
rename from phatk2_2.cl
rename to phatk110816.cl
index fd1ef737..e9e913aa 100644
--- a/phatk2_2.cl
+++ b/phatk110816.cl
@@ -388,46 +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));
 	
-	uint nonce = 0;
+#define NFLAG (0xFFEUL)
+
 #ifdef VECTORS4
 	if (v.x == g.x)
 	{
-		nonce = W[3].x;
+		output[MAXBUFFERS] = output[NFLAG & W[3].x] = W[3].x;
 	}
 	if (v.y == g.y)
 	{
-		nonce = W[3].y;
+		output[MAXBUFFERS] = output[NFLAG & W[3].y] = W[3].y;
 	}
 	if (v.z == g.z)
 	{
-		nonce = W[3].z;
+		output[MAXBUFFERS] = output[NFLAG & W[3].z] = W[3].z;
 	}
 	if (v.w == g.w)
 	{
-		nonce = W[3].w;
+		output[MAXBUFFERS] = output[NFLAG & W[3].w] = W[3].w;
 	}
 #else
 	#ifdef VECTORS2
 		if (v.x == g.x)
 		{
-			nonce = W[3].x;
+			output[MAXBUFFERS] = output[NFLAG & W[3].x] = W[3].x;
 		}
 		if (v.y == g.y)
 		{
-			nonce = W[3].y;
+			output[MAXBUFFERS] = output[NFLAG & W[3].y] = W[3].y;
 		}
 	#else
 		if (v == g)
 		{
-			nonce = W[3];
+			output[MAXBUFFERS] = output[NFLAG & W[3]] = W[3];
 		}
 	#endif
 #endif
-	if(nonce)
-	{
-		//Faster to shift the nonce by 2 due to 4-DWORD addressing and does not add more collisions
-		output[MAXBUFFERS] = nonce;
-		output[get_local_id(0)] = nonce;
-	}
 }