From d60fc1c5ca9d4ad41246494279043cc3824df217 Mon Sep 17 00:00:00 2001 From: Philipp Otterbein Date: Tue, 29 Nov 2016 18:44:45 +0100 Subject: [PATCH] ethash DAG fixes --- algorithm.c | 303 +++++++++++++++++++--------------------- algorithm.h | 1 - algorithm/ethash.c | 189 ++++++++++++------------- algorithm/ethash.h | 4 +- algorithm/ethgencache.c | 60 ++++---- driver-opencl.c | 9 ++ miner.h | 35 ++++- ocl.c | 40 +++--- ocl.h | 3 - sgminer.c | 182 ++++++++---------------- util.c | 12 +- 11 files changed, 395 insertions(+), 443 deletions(-) diff --git a/algorithm.c b/algorithm.c index 6d762a5f..cbbc62c2 100644 --- a/algorithm.c +++ b/algorithm.c @@ -8,6 +8,7 @@ */ #include "algorithm.h" +#include "findnonce.h" #include "sph/sph_sha2.h" #include "ocl.h" #include "ocl/build_kernel.h" @@ -118,13 +119,8 @@ static void append_scrypt_compiler_options(struct _build_kernel_data *data, stru strcat(data->binary_filename, buf); } -extern uint32_t EthereumEpochNumber; - static void append_ethash_compiler_options(struct _build_kernel_data *data, struct cgpu_info *cgpu, struct _algorithm_t *algorithm) { - //char buf[255]; - //sprintf(buf, " -D DAG_SIZE=%lluUL ", EthGetDAGSize(EthereumEpochNumber) / 128); - //strcat(data->compiler_options, buf); } static void append_neoscrypt_compiler_options(struct _build_kernel_data *data, struct cgpu_info *cgpu, struct _algorithm_t *algorithm) @@ -935,141 +931,154 @@ static cl_int queue_pluck_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_un static cl_int queue_blake_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; - cl_ulong le_target; - - le_target = *(cl_ulong *)(blk->work->device_target + 24); - flip80(clState->cldata, blk->work->data); - status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL, NULL); - - CL_SET_ARG(clState->outputBuffer); - CL_SET_ARG(blk->work->blk.ctx_a); - CL_SET_ARG(blk->work->blk.ctx_b); - CL_SET_ARG(blk->work->blk.ctx_c); - CL_SET_ARG(blk->work->blk.ctx_d); - CL_SET_ARG(blk->work->blk.ctx_e); - CL_SET_ARG(blk->work->blk.ctx_f); - CL_SET_ARG(blk->work->blk.ctx_g); - CL_SET_ARG(blk->work->blk.ctx_h); - - CL_SET_ARG(blk->work->blk.cty_a); - CL_SET_ARG(blk->work->blk.cty_b); - CL_SET_ARG(blk->work->blk.cty_c); + cl_kernel *kernel = &clState->kernel; + unsigned int num = 0; + cl_int status = 0; + cl_ulong le_target; - return status; + le_target = *(cl_ulong *)(blk->work->device_target + 24); + flip80(clState->cldata, blk->work->data); + status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL, NULL); + + CL_SET_ARG(clState->outputBuffer); + CL_SET_ARG(blk->work->blk.ctx_a); + CL_SET_ARG(blk->work->blk.ctx_b); + CL_SET_ARG(blk->work->blk.ctx_c); + CL_SET_ARG(blk->work->blk.ctx_d); + CL_SET_ARG(blk->work->blk.ctx_e); + CL_SET_ARG(blk->work->blk.ctx_f); + CL_SET_ARG(blk->work->blk.ctx_g); + CL_SET_ARG(blk->work->blk.ctx_h); + + CL_SET_ARG(blk->work->blk.cty_a); + CL_SET_ARG(blk->work->blk.cty_b); + CL_SET_ARG(blk->work->blk.cty_c); + + return status; } -extern cglock_t EthCacheLock[2]; -extern uint8_t* EthCache[2]; extern pthread_mutex_t eth_nonce_lock; extern uint32_t eth_nonce; +static const int eth_future_epochs = 6; +extern struct pool *currentpool; static cl_int queue_ethash_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_unused cl_uint threads) { - cl_kernel *kernel; - unsigned int num = 0; - cl_int status = 0; - cl_ulong le_target; - cl_uint HighNonce, Isolate = 0xFFFFFFFFUL; - cl_ulong DAGSize = EthGetDAGSize(blk->work->EpochNumber); - size_t DAGItems = (size_t) (DAGSize / 64); - - le_target = *(cl_ulong *)(blk->work->device_target + 24); - - // DO NOT flip80. - status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 32, blk->work->data, 0, NULL, NULL); - if (clState->EpochNumber != blk->work->EpochNumber) - { - clState->EpochNumber = blk->work->EpochNumber; - cl_ulong CacheSize = EthGetCacheSize(blk->work->EpochNumber); - cl_event DAGGenEvent; - - applog(LOG_DEBUG, "DAG being regenerated."); - if (clState->EthCache) - clReleaseMemObject(clState->EthCache); - if (clState->DAG) - clReleaseMemObject(clState->DAG); - - clState->DAG = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, DAGSize, NULL, &status); - if (status != CL_SUCCESS) - { - applog(LOG_ERR, "Error %d: Creating the DAG buffer.", status); - return(status); - } - - clState->EthCache = clCreateBuffer(clState->context, CL_MEM_READ_ONLY, CacheSize, NULL, &status); - - int idx = blk->work->EpochNumber % 2; - cg_ilock(&EthCacheLock[idx]); - bool update = (EthCache[idx] == NULL || *(uint32_t*) EthCache[idx] != blk->work->EpochNumber); - if (update) - { - cg_ulock(&EthCacheLock[idx]); - EthCache[idx] = realloc(EthCache[idx], sizeof(uint8_t) * CacheSize + 64); - *(uint32_t*) EthCache[idx] = blk->work->EpochNumber; - EthGenerateCache(EthCache[idx] + 64, blk->work->seedhash, CacheSize); - } - else - cg_dlock(&EthCacheLock[idx]); - - if (status == CL_SUCCESS) - status = clEnqueueWriteBuffer(clState->commandQueue, clState->EthCache, true, 0, sizeof(cl_uchar) * CacheSize, EthCache[idx] + 64, 0, NULL, NULL); - - if (update) - cg_wunlock(&EthCacheLock[idx]); - else - cg_runlock(&EthCacheLock[idx]); - - if (status != CL_SUCCESS) - { - applog(LOG_ERR, "Error %d: Creating the cache buffer and/or writing to it.", status); - return(status); - } - - // enqueue DAG gen kernel - kernel = &clState->GenerateDAG; - - cl_uint zero = 0; - cl_uint CacheSize64 = CacheSize / 64; - - CL_SET_ARG(zero); - CL_SET_ARG(clState->EthCache); - CL_SET_ARG(clState->DAG); - CL_SET_ARG(CacheSize64); - CL_SET_ARG(Isolate); - - status |= clEnqueueNDRangeKernel(clState->commandQueue, clState->GenerateDAG, 1, NULL, &DAGItems, NULL, 0, NULL, &DAGGenEvent); - status |= clWaitForEvents(1, &DAGGenEvent); - clReleaseEvent(DAGGenEvent); - - if(status != CL_SUCCESS) - { - applog(LOG_ERR, "Error %d: Setting args for the DAG kernel and/or executing it.", status); - return(status); - } - } - - mutex_lock(ð_nonce_lock); - HighNonce = eth_nonce++; - blk->work->Nonce = (cl_ulong) HighNonce << 32; - mutex_unlock(ð_nonce_lock); - - num = 0; - kernel = &clState->kernel; - - // Not nodes now (64 bytes), but DAG entries (128 bytes) - cl_uint ItemsArg = DAGItems >> 1; - - CL_SET_ARG(clState->outputBuffer); - CL_SET_ARG(clState->CLbuffer0); - CL_SET_ARG(clState->DAG); - CL_SET_ARG(ItemsArg); - CL_SET_ARG(blk->work->Nonce); - CL_SET_ARG(le_target); - CL_SET_ARG(Isolate); - - return(status); + struct pool *pool = blk->work->pool; + cl_kernel *kernel; + unsigned int num = 0; + cl_int status = 0; + cl_ulong le_target; + cl_uint HighNonce, Isolate = UINT32_MAX; + + eth_dag_t *dag = &blk->work->thr->cgpu->eth_dag; + cg_ilock(&dag->lock); + cg_ilock(&pool->data_lock); + if (pool->eth_cache.disabled || pool->eth_cache.dag_cache == NULL || pool->algorithm.type != ALGO_ETHASH || blk->work->eth_epoch == UINT32_MAX) { + blk->work->pool = currentpool; + cg_iunlock(&pool->data_lock); + cg_iunlock(&dag->lock); + cgsleep_ms(200); + applog(LOG_DEBUG, "THR[%d]: stop ETHASH mining", blk->work->thr_id); + return 1; + } + if (dag->current_epoch != blk->work->eth_epoch) { + cl_ulong CacheSize = EthGetCacheSize(blk->work->eth_epoch); + 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); + 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); + } + 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; + } + } + else + cg_dlock(&pool->data_lock); + + applog(LOG_DEBUG, "DAG being regenerated."); + cl_mem eth_cache = clCreateBuffer(clState->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, CacheSize, pool->eth_cache.dag_cache, &status); + cg_runlock(&pool->data_lock); + if (status != CL_SUCCESS) { + clReleaseMemObject(eth_cache); + cg_wunlock(&dag->lock); + applog(LOG_ERR, "Error %d: Creating the ethash cache buffer failed.", status); + return status; + } + + // enqueue DAG gen kernel + kernel = &clState->GenerateDAG; + + cl_uint zero = 0; + cl_uint CacheSize64 = CacheSize / 64; + + CL_SET_ARG(zero); + CL_SET_ARG(eth_cache); + CL_SET_ARG(dag->dag_buffer); + CL_SET_ARG(CacheSize64); + CL_SET_ARG(Isolate); + + cl_ulong DAGSize = EthGetDAGSize(blk->work->eth_epoch); + size_t DAGItems = (size_t) (DAGSize / 64); + status |= clEnqueueNDRangeKernel(clState->commandQueue, clState->GenerateDAG, 1, NULL, &DAGItems, NULL, 0, NULL, NULL); + clFinish(clState->commandQueue); + + clReleaseMemObject(eth_cache); + if (status != CL_SUCCESS) { + cg_wunlock(&dag->lock); + applog(LOG_ERR, "Error %d: Setting args for the DAG kernel and/or executing it.", status); + return status; + } + dag->current_epoch = blk->work->eth_epoch; + cg_dwlock(&dag->lock); + } + else { + cg_dlock(&dag->lock); + cg_iunlock(&pool->data_lock); + } + + memcpy(&le_target, blk->work->device_target + 24, 8); + mutex_lock(ð_nonce_lock); + HighNonce = eth_nonce++; + blk->work->Nonce = (cl_ulong) HighNonce << 32; + mutex_unlock(ð_nonce_lock); + + num = 0; + kernel = &clState->kernel; + + // Not nodes now (64 bytes), but DAG entries (128 bytes) + cl_ulong DAGSize = EthGetDAGSize(blk->work->eth_epoch); + cl_uint ItemsArg = DAGSize / 128; + + // 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); + CL_SET_ARG(ItemsArg); + CL_SET_ARG(blk->work->Nonce); + CL_SET_ARG(le_target); + CL_SET_ARG(Isolate); + + if (status != CL_SUCCESS) + cg_runlock(&dag->lock); + return status; } static void append_equihash_compiler_options(struct _build_kernel_data *data, struct cgpu_info *cgpu, struct _algorithm_t *algorithm) @@ -1269,12 +1278,10 @@ void copy_algorithm_settings(algorithm_t* dest, const char* algo) } } -static const char *lookup_algorithm_alias(const char *lookup_alias, uint8_t *nfactor, uint8_t *kfactor) +static const char *lookup_algorithm_alias(const char *lookup_alias, uint8_t *nfactor) { #define ALGO_ALIAS_NF(alias, name, nf) \ if (strcasecmp(alias, lookup_alias) == 0) { *nfactor = nf; return name; } -#define ALGO_ALIAS_NFK(alias, name, nf, kf) \ - if (strcasecmp(alias, lookup_alias) == 0) { *kfactor = kf, *nfactor = nf; return name; } #define ALGO_ALIAS(alias, name) \ if (strcasecmp(alias, lookup_alias) == 0) return name; @@ -1302,11 +1309,10 @@ static const char *lookup_algorithm_alias(const char *lookup_alias, uint8_t *nfa ALGO_ALIAS("lyra2v2", "lyra2rev2"); ALGO_ALIAS("blakecoin", "blake256r8"); ALGO_ALIAS("blake", "blake256r14"); - ALGO_ALIAS_NFK("zcash", "equihash", 200, 9); + ALGO_ALIAS("zcash", "equihash"); #undef ALGO_ALIAS #undef ALGO_ALIAS_NF -#undef ALGO_ALIAS_NFK return NULL; } @@ -1317,16 +1323,11 @@ void set_algorithm(algorithm_t* algo, const char* newname_alias) //load previous algorithm nfactor in case nfactor was applied before algorithm... or default to 10 uint8_t old_nfactor = ((algo->nfactor) ? algo->nfactor : 0); - - //load previous algorithm kfactor in case kfactor was applied before algorithm... or default to 9 - uint8_t old_kfactor = ((algo->kfactor) ? algo->kfactor : 0); - //load previous kernel file name if was applied before algorithm... const char *kernelfile = algo->kernelfile; uint8_t nfactor = 10; - uint8_t kfactor = 9; - if (!(newname = lookup_algorithm_alias(newname_alias, &nfactor, &kfactor))) { + if (!(newname = lookup_algorithm_alias(newname_alias, &nfactor))) { newname = newname_alias; } @@ -1339,13 +1340,6 @@ void set_algorithm(algorithm_t* algo, const char* newname_alias) set_algorithm_nfactor(algo, nfactor); - // use old kfactor if it was previously set and is different than the one set by alias - if ((old_kfactor > 0) && (old_kfactor != kfactor)) { - kfactor = old_kfactor; - } - - set_algorithm_kfactor(algo, kfactor); - //reapply kernelfile if was set if (!empty_string(kernelfile)) { algo->kernelfile = kernelfile; @@ -1377,12 +1371,7 @@ void set_algorithm_nfactor(algorithm_t* algo, const uint8_t nfactor) } } -void set_algorithm_kfactor(algorithm_t* algo, const uint8_t kfactor) -{ - algo->kfactor = kfactor; -} - bool cmp_algorithm(const algorithm_t* algo1, const algorithm_t* algo2) { - return (!safe_cmp(algo1->name, algo2->name) && !safe_cmp(algo1->kernelfile, algo2->kernelfile) && (algo1->nfactor == algo2->nfactor) && (algo1->kfactor == algo2->kfactor)); + return (!safe_cmp(algo1->name, algo2->name) && !safe_cmp(algo1->kernelfile, algo2->kernelfile) && (algo1->nfactor == algo2->nfactor)); } diff --git a/algorithm.h b/algorithm.h index 397937aa..98db4ee5 100644 --- a/algorithm.h +++ b/algorithm.h @@ -61,7 +61,6 @@ typedef struct _algorithm_t { const char *kernelfile; /* alternate kernel file */ uint32_t n; /* N (CPU/Memory tradeoff parameter) */ uint8_t nfactor; /* Factor of N above (n = 2^nfactor) */ - uint8_t kfactor; /* Factor of N above (n = 2^nfactor) */ double diff_multiplier1; double diff_multiplier2; double share_diff_multiplier; diff --git a/algorithm/ethash.c b/algorithm/ethash.c index 40fb8083..b9606d54 100644 --- a/algorithm/ethash.c +++ b/algorithm/ethash.c @@ -1,136 +1,127 @@ #include #include "config.h" -#include "miner.h" #include "algorithm/ethash.h" #include "algorithm/eth-sha3.h" -#define FNV_PRIME 0x01000193 +#define FNV_PRIME 0x01000193 -#define fnv(x, y) (((x) * FNV_PRIME) ^ (y)) -#define fnv_reduce(v) fnv(fnv(fnv((v)[0], (v)[1]), (v)[2]), (v)[3]) +#define fnv(x, y) (((x) * FNV_PRIME) ^ (y)) +#define fnv_reduce(v) fnv(fnv(fnv((v)[0], (v)[1]), (v)[2]), (v)[3]) #define ETHEREUM_EPOCH_LENGTH 30000UL -extern cglock_t EthCacheLock[2]; -extern uint8_t* EthCache[2]; typedef struct _DAG128 { - uint32_t Columns[32]; + uint32_t Columns[32]; } DAG128; typedef union _Node { - uint8_t bytes[16 * 4]; - uint32_t words[16]; - uint64_t double_words[16 / 2]; + uint8_t bytes[16 * 4]; + uint32_t words[16]; + uint64_t double_words[16 / 2]; } Node; uint32_t EthCalcEpochNumber(uint8_t *SeedHash) { - uint8_t TestSeedHash[32] = { 0 }; - - for(int Epoch = 0; Epoch < 2048; ++Epoch) - { - SHA3_256(TestSeedHash, TestSeedHash, 32); - if(!memcmp(TestSeedHash, SeedHash, 32)) return(Epoch + 1); - } - - applog(LOG_ERR, "Error on epoch calculation."); - - return(0UL); + uint8_t TestSeedHash[32] = { 0 }; + + for(int Epoch = 0; Epoch < 2048; ++Epoch) { + SHA3_256(TestSeedHash, TestSeedHash, 32); + if(!memcmp(TestSeedHash, SeedHash, 32)) return(Epoch + 1); + } + + applog(LOG_ERR, "Error on epoch calculation."); + + return 0UL; } Node CalcDAGItem(const Node *CacheInputNodes, uint32_t NodeCount, uint32_t NodeIdx) { - Node DAGNode = CacheInputNodes[NodeIdx % NodeCount]; - - DAGNode.words[0] ^= NodeIdx; + Node DAGNode = CacheInputNodes[NodeIdx % NodeCount]; + + DAGNode.words[0] ^= NodeIdx; - SHA3_512(DAGNode.bytes, DAGNode.bytes, sizeof(Node)); - - for(uint32_t i = 0; i < 256; ++i) - { - uint32_t parent_index = fnv(NodeIdx ^ i, DAGNode.words[i % 16]) % NodeCount; - Node const *parent = CacheInputNodes + parent_index; //&cache_nodes[parent_index]; - - for(int i = 0; i < 16; ++i) - { - DAGNode.words[i] *= FNV_PRIME; - DAGNode.words[i] ^= parent->words[i]; - } - } + SHA3_512(DAGNode.bytes, DAGNode.bytes, sizeof(Node)); + + for(uint32_t i = 0; i < 256; ++i) { + uint32_t parent_index = fnv(NodeIdx ^ i, DAGNode.words[i % 16]) % NodeCount; + Node const *parent = CacheInputNodes + parent_index; //&cache_nodes[parent_index]; + + for(int i = 0; i < 16; ++i) { + DAGNode.words[i] *= FNV_PRIME; + DAGNode.words[i] ^= parent->words[i]; + } + } - SHA3_512(DAGNode.bytes, DAGNode.bytes, sizeof(Node)); - - return(DAGNode); + SHA3_512(DAGNode.bytes, DAGNode.bytes, sizeof(Node)); + + return DAGNode; } // OutHash & MixHash MUST have 32 bytes allocated (at least) void LightEthash(uint8_t *restrict OutHash, uint8_t *restrict MixHash, const uint8_t *restrict HeaderPoWHash, const Node *Cache, const uint64_t EpochNumber, const uint64_t Nonce) { - uint32_t MixState[32], TmpBuf[24], NodeCount = EthGetCacheSize(EpochNumber) / sizeof(Node); - uint64_t DagSize; - Node *EthCache = Cache; - - // Initial hash - append nonce to header PoW hash and - // run it through SHA3 - this becomes the initial value - // for the mixing state buffer. The init value is used - // later for the final hash, and is therefore saved. - memcpy(TmpBuf, HeaderPoWHash, 32UL); - memcpy(TmpBuf + 8UL, &Nonce, 8UL); - sha3_512((uint8_t *)TmpBuf, 64UL, (uint8_t *)TmpBuf, 40UL); - - memcpy(MixState, TmpBuf, 64UL); - - // The other half of the state is filled by simply - // duplicating the first half of its initial value. - memcpy(MixState + 16UL, MixState, 64UL); - - DagSize = EthGetDAGSize(EpochNumber) / (sizeof(Node) << 1); - - // Main mix of Ethash - for(uint32_t i = 0, Init0 = MixState[0], MixValue = MixState[0]; i < 64; ++i) - { - uint32_t row = fnv(Init0 ^ i, MixValue) % DagSize; - Node DAGSliceNodes[2]; - DAGSliceNodes[0] = CalcDAGItem(EthCache, NodeCount, row << 1); - DAGSliceNodes[1] = CalcDAGItem(EthCache, NodeCount, (row << 1) + 1); - DAG128 *DAGSlice = (DAG128 *)DAGSliceNodes; - - for(uint32_t col = 0; col < 32; ++col) - { - MixState[col] = fnv(MixState[col], DAGSlice->Columns[col]); - MixValue = col == ((i + 1) & 0x1F) ? MixState[col] : MixValue; - } - } - - // The reducing of the mix state directly into where - // it will be hashed to produce the final hash. Note - // that the initial hash is still in the first 64 - // bytes of TmpBuf - we're appending the mix hash. - for(int i = 0; i < 8; ++i) TmpBuf[i + 16] = fnv_reduce(MixState + (i << 2)); - - memcpy(MixHash, TmpBuf + 16, 32UL); - - // Hash the initial hash and the mix hash concatenated - // to get the final proof-of-work hash that is our output. - sha3_256(OutHash, 32UL, (uint8_t *)TmpBuf, 96UL); + uint32_t MixState[32], TmpBuf[24], NodeCount = EthGetCacheSize(EpochNumber) / sizeof(Node); + uint64_t DagSize; + + // Initial hash - append nonce to header PoW hash and + // run it through SHA3 - this becomes the initial value + // for the mixing state buffer. The init value is used + // later for the final hash, and is therefore saved. + memcpy(TmpBuf, HeaderPoWHash, 32UL); + memcpy(TmpBuf + 8UL, &Nonce, 8UL); + sha3_512((uint8_t *)TmpBuf, 64UL, (uint8_t *)TmpBuf, 40UL); + + memcpy(MixState, TmpBuf, 64UL); + + // The other half of the state is filled by simply + // duplicating the first half of its initial value. + memcpy(MixState + 16UL, MixState, 64UL); + + DagSize = EthGetDAGSize(EpochNumber) / (sizeof(Node) << 1); + + // Main mix of Ethash + for(uint32_t i = 0, Init0 = MixState[0], MixValue = MixState[0]; i < 64; ++i) { + uint32_t row = fnv(Init0 ^ i, MixValue) % DagSize; + Node DAGSliceNodes[2]; + DAGSliceNodes[0] = CalcDAGItem(Cache, NodeCount, row << 1); + DAGSliceNodes[1] = CalcDAGItem(Cache, NodeCount, (row << 1) + 1); + DAG128 *DAGSlice = (DAG128 *)DAGSliceNodes; + + for(uint32_t col = 0; col < 32; ++col) { + MixState[col] = fnv(MixState[col], DAGSlice->Columns[col]); + MixValue = col == ((i + 1) & 0x1F) ? MixState[col] : MixValue; + } + } + + // The reducing of the mix state directly into where + // it will be hashed to produce the final hash. Note + // that the initial hash is still in the first 64 + // bytes of TmpBuf - we're appending the mix hash. + for(int i = 0; i < 8; ++i) + TmpBuf[i + 16] = fnv_reduce(MixState + (i << 2)); + + memcpy(MixHash, TmpBuf + 16, 32UL); + + // Hash the initial hash and the mix hash concatenated + // to get the final proof-of-work hash that is our output. + sha3_256(OutHash, 32UL, (uint8_t *)TmpBuf, 96UL); } void ethash_regenhash(struct work *work) { - work->Nonce += *((uint32_t *)(work->data + 32)); - applog(LOG_DEBUG, "Regenhash: First qword of input: 0x%016llX.", work->Nonce); - int idx = work->EpochNumber % 2; - cg_rlock(&EthCacheLock[idx]); - LightEthash(work->hash, work->mixhash, work->data, EthCache[idx] + 64, work->EpochNumber, work->Nonce); - cg_runlock(&EthCacheLock[idx]); - - char *DbgHash = bin2hex(work->hash, 32); - - applog(LOG_DEBUG, "Regenhash result: %s.", DbgHash); - applog(LOG_DEBUG, "Last ulong: 0x%016llX.", bswap_64(*((uint64_t *)(work->hash + 0)))); - free(DbgHash); + work->Nonce += *((uint32_t *)(work->data + 32)); + applog(LOG_DEBUG, "Regenhash: First qword of input: 0x%016llX.", work->Nonce); + cg_rlock(&work->pool->data_lock); + LightEthash(work->hash, work->mixhash, work->data, work->pool->eth_cache.dag_cache, work->eth_epoch, work->Nonce); + cg_runlock(&work->pool->data_lock); + + char *DbgHash = bin2hex(work->hash, 32); + + applog(LOG_DEBUG, "Regenhash result: %s.", DbgHash); + applog(LOG_DEBUG, "Last ulong: 0x%016llX.", bswap_64(*((uint64_t *)(work->hash + 0)))); + free(DbgHash); } diff --git a/algorithm/ethash.h b/algorithm/ethash.h index 869b1cda..36556e06 100644 --- a/algorithm/ethash.h +++ b/algorithm/ethash.h @@ -2,6 +2,8 @@ #define __ETHASH_H #include +#include "miner.h" + static const uint64_t dag_sizes[2048] = { @@ -752,7 +754,7 @@ static const uint64_t cache_sizes[2048] = #define EthGetDAGSize(EpochNum) dag_sizes[EpochNum] struct work; -void EthGenerateCache(void *cache_nodes_in, const uint8_t *seedhash, uint64_t cache_size); +void eth_gen_cache(struct pool *); void ethash_regenhash(struct work *work); uint32_t EthCalcEpochNumber(uint8_t *SeedHash); diff --git a/algorithm/ethgencache.c b/algorithm/ethgencache.c index f60434a9..96ae1ce3 100644 --- a/algorithm/ethgencache.c +++ b/algorithm/ethgencache.c @@ -1,42 +1,46 @@ #include +#include "miner.h" #include "sph/sph_keccak.h" #include "algorithm/ethash.h" #include "algorithm/eth-sha3.h" typedef union node { - uint8_t bytes[16 * 4]; - uint32_t words[16]; - uint64_t double_words[16 / 2]; + uint8_t bytes[16 * 4]; + uint32_t words[16]; + uint64_t double_words[16 / 2]; } node; // Output (cache_nodes) MUST have at least cache_size bytes -void EthGenerateCache(void *cache_nodes_in, const uint8_t *seedhash, uint64_t cache_size) +static void EthGenerateCache(uint8_t *cache_nodes_in, const uint8_t *seedhash, uint64_t cache_size) { - uint32_t const num_nodes = (uint32_t)(cache_size / sizeof(node)); - node *cache_nodes = (node *)cache_nodes_in; - - SHA3_512(cache_nodes[0].bytes, seedhash, 32); - - for(uint32_t i = 1; i != num_nodes; ++i) - { - SHA3_512(cache_nodes[i].bytes, cache_nodes[i - 1].bytes, 64); - } + uint32_t const num_nodes = (uint32_t)(cache_size / sizeof(node)); + node *cache_nodes = (node *)cache_nodes_in; + + SHA3_512(cache_nodes[0].bytes, seedhash, 32); + + for(uint32_t i = 1; i != num_nodes; ++i) { + SHA3_512(cache_nodes[i].bytes, cache_nodes[i - 1].bytes, 64); + } - for(uint32_t j = 0; j < 3; j++) // this one can be unrolled entirely, ETHASH_CACHE_ROUNDS is constant - { - for(uint32_t i = 0; i != num_nodes; i++) - { - uint32_t const idx = cache_nodes[i].words[0] % num_nodes; - node data; - data = cache_nodes[(num_nodes - 1 + i) % num_nodes]; - for(uint32_t w = 0; w != 16; ++w) // this one can be unrolled entirely as well - { - data.words[w] ^= cache_nodes[idx].words[w]; - } - - SHA3_512(cache_nodes[i].bytes, data.bytes, sizeof(data)); - } - } + for(uint32_t j = 0; j < 3; j++) { // this one can be unrolled entirely, ETHASH_CACHE_ROUNDS is constant + for(uint32_t i = 0; i != num_nodes; i++) { + uint32_t const idx = cache_nodes[i].words[0] % num_nodes; + node data; + data = cache_nodes[(num_nodes - 1 + i) % num_nodes]; + for(uint32_t w = 0; w != 16; ++w) { // this one can be unrolled entirely as well + data.words[w] ^= cache_nodes[idx].words[w]; + } + + SHA3_512(cache_nodes[i].bytes, data.bytes, sizeof(data)); + } + } +} + + +void eth_gen_cache(struct pool *pool) { + size_t cache_size = EthGetCacheSize(pool->eth_cache.current_epoch); + pool->eth_cache.dag_cache = realloc(pool->eth_cache.dag_cache, cache_size); + EthGenerateCache(pool->eth_cache.dag_cache, pool->eth_cache.seed_hash, cache_size); } diff --git a/driver-opencl.c b/driver-opencl.c index 5e1dff47..9d504760 100644 --- a/driver-opencl.c +++ b/driver-opencl.c @@ -1469,9 +1469,12 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work, status = thrdata->queue_kernel_parameters(clState, &work->blk, globalThreads[0]); if (unlikely(status != CL_SUCCESS)) { + if (status > 0) + return 0; applog(LOG_ERR, "Error: clSetKernelArg of all params failed."); return -1; } + // if (algorithm.type == ALGO_ETHASH) read lock gpu->eth_dag.lock has to be released if (clState->goffset) p_global_work_offset = (size_t *)&work->blk.nonce; @@ -1479,6 +1482,8 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work, status = clEnqueueNDRangeKernel(clState->commandQueue, clState->kernel, 1, p_global_work_offset, globalThreads, localThreads, 0, NULL, NULL); if (unlikely(status != CL_SUCCESS)) { + if (work->pool->algorithm.type == ALGO_ETHASH) + cg_runlock(&gpu->eth_dag.lock); applog(LOG_ERR, "Error %d: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)", status); return -1; } @@ -1495,6 +1500,8 @@ 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)) { + if (work->pool->algorithm.type == ALGO_ETHASH) + cg_runlock(&gpu->eth_dag.lock); applog(LOG_ERR, "Error: clEnqueueReadBuffer failed error %d. (clEnqueueReadBuffer)", status); return -1; } @@ -1506,6 +1513,8 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work, /* This finish flushes the readbuffer set with CL_FALSE in clEnqueueReadBuffer */ clFinish(clState->commandQueue); + if (work->pool->algorithm.type == ALGO_ETHASH) + cg_runlock(&gpu->eth_dag.lock); /* found entry is used as a counter to say how many nonces exist */ if (thrdata->res[found]) { diff --git a/miner.h b/miner.h index 2d2b0373..5743f927 100644 --- a/miner.h +++ b/miner.h @@ -527,8 +527,7 @@ struct sgminer_pool_stats { uint64_t net_bytes_received; }; -typedef struct _gpu_sysfs_info -{ +typedef struct _gpu_sysfs_info { char *HWMonPath; uint32_t MinFanSpeed; uint32_t MaxFanSpeed; @@ -539,6 +538,24 @@ typedef struct _gpu_sysfs_info float LastTemp; } gpu_sysfs_info; +struct _eth_dag_t; +typedef struct _eth_cache_t { + uint8_t seed_hash[32]; + uint8_t *dag_cache; + struct _eth_dag_t **dags; + uint32_t current_epoch; + uint32_t nDevs; + bool disabled; +} eth_cache_t; + +typedef struct _eth_dag_t { + cglock_t lock; + cl_mem dag_buffer; + struct pool *pool; + uint32_t current_epoch; + uint32_t max_epoch; +} eth_dag_t; + struct cgpu_info { int sgminer_id; struct device_drv *drv; @@ -625,6 +642,7 @@ struct cgpu_info { int dev_throttle_count; struct sgminer_stats sgminer_stats; + eth_dag_t eth_dag; bool shutdown; @@ -851,6 +869,7 @@ extern void api_initlock(void *lock, enum cglock_typ typ, const char *file, cons #define cglock_init(_lock) _cglock_init(_lock, __FILE__, __func__, __LINE__) #define cg_rlock(_lock) _cg_rlock(_lock, __FILE__, __func__, __LINE__) #define cg_ilock(_lock) _cg_ilock(_lock, __FILE__, __func__, __LINE__) +#define cg_iunlock(_lock) _cg_iunlock(_lock, __FILE__, __func__, __LINE__) #define cg_ulock(_lock) _cg_ulock(_lock, __FILE__, __func__, __LINE__) #define cg_wlock(_lock) _cg_wlock(_lock, __FILE__, __func__, __LINE__) #define cg_dwlock(_lock) _cg_dwlock(_lock, __FILE__, __func__, __LINE__) @@ -995,6 +1014,12 @@ static inline void _cg_ilock(cglock_t *lock, const char *file, const char *func, _mutex_lock(&lock->mutex, file, func, line); } +/* Unlock intermediate lock - behaves like a mutex. */ +static inline void _cg_iunlock(cglock_t *lock, const char *file, const char *func, const int line) +{ + _mutex_unlock_noyield(&lock->mutex, file, func, line); +} + /* Upgrade intermediate variant to a write lock */ static inline void _cg_ulock(cglock_t *lock, const char *file, const char *func, const int line) { @@ -1318,8 +1343,7 @@ struct pool { int quota_gcd; int quota_used; int works; - uint8_t SeedHash[32]; - uint32_t EpochNumber; + eth_cache_t eth_cache; uint8_t Target[32]; uint8_t EthWork[32]; uint8_t NetDiff[32]; @@ -1468,7 +1492,6 @@ struct work { unsigned char data[168]; unsigned char midstate[32]; unsigned char target[32]; - unsigned char seedhash[32]; unsigned char hash[32]; unsigned char mixhash[32]; @@ -1477,7 +1500,7 @@ struct work { double share_diff; double network_diff; - uint32_t EpochNumber; + uint32_t eth_epoch; uint64_t Nonce; unsigned char equihash_data[1487]; diff --git a/ocl.c b/ocl.c index 3f5e60cf..324eda3c 100644 --- a/ocl.c +++ b/ocl.c @@ -813,10 +813,6 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg return NULL; } - // Load kernels - applog(LOG_NOTICE, "Initialising kernel %s with params N=%d, K=%d", - filename, algorithm->nfactor, algorithm->kfactor); - clState->n_extra_kernels = algorithm->n_extra_kernels; if (clState->n_extra_kernels > 0) { unsigned int i; @@ -836,22 +832,30 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg } - if(algorithm->type == ALGO_ETHASH) - { - clState->GenerateDAG = clCreateKernel(clState->program, "GenerateDAG", &status); + if (algorithm->type == ALGO_ETHASH) { + clState->GenerateDAG = clCreateKernel(clState->program, "GenerateDAG", &status); - if(status != CL_SUCCESS) - { - applog(LOG_ERR, "Error %d while creating DAG generation kernel.", status); - return(NULL); - } + if (status != CL_SUCCESS) { + applog(LOG_ERR, "Error %d while creating DAG generation kernel.", status); + return NULL; + } } size_t bufsize; size_t buf1size; size_t buf3size; size_t buf2size; - size_t readbufsize = (algorithm->type == ALGO_CRE) ? 168 : 128; + size_t readbufsize; + switch (algorithm->type) { + case ALGO_CRE: + readbufsize = 168; + break; + case ALGO_ETHASH: + readbufsize = 32; + break; + default: + readbufsize = 128; + } if (algorithm->rw_buffer_size < 0) { // calc buffer size for neoscrypt @@ -974,11 +978,6 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg } } - if(algorithm->type == ALGO_ETHASH) - { - readbufsize = 32UL; - clState->DAG = clState->EthCache = NULL; - } applog(LOG_DEBUG, "Using read buffer sized %lu", (unsigned long)readbufsize); clState->CLbuffer0 = clCreateBuffer(clState->context, CL_MEM_READ_ONLY, readbufsize, NULL, &status); if (status != CL_SUCCESS) { @@ -988,8 +987,9 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg clState->devid = cgpu->device_id; - applog(LOG_DEBUG, "Using output buffer sized %lu", BUFFERSIZE); - clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_WRITE_ONLY, BUFFERSIZE, NULL, &status); + size_t buffersize = MAX(sizeof(sols_t), BUFFERSIZE); + applog(LOG_DEBUG, "Using output buffer sized %lu", buffersize); + clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_WRITE_ONLY, buffersize, NULL, &status); if (status != CL_SUCCESS) { applog(LOG_ERR, "Error %d: clCreateBuffer (outputBuffer)", status); return NULL; diff --git a/ocl.h b/ocl.h index a07c1f75..3ad6f6ef 100644 --- a/ocl.h +++ b/ocl.h @@ -22,9 +22,6 @@ typedef struct __clState { cl_program program; cl_mem outputBuffer; cl_mem CLbuffer0; - cl_mem DAG; - cl_mem EthCache; - cl_uint EpochNumber; cl_mem MidstateBuf; cl_mem padbuffer8; cl_mem buffer1; diff --git a/sgminer.c b/sgminer.c index b202bfd6..73d81dca 100644 --- a/sgminer.c +++ b/sgminer.c @@ -164,9 +164,6 @@ static bool opt_morenotices; uint8_t entropy[32]; uint32_t eth_nonce; pthread_mutex_t eth_nonce_lock; -uint32_t EthereumEpochNumber = 0; -cglock_t EthCacheLock[2]; -uint8_t* EthCache[2]; bool opt_autofan; bool opt_autoengine; bool opt_noadl; @@ -268,7 +265,7 @@ unsigned int local_work; unsigned int total_go, total_ro; struct pool **pools; -static struct pool *currentpool = NULL; +struct pool *currentpool = NULL; struct strategies strategies[] = { { "Failover" }, @@ -348,6 +345,36 @@ struct schedtime schedstart; struct schedtime schedstop; bool sched_paused; +static void set_current_pool(struct pool *pool) { + applog(LOG_DEBUG, "Trying to set current pool..."); + bool free_dag = (currentpool != NULL && currentpool->algorithm.type == ALGO_ETHASH && pool->algorithm.type != ALGO_ETHASH); + if (free_dag) { + cg_wlock(¤tpool->data_lock); + eth_cache_t *cache = ¤tpool->eth_cache; + for (int i = 0; i < cache->nDevs; i++) { + cg_wlock(&cache->dags[i]->lock); + if (cache->dags[i]->dag_buffer != NULL) + clReleaseMemObject(cache->dags[i]->dag_buffer); + cache->dags[i]->dag_buffer = NULL; + cache->dags[i]->pool = NULL; + cache->dags[i]->max_epoch = 0; + cache->dags[i]->current_epoch = UINT32_MAX; + cg_wunlock(&cache->dags[i]->lock); + } + free(cache->dags); + cache->dags = NULL; + cache->nDevs = 0; + cache->disabled = true; + cg_wunlock(¤tpool->data_lock); + } + currentpool = pool; + + cg_wlock(¤tpool->data_lock); + if (currentpool->algorithm.type == ALGO_ETHASH) + currentpool->eth_cache.disabled = false; + cg_wunlock(¤tpool->data_lock); +} + static bool time_before(struct tm *tm1, struct tm *tm2) { if (tm1->tm_hour < tm2->tm_hour) @@ -2122,9 +2149,8 @@ static double get_work_blockdiff(const struct work *work) diff64 = bswap_64(((uint64_t)(be32toh(*((uint32_t *)(work->data + 72))) & 0xFFFFFF00)) << 8); numerator = (double)work->pool->algorithm.diff_numerator; } - if(work->pool->algorithm.type == ALGO_ETHASH) - { - return(work->network_diff); + if (work->pool->algorithm.type == ALGO_ETHASH) { + return work->network_diff; } else { uint8_t pow = work->data[72]; @@ -2483,13 +2509,13 @@ static bool work_decode(struct pool *pool, struct work *work, json_t *val) } -bool parse_diff_ethash(char* Target, char* TgtStr); +bool parse_diff_ethash(char* Target, const char* TgtStr); static bool work_decode_eth(struct pool *pool, struct work *work, json_t *val, json_t *ethval2) { int i; bool ret = false; uint8_t EthWork[32], SeedHash[32], Target[32]; - char *EthWorkStr, *SeedHashStr, *TgtStr, *BlockHeightStr, *NetDiffStr, FinalNetDiffStr[65]; + const char *EthWorkStr, *SeedHashStr, *TgtStr, *BlockHeightStr, *NetDiffStr, FinalNetDiffStr[65]; cgtime(&pool->tv_lastwork); @@ -2554,18 +2580,24 @@ static bool work_decode_eth(struct pool *pool, struct work *work, json_t *val, j else if(!hex2bin(FinalNetDiffStr, NetDiffStr + 2, 32UL)) return(false); */ - if (memcmp(pool->SeedHash, SeedHash, 32)) { - pool->EpochNumber = EthCalcEpochNumber(SeedHash); - memcpy(pool->SeedHash, SeedHash, 32); + cg_ilock(&pool->data_lock); + if (memcmp(pool->eth_cache.seed_hash, SeedHash, 32)) { + cg_ulock(&pool->data_lock); + pool->eth_cache.current_epoch = EthCalcEpochNumber(SeedHash); + memcpy(pool->eth_cache.seed_hash, SeedHash, 32); + eth_gen_cache(pool); + cg_dwlock(&pool->data_lock); } + else + cg_dlock(&pool->data_lock); + //work->height = strtoul(BlockHeightStr + 2, NULL, 16) / 30000UL; + work->eth_epoch = pool->eth_cache.current_epoch; + cg_runlock(&pool->data_lock); memcpy(work->data, EthWork, 32); - memcpy(work->seedhash, pool->SeedHash, 32); swab256(work->target, Target); //work->network_diff = eth2pow256 / le256todouble(FinalNetDiffStr); - //work->EpochNumber = strtoul(BlockHeightStr + 2, NULL, 16) / 30000UL; - work->EpochNumber = pool->EpochNumber; cgtime(&work->tv_staged); ret = true; @@ -3478,8 +3510,9 @@ static bool submit_upstream_work(struct work *work, CURL *curl, char *curl_err_s return rc; } -const char eth_getwork_rpc[] = "{\"jsonrpc\":\"2.0\",\"method\":\"eth_getWork\",\"params\":[],\"id\":1}"; -const char eth_gethighestblock_rpc[] = "{\"jsonrpc\":\"2.0\",\"method\":\"eth_getBlockByNumber\",\"params\":[\"latest\", false],\"id\":1}"; + +char eth_getwork_rpc[] = "{\"jsonrpc\":\"2.0\",\"method\":\"eth_getWork\",\"params\":[],\"id\":1}"; +char eth_gethighestblock_rpc[] = "{\"jsonrpc\":\"2.0\",\"method\":\"eth_getBlockByNumber\",\"params\":[\"latest\", false],\"id\":1}"; static bool get_upstream_work(struct work *work, CURL *curl, char *curl_err_str) { @@ -4436,7 +4469,7 @@ void __switch_pools(struct pool *selected, bool saveprio) break; } - currentpool = pools[pool_no]; + set_current_pool(pools[pool_no]); pool = currentpool; on_backup_pool = pool->backup; cg_wunlock(&control_lock); @@ -6481,10 +6514,9 @@ static void gen_stratum_work_eth(struct pool *pool, struct work *work) applog(LOG_DEBUG, "[THR%d] gen_stratum_work() - algorithm = %s", work->thr_id, pool->algorithm.name); cg_rlock(&pool->data_lock); - work->EpochNumber = pool->EpochNumber; + work->eth_epoch = pool->eth_cache.current_epoch; work->job_id = strdup(pool->swork.job_id); memcpy(work->data, pool->EthWork, 32); - memcpy(work->seedhash, pool->SeedHash, 32); memcpy(work->target, pool->Target, 32); work->sdiff = pool->swork.diff; work->work_difficulty = pool->swork.diff; @@ -6726,110 +6758,13 @@ static void apply_initial_gpu_settings(struct pool *pool) rd_lock(&mining_thr_lock); apply_switcher_options(options, pool); - -/* - //reset devices - opt_devs_enabled = 0; - for (i = 0; i < MAX_DEVICES; i++) - devices_enabled[i] = false; - - //assign pool devices if any - if(!empty_string((opt = get_pool_setting(pool->devices, ((!empty_string(default_profile.devices))?default_profile.devices:"all"))))) { - set_devices((char *)opt); - } - - //lookup gap - if(!empty_string((opt = get_pool_setting(pool->lookup_gap, default_profile.lookup_gap)))) - set_lookup_gap((char *)opt); - - //set intensity - if(!empty_string((opt = get_pool_setting(pool->rawintensity, default_profile.rawintensity)))) { - set_rawintensity((char *)opt); - } - else if(!empty_string((opt = get_pool_setting(pool->xintensity, default_profile.xintensity)))) { - set_xintensity((char *)opt); - } - else if(!empty_string((opt = get_pool_setting(pool->intensity, ((!empty_string(default_profile.intensity))?default_profile.intensity:"8"))))) { - set_intensity((char *)opt); - } - - //shaders - if(!empty_string((opt = get_pool_setting(pool->shaders, default_profile.shaders)))) - set_shaders((char *)opt); - - //thread-concurrency - // neoscrypt - if not specified set TC to 0 so that TC will be calculated by intensity settings - if (pool->algorithm.type == ALGO_NEOSCRYPT) { - opt = ((empty_string(pool->thread_concurrency))?"0":get_pool_setting(pool->thread_concurrency, default_profile.thread_concurrency)); - } - // otherwise use pool/profile setting or default to default profile setting - else { - opt = get_pool_setting(pool->thread_concurrency, default_profile.thread_concurrency); - } - - if (!empty_string(opt)) { - set_thread_concurrency(opt); - } - - //worksize - if(!empty_string((opt = get_pool_setting(pool->worksize, default_profile.worksize)))) - set_worksize(opt); -*/ + //manually apply algorithm for (i = 0; i < nDevs; i++) { applog(LOG_DEBUG, "Set GPU %d to %s", i, isnull(pool->algorithm.name, "")); gpus[i].algorithm = pool->algorithm; } -/* - #ifdef HAVE_ADL - options = APPLY_ENGINE | APPLY_MEMCLOCK | APPLY_FANSPEED | APPLY_POWERTUNE | APPLY_VDDC; - - //GPU clock - if(!empty_string((opt = get_pool_setting(pool->gpu_engine, default_profile.gpu_engine)))) - set_gpu_engine((char *)opt); - else - options ^= APPLY_ENGINE; - - //GPU memory clock - if(!empty_string((opt = get_pool_setting(pool->gpu_memclock, default_profile.gpu_memclock)))) - set_gpu_memclock((char *)opt); - else - options ^= APPLY_MEMCLOCK; - - //GPU fans - if(!empty_string((opt = get_pool_setting(pool->gpu_fan, default_profile.gpu_fan)))) - set_gpu_fan((char *)opt); - else - options ^= APPLY_FANSPEED; - - //GPU powertune - if(!empty_string((opt = get_pool_setting(pool->gpu_powertune, default_profile.gpu_powertune)))) - set_gpu_powertune((char *)opt); - else - options ^= APPLY_POWERTUNE; - - //GPU vddc - if(!empty_string((opt = get_pool_setting(pool->gpu_vddc, default_profile.gpu_vddc)))) - set_gpu_vddc((char *)opt); - else - options ^= APPLY_VDDC; - - //apply gpu settings - for (i = 0; i < nDevs; i++) - { - if(opt_isset(options, APPLY_ENGINE)) - set_engineclock(i, gpus[i].min_engine); - if(opt_isset(options, APPLY_MEMCLOCK)) - set_memoryclock(i, gpus[i].gpu_memclock); - if(opt_isset(options, APPLY_FANSPEED)) - set_fanspeed(i, gpus[i].min_fan); - if(opt_isset(options, APPLY_POWERTUNE)) - set_powertune(i, gpus[i].gpu_powertune); - if(opt_isset(options, APPLY_VDDC)) - set_vddc(i, gpus[i].gpu_vddc); - } - #endif*/ rd_unlock(&mining_thr_lock); @@ -7522,6 +7457,7 @@ struct work *get_work(struct thr_info *thr, const int thr_id) applog(LOG_DEBUG, "[THR%d] Got work from get queue", thr_id); work->thr_id = thr_id; + work->thr = thr; thread_reportin(thr); work->mined = true; work->device_diff = MIN(thr->cgpu->drv->max_diff, work->work_difficulty); @@ -8644,7 +8580,7 @@ static void *test_pool_thread(void *arg) cg_wlock(&control_lock); if (!pools_active) { - currentpool = pool; + set_current_pool(pool); if (pool->pool_no != 0) first_pool = true; pools_active = true; @@ -9027,6 +8963,9 @@ bool add_cgpu(struct cgpu_info *cgpu) devices[total_devices++] = cgpu; wr_unlock(&devices_lock); + cgpu->eth_dag.current_epoch = 0xffffffffU; + cglock_init(&cgpu->eth_dag.lock); + adjust_mostdevs(); return true; } @@ -9225,9 +9164,6 @@ int main(int argc, char *argv[]) initial_args[i] = (const char *)strdup(argv[i]); initial_args[argc] = NULL; - EthCache[0] = EthCache[1] = NULL; - cglock_init(&EthCacheLock[0]); - cglock_init(&EthCacheLock[1]); mutex_init(ð_nonce_lock); #ifdef WIN32 rand_s(ð_nonce); @@ -9494,7 +9430,7 @@ int main(int argc, char *argv[]) } } /* Set the currentpool to pool 0 */ - currentpool = pools[0]; + set_current_pool(pools[0]); #ifdef HAVE_SYSLOG_H if (use_syslog) diff --git a/util.c b/util.c index 8c64a38b..0d21a3d0 100644 --- a/util.c +++ b/util.c @@ -1820,7 +1820,7 @@ static bool parse_notify(struct pool *pool, json_t *val) } -bool parse_diff_ethash(char* Target, char* TgtStr) +bool parse_diff_ethash(char* Target, const char* TgtStr) { bool ret = false; int len = strlen(TgtStr); @@ -1893,10 +1893,11 @@ static bool parse_notify_ethash(struct pool *pool, json_t *val) free(pool->swork.job_id); pool->swork.job_id = strdup(job_id); pool->swork.clean = clean; - - if (memcmp(pool->SeedHash, SeedHash, 32)) { - pool->EpochNumber = EthCalcEpochNumber(SeedHash); - memcpy(pool->SeedHash, SeedHash, 32); + + if (memcmp(pool->eth_cache.seed_hash, SeedHash, 32)) { + pool->eth_cache.current_epoch = EthCalcEpochNumber(SeedHash); + memcpy(pool->eth_cache.seed_hash, SeedHash, 32); + eth_gen_cache(pool); } memcpy(pool->EthWork, EthWork, 32); @@ -1910,6 +1911,7 @@ static bool parse_notify_ethash(struct pool *pool, json_t *val) pool->diff1 = eth2pow256 / le256todouble(pool->NetDiff); } pool->getwork_requested++; + //pool->eth_cache.disabled = false; cg_wunlock(&pool->data_lock);