igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity
igneum/proto-cuda/nvrtc/worker.cpp
igneum-labs 27bcb81494 Variant racing: the NVRTC and Metal workers compile several kernel variants at every prepare and keep the fastest for the hour 
the project lead, 4 Oct 2026 evening: "we need to make our miner better than anything else can be". The compile-ahead pipeline
built one kernel per program; it now builds a catalogue (unroll 2 or 8; the dataset load path __ldg, __ldcg, __ldcs
on NVIDIA; a register budget by --maxrregcount or __launch_bounds__, max_total_threads_per_threadgroup on Apple;
2, 4, 8 warps per block or 64 to 256 threads per threadgroup; combinations), self-tests each against the pack's
vector warps (NVIDIA) or the base kernel over 2^16 nonces (Metal), bit for bit or out, and times each for about two
seconds with the job loop paused (one mutex, mining resumes between variants). Base is the pack's text as shipped,
always first, never discarded; a race has a budget (default 120 s against the 600-DAA lead) and keeps the best so
far when it runs out, so the swap is never delayed. One line per race: variants, MH/s each, winner, gain, time.
A tuning file (--tuning, IGNEUM_TUNING_FILE) pins a variant or orders the candidates per card model.

NVIDIA: textual rewrites on the pack's own kernel_bound.cu with exact anchors from igneum-pow's emitter, so the
pack format, the miner and igneum-pow are untouched and old packs race. --race --pack <dir> runs the race alone.
Under IGNEUM_EMU the race is off (the stand-in checks the handed-over text is the pack's). Metal: the MSL hooks in
generateMSL, a lock-protected kernel slot per program, a pair compiled inline races after its first job,
--race-test --seed --day runs the race alone with a table. OpenCL is not raced yet. docs/design/miner-tuning.md.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
2026-10-04 19:08:45 +00:00

1149 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;
}
}
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