Skip to content

Commit 454ec51

Browse files
luminousminingluminousmining
authored
Merge pull request #184 from yuzi-co/feat/kheavyhash-opencl-common
refactor(kheavyhash): reuse common OpenCL helpers + add resolver tests
2 parents 4e5b03c + 20ce8be commit 454ec51

7 files changed

Lines changed: 266 additions & 106 deletions

File tree

‎sources/algo/kheavyhash/opencl/CMakeLists.txt‎

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,7 @@ file(MAKE_DIRECTORY ${OUT_KHEAVYHASH})
44

55
set(KHEAVYHASH_FILES
66
kheavyhash.cl
7+
kheavyhash_result.cl
78
)
89

910
foreach(FILE ${KHEAVYHASH_FILES})
@@ -18,6 +19,7 @@ endforeach()
1819
add_custom_target(copy_kheavyhash_opencl_files ALL
1920
DEPENDS
2021
${OUT_KHEAVYHASH}/kheavyhash.cl
22+
${OUT_KHEAVYHASH}/kheavyhash_result.cl
2123
)
2224

2325
if (BUILD_EXE_UNIT_TEST AND BUILD_AMD)
@@ -26,5 +28,7 @@ if (BUILD_EXE_UNIT_TEST AND BUILD_AMD)
2628
${CMAKE_CURRENT_SOURCE_DIR}/../tests
2729
${OpenCL_INCLUDE_DIRS})
2830
target_compile_definitions(${UNIT_TEST_EXE}
29-
PRIVATE KH_CL_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kheavyhash.cl")
31+
PRIVATE KH_CL_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kheavyhash.cl"
32+
KH_CL_DIR="${CMAKE_CURRENT_SOURCE_DIR}"
33+
KH_CL_COMMON_DIR="${CMAKE_CURRENT_SOURCE_DIR}/../../../common/opencl")
3034
endif()

‎sources/algo/kheavyhash/opencl/kheavyhash.cl‎

Lines changed: 38 additions & 92 deletions
Original file line numberDiff line numberDiff line change
@@ -44,35 +44,6 @@ __constant int PI_LANE[24] = { 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
4444
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 };
4545

4646

47-
inline ulong rotl64(ulong const x, int const k)
48-
{
49-
return (x << k) | (x >> (64 - k));
50-
}
51-
52-
53-
inline ulong loadLe64(uchar const* p)
54-
{
55-
ulong v = 0;
56-
for (int b = 0; b < 8; ++b)
57-
{
58-
v |= ((ulong)p[b]) << (8 * b);
59-
}
60-
return v;
61-
}
62-
63-
64-
inline void storeLe256(ulong const* state, uchar* out)
65-
{
66-
for (int w = 0; w < 4; ++w)
67-
{
68-
for (int b = 0; b < 8; ++b)
69-
{
70-
out[w * 8 + b] = (uchar)((state[w] >> (8 * b)) & 0xFF);
71-
}
72-
}
73-
}
74-
75-
7647
void keccakF1600(ulong* a)
7748
{
7849
for (int round = 0; round < 24; ++round)
@@ -85,7 +56,7 @@ void keccakF1600(ulong* a)
8556
}
8657
for (int i = 0; i < 5; ++i)
8758
{
88-
ulong const t = bc[(i + 4) % 5] ^ rotl64(bc[(i + 1) % 5], 1);
59+
ulong const t = bc[(i + 4) % 5] ^ rol_u64(bc[(i + 1) % 5], 1);
8960
for (int j = 0; j < 25; j += 5)
9061
{
9162
a[j + i] ^= t;
@@ -98,7 +69,7 @@ void keccakF1600(ulong* a)
9869
{
9970
int const j = PI_LANE[i];
10071
ulong const tmp = a[j];
101-
a[j] = rotl64(t, ROTATIONS[i]);
72+
a[j] = rol_u64(t, ROTATIONS[i]);
10273
t = tmp;
10374
}
10475

@@ -131,12 +102,12 @@ void powHash(uchar const* prePowHash, ulong const timestamp, ulong const nonce,
131102
}
132103
for (int w = 0; w < 4; ++w)
133104
{
134-
state[w] ^= loadLe64(prePowHash + w * 8);
105+
state[w] ^= load_le_u64(prePowHash + w * 8);
135106
}
136107
state[4] ^= timestamp;
137108
state[9] ^= nonce;
138109
keccakF1600(state);
139-
storeLe256(state, out);
110+
store_le_u256(state, out);
140111
}
141112

142113

@@ -150,10 +121,10 @@ void kHeavyHash(uchar const* input, uchar* out)
150121
}
151122
for (int w = 0; w < 4; ++w)
152123
{
153-
state[w] ^= loadLe64(input + w * 8);
124+
state[w] ^= load_le_u64(input + w * 8);
154125
}
155126
keccakF1600(state);
156-
storeLe256(state, out);
127+
store_le_u256(state, out);
157128
}
158129

