igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity
igneum/proto-cuda/nvrtc/worker.cpp
igneum-labs cb3bc0efe7 hot table (Counter ASIC 2.0 layer 5), measured and not adopted, on the ca2-v3 composed class (squash of tag ca2-cache-history-2026-10-05) 
LoadClass::hot (Option<HotClass { mb, k, added }>) beside mix, slots, scratch, mixer_mult, growth and era; V3_CLASS = { era: None, hot: None, ..LoadClass::MX4 } (hot stays None: the hot code is behind the flag, measured and not adopted). Op::Hot, the hot slots drawn after the scratch slots, no width roll (v2_loads allows the added form's extra slots), the id suffix hot/<S><k>[added], the hot parsing inside parse_loads, the era branch first in name(). HotTable under seed_words("igneum-hot/" || epoch seed) with the cache chain and tag umHT, read at H[mulhi(src, HOT_WORDS)]; DatasetSource::hot attached by new_class_day and from_seed_bytes_class; the acceptance stand-in dataset_elem(idx, S[2], S[3]); the three emitters (hot argument after the init words, ht_segment and igneum_hot_fill beside the layout-aware cores); packfile.h hot fields beside class, era, attempt and mixerMult; OpenCL host, Metal packbench and NVRTC worker fill H on the device and self-test it. Eight packs under proto-cuda/packs-ca2-hot (replaced hot32k4 hot64k4 hot96k4 hot64k2 hot64k8, added hot32k4a hot64k4a hot96k4a) re-exported on the merged crate: vectors unchanged, program.h and program.json carry the mixer fields. docs/plans/hot-table.md (design, spec text, per-tier budget, chip model, Mac and PC measurements, the decision: layer 5 out of v3, the 3.0 note); bench-log entry and addenda with job ids and worker sha256s; the two PC playbooks.

Checks on this commit: cargo test --release 53 + 19 pass (the pinned v2, mx4, era, readwidth and hot packs); the pinned packs under proto-cuda/packs, packs-ca2-mixer, packs-ca2-era and packs-readwidth untouched; Metal packbench and Apple OpenCL --bench-pack on all eight hot packs (run lock, 2^20 at base 0): 96/96 lanes, hot table head, last line and FNV PASS, one fingerprint per pack on both harnesses, equal to the fingerprints before the rebase (hot32k4 679e5e83378d3790, hot64k4 d4c9e456b039fdef, hot96k4 7c98eceffee9fd73, hot64k2 f43b10a95879b8e5, hot64k8 c11309d743be9392, hot32k4a afb700b2d997c847, hot64k4a ba214baa9c1a9e85, hot96k4a 29e1916aed6deff5). No era or mixer behaviour changed: every resolution kept the ca2-v3 side and appended the hot branch.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
2026-10-05 21:57:39 +00:00

1506 lines
91 KiB
C++
Raw Blame History

// igneum-worker-cuda: the one-click NVIDIA worker for igneum-miner --worker. 4 October 2026.
//
// Nothing to install but the NVIDIA driver. The pack's kernels (kernel.cu: cache fill and dataset build;
// kernel_bound.cu: the header-bound hash) are compiled at run time by NVRTC, the toolkit's runtime compiler, which
// ships next to this exe as nvrtc64_120_0.dll plus nvrtc-builtins64_128.dll (NVIDIA's redistributable, see
// THIRD-PARTY.md). The GPU is driven through the driver API in nvcuda.dll, which every NVIDIA driver installs. Both
// libraries are loaded with LoadLibrary/GetProcAddress (cuda_api.h), so no import library is linked and the exe is
// cross-compiled on the Mac with mingw (build-windows.sh). Plain C++17 otherwise.
//
// The source handed to NVRTC is the pack's own text: kernel.cu and kernel_bound.cu up to the host-side launch
// wrappers (which nvcc compiles for the host and NVRTC has no use for), with the pack's program.h and memhard.h as
// named headers, byte for byte. Two stub headers stand in for <cuda_runtime.h> and <cstdint>, which nvcc takes from
// the toolkit. emu/test.sh checks the equality on the Mac. The memory-hard core is therefore the same text host.cu
// compiles, and the miner's CPU re-check of every found nonce covers the rest.
//
// Worker protocol (the same lines as proto-cuda/host.cu --serve and proto-opencl/host.c --serve):
// stdin: job <job_id> <header_prehash_hex 64> <target_hex 16> <nonce_start u64> <nonce_count u64> <epoch_seed_hex 64> <day_seed_hex>
// prepare <epoch_seed_hex 64> <day_seed_hex> <pack_dir> compile that pack in the background, build its cache
// and dataset, self-test it; a job on it then switches
// quit
// stdout: ready cuda <device> pack <seed string> dataset-log2 N batch B regs R prepare 1 path nvrtc ...
// found <job_id> <nonce u64> <hash_hex 16>
// done <job_id> <hashes> <ms>
// error <job_id> <text>
// need <epoch_seed_hex> <day_seed_hex> before the mismatch error: the pair this worker lacks (the miner prepares it)
// prepared <epoch_seed_hex> <day_seed_hex> <ms> ... | prepare-failed <epoch_seed_hex> <day_seed_hex> <text>
// info ...
// The first pack comes from --pack <dir> (igneum-miner export-pack writes it; the launcher passes it). Every pack is
// self-tested before it serves a job: cache head, last line and FNV-1a 64, dataset head, last word and 64 samples,
// and the three vector warps of vectors.h through the bound kernel with the pack's own seed words. A pack that fails
// is refused.
//
// Variant racing (4 October 2026, evening; docs/design/miner-tuning.md): every pack's bound kernel is compiled in
// several variants (loop unrolling, the dataset load path: plain, __ldg, __ldcg, __ldcs; a register budget through
// -maxrregcount or __launch_bounds__; threads per block), each self-tested against the pack's vectors (bit-exact or
// discarded) and run for about two seconds on the card; the fastest serves the hour. A race runs inside the prepare
// (the hourly compile-ahead, one lead before the boundary) and never delays the swap: it has a time budget, "base"
// (the pack's text as shipped) is always the first entry, and a prepare that runs out of budget keeps the best so far.
// While a variant is timed the job loop pauses (one mutex): the numbers are exclusive, mining resumes between
// variants. One line per race: `race <epoch16> device <name> ... variants N a=MH/s b=MH/s ... winner <name> <MH/s>
// gain <pct> ...`. A tuning file (--tuning, or IGNEUM_TUNING_FILE from the app) may pin a variant for this card
// model or order the candidates; the app's over-the-air manifest carries it (fleet learning). Under IGNEUM_EMU the
// race is off (the stand-in checks that the handed-over text equals the pack's).
//
// Usage: igneum-worker-cuda --serve --pack <dir> [--device D] [--batch-log2 22] [--block-warps 1] [--arch sm_120|auto]
// [--race on|off|<name,name,...>] [--race-bench-ms 2000] [--race-budget-s 120] [--race-rounds 1]
// [--variant <name>] [--tuning <file>]
// igneum-worker-cuda --check --pack <dir> [--device D] compile, build, self-test, print timings, exit 0/1
// igneum-worker-cuda --race --pack <dir> [--device D] [--race-rounds 3] the race alone: one line per variant, exit 0/1
#include <cstdint>
#include <cstdarg>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <chrono>
#include <string>
#include <vector>
#include <thread>
#include <atomic>
#include <mutex>
#include <algorithm>
#include <iostream>
#include "cuda_api.h"
#include "packfile.h"
#ifdef _WIN32
#define WIN32_LEAN_AND_MEAN
#include <windows.h>
#else
#include <dlfcn.h>
#include <dirent.h>
#endif
static const char* WORKER_VERSION = "1.0 (4 October 2026)";
// ---------------------------------------------------------------------------------------------
// Helpers
static double wallMs() {
using namespace std::chrono;
return duration<double, std::milli>(steady_clock::now().time_since_epoch()).count();
}
static void emit(const std::string& s) { std::fputs(s.c_str(), stdout); std::fputc('\n', stdout); std::fflush(stdout); }
static void info(const std::string& s) { emit("info " + s); }
static std::string fmt(const char* f, ...) {
char buf[2048];
va_list ap;
va_start(ap, f);
vsnprintf(buf, sizeof(buf), f, ap);
va_end(ap);
return buf;
}
static std::string readText(const std::string& path, bool& ok) {
size_t n = 0;
char* b = pf_read_file(path.c_str(), &n);
if (!b) { ok = false; return ""; }
std::string s(b, n);
free(b);
ok = true;
return s;
}
static std::string exeDir() {
#ifdef _WIN32
char buf[MAX_PATH];
DWORD n = GetModuleFileNameA(nullptr, buf, MAX_PATH);
std::string p(buf, n);
size_t i = p.find_last_of("\\/");
return i == std::string::npos ? "." : p.substr(0, i);
#else
return ".";
#endif
}
// ---------------------------------------------------------------------------------------------
// Loading the two libraries
static void* libOpen(const std::string& name) {
#ifdef _WIN32
return (void*)LoadLibraryA(name.c_str());
#else
return dlopen(name.c_str(), RTLD_NOW);
#endif
}
static void* libSym(void* lib, const char* name) {
#ifdef _WIN32
return (void*)GetProcAddress((HMODULE)lib, name);
#else
return dlsym(lib, name);
#endif
}
#define LOAD_SYM(table, field, name) do { table.field = (decltype(table.field))libSym(lib, name); if (!table.field) { missing += std::string(missing.empty() ? "" : ", ") + name; } } while (0)
static bool loadDriver(Drv& d, std::string& err, std::string& libName) {
#ifdef IGNEUM_EMU
emu_fill_driver(d); libName = "emulation (host threads, no GPU)"; (void)err; return true;
#else
#ifdef _WIN32
const char* names[] = { "nvcuda.dll" };
#else
const char* names[] = { "libcuda.so.1", "libcuda.so" };
#endif
void* lib = nullptr;
for (const char* n : names) { lib = libOpen(n); if (lib) { libName = n; break; } }
if (!lib) { err = "the CUDA driver library (nvcuda.dll) is not installed: install or update the NVIDIA driver"; return false; }
std::string missing;
LOAD_SYM(d, init, "cuInit");
LOAD_SYM(d, driverGetVersion, "cuDriverGetVersion");
LOAD_SYM(d, deviceGetCount, "cuDeviceGetCount");
LOAD_SYM(d, deviceGet, "cuDeviceGet");
LOAD_SYM(d, deviceGetName, "cuDeviceGetName");
LOAD_SYM(d, deviceGetAttribute, "cuDeviceGetAttribute");
LOAD_SYM(d, deviceTotalMem, "cuDeviceTotalMem_v2");
LOAD_SYM(d, primaryCtxSetFlags, "cuDevicePrimaryCtxSetFlags_v2");
LOAD_SYM(d, primaryCtxRetain, "cuDevicePrimaryCtxRetain");
LOAD_SYM(d, primaryCtxRelease, "cuDevicePrimaryCtxRelease_v2");
LOAD_SYM(d, ctxSetCurrent, "cuCtxSetCurrent");
LOAD_SYM(d, ctxSynchronize, "cuCtxSynchronize");
LOAD_SYM(d, memGetInfo, "cuMemGetInfo_v2");
LOAD_SYM(d, memAlloc, "cuMemAlloc_v2");
LOAD_SYM(d, memFree, "cuMemFree_v2");
LOAD_SYM(d, memcpyDtoH, "cuMemcpyDtoH_v2");
LOAD_SYM(d, moduleLoadData, "cuModuleLoadData");
LOAD_SYM(d, moduleUnload, "cuModuleUnload");
LOAD_SYM(d, moduleGetFunction, "cuModuleGetFunction");
LOAD_SYM(d, launchKernel, "cuLaunchKernel");
LOAD_SYM(d, streamCreate, "cuStreamCreate");
LOAD_SYM(d, streamSynchronize, "cuStreamSynchronize");
LOAD_SYM(d, streamDestroy, "cuStreamDestroy_v2");
LOAD_SYM(d, funcGetAttribute, "cuFuncGetAttribute");
LOAD_SYM(d, occupancy, "cuOccupancyMaxActiveBlocksPerMultiprocessor");
LOAD_SYM(d, getErrorString, "cuGetErrorString");
LOAD_SYM(d, getErrorName, "cuGetErrorName");
if (!missing.empty()) { err = "the driver library lacks " + missing + " (driver too old; CUDA 11 or newer is needed)"; return false; }
return true;
#endif
}
#ifdef _WIN32
// nvrtc64_<major>0_0.dll next to the exe (any major), then a toolkit on PATH. IGNEUM_NVRTC_DLL overrides.
static std::vector<std::string> nvrtcCandidates() {
std::vector<std::string> v;
if (const char* o = std::getenv("IGNEUM_NVRTC_DLL")) v.push_back(o);
std::string dir = exeDir();
WIN32_FIND_DATAA fd;
HANDLE h = FindFirstFileA((dir + "\\nvrtc64_*_0.dll").c_str(), &fd);
if (h != INVALID_HANDLE_VALUE) {
do { std::string n = fd.cFileName; if (n.find(".alt.") == std::string::npos) v.push_back(dir + "\\" + n); } while (FindNextFileA(h, &fd));
FindClose(h);
}
v.push_back("nvrtc64_120_0.dll");
v.push_back("nvrtc64_130_0.dll");
if (const char* cp = std::getenv("CUDA_PATH")) { v.push_back(std::string(cp) + "\\bin\\nvrtc64_120_0.dll"); v.push_back(std::string(cp) + "\\bin\\nvrtc64_130_0.dll"); }
return v;
}
#endif
static bool loadNvrtc(Rtc& r, std::string& err, std::string& libName) {
#ifdef IGNEUM_EMU
emu_fill_nvrtc(r); libName = "emulation (source recorded and checked, nothing compiled)"; (void)err; return true;
#else
void* lib = nullptr;
#ifdef _WIN32
for (const std::string& n : nvrtcCandidates()) { lib = libOpen(n); if (lib) { libName = n; break; } }
if (!lib) { err = "nvrtc64_120_0.dll (and nvrtc-builtins64_128.dll) must sit next to " + exeDir() + "\\igneum-worker-cuda.exe; they are in the package"; return false; }
#else
const char* names[] = { "libnvrtc.so.12", "libnvrtc.so" };
for (const char* n : names) { lib = libOpen(n); if (lib) { libName = n; break; } }
if (!lib) { err = "libnvrtc.so.12 not found"; return false; }
#endif
std::string missing;
LOAD_SYM(r, version, "nvrtcVersion");
LOAD_SYM(r, createProgram, "nvrtcCreateProgram");
LOAD_SYM(r, destroyProgram, "nvrtcDestroyProgram");
LOAD_SYM(r, compileProgram, "nvrtcCompileProgram");
LOAD_SYM(r, getProgramLogSize, "nvrtcGetProgramLogSize");
LOAD_SYM(r, getProgramLog, "nvrtcGetProgramLog");
LOAD_SYM(r, getPTXSize, "nvrtcGetPTXSize");
LOAD_SYM(r, getPTX, "nvrtcGetPTX");
LOAD_SYM(r, getCUBINSize, "nvrtcGetCUBINSize");
LOAD_SYM(r, getCUBIN, "nvrtcGetCUBIN");
LOAD_SYM(r, addNameExpression, "nvrtcAddNameExpression");
LOAD_SYM(r, getLoweredName, "nvrtcGetLoweredName");
LOAD_SYM(r, getErrorString, "nvrtcGetErrorString");
if (!missing.empty()) { err = "the NVRTC library lacks " + missing; return false; }
r.getNumSupportedArchs = (decltype(r.getNumSupportedArchs))libSym(lib, "nvrtcGetNumSupportedArchs");
r.getSupportedArchs = (decltype(r.getSupportedArchs))libSym(lib, "nvrtcGetSupportedArchs");
return true;
#endif
}
// ---------------------------------------------------------------------------------------------
// The device context
struct Ctx {
Drv drv;
Rtc rtc;
// read-width experiment, variant 5 (5 October 2026): persistent warps and their scratch; --warps caps the launch
int warps = 0; // 0 = the resident capacity from the occupancy query, rounded down to a power of two
uint32_t salt = 1; // the running per-unit tag salt (+= units per launch)
int batches = 5; // --bench: timed dispatches
CUdevice dev = 0;
CUcontext ctx = nullptr;
std::string name;
int major = 0, minor = 0, sms = 0, driverVersion = 0, rtcMajor = 0, rtcMinor = 0;
std::vector<int> rtcArchs; // what NVRTC can target (empty when the query is unavailable)
std::string archOpt; // "sm_120" or "compute_120": what the packs are compiled for
bool ptx = false; // true when archOpt is compute_XY (PTX, driver JIT)
std::string why; // how archOpt was chosen
int blockWarps = 1;
// variant racing (see the header): which variants, how long each is timed, the budget of a race, rounds
std::string race = "on"; // on | off | comma list of variant names
int raceBenchMs = 2000, raceBudgetS = 120, raceRounds = 1, batchLog2 = 22;
std::string pinned; // --variant: use this variant, no race
std::string tuning; // the tuning file's text ("" = none)
std::string err(CUresult r) { const char* s = nullptr; if (drv.getErrorString) drv.getErrorString(r, &s); return s ? s : "CUDA driver error"; }
};
#define DRV_CHECK(c, call, what) do { CUresult r_ = (call); if (r_ != CUDA_SUCCESS) { err = std::string(what) + ": " + (c).err(r_); return false; } } while (0)
static bool openDevice(Ctx& c, int device, const std::string& archArg, std::string& err) {
DRV_CHECK(c, c.drv.init(0), "cuInit");
int count = 0;
DRV_CHECK(c, c.drv.deviceGetCount(&count), "cuDeviceGetCount");
if (count == 0) { err = "no CUDA device"; return false; }
if (device < 0 || device >= count) { err = fmt("device %d out of range (%d devices)", device, count); return false; }
DRV_CHECK(c, c.drv.deviceGet(&c.dev, device), "cuDeviceGet");
char name[256] = {0};
DRV_CHECK(c, c.drv.deviceGetName(name, 255, c.dev), "cuDeviceGetName");
c.name = name;
for (char& ch : c.name) if (ch == ' ') ch = '_';
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, c.dev), "compute capability major");
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, c.dev), "compute capability minor");
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.sms, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, c.dev), "multiprocessor count");
c.drv.driverGetVersion(&c.driverVersion);
// Blocking sync, set before the context exists: the host thread sleeps in cuStreamSynchronize instead of spinning
// (one full core per worker at the default spin schedule, measured on the RTX 5090 with eight workers, 3 Oct 2026).
c.drv.primaryCtxSetFlags(c.dev, CU_CTX_SCHED_BLOCKING_SYNC);
DRV_CHECK(c, c.drv.primaryCtxRetain(&c.ctx, c.dev), "cuDevicePrimaryCtxRetain");
DRV_CHECK(c, c.drv.ctxSetCurrent(c.ctx), "cuCtxSetCurrent");
c.rtc.version(&c.rtcMajor, &c.rtcMinor);
if (c.rtc.getNumSupportedArchs && c.rtc.getSupportedArchs) {
int n = 0;
if (c.rtc.getNumSupportedArchs(&n) == NVRTC_SUCCESS && n > 0 && n < 256) { c.rtcArchs.assign((size_t)n, 0); if (c.rtc.getSupportedArchs(c.rtcArchs.data()) != NVRTC_SUCCESS) c.rtcArchs.clear(); }
}
// The target: the device's own SASS (sm_XY) when this NVRTC knows the architecture, else PTX for the newest
// architecture it knows below the device's, which the driver JIT-compiles forward. --arch overrides.
int cc = c.major * 10 + c.minor;
if (archArg != "auto" && !archArg.empty()) {
c.archOpt = archArg; c.ptx = archArg.rfind("compute_", 0) == 0; c.why = "--arch";
} else if (c.rtcArchs.empty()) {
c.archOpt = fmt("sm_%d", cc); c.why = "the device's architecture (NVRTC did not list its targets)";
} else {
bool known = false; int best = 0;
for (int a : c.rtcArchs) { if (a == cc) known = true; if (a <= cc && a > best) best = a; }
if (known) { c.archOpt = fmt("sm_%d", cc); c.why = "the device's architecture, listed by NVRTC"; }
else if (best > 0) { c.archOpt = fmt("compute_%d", best); c.ptx = true; c.why = fmt("this NVRTC does not know sm_%d; PTX for compute_%d, JIT-compiled by the driver", cc, best); }
else { c.archOpt = fmt("compute_%d", c.rtcArchs.front()); c.ptx = true; c.why = fmt("this NVRTC knows nothing at or below sm_%d; PTX for its oldest target", cc); }
}
return true;
}
// ---------------------------------------------------------------------------------------------
// NVRTC: compile one of the pack's kernel files
// The pack's kernel files end in host-side launch wrappers (cudaError_t igneum_launch_* with <<< >>> launches) that
// nvcc compiles for the host. NVRTC compiles device code only, so the text is cut there. The cut is checked: nothing
// device-side may follow it.
static bool deviceOnly(const std::string& text, std::string& out, std::string& err) {
size_t cut = text.find("\n// Host-side launch wrappers");
if (cut == std::string::npos) cut = text.find("\ncudaError_t ");
if (cut == std::string::npos) { out = text; return true; }
std::string tail = text.substr(cut + 1);
if (tail.find("__global__") != std::string::npos || tail.find("__device__") != std::string::npos) { err = "device code after the host launch wrappers; the pack layout is not the one this worker knows"; return false; }
out = text.substr(0, cut + 1);
return true;
}
static const char* STUB_CUDA_RUNTIME =
"// igneum-worker-cuda: stand-in for <cuda_runtime.h> under NVRTC, which has the device built-ins already\n"
"#pragma once\n"
"#ifndef __CUDACC_RTC__\n#error \"this stub is for NVRTC only\"\n#endif\n"
"#ifdef __SIZE_TYPE__\ntypedef __SIZE_TYPE__ size_t;\n#elif defined(__LP64__) || defined(_LP64)\ntypedef unsigned long size_t;\n#else\ntypedef unsigned long long size_t;\n#endif\n";
static const char* STUB_CSTDINT =
"// igneum-worker-cuda: stand-in for <cstdint> under NVRTC (the fixed-width types the packs use)\n"
"#pragma once\n"
"typedef signed char int8_t; typedef unsigned char uint8_t; typedef short int16_t; typedef unsigned short uint16_t;\n"
"typedef int int32_t; typedef unsigned int uint32_t;\n"
"#if defined(__LP64__) || defined(_LP64)\ntypedef long int64_t; typedef unsigned long uint64_t;\n"
"#else\ntypedef long long int64_t; typedef unsigned long long uint64_t;\n#endif\n";
struct Compiled {
std::vector<char> image;
std::vector<std::string> lowered;
double ms = 0;
std::string log;
};
static bool rtcCompile(Ctx& c, const std::string& src, const char* name, const std::string& programH, const std::string& memhardH,
const std::vector<std::string>& nameExprs, Compiled& out, std::string& err, const std::vector<std::string>& extraOpts = {}) {
double t0 = wallMs();
const char* headers[4] = { STUB_CUDA_RUNTIME, STUB_CSTDINT, programH.c_str(), memhardH.c_str() };
const char* names[4] = { "cuda_runtime.h", "cstdint", "program.h", "memhard.h" };
nvrtcProgram prog = nullptr;
nvrtcResult r = c.rtc.createProgram(&prog, src.c_str(), name, 4, headers, names);
if (r != NVRTC_SUCCESS) { err = std::string("nvrtcCreateProgram: ") + c.rtc.getErrorString(r); return false; }
for (const std::string& e : nameExprs) {
r = c.rtc.addNameExpression(prog, e.c_str());
if (r != NVRTC_SUCCESS) { err = "nvrtcAddNameExpression " + e + ": " + c.rtc.getErrorString(r); c.rtc.destroyProgram(&prog); return false; }
}
std::string archOpt = "--gpu-architecture=" + c.archOpt;
// -default-device: NVRTC rejects unannotated functions as host code (nvcc treats them as host and discards them);
// the pack headers (program.h, memhard.h) carry plain inline helpers, so every unannotated function is device code here.
std::vector<const char*> opts = { archOpt.c_str(), "--std=c++17", "-default-device" };
for (const std::string& o : extraOpts) opts.push_back(o.c_str());
r = c.rtc.compileProgram(prog, (int)opts.size(), opts.data());
{
size_t logSize = 0;
if (c.rtc.getProgramLogSize(prog, &logSize) == NVRTC_SUCCESS && logSize > 1) {
std::vector<char> log(logSize);
c.rtc.getProgramLog(prog, log.data());
out.log.assign(log.data(), logSize - 1);
}
}
if (r != NVRTC_SUCCESS) {
std::string one;
for (char ch : out.log) { if (ch == '\n' || ch == '\r') { if (one.size() && one.back() != '|') one += " | "; } else one += ch; if (one.size() > 600) break; }
err = std::string("nvrtcCompileProgram ") + name + " for " + c.archOpt + ": " + c.rtc.getErrorString(r) + ": " + one;
c.rtc.destroyProgram(&prog);
return false;
}
for (const std::string& e : nameExprs) {
const char* lowered = nullptr;
r = c.rtc.getLoweredName(prog, e.c_str(), &lowered);
if (r != NVRTC_SUCCESS || !lowered) { err = "nvrtcGetLoweredName " + e + ": " + c.rtc.getErrorString(r); c.rtc.destroyProgram(&prog); return false; }
out.lowered.push_back(lowered);
}
size_t n = 0;
if (c.ptx) {
r = c.rtc.getPTXSize(prog, &n);
if (r == NVRTC_SUCCESS) { out.image.resize(n); r = c.rtc.getPTX(prog, out.image.data()); }
} else {
r = c.rtc.getCUBINSize(prog, &n);
if (r == NVRTC_SUCCESS) { out.image.resize(n); r = c.rtc.getCUBIN(prog, out.image.data()); }
}
c.rtc.destroyProgram(&prog);
if (r != NVRTC_SUCCESS || n == 0) { err = std::string(c.ptx ? "nvrtcGetPTX" : "nvrtcGetCUBIN") + ": " + c.rtc.getErrorString(r); return false; }
out.ms = wallMs() - t0;
return true;
}
// ---------------------------------------------------------------------------------------------
// A resident pair: one pack compiled, its cache and dataset on the device, self-tested
static bool hexEq(const std::string& a, const std::string& b) {
if (a.size() != b.size()) return false;
for (size_t i = 0; i < a.size(); ++i) if (std::tolower((unsigned char)a[i]) != std::tolower((unsigned char)b[i])) return false;
return true;
}
struct Pair;
// 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.
static bool pairIs(const Pair* p, const std::string& epochHex, const std::string& dayHex);
struct Pair {
std::string dir, epochHex, dayHex, seedString;
uint32_t sw[8] = {0}, kw[8] = {0};
uint32_t datasetLog2 = 0, words = 0, cacheWords = 0, cacheSegments = 0;
CUmodule modKernel = nullptr, modBound = nullptr;
CUfunction fCacheFill = nullptr, fBuild = nullptr, fHashBound = nullptr;
CUdeviceptr cache = 0, ds = 0;
double compileMs = 0, cacheMs = 0, dsMs = 0, checkMs = 0;
std::string check;
bool checkPass = false, checked = false;
int regs = 0, blocksPerSM = 0;
int blockWarps = 1; // threads per block = 32 x this (the winning variant's, else the worker's default)
std::string variant = "base"; // the bound kernel in service: a variant name (see allVariants)
std::string raceLine; // the race's one-line report, emitted by the main thread with "prepared"
double raceMs = 0;
// read-width experiment (5 October 2026): the pack's load class and, for variant 5, the persistent-warp scratch
std::string loadClass = "v2";
std::string programClass = "v2", eraHex; // Counter ASIC 2.0: the pack's class and era seed (packfile.h)
uint32_t loadsPerHash = 128, bytesPerHash = 512, scratchOps = 0;
bool persistent = false;
CUdeviceptr scratch = 0;
int warps = 0; // persistent warps launched (the arena holds this many)
int residentWarps = 0; // the occupancy query's capacity: blocks/SM x warps/block x SMs
size_t scratchBytes = 0;
// hot-table experiment (5 October 2026, docs/plans/hot-table.md): the epoch's hot table, filled on the device by the
// pack's igneum_hot_fill, the argument after the init words
uint32_t hotMb = 0, hotWords = 0, hotSegments = 0, hotSlots = 0;
CUfunction fHotFill = nullptr;
CUdeviceptr hot = 0;
double hotMs = 0;
};
static bool pairIs(const Pair* p, const std::string& epochHex, const std::string& dayHex) {
return p && hexEq(p->epochHex, epochHex) && hexEq(p->dayHex, dayHex);
}
// Counter ASIC 2.0: a job that names a class (and an era) belongs to a pair of that class (and era) only, so a pack
// of the old class for the same seeds is not this job's pair and the prepared pack of the right class wins.
static bool pairIsClass(const Pair* p, const std::string& epochHex, const std::string& dayHex, const std::string& cls, const std::string& era) {
char why[256];
return pairIs(p, epochHex, dayHex) && pf_pack_class_ok(p->programClass.c_str(), p->eraHex.c_str(), cls.c_str(), era.c_str(), 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(std::vector<std::string>& f, std::string& cls, std::string& era) {
while (!f.empty()) {
char c[8] = {0}, e[65] = {0};
if (!pf_class_token(f.back().c_str(), c, sizeof(c), e, sizeof(e))) break;
if (c[0]) cls = c;
if (e[0]) era = e;
f.pop_back();
}
}
// The job loop and a race take turns on the card: a variant is timed with no job running (exclusive numbers), and
// mining resumes between variants. Held per chunk by the job loop, per variant by the race.
static std::mutex gpuMutex;
static void releasePair(Ctx& c, Pair* p) {
if (!p) return;
if (p->ds) c.drv.memFree(p->ds);
if (p->cache) c.drv.memFree(p->cache);
if (p->scratch) c.drv.memFree(p->scratch);
if (p->hot) c.drv.memFree(p->hot);
if (p->modBound) c.drv.moduleUnload(p->modBound);
if (p->modKernel) c.drv.moduleUnload(p->modKernel);
delete p;
}
struct IgneumInitWordsArg { uint32_t w[8]; };
// `block` threads per block (32 x warps); `nonces` must be a multiple of it.
static bool launchHash(Ctx& c, Pair* p, CUdeviceptr out, uint32_t baseNonce, const uint32_t iw[8], uint32_t nonces, uint32_t block, CUstream s, std::string& err) {
uint32_t mask = p->words - 1u;
IgneumInitWordsArg a; std::memcpy(a.w, iw, 32);
if (p->persistent) {
// Variant 5: N persistent warps over nonces / 32 units; the arena was sized for p->warps warps in buildPair.
uint32_t units = nonces / 32u, warps = (uint32_t)p->warps;
if (warps > units) warps = units;
while (warps > 1u && units % warps != 0u) warps >>= 1;
if (block != 32u) { err = "a variant-5 pack runs one warp per block (--block-warps 1)"; return false; }
uint32_t salt = c.salt; c.salt += units;
// the hot table (when the pack has one) sits between the init words and the scratch triple
void* args[9] = { &p->ds, &out, &baseNonce, &mask, &a, &p->scratch, &units, &salt, nullptr };
if (p->hot) { args[5] = &p->hot; args[6] = &p->scratch; args[7] = &units; args[8] = &salt; }
DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, warps, 1, 1, 32, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound (persistent)");
return true;
}
void* args[6] = { &p->ds, &out, &baseNonce, &mask, &a, &p->hot };
DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, nonces / block, 1, 1, block, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound");
return true;
}
// ---------------------------------------------------------------------------------------------
// Variant racing
struct Variant {
std::string name;
int unroll = 0; // 0: the iteration loop as emitted; N: "#pragma unroll N" before it (8 = fully unrolled)
int load = 0; // 0: plain ds[i]; 1: __ldg (read-only data path); 2: __ldcg (L2 only, no L1); 3: __ldcs (streaming)
int maxrreg = 0; // 0: none; N: --maxrregcount=N (registers per thread, occupancy against spills)
int blockWarps = 0; // 0: the worker's --block-warps; N: 32 x N threads per block
int minBlocks = 0; // N > 0: __launch_bounds__(32 x blockWarps, N) (the compiler fits N blocks per SM)
};
// The catalogue. Names are stable: the tuning file and the fleet records use them. "base" is the pack's text as
// shipped with the worker's default block and is always the first entry of a race.
static std::vector<Variant> allVariants() {
std::vector<Variant> v;
auto add = [&](const char* n, int unroll, int load, int maxrreg, int bw, int minBlocks) { Variant x; x.name = n; x.unroll = unroll; x.load = load; x.maxrreg = maxrreg; x.blockWarps = bw; x.minBlocks = minBlocks; v.push_back(x); };
add("base", 0, 0, 0, 0, 0);
add("w2", 0, 0, 0, 2, 0);
add("w4", 0, 0, 0, 4, 0);
add("w8", 0, 0, 0, 8, 0);
add("u2", 2, 0, 0, 0, 0);
add("u8", 8, 0, 0, 0, 0);
add("ldg", 0, 1, 0, 0, 0);
add("ldcg", 0, 2, 0, 0, 0);
add("ldcs", 0, 3, 0, 0, 0);
add("r32", 0, 0, 32, 0, 0);
add("r64", 0, 0, 64, 0, 0);
add("lb4-w4", 0, 0, 0, 4, 4);
add("lb8-w2", 0, 0, 0, 2, 8);
add("u2-ldg", 2, 1, 0, 0, 0);
add("u2-w4", 2, 0, 0, 4, 0);
add("ldg-w4", 0, 1, 0, 4, 0);
add("ldcg-w4", 0, 2, 0, 4, 0);
return v;
}
static const Variant* findVariant(const std::vector<Variant>& all, const std::string& name) {
for (const Variant& v : all) if (v.name == name) return &v;
return nullptr;
}
// The variant's source: the pack's bound-kernel text with the variant's rewrites. Every rewrite has an exact anchor
// in the text igneum-pow emits; a text without the anchor refuses the variant (why), it is never guessed.
static bool variantSource(const std::string& base, const Variant& v, int blockWarps, std::string& out, std::string& why) {
out = base;
if (v.unroll > 0) {
const char* anchor = "\n for (uint32_t it = 0u; it < ";
size_t p = out.find(anchor);
if (p == std::string::npos) { why = "no iteration loop in the bound kernel text"; return false; }
out.insert(p + 1, fmt("#pragma unroll %d\n", v.unroll));
}
if (v.load > 0) {
const char* fn = v.load == 1 ? "__ldg" : v.load == 2 ? "__ldcg" : "__ldcs";
size_t body = out.find("igneum_hash_bound(");
if (body == std::string::npos) { why = "no igneum_hash_bound in the text"; return false; }
size_t p = body; int n = 0;
while ((p = out.find(" ^ ds[", p)) != std::string::npos) {
size_t close = out.find(']', p);
if (close == std::string::npos) { why = "an unterminated dataset load"; return false; }
out.insert(close + 1, ")"); // " ^ ds[idx]" -> " ^ __ldg(&ds[idx])"
out.insert(p + 3, std::string(fn) + "(&");
p += 6; ++n;
}
if (n == 0) { why = "no dataset loads in the bound kernel"; return false; }
}
if (v.minBlocks > 0) {
const char* a = "__global__ void igneum_hash_bound(";
size_t p = out.find(a);
if (p == std::string::npos) { why = "no kernel declaration anchor"; return false; }
out.replace(p, std::strlen(a), fmt("__global__ void __launch_bounds__(%d, %d) igneum_hash_bound(", 32 * blockWarps, v.minBlocks));
}
return true;
}
// The tuning file: {"cards": {"<device name as this worker prints it>": {"variant": "u2-ldg", "race": false,
// "candidates": ["u2-ldg", "ldg", "base"]}}, ...}. Read with plain string scanning (no JSON library in this exe);
// a file that does not parse means no tuning. Keys and names are [A-Za-z0-9_.-].
struct Tuning {
bool found = false;
std::string variant; // pinned variant ("" = none)
bool race = true; // false: use the pinned variant without a race
std::vector<std::string> candidates;
};
static std::string jsonStringAfter(const std::string& t, size_t from, const char* key, size_t limit) {
size_t k = t.find(std::string("\"") + key + "\"", from);
if (k == std::string::npos || k > limit) return "";
size_t q = t.find('"', t.find(':', k) + 1);
if (q == std::string::npos) return "";
size_t e = t.find('"', q + 1);
return e == std::string::npos ? "" : t.substr(q + 1, e - q - 1);
}
static Tuning readTuning(const std::string& text, const std::string& device) {
Tuning tu;
if (text.empty()) return tu;
size_t cards = text.find("\"cards\"");
if (cards == std::string::npos) return tu;
size_t k = text.find("\"" + device + "\"", cards);
if (k == std::string::npos) return tu;
size_t open = text.find('{', k);
if (open == std::string::npos) return tu;
size_t close = open; int depth = 0;
for (; close < text.size(); ++close) { if (text[close] == '{') ++depth; else if (text[close] == '}' && --depth == 0) break; }
if (close >= text.size()) return tu;
tu.found = true;
tu.variant = jsonStringAfter(text, open, "variant", close);
size_t r = text.find("\"race\"", open);
if (r != std::string::npos && r < close) { size_t c = text.find(':', r); tu.race = text.compare(text.find_first_not_of(" \t\r\n", c + 1), 5, "false") != 0; }
size_t cand = text.find("\"candidates\"", open);
if (cand != std::string::npos && cand < close) {
size_t a = text.find('[', cand), b = text.find(']', a == std::string::npos ? cand : a);
if (a != std::string::npos && b != std::string::npos && b < close) {
size_t i = a;
while ((i = text.find('"', i + 1)) != std::string::npos && i < b) { size_t e = text.find('"', i + 1); if (e == std::string::npos || e > b) break; tu.candidates.push_back(text.substr(i + 1, e - i - 1)); i = e; }
}
}
return tu;
}
struct RaceEntry {
Variant v;
int blockWarps = 1; // the block this entry runs with
Compiled cb;
CUmodule mod = nullptr;
CUfunction fn = nullptr;
int regs = 0, blocksPerSM = 0;
double mhs = 0; // best round
bool ok = false; // compiled, loaded, self-tested
std::string note; // why not, or a detail
std::string src;
};
// One timed window on the card for an entry: the pack's vector warps (bit-exact or the entry is out), then
// launches of `batch` nonces until benchMs elapsed (the first launch warms up and is not counted). Holds gpuMutex.
static bool raceTime(Ctx& c, Pair* p, const PfPack& pk, RaceEntry& e, CUdeviceptr dOut, uint32_t batch, int benchMs, CUstream s, bool selfTest) {
std::lock_guard<std::mutex> hold(gpuMutex);
CUfunction keep = p->fHashBound;
p->fHashBound = e.fn;
std::string err;
uint32_t block = 32u * (uint32_t)e.blockWarps;
bool ok = true;
if (selfTest && pk.haveVectors) {
std::vector<uint64_t> vec(32);
for (int w = 0; w < pk.vecWarps && ok; ++w) {
if (!launchHash(c, p, dOut, pk.vecBase[w], p->sw, block, block, s, err)) { e.note = "launch: " + err; ok = false; break; }
CUresult r = c.drv.streamSynchronize(s);
if (r == CUDA_SUCCESS) r = c.drv.memcpyDtoH(vec.data(), dOut, 32u * 8u);
if (r != CUDA_SUCCESS) { e.note = "vector warp: " + c.err(r); ok = false; break; }
for (int l = 0; l < 32; ++l) if (vec[(size_t)l] != pk.vecOut[w][l]) { e.note = fmt("vector warp %d lane %d: device %016llx expected %016llx (discarded)", w, l, (unsigned long long)vec[(size_t)l], (unsigned long long)pk.vecOut[w][l]); ok = false; break; }
}
}
if (ok) {
uint32_t iw[8]; std::memcpy(iw, p->sw, 32);
uint32_t n = batch - (batch % block);
if (n == 0) n = block;
double t0 = 0; uint64_t hashes = 0; int launches = 0;
while (true) {
if (!launchHash(c, p, dOut, 0x10000000u + (uint32_t)launches * n, iw, n, block, s, err)) { e.note = "launch: " + err; ok = false; break; }
CUresult r = c.drv.streamSynchronize(s);
if (r != CUDA_SUCCESS) { e.note = "bench: " + c.err(r); ok = false; break; }
double now = wallMs();
if (launches == 0) t0 = now; else hashes += n;
++launches;
if (launches >= 3 && now - t0 >= benchMs) { double mhs = (double)hashes / (now - t0) / 1000.0; if (mhs > e.mhs) e.mhs = mhs; break; }
}
}
p->fHashBound = keep;
return ok;
}
// Races the bound kernel of `p` (its cache and dataset are built, its base kernel self-tested) and installs the
// winner: p->modBound, fHashBound, regs, blockWarps, variant. The pair keeps serving its base kernel if every other
// entry fails. `boundDev`, `programH`, `memhardH` are the texts the base was compiled from. Sets p->raceLine.
static void racePair(Ctx& c, Pair* p, const PfPack& pk, const std::string& boundDev, const std::string& programH, const std::string& memhardH, CUstream s) {
double t0 = wallMs();
std::vector<Variant> all = allVariants();
Tuning tu = readTuning(c.tuning, c.name);
std::string pinned = !c.pinned.empty() ? c.pinned : (tu.found && !tu.race ? tu.variant : "");
// the order: base first, then the pinned or tuned candidates, then the rest (or the --race list only)
std::vector<Variant> order;
auto push = [&](const std::string& n) { const Variant* v = findVariant(all, n); if (v && !findVariant(order, n)) order.push_back(*v); };
push("base");
if (!pinned.empty()) push(pinned);
else {
for (const std::string& n : tu.candidates) push(n);
if (c.race != "on" && c.race != "off") { std::string rest = c.race; size_t i = 0; while (i <= rest.size()) { size_t j = rest.find(',', i); if (j == std::string::npos) j = rest.size(); if (j > i) push(rest.substr(i, j - i)); i = j + 1; } }
else if (c.race == "on") for (const Variant& v : all) push(v.name);
}
#ifdef IGNEUM_EMU
order.resize(1); // the stand-in checks that the handed-over text is the pack's; no rewrites under emulation
#endif
if (order.size() < 2) {
// nothing to race against base (emulation, or --race with no known name): no timing, the base kernel serves
p->blockWarps = c.blockWarps; p->variant = "base"; p->raceMs = wallMs() - t0;
p->raceLine = fmt("race %.16s device %s variants 1 base only, no race (%s)", p->epochHex.c_str(), c.name.c_str(),
#ifdef IGNEUM_EMU
"emulation");
#else
"no other variant named");
#endif
return;
}
const bool pinnedOnly = !pinned.empty() && order.size() == 2;
uint32_t batch = 1u << c.batchLog2;
int benchMs = c.raceBenchMs;
double deadline = t0 + c.raceBudgetS * 1000.0;
std::vector<RaceEntry> entries;
for (const Variant& v : order) {
RaceEntry e; e.v = v; e.blockWarps = v.blockWarps > 0 ? v.blockWarps : c.blockWarps;
if (v.name == "base") { e.mod = p->modBound; e.fn = p->fHashBound; e.regs = p->regs; e.blocksPerSM = p->blocksPerSM; e.ok = true; }
else if (!variantSource(boundDev, v, e.blockWarps, e.src, e.note)) e.ok = false;
else e.ok = true; // compiled below
entries.push_back(std::move(e));
}
// Compile the variants, up to four at a time (NVRTC is thread-safe; the compile is CPU work)
{
std::vector<size_t> todo;
for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok) todo.push_back(i);
size_t next = 0;
std::mutex m;
auto work = [&]() {
while (true) {
size_t i;
{ std::lock_guard<std::mutex> g(m); if (next >= todo.size() || wallMs() > deadline - benchMs) return; i = todo[next++]; }
RaceEntry& e = entries[i];
std::vector<std::string> extra;
if (e.v.maxrreg > 0) extra.push_back(fmt("--maxrregcount=%d", e.v.maxrreg));
std::string err;
if (!rtcCompile(c, e.src, "kernel_bound.cu", programH, memhardH, { "igneum_hash_bound" }, e.cb, err, extra)) { e.ok = false; e.note = "compile: " + err.substr(0, 200); }
}
};
int threads = (int)std::min<size_t>(4, std::max<size_t>(1, todo.size()));
std::vector<std::thread> ts;
for (int t = 0; t < threads; ++t) ts.emplace_back(work);
for (std::thread& t : ts) t.join();
for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok && entries[i].cb.image.empty()) { entries[i].ok = false; entries[i].note = "not compiled: the race budget ran out"; }
}
// Load the modules (the context is current on this thread)
for (size_t i = 1; i < entries.size(); ++i) {
RaceEntry& e = entries[i];
if (!e.ok) continue;
CUresult r = c.drv.moduleLoadData(&e.mod, e.cb.image.data());
if (r != CUDA_SUCCESS) { e.ok = false; e.note = "cuModuleLoadData: " + c.err(r); e.mod = nullptr; continue; }
if (c.drv.moduleGetFunction(&e.fn, e.mod, e.cb.lowered[0].c_str()) != CUDA_SUCCESS) { e.ok = false; e.note = "function not in the module"; continue; }
c.drv.funcGetAttribute(&e.regs, CU_FUNC_ATTRIBUTE_NUM_REGS, e.fn);
c.drv.occupancy(&e.blocksPerSM, e.fn, 32 * e.blockWarps, 0);
}
double compileMs = wallMs() - t0;
// Time them: rounds over the entries, interleaved, best per entry. A pinned variant is only self-tested.
CUdeviceptr dOut = 0;
std::string benchErr;
if (c.drv.memAlloc(&dOut, (size_t)batch * 8u) != CUDA_SUCCESS) { benchErr = "cuMemAlloc for the race"; for (RaceEntry& e : entries) if (e.v.name != "base") e.ok = false; }
int rounds = pinnedOnly ? 1 : std::max(1, c.raceRounds);
for (int round = 0; round < rounds && benchErr.empty(); ++round) {
for (size_t i = 0; i < entries.size(); ++i) {
RaceEntry& e = entries[i];
if (!e.ok) continue;
if (i > 0 && round == 0 && wallMs() > deadline) { e.ok = false; e.note = "not timed: the race budget ran out"; continue; }
if (pinnedOnly && i == 0) continue;
if (!raceTime(c, p, pk, e, dOut, batch, pinnedOnly ? 0 : benchMs, s, round == 0)) e.ok = false;
// the mutex is not fair: give the job loop the card between windows (measured on the Mac, 4 October
// 2026: without this a queued job waited the whole race, 36 s)
std::this_thread::sleep_for(std::chrono::milliseconds(150));
}
}
if (dOut) c.drv.memFree(dOut);
// The winner: the fastest entry; base keeps its place unless a variant is at least 0.5% faster (noise guard).
size_t win = 0;
if (pinnedOnly && entries.size() == 2 && entries[1].ok) win = 1;
else for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok && entries[i].mhs > entries[win].mhs * (win == 0 ? 1.005 : 1.0)) win = i;
double baseMhs = entries[0].mhs, winMhs = entries[win].mhs;
if (win != 0) {
RaceEntry& w = entries[win];
c.drv.moduleUnload(p->modBound);
p->modBound = w.mod; p->fHashBound = w.fn; p->regs = w.regs; p->blocksPerSM = w.blocksPerSM; p->blockWarps = w.blockWarps; p->variant = w.v.name;
w.mod = nullptr;
} else {
p->blockWarps = c.blockWarps; p->variant = "base";
}
for (size_t i = 1; i < entries.size(); ++i) if (entries[i].mod) c.drv.moduleUnload(entries[i].mod);
p->raceMs = wallMs() - t0;
// The one line. Variants in race order: name=MH/s (regs), or name=- (why).
uint32_t loads = 0, wide = 0;
pf_define_u32(programH.c_str(), "IGNEUM_LOADS_PER_HASH", &loads);
pf_define_u32(programH.c_str(), "IGNEUM_WIDE_LOADS_PER_HASH", &wide);
std::string line = fmt("race %.16s device %s driver %d.%d arch %s loads %u wide %u variants %zu", p->epochHex.c_str(), c.name.c_str(), c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.archOpt.c_str(), loads, wide, entries.size());
for (const RaceEntry& e : entries) {
if (e.ok && (e.mhs > 0 || pinnedOnly)) line += fmt(" %s=%.3f/%dr/%dw", e.v.name.c_str(), e.mhs, e.regs, e.blockWarps);
else line += fmt(" %s=-", e.v.name.c_str());
}
line += fmt(" winner %s %.3f base %.3f gain %+.2f%% compile %.0f bench %.0f total %.0f ms%s%s", p->variant.c_str(), winMhs, baseMhs, baseMhs > 0 ? (winMhs / baseMhs - 1.0) * 100.0 : 0.0, compileMs, p->raceMs - compileMs, p->raceMs,
pinnedOnly ? " pinned by tuning" : (tu.found ? " tuned order" : ""), benchErr.empty() ? "" : (" " + benchErr).c_str());
for (const RaceEntry& e : entries) if (!e.ok && !e.note.empty()) line += " | " + e.v.name + ": " + e.note;
p->raceLine = line;
}
// Compiles the pack in `dir`, builds its cache and dataset on stream `s`, runs the self-test, races the variants
// (`race`). Returns the pair or null with `err` set. Runs on the main thread for --pack and --check, on the prepare
// thread for `prepare`.
static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string& err, bool race) {
PfPack pk;
char perr[512];
if (!pf_load(dir.c_str(), &pk, perr, sizeof(perr))) { err = std::string("pack ") + dir + ": " + perr; return nullptr; }
bool ok1, ok2, ok3, ok4;
std::string kernelCu = readText(dir + "/kernel.cu", ok1), boundCu = readText(dir + "/kernel_bound.cu", ok2);
std::string programH = readText(dir + "/program.h", ok3), memhardH = readText(dir + "/memhard.h", ok4);
if (!ok1 || !ok2 || !ok3 || !ok4) { err = "pack " + dir + " lacks kernel.cu, kernel_bound.cu, program.h or memhard.h"; return nullptr; }
std::string kernelDev, boundDev;
if (!deviceOnly(kernelCu, kernelDev, err) || !deviceOnly(boundCu, boundDev, err)) { err = "pack " + dir + ": " + err; return nullptr; }
Pair* p = new Pair();
p->dir = dir; p->epochHex = pk.epochHex; p->dayHex = pk.dayHex; p->seedString = pk.seedString;
std::memcpy(p->sw, pk.seedw, 32); std::memcpy(p->kw, pk.keyw, 32);
p->datasetLog2 = pk.datasetLog2; p->words = 1u << pk.datasetLog2; p->cacheWords = 1u << pk.cacheLog2Words; p->cacheSegments = pk.cacheSegments;
// Compile
Compiled ck, cb;
std::vector<std::string> kernelNames = { "igneum_cache_fill", "igneum_build" };
if (pk.hotMb) kernelNames.push_back("igneum_hot_fill");
if (!rtcCompile(c, kernelDev, "kernel.cu", programH, memhardH, kernelNames, ck, err)) { releasePair(c, p); return nullptr; }
if (!rtcCompile(c, boundDev, "kernel_bound.cu", programH, memhardH, { "igneum_hash_bound" }, cb, err)) { releasePair(c, p); return nullptr; }
p->compileMs = ck.ms + cb.ms;
// Load
{
CUresult r = c.drv.moduleLoadData(&p->modKernel, ck.image.data());
if (r != CUDA_SUCCESS) { err = "cuModuleLoadData kernel.cu (" + c.archOpt + "): " + c.err(r); releasePair(c, p); return nullptr; }
r = c.drv.moduleLoadData(&p->modBound, cb.image.data());
if (r != CUDA_SUCCESS) { err = "cuModuleLoadData kernel_bound.cu (" + c.archOpt + "): " + c.err(r); releasePair(c, p); return nullptr; }
if (c.drv.moduleGetFunction(&p->fCacheFill, p->modKernel, ck.lowered[0].c_str()) != CUDA_SUCCESS) { err = "igneum_cache_fill (" + ck.lowered[0] + ") not in the module"; releasePair(c, p); return nullptr; }
if (c.drv.moduleGetFunction(&p->fBuild, p->modKernel, ck.lowered[1].c_str()) != CUDA_SUCCESS) { err = "igneum_build (" + ck.lowered[1] + ") not in the module"; releasePair(c, p); return nullptr; }
if (c.drv.moduleGetFunction(&p->fHashBound, p->modBound, cb.lowered[0].c_str()) != CUDA_SUCCESS) { err = "igneum_hash_bound (" + cb.lowered[0] + ") not in the module"; releasePair(c, p); return nullptr; }
if (pk.hotMb && c.drv.moduleGetFunction(&p->fHotFill, p->modKernel, ck.lowered[2].c_str()) != CUDA_SUCCESS) { err = "igneum_hot_fill (" + ck.lowered[2] + ") not in the module"; releasePair(c, p); return nullptr; }
c.drv.funcGetAttribute(&p->regs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->fHashBound);
c.drv.occupancy(&p->blocksPerSM, p->fHashBound, 32 * c.blockWarps, 0);
}
p->loadClass = pk.loadClass; p->loadsPerHash = pk.loadsPerHash; p->bytesPerHash = pk.bytesPerHash; p->scratchOps = pk.scratchOps;
p->programClass = pk.programClass; p->eraHex = pk.eraHex;
p->persistent = pk.persistent != 0;
p->hotMb = pk.hotMb; p->hotWords = pk.hotWords; p->hotSegments = pk.hotSegments; p->hotSlots = pk.hotSlots;
p->residentWarps = p->blocksPerSM * c.blockWarps * c.sms;
size_t scratchBytes = 0, hotBytes = (size_t)pk.hotWords * 4u;
if (p->persistent) {
// Variant 5: one arena per launched warp. The launch is the resident capacity (the occupancy query), rounded
// down to a power of two so it divides every batch, or --warps; the allocation cannot change the occupancy
// (registers and shared memory decide it), and the number is re-queried after the allocation below to show it.
if (c.blockWarps != 1) { err = "a variant-5 pack runs one warp per block: use --block-warps 1"; releasePair(c, p); return nullptr; }
int w = c.warps > 0 ? c.warps : p->residentWarps;
int pw = 1; while (pw * 2 <= w) pw *= 2;
p->warps = pw;
scratchBytes = (size_t)p->warps * 32u * (size_t)pk.scratchWordsPerLane * 4u;
}
// Cache
double t0 = wallMs();
size_t cacheBytes = (size_t)p->cacheWords * 4u, dsBytes = (size_t)p->words * 4u;
{
size_t freeB = 0, totalB = 0;
if (c.drv.memGetInfo(&freeB, &totalB) == CUDA_SUCCESS && freeB < cacheBytes + dsBytes + scratchBytes + hotBytes + (64u << 20)) {
err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu + scratch %llu + hot %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes + scratchBytes + hotBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20), (unsigned long long)(scratchBytes >> 20), (unsigned long long)(hotBytes >> 20));
releasePair(c, p); return nullptr;
}
}
if (p->persistent) {
CUresult r = c.drv.memAlloc(&p->scratch, scratchBytes);
if (r != CUDA_SUCCESS) { err = "cuMemAlloc scratch: " + c.err(r); p->scratch = 0; releasePair(c, p); return nullptr; }
p->scratchBytes = scratchBytes;
int after = 0;
c.drv.occupancy(&after, p->fHashBound, 32 * c.blockWarps, 0);
info(fmt("variant 5: %d persistent warps (resident capacity %d = %d blocks/SM x %d warps/block x %d SMs; occupancy query after the allocation %d blocks/SM), scratch %llu MiB (%u KiB per warp)",
p->warps, p->residentWarps, p->blocksPerSM, c.blockWarps, c.sms, after, (unsigned long long)(scratchBytes >> 20), pk.scratchWordsPerLane * 4u * 32u / 1024u));
}
{
CUresult r = c.drv.memAlloc(&p->cache, cacheBytes);
if (r != CUDA_SUCCESS) { err = "cuMemAlloc cache: " + c.err(r); p->cache = 0; releasePair(c, p); return nullptr; }
uint32_t nSeg = p->cacheSegments, block = 256u, grid = (nSeg + block - 1u) / block;
void* args[2] = { &p->cache, &nSeg };
r = c.drv.launchKernel(p->fCacheFill, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
if (r != CUDA_SUCCESS) { err = "cache fill: " + c.err(r); releasePair(c, p); return nullptr; }
}
p->cacheMs = wallMs() - t0;
// Dataset
t0 = wallMs();
{
CUresult r = c.drv.memAlloc(&p->ds, dsBytes);
if (r != CUDA_SUCCESS) { err = "cuMemAlloc dataset: " + c.err(r); p->ds = 0; releasePair(c, p); return nullptr; }
uint32_t nItems = p->words / 16u, block = 256u, grid = (nItems + block - 1u) / block;
void* args[3] = { &p->ds, &p->cache, &nItems };
r = c.drv.launchKernel(p->fBuild, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
if (r != CUDA_SUCCESS) { err = "dataset build: " + c.err(r); releasePair(c, p); return nullptr; }
}
p->dsMs = wallMs() - t0;
// Hot table (hot-table experiment): filled from the epoch seed by the pack's own kernel, never shipped
if (pk.hotMb) {
t0 = wallMs();
CUresult r = c.drv.memAlloc(&p->hot, hotBytes);
if (r != CUDA_SUCCESS) { err = "cuMemAlloc hot table: " + c.err(r); p->hot = 0; releasePair(c, p); return nullptr; }
uint32_t nSeg = pk.hotSegments, block = 256u, grid = (nSeg + block - 1u) / block;
void* args[2] = { &p->hot, &nSeg };
r = c.drv.launchKernel(p->fHotFill, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
if (r != CUDA_SUCCESS) { err = "hot table fill: " + c.err(r); releasePair(c, p); return nullptr; }
p->hotMs = wallMs() - t0;
}
// Self-test against vectors.h
t0 = wallMs();
if (!pk.haveVectors) {
p->checked = false; p->checkPass = true;
p->check = "self-test skipped (no vectors.h in the pack); the miner's CPU re-check covers every found nonce";
} else {
uint32_t cacheHead[16], cacheLast[16], dsHead[16], dsLast = 0;
std::vector<uint32_t> samples((size_t)(pk.nSamples > 0 ? pk.nSamples : 1), 0u);
std::vector<uint64_t> vec((size_t)pk.vecWarps * 32u, 0ull);
std::vector<uint32_t> whole(p->cacheWords);
CUresult r = c.drv.memcpyDtoH(whole.data(), p->cache, cacheBytes);
if (r != CUDA_SUCCESS) { err = "cuMemcpyDtoH cache: " + c.err(r); releasePair(c, p); return nullptr; }
std::memcpy(cacheHead, whole.data(), 64);
std::memcpy(cacheLast, whole.data() + p->cacheWords - 16u, 64);
uint64_t fnv = pf_fnv1a64(whole.data(), cacheBytes);
whole.clear(); whole.shrink_to_fit();
if ((r = c.drv.memcpyDtoH(dsHead, p->ds, 64)) != CUDA_SUCCESS) { err = "cuMemcpyDtoH dataset head: " + c.err(r); releasePair(c, p); return nullptr; }
if (pk.dsLastIndex < p->words) r = c.drv.memcpyDtoH(&dsLast, p->ds + (CUdeviceptr)pk.dsLastIndex * 4u, 4);
for (int i = 0; i < pk.nSamples && r == CUDA_SUCCESS; ++i) if (pk.sampleIdx[i] < p->words) r = c.drv.memcpyDtoH(&samples[(size_t)i], p->ds + (CUdeviceptr)pk.sampleIdx[i] * 4u, 4);
if (r != CUDA_SUCCESS) { err = "cuMemcpyDtoH dataset words: " + c.err(r); releasePair(c, p); return nullptr; }
CUdeviceptr out = 0;
if ((r = c.drv.memAlloc(&out, 32u * (size_t)c.blockWarps * 8u)) != CUDA_SUCCESS) { err = "cuMemAlloc vector out: " + c.err(r); releasePair(c, p); return nullptr; }
for (int w = 0; w < pk.vecWarps; ++w) {
// One block of 32 x block-warps lanes; the vector warp is its first 32 lanes (lane nonce = base + gid)
if (!launchHash(c, p, out, pk.vecBase[w], p->sw, 32u * (uint32_t)c.blockWarps, 32u * (uint32_t)c.blockWarps, s, err)) { c.drv.memFree(out); releasePair(c, p); return nullptr; }
if ((r = c.drv.streamSynchronize(s)) != CUDA_SUCCESS || (r = c.drv.memcpyDtoH(&vec[(size_t)w * 32u], out, 32u * 8u)) != CUDA_SUCCESS) { err = "vector warp: " + c.err(r); c.drv.memFree(out); releasePair(c, p); return nullptr; }
}
c.drv.memFree(out);
uint32_t hotHead[16] = {0}, hotLast[16] = {0};
uint64_t hotFnv = 0;
if (p->hot) {
std::vector<uint32_t> hw(p->hotWords);
if ((r = c.drv.memcpyDtoH(hw.data(), p->hot, hotBytes)) != CUDA_SUCCESS) { err = "cuMemcpyDtoH hot table: " + c.err(r); releasePair(c, p); return nullptr; }
std::memcpy(hotHead, hw.data(), 64);
std::memcpy(hotLast, hw.data() + p->hotWords - 16u, 64);
hotFnv = pf_fnv1a64(hw.data(), hotBytes);
}
char line[1024];
p->checkPass = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples.data(), vec.data(), p->hot ? hotHead : nullptr, p->hot ? hotLast : nullptr, hotFnv, line, sizeof(line)) != 0;
p->checked = true;
p->check = line;
}
p->checkMs = wallMs() - t0;
if (!p->checkPass) { err = p->check; releasePair(c, p); return nullptr; }
p->blockWarps = c.blockWarps;
if (race && c.race != "off") racePair(c, p, pk, boundDev, programH, memhardH, s);
return p;
}
static std::string pairSummary(const Pair* p) {
return fmt("nvrtc %.0f cache %.0f dataset %.0f hot %.0f check %.0f race %.0f ms variant %s; %s", p->compileMs, p->cacheMs, p->dsMs, p->hotMs, p->checkMs, p->raceMs, p->variant.c_str(), p->check.c_str());
}
// ---------------------------------------------------------------------------------------------
// Prepare, on its own thread
struct PrepareTask {
std::string epochHex, dayHex, dir, error;
std::string wantClass, wantEra; // the class and era the prepare line named (empty: any)
std::atomic<bool> done{false};
Pair* result = nullptr;
double t0 = 0;
std::thread thread;
};
static void prepareRun(Ctx* c, PrepareTask* t) {
CUstream s = nullptr;
std::string err;
if (c->drv.ctxSetCurrent(c->ctx) != CUDA_SUCCESS) { t->error = "cuCtxSetCurrent on the prepare thread"; t->done = true; return; }
if (c->drv.streamCreate(&s, CU_STREAM_NON_BLOCKING) != CUDA_SUCCESS) { t->error = "cuStreamCreate on the prepare thread"; t->done = true; return; }
Pair* p = buildPair(*c, t->dir, s, err, true);
c->drv.streamDestroy(s);
if (p && (p->epochHex != t->epochHex || p->dayHex != t->dayHex)) {
err = "the pack in " + t->dir + " is for epoch " + p->epochHex.substr(0, 16) + " day " + p->dayHex + ", not the prepared seeds";
releasePair(*c, p); p = nullptr;
}
char why[256];
if (p && !pf_pack_class_ok(p->programClass.c_str(), p->eraHex.c_str(), t->wantClass.c_str(), t->wantEra.c_str(), why, sizeof(why))) {
err = "pack " + t->dir + ": " + why;
releasePair(*c, p); p = nullptr;
}
t->result = p;
t->error = err;
t->done = true;
}
// ---------------------------------------------------------------------------------------------
// Serve
struct Options {
bool serve = false, check = false, raceOnly = false;
bool bench = false, memprobe = false; // read-width experiment (5 October 2026)
int batches = 5, warps = 0, probeMib = 0;
int device = 0, batchLog2 = 22, blockWarps = 1;
std::string pack, arch = "auto";
std::string race = "on", pinned, tuningPath;
int raceBenchMs = 2000, raceBudgetS = 120, raceRounds = 0; // rounds 0 = 1 in --serve, 3 in --race
};
static void usage() {
std::printf("igneum-worker-cuda %s\n"
" --serve --pack <dir> GPU worker for igneum-miner --worker: jobs on stdin, found/done lines on stdout\n"
" --check --pack <dir> compile the pack, build its cache and dataset, self-test, print timings, exit 0 or 1\n"
" --device D CUDA device index (default 0)\n"
" --batch-log2 B nonces per dispatch = 2^B (default 22)\n"
" --block-warps W warps per thread block (default 1)\n"
" --arch sm_XY|compute_XY|auto NVRTC target (default auto: the device's architecture)\n"
" --bench --pack <dir> read-width experiment: build and self-test the pack, time --batches dispatches of 2^B nonces,\n"
" print the 2^B fingerprint at base nonce 0 (one RESULT line); a variant-5 pack runs --warps persistent warps\n"
" --memprobe [--probe-mib N] no pack: dependent random 4, 16 and 64-byte reads, independent reads, a coalesced stream and an\n"
" integer chain at 4, 64 and 1024 MiB (the same table as igneum-worker-opencl --memprobe)\n"
" --batches N --bench: timed dispatches (default 5)\n"
" --warps N --bench on a variant-5 pack: persistent warps (default: the occupancy capacity, rounded down to a power of two)\n"
" --race --pack <dir> the variant race alone (3 rounds): one line per variant, the race line, exit 0 or 1\n"
" --race on|off|a,b,c in --serve: race every variant (default), none, or these names\n"
" --race-bench-ms N timed window per variant (default 2000)\n"
" --race-budget-s N a race stops compiling and timing after this (default 120; base is kept)\n"
" --race-rounds N interleaved rounds, best per variant (default 1 in --serve, 3 in --race)\n"
" --variant <name> use this variant without a race (also from the tuning file)\n"
" --tuning <file> the per-card tuning file (default: IGNEUM_TUNING_FILE from the environment)\n", WORKER_VERSION);
}
static Options parseArgs(int argc, char** argv) {
Options o;
for (int i = 1; i < argc; ++i) {
std::string a = argv[i];
auto next = [&]() -> std::string { if (i + 1 >= argc) { usage(); std::exit(2); } return argv[++i]; };
if (a == "--serve") o.serve = true;
else if (a == "--check") o.check = true;
else if (a == "--bench") o.bench = true;
else if (a == "--memprobe") o.memprobe = true;
else if (a == "--batches") o.batches = std::atoi(next().c_str());
else if (a == "--warps") o.warps = std::atoi(next().c_str());
else if (a == "--probe-mib") o.probeMib = std::atoi(next().c_str());
else if (a == "--race" && (i + 1 >= argc || std::string(argv[i + 1]).rfind("--", 0) == 0)) o.raceOnly = true;
else if (a == "--race") o.race = next();
else if (a == "--race-bench-ms") o.raceBenchMs = std::atoi(next().c_str());
else if (a == "--race-budget-s") o.raceBudgetS = std::atoi(next().c_str());
else if (a == "--race-rounds") o.raceRounds = std::atoi(next().c_str());
else if (a == "--variant") o.pinned = next();
else if (a == "--tuning") o.tuningPath = next();
else if (a == "--pack") o.pack = next();
else if (a == "--device") o.device = std::atoi(next().c_str());
else if (a == "--batch-log2") o.batchLog2 = std::atoi(next().c_str());
else if (a == "--block-warps") o.blockWarps = std::atoi(next().c_str());
else if (a == "--arch") o.arch = next();
else if (a == "--no-prepare") { /* accepted for symmetry with the other workers; prepare is always on here */ }
else if (a == "-h" || a == "--help") { usage(); std::exit(0); }
else { std::printf("unknown argument %s\n", argv[i]); usage(); std::exit(2); }
}
if (o.batchLog2 < 10 || o.batchLog2 > 28) { std::printf("--batch-log2 must be between 10 and 28\n"); std::exit(2); }
if (o.blockWarps < 1 || o.blockWarps > 32) { std::printf("--block-warps must be between 1 and 32\n"); std::exit(2); }
if (!o.serve && !o.check && !o.raceOnly && !o.bench && !o.memprobe) { usage(); std::exit(2); }
if (o.raceBenchMs < 200 || o.raceBenchMs > 20000) { std::printf("--race-bench-ms must be between 200 and 20000\n"); std::exit(2); }
if (o.raceBudgetS < 5 || o.raceBudgetS > 540) { std::printf("--race-budget-s must be between 5 and 540 (the prepare lead is 600 DAA)\n"); std::exit(2); }
if (o.raceRounds == 0) o.raceRounds = o.raceOnly ? 3 : 1;
if (o.tuningPath.empty()) if (const char* t = std::getenv("IGNEUM_TUNING_FILE")) o.tuningPath = t;
if (o.pack.empty() && !o.memprobe) { std::printf("--pack <dir> is required (igneum-miner export-pack <node> <dir> writes one)\n"); std::exit(2); }
while (o.pack.size() > 1 && (o.pack.back() == '/' || o.pack.back() == '\\')) o.pack.pop_back();
return o;
}
static std::vector<std::string> split(const std::string& line) {
std::vector<std::string> f;
size_t i = 0;
while (i < line.size()) {
while (i < line.size() && (line[i] == ' ' || line[i] == '\t' || line[i] == '\r')) ++i;
size_t j = i;
while (j < line.size() && line[j] != ' ' && line[j] != '\t' && line[j] != '\r') ++j;
if (j > i) f.push_back(line.substr(i, j - i));
i = j;
}
return f;
}
// A pack directory for the given seeds under `root` (one subdirectory per pack, each with seeds.txt), or "".
static std::string findPackFor(const std::string& root, const std::string& epochHex, const std::string& dayHex) {
if (root.empty()) return "";
std::vector<std::string> names;
#ifdef _WIN32
WIN32_FIND_DATAA fd;
HANDLE h = FindFirstFileA((root + "\\*").c_str(), &fd);
if (h == INVALID_HANDLE_VALUE) return "";
do { if ((fd.dwFileAttributes & FILE_ATTRIBUTE_DIRECTORY) && fd.cFileName[0] != '.') names.push_back(root + "\\" + fd.cFileName); } while (FindNextFileA(h, &fd));
FindClose(h);
#else
DIR* d = opendir(root.c_str());
if (!d) return "";
while (dirent* e = readdir(d)) if (e->d_name[0] != '.') names.push_back(root + "/" + e->d_name);
closedir(d);
#endif
for (const std::string& dir : names) {
bool ok = false;
std::string seeds = readText(dir + "/seeds.txt", ok);
if (!ok) continue;
char e[65] = {0}, dd[PF_HEX_CAP] = {0};
if (pf_seeds_line(seeds.c_str(), "epoch_seed_hex", e, sizeof(e)) && pf_seeds_line(seeds.c_str(), "day_seed_hex", dd, sizeof(dd)) && epochHex == e && dayHex == dd) return dir;
}
return "";
}
static std::string parentDir(const std::string& p) { size_t i = p.find_last_of("/\\"); return i == std::string::npos ? "." : p.substr(0, i); }
static int runServe(Ctx& c, const Options& o, Pair* cur) {
const uint32_t batch = 1u << o.batchLog2;
CUdeviceptr dOut = 0;
std::string err;
if (c.drv.memAlloc(&dOut, (size_t)batch * 8u) != CUDA_SUCCESS) { emit("error 0 cuMemAlloc out buffer"); return 2; }
std::vector<uint64_t> hOut(batch);
Pair* prepared = nullptr;
Pair* old = nullptr;
PrepareTask* task = nullptr;
std::string prepareRoot; // the parent of the last prepare's pack directory: where the miner writes its packs
emit(fmt("ready cuda %s pack %s dataset-log2 %u batch %u regs %d prepare 1 path nvrtc %d.%d driver %d.%d arch %s variant %s race %s worker %s",
c.name.c_str(), cur->seedString.c_str(), cur->datasetLog2, batch, cur->regs, c.rtcMajor, c.rtcMinor, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.archOpt.c_str(), cur->variant.c_str(), c.race.c_str(), WORKER_VERSION));
info(fmt("first pack %s: %s", cur->dir.c_str(), pairSummary(cur).c_str()));
if (!cur->raceLine.empty()) emit(cur->raceLine);
std::string line;
while (std::getline(std::cin, line)) {
if (line == "quit") break;
std::vector<std::string> f = split(line);
if (f.empty()) continue;
// A finished prepare is reported here, between lines
if (task && task->done) {
task->thread.join();
if (task->result) {
if (prepared) releasePair(c, prepared);
prepared = task->result;
if (!prepared->raceLine.empty()) emit(prepared->raceLine);
emit(fmt("prepared %s %s %.1f %s resident 2 programs 2 datasets", prepared->epochHex.c_str(), prepared->dayHex.c_str(), wallMs() - task->t0, pairSummary(prepared).c_str()));
} else {
emit(fmt("prepare-failed %s %s %s", task->epochHex.c_str(), task->dayHex.c_str(), task->error.c_str()));
}
delete task; task = nullptr;
}
std::string wantClass, wantEra;
takeClassTokens(f, wantClass, wantEra);
if (f[0] == "prepare") {
if (f.size() < 4) { emit(fmt("prepare-failed %s %s a pack directory is needed as the third field (igneum-miner --prepare-packs <dir>)", f.size() > 1 ? f[1].c_str() : "0", f.size() > 2 ? f[2].c_str() : "0")); continue; }
if (f[1].size() != 64) { emit(fmt("prepare-failed %s %s bad field (epoch_seed 64 hex, day_seed hex)", f[1].c_str(), f[2].c_str())); continue; }
if (task) { emit(fmt("prepare-failed %s %s a prepare is still running", f[1].c_str(), f[2].c_str())); continue; }
if (prepared && prepared->epochHex == f[1] && prepared->dayHex == f[2]) { emit(fmt("prepared %s %s 0 (already resident)", f[1].c_str(), f[2].c_str())); continue; }
if (cur->epochHex == f[1] && cur->dayHex == f[2]) { emit(fmt("prepared %s %s 0 (already the current pair)", f[1].c_str(), f[2].c_str())); continue; }
std::string dir = f[3];
for (size_t i = 4; i < f.size(); ++i) dir += " " + f[i]; // a directory with spaces arrives as several fields
prepareRoot = parentDir(dir);
task = new PrepareTask();
task->epochHex = f[1]; task->dayHex = f[2]; task->dir = dir; task->t0 = wallMs();
task->wantClass = wantClass; task->wantEra = wantEra;
task->thread = std::thread(prepareRun, &c, task);
info(fmt("prepare started for epoch %.16s day %s from %s (NVRTC %s in the background)", f[1].c_str(), f[2].c_str(), dir.c_str(), c.archOpt.c_str()));
continue;
}
if (f[0] != "job") { info("ignored: " + line); continue; }
std::string jobId = f.size() > 1 ? f[1] : "0";
if (f.size() < 8) { emit("error " + jobId + " malformed job line (need 7 fields after job)"); continue; }
uint8_t prehash[32], epochSeed[32], daySeed[256];
size_t pl = 0, el = 0, dl = 0;
unsigned long long target = 0, nonceStart = 0, nonceCount = 0;
if (!pf_unhex(f[2].c_str(), prehash, 32, &pl) || pl != 32 || std::sscanf(f[3].c_str(), "%llx", &target) != 1 ||
std::sscanf(f[4].c_str(), "%llu", &nonceStart) != 1 || std::sscanf(f[5].c_str(), "%llu", &nonceCount) != 1 ||
!pf_unhex(f[6].c_str(), epochSeed, 32, &el) || el != 32 || !pf_unhex(f[7].c_str(), daySeed, sizeof(daySeed), &dl)) {
emit("error " + jobId + " bad field (prehash 64 hex, target 16 hex, nonce_start u64, nonce_count u64, epoch_seed 64 hex, day_seed hex)"); continue;
}
if (nonceCount == 0 || nonceCount % 32 != 0 || (nonceStart & 31) != 0) { emit("error " + jobId + " nonce_start must be 32-aligned and nonce_count a non-zero multiple of 32"); continue; }
uint32_t sw[8], kw[8];
pf_seed_words_from_bytes(epochSeed, 32, sw);
pf_seed_words_from_bytes(daySeed, dl, kw);
double t0 = wallMs();
bool switched = false;
if (!pairIsClass(cur, f[6], f[7], wantClass, wantEra) && !task && !pairIsClass(prepared, f[6], f[7], wantClass, wantEra)) {
// Self-heal: a job on seeds this worker has no pair for and no prepare in flight (a prepare failed, or
// the miner never sent one). The miner writes a pack per pair under its --prepare-packs root; find it by
// seeds.txt and build it now, in the foreground. The miner only re-sends prepare for the pair after this one.
std::string dir = findPackFor(prepareRoot, f[6], f[7]);
if (dir.empty()) dir = findPackFor(parentDir(o.pack) + "/prepare", f[6], f[7]);
if (dir.empty()) dir = findPackFor(parentDir(o.pack), f[6], f[7]);
if (!dir.empty()) {
info(fmt("job %s is for epoch %.16s day %s, which is not resident; building its pack %s now (foreground)", jobId.c_str(), f[6].c_str(), f[7].c_str(), dir.c_str()));
std::string berr;
Pair* p = buildPair(c, dir, nullptr, berr, true);
if (p && (p->epochHex != f[6] || p->dayHex != f[7])) { berr = "the pack in " + dir + " is for other seeds"; releasePair(c, p); p = nullptr; }
if (p) { if (prepared) releasePair(c, prepared); prepared = p; info(fmt("built %s: %s", dir.c_str(), pairSummary(p).c_str())); if (!p->raceLine.empty()) emit(p->raceLine); }
else emit("error " + jobId + " could not build " + dir + ": " + berr);
}
}
if (!pairIsClass(cur, f[6], f[7], wantClass, wantEra)) {
char why[256];
if (pairIsClass(prepared, f[6], f[7], wantClass, wantEra)) {
if (old) releasePair(c, old);
old = cur; cur = prepared; prepared = nullptr; switched = true;
info(fmt("switched to the prepared pair epoch %.16s day %s (class %s) in %.2f ms", cur->epochHex.c_str(), cur->dayHex.c_str(), cur->programClass.c_str(), wallMs() - t0));
} else if (pairIs(cur, f[6], f[7]) && !pf_pack_class_ok(cur->programClass.c_str(), cur->eraHex.c_str(), wantClass.c_str(), wantEra.c_str(), why, sizeof(why))) {
// Counter ASIC 2.0: right seeds, wrong class or era. The miner prepares the pair again from a pack of
// the class the chain is on (the `need` line), and the pack of the other class is never mined.
emit(fmt("need %s %s", f[6].c_str(), f[7].c_str()));
emit(fmt("error %s pack %s: %s", jobId.c_str(), cur->dir.c_str(), why));
continue;
} else if (!hexEq(cur->epochHex, f[6])) {
emit(fmt("need %s %s", f[6].c_str(), f[7].c_str())); // the miner prepares this pair (4 October 2026)
emit(fmt("error %s epoch seed mismatch: this worker holds epoch %.16s (program words %08x %08x ...)%s, the job is for epoch %.16s (bare seed words %08x %08x ...); send prepare with a pack directory",
jobId.c_str(), cur->epochHex.c_str(), cur->sw[0], cur->sw[1], prepared ? " plus one prepared pair" : "", f[6].c_str(), sw[0], sw[1]));
continue;
} else {
emit(fmt("need %s %s", f[6].c_str(), f[7].c_str()));
emit(fmt("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",
jobId.c_str(), cur->kw[0], cur->kw[1], f[7].c_str(), kw[0], kw[1]));
continue;
}
}
uint64_t remaining = nonceCount, hashes = 0;
uint32_t hi = (uint32_t)(nonceStart >> 32), lo = (uint32_t)nonceStart;
bool failed = false;
while (remaining > 0) {
uint64_t room = (uint64_t)(0xffffffffu - lo) + 1ull;
uint64_t chunk64 = remaining < batch ? remaining : batch;
if (chunk64 > room) chunk64 = room;
uint32_t chunk = (uint32_t)chunk64;
uint32_t iw[8];
{
uint8_t b[49];
std::memcpy(b, "igneum-block/", 13);
std::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);
pf_seed_words_from_bytes(b, 49, iw);
}
// A chunk that is not a multiple of the block is finished one 32-lane block at a time. The block is the
// pair's (its winning variant's). The mutex gives a race its exclusive windows between chunks.
std::lock_guard<std::mutex> hold(gpuMutex);
uint32_t block = 32u * (uint32_t)cur->blockWarps;
uint32_t main = chunk - (chunk % block);
CUresult r = CUDA_SUCCESS;
if (main > 0 && !launchHash(c, cur, dOut, lo, iw, main, block, nullptr, err)) { emit("error " + jobId + " dispatch failed: " + err); failed = true; break; }
if (main < chunk) {
for (uint32_t off = main; off < chunk && !failed; off += 32u) if (!launchHash(c, cur, dOut + (CUdeviceptr)off * 8u, lo + off, iw, 32u, 32u, nullptr, err)) { emit("error " + jobId + " dispatch failed: " + err); failed = true; }
if (failed) break;
}
r = c.drv.streamSynchronize(nullptr);
if (r == CUDA_SUCCESS) r = c.drv.memcpyDtoH(hOut.data(), dOut, (size_t)chunk * 8u);
if (r != CUDA_SUCCESS) { emit("error " + jobId + " dispatch failed: " + c.err(r)); failed = true; break; }
for (uint32_t i = 0; i < chunk; ++i) if (hOut[i] <= target) {
uint64_t nonce = ((uint64_t)hi << 32) | (uint64_t)(uint32_t)(lo + i);
std::printf("found %s %llu %016llx\n", jobId.c_str(), (unsigned long long)nonce, (unsigned long long)hOut[i]);
}
std::fflush(stdout);
hashes += chunk;
remaining -= chunk;
if (chunk64 == room) { hi += 1u; lo = 0u; } else lo += chunk;
}
if (failed) continue;
emit(fmt("done %s %llu %.2f", jobId.c_str(), (unsigned long long)hashes, wallMs() - t0));
if (switched && old) { releasePair(c, old); old = nullptr; info("dropped the previous pair (its program, cache and dataset)"); }
}
if (task) { task->thread.join(); if (task->result) releasePair(c, task->result); delete task; }
c.drv.memFree(dOut);
if (old) releasePair(c, old);
if (prepared) releasePair(c, prepared);
releasePair(c, cur);
return 0;
}
// ---------------------------------------------------------------------------------------------
// Main
// ---------------------------------------------------------------------------------------------
// Read-width experiment (5 October 2026, docs/plans/read-width.md): --bench and --memprobe
static uint64_t fnv1a64Bytes(const void* p, size_t n) {
const uint8_t* b = (const uint8_t*)p;
uint64_t h = 0xcbf29ce484222325ull;
for (size_t i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
return h;
}
// --bench: the pair is built and self-tested (vectors through the bound kernel with the seed words); then a warm-up
// dispatch at base nonce 0 (fingerprinted) and --batches timed dispatches of 2^B nonces, wall time around
// cuStreamSynchronize (the driver API path loads no event symbols; a 2^24 dispatch is 60 to 900 ms on the cards here,
// so the launch overhead is under 1 percent).
static int runBench(Ctx& c, const Options& o, Pair* p) {
uint32_t nonces = 1u << o.batchLog2, block = 32u * (uint32_t)c.blockWarps;
if (p->persistent) { uint32_t unit = 32u * (uint32_t)p->warps; nonces = (nonces / unit) * unit; if (nonces == 0) nonces = unit; }
CUdeviceptr dOut = 0;
std::string err;
if (c.drv.memAlloc(&dOut, (size_t)nonces * 8u) != CUDA_SUCCESS) { std::printf("FAIL: cuMemAlloc out\n"); return 2; }
std::vector<uint64_t> hOut(nonces);
double sum = 0, warm = 0;
uint64_t fp = 0;
for (int b = -1; b < o.batches; ++b) {
double t0 = wallMs();
if (!launchHash(c, p, dOut, (uint32_t)(b + 1) * nonces, p->sw, nonces, block, nullptr, err)) { std::printf("FAIL: %s\n", err.c_str()); return 2; }
CUresult r = c.drv.streamSynchronize(nullptr);
if (r != CUDA_SUCCESS) { std::printf("FAIL: dispatch %d: %s\n", b, c.err(r).c_str()); return 2; }
double ms = wallMs() - t0;
if (b < 0) {
warm = ms;
if (c.drv.memcpyDtoH(hOut.data(), dOut, (size_t)nonces * 8u) != CUDA_SUCCESS) { std::printf("FAIL: read-back\n"); return 2; }
fp = fnv1a64Bytes(hOut.data(), (size_t)nonces * 8u);
} else sum += ms;
}
c.drv.memFree(dOut);
std::string dev = c.name; for (char& ch : dev) if (ch == ' ') ch = '_';
std::printf("warm-up dispatch (base 0): %.2f ms; %d timed dispatches of %u nonces: mean %.2f ms\n", warm, o.batches, nonces, sum / o.batches);
std::printf("RESULT pack=%s class=%s device=%s arch=%s regs=%d blocks_per_sm=%d warps=%d resident=%d 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=wall\n",
p->dir.c_str(), p->loadClass.c_str(), dev.c_str(), c.archOpt.c_str(), p->regs, p->blocksPerSM, p->warps, p->residentWarps, (unsigned long long)(p->scratchBytes >> 20), p->hotMb, p->hotSlots, p->hotMs, nonces, o.batches,
p->checked ? (p->checkPass ? "PASS" : "FAIL") : "skipped", (unsigned long long)fp, (double)nonces * (double)o.batches / (sum / 1000.0) / 1e6,
p->loadsPerHash, p->bytesPerHash, p->scratchOps * 8u);
return 0;
}
// --memprobe: the OpenCL worker's table (proto-opencl/host.c, 5 October 2026) in CUDA C through NVRTC, so the two
// vendors are probed with the same access patterns: a dependent chain of random 4-byte reads (the hash's pattern),
// eight independent chains per lane, dependent random 16-byte and 64-byte reads, a coalesced stream and an integer
// chain, at 4, 64 and 1024 MiB. Wall time around cuStreamSynchronize, best of 3, a fresh seed per repetition.
static const char* PROBE_CUDA =
"#include <cstdint>\n"
"__device__ __forceinline__ uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
"extern \"C\" __global__ void probe_fill(uint32_t* ds, uint32_t n) { uint32_t i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) ds[i] = pm_mix(i ^ 0x9E3779B9u); }\n"
"extern \"C\" __global__ void probe_chase(const uint32_t* ds, uint32_t mask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
" for (uint32_t s = 0u; s < steps; ++s) x = ds[x & mask] ^ (x * 0x9E3779B1u + s);\n"
" out[g] = x;\n"
"}\n"
"extern \"C\" __global__ void probe_indep(const uint32_t* ds, uint32_t mask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;\n"
" uint32_t 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"
" uint32_t 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 (uint32_t 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"
"extern \"C\" __global__ void probe_line16(const uint4* ds, uint32_t vecMask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
" for (uint32_t s = 0u; s < steps; ++s) { uint4 a = ds[x & vecMask]; x = (a.x ^ a.y ^ a.z ^ a.w) ^ (x * 0x9E3779B1u + s); }\n"
" out[g] = x;\n"
"}\n"
"extern \"C\" __global__ void probe_line(const uint4* ds, uint32_t lineMask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
" for (uint32_t s = 0u; s < steps; ++s) { uint32_t l = (x & lineMask) * 4u; uint4 a = ds[l], b = ds[l + 1u], c = ds[l + 2u], d = ds[l + 3u]; x = (a.x ^ b.y ^ c.z ^ d.w) ^ (x * 0x9E3779B1u + s); }\n"
" out[g] = x;\n"
"}\n"
"extern \"C\" __global__ void probe_stream(const uint4* ds, uint32_t perLane, uint32_t* out) {\n"
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x, n = gridDim.x * blockDim.x; uint4 acc = make_uint4(0u, 0u, 0u, 0u);\n"
" for (uint32_t s = 0u; s < perLane; ++s) { uint4 v = ds[s * n + g]; acc.x ^= v.x; acc.y ^= v.y; acc.z ^= v.z; acc.w ^= v.w; }\n"
" out[g] = acc.x ^ acc.y ^ acc.z ^ acc.w;\n"
"}\n"
"extern \"C\" __global__ void probe_alu(uint32_t steps, uint32_t seed, uint32_t* out) {\n"
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n"
" for (uint32_t s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + ((y << 7u) | (y >> 25u)); y = (y ^ x) + s; }\n"
" out[g] = x ^ y;\n"
"}\n";
static double probeLaunch(Ctx& c, CUfunction f, size_t lanes, size_t local, int reps, int seedArg, uint32_t seed, void** args) {
double best = -1;
for (int r = 0; r < reps; ++r) {
uint32_t s = seed + (uint32_t)r * 0x9E3779B9u;
if (seedArg >= 0) args[seedArg] = &s;
double t0 = wallMs();
if (c.drv.launchKernel(f, (unsigned)(lanes / local), 1, 1, (unsigned)local, 1, 1, 0, nullptr, args, nullptr) != CUDA_SUCCESS) return -1;
if (c.drv.streamSynchronize(nullptr) != CUDA_SUCCESS) return -1;
double ms = wallMs() - t0;
if (best < 0 || ms < best) best = ms;
}
return best;
}
static int runMemprobe(Ctx& c, const Options& o) {
Compiled cp;
std::string err;
if (!rtcCompile(c, PROBE_CUDA, "probe.cu", "", "", {}, cp, err)) { std::printf("memprobe: build FAILED: %s\n", err.c_str()); return 2; }
CUmodule mod = nullptr;
if (c.drv.moduleLoadData(&mod, cp.image.data()) != CUDA_SUCCESS) { std::printf("memprobe: cuModuleLoadData failed\n"); return 2; }
CUfunction kFill, kChase, kIndep, kLine16, kLine, kStream, kAlu;
const char* names[7] = { "probe_fill", "probe_chase", "probe_indep", "probe_line16", "probe_line", "probe_stream", "probe_alu" };
CUfunction* fns[7] = { &kFill, &kChase, &kIndep, &kLine16, &kLine, &kStream, &kAlu };
for (int i = 0; i < 7; ++i) if (c.drv.moduleGetFunction(fns[i], mod, names[i]) != CUDA_SUCCESS) { std::printf("memprobe: %s not in the module\n", names[i]); return 2; }
int sizes[3] = { 4, 64, 1024 }, nSizes = 3;
if (o.probeMib > 0) { sizes[0] = o.probeMib; nSizes = 1; }
const size_t lanesList[8] = { 256, 1024, 1u << 12, 1u << 14, 1u << 16, 1u << 18, 1u << 20, 1u << 22 };
const size_t groups[2] = { 32, 256 };
const uint32_t STEPS = 256u, ALU_STEPS = 4096u;
const size_t maxLanes = 1u << 22;
CUdeviceptr dOut = 0;
if (c.drv.memAlloc(&dOut, maxLanes * 4u) != CUDA_SUCCESS) { std::printf("memprobe: cuMemAlloc out\n"); return 2; }
std::printf("memprobe on %s (sm_%d%d, %d SMs, driver %d.%d, NVRTC %d.%d), wall time around cuStreamSynchronize\n", c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor);
std::printf("| probe | MiB | block | lanes in flight | steps per lane | best ms | G loads/s | ns per dependent load |\n|---|---|---|---|---|---|---|---|\n");
for (int si = 0; si < nSizes; ++si) {
int mib = sizes[si];
uint64_t bytes = (uint64_t)mib << 20;
uint32_t words = (uint32_t)(bytes / 4ull), mask = words - 1u, n = words;
CUdeviceptr dDs = 0;
if (c.drv.memAlloc(&dDs, (size_t)bytes) != CUDA_SUCCESS) { std::printf("| chase | %d | skipped: cuMemAlloc failed | | | | | |\n", mib); continue; }
{ void* a[2] = { &dDs, &n }; probeLaunch(c, kFill, ((size_t)words + 255) / 256 * 256, 256, 1, -1, 0, a); }
for (int gi = 0; gi < 2; ++gi) {
size_t local = groups[gi];
for (int li = 0; li < 8; ++li) {
size_t lanes = lanesList[li];
if (lanes < local) continue;
uint32_t seed = 0x1234567u + (uint32_t)li * 977u, steps = STEPS;
void* a[5] = { &dDs, &mask, &steps, &seed, &dOut };
double ms = probeLaunch(c, kChase, lanes, local, 3, 3, seed, a);
std::printf("| chase | %d | %zu | %zu | %u | %.3f | %.3f | %.0f |\n", mib, local, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, ms * 1e6 / STEPS);
std::fflush(stdout);
}
}
for (size_t lanes = 1u << 16; lanes <= maxLanes; lanes <<= 2) {
uint32_t seed = 0x7654321u, steps = STEPS;
void* a[5] = { &dDs, &mask, &steps, &seed, &dOut };
double ms = probeLaunch(c, kIndep, lanes, 256, 3, 3, seed, a);
std::printf("| indep x8 | %d | 256 | %zu | %u | %.3f | %.3f | (8 loads in flight per lane) |\n", mib, lanes, STEPS, ms, (double)lanes * 8.0 * STEPS / (ms / 1000.0) / 1e9);
}
for (size_t lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
uint32_t vecMask = (words / 4u) - 1u, seed = 0x2718281u, steps = STEPS;
void* a[5] = { &dDs, &vecMask, &steps, &seed, &dOut };
double ms = probeLaunch(c, kLine16, lanes, 256, 3, 3, seed, a);
std::printf("| line 16 B | %d | 256 | %zu | %u | %.3f | %.3f G reads/s | %.1f GB/s in 16 B reads |\n", mib, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, (double)lanes * STEPS * 16.0 / (ms / 1000.0) / 1e9);
}
for (size_t lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
uint32_t lineMask = (words / 16u) - 1u, seed = 0x3141592u, steps = STEPS;
void* a[5] = { &dDs, &lineMask, &steps, &seed, &dOut };
double ms = probeLaunch(c, kLine, lanes, 256, 3, 3, seed, a);
std::printf("| line 64 B | %d | 256 | %zu | %u | %.3f | %.3f G lines/s | %.1f GB/s in lines |\n", mib, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, (double)lanes * STEPS * 64.0 / (ms / 1000.0) / 1e9);
}
{
size_t lanes = 1u << 20;
uint32_t perLane = (uint32_t)((uint64_t)words / 4ull / (uint64_t)lanes);
if (perLane == 0) { perLane = 1; lanes = (size_t)words / 4u; }
double bytesRead = (double)perLane * (double)lanes * 16.0;
void* a[3] = { &dDs, &perLane, &dOut };
double ms = probeLaunch(c, kStream, lanes, 256, 3, -1, 0, a);
std::printf("| stream | %d | 256 | %zu | %u | %.3f | %.1f GB/s coalesced | (%.0f MiB read once) |\n", mib, lanes, perLane, ms, bytesRead / (ms / 1000.0) / 1e9, bytesRead / 1048576.0);
}
c.drv.memFree(dDs);
std::fflush(stdout);
}
{
size_t lanes = 1u << 20;
uint32_t seed = 0x2468aceu, steps = ALU_STEPS;
void* a[3] = { &steps, &seed, &dOut };
double ms = probeLaunch(c, kAlu, lanes, 256, 3, 1, seed, a);
double ops = (double)lanes * ALU_STEPS * 5.0;
std::printf("| alu | 0 | 256 | %zu | %u | %.3f | %.1f G int ops/s | %.3f G steps/s per SM (approximate: 5 ops per step counted) |\n", lanes, ALU_STEPS, ms, ops / (ms / 1000.0) / 1e9, (double)lanes * ALU_STEPS / (ms / 1000.0) / 1e9 / (c.sms ? c.sms : 1));
}
c.drv.memFree(dOut);
c.drv.moduleUnload(mod);
std::printf("memprobe: done\n");
return 0;
}
int main(int argc, char** argv) {
Options o = parseArgs(argc, argv);
Ctx c;
c.blockWarps = o.blockWarps;
c.race = o.race; c.raceBenchMs = o.raceBenchMs; c.raceBudgetS = o.raceBudgetS; c.raceRounds = o.raceRounds; c.batchLog2 = o.batchLog2; c.pinned = o.pinned;
c.warps = o.warps; c.batches = o.batches;
if (o.bench || o.memprobe) c.race = "off";
if (!o.tuningPath.empty()) { bool ok = false; c.tuning = readText(o.tuningPath, ok); if (!ok) c.tuning.clear(); }
std::string err, drvLib, rtcLib;
if (!loadDriver(c.drv, err, drvLib)) { emit("error 0 " + err); return 2; }
if (!loadNvrtc(c.rtc, err, rtcLib)) { emit("error 0 " + err); return 2; }
if (!openDevice(c, o.device, o.arch, err)) { emit("error 0 " + err); return 2; }
info(fmt("igneum-worker-cuda %s: device %d %s (sm_%d%d, %d SMs), driver %d.%d from %s, NVRTC %d.%d from %s, target %s (%s)",
WORKER_VERSION, o.device, c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, drvLib.c_str(), c.rtcMajor, c.rtcMinor, rtcLib.c_str(), c.archOpt.c_str(), c.why.c_str()));
if (!c.tuning.empty()) info(fmt("tuning file %s (%zu bytes): %s", o.tuningPath.c_str(), c.tuning.size(), readTuning(c.tuning, c.name).found ? "has an entry for this card" : "no entry for this card"));
if (o.memprobe) { int rc = runMemprobe(c, o); c.drv.primaryCtxRelease(c.dev); return rc; }
double t0 = wallMs();
Pair* cur = buildPair(c, o.pack, nullptr, err, !o.check && !o.bench);
if (!cur) { emit("error 0 " + err); return 1; }
if (o.bench) {
std::printf("pack %s on %s: %s\n", o.pack.c_str(), c.name.c_str(), pairSummary(cur).c_str());
int rc = runBench(c, o, cur);
releasePair(c, cur);
c.drv.primaryCtxRelease(c.dev);
return rc;
}
if (o.raceOnly) {
std::printf("race %s on %s (%s, %d SMs, driver %d.%d, NVRTC %d.%d, %s): %s\n", o.pack.c_str(), c.name.c_str(), c.archOpt.c_str(), c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor, c.why.c_str(), pairSummary(cur).c_str());
std::printf("%s\n", cur->raceLine.c_str());
std::printf("winner %s: %d registers, %d blocks/SM at %d warp(s)/block\n", cur->variant.c_str(), cur->regs, cur->blocksPerSM, cur->blockWarps);
releasePair(c, cur);
return 0;
}
if (o.check) {
std::printf("check PASS %s in %.0f ms: %s\n", o.pack.c_str(), wallMs() - t0, pairSummary(cur).c_str());
std::printf(" epoch %s day %s, dataset 2^%u words, cache 2^%u words in %u segments, %d registers, %d blocks/SM at %d warp(s)/block, target %s\n",
cur->epochHex.c_str(), cur->dayHex.c_str(), cur->datasetLog2, (unsigned)__builtin_ctz(cur->cacheWords), cur->cacheSegments, cur->regs, cur->blocksPerSM, c.blockWarps, c.archOpt.c_str());
releasePair(c, cur);
return 0;
}
int rc = runServe(c, o, cur);
c.drv.primaryCtxRelease(c.dev);
return rc;
}
Reference in a new issue View git blame Copy permalink