From 5c4df1309acdd7b7b07c2a221bb4ef20529bc6a8 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Tue, 14 Feb 2012 23:28:46 +1100
Subject: [PATCH] Import PreVal4 and PreVal0 into poclbm kernel.

---
 device-gpu.c    | 4 +++-
 poclbm120213.cl | 9 ++++-----
 2 files changed, 7 insertions(+), 6 deletions(-)

diff --git a/device-gpu.c b/device-gpu.c
index ae054478..fca9d11d 100644
--- a/device-gpu.c
+++ b/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);
 
diff --git a/poclbm120213.cl b/poclbm120213.cl
index 6d103166..5785a7a3 100644
--- a/poclbm120213.cl
+++ b/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);