39ecd7c52b
LoadClass::hot (Option<HotClass { mb, k, added }>) beside mix, slots, scratch, mixer_mult, growth and era; V3_CLASS = { era: None, hot: None, ..LoadClass::MX4 } (hot stays None: the hot code is behind the flag, measured and not adopted). Op::Hot, the hot slots drawn after the scratch slots, no width roll (v2_loads allows the added form's extra slots), the id suffix hot/<S><k>[added], the hot parsing inside parse_loads, the era branch first in name(). HotTable under seed_words("igneum-hot/" || epoch seed) with the cache chain and tag umHT, read at H[mulhi(src, HOT_WORDS)]; DatasetSource::hot attached by new_class_day and from_seed_bytes_class; the acceptance stand-in dataset_elem(idx, S[2], S[3]); the three emitters (hot argument after the init words, ht_segment and igneum_hot_fill beside the layout-aware cores); packfile.h hot fields beside class, era, attempt and mixerMult; OpenCL host, Metal packbench and NVRTC worker fill H on the device and self-test it. Eight packs under proto-cuda/packs-ca2-hot (replaced hot32k4 hot64k4 hot96k4 hot64k2 hot64k8, added hot32k4a hot64k4a hot96k4a) re-exported on the merged crate: vectors unchanged, program.h and program.json carry the mixer fields. docs/plans/hot-table.md (design, spec text, per-tier budget, chip model, Mac and PC measurements, the decision: layer 5 out of v3, the 3.0 note); bench-log entry and addenda with job ids and worker sha256s; the two PC playbooks.
Checks on this commit: cargo test --release 53 + 19 pass (the pinned v2, mx4, era, readwidth and hot packs); the pinned packs under proto-cuda/packs, packs-ca2-mixer, packs-ca2-era and packs-readwidth untouched; Metal packbench and Apple OpenCL --bench-pack on all eight hot packs (run lock, 2^20 at base 0): 96/96 lanes, hot table head, last line and FNV PASS, one fingerprint per pack on both harnesses, equal to the fingerprints before the rebase (hot32k4 679e5e83378d3790, hot64k4 d4c9e456b039fdef, hot96k4 7c98eceffee9fd73, hot64k2 f43b10a95879b8e5, hot64k8 c11309d743be9392, hot32k4a afb700b2d997c847, hot64k4a ba214baa9c1a9e85, hot96k4a 29e1916aed6deff5). No era or mixer behaviour changed: every resolution kept the ca2-v3 side and appended the hot branch.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
2205 lines
133 KiB
C
2205 lines
133 KiB
C
// igneum-bench-cl: OpenCL host program for Igneum's random-program proof-of-work test harness.
|
|
// The portable third path after Apple Metal (proto-metal) and NVIDIA CUDA (proto-cuda): it runs on AMD (Windows
|
|
// and Linux), NVIDIA, Intel and, as a correctness check only, on Apple's deprecated OpenCL 1.2 runtime.
|
|
//
|
|
// TEST HARNESS ONLY. No pool, no network, no wallet, no mining protocol. It fills the dataset on the device,
|
|
// checks the device against vectors produced on the Mac (proto-metal), and times the kernel.
|
|
//
|
|
// C99 plus the OpenCL 1.2 API, nothing else. The kernels are compiled from packs/<seed>/kernel.cl at runtime.
|
|
// The pack's program.h, vectors.h and (memory-hard packs) memhard.h are included at compile time; memhard.h is the
|
|
// host reference that fills the cache on one thread and derives dataset words for the self-test.
|
|
//
|
|
// Build: see README.md (macOS -framework OpenCL, Linux -lOpenCL, Windows cl.exe + OpenCL.lib), or build.sh / build.bat.
|
|
|
|
#define _CRT_SECURE_NO_WARNINGS
|
|
#define CL_TARGET_OPENCL_VERSION 120
|
|
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
|
|
#if defined(__APPLE__) && !defined(IGNEUM_KHR_HEADERS)
|
|
#include <OpenCL/opencl.h>
|
|
#else
|
|
#include <CL/cl.h>
|
|
#endif
|
|
#ifdef IGNEUM_CL_DYNAMIC
|
|
#include "cl_dynamic.h" /* Windows one-click build: OpenCL.dll loaded at run time, no import library */
|
|
#endif
|
|
|
|
#include <stdint.h>
|
|
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
#include <string.h>
|
|
#ifdef _WIN32
|
|
#define WIN32_LEAN_AND_MEAN
|
|
#include <windows.h>
|
|
#define strtok_r strtok_s
|
|
#else
|
|
#include <time.h>
|
|
#include <dlfcn.h>
|
|
#include <pthread.h>
|
|
#endif
|
|
|
|
#define IGNEUM_NO_CUDA
|
|
#include "program.h"
|
|
#include "vectors.h"
|
|
#include "../proto-cuda/nvrtc/packfile.h" /* --pack: a pack read at run time (generic serve mode, 4 October 2026) */
|
|
|
|
#ifndef IGNEUM_DATASET_MODE
|
|
#define IGNEUM_DATASET_MODE 0
|
|
#endif
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
#include "memhard.h"
|
|
#endif
|
|
#ifndef IGNEUM_KERNEL_PATH
|
|
#define IGNEUM_KERNEL_PATH "kernel.cl"
|
|
#endif
|
|
|
|
// Sub-group query constants (cl_khr_subgroups / OpenCL 2.1). Spelled out because OpenCL 1.2 headers lack them.
|
|
#define IG_CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE 0x2033
|
|
#define IG_CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE 0x2034
|
|
// Vendor device attributes (cl_amd_device_attribute_query, cl_nv_device_attribute_query).
|
|
#define IG_CL_DEVICE_WAVEFRONT_WIDTH_AMD 0x4043
|
|
#define IG_CL_DEVICE_WARP_SIZE_NV 0x4003
|
|
|
|
typedef cl_int (CL_API_CALL *ig_pfn_subgroup_info)(cl_kernel, cl_device_id, cl_uint, size_t, const void*, size_t, void*, size_t*);
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Errors and timing
|
|
|
|
static const char* clErrName(cl_int e) {
|
|
switch (e) {
|
|
case CL_SUCCESS: return "CL_SUCCESS";
|
|
case CL_DEVICE_NOT_FOUND: return "CL_DEVICE_NOT_FOUND";
|
|
case CL_DEVICE_NOT_AVAILABLE: return "CL_DEVICE_NOT_AVAILABLE";
|
|
case CL_COMPILER_NOT_AVAILABLE: return "CL_COMPILER_NOT_AVAILABLE";
|
|
case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
|
|
case CL_OUT_OF_RESOURCES: return "CL_OUT_OF_RESOURCES";
|
|
case CL_OUT_OF_HOST_MEMORY: return "CL_OUT_OF_HOST_MEMORY";
|
|
case CL_PROFILING_INFO_NOT_AVAILABLE: return "CL_PROFILING_INFO_NOT_AVAILABLE";
|
|
case CL_BUILD_PROGRAM_FAILURE: return "CL_BUILD_PROGRAM_FAILURE";
|
|
case CL_INVALID_VALUE: return "CL_INVALID_VALUE";
|
|
case CL_INVALID_DEVICE: return "CL_INVALID_DEVICE";
|
|
case CL_INVALID_CONTEXT: return "CL_INVALID_CONTEXT";
|
|
case CL_INVALID_QUEUE_PROPERTIES: return "CL_INVALID_QUEUE_PROPERTIES";
|
|
case CL_INVALID_COMMAND_QUEUE: return "CL_INVALID_COMMAND_QUEUE";
|
|
case CL_INVALID_MEM_OBJECT: return "CL_INVALID_MEM_OBJECT";
|
|
case CL_INVALID_BUFFER_SIZE: return "CL_INVALID_BUFFER_SIZE";
|
|
case CL_INVALID_BUILD_OPTIONS: return "CL_INVALID_BUILD_OPTIONS";
|
|
case CL_INVALID_PROGRAM: return "CL_INVALID_PROGRAM";
|
|
case CL_INVALID_PROGRAM_EXECUTABLE: return "CL_INVALID_PROGRAM_EXECUTABLE";
|
|
case CL_INVALID_KERNEL_NAME: return "CL_INVALID_KERNEL_NAME";
|
|
case CL_INVALID_KERNEL: return "CL_INVALID_KERNEL";
|
|
case CL_INVALID_ARG_INDEX: return "CL_INVALID_ARG_INDEX";
|
|
case CL_INVALID_ARG_VALUE: return "CL_INVALID_ARG_VALUE";
|
|
case CL_INVALID_ARG_SIZE: return "CL_INVALID_ARG_SIZE";
|
|
case CL_INVALID_KERNEL_ARGS: return "CL_INVALID_KERNEL_ARGS";
|
|
case CL_INVALID_WORK_DIMENSION: return "CL_INVALID_WORK_DIMENSION";
|
|
case CL_INVALID_WORK_GROUP_SIZE: return "CL_INVALID_WORK_GROUP_SIZE";
|
|
case CL_INVALID_WORK_ITEM_SIZE: return "CL_INVALID_WORK_ITEM_SIZE";
|
|
case CL_INVALID_GLOBAL_OFFSET: return "CL_INVALID_GLOBAL_OFFSET";
|
|
case CL_INVALID_EVENT: return "CL_INVALID_EVENT";
|
|
case CL_INVALID_OPERATION: return "CL_INVALID_OPERATION";
|
|
case CL_INVALID_GLOBAL_WORK_SIZE: return "CL_INVALID_GLOBAL_WORK_SIZE";
|
|
case CL_INVALID_PLATFORM: return "CL_INVALID_PLATFORM";
|
|
default: return "(other)";
|
|
}
|
|
}
|
|
|
|
static void clFail(cl_int e, const char* what, int line) {
|
|
fprintf(stderr, "OpenCL error: %s (%d)\n at host.c:%d\n in %s\n", clErrName(e), (int)e, line, what);
|
|
exit(2);
|
|
}
|
|
#define CL_CHECK(call) do { cl_int err_ = (call); if (err_ != CL_SUCCESS) clFail(err_, #call, __LINE__); } while (0)
|
|
#define CL_CHECK_ERR(err_, what) do { if ((err_) != CL_SUCCESS) clFail((err_), what, __LINE__); } while (0)
|
|
|
|
static double wallMs(void) {
|
|
#ifdef _WIN32
|
|
LARGE_INTEGER f, c;
|
|
QueryPerformanceFrequency(&f);
|
|
QueryPerformanceCounter(&c);
|
|
return (double)c.QuadPart * 1000.0 / (double)f.QuadPart;
|
|
#else
|
|
struct timespec ts;
|
|
clock_gettime(CLOCK_MONOTONIC, &ts);
|
|
return (double)ts.tv_sec * 1000.0 + (double)ts.tv_nsec / 1e6;
|
|
#endif
|
|
}
|
|
|
|
// Event profiling. A runtime that cannot report timestamps (CL_PROFILING_INFO_NOT_AVAILABLE) gives -1 and the
|
|
// harness switches the rate to wall time instead of stopping; the count of such events is reported.
|
|
static int gProfilingFailures = 0;
|
|
static double eventMs(cl_event e) {
|
|
cl_ulong t0 = 0, t1 = 0;
|
|
if (clGetEventProfilingInfo(e, CL_PROFILING_COMMAND_START, sizeof(t0), &t0, NULL) != CL_SUCCESS ||
|
|
clGetEventProfilingInfo(e, CL_PROFILING_COMMAND_END, sizeof(t1), &t1, NULL) != CL_SUCCESS) { ++gProfilingFailures; return -1.0; }
|
|
return (double)(t1 - t0) / 1e6;
|
|
}
|
|
|
|
static double spanMs(cl_event first, cl_event last) {
|
|
cl_ulong t0 = 0, t1 = 0;
|
|
if (clGetEventProfilingInfo(first, CL_PROFILING_COMMAND_START, sizeof(t0), &t0, NULL) != CL_SUCCESS ||
|
|
clGetEventProfilingInfo(last, CL_PROFILING_COMMAND_END, sizeof(t1), &t1, NULL) != CL_SUCCESS) { ++gProfilingFailures; return -1.0; }
|
|
return (double)(t1 - t0) / 1e6;
|
|
}
|
|
|
|
static void* loadSym(const char* name) {
|
|
#ifdef _WIN32
|
|
HMODULE m = GetModuleHandleA("OpenCL.dll");
|
|
return m ? (void*)GetProcAddress(m, name) : NULL;
|
|
#else
|
|
return dlsym(RTLD_DEFAULT, name);
|
|
#endif
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Host reference
|
|
|
|
static const uint32_t SEEDW[8] = IGNEUM_SEEDW_INIT;
|
|
|
|
#if IGNEUM_DATASET_MODE == 0
|
|
// Same closed form as ds_elem in kernel.cl and datasetElem in proto-metal/main.swift.
|
|
static uint32_t host_ds_elem(uint32_t i, uint32_t d0, uint32_t d1) {
|
|
uint32_t x = i ^ d0;
|
|
x *= 0x9E3779B1u; x ^= x >> 15;
|
|
x += d1;
|
|
x *= 0x85EBCA77u; x ^= x >> 13;
|
|
x *= 0xC2B2AE3Du; x ^= x >> 16;
|
|
return x;
|
|
}
|
|
#else
|
|
static uint64_t fnv1a64(const void* p, size_t n) {
|
|
const uint8_t* b = (const uint8_t*)p;
|
|
uint64_t h = 0xcbf29ce484222325ull;
|
|
size_t i;
|
|
for (i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
|
|
return h;
|
|
}
|
|
static const uint32_t CACHE_WORDS_HOST = 1u << IGNEUM_CACHE_LOG2_WORDS;
|
|
static uint32_t* hCache = NULL;
|
|
/* dataset[w] through the pack's own mh_word (memhard.h), which carries the pack's item-to-word layout (era layout,
|
|
* 5 October 2026; the former w >> 4 / w & 15 here failed the random points of every interleaved pack). */
|
|
static uint32_t host_ds_word(uint32_t w) {
|
|
return mh_word(hCache, w);
|
|
}
|
|
#endif
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Options
|
|
|
|
typedef struct {
|
|
int datasetMib;
|
|
int batchLog2;
|
|
int batches;
|
|
int groupWarps;
|
|
int sweep;
|
|
int device; // flat index into the enumerated list, -1 = first GPU
|
|
int exchange; // 0 auto, 1 force local-memory fallback, 2 force sub-group shuffles
|
|
int list;
|
|
int timeWall; // 1 = rate from wall time, 0 = from device event profiling, -1 = auto (wall on the Apple platform)
|
|
const char* kernelPath;
|
|
const char* extraOpts;
|
|
int serve; // --serve: GPU worker for igneum-miner --worker (jobs on stdin), 3 October 2026
|
|
int noPrepare; // --no-prepare: serve without the prepare command (ready line says "prepare 0"), to test the miner's fallback
|
|
int kernelGiven; // --kernel was passed
|
|
const char* vendor; // --vendor S: pick the first GPU whose vendor string contains S (default: first GPU of any vendor)
|
|
const char* packDir; // --pack D (serve only): generic mode, the pack is read from D at run time; the compiled-in pack is
|
|
// then only the build-time placeholder of the prebuilt exe (4 October 2026)
|
|
int readback; // --readback: 0 select (a GPU-side pass reads back only the hits and 34 sentinel words), 1 full
|
|
// (every output word comes back, 8 bytes per nonce, the path before 5 October 2026)
|
|
int memprobe; // --memprobe: dependent-load latency and throughput, independent-load throughput and an ALU
|
|
// chain on the chosen device, no pack needed (5 October 2026, the 9070 XT on the eGPU)
|
|
int probeMib; // --probe-mib N: --memprobe at that one buffer size only (default 0 = 4, 64 and 1024 MiB)
|
|
int benchPack; // --bench-pack: with --pack D, build and self-test the pack at run time (as --serve does) and time
|
|
// igneum_hash_bound with the pack's seed words as init words; read-width experiment, 5 October 2026
|
|
int warps; // --warps N: persistent warps for a variant-5 pack (IGNEUM_PERSISTENT_WARPS); 0 = 2048
|
|
} Options;
|
|
|
|
static int packMib(void) { return (int)(((1ull << IGNEUM_DATASET_LOG2) * 4ull) >> 20); }
|
|
|
|
static void usage(void) {
|
|
printf(
|
|
"igneum-bench-cl [--list] [--device D] [--dataset-mib N] [--sweep] [--batch-log2 24] [--batches 5] [--group-warps 1]\n"
|
|
" [--exchange auto|local|subgroup] [--kernel path/to/kernel.cl] [--build-opts \"...\"]\n"
|
|
" --list print every OpenCL platform and device, then exit\n"
|
|
" --device D device index from the list (default: the first GPU, else device 0)\n"
|
|
" --dataset-mib N dataset size in MiB, power of two (default 1024; vectors are only checked at %d MiB)\n"
|
|
" --sweep run 4, 64, 256, 512 and 1024 MiB in sequence (same sweep as the Mac and the CUDA harness)\n"
|
|
" --batch-log2 B nonces per batch = 2^B (default 24)\n"
|
|
" --batches N timed batches after one warm-up batch (default 5)\n"
|
|
" --group-warps W 32-lane units per work-group, 1..8 (default 1 = one work-group per unit; sub-group shuffles need 1)\n"
|
|
" --exchange M auto (default): sub-group shuffles when the device has them and its sub-group size is 32, else local memory\n"
|
|
" local: force the local-memory exchange; subgroup: require sub-group shuffles or fail\n"
|
|
" --kernel P path to the pack's kernel.cl (default: the path compiled in, %s)\n"
|
|
" --build-opts S extra options appended to clBuildProgram (for example \"-cl-std=CL2.0\")\n"
|
|
" --time T event (default): hashes/s from device event profiling, like cudaEvent time; wall: from host wall time.\n"
|
|
" Apple's OpenCL runtime reports unusable event timestamps, so wall is the default on the Apple platform.\n"
|
|
" --vendor S choose the first GPU whose vendor string contains S (for example \"Advanced Micro Devices\"); fails if none\n"
|
|
" --serve GPU worker for igneum-miner --worker: reads \"job ...\" lines on stdin, prints found/done lines.\n"
|
|
" --no-prepare with --serve: no prepare support (the miner then falls back to exit 42 at a seed change).\n"
|
|
" Builds the pack's kernel_bound.cl (next to the compiled-in kernel.cl) unless --kernel says otherwise.\n"
|
|
" --pack D with --serve: serve the pack in directory D (program.h, seeds.txt, vectors.h, kernel_bound.cl), whatever\n"
|
|
" pack this exe was built against; it is self-tested against its vectors.h first (the one-click worker)\n"
|
|
" --readback M with --serve: select (default) reads back only the hits and 34 sentinel words of each dispatch through a\n"
|
|
" GPU-side pass; full reads back every output (8 bytes per nonce). IGNEUM_READBACK=full does the same.\n"
|
|
" --bench-pack with --pack D: read the pack at run time, build and self-test it, time its bound kernel (one exe, any pack)\n"
|
|
" --warps N persistent warps for a variant-5 pack (a 1 MiB scratch each; default 2048; the batch rounds to 32 x N)\n"
|
|
" --memprobe no pack: dependent random loads (latency and throughput against lanes in flight), independent random\n"
|
|
" loads and an ALU chain on the chosen device, at 4, 64 and 1024 MiB (--probe-mib N for one size)\n", packMib(), IGNEUM_KERNEL_PATH);
|
|
}
|
|
|
|
static int isPow2(long long v) { return v > 0 && (v & (v - 1)) == 0; }
|
|
static int log2u32(uint32_t v) { int n = 0; while (v > 1u) { v >>= 1; ++n; } return n; }
|
|
|
|
static Options parseArgs(int argc, char** argv) {
|
|
Options o;
|
|
int i;
|
|
o.datasetMib = 1024; o.batchLog2 = 24; o.batches = 5; o.groupWarps = 1; o.sweep = 0; o.device = -1;
|
|
o.exchange = 0; o.list = 0; o.timeWall = -1; o.kernelPath = IGNEUM_KERNEL_PATH; o.extraOpts = ""; o.serve = 0; o.noPrepare = 0; o.kernelGiven = 0; o.vendor = NULL; o.packDir = NULL;
|
|
o.readback = (getenv("IGNEUM_READBACK") && strcmp(getenv("IGNEUM_READBACK"), "full") == 0) ? 1 : 0; o.memprobe = 0; o.probeMib = 0; o.benchPack = 0; o.warps = 0;
|
|
for (i = 1; i < argc; ++i) {
|
|
const char* a = argv[i];
|
|
int needs = (strcmp(a, "--dataset-mib") == 0 || strcmp(a, "--batch-log2") == 0 || strcmp(a, "--batches") == 0 ||
|
|
strcmp(a, "--group-warps") == 0 || strcmp(a, "--device") == 0 || strcmp(a, "--exchange") == 0 ||
|
|
strcmp(a, "--kernel") == 0 || strcmp(a, "--build-opts") == 0 || strcmp(a, "--time") == 0);
|
|
if (needs && i + 1 >= argc) { usage(); exit(2); }
|
|
if (strcmp(a, "--dataset-mib") == 0) o.datasetMib = atoi(argv[++i]);
|
|
else if (strcmp(a, "--batch-log2") == 0) o.batchLog2 = atoi(argv[++i]);
|
|
else if (strcmp(a, "--batches") == 0) o.batches = atoi(argv[++i]);
|
|
else if (strcmp(a, "--group-warps") == 0) o.groupWarps = atoi(argv[++i]);
|
|
else if (strcmp(a, "--device") == 0) o.device = atoi(argv[++i]);
|
|
else if (strcmp(a, "--kernel") == 0) { o.kernelPath = argv[++i]; o.kernelGiven = 1; }
|
|
else if (strcmp(a, "--serve") == 0) o.serve = 1;
|
|
else if (strcmp(a, "--no-prepare") == 0) o.noPrepare = 1;
|
|
else if (strcmp(a, "--vendor") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.vendor = argv[++i]; }
|
|
else if (strcmp(a, "--pack") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.packDir = argv[++i]; }
|
|
else if (strcmp(a, "--readback") == 0) {
|
|
const char* m;
|
|
if (i + 1 >= argc) { usage(); exit(2); }
|
|
m = argv[++i];
|
|
if (strcmp(m, "select") == 0) o.readback = 0;
|
|
else if (strcmp(m, "full") == 0) o.readback = 1;
|
|
else { printf("--readback must be select or full\n"); exit(2); }
|
|
}
|
|
else if (strcmp(a, "--memprobe") == 0) o.memprobe = 1;
|
|
else if (strcmp(a, "--bench-pack") == 0) o.benchPack = 1;
|
|
else if (strcmp(a, "--warps") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.warps = atoi(argv[++i]); }
|
|
else if (strcmp(a, "--probe-mib") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.probeMib = atoi(argv[++i]); }
|
|
else if (strcmp(a, "--build-opts") == 0) o.extraOpts = argv[++i];
|
|
else if (strcmp(a, "--time") == 0) {
|
|
const char* m = argv[++i];
|
|
if (strcmp(m, "event") == 0) o.timeWall = 0;
|
|
else if (strcmp(m, "wall") == 0) o.timeWall = 1;
|
|
else { printf("--time must be event or wall\n"); exit(2); }
|
|
}
|
|
else if (strcmp(a, "--exchange") == 0) {
|
|
const char* m = argv[++i];
|
|
if (strcmp(m, "auto") == 0) o.exchange = 0;
|
|
else if (strcmp(m, "local") == 0) o.exchange = 1;
|
|
else if (strcmp(m, "subgroup") == 0) o.exchange = 2;
|
|
else { printf("--exchange must be auto, local or subgroup\n"); exit(2); }
|
|
}
|
|
else if (strcmp(a, "--sweep") == 0) o.sweep = 1;
|
|
else if (strcmp(a, "--list") == 0) o.list = 1;
|
|
else if (strcmp(a, "-h") == 0 || strcmp(a, "--help") == 0) { usage(); exit(0); }
|
|
else { printf("unknown argument %s\n", a); usage(); exit(2); }
|
|
}
|
|
if (!isPow2(o.datasetMib) || o.datasetMib < 1 || o.datasetMib > 16384) { printf("--dataset-mib must be a power of two between 1 and 16384\n"); exit(2); }
|
|
if (o.batchLog2 < 10 || o.batchLog2 > 28) { printf("--batch-log2 must be between 10 and 28\n"); exit(2); }
|
|
if (o.batches < 1) { printf("--batches must be at least 1\n"); exit(2); }
|
|
if (o.groupWarps < 1 || o.groupWarps > 8) { printf("--group-warps must be between 1 and 8\n"); exit(2); }
|
|
return o;
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Devices
|
|
|
|
typedef struct {
|
|
cl_platform_id platform;
|
|
cl_device_id device;
|
|
char platformName[256], platformVersion[256];
|
|
char name[256], vendor[256], version[256], driver[256], cVersion[256];
|
|
char* extensions;
|
|
cl_device_type type;
|
|
cl_uint computeUnits, clockMHz;
|
|
cl_ulong globalMem, maxAlloc, localMem;
|
|
size_t maxWorkGroup;
|
|
int cMajor, cMinor; // OpenCL C version
|
|
int dMajor, dMinor; // device (platform profile) version
|
|
cl_uint amdWavefront, nvWarp; // 0 if not reported
|
|
int dupOf; // index of the same card on a newer platform of the same vendor, -1 if none (5 October 2026)
|
|
} DeviceInfo;
|
|
|
|
/* The driver version as a number for ordering ("3683.0 (PAL,LC)" -> 3683.0; "617.14" -> 617.14; 0 when unreadable). */
|
|
static double driverNumber(const char* driver) {
|
|
const char* p = driver;
|
|
while (*p && (*p < '0' || *p > '9')) ++p;
|
|
return *p ? strtod(p, NULL) : 0.0;
|
|
}
|
|
|
|
/* Two AMD platforms are registered after a driver update on Windows (PC 1, 5 October 2026: 32.0.21042 and 32.0.32015,
|
|
* OpenCL driver strings 3652.0 and 3683.0): every card is listed twice, the app ran two workers on one 9070 XT, and
|
|
* each got half. The same card on the same-named platform with a different platform version is the one card; the
|
|
* entry whose driver number is lower is the duplicate. Two real cards of one model sit on the SAME platform and are
|
|
* never folded. Indices stay flat (the app passes them back as --device). Returns how many duplicates were marked. */
|
|
static int markDuplicates(DeviceInfo* list, int n) {
|
|
int i, j, marked = 0;
|
|
for (i = 0; i < n; ++i) list[i].dupOf = -1;
|
|
for (i = 0; i < n; ++i) {
|
|
if (list[i].dupOf >= 0) continue;
|
|
for (j = i + 1; j < n; ++j) {
|
|
int older;
|
|
if (list[j].dupOf >= 0) continue;
|
|
if (strcmp(list[i].platformName, list[j].platformName) != 0) continue;
|
|
if (strcmp(list[i].platformVersion, list[j].platformVersion) == 0) continue;
|
|
if (strcmp(list[i].name, list[j].name) != 0 || strcmp(list[i].vendor, list[j].vendor) != 0) continue;
|
|
if (list[i].globalMem != list[j].globalMem || list[i].computeUnits != list[j].computeUnits) continue;
|
|
older = driverNumber(list[j].driver) < driverNumber(list[i].driver) ? j : i;
|
|
if (older == j) { list[j].dupOf = i; }
|
|
else { list[i].dupOf = j; }
|
|
++marked;
|
|
if (older == i) break; /* i itself is the duplicate; j stays the real one */
|
|
}
|
|
}
|
|
/* three registrations of one card: every duplicate points at the one that stays, not at another duplicate */
|
|
for (i = 0; i < n; ++i) {
|
|
int k = list[i].dupOf, hops = 0;
|
|
while (k >= 0 && list[k].dupOf >= 0 && hops++ < n) k = list[k].dupOf;
|
|
if (list[i].dupOf >= 0) list[i].dupOf = k;
|
|
}
|
|
return marked;
|
|
}
|
|
|
|
static void devStr(cl_device_id d, cl_device_info what, char* out, size_t n) {
|
|
out[0] = 0;
|
|
clGetDeviceInfo(d, what, n - 1, out, NULL);
|
|
out[n - 1] = 0;
|
|
}
|
|
|
|
static const char* typeName(cl_device_type t) {
|
|
if (t & CL_DEVICE_TYPE_GPU) return "GPU";
|
|
if (t & CL_DEVICE_TYPE_CPU) return "CPU";
|
|
if (t & CL_DEVICE_TYPE_ACCELERATOR) return "accelerator";
|
|
return "other";
|
|
}
|
|
|
|
static int enumerateDevices(DeviceInfo** outList) {
|
|
cl_uint np = 0, p;
|
|
cl_platform_id plats[16];
|
|
DeviceInfo* list = NULL;
|
|
int n = 0;
|
|
cl_int e = clGetPlatformIDs(16, plats, &np);
|
|
if (e != CL_SUCCESS || np == 0) { *outList = NULL; return 0; }
|
|
for (p = 0; p < np; ++p) {
|
|
cl_uint nd = 0, d;
|
|
cl_device_id devs[32];
|
|
char pname[256] = {0}, pver[256] = {0};
|
|
clGetPlatformInfo(plats[p], CL_PLATFORM_NAME, sizeof(pname) - 1, pname, NULL);
|
|
clGetPlatformInfo(plats[p], CL_PLATFORM_VERSION, sizeof(pver) - 1, pver, NULL);
|
|
if (clGetDeviceIDs(plats[p], CL_DEVICE_TYPE_ALL, 32, devs, &nd) != CL_SUCCESS) continue;
|
|
for (d = 0; d < nd; ++d) {
|
|
DeviceInfo di;
|
|
size_t extLen = 0;
|
|
memset(&di, 0, sizeof(di));
|
|
di.platform = plats[p]; di.device = devs[d];
|
|
strncpy(di.platformName, pname, 255); strncpy(di.platformVersion, pver, 255);
|
|
devStr(devs[d], CL_DEVICE_NAME, di.name, sizeof(di.name));
|
|
devStr(devs[d], CL_DEVICE_VENDOR, di.vendor, sizeof(di.vendor));
|
|
devStr(devs[d], CL_DEVICE_VERSION, di.version, sizeof(di.version));
|
|
devStr(devs[d], CL_DRIVER_VERSION, di.driver, sizeof(di.driver));
|
|
devStr(devs[d], CL_DEVICE_OPENCL_C_VERSION, di.cVersion, sizeof(di.cVersion));
|
|
clGetDeviceInfo(devs[d], CL_DEVICE_EXTENSIONS, 0, NULL, &extLen);
|
|
di.extensions = (char*)calloc(extLen + 1, 1);
|
|
if (extLen) clGetDeviceInfo(devs[d], CL_DEVICE_EXTENSIONS, extLen, di.extensions, NULL);
|
|
clGetDeviceInfo(devs[d], CL_DEVICE_TYPE, sizeof(di.type), &di.type, NULL);
|
|
clGetDeviceInfo(devs[d], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(di.computeUnits), &di.computeUnits, NULL);
|
|
clGetDeviceInfo(devs[d], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(di.clockMHz), &di.clockMHz, NULL);
|
|
clGetDeviceInfo(devs[d], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(di.globalMem), &di.globalMem, NULL);
|
|
clGetDeviceInfo(devs[d], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(di.maxAlloc), &di.maxAlloc, NULL);
|
|
clGetDeviceInfo(devs[d], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(di.localMem), &di.localMem, NULL);
|
|
clGetDeviceInfo(devs[d], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(di.maxWorkGroup), &di.maxWorkGroup, NULL);
|
|
if (sscanf(di.cVersion, "OpenCL C %d.%d", &di.cMajor, &di.cMinor) != 2) { di.cMajor = 1; di.cMinor = 2; }
|
|
if (sscanf(di.version, "OpenCL %d.%d", &di.dMajor, &di.dMinor) != 2) { di.dMajor = 1; di.dMinor = 2; }
|
|
if (strstr(di.extensions, "cl_amd_device_attribute_query"))
|
|
clGetDeviceInfo(devs[d], IG_CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof(di.amdWavefront), &di.amdWavefront, NULL);
|
|
if (strstr(di.extensions, "cl_nv_device_attribute_query"))
|
|
clGetDeviceInfo(devs[d], IG_CL_DEVICE_WARP_SIZE_NV, sizeof(di.nvWarp), &di.nvWarp, NULL);
|
|
list = (DeviceInfo*)realloc(list, sizeof(DeviceInfo) * (size_t)(n + 1));
|
|
list[n++] = di;
|
|
}
|
|
}
|
|
if (n) markDuplicates(list, n);
|
|
*outList = list;
|
|
return n;
|
|
}
|
|
|
|
static void printDevice(int idx, const DeviceInfo* d, int chosen) {
|
|
const char* subExt = strstr(d->extensions, "cl_khr_subgroup_shuffle") ? "cl_khr_subgroup_shuffle" :
|
|
strstr(d->extensions, "cl_intel_subgroups") ? "cl_intel_subgroups" :
|
|
strstr(d->extensions, "cl_khr_subgroups") ? "cl_khr_subgroups (no shuffle extension)" : "none";
|
|
if (d->dupOf >= 0) {
|
|
/* Hidden from the app's card list: its parser takes only lines that start with "[" (detect.rs). The index
|
|
* is still valid for --device, so a run on the older platform stays possible for a comparison. */
|
|
printf(" dup [%d] %s | %s (%s): the same card as [%d] on an older platform (driver %s); hidden, use [%d]\n",
|
|
idx, d->name, d->platformName, d->platformVersion, d->dupOf, d->driver, d->dupOf);
|
|
return;
|
|
}
|
|
printf("%s[%d] %s | %s (%s)\n", chosen ? "*" : " ", idx, d->name, d->platformName, d->platformVersion);
|
|
printf(" %s, vendor %s, driver %s, %s, %u compute units, %u MHz\n", typeName(d->type), d->vendor, d->driver, d->cVersion, d->computeUnits, d->clockMHz);
|
|
printf(" global %llu MiB, max alloc %llu MiB, local %llu KiB, max work-group %llu, sub-group extension: %s",
|
|
(unsigned long long)(d->globalMem >> 20), (unsigned long long)(d->maxAlloc >> 20), (unsigned long long)(d->localMem >> 10),
|
|
(unsigned long long)d->maxWorkGroup, subExt);
|
|
if (d->amdWavefront) printf(", AMD wavefront width %u", d->amdWavefront);
|
|
if (d->nvWarp) printf(", NVIDIA warp size %u", d->nvWarp);
|
|
printf("\n");
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Program build
|
|
|
|
static char* readFile(const char* path, size_t* len) {
|
|
FILE* f = fopen(path, "rb");
|
|
char* buf;
|
|
long n;
|
|
if (!f) return NULL;
|
|
fseek(f, 0, SEEK_END); n = ftell(f); fseek(f, 0, SEEK_SET);
|
|
if (n < 0) { fclose(f); return NULL; }
|
|
buf = (char*)malloc((size_t)n + 1);
|
|
if (fread(buf, 1, (size_t)n, f) != (size_t)n) { fclose(f); free(buf); return NULL; }
|
|
buf[n] = 0;
|
|
fclose(f);
|
|
*len = (size_t)n;
|
|
return buf;
|
|
}
|
|
|
|
typedef struct {
|
|
cl_context ctx;
|
|
cl_command_queue q;
|
|
cl_program prog;
|
|
cl_kernel kHash, kCacheFill, kBuild, kFill;
|
|
cl_kernel kHashBound; // igneum_hash_bound (serve mode; NULL when the source has none)
|
|
int exchange; // 0 local memory, 1 khr sub-group shuffle, 2 intel
|
|
size_t subGroupSize; // as queried for a 32-item work-group, 0 if not queried
|
|
char exchangeNote[512];
|
|
char buildOptions[512];
|
|
int groupSize; // work-group size the program was built for (IGNEUM_GROUP = 32 x group-warps)
|
|
} Device;
|
|
|
|
static const char* exchangeName(int m) { return m == 1 ? "sub_group_shuffle_xor (cl_khr_subgroup_shuffle)" : m == 2 ? "intel_sub_group_shuffle_xor (cl_intel_subgroups)" : "local-memory exchange with barrier"; }
|
|
|
|
// Returns 0 on success, 1 on build failure (log printed).
|
|
static int buildProgram(Device* dv, const DeviceInfo* di, const char* src, size_t srcLen, int exchangeMode, int groupSize, const char* extra) {
|
|
cl_int err = 0;
|
|
const char* std;
|
|
// The sub-group built-ins need OpenCL C 2.0 or 3.0. OpenCL 3.0 devices may report "OpenCL C 1.2" as the default
|
|
// CL_DEVICE_OPENCL_C_VERSION while supporting 3.0 (the 3.0 API lists all versions; the 1.2 API cannot ask), so
|
|
// the device version counts as well. The local-memory variant is always built as OpenCL C 1.2, the same text everywhere.
|
|
int major = di->cMajor > di->dMajor ? di->cMajor : di->dMajor;
|
|
if (exchangeMode == 0) std = "-cl-std=CL1.2";
|
|
else if (major >= 3) std = "-cl-std=CL3.0";
|
|
else if (major >= 2) std = "-cl-std=CL2.0";
|
|
else std = "-cl-std=CL1.2";
|
|
snprintf(dv->buildOptions, sizeof(dv->buildOptions), "%s -D IGNEUM_GROUP=%d -D IGNEUM_EXCHANGE=%d %s", std, groupSize, exchangeMode, extra);
|
|
dv->prog = clCreateProgramWithSource(dv->ctx, 1, &src, &srcLen, &err);
|
|
CL_CHECK_ERR(err, "clCreateProgramWithSource");
|
|
err = clBuildProgram(dv->prog, 1, &di->device, dv->buildOptions, NULL, NULL);
|
|
if (err != CL_SUCCESS) {
|
|
size_t logLen = 0;
|
|
char* log;
|
|
clGetProgramBuildInfo(dv->prog, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
|
|
log = (char*)calloc(logLen + 1, 1);
|
|
if (logLen) clGetProgramBuildInfo(dv->prog, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
|
|
printf("build FAILED (%s) with options \"%s\"\n--- build log ---\n%s\n--- end of build log ---\n", clErrName(err), dv->buildOptions, log);
|
|
free(log);
|
|
clReleaseProgram(dv->prog); dv->prog = NULL;
|
|
return 1;
|
|
}
|
|
dv->kHash = clCreateKernel(dv->prog, "igneum_hash", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_hash");
|
|
dv->kHashBound = clCreateKernel(dv->prog, "igneum_hash_bound", &err);
|
|
if (err != CL_SUCCESS) dv->kHashBound = NULL; /* kernel.cl without the bound kernel: fine outside --serve */
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
dv->kCacheFill = clCreateKernel(dv->prog, "igneum_cache_fill", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_cache_fill");
|
|
dv->kBuild = clCreateKernel(dv->prog, "igneum_build", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_build");
|
|
#else
|
|
dv->kFill = clCreateKernel(dv->prog, "igneum_fill", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_fill");
|
|
#endif
|
|
return 0;
|
|
}
|
|
|
|
static void releaseProgram(Device* dv) {
|
|
if (dv->kHash) clReleaseKernel(dv->kHash);
|
|
if (dv->kHashBound) clReleaseKernel(dv->kHashBound);
|
|
if (dv->kCacheFill) clReleaseKernel(dv->kCacheFill);
|
|
if (dv->kBuild) clReleaseKernel(dv->kBuild);
|
|
if (dv->kFill) clReleaseKernel(dv->kFill);
|
|
if (dv->prog) clReleaseProgram(dv->prog);
|
|
dv->kHash = dv->kCacheFill = dv->kBuild = dv->kFill = dv->kHashBound = NULL; dv->prog = NULL;
|
|
}
|
|
|
|
// Sub-group size of kernel k for a work-group of `local` items. 0 if the query is unavailable (reason in *why).
|
|
static size_t querySubGroupSizeOf(cl_kernel k, const DeviceInfo* di, size_t local, char* why, size_t whyLen) {
|
|
ig_pfn_subgroup_info fn = (ig_pfn_subgroup_info)clGetExtensionFunctionAddressForPlatform(di->platform, "clGetKernelSubGroupInfoKHR");
|
|
const char* via = "clGetKernelSubGroupInfoKHR";
|
|
size_t sg = 0;
|
|
cl_int e;
|
|
if (!fn) { fn = (ig_pfn_subgroup_info)loadSym("clGetKernelSubGroupInfo"); via = "clGetKernelSubGroupInfo (OpenCL 2.1 core, through the loader)"; }
|
|
if (!fn) { snprintf(why, whyLen, "the sub-group size could not be queried (neither clGetKernelSubGroupInfoKHR nor clGetKernelSubGroupInfo is available)"); return 0; }
|
|
e = fn(k, di->device, IG_CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE, sizeof(local), &local, sizeof(sg), &sg, NULL);
|
|
if (e != CL_SUCCESS) { snprintf(why, whyLen, "the sub-group size query failed: %s returned %s (%d)", via, clErrName(e), (int)e); return 0; }
|
|
snprintf(why, whyLen, "queried through %s", via);
|
|
return sg;
|
|
}
|
|
static size_t querySubGroupSize(const Device* dv, const DeviceInfo* di, size_t local, char* why, size_t whyLen) {
|
|
return querySubGroupSizeOf(dv->kHash, di, local, why, whyLen);
|
|
}
|
|
|
|
/* The compiled kernel as the driver sees it, on every exchange path (5 October 2026, the 9070 XT question): the
|
|
* work-group limit, the preferred multiple (the wave width the compiler chose), local memory, private memory (scratch:
|
|
* anything above 0 means spilled registers, which on AMD costs a memory round trip per spill), and the sub-group size
|
|
* for the built work-group size (wave32 or wave64 on RDNA; the local-memory path never queried it before). */
|
|
#define IG_CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 0x11B3
|
|
#define IG_CL_KERNEL_PRIVATE_MEM_SIZE 0x11B4
|
|
static void printKernelInfo(const DeviceInfo* di, cl_kernel k, const char* kernelName, int groupSize, const char* prefix) {
|
|
size_t wg = 0, mult = 0;
|
|
cl_ulong lmem = 0, pmem = 0;
|
|
char why[256];
|
|
size_t sg;
|
|
clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL);
|
|
clGetKernelWorkGroupInfo(k, di->device, IG_CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(mult), &mult, NULL);
|
|
clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(lmem), &lmem, NULL);
|
|
clGetKernelWorkGroupInfo(k, di->device, IG_CL_KERNEL_PRIVATE_MEM_SIZE, sizeof(pmem), &pmem, NULL);
|
|
sg = querySubGroupSizeOf(k, di, (size_t)groupSize, why, sizeof(why));
|
|
printf("%skernel: %s max work-group %llu, preferred multiple %llu, local memory %llu bytes, private memory %llu bytes%s, work-group %d, sub-group size %llu (%s)%s\n",
|
|
prefix, kernelName, (unsigned long long)wg, (unsigned long long)mult, (unsigned long long)lmem, (unsigned long long)pmem,
|
|
pmem ? " (SPILLED: registers in scratch memory)" : "", groupSize, (unsigned long long)sg, why,
|
|
(sg > (size_t)groupSize) ? " (the work-group fills only part of a wave: see --group-warps)" : "");
|
|
fflush(stdout);
|
|
}
|
|
|
|
// Decide the exchange implementation and build. See WAVEFRONT.md for the rule.
|
|
static void setupProgram(Device* dv, const DeviceInfo* di, const Options* o, const char* src, size_t srcLen) {
|
|
int groupSize = 32 * o->groupWarps;
|
|
int want = 0;
|
|
dv->groupSize = groupSize;
|
|
const char* why = "";
|
|
if (strstr(di->extensions, "cl_khr_subgroup_shuffle")) want = 1;
|
|
else if (strstr(di->extensions, "cl_intel_subgroups")) want = 2;
|
|
else why = "device lists no sub-group shuffle extension";
|
|
if (o->exchange == 1) { want = 0; why = "forced by --exchange local"; }
|
|
if (want != 0 && o->groupWarps != 1) {
|
|
if (o->exchange == 2) { printf("FAIL: --exchange subgroup needs --group-warps 1 (one work-group = one 32-lane unit)\n"); exit(2); }
|
|
want = 0; why = "--group-warps is not 1, so a work-group is not one 32-lane unit";
|
|
}
|
|
if (want == 0 && o->exchange == 2) { printf("FAIL: --exchange subgroup requested but %s\n", why); exit(2); }
|
|
|
|
if (want != 0) {
|
|
if (buildProgram(dv, di, src, srcLen, want, groupSize, o->extraOpts) != 0) {
|
|
if (o->exchange == 2) { printf("FAIL: the sub-group variant did not compile\n"); exit(2); }
|
|
want = 0; why = "the sub-group variant did not compile (log above)";
|
|
} else {
|
|
static char qwhy[256];
|
|
size_t sg = querySubGroupSize(dv, di, 32, qwhy, sizeof(qwhy));
|
|
if (sg == 0) {
|
|
// Second choice: a probe kernel from the same build reports get_sub_group_size() for a 32-item work-group.
|
|
// Weaker than the per-kernel query (a compiler may pick the wave width per kernel), and the note says so.
|
|
cl_int perr = 0;
|
|
cl_kernel kp = clCreateKernel(dv->prog, "igneum_probe_subgroup", &perr);
|
|
if (perr == CL_SUCCESS) {
|
|
cl_uint probe[2] = { 0u, 0u };
|
|
cl_mem pb = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, sizeof(probe), NULL, &perr);
|
|
size_t g = 32, l = 32;
|
|
if (perr == CL_SUCCESS && clSetKernelArg(kp, 0, sizeof(cl_mem), &pb) == CL_SUCCESS &&
|
|
clEnqueueNDRangeKernel(dv->q, kp, 1, NULL, &g, &l, 0, NULL, NULL) == CL_SUCCESS &&
|
|
clEnqueueReadBuffer(dv->q, pb, CL_TRUE, 0, sizeof(probe), probe, 0, NULL, NULL) == CL_SUCCESS && probe[0] != 0u) {
|
|
static char pwhy[400];
|
|
snprintf(pwhy, sizeof(pwhy), "%s; probe kernel reports get_sub_group_size() %u and %u sub-group(s) for a 32-item work-group (weaker than the per-kernel query)", qwhy, probe[0], probe[1]);
|
|
sg = probe[0];
|
|
strncpy(qwhy, pwhy, sizeof(qwhy) - 1); qwhy[sizeof(qwhy) - 1] = 0;
|
|
}
|
|
if (pb) clReleaseMemObject(pb);
|
|
clReleaseKernel(kp);
|
|
}
|
|
}
|
|
dv->subGroupSize = sg;
|
|
if (sg == 32) {
|
|
dv->exchange = want;
|
|
snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s, sub-group size %llu for a 32-item work-group (%s)", exchangeName(want), (unsigned long long)sg, qwhy);
|
|
return;
|
|
}
|
|
releaseProgram(dv);
|
|
if (sg == 0) why = qwhy;
|
|
else why = "the sub-group size for a 32-item work-group is not 32";
|
|
if (o->exchange == 2) { printf("FAIL: --exchange subgroup requested but %s (sub-group size %llu)\n", why, (unsigned long long)sg); exit(2); }
|
|
}
|
|
}
|
|
if (buildProgram(dv, di, src, srcLen, 0, groupSize, o->extraOpts) != 0) { printf("FAIL: kernel.cl did not compile\n"); exit(2); }
|
|
dv->exchange = 0;
|
|
if (dv->subGroupSize) snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s (%s; queried sub-group size %llu)", exchangeName(0), why, (unsigned long long)dv->subGroupSize);
|
|
else snprintf(dv->exchangeNote, sizeof(dv->exchangeNote), "%s (%s)", exchangeName(0), why);
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Launch helpers
|
|
|
|
static size_t kernelMaxLocal(const Device* dv, cl_kernel k, const DeviceInfo* di, size_t want) {
|
|
size_t wg = 0;
|
|
(void)dv;
|
|
if (clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL) != CL_SUCCESS || wg == 0) wg = want;
|
|
if (wg > di->maxWorkGroup) wg = di->maxWorkGroup;
|
|
return wg < want ? wg : want;
|
|
}
|
|
|
|
/* Object accounting for the serve loop's stats line (4 October 2026, after the gfx1036 fault): every event and buffer
|
|
* the serve path creates or releases is counted here, so a leak shows as a growing "live" count long before a runtime
|
|
* limit is hit. The bench path does not count. */
|
|
static unsigned long gEvCreated = 0, gEvReleased = 0, gMemCreated = 0, gMemReleased = 0;
|
|
static void countRelease(cl_event ev) { clReleaseEvent(ev); ++gEvReleased; }
|
|
|
|
static cl_event launch1D(const Device* dv, cl_kernel k, size_t global, size_t local) {
|
|
cl_event ev = NULL;
|
|
size_t g = ((global + local - 1) / local) * local;
|
|
CL_CHECK(clEnqueueNDRangeKernel(dv->q, k, 1, NULL, &g, &local, 0, NULL, &ev));
|
|
++gEvCreated;
|
|
return ev;
|
|
}
|
|
|
|
static cl_event launchHash(const Device* dv, cl_mem ds, cl_mem out, cl_uint baseNonce, cl_uint mask, size_t nonces, size_t groupSize) {
|
|
CL_CHECK(clSetKernelArg(dv->kHash, 0, sizeof(cl_mem), &ds));
|
|
CL_CHECK(clSetKernelArg(dv->kHash, 1, sizeof(cl_mem), &out));
|
|
CL_CHECK(clSetKernelArg(dv->kHash, 2, sizeof(cl_uint), &baseNonce));
|
|
CL_CHECK(clSetKernelArg(dv->kHash, 3, sizeof(cl_uint), &mask));
|
|
return launch1D(dv, dv->kHash, nonces, groupSize);
|
|
}
|
|
|
|
static void readWords(const Device* dv, cl_mem buf, size_t wordIndex, size_t nWords, uint32_t* dst) {
|
|
CL_CHECK(clEnqueueReadBuffer(dv->q, buf, CL_TRUE, wordIndex * 4u, nWords * 4u, dst, 0, NULL, NULL));
|
|
}
|
|
|
|
static int compareWarp(const uint64_t* got, const uint64_t* want, uint32_t base, const char* how) {
|
|
int bad = 0, first = -1, l;
|
|
for (l = 0; l < 32; ++l) if (got[l] != want[l]) { if (first < 0) first = l; ++bad; }
|
|
if (bad == 0) printf("verify warp base %u (nonces %u..%u) %s: PASS\n", base, base, base + 31u, how);
|
|
else printf("verify warp base %u (nonces %u..%u) %s: FAIL %d of 32 lanes differ, first lane %d: device=%016llx expected=%016llx\n",
|
|
base, base, base + 31u, how, bad, first, (unsigned long long)got[first], (unsigned long long)want[first]);
|
|
return bad == 0;
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Cache (memory-hard packs)
|
|
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
static cl_mem gCache = NULL;
|
|
static double gCacheFillFirstMs = 0, gCacheFillSecondMs = 0, gCacheHostMs = 0;
|
|
static int gCachePass = 0;
|
|
|
|
static int setupCache(Device* dv, const DeviceInfo* di) {
|
|
size_t bytes = (size_t)CACHE_WORDS_HOST * 4u;
|
|
cl_int err = 0;
|
|
cl_uint nSeg = IGNEUM_CACHE_SEGMENTS;
|
|
size_t local;
|
|
int pass;
|
|
uint32_t* dev;
|
|
uint32_t seg;
|
|
double h0;
|
|
int same, fnvOk, headOk, lastOk;
|
|
uint64_t fnv;
|
|
gCache = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, bytes, NULL, &err);
|
|
CL_CHECK_ERR(err, "clCreateBuffer cache");
|
|
local = kernelMaxLocal(dv, dv->kCacheFill, di, 256);
|
|
for (pass = 0; pass < 2; ++pass) {
|
|
cl_event ev;
|
|
CL_CHECK(clSetKernelArg(dv->kCacheFill, 0, sizeof(cl_mem), &gCache));
|
|
CL_CHECK(clSetKernelArg(dv->kCacheFill, 1, sizeof(cl_uint), &nSeg));
|
|
ev = launch1D(dv, dv->kCacheFill, nSeg, local);
|
|
CL_CHECK(clFinish(dv->q));
|
|
if (pass == 0) gCacheFillFirstMs = eventMs(ev); else gCacheFillSecondMs = eventMs(ev);
|
|
countRelease(ev);
|
|
}
|
|
printf("cache fill (device): %.2f ms first, %.2f ms second (%u chains x %u ChaCha blocks, %u MiB, work-group %llu)\n",
|
|
gCacheFillFirstMs, gCacheFillSecondMs, (unsigned)IGNEUM_CACHE_SEGMENTS, 1u << IGNEUM_CACHE_SEGMENT_LOG2_LINES,
|
|
(unsigned)(bytes >> 20), (unsigned long long)local);
|
|
|
|
hCache = (uint32_t*)malloc(bytes);
|
|
h0 = wallMs();
|
|
for (seg = 0; seg < IGNEUM_CACHE_SEGMENTS; ++seg) mh_cache_segment(hCache, seg);
|
|
gCacheHostMs = wallMs() - h0;
|
|
printf("cache fill (host, one thread): %.1f ms\n", gCacheHostMs);
|
|
|
|
dev = (uint32_t*)malloc(bytes);
|
|
readWords(dv, gCache, 0, CACHE_WORDS_HOST, dev);
|
|
same = memcmp(dev, hCache, bytes) == 0;
|
|
fnv = fnv1a64(hCache, bytes);
|
|
fnvOk = (fnv == IGNEUM_CACHE_FNV64);
|
|
headOk = memcmp(hCache, IGNEUM_CACHE_HEAD, 64) == 0;
|
|
lastOk = memcmp(hCache + CACHE_WORDS_HOST - 16u, IGNEUM_CACHE_LAST, 64) == 0;
|
|
if (!same) {
|
|
uint32_t i;
|
|
for (i = 0; i < CACHE_WORDS_HOST; ++i) if (dev[i] != hCache[i]) {
|
|
printf(" cache[%u]: device 0x%08x host 0x%08x (first difference)\n", i, dev[i], hCache[i]); break;
|
|
}
|
|
}
|
|
free(dev);
|
|
gCachePass = same && fnvOk && headOk && lastOk;
|
|
printf("cache check: %s (device == host all %u words %s, host FNV-1a 64 %016llx vs Mac %016llx %s, head 16 vs Mac %s, last line vs Mac %s)\n",
|
|
gCachePass ? "PASS" : "FAIL", CACHE_WORDS_HOST, same ? "PASS" : "FAIL",
|
|
(unsigned long long)fnv, (unsigned long long)IGNEUM_CACHE_FNV64, fnvOk ? "PASS" : "FAIL",
|
|
headOk ? "PASS" : "FAIL", lastOk ? "PASS" : "FAIL");
|
|
return gCachePass;
|
|
}
|
|
#endif
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// One dataset size: fill or build, self-test, vectors, bench
|
|
|
|
typedef struct {
|
|
int mib;
|
|
uint32_t words;
|
|
double fillFirstMs, fillSecondMs;
|
|
int dsPass;
|
|
int vecChecked, vecPass;
|
|
double gpuMs, wallMsTimed;
|
|
double hashesPerSec, gbps;
|
|
} SizeResult;
|
|
|
|
static SizeResult runSize(Device* dv, const DeviceInfo* di, const Options* o, int mib, cl_mem dOut, uint32_t nonces) {
|
|
SizeResult r;
|
|
uint64_t bytes = (uint64_t)mib << 20;
|
|
uint32_t mask;
|
|
int atPackSize;
|
|
cl_int err = 0;
|
|
cl_mem dDs;
|
|
int pass;
|
|
size_t groupSize = 32u * (size_t)o->groupWarps;
|
|
uint64_t got[32];
|
|
double w0, w1, t0, t1, total;
|
|
cl_event* evs;
|
|
int b;
|
|
memset(&r, 0, sizeof(r));
|
|
r.mib = mib;
|
|
r.words = (uint32_t)(bytes / 4ull);
|
|
mask = r.words - 1u;
|
|
atPackSize = (r.words == (1u << IGNEUM_DATASET_LOG2));
|
|
printf("\n=== dataset %d MiB (2^%d words, mask 0x%08x)%s ===\n", mib, log2u32(r.words), mask,
|
|
atPackSize ? "" : " [not the pack size: vectors skipped, dataset head and random points still checked]");
|
|
if ((uint64_t)di->maxAlloc < bytes) {
|
|
printf("FAIL: CL_DEVICE_MAX_MEM_ALLOC_SIZE is %llu MiB, the dataset needs %d MiB in one buffer\n", (unsigned long long)(di->maxAlloc >> 20), mib);
|
|
exit(2);
|
|
}
|
|
dDs = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)bytes, NULL, &err);
|
|
CL_CHECK_ERR(err, "clCreateBuffer dataset");
|
|
|
|
// Fill (or build) twice: the Mac showed a first-touch cost on the first fill of a process.
|
|
for (pass = 0; pass < 2; ++pass) {
|
|
cl_event ev;
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
cl_uint nItems = r.words / 16u;
|
|
size_t local = kernelMaxLocal(dv, dv->kBuild, di, 256);
|
|
CL_CHECK(clSetKernelArg(dv->kBuild, 0, sizeof(cl_mem), &dDs));
|
|
CL_CHECK(clSetKernelArg(dv->kBuild, 1, sizeof(cl_mem), &gCache));
|
|
CL_CHECK(clSetKernelArg(dv->kBuild, 2, sizeof(cl_uint), &nItems));
|
|
ev = launch1D(dv, dv->kBuild, nItems, local);
|
|
#else
|
|
cl_uint n = r.words, d0 = IGNEUM_DAY0, d1 = IGNEUM_DAY1;
|
|
size_t local = kernelMaxLocal(dv, dv->kFill, di, 256);
|
|
CL_CHECK(clSetKernelArg(dv->kFill, 0, sizeof(cl_mem), &dDs));
|
|
CL_CHECK(clSetKernelArg(dv->kFill, 1, sizeof(cl_uint), &n));
|
|
CL_CHECK(clSetKernelArg(dv->kFill, 2, sizeof(cl_uint), &d0));
|
|
CL_CHECK(clSetKernelArg(dv->kFill, 3, sizeof(cl_uint), &d1));
|
|
ev = launch1D(dv, dv->kFill, n, local);
|
|
#endif
|
|
CL_CHECK(clFinish(dv->q));
|
|
if (pass == 0) r.fillFirstMs = eventMs(ev); else r.fillSecondMs = eventMs(ev);
|
|
countRelease(ev);
|
|
}
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
printf("dataset build (memory-hard, from the cache): %.2f ms first, %.2f ms second -> %.1f M items/s, %.2f G cache-line reads/s (second, device time)\n",
|
|
r.fillFirstMs, r.fillSecondMs, (double)(r.words / 16u) / 1e6 / (r.fillSecondMs / 1000.0),
|
|
(double)(r.words / 16u) * (double)IGNEUM_ITEM_ROUNDS / 1e9 / (r.fillSecondMs / 1000.0));
|
|
#else
|
|
printf("dataset fill: %.2f ms first, %.2f ms second -> %.0f GB/s write (second, device time)\n",
|
|
r.fillFirstMs, r.fillSecondMs, (double)bytes / 1e9 / (r.fillSecondMs / 1000.0));
|
|
#endif
|
|
|
|
// Dataset self-test: head 16 (any size), element [MASK] (pack size only), 64 pseudo-random points vs the host
|
|
// formula (closed form) or the host derivation from the host cache (memory-hard), and the Mac's 64 sampled words.
|
|
{
|
|
uint32_t head[16];
|
|
int badHead = 0, i, k, badRnd = 0, badSample = 0, nSample = 0, lastOk = 1;
|
|
const char* lastText = "skipped";
|
|
uint64_t s = 0x9E3779B97F4A7C15ull ^ (uint64_t)r.words;
|
|
readWords(dv, dDs, 0, 16, head);
|
|
for (i = 0; i < 16; ++i) if (head[i] != IGNEUM_DS_HEAD[i]) {
|
|
if (badHead == 0) printf(" dataset[%d] = 0x%08x, expected 0x%08x\n", i, head[i], IGNEUM_DS_HEAD[i]);
|
|
++badHead;
|
|
}
|
|
if (atPackSize) {
|
|
uint32_t last = 0;
|
|
readWords(dv, dDs, IGNEUM_DS_LAST_INDEX, 1, &last);
|
|
lastOk = (last == IGNEUM_DS_LAST);
|
|
lastText = lastOk ? "PASS" : "FAIL";
|
|
if (!lastOk) printf(" dataset[%u] = 0x%08x, expected 0x%08x\n", IGNEUM_DS_LAST_INDEX, last, IGNEUM_DS_LAST);
|
|
}
|
|
for (k = 0; k < 64; ++k) {
|
|
uint64_t z;
|
|
uint32_t idx, v = 0, want;
|
|
s += 0x9E3779B97F4A7C15ull;
|
|
z = s;
|
|
z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ull;
|
|
z = (z ^ (z >> 27)) * 0x94D049BB133111EBull;
|
|
z ^= z >> 31;
|
|
idx = (uint32_t)z & mask;
|
|
readWords(dv, dDs, idx, 1, &v);
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
want = host_ds_word(idx);
|
|
#else
|
|
want = host_ds_elem(idx, IGNEUM_DAY0, IGNEUM_DAY1);
|
|
#endif
|
|
if (v != want) {
|
|
if (badRnd == 0) printf(" dataset[%u] = 0x%08x, host %s 0x%08x\n", idx, v, IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", want);
|
|
++badRnd;
|
|
}
|
|
}
|
|
#ifdef IGNEUM_DS_SAMPLES
|
|
for (k = 0; k < IGNEUM_DS_SAMPLES; ++k) {
|
|
uint32_t v = 0;
|
|
if (IGNEUM_DS_SAMPLE_INDEX[k] > mask) continue;
|
|
++nSample;
|
|
readWords(dv, dDs, IGNEUM_DS_SAMPLE_INDEX[k], 1, &v);
|
|
if (v != IGNEUM_DS_SAMPLE_VALUE[k]) {
|
|
if (badSample == 0) printf(" dataset[%u] = 0x%08x, Mac 0x%08x\n", IGNEUM_DS_SAMPLE_INDEX[k], v, IGNEUM_DS_SAMPLE_VALUE[k]);
|
|
++badSample;
|
|
}
|
|
}
|
|
#endif
|
|
r.dsPass = (badHead == 0 && lastOk && badRnd == 0 && badSample == 0);
|
|
printf("dataset self-test: %s (head 16 vs Mac %s, element [MASK] vs Mac %s, 64 random points vs host %s %s, %d Mac samples %s)\n",
|
|
r.dsPass ? "PASS" : "FAIL", badHead == 0 ? "PASS" : "FAIL", lastText,
|
|
IGNEUM_DATASET_MODE == 1 ? "derivation" : "formula", badRnd == 0 ? "PASS" : "FAIL",
|
|
nSample, nSample == 0 ? "none in pack" : (badSample == 0 ? "PASS" : "FAIL"));
|
|
}
|
|
|
|
// Vectors, standalone: one 32-item work-group per base nonce, exactly like the Mac cross-check.
|
|
if (atPackSize) {
|
|
int w;
|
|
r.vecChecked = 1; r.vecPass = 1;
|
|
for (w = 0; w < IGNEUM_VEC_WARPS; ++w) {
|
|
cl_event ev;
|
|
if (o->groupWarps != 1) {
|
|
// reqd_work_group_size pins the hash kernel to 32 x group-warps items; run one full group and read its first 32.
|
|
ev = launchHash(dv, dDs, dOut, IGNEUM_VEC_BASE[w], mask, groupSize, groupSize);
|
|
} else {
|
|
ev = launchHash(dv, dDs, dOut, IGNEUM_VEC_BASE[w], mask, 32u, 32u);
|
|
}
|
|
CL_CHECK(clFinish(dv->q));
|
|
countRelease(ev);
|
|
CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, sizeof(got), got, 0, NULL, NULL));
|
|
r.vecPass = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], o->groupWarps == 1 ? "standalone, 1 unit/work-group" : "standalone, first unit of one work-group") && r.vecPass;
|
|
}
|
|
} else {
|
|
printf("vectors: skipped (the pack's vectors are for %d MiB)\n", packMib());
|
|
}
|
|
|
|
// Warm-up batch at base nonce 0. With the default 2^24 nonces it contains all three vector warps,
|
|
// so the bench configuration itself (work-group = 32 x group-warps) is also checked bit for bit.
|
|
w0 = wallMs();
|
|
{
|
|
cl_event ev = launchHash(dv, dDs, dOut, 0u, mask, nonces, groupSize);
|
|
CL_CHECK(clFinish(dv->q));
|
|
countRelease(ev);
|
|
}
|
|
w1 = wallMs();
|
|
printf("warm-up batch: %u hashes in %.2f ms wall\n", nonces, w1 - w0);
|
|
if (atPackSize) {
|
|
// Fingerprint of every output in the batch, so two runs (two devices, two exchange paths, the CPU emulator at the
|
|
// same --batch-log2) can be compared for all nonces, not only the vector warps.
|
|
uint64_t* all = (uint64_t*)malloc((size_t)nonces * sizeof(uint64_t));
|
|
uint64_t fp = 0xcbf29ce484222325ull;
|
|
size_t i;
|
|
CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)nonces * sizeof(uint64_t), all, 0, NULL, NULL));
|
|
for (i = 0; i < (size_t)nonces * 8u; ++i) { fp ^= ((const uint8_t*)all)[i]; fp *= 0x100000001b3ull; }
|
|
free(all);
|
|
printf("batch fingerprint (FNV-1a 64 of 2^%d outputs at base nonce 0): %016llx\n", o->batchLog2, (unsigned long long)fp);
|
|
}
|
|
if (atPackSize) {
|
|
char how[64];
|
|
int w;
|
|
snprintf(how, sizeof(how), "in batch, %d unit(s)/work-group", o->groupWarps);
|
|
for (w = 0; w < IGNEUM_VEC_WARPS; ++w) {
|
|
if ((uint64_t)IGNEUM_VEC_BASE[w] + 32ull > (uint64_t)nonces) {
|
|
printf("verify warp base %u in batch: skipped (batch has %u nonces)\n", IGNEUM_VEC_BASE[w], nonces);
|
|
continue;
|
|
}
|
|
CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, (size_t)IGNEUM_VEC_BASE[w] * 8u, sizeof(got), got, 0, NULL, NULL));
|
|
r.vecPass = compareWarp(got, IGNEUM_VEC_OUT[w], IGNEUM_VEC_BASE[w], how) && r.vecPass;
|
|
}
|
|
}
|
|
|
|
// Timed batches. Base nonces (b * nonces) mod 2^32, as in the Metal and CUDA runs. Device time is the span from
|
|
// the first batch's start to the last batch's end from event profiling, like cudaEvent elapsed time.
|
|
evs = (cl_event*)calloc((size_t)o->batches, sizeof(cl_event));
|
|
t0 = wallMs();
|
|
for (b = 1; b <= o->batches; ++b) {
|
|
uint32_t base = (uint32_t)((uint64_t)b * (uint64_t)nonces);
|
|
evs[b - 1] = launchHash(dv, dDs, dOut, base, mask, nonces, groupSize);
|
|
}
|
|
CL_CHECK(clFinish(dv->q));
|
|
t1 = wallMs();
|
|
total = (double)nonces * (double)o->batches;
|
|
r.gpuMs = spanMs(evs[0], evs[o->batches - 1]);
|
|
for (b = 0; b < o->batches; ++b) clReleaseEvent(evs[b]);
|
|
free(evs);
|
|
r.wallMsTimed = t1 - t0;
|
|
if (r.gpuMs < 0 && !o->timeWall) { printf("NOTE: event profiling unavailable on this runtime; the rate uses wall time\n"); ((Options*)o)->timeWall = 1; }
|
|
r.hashesPerSec = total / ((o->timeWall ? r.wallMsTimed : r.gpuMs) / 1000.0);
|
|
r.gbps = r.hashesPerSec * (double)IGNEUM_LOADS_PER_HASH * 4.0 / 1e9;
|
|
printf("timed: %d batches x %u hashes = %.0f hashes (rate below from %s time)\n", o->batches, nonces, total, o->timeWall ? "wall" : "device event");
|
|
printf(" device %.2f ms -> %.3f Mhash/s%s\n", r.gpuMs, total / (r.gpuMs / 1000.0) / 1e6, o->timeWall ? " (event profiling, not used for the rate on this platform)" : "");
|
|
printf(" wall %.2f ms -> %.3f Mhash/s\n", r.wallMsTimed, total / (r.wallMsTimed / 1000.0) / 1e6);
|
|
printf(" rate %.3f Mhash/s (%.0f hashes/s), %.2f GB/s useful (loads x 4 B)\n", r.hashesPerSec / 1e6, r.hashesPerSec, r.gbps);
|
|
|
|
CL_CHECK(clReleaseMemObject(dDs));
|
|
return r;
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Main
|
|
|
|
|
|
/* ---------------------------------------------------------------------------------------------
|
|
* Serve mode: GPU worker for igneum-miner --worker (3 October 2026)
|
|
*
|
|
* Protocol, one line each. Only "found", "done" and "error" are parsed by the miner; every other line is logged.
|
|
* stdin: job <job_id> <header_prehash_hex 64> <target_hex 16> <nonce_start u64> <nonce_count u64> <epoch_seed_hex 64> <day_seed_hex>
|
|
* quit
|
|
* stdout: ready opencl <device>
|
|
* found <job_id> <nonce u64> <hash_hex 16> every nonce whose 64-bit hash is <= target
|
|
* done <job_id> <hashes> <ms> end of the job (wall ms)
|
|
* error <job_id> <text>
|
|
* prepare <epoch_seed_hex 64> <day_seed_hex> <pack_dir> build <pack_dir>/kernel_bound.cl (the pack the miner wrote
|
|
* for those seeds) plus its cache and dataset in the background
|
|
* stdout: ready opencl <device> ... prepare 1
|
|
* prepared <epoch_seed_hex> <day_seed_hex> <ms> ... the pair is resident; a job on it switches instantly
|
|
* prepare-failed <epoch_seed_hex> <day_seed_hex> <text>
|
|
* The pack's program.h is compiled in and kernel_bound.cl is built at runtime, so at start this worker serves exactly
|
|
* one epoch seed and one day seed: the pack's. The next pair arrives through `prepare`: the miner writes the pack for the
|
|
* prepared seeds (igneum-miner --prepare-packs <dir>) and names its directory; a background thread builds that pack's
|
|
* kernel_bound.cl with the same build options, fills its cache and builds its dataset on a second queue while jobs on
|
|
* the current pair keep running (at most two pairs resident: the current one and the prepared one; the old pair is
|
|
* released after the first job on the new one). A job for seeds that are neither the current nor the prepared pair is
|
|
* answered with an error naming both. Without prepare, re-export the pack with `igneum-miner export-pack <node> <dir>`
|
|
* and rebuild. A prepared pair's cache is not cross-checked against a host fill (memhard.h is compiled in for the
|
|
* original day); the miner's CPU re-check of every found nonce covers it. The init words of a dispatch are
|
|
* seed_words_from_bytes("igneum-block/" || prehash || nonce_hi_le32), written to a small buffer that is the fifth
|
|
* argument of igneum_hash_bound; the lane nonce is baseNonce + gid as in the bench kernel. The exchange rule of
|
|
* WAVEFRONT.md applies unchanged (the bound kernel has the same body and the same IGNEUM_EXCHANGE build).
|
|
*/
|
|
|
|
static void seedWordsFromBytes(const uint8_t* b, size_t n, uint32_t out[8]) {
|
|
uint64_t salt;
|
|
for (salt = 0; salt < 4; ++salt) {
|
|
uint64_t h = 0xcbf29ce484222325ull ^ (salt * 0x9E3779B97F4A7C15ull);
|
|
size_t i;
|
|
for (i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
|
|
h ^= h >> 33; h *= 0xff51afd7ed558ccdull; h ^= h >> 33;
|
|
out[2 * salt] = (uint32_t)h;
|
|
out[2 * salt + 1] = (uint32_t)(h >> 32);
|
|
}
|
|
}
|
|
|
|
static int unhexBuf(const char* s, uint8_t* out, size_t cap, size_t* len) {
|
|
size_t n = strlen(s), i;
|
|
if (n % 2 || n / 2 > cap) return 0;
|
|
for (i = 0; i < n; i += 2) {
|
|
unsigned v = 0;
|
|
char two[3] = { s[i], s[i + 1], 0 };
|
|
if (sscanf(two, "%2x", &v) != 1) return 0;
|
|
out[i / 2] = (uint8_t)v;
|
|
}
|
|
*len = n / 2;
|
|
return 1;
|
|
}
|
|
|
|
/* Generic serve mode (--pack, 4 October 2026): the sizes and seeds come from the pack directory read at run time, not
|
|
* from the compiled-in program.h, so one prebuilt exe serves every pack. The compiled-in values are the defaults. */
|
|
static PfPack gPack;
|
|
static int gGeneric = 0;
|
|
/* Variant 5 of the read-width experiment (5 October 2026): the scratch arena of a persistent-warp pack, its warp count
|
|
* and the running tag salt; set by --bench-pack before the self-test. Serve mode does not support these packs. */
|
|
static cl_mem gScratch = NULL;
|
|
static cl_uint gScratchWarps = 0;
|
|
static cl_uint gSalt = 1;
|
|
/* Sets the three extra arguments of a variant-5 kernel (after the five of igneum_hash_bound) for `units` units. */
|
|
static cl_int setScratchArgs(cl_kernel k, cl_uint firstArg, cl_uint units) {
|
|
cl_int e = clSetKernelArg(k, firstArg, sizeof(cl_mem), &gScratch);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(k, firstArg + 1, sizeof(cl_uint), &units);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(k, firstArg + 2, sizeof(cl_uint), &gSalt);
|
|
gSalt += units;
|
|
return e;
|
|
}
|
|
static uint32_t gServeWords = 1u << IGNEUM_DATASET_LOG2;
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
static uint32_t gServeCacheWords = 1u << IGNEUM_CACHE_LOG2_WORDS;
|
|
static uint32_t gServeSegments = IGNEUM_CACHE_SEGMENTS;
|
|
#endif
|
|
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
/* One resident (program, cache, dataset) triple for a seed pair. The first one is the compiled-in pack on the main
|
|
* queue (or, with --pack, the directory's pack); prepared ones are built from a pack directory on their own queue. */
|
|
typedef struct {
|
|
char epochHex[65];
|
|
char dayHex[512];
|
|
uint32_t sw[8], kw[8]; /* seed words and key words, for job matching */
|
|
cl_program prog;
|
|
cl_kernel kHashBound, kCacheFill, kBuild;
|
|
cl_mem cache, ds;
|
|
/* hot-table experiment (5 October 2026, docs/plans/hot-table.md): the epoch's table, filled on the device by the
|
|
* pack's igneum_hot_fill, the argument after the init words; hotWords 0 for a pack without one */
|
|
cl_kernel kHotFill;
|
|
cl_mem hot;
|
|
uint32_t hotWords, hotSegments, hotMb, hotSlots;
|
|
double hotMs;
|
|
double buildMs, cacheMs, datasetMs, checkMs;
|
|
char check[1024]; /* the self-test verdict (packfile.h), one line */
|
|
int checked;
|
|
char programClass[8]; /* Counter ASIC 2.0: the pack's class ("v2" or "v3") and era seed, from packfile.h */
|
|
char eraHex[65];
|
|
} ServePair;
|
|
|
|
static int hexEq(const char* a, const char* b) {
|
|
size_t i;
|
|
if (strlen(a) != strlen(b)) return 0;
|
|
for (i = 0; a[i]; ++i) if (tolower((unsigned char)a[i]) != tolower((unsigned char)b[i])) return 0;
|
|
return 1;
|
|
}
|
|
|
|
/* A pair is the pair of a job when the job's seeds (the hex the node sent) are the pair's seeds. The derived seed
|
|
* words are no identity: a retried program's words are its attempt's words, not the bare seed's (packfile.h,
|
|
* 5 October 2026), so comparing words refused every job of a retried program. The compiled-in placeholder pack
|
|
* (no --pack, no prepared pair) has no seed hex; it keeps the word comparison. */
|
|
static int pairIs(const ServePair* p, const char* epochHex, const char* dayHex, const uint32_t sw[8], const uint32_t kw[8]) {
|
|
if (!p) return 0;
|
|
if (p->epochHex[0]) return hexEq(p->epochHex, epochHex) && hexEq(p->dayHex, dayHex);
|
|
return memcmp(sw, p->sw, 32) == 0 && memcmp(kw, p->kw, 32) == 0;
|
|
}
|
|
|
|
/* Counter ASIC 2.0: a job that names a class (and an era) belongs to a pair of that class (and era) only. */
|
|
static int pairIsClass(const ServePair* p, const char* epochHex, const char* dayHex, const uint32_t sw[8], const uint32_t kw[8], const char* cls, const char* era) {
|
|
char why[256];
|
|
return pairIs(p, epochHex, dayHex, sw, kw) && pf_pack_class_ok(p->programClass[0] ? p->programClass : "v2", p->eraHex, cls, era, why, sizeof(why));
|
|
}
|
|
|
|
/* The trailing `class=` and `era=` tokens of a job or prepare line (absent on every class v2 line), removed from f. */
|
|
static void takeClassTokens(char** f, int* nf, char* cls, size_t clsCap, char* era, size_t eraCap) {
|
|
cls[0] = 0; era[0] = 0;
|
|
while (*nf > 0 && pf_class_token(f[*nf - 1], cls, clsCap, era, eraCap)) --*nf;
|
|
}
|
|
|
|
static void releasePair(ServePair* p) {
|
|
if (!p) return;
|
|
if (p->ds) { clReleaseMemObject(p->ds); ++gMemReleased; }
|
|
if (p->cache) { clReleaseMemObject(p->cache); ++gMemReleased; }
|
|
if (p->hot) { clReleaseMemObject(p->hot); ++gMemReleased; }
|
|
if (p->kHotFill) clReleaseKernel(p->kHotFill);
|
|
if (p->kHashBound) clReleaseKernel(p->kHashBound);
|
|
if (p->kCacheFill) clReleaseKernel(p->kCacheFill);
|
|
if (p->kBuild) clReleaseKernel(p->kBuild);
|
|
if (p->prog) clReleaseProgram(p->prog);
|
|
free(p);
|
|
}
|
|
|
|
/* The prepare request and its result, handed between the main loop and the prepare thread. */
|
|
typedef struct {
|
|
Device* dv;
|
|
const DeviceInfo* di;
|
|
char epochHex[65];
|
|
char dayHex[512];
|
|
char packDir[1024];
|
|
char wantClass[8]; /* the class and era the prepare line named (empty: any) */
|
|
char wantEra[65];
|
|
uint32_t words, cacheWords, segments;
|
|
char error[512];
|
|
ServePair* result; /* set by the thread on success */
|
|
volatile int done; /* 1 when the thread has finished (success or failure) */
|
|
double t0, doneAt;
|
|
} PrepareTask;
|
|
|
|
static void prepareFail(PrepareTask* t, const char* what, cl_int err) {
|
|
snprintf(t->error, sizeof(t->error), "%s (%s)", what, clErrName(err));
|
|
}
|
|
|
|
/* The self-test of packfile.h on a pair: cache head, last line and FNV-1a 64 over the whole cache, dataset head, last
|
|
* word and samples, and the vector warps through igneum_hash_bound with the pack's own seed words as init words (that
|
|
* is igneum_hash of kernel.cl). The verdict goes to p->check. A pack directory without a readable vectors.h is not
|
|
* tested (p->checked stays 0: the miner's CPU re-check still covers every found nonce); a FAIL is an error. */
|
|
static int pairSelfTest(Device* dv, const DeviceInfo* di, cl_command_queue q, ServePair* p, uint32_t words, uint32_t cacheWords, const char* packDir, char* err, size_t errCap) {
|
|
PfPack pk;
|
|
char perr[256];
|
|
uint32_t cacheHead[16], cacheLast[16], dsHead[16], dsLast = 0;
|
|
uint32_t hotHead[16], hotLast[16];
|
|
uint64_t hotFnv = 0;
|
|
uint32_t samples[PF_MAX_SAMPLES];
|
|
uint64_t vec[PF_MAX_WARPS * 32];
|
|
uint64_t fnv;
|
|
uint32_t* whole;
|
|
cl_mem out = NULL, init = NULL;
|
|
cl_int e = 0;
|
|
int w, i, ok;
|
|
cl_uint extra = 5; /* the first argument after the five of igneum_hash_bound: the hot table, then the scratch triple */
|
|
size_t local = (size_t)dv->groupSize, g = (size_t)dv->groupSize; /* one work-group; its first 32 lanes are the warp */
|
|
double tb = wallMs();
|
|
(void)di;
|
|
p->checked = 0;
|
|
perr[0] = 0;
|
|
if (!packDir || !packDir[0] || !pf_load(packDir, &pk, perr, sizeof(perr)) || !pk.haveVectors) {
|
|
snprintf(p->check, sizeof(p->check), "self-test skipped (%s); the miner's CPU re-check covers every found nonce", (packDir && packDir[0]) ? (perr[0] ? perr : "no vectors.h in the pack") : "no pack directory");
|
|
return 1;
|
|
}
|
|
whole = (uint32_t*)malloc((size_t)cacheWords * 4u);
|
|
if (!whole) { snprintf(err, errCap, "self-test: no host memory for the cache read-back"); return 0; }
|
|
e = clEnqueueReadBuffer(q, p->cache, CL_TRUE, 0, (size_t)cacheWords * 4u, whole, 0, NULL, NULL);
|
|
if (e != CL_SUCCESS) { free(whole); snprintf(err, errCap, "self-test: cache read-back (%s)", clErrName(e)); return 0; }
|
|
memcpy(cacheHead, whole, 64);
|
|
memcpy(cacheLast, whole + cacheWords - 16u, 64);
|
|
fnv = pf_fnv1a64(whole, (size_t)cacheWords * 4u);
|
|
free(whole);
|
|
e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, 0, 64, dsHead, 0, NULL, NULL);
|
|
if (e == CL_SUCCESS && pk.dsLastIndex < words) e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, (size_t)pk.dsLastIndex * 4u, 4, &dsLast, 0, NULL, NULL);
|
|
for (i = 0; i < pk.nSamples && e == CL_SUCCESS; ++i) { samples[i] = 0; if (pk.sampleIdx[i] < words) e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, (size_t)pk.sampleIdx[i] * 4u, 4, &samples[i], 0, NULL, NULL); }
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "self-test: dataset read-back (%s)", clErrName(e)); return 0; }
|
|
if (p->hot) {
|
|
uint32_t* hw = (uint32_t*)malloc((size_t)p->hotWords * 4u);
|
|
if (!hw) { snprintf(err, errCap, "self-test: no host memory for the hot table read-back"); return 0; }
|
|
e = clEnqueueReadBuffer(q, p->hot, CL_TRUE, 0, (size_t)p->hotWords * 4u, hw, 0, NULL, NULL);
|
|
if (e != CL_SUCCESS) { free(hw); snprintf(err, errCap, "self-test: hot table read-back (%s)", clErrName(e)); return 0; }
|
|
memcpy(hotHead, hw, 64);
|
|
memcpy(hotLast, hw + p->hotWords - 16u, 64);
|
|
hotFnv = pf_fnv1a64(hw, (size_t)p->hotWords * 4u);
|
|
free(hw);
|
|
}
|
|
out = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, g * sizeof(uint64_t), NULL, &e);
|
|
if (e == CL_SUCCESS) { ++gMemCreated; init = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 32, pk.seedw, &e); }
|
|
if (e == CL_SUCCESS) ++gMemCreated;
|
|
if (pk.persistent) {
|
|
if (!gScratch) { snprintf(err, errCap, "self-test: a variant-5 pack (persistent warps) needs --bench-pack (serve mode does not carry a scratch)"); return 0; }
|
|
g = 32; local = 32; /* one persistent warp runs the one unit */
|
|
}
|
|
for (w = 0; w < pk.vecWarps && e == CL_SUCCESS; ++w) {
|
|
cl_uint base = pk.vecBase[w], mask = words - 1u;
|
|
e = clSetKernelArg(p->kHashBound, 0, sizeof(cl_mem), &p->ds);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 1, sizeof(cl_mem), &out);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 2, sizeof(cl_uint), &base);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 3, sizeof(cl_uint), &mask);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 4, sizeof(cl_mem), &init);
|
|
if (e == CL_SUCCESS && p->hot) { e = clSetKernelArg(p->kHashBound, 5, sizeof(cl_mem), &p->hot); extra = 6; }
|
|
if (e == CL_SUCCESS && pk.persistent) e = setScratchArgs(p->kHashBound, extra, 1u);
|
|
if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kHashBound, 1, NULL, &g, &local, 0, NULL, NULL);
|
|
if (e == CL_SUCCESS) e = clEnqueueReadBuffer(q, out, CL_TRUE, 0, 32 * sizeof(uint64_t), &vec[w * 32], 0, NULL, NULL);
|
|
}
|
|
if (out) { clReleaseMemObject(out); ++gMemReleased; }
|
|
if (init) { clReleaseMemObject(init); ++gMemReleased; }
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "self-test: vector warp (%s)", clErrName(e)); return 0; }
|
|
ok = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples, vec, p->hot ? hotHead : NULL, p->hot ? hotLast : NULL, hotFnv, p->check, sizeof(p->check));
|
|
p->checked = 1;
|
|
p->checkMs = wallMs() - tb;
|
|
if (!ok) { snprintf(err, errCap, "%s", p->check); return 0; }
|
|
return 1;
|
|
}
|
|
|
|
/* Fills the pair's cache and builds its dataset on queue q (the pair's kernels must exist), then self-tests it.
|
|
* Shared by the first pair of --pack and by every prepared pair. Returns 1, or 0 with err set. */
|
|
static int pairBuffers(Device* dv, const DeviceInfo* di, cl_command_queue q, ServePair* p, uint32_t words, uint32_t cacheWords, uint32_t segments, const char* packDir, char* err, size_t errCap) {
|
|
cl_int e = 0;
|
|
cl_uint nSeg = segments, nItems = words / 16u;
|
|
size_t local, g, bytes = (size_t)cacheWords * 4u;
|
|
double tb = wallMs();
|
|
p->cache = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, bytes, NULL, &e);
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer cache (%s)", clErrName(e)); return 0; }
|
|
++gMemCreated;
|
|
local = kernelMaxLocal(dv, p->kCacheFill, di, 256);
|
|
g = ((nSeg + local - 1) / local) * local;
|
|
e = clSetKernelArg(p->kCacheFill, 0, sizeof(cl_mem), &p->cache);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(p->kCacheFill, 1, sizeof(cl_uint), &nSeg);
|
|
if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kCacheFill, 1, NULL, &g, &local, 0, NULL, NULL);
|
|
if (e == CL_SUCCESS) e = clFinish(q);
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "cache fill (%s)", clErrName(e)); return 0; }
|
|
p->cacheMs = wallMs() - tb;
|
|
tb = wallMs();
|
|
p->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &e);
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer dataset (%s)", clErrName(e)); return 0; }
|
|
++gMemCreated;
|
|
local = kernelMaxLocal(dv, p->kBuild, di, 256);
|
|
g = ((nItems + local - 1) / local) * local;
|
|
e = clSetKernelArg(p->kBuild, 0, sizeof(cl_mem), &p->ds);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(p->kBuild, 1, sizeof(cl_mem), &p->cache);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(p->kBuild, 2, sizeof(cl_uint), &nItems);
|
|
if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kBuild, 1, NULL, &g, &local, 0, NULL, NULL);
|
|
if (e == CL_SUCCESS) e = clFinish(q);
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "dataset build (%s)", clErrName(e)); return 0; }
|
|
p->datasetMs = wallMs() - tb;
|
|
if (p->kHotFill && p->hotWords) {
|
|
/* hot-table experiment: the epoch's table from the pack's own fill kernel (one work-item per segment) */
|
|
cl_uint nSegH = p->hotSegments;
|
|
tb = wallMs();
|
|
p->hot = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)p->hotWords * 4u, NULL, &e);
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer hot table (%s)", clErrName(e)); return 0; }
|
|
++gMemCreated;
|
|
local = kernelMaxLocal(dv, p->kHotFill, di, 256);
|
|
g = ((nSegH + local - 1) / local) * local;
|
|
e = clSetKernelArg(p->kHotFill, 0, sizeof(cl_mem), &p->hot);
|
|
if (e == CL_SUCCESS) e = clSetKernelArg(p->kHotFill, 1, sizeof(cl_uint), &nSegH);
|
|
if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kHotFill, 1, NULL, &g, &local, 0, NULL, NULL);
|
|
if (e == CL_SUCCESS) e = clFinish(q);
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "hot table fill (%s)", clErrName(e)); return 0; }
|
|
p->hotMs = wallMs() - tb;
|
|
}
|
|
return pairSelfTest(dv, di, q, p, words, cacheWords, packDir, err, errCap);
|
|
}
|
|
|
|
/* The hot-table kernel and sizes of a pair whose program came from a pack directory (none when the pack has no hot table). */
|
|
static int pairHotKernel(ServePair* p, const char* packDir, char* err, size_t errCap) {
|
|
PfPack pk;
|
|
char perr[256];
|
|
cl_int e = 0;
|
|
if (!packDir || !packDir[0] || !pf_load(packDir, &pk, perr, sizeof(perr)) || !pk.hotMb) return 1;
|
|
p->kHotFill = clCreateKernel(p->prog, "igneum_hot_fill", &e);
|
|
if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateKernel igneum_hot_fill (%s): the pack says IGNEUM_HOT_MB %u but its kernel source has no hot fill", clErrName(e), pk.hotMb); return 0; }
|
|
p->hotWords = pk.hotWords; p->hotSegments = pk.hotSegments; p->hotMb = pk.hotMb; p->hotSlots = pk.hotSlots;
|
|
return 1;
|
|
}
|
|
|
|
/* Builds the pair for a prepare request. Runs on its own thread with its own command queue. */
|
|
static void prepareRun(PrepareTask* t) {
|
|
cl_int err = 0;
|
|
size_t srcLen = 0;
|
|
char path[1200];
|
|
char* src;
|
|
ServePair* p = (ServePair*)calloc(1, sizeof(ServePair));
|
|
cl_command_queue q = NULL;
|
|
double tb;
|
|
strncpy(p->epochHex, t->epochHex, 64); p->epochHex[64] = 0;
|
|
strncpy(p->dayHex, t->dayHex, sizeof(p->dayHex) - 1);
|
|
{
|
|
/* Counter ASIC 2.0: the pack's class and era are read before anything is built; a pack of another class
|
|
* than the prepare line names is refused here, so the miner exports one of the right class */
|
|
PfPack pk; char perr[512]; char why[256];
|
|
if (!pf_load(t->packDir, &pk, perr, sizeof(perr))) { snprintf(t->error, sizeof(t->error), "pack %s: %s", t->packDir, perr); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
strncpy(p->programClass, pk.programClass, sizeof(p->programClass) - 1); strncpy(p->eraHex, pk.eraHex, sizeof(p->eraHex) - 1);
|
|
if (!pf_pack_class_ok(p->programClass, p->eraHex, t->wantClass, t->wantEra, why, sizeof(why))) { snprintf(t->error, sizeof(t->error), "pack %s: %s", t->packDir, why); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
}
|
|
snprintf(path, sizeof(path), "%s/kernel_bound.cl", t->packDir);
|
|
src = readFile(path, &srcLen);
|
|
if (!src) { snprintf(t->error, sizeof(t->error), "cannot read %s", path); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
tb = wallMs();
|
|
p->prog = clCreateProgramWithSource(t->dv->ctx, 1, (const char**)&src, &srcLen, &err);
|
|
free(src);
|
|
if (err != CL_SUCCESS) { prepareFail(t, "clCreateProgramWithSource", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
err = clBuildProgram(p->prog, 1, &t->di->device, t->dv->buildOptions, NULL, NULL);
|
|
if (err != CL_SUCCESS) {
|
|
size_t logLen = 0;
|
|
char* log;
|
|
clGetProgramBuildInfo(p->prog, t->di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
|
|
log = (char*)calloc(logLen + 1, 1);
|
|
if (logLen) clGetProgramBuildInfo(p->prog, t->di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
|
|
snprintf(t->error, sizeof(t->error), "clBuildProgram failed (%s): %.300s", clErrName(err), log);
|
|
free(log); releasePair(p); t->done = 1; return;
|
|
}
|
|
p->buildMs = wallMs() - tb;
|
|
p->kHashBound = clCreateKernel(p->prog, "igneum_hash_bound", &err);
|
|
if (err != CL_SUCCESS) { prepareFail(t, "clCreateKernel igneum_hash_bound (is this a kernel_bound.cl?)", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
p->kCacheFill = clCreateKernel(p->prog, "igneum_cache_fill", &err);
|
|
if (err != CL_SUCCESS) { prepareFail(t, "clCreateKernel igneum_cache_fill", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
p->kBuild = clCreateKernel(p->prog, "igneum_build", &err);
|
|
if (err != CL_SUCCESS) { prepareFail(t, "clCreateKernel igneum_build", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
if (!pairHotKernel(p, t->packDir, t->error, sizeof(t->error))) { releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
q = clCreateCommandQueue(t->dv->ctx, t->di->device, 0, &err);
|
|
if (err != CL_SUCCESS) { prepareFail(t, "clCreateCommandQueue (prepare)", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
if (!pairBuffers(t->dv, t->di, q, p, t->words, t->cacheWords, t->segments, t->packDir, t->error, sizeof(t->error))) { clReleaseCommandQueue(q); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; }
|
|
clReleaseCommandQueue(q);
|
|
t->result = p;
|
|
t->doneAt = wallMs();
|
|
t->done = 1;
|
|
}
|
|
|
|
#ifdef _WIN32
|
|
static DWORD WINAPI prepareThreadMain(LPVOID arg) { prepareRun((PrepareTask*)arg); return 0; }
|
|
static int startPrepareThread(PrepareTask* t) { HANDLE h = CreateThread(NULL, 0, prepareThreadMain, t, 0, NULL); if (!h) return 0; CloseHandle(h); return 1; }
|
|
#else
|
|
static void* prepareThreadMain(void* arg) { prepareRun((PrepareTask*)arg); return NULL; }
|
|
static int startPrepareThread(PrepareTask* t) { pthread_t th; if (pthread_create(&th, NULL, prepareThreadMain, t) != 0) return 0; pthread_detach(th); return 1; }
|
|
#endif
|
|
#endif
|
|
|
|
/* --bench-pack (read-width experiment, 5 October 2026): the pack in --pack is built and self-tested exactly as the
|
|
* first pair of --serve (pairBuffers: cache, dataset, cache FNV, dataset words, the vector warps through
|
|
* igneum_hash_bound with the pack's seed words), then the bound kernel is timed over --batches dispatches of
|
|
* 2^--batch-log2 nonces with device event time, and the 2^B outputs at base nonce 0 are fingerprinted (FNV-1a 64) so
|
|
* the same pack can be compared bit for bit across vendors. A variant-5 pack is launched as --warps persistent warps
|
|
* with a 1 MiB scratch each. One line per run starts with RESULT. */
|
|
static int runBenchPack(Device* dv, const DeviceInfo* di, const Options* o) {
|
|
#if IGNEUM_DATASET_MODE != 1
|
|
(void)dv; (void)di; (void)o;
|
|
printf("FAIL: --bench-pack needs a memory-hard placeholder pack\n");
|
|
return 2;
|
|
#else
|
|
const uint32_t words = gServeWords, mask = words - 1u;
|
|
uint32_t nonces = 1u << o->batchLog2;
|
|
size_t groupSize = 32 * (size_t)o->groupWarps, g;
|
|
cl_int err = 0;
|
|
cl_mem dOut, dInit;
|
|
uint64_t* hOut;
|
|
ServePair* cur;
|
|
char perr[512], devName[256];
|
|
double t0 = wallMs(), sum = 0, warmMs = 0;
|
|
uint64_t fp = 0;
|
|
int b, k;
|
|
cl_uint warps = (cl_uint)(o->warps > 0 ? o->warps : 2048), units = nonces / 32u;
|
|
if (!dv->kHashBound) { printf("FAIL: the kernel source has no igneum_hash_bound\n"); return 2; }
|
|
if (gPack.persistent) {
|
|
size_t arena;
|
|
if (o->groupWarps != 1) { printf("FAIL: a variant-5 pack needs --group-warps 1 (one warp per work-group: the loop trip count must be uniform)\n"); return 2; }
|
|
while (warps > 1 && units % warps != 0) warps >>= 1;
|
|
arena = (size_t)warps * 32u * (size_t)gPack.scratchWordsPerLane * 4u;
|
|
if ((uint64_t)arena > di->maxAlloc) { printf("FAIL: scratch arena %llu MiB exceeds the device's max alloc %llu MiB; lower --warps\n", (unsigned long long)(arena >> 20), (unsigned long long)(di->maxAlloc >> 20)); return 2; }
|
|
gScratch = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, arena, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer scratch");
|
|
gScratchWarps = warps;
|
|
printf("variant 5: %u persistent warps (%u x 32 work-items, work-group 32), scratch arena %llu MiB, %u units per dispatch, lazy tagged fill\n", warps, warps, (unsigned long long)(arena >> 20), units);
|
|
}
|
|
cur = (ServePair*)calloc(1, sizeof(ServePair));
|
|
cur->kHashBound = dv->kHashBound; cur->kCacheFill = dv->kCacheFill; cur->kBuild = dv->kBuild; cur->prog = dv->prog;
|
|
dv->kHashBound = dv->kCacheFill = dv->kBuild = NULL; dv->prog = NULL;
|
|
memcpy(cur->sw, gPack.seedw, 32); memcpy(cur->kw, gPack.keyw, 32);
|
|
if (!pairHotKernel(cur, o->packDir, perr, sizeof(perr))) { printf("FAIL: pack %s: %s\n", o->packDir, perr); return 1; }
|
|
if (!pairBuffers(dv, di, dv->q, cur, words, gServeCacheWords, gServeSegments, o->packDir, perr, sizeof(perr))) { printf("FAIL: pack %s: %s\n", o->packDir, perr); return 1; }
|
|
printf("pack %s: cache %.0f dataset %.0f hot %.0f check %.0f ms (%.0f ms in all); %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->hotMs, cur->checkMs, wallMs() - t0, cur->check);
|
|
printKernelInfo(di, cur->kHashBound, "igneum_hash_bound", (int)groupSize, "");
|
|
dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)nonces * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer out");
|
|
dInit = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 32, cur->sw, &err); CL_CHECK_ERR(err, "clCreateBuffer init words");
|
|
hOut = (uint64_t*)malloc((size_t)nonces * sizeof(uint64_t));
|
|
strncpy(devName, di->name, 255); devName[255] = 0;
|
|
for (k = 0; devName[k]; ++k) if (devName[k] == ' ') devName[k] = '_';
|
|
g = gPack.persistent ? (size_t)warps * 32u : (size_t)nonces;
|
|
for (b = -1; b < o->batches; ++b) {
|
|
cl_uint base = (cl_uint)((uint32_t)(b + 1) * nonces);
|
|
cl_event ev = NULL;
|
|
double ms;
|
|
CL_CHECK(clSetKernelArg(cur->kHashBound, 0, sizeof(cl_mem), &cur->ds));
|
|
CL_CHECK(clSetKernelArg(cur->kHashBound, 1, sizeof(cl_mem), &dOut));
|
|
CL_CHECK(clSetKernelArg(cur->kHashBound, 2, sizeof(cl_uint), &base));
|
|
CL_CHECK(clSetKernelArg(cur->kHashBound, 3, sizeof(cl_uint), &mask));
|
|
CL_CHECK(clSetKernelArg(cur->kHashBound, 4, sizeof(cl_mem), &dInit));
|
|
if (cur->hot) CL_CHECK(clSetKernelArg(cur->kHashBound, 5, sizeof(cl_mem), &cur->hot));
|
|
if (gPack.persistent) CL_CHECK(setScratchArgs(cur->kHashBound, cur->hot ? 6 : 5, units));
|
|
{
|
|
double w0 = wallMs();
|
|
CL_CHECK(clEnqueueNDRangeKernel(dv->q, cur->kHashBound, 1, NULL, &g, &groupSize, 0, NULL, &ev));
|
|
CL_CHECK(clWaitForEvents(1, &ev));
|
|
ms = o->timeWall ? wallMs() - w0 : eventMs(ev);
|
|
if (ms < 0) ms = wallMs() - w0;
|
|
clReleaseEvent(ev);
|
|
}
|
|
if (b < 0) {
|
|
warmMs = ms;
|
|
CL_CHECK(clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)nonces * sizeof(uint64_t), hOut, 0, NULL, NULL));
|
|
fp = pf_fnv1a64((const uint32_t*)hOut, (size_t)nonces * 8u);
|
|
} else sum += ms;
|
|
}
|
|
printf("warm-up dispatch (base 0): %.2f ms; %d timed dispatches of 2^%d nonces: mean %.2f ms\n", warmMs, o->batches, o->batchLog2, sum / o->batches);
|
|
printf("RESULT pack=%s class=%s device=%s platform=%s group=%d warps=%u arena_mib=%llu hot_mib=%u hot_slots=%u hot_fill_ms=%.2f nonces=%u batches=%d check=%s fingerprint=%016llx mhs=%.3f loads=%u bytes=%u scratch_ops=%u time=%s\n",
|
|
o->packDir, gPack.loadClass, devName, strcmp(di->platformName, "Apple") == 0 ? "Apple" : "other", (int)groupSize, gPack.persistent ? warps : 0u,
|
|
gPack.persistent ? (unsigned long long)(((size_t)warps * 32u * gPack.scratchWordsPerLane * 4u) >> 20) : 0ull, cur->hotMb, cur->hotSlots, cur->hotMs, nonces, o->batches,
|
|
cur->checked ? "PASS" : "skipped", (unsigned long long)fp, (double)nonces * (double)o->batches / (sum / 1000.0) / 1e6,
|
|
gPack.loadsPerHash, gPack.bytesPerHash, gPack.scratchOps * 8u, o->timeWall ? "wall" : "event");
|
|
free(hOut);
|
|
clReleaseMemObject(dOut); clReleaseMemObject(dInit);
|
|
if (gScratch) clReleaseMemObject(gScratch);
|
|
releasePair(cur);
|
|
return 0;
|
|
#endif
|
|
}
|
|
|
|
static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
|
|
#if IGNEUM_DATASET_MODE != 1
|
|
(void)dv; (void)di; (void)o;
|
|
printf("error 0 --serve needs a memory-hard pack (IGNEUM_DATASET_MODE 1)\n"); fflush(stdout);
|
|
return 2;
|
|
#else
|
|
static const uint32_t KEYW[8] = IGNEUM_KEY_INIT;
|
|
const uint32_t words = gServeWords;
|
|
const uint32_t mask = words - 1u;
|
|
const uint32_t batch = 1u << (o->batchLog2 == 24 ? 22 : o->batchLog2); /* 2^22 nonces per dispatch by default */
|
|
size_t groupSize = 32 * (size_t)o->groupWarps;
|
|
cl_int err = 0;
|
|
cl_mem dOut, dInit;
|
|
cl_uint nItems = words / 16u;
|
|
uint64_t* hOut;
|
|
char devName[256];
|
|
char line[2048];
|
|
size_t k;
|
|
ServePair* cur; /* the pair jobs run on */
|
|
ServePair* prepared = NULL; /* the pair the last prepare built, until a job switches to it */
|
|
ServePair* old = NULL; /* the previous pair, released after the first job on the new one */
|
|
PrepareTask* task = NULL; /* the prepare in flight */
|
|
if (!dv->kHashBound) { printf("error 0 the kernel source has no igneum_hash_bound (build from the pack's kernel_bound.cl, or pass --kernel)\n"); fflush(stdout); return 2; }
|
|
cur = (ServePair*)calloc(1, sizeof(ServePair));
|
|
cur->kHashBound = dv->kHashBound; cur->kCacheFill = dv->kCacheFill; cur->kBuild = dv->kBuild; cur->prog = dv->prog;
|
|
dv->kHashBound = dv->kCacheFill = dv->kBuild = NULL; dv->prog = NULL; /* owned by the pair now */
|
|
if (gGeneric) {
|
|
/* --pack: the first pair is the directory's pack, built and self-tested exactly like a prepared one */
|
|
char perr[512];
|
|
double t0 = wallMs();
|
|
memcpy(cur->sw, gPack.seedw, 32); memcpy(cur->kw, gPack.keyw, 32);
|
|
strncpy(cur->epochHex, gPack.epochHex, 64); cur->epochHex[64] = 0;
|
|
strncpy(cur->dayHex, gPack.dayHex, sizeof(cur->dayHex) - 1);
|
|
strncpy(cur->programClass, gPack.programClass, sizeof(cur->programClass) - 1); strncpy(cur->eraHex, gPack.eraHex, sizeof(cur->eraHex) - 1);
|
|
if (!pairHotKernel(cur, o->packDir, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o->packDir, perr); fflush(stdout); return 1; }
|
|
if (!pairBuffers(dv, di, dv->q, cur, words, gServeCacheWords, gServeSegments, o->packDir, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o->packDir, perr); fflush(stdout); return 1; }
|
|
printf("info first pack %s: cache %.0f dataset %.0f check %.0f ms (%.0f ms in all); %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->checkMs, wallMs() - t0, cur->check);
|
|
} else {
|
|
if (!setupCache(dv, di)) { printf("error 0 cache check failed (device cache differs from the host cache or the pack's FNV)\n"); fflush(stdout); return 1; }
|
|
memcpy(cur->sw, SEEDW, 32); memcpy(cur->kw, KEYW, 32);
|
|
cur->cache = gCache; gCache = NULL;
|
|
cur->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer dataset");
|
|
CL_CHECK(clSetKernelArg(cur->kBuild, 0, sizeof(cl_mem), &cur->ds));
|
|
CL_CHECK(clSetKernelArg(cur->kBuild, 1, sizeof(cl_mem), &cur->cache));
|
|
CL_CHECK(clSetKernelArg(cur->kBuild, 2, sizeof(cl_uint), &nItems));
|
|
countRelease(launch1D(dv, cur->kBuild, nItems, kernelMaxLocal(dv, cur->kBuild, di, 256)));
|
|
CL_CHECK(clFinish(dv->q));
|
|
}
|
|
dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)batch * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer out");
|
|
dInit = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY, 32, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer init words");
|
|
gMemCreated += 2;
|
|
hOut = (uint64_t*)malloc((size_t)batch * sizeof(uint64_t));
|
|
strncpy(devName, di->name, 255); devName[255] = 0;
|
|
for (k = 0; devName[k]; ++k) if (devName[k] == ' ') devName[k] = '_';
|
|
printKernelInfo(di, cur->kHashBound, "igneum_hash_bound", (int)groupSize, "info ");
|
|
/* The select pass (5 October 2026). Before it every dispatch read back 8 bytes per nonce (16 MiB for a 2^21-nonce
|
|
* job) over the bus and scanned them on the host; on an eGPU over USB4 that is a measurable part of every job.
|
|
* Now a tiny kernel built here (no pack involved) writes the hits (index, hash) behind an atomic counter plus 34
|
|
* sentinel words (the first 32 outputs, the middle and the last), and the host reads back a few hundred bytes.
|
|
* The fault detectors read the sentinels; the found lines are printed in nonce order from the sorted hits. If a
|
|
* chunk has more hits than the table holds (a target that loose is a test, not a block), the chunk falls back to the
|
|
* full read. --readback full or IGNEUM_READBACK=full keeps the old path for a comparison. */
|
|
#define IG_MAX_HITS 256u
|
|
cl_program selProg = NULL;
|
|
cl_kernel kSelect = NULL;
|
|
cl_mem dCount = NULL, dHits = NULL, dSentinel = NULL;
|
|
size_t selLocal = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
|
|
uint32_t hCount = 0;
|
|
uint64_t hHits[IG_MAX_HITS * 2];
|
|
uint64_t hSentinel[34];
|
|
unsigned long long bytesUp = 0, bytesDown = 0;
|
|
double kernelMsSum = 0, selectMsSum = 0, readMsSum = 0, scanMsSum = 0;
|
|
unsigned long fullFallbacks = 0;
|
|
if (!o->readback) {
|
|
static const char* SELECT_SRC =
|
|
"__kernel void igneum_select(__global const ulong* out, uint n, ulong target, volatile __global uint* count,\n"
|
|
" __global ulong* hits, uint maxHits, __global ulong* sentinel) {\n"
|
|
" uint i = (uint)get_global_id(0);\n"
|
|
" if (i < n) {\n"
|
|
" ulong h = out[i];\n"
|
|
" if (h <= target) { uint k = atomic_inc(count); if (k < maxHits) { hits[2u * k] = (ulong)i; hits[2u * k + 1u] = h; } }\n"
|
|
" if (i < 32u) sentinel[i] = h;\n"
|
|
" if (i == 0u) { sentinel[32] = out[n / 2u]; sentinel[33] = out[n - 1u]; }\n"
|
|
" }\n"
|
|
"}\n";
|
|
size_t selLen = strlen(SELECT_SRC);
|
|
selProg = clCreateProgramWithSource(dv->ctx, 1, &SELECT_SRC, &selLen, &err); CL_CHECK_ERR(err, "clCreateProgramWithSource select");
|
|
err = clBuildProgram(selProg, 1, &di->device, "-cl-std=CL1.2", NULL, NULL);
|
|
if (err != CL_SUCCESS) {
|
|
size_t logLen = 0; char* log;
|
|
clGetProgramBuildInfo(selProg, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
|
|
log = (char*)calloc(logLen + 1, 1);
|
|
if (logLen) clGetProgramBuildInfo(selProg, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
|
|
printf("info the select pass did not build (%s): %.300s; using the full read-back\n", clErrName(err), log);
|
|
free(log); clReleaseProgram(selProg); selProg = NULL; ((Options*)o)->readback = 1;
|
|
} else {
|
|
kSelect = clCreateKernel(selProg, "igneum_select", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_select");
|
|
dCount = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, 4, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer count");
|
|
dHits = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, IG_MAX_HITS * 2 * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer hits");
|
|
dSentinel = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, 34 * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer sentinel");
|
|
gMemCreated += 3;
|
|
{
|
|
size_t wg = 0;
|
|
if (clGetKernelWorkGroupInfo(kSelect, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL) == CL_SUCCESS && wg && wg < selLocal) selLocal = wg;
|
|
}
|
|
}
|
|
}
|
|
printf("info readback %s (per dispatch of %u nonces: %s)\n", o->readback ? "full" : "select",
|
|
batch, o->readback ? "8 bytes per nonce come back and the host scans them" : "the hits and 34 sentinel words come back; the GPU scans");
|
|
printf("ready opencl %s platform %s pack %s dataset-log2 %d batch %u exchange %d prepare %d path %s\n", devName, di->platformName,
|
|
gGeneric ? gPack.seedString : IGNEUM_SEED_STRING, log2u32(words), batch, dv->exchange, o->noPrepare ? 0 : 1, gGeneric ? "prebuilt-generic" : "compiled-in");
|
|
fflush(stdout);
|
|
|
|
/* Fault detection (4 October 2026, after the gfx1036 run that completed 2,000 jobs a second with no hash after
|
|
* 600 s): any OpenCL error in the job path is fatal (exit 3, the miner restarts the worker), the dispatch event
|
|
* must report CL_COMPLETE, a chunk that runs 20x faster per nonce than the running mean is a fault, and two
|
|
* consecutive chunks with the same output signature (different nonces or init words give different outputs) mean
|
|
* the kernel did not run. A stats line every 200 jobs shows the live event and buffer counts. */
|
|
long gFaultTestChunk = getenv("IGNEUM_FAULT_TEST") ? atol(getenv("IGNEUM_FAULT_TEST")) : -1;
|
|
double meanNsPerNonce = 0; /* running mean of wall ns per nonce over the chunks so far */
|
|
unsigned long chunksSeen = 0, jobsSeen = 0;
|
|
unsigned long long noncesSeen = 0; /* for the mean chunk in the stats line */
|
|
uint64_t prevSig = 0; int havePrevSig = 0;
|
|
double jobMsSum = 0;
|
|
#define SERVE_FATAL(jobIdStr, fmt, ...) do { printf("error %s worker fault: " fmt "; exiting 3 so the miner restarts the worker\n", jobIdStr, __VA_ARGS__); fflush(stdout); exit(3); } while (0)
|
|
#define SERVE_CHECK(jobIdStr, call) do { cl_int e_ = (call); if (e_ != CL_SUCCESS) SERVE_FATAL(jobIdStr, "%s failed with %s (%d)", #call, clErrName(e_), (int)e_); } while (0)
|
|
while (fgets(line, sizeof(line), stdin)) {
|
|
char* f[12]; /* job + 7 fields, plus the class= and era= tokens of a class v3 line */
|
|
int nf = 0;
|
|
char* tok;
|
|
char* save = NULL;
|
|
char jobId[64];
|
|
char wantClass[8], wantEra[65];
|
|
uint8_t prehash[32], epochSeed[32], daySeed[256];
|
|
size_t prehashLen = 0, epochLen = 0, dayLen = 0;
|
|
unsigned long long target = 0, nonceStart = 0, nonceCount = 0;
|
|
uint32_t sw[8], kw[8];
|
|
uint64_t remaining, hashes = 0;
|
|
uint32_t hi, lo;
|
|
double t0;
|
|
int failed = 0, switched = 0;
|
|
line[strcspn(line, "\r\n")] = 0;
|
|
/* A finished prepare is reported here, between lines (the thread never prints) */
|
|
if (task && task->done) {
|
|
if (task->result) {
|
|
ServePair* p = task->result;
|
|
{
|
|
uint8_t eb[32], db[256]; size_t el = 0, dl = 0;
|
|
if (unhexBuf(p->epochHex, eb, 32, &el) && el == 32) seedWordsFromBytes(eb, 32, p->sw);
|
|
if (unhexBuf(p->dayHex, db, sizeof(db), &dl)) seedWordsFromBytes(db, dl, p->kw);
|
|
}
|
|
if (prepared) releasePair(prepared);
|
|
prepared = p;
|
|
printf("prepared %s %s %.1f build %.1f cache %.1f dataset %.1f check %.1f resident 2 programs 2 datasets; %s\n", p->epochHex, p->dayHex, task->doneAt - task->t0, p->buildMs, p->cacheMs, p->datasetMs, p->checkMs, p->check);
|
|
} else {
|
|
printf("prepare-failed %s %s %s\n", task->epochHex, task->dayHex, task->error);
|
|
}
|
|
fflush(stdout);
|
|
free(task); task = NULL;
|
|
}
|
|
for (tok = strtok_r(line, " ", &save); tok && nf < 12; tok = strtok_r(NULL, " ", &save)) f[nf++] = tok;
|
|
if (nf == 0) continue;
|
|
if (strcmp(f[0], "quit") == 0) break;
|
|
takeClassTokens(f, &nf, wantClass, sizeof(wantClass), wantEra, sizeof(wantEra));
|
|
if (strcmp(f[0], "prepare") == 0) {
|
|
if (o->noPrepare) { printf("info ignored (started with --no-prepare): prepare\n"); fflush(stdout); continue; }
|
|
if (nf < 4) { printf("prepare-failed %s %s this ahead-of-time worker needs a pack directory as the third field (igneum-miner --prepare-packs <dir>)\n", nf > 1 ? f[1] : "0", nf > 2 ? f[2] : "0"); fflush(stdout); continue; }
|
|
if (strlen(f[1]) != 64 || strlen(f[2]) >= 500) { printf("prepare-failed %s %s bad field (epoch_seed 64 hex, day_seed hex)\n", f[1], f[2]); fflush(stdout); continue; }
|
|
if (task) { printf("prepare-failed %s %s a prepare is still running\n", f[1], f[2]); fflush(stdout); continue; }
|
|
if (prepared && strcmp(prepared->epochHex, f[1]) == 0 && strcmp(prepared->dayHex, f[2]) == 0) { printf("prepared %s %s 0 (already resident)\n", f[1], f[2]); fflush(stdout); continue; }
|
|
task = (PrepareTask*)calloc(1, sizeof(PrepareTask));
|
|
task->dv = dv; task->di = di; task->words = words; task->cacheWords = gServeCacheWords; task->segments = gServeSegments; task->t0 = wallMs();
|
|
strncpy(task->epochHex, f[1], 64); strncpy(task->dayHex, f[2], sizeof(task->dayHex) - 1); strncpy(task->packDir, f[3], sizeof(task->packDir) - 1);
|
|
strncpy(task->wantClass, wantClass, sizeof(task->wantClass) - 1); strncpy(task->wantEra, wantEra, sizeof(task->wantEra) - 1);
|
|
if (!startPrepareThread(task)) { printf("prepare-failed %s %s cannot start the prepare thread\n", f[1], f[2]); fflush(stdout); free(task); task = NULL; continue; }
|
|
printf("info prepare started for epoch %.16s day %s from %s (builds in the background)\n", f[1], f[2], f[3]); fflush(stdout);
|
|
continue;
|
|
}
|
|
if (strcmp(f[0], "job") != 0) { printf("info ignored line\n"); fflush(stdout); continue; }
|
|
strncpy(jobId, nf > 1 ? f[1] : "0", 63); jobId[63] = 0;
|
|
if (nf < 8) { printf("error %s malformed job line (need 7 fields after job)\n", jobId); fflush(stdout); continue; }
|
|
if (!unhexBuf(f[2], prehash, 32, &prehashLen) || prehashLen != 32 || sscanf(f[3], "%llx", &target) != 1 ||
|
|
sscanf(f[4], "%llu", &nonceStart) != 1 || sscanf(f[5], "%llu", &nonceCount) != 1 ||
|
|
!unhexBuf(f[6], epochSeed, 32, &epochLen) || epochLen != 32 || !unhexBuf(f[7], daySeed, sizeof(daySeed), &dayLen)) {
|
|
printf("error %s bad field (prehash 64 hex, target 16 hex, nonce_start u64, nonce_count u64, epoch_seed 64 hex, day_seed hex)\n", jobId); fflush(stdout); continue;
|
|
}
|
|
if (nonceCount == 0 || nonceCount % 32 != 0 || (nonceStart & 31) != 0) { printf("error %s nonce_start must be 32-aligned and nonce_count a non-zero multiple of 32\n", jobId); fflush(stdout); continue; }
|
|
seedWordsFromBytes(epochSeed, 32, sw);
|
|
seedWordsFromBytes(daySeed, dayLen, kw);
|
|
t0 = wallMs();
|
|
if (!pairIsClass(cur, f[6], f[7], sw, kw, wantClass, wantEra)) {
|
|
char why[256];
|
|
if (pairIsClass(prepared, f[6], f[7], sw, kw, wantClass, wantEra)) {
|
|
/* The prepared pair: switch now, release the old one after this job */
|
|
if (old) releasePair(old);
|
|
old = cur; cur = prepared; prepared = NULL; switched = 1;
|
|
printf("info switched to the prepared pair epoch %.16s day %s (class %s) in %.2f ms\n", cur->epochHex, cur->dayHex, cur->programClass[0] ? cur->programClass : "v2", wallMs() - t0); fflush(stdout);
|
|
} else if (pairIs(cur, f[6], f[7], sw, kw) && !pf_pack_class_ok(cur->programClass[0] ? cur->programClass : "v2", cur->eraHex, wantClass, wantEra, why, sizeof(why))) {
|
|
/* Counter ASIC 2.0: right seeds, wrong class or era; the miner prepares the pair from a pack of the
|
|
* class the chain is on (the need line), and this pack is never mined */
|
|
printf("need %s %s\n", f[6], f[7]);
|
|
printf("error %s pack %s: %s\n", jobId, o->packDir ? o->packDir : "(compiled-in)", why);
|
|
fflush(stdout); continue;
|
|
} else if (cur->epochHex[0] ? !hexEq(cur->epochHex, f[6]) : memcmp(sw, cur->sw, 32) != 0) {
|
|
printf("need %s %s\n", f[6], f[7]); /* the miner prepares this pair (4 October 2026) */
|
|
printf("error %s epoch seed mismatch: this worker holds %s%s (program words %08x %08x ...)%s, the job is for epoch %.16s (bare seed words %08x %08x ...); send prepare with a pack directory, or run igneum-miner export-pack and rebuild\n",
|
|
jobId, cur->epochHex[0] ? "prepared epoch " : "pack \"" IGNEUM_SEED_STRING "\"", cur->epochHex[0] ? cur->epochHex : "", cur->sw[0], cur->sw[1], prepared ? " plus one prepared pair" : "", f[6], sw[0], sw[1]);
|
|
fflush(stdout); continue;
|
|
} else {
|
|
printf("need %s %s\n", f[6], f[7]);
|
|
printf("error %s day seed mismatch: this worker's cache is for key %08x %08x ..., the job's day seed %s gives %08x %08x ...; send prepare with a pack directory, or run igneum-miner export-pack and rebuild\n",
|
|
jobId, cur->kw[0], cur->kw[1], f[7], kw[0], kw[1]);
|
|
fflush(stdout); continue;
|
|
}
|
|
}
|
|
remaining = nonceCount;
|
|
hi = (uint32_t)(nonceStart >> 32); lo = (uint32_t)nonceStart;
|
|
while (remaining > 0) {
|
|
uint64_t room = (uint64_t)(0xffffffffu - lo) + 1ull;
|
|
uint64_t chunk64 = remaining < batch ? remaining : batch;
|
|
uint32_t chunk, i;
|
|
uint32_t iw[8];
|
|
uint8_t b[49];
|
|
cl_uint baseNonce, maskArg = mask;
|
|
cl_event ev;
|
|
if (chunk64 > room) chunk64 = room;
|
|
chunk = (uint32_t)chunk64;
|
|
memcpy(b, "igneum-block/", 13);
|
|
memcpy(b + 13, prehash, 32);
|
|
b[45] = (uint8_t)hi; b[46] = (uint8_t)(hi >> 8); b[47] = (uint8_t)(hi >> 16); b[48] = (uint8_t)(hi >> 24);
|
|
seedWordsFromBytes(b, 49, iw);
|
|
{
|
|
double c0 = wallMs(), cms, kernelMs = -1.0, selectMs = 0.0, readMs = 0.0, scanMs = 0.0, r0;
|
|
cl_int status = 0;
|
|
size_t g = ((chunk + groupSize - 1) / groupSize) * groupSize;
|
|
uint64_t sig;
|
|
int useSelect = (kSelect != NULL);
|
|
uint32_t nHits = 0;
|
|
SERVE_CHECK(jobId, clEnqueueWriteBuffer(dv->q, dInit, CL_TRUE, 0, 32, iw, 0, NULL, NULL));
|
|
bytesUp += 32;
|
|
if (useSelect) { hCount = 0; SERVE_CHECK(jobId, clEnqueueWriteBuffer(dv->q, dCount, CL_TRUE, 0, 4, &hCount, 0, NULL, NULL)); bytesUp += 4; }
|
|
baseNonce = lo;
|
|
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 0, sizeof(cl_mem), &cur->ds));
|
|
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 1, sizeof(cl_mem), &dOut));
|
|
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 2, sizeof(cl_uint), &baseNonce));
|
|
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 3, sizeof(cl_uint), &maskArg));
|
|
SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 4, sizeof(cl_mem), &dInit));
|
|
if (cur->hot) SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 5, sizeof(cl_mem), &cur->hot));
|
|
ev = NULL;
|
|
if (gFaultTestChunk >= 0 && (long)chunksSeen >= gFaultTestChunk) {
|
|
/* IGNEUM_FAULT_TEST=N (test only): from chunk N on, behave like a runtime that answers every call with
|
|
* success and runs nothing, so the detectors below are exercised on a healthy device. */
|
|
} else {
|
|
SERVE_CHECK(jobId, clEnqueueNDRangeKernel(dv->q, cur->kHashBound, 1, NULL, &g, &groupSize, 0, NULL, &ev));
|
|
++gEvCreated;
|
|
/* Wait on the dispatch event, not clFinish: the runtime can sleep the thread on an event, where clFinish
|
|
* on some drivers spins one core for the whole dispatch. The event must then say CL_COMPLETE: a runtime
|
|
* that has dropped its device answers the wait at once with a negative status. */
|
|
err = clWaitForEvents(1, &ev);
|
|
if (err != CL_SUCCESS) { countRelease(ev); SERVE_FATAL(jobId, "clWaitForEvents on the dispatch returned %s (%d)", clErrName(err), (int)err); }
|
|
err = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(status), &status, NULL);
|
|
kernelMs = eventMs(ev);
|
|
countRelease(ev);
|
|
if (err != CL_SUCCESS) SERVE_FATAL(jobId, "clGetEventInfo on the dispatch returned %s (%d)", clErrName(err), (int)err);
|
|
if (status != CL_COMPLETE) SERVE_FATAL(jobId, "the dispatch event ended with status %d, not CL_COMPLETE (a device reset or a lost context)", (int)status);
|
|
if (useSelect) {
|
|
cl_event evs = NULL;
|
|
cl_uint nArg = chunk, maxArg = IG_MAX_HITS;
|
|
cl_ulong tArg = (cl_ulong)target;
|
|
size_t gs = ((chunk + selLocal - 1) / selLocal) * selLocal;
|
|
SERVE_CHECK(jobId, clSetKernelArg(kSelect, 0, sizeof(cl_mem), &dOut));
|
|
SERVE_CHECK(jobId, clSetKernelArg(kSelect, 1, sizeof(cl_uint), &nArg));
|
|
SERVE_CHECK(jobId, clSetKernelArg(kSelect, 2, sizeof(cl_ulong), &tArg));
|
|
SERVE_CHECK(jobId, clSetKernelArg(kSelect, 3, sizeof(cl_mem), &dCount));
|
|
SERVE_CHECK(jobId, clSetKernelArg(kSelect, 4, sizeof(cl_mem), &dHits));
|
|
SERVE_CHECK(jobId, clSetKernelArg(kSelect, 5, sizeof(cl_uint), &maxArg));
|
|
SERVE_CHECK(jobId, clSetKernelArg(kSelect, 6, sizeof(cl_mem), &dSentinel));
|
|
SERVE_CHECK(jobId, clEnqueueNDRangeKernel(dv->q, kSelect, 1, NULL, &gs, &selLocal, 0, NULL, &evs));
|
|
++gEvCreated;
|
|
err = clWaitForEvents(1, &evs);
|
|
if (err != CL_SUCCESS) { countRelease(evs); SERVE_FATAL(jobId, "clWaitForEvents on the select pass returned %s (%d)", clErrName(err), (int)err); }
|
|
selectMs = eventMs(evs);
|
|
countRelease(evs);
|
|
}
|
|
}
|
|
r0 = wallMs();
|
|
if (useSelect) {
|
|
SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dCount, CL_TRUE, 0, 4, &hCount, 0, NULL, NULL));
|
|
bytesDown += 4;
|
|
nHits = hCount;
|
|
if (nHits > IG_MAX_HITS) {
|
|
/* more hits than the table holds: this chunk takes the full path (the test target case) */
|
|
++fullFallbacks;
|
|
useSelect = 0;
|
|
} else {
|
|
if (nHits) { SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dHits, CL_TRUE, 0, (size_t)nHits * 2 * sizeof(uint64_t), hHits, 0, NULL, NULL)); bytesDown += (unsigned long long)nHits * 16; }
|
|
SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dSentinel, CL_TRUE, 0, 34 * sizeof(uint64_t), hSentinel, 0, NULL, NULL));
|
|
bytesDown += 34 * 8;
|
|
}
|
|
}
|
|
if (!useSelect) {
|
|
SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)chunk * sizeof(uint64_t), hOut, 0, NULL, NULL));
|
|
bytesDown += (unsigned long long)chunk * 8;
|
|
}
|
|
readMs = wallMs() - r0;
|
|
cms = wallMs() - c0;
|
|
/* Plausibility: wall time per nonce against the running mean (the first chunk sets it; a chunk is a full
|
|
* batch except at the 32-bit boundary, so per nonce is the comparable unit). 20x faster = the kernel did not run. */
|
|
{
|
|
double ns = cms * 1e6 / (double)chunk;
|
|
if (chunksSeen >= 4 && meanNsPerNonce > 0 && ns * 20.0 < meanNsPerNonce)
|
|
SERVE_FATAL(jobId, "%u nonces reported complete in %.3f ms, %.0fx faster than the running mean of %.2f ms per million (the runtime is not running the kernel)", chunk, cms, meanNsPerNonce / ns, meanNsPerNonce / 1e3);
|
|
meanNsPerNonce = chunksSeen == 0 ? ns : meanNsPerNonce + (ns - meanNsPerNonce) / (double)(chunksSeen + 1 < 64 ? chunksSeen + 1 : 64);
|
|
++chunksSeen; noncesSeen += chunk;
|
|
}
|
|
/* Stale output: the first, middle and last words plus an FNV of the first 32. Two chunks never share it.
|
|
* On the select path the same 34 words come from the sentinel buffer the select pass wrote. */
|
|
if (useSelect) sig = fnv1a64(hSentinel, 32 * sizeof(uint64_t)) ^ hSentinel[32] ^ hSentinel[33];
|
|
else sig = fnv1a64(hOut, 32 * sizeof(uint64_t)) ^ hOut[chunk / 2] ^ hOut[chunk - 1];
|
|
if (chunk >= 64) {
|
|
if (havePrevSig && sig == prevSig) SERVE_FATAL(jobId, "the output buffer is unchanged since the previous dispatch (signature %016llx): the kernel did not run", (unsigned long long)sig);
|
|
prevSig = sig; havePrevSig = 1;
|
|
}
|
|
r0 = wallMs();
|
|
if (useSelect) {
|
|
/* nonce order, as the full scan printed them (the miner submits the first found of a job) */
|
|
uint32_t a, b;
|
|
for (a = 1; a < nHits; ++a) {
|
|
uint64_t ki = hHits[2 * a], kh = hHits[2 * a + 1];
|
|
for (b = a; b > 0 && hHits[2 * (b - 1)] > ki; --b) { hHits[2 * b] = hHits[2 * (b - 1)]; hHits[2 * b + 1] = hHits[2 * (b - 1) + 1]; }
|
|
hHits[2 * b] = ki; hHits[2 * b + 1] = kh;
|
|
}
|
|
for (a = 0; a < nHits; ++a) {
|
|
uint32_t idx = (uint32_t)hHits[2 * a];
|
|
unsigned long long nonce = ((unsigned long long)hi << 32) | (unsigned long long)(uint32_t)(lo + idx);
|
|
printf("found %s %llu %016llx\n", jobId, nonce, (unsigned long long)hHits[2 * a + 1]);
|
|
}
|
|
} else {
|
|
for (i = 0; i < chunk; ++i) if (hOut[i] <= target) {
|
|
unsigned long long nonce = ((unsigned long long)hi << 32) | (unsigned long long)(uint32_t)(lo + i);
|
|
printf("found %s %llu %016llx\n", jobId, nonce, (unsigned long long)hOut[i]);
|
|
}
|
|
}
|
|
scanMs = wallMs() - r0;
|
|
if (kernelMs > 0) kernelMsSum += kernelMs;
|
|
selectMsSum += selectMs; readMsSum += readMs; scanMsSum += scanMs;
|
|
}
|
|
fflush(stdout);
|
|
hashes += chunk;
|
|
remaining -= chunk;
|
|
if (chunk64 == room) { hi += 1u; lo = 0u; } else lo += chunk;
|
|
}
|
|
if (failed) continue;
|
|
printf("done %s %llu %.2f\n", jobId, (unsigned long long)hashes, wallMs() - t0);
|
|
jobMsSum += wallMs() - t0;
|
|
++jobsSeen;
|
|
if (jobsSeen % 200 == 0) {
|
|
printf("info stats jobs %lu chunks %lu mean job %.1f ms mean %.2f ms per million nonces; events created %lu released %lu live %lu; buffers created %lu released %lu live %lu\n",
|
|
jobsSeen, chunksSeen, jobMsSum / (double)jobsSeen, meanNsPerNonce / 1e3, gEvCreated, gEvReleased, gEvCreated - gEvReleased, gMemCreated, gMemReleased, gMemCreated - gMemReleased);
|
|
/* The transfer and time budget per chunk (5 October 2026): bytes up (init words, the count reset), bytes
|
|
* down (hits and sentinels, or 8 bytes per nonce on the full path), and the mean device time of the hash
|
|
* kernel (event profiling), the select pass, the blocking read-back and the host scan. */
|
|
printf("info transfers per chunk: up %.0f B, down %.0f B (%s, %lu full fallbacks); mean per chunk: kernel %.2f ms (device), select %.3f ms, read-back %.2f ms, scan %.2f ms, chunk wall %.2f ms\n",
|
|
(double)bytesUp / (double)chunksSeen, (double)bytesDown / (double)chunksSeen, kSelect ? "select" : "full", fullFallbacks,
|
|
kernelMsSum / (double)chunksSeen, selectMsSum / (double)chunksSeen, readMsSum / (double)chunksSeen, scanMsSum / (double)chunksSeen, meanNsPerNonce * ((double)noncesSeen / (double)chunksSeen) / 1e6);
|
|
if (gEvCreated - gEvReleased > 16 || gMemCreated - gMemReleased > 12) SERVE_FATAL(jobId, "object leak: %lu events and %lu buffers live after %lu jobs", gEvCreated - gEvReleased, gMemCreated - gMemReleased, jobsSeen);
|
|
}
|
|
if (switched && old) { releasePair(old); old = NULL; printf("info dropped the previous pair (its program, cache and dataset)\n"); }
|
|
fflush(stdout);
|
|
}
|
|
free(hOut);
|
|
clReleaseMemObject(dInit);
|
|
clReleaseMemObject(dOut);
|
|
gMemReleased += 2;
|
|
if (kSelect) { clReleaseKernel(kSelect); clReleaseMemObject(dCount); clReleaseMemObject(dHits); clReleaseMemObject(dSentinel); gMemReleased += 3; }
|
|
if (selProg) clReleaseProgram(selProg);
|
|
if (chunksSeen) printf("info transfers total: %lu chunks, up %llu B, down %llu B (%s, %lu full fallbacks)\n", chunksSeen, bytesUp, bytesDown, kSelect ? "select" : "full", fullFallbacks);
|
|
if (old) releasePair(old);
|
|
if (prepared) releasePair(prepared);
|
|
releasePair(cur);
|
|
return 0;
|
|
#endif
|
|
}
|
|
|
|
/* ---------------------------------------------------------------------------------------------
|
|
* --memprobe (5 October 2026): what the device itself can do with the access pattern of the hash, with no pack and no
|
|
* program. Three kernels built from the text below:
|
|
* chase one dependent random 4-byte load per step per lane (the next address comes from the loaded word), so the
|
|
* time per step at a small lane count is the loaded-latency of one random read, and the loads per second at a
|
|
* large lane count is the device's random-read throughput for a dependent chain (the hash is 128 of these)
|
|
* indep eight independent chains per lane: the throughput when latency is hidden inside one lane
|
|
* alu an integer multiply-add-rotate chain with no memory: the achieved integer rate, which moves with the clock
|
|
* Sizes 4 MiB (cache-resident), 64 MiB (the last-level cache on RDNA 3 and 4 is 64 MiB or more, approximate) and
|
|
* 1024 MiB (the dataset size: DRAM). Times from event profiling (wall on Apple). Every number is printed with the
|
|
* configuration that produced it. Rates are G loads/s = 1e9 loads per second; ns per load = time / steps.
|
|
*/
|
|
static const char* PROBE_SRC =
|
|
"static inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
|
|
"__kernel void probe_fill(__global uint* ds, uint n) { uint i = (uint)get_global_id(0); if (i < n) ds[i] = pm_mix(i ^ 0x9E3779B9u); }\n"
|
|
"__kernel void probe_chase(__global const uint* ds, uint mask, uint steps, uint seed, __global uint* out) {\n"
|
|
" uint x = pm_mix((uint)get_global_id(0) ^ seed);\n"
|
|
" for (uint s = 0u; s < steps; ++s) x = ds[x & mask] ^ (x * 0x9E3779B1u + s);\n"
|
|
" out[get_global_id(0)] = x;\n"
|
|
"}\n"
|
|
"__kernel void probe_indep(__global const uint* ds, uint mask, uint steps, uint seed, __global uint* out) {\n"
|
|
" uint g = (uint)get_global_id(0);\n"
|
|
" uint x0 = pm_mix(g * 8u ^ seed), x1 = pm_mix((g * 8u + 1u) ^ seed), x2 = pm_mix((g * 8u + 2u) ^ seed), x3 = pm_mix((g * 8u + 3u) ^ seed);\n"
|
|
" uint x4 = pm_mix((g * 8u + 4u) ^ seed), x5 = pm_mix((g * 8u + 5u) ^ seed), x6 = pm_mix((g * 8u + 6u) ^ seed), x7 = pm_mix((g * 8u + 7u) ^ seed);\n"
|
|
" for (uint s = 0u; s < steps; ++s) {\n"
|
|
" x0 = ds[x0 & mask] ^ (x0 * 0x9E3779B1u + s); x1 = ds[x1 & mask] ^ (x1 * 0x9E3779B1u + s);\n"
|
|
" x2 = ds[x2 & mask] ^ (x2 * 0x9E3779B1u + s); x3 = ds[x3 & mask] ^ (x3 * 0x9E3779B1u + s);\n"
|
|
" x4 = ds[x4 & mask] ^ (x4 * 0x9E3779B1u + s); x5 = ds[x5 & mask] ^ (x5 * 0x9E3779B1u + s);\n"
|
|
" x6 = ds[x6 & mask] ^ (x6 * 0x9E3779B1u + s); x7 = ds[x7 & mask] ^ (x7 * 0x9E3779B1u + s);\n"
|
|
" }\n"
|
|
" out[g] = x0 ^ x1 ^ x2 ^ x3 ^ x4 ^ x5 ^ x6 ^ x7;\n"
|
|
"}\n"
|
|
"__kernel void probe_line16(__global const uint4* ds, uint vecMask, uint steps, uint seed, __global uint* out) {\n"
|
|
" uint x = pm_mix((uint)get_global_id(0) ^ seed);\n"
|
|
" for (uint s = 0u; s < steps; ++s) { uint4 a = ds[x & vecMask]; x = (a.x ^ a.y ^ a.z ^ a.w) ^ (x * 0x9E3779B1u + s); }\n"
|
|
" out[get_global_id(0)] = x;\n"
|
|
"}\n"
|
|
"__kernel void probe_line(__global const uint4* ds, uint lineMask, uint steps, uint seed, __global uint* out) {\n"
|
|
" uint x = pm_mix((uint)get_global_id(0) ^ seed);\n"
|
|
" for (uint s = 0u; s < steps; ++s) {\n"
|
|
" uint l = (x & lineMask) * 4u;\n"
|
|
" uint4 a = ds[l], b = ds[l + 1u], c = ds[l + 2u], d = ds[l + 3u];\n"
|
|
" x = (a.x ^ b.y ^ c.z ^ d.w) ^ (x * 0x9E3779B1u + s);\n"
|
|
" }\n"
|
|
" out[get_global_id(0)] = x;\n"
|
|
"}\n"
|
|
"__kernel void probe_stream(__global const uint4* ds, uint perLane, __global uint* out) {\n"
|
|
" uint g = (uint)get_global_id(0), n = (uint)get_global_size(0);\n"
|
|
" uint4 acc = (uint4)(0u, 0u, 0u, 0u);\n"
|
|
" for (uint s = 0u; s < perLane; ++s) acc ^= ds[s * n + g];\n"
|
|
" out[g] = acc.x ^ acc.y ^ acc.z ^ acc.w;\n"
|
|
"}\n"
|
|
"__kernel void probe_alu(uint steps, uint seed, __global uint* out) {\n"
|
|
" uint g = (uint)get_global_id(0); uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n"
|
|
" for (uint s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + rotate(y, 7u); y = (y ^ x) + s; }\n"
|
|
" out[g] = x ^ y;\n"
|
|
"}\n";
|
|
|
|
/* One timed launch, best of `reps`, in ms (event time, or wall when the platform's events are unusable). */
|
|
static double probeLaunch(Device* dv, const Options* o, cl_kernel k, size_t global, size_t local, int reps, int seedArg, cl_uint seed) {
|
|
double best = -1.0;
|
|
int r;
|
|
for (r = 0; r < reps; ++r) {
|
|
cl_event ev = NULL;
|
|
double w0 = wallMs(), ms;
|
|
/* a fresh seed per repetition: a replay of the same addresses would be served from the last-level cache
|
|
* (65,536 reads x 64 B = 4 MiB fits any of them) and read as DRAM latency (seen on the 9070 XT, 5 October) */
|
|
if (seedArg >= 0) { cl_uint s = seed + (cl_uint)r * 0x9E3779B9u; CL_CHECK(clSetKernelArg(k, (cl_uint)seedArg, sizeof(cl_uint), &s)); }
|
|
CL_CHECK(clEnqueueNDRangeKernel(dv->q, k, 1, NULL, &global, &local, 0, NULL, &ev));
|
|
CL_CHECK(clWaitForEvents(1, &ev));
|
|
ms = o->timeWall ? wallMs() - w0 : eventMs(ev);
|
|
if (ms < 0) ms = wallMs() - w0;
|
|
clReleaseEvent(ev);
|
|
if (best < 0 || ms < best) best = ms;
|
|
}
|
|
return best;
|
|
}
|
|
|
|
static int runMemprobe(Device* dv, const DeviceInfo* di, const Options* o) {
|
|
cl_int err = 0;
|
|
cl_program prog;
|
|
cl_kernel kFill, kChase, kIndep, kAlu, kLine, kStream, kLine16;
|
|
size_t srcLen = strlen(PROBE_SRC);
|
|
int sizes[3] = { 4, 64, 1024 }, nSizes = 3, si;
|
|
size_t lanesList[8] = { 256, 1024, 1u << 12, 1u << 14, 1u << 16, 1u << 18, 1u << 20, 1u << 22 };
|
|
const int nLanes = 8;
|
|
size_t groups[2] = { 32, 256 };
|
|
const cl_uint STEPS = 256u, ALU_STEPS = 4096u;
|
|
const size_t maxLanes = 1u << 22;
|
|
cl_mem dOut;
|
|
if (o->probeMib > 0) { sizes[0] = o->probeMib; nSizes = 1; }
|
|
prog = clCreateProgramWithSource(dv->ctx, 1, &PROBE_SRC, &srcLen, &err); CL_CHECK_ERR(err, "clCreateProgramWithSource probe");
|
|
err = clBuildProgram(prog, 1, &di->device, "-cl-std=CL1.2", NULL, NULL);
|
|
if (err != CL_SUCCESS) {
|
|
size_t logLen = 0; char* log;
|
|
clGetProgramBuildInfo(prog, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
|
|
log = (char*)calloc(logLen + 1, 1);
|
|
if (logLen) clGetProgramBuildInfo(prog, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
|
|
printf("memprobe: build FAILED (%s)\n--- build log ---\n%s\n--- end of build log ---\n", clErrName(err), log);
|
|
free(log);
|
|
return 2;
|
|
}
|
|
kFill = clCreateKernel(prog, "probe_fill", &err); CL_CHECK_ERR(err, "probe_fill");
|
|
kChase = clCreateKernel(prog, "probe_chase", &err); CL_CHECK_ERR(err, "probe_chase");
|
|
kIndep = clCreateKernel(prog, "probe_indep", &err); CL_CHECK_ERR(err, "probe_indep");
|
|
kAlu = clCreateKernel(prog, "probe_alu", &err); CL_CHECK_ERR(err, "probe_alu");
|
|
kLine = clCreateKernel(prog, "probe_line", &err); CL_CHECK_ERR(err, "probe_line");
|
|
kLine16 = clCreateKernel(prog, "probe_line16", &err); CL_CHECK_ERR(err, "probe_line16");
|
|
kStream = clCreateKernel(prog, "probe_stream", &err); CL_CHECK_ERR(err, "probe_stream");
|
|
dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, maxLanes * 4u, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer probe out");
|
|
printf("memprobe on [%s] %s, driver %s, %u compute units, %u MHz, %s time\n", di->platformName, di->name, di->driver, di->computeUnits, di->clockMHz, o->timeWall ? "wall" : "device event");
|
|
printKernelInfo(di, kChase, "probe_chase", 32, "memprobe ");
|
|
printKernelInfo(di, kChase, "probe_chase", 256, "memprobe ");
|
|
printf("| probe | MiB | work-group | lanes in flight | steps per lane | best ms | G loads/s | ns per dependent load |\n|---|---|---|---|---|---|---|---|\n");
|
|
for (si = 0; si < nSizes; ++si) {
|
|
int mib = sizes[si];
|
|
uint64_t bytes = (uint64_t)mib << 20;
|
|
cl_uint words = (cl_uint)(bytes / 4ull), mask = words - 1u, n = words;
|
|
cl_mem dDs;
|
|
size_t gi, li;
|
|
size_t fillLocal = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
|
|
size_t fillGlobal = ((size_t)words + fillLocal - 1) / fillLocal * fillLocal;
|
|
if ((uint64_t)di->maxAlloc < bytes) { printf("| chase | %d | skipped: max alloc %llu MiB | | | | | |\n", mib, (unsigned long long)(di->maxAlloc >> 20)); continue; }
|
|
dDs = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)bytes, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer probe dataset");
|
|
CL_CHECK(clSetKernelArg(kFill, 0, sizeof(cl_mem), &dDs));
|
|
CL_CHECK(clSetKernelArg(kFill, 1, sizeof(cl_uint), &n));
|
|
probeLaunch(dv, o, kFill, fillGlobal, fillLocal, 1, -1, 0u);
|
|
for (gi = 0; gi < 2; ++gi) {
|
|
size_t local = groups[gi];
|
|
if (local > di->maxWorkGroup) continue;
|
|
for (li = 0; li < (size_t)nLanes; ++li) {
|
|
size_t lanes = lanesList[li];
|
|
cl_uint seed = (cl_uint)(0x1234567u + (cl_uint)li * 977u);
|
|
double ms;
|
|
if (lanes < local) continue;
|
|
CL_CHECK(clSetKernelArg(kChase, 0, sizeof(cl_mem), &dDs));
|
|
CL_CHECK(clSetKernelArg(kChase, 1, sizeof(cl_uint), &mask));
|
|
CL_CHECK(clSetKernelArg(kChase, 2, sizeof(cl_uint), &STEPS));
|
|
CL_CHECK(clSetKernelArg(kChase, 3, sizeof(cl_uint), &seed));
|
|
CL_CHECK(clSetKernelArg(kChase, 4, sizeof(cl_mem), &dOut));
|
|
ms = probeLaunch(dv, o, kChase, lanes, local, 3, 3, seed);
|
|
printf("| chase | %d | %llu | %llu | %u | %.3f | %.3f | %.0f |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
|
|
(double)lanes * (double)STEPS / (ms / 1000.0) / 1e9, ms * 1e6 / (double)STEPS);
|
|
fflush(stdout);
|
|
}
|
|
}
|
|
{
|
|
size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
|
|
size_t lanes;
|
|
for (lanes = 1u << 16; lanes <= maxLanes; lanes <<= 2) {
|
|
cl_uint seed = 0x7654321u;
|
|
double ms;
|
|
CL_CHECK(clSetKernelArg(kIndep, 0, sizeof(cl_mem), &dDs));
|
|
CL_CHECK(clSetKernelArg(kIndep, 1, sizeof(cl_uint), &mask));
|
|
CL_CHECK(clSetKernelArg(kIndep, 2, sizeof(cl_uint), &STEPS));
|
|
CL_CHECK(clSetKernelArg(kIndep, 3, sizeof(cl_uint), &seed));
|
|
CL_CHECK(clSetKernelArg(kIndep, 4, sizeof(cl_mem), &dOut));
|
|
ms = probeLaunch(dv, o, kIndep, lanes, local, 3, 3, seed);
|
|
printf("| indep x8 | %d | %llu | %llu | %u | %.3f | %.3f | (8 loads in flight per lane) |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
|
|
(double)lanes * 8.0 * (double)STEPS / (ms / 1000.0) / 1e9);
|
|
fflush(stdout);
|
|
}
|
|
}
|
|
{
|
|
/* Random 16-byte reads (one uint4) in a dependent chain: the W = 16 width of the read-width experiment. */
|
|
size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
|
|
size_t lanes;
|
|
cl_uint vecMask = (words / 4u) - 1u;
|
|
for (lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
|
|
cl_uint seed = 0x2718281u;
|
|
double ms;
|
|
CL_CHECK(clSetKernelArg(kLine16, 0, sizeof(cl_mem), &dDs));
|
|
CL_CHECK(clSetKernelArg(kLine16, 1, sizeof(cl_uint), &vecMask));
|
|
CL_CHECK(clSetKernelArg(kLine16, 2, sizeof(cl_uint), &STEPS));
|
|
CL_CHECK(clSetKernelArg(kLine16, 3, sizeof(cl_uint), &seed));
|
|
CL_CHECK(clSetKernelArg(kLine16, 4, sizeof(cl_mem), &dOut));
|
|
ms = probeLaunch(dv, o, kLine16, lanes, local, 3, 3, seed);
|
|
printf("| line 16 B | %d | %llu | %llu | %u | %.3f | %.3f G reads/s | %.1f GB/s in 16 B reads |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
|
|
(double)lanes * (double)STEPS / (ms / 1000.0) / 1e9, (double)lanes * (double)STEPS * 16.0 / (ms / 1000.0) / 1e9);
|
|
fflush(stdout);
|
|
}
|
|
}
|
|
{
|
|
/* Random 64-byte lines (16 words, four uint4 loads) in a dependent chain: lines per second against the
|
|
* 4-byte chase above says what one random 4-byte read costs the memory system. If the two rates are
|
|
* equal, every 4-byte read fetches a whole line; if lines/s is a quarter of loads/s, reads cost a 16-byte
|
|
* sector. */
|
|
size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
|
|
size_t lanes;
|
|
cl_uint lineMask = (words / 16u) - 1u;
|
|
for (lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
|
|
cl_uint seed = 0x3141592u;
|
|
double ms;
|
|
CL_CHECK(clSetKernelArg(kLine, 0, sizeof(cl_mem), &dDs));
|
|
CL_CHECK(clSetKernelArg(kLine, 1, sizeof(cl_uint), &lineMask));
|
|
CL_CHECK(clSetKernelArg(kLine, 2, sizeof(cl_uint), &STEPS));
|
|
CL_CHECK(clSetKernelArg(kLine, 3, sizeof(cl_uint), &seed));
|
|
CL_CHECK(clSetKernelArg(kLine, 4, sizeof(cl_mem), &dOut));
|
|
ms = probeLaunch(dv, o, kLine, lanes, local, 3, 3, seed);
|
|
printf("| line 64 B | %d | %llu | %llu | %u | %.3f | %.3f G lines/s | %.1f GB/s in lines |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
|
|
(double)lanes * (double)STEPS / (ms / 1000.0) / 1e9, (double)lanes * (double)STEPS * 64.0 / (ms / 1000.0) / 1e9);
|
|
fflush(stdout);
|
|
}
|
|
}
|
|
{
|
|
/* Coalesced read of the whole buffer (uint4 per lane per step, consecutive lanes consecutive addresses):
|
|
* the sequential bandwidth. Against the card's rated figure this says whether the memory clock is in its
|
|
* full state; a card parked in a middle memory state shows about half (approximate). */
|
|
size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
|
|
size_t lanes = 1u << 20;
|
|
cl_uint perLane = (cl_uint)((uint64_t)words / 4ull / (uint64_t)lanes);
|
|
double ms, bytes = (double)perLane * (double)lanes * 16.0;
|
|
if (perLane == 0) { perLane = 1; lanes = (size_t)words / 4u; bytes = (double)lanes * 16.0; }
|
|
CL_CHECK(clSetKernelArg(kStream, 0, sizeof(cl_mem), &dDs));
|
|
CL_CHECK(clSetKernelArg(kStream, 1, sizeof(cl_uint), &perLane));
|
|
CL_CHECK(clSetKernelArg(kStream, 2, sizeof(cl_mem), &dOut));
|
|
ms = probeLaunch(dv, o, kStream, lanes, local, 3, -1, 0u);
|
|
printf("| stream | %d | %llu | %llu | %u | %.3f | %.1f GB/s coalesced | (%.0f MiB read once) |\n", mib, (unsigned long long)local, (unsigned long long)lanes, perLane, ms, bytes / (ms / 1000.0) / 1e9, bytes / 1048576.0);
|
|
fflush(stdout);
|
|
}
|
|
clReleaseMemObject(dDs);
|
|
}
|
|
{
|
|
size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
|
|
size_t lanes = 1u << 20;
|
|
cl_uint seed = 0x2468aceu;
|
|
double ms, ops;
|
|
CL_CHECK(clSetKernelArg(kAlu, 0, sizeof(cl_uint), &ALU_STEPS));
|
|
CL_CHECK(clSetKernelArg(kAlu, 1, sizeof(cl_uint), &seed));
|
|
CL_CHECK(clSetKernelArg(kAlu, 2, sizeof(cl_mem), &dOut));
|
|
ms = probeLaunch(dv, o, kAlu, lanes, local, 3, 1, seed);
|
|
ops = (double)lanes * (double)ALU_STEPS * 5.0; /* mul, add, rotate, xor, add per step */
|
|
printf("| alu | 0 | %llu | %llu | %u | %.3f | %.1f G int ops/s | %.3f G steps/s per compute unit (approximate: 5 ops per step counted) |\n",
|
|
(unsigned long long)local, (unsigned long long)lanes, ALU_STEPS, ms, ops / (ms / 1000.0) / 1e9,
|
|
(double)lanes * (double)ALU_STEPS / (ms / 1000.0) / 1e9 / (double)(di->computeUnits ? di->computeUnits : 1));
|
|
}
|
|
clReleaseMemObject(dOut);
|
|
clReleaseKernel(kFill); clReleaseKernel(kChase); clReleaseKernel(kIndep); clReleaseKernel(kAlu); clReleaseKernel(kLine); clReleaseKernel(kStream); clReleaseKernel(kLine16);
|
|
clReleaseProgram(prog);
|
|
printf("memprobe: done\n");
|
|
return 0;
|
|
}
|
|
|
|
int main(int argc, char** argv) {
|
|
Options o = parseArgs(argc, argv);
|
|
DeviceInfo* devs = NULL;
|
|
int nDev, i, chosen = -1;
|
|
DeviceInfo* di;
|
|
Device dv;
|
|
cl_int err = 0;
|
|
size_t srcLen = 0;
|
|
char* src;
|
|
uint32_t nonces;
|
|
cl_mem dOut;
|
|
int sizes[5], nSizes = 0;
|
|
SizeResult results[5];
|
|
int cachePass = 1, overall, anyVec = 0;
|
|
|
|
#ifdef IGNEUM_CL_DYNAMIC
|
|
if (!ig_cl_load()) { printf("%s %s\n", o.serve ? "error 0" : "FAIL:", ig_cl_error); fflush(stdout); return 2; }
|
|
#endif
|
|
if (o.packDir) {
|
|
/* Generic serve mode: the pack directory replaces the compiled-in pack (4 October 2026) */
|
|
static char boundPath[1200];
|
|
char perr[512];
|
|
size_t n = strlen(o.packDir);
|
|
if (!o.serve && !o.benchPack) { printf("FAIL: --pack goes with --serve or --bench-pack (the plain bench runs the compiled-in pack)\n"); return 2; }
|
|
if (n > 1 && (o.packDir[n - 1] == '/' || o.packDir[n - 1] == '\\')) ((char*)o.packDir)[n - 1] = 0;
|
|
if (!pf_load(o.packDir, &gPack, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o.packDir, perr); fflush(stdout); return 2; }
|
|
gGeneric = 1;
|
|
gServeWords = 1u << gPack.datasetLog2;
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
gServeCacheWords = 1u << gPack.cacheLog2Words;
|
|
gServeSegments = gPack.cacheSegments;
|
|
#endif
|
|
if (!o.kernelGiven) { snprintf(boundPath, sizeof(boundPath), "%s/kernel_bound.cl", o.packDir); o.kernelPath = boundPath; o.kernelGiven = 1; }
|
|
}
|
|
printf("igneum-bench-cl pack \"%s\" (test harness: no pool, no network, no wallet)%s\n", gGeneric ? gPack.seedString : IGNEUM_SEED_STRING, gGeneric ? " [--pack: generic serve mode]" : "");
|
|
nDev = enumerateDevices(&devs);
|
|
if (nDev == 0) { printf("%s no OpenCL platform or device found (is the GPU driver installed? it provides OpenCL)\n", o.serve ? "error 0" : "FAIL:"); fflush(stdout); return 2; }
|
|
if (o.device >= 0) {
|
|
if (o.device >= nDev) { printf("FAIL: --device %d out of range (%d devices)\n", o.device, nDev); return 2; }
|
|
chosen = o.device;
|
|
} else if (o.vendor) {
|
|
for (i = 0; i < nDev; ++i) if ((devs[i].type & CL_DEVICE_TYPE_GPU) && devs[i].dupOf < 0 && strstr(devs[i].vendor, o.vendor)) { chosen = i; break; }
|
|
if (chosen < 0) {
|
|
printf("OpenCL devices (%d):\n", nDev);
|
|
for (i = 0; i < nDev; ++i) printDevice(i, &devs[i], 0);
|
|
printf("FAIL: no GPU whose vendor contains \"%s\"\n", o.vendor);
|
|
return 2;
|
|
}
|
|
} else {
|
|
for (i = 0; i < nDev; ++i) if ((devs[i].type & CL_DEVICE_TYPE_GPU) && devs[i].dupOf < 0) { chosen = i; break; }
|
|
if (chosen < 0) chosen = 0;
|
|
}
|
|
printf("OpenCL devices (%d):\n", nDev);
|
|
for (i = 0; i < nDev; ++i) printDevice(i, &devs[i], i == chosen && !o.list);
|
|
{
|
|
int dups = 0;
|
|
for (i = 0; i < nDev; ++i) if (devs[i].dupOf >= 0) ++dups;
|
|
if (dups) printf("platforms: %d device(s) hidden as the same card on an older platform of the same vendor (an old driver's OpenCL registration is still present)\n", dups);
|
|
}
|
|
if (o.list) return 0;
|
|
di = &devs[chosen];
|
|
printf("using device [%d] %s\n", chosen, di->name);
|
|
if (di->dupOf >= 0) printf("NOTE: --device %d is the older platform's listing of the card [%d] (driver %s); results on it are for comparison only\n", chosen, di->dupOf, di->driver);
|
|
if (o.timeWall < 0) o.timeWall = (strcmp(di->platformName, "Apple") == 0) ? 1 : 0;
|
|
if (o.timeWall) printf("timing: host wall time (Apple's OpenCL event timestamps are not usable; the rate is still a device rate, see README.md)\n");
|
|
else printf("timing: device event profiling (CL_PROFILING_COMMAND_START/END), like cudaEvent elapsed time\n");
|
|
|
|
memset(&dv, 0, sizeof(dv));
|
|
dv.ctx = clCreateContext(NULL, 1, &di->device, NULL, NULL, &err);
|
|
CL_CHECK_ERR(err, "clCreateContext");
|
|
dv.q = clCreateCommandQueue(dv.ctx, di->device, CL_QUEUE_PROFILING_ENABLE, &err);
|
|
CL_CHECK_ERR(err, "clCreateCommandQueue");
|
|
|
|
if (o.memprobe) {
|
|
int rc = runMemprobe(&dv, di, &o);
|
|
clReleaseCommandQueue(dv.q);
|
|
clReleaseContext(dv.ctx);
|
|
return rc;
|
|
}
|
|
if (o.benchPack && !o.packDir) { printf("FAIL: --bench-pack needs --pack <dir>\n"); return 2; }
|
|
if (o.serve && !o.kernelGiven) {
|
|
/* The bound kernel lives next to the compiled-in kernel.cl as kernel_bound.cl (packs from igneum-pow or igneum-miner export-pack). */
|
|
static char boundPath[1024];
|
|
size_t n = strlen(o.kernelPath);
|
|
if (n >= 9 && strcmp(o.kernelPath + n - 9, "kernel.cl") == 0) {
|
|
snprintf(boundPath, sizeof(boundPath), "%.*skernel_bound.cl", (int)(n - 9), o.kernelPath);
|
|
o.kernelPath = boundPath;
|
|
}
|
|
}
|
|
src = readFile(o.kernelPath, &srcLen);
|
|
if (!src) { printf("FAIL: cannot read kernel source %s (run from proto-opencl/ or pass --kernel)\n", o.kernelPath); return 2; }
|
|
printf("kernel source: %s (%llu bytes)\n", o.kernelPath, (unsigned long long)srcLen);
|
|
setupProgram(&dv, di, &o, src, srcLen);
|
|
free(src);
|
|
printf("build options: %s\n", dv.buildOptions);
|
|
printf("exchange: %s\n", dv.exchangeNote);
|
|
if (o.serve) return runServe(&dv, di, &o);
|
|
if (o.benchPack) return runBenchPack(&dv, di, &o);
|
|
printKernelInfo(di, dv.kHash, "igneum_hash", dv.groupSize, "");
|
|
printf("program: %d instructions x %d iterations, loads/hash %d, op mix %s\n",
|
|
IGNEUM_INSTR_COUNT, IGNEUM_ITERATIONS, IGNEUM_LOADS_PER_HASH, IGNEUM_OP_MIX);
|
|
printf("seed words: %08x %08x %08x %08x %08x %08x %08x %08x\n",
|
|
SEEDW[0], SEEDW[1], SEEDW[2], SEEDW[3], SEEDW[4], SEEDW[5], SEEDW[6], SEEDW[7]);
|
|
printf("day \"%s\" (d0 0x%08x, d1 0x%08x), pack dataset 2^%d words = %d MiB\n",
|
|
IGNEUM_DAY_STRING, IGNEUM_DAY0, IGNEUM_DAY1, IGNEUM_DATASET_LOG2, packMib());
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
#ifndef IGNEUM_MIXER_MULT
|
|
#define IGNEUM_MIXER_MULT 1
|
|
#endif
|
|
printf("dataset construction: memory-hard (%u MiB ChaCha cache, %d dependent cache reads per 64-byte item, mixer x%d; proto-metal/MEMHARD.md)\n",
|
|
(unsigned)(((uint64_t)CACHE_WORDS_HOST * 4u) >> 20), IGNEUM_ITEM_ROUNDS, IGNEUM_MIXER_MULT);
|
|
cachePass = setupCache(&dv, di);
|
|
#else
|
|
printf("dataset construction: closed-form ds_elem (the original prototype dataset, not memory-hard)\n");
|
|
#endif
|
|
|
|
nonces = 1u << o.batchLog2;
|
|
if (nonces % (32u * (uint32_t)o.groupWarps) != 0u) {
|
|
printf("FAIL: 2^%d nonces is not a multiple of %d items per work-group\n", o.batchLog2, 32 * o.groupWarps);
|
|
return 2;
|
|
}
|
|
dOut = clCreateBuffer(dv.ctx, CL_MEM_READ_WRITE, (size_t)nonces * sizeof(uint64_t), NULL, &err);
|
|
CL_CHECK_ERR(err, "clCreateBuffer out");
|
|
|
|
if (o.sweep) { sizes[0] = 4; sizes[1] = 64; sizes[2] = 256; sizes[3] = 512; sizes[4] = 1024; nSizes = 5; }
|
|
else { sizes[0] = o.datasetMib; nSizes = 1; }
|
|
for (i = 0; i < nSizes; ++i) results[i] = runSize(&dv, di, &o, sizes[i], dOut, nonces);
|
|
CL_CHECK(clReleaseMemObject(dOut));
|
|
|
|
printf("\n=== summary (%s, %s, pack %s, batch 2^%d x %d, %d unit(s)/work-group, exchange %s, %s time) ===\n",
|
|
di->name, di->platformName, IGNEUM_SEED_STRING, o.batchLog2, o.batches, o.groupWarps, exchangeName(dv.exchange), o.timeWall ? "wall" : "device-event");
|
|
printf("| dataset MiB | %s ms (second) | Mhash/s | GB/s useful | random loads/s (G) | loads/hash | dataset self-test | vectors |\n",
|
|
IGNEUM_DATASET_MODE == 1 ? "build" : "fill");
|
|
printf("|---|---|---|---|---|---|---|---|\n");
|
|
overall = cachePass;
|
|
for (i = 0; i < nSizes; ++i) {
|
|
const SizeResult* r = &results[i];
|
|
overall = overall && r->dsPass && (!r->vecChecked || r->vecPass);
|
|
anyVec = anyVec || r->vecChecked;
|
|
printf("| %d | %.2f | %.3f | %.2f | %.2f | %d | %s | %s |\n",
|
|
r->mib, r->fillSecondMs, r->hashesPerSec / 1e6, r->gbps,
|
|
r->hashesPerSec * (double)IGNEUM_LOADS_PER_HASH / 1e9, IGNEUM_LOADS_PER_HASH,
|
|
r->dsPass ? "PASS" : "FAIL",
|
|
r->vecChecked ? (r->vecPass ? "PASS (3 warps, standalone and in batch)" : "FAIL") : "skipped (not pack size)");
|
|
}
|
|
#if IGNEUM_DATASET_MODE == 1
|
|
printf("cache: device fill %.2f ms (second), host fill %.1f ms one thread, cache check %s\n", gCacheFillSecondMs, gCacheHostMs, cachePass ? "PASS" : "FAIL");
|
|
if (gCache) clReleaseMemObject(gCache);
|
|
free(hCache);
|
|
#endif
|
|
if (gProfilingFailures) printf("NOTE: %d event profiling queries failed on this runtime (timings printed as -1.00 ms are unavailable)\n", gProfilingFailures);
|
|
printf("exchange: %s\n", dv.exchangeNote);
|
|
if (!anyVec) printf("NOTE: no vectors were checked. Run at %d MiB (the default) to verify against the Mac.\n", packMib());
|
|
printf("OVERALL: %s\n", overall ? "PASS" : "FAIL");
|
|
|
|
releaseProgram(&dv);
|
|
clReleaseCommandQueue(dv.q);
|
|
clReleaseContext(dv.ctx);
|
|
for (i = 0; i < nDev; ++i) free(devs[i].extensions);
|
|
free(devs);
|
|
return overall ? 0 : 1;
|
|
}
|