4714b510f2
After PC 2's gfx1036 (prebuilt-generic path) completed 2,000 jobs a second with no hash from 600 s on: every OpenCL call in host.c's job path is now fatal on error (exit 3, the miner restarts the worker), the dispatch event must read CL_COMPLETE, a chunk 20x faster per nonce than the running mean or an output buffer unchanged since the previous dispatch is a fault, and a stats line every 200 jobs carries the live event and buffer counts (a leak over 16 events or 12 buffers is fatal too). IGNEUM_FAULT_TEST=N exercises the detectors on a healthy device (verified on Apple OpenCL: the stale-output guard fires on the chunk after the injected fault). The launcher shows 'worker fault' and 'restarting' for the card on the miner's WORKER FAULT line and drops the last rate. The emulation's NVRTC stand-in now applies NVRTC's execution-space rule (program.h(46) igneum_launch_* declarations are host code unless -default-device or -DIGNEUM_NO_CUDA), which is what the RTX 5090 reported; test.sh checks the rejection. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
1507 lines
85 KiB
C
1507 lines
85 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;
|
|
static uint32_t host_ds_word(uint32_t w) {
|
|
uint32_t s[16];
|
|
mh_item(hCache, w >> 4u, s);
|
|
return s[w & 15u];
|
|
}
|
|
#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)
|
|
} 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", 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;
|
|
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, "--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
|
|
} DeviceInfo;
|
|
|
|
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;
|
|
}
|
|
}
|
|
*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";
|
|
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 igneum_hash for a work-group of `local` items. 0 if the query is unavailable (reason in *why).
|
|
static size_t querySubGroupSize(const Device* dv, 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(dv->kHash, 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;
|
|
}
|
|
|
|
// 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;
|
|
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;
|
|
double buildMs, cacheMs, datasetMs, checkMs;
|
|
char check[1024]; /* the self-test verdict (packfile.h), one line */
|
|
int checked;
|
|
} ServePair;
|
|
|
|
static void releasePair(ServePair* p) {
|
|
if (!p) return;
|
|
if (p->ds) { clReleaseMemObject(p->ds); ++gMemReleased; }
|
|
if (p->cache) { clReleaseMemObject(p->cache); ++gMemReleased; }
|
|
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];
|
|
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 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;
|
|
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; }
|
|
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;
|
|
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) 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->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;
|
|
return pairSelfTest(dv, di, q, p, words, cacheWords, packDir, err, errCap);
|
|
}
|
|
|
|
/* 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);
|
|
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; }
|
|
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
|
|
|
|
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);
|
|
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] = '_';
|
|
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;
|
|
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[9];
|
|
int nf = 0;
|
|
char* tok;
|
|
char* save = NULL;
|
|
char jobId[64];
|
|
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 < 9; tok = strtok_r(NULL, " ", &save)) f[nf++] = tok;
|
|
if (nf == 0) continue;
|
|
if (strcmp(f[0], "quit") == 0) break;
|
|
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);
|
|
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 (memcmp(sw, cur->sw, 32) != 0 || memcmp(kw, cur->kw, 32) != 0) {
|
|
if (prepared && memcmp(sw, prepared->sw, 32) == 0 && memcmp(kw, prepared->kw, 32) == 0) {
|
|
/* 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 in %.2f ms\n", cur->epochHex, cur->dayHex, wallMs() - t0); fflush(stdout);
|
|
} else if (memcmp(sw, cur->sw, 32) != 0) {
|
|
printf("error %s epoch seed mismatch: this worker holds %s%s (seed words %08x %08x ...)%s, the job's epoch seed %.16s gives %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("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;
|
|
cl_int status = 0;
|
|
size_t g = ((chunk + groupSize - 1) / groupSize) * groupSize;
|
|
uint64_t sig;
|
|
SERVE_CHECK(jobId, clEnqueueWriteBuffer(dv->q, dInit, CL_TRUE, 0, 32, iw, 0, NULL, NULL));
|
|
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));
|
|
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);
|
|
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);
|
|
}
|
|
SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)chunk * sizeof(uint64_t), hOut, 0, NULL, NULL));
|
|
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;
|
|
}
|
|
/* Stale output: the first, middle and last words plus an FNV of the first 32. Two chunks never share it. */
|
|
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;
|
|
}
|
|
}
|
|
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]);
|
|
}
|
|
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);
|
|
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 (old) releasePair(old);
|
|
if (prepared) releasePair(prepared);
|
|
releasePair(cur);
|
|
return 0;
|
|
#endif
|
|
}
|
|
|
|
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) { printf("FAIL: --pack goes with --serve (the 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) && 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) { 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);
|
|
if (o.list) return 0;
|
|
di = &devs[chosen];
|
|
printf("using device [%d] %s\n", chosen, di->name);
|
|
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.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);
|
|
{
|
|
size_t wg = 0;
|
|
cl_ulong lmem = 0;
|
|
clGetKernelWorkGroupInfo(dv.kHash, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL);
|
|
clGetKernelWorkGroupInfo(dv.kHash, di->device, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(lmem), &lmem, NULL);
|
|
printf("kernel: igneum_hash max work-group %llu, local memory %llu bytes, work-group %d x 32\n",
|
|
(unsigned long long)wg, (unsigned long long)lmem, o.groupWarps);
|
|
}
|
|
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
|
|
printf("dataset construction: memory-hard (256 MiB ChaCha cache, %d dependent cache reads per 64-byte item; proto-metal/MEMHARD.md)\n", IGNEUM_ITEM_ROUNDS);
|
|
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;
|
|
}
|