igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity
igneum/proto-cuda/nvrtc/worker.cpp

1152 lines
67 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;
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
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;
};
// 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->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);
void* args[5] = { &p->ds, &out, &baseNonce, &mask, &a };
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
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;
if (!rtcCompile(c, kernelDev, "kernel.cu", programH, memhardH, { "igneum_cache_fill", "igneum_build" }, 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; }
c.drv.funcGetAttribute(&p->regs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->fHashBound);
c.drv.occupancy(&p->blocksPerSM, p->fHashBound, 32 * c.blockWarps, 0);
}
// 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 + (64u << 20)) {
err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20));
releasePair(c, p); return nullptr;
}
}
{
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;
// 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);
char line[1024];
p->checkPass = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples.data(), vec.data(), 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 check %.0f race %.0f ms variant %s; %s", p->compileMs, p->cacheMs, p->dsMs, 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::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;
}
t->result = p;
t->error = err;
t->done = true;
}
// ---------------------------------------------------------------------------------------------
// Serve
struct Options {
bool serve = false, check = false, raceOnly = false;
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"
" --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 == "--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) { 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()) { 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;
}
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->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 ((std::memcmp(sw, cur->sw, 32) != 0 || std::memcmp(kw, cur->kw, 32) != 0) && !task &&
!(prepared && std::memcmp(sw, prepared->sw, 32) == 0 && std::memcmp(kw, prepared->kw, 32) == 0)) {
// 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 (std::memcmp(sw, cur->sw, 32) != 0 || std::memcmp(kw, cur->kw, 32) != 0) {
if (prepared && std::memcmp(sw, prepared->sw, 32) == 0 && std::memcmp(kw, prepared->kw, 32) == 0) {
if (old) releasePair(c, old);
old = cur; cur = prepared; prepared = nullptr; switched = true;
info(fmt("switched to the prepared pair epoch %.16s day %s in %.2f ms", cur->epochHex.c_str(), cur->dayHex.c_str(), wallMs() - t0));
} else if (std::memcmp(sw, cur->sw, 32) != 0) {
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 (seed words %08x %08x ...)%s, the job's epoch seed %.16s gives %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
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;
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"));
double t0 = wallMs();
Pair* cur = buildPair(c, o.pack, nullptr, err, !o.check);
if (!cur) { emit("error 0 " + err); return 1; }
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