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

Merge pull request #185 from Diapolo/diakgcn

Browse Source 
add goffset support for diakgcn with -v 1 and update kernel version
This commit is contained in:
Con Kolivas 2012-04-27 06:25:38 -07:00
commit 5d99f5d3e8
4 changed files with 19 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
configure.ac
View File

@ -367,7 +367,7 @@ AC_DEFINE_UNQUOTED([CGMINER_PREFIX], ["$prefix/bin"], [Path to cgminer install])
AC_DEFINE_UNQUOTED([PHATK_KERNNAME], ["phatk120223"], [Filename for phatk kernel])
AC_DEFINE_UNQUOTED([POCLBM_KERNNAME], ["poclbm120327"], [Filename for poclbm kernel])
AC_DEFINE_UNQUOTED([DIAKGCN_KERNNAME], ["diakgcn120223"], [Filename for diakgcn kernel])
AC_DEFINE_UNQUOTED([DIAKGCN_KERNNAME], ["diakgcn120427"], [Filename for diakgcn kernel])
AC_DEFINE_UNQUOTED([DIABLO_KERNNAME], ["diablo120328"], [Filename for diablo kernel])

10
diakgcn120223.cl → diakgcn120427.cl
View File

@ -1,4 +1,4 @@
// DiaKGCN 16-03-2012 - OpenCL kernel by Diapolo
// DiaKGCN 27-04-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.
// The kernel was rewritten by me (Diapolo) and is still public-domain!
@ -33,7 +33,9 @@
__kernel
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
void search(
#ifndef GOFFSET
const u base,
#endif
const uint PreVal0, const uint PreVal4,
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,
@ -56,7 +58,11 @@ __kernel
#elif defined VECTORS2
const u nonce = (uint)(get_local_id(0)) * 2U + (uint)(get_group_id(0)) * (uint)(WORKVEC) + base;
#else
const u nonce = (uint)(get_local_id(0)) + (uint)(get_group_id(0)) * (uint)(WORKSIZE) + base;
#ifdef GOFFSET
const u nonce = (uint)(get_global_id(0));
#else
const u nonce = (uint)(get_local_id(0)) + (uint)(get_group_id(0)) * (uint)(WORKSIZE) + base;
#endif
#endif
V[0] = PreVal0 + nonce;

16
driver-opencl.c
View File

@ -879,15 +879,17 @@ static cl_int queue_diakgcn_kernel(_clState *clState, dev_blk_ctx *blk,
__maybe_unused cl_uint threads)
{
cl_kernel *kernel = &clState->kernel;
cl_uint vwidth = clState->vwidth;
unsigned int i, num = 0;
unsigned int num = 0;
cl_int status = 0;
uint *nonces;
nonces = alloca(sizeof(uint) * vwidth);
for (i = 0; i < vwidth; i++)
nonces[i] = blk->nonce + i;
CL_SET_VARG(vwidth, nonces);
if (!clState->goffset) {
cl_uint vwidth = clState->vwidth;
uint *nonces = alloca(sizeof(uint) * vwidth);
unsigned int i;
for (i = 0; i < vwidth; i++)
nonces[i] = blk->nonce + i;
CL_SET_VARG(vwidth, nonces);
}
CL_SET_BLKARG(PreVal0);
CL_SET_BLKARG(PreVal4_2);

2
ocl.c
View File

@ -421,7 +421,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
gpus[gpu].vwidth = preferred_vwidth;
}
if ((clState->chosen_kernel == KL_POCLBM || clState->chosen_kernel == KL_DIABLO) &&
if ((clState->chosen_kernel == KL_POCLBM || clState->chosen_kernel == KL_DIABLO || clState->chosen_kernel == KL_DIAKGCN) &&
clState->vwidth == 1 && clState->hasOpenCL11plus)
clState->goffset = true;