Skip to content

Double queue buffering OpenCL for AMD device #108

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
luminousmining merged 7 commits into from
Feb 10, 2026
Merged

Double queue buffering OpenCL for AMD device #108

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
408a007
start bench lm4
luminousmining Feb 8, 2026
2e1e866
AMD: Add double queue buffering
luminousmining Feb 9, 2026
527b307
Get correct DAG context in benchmark, crosslane can use LDS
luminousmining Feb 10, 2026
284630c
Get correct DAG context in benchmark, crosslane can use LDS
luminousmining Feb 10, 2026
f4f7b52
split init context dag and build lightCache
luminousmining Feb 10, 2026
d316c32
cleanup
luminousmining Feb 10, 2026
a38e7ed
clean up print
luminousmining Feb 10, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 8 additions & 0 deletions sources/algo/ethash/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,14 @@ if (BUILD_EXE_UNIT_TEST)
)
endif()

if (BUILD_EXE_MINER)
target_sources(${BENCH_EXE} PUBLIC
${HEADERS}
${SOURCES}
)
endif()


add_subdirectory(opencl)
add_subdirectory(tests)

Expand Down
39 changes: 17 additions & 22 deletions sources/algo/ethash/ethash.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,8 +71,7 @@ void algo::ethash::initializeDagContext(
uint64_t const dagCountItemsGrowth,
uint64_t const dagCountItemsInit,
uint32_t const lightCacheCountItemsGrowth,
uint32_t const lightCacheCountItemsInit,
bool const buildOnCPU)
uint32_t const lightCacheCountItemsInit)
{
////////////////////////////////////////////////////////////////////////////
context.epoch = castU32(currentEpoch);
Expand Down Expand Up @@ -115,12 +114,23 @@ void algo::ethash::initializeDagContext(
seed = algo::keccak(seed);
}
algo::copyHash(context.originalSeedCache, seed);
}


void algo::ethash::buildLightCache(
algo::DagContext& context,
bool const buildOnCPU)
{
////////////////////////////////////////////////////////////////////////////
algo::hash512 item
{
algo::keccak<algo::hash512, algo::hash256>(context.originalSeedCache)
};

////////////////////////////////////////////////////////////////////////////
if (false == buildOnCPU)
{
algo::hash512 const hashedSeed{ algo::keccak<algo::hash512, algo::hash256>(seed) };
algo::copyHash(context.hashedSeedCache, hashedSeed);
algo::copyHash(context.hashedSeedCache, item);
return;
}

Expand All @@ -133,27 +143,12 @@ void algo::ethash::initializeDagContext(
logErr() << "Cannot alloc context data";
return;
}

////////////////////////////////////////////////////////////////////////////
context.lightCache.hash = castPtrHash512(context.data + algo::LEN_HASH_512);
context.lightCache.hash[0] = item;

////////////////////////////////////////////////////////////////////////////
if (true == buildOnCPU)
{
logInfo() << "Building light cache on CPU";
common::ChronoGuard chrono{ "Built light cache", common::CHRONO_UNIT::MS };
buildLightCache(context, seed);
}
}


void algo::ethash::buildLightCache(
algo::DagContext& context,
algo::hash256 const& seed)
{
////////////////////////////////////////////////////////////////////////////
algo::hash512 item{ algo::keccak<algo::hash512, algo::hash256>(seed) };
context.lightCache.hash[0] = item;
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)
Expand Down
5 changes: 2 additions & 3 deletions sources/algo/ethash/ethash.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,11 +40,10 @@ namespace algo
uint64_t const dagCountItemsGrowth,
uint64_t const dagCountItemsInit,
uint32_t const lightCacheCountItemsGrowth,
uint32_t const lightCacheCountItemsInit,
bool const buildOnCPU);
uint32_t const lightCacheCountItemsInit);
void freeDagContext(algo::DagContext& context);
void buildLightCache(algo::DagContext& context,
algo::hash256 const& seed);
bool const buildOnCPU);
#endif
}
}
4 changes: 2 additions & 2 deletions sources/algo/ethash/tests/ethash.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,9 +63,9 @@ TEST_F(EthashTest, lightCacheBuild)
dagCountItemsGrowth,
dagCountItemsInit,
lightCacheCountItemsGrowth,
lightCacheCountItemsInit,
true
lightCacheCountItemsInit
);
algo::ethash::buildLightCache(context, true);

