Navigation Menu

Skip to content
This repository has been archived by the owner on Jan 9, 2021. It is now read-only.

Commit

Permalink
Fix indents
Browse files Browse the repository at this point in the history
  • Loading branch information
@bitbandi
bitbandi committed Jan 20, 2017
1 parent 41a4798 commit ade8905
Show file tree
Hide file tree
Showing 4 changed files with 294 additions and 295 deletions.
150 changes: 75 additions & 75 deletions algorithm.c
View file
Expand Up @@ -986,26 +986,26 @@ static cl_int queue_ethash_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u
cg_ulock(&dag->lock);
if (dag->dag_buffer == NULL || blk->work->eth_epoch > dag->max_epoch) {
if (dag->dag_buffer != NULL) {
cg_dlock(&pool->data_lock);
cg_dlock(&pool->data_lock);
clReleaseMemObject(dag->dag_buffer);
}
else {
cg_ulock(&pool->data_lock);
int size = ++pool->eth_cache.nDevs;
pool->eth_cache.dags = (eth_dag_t **) realloc(pool->eth_cache.dags, sizeof(void*) * size);
pool->eth_cache.dags[size-1] = dag;
dag->pool = pool;
cg_dwlock(&pool->data_lock);
cg_ulock(&pool->data_lock);
int size = ++pool->eth_cache.nDevs;
pool->eth_cache.dags = (eth_dag_t **) realloc(pool->eth_cache.dags, sizeof(void*) * size);
pool->eth_cache.dags[size-1] = dag;
dag->pool = pool;
cg_dwlock(&pool->data_lock);
}
dag->max_epoch = blk->work->eth_epoch + eth_future_epochs;
dag->dag_buffer = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, EthGetDAGSize(dag->max_epoch), NULL, &status);
if (status != CL_SUCCESS) {
cg_runlock(&pool->data_lock);
dag->max_epoch = 0;
dag->dag_buffer = NULL;
cg_wunlock(&dag->lock);
applog(LOG_ERR, "Error %d: Creating the DAG buffer failed.", status);
return status;
cg_runlock(&pool->data_lock);
dag->max_epoch = 0;
dag->dag_buffer = NULL;
cg_wunlock(&dag->lock);
applog(LOG_ERR, "Error %d: Creating the DAG buffer failed.", status);
return status;
}
}
else
Expand Down Expand Up @@ -1067,7 +1067,7 @@ static cl_int queue_ethash_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u

// DO NOT flip80.
status |= clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 32, blk->work->data, 0, NULL, NULL);

CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(clState->CLbuffer0);
CL_SET_ARG(dag->dag_buffer);
Expand All @@ -1082,68 +1082,68 @@ static cl_int queue_ethash_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u
}

static void append_equihash_compiler_options(struct _build_kernel_data *data, struct cgpu_info *cgpu, struct _algorithm_t *algorithm)
{
{
strcat(data->compiler_options, "");
}

static cl_int queue_cryptonight_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_unused cl_uint threads)
{
cl_kernel *kernel = &clState->kernel;
unsigned int num = 0;
cl_int status = 0, tgt32 = (blk->work->XMRTarget);
cl_ulong le_target = ((cl_ulong)(blk->work->XMRTarget));

//le_target = *(cl_ulong *)(blk->work->device_target + 24);
memcpy(clState->cldata, blk->work->data, 76);
status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 76, clState->cldata , 0, NULL, NULL);
CL_SET_ARG(clState->CLbuffer0);
CL_SET_ARG(clState->Scratchpads);
CL_SET_ARG(clState->States);
num = 0;
kernel = clState->extra_kernels;
CL_SET_ARG(clState->Scratchpads);
CL_SET_ARG(clState->States);
num = 0;
CL_NEXTKERNEL_SET_ARG(clState->Scratchpads);
CL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[0]);
CL_SET_ARG(clState->BranchBuffer[1]);
CL_SET_ARG(clState->BranchBuffer[2]);
CL_SET_ARG(clState->BranchBuffer[3]);
num = 0;
CL_NEXTKERNEL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[0]);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(tgt32);
// last to be set in driver-opencl.c
num = 0;
CL_NEXTKERNEL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[1]);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(tgt32);
num = 0;
CL_NEXTKERNEL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[2]);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(tgt32);
num = 0;
CL_NEXTKERNEL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[3]);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(tgt32);
return(status);
cl_kernel *kernel = &clState->kernel;
unsigned int num = 0;
cl_int status = 0, tgt32 = (blk->work->XMRTarget);
cl_ulong le_target = ((cl_ulong)(blk->work->XMRTarget));

