From 0bff6a90f8bb49052c5c63833d420824055bbdae Mon Sep 17 00:00:00 2001 From: luminousminer Date: Wed, 11 Feb 2026 18:27:49 +0100 Subject: [PATCH 01/15] 1/x: lightcache compute with 1 thread and shared --- sources/algo/dag_context.hpp | 1 - sources/algo/ethash/ethash.cpp | 90 +++++++++++-------- sources/algo/ethash/ethash.hpp | 3 +- sources/algo/ethash/tests/ethash.cpp | 2 +- sources/benchmark/amd/kawpow.cpp | 2 +- sources/resolver/amd/etchash.cpp | 5 +- sources/resolver/amd/ethash.cpp | 2 +- sources/resolver/amd/progpow.cpp | 2 +- sources/resolver/cpu/progpow.cpp | 2 +- sources/resolver/mocker.cpp | 5 +- sources/resolver/nvidia/etchash.cpp | 5 +- sources/resolver/nvidia/ethash.cpp | 5 +- sources/resolver/nvidia/progpow.cpp | 5 +- .../resolver/nvidia/tests/autolykos_v2.cpp | 2 +- sources/resolver/nvidia/tests/blake3.cpp | 2 +- sources/resolver/nvidia/tests/ethash.cpp | 2 +- sources/resolver/nvidia/tests/firopow.cpp | 2 +- sources/resolver/nvidia/tests/kawpow.cpp | 2 +- sources/resolver/nvidia/tests/meowpow.cpp | 2 +- .../resolver/nvidia/tests/progpow_quai.cpp | 2 +- sources/resolver/nvidia/tests/progpow_z.cpp | 2 +- sources/resolver/tests/amd.hpp | 6 +- sources/resolver/tests/nvidia.hpp | 15 ++-- 23 files changed, 100 insertions(+), 66 deletions(-) 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/ethash.cpp b/sources/algo/ethash/ethash.cpp index fa7b5df..b55148a 100644 --- a/sources/algo/ethash/ethash.cpp +++ b/sources/algo/ethash/ethash.cpp @@ -10,6 +10,10 @@ #include +boost::mutex mtxDagContext{}; +algo::DagContext localDagContext{}; + + int32_t algo::ethash::findEpoch( algo::hash256 const& seedHash, uint32_t const maxEpoch) @@ -59,8 +63,8 @@ int32_t algo::ethash::findEpoch( void algo::ethash::freeDagContext(algo::DagContext& context) { //////////////////////////////////////////////////////////////////////////// - SAFE_DELETE_ARRAY(context.data); - context.lightCache.hash = nullptr; + SAFE_DELETE_ARRAY(localDagContext.data); + localDagContext.lightCache.hash = nullptr; } @@ -74,11 +78,14 @@ void algo::ethash::initializeDagContext( uint32_t const lightCacheCountItemsInit) { //////////////////////////////////////////////////////////////////////////// - context.epoch = castU32(currentEpoch); - if ( castU32(context.epoch) > maxEpoch + UNIQUE_LOCK(mtxDagContext); + + //////////////////////////////////////////////////////////////////////////// + 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; } @@ -97,80 +104,87 @@ void algo::ethash::initializeDagContext( 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); + + //////////////////////////////////////////////////////////////////////////// + algo::hash512 seedHash{ algo::keccak(seed) }; + algo::copyHash(localDagContext.hashedSeedCache, seedHash); + + //////////////////////////////////////////////////////////////////////////// + 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; + algo::copyHash(context.hashedSeedCache, localDagContext.hashedSeedCache); } void algo::ethash::buildLightCache( - algo::DagContext& context, - bool const buildOnCPU) + algo::DagContext& context) { //////////////////////////////////////////////////////////////////////////// - algo::hash512 item - { - algo::keccak(context.originalSeedCache) - }; + UNIQUE_LOCK(mtxDagContext); //////////////////////////////////////////////////////////////////////////// - if (false == buildOnCPU) - { - algo::copyHash(context.hashedSeedCache, item); - return; - } - - //////////////////////////////////////////////////////////////////////////// - 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) + 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 }; //////////////////////////////////////////////////////////////////////////// - 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); } } + + //////////////////////////////////////////////////////////////////////////// + context.data = localDagContext.data; + for (uint64_t i{ 0ull }; i < localDagContext.lightCache.numberItem; ++i) + { + algo::copyHash(context.lightCache.hash[i], localDagContext.lightCache.hash[i]); + } } diff --git a/sources/algo/ethash/ethash.hpp b/sources/algo/ethash/ethash.hpp index 4156814..1ff1858 100644 --- a/sources/algo/ethash/ethash.hpp +++ b/sources/algo/ethash/ethash.hpp @@ -42,8 +42,7 @@ namespace algo uint32_t const lightCacheCountItemsGrowth, uint32_t const lightCacheCountItemsInit); void freeDagContext(algo::DagContext& context); - void buildLightCache(algo::DagContext& context, - bool const buildOnCPU); + void buildLightCache(algo::DagContext& context); #endif } } diff --git a/sources/algo/ethash/tests/ethash.cpp b/sources/algo/ethash/tests/ethash.cpp index 2a32b79..a41eb2e 100644 --- a/sources/algo/ethash/tests/ethash.cpp +++ b/sources/algo/ethash/tests/ethash.cpp @@ -65,7 +65,7 @@ TEST_F(EthashTest, lightCacheBuild) lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - algo::ethash::buildLightCache(context, true); + algo::ethash::buildLightCache(context); // light cache size ASSERT_EQ(1411061ull, context.lightCache.numberItem); diff --git a/sources/benchmark/amd/kawpow.cpp b/sources/benchmark/amd/kawpow.cpp index 8b4a408..0ebc04e 100644 --- a/sources/benchmark/amd/kawpow.cpp +++ b/sources/benchmark/amd/kawpow.cpp @@ -60,7 +60,7 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow() algo::ethash::LIGHT_CACHE_COUNT_ITEMS_GROWTH, algo::ethash::LIGHT_CACHE_COUNT_ITEMS_INIT ); - algo::ethash::buildLightCache(dagContext, true); + algo::ethash::buildLightCache(dagContext); //////////////////////////////////////////////////////////////////////////// dagCache.setSize(dagContext.dagCache.size); diff --git a/sources/resolver/amd/etchash.cpp b/sources/resolver/amd/etchash.cpp index 8ea9fcd..1c8e43b 100644 --- a/sources/resolver/amd/etchash.cpp +++ b/sources/resolver/amd/etchash.cpp @@ -28,7 +28,10 @@ bool resolver::ResolverAmdEtchash::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); + if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + { + algo::ethash::buildLightCache(context); + } if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull diff --git a/sources/resolver/amd/ethash.cpp b/sources/resolver/amd/ethash.cpp index 7f82ea0..62eacdb 100644 --- a/sources/resolver/amd/ethash.cpp +++ b/sources/resolver/amd/ethash.cpp @@ -39,7 +39,7 @@ bool resolver::ResolverAmdEthash::updateContext( lightCacheCountItemsInit ); // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context, true); + algo::ethash::buildLightCache(context); if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull diff --git a/sources/resolver/amd/progpow.cpp b/sources/resolver/amd/progpow.cpp index b89691a..2631230 100644 --- a/sources/resolver/amd/progpow.cpp +++ b/sources/resolver/amd/progpow.cpp @@ -35,7 +35,7 @@ bool resolver::ResolverAmdProgPOW::updateContext( lightCacheCountItemsInit ); // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context, true); + algo::ethash::buildLightCache(context); if ( 0ull == context.lightCache.numberItem || 0ull == context.lightCache.size diff --git a/sources/resolver/cpu/progpow.cpp b/sources/resolver/cpu/progpow.cpp index f1bab19..11e353e 100644 --- a/sources/resolver/cpu/progpow.cpp +++ b/sources/resolver/cpu/progpow.cpp @@ -28,7 +28,7 @@ bool resolver::ResolverCpuProgPOW::updateContext(stratum::StratumJobInfo const& lightCacheCountItemsInit ); // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context, true); + algo::ethash::buildLightCache(context); if ( 0ull == context.lightCache.numberItem || 0ull == context.lightCache.size diff --git a/sources/resolver/mocker.cpp b/sources/resolver/mocker.cpp index f0e700e..04847fd 100644 --- a/sources/resolver/mocker.cpp +++ b/sources/resolver/mocker.cpp @@ -50,7 +50,10 @@ bool resolver::ResolverMocker::updateMemory( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); + if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + { + algo::ethash::buildLightCache(context); + } /////////////////////////////////////////////////////////////////////////// boost::this_thread::sleep_for(WAIT_UPDATE_MEMORY); diff --git a/sources/resolver/nvidia/etchash.cpp b/sources/resolver/nvidia/etchash.cpp index 865632c..9b8e16a 100644 --- a/sources/resolver/nvidia/etchash.cpp +++ b/sources/resolver/nvidia/etchash.cpp @@ -26,7 +26,10 @@ bool resolver::ResolverNvidiaEtchash::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); + if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + { + algo::ethash::buildLightCache(context); + } /////////////////////////////////////////////////////////////////////////// if ( context.lightCache.numberItem == 0ull diff --git a/sources/resolver/nvidia/ethash.cpp b/sources/resolver/nvidia/ethash.cpp index e9bd014..057210e 100644 --- a/sources/resolver/nvidia/ethash.cpp +++ b/sources/resolver/nvidia/ethash.cpp @@ -31,7 +31,10 @@ bool resolver::ResolverNvidiaEthash::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); + if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + { + algo::ethash::buildLightCache(context); + } if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index 460cfc4..9f417b6 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -33,7 +33,10 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - algo::ethash::buildLightCache(context, config.deviceAlgorithm.ethashBuildLightCacheCPU); + if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + { + algo::ethash::buildLightCache(context); + } if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull 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..72ddd59 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) diff --git a/sources/resolver/nvidia/tests/firopow.cpp b/sources/resolver/nvidia/tests/firopow.cpp index 046d8a2..9545022 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) diff --git a/sources/resolver/nvidia/tests/kawpow.cpp b/sources/resolver/nvidia/tests/kawpow.cpp index bceedab..acf6d82 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( diff --git a/sources/resolver/nvidia/tests/meowpow.cpp b/sources/resolver/nvidia/tests/meowpow.cpp index df296c8..4148306 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) diff --git a/sources/resolver/nvidia/tests/progpow_quai.cpp b/sources/resolver/nvidia/tests/progpow_quai.cpp index 76a97a2..05a4e9f 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) 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/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; } From 6933b157a2e10878cd416f34ab272bcec5b0e0e6 Mon Sep 17 00:00:00 2001 From: luminousminer Date: Wed, 11 Feb 2026 18:29:38 +0100 Subject: [PATCH 02/15] sharing ptr instead of copy data --- sources/algo/ethash/ethash.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/sources/algo/ethash/ethash.cpp b/sources/algo/ethash/ethash.cpp index b55148a..85657df 100644 --- a/sources/algo/ethash/ethash.cpp +++ b/sources/algo/ethash/ethash.cpp @@ -183,8 +183,5 @@ void algo::ethash::buildLightCache( //////////////////////////////////////////////////////////////////////////// context.data = localDagContext.data; - for (uint64_t i{ 0ull }; i < localDagContext.lightCache.numberItem; ++i) - { - algo::copyHash(context.lightCache.hash[i], localDagContext.lightCache.hash[i]); - } + context.lightCache.hash = localDagContext.lightCache.hash; } From 373f43a9d87ad701ec38e115bebd4c9093a5e85b Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 11:33:59 +0100 Subject: [PATCH 03/15] add parameter --tool_mocker_resolver_count --- CMakeLists.txt | 3 +++ PARAMETERS.md | 7 +++++++ sources/algo/ethash/ethash.cpp | 15 +++++++++++++-- sources/algo/ethash/ethash.hpp | 2 +- sources/common/cli/cli.hpp | 4 ++++ sources/common/cli/cli_tools.cpp | 17 +++++++++++++++++ sources/common/config.cpp | 7 +++++++ sources/common/config.hpp | 12 ++++++++++++ sources/device/device_manager.cpp | 8 ++++++-- sources/resolver/mocker.cpp | 26 +++++++++++++++----------- sources/resolver/nvidia/progpow.cpp | 26 +++++++++++++++----------- 11 files changed, 100 insertions(+), 27 deletions(-) create mode 100644 sources/common/cli/cli_tools.cpp 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/ethash/ethash.cpp b/sources/algo/ethash/ethash.cpp index 85657df..28012ce 100644 --- a/sources/algo/ethash/ethash.cpp +++ b/sources/algo/ethash/ethash.cpp @@ -68,7 +68,7 @@ void algo::ethash::freeDagContext(algo::DagContext& context) } -void algo::ethash::initializeDagContext( +bool algo::ethash::initializeDagContext( algo::DagContext& context, uint64_t const currentEpoch, uint32_t const maxEpoch, @@ -80,13 +80,20 @@ void algo::ethash::initializeDagContext( //////////////////////////////////////////////////////////////////////////// UNIQUE_LOCK(mtxDagContext); + //////////////////////////////////////////////////////////////////////////// + if (localDagContext.epoch == castU32(currentEpoch)) + { + logInfo() << "Skip epoch: " << currentEpoch; + return false; + } + //////////////////////////////////////////////////////////////////////////// localDagContext.epoch = castU32(currentEpoch); if ( castU32(localDagContext.epoch) > maxEpoch && algo::ethash::EIP1057_MAX_EPOCH_NUMER != maxEpoch) { logErr() << "context.epoch: " << localDagContext.epoch << " | maxEpoch: " << maxEpoch; - return; + return false; } //////////////////////////////////////////////////////////////////////////// @@ -132,6 +139,10 @@ void algo::ethash::initializeDagContext( context.dagCache.numberItem = localDagContext.dagCache.numberItem; context.dagCache.size = localDagContext.dagCache.size; algo::copyHash(context.hashedSeedCache, localDagContext.hashedSeedCache); + + //////////////////////////////////////////////////////////////////////////// + logInfo() << "Need build"; + return true; } diff --git a/sources/algo/ethash/ethash.hpp b/sources/algo/ethash/ethash.hpp index 1ff1858..568fb46 100644 --- a/sources/algo/ethash/ethash.hpp +++ b/sources/algo/ethash/ethash.hpp @@ -34,7 +34,7 @@ namespace algo #if !defined(__LIB_CUDA) int32_t findEpoch(algo::hash256 const& seedHash, uint32_t const maxEpoch); - void initializeDagContext(algo::DagContext& context, + bool initializeDagContext(algo::DagContext& context, uint64_t const currentEpoch, uint32_t const maxEpoch, uint64_t const dagCountItemsGrowth, diff --git a/sources/common/cli/cli.hpp b/sources/common/cli/cli.hpp index 5c2adda..ed1e5f8 100644 --- a/sources/common/cli/cli.hpp +++ b/sources/common/cli/cli.hpp @@ -90,6 +90,10 @@ namespace common #endif bool isCpuEnable() const; +#if defined(TOOLS_ENABLE) && defined(TOOL_MOCKER) + uint32_t getMockerResolverCount() const; +#endif + // Device settings custom std::vector getDevicesDisable() const; customTupleStr getCustomHost() const; diff --git a/sources/common/cli/cli_tools.cpp b/sources/common/cli/cli_tools.cpp new file mode 100644 index 0000000..3268fb9 --- /dev/null +++ b/sources/common/cli/cli_tools.cpp @@ -0,0 +1,17 @@ +#if defined(TOOLS_ENABLE) && defined(TOOL_MOCKER) + +#include + + +uint32_t common::Cli::getMockerResolverCount() const +{ + if (true == contains("tool_mocker_resolver_count")) + { + return params["tool_mocker_resolver_count"].as(); + } + + return 8u; +} + + +#endif // TOOLS_ENABLE && TOOL_MOCKER diff --git a/sources/common/config.cpp b/sources/common/config.cpp index 6e7001b..e71e4c5 100644 --- a/sources/common/config.cpp +++ b/sources/common/config.cpp @@ -481,6 +481,13 @@ bool common::Config::loadCli(int argc, char** argv) { api.port = ApiPort; } + + //////////////////////////////////////////////////////////////////////// + // TOOL MOCKER + //////////////////////////////////////////////////////////////////////// +#if defined(TOOLS_ENABLE) && defined(TOOL_MOCKER) + toolConfigs.mockerResolverCount = cli.getMockerResolverCount(); +#endif } catch(std::exception const& e) { diff --git a/sources/common/config.hpp b/sources/common/config.hpp index e70f2cc..461cba3 100644 --- a/sources/common/config.hpp +++ b/sources/common/config.hpp @@ -83,6 +83,15 @@ namespace common uint32_t port{ 8080u }; }; +#if defined(TOOLS_ENABLE) + struct ToolConfig + { +#if defined(TOOL_MOCKER) + std::optional mockerResolverCount{ 8u }; +#endif + }; +#endif + common::PROFILE profile{ common::PROFILE::STANDARD }; common::Cli cli{}; LogConfig log{}; @@ -97,6 +106,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/device/device_manager.cpp b/sources/device/device_manager.cpp index 10df18d..94af2d7 100644 --- a/sources/device/device_manager.cpp +++ b/sources/device/device_manager.cpp @@ -285,14 +285,18 @@ 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() }; + + /////////////////////////////////////////////////////////////////////////// + for (uint32_t i{ 0u }; i < config.toolConfigs.mockerResolverCount; ++i) { device::DeviceMocker* device{ NEW(device::DeviceMocker) }; if (nullptr == device) { return false; } - device->id = 1000 + i; + device->id = 1000u + i; device->deviceType = device::DEVICE_TYPE::MOCKER; devices.push_back(device); } diff --git a/sources/resolver/mocker.cpp b/sources/resolver/mocker.cpp index 04847fd..d4ac5b2 100644 --- a/sources/resolver/mocker.cpp +++ b/sources/resolver/mocker.cpp @@ -40,17 +40,21 @@ bool resolver::ResolverMocker::updateMemory( 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 - ( - context, - jobInfo.epoch, - maxEpoch, - dagCountItemsGrowth, - dagCountItemsInit, - lightCacheCountItemsGrowth, - lightCacheCountItemsInit - ); - if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + bool const needBuildLightCache + { + algo::ethash::initializeDagContext + ( + context, + jobInfo.epoch, + maxEpoch, + dagCountItemsGrowth, + dagCountItemsInit, + lightCacheCountItemsGrowth, + lightCacheCountItemsInit + ) + }; + if ( true == needBuildLightCache + && true == config.deviceAlgorithm.ethashBuildLightCacheCPU) { algo::ethash::buildLightCache(context); } diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index 9f417b6..e5da5ed 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -23,17 +23,21 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( common::Config& config{ common::Config::instance() }; //////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext - ( - context, - jobInfo.epoch, - maxEpoch, - dagCountItemsGrowth, - dagCountItemsInit, - lightCacheCountItemsGrowth, - lightCacheCountItemsInit - ); - if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + bool const updated + { + algo::ethash::initializeDagContext + ( + context, + jobInfo.epoch, + maxEpoch, + dagCountItemsGrowth, + dagCountItemsInit, + lightCacheCountItemsGrowth, + lightCacheCountItemsInit + ) + }; + if ( true == updated + && true == config.deviceAlgorithm.ethashBuildLightCacheCPU) { algo::ethash::buildLightCache(context); } From 026485ea9899037221172597c21baa1533ae58ff Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 16:44:27 +0100 Subject: [PATCH 04/15] 2/x: lightcache compute with 1 thread and shared --- sources/algo/ethash/ethash.cpp | 41 +++++++++++++++++----------- sources/algo/ethash/ethash.hpp | 17 ++++++------ sources/algo/ethash/tests/ethash.cpp | 3 +- sources/benchmark/amd/kawpow.cpp | 3 +- sources/common/log/log.hpp | 9 ++---- sources/resolver/amd/etchash.cpp | 6 +--- sources/resolver/amd/ethash.cpp | 4 +-- sources/resolver/amd/progpow.cpp | 4 +-- sources/resolver/cpu/progpow.cpp | 4 +-- sources/resolver/mocker.cpp | 33 ++++++++-------------- sources/resolver/nvidia/etchash.cpp | 6 +--- sources/resolver/nvidia/ethash.cpp | 9 +----- sources/resolver/nvidia/progpow.cpp | 30 ++++++++------------ 13 files changed, 69 insertions(+), 100 deletions(-) diff --git a/sources/algo/ethash/ethash.cpp b/sources/algo/ethash/ethash.cpp index 28012ce..229b852 100644 --- a/sources/algo/ethash/ethash.cpp +++ b/sources/algo/ethash/ethash.cpp @@ -68,7 +68,7 @@ void algo::ethash::freeDagContext(algo::DagContext& context) } -bool algo::ethash::initializeDagContext( +void algo::ethash::buildContext( algo::DagContext& context, uint64_t const currentEpoch, uint32_t const maxEpoch, @@ -80,11 +80,15 @@ bool algo::ethash::initializeDagContext( //////////////////////////////////////////////////////////////////////////// UNIQUE_LOCK(mtxDagContext); + /////////////////////////////////////////////////////////////////////////// + common::Config& config{ common::Config::instance() }; + //////////////////////////////////////////////////////////////////////////// if (localDagContext.epoch == castU32(currentEpoch)) { logInfo() << "Skip epoch: " << currentEpoch; - return false; + copyContext(context); + return; } //////////////////////////////////////////////////////////////////////////// @@ -93,7 +97,7 @@ bool algo::ethash::initializeDagContext( && algo::ethash::EIP1057_MAX_EPOCH_NUMER != maxEpoch) { logErr() << "context.epoch: " << localDagContext.epoch << " | maxEpoch: " << maxEpoch; - return false; + return; } //////////////////////////////////////////////////////////////////////////// @@ -133,25 +137,19 @@ bool algo::ethash::initializeDagContext( algo::copyHash(localDagContext.hashedSeedCache, seedHash); //////////////////////////////////////////////////////////////////////////// - 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; - algo::copyHash(context.hashedSeedCache, localDagContext.hashedSeedCache); + if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + { + algo::ethash::buildLightCacheOnCPU(context); + } //////////////////////////////////////////////////////////////////////////// - logInfo() << "Need build"; - return true; + copyContext(context); } -void algo::ethash::buildLightCache( +void algo::ethash::buildLightCacheOnCPU( algo::DagContext& context) { - //////////////////////////////////////////////////////////////////////////// - UNIQUE_LOCK(mtxDagContext); - //////////////////////////////////////////////////////////////////////////// algo::hash512 item{ localDagContext.hashedSeedCache }; size_t const dataLength{ algo::LEN_HASH_512 + localDagContext.lightCache.size }; @@ -191,8 +189,19 @@ void algo::ethash::buildLightCache( localDagContext.lightCache.hash[i] = algo::keccak(xored); } } +} - //////////////////////////////////////////////////////////////////////////// + +void algo::ethash::copyContext(algo::DagContext& 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 568fb46..ae5f7ee 100644 --- a/sources/algo/ethash/ethash.hpp +++ b/sources/algo/ethash/ethash.hpp @@ -34,15 +34,16 @@ namespace algo #if !defined(__LIB_CUDA) int32_t findEpoch(algo::hash256 const& seedHash, uint32_t const maxEpoch); - bool 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 buildContext(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 buildLightCacheOnCPU(algo::DagContext& context); void freeDagContext(algo::DagContext& context); - void buildLightCache(algo::DagContext& context); + void copyContext(algo::DagContext& context); #endif } } diff --git a/sources/algo/ethash/tests/ethash.cpp b/sources/algo/ethash/tests/ethash.cpp index a41eb2e..6bc8980 100644 --- a/sources/algo/ethash/tests/ethash.cpp +++ b/sources/algo/ethash/tests/ethash.cpp @@ -55,7 +55,7 @@ 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::buildContext ( context, 561ull, @@ -65,7 +65,6 @@ TEST_F(EthashTest, lightCacheBuild) lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - algo::ethash::buildLightCache(context); // light cache size ASSERT_EQ(1411061ull, context.lightCache.numberItem); diff --git a/sources/benchmark/amd/kawpow.cpp b/sources/benchmark/amd/kawpow.cpp index 0ebc04e..7b89319 100644 --- a/sources/benchmark/amd/kawpow.cpp +++ b/sources/benchmark/amd/kawpow.cpp @@ -50,7 +50,7 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow() //////////////////////////////////////////////////////////////////////////// algo::DagContext dagContext{}; - algo::ethash::initializeDagContext + algo::ethash::buildContext ( dagContext, epoch, @@ -60,7 +60,6 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow() algo::ethash::LIGHT_CACHE_COUNT_ITEMS_GROWTH, algo::ethash::LIGHT_CACHE_COUNT_ITEMS_INIT ); - algo::ethash::buildLightCache(dagContext); //////////////////////////////////////////////////////////////////////////// dagCache.setSize(dagContext.dagCache.size); 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/resolver/amd/etchash.cpp b/sources/resolver/amd/etchash.cpp index 1c8e43b..7f41dc1 100644 --- a/sources/resolver/amd/etchash.cpp +++ b/sources/resolver/amd/etchash.cpp @@ -18,7 +18,7 @@ bool resolver::ResolverAmdEtchash::updateContext( common::Config& config{ common::Config::instance() }; /////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::buildContext ( context, jobInfo.epoch, @@ -28,10 +28,6 @@ bool resolver::ResolverAmdEtchash::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) - { - algo::ethash::buildLightCache(context); - } if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull diff --git a/sources/resolver/amd/ethash.cpp b/sources/resolver/amd/ethash.cpp index 62eacdb..6261dd4 100644 --- a/sources/resolver/amd/ethash.cpp +++ b/sources/resolver/amd/ethash.cpp @@ -28,7 +28,7 @@ bool resolver::ResolverAmdEthash::updateContext( common::Config& config{ common::Config::instance() }; /////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::buildContext ( context, jobInfo.epoch, @@ -38,8 +38,6 @@ bool resolver::ResolverAmdEthash::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context); if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull diff --git a/sources/resolver/amd/progpow.cpp b/sources/resolver/amd/progpow.cpp index 2631230..f089912 100644 --- a/sources/resolver/amd/progpow.cpp +++ b/sources/resolver/amd/progpow.cpp @@ -24,7 +24,7 @@ bool resolver::ResolverAmdProgPOW::updateContext( common::Config& config{ common::Config::instance() }; /////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::buildContext ( context, jobInfo.epoch, @@ -34,8 +34,6 @@ bool resolver::ResolverAmdProgPOW::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context); if ( 0ull == context.lightCache.numberItem || 0ull == context.lightCache.size diff --git a/sources/resolver/cpu/progpow.cpp b/sources/resolver/cpu/progpow.cpp index 11e353e..b89ccf5 100644 --- a/sources/resolver/cpu/progpow.cpp +++ b/sources/resolver/cpu/progpow.cpp @@ -17,7 +17,7 @@ resolver::ResolverCpuProgPOW::~ResolverCpuProgPOW() bool resolver::ResolverCpuProgPOW::updateContext(stratum::StratumJobInfo const& jobInfo) { //////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::buildContext ( context, jobInfo.epoch, @@ -27,8 +27,6 @@ bool resolver::ResolverCpuProgPOW::updateContext(stratum::StratumJobInfo const& lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - // TODO: config.deviceAlgorithm.ethashBuildLightCacheCPU - algo::ethash::buildLightCache(context); if ( 0ull == context.lightCache.numberItem || 0ull == context.lightCache.size diff --git a/sources/resolver/mocker.cpp b/sources/resolver/mocker.cpp index d4ac5b2..c070511 100644 --- a/sources/resolver/mocker.cpp +++ b/sources/resolver/mocker.cpp @@ -27,9 +27,6 @@ void resolver::ResolverMocker::overrideOccupancy( bool resolver::ResolverMocker::updateMemory( [[maybe_unused]] stratum::StratumJobInfo const& jobInfo) { - /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; - /////////////////////////////////////////////////////////////////////////// algo::DagContext context{}; uint32_t maxEpoch{ algo::ethash::MAX_EPOCH_NUMBER }; @@ -40,24 +37,18 @@ bool resolver::ResolverMocker::updateMemory( 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 }; - bool const needBuildLightCache - { - algo::ethash::initializeDagContext - ( - context, - jobInfo.epoch, - maxEpoch, - dagCountItemsGrowth, - dagCountItemsInit, - lightCacheCountItemsGrowth, - lightCacheCountItemsInit - ) - }; - if ( true == needBuildLightCache - && true == config.deviceAlgorithm.ethashBuildLightCacheCPU) - { - algo::ethash::buildLightCache(context); - } + __TRACE_RESOLVER(); + algo::ethash::buildContext + ( + context, + jobInfo.epoch, + maxEpoch, + dagCountItemsGrowth, + dagCountItemsInit, + lightCacheCountItemsGrowth, + lightCacheCountItemsInit + ); + __TRACE_RESOLVER(); /////////////////////////////////////////////////////////////////////////// boost::this_thread::sleep_for(WAIT_UPDATE_MEMORY); diff --git a/sources/resolver/nvidia/etchash.cpp b/sources/resolver/nvidia/etchash.cpp index 9b8e16a..e6b7fa1 100644 --- a/sources/resolver/nvidia/etchash.cpp +++ b/sources/resolver/nvidia/etchash.cpp @@ -16,7 +16,7 @@ bool resolver::ResolverNvidiaEtchash::updateContext( common::Config& config{ common::Config::instance() }; /////////////////////////////////////////////////////////////////////////// - algo::ethash::initializeDagContext + algo::ethash::buildContext ( context, jobInfo.epoch, @@ -26,10 +26,6 @@ bool resolver::ResolverNvidiaEtchash::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) - { - algo::ethash::buildLightCache(context); - } /////////////////////////////////////////////////////////////////////////// if ( context.lightCache.numberItem == 0ull diff --git a/sources/resolver/nvidia/ethash.cpp b/sources/resolver/nvidia/ethash.cpp index 057210e..c839c73 100644 --- a/sources/resolver/nvidia/ethash.cpp +++ b/sources/resolver/nvidia/ethash.cpp @@ -18,10 +18,7 @@ resolver::ResolverNvidiaEthash::~ResolverNvidiaEthash() bool resolver::ResolverNvidiaEthash::updateContext( stratum::StratumJobInfo const& jobInfo) { - /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; - - algo::ethash::initializeDagContext + algo::ethash::buildContext ( context, jobInfo.epoch, @@ -31,10 +28,6 @@ bool resolver::ResolverNvidiaEthash::updateContext( lightCacheCountItemsGrowth, lightCacheCountItemsInit ); - if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) - { - algo::ethash::buildLightCache(context); - } if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index e5da5ed..66a8a17 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -23,24 +23,18 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( common::Config& config{ common::Config::instance() }; //////////////////////////////////////////////////////////////////////////// - bool const updated - { - algo::ethash::initializeDagContext - ( - context, - jobInfo.epoch, - maxEpoch, - dagCountItemsGrowth, - dagCountItemsInit, - lightCacheCountItemsGrowth, - lightCacheCountItemsInit - ) - }; - if ( true == updated - && true == config.deviceAlgorithm.ethashBuildLightCacheCPU) - { - algo::ethash::buildLightCache(context); - } + __TRACE_RESOLVER(); + algo::ethash::buildContext + ( + context, + jobInfo.epoch, + maxEpoch, + dagCountItemsGrowth, + dagCountItemsInit, + lightCacheCountItemsGrowth, + lightCacheCountItemsInit + ); + __TRACE_RESOLVER(); if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull From 8b04caf2557691888a11025f33cd6c36d62de58f Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 18:13:41 +0100 Subject: [PATCH 05/15] Create singleton to context ETHASH --- sources/algo/ethash/ethash.cpp | 55 ++++++++++++------- sources/algo/ethash/ethash.hpp | 41 ++++++++++---- sources/algo/ethash/tests/ethash.cpp | 28 +++++----- sources/benchmark/amd/kawpow.cpp | 10 ++-- sources/common/cli/cli.hpp | 1 + sources/common/cli/cli_tools.cpp | 10 ++++ sources/common/config.cpp | 1 + sources/common/config.hpp | 1 + sources/resolver/amd/etchash.cpp | 9 ++- sources/resolver/amd/ethash.cpp | 11 ++-- sources/resolver/amd/progpow.cpp | 11 ++-- sources/resolver/amd/tests/ethash.cpp | 2 +- sources/resolver/amd/tests/firopow.cpp | 2 +- sources/resolver/amd/tests/kawpow.cpp | 2 +- sources/resolver/amd/tests/meowpow.cpp | 2 +- sources/resolver/amd/tests/progpow_quai.cpp | 2 +- sources/resolver/amd/tests/progpow_z.cpp | 2 +- sources/resolver/cpu/progpow.cpp | 6 +- sources/resolver/mocker.cpp | 24 +++++--- sources/resolver/nvidia/etchash.cpp | 6 +- sources/resolver/nvidia/ethash.cpp | 15 ++++- sources/resolver/nvidia/progpow.cpp | 14 +++-- sources/resolver/nvidia/tests/ethash.cpp | 2 +- sources/resolver/nvidia/tests/firopow.cpp | 2 +- sources/resolver/nvidia/tests/kawpow.cpp | 28 +++++++--- sources/resolver/nvidia/tests/meowpow.cpp | 2 +- .../resolver/nvidia/tests/progpow_quai.cpp | 2 +- sources/stratum/ethash.cpp | 2 +- sources/stratum/progpow.cpp | 4 +- 29 files changed, 188 insertions(+), 109 deletions(-) diff --git a/sources/algo/ethash/ethash.cpp b/sources/algo/ethash/ethash.cpp index 229b852..505416e 100644 --- a/sources/algo/ethash/ethash.cpp +++ b/sources/algo/ethash/ethash.cpp @@ -10,13 +10,16 @@ #include -boost::mutex mtxDagContext{}; -algo::DagContext localDagContext{}; +algo::ethash::ContextGenerator& algo::ethash::ContextGenerator::instance() +{ + static algo::ethash::ContextGenerator handler{}; + return handler; +} -int32_t algo::ethash::findEpoch( +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 }; @@ -60,34 +63,39 @@ int32_t algo::ethash::findEpoch( } -void algo::ethash::freeDagContext(algo::DagContext& context) +void algo::ethash::ContextGenerator::free(algo::ALGORITHM const algorithm) { //////////////////////////////////////////////////////////////////////////// - SAFE_DELETE_ARRAY(localDagContext.data); - localDagContext.lightCache.hash = nullptr; + UNIQUE_LOCK(mtxContexts); + + //////////////////////////////////////////////////////////////////////////// + algo::DagContext& context{ contexts[algorithm] }; + SAFE_DELETE_ARRAY(context.data); + context.lightCache.hash = nullptr; } -void algo::ethash::buildContext( +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) { //////////////////////////////////////////////////////////////////////////// - UNIQUE_LOCK(mtxDagContext); + UNIQUE_LOCK(mtxContexts); /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; + algo::DagContext& localDagContext{ contexts[algorithm] }; //////////////////////////////////////////////////////////////////////////// if (localDagContext.epoch == castU32(currentEpoch)) { - logInfo() << "Skip epoch: " << currentEpoch; - copyContext(context); + copyContext(context, algorithm); return; } @@ -137,19 +145,22 @@ void algo::ethash::buildContext( algo::copyHash(localDagContext.hashedSeedCache, seedHash); //////////////////////////////////////////////////////////////////////////// - if (true == config.deviceAlgorithm.ethashBuildLightCacheCPU) + if (true == buildOnCPU) { - algo::ethash::buildLightCacheOnCPU(context); + buildLightCache(algorithm); } //////////////////////////////////////////////////////////////////////////// - copyContext(context); + copyContext(context, algorithm); } -void algo::ethash::buildLightCacheOnCPU( - algo::DagContext& context) +void algo::ethash::ContextGenerator::buildLightCache( + algo::ALGORITHM const algorithm) { + //////////////////////////////////////////////////////////////////////////// + algo::DagContext& localDagContext{ contexts[algorithm] }; + //////////////////////////////////////////////////////////////////////////// algo::hash512 item{ localDagContext.hashedSeedCache }; size_t const dataLength{ algo::LEN_HASH_512 + localDagContext.lightCache.size }; @@ -192,8 +203,13 @@ void algo::ethash::buildLightCacheOnCPU( } -void algo::ethash::copyContext(algo::DagContext& context) +void algo::ethash::ContextGenerator::copyContext( + algo::DagContext& context, + algo::ALGORITHM const algorithm) { + //////////////////////////////////////////////////////////////////////////// + algo::DagContext& localDagContext{ contexts[algorithm] }; + context.epoch = localDagContext.epoch; context.lightCache.numberItem = localDagContext.lightCache.numberItem; context.lightCache.size = localDagContext.lightCache.size; @@ -204,4 +220,3 @@ void algo::ethash::copyContext(algo::DagContext& context) algo::copyHash(context.hashedSeedCache, localDagContext.hashedSeedCache); } - diff --git a/sources/algo/ethash/ethash.hpp b/sources/algo/ethash/ethash.hpp index ae5f7ee..d2e5967 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,32 @@ 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 buildContext(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 buildLightCacheOnCPU(algo::DagContext& context); - void freeDagContext(algo::DagContext& context); - void copyContext(algo::DagContext& context); + 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: + boost::mutex mtxContexts{}; + std::map contexts{}; + + ContextGenerator() = default; + void buildLightCache(algo::ALGORITHM const algorithm); + void copyContext(algo::DagContext& context, + algo::ALGORITHM const algorithm); + }; #endif } } diff --git a/sources/algo/ethash/tests/ethash.cpp b/sources/algo/ethash/tests/ethash.cpp index 6bc8980..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,15 +55,17 @@ TEST_F(EthashTest, lightCacheBuild) uint32_t dagCountItemsGrowth{ algo::ethash::DAG_COUNT_ITEMS_GROWTH }; uint32_t dagCountItemsInit{ algo::ethash::DAG_COUNT_ITEMS_INIT }; - algo::ethash::buildContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::ETHASH, context, 561ull, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true ); // light cache size @@ -110,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/benchmark/amd/kawpow.cpp b/sources/benchmark/amd/kawpow.cpp index 7b89319..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,15 +50,17 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow() //////////////////////////////////////////////////////////////////////////// algo::DagContext dagContext{}; - algo::ethash::buildContext + 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*/ ); //////////////////////////////////////////////////////////////////////////// @@ -227,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.hpp b/sources/common/cli/cli.hpp index ed1e5f8..0af677d 100644 --- a/sources/common/cli/cli.hpp +++ b/sources/common/cli/cli.hpp @@ -92,6 +92,7 @@ namespace common #if defined(TOOLS_ENABLE) && defined(TOOL_MOCKER) uint32_t getMockerResolverCount() const; + uint32_t getMockerResolverUpdateMemorySleep() const; #endif // Device settings custom diff --git a/sources/common/cli/cli_tools.cpp b/sources/common/cli/cli_tools.cpp index 3268fb9..7813bb7 100644 --- a/sources/common/cli/cli_tools.cpp +++ b/sources/common/cli/cli_tools.cpp @@ -14,4 +14,14 @@ uint32_t common::Cli::getMockerResolverCount() const } +uint32_t common::Cli::getMockerResolverUpdateMemorySleep() const +{ + if (true == contains("tool_mocker_resolver_update_memory_sleep")) + { + return params["tool_mocker_resolver_update_memory_sleep"].as(); + } + + return 1000u; +} + #endif // TOOLS_ENABLE && TOOL_MOCKER diff --git a/sources/common/config.cpp b/sources/common/config.cpp index e71e4c5..b0a2877 100644 --- a/sources/common/config.cpp +++ b/sources/common/config.cpp @@ -487,6 +487,7 @@ bool common::Config::loadCli(int argc, char** argv) //////////////////////////////////////////////////////////////////////// #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 461cba3..fd29df0 100644 --- a/sources/common/config.hpp +++ b/sources/common/config.hpp @@ -88,6 +88,7 @@ namespace common { #if defined(TOOL_MOCKER) std::optional mockerResolverCount{ 8u }; + std::optional mockerResolverUpdateMemorySleep{ 8u }; #endif }; #endif diff --git a/sources/resolver/amd/etchash.cpp b/sources/resolver/amd/etchash.cpp index 7f41dc1..6f30295 100644 --- a/sources/resolver/amd/etchash.cpp +++ b/sources/resolver/amd/etchash.cpp @@ -15,18 +15,17 @@ bool resolver::ResolverAmdEtchash::updateContext( stratum::StratumJobInfo const& jobInfo) { /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; - - /////////////////////////////////////////////////////////////////////////// - algo::ethash::buildContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::ETCHASH, context, jobInfo.epoch, algo::ethash::EIP1099_MAX_EPOCH_NUMBER, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true /*config.deviceAlgorithm.ethashBuildLightCacheCPU*/ ); if ( context.lightCache.numberItem == 0ull diff --git a/sources/resolver/amd/ethash.cpp b/sources/resolver/amd/ethash.cpp index 6261dd4..e11b88b 100644 --- a/sources/resolver/amd/ethash.cpp +++ b/sources/resolver/amd/ethash.cpp @@ -25,18 +25,17 @@ bool resolver::ResolverAmdEthash::updateContext( stratum::StratumJobInfo const& jobInfo) { /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; - - /////////////////////////////////////////////////////////////////////////// - algo::ethash::buildContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::ETHASH, context, jobInfo.epoch, algo::ethash::MAX_EPOCH_NUMBER, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true /*config.deviceAlgorithm.ethashBuildLightCacheCPU*/ ); if ( context.lightCache.numberItem == 0ull @@ -124,7 +123,7 @@ bool resolver::ResolverAmdEthash::updateMemory( } //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(context); + algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::ETHASH); //////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/amd/progpow.cpp b/sources/resolver/amd/progpow.cpp index f089912..95d3435 100644 --- a/sources/resolver/amd/progpow.cpp +++ b/sources/resolver/amd/progpow.cpp @@ -21,18 +21,17 @@ bool resolver::ResolverAmdProgPOW::updateContext( stratum::StratumJobInfo const& jobInfo) { /////////////////////////////////////////////////////////////////////////// - common::Config& config{ common::Config::instance() }; - - /////////////////////////////////////////////////////////////////////////// - algo::ethash::buildContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::PROGPOW, context, jobInfo.epoch, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true /*config.deviceAlgorithm.ethashBuildLightCacheCPU*/ ); if ( 0ull == context.lightCache.numberItem @@ -112,7 +111,7 @@ bool resolver::ResolverAmdProgPOW::updateMemory( } //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(context); + algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::PROGPOW); //////////////////////////////////////////////////////////////////////////// return true; 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/progpow.cpp b/sources/resolver/cpu/progpow.cpp index b89ccf5..203b78a 100644 --- a/sources/resolver/cpu/progpow.cpp +++ b/sources/resolver/cpu/progpow.cpp @@ -17,15 +17,17 @@ resolver::ResolverCpuProgPOW::~ResolverCpuProgPOW() bool resolver::ResolverCpuProgPOW::updateContext(stratum::StratumJobInfo const& jobInfo) { //////////////////////////////////////////////////////////////////////////// - algo::ethash::buildContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::PROGPOW, context, jobInfo.epoch, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + true ); if ( 0ull == context.lightCache.numberItem diff --git a/sources/resolver/mocker.cpp b/sources/resolver/mocker.cpp index c070511..7113c4c 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 }; @@ -27,31 +26,38 @@ void resolver::ResolverMocker::overrideOccupancy( bool resolver::ResolverMocker::updateMemory( [[maybe_unused]] stratum::StratumJobInfo const& jobInfo) { + /////////////////////////////////////////////////////////////////////////// + common::Config& config{ common::Config::instance() }; + /////////////////////////////////////////////////////////////////////////// algo::DagContext context{}; 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 }; - __TRACE_RESOLVER(); - algo::ethash::buildContext + + /////////////////////////////////////////////////////////////////////////// + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::PROGPOW, context, jobInfo.epoch, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + config.deviceAlgorithm.ethashBuildLightCacheCPU ); - __TRACE_RESOLVER(); /////////////////////////////////////////////////////////////////////////// - boost::this_thread::sleep_for(WAIT_UPDATE_MEMORY); + 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/etchash.cpp b/sources/resolver/nvidia/etchash.cpp index e6b7fa1..ef4a0bb 100644 --- a/sources/resolver/nvidia/etchash.cpp +++ b/sources/resolver/nvidia/etchash.cpp @@ -16,15 +16,17 @@ bool resolver::ResolverNvidiaEtchash::updateContext( common::Config& config{ common::Config::instance() }; /////////////////////////////////////////////////////////////////////////// - algo::ethash::buildContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::ETCHASH, context, jobInfo.epoch, algo::ethash::EIP1099_MAX_EPOCH_NUMBER, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + config.deviceAlgorithm.ethashBuildLightCacheCPU ); /////////////////////////////////////////////////////////////////////////// diff --git a/sources/resolver/nvidia/ethash.cpp b/sources/resolver/nvidia/ethash.cpp index c839c73..c6ee9bb 100644 --- a/sources/resolver/nvidia/ethash.cpp +++ b/sources/resolver/nvidia/ethash.cpp @@ -18,17 +18,24 @@ resolver::ResolverNvidiaEthash::~ResolverNvidiaEthash() bool resolver::ResolverNvidiaEthash::updateContext( stratum::StratumJobInfo const& jobInfo) { - algo::ethash::buildContext + /////////////////////////////////////////////////////////////////////////// + common::Config& config{ common::Config::instance() }; + + /////////////////////////////////////////////////////////////////////////// + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::ETHASH, context, jobInfo.epoch, algo::ethash::MAX_EPOCH_NUMBER, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + config.deviceAlgorithm.ethashBuildLightCacheCPU ); + /////////////////////////////////////////////////////////////////////////// if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull || context.dagCache.numberItem == 0ull @@ -46,6 +53,7 @@ bool resolver::ResolverNvidiaEthash::updateContext( return false; } + /////////////////////////////////////////////////////////////////////////// uint64_t const totalMemoryNeeded{ (context.dagCache.size + context.lightCache.size) }; if ( 0ull != deviceMemoryAvailable && totalMemoryNeeded >= deviceMemoryAvailable) @@ -56,6 +64,7 @@ bool resolver::ResolverNvidiaEthash::updateContext( return false; } + /////////////////////////////////////////////////////////////////////////// return true; } @@ -112,7 +121,7 @@ bool resolver::ResolverNvidiaEthash::updateMemory( CU_SAFE_DELETE(parameters.lightCache); //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(context); + algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::ETCHASH); //////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index 66a8a17..0b226b9 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -23,19 +23,20 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( common::Config& config{ common::Config::instance() }; //////////////////////////////////////////////////////////////////////////// - __TRACE_RESOLVER(); - algo::ethash::buildContext + algo::ethash::ContextGenerator::instance().build ( + algo::ALGORITHM::PROGPOW, context, jobInfo.epoch, maxEpoch, dagCountItemsGrowth, dagCountItemsInit, lightCacheCountItemsGrowth, - lightCacheCountItemsInit + lightCacheCountItemsInit, + config.deviceAlgorithm.ethashBuildLightCacheCPU ); - __TRACE_RESOLVER(); + //////////////////////////////////////////////////////////////////////////// if ( context.lightCache.numberItem == 0ull || context.lightCache.size == 0ull || context.dagCache.numberItem == 0ull @@ -53,7 +54,7 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( return false; } - + //////////////////////////////////////////////////////////////////////////// uint64_t const totalMemoryNeeded{ context.dagCache.size + context.lightCache.size }; if ( 0ull != deviceMemoryAvailable && totalMemoryNeeded >= deviceMemoryAvailable) @@ -65,6 +66,7 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( return false; } + //////////////////////////////////////////////////////////////////////////// return true; } @@ -126,7 +128,7 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( CU_SAFE_DELETE(parameters.lightCache); //////////////////////////////////////////////////////////////////////////// - algo::ethash::freeDagContext(context); + algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::PROGPOW); //////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/nvidia/tests/ethash.cpp b/sources/resolver/nvidia/tests/ethash.cpp index 72ddd59..b7221ed 100644 --- a/sources/resolver/nvidia/tests/ethash.cpp +++ b/sources/resolver/nvidia/tests/ethash.cpp @@ -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 9545022..5c3a64d 100644 --- a/sources/resolver/nvidia/tests/firopow.cpp +++ b/sources/resolver/nvidia/tests/firopow.cpp @@ -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 acf6d82..66877c2 100644 --- a/sources/resolver/nvidia/tests/kawpow.cpp +++ b/sources/resolver/nvidia/tests/kawpow.cpp @@ -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 4148306..56630b9 100644 --- a/sources/resolver/nvidia/tests/meowpow.cpp +++ b/sources/resolver/nvidia/tests/meowpow.cpp @@ -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 05a4e9f..45323cb 100644 --- a/sources/resolver/nvidia/tests/progpow_quai.cpp +++ b/sources/resolver/nvidia/tests/progpow_quai.cpp @@ -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/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..c8d44a8 100644 --- a/sources/stratum/progpow.cpp +++ b/sources/stratum/progpow.cpp @@ -340,7 +340,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) { @@ -413,7 +413,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) { From 9855e66dd5d24ed50b6c6b2ac40161a768536e50 Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 18:15:20 +0100 Subject: [PATCH 06/15] Create singleton to context ETHASH --- sources/algo/progpow/cuda/dag.cuh | 1 + 1 file changed, 1 insertion(+) diff --git a/sources/algo/progpow/cuda/dag.cuh b/sources/algo/progpow/cuda/dag.cuh index c756fec..0b0bafb 100644 --- a/sources/algo/progpow/cuda/dag.cuh +++ b/sources/algo/progpow/cuda/dag.cuh @@ -198,6 +198,7 @@ bool progpowBuildDag( for (uint32_t i{ 0u }; i < dagNumberItems; i += itemByKernel) { + printf("%u/%u\n", i, dagNumberItems); kernelProgpowBuildDag<<<512, 128, 0, stream>>>(i, loop); CUDA_ER(cudaStreamSynchronize(stream)); CUDA_ER(cudaGetLastError()); From b55e007267bdd31b7fc3aa10c0e13b2ee34a880f Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 18:17:53 +0100 Subject: [PATCH 07/15] DEBUG print dag item building --- sources/algo/progpow/cuda/dag.cuh | 5 +++-- sources/algo/progpow/cuda/progpow.cuh | 2 +- sources/resolver/nvidia/progpow.cpp | 2 +- 3 files changed, 5 insertions(+), 4 deletions(-) diff --git a/sources/algo/progpow/cuda/dag.cuh b/sources/algo/progpow/cuda/dag.cuh index 0b0bafb..7585e65 100644 --- a/sources/algo/progpow/cuda/dag.cuh +++ b/sources/algo/progpow/cuda/dag.cuh @@ -191,14 +191,15 @@ __host__ bool progpowBuildDag( cudaStream_t stream, uint32_t const dagItemParents, - uint32_t const dagNumberItems) + uint32_t const dagNumberItems, + uint32_t const deviceID) { uint32_t const itemByKernel{ 65536u }; uint32_t const loop{ dagItemParents / 4u / 4u }; for (uint32_t i{ 0u }; i < dagNumberItems; i += itemByKernel) { - printf("%u/%u\n", i, dagNumberItems); + printf("[%u] BUILDING DAG ITEM(%u/%u)\n", deviceID, i, dagNumberItems); kernelProgpowBuildDag<<<512, 128, 0, stream>>>(i, loop); CUDA_ER(cudaStreamSynchronize(stream)); CUDA_ER(cudaGetLastError()); diff --git a/sources/algo/progpow/cuda/progpow.cuh b/sources/algo/progpow/cuda/progpow.cuh index 8c2d6a1..701013d 100644 --- a/sources/algo/progpow/cuda/progpow.cuh +++ b/sources/algo/progpow/cuda/progpow.cuh @@ -16,4 +16,4 @@ bool progpowUpdateConstants(uint32_t const* const headerSrc, uint32_t* const headerDst); bool progpowBuildDag(cudaStream_t stream, uint32_t const dagItemParents, - uint32_t const dagNumberItems); + uint32_t const dagNumberItems, uint32_t const deviceID); diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index 0b226b9..933bc63 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -117,7 +117,7 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( chrono.start(); if (false == progpowBuildDag(cuStream[currentIndexStream], dagItemParents, - castU32(context.dagCache.numberItem))) + castU32(context.dagCache.numberItem), deviceId)) { return false; } From 59ad972f9391857a894e0d1f72063ea0a938d19a Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 18:36:33 +0100 Subject: [PATCH 08/15] remove printf during build DAG item, tools mocker use std::optional flags --- sources/algo/progpow/cuda/dag.cuh | 1 - sources/common/cli/cli.hpp | 4 ++-- sources/common/cli/cli_tools.cpp | 8 ++++---- sources/device/device_manager.cpp | 18 +++++++++++------- sources/resolver/mocker.cpp | 15 +++++++++------ 5 files changed, 26 insertions(+), 20 deletions(-) diff --git a/sources/algo/progpow/cuda/dag.cuh b/sources/algo/progpow/cuda/dag.cuh index 7585e65..18d43ee 100644 --- a/sources/algo/progpow/cuda/dag.cuh +++ b/sources/algo/progpow/cuda/dag.cuh @@ -199,7 +199,6 @@ bool progpowBuildDag( for (uint32_t i{ 0u }; i < dagNumberItems; i += itemByKernel) { - printf("[%u] BUILDING DAG ITEM(%u/%u)\n", deviceID, i, dagNumberItems); kernelProgpowBuildDag<<<512, 128, 0, stream>>>(i, loop); CUDA_ER(cudaStreamSynchronize(stream)); CUDA_ER(cudaGetLastError()); diff --git a/sources/common/cli/cli.hpp b/sources/common/cli/cli.hpp index 0af677d..dffa210 100644 --- a/sources/common/cli/cli.hpp +++ b/sources/common/cli/cli.hpp @@ -91,8 +91,8 @@ namespace common bool isCpuEnable() const; #if defined(TOOLS_ENABLE) && defined(TOOL_MOCKER) - uint32_t getMockerResolverCount() const; - uint32_t getMockerResolverUpdateMemorySleep() const; + std::optional getMockerResolverCount() const; + std::optional getMockerResolverUpdateMemorySleep() const; #endif // Device settings custom diff --git a/sources/common/cli/cli_tools.cpp b/sources/common/cli/cli_tools.cpp index 7813bb7..b1435e6 100644 --- a/sources/common/cli/cli_tools.cpp +++ b/sources/common/cli/cli_tools.cpp @@ -3,25 +3,25 @@ #include -uint32_t common::Cli::getMockerResolverCount() const +std::optional common::Cli::getMockerResolverCount() const { if (true == contains("tool_mocker_resolver_count")) { return params["tool_mocker_resolver_count"].as(); } - return 8u; + return std::nullopt; } -uint32_t common::Cli::getMockerResolverUpdateMemorySleep() const +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 1000u; + return std::nullopt; } #endif // TOOLS_ENABLE && TOOL_MOCKER diff --git a/sources/device/device_manager.cpp b/sources/device/device_manager.cpp index 94af2d7..5b6b62d 100644 --- a/sources/device/device_manager.cpp +++ b/sources/device/device_manager.cpp @@ -289,16 +289,20 @@ bool device::DeviceManager::initializeMocker() common::Config const& config{ common::Config::instance() }; /////////////////////////////////////////////////////////////////////////// - for (uint32_t i{ 0u }; i < config.toolConfigs.mockerResolverCount; ++i) + 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 = 1000u + i; - device->deviceType = device::DEVICE_TYPE::MOCKER; - devices.push_back(device); } return true; } diff --git a/sources/resolver/mocker.cpp b/sources/resolver/mocker.cpp index 7113c4c..da1c767 100644 --- a/sources/resolver/mocker.cpp +++ b/sources/resolver/mocker.cpp @@ -52,12 +52,15 @@ bool resolver::ResolverMocker::updateMemory( ); /////////////////////////////////////////////////////////////////////////// - resolverInfo() - << "Update memory force waiting: " - << *config.toolConfigs.mockerResolverUpdateMemorySleep - << "ms"; - boost::chrono::milliseconds const ms{ *config.toolConfigs.mockerResolverUpdateMemorySleep }; - boost::this_thread::sleep_for(ms); + 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; From 56c7a48af939bb913f410c957d328647170385a1 Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 21:18:03 +0100 Subject: [PATCH 09/15] Assign algorithm in Resolver --- sources/resolver/amd/autolykos_v2.cpp | 7 +++++++ sources/resolver/amd/autolykos_v2.hpp | 2 +- sources/resolver/amd/etchash.cpp | 9 ++++++++- sources/resolver/amd/etchash.hpp | 2 +- sources/resolver/amd/ethash.cpp | 12 +++++++++++- sources/resolver/amd/ethash.hpp | 2 +- sources/resolver/amd/evrprogpow.cpp | 5 +++++ sources/resolver/amd/firopow.cpp | 5 +++++ sources/resolver/amd/kawpow.cpp | 5 +++++ sources/resolver/amd/meowpow.cpp | 6 ++++++ sources/resolver/amd/progpow.cpp | 12 +++++++++++- sources/resolver/amd/progpow.hpp | 2 +- sources/resolver/amd/progpow_quai.cpp | 5 +++++ sources/resolver/cpu/kawpow.cpp | 6 ++++++ sources/resolver/cpu/progpow.cpp | 12 +++++++++++- sources/resolver/cpu/progpow.hpp | 2 +- sources/resolver/mocker.cpp | 2 +- sources/resolver/nvidia/autolykos_v2.cpp | 7 +++++++ sources/resolver/nvidia/autolykos_v2.hpp | 2 +- sources/resolver/nvidia/blake3.cpp | 7 +++++++ sources/resolver/nvidia/blake3.hpp | 2 +- sources/resolver/nvidia/etchash.cpp | 9 ++++++++- sources/resolver/nvidia/etchash.hpp | 2 +- sources/resolver/nvidia/ethash.cpp | 12 +++++++++++- sources/resolver/nvidia/ethash.hpp | 2 +- sources/resolver/nvidia/evrprogpow.cpp | 5 +++++ sources/resolver/nvidia/firopow.cpp | 5 +++++ sources/resolver/nvidia/kawpow.cpp | 6 ++++++ sources/resolver/nvidia/meowpow.cpp | 6 ++++++ sources/resolver/nvidia/progpow.cpp | 17 ++++++++++++----- sources/resolver/nvidia/progpow.hpp | 2 +- sources/resolver/nvidia/progpow_quai.cpp | 3 +++ sources/resolver/resolver.hpp | 10 ++++++---- 33 files changed, 167 insertions(+), 26 deletions(-) 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 6f30295..63ab816 100644 --- a/sources/resolver/amd/etchash.cpp +++ b/sources/resolver/amd/etchash.cpp @@ -11,13 +11,20 @@ #include +resolver::ResolverAmdEtchash::ResolverAmdEtchash(): + resolver::ResolverAmdEthash() +{ + algorithm = algo::ALGORITHM::ETHASH; +} + + bool resolver::ResolverAmdEtchash::updateContext( stratum::StratumJobInfo const& jobInfo) { /////////////////////////////////////////////////////////////////////////// algo::ethash::ContextGenerator::instance().build ( - algo::ALGORITHM::ETCHASH, + algorithm, context, jobInfo.epoch, algo::ethash::EIP1099_MAX_EPOCH_NUMBER, 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 e11b88b..e7d493c 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(); @@ -27,7 +37,7 @@ bool resolver::ResolverAmdEthash::updateContext( /////////////////////////////////////////////////////////////////////////// algo::ethash::ContextGenerator::instance().build ( - algo::ALGORITHM::ETHASH, + algorithm, context, jobInfo.epoch, algo::ethash::MAX_EPOCH_NUMBER, 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 95d3435..8655eda 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(); @@ -23,7 +33,7 @@ bool resolver::ResolverAmdProgPOW::updateContext( /////////////////////////////////////////////////////////////////////////// algo::ethash::ContextGenerator::instance().build ( - algo::ALGORITHM::PROGPOW, + algorithm, context, jobInfo.epoch, maxEpoch, 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/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 203b78a..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); @@ -19,7 +29,7 @@ bool resolver::ResolverCpuProgPOW::updateContext(stratum::StratumJobInfo const& //////////////////////////////////////////////////////////////////////////// algo::ethash::ContextGenerator::instance().build ( - algo::ALGORITHM::PROGPOW, + algorithm, context, jobInfo.epoch, maxEpoch, 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 da1c767..7db4663 100644 --- a/sources/resolver/mocker.cpp +++ b/sources/resolver/mocker.cpp @@ -40,7 +40,7 @@ bool resolver::ResolverMocker::updateMemory( /////////////////////////////////////////////////////////////////////////// algo::ethash::ContextGenerator::instance().build ( - algo::ALGORITHM::PROGPOW, + algorithm, context, jobInfo.epoch, maxEpoch, 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 ef4a0bb..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) { @@ -18,7 +25,7 @@ bool resolver::ResolverNvidiaEtchash::updateContext( /////////////////////////////////////////////////////////////////////////// algo::ethash::ContextGenerator::instance().build ( - algo::ALGORITHM::ETCHASH, + algorithm, context, jobInfo.epoch, algo::ethash::EIP1099_MAX_EPOCH_NUMBER, 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 c6ee9bb..27bc0b7 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); @@ -24,7 +34,7 @@ bool resolver::ResolverNvidiaEthash::updateContext( /////////////////////////////////////////////////////////////////////////// algo::ethash::ContextGenerator::instance().build ( - algo::ALGORITHM::ETHASH, + algorithm, context, jobInfo.epoch, algo::ethash::MAX_EPOCH_NUMBER, 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 933bc63..14197bf 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); @@ -25,7 +35,7 @@ bool resolver::ResolverNvidiaProgPOW::updateContext( //////////////////////////////////////////////////////////////////////////// algo::ethash::ContextGenerator::instance().build ( - algo::ALGORITHM::PROGPOW, + algorithm, context, jobInfo.epoch, maxEpoch, @@ -109,7 +119,7 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( } else { - + algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::PROGPOW); } //////////////////////////////////////////////////////////////////////////// @@ -127,9 +137,6 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( //////////////////////////////////////////////////////////////////////////// CU_SAFE_DELETE(parameters.lightCache); - //////////////////////////////////////////////////////////////////////////// - algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::PROGPOW); - //////////////////////////////////////////////////////////////////////////// return true; } 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/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; From 4a03f46d78ad36970ba65731280199540d68ff3c Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 21:22:52 +0100 Subject: [PATCH 10/15] remove TRACE debug --- sources/stratum/progpow.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/sources/stratum/progpow.cpp b/sources/stratum/progpow.cpp index c8d44a8..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()) @@ -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; From 0c9c22dc366aa8ed2d913047491e55b5462acdab Mon Sep 17 00:00:00 2001 From: luminousminer Date: Thu, 12 Feb 2026 21:29:28 +0100 Subject: [PATCH 11/15] free context using algorithm resolver type --- sources/algo/ethash/ethash.cpp | 2 +- sources/resolver/amd/ethash.cpp | 2 +- sources/resolver/amd/progpow.cpp | 2 +- sources/resolver/nvidia/ethash.cpp | 2 +- sources/resolver/nvidia/progpow.cpp | 3 ++- 5 files changed, 6 insertions(+), 5 deletions(-) diff --git a/sources/algo/ethash/ethash.cpp b/sources/algo/ethash/ethash.cpp index 505416e..03baf83 100644 --- a/sources/algo/ethash/ethash.cpp +++ b/sources/algo/ethash/ethash.cpp @@ -93,7 +93,7 @@ void algo::ethash::ContextGenerator::build( algo::DagContext& localDagContext{ contexts[algorithm] }; //////////////////////////////////////////////////////////////////////////// - if (localDagContext.epoch == castU32(currentEpoch)) + if (localDagContext.epoch == cast32(currentEpoch)) { copyContext(context, algorithm); return; diff --git a/sources/resolver/amd/ethash.cpp b/sources/resolver/amd/ethash.cpp index e7d493c..9226d98 100644 --- a/sources/resolver/amd/ethash.cpp +++ b/sources/resolver/amd/ethash.cpp @@ -133,7 +133,7 @@ bool resolver::ResolverAmdEthash::updateMemory( } //////////////////////////////////////////////////////////////////////////// - algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::ETHASH); + algo::ethash::ContextGenerator::instance().free(algorithm); //////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/amd/progpow.cpp b/sources/resolver/amd/progpow.cpp index 8655eda..ea6455a 100644 --- a/sources/resolver/amd/progpow.cpp +++ b/sources/resolver/amd/progpow.cpp @@ -121,7 +121,7 @@ bool resolver::ResolverAmdProgPOW::updateMemory( } //////////////////////////////////////////////////////////////////////////// - algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::PROGPOW); + algo::ethash::ContextGenerator::instance().free(algorithm); //////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/nvidia/ethash.cpp b/sources/resolver/nvidia/ethash.cpp index 27bc0b7..02b6e44 100644 --- a/sources/resolver/nvidia/ethash.cpp +++ b/sources/resolver/nvidia/ethash.cpp @@ -131,7 +131,7 @@ bool resolver::ResolverNvidiaEthash::updateMemory( CU_SAFE_DELETE(parameters.lightCache); //////////////////////////////////////////////////////////////////////////// - algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::ETCHASH); + algo::ethash::ContextGenerator::instance().free(algorithm); //////////////////////////////////////////////////////////////////////////// return true; diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index 14197bf..5b156d5 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -114,12 +114,13 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( } chrono.stop(); resolverInfo() << "Light Cache built in " << chrono.elapsed(common::CHRONO_UNIT::MS) << "ms"; + //////////////////////////////////////////////////////////////////////////// CU_SAFE_DELETE(parameters.seedCache); } else { - algo::ethash::ContextGenerator::instance().free(algo::ALGORITHM::PROGPOW); + algo::ethash::ContextGenerator::instance().free(algorithm); } //////////////////////////////////////////////////////////////////////////// From 1641eb528524e117edb56c5096a8cba6cc4b3d25 Mon Sep 17 00:00:00 2001 From: luminousminer Date: Fri, 13 Feb 2026 10:40:47 +0100 Subject: [PATCH 12/15] protect context with counter --- sources/algo/ethash/cuda/memory.cuh | 14 ++- sources/algo/ethash/ethash.cpp | 35 +++++-- sources/algo/ethash/ethash.hpp | 8 +- sources/algo/progpow/cuda/dag.cuh | 137 +++++++++++++------------- sources/algo/progpow/cuda/memory.cuh | 28 ++++-- sources/algo/progpow/cuda/progpow.cuh | 2 +- sources/resolver/nvidia/ethash.cpp | 8 +- sources/resolver/nvidia/progpow.cpp | 10 +- 8 files changed, 151 insertions(+), 91 deletions(-) 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 03baf83..91d5a77 100644 --- a/sources/algo/ethash/ethash.cpp +++ b/sources/algo/ethash/ethash.cpp @@ -69,9 +69,22 @@ void algo::ethash::ContextGenerator::free(algo::ALGORITHM const algorithm) UNIQUE_LOCK(mtxContexts); //////////////////////////////////////////////////////////////////////////// - algo::DagContext& context{ contexts[algorithm] }; - SAFE_DELETE_ARRAY(context.data); - context.lightCache.hash = nullptr; + 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]; } @@ -90,7 +103,11 @@ void algo::ethash::ContextGenerator::build( UNIQUE_LOCK(mtxContexts); /////////////////////////////////////////////////////////////////////////// - algo::DagContext& localDagContext{ contexts[algorithm] }; + algo::ethash::ContextGenerator::DagContextShare& contextShare{ getContext(algorithm) }; + algo::DagContext& localDagContext{ contextShare.context }; + + //////////////////////////////////////////////////////////////////////////// + ++contextShare.count; //////////////////////////////////////////////////////////////////////////// if (localDagContext.epoch == cast32(currentEpoch)) @@ -120,7 +137,7 @@ void algo::ethash::ContextGenerator::build( } //////////////////////////////////////////////////////////////////////////// - uint64_t lightCacheNumItemsUpperBound { castU64(epochEIP) }; + uint64_t lightCacheNumItemsUpperBound{ castU64(epochEIP) }; lightCacheNumItemsUpperBound *= lightCacheCountItemsGrowth; lightCacheNumItemsUpperBound += lightCacheCountItemsInit; localDagContext.lightCache.numberItem = algo::largestPrime(lightCacheNumItemsUpperBound); @@ -159,7 +176,8 @@ void algo::ethash::ContextGenerator::buildLightCache( algo::ALGORITHM const algorithm) { //////////////////////////////////////////////////////////////////////////// - algo::DagContext& localDagContext{ contexts[algorithm] }; + algo::ethash::ContextGenerator::DagContextShare& contextShare{ getContext(algorithm) }; + algo::DagContext& localDagContext{ contextShare.context }; //////////////////////////////////////////////////////////////////////////// algo::hash512 item{ localDagContext.hashedSeedCache }; @@ -175,7 +193,7 @@ void algo::ethash::ContextGenerator::buildLightCache( //////////////////////////////////////////////////////////////////////////// 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{ 0ull }; i < localDagContext.lightCache.numberItem; ++i) @@ -208,7 +226,8 @@ void algo::ethash::ContextGenerator::copyContext( algo::ALGORITHM const algorithm) { //////////////////////////////////////////////////////////////////////////// - algo::DagContext& localDagContext{ contexts[algorithm] }; + algo::ethash::ContextGenerator::DagContextShare& contextShare{ getContext(algorithm) }; + algo::DagContext& localDagContext{ contextShare.context }; context.epoch = localDagContext.epoch; context.lightCache.numberItem = localDagContext.lightCache.numberItem; diff --git a/sources/algo/ethash/ethash.hpp b/sources/algo/ethash/ethash.hpp index d2e5967..12f7215 100644 --- a/sources/algo/ethash/ethash.hpp +++ b/sources/algo/ethash/ethash.hpp @@ -53,10 +53,16 @@ namespace algo bool const buildOnCPU); private: + struct DagContextShare + { + uint32_t count{ 0u }; + algo::DagContext context; + }; boost::mutex mtxContexts{}; - std::map contexts{}; + 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); diff --git a/sources/algo/progpow/cuda/dag.cuh b/sources/algo/progpow/cuda/dag.cuh index 18d43ee..a53656e 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) { - 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); } @@ -191,15 +195,16 @@ __host__ bool progpowBuildDag( cudaStream_t stream, uint32_t const dagItemParents, - uint32_t const dagNumberItems, - uint32_t const deviceID) + 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) { - 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/algo/progpow/cuda/progpow.cuh b/sources/algo/progpow/cuda/progpow.cuh index 701013d..8c2d6a1 100644 --- a/sources/algo/progpow/cuda/progpow.cuh +++ b/sources/algo/progpow/cuda/progpow.cuh @@ -16,4 +16,4 @@ bool progpowUpdateConstants(uint32_t const* const headerSrc, uint32_t* const headerDst); bool progpowBuildDag(cudaStream_t stream, uint32_t const dagItemParents, - uint32_t const dagNumberItems, uint32_t const deviceID); + uint32_t const dagNumberItems); diff --git a/sources/resolver/nvidia/ethash.cpp b/sources/resolver/nvidia/ethash.cpp index 02b6e44..c5c9742 100644 --- a/sources/resolver/nvidia/ethash.cpp +++ b/sources/resolver/nvidia/ethash.cpp @@ -104,12 +104,18 @@ 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"; + } + else + { + algo::ethash::ContextGenerator::instance().free(algorithm); } //////////////////////////////////////////////////////////////////////////// diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index 5b156d5..02fa156 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -113,22 +113,22 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( return false; } chrono.stop(); - resolverInfo() << "Light Cache built in " << chrono.elapsed(common::CHRONO_UNIT::MS) << "ms"; - - //////////////////////////////////////////////////////////////////////////// - CU_SAFE_DELETE(parameters.seedCache); + resolverInfo() << "Light Cache built on GPU in " << chrono.elapsed(common::CHRONO_UNIT::MS) << "ms"; } else { algo::ethash::ContextGenerator::instance().free(algorithm); } + //////////////////////////////////////////////////////////////////////////// + CU_SAFE_DELETE(parameters.seedCache); + //////////////////////////////////////////////////////////////////////////// resolverInfo() << "Building DAG"; chrono.start(); if (false == progpowBuildDag(cuStream[currentIndexStream], dagItemParents, - castU32(context.dagCache.numberItem), deviceId)) + castU32(context.dagCache.numberItem))) { return false; } From 350c33e428cdd899b421861b9b8213891895eaaa Mon Sep 17 00:00:00 2001 From: luminousminer Date: Fri, 13 Feb 2026 14:12:36 +0100 Subject: [PATCH 13/15] Free SeedCache if lightCache build from GPU --- sources/algo/ethash/result.hpp | 8 ++-- sources/resolver/nvidia/ethash.cpp | 13 ++++-- sources/resolver/nvidia/progpow.cpp | 62 ++++++++++++++--------------- 3 files changed, 44 insertions(+), 39 deletions(-) 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/resolver/nvidia/ethash.cpp b/sources/resolver/nvidia/ethash.cpp index c5c9742..de43b21 100644 --- a/sources/resolver/nvidia/ethash.cpp +++ b/sources/resolver/nvidia/ethash.cpp @@ -103,6 +103,7 @@ bool resolver::ResolverNvidiaEthash::updateMemory( //////////////////////////////////////////////////////////////////////////// if (false == config.deviceAlgorithm.ethashBuildLightCacheCPU) { + /////////////////////////////////////////////////////////////////////// resolverInfo() << "Building light cache on GPU"; chrono.start(); if (false == ethashBuildLightCache(cuStream[currentIndexStream], @@ -112,15 +113,15 @@ bool resolver::ResolverNvidiaEthash::updateMemory( } chrono.start(); resolverInfo() << "Built light cache on GPU in " << chrono.elapsed(common::CHRONO_UNIT::MS) << "ms"; + + /////////////////////////////////////////////////////////////////////// + CU_SAFE_DELETE(parameters.seedCache); } else { algo::ethash::ContextGenerator::instance().free(algorithm); } - //////////////////////////////////////////////////////////////////////////// - CU_SAFE_DELETE(parameters.seedCache); - //////////////////////////////////////////////////////////////////////////// resolverInfo() << "Building DAG"; chrono.start(); @@ -147,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)) @@ -157,6 +159,7 @@ bool resolver::ResolverNvidiaEthash::updateConstants( //////////////////////////////////////////////////////////////////////////// overrideOccupancy(128u, 8192u); + //////////////////////////////////////////////////////////////////////////// return true; } @@ -164,6 +167,7 @@ bool resolver::ResolverNvidiaEthash::updateConstants( bool resolver::ResolverNvidiaEthash::executeSync( stratum::StratumJobInfo const& jobInfo) { + //////////////////////////////////////////////////////////////////////////// ethashSearch(cuStream[currentIndexStream], ¶meters.resultCache[currentIndexStream], blocks, @@ -172,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) { @@ -194,6 +199,7 @@ bool resolver::ResolverNvidiaEthash::executeSync( resultCache->count = 0u; } + //////////////////////////////////////////////////////////////////////////// return true; } @@ -240,6 +246,7 @@ bool resolver::ResolverNvidiaEthash::executeAsync( //////////////////////////////////////////////////////////////////////////// swapIndexStream(); + //////////////////////////////////////////////////////////////////////////// return true; } diff --git a/sources/resolver/nvidia/progpow.cpp b/sources/resolver/nvidia/progpow.cpp index 02fa156..ee18371 100644 --- a/sources/resolver/nvidia/progpow.cpp +++ b/sources/resolver/nvidia/progpow.cpp @@ -105,6 +105,7 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( //////////////////////////////////////////////////////////////////////////// if (false == config.deviceAlgorithm.ethashBuildLightCacheCPU) { + /////////////////////////////////////////////////////////////////////// resolverInfo() << "Building LightCache on GPU"; chrono.start(); if (false == progpowBuildLightCache(cuStream[currentIndexStream], @@ -114,15 +115,15 @@ bool resolver::ResolverNvidiaProgPOW::updateMemory( } chrono.stop(); 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); } - //////////////////////////////////////////////////////////////////////////// - CU_SAFE_DELETE(parameters.seedCache); - //////////////////////////////////////////////////////////////////////////// resolverInfo() << "Building DAG"; chrono.start(); @@ -355,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()); @@ -378,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]; @@ -420,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 @@ -443,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]; @@ -467,6 +464,7 @@ bool resolver::ResolverNvidiaProgPOW::executeAsync( return true; } + void resolver::ResolverNvidiaProgPOW::submit( stratum::Stratum* const stratum) { @@ -474,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]; } @@ -555,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]; } From 3e8cc71d5e28f83799e745acf96db278508c3fde Mon Sep 17 00:00:00 2001 From: luminousminer Date: Fri, 13 Feb 2026 14:14:22 +0100 Subject: [PATCH 14/15] cuda coding style --- sources/algo/progpow/cuda/dag.cuh | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/sources/algo/progpow/cuda/dag.cuh b/sources/algo/progpow/cuda/dag.cuh index a53656e..ad0ec50 100644 --- a/sources/algo/progpow/cuda/dag.cuh +++ b/sources/algo/progpow/cuda/dag.cuh @@ -15,7 +15,7 @@ uint4 doCopy( //////////////////////////////////////////////////////////////////////////// #pragma unroll - for (uint32_t i{ 0u }; i < DAG_HASH_U4_SIZE; ++i) + for (uint32_t i = 0u; i < DAG_HASH_U4_SIZE; ++i) { uint4 const tmp_src = src[i]; uint32_t const x = reg_load(tmp_src.x, lane_id, DAG_PARRALLEL_LANE); @@ -182,10 +182,10 @@ void kernelProgpowBuildDag( } //////////////////////////////////////////////////////////////////////////// - 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 }; + 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); } @@ -197,12 +197,12 @@ bool progpowBuildDag( uint32_t const dagItemParents, uint32_t const dagNumberItems) { - 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 }; + 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<<>>(i, loop); CUDA_ER(cudaStreamSynchronize(stream)); From aa4033bc1d87aec2f4f066c70181e91c08d2c838 Mon Sep 17 00:00:00 2001 From: luminousminer Date: Fri, 13 Feb 2026 14:37:03 +0100 Subject: [PATCH 15/15] add parameters --log_new_job --- sources/common/cli/cli.cpp | 7 +++++++ sources/common/cli/cli.hpp | 3 ++- sources/common/cli/cli_log.cpp | 11 +++++++++++ sources/common/config.cpp | 3 +++ sources/common/config.hpp | 1 + sources/device/device_manager.cpp | 8 ++++++-- 6 files changed, 30 insertions(+), 3 deletions(-) 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 dffa210..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; 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/config.cpp b/sources/common/config.cpp index b0a2877..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 //////////////////////////////////////////////////////////////////////// diff --git a/sources/common/config.hpp b/sources/common/config.hpp index fd29df0..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 diff --git a/sources/device/device_manager.cpp b/sources/device/device_manager.cpp index 5b6b62d..1de3593 100644 --- a/sources/device/device_manager.cpp +++ b/sources/device/device_manager.cpp @@ -509,6 +509,7 @@ void device::DeviceManager::onUpdateJob( UNIQUE_LOCK(mtxJobInfo); //////////////////////////////////////////////////////////////////////////// + auto const& config{ common::Config::instance() }; stratum::StratumJobInfo& jobInfo{ jobInfos[stratumUUID] }; //////////////////////////////////////////////////////////////////////////// @@ -551,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); } }