diff --git a/CMakeLists.txt b/CMakeLists.txt index aa0d965..f18582b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -281,6 +281,9 @@ endif() if (TOOL_MOCKER) add_compile_definitions(TOOL_MOCKER) endif() +if (TOOL_TRACE_MEMORY OR TOOL_MOCKER) + add_compile_definitions(TOOLS_ENABLE) +endif() ################################################################################ diff --git a/PARAMETERS.md b/PARAMETERS.md index 8b12667..b9be8ec 100644 --- a/PARAMETERS.md +++ b/PARAMETERS.md @@ -105,3 +105,10 @@ N/A : No default value is set. |---------|--------|--------------|-------------|---------| | `--sm_wallet` | ✅ | N/A | Assign a wallet with a coin. | `--sm_wallet=COIN:WALLET` | | `--sm_pool` | ✅ | N/A | Assign a pool with a coin. | `--sm_pool=COIN@POOL_URL:POOL_PORT` | + + +## Tools Developer + +| Parameter | Optional | Default Value | Description | Example | +|---------|--------|--------------|-------------|---------| +| `--tool_mocker_resolver_count` | ✅ | 8 | Defines the number of fake devices that will mine / create fake latency. | `--tool_mocker_resolver_count=8` | diff --git a/sources/algo/dag_context.hpp b/sources/algo/dag_context.hpp index c546a04..cf47f69 100644 --- a/sources/algo/dag_context.hpp +++ b/sources/algo/dag_context.hpp @@ -21,7 +21,6 @@ namespace algo uint64_t period{ 0ull }; HashContext lightCache{}; HashContext dagCache{}; - algo::hash256 originalSeedCache{}; algo::hash512 hashedSeedCache{}; }; } diff --git a/sources/algo/ethash/cuda/memory.cuh b/sources/algo/ethash/cuda/memory.cuh index c5d218d..0e90bc5 100644 --- a/sources/algo/ethash/cuda/memory.cuh +++ b/sources/algo/ethash/cuda/memory.cuh @@ -31,7 +31,12 @@ bool ethashInitMemory( } //////////////////////////////////////////////////////////////////////////// - CU_ALLOC(¶ms.seedCache, algo::LEN_HASH_512); + if (true == buildLightCacheOnGPU) + { + CU_ALLOC(¶ms.seedCache, algo::LEN_HASH_512); + } + + //////////////////////////////////////////////////////////////////////////// CU_ALLOC(¶ms.lightCache, context.lightCache.size); CU_ALLOC(¶ms.dagCache, context.dagCache.size); CU_ALLOC_HOST(¶ms.resultCache, sizeof(algo::ethash::Result) * 2u); @@ -47,10 +52,15 @@ bool ethashInitMemory( params.resultCache[1].nonces[i] = 0ull; } + //////////////////////////////////////////////////////////////////////////// + if (true == buildLightCacheOnGPU) + { + IS_NULL(params.seedCache); + } + //////////////////////////////////////////////////////////////////////////// IS_NULL(params.lightCache); IS_NULL(params.dagCache); - IS_NULL(params.seedCache); IS_NULL(params.resultCache); //////////////////////////////////////////////////////////////////////////// diff --git a/sources/algo/ethash/ethash.cpp b/sources/algo/ethash/ethash.cpp index fa7b5df..91d5a77 100644 --- a/sources/algo/ethash/ethash.cpp +++ b/sources/algo/ethash/ethash.cpp @@ -10,9 +10,16 @@ #include -int32_t algo::ethash::findEpoch( +algo::ethash::ContextGenerator& algo::ethash::ContextGenerator::instance() +{ + static algo::ethash::ContextGenerator handler{}; + return handler; +} + + +int32_t algo::ethash::ContextGenerator::findEpoch( algo::hash256 const& seedHash, - uint32_t const maxEpoch) + uint32_t const maxEpoch) const { //////////////////////////////////////////////////////////////////////////// static thread_local int32_t cachedEpochNumber{ 0 }; @@ -56,29 +63,65 @@ int32_t algo::ethash::findEpoch( } -void algo::ethash::freeDagContext(algo::DagContext& context) +void algo::ethash::ContextGenerator::free(algo::ALGORITHM const algorithm) { //////////////////////////////////////////////////////////////////////////// - SAFE_DELETE_ARRAY(context.data); - context.lightCache.hash = nullptr; + UNIQUE_LOCK(mtxContexts); + + //////////////////////////////////////////////////////////////////////////// + algo::ethash::ContextGenerator::DagContextShare& contextShare{ getContext(algorithm) }; + + //////////////////////////////////////////////////////////////////////////// + --contextShare.count; + if (0u == contextShare.count) + { + SAFE_DELETE_ARRAY(contextShare.context.data); + contextShare.context.lightCache.hash = nullptr; + } +} + + +algo::ethash::ContextGenerator::DagContextShare& algo::ethash::ContextGenerator::getContext( + algo::ALGORITHM const algorithm) +{ + return contexts[algorithm]; } -void algo::ethash::initializeDagContext( +void algo::ethash::ContextGenerator::build( + algo::ALGORITHM const algorithm, algo::DagContext& context, uint64_t const currentEpoch, uint32_t const maxEpoch, uint64_t const dagCountItemsGrowth, uint64_t const dagCountItemsInit, uint32_t const lightCacheCountItemsGrowth, - uint32_t const lightCacheCountItemsInit) + uint32_t const lightCacheCountItemsInit, + bool const buildOnCPU) { //////////////////////////////////////////////////////////////////////////// - context.epoch = castU32(currentEpoch); - if ( castU32(context.epoch) > maxEpoch + UNIQUE_LOCK(mtxContexts); + + /////////////////////////////////////////////////////////////////////////// + algo::ethash::ContextGenerator::DagContextShare& contextShare{ getContext(algorithm) }; + algo::DagContext& localDagContext{ contextShare.context }; + + //////////////////////////////////////////////////////////////////////////// + ++contextShare.count; + + //////////////////////////////////////////////////////////////////////////// + if (localDagContext.epoch == cast32(currentEpoch)) + { + copyContext(context, algorithm); + return; + } + + //////////////////////////////////////////////////////////////////////////// + localDagContext.epoch = castU32(currentEpoch); + if ( castU32(localDagContext.epoch) > maxEpoch && algo::ethash::EIP1057_MAX_EPOCH_NUMER != maxEpoch) { - logErr() << "context.epoch: " << context.epoch << " | maxEpoch: " << maxEpoch; + logErr() << "context.epoch: " << localDagContext.epoch << " | maxEpoch: " << maxEpoch; return; } @@ -94,83 +137,105 @@ void algo::ethash::initializeDagContext( } //////////////////////////////////////////////////////////////////////////// - uint64_t lightCacheNumItemsUpperBound { castU64(epochEIP) }; + uint64_t lightCacheNumItemsUpperBound{ castU64(epochEIP) }; lightCacheNumItemsUpperBound *= lightCacheCountItemsGrowth; lightCacheNumItemsUpperBound += lightCacheCountItemsInit; - context.lightCache.numberItem = algo::largestPrime(lightCacheNumItemsUpperBound); - context.lightCache.size = context.lightCache.numberItem * algo::LEN_HASH_512; + localDagContext.lightCache.numberItem = algo::largestPrime(lightCacheNumItemsUpperBound); + localDagContext.lightCache.size = localDagContext.lightCache.numberItem * algo::LEN_HASH_512; //////////////////////////////////////////////////////////////////////////// uint64_t numberItemUpperBound{ castU64(epochEIP) }; numberItemUpperBound *= dagCountItemsGrowth; numberItemUpperBound += dagCountItemsInit; - context.dagCache.numberItem = algo::largestPrime(numberItemUpperBound); - context.dagCache.size = context.dagCache.numberItem * algo::LEN_HASH_1024; + localDagContext.dagCache.numberItem = algo::largestPrime(numberItemUpperBound); + localDagContext.dagCache.size = localDagContext.dagCache.numberItem * algo::LEN_HASH_1024; //////////////////////////////////////////////////////////////////////////// algo::hash256 seed{}; - for (int32_t i{ 0 }; i < context.epoch; ++i) + for (int32_t i{ 0 }; i < localDagContext.epoch; ++i) { seed = algo::keccak(seed); } - algo::copyHash(context.originalSeedCache, seed); -} - -void algo::ethash::buildLightCache( - algo::DagContext& context, - bool const buildOnCPU) -{ //////////////////////////////////////////////////////////////////////////// - algo::hash512 item - { - algo::keccak(context.originalSeedCache) - }; + algo::hash512 seedHash{ algo::keccak(seed) }; + algo::copyHash(localDagContext.hashedSeedCache, seedHash); //////////////////////////////////////////////////////////////////////////// - if (false == buildOnCPU) + if (true == buildOnCPU) { - algo::copyHash(context.hashedSeedCache, item); - return; + buildLightCache(algorithm); } //////////////////////////////////////////////////////////////////////////// - size_t const dataLength{ algo::LEN_HASH_512 + context.lightCache.size }; - context.data = NEW_ARRAY(char, dataLength); - std::memset(context.data, 0, sizeof(char) * dataLength); - if (nullptr == context.data) + copyContext(context, algorithm); +} + + +void algo::ethash::ContextGenerator::buildLightCache( + algo::ALGORITHM const algorithm) +{ + //////////////////////////////////////////////////////////////////////////// + algo::ethash::ContextGenerator::DagContextShare& contextShare{ getContext(algorithm) }; + algo::DagContext& localDagContext{ contextShare.context }; + + //////////////////////////////////////////////////////////////////////////// + algo::hash512 item{ localDagContext.hashedSeedCache }; + size_t const dataLength{ algo::LEN_HASH_512 + localDagContext.lightCache.size }; + localDagContext.data = NEW_ARRAY(char, dataLength); + std::memset(localDagContext.data, 0, sizeof(char) * dataLength); + if (nullptr == localDagContext.data) { logErr() << "Cannot alloc context data"; return; } - context.lightCache.hash = castPtrHash512(context.data + algo::LEN_HASH_512); - context.lightCache.hash[0] = item; + localDagContext.lightCache.hash = castPtrHash512(localDagContext.data + algo::LEN_HASH_512); //////////////////////////////////////////////////////////////////////////// logInfo() << "Building light cache on CPU"; - common::ChronoGuard chrono{ "Built light cache", common::CHRONO_UNIT::MS }; + common::ChronoGuard chrono{ "Built light cache on CPU", common::CHRONO_UNIT::MS }; //////////////////////////////////////////////////////////////////////////// - for (uint64_t i{ 1ull }; i < context.lightCache.numberItem; ++i) + for (uint64_t i{ 0ull }; i < localDagContext.lightCache.numberItem; ++i) { + localDagContext.lightCache.hash[i] = item; item = algo::keccak(item); - context.lightCache.hash[i] = item; } //////////////////////////////////////////////////////////////////////////// - uint32_t const numberItemu32 { castU32(context.lightCache.numberItem) }; + uint32_t const numberItemu32{ castU32(localDagContext.lightCache.numberItem) }; for (uint64_t round{ 0ull }; round < algo::ethash::LIGHT_CACHE_ROUNDS; ++round) { - for (uint64_t i{ 0ull }; i < context.lightCache.numberItem; ++i) + for (uint64_t i{ 0ull }; i < localDagContext.lightCache.numberItem; ++i) { - uint32_t const fi{ context.lightCache.hash[i].word32[0] % numberItemu32 }; + uint32_t const fi{ localDagContext.lightCache.hash[i].word32[0] % numberItemu32 }; uint32_t const si{ (numberItemu32 + (castU32(i) - 1u)) % numberItemu32 }; - algo::hash512 const& firstCache{ context.lightCache.hash[fi] }; - algo::hash512 const& secondCache{ context.lightCache.hash[si] }; + algo::hash512 const& firstCache{ localDagContext.lightCache.hash[fi] }; + algo::hash512 const& secondCache{ localDagContext.lightCache.hash[si] }; algo::hash512 const xored = algo::hashXor(firstCache, secondCache); - context.lightCache.hash[i] = algo::keccak(xored); + localDagContext.lightCache.hash[i] = algo::keccak(xored); } } } + + +void algo::ethash::ContextGenerator::copyContext( + algo::DagContext& context, + algo::ALGORITHM const algorithm) +{ + //////////////////////////////////////////////////////////////////////////// + algo::ethash::ContextGenerator::DagContextShare& contextShare{ getContext(algorithm) }; + algo::DagContext& localDagContext{ contextShare.context }; + + context.epoch = localDagContext.epoch; + context.lightCache.numberItem = localDagContext.lightCache.numberItem; + context.lightCache.size = localDagContext.lightCache.size; + context.dagCache.numberItem = localDagContext.dagCache.numberItem; + context.dagCache.size = localDagContext.dagCache.size; + context.data = localDagContext.data; + context.lightCache.hash = localDagContext.lightCache.hash; + + algo::copyHash(context.hashedSeedCache, localDagContext.hashedSeedCache); +} diff --git a/sources/algo/ethash/ethash.hpp b/sources/algo/ethash/ethash.hpp index 4156814..12f7215 100644 --- a/sources/algo/ethash/ethash.hpp +++ b/sources/algo/ethash/ethash.hpp @@ -1,7 +1,10 @@ #pragma once #if !defined(__LIB_CUDA) +#include + #include +#include #include #endif @@ -32,18 +35,38 @@ namespace algo constexpr uint64_t DAG_COUNT_ITEMS_GROWTH{ algo::ethash::DAG_GROWTH / algo::LEN_HASH_1024 }; #if !defined(__LIB_CUDA) - int32_t findEpoch(algo::hash256 const& seedHash, - uint32_t const maxEpoch); - void initializeDagContext(algo::DagContext& context, - uint64_t const currentEpoch, - uint32_t const maxEpoch, - uint64_t const dagCountItemsGrowth, - uint64_t const dagCountItemsInit, - uint32_t const lightCacheCountItemsGrowth, - uint32_t const lightCacheCountItemsInit); - void freeDagContext(algo::DagContext& context); - void buildLightCache(algo::DagContext& context, - bool const buildOnCPU); + struct ContextGenerator + { + public: + static ContextGenerator& instance(); + int32_t findEpoch(algo::hash256 const& seedHash, + uint32_t const maxEpoch) const; + void free(algo::ALGORITHM const algorithm); + void build(algo::ALGORITHM const algorithm, + algo::DagContext& context, + uint64_t const currentEpoch, + uint32_t const maxEpoch, + uint64_t const dagCountItemsGrowth, + uint64_t const dagCountItemsInit, + uint32_t const lightCacheCountItemsGrowth, + uint32_t const lightCacheCountItemsInit, + bool const buildOnCPU); + + private: + struct DagContextShare + { + uint32_t count{ 0u }; + algo::DagContext context; + }; + boost::mutex mtxContexts{}; + std::map contexts{}; + + ContextGenerator() = default; + DagContextShare& getContext(algo::ALGORITHM const algorithm); + void buildLightCache(algo::ALGORITHM const algorithm); + void copyContext(algo::DagContext& context, + algo::ALGORITHM const algorithm); + }; #endif } } diff --git a/sources/algo/ethash/result.hpp b/sources/algo/ethash/result.hpp index 708d9a2..6e66d33 100644 --- a/sources/algo/ethash/result.hpp +++ b/sources/algo/ethash/result.hpp @@ -21,10 +21,10 @@ namespace algo struct ResultShare { std::string jobId{}; - bool found { false }; - uint32_t extraNonceSize { 0u }; - uint32_t count { 0u }; - uint64_t nonces[algo::ethash::MAX_RESULT] { 0ull, 0ull, 0ull, 0ull }; + bool found{ false }; + uint32_t extraNonceSize{ 0u }; + uint32_t count{ 0u }; + uint64_t nonces[algo::ethash::MAX_RESULT]{ 0ull, 0ull, 0ull, 0ull }; }; #endif } diff --git a/sources/algo/ethash/tests/ethash.cpp b/sources/algo/ethash/tests/ethash.cpp index 2a32b79..5bc187c 100644 --- a/sources/algo/ethash/tests/ethash.cpp +++ b/sources/algo/ethash/tests/ethash.cpp @@ -33,16 +33,16 @@ TEST_F(EthashTest, epoch) { int32_t const maxEpochNumber { cast32(algo::ethash::MAX_EPOCH_NUMBER) }; - EXPECT_EQ(0, algo::ethash::findEpoch(hashes[0], maxEpochNumber)); - EXPECT_EQ(1, algo::ethash::findEpoch(hashes[1], maxEpochNumber)); - EXPECT_EQ(171, algo::ethash::findEpoch(hashes[2], maxEpochNumber)); - EXPECT_EQ(604, algo::ethash::findEpoch(hashes[3], maxEpochNumber)); - EXPECT_EQ(588, algo::ethash::findEpoch(hashes[4], maxEpochNumber)); - EXPECT_EQ(2048, algo::ethash::findEpoch(hashes[5], maxEpochNumber)); - EXPECT_EQ(29998, algo::ethash::findEpoch(hashes[6], maxEpochNumber)); - EXPECT_EQ(29999, algo::ethash::findEpoch(hashes[7], maxEpochNumber)); - EXPECT_EQ(maxEpochNumber - 1, algo::ethash::findEpoch(hashes[8], maxEpochNumber)); - EXPECT_EQ(maxEpochNumber, algo::ethash::findEpoch(hashes[9], maxEpochNumber)); + EXPECT_EQ(0, algo::ethash::ContextGenerator::instance().findEpoch(hashes[0], maxEpochNumber)); + EXPECT_EQ(1, algo::ethash::ContextGenerator::instance().findEpoch(hashes[1], maxEpochNumber)); + EXPECT_EQ(171, algo::ethash::ContextGenerator::instance().findEpoch(hashes[2], maxEpochNumber)); + EXPECT_EQ(604, algo::ethash::ContextGenerator::instance().findEpoch(hashes[3], maxEpochNumber)); + EXPECT_EQ(588, algo::ethash::ContextGenerator::instance().findEpoch(hashes[4], maxEpochNumber)); + EXPECT_EQ(2048, algo::ethash::ContextGenerator::instance().findEpoch(hashes[5], maxEpochNumber)); + EXPECT_EQ(29998, algo::ethash::ContextGenerator::instance().findEpoch(hashes[6], maxEpochNumber)); + EXPECT_EQ(29999, algo::ethash::ContextGenerator::instance().findEpoch(hashes[7], maxEpochNumber)); + EXPECT_EQ(maxEpochNumber - 1, algo::ethash::ContextGenerator::instance().findEpoch(hashes[8], maxEpochNumber)); + EXPECT_EQ(maxEpochNumber, algo::ethash::ContextGenerator::instance().findEpoch(hashes[9], maxEpochNumber)); } @@ -55,17 +55,18 @@ TEST_F(EthashTest, lightCacheBuild) uint32_t dagCountItemsGrowth{ algo::ethash::DAG_COUNT_ITEMS_GROWTH }; uint32_t dagCountItemsInit{ algo::ethash::DAG_COUNT_ITEMS_INIT }; - algo::ethash::initializeDagContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::ETHASH, context, 561ull, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true ); - algo::ethash::buildLightCache(context, true); // light cache size ASSERT_EQ(1411061ull, context.lightCache.numberItem); @@ -111,5 +112,5 @@ TEST_F(EthashTest, lightCacheBuild) ASSERT_EQ(762008954u, context.lightCache.hash[1000008].word32[0]); ASSERT_EQ(311220328u, context.lightCache.hash[1000009].word32[0]); - algo::ethash::freeDagContext(context); + algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::ETHASH); } diff --git a/sources/algo/progpow/cuda/dag.cuh b/sources/algo/progpow/cuda/dag.cuh index c756fec..ad0ec50 100644 --- a/sources/algo/progpow/cuda/dag.cuh +++ b/sources/algo/progpow/cuda/dag.cuh @@ -5,54 +5,60 @@ __device__ __forceinline__ -void doCopy( - uint4* __restrict__ const dst, +uint4 doCopy( uint4* __restrict__ const src, - uint32_t const workerId, - uint32_t const laneId) + uint32_t const worker_id, + uint32_t const lane_id) { + //////////////////////////////////////////////////////////////////////////// + uint4 item; + //////////////////////////////////////////////////////////////////////////// #pragma unroll - for (uint32_t i{ 0u }; i < DAG_HASH_U4_SIZE; ++i) + for (uint32_t i = 0u; i < DAG_HASH_U4_SIZE; ++i) { - uint32_t const x{ reg_load(src[i].x, laneId, DAG_PARRALLEL_LANE) }; - uint32_t const y{ reg_load(src[i].y, laneId, DAG_PARRALLEL_LANE) }; - uint32_t const z{ reg_load(src[i].z, laneId, DAG_PARRALLEL_LANE) }; - uint32_t const w{ reg_load(src[i].w, laneId, DAG_PARRALLEL_LANE) }; + uint4 const tmp_src = src[i]; + uint32_t const x = reg_load(tmp_src.x, lane_id, DAG_PARRALLEL_LANE); + uint32_t const y = reg_load(tmp_src.y, lane_id, DAG_PARRALLEL_LANE); + uint32_t const z = reg_load(tmp_src.z, lane_id, DAG_PARRALLEL_LANE); + uint32_t const w = reg_load(tmp_src.w, lane_id, DAG_PARRALLEL_LANE); - if (i == workerId) + if (i == worker_id) { - dst->x = x; - dst->y = y; - dst->z = z; - dst->w = w; + item.x = x; + item.y = y; + item.z = z; + item.w = w; } } + + //////////////////////////////////////////////////////////////////////////// + return item; } __device__ __forceinline__ void buildItemFromCache( uint4* const __restrict__ hash, - uint32_t const workerId, - uint32_t const cacheIndex) + uint32_t const worker_id, + uint32_t const cache_index) { //////////////////////////////////////////////////////////////////////////// - uint32_t const cacheIndexLimit{ (cacheIndex % d_light_number_item) * DAG_HASH_U4_SIZE }; + uint32_t const cache_index_limit = (cache_index % d_light_number_item) * DAG_HASH_U4_SIZE; //////////////////////////////////////////////////////////////////////////// #pragma unroll - for (uint32_t laneId{ 0u }; laneId < DAG_PARRALLEL_LANE; ++laneId) + for (uint32_t lane_id = 0u; lane_id < DAG_PARRALLEL_LANE; ++lane_id) { //////////////////////////////////////////////////////////////////////// - uint32_t const tmpCacheIndex{ reg_load(cacheIndexLimit, laneId, DAG_PARRALLEL_LANE) }; - uint32_t const index{ tmpCacheIndex + workerId }; - uint4 cache16Bytes = d_light_cache[index]; + uint32_t const tmp_cache_index = reg_load(cache_index_limit, lane_id, DAG_PARRALLEL_LANE); + uint32_t const index = tmp_cache_index + worker_id; + uint4 const cache16Bytes = d_light_cache[index]; //////////////////////////////////////////////////////////////////////// uint4 cache64Bytes[DAG_HASH_U4_SIZE]; #pragma unroll - for (uint32_t i{ 0u }; i < DAG_HASH_U4_SIZE; ++i) + for (uint32_t i = 0u; i < DAG_HASH_U4_SIZE; ++i) { cache64Bytes[i].x = reg_load(cache16Bytes.x, i, DAG_PARRALLEL_LANE); cache64Bytes[i].y = reg_load(cache16Bytes.y, i, DAG_PARRALLEL_LANE); @@ -61,10 +67,10 @@ void buildItemFromCache( } //////////////////////////////////////////////////////////////////////// - if (laneId == workerId) + if (lane_id == worker_id) { #pragma unroll - for (uint32_t i{ 0u }; i < DAG_HASH_U4_SIZE; ++i) + for (uint32_t i = 0u; i < DAG_HASH_U4_SIZE; ++i) { fnv1(&hash[i], &cache64Bytes[i]); } @@ -75,29 +81,29 @@ void buildItemFromCache( __device__ __forceinline__ void buildItem( - bool is_access, - uint32_t const workerId, - uint32_t const dagIndex, - uint32_t const loopIndex, - uint32_t const cacheIndex) + bool const is_access, + uint32_t const worker_id, + uint32_t const dag_index, + uint32_t const loop_index, + uint32_t const cache_index) { //////////////////////////////////////////////////////////////////////////// uint4 hash[DAG_HASH_U4_SIZE]; - uint32_t const cacheStartIndex{ (cacheIndex % d_light_number_item) * DAG_HASH_U4_SIZE }; + uint32_t const cache_start_index = (cache_index % d_light_number_item) * DAG_HASH_U4_SIZE; //////////////////////////////////////////////////////////////////////////// #pragma unroll - for (uint32_t lane_id{ 0u }; lane_id < DAG_PARRALLEL_LANE; ++lane_id) + for (uint32_t lane_id = 0u; lane_id < DAG_PARRALLEL_LANE; ++lane_id) { //////////////////////////////////////////////////////////////////////// - uint32_t const tmpCacheStartIndex{ reg_load(cacheStartIndex, lane_id, DAG_PARRALLEL_LANE) }; - uint32_t const tmpCacheIndex{ tmpCacheStartIndex + workerId }; - uint4 const tmp_cache{ d_light_cache[tmpCacheIndex] }; + uint32_t const tmp_cache_start_index = reg_load(cache_start_index, lane_id, DAG_PARRALLEL_LANE); + uint32_t const tmp_cache_index = tmp_cache_start_index + worker_id; + uint4 const tmp_cache = d_light_cache[tmp_cache_index]; uint4 tmp_hash[DAG_HASH_U4_SIZE]; //////////////////////////////////////////////////////////////////////// #pragma unroll - for (uint32_t i{ 0u }; i < DAG_PARRALLEL_LANE; ++i) + for (uint32_t i = 0u; i < DAG_PARRALLEL_LANE; ++i) { tmp_hash[i].x = reg_load(tmp_cache.x, i, DAG_PARRALLEL_LANE); tmp_hash[i].y = reg_load(tmp_cache.y, i, DAG_PARRALLEL_LANE); @@ -106,10 +112,10 @@ void buildItem( } //////////////////////////////////////////////////////////////////////// - if (lane_id == workerId) + if (lane_id == worker_id) { #pragma unroll - for (uint32_t i{ 0u }; i < DAG_HASH_U4_SIZE; ++i) + for (uint32_t i = 0u; i < DAG_HASH_U4_SIZE; ++i) { hash[i].x = tmp_hash[i].x; hash[i].y = tmp_hash[i].y; @@ -118,23 +124,23 @@ void buildItem( } } } - hash[0].x ^= cacheIndex; + hash[0].x ^= cache_index; //////////////////////////////////////////////////////////////////////////// keccak_f1600(hash); //////////////////////////////////////////////////////////////////////////// - uint32_t j{ 0u }; + uint32_t j = 0u; #pragma unroll - for (uint32_t y{ 0u }; y < loopIndex; ++y) + for (uint32_t y = 0u; y < loop_index; ++y) { #pragma unroll - for (uint32_t i{ 0u }; i < DAG_HASH_U4_SIZE; ++i) + for (uint32_t i = 0u; i < DAG_HASH_U4_SIZE; ++i) { - buildItemFromCache(hash, workerId, fnv1(cacheIndex ^ j, hash[i].x)); - buildItemFromCache(hash, workerId, fnv1(cacheIndex ^ (j + 1u), hash[i].y)); - buildItemFromCache(hash, workerId, fnv1(cacheIndex ^ (j + 2u), hash[i].z)); - buildItemFromCache(hash, workerId, fnv1(cacheIndex ^ (j + 3u), hash[i].w)); + buildItemFromCache(hash, worker_id, fnv1(cache_index ^ j, hash[i].x)); + buildItemFromCache(hash, worker_id, fnv1(cache_index ^ (j + 1u), hash[i].y)); + buildItemFromCache(hash, worker_id, fnv1(cache_index ^ (j + 2u), hash[i].z)); + buildItemFromCache(hash, worker_id, fnv1(cache_index ^ (j + 3u), hash[i].w)); j += 4u; } @@ -144,15 +150,14 @@ void buildItem( keccak_f1600(hash); //////////////////////////////////////////////////////////////////////////// - uint4 itemHash; #pragma unroll - for (uint32_t laneId{ 0u }; laneId < DAG_HASH_U4_SIZE; ++laneId) + for (uint32_t lane_id = 0u; lane_id < DAG_HASH_U4_SIZE; ++lane_id) { - doCopy(&itemHash, hash, workerId, laneId); - uint32_t const tmpDagIndex{ reg_load(dagIndex, laneId, DAG_PARRALLEL_LANE) }; + uint4 const itemHash = doCopy(hash, worker_id, lane_id); + uint32_t const tmpdag_index = reg_load(dag_index, lane_id, DAG_PARRALLEL_LANE); if (true == is_access) { - uint32_t const index{ tmpDagIndex + workerId }; + uint32_t const index = tmpdag_index + worker_id; d_dag[index] = itemHash; } } @@ -161,15 +166,14 @@ void buildItem( __global__ void kernelProgpowBuildDag( - uint32_t const startIndex, + uint32_t const start_index, uint32_t const loop) { //////////////////////////////////////////////////////////////////////////// - uint32_t const threadId{ (blockIdx.x * blockDim.x) + threadIdx.x }; - uint32_t const workerId{ threadId & 3u }; - uint32_t const dagAccessWorker{ (threadId / 4u) + startIndex }; - uint32_t const dagAccessIndex{ threadId + startIndex }; - uint32_t const dagIndex{ threadId + startIndex }; + uint32_t const threadId = (blockIdx.x * blockDim.x) + threadIdx.x; + uint32_t const worker_id = threadId & 3u; + uint32_t const dagAccessWorker = (threadId / 4u) + start_index; + uint32_t const dag_index = threadId + start_index; //////////////////////////////////////////////////////////////////////////// if (dagAccessWorker >= d_dag_number_item) @@ -178,12 +182,12 @@ void kernelProgpowBuildDag( } //////////////////////////////////////////////////////////////////////////// - uint32_t const cacheIndex{ dagIndex * 2u }; - uint32_t const dagIndex1{ dagIndex * 8u }; - uint32_t const dagIndex2{ dagIndex1 + 4u }; - bool const isAccess{ dagAccessIndex < d_dag_number_item }; - buildItem(isAccess, workerId, dagIndex1, loop, cacheIndex); - buildItem(isAccess, workerId, dagIndex2, loop, cacheIndex + 1u); + uint32_t const cache_index = dag_index * 2u; + uint32_t const dag_index_1 = dag_index * 8u; + uint32_t const dag_index_2 = dag_index_1 + 4u; + bool const isAccess = dag_index < d_dag_number_item; + buildItem(isAccess, worker_id, dag_index_1, loop, cache_index); + buildItem(isAccess, worker_id, dag_index_2, loop, cache_index + 1u); } @@ -193,12 +197,14 @@ bool progpowBuildDag( uint32_t const dagItemParents, uint32_t const dagNumberItems) { - uint32_t const itemByKernel{ 65536u }; - uint32_t const loop{ dagItemParents / 4u / 4u }; + uint32_t const threads = 128u; + uint32_t const blocks = 512u; + uint32_t const itemByKernel = threads * blocks; + uint32_t const loop = dagItemParents / DAG_PARRALLEL_LANE / DAG_HASH_U4_SIZE; - for (uint32_t i{ 0u }; i < dagNumberItems; i += itemByKernel) + for (uint32_t i = 0u; i < dagNumberItems; i += itemByKernel) { - kernelProgpowBuildDag<<<512, 128, 0, stream>>>(i, loop); + kernelProgpowBuildDag<<>>(i, loop); CUDA_ER(cudaStreamSynchronize(stream)); CUDA_ER(cudaGetLastError()); } diff --git a/sources/algo/progpow/cuda/memory.cuh b/sources/algo/progpow/cuda/memory.cuh index c338095..e8dc3f0 100644 --- a/sources/algo/progpow/cuda/memory.cuh +++ b/sources/algo/progpow/cuda/memory.cuh @@ -15,6 +15,7 @@ bool progpowFreeMemory( CU_SAFE_DELETE(params.dagCache); CU_SAFE_DELETE(params.headerCache); CU_SAFE_DELETE_HOST(params.resultCache); + return true; } @@ -32,12 +33,23 @@ bool progpowInitMemory( } //////////////////////////////////////////////////////////////////////////// - CU_ALLOC(¶ms.seedCache, algo::LEN_HASH_512); + if (true == buildLightCacheOnGPU) + { + CU_ALLOC(¶ms.seedCache, algo::LEN_HASH_512); + } + + //////////////////////////////////////////////////////////////////////////// CU_ALLOC(¶ms.lightCache, context.lightCache.size); CU_ALLOC(¶ms.dagCache, context.dagCache.size); CU_ALLOC(¶ms.headerCache, sizeof(uint32_t) * algo::LEN_HASH_256_WORD_32); CU_ALLOC_HOST(¶ms.resultCache, sizeof(algo::progpow::Result) * 2u); + //////////////////////////////////////////////////////////////////////////// + if (true == buildLightCacheOnGPU) + { + IS_NULL(params.seedCache); + } + //////////////////////////////////////////////////////////////////////////// IS_NULL(params.lightCache); IS_NULL(params.dagCache); @@ -49,11 +61,11 @@ bool progpowInitMemory( params.resultCache[0].found = false; params.resultCache[1].count = 0u; params.resultCache[1].found = false; - for (uint32_t i{ 0u }; i < 4u; ++i) + for (uint32_t i{ 0u }; i < algo::progpow::MAX_RESULT; ++i) { params.resultCache[0].nonces[i] = 0ull; params.resultCache[1].nonces[i] = 0ull; - for (uint32_t x{ 0u }; x < 8u; ++x) + for (uint32_t x{ 0u }; x < algo::LEN_HASH_256_WORD_32; ++x) { params.resultCache[0].hash[i][x] = 0u; params.resultCache[1].hash[i][x] = 0u; @@ -118,10 +130,12 @@ bool progpowUpdateConstants( uint32_t const* const hostBuffer, uint32_t* const deviceBuffer) { - CUDA_ER(cudaMemcpy(deviceBuffer, - hostBuffer, - algo::LEN_HASH_256_WORD_32 * sizeof(uint32_t), - cudaMemcpyHostToDevice)); + CUDA_ER( + cudaMemcpy( + deviceBuffer, + hostBuffer, + algo::LEN_HASH_256_WORD_32 * sizeof(uint32_t), + cudaMemcpyHostToDevice)); return true; } diff --git a/sources/benchmark/amd/kawpow.cpp b/sources/benchmark/amd/kawpow.cpp index 8b4a408..8bc00af 100644 --- a/sources/benchmark/amd/kawpow.cpp +++ b/sources/benchmark/amd/kawpow.cpp @@ -36,7 +36,7 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow() { algo::toHash256("7c4fb8a5d141973b69b521ce76b0dc50f0d2834d817c7f8310a6ab5becc6bb0c") }; - int32_t const epoch{ algo::ethash::findEpoch(seedHash, algo::ethash::EPOCH_LENGTH) }; + int32_t const epoch{ algo::ethash::ContextGenerator::instance().findEpoch(seedHash, algo::ethash::EPOCH_LENGTH) }; //////////////////////////////////////////////////////////////////////////// common::opencl::Buffer lightCache{ CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY }; @@ -50,17 +50,18 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow() //////////////////////////////////////////////////////////////////////////// algo::DagContext dagContext{}; - algo::ethash::initializeDagContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::KAWPOW, dagContext, epoch, algo::ethash::MAX_EPOCH_NUMBER, algo::ethash::DAG_COUNT_ITEMS_GROWTH, algo::ethash::DAG_COUNT_ITEMS_INIT, algo::ethash::LIGHT_CACHE_COUNT_ITEMS_GROWTH, - algo::ethash::LIGHT_CACHE_COUNT_ITEMS_INIT + algo::ethash::LIGHT_CACHE_COUNT_ITEMS_INIT, + true /*config.deviceAlgorithm.ethashBuildLightCacheCPU*/ ); - algo::ethash::buildLightCache(dagContext, true); //////////////////////////////////////////////////////////////////////////// dagCache.setSize(dagContext.dagCache.size); @@ -228,7 +229,7 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow() } //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(dagContext); + algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::KAWPOW); //////////////////////////////////////////////////////////////////////////// dagCache.free(); diff --git a/sources/common/cli/cli.cpp b/sources/common/cli/cli.cpp index 4070392..aa6d584 100644 --- a/sources/common/cli/cli.cpp +++ b/sources/common/cli/cli.cpp @@ -47,6 +47,13 @@ common::Cli::Cli() "[OPTIONAL] Set the time interval (in milliseconds) between logs of information about the hashrate.\n" "--log_interval_hash=10000" ) + ( + "log_new_job", + value(), + "[OPTIONAL] Show log when receive new job.\n" + "Default value is true.\n" + "--log_new_job=true" + ) // Environment ( diff --git a/sources/common/cli/cli.hpp b/sources/common/cli/cli.hpp index 5c2adda..1f2f54f 100644 --- a/sources/common/cli/cli.hpp +++ b/sources/common/cli/cli.hpp @@ -49,6 +49,7 @@ namespace common std::optional getLevelLog() const; std::optional getLogFilenaName() const; std::optional getLogIntervalHashStats() const; + bool isLogNewJob() const; // Environment std::optional getEnvironmentCudaLazy() const; @@ -58,7 +59,7 @@ namespace common std::optional getEnvironmentGpuSingleAllocPercent() const; // Common - std::optional getPricekWH() const; + std::optional getPricekWH() const; // Pool Connection std::optional getHost() const; @@ -90,6 +91,11 @@ namespace common #endif bool isCpuEnable() const; +#if defined(TOOLS_ENABLE) && defined(TOOL_MOCKER) + std::optional getMockerResolverCount() const; + std::optional getMockerResolverUpdateMemorySleep() const; +#endif + // Device settings custom std::vector getDevicesDisable() const; customTupleStr getCustomHost() const; diff --git a/sources/common/cli/cli_log.cpp b/sources/common/cli/cli_log.cpp index 4418ad9..81b96c7 100644 --- a/sources/common/cli/cli_log.cpp +++ b/sources/common/cli/cli_log.cpp @@ -48,3 +48,14 @@ std::optional common::Cli::getLogIntervalHashStats() const return std::nullopt; } + + +bool common::Cli::isLogNewJob() const +{ + if (true == contains("log_new_job")) + { + return params["log_new_job"].as(); + } + + return true; +} diff --git a/sources/common/cli/cli_tools.cpp b/sources/common/cli/cli_tools.cpp new file mode 100644 index 0000000..b1435e6 --- /dev/null +++ b/sources/common/cli/cli_tools.cpp @@ -0,0 +1,27 @@ +#if defined(TOOLS_ENABLE) && defined(TOOL_MOCKER) + +#include + + +std::optional common::Cli::getMockerResolverCount() const +{ + if (true == contains("tool_mocker_resolver_count")) + { + return params["tool_mocker_resolver_count"].as(); + } + + return std::nullopt; +} + + +std::optional common::Cli::getMockerResolverUpdateMemorySleep() const +{ + if (true == contains("tool_mocker_resolver_update_memory_sleep")) + { + return params["tool_mocker_resolver_update_memory_sleep"].as(); + } + + return std::nullopt; +} + +#endif // TOOLS_ENABLE && TOOL_MOCKER diff --git a/sources/common/config.cpp b/sources/common/config.cpp index 6e7001b..f009a09 100644 --- a/sources/common/config.cpp +++ b/sources/common/config.cpp @@ -110,6 +110,9 @@ bool common::Config::loadCli(int argc, char** argv) log.intervalHashStats = common::min_limit(*intervalHashStats, 100u); } + logInfo() << "cli.isLogNewJob(): " << cli.isLogNewJob(); + log.showNewJob = cli.isLogNewJob(); + //////////////////////////////////////////////////////////////////////// // ENVIRONMENT //////////////////////////////////////////////////////////////////////// @@ -481,6 +484,14 @@ bool common::Config::loadCli(int argc, char** argv) { api.port = ApiPort; } + + //////////////////////////////////////////////////////////////////////// + // TOOL MOCKER + //////////////////////////////////////////////////////////////////////// +#if defined(TOOLS_ENABLE) && defined(TOOL_MOCKER) + toolConfigs.mockerResolverCount = cli.getMockerResolverCount(); + toolConfigs.mockerResolverUpdateMemorySleep = cli.getMockerResolverUpdateMemorySleep(); +#endif } catch(std::exception const& e) { diff --git a/sources/common/config.hpp b/sources/common/config.hpp index e70f2cc..6d1324e 100644 --- a/sources/common/config.hpp +++ b/sources/common/config.hpp @@ -76,6 +76,7 @@ namespace common common::TYPELOG level{ common::TYPELOG::__INFO }; std::string file{}; std::optional intervalHashStats{}; + bool showNewJob{ true }; }; struct ApiConfig @@ -83,6 +84,16 @@ namespace common uint32_t port{ 8080u }; }; +#if defined(TOOLS_ENABLE) + struct ToolConfig + { +#if defined(TOOL_MOCKER) + std::optional mockerResolverCount{ 8u }; + std::optional mockerResolverUpdateMemorySleep{ 8u }; +#endif + }; +#endif + common::PROFILE profile{ common::PROFILE::STANDARD }; common::Cli cli{}; LogConfig log{}; @@ -97,6 +108,9 @@ namespace common PoolConfig nvidiaSetting{}; ApiConfig api{}; CommonConfig common{}; +#if defined(TOOLS_ENABLE) + ToolConfig toolConfigs{}; +#endif static Config& instance(); bool load(int argc, char** argv); diff --git a/sources/common/log/log.hpp b/sources/common/log/log.hpp index 74f9aaf..f1965ed 100644 --- a/sources/common/log/log.hpp +++ b/sources/common/log/log.hpp @@ -75,9 +75,6 @@ namespace common #define resolverTrace() common::Logger(__FUNCTION__, __LINE__, common::TYPELOG::__TRACE) << "Device[" << deviceId << "]: " #define resolverDebug() common::Logger(__FUNCTION__, __LINE__, common::TYPELOG::__DEBUG) << "Device[" << deviceId << "]: " -#define __TRACE() { logTrace() << __FUNCTION__ << ":" << __LINE__; } -#define __TRACE_DEVICE() \ - { \ - logTrace() << __FUNCTION__ << ":" << __LINE__ \ - << ": device[" << id << "]"; \ - } +#define __TRACE() { logTrace() << __FUNCTION__ << ":" << __LINE__; } +#define __TRACE_DEVICE() { deviceTrace() << __FUNCTION__ << ":" << __LINE__; } +#define __TRACE_RESOLVER() { resolverTrace() << __FUNCTION__ << ":" << __LINE__; } diff --git a/sources/device/device_manager.cpp b/sources/device/device_manager.cpp index 10df18d..1de3593 100644 --- a/sources/device/device_manager.cpp +++ b/sources/device/device_manager.cpp @@ -285,16 +285,24 @@ bool device::DeviceManager::initializeStratum( #if defined(TOOL_MOCKER) bool device::DeviceManager::initializeMocker() { - for (uint32_t i = 0; i < 100u; ++i) + /////////////////////////////////////////////////////////////////////////// + common::Config const& config{ common::Config::instance() }; + + /////////////////////////////////////////////////////////////////////////// + if (std::nullopt != config.toolConfigs.mockerResolverCount) { - device::DeviceMocker* device{ NEW(device::DeviceMocker) }; - if (nullptr == device) + uint32_t const mockerCount{ *config.toolConfigs.mockerResolverCount }; + for (uint32_t i{ 0u }; i < mockerCount; ++i) { - return false; + device::DeviceMocker* device{ NEW(device::DeviceMocker) }; + if (nullptr == device) + { + return false; + } + device->id = 1000u + i; + device->deviceType = device::DEVICE_TYPE::MOCKER; + devices.push_back(device); } - device->id = 1000 + i; - device->deviceType = device::DEVICE_TYPE::MOCKER; - devices.push_back(device); } return true; } @@ -501,6 +509,7 @@ void device::DeviceManager::onUpdateJob( UNIQUE_LOCK(mtxJobInfo); //////////////////////////////////////////////////////////////////////////// + auto const& config{ common::Config::instance() }; stratum::StratumJobInfo& jobInfo{ jobInfos[stratumUUID] }; //////////////////////////////////////////////////////////////////////////// @@ -543,11 +552,14 @@ void device::DeviceManager::onUpdateJob( { jobInfo.copy(newJobInfo); jobInfo.gapNonce /= devices.size(); + if (true == config.log.showNewJob) + { #if defined(_DEBUG) - logInfo() << jobInfo; + logInfo() << jobInfo; #else - logInfo() << "New Job[" << jobInfo.jobID << "]"; + logInfo() << "New Job[" << jobInfo.jobID << "]"; #endif + } updateDevice(stratumUUID, updateMemory, updateConstants); } } diff --git a/sources/resolver/amd/autolykos_v2.cpp b/sources/resolver/amd/autolykos_v2.cpp index d644e51..802e021 100644 --- a/sources/resolver/amd/autolykos_v2.cpp +++ b/sources/resolver/amd/autolykos_v2.cpp @@ -11,6 +11,13 @@ #include +resolver::ResolverAmdAutolykosV2::ResolverAmdAutolykosV2(): + resolver::ResolverAmd() +{ + algorithm = algo::ALGORITHM::AUTOLYKOS_V2; +} + + resolver::ResolverAmdAutolykosV2::~ResolverAmdAutolykosV2() { parameters.BHashes.free(); diff --git a/sources/resolver/amd/autolykos_v2.hpp b/sources/resolver/amd/autolykos_v2.hpp index b480450..39efa96 100644 --- a/sources/resolver/amd/autolykos_v2.hpp +++ b/sources/resolver/amd/autolykos_v2.hpp @@ -15,7 +15,7 @@ namespace resolver class ResolverAmdAutolykosV2 : public resolver::ResolverAmd { public: - ResolverAmdAutolykosV2() = default; + ResolverAmdAutolykosV2(); ~ResolverAmdAutolykosV2(); bool updateMemory(stratum::StratumJobInfo const& jobInfo) final; diff --git a/sources/resolver/amd/etchash.cpp b/sources/resolver/amd/etchash.cpp index 8ea9fcd..63ab816 100644 --- a/sources/resolver/amd/etchash.cpp +++ b/sources/resolver/amd/etchash.cpp @@ -11,24 +11,29 @@ #include +resolver::ResolverAmdEtchash::ResolverAmdEtchash(): + resolver::ResolverAmdEthash() +{ + algorithm = algo::ALGORITHM::ETHASH; +} + + bool resolver::ResolverAmdEtchash::updateContext( stratum::StratumJobInfo const& jobInfo) { /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; - - /////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::ContextGenerator::instance().build ( + algorithm, context, jobInfo.epoch, algo::ethash::EIP1099_MAX_EPOCH_NUMBER, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true /*config.deviceAlgorithm.ethashBuildLightCacheCPU*/ ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull diff --git a/sources/resolver/amd/etchash.hpp b/sources/resolver/amd/etchash.hpp index 8e2d062..29a7f7a 100644 --- a/sources/resolver/amd/etchash.hpp +++ b/sources/resolver/amd/etchash.hpp @@ -14,7 +14,7 @@ namespace resolver class ResolverAmdEtchash : public resolver::ResolverAmdEthash { public: - ResolverAmdEtchash() = default; + ResolverAmdEtchash(); ~ResolverAmdEtchash() = default; protected: diff --git a/sources/resolver/amd/ethash.cpp b/sources/resolver/amd/ethash.cpp index 7f82ea0..9226d98 100644 --- a/sources/resolver/amd/ethash.cpp +++ b/sources/resolver/amd/ethash.cpp @@ -11,6 +11,16 @@ #include +resolver::ResolverAmdEthash::ResolverAmdEthash(): + resolver::ResolverAmd() +{ + if (algorithm == algo::ALGORITHM::UNKNOWN) + { + algorithm = algo::ALGORITHM::ETHASH; + } +} + + resolver::ResolverAmdEthash::~ResolverAmdEthash() { parameters.lightCache.free(); @@ -25,21 +35,18 @@ bool resolver::ResolverAmdEthash::updateContext( stratum::StratumJobInfo const& jobInfo) { /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; - - /////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::ContextGenerator::instance().build ( + algorithm, context, jobInfo.epoch, algo::ethash::MAX_EPOCH_NUMBER, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true /*config.deviceAlgorithm.ethashBuildLightCacheCPU*/ ); - // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context, true); if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull @@ -126,7 +133,7 @@ bool resolver::ResolverAmdEthash::updateMemory( } //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(context); + algo::ethash::ContextGenerator::instance().free(algorithm); //////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/amd/ethash.hpp b/sources/resolver/amd/ethash.hpp index 7e3920b..0ece2fe 100644 --- a/sources/resolver/amd/ethash.hpp +++ b/sources/resolver/amd/ethash.hpp @@ -16,7 +16,7 @@ namespace resolver class ResolverAmdEthash : public resolver::ResolverAmd { public: - ResolverAmdEthash() = default; + ResolverAmdEthash(); ~ResolverAmdEthash(); bool updateMemory(stratum::StratumJobInfo const& jobInfo) final; diff --git a/sources/resolver/amd/evrprogpow.cpp b/sources/resolver/amd/evrprogpow.cpp index 9b5d35c..8e791b5 100644 --- a/sources/resolver/amd/evrprogpow.cpp +++ b/sources/resolver/amd/evrprogpow.cpp @@ -6,11 +6,16 @@ resolver::ResolverAmdEvrprogPOW::ResolverAmdEvrprogPOW(): resolver::ResolverAmdProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::EVRPROGPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::MAX_EPOCH_NUMBER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; dagCountItemsInit = algo::evrprogpow::DAG_COUNT_ITEMS_INIT; + /////////////////////////////////////////////////////////////////////////// // KawPow progpowVersion = algo::progpow::VERSION::EVRPROGPOW; dagItemParents = algo::evrprogpow::DAG_ITEM_PARENTS; diff --git a/sources/resolver/amd/firopow.cpp b/sources/resolver/amd/firopow.cpp index d80a180..896cbff 100644 --- a/sources/resolver/amd/firopow.cpp +++ b/sources/resolver/amd/firopow.cpp @@ -6,11 +6,16 @@ resolver::ResolverAmdFiroPOW::ResolverAmdFiroPOW(): resolver::ResolverAmdProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::FIROPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::MAX_EPOCH_NUMBER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; dagCountItemsInit = algo::firopow::DAG_COUNT_ITEMS_INIT; + /////////////////////////////////////////////////////////////////////////// // KawPow progpowVersion = algo::progpow::VERSION::FIROPOW; dagItemParents = algo::firopow::DAG_ITEM_PARENTS; diff --git a/sources/resolver/amd/kawpow.cpp b/sources/resolver/amd/kawpow.cpp index 3b06a07..8bc36f7 100644 --- a/sources/resolver/amd/kawpow.cpp +++ b/sources/resolver/amd/kawpow.cpp @@ -6,11 +6,16 @@ resolver::ResolverAmdKawPOW::ResolverAmdKawPOW(): resolver::ResolverAmdProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::KAWPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::MAX_EPOCH_NUMBER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; dagCountItemsInit = algo::ethash::DAG_COUNT_ITEMS_INIT; + /////////////////////////////////////////////////////////////////////////// // KawPow progpowVersion = algo::progpow::VERSION::KAWPOW; dagItemParents = algo::kawpow::DAG_ITEM_PARENTS; diff --git a/sources/resolver/amd/meowpow.cpp b/sources/resolver/amd/meowpow.cpp index 12b089d..215e1de 100644 --- a/sources/resolver/amd/meowpow.cpp +++ b/sources/resolver/amd/meowpow.cpp @@ -8,15 +8,21 @@ resolver::ResolverAmdMeowPOW::ResolverAmdMeowPOW(): resolver::ResolverAmdProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::MEOWPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::EIP1057_MAX_EPOCH_NUMER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; + /////////////////////////////////////////////////////////////////////////// // ProgPow progpowVersion = algo::progpow::VERSION::MEOWPOW; regs = algo::meowpow::REGS; moduleSource = algo::meowpow::MODULE_SOURCE; + /////////////////////////////////////////////////////////////////////////// // MeowPow dagItemParents = algo::meowpow::DAG_ITEM_PARENTS; countCache = algo::meowpow::COUNT_CACHE; diff --git a/sources/resolver/amd/progpow.cpp b/sources/resolver/amd/progpow.cpp index b89691a..ea6455a 100644 --- a/sources/resolver/amd/progpow.cpp +++ b/sources/resolver/amd/progpow.cpp @@ -8,6 +8,16 @@ #include +resolver::ResolverAmdProgPOW::ResolverAmdProgPOW(): + resolver::ResolverAmd() +{ + if (algorithm == algo::ALGORITHM::UNKNOWN) + { + algorithm = algo::ALGORITHM::PROGPOW; + } +} + + resolver::ResolverAmdProgPOW::~ResolverAmdProgPOW() { parameters.lightCache.free(); @@ -21,21 +31,18 @@ bool resolver::ResolverAmdProgPOW::updateContext( stratum::StratumJobInfo const& jobInfo) { /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; - - /////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::ContextGenerator::instance().build ( + algorithm, context, jobInfo.epoch, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true /*config.deviceAlgorithm.ethashBuildLightCacheCPU*/ ); - // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context, true); if ( 0ull == context.lightCache.numberItem || 0ull == context.lightCache.size @@ -114,7 +121,7 @@ bool resolver::ResolverAmdProgPOW::updateMemory( } //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(context); + algo::ethash::ContextGenerator::instance().free(algorithm); //////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/amd/progpow.hpp b/sources/resolver/amd/progpow.hpp index 6d9ca77..3f34bad 100644 --- a/sources/resolver/amd/progpow.hpp +++ b/sources/resolver/amd/progpow.hpp @@ -17,7 +17,7 @@ namespace resolver class ResolverAmdProgPOW : public resolver::ResolverAmd { public: - ResolverAmdProgPOW() = default; + ResolverAmdProgPOW(); ~ResolverAmdProgPOW(); bool updateMemory(stratum::StratumJobInfo const& jobInfo) final; diff --git a/sources/resolver/amd/progpow_quai.cpp b/sources/resolver/amd/progpow_quai.cpp index b526a7b..d995fc0 100644 --- a/sources/resolver/amd/progpow_quai.cpp +++ b/sources/resolver/amd/progpow_quai.cpp @@ -5,12 +5,17 @@ resolver::ResolverAmdProgpowQuai::ResolverAmdProgpowQuai(): resolver::ResolverAmdProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::PROGPOWQUAI; + + /////////////////////////////////////////////////////////////////////////// //Ethash dagItemParents = algo::progpow_quai::DAG_ITEM_PARENTS; dagCountItemsGrowth = algo::progpow_quai::DAG_COUNT_ITEMS_GROWTH; dagCountItemsInit = algo::progpow_quai::DAG_COUNT_ITEMS_INIT; lightCacheCountItemsGrowth = algo::progpow_quai::LIGHT_CACHE_COUNT_ITEMS_GROWTH; + /////////////////////////////////////////////////////////////////////////// // ProgpowQuai progpowVersion = algo::progpow::VERSION::PROGPOWQUAI; countCache = algo::progpow_quai::COUNT_CACHE; diff --git a/sources/resolver/amd/tests/ethash.cpp b/sources/resolver/amd/tests/ethash.cpp index 2c616bc..85f64ef 100644 --- a/sources/resolver/amd/tests/ethash.cpp +++ b/sources/resolver/amd/tests/ethash.cpp @@ -45,7 +45,7 @@ struct ResolverEthashAmdTest : public testing::Test jobInfo.seedHash = algo::toHash256("3c77b17f5e89ebc92dc434fd7d631d80e5b522f07c9a452bc735c05c152381a1"); jobInfo.boundary = algo::toHash256("0000000225c17d04dad2a2a5b11fb33e22c8af8733a66744290d88c89181a499"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::ethash::MAX_EPOCH_NUMBER); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::ethash::MAX_EPOCH_NUMBER); } }; diff --git a/sources/resolver/amd/tests/firopow.cpp b/sources/resolver/amd/tests/firopow.cpp index fb33940..7d6b3b5 100644 --- a/sources/resolver/amd/tests/firopow.cpp +++ b/sources/resolver/amd/tests/firopow.cpp @@ -46,7 +46,7 @@ struct ResolverFiropowAmdTest : public testing::Test jobInfo.seedHash = algo::toHash256("ac88bf0324754ae04cffe412accbbd72e534acd928bdcbe95239f660667ae26d"); jobInfo.boundary = algo::toHash256("00000003fffc0000000000000000000000000000000000000000000000000000"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::firopow::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::firopow::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::firopow::MAX_PERIOD; } diff --git a/sources/resolver/amd/tests/kawpow.cpp b/sources/resolver/amd/tests/kawpow.cpp index 705645f..6a657e1 100644 --- a/sources/resolver/amd/tests/kawpow.cpp +++ b/sources/resolver/amd/tests/kawpow.cpp @@ -49,7 +49,7 @@ struct ResolverKawpowAmdTest : public testing::Test jobInfo.seedHash = algo::toHash256("7c4fb8a5d141973b69b521ce76b0dc50f0d2834d817c7f8310a6ab5becc6bb0c"); jobInfo.boundary = algo::toHash256("00000000ffff0000000000000000000000000000000000000000000000000000"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::progpow::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::progpow::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::kawpow::MAX_PERIOD; } diff --git a/sources/resolver/amd/tests/meowpow.cpp b/sources/resolver/amd/tests/meowpow.cpp index 3547ce2..e9e85cf 100644 --- a/sources/resolver/amd/tests/meowpow.cpp +++ b/sources/resolver/amd/tests/meowpow.cpp @@ -49,7 +49,7 @@ struct ResolverMeowpowAmdTest : public testing::Test jobInfo.seedHash = algo::toHash256("cfa3e37c459ebd9b4138bd2141a52d89f6f8f671ecf91456f5a29176eb132fc0"); jobInfo.boundary = algo::toHash256("0000000500000000000000000000000000000000000000000000000000000000"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::progpow::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::progpow::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::meowpow::MAX_PERIOD; } }; diff --git a/sources/resolver/amd/tests/progpow_quai.cpp b/sources/resolver/amd/tests/progpow_quai.cpp index 4590858..d35cc44 100644 --- a/sources/resolver/amd/tests/progpow_quai.cpp +++ b/sources/resolver/amd/tests/progpow_quai.cpp @@ -49,7 +49,7 @@ struct ResolverProgpowQuaiAmdTest : public testing::Test jobInfo.seedHash = algo::toHash256("0000000000000000000000000000000000000000000000000000000000000000"); jobInfo.boundary = algo::toHash256("0000000225c17d04dad2965cc5a02a23e254c0c3f75d9178046aeb27ce1ca574"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::progpow_quai::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::progpow_quai::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::progpow_quai::MAX_PERIOD; } }; diff --git a/sources/resolver/amd/tests/progpow_z.cpp b/sources/resolver/amd/tests/progpow_z.cpp index e63b19c..c2ffd6b 100644 --- a/sources/resolver/amd/tests/progpow_z.cpp +++ b/sources/resolver/amd/tests/progpow_z.cpp @@ -49,7 +49,7 @@ struct ResolverProgpowZAmdTest : public testing::Test jobInfo.seedHash = algo::toHash256("0x71a56feffb6f10ea9d76e1a9464eb0abd86e4349ae98fb794923a65b650282a3"); jobInfo.boundary = algo::toHash256("0x00000002dd01fc067918c87bb22ae2da831babaff47cc1f55b39398a682631f0"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::ethash::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::ethash::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::progpow::v_0_9_2::MAX_PERIOD; } }; diff --git a/sources/resolver/cpu/kawpow.cpp b/sources/resolver/cpu/kawpow.cpp index a9fc3fa..bb6edc2 100644 --- a/sources/resolver/cpu/kawpow.cpp +++ b/sources/resolver/cpu/kawpow.cpp @@ -6,14 +6,20 @@ resolver::ResolverCpuKawPOW::ResolverCpuKawPOW(): resolver::ResolverCpuProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::KAWPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::MAX_EPOCH_NUMBER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; dagCountItemsInit = algo::ethash::DAG_COUNT_ITEMS_INIT; + /////////////////////////////////////////////////////////////////////////// // ProgPow progpowVersion = algo::progpow::VERSION::KAWPOW; + /////////////////////////////////////////////////////////////////////////// // KawPow dagItemParents = algo::kawpow::DAG_ITEM_PARENTS; countCache = algo::kawpow::COUNT_CACHE; diff --git a/sources/resolver/cpu/progpow.cpp b/sources/resolver/cpu/progpow.cpp index f1bab19..4a7b505 100644 --- a/sources/resolver/cpu/progpow.cpp +++ b/sources/resolver/cpu/progpow.cpp @@ -6,6 +6,16 @@ #include +resolver::ResolverCpuProgPOW::ResolverCpuProgPOW(): + resolver::ResolverCpu() +{ + if (algorithm == algo::ALGORITHM::UNKNOWN) + { + algorithm = algo::ALGORITHM::PROGPOW; + } +} + + resolver::ResolverCpuProgPOW::~ResolverCpuProgPOW() { SAFE_DELETE_ARRAY(parameters.headerCache); @@ -17,18 +27,18 @@ resolver::ResolverCpuProgPOW::~ResolverCpuProgPOW() bool resolver::ResolverCpuProgPOW::updateContext(stratum::StratumJobInfo const& jobInfo) { //////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::ContextGenerator::instance().build ( + algorithm, context, jobInfo.epoch, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true ); - // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context, true); if ( 0ull == context.lightCache.numberItem || 0ull == context.lightCache.size diff --git a/sources/resolver/cpu/progpow.hpp b/sources/resolver/cpu/progpow.hpp index 35eb73e..05cefa6 100644 --- a/sources/resolver/cpu/progpow.hpp +++ b/sources/resolver/cpu/progpow.hpp @@ -14,7 +14,7 @@ namespace resolver class ResolverCpuProgPOW : public resolver::ResolverCpu { public: - ResolverCpuProgPOW() = default; + ResolverCpuProgPOW(); virtual ~ResolverCpuProgPOW(); bool updateMemory(stratum::StratumJobInfo const& jobInfo) final; diff --git a/sources/resolver/mocker.cpp b/sources/resolver/mocker.cpp index f0e700e..7db4663 100644 --- a/sources/resolver/mocker.cpp +++ b/sources/resolver/mocker.cpp @@ -11,7 +11,6 @@ #include -static constexpr boost::chrono::milliseconds WAIT_UPDATE_MEMORY{ 40000 }; static constexpr boost::chrono::milliseconds WAIT_UPDATE_CONSTANT{ 2000 }; static constexpr boost::chrono::milliseconds WAIT_UPDATE_EXECUTE_SYNC{ 1000 }; static constexpr boost::chrono::milliseconds WAIT_UPDATE_EXECUTE_ASYNC{ 1000 }; @@ -35,25 +34,33 @@ bool resolver::ResolverMocker::updateMemory( uint32_t maxEpoch{ algo::ethash::MAX_EPOCH_NUMBER }; uint32_t lightCacheCountItemsGrowth{ algo::ethash::LIGHT_CACHE_COUNT_ITEMS_GROWTH }; uint32_t lightCacheCountItemsInit{ algo::ethash::LIGHT_CACHE_COUNT_ITEMS_INIT }; - uint32_t dagItemParents{ algo::ethash::DAG_ITEM_PARENTS }; uint32_t dagCountItemsGrowth{ algo::ethash::DAG_COUNT_ITEMS_GROWTH }; uint32_t dagCountItemsInit{ algo::ethash::DAG_COUNT_ITEMS_INIT }; - uint32_t countCache{ algo::progpow::v_0_9_3::COUNT_CACHE }; - uint32_t countMath{ algo::progpow::v_0_9_3::COUNT_MATH }; - algo::ethash::initializeDagContext + + /////////////////////////////////////////////////////////////////////////// + algo::ethash::ContextGenerator::instance().build ( + algorithm, context, jobInfo.epoch, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + config.deviceAlgorithm.ethashBuildLightCacheCPU ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); /////////////////////////////////////////////////////////////////////////// - boost::this_thread::sleep_for(WAIT_UPDATE_MEMORY); + if (std::nullopt != config.toolConfigs.mockerResolverUpdateMemorySleep) + { + resolverInfo() + << "Update memory force waiting: " + << *config.toolConfigs.mockerResolverUpdateMemorySleep + << "ms"; + boost::chrono::milliseconds const ms{ *config.toolConfigs.mockerResolverUpdateMemorySleep }; + boost::this_thread::sleep_for(ms); + } /////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/nvidia/autolykos_v2.cpp b/sources/resolver/nvidia/autolykos_v2.cpp index da0c14a..70b52f6 100644 --- a/sources/resolver/nvidia/autolykos_v2.cpp +++ b/sources/resolver/nvidia/autolykos_v2.cpp @@ -8,6 +8,13 @@ #include +resolver::ResolverNvidiaAutolykosV2::ResolverNvidiaAutolykosV2(): + resolver::ResolverNvidia() +{ + algorithm = algo::ALGORITHM::AUTOLYKOS_V2; +} + + resolver::ResolverNvidiaAutolykosV2::~ResolverNvidiaAutolykosV2() { autolykosv2FreeMemory(parameters); diff --git a/sources/resolver/nvidia/autolykos_v2.hpp b/sources/resolver/nvidia/autolykos_v2.hpp index fefb69f..b7e62f6 100644 --- a/sources/resolver/nvidia/autolykos_v2.hpp +++ b/sources/resolver/nvidia/autolykos_v2.hpp @@ -14,7 +14,7 @@ namespace resolver class ResolverNvidiaAutolykosV2 : public resolver::ResolverNvidia { public: - ResolverNvidiaAutolykosV2() = default; + ResolverNvidiaAutolykosV2(); virtual ~ResolverNvidiaAutolykosV2(); bool updateMemory(stratum::StratumJobInfo const& jobInfo) final; diff --git a/sources/resolver/nvidia/blake3.cpp b/sources/resolver/nvidia/blake3.cpp index 9134c0d..8870148 100644 --- a/sources/resolver/nvidia/blake3.cpp +++ b/sources/resolver/nvidia/blake3.cpp @@ -6,6 +6,13 @@ #include +resolver::ResolverNvidiaBlake3::ResolverNvidiaBlake3(): + resolver::ResolverNvidia() +{ + algorithm = algo::ALGORITHM::BLAKE3; +} + + resolver::ResolverNvidiaBlake3::~ResolverNvidiaBlake3() { blake3FreeMemory(parameters); diff --git a/sources/resolver/nvidia/blake3.hpp b/sources/resolver/nvidia/blake3.hpp index 61b537e..a7359b8 100644 --- a/sources/resolver/nvidia/blake3.hpp +++ b/sources/resolver/nvidia/blake3.hpp @@ -14,7 +14,7 @@ namespace resolver class ResolverNvidiaBlake3 : public resolver::ResolverNvidia { public: - ResolverNvidiaBlake3() = default; + ResolverNvidiaBlake3(); ~ResolverNvidiaBlake3(); bool updateMemory(stratum::StratumJobInfo const& jobInfo) final; diff --git a/sources/resolver/nvidia/etchash.cpp b/sources/resolver/nvidia/etchash.cpp index 865632c..8d86fa4 100644 --- a/sources/resolver/nvidia/etchash.cpp +++ b/sources/resolver/nvidia/etchash.cpp @@ -9,6 +9,13 @@ #include +resolver::ResolverNvidiaEtchash::ResolverNvidiaEtchash(): + resolver::ResolverNvidiaEthash() +{ + algorithm = algo::ALGORITHM::ETCHASH; +} + + bool resolver::ResolverNvidiaEtchash::updateContext( stratum::StratumJobInfo const& jobInfo) { @@ -16,17 +23,18 @@ bool resolver::ResolverNvidiaEtchash::updateContext( common::Config& config{ common::Config::instance() }; /////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::ContextGenerator::instance().build ( + algorithm, context, jobInfo.epoch, algo::ethash::EIP1099_MAX_EPOCH_NUMBER, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + config.deviceAlgorithm.ethashBuildLightCacheCPU ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); /////////////////////////////////////////////////////////////////////////// if ( context.lightCache.numberItem == 0ull diff --git a/sources/resolver/nvidia/etchash.hpp b/sources/resolver/nvidia/etchash.hpp index 89e8eff..61244a0 100644 --- a/sources/resolver/nvidia/etchash.hpp +++ b/sources/resolver/nvidia/etchash.hpp @@ -14,7 +14,7 @@ namespace resolver class ResolverNvidiaEtchash : public resolver::ResolverNvidiaEthash { public: - ResolverNvidiaEtchash() = default; + ResolverNvidiaEtchash(); ~ResolverNvidiaEtchash() = default; protected: diff --git a/sources/resolver/nvidia/ethash.cpp b/sources/resolver/nvidia/ethash.cpp index e9bd014..de43b21 100644 --- a/sources/resolver/nvidia/ethash.cpp +++ b/sources/resolver/nvidia/ethash.cpp @@ -9,6 +9,16 @@ #include +resolver::ResolverNvidiaEthash::ResolverNvidiaEthash(): + resolver::ResolverNvidia() +{ + if (algorithm == algo::ALGORITHM::UNKNOWN) + { + algorithm = algo::ALGORITHM::ETHASH; + } +} + + resolver::ResolverNvidiaEthash::~ResolverNvidiaEthash() { ethashFreeMemory(parameters); @@ -21,18 +31,21 @@ bool resolver::ResolverNvidiaEthash::updateContext( /////////////////////////////////////////////////////////////////////////// common::Config& config{ common::Config::instance() }; - algo::ethash::initializeDagContext + /////////////////////////////////////////////////////////////////////////// + algo::ethash::ContextGenerator::instance().build ( + algorithm, context, jobInfo.epoch, algo::ethash::MAX_EPOCH_NUMBER, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + config.deviceAlgorithm.ethashBuildLightCacheCPU ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); + /////////////////////////////////////////////////////////////////////////// if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull || context.dagCache.numberItem == 0ull @@ -50,6 +63,7 @@ bool resolver::ResolverNvidiaEthash::updateContext( return false; } + /////////////////////////////////////////////////////////////////////////// uint64_t const totalMemoryNeeded{ (context.dagCache.size + context.lightCache.size) }; if ( 0ull != deviceMemoryAvailable && totalMemoryNeeded >= deviceMemoryAvailable) @@ -60,6 +74,7 @@ bool resolver::ResolverNvidiaEthash::updateContext( return false; } + /////////////////////////////////////////////////////////////////////////// return true; } @@ -88,17 +103,24 @@ bool resolver::ResolverNvidiaEthash::updateMemory( //////////////////////////////////////////////////////////////////////////// if (false == config.deviceAlgorithm.ethashBuildLightCacheCPU) { + /////////////////////////////////////////////////////////////////////// resolverInfo() << "Building light cache on GPU"; - common::ChronoGuard chronoCPU{ "Built light cache", common::CHRONO_UNIT::MS }; + chrono.start(); if (false == ethashBuildLightCache(cuStream[currentIndexStream], parameters.seedCache)) { return false; } - } + chrono.start(); + resolverInfo() << "Built light cache on GPU in " << chrono.elapsed(common::CHRONO_UNIT::MS) << "ms"; - //////////////////////////////////////////////////////////////////////////// - CU_SAFE_DELETE(parameters.seedCache); + /////////////////////////////////////////////////////////////////////// + CU_SAFE_DELETE(parameters.seedCache); + } + else + { + algo::ethash::ContextGenerator::instance().free(algorithm); + } //////////////////////////////////////////////////////////////////////////// resolverInfo() << "Building DAG"; @@ -116,7 +138,7 @@ bool resolver::ResolverNvidiaEthash::updateMemory( CU_SAFE_DELETE(parameters.lightCache); //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(context); + algo::ethash::ContextGenerator::instance().free(algorithm); //////////////////////////////////////////////////////////////////////////// return true; @@ -126,6 +148,7 @@ bool resolver::ResolverNvidiaEthash::updateMemory( bool resolver::ResolverNvidiaEthash::updateConstants( stratum::StratumJobInfo const& jobInfo) { + //////////////////////////////////////////////////////////////////////////// uint32_t const* const header { jobInfo.headerHash.word32 }; uint64_t const boundary { jobInfo.boundaryU64 }; if (false == ethashUpdateConstants(header, boundary)) @@ -136,6 +159,7 @@ bool resolver::ResolverNvidiaEthash::updateConstants( //////////////////////////////////////////////////////////////////////////// overrideOccupancy(128u, 8192u); + //////////////////////////////////////////////////////////////////////////// return true; } @@ -143,6 +167,7 @@ bool resolver::ResolverNvidiaEthash::updateConstants( bool resolver::ResolverNvidiaEthash::executeSync( stratum::StratumJobInfo const& jobInfo) { + //////////////////////////////////////////////////////////////////////////// ethashSearch(cuStream[currentIndexStream], ¶meters.resultCache[currentIndexStream], blocks, @@ -151,6 +176,7 @@ bool resolver::ResolverNvidiaEthash::executeSync( CUDA_ER(cudaStreamSynchronize(cuStream[currentIndexStream])); CUDA_ER(cudaGetLastError()); + //////////////////////////////////////////////////////////////////////////// algo::ethash::Result* resultCache{ ¶meters.resultCache[currentIndexStream] }; if (true == resultCache->found) { @@ -173,6 +199,7 @@ bool resolver::ResolverNvidiaEthash::executeSync( resultCache->count = 0u; } + //////////////////////////////////////////////////////////////////////////// return true; } @@ -219,6 +246,7 @@ bool resolver::ResolverNvidiaEthash::executeAsync( //////////////////////////////////////////////////////////////////////////// swapIndexStream(); + //////////////////////////////////////////////////////////////////////////// return true; } diff --git a/sources/resolver/nvidia/ethash.hpp b/sources/resolver/nvidia/ethash.hpp index d19a8d0..cfec87c 100644 --- a/sources/resolver/nvidia/ethash.hpp +++ b/sources/resolver/nvidia/ethash.hpp @@ -15,7 +15,7 @@ namespace resolver class ResolverNvidiaEthash : public resolver::ResolverNvidia { public: - ResolverNvidiaEthash() = default; + ResolverNvidiaEthash(); ~ResolverNvidiaEthash(); bool updateMemory(stratum::StratumJobInfo const& jobInfo) final; diff --git a/sources/resolver/nvidia/evrprogpow.cpp b/sources/resolver/nvidia/evrprogpow.cpp index 1baa419..2fff3d7 100644 --- a/sources/resolver/nvidia/evrprogpow.cpp +++ b/sources/resolver/nvidia/evrprogpow.cpp @@ -6,11 +6,16 @@ resolver::ResolverNvidiaEvrprogPOW::ResolverNvidiaEvrprogPOW(): resolver::ResolverNvidiaProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::EVRPROGPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::MAX_EPOCH_NUMBER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; dagCountItemsInit = algo::evrprogpow::DAG_COUNT_ITEMS_INIT; + /////////////////////////////////////////////////////////////////////////// // EvrprogPow progpowVersion = algo::progpow::VERSION::EVRPROGPOW; dagItemParents = algo::evrprogpow::DAG_ITEM_PARENTS; diff --git a/sources/resolver/nvidia/firopow.cpp b/sources/resolver/nvidia/firopow.cpp index 54f0e89..bfc3d22 100644 --- a/sources/resolver/nvidia/firopow.cpp +++ b/sources/resolver/nvidia/firopow.cpp @@ -6,11 +6,16 @@ resolver::ResolverNvidiaFiroPOW::ResolverNvidiaFiroPOW(): resolver::ResolverNvidiaProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::FIROPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::MAX_EPOCH_NUMBER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; dagCountItemsInit = algo::firopow::DAG_COUNT_ITEMS_INIT; + /////////////////////////////////////////////////////////////////////////// // FiroPow progpowVersion = algo::progpow::VERSION::FIROPOW; dagItemParents = algo::firopow::DAG_ITEM_PARENTS; diff --git a/sources/resolver/nvidia/kawpow.cpp b/sources/resolver/nvidia/kawpow.cpp index 9580f56..2754f0e 100644 --- a/sources/resolver/nvidia/kawpow.cpp +++ b/sources/resolver/nvidia/kawpow.cpp @@ -6,14 +6,20 @@ resolver::ResolverNvidiaKawPOW::ResolverNvidiaKawPOW(): resolver::ResolverNvidiaProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::KAWPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::MAX_EPOCH_NUMBER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; dagCountItemsInit = algo::ethash::DAG_COUNT_ITEMS_INIT; + /////////////////////////////////////////////////////////////////////////// // ProgPow progpowVersion = algo::progpow::VERSION::KAWPOW; + /////////////////////////////////////////////////////////////////////////// // KawPow dagItemParents = algo::kawpow::DAG_ITEM_PARENTS; countCache = algo::kawpow::COUNT_CACHE; diff --git a/sources/resolver/nvidia/meowpow.cpp b/sources/resolver/nvidia/meowpow.cpp index a43779b..27e1c1b 100644 --- a/sources/resolver/nvidia/meowpow.cpp +++ b/sources/resolver/nvidia/meowpow.cpp @@ -6,15 +6,21 @@ resolver::ResolverNvidiaMeowPOW::ResolverNvidiaMeowPOW(): resolver::ResolverNvidiaProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::MEOWPOW; + + /////////////////////////////////////////////////////////////////////////// // Ethash maxEpoch = algo::ethash::EIP1057_MAX_EPOCH_NUMER; dagCountItemsGrowth = algo::ethash::DAG_COUNT_ITEMS_GROWTH; + /////////////////////////////////////////////////////////////////////////// // ProgPow progpowVersion = algo::progpow::VERSION::MEOWPOW; regs = algo::meowpow::REGS; moduleSource = algo::meowpow::MODULE_SOURCE; + /////////////////////////////////////////////////////////////////////////// // MeowPow dagItemParents = algo::meowpow::DAG_ITEM_PARENTS; countCache = algo::meowpow::COUNT_CACHE; diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index 460cfc4..ee18371 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -10,6 +10,16 @@ #include +resolver::ResolverNvidiaProgPOW::ResolverNvidiaProgPOW(): + resolver::ResolverNvidia() +{ + if (algorithm == algo::ALGORITHM::UNKNOWN) + { + algorithm = algo::ALGORITHM::PROGPOWQUAI; + } +} + + resolver::ResolverNvidiaProgPOW::~ResolverNvidiaProgPOW() { progpowFreeMemory(parameters); @@ -23,18 +33,20 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( common::Config& config{ common::Config::instance() }; //////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::ContextGenerator::instance().build ( + algorithm, context, jobInfo.epoch, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + config.deviceAlgorithm.ethashBuildLightCacheCPU ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); + //////////////////////////////////////////////////////////////////////////// if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull || context.dagCache.numberItem == 0ull @@ -52,7 +64,7 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( return false; } - + //////////////////////////////////////////////////////////////////////////// uint64_t const totalMemoryNeeded{ context.dagCache.size + context.lightCache.size }; if ( 0ull != deviceMemoryAvailable && totalMemoryNeeded >= deviceMemoryAvailable) @@ -64,6 +76,7 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( return false; } + //////////////////////////////////////////////////////////////////////////// return true; } @@ -92,6 +105,7 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( //////////////////////////////////////////////////////////////////////////// if (false == config.deviceAlgorithm.ethashBuildLightCacheCPU) { + /////////////////////////////////////////////////////////////////////// resolverInfo() << "Building LightCache on GPU"; chrono.start(); if (false == progpowBuildLightCache(cuStream[currentIndexStream], @@ -100,13 +114,14 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( return false; } chrono.stop(); - resolverInfo() << "Light Cache built in " << chrono.elapsed(common::CHRONO_UNIT::MS) << "ms"; - //////////////////////////////////////////////////////////////////////////// + resolverInfo() << "Light Cache built on GPU in " << chrono.elapsed(common::CHRONO_UNIT::MS) << "ms"; + + /////////////////////////////////////////////////////////////////////// CU_SAFE_DELETE(parameters.seedCache); } else { - + algo::ethash::ContextGenerator::instance().free(algorithm); } //////////////////////////////////////////////////////////////////////////// @@ -124,9 +139,6 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( //////////////////////////////////////////////////////////////////////////// CU_SAFE_DELETE(parameters.lightCache); - //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(context); - //////////////////////////////////////////////////////////////////////////// return true; } @@ -344,13 +356,15 @@ bool resolver::ResolverNvidiaProgPOW::executeSync( }; //////////////////////////////////////////////////////////////////////////// - CU_ER(cuLaunchKernel(kernelGenerator.cuFunction, - blocks, 1u, 1u, - threads, 1u, 1u, - 0u, - cuStream[currentIndexStream], - arguments, - nullptr)); + CU_ER( + cuLaunchKernel( + kernelGenerator.cuFunction, + blocks, 1u, 1u, + threads, 1u, 1u, + 0u, + cuStream[currentIndexStream], + arguments, + nullptr)); CUDA_ER(cudaStreamSynchronize(cuStream[currentIndexStream])); CUDA_ER(cudaGetLastError()); @@ -367,13 +381,9 @@ bool resolver::ResolverNvidiaProgPOW::executeSync( resultShare.jobId = jobInfo.jobIDStr; resultShare.extraNonceSize = jobInfo.extraNonceSize; - for (uint32_t i { 0u }; i < count; ++i) + for (uint32_t i{ 0u }; i < count; ++i) { resultShare.nonces[i] = result->nonces[i]; - } - - for (uint32_t i { 0u }; i < count; ++i) - { for (uint32_t j{ 0u }; j < algo::LEN_HASH_256_WORD_32; ++j) { resultShare.hash[i][j] = result->hash[i][j]; @@ -409,17 +419,19 @@ bool resolver::ResolverNvidiaProgPOW::executeAsync( ¶meters.dagCache, &result }; - CU_ER(cuLaunchKernel(kernelGenerator.cuFunction, - blocks, 1u, 1u, - threads, 1u, 1u, - 0u, - cuStream[currentIndexStream], - arguments, - nullptr)); + CU_ER( + cuLaunchKernel( + kernelGenerator.cuFunction, + blocks, 1u, 1u, + threads, 1u, 1u, + 0u, + cuStream[currentIndexStream], + arguments, + nullptr)); //////////////////////////////////////////////////////////////////////////// swapIndexStream(); - algo::progpow::Result* resultCache { ¶meters.resultCache[currentIndexStream] }; + algo::progpow::Result* resultCache{ ¶meters.resultCache[currentIndexStream] }; if (true == resultCache->found) { uint32_t const count @@ -432,13 +444,9 @@ bool resolver::ResolverNvidiaProgPOW::executeAsync( resultShare.jobId = jobInfo.jobIDStr; resultShare.extraNonceSize = jobInfo.extraNonceSize; - for (uint32_t i { 0u }; i < count; ++i) + for (uint32_t i{ 0u }; i < count; ++i) { resultShare.nonces[i] = resultCache->nonces[i]; - } - - for (uint32_t i { 0u }; i < count; ++i) - { for (uint32_t j{ 0u }; j < algo::LEN_HASH_256_WORD_32; ++j) { resultShare.hash[i][j] = resultCache->hash[i][j]; @@ -456,6 +464,7 @@ bool resolver::ResolverNvidiaProgPOW::executeAsync( return true; } + void resolver::ResolverNvidiaProgPOW::submit( stratum::Stratum* const stratum) { @@ -463,10 +472,10 @@ void resolver::ResolverNvidiaProgPOW::submit( { if (false == isStale(resultShare.jobId)) { - for (uint32_t i { 0u }; i < resultShare.count; ++i) + for (uint32_t i{ 0u }; i < resultShare.count; ++i) { uint32_t hash[algo::LEN_HASH_256_WORD_32]{}; - for (uint32_t j { 0u }; j < algo::LEN_HASH_256_WORD_32; ++j) + for (uint32_t j{ 0u }; j < algo::LEN_HASH_256_WORD_32; ++j) { hash[j] = resultShare.hash[i][j]; } @@ -544,10 +553,10 @@ void resolver::ResolverNvidiaProgPOW::submit( { if (false == isStale(resultShare.jobId)) { - for (uint32_t i { 0u }; i < resultShare.count; ++i) + for (uint32_t i{ 0u }; i < resultShare.count; ++i) { uint32_t hash[algo::LEN_HASH_256_WORD_32]{}; - for (uint32_t j { 0u }; j < algo::LEN_HASH_256_WORD_32; ++j) + for (uint32_t j{ 0u }; j < algo::LEN_HASH_256_WORD_32; ++j) { hash[j] = resultShare.hash[i][j]; } diff --git a/sources/resolver/nvidia/progpow.hpp b/sources/resolver/nvidia/progpow.hpp index 2113e88..7fc6c89 100644 --- a/sources/resolver/nvidia/progpow.hpp +++ b/sources/resolver/nvidia/progpow.hpp @@ -17,7 +17,7 @@ namespace resolver class ResolverNvidiaProgPOW : public resolver::ResolverNvidia { public: - ResolverNvidiaProgPOW() = default; + ResolverNvidiaProgPOW(); virtual ~ResolverNvidiaProgPOW(); bool updateMemory(stratum::StratumJobInfo const& jobInfo) final; diff --git a/sources/resolver/nvidia/progpow_quai.cpp b/sources/resolver/nvidia/progpow_quai.cpp index 446f0f0..e78ba93 100644 --- a/sources/resolver/nvidia/progpow_quai.cpp +++ b/sources/resolver/nvidia/progpow_quai.cpp @@ -5,6 +5,9 @@ resolver::ResolverNvidiaProgpowQuai::ResolverNvidiaProgpowQuai(): resolver::ResolverNvidiaProgPOW() { + /////////////////////////////////////////////////////////////////////////// + algorithm = algo::ALGORITHM::PROGPOWQUAI; + //Ethash dagItemParents = algo::progpow_quai::DAG_ITEM_PARENTS; dagCountItemsGrowth = algo::progpow_quai::DAG_COUNT_ITEMS_GROWTH; diff --git a/sources/resolver/nvidia/tests/autolykos_v2.cpp b/sources/resolver/nvidia/tests/autolykos_v2.cpp index 2df39ad..c1b05d3 100644 --- a/sources/resolver/nvidia/tests/autolykos_v2.cpp +++ b/sources/resolver/nvidia/tests/autolykos_v2.cpp @@ -31,7 +31,7 @@ struct ResolverAutolykosv2NvidiaTest : public testing::Test ~ResolverAutolykosv2NvidiaTest() { - resolver::tests::cleanUpCuda(); + resolver::tests::cleanUpCuda(properties); } void initializeJob(uint64_t const nonce) diff --git a/sources/resolver/nvidia/tests/blake3.cpp b/sources/resolver/nvidia/tests/blake3.cpp index 49a1989..5e69493 100644 --- a/sources/resolver/nvidia/tests/blake3.cpp +++ b/sources/resolver/nvidia/tests/blake3.cpp @@ -31,7 +31,7 @@ struct ResolverBlake3NvidiaTest : public testing::Test ~ResolverBlake3NvidiaTest() { - resolver::tests::cleanUpCuda(); + resolver::tests::cleanUpCuda(properties); } void initializeJob(uint64_t const nonce) diff --git a/sources/resolver/nvidia/tests/ethash.cpp b/sources/resolver/nvidia/tests/ethash.cpp index 4ace9a3..b7221ed 100644 --- a/sources/resolver/nvidia/tests/ethash.cpp +++ b/sources/resolver/nvidia/tests/ethash.cpp @@ -40,7 +40,7 @@ struct ResolverEthashNvidiaTest : public testing::Test ~ResolverEthashNvidiaTest() { - resolver::tests::cleanUpCuda(); + resolver::tests::cleanUpCuda(properties); } void initializeJob(uint64_t const nonce) @@ -50,7 +50,7 @@ struct ResolverEthashNvidiaTest : public testing::Test jobInfo.seedHash = algo::toHash256("3c77b17f5e89ebc92dc434fd7d631d80e5b522f07c9a452bc735c05c152381a1"); jobInfo.boundary = algo::toHash256("0000000225c17d04dad2a2a5b11fb33e22c8af8733a66744290d88c89181a499"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::ethash::MAX_EPOCH_NUMBER); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::ethash::MAX_EPOCH_NUMBER); } }; diff --git a/sources/resolver/nvidia/tests/firopow.cpp b/sources/resolver/nvidia/tests/firopow.cpp index 046d8a2..5c3a64d 100644 --- a/sources/resolver/nvidia/tests/firopow.cpp +++ b/sources/resolver/nvidia/tests/firopow.cpp @@ -30,7 +30,7 @@ struct ResolverFiropowNvidiaTest : public testing::Test ~ResolverFiropowNvidiaTest() { - resolver::tests::cleanUpCuda(); + resolver::tests::cleanUpCuda(properties); } void initializeJob(uint64_t const nonce) @@ -41,7 +41,7 @@ struct ResolverFiropowNvidiaTest : public testing::Test jobInfo.seedHash = algo::toHash256("ac88bf0324754ae04cffe412accbbd72e534acd928bdcbe95239f660667ae26d"); jobInfo.boundary = algo::toHash256("00000003fffc0000000000000000000000000000000000000000000000000000"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::firopow::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::firopow::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::firopow::MAX_PERIOD; } }; diff --git a/sources/resolver/nvidia/tests/kawpow.cpp b/sources/resolver/nvidia/tests/kawpow.cpp index bceedab..66877c2 100644 --- a/sources/resolver/nvidia/tests/kawpow.cpp +++ b/sources/resolver/nvidia/tests/kawpow.cpp @@ -45,7 +45,7 @@ struct ResolverKawpowNvidiaTest : public testing::Test config.occupancy.blocks = std::nullopt; //////////////////////////////////////////////////////////////////////////// - resolver::tests::cleanUpCuda(); + resolver::tests::cleanUpCuda(properties); } void setOccupancy( @@ -76,7 +76,7 @@ struct ResolverKawpowNvidiaTest : public testing::Test jobInfo.seedHash = algo::toHash256("7c4fb8a5d141973b69b521ce76b0dc50f0d2834d817c7f8310a6ab5becc6bb0c"); jobInfo.boundary = algo::toHash256("00000000ffff0000000000000000000000000000000000000000000000000000"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::progpow::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::progpow::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::kawpow::MAX_PERIOD; } }; @@ -88,18 +88,21 @@ TEST_F(ResolverKawpowNvidiaTest, findNonceWithLightCacheGPU) common::Config& config{ common::Config::instance() }; config.deviceAlgorithm.ethashBuildLightCacheCPU = false; + //////////////////////////////////////////////////////////////////////////// initializeJob(0xce00000017f87f70); + //////////////////////////////////////////////////////////////////////////// ASSERT_NE(nullptr, resolver.cuStream[0]); ASSERT_TRUE(resolver.updateMemory(jobInfo)); ASSERT_TRUE(resolver.updateConstants(jobInfo)); ASSERT_TRUE(resolver.executeSync(jobInfo)); - resolver.submit(&stratum); + //////////////////////////////////////////////////////////////////////////// + resolver.submit(&stratum); ASSERT_FALSE(stratum.paramSubmit.empty()); - std::string const nonceStr { stratum.paramSubmit[1].as_string().c_str() }; + //////////////////////////////////////////////////////////////////////////// using namespace std::string_literals; EXPECT_EQ("0xce00000017f87f7a"s, nonceStr); } @@ -107,18 +110,25 @@ TEST_F(ResolverKawpowNvidiaTest, findNonceWithLightCacheGPU) TEST_F(ResolverKawpowNvidiaTest, findNonceWithLightCacheCPU) { + //////////////////////////////////////////////////////////////////////////// + common::Config& config{ common::Config::instance() }; + config.deviceAlgorithm.ethashBuildLightCacheCPU = true; + + //////////////////////////////////////////////////////////////////////////// initializeJob(0xce00000017f87f70); + //////////////////////////////////////////////////////////////////////////// ASSERT_NE(nullptr, resolver.cuStream[0]); ASSERT_TRUE(resolver.updateMemory(jobInfo)); ASSERT_TRUE(resolver.updateConstants(jobInfo)); ASSERT_TRUE(resolver.executeSync(jobInfo)); - resolver.submit(&stratum); + //////////////////////////////////////////////////////////////////////////// + resolver.submit(&stratum); ASSERT_FALSE(stratum.paramSubmit.empty()); - std::string const nonceStr { stratum.paramSubmit[1].as_string().c_str() }; + //////////////////////////////////////////////////////////////////////////// using namespace std::string_literals; EXPECT_EQ("0xce00000017f87f7a"s, nonceStr); } @@ -137,9 +147,9 @@ TEST_F(ResolverKawpowNvidiaTest, findNonceInternalLoop) ASSERT_TRUE(resolver.updateMemory(jobInfo)); ASSERT_TRUE(resolver.updateConstants(jobInfo)); ASSERT_TRUE(resolver.executeSync(jobInfo)); - resolver.submit(&stratum); //////////////////////////////////////////////////////////////////////////// + resolver.submit(&stratum); ASSERT_FALSE(stratum.paramSubmit.empty()); //////////////////////////////////////////////////////////////////////////// @@ -162,9 +172,9 @@ TEST_F(ResolverKawpowNvidiaTest, findNonceLittleOccupancy) ASSERT_TRUE(resolver.updateMemory(jobInfo)); ASSERT_TRUE(resolver.updateConstants(jobInfo)); ASSERT_TRUE(resolver.executeSync(jobInfo)); - resolver.submit(&stratum); //////////////////////////////////////////////////////////////////////////// + resolver.submit(&stratum); ASSERT_FALSE(stratum.paramSubmit.empty()); //////////////////////////////////////////////////////////////////////////// @@ -186,9 +196,9 @@ TEST_F(ResolverKawpowNvidiaTest, aroundFindNonce) ASSERT_TRUE(resolver.updateMemory(jobInfo)); ASSERT_TRUE(resolver.updateConstants(jobInfo)); ASSERT_TRUE(resolver.executeSync(jobInfo)); - resolver.submit(&stratum); //////////////////////////////////////////////////////////////////////////// + resolver.submit(&stratum); ASSERT_FALSE(stratum.paramSubmit.empty()); //////////////////////////////////////////////////////////////////////////// @@ -210,6 +220,8 @@ TEST_F(ResolverKawpowNvidiaTest, notFindNonce) ASSERT_TRUE(resolver.updateMemory(jobInfo)); ASSERT_TRUE(resolver.updateConstants(jobInfo)); ASSERT_TRUE(resolver.executeSync(jobInfo)); + + //////////////////////////////////////////////////////////////////////////// resolver.submit(&stratum); //////////////////////////////////////////////////////////////////////////// diff --git a/sources/resolver/nvidia/tests/meowpow.cpp b/sources/resolver/nvidia/tests/meowpow.cpp index df296c8..56630b9 100644 --- a/sources/resolver/nvidia/tests/meowpow.cpp +++ b/sources/resolver/nvidia/tests/meowpow.cpp @@ -30,7 +30,7 @@ struct ResolverMeowpowNvidiaTest : public testing::Test ~ResolverMeowpowNvidiaTest() { - resolver::tests::cleanUpCuda(); + resolver::tests::cleanUpCuda(properties); } void initializeJob(uint64_t const nonce) @@ -41,7 +41,7 @@ struct ResolverMeowpowNvidiaTest : public testing::Test jobInfo.seedHash = algo::toHash256("cfa3e37c459ebd9b4138bd2141a52d89f6f8f671ecf91456f5a29176eb132fc0"); jobInfo.boundary = algo::toHash256("0000000500000000000000000000000000000000000000000000000000000000"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::progpow::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::progpow::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::meowpow::MAX_PERIOD; } }; diff --git a/sources/resolver/nvidia/tests/progpow_quai.cpp b/sources/resolver/nvidia/tests/progpow_quai.cpp index 76a97a2..45323cb 100644 --- a/sources/resolver/nvidia/tests/progpow_quai.cpp +++ b/sources/resolver/nvidia/tests/progpow_quai.cpp @@ -30,7 +30,7 @@ struct ResolverProgpowQuaiNvidiaTest : public testing::Test ~ResolverProgpowQuaiNvidiaTest() { - resolver::tests::cleanUpCuda(); + resolver::tests::cleanUpCuda(properties); } void initializeJob(uint64_t const nonce) @@ -41,7 +41,7 @@ struct ResolverProgpowQuaiNvidiaTest : public testing::Test jobInfo.seedHash = algo::toHash256("0000000000000000000000000000000000000000000000000000000000000000"); jobInfo.boundary = algo::toHash256("0000000225c17d04dad2965cc5a02a23e254c0c3f75d9178046aeb27ce1ca574"); jobInfo.boundaryU64 = algo::toUINT64(jobInfo.boundary); - jobInfo.epoch = algo::ethash::findEpoch(jobInfo.seedHash, algo::progpow_quai::EPOCH_LENGTH); + jobInfo.epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, algo::progpow_quai::EPOCH_LENGTH); jobInfo.period = jobInfo.blockNumber / algo::progpow_quai::MAX_PERIOD; } }; diff --git a/sources/resolver/nvidia/tests/progpow_z.cpp b/sources/resolver/nvidia/tests/progpow_z.cpp index bebbbd4..cf2db6c 100644 --- a/sources/resolver/nvidia/tests/progpow_z.cpp +++ b/sources/resolver/nvidia/tests/progpow_z.cpp @@ -30,7 +30,7 @@ struct ResolverProgpowZNvidiaTest : public testing::Test ~ResolverProgpowZNvidiaTest() { - resolver::tests::cleanUpCuda(); + resolver::tests::cleanUpCuda(properties); } void initializeJob(uint64_t const nonce) diff --git a/sources/resolver/resolver.hpp b/sources/resolver/resolver.hpp index 812ca61..dd92e17 100644 --- a/sources/resolver/resolver.hpp +++ b/sources/resolver/resolver.hpp @@ -1,5 +1,6 @@ #pragma once +#include #include #include #include @@ -10,10 +11,11 @@ namespace resolver class Resolver { public: - uint32_t currentIndexStream{ 0u }; - uint32_t deviceId{ 0u }; - std::string jobId{}; - uint64_t deviceMemoryAvailable{ 0ull }; + algo::ALGORITHM algorithm{ algo::ALGORITHM::UNKNOWN }; + uint32_t currentIndexStream{ 0u }; + uint32_t deviceId{ 0u }; + std::string jobId{}; + uint64_t deviceMemoryAvailable{ 0ull }; Resolver() = default; virtual ~Resolver() = default; diff --git a/sources/resolver/tests/amd.hpp b/sources/resolver/tests/amd.hpp index 75793e4..da926c5 100644 --- a/sources/resolver/tests/amd.hpp +++ b/sources/resolver/tests/amd.hpp @@ -17,7 +17,8 @@ namespace resolver cl::CommandQueue clQueue{ nullptr }; }; - inline uint32_t getDeviceCount() + inline + uint32_t getDeviceCount() { uint32_t count { 0u }; std::vector cldevices{}; @@ -39,7 +40,8 @@ namespace resolver return count; } - inline cl::Device getDevice(uint32_t const index) + inline + cl::Device getDevice(uint32_t const index) { std::vector cldevices{}; std::vector platforms{}; diff --git a/sources/resolver/tests/nvidia.hpp b/sources/resolver/tests/nvidia.hpp index 8911561..a20057f 100644 --- a/sources/resolver/tests/nvidia.hpp +++ b/sources/resolver/tests/nvidia.hpp @@ -9,23 +9,28 @@ namespace resolver { struct Properties { - uint32_t cuIndex { 0u }; + uint32_t cuIndex{ 0u }; CUdevice cuDevice; CUcontext cuContext{ nullptr }; cudaStream_t cuStream{ nullptr }; cudaDeviceProp cuProperties{}; }; - inline bool cleanUpCuda() + inline + bool cleanUpCuda(resolver::tests::Properties& properties) { - CUDA_ER(cudaDeviceReset()); + if (nullptr != properties.cuContext) + { + CUDA_ER(cudaDeviceReset()); + } return true; } - inline bool initializeCuda(resolver::tests::Properties& properties, + inline + bool initializeCuda(resolver::tests::Properties& properties, uint32_t const index = 0u) { - if (false == cleanUpCuda()) + if (false == cleanUpCuda(properties)) { return false; } diff --git a/sources/stratum/ethash.cpp b/sources/stratum/ethash.cpp index 57319ae..81f0419 100644 --- a/sources/stratum/ethash.cpp +++ b/sources/stratum/ethash.cpp @@ -87,7 +87,7 @@ void stratum::StratumEthash::onMiningNotify( jobInfo.jobID = algo::toHash256(jobID); //////////////////////////////////////////////////////////////////////////// - int32_t const epoch{ algo::ethash::findEpoch(jobInfo.seedHash, maxEpochNumber) }; + int32_t const epoch{ algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, maxEpochNumber) }; if (-1 != epoch) { jobInfo.epoch = epoch; diff --git a/sources/stratum/progpow.cpp b/sources/stratum/progpow.cpp index de71e67..67139e7 100644 --- a/sources/stratum/progpow.cpp +++ b/sources/stratum/progpow.cpp @@ -191,7 +191,6 @@ void stratum::StratumProgPOW::onResponseEthProxy(boost::json::object const& root { case stratum::ETHPROXY_ID::SUBMITLOGIN: { - __TRACE(); //////////////////////////////////////////////////////////////////// if ( false == root.contains("error") || true == root.at("error").is_null()) @@ -340,7 +339,7 @@ void stratum::StratumProgPOW::onMiningNotify( } else { - epoch = algo::ethash::findEpoch(jobInfo.seedHash, maxEthashEpoch); + epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, maxEthashEpoch); } if (-1 != epoch) { @@ -372,20 +371,14 @@ void stratum::StratumProgPOW::onMiningNotify( } case stratum::STRATUM_TYPE::ETHPROXY: { - __TRACE(); //////////////////////////////////////////////////////////////////// auto const& params{ root.at("result").as_array() }; //////////////////////////////////////////////////////////////////// - __TRACE(); std::string const jobID{ common::boostGetString(params, 0) }; - __TRACE(); jobInfo.headerHash = algo::toHash256(common::boostGetString(params, 0)); - __TRACE(); jobInfo.seedHash = algo::toHash256(common::boostGetString(params, 1)); - __TRACE(); jobInfo.boundary = algo::toHash256(common::boostGetString(params, 2)); - __TRACE(); std::string const blockNumber{ common::boostGetString(params, 3) }; logInfo() << "blockNumber: " << blockNumber; @@ -413,7 +406,7 @@ void stratum::StratumProgPOW::onMiningNotify( } else { - epoch = algo::ethash::findEpoch(jobInfo.seedHash, maxEthashEpoch); + epoch = algo::ethash::ContextGenerator::instance().findEpoch(jobInfo.seedHash, maxEthashEpoch); } if (-1 != epoch) {