// light cache size
ASSERT_EQ(1411061ull, context.lightCache.numberItem);
Expand Down
89 changes: 49 additions & 40 deletions sources/algo/progpow/opencl/progpow.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -122,60 +122,69 @@ void progpow_search(
uint const thread_id = get_global_id(0) + (get_global_id(1) * GROUP_SIZE);
uint const lane_id = thread_id % LANES;
uint const worker_group = get_global_id(0) / LANES;
ulong const nonce = start_nonce + thread_id;
ulong nonce = start_nonce + thread_id;
uint const index_share_seed = get_global_id(0) / BATCH_GROUP_LANE;

////////////////////////////////////////////////////////////////////////
initialize_header(dag, header_dag, (thread_id % GROUP_SIZE));
ulong const seed = initialize_seed(header, state_mix, nonce);

#if defined(INTERNAL_LOOP)
__attribute__((opencl_unroll_hint(1)))
for (uint l_id = 0u; l_id < LANES; ++l_id)
for (uint i = 0; i < INTERNAL_LOOP; ++i)
{
////////////////////////////////////////////////////////////////////////
if (l_id == lane_id)
nonce += (i * TOTAL_THREADS);
#endif
__attribute__((opencl_unroll_hint(1)))
for (uint l_id = 0u; l_id < LANES; ++l_id)
{
share_msb_lsb[index_share_seed] = seed;
////////////////////////////////////////////////////////////////////////
if (l_id == lane_id)
{
share_msb_lsb[index_share_seed] = seed;
}
barrier(CLK_LOCAL_MEM_FENCE);

////////////////////////////////////////////////////////////////////////
ulong const seedShare = share_msb_lsb[index_share_seed];

////////////////////////////////////////////////////////////////////////
fill_hash(hash, lane_id, seedShare);
loop_math(dag, share_hash0, header_dag, hash, lane_id, worker_group);
reduce_hash(
share_fnv1a,
hash,
digest,
worker_group,
l_id == lane_id);
}
barrier(CLK_LOCAL_MEM_FENCE);

////////////////////////////////////////////////////////////////////////
ulong const seedShare = share_msb_lsb[index_share_seed];

////////////////////////////////////////////////////////////////////////
fill_hash(hash, lane_id, seedShare);
loop_math(dag, share_hash0, header_dag, hash, lane_id, worker_group);
reduce_hash(
share_fnv1a,
hash,
digest,
worker_group,
l_id == lane_id);
}
#if defined(__KERNEL_PROGPOW)
ulong const bytes_result = is_valid(header, digest, seed);
#else
ulong const bytes_result = is_valid(state_mix, digest);
#endif

////////////////////////////////////////////////////////////////////////
#if defined(__KERNEL_PROGPOW)
ulong const bytes_result = is_valid(header, digest, seed);
#else
ulong const bytes_result = is_valid(state_mix, digest);
#endif

if (bytes_result <= boundary)
{
uint const index = atomic_inc(&result->count);
if (index < MAX_RESULT)
if (bytes_result <= boundary)
{
result->found = true;
result->nonces[index] = nonce;

result->hash[index][0] = digest[0];
result->hash[index][1] = digest[1];
result->hash[index][2] = digest[2];
result->hash[index][3] = digest[3];
result->hash[index][4] = digest[4];
result->hash[index][5] = digest[5];
result->hash[index][6] = digest[6];
result->hash[index][7] = digest[7];
uint const index = atomic_inc(&result->count);
if (index < MAX_RESULT)
{
result->found = true;
result->nonces[index] = nonce;

result->hash[index][0] = digest[0];
result->hash[index][1] = digest[1];
result->hash[index][2] = digest[2];
result->hash[index][3] = digest[3];
result->hash[index][4] = digest[4];
result->hash[index][5] = digest[5];
result->hash[index][6] = digest[6];
result->hash[index][7] = digest[7];
}
}
#if defined(INTERNAL_LOOP)
}
#endif
}
57 changes: 41 additions & 16 deletions sources/benchmark/amd/kawpow.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,8 @@

