1549 lines
88 KiB
C
1549 lines
88 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
|
|
// Where the card sits on the PCI bus, so the app can tell one physical card listed by two OpenCL platforms (two
|
|
// AMD ICDs after a driver upgrade, PC 1 on 5 October 2026) from two cards: CL_DEVICE_TOPOLOGY_AMD and the NVIDIA pair.
|
|
#define IG_CL_DEVICE_TOPOLOGY_AMD 0x4037
|
|
#define IG_CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD 1
|
|
#define IG_CL_DEVICE_PCI_BUS_ID_NV 0x4008
|
|
#define IG_CL_DEVICE_PCI_SLOT_ID_NV 0x4009
|
|
typedef union {
|
|
struct { cl_uint type; cl_uint data[5]; } raw;
|
|
struct { cl_uint type; cl_char unused[17]; cl_char bus; cl_char device; cl_char function; } pcie;
|
|
} ig_topology_amd;
|
|
|
|
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
|
|
char pci[32]; // "01:00.0" (bus:device.function) when the vendor extension reports it, else ""
|
|
} 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")) {
|
|
ig_topology_amd topo;
|
|
memset(&topo, 0, sizeof(topo));
|
|
clGetDeviceInfo(devs[d], IG_CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof(di.amdWavefront), &di.amdWavefront, NULL);
|
|
if (clGetDeviceInfo(devs[d], IG_CL_DEVICE_TOPOLOGY_AMD, sizeof(topo), &topo, NULL) == CL_SUCCESS && topo.raw.type == IG_CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD)
|
|
snprintf(di.pci, sizeof(di.pci), "%02x:%02x.%x", (unsigned)(unsigned char)topo.pcie.bus, (unsigned)(unsigned char)topo.pcie.device, (unsigned)(unsigned char)topo.pcie.function);
|
|
}
|
|
if (strstr(di.extensions, "cl_nv_device_attribute_query")) {
|
|
cl_uint bus = 0, slot = 0;
|
|
clGetDeviceInfo(devs[d], IG_CL_DEVICE_WARP_SIZE_NV, sizeof(di.nvWarp), &di.nvWarp, NULL);
|
|
if (clGetDeviceInfo(devs[d], IG_CL_DEVICE_PCI_BUS_ID_NV, sizeof(bus), &bus, NULL) == CL_SUCCESS && clGetDeviceInfo(devs[d], IG_CL_DEVICE_PCI_SLOT_ID_NV, sizeof(slot), &slot, NULL) == CL_SUCCESS)
|
|
snprintf(di.pci, sizeof(di.pci), "%02x:%02x.0", (unsigned)(bus & 0xff), (unsigned)(slot & 0xff));
|
|
}
|
|
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);
|
|
// the app's detect.rs reads this line: type, vendor, driver, the compute units, and the PCI address when known
|
|
printf(" %s, vendor %s, driver %s, %s, %u compute units, %u MHz", typeName(d->type), d->vendor, d->driver, d->cVersion, d->computeUnits, d->clockMHz);
|
|
if (d->pci[0]) printf(", pci %s", d->pci);
|
|
printf("\n");
|
|
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 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;
|
|
}
|
|
|
|
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 (!pairIs(cur, f[6], f[7], sw, kw)) {
|
|
if (pairIs(prepared, f[6], f[7], sw, kw)) {
|
|
/* 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 (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;
|
|
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;
|
|
}
|