igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity
igneum/proto-cuda/nvrtc/worker.cpp
igneum-labs 734dbd134f NVRTC worker: -default-device (NVRTC rejects unannotated pack helpers as host code; found on the first real card) 
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
2026-10-04 10:12:35 +00:00

793 lines
44 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>
// 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("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("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