#include <CL/opencl.hpp>

#include <algo/dag_context.hpp>
#include <algo/ethash/ethash.hpp>
#include <algo/hash.hpp>
#include <algo/hash_utils.hpp>
#include <algo/progpow/progpow.hpp>
Expand All @@ -19,14 +21,18 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow()

////////////////////////////////////////////////////////////////////////////
bool dagInitialized{ false };
uint64_t const dagItems{ 16777213ull };
uint64_t const dagItemsKawpow{ dagItems / 2ull };
auto const header
algo::hash256 const headerHash
{
algo::toHash256("71c967486cb3b70d5dfcb2ebd8eeef138453637cacbf3ccb580a41a7e96986bb")
};
algo::hash256 const seedHash
{
algo::toHash256("7c4fb8a5d141973b69b521ce76b0dc50f0d2834d817c7f8310a6ab5becc6bb0c")
};
int32_t const epoch{ algo::ethash::findEpoch(seedHash, algo::ethash::EPOCH_LENGTH) };

////////////////////////////////////////////////////////////////////////////
common::opencl::Buffer<algo::hash512> lightCache { CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY };
common::opencl::Buffer<algo::hash1024> dagCache{ CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS };
common::opencl::BufferMapped<uint32_t> headerCache
{
Expand All @@ -36,14 +42,29 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow()
common::opencl::BufferMapped<t_result> resultCache{ CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR };

////////////////////////////////////////////////////////////////////////////
dagCache.setSize(dagItems * algo::LEN_HASH_1024);
algo::DagContext dagContext{};
algo::ethash::initializeDagContext
(
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
);

////////////////////////////////////////////////////////////////////////////
dagCache.setSize(dagContext.dagCache.size);

////////////////////////////////////////////////////////////////////////////
dagCache.alloc(propertiesAmd.clContext);
headerCache.alloc(&propertiesAmd.clQueue, propertiesAmd.clContext);
resultCache.alloc(&propertiesAmd.clQueue, propertiesAmd.clContext);

////////////////////////////////////////////////////////////////////////////
if (false == headerCache.setBufferDevice(&propertiesAmd.clQueue,
(uint32_t*)header.word32))
(uint32_t*)headerHash.word32))
{
logErr() << "Fail to copy header in cache";
}
Expand All @@ -66,7 +87,7 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow()
///////////////////////////////////////////////////////////////////
auto& clKernel{ generator.clKernel };
OPENCL_ER(clKernel.setArg(0u, *dagCache.getBuffer()));
OPENCL_ER(clKernel.setArg(1u, castU32(dagItems)));
OPENCL_ER(clKernel.setArg(1u, castU32(dagContext.dagCache.numberItem)));

