From 49c28b3929866a3e26668ec5000e0dfd41dcb6e8 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Wed, 15 Feb 2012 00:00:05 +1100
Subject: [PATCH] Use PreVal4addT1 instead of PreVal4 in poclbm kernel.

---
 device-gpu.c    | 2 +-
 poclbm120213.cl | 5 ++---
 2 files changed, 3 insertions(+), 4 deletions(-)

diff --git a/device-gpu.c b/device-gpu.c
index fca9d11d..c42b6028 100644
--- a/device-gpu.c
+++ b/device-gpu.c
@@ -692,7 +692,7 @@ static cl_int queue_poclbm_kernel(_clState *clState, dev_blk_ctx *blk)
 	CL_SET_BLKARG(B1addK6);
 	CL_SET_BLKARG(W16addK16);
 	CL_SET_BLKARG(W17addK17);
-	CL_SET_BLKARG(PreVal4);
+	CL_SET_BLKARG(PreVal4addT1);
 	CL_SET_BLKARG(PreVal0);
 
 	CL_SET_ARG(clState->outputBuffer);
diff --git a/poclbm120213.cl b/poclbm120213.cl
index 5785a7a3..78c928c9 100644
--- a/poclbm120213.cl
+++ b/poclbm120213.cl
@@ -76,7 +76,7 @@ __kernel void search(const uint state0, const uint state1, const uint state2, co
 						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,
+						const uint PreVal4addT1, const uint Preval0,
 						__global uint * output)
 {
 	u W[24];
@@ -90,7 +90,6 @@ __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]=PreVal4+nonce;
 
 Vals[0]=Preval0+nonce;
 
@@ -100,7 +99,7 @@ Vals[3]+=D1A;
 
 Vals[7]=Vals[3];
 Vals[7]+=h1;
-Vals[4]+=fcty_e2;
+Vals[4]=PreVal4addT1+nonce;
 Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
 
 Vals[2]=C1addK5;