Browse Source

Move to cgminer style buffer return and file naming convention and fix a compiler warning.

Con Kolivas 14 years ago
parent
commit
c40f51c7c1
3 changed files with 14 additions and 19 deletions
  1. 2 2
      Makefile.am
  2. 3 3
      ocl.c
  3. 9 14
      phatk110816.cl

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

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

+ 9 - 14
phatk2_2.cl → 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;
-	}
 }