//le_target = *(cl_ulong *)(blk->work->device_target + 24);
memcpy(clState->cldata, blk->work->data, 76);

status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 76, clState->cldata , 0, NULL, NULL);

CL_SET_ARG(clState->CLbuffer0);
CL_SET_ARG(clState->Scratchpads);
CL_SET_ARG(clState->States);

num = 0;
kernel = clState->extra_kernels;
CL_SET_ARG(clState->Scratchpads);
CL_SET_ARG(clState->States);

num = 0;
CL_NEXTKERNEL_SET_ARG(clState->Scratchpads);
CL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[0]);
CL_SET_ARG(clState->BranchBuffer[1]);
CL_SET_ARG(clState->BranchBuffer[2]);
CL_SET_ARG(clState->BranchBuffer[3]);

num = 0;
CL_NEXTKERNEL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[0]);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(tgt32);

// last to be set in driver-opencl.c

num = 0;
CL_NEXTKERNEL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[1]);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(tgt32);


num = 0;
CL_NEXTKERNEL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[2]);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(tgt32);


num = 0;
CL_NEXTKERNEL_SET_ARG(clState->States);
CL_SET_ARG(clState->BranchBuffer[3]);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(tgt32);

return(status);
}


Expand All @@ -1160,7 +1160,7 @@ static cl_int queue_equihash_kernel(_clState *clState, dev_blk_ctx *blk, __maybe
status = clEnqueueWriteBuffer(clState->commandQueue, clState->MidstateBuf, CL_TRUE, 0, sizeof(mid_hash), mid_hash, 0, NULL, NULL);
uint32_t dbg[2] = {0};
status |= clEnqueueWriteBuffer(clState->commandQueue, clState->padbuffer8, CL_TRUE, 0, sizeof(dbg), &dbg, 0, NULL, NULL);

cl_mem rowCounters[2] = {clState->buffer2, clState->buffer3};
for (int round = 0; round < PARAM_K; round++) {
size_t global_ws = RC_SIZE;
Expand All @@ -1173,7 +1173,7 @@ static cl_int queue_equihash_kernel(_clState *clState, dev_blk_ctx *blk, __maybe
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(clState->CLbuffer0);
status |= clEnqueueNDRangeKernel(clState->commandQueue, *kernel, 1, NULL, &global_ws, &local_ws, 0, NULL, NULL);

kernel = &clState->extra_kernels[1 + round];
if (!round) {
worksize = LOCAL_WORK_SIZE_ROUND0;
Expand Down Expand Up @@ -1295,9 +1295,9 @@ static algorithm_settings_t algos[] = {
{ "ethash-genoil", ALGO_ETHASH, "", (1ULL << 32), (1ULL << 32), 1, 0, 0, 0xFF, 0xFFFF000000000000ULL, 0x00000000UL, 0, 128, 0, ethash_regenhash, NULL, queue_ethash_kernel, gen_hash, append_ethash_compiler_options },

{ "cryptonight", ALGO_CRYPTONIGHT, "", (1ULL << 32), (1ULL << 32), (1ULL << 32), 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 6, 0, 0, cryptonight_regenhash, NULL, queue_cryptonight_kernel, gen_hash, NULL },

{ "equihash", ALGO_EQUIHASH, "", 1, (1ULL << 28), (1ULL << 28), 0, 0, 0x20000, 0xFFFF000000000000ULL, 0x00000000UL, 0, 128, 0, equihash_regenhash, NULL, queue_equihash_kernel, gen_hash, append_equihash_compiler_options },

// Terminator (do not remove)
{ NULL, ALGO_UNK, "", 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, NULL, NULL, NULL, NULL }
};
Expand Down
56 changes: 28 additions & 28 deletions driver-opencl.c
View file
Expand Up @@ -1211,7 +1211,7 @@ static void opencl_detect(void)
cgpu->algorithm = default_profile.algorithm;
add_cgpu(cgpu);
}

if(!init_sysfs_hwcontrols(nDevs) && !opt_noadl) {
init_adl(nDevs);
}
Expand Down Expand Up @@ -1424,14 +1424,14 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
uint8_t prev_hash[32];
size_t txns;
bool stale = true;

if (work->getwork_mode != GETWORK_MODE_STRATUM) {
cg_rlock(&work->pool->gbt_lock);
txns = work->pool->gbt_txns;
memcpy(prev_hash, work->pool->previousblockhash, 32);
cg_runlock(&work->pool->gbt_lock);
}

thrdata->res = realloc(thrdata->res, length);
sols_t *sols = (sols_t*) thrdata->res;
work->thr = thr;
Expand All @@ -1450,22 +1450,22 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
applog(LOG_ERR, "Error %d: Reading result buffer for ALGO_EQUIHASH failed. (clEnqueueReadBuffer)", status);
return -1;
}

// increase nonce
work->blk.nonce++;
if (work->getwork_mode == GETWORK_MODE_STRATUM)
*(uint16_t*)(work->equihash_data + 108 + strlen(work->nonce1) / 2) += 1;
else {
*(uint64_t*)(work->equihash_data + 108) += 1;

cg_rlock(&work->pool->gbt_lock);
stale = (work->pool->gbt_txns != txns) || (memcmp(prev_hash, work->pool->previousblockhash, 32) != 0);
cg_runlock(&work->pool->gbt_lock);
}
} while (!stale && (time(NULL) - t0) < 2);
return ret;
}

