Commit 5c4df1309acdd7b7b07c2a221bb4ef20529bc6a8
2012-02-14T23:28:46
Import PreVal4 and PreVal0 into poclbm kernel.
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51
diff --git a/device-gpu.c b/device-gpu.c
index ae05447..fca9d11 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 6d10316..5785a7a 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);