159130

@@ -257,20 +228,6 @@ __kernel void test_heavy_hash(__global ushort const* matrix,
257228

258229

259230

260-
// Result buffer shared with the host (mirrors algo::ethash/blake3 Result).
261-
// MAX_RESULT is overridable by the host kernel generator (addDefine).
262-
#ifndef MAX_RESULT
263-
#define MAX_RESULT 4
264-
#endif
265-
266-
typedef struct __attribute__((aligned(8)))
267-
{
268-
uchar found;
269-
uint count;
270-
ulong nonces[MAX_RESULT];
271-
} Result;
272-
273-
274231
// Real mining kernel: each work-item tries nonce = startNonce + global_id(0).
275232
// On a hit (pow <= target, little-endian) it publishes its nonce into result.
276233

@@ -281,17 +238,6 @@ typedef struct __attribute__((aligned(8)))
281238
#define KH_MATRIX_ELEMS (KH_MATRIX_N * KH_MATRIX_N)
282239

283240

284-
inline void publishHit(__global Result* result, ulong const nonce)
285-
{
286-
uint const idx = atomic_inc(&result->count);
287-
result->found = 1;
288-
if (idx < MAX_RESULT)
289-
{
290-
result->nonces[idx] = nonce;
291-
}
292-
}
293-
294-
295241
#define KH_MATRIX_WORDS (KH_MATRIX_ELEMS / 4)
296242

297243

@@ -349,11 +295,11 @@ inline void khTheta(ulong* a)
349295
ulong const c2 = a[2] ^ a[7] ^ a[12] ^ a[17] ^ a[22];
350296
ulong const c3 = a[3] ^ a[8] ^ a[13] ^ a[18] ^ a[23];
351297
ulong const c4 = a[4] ^ a[9] ^ a[14] ^ a[19] ^ a[24];
352-
ulong const d0 = c4 ^ rotl64(c1, 1);
353-
ulong const d1 = c0 ^ rotl64(c2, 1);
354-
ulong const d2 = c1 ^ rotl64(c3, 1);
355-
ulong const d3 = c2 ^ rotl64(c4, 1);
356-
ulong const d4 = c3 ^ rotl64(c0, 1);
298+
ulong const d0 = c4 ^ rol_u64(c1, 1);
299+
ulong const d1 = c0 ^ rol_u64(c2, 1);
300+
ulong const d2 = c1 ^ rol_u64(c3, 1);
301+
ulong const d3 = c2 ^ rol_u64(c4, 1);
302+
ulong const d4 = c3 ^ rol_u64(c0, 1);
357303
a[0] ^= d0; a[5] ^= d0; a[10] ^= d0; a[15] ^= d0; a[20] ^= d0;
358304
a[1] ^= d1; a[6] ^= d1; a[11] ^= d1; a[16] ^= d1; a[21] ^= d1;
359305
a[2] ^= d2; a[7] ^= d2; a[12] ^= d2; a[17] ^= d2; a[22] ^= d2;
@@ -366,30 +312,30 @@ inline void khRhoPi(ulong* a)
366312
{
367313
ulong t = a[1];
368314
ulong tmp;
369-
tmp = a[10]; a[10] = rotl64(t, 1); t = tmp;
370-
tmp = a[7]; a[7] = rotl64(t, 3); t = tmp;
371-
tmp = a[11]; a[11] = rotl64(t, 6); t = tmp;
372-
tmp = a[17]; a[17] = rotl64(t, 10); t = tmp;
373-
tmp = a[18]; a[18] = rotl64(t, 15); t = tmp;
374-
tmp = a[3]; a[3] = rotl64(t, 21); t = tmp;
375-
tmp = a[5]; a[5] = rotl64(t, 28); t = tmp;
376-
tmp = a[16]; a[16] = rotl64(t, 36); t = tmp;
377-
tmp = a[8]; a[8] = rotl64(t, 45); t = tmp;
378-
tmp = a[21]; a[21] = rotl64(t, 55); t = tmp;
379-
tmp = a[24]; a[24] = rotl64(t, 2); t = tmp;
380-
tmp = a[4]; a[4] = rotl64(t, 14); t = tmp;
381-
tmp = a[15]; a[15] = rotl64(t, 27); t = tmp;
382-
tmp = a[23]; a[23] = rotl64(t, 41); t = tmp;
383-
tmp = a[19]; a[19] = rotl64(t, 56); t = tmp;
384-
tmp = a[13]; a[13] = rotl64(t, 8); t = tmp;
385-
tmp = a[12]; a[12] = rotl64(t, 25); t = tmp;
386-
tmp = a[2]; a[2] = rotl64(t, 43); t = tmp;
387-
tmp = a[20]; a[20] = rotl64(t, 62); t = tmp;
388-
tmp = a[14]; a[14] = rotl64(t, 18); t = tmp;
389-
tmp = a[22]; a[22] = rotl64(t, 39); t = tmp;
390-
tmp = a[9]; a[9] = rotl64(t, 61); t = tmp;
391-
tmp = a[6]; a[6] = rotl64(t, 20); t = tmp;
392-
tmp = a[1]; a[1] = rotl64(t, 44); t = tmp;
315+
tmp = a[10]; a[10] = rol_u64(t, 1); t = tmp;
316+
tmp = a[7]; a[7] = rol_u64(t, 3); t = tmp;
317+
tmp = a[11]; a[11] = rol_u64(t, 6); t = tmp;
318+
tmp = a[17]; a[17] = rol_u64(t, 10); t = tmp;
319+
tmp = a[18]; a[18] = rol_u64(t, 15); t = tmp;
320+
tmp = a[3]; a[3] = rol_u64(t, 21); t = tmp;
321+
tmp = a[5]; a[5] = rol_u64(t, 28); t = tmp;
322+
tmp = a[16]; a[16] = rol_u64(t, 36); t = tmp;
323+
tmp = a[8]; a[8] = rol_u64(t, 45); t = tmp;
324+
tmp = a[21]; a[21] = rol_u64(t, 55); t = tmp;
325+
tmp = a[24]; a[24] = rol_u64(t, 2); t = tmp;
326+
tmp = a[4]; a[4] = rol_u64(t, 14); t = tmp;
327+
tmp = a[15]; a[15] = rol_u64(t, 27); t = tmp;
328+
tmp = a[23]; a[23] = rol_u64(t, 41); t = tmp;
329+
tmp = a[19]; a[19] = rol_u64(t, 56); t = tmp;
330+
tmp = a[13]; a[13] = rol_u64(t, 8); t = tmp;
331+
tmp = a[12]; a[12] = rol_u64(t, 25); t = tmp;
332+
tmp = a[2]; a[2] = rol_u64(t, 43); t = tmp;
333+
tmp = a[20]; a[20] = rol_u64(t, 62); t = tmp;
334+
tmp = a[14]; a[14] = rol_u64(t, 18); t = tmp;
335+
tmp = a[22]; a[22] = rol_u64(t, 39); t = tmp;
336+
tmp = a[9]; a[9] = rol_u64(t, 61); t = tmp;
337+
tmp = a[6]; a[6] = rol_u64(t, 20); t = tmp;
338+
tmp = a[1]; a[1] = rol_u64(t, 44); t = tmp;
393339
}
394340

395341

@@ -458,7 +404,7 @@ __kernel void search(__global ushort const* matrix,
458404
}
459405
for (int w = 0; w < 4; ++w)
460406
{
461-
ms[w] ^= loadLe64(pre + w * 8);
407+
ms[w] ^= load_le_u64(pre + w * 8);
462408
}
463409
ms[4] ^= timestamp; // nonce (ms[9]) intentionally NOT folded in here
464410
khTheta(ms);
@@ -483,12 +429,12 @@ __kernel void search(__global ushort const* matrix,
483429
{
484430
st[i] = powMid[i];
485431
}
486-
ulong const nr = rotl64(nonce, 1);
432+
ulong const nr = rol_u64(nonce, 1);
487433
st[0] ^= nonce; st[5] ^= nonce; st[10] ^= nonce; st[15] ^= nonce; st[20] ^= nonce; st[9] ^= nonce;
488434
st[3] ^= nr; st[8] ^= nr; st[13] ^= nr; st[18] ^= nr; st[23] ^= nr;
489435
keccakF1600FromTheta1(st);
490436
uchar h1[32];
491-
storeLe256(st, h1);
437+
store_le_u256(st, h1);
492438

493439
uchar product[32];
494440
matmulDot(matU, h1, product);

‎sources/algo/kheavyhash/opencl/kheavyhash_result.cl‎

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,18 @@
1+
typedef struct __attribute__((aligned(8)))
2+
{
3+
uint found;
4+
uint count;
5+
ulong nonces[MAX_RESULT];
6+
} Result;
7+
8+
9+
inline
10+
void publishHit(__global Result* result, ulong const nonce)
11+
{
12+
uint const idx = atomic_inc(&result->count);
13+
result->found = 1;
14+
if (idx < MAX_RESULT)
15+
{
16+
result->nonces[idx] = nonce;
17+
}
18+
}

‎sources/algo/kheavyhash/opencl/tests/opencl_kat_test.cpp‎

Lines changed: 27 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -46,13 +46,11 @@ static std::vector<uint16_t> flatten(Matrix const& matrix)
4646
}
4747

4848

49-
// Resolve the shipped .cl: env override (fast iteration on the rig) first, then the
50-
// compile-time in-tree source path (POCL/dev), then the deployed path next to the
51-
// binary. First that opens wins, so the same unit_test runs in the dev harness and on
52-
// a real GPU. Returns "" if none open.
53-
static std::string resolveKernelPath()
49+
// Resolve a .cl by trying each candidate in order; first that opens wins, so the same
50+
// unit_test runs in the dev harness (in-tree source) and on a real GPU (deployed next
51+
// to the binary). Returns "" if none open.
52+
static std::string firstReadable(std::initializer_list<char const*> candidates)
5453
{
55-
char const* const candidates[3]{ std::getenv("KH_CL_PATH"), KH_CL_PATH, "kernel/kheavyhash/kheavyhash.cl" };
5654
for (char const* const cand : candidates)
5755
{
5856
if (nullptr != cand)
@@ -68,6 +66,20 @@ static std::string resolveKernelPath()
6866
}
6967

7068

69+
// The search kernel is split across shared helpers (rol_u64, load/store LE), the
70+
// per-algo Result struct, and the kheavyhash body — appended in this order so each
71+
// definition precedes its use, mirroring ResolverAmdKHeavyHash::buildSearch().
72+
static std::array<std::string, 4> resolveKernelPaths()
73+
{
74+
return {
75+
firstReadable({ KH_CL_COMMON_DIR "/rotate_byte.cl", "kernel/common/rotate_byte.cl" }),
76+
firstReadable({ KH_CL_COMMON_DIR "/load_store_le.cl", "kernel/common/load_store_le.cl" }),
77+
firstReadable({ KH_CL_DIR "/kheavyhash_result.cl", "kernel/kheavyhash/kheavyhash_result.cl" }),
78+
firstReadable({ std::getenv("KH_CL_PATH"), KH_CL_PATH, "kernel/kheavyhash/kheavyhash.cl" }),
79+
};
80+
}
81+
82+
7183
// One built program, shared by every test; each test pulls its kernel by name.
7284
class OpenClKat : public ::testing::Test
7385
{
@@ -93,14 +105,17 @@ class OpenClKat : public ::testing::Test
93105
context = cl::Context(device);
94106
queue = cl::CommandQueue(context, device);
95107

96-
std::string const clPath{ resolveKernelPath() };
97-
ASSERT_FALSE(clPath.empty()) << "cannot open kernel source (tried env KH_CL_PATH, " << KH_CL_PATH
98-
<< ", kernel/kheavyhash/kheavyhash.cl)";
108+
std::array<std::string, 4> const clPaths{ resolveKernelPaths() };
99109

100110
generator.clear();
101111
generator.setKernelName("search");
102112
generator.addDefine("MAX_RESULT", algo::kheavyhash::MAX_RESULT);
103-
ASSERT_TRUE(generator.appendFile(clPath)) << "cannot append kernel source: " << clPath;
113+
for (std::string const& clPath : clPaths)
114+
{
115+
ASSERT_FALSE(clPath.empty()) << "cannot open a kernel source (tried env KH_CL_PATH, the in-tree "
116+
"source dirs, and the deployed kernel/ tree)";
117+
ASSERT_TRUE(generator.appendFile(clPath)) << "cannot append kernel source: " << clPath;
118+
}
104119
ASSERT_TRUE(generator.build(&device, &context)) << "kernel build failed (see build log above)";
105120
}
106121
catch (cl::Error const& clErr)
@@ -233,10 +248,10 @@ TEST_F(OpenClKat, heavyHashMatchesReference)
233248
class SearchKernel : public OpenClKat, public ::testing::WithParamInterface<char const*>
234249
{
235250
protected:
236-
// Mirrors the kernel's Result struct (and algo::ethash::Result layout).
251+
// Mirrors the kernel's Result struct (algo::kheavyhash::Result layout).
237252
struct alignas(8) Result
238253
{
239-
uint8_t found{ 0u };
254+
bool found{ false };
240255
uint32_t count{ 0u };
241256
uint64_t nonces[algo::kheavyhash::MAX_RESULT]{ 0ull, 0ull, 0ull, 0ull };
242257
};

‎sources/resolver/amd/kheavyhash.cpp‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -126,7 +126,10 @@ bool resolver::ResolverAmdKHeavyHash::buildSearch()
126126
kernelGenerator.addDefine("MAX_RESULT", algo::kheavyhash::MAX_RESULT);
127127

128128
////////////////////////////////////////////////////////////////////////////
129-
if (false == kernelGenerator.appendFile("kernel/kheavyhash/kheavyhash.cl"))
129+
if (false == kernelGenerator.appendFile("kernel/common/rotate_byte.cl")
130+
|| false == kernelGenerator.appendFile("kernel/common/load_store_le.cl")
131+
|| false == kernelGenerator.appendFile("kernel/kheavyhash/kheavyhash_result.cl")
132+
|| false == kernelGenerator.appendFile("kernel/kheavyhash/kheavyhash.cl"))
130133
{
131134
return false;
132135
}

0 commit comments

Comments
 (0)