status = thrdata->queue_kernel_parameters(clState, &work->blk, globalThreads[0]);
if (unlikely(status != CL_SUCCESS)) {
if (status > 0)
Expand All @@ -1474,21 +1474,21 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
return -1;
}
// if (algorithm.type == ALGO_ETHASH) read lock gpu->eth_dag.lock has to be released

if(gpu->algorithm.type == ALGO_CRYPTONIGHT) {
mutex_lock(&work->pool->XMRGlobalNonceLock);
work->blk.nonce = work->pool->XMRGlobalNonce;
work->pool->XMRGlobalNonce += gpu->max_hashes;
mutex_unlock(&work->pool->XMRGlobalNonceLock);
}

if (clState->goffset)
p_global_work_offset = (size_t *)&work->blk.nonce;

if (gpu->algorithm.type == ALGO_CRYPTONIGHT) {
size_t GlobalThreads = *globalThreads, Nonce[2] = { (size_t)work->blk.nonce, 1}, gthreads[2] = { *globalThreads, 8 }, lthreads[2] = { *localThreads, 8 };
size_t BranchBufCount[4] = { 0, 0, 0, 0 };

for (int i = 0; i < 4; ++i) {
cl_uint zero = 0;

Expand All @@ -1499,63 +1499,63 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
return -1;
}
}

clFinish(clState->commandQueue);

// Main CN P0
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->kernel, 2, Nonce, gthreads, lthreads, 0, NULL, NULL);

if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error %d while attempting to enqueue kernel 0.", status);
return -1;
}

// Main CN P1
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[0], 1, p_global_work_offset, globalThreads, localThreads, 0, NULL, NULL);

if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error %d while attempting to enqueue kernel 1.", status);
return -1;
}

// Main CN P2
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[1], 2, Nonce, gthreads, lthreads, 0, NULL, NULL);

if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error %d while attempting to enqueue kernel 2.", status);
return -1;
}

// Read BranchBuf counters

for (int i = 0; i < 4; ++i) {
status = clEnqueueReadBuffer(clState->commandQueue, clState->BranchBuffer[i], CL_FALSE, sizeof(cl_uint) * GlobalThreads, sizeof(cl_uint), BranchBufCount + i, 0, NULL, NULL);

if(status != CL_SUCCESS) {
applog(LOG_ERR, "Error %d while attempting to read branch buffer counter %d.", status, i);
return(-1);
}
}

clFinish(clState->commandQueue);

for (int i = 0; i < 4; ++i) {
if(BranchBufCount[i]) {
cl_ulong tmp = BranchBufCount[i];

// Threads
status = clSetKernelArg(clState->extra_kernels[i + 2], 4, sizeof(cl_ulong), &tmp);

if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error %d while attempting to set argument 4 for kernel %d.", status, i + 2);
return -1;
}

// Make it a multiple of the local worksize (some drivers will otherwise shit a brick)
BranchBufCount[i] += (clState->wsize - (BranchBufCount[i] & (clState->wsize - 1)));

status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[i + 2], 1, p_global_work_offset, BranchBufCount + i, localThreads, 0, NULL, NULL);

if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error %d while attempting to enqueue kernel %d.", status, i + 2);
return -1;
Expand Down Expand Up @@ -1583,7 +1583,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
}
}
}

status = clEnqueueReadBuffer(clState->commandQueue, clState->outputBuffer, CL_FALSE, 0,
buffersize, thrdata->res, 0, NULL, NULL);
if (unlikely(status != CL_SUCCESS)) {
Expand Down

0 comments on commit ade8905

Please sign in to comment.