///////////////////////////////////////////////////////////////////
OPENCL_ER(
Expand All @@ -87,10 +108,10 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow()

///////////////////////////////////////////////////////////////////////////
auto benchKawpow = [&](std::string const& kernelName,
uint32_t const loop,
uint32_t const groupSize,
uint32_t const workerGroupCount,
uint32_t const workItemCollaborate,
uint32_t const loop) -> bool
uint32_t const workItemCollaborate) -> bool
{
///////////////////////////////////////////////////////////////////////
common::KernelGeneratorOpenCL generator{};
Expand All @@ -113,7 +134,7 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow()
generator.addDefine("MODULE_CACHE", algo::progpow::MODULE_CACHE);
generator.addDefine("HEADER_ITEM_BY_THREAD", algo::progpow::MODULE_CACHE / groupSize);
generator.addDefine("COUNT_DAG", algo::progpow::COUNT_DAG);
generator.addDefine("DAG_SIZE", castU32(dagItemsKawpow));
generator.addDefine("DAG_SIZE", castU32(dagContext.dagCache.numberItem / 2ull));
generator.addDefine("BATCH_GROUP_LANE", batchGroupLane);
generator.addDefine("SHARE_SEED_SIZE", batchGroupLane);
generator.addDefine("SHARE_HASH0_SIZE", batchGroupLane);
Expand Down Expand Up @@ -145,10 +166,10 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow()

///////////////////////////////////////////////////////////////////////
auto& clKernel{ generator.clKernel };
OPENCL_ER(clKernel.setArg(0u, 0ull));
OPENCL_ER(clKernel.setArg(1u, *headerCache.getBuffer()));
OPENCL_ER(clKernel.setArg(2u, *dagCache.getBuffer()));
OPENCL_ER(clKernel.setArg(3u, *resultCache.getBuffer()));
OPENCL_ER(clKernel.setArg(0u, *dagCache.getBuffer()));
OPENCL_ER(clKernel.setArg(1u, *resultCache.getBuffer()));
OPENCL_ER(clKernel.setArg(2u, *headerCache.getBuffer()));
OPENCL_ER(clKernel.setArg(3u, 0ull));

///////////////////////////////////////////////////////////////////////
setGrid(groupSize, workerGroupCount);
Expand Down Expand Up @@ -177,11 +198,15 @@ bool benchmark::BenchmarkWorkflow::runAmdKawpow()
////////////////////////////////////////////////////////////////////////////
if (true == dagInitialized)
{
benchKawpow("kawpow_lm1", 256u, 1024u, algo::progpow::LANES, 1u); // Parallele + LDS
benchKawpow("kawpow_lm2", 256u, 1024u, algo::progpow::LANES, 1u); // Parallel + crosslane
benchKawpow("kawpow_lm3", 256u, 1024u, algo::progpow::LANES, 1u); // Parallel + crosslane + LDS header
benchKawpow("kawpow_lm1", 1u, 256u, 1024u, algo::progpow::LANES); // Parallele + LDS
benchKawpow("kawpow_lm2", 1u, 256u, 1024u, algo::progpow::LANES); // Parallel + crosslane
benchKawpow("kawpow_lm3", 1u, 256u, 1024u, algo::progpow::LANES); // Parallel + crosslane + LDS header
benchKawpow("kawpow_lm4", 1u, 256u, 1024u, algo::progpow::LANES); // Parallel + crosslane + LDS header
}

////////////////////////////////////////////////////////////////////////////
algo::ethash::freeDagContext(dagContext);

////////////////////////////////////////////////////////////////////////////
dagCache.free();
headerCache.free();
Expand Down
2 changes: 2 additions & 0 deletions sources/benchmark/opencl/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@ set(BENCHMARK_OPENCL_FILES
kawpow/kawpow_lm1.cl
kawpow/kawpow_lm2.cl
kawpow/kawpow_lm3.cl
kawpow/kawpow_lm4.cl
kawpow/sequence_dynamic.cl
kawpow/sequence_dynamic_local.cl
)
Expand All @@ -28,6 +29,7 @@ add_custom_target(copy_benchmark_opencl ALL
${OUT_COMMON}/kawpow/kawpow_lm1.cl
${OUT_COMMON}/kawpow/kawpow_lm2.cl
${OUT_COMMON}/kawpow/kawpow_lm3.cl
${OUT_COMMON}/kawpow/kawpow_lm4.cl
${OUT_COMMON}/kawpow/sequence_dynamic.cl
${OUT_COMMON}/kawpow/sequence_dynamic_local.cl
)
Loading
Loading