igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity
igneum/proto-cuda/nvrtc/worker.cpp
igneum-labs 4d907d8a14 Workers answer a job they have no pair for with a need line; the miner prepares the current pair (PC 2 stuck on the previous epoch) 
Root cause from the uploads (bench-log entry): the app exported the pack while its node was in IBD inside the previous
epoch, the OpenCL worker started after the boundary with no next epoch within lead, so no prepare was ever sent and
every job was a seed mismatch; the CUDA worker on the same PC had swapped correctly. Both workers now print
need <epoch> <day> before the error; the devnet-v4 miner (3bfe346f) prepares the current pair on a need line or three
mismatches, exits 42 for a worker without prepare support, and restarts a ready worker that completes no job for
60 s with jobs queued. Package rebuilt with the guarded miner (ship build on dc749905), payload inputs published.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
2026-10-04 13:32:58 +00:00

796 lines
45 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.
//
// Usage: igneum-worker-cuda --serve --pack <dir> [--device D] [--batch-log2 22] [--block-warps 1] [--arch sm_120|auto]
// igneum-worker-cuda --check --pack <dir> [--device D] compile, build, self-test, print timings, exit 0/1
#include <cstdint>
#include <cstdarg>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <chrono>
#include <string>
#include <vector>
#include <thread>
#include <atomic>
#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;
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) {
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.
const char* opts[3] = { archOpt.c_str(), "--std=c++17", "-default-device" };
r = c.rtc.compileProgram(prog, 3, opts);
{
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;
};
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;
}
// Compiles the pack in `dir`, builds its cache and dataset on stream `s`, runs the self-test. 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) {
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; }
return p;
}
static std::string pairSummary(const Pair* p) {
return fmt("nvrtc %.0f cache %.0f dataset %.0f check %.0f ms; %s", p->compileMs, p->cacheMs, p->dsMs, p->checkMs, 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);
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;
int device = 0, batchLog2 = 22, blockWarps = 1;
std::string pack, arch = "auto";
};
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", 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 == "--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) { usage(); std::exit(2); }
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 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(), WORKER_VERSION));
info(fmt("first pack %s: %s", cur->dir.c_str(), pairSummary(cur).c_str()));
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;
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);
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())); }
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
uint32_t block = 32u * (uint32_t)o.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;
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()));
double t0 = wallMs();
Pair* cur = buildPair(c, o.pack, nullptr, err);
if (!cur) { emit("error 0 " + err); return 1; }
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