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