Browse Source

Import PreVal4 and PreVal0 into poclbm kernel.

Con Kolivas 14 years ago
parent
commit
5c4df1309a
2 changed files with 7 additions and 6 deletions
  1. 3 1
      device-gpu.c
  2. 4 5
      poclbm120213.cl

+ 3 - 1
device-gpu.c

@@ -685,13 +685,15 @@ static cl_int queue_poclbm_kernel(_clState *clState, dev_blk_ctx *blk)
 	CL_SET_BLKARG(fW3);
 	CL_SET_BLKARG(fW15);
 	CL_SET_BLKARG(fW01r);
-	CL_SET_BLKARG(fcty_e);
+
 	CL_SET_BLKARG(fcty_e2);
 	CL_SET_BLKARG(D1A);
 	CL_SET_BLKARG(C1addK5);
 	CL_SET_BLKARG(B1addK6);
 	CL_SET_BLKARG(W16addK16);
 	CL_SET_BLKARG(W17addK17);
+	CL_SET_BLKARG(PreVal4);
+	CL_SET_BLKARG(PreVal0);
 
 	CL_SET_ARG(clState->outputBuffer);
 

+ 4 - 5
poclbm120213.cl

@@ -73,9 +73,10 @@ __kernel void search(const uint state0, const uint state1, const uint state2, co
 						const uint f1, const uint g1, const uint h1,
 						const u base,
 						const uint fw0, const uint fw1, const uint fw2, const uint fw3, const uint fw15, const uint fw01r,
-						const uint fcty_e, const uint fcty_e2,
+						const uint fcty_e2,
 						const uint D1A, const uint C1addK5, const uint B1addK6,
 						const uint W16addK16, const uint W17addK17,
+						const uint PreVal4, const uint Preval0,
 						__global uint * output)
 {
 	u W[24];
@@ -89,11 +90,9 @@ __kernel void search(const uint state0, const uint state1, const uint state2, co
 	const u nonce = base + get_local_id(0) + get_group_id(0) * (WORKSIZE);
 #endif
 
-Vals[4]=fcty_e;
-Vals[4]+=nonce;
+Vals[4]=PreVal4+nonce;
 
-Vals[0]=Vals[4];
-Vals[0]+=state0;
+Vals[0]=Preval0+nonce;
 
 Vals[3]=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
 Vals[3]+=ch(Vals[0],b1,c1);