GOSTSec/sgminer
1
0
Fork 0
mirror of https://github.com/GOSTSec/sgminer synced 2025-02-02 10:04:33 +00:00
Code Issues Releases Wiki Activity

tried to fix vector offset with diakgcn

Browse Source
This commit is contained in:
Philip Kaufmann 2012-02-22 08:53:57 +01:00
parent 26c59fbf0f
commit bb51b628d3
2 changed files with 4 additions and 31 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
device-gpu.c
View File

@ -796,8 +796,7 @@ static cl_int queue_phatk_kernel(_clState *clState, dev_blk_ctx *blk,
return status; return status;
} }
static cl_int queue_diakgcn_kernel(_clState *clState, dev_blk_ctx *blk, static cl_int queue_diakgcn_kernel(_clState *clState, dev_blk_ctx *blk, cl_uint threads)
__maybe_unused cl_uint threads)
{ {
cl_kernel *kernel = &clState->kernel; cl_kernel *kernel = &clState->kernel;
cl_uint vwidth = clState->vwidth; cl_uint vwidth = clState->vwidth;
@ -807,7 +806,7 @@ static cl_int queue_diakgcn_kernel(_clState *clState, dev_blk_ctx *blk,
nonces = alloca(sizeof(uint) * vwidth); nonces = alloca(sizeof(uint) * vwidth);
for (i = 0; i < vwidth; i++) for (i = 0; i < vwidth; i++)
nonces[i] = blk->nonce + i; nonces[i] = blk->nonce + (i * threads);
CL_SET_VARG(vwidth, nonces); CL_SET_VARG(vwidth, nonces);
CL_SET_BLKARG(PreVal0); CL_SET_BLKARG(PreVal0);

30
diakgcn120222.cl
View File

@ -1,4 +1,4 @@
// DiaKGCN 20-02-2012 - OpenCL kernel by Diapolo // DiaKGCN 22-02-2012 - OpenCL kernel by Diapolo
// //
// Parts and / or ideas for this kernel are based upon the public-domain poclbm project, the phatk kernel by Phateus and the DiabloMiner kernel by DiabloD3. // Parts and / or ideas for this kernel are based upon the public-domain poclbm project, the phatk kernel by Phateus and the DiabloMiner kernel by DiabloD3.
// The kernel was rewritten by me (Diapolo) and is still public-domain! // The kernel was rewritten by me (Diapolo) and is still public-domain!
@ -35,9 +35,7 @@
__kernel __kernel
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
void search( void search(
#ifndef GOFFSET
const u base, const u base,
#endif
const uint PreVal0, const uint PreVal4, const uint PreVal0, const uint PreVal4,
const uint H1, const uint D1A, const uint B1, const uint C1, const uint H1, const uint D1A, const uint B1, const uint C1,
const uint F1, const uint G1, const uint C1addK5, const uint B1addK6, const uint PreVal0addK7, const uint F1, const uint G1, const uint C1addK5, const uint B1addK6, const uint PreVal0addK7,
@ -55,31 +53,7 @@ __kernel
u V[8]; u V[8];
u W[16]; u W[16];
#ifdef VECTORS8 const u nonce = base + (uint)get_global_id(0);
#ifdef GOFFSET
const u nonce = ((uint)get_global_id(0) << 3) + (u)(0, 1, 2, 3, 4, 5, 6, 7);
#else
const u nonce = ((uint)get_group_id(0) * (uint)get_local_size(0) << 3) + ((uint)get_local_id(0) << 3) + base;
#endif
#elif defined VECTORS4
#ifdef GOFFSET
const u nonce = ((uint)get_global_id(0) << 2) + (u)(0, 1, 2, 3);
#else
const u nonce = ((uint)get_group_id(0) * (uint)get_local_size(0) << 2) + ((uint)get_local_id(0) << 2) + base;
#endif
#elif defined VECTORS2
#ifdef GOFFSET
const u nonce = ((uint)get_global_id(0) << 1) + (u)(0, 1);
#else
const u nonce = ((uint)get_group_id(0) * (uint)get_local_size(0) << 1) + ((uint)get_local_id(0) << 1) + base;
#endif
#else
#ifdef GOFFSET
const u nonce = (uint)get_global_id(0);
#else
const u nonce = ((uint)get_group_id(0) * (uint)get_local_size(0)) + (uint)get_local_id(0) + base;
#endif
#endif
V[0] = PreVal0 + nonce; V[0] = PreVal0 + nonce;
V[1] = B1; V[1] = B1;