GOSTSec/sgminer
1
0
Fork 0
mirror of https://github.com/GOSTSec/sgminer synced 2025-01-23 04:54:26 +00:00
Code Issues Releases Wiki Activity

Use local and group id on poclbm kernel as well.

Browse Source
This commit is contained in:
Con Kolivas 2012-02-13 12:51:24 +11:00
parent b2b5083bda
commit 60f8ccb313
2 changed files with 13 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
device-gpu.c
View File

@ -651,9 +651,11 @@ static _clState *clStates[MAX_GPUDEVICES];
static cl_int queue_poclbm_kernel(_clState *clState, dev_blk_ctx *blk) static cl_int queue_poclbm_kernel(_clState *clState, dev_blk_ctx *blk)
{ {
cl_uint vwidth = clState->preferred_vwidth;
cl_kernel *kernel = &clState->kernel; cl_kernel *kernel = &clState->kernel;
unsigned int i, num = 0;
cl_int status = 0; cl_int status = 0;
int num = 0; uint *nonces;
CL_SET_BLKARG(ctx_a); CL_SET_BLKARG(ctx_a);
CL_SET_BLKARG(ctx_b); CL_SET_BLKARG(ctx_b);
@ -663,13 +665,18 @@ static cl_int queue_poclbm_kernel(_clState *clState, dev_blk_ctx *blk)
CL_SET_BLKARG(ctx_f); CL_SET_BLKARG(ctx_f);
CL_SET_BLKARG(ctx_g); CL_SET_BLKARG(ctx_g);
CL_SET_BLKARG(ctx_h); CL_SET_BLKARG(ctx_h);
CL_SET_BLKARG(cty_b); CL_SET_BLKARG(cty_b);
CL_SET_BLKARG(cty_c); CL_SET_BLKARG(cty_c);
CL_SET_BLKARG(cty_d); CL_SET_BLKARG(cty_d);
CL_SET_BLKARG(cty_f); CL_SET_BLKARG(cty_f);
CL_SET_BLKARG(cty_g); CL_SET_BLKARG(cty_g);
CL_SET_BLKARG(cty_h); CL_SET_BLKARG(cty_h);
CL_SET_BLKARG(nonce);
nonces = alloca(sizeof(uint) * vwidth);
for (i = 0; i < vwidth; i++)
nonces[i] = blk->nonce + i;
CL_SET_VARG(vwidth, nonces);
CL_SET_BLKARG(fW0); CL_SET_BLKARG(fW0);
CL_SET_BLKARG(fW1); CL_SET_BLKARG(fW1);

8
poclbm120210.cl
View File

@ -71,7 +71,7 @@ __kernel void search(const uint state0, const uint state1, const uint state2, co
const uint state4, const uint state5, const uint state6, const uint state7, const uint state4, const uint state5, const uint state6, const uint state7,
const uint b1, const uint c1, const uint d1, const uint b1, const uint c1, const uint d1,
const uint f1, const uint g1, const uint h1, const uint f1, const uint g1, const uint h1,
const uint base, 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 fw0, const uint fw1, const uint fw2, const uint fw3, const uint fw15, const uint fw01r, const uint fcty_e, const uint fcty_e2,
__global uint * output) __global uint * output)
{ {
@ -80,11 +80,11 @@ __kernel void search(const uint state0, const uint state1, const uint state2, co
u nonce; u nonce;
#ifdef VECTORS4 #ifdef VECTORS4
nonce = base + (get_global_id(0)<<2) + (uint4)(0, 1, 2, 3); nonce = base + (uint)(get_local_id(0)) * 4u + (uint)(get_group_id(0)) * (WORKSIZE * 4u);
#elif defined VECTORS2 #elif defined VECTORS2
nonce = base + (get_global_id(0)<<1) + (uint2)(0, 1); nonce = base + (uint)(get_local_id(0)) * 2u + (uint)(get_group_id(0)) * (WORKSIZE * 2u);
#else #else
nonce = base + get_global_id(0); nonce = base + get_local_id(0) + get_group_id(0) * (WORKSIZE);
#endif #endif
W[20] = fcty_e + nonce; W[20] = fcty_e + nonce;