From f535abb59d417596b55df291c1e72ca49fd4a4a9 Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Sun, 4 Oct 2026 09:55:05 +0000 Subject: [PATCH] One-click CUDA worker: driver API + NVRTC loaded at run time, pack read at run time, Mac emulation check proto-cuda/nvrtc/worker.cpp serves igneum-miner's worker protocol with nothing installed but the NVIDIA driver: nvcuda.dll and the redistributable nvrtc64_120_0.dll are opened with LoadLibrary (cuda_api.h), the pack's kernel.cu and kernel_bound.cu are handed to NVRTC byte for byte up to the host launch wrappers with program.h and memhard.h as named headers, every pack is self-tested against its vectors.h (cache head, last line, FNV-1a 64, dataset head, last word, 64 samples, 96 vector lanes) before it serves a job, prepare runs on a thread for the hourly swap, and a job on seeds without a pair makes the worker find the miner's pack by seeds.txt and build it. packfile.h (C99) reads a pack directory. fetch-redist.sh verifies and stages the NVIDIA 12.8.93 redistributables and the Khronos headers (THIRD-PARTY.md records URLs, hashes and the EULA clause); build-windows.sh cross-compiles with mingw (static, KERNEL32 + Universal CRT only). emu/: the two libraries as host functions, the pack kernels on host threads, the NVRTC source compared with the pack files; test.sh PASS on the Mac (two real packs, swap, self-heal, 17 sampled hashes equal igneum-pow hash-bound). Not run here: the real NVRTC compile and driver load (RTX 5090 PC). Co-Authored-By: Claude Fable 5.1 --- .gitignore | 8 + proto-cuda/nvrtc/THIRD-PARTY.md | 62 +++ proto-cuda/nvrtc/build-windows.sh | 31 ++ proto-cuda/nvrtc/cuda_api.h | 61 +++ proto-cuda/nvrtc/emu/emu_backend.cpp | 213 ++++++++ proto-cuda/nvrtc/emu/serve-check.sh | 62 +++ proto-cuda/nvrtc/emu/test.sh | 61 +++ proto-cuda/nvrtc/fetch-redist.sh | 52 ++ proto-cuda/nvrtc/packfile.h | 350 ++++++++++++ proto-cuda/nvrtc/worker.cpp | 791 +++++++++++++++++++++++++++ 10 files changed, 1691 insertions(+) create mode 100644 proto-cuda/nvrtc/THIRD-PARTY.md create mode 100755 proto-cuda/nvrtc/build-windows.sh create mode 100644 proto-cuda/nvrtc/cuda_api.h create mode 100644 proto-cuda/nvrtc/emu/emu_backend.cpp create mode 100755 proto-cuda/nvrtc/emu/serve-check.sh create mode 100755 proto-cuda/nvrtc/emu/test.sh create mode 100755 proto-cuda/nvrtc/fetch-redist.sh create mode 100644 proto-cuda/nvrtc/packfile.h create mode 100644 proto-cuda/nvrtc/worker.cpp diff --git a/.gitignore b/.gitignore index 3587bb2c..f4d63d7d 100644 --- a/.gitignore +++ b/.gitignore @@ -16,3 +16,11 @@ sim/__pycache__/ *.pyc brand/Unbounded.ttf .vercel +# one-click worker: fetched third-party pieces and the Mac emulation build +proto-cuda/nvrtc/redist/ +proto-cuda/nvrtc/emu/build/ +proto-cuda/nvrtc/*.exe +proto-opencl/*.exe +proto-opencl/igneum-bench-cl-generic-test +proto-opencl/generic-test*.log +proto-cuda/nvrtc/emu/test-run.log diff --git a/proto-cuda/nvrtc/THIRD-PARTY.md b/proto-cuda/nvrtc/THIRD-PARTY.md new file mode 100644 index 00000000..55bb8a1c --- /dev/null +++ b/proto-cuda/nvrtc/THIRD-PARTY.md @@ -0,0 +1,62 @@ +# Third-party pieces of the one-click Windows workers + +Recorded 4 October 2026. Nothing listed here is committed to this repository; `fetch-redist.sh` downloads every +file into `~/.cache/igneum/`, checks its SHA-256 against the values below and stages it under `proto-cuda/nvrtc/redist/` +(gitignored). `build-windows.sh` compiles against the headers; `windows-app/make-package.sh` ships the two NVRTC DLLs +and the two licence texts in the zip. Nothing else third-party is in the package. + +## Shipped in the package + +| File in the zip | Source | Version | Size (bytes) | SHA-256 | +|---|---|---|---|---| +| `nvrtc64_120_0.dll` | NVIDIA CUDA redistributable, package `cuda_nvrtc` (NVIDIA Runtime Compilation Library) | 12.8.93 (CUDA 12.8 Update 1, released 2025-03-06) | 86,728,192 | `02d2d7ef4690bf1a` (first 16 hex; the archive's hash is below) | +| `nvrtc-builtins64_128.dll` | same package | 12.8.93 | 6,356,480 | `a7afa2a40cbd3f92` (first 16 hex) | +| `LICENSE-NVIDIA-CUDA-EULA.txt` | the `LICENSE` file inside the same archive (NVIDIA Software License Agreement and CUDA Supplement) | | 63,021 | `e2c71babfd18a8e6` (first 16 hex) | +| `LICENSE-Khronos-OpenCL-Headers.txt` | Khronos OpenCL-Headers `LICENSE` (Apache License 2.0); the headers themselves are build-only | v2026.05.29 | 11,358 | `cfc7749b96f63bd31c3c42b5c471bf756814053e847c10f3eb003417bc523d30` | + +Archive the two DLLs come from (verified by `fetch-redist.sh`): + +| Archive | URL | SHA-256 (from NVIDIA's `redistrib_12.8.1.json`) | +|---|---|---| +| `cuda_nvrtc-windows-x86_64-12.8.93-archive.zip` (305,588,898 bytes) | `https://developer.download.nvidia.com/compute/cuda/redist/cuda_nvrtc/windows-x86_64/cuda_nvrtc-windows-x86_64-12.8.93-archive.zip` | `a63302a077f0248a743a1a7caa7dbd80d0fac56c6cfa9c41fa05fac9b7e5eda5` | + +The archive also holds `nvrtc64_120_0.alt.dll` and static libraries; they are not shipped and not used. + +## Why the two DLLs may be redistributed + +The NVIDIA Software License Agreement (the `LICENSE` file above, section 2.2 "Distribution": "The portions of the SDK that +are distributable under the Agreement are listed in Attachment A") lists in Attachment A, under "NVIDIA Runtime +Compilation Library and Header": + +> All: nvrtc.h +> Windows: nvrtc.dll, nvrtc-builtins.dll + +and the preamble of Attachment A covers the versioned file names ("including certain variations of these files that have +version number or architecture specific information embedded in the file name"), which is what `nvrtc64_120_0.dll` and +`nvrtc-builtins64_128.dll` are. Section 2.1 limits the licence to applications "for use in systems with NVIDIA GPUs", +which is the case: the worker runs on NVIDIA GPUs only. The DLLs are shipped unmodified next to the exe. The licence text +is in the package so the user can read it. + +## Build-only (compiled against, never shipped) + +| File | Source | Version | SHA-256 | +|---|---|---|---| +| `cuda.h` (the driver API header) | NVIDIA CUDA redistributable, package `cuda_cudart`, archive `cuda_cudart-windows-x86_64-12.8.90-archive.zip` (3,037,735 bytes) at `https://developer.download.nvidia.com/compute/cuda/redist/cuda_cudart/windows-x86_64/cuda_cudart-windows-x86_64-12.8.90-archive.zip`, archive SHA-256 `4a39058fd8519444a81cfc7ae055d136f48d1a31ffa41ae255b35b2edd61e13b` | 12.8.90 | `167e383a3791226c` (first 16 hex) | +| `nvrtc.h` | the `cuda_nvrtc` archive above (Attachment A lists it as redistributable too) | 12.8.93 | `85d79ff186392c1b` (first 16 hex) | +| `CL/cl.h` | `https://raw.githubusercontent.com/KhronosGroup/OpenCL-Headers/v2026.05.29/CL/cl.h` | v2026.05.29 | `16d09614cd7eef73b4094089cc4ce2af777181d3ffe9613473fb6d552234a4d1` | +| `CL/cl_platform.h` | same tag | | `b125cece6fe41f2e6690d88b24652c151cb6654da13b42180d69d5bd9ed526d7` | +| `CL/cl_version.h` | same tag | | `1e571dd68566330dd066cf778e07909a3cabeb2b969c5578c60f14baba7df1af` | +| `CL/cl_ext.h` | same tag (fetched for completeness; `host.c` includes `cl.h` only) | | `fe8916fd98b739cbc47ecab387bee2d258fa939c1bd2e10415c6f8f8b7055968` | + +`cuda.h` is not in Attachment A, so it is used only to compile the worker on the Mac and is not in the package. The +worker never links a CUDA import library: `nvcuda.dll` (installed by every NVIDIA driver) and `nvrtc64_120_0.dll` are +opened with `LoadLibrary` and their entry points fetched with `GetProcAddress` (`cuda_api.h`, `worker.cpp`). The +OpenCL worker does the same with `OpenCL.dll`, the Khronos ICD loader that the AMD, NVIDIA and Intel drivers install +(`proto-opencl/cl_dynamic.h`). + +## What the user's PC needs + +An NVIDIA driver (for the CUDA worker; the driver API is in `nvcuda.dll`, and the driver must be new enough for +CUDA 12.x) or an AMD, NVIDIA or Intel driver with OpenCL (for the OpenCL worker). Windows 10 or 11 (the exes import +only `KERNEL32.dll` and the Universal CRT `api-ms-win-crt-*` libraries, which both carry). No CUDA Toolkit, no Visual +Studio, no SDK. diff --git a/proto-cuda/nvrtc/build-windows.sh b/proto-cuda/nvrtc/build-windows.sh new file mode 100755 index 00000000..0cad5564 --- /dev/null +++ b/proto-cuda/nvrtc/build-windows.sh @@ -0,0 +1,31 @@ +#!/usr/bin/env bash +# Cross-compiles the two one-click Windows workers on the Mac with Homebrew mingw-w64 (x86_64-w64-mingw32-gcc/g++): +# proto-cuda/nvrtc/igneum-worker-cuda.exe driver API + NVRTC, both loaded at run time (cuda_api.h); needs cuda.h and +# nvrtc.h from fetch-redist.sh to compile, no import library +# proto-opencl/igneum-worker-opencl.exe host.c with IGNEUM_CL_DYNAMIC (OpenCL.dll loaded at run time) and the generic +# --pack serve mode; the pack it is built against is only a placeholder +# Static libgcc/libstdc++/winpthread, so each exe needs only KERNEL32 and the Universal CRT (Windows 10 and 11 have it). +# Usage: build-windows.sh (then windows-app/make-package.sh ships them with the NVRTC DLLs) +set -euo pipefail +HERE="$(cd "$(dirname "$0")" && pwd)" +ROOT="$(cd "$HERE/../.." && pwd)" +RED="$HERE/redist" +[ -f "$RED/include/cuda.h" ] && [ -f "$RED/include/nvrtc.h" ] && [ -f "$RED/include/CL/cl.h" ] || { echo "run $HERE/fetch-redist.sh first" >&2; exit 1; } +CXX=x86_64-w64-mingw32-g++ +CC=x86_64-w64-mingw32-gcc +STRIP=x86_64-w64-mingw32-strip +PLACEHOLDER="$ROOT/proto-cuda/packs/igneum-devnet-v4-epoch0" + +echo "== igneum-worker-cuda.exe" +"$CXX" -std=c++17 -O2 -Wall -Wextra -static -I "$HERE" -I "$RED/include" -o "$HERE/igneum-worker-cuda.exe" "$HERE/worker.cpp" +"$STRIP" "$HERE/igneum-worker-cuda.exe" +echo "== igneum-worker-opencl.exe" +"$CC" -std=c99 -O2 -Wall -Wextra -Wno-stringop-truncation -Wno-format-truncation -static -DIGNEUM_CL_DYNAMIC -DCL_TARGET_OPENCL_VERSION=120 \ + -I "$RED/include" -I "$PLACEHOLDER" -DIGNEUM_KERNEL_PATH='"kernel_bound.cl"' \ + -o "$ROOT/proto-opencl/igneum-worker-opencl.exe" "$ROOT/proto-opencl/host.c" +"$STRIP" "$ROOT/proto-opencl/igneum-worker-opencl.exe" +for exe in "$HERE/igneum-worker-cuda.exe" "$ROOT/proto-opencl/igneum-worker-opencl.exe"; do + printf '%s: %d bytes, imports:' "$(basename "$exe")" "$(stat -f %z "$exe")" + x86_64-w64-mingw32-objdump -p "$exe" | sed -n 's/^[[:space:]]*DLL Name: //p' | tr '\n' ' ' + echo +done diff --git a/proto-cuda/nvrtc/cuda_api.h b/proto-cuda/nvrtc/cuda_api.h new file mode 100644 index 00000000..e16e8c6e --- /dev/null +++ b/proto-cuda/nvrtc/cuda_api.h @@ -0,0 +1,61 @@ +// cuda_api.h: the slice of the CUDA driver API and of NVRTC that igneum-worker-cuda calls, as function pointers. +// worker.cpp fills them from nvcuda.dll (the driver) and nvrtc64_*.dll (shipped next to the exe) with +// LoadLibrary/GetProcAddress, so no import library is linked and the exe cross-compiles with mingw on the Mac. +// The emulation build (emu/emu_backend.cpp) fills the same table with host functions. 4 October 2026. +#pragma once +#include +#include + +struct Drv { + decltype(&cuInit) init = nullptr; + decltype(&cuDriverGetVersion) driverGetVersion = nullptr; + decltype(&cuDeviceGetCount) deviceGetCount = nullptr; + decltype(&cuDeviceGet) deviceGet = nullptr; + decltype(&cuDeviceGetName) deviceGetName = nullptr; + decltype(&cuDeviceGetAttribute) deviceGetAttribute = nullptr; + decltype(&cuDeviceTotalMem_v2) deviceTotalMem = nullptr; + decltype(&cuDevicePrimaryCtxSetFlags_v2) primaryCtxSetFlags = nullptr; + decltype(&cuDevicePrimaryCtxRetain) primaryCtxRetain = nullptr; + decltype(&cuDevicePrimaryCtxRelease_v2) primaryCtxRelease = nullptr; + decltype(&cuCtxSetCurrent) ctxSetCurrent = nullptr; + decltype(&cuCtxSynchronize) ctxSynchronize = nullptr; + decltype(&cuMemGetInfo_v2) memGetInfo = nullptr; + decltype(&cuMemAlloc_v2) memAlloc = nullptr; + decltype(&cuMemFree_v2) memFree = nullptr; + decltype(&cuMemcpyDtoH_v2) memcpyDtoH = nullptr; + decltype(&cuModuleLoadData) moduleLoadData = nullptr; + decltype(&cuModuleUnload) moduleUnload = nullptr; + decltype(&cuModuleGetFunction) moduleGetFunction = nullptr; + decltype(&cuLaunchKernel) launchKernel = nullptr; + decltype(&cuStreamCreate) streamCreate = nullptr; + decltype(&cuStreamSynchronize) streamSynchronize = nullptr; + decltype(&cuStreamDestroy_v2) streamDestroy = nullptr; + decltype(&cuFuncGetAttribute) funcGetAttribute = nullptr; + decltype(&cuOccupancyMaxActiveBlocksPerMultiprocessor) occupancy = nullptr; + decltype(&cuGetErrorString) getErrorString = nullptr; + decltype(&cuGetErrorName) getErrorName = nullptr; +}; + +struct Rtc { + decltype(&nvrtcVersion) version = nullptr; + decltype(&nvrtcGetNumSupportedArchs) getNumSupportedArchs = nullptr; // CUDA 11.2 and newer; optional + decltype(&nvrtcGetSupportedArchs) getSupportedArchs = nullptr; + decltype(&nvrtcCreateProgram) createProgram = nullptr; + decltype(&nvrtcDestroyProgram) destroyProgram = nullptr; + decltype(&nvrtcCompileProgram) compileProgram = nullptr; + decltype(&nvrtcGetProgramLogSize) getProgramLogSize = nullptr; + decltype(&nvrtcGetProgramLog) getProgramLog = nullptr; + decltype(&nvrtcGetPTXSize) getPTXSize = nullptr; + decltype(&nvrtcGetPTX) getPTX = nullptr; + decltype(&nvrtcGetCUBINSize) getCUBINSize = nullptr; + decltype(&nvrtcGetCUBIN) getCUBIN = nullptr; + decltype(&nvrtcAddNameExpression) addNameExpression = nullptr; + decltype(&nvrtcGetLoweredName) getLoweredName = nullptr; + decltype(&nvrtcGetErrorString) getErrorString = nullptr; +}; + +#ifdef IGNEUM_EMU +// emu/emu_backend.cpp: host functions in place of the two libraries (the pack's kernels run on host threads) +void emu_fill_driver(Drv& d); +void emu_fill_nvrtc(Rtc& r); +#endif diff --git a/proto-cuda/nvrtc/emu/emu_backend.cpp b/proto-cuda/nvrtc/emu/emu_backend.cpp new file mode 100644 index 00000000..995998de --- /dev/null +++ b/proto-cuda/nvrtc/emu/emu_backend.cpp @@ -0,0 +1,213 @@ +// emu_backend.cpp: the CUDA driver API and NVRTC replaced by host functions, for checking igneum-worker-cuda on a +// machine without an NVIDIA GPU (the Mac). Not part of the deliverable. 4 October 2026. +// +// What it proves: the worker's protocol, job walk, init words, pack reading, self-test, prepare thread and swap, +// with the pack's real kernel text running on host threads through proto-cuda/emu (the same shim the host.cu +// emulation uses, 32 threads per warp, a barrier inside __shfl_xor_sync). What it cannot prove: that NVRTC accepts +// the text and that the driver runs it; that is the RTX 5090 run (windows-app/TEST.md). +// +// The NVRTC stand-in records every program the worker creates and checks it against the pack on disk: the source +// must be the pack's kernel file up to the host launch wrappers, byte for byte, the dropped tail must hold host +// wrappers only, and program.h and memhard.h must be the pack's files byte for byte. Two packs may be compiled in +// (IGNEUM_EMU_PACK and IGNEUM_EMU_PACK2, the directories), so a prepare and the swap run with real second-pack code. +#include +#include +#include +#include +#include +#include +#include +#include + +#include "../cuda_api.h" // the real cuda.h / nvrtc.h types +#include "cuda_runtime.h" // the shim (proto-cuda/emu): emu_launch and the kernel symbols' world + +namespace emu_pack_a { +struct IgneumInitWords { uint32_t w[8]; }; // the same definition as the pack's kernel_bound.cu, inside its namespace +void igneum_cache_fill(uint32_t* cache, uint32_t nSegments); +void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems); +void igneum_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw); +} +#ifdef IGNEUM_EMU_TWO_PACKS +namespace emu_pack_b { +struct IgneumInitWords { uint32_t w[8]; }; +void igneum_cache_fill(uint32_t* cache, uint32_t nSegments); +void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems); +void igneum_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw); +} +#endif + +static std::string readAll(const std::string& path) { + FILE* f = std::fopen(path.c_str(), "rb"); + if (!f) return std::string(); + std::string s; + char buf[65536]; + size_t n; + while ((n = std::fread(buf, 1, sizeof(buf), f)) > 0) s.append(buf, n); + std::fclose(f); + return s; +} + +// --------------------------------------------------------------------------------------------- +// NVRTC stand-in + +struct EmuProg { + std::string src, name, log; + std::map headers; + std::list exprs; + int pack = 0; // 1 = IGNEUM_EMU_PACK, 2 = IGNEUM_EMU_PACK2, 0 = unknown + bool checked = false, ok = false, compiled = false; +}; + +static const char* packDir(int which) { + const char* e = std::getenv(which == 1 ? "IGNEUM_EMU_PACK" : "IGNEUM_EMU_PACK2"); + return e ? e : ""; +} + +// The source check. Fills p->pack, p->ok, p->log. +static void checkSource(EmuProg* p) { + p->checked = true; + for (int which = 1; which <= 2; ++which) { + std::string dir = packDir(which); + if (dir.empty()) continue; + std::string file = readAll(dir + "/" + p->name); + if (file.empty() || file.size() < p->src.size() || file.compare(0, p->src.size(), p->src) != 0) continue; + std::string tail = file.substr(p->src.size()); + bool tailOk = tail.empty() || tail.rfind("// Host-side launch wrappers", 0) == 0 || tail.rfind("cudaError_t ", 0) == 0; + tailOk = tailOk && tail.find("__global__") == std::string::npos && tail.find("__device__") == std::string::npos; + if (!tailOk) { p->log = "emu-nvrtc: the dropped tail of " + p->name + " is not the host launch wrappers"; p->ok = false; p->pack = which; return; } + std::string ph = readAll(dir + "/program.h"), mh = readAll(dir + "/memhard.h"); + if (!p->headers.count("program.h") || p->headers["program.h"] != ph) { p->log = "emu-nvrtc: program.h handed to NVRTC differs from the pack's"; p->ok = false; p->pack = which; return; } + if (!p->headers.count("memhard.h") || p->headers["memhard.h"] != mh) { p->log = "emu-nvrtc: memhard.h handed to NVRTC differs from the pack's"; p->ok = false; p->pack = which; return; } + if (!p->headers.count("cuda_runtime.h") || !p->headers.count("cstdint")) { p->log = "emu-nvrtc: the stub headers cuda_runtime.h and cstdint were not handed over"; p->ok = false; p->pack = which; return; } + p->pack = which; p->ok = true; + std::printf("info emu-nvrtc: %s source check PASS for pack %d: %zu bytes handed over = the pack file's first %zu of %zu bytes; dropped tail = host launch wrappers only (%zu bytes); program.h (%zu bytes) and memhard.h (%zu bytes) byte-identical\n", + p->name.c_str(), which, p->src.size(), p->src.size(), file.size(), tail.size(), ph.size(), mh.size()); + std::fflush(stdout); + return; + } + p->log = "emu-nvrtc: the source of " + p->name + " is not a prefix of any pack's file (IGNEUM_EMU_PACK / IGNEUM_EMU_PACK2)"; + p->ok = false; +} + +static nvrtcResult e_version(int* major, int* minor) { *major = 12; *minor = 8; return NVRTC_SUCCESS; } +static int ARCHS[] = { 50, 52, 53, 60, 61, 62, 70, 72, 75, 80, 86, 87, 89, 90, 100, 101, 120 }; +static nvrtcResult e_numArchs(int* n) { *n = (int)(sizeof(ARCHS) / sizeof(ARCHS[0])); return NVRTC_SUCCESS; } +static nvrtcResult e_archs(int* out) { for (size_t i = 0; i < sizeof(ARCHS) / sizeof(ARCHS[0]); ++i) out[i] = ARCHS[i]; return NVRTC_SUCCESS; } +static nvrtcResult e_create(nvrtcProgram* prog, const char* src, const char* name, int numHeaders, const char* const* headers, const char* const* includeNames) { + EmuProg* p = new EmuProg(); + p->src = src ? src : ""; p->name = name ? name : ""; + for (int i = 0; i < numHeaders; ++i) p->headers[includeNames[i]] = headers[i]; + *prog = (nvrtcProgram)p; + return NVRTC_SUCCESS; +} +static nvrtcResult e_destroy(nvrtcProgram* prog) { delete (EmuProg*)*prog; *prog = nullptr; return NVRTC_SUCCESS; } +static nvrtcResult e_compile(nvrtcProgram prog, int numOptions, const char* const* options) { + EmuProg* p = (EmuProg*)prog; + bool arch = false, std17 = false; + for (int i = 0; i < numOptions; ++i) { std::string o = options[i]; if (o.rfind("--gpu-architecture=", 0) == 0) arch = true; if (o == "--std=c++17") std17 = true; } + if (!arch || !std17) { p->log = "emu-nvrtc: expected --gpu-architecture=... and --std=c++17"; return NVRTC_ERROR_COMPILATION; } + checkSource(p); + if (!p->ok) return NVRTC_ERROR_COMPILATION; + p->compiled = true; + return NVRTC_SUCCESS; +} +static nvrtcResult e_logSize(nvrtcProgram prog, size_t* n) { *n = ((EmuProg*)prog)->log.size() + 1; return NVRTC_SUCCESS; } +static nvrtcResult e_log(nvrtcProgram prog, char* out) { std::strcpy(out, ((EmuProg*)prog)->log.c_str()); return NVRTC_SUCCESS; } +static std::string image(EmuProg* p) { return std::string("EMU-IMAGE:") + (p->pack == 2 ? "B" : "A"); } +static nvrtcResult e_imgSize(nvrtcProgram prog, size_t* n) { *n = image((EmuProg*)prog).size() + 1; return NVRTC_SUCCESS; } +static nvrtcResult e_img(nvrtcProgram prog, char* out) { std::strcpy(out, image((EmuProg*)prog).c_str()); return NVRTC_SUCCESS; } +static nvrtcResult e_addName(nvrtcProgram prog, const char* e) { ((EmuProg*)prog)->exprs.push_back(e); return NVRTC_SUCCESS; } +static nvrtcResult e_lowered(nvrtcProgram prog, const char* e, const char** out) { + EmuProg* p = (EmuProg*)prog; + for (const std::string& s : p->exprs) if (s == e) { *out = s.c_str(); return NVRTC_SUCCESS; } + return NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID; +} +static const char* e_errstr(nvrtcResult r) { return r == NVRTC_SUCCESS ? "NVRTC_SUCCESS" : r == NVRTC_ERROR_COMPILATION ? "NVRTC_ERROR_COMPILATION" : "NVRTC_ERROR (emulated)"; } + +void emu_fill_nvrtc(Rtc& r) { + r.version = e_version; r.getNumSupportedArchs = e_numArchs; r.getSupportedArchs = e_archs; + r.createProgram = e_create; r.destroyProgram = e_destroy; r.compileProgram = e_compile; + r.getProgramLogSize = e_logSize; r.getProgramLog = e_log; + r.getPTXSize = e_imgSize; r.getPTX = e_img; r.getCUBINSize = e_imgSize; r.getCUBIN = e_img; + r.addNameExpression = e_addName; r.getLoweredName = e_lowered; r.getErrorString = e_errstr; +} + +// --------------------------------------------------------------------------------------------- +// Driver API stand-in + +static CUresult d_init(unsigned) { return CUDA_SUCCESS; } +static CUresult d_driverVersion(int* v) { *v = 12080; return CUDA_SUCCESS; } +static CUresult d_count(int* c) { *c = 1; return CUDA_SUCCESS; } +static CUresult d_get(CUdevice* d, int i) { *d = i; return CUDA_SUCCESS; } +static CUresult d_name(char* out, int len, CUdevice) { std::snprintf(out, (size_t)len, "CPU emulation shim (not a GPU)"); return CUDA_SUCCESS; } +static CUresult d_attr(int* v, CUdevice_attribute a, CUdevice) { + switch (a) { + case CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR: *v = 12; break; + case CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR: *v = 0; break; + case CU_DEVICE_ATTRIBUTE_WARP_SIZE: *v = 32; break; + default: *v = 0; break; + } + return CUDA_SUCCESS; +} +static CUresult d_totalMem(size_t* b, CUdevice) { *b = 16ull << 30; return CUDA_SUCCESS; } +static CUresult d_ctxFlags(CUdevice, unsigned) { return CUDA_SUCCESS; } +static int gCtx = 0; +static CUresult d_ctxRetain(CUcontext* c, CUdevice) { *c = (CUcontext)&gCtx; return CUDA_SUCCESS; } +static CUresult d_ctxRelease(CUdevice) { return CUDA_SUCCESS; } +static CUresult d_ctxSet(CUcontext) { return CUDA_SUCCESS; } +static CUresult d_ctxSync() { return CUDA_SUCCESS; } +static CUresult d_memInfo(size_t* f, size_t* t) { *f = 8ull << 30; *t = 16ull << 30; return CUDA_SUCCESS; } +static CUresult d_alloc(CUdeviceptr* p, size_t n) { void* m = std::malloc(n); if (!m) return CUDA_ERROR_OUT_OF_MEMORY; *p = (CUdeviceptr)(uintptr_t)m; return CUDA_SUCCESS; } +static CUresult d_free(CUdeviceptr p) { std::free((void*)(uintptr_t)p); return CUDA_SUCCESS; } +static CUresult d_dtoh(void* dst, CUdeviceptr src, size_t n) { std::memcpy(dst, (const void*)(uintptr_t)src, n); return CUDA_SUCCESS; } +static CUresult d_modLoad(CUmodule* m, const void* img) { + const char* s = (const char*)img; + if (std::strncmp(s, "EMU-IMAGE:", 10) != 0) return CUDA_ERROR_INVALID_IMAGE; + *m = (CUmodule)(uintptr_t)(s[10] == 'B' ? 2 : 1); + return CUDA_SUCCESS; +} +static CUresult d_modUnload(CUmodule) { return CUDA_SUCCESS; } +static CUresult d_getFn(CUfunction* f, CUmodule m, const char* name) { + int pack = (int)(uintptr_t)m, idx; + if (std::strcmp(name, "igneum_cache_fill") == 0) idx = 1; + else if (std::strcmp(name, "igneum_build") == 0) idx = 2; + else if (std::strcmp(name, "igneum_hash_bound") == 0) idx = 3; + else return CUDA_ERROR_NOT_FOUND; +#ifndef IGNEUM_EMU_TWO_PACKS + if (pack == 2) return CUDA_ERROR_NOT_FOUND; +#endif + *f = (CUfunction)(uintptr_t)((pack - 1) * 3 + idx); + return CUDA_SUCCESS; +} +template static T arg(void** params, int i) { return *(T*)params[i]; } +template static T* dptr(void** params, int i) { return (T*)(uintptr_t)(*(CUdeviceptr*)params[i]); } +static CUresult d_launch(CUfunction f, unsigned gx, unsigned, unsigned, unsigned bx, unsigned, unsigned, unsigned, CUstream, void** params, void**) { + switch ((int)(uintptr_t)f) { + case 1: emu_launch(emu_pack_a::igneum_cache_fill, gx, bx, dptr(params, 0), arg(params, 1)); return CUDA_SUCCESS; + case 2: emu_launch(emu_pack_a::igneum_build, gx, bx, dptr(params, 0), dptr(params, 1), arg(params, 2)); return CUDA_SUCCESS; + case 3: emu_launch(emu_pack_a::igneum_hash_bound, gx, bx, dptr(params, 0), dptr(params, 1), arg(params, 2), arg(params, 3), arg(params, 4)); return CUDA_SUCCESS; +#ifdef IGNEUM_EMU_TWO_PACKS + case 4: emu_launch(emu_pack_b::igneum_cache_fill, gx, bx, dptr(params, 0), arg(params, 1)); return CUDA_SUCCESS; + case 5: emu_launch(emu_pack_b::igneum_build, gx, bx, dptr(params, 0), dptr(params, 1), arg(params, 2)); return CUDA_SUCCESS; + case 6: emu_launch(emu_pack_b::igneum_hash_bound, gx, bx, dptr(params, 0), dptr(params, 1), arg(params, 2), arg(params, 3), arg(params, 4)); return CUDA_SUCCESS; +#endif + default: return CUDA_ERROR_INVALID_HANDLE; + } +} +static CUresult d_streamCreate(CUstream* s, unsigned) { *s = (CUstream)&gCtx; return CUDA_SUCCESS; } +static CUresult d_streamSync(CUstream) { return CUDA_SUCCESS; } +static CUresult d_streamDestroy(CUstream) { return CUDA_SUCCESS; } +static CUresult d_funcAttr(int* v, CUfunction_attribute, CUfunction) { *v = 0; return CUDA_SUCCESS; } +static CUresult d_occ(int* n, CUfunction, int, size_t) { *n = 0; return CUDA_SUCCESS; } +static CUresult d_errStr(CUresult r, const char** s) { *s = r == CUDA_SUCCESS ? "no error" : "emulated driver error"; return CUDA_SUCCESS; } +static CUresult d_errName(CUresult r, const char** s) { *s = r == CUDA_SUCCESS ? "CUDA_SUCCESS" : "CUDA_ERROR (emulated)"; return CUDA_SUCCESS; } + +void emu_fill_driver(Drv& d) { + d.init = d_init; d.driverGetVersion = d_driverVersion; d.deviceGetCount = d_count; d.deviceGet = d_get; d.deviceGetName = d_name; + d.deviceGetAttribute = d_attr; d.deviceTotalMem = d_totalMem; d.primaryCtxSetFlags = d_ctxFlags; d.primaryCtxRetain = d_ctxRetain; + d.primaryCtxRelease = d_ctxRelease; d.ctxSetCurrent = d_ctxSet; d.ctxSynchronize = d_ctxSync; d.memGetInfo = d_memInfo; + d.memAlloc = d_alloc; d.memFree = d_free; d.memcpyDtoH = d_dtoh; d.moduleLoadData = d_modLoad; d.moduleUnload = d_modUnload; + d.moduleGetFunction = d_getFn; d.launchKernel = d_launch; d.streamCreate = d_streamCreate; d.streamSynchronize = d_streamSync; + d.streamDestroy = d_streamDestroy; d.funcGetAttribute = d_funcAttr; d.occupancy = d_occ; d.getErrorString = d_errStr; d.getErrorName = d_errName; +} diff --git a/proto-cuda/nvrtc/emu/serve-check.sh b/proto-cuda/nvrtc/emu/serve-check.sh new file mode 100755 index 00000000..557ddc95 --- /dev/null +++ b/proto-cuda/nvrtc/emu/serve-check.sh @@ -0,0 +1,62 @@ +#!/usr/bin/env bash +# The protocol check shared by the CUDA emulation (test.sh) and the OpenCL generic worker (proto-opencl/test-generic.sh): +# drives one worker through two jobs on pack A (one across the 32-bit nonce boundary), a background prepare of pack B, +# the swap, a refused job on A after the swap, and compares a sample of the found hashes with igneum-pow hash-bound. +# Usage: serve-check.sh +# The worker command is run as given plus the stdin script; it must already carry --serve --pack etc. +set -euo pipefail +PACK_A="$1"; PACK_B="$2"; LOG="$3"; shift 3 +HERE="$(cd "$(dirname "$0")" && pwd)" +ROOT="$(cd "$HERE/../../.." && pwd)" +POW="$ROOT/igneum-pow/target/release/igneum-pow" +seedline() { sed -n "s/^$2 //p" "$1/seeds.txt"; } +A_EPOCH="$(seedline "$PACK_A" epoch_seed_hex)"; A_DAY="$(seedline "$PACK_A" day_seed_hex)" +B_EPOCH="$(seedline "$PACK_B" epoch_seed_hex)"; B_DAY="$(seedline "$PACK_B" day_seed_hex)" +WAIT="${IGNEUM_PREPARE_WAIT:-40}" +PRE="0000000000000000000000000000000000000000000000000000000000000001" +{ + echo "job 1 $PRE ffffffffffffffff 0 64 $A_EPOCH $A_DAY" + echo "job 2 $PRE ffffffffffffffff 4294967264 64 $A_EPOCH $A_DAY" + echo "prepare $B_EPOCH $B_DAY $PACK_B" + sleep "$WAIT" + echo "job 3 $PRE ffffffffffffffff 128 32 $A_EPOCH $A_DAY" + sleep 1 + echo "job 4 $PRE ffffffffffffffff 0 32 $B_EPOCH $B_DAY" + echo "job 5 $PRE ffffffffffffffff 96 32 $A_EPOCH $A_DAY" + echo "quit" +} | "$@" | tee "$LOG" + +echo "== verdict" +bad=0; n=0 +check() { # job nonce epoch day + want="$("$POW" hash-bound --prehash "$PRE" --nonce "$2" --epoch-hex "$3" --day-hex "$4" | tail -1)" + got="$(grep "^found $1 $2 " "$LOG" | awk '{print $4}')" + n=$((n + 1)) + if [ "$want" != "$got" ]; then echo "MISMATCH job $1 nonce $2: worker $got cpu $want"; bad=$((bad + 1)); fi +} +grep -q '^ready .* prepare 1 ' "$LOG" || { echo "FAIL: no ready line with prepare 1"; exit 1; } +[ "$(grep -c '^found 1 ' "$LOG")" = 64 ] || { echo "FAIL: job 1 should find 64"; exit 1; } +[ "$(grep -c '^found 2 ' "$LOG")" = 64 ] || { echo "FAIL: job 2 should find 64"; exit 1; } +[ "$(grep -c '^found 3 ' "$LOG")" = 32 ] || { echo "FAIL: job 3 should find 32"; exit 1; } +[ "$(grep -c '^found 4 ' "$LOG")" = 32 ] || { echo "FAIL: job 4 should find 32 (was pack B prepared in time? IGNEUM_PREPARE_WAIT=$WAIT s)"; exit 1; } +grep -q "^prepared $B_EPOCH $B_DAY .*self-test PASS" "$LOG" || { echo "FAIL: no prepared line with self-test PASS"; exit 1; } +grep -q '^info switched to the prepared pair' "$LOG" || { echo "FAIL: no swap"; exit 1; } +# Job 5 is on pack A after the swap dropped it: an ahead-of-time worker refuses it (OpenCL host.c), the CUDA worker +# self-heals by finding pack A next to its first pack and building it again in the foreground (32 found, 5 done lines). +if grep -q '^error 5 epoch seed mismatch' "$LOG"; then + job5="refused (no pair for its seeds)" + [ "$(grep -c '^done ' "$LOG")" = 4 ] || { echo "FAIL: 4 done lines expected"; exit 1; } +elif grep -q '^info job 5 is for epoch .* building its pack' "$LOG" && [ "$(grep -c '^found 5 ' "$LOG")" = 32 ]; then + job5="served after the worker rebuilt pack A by itself (self-heal)" + [ "$(grep -c '^done ' "$LOG")" = 5 ] || { echo "FAIL: 5 done lines expected"; exit 1; } + for x in 96 127; do check 5 $x "$A_EPOCH" "$A_DAY"; done 2>/dev/null || true +else + echo "FAIL: job 5 was neither refused nor self-healed"; exit 1 +fi +grep -q '^found 2 4294967295 ' "$LOG" && grep -q '^found 2 4294967296 ' "$LOG" || { echo "FAIL: the 32-bit boundary nonces are missing"; exit 1; } +for x in 0 1 2 31 32 63; do check 1 $x "$A_EPOCH" "$A_DAY"; done +for x in 4294967264 4294967295 4294967296 4294967327; do check 2 $x "$A_EPOCH" "$A_DAY"; done +for x in 128 159; do check 3 $x "$A_EPOCH" "$A_DAY"; done +for x in 0 7 31; do check 4 $x "$B_EPOCH" "$B_DAY"; done +[ "$bad" = 0 ] || { echo "FAIL: $bad of $n sampled hashes differ from igneum-pow"; exit 1; } +echo "PASS: ready + prepare 1; 64 + 64 + 32 found on pack A, prepared B with self-test PASS, swapped, 32 found on B, job 5 $job5; $n sampled hashes (both packs, both sides of the 32-bit nonce boundary) equal igneum-pow hash-bound" diff --git a/proto-cuda/nvrtc/emu/test.sh b/proto-cuda/nvrtc/emu/test.sh new file mode 100755 index 00000000..7c30bed6 --- /dev/null +++ b/proto-cuda/nvrtc/emu/test.sh @@ -0,0 +1,61 @@ +#!/usr/bin/env bash +# Checks igneum-worker-cuda on the Mac (no NVIDIA GPU): the worker is built with the emulation backend, so the CUDA +# driver API and NVRTC are host functions and the packs' real kernel text runs on host threads (proto-cuda/emu shim). +# What is checked: (1) the source handed to NVRTC equals the pack's kernel files up to the host launch wrappers, with +# program.h and memhard.h byte-identical (the backend prints a PASS line per file); (2) --check self-tests pack A +# against its vectors.h; (3) --serve: jobs, a job across the 32-bit nonce boundary, a prepare of pack B in the +# background, the swap, the mismatch error after the swap; (4) every found hash of a sample of nonces equals +# igneum-pow hash-bound (the Rust CPU reference, the same code the miner re-checks shares with). +# What it cannot check: that the real NVRTC compiles the text and that the real driver runs it (windows-app/TEST.md). +# Usage: emu/test.sh [pack A dir] (default proto-cuda/packs/igneum-devnet-v4-epoch0; needs fetch-redist.sh first) +set -euo pipefail +HERE="$(cd "$(dirname "$0")" && pwd)" +NV="$(cd "$HERE/.." && pwd)" +ROOT="$(cd "$NV/../.." && pwd)" +EMU="$ROOT/proto-cuda/emu" +INC="${IGNEUM_CUDA_INC:-$NV/redist/include}" +POW="$ROOT/igneum-pow/target/release/igneum-pow" +OUT="$HERE/build" +PACK_A_SRC="${1:-$ROOT/proto-cuda/packs/igneum-devnet-v4-epoch0}" +[ -f "$INC/cuda.h" ] && [ -f "$INC/nvrtc.h" ] || { echo "no cuda.h / nvrtc.h under $INC: run $NV/fetch-redist.sh first" >&2; exit 1; } +[ -x "$POW" ] || (cd "$ROOT/igneum-pow" && cargo build --release 2>&1 | tail -2) +CXX="${CXX:-c++}" +rm -rf "$OUT"; mkdir -p "$OUT" + +# Pack A: a copy with seeds.txt as igneum-miner export-pack writes it. Pack B: a different epoch and day, exported now. +cp -R "$PACK_A_SRC" "$OUT/pack-a" +A_EPOCH="$(sed -n 's/^#define IGNEUM_SEED_BYTES_HEX "\(.*\)"/\1/p' "$OUT/pack-a/program.h")" +A_DAY="$(sed -n 's/^#define IGNEUM_DAY_BYTES_HEX "\(.*\)"/\1/p' "$OUT/pack-a/program.h")" +[ -n "$A_EPOCH" ] && [ -n "$A_DAY" ] || { echo "pack A has no byte seeds in program.h" >&2; exit 1; } +printf 'epoch_seed_hex %s\nday_seed_hex %s\nepoch_index 0\nday_index 0\ndaa_score 0\n' "$A_EPOCH" "$A_DAY" > "$OUT/pack-a/seeds.txt" +B_EPOCH="5e5d0c3b2a19f8e7d6c5b4a3928170f6e5d4c3b2a1908f7e6d5c4b3a2918f7e6" +B_DAY="69676e65756d2d6461792ffa51000000000000" +"$POW" export --epoch-hex "$B_EPOCH" --day-hex "$B_DAY" --out "$OUT/pack-b" > "$OUT/export-b.log"; head -2 "$OUT/export-b.log" +printf 'epoch_seed_hex %s\nday_seed_hex %s\nday_index 1\n' "$B_EPOCH" "$B_DAY" > "$OUT/pack-b/seeds.txt" + +# The kernels as the emulation runs them: the launch syntax rewritten (as emu/emu.sh does), each pack in its own namespace. +emu_kernel() { # pack dir, namespace, file stem + { echo '#include '; echo '#include '; echo "namespace $2 {" + sed -E 's/([A-Za-z_0-9]+)<<<([^,]+), ([^>]+)>>>\(/emu_launch(\1, \2, \3, /' "$1/$3.cu" + echo "}"; } > "$OUT/$3_$2.cpp" + "$CXX" -std=c++17 -O2 -w -I "$EMU" -I "$1" -c "$OUT/$3_$2.cpp" -o "$OUT/$3_$2.o" +} +emu_kernel "$OUT/pack-a" emu_pack_a kernel +emu_kernel "$OUT/pack-a" emu_pack_a kernel_bound +emu_kernel "$OUT/pack-b" emu_pack_b kernel +emu_kernel "$OUT/pack-b" emu_pack_b kernel_bound +"$CXX" -std=c++17 -O2 -Wall -Wextra -DIGNEUM_EMU -I "$NV" -I "$INC" -c "$NV/worker.cpp" -o "$OUT/worker.o" +"$CXX" -std=c++17 -O2 -Wall -Wextra -DIGNEUM_EMU -DIGNEUM_EMU_TWO_PACKS -I "$NV" -I "$EMU" -I "$INC" -c "$HERE/emu_backend.cpp" -o "$OUT/emu_backend.o" +"$CXX" -std=c++17 -O2 -w -I "$EMU" -c "$EMU/shim.cpp" -o "$OUT/shim.o" +"$CXX" -o "$OUT/igneum-worker-cuda-emu" "$OUT"/*.o -pthread +echo "built $OUT/igneum-worker-cuda-emu (CPU emulation, not a GPU build)" +export IGNEUM_EMU_PACK="$OUT/pack-a" IGNEUM_EMU_PACK2="$OUT/pack-b" + +echo "== --check pack A" +"$OUT/igneum-worker-cuda-emu" --check --pack "$OUT/pack-a" | tee "$OUT/check-a.log" +grep -q '^check PASS' "$OUT/check-a.log" +grep -c 'source check PASS' "$OUT/check-a.log" | grep -qx 2 + +echo "== --serve: two jobs on A, prepare B, a job on B (swap), a job on A after the swap (error), quit" +"$HERE/serve-check.sh" "$OUT/pack-a" "$OUT/pack-b" "$OUT/serve.log" "$OUT/igneum-worker-cuda-emu" --serve --pack "$OUT/pack-a" --batch-log2 13 +echo "NVRTC source check PASS for $(grep -c 'source check PASS' "$OUT/serve.log") files in --serve, $(grep -c 'source check PASS' "$OUT/check-a.log") in --check" diff --git a/proto-cuda/nvrtc/fetch-redist.sh b/proto-cuda/nvrtc/fetch-redist.sh new file mode 100755 index 00000000..eac6ada2 --- /dev/null +++ b/proto-cuda/nvrtc/fetch-redist.sh @@ -0,0 +1,52 @@ +#!/usr/bin/env bash +# Fetches the third-party pieces the one-click workers are built and packaged with, verifies them, and stages them +# under proto-cuda/nvrtc/redist/ (gitignored). Nothing third-party is committed. THIRD-PARTY.md records every file. +# NVIDIA cuda_nvrtc 12.8.93 (Windows x86_64): nvrtc64_120_0.dll and nvrtc-builtins64_128.dll (shipped), nvrtc.h (build), LICENSE +# NVIDIA cuda_cudart 12.8.90 (Windows x86_64): cuda.h (build only, the driver API header; not shipped) +# Khronos OpenCL-Headers v2026.05.29: CL/cl.h, cl_platform.h, cl_version.h, cl_ext.h (build only), LICENSE (Apache 2.0) +# The archives are cached in ~/.cache/igneum/ (the nvrtc archive is 305 MB). Usage: fetch-redist.sh +set -euo pipefail +HERE="$(cd "$(dirname "$0")" && pwd)" +OUT="$HERE/redist" +CACHE="${IGNEUM_REDIST_CACHE:-$HOME/.cache/igneum}" +BASE="https://developer.download.nvidia.com/compute/cuda/redist" +NVRTC_REL="cuda_nvrtc/windows-x86_64/cuda_nvrtc-windows-x86_64-12.8.93-archive.zip" +NVRTC_SHA="a63302a077f0248a743a1a7caa7dbd80d0fac56c6cfa9c41fa05fac9b7e5eda5" +CUDART_REL="cuda_cudart/windows-x86_64/cuda_cudart-windows-x86_64-12.8.90-archive.zip" +CUDART_SHA="4a39058fd8519444a81cfc7ae055d136f48d1a31ffa41ae255b35b2edd61e13b" +KHR_TAG="v2026.05.29" +KHR_BASE="https://raw.githubusercontent.com/KhronosGroup/OpenCL-Headers/$KHR_TAG" +# sha256 of the Khronos files at that tag, recorded 4 October 2026 (file sha256, one per line; bash 3 has no maps) +KHR_FILES="CL/cl.h 16d09614cd7eef73b4094089cc4ce2af777181d3ffe9613473fb6d552234a4d1 +CL/cl_platform.h b125cece6fe41f2e6690d88b24652c151cb6654da13b42180d69d5bd9ed526d7 +CL/cl_version.h 1e571dd68566330dd066cf778e07909a3cabeb2b969c5578c60f14baba7df1af +CL/cl_ext.h fe8916fd98b739cbc47ecab387bee2d258fa939c1bd2e10415c6f8f8b7055968 +LICENSE cfc7749b96f63bd31c3c42b5c471bf756814053e847c10f3eb003417bc523d30" + +sha() { shasum -a 256 "$1" | cut -d' ' -f1; } +fetch() { # url dest sha + if [ ! -f "$2" ] || [ "$(sha "$2")" != "$3" ]; then + echo "downloading $(basename "$2")" + curl -sS -L -m 3600 -o "$2.part" "$1" && mv "$2.part" "$2" + fi + [ "$(sha "$2")" = "$3" ] || { echo "sha256 mismatch for $2" >&2; exit 1; } +} + +mkdir -p "$CACHE/cuda-redist" "$CACHE/opencl-headers/CL" "$OUT/include/CL" "$OUT/bin" +fetch "$BASE/$NVRTC_REL" "$CACHE/cuda-redist/$(basename "$NVRTC_REL")" "$NVRTC_SHA" +fetch "$BASE/$CUDART_REL" "$CACHE/cuda-redist/$(basename "$CUDART_REL")" "$CUDART_SHA" +echo "$KHR_FILES" | while read -r f h; do fetch "$KHR_BASE/$f" "$CACHE/opencl-headers/$f" "$h"; done + +X="$(mktemp -d)" +unzip -q "$CACHE/cuda-redist/$(basename "$NVRTC_REL")" -d "$X" +unzip -q "$CACHE/cuda-redist/$(basename "$CUDART_REL")" -d "$X" +N="$X/cuda_nvrtc-windows-x86_64-12.8.93-archive" +C="$X/cuda_cudart-windows-x86_64-12.8.90-archive" +cp "$N/bin/nvrtc64_120_0.dll" "$N/bin/nvrtc-builtins64_128.dll" "$OUT/bin/" +cp "$N/include/nvrtc.h" "$C/include/cuda.h" "$OUT/include/" +cp "$N/LICENSE" "$OUT/LICENSE-NVIDIA-CUDA-EULA.txt" +cp "$CACHE"/opencl-headers/CL/*.h "$OUT/include/CL/" +cp "$CACHE/opencl-headers/LICENSE" "$OUT/LICENSE-Khronos-OpenCL-Headers.txt" +rm -rf "$X" +echo "staged in $OUT:" +(cd "$OUT" && find . -type f | sort | while read -r f; do printf ' %10d %s %s\n' "$(stat -f %z "$f")" "$(sha "$f" | cut -c1-16)" "$f"; done) diff --git a/proto-cuda/nvrtc/packfile.h b/proto-cuda/nvrtc/packfile.h new file mode 100644 index 00000000..fa407f93 --- /dev/null +++ b/proto-cuda/nvrtc/packfile.h @@ -0,0 +1,350 @@ +// packfile.h: read a program pack at run time (C99, header only). 4 October 2026. +// +// The one-click workers (proto-cuda/nvrtc/worker.cpp and proto-opencl/host.c with --pack) are built once and serve +// any pack, so the constants that host.cu and host.c take from program.h and vectors.h at compile time are read +// here from the same files at run time: the dataset and cache sizes, the seed words, the seeds as hex (seeds.txt, +// written by igneum-miner export-pack and --prepare-packs; program.h carries the same bytes), and the self-test +// values (cache head, last line and FNV-1a 64; dataset head, last word and 64 samples; the three vector warps). +// No JSON parser: the headers are scanned for "#define NAME value" and "NAME[..] = { numbers }". Comments are skipped. +// +// Included by C (proto-opencl/host.c) and C++ (worker.cpp). Everything is static. +#ifndef IGNEUM_PACKFILE_H +#define IGNEUM_PACKFILE_H +#include +#include +#include +#include +#include + +#define PF_MAX_WARPS 8 +#define PF_MAX_SAMPLES 256 +#define PF_HEX_CAP 520 + +typedef struct { + // program.h + uint32_t datasetLog2, cacheLog2Words, cacheSegments, datasetMode, generator; + uint32_t seedw[8], keyw[8]; + char seedString[600]; + // seeds.txt (or program.h): the seeds as the worker protocol carries them + char epochHex[65]; + char dayHex[PF_HEX_CAP]; + // vectors.h + int haveVectors; + int vecWarps; + uint32_t vecBase[PF_MAX_WARPS]; + uint64_t vecOut[PF_MAX_WARPS][32]; + uint32_t dsHead[16]; + uint32_t dsLastIndex, dsLast; + int nSamples; + uint32_t sampleIdx[PF_MAX_SAMPLES], sampleVal[PF_MAX_SAMPLES]; + uint32_t cacheHead[16], cacheLast[16]; + uint64_t cacheFnv; +} PfPack; + +static char* pf_read_file(const char* path, size_t* len) { + FILE* f = fopen(path, "rb"); + char* buf; + long n; + if (!f) return NULL; + fseek(f, 0, SEEK_END); n = ftell(f); fseek(f, 0, SEEK_SET); + if (n < 0) { fclose(f); return NULL; } + buf = (char*)malloc((size_t)n + 1); + if (!buf) { fclose(f); return NULL; } + if (fread(buf, 1, (size_t)n, f) != (size_t)n) { fclose(f); free(buf); return NULL; } + buf[n] = 0; + fclose(f); + if (len) *len = (size_t)n; + return buf; +} + +static char* pf_read_pack_file(const char* dir, const char* name, size_t* len) { + char path[2048]; + snprintf(path, sizeof(path), "%s/%s", dir, name); + return pf_read_file(path, len); +} + +// The text after "#define NAME" (NAME followed by a space or tab) at the start of a line, or NULL. +static const char* pf_find_define(const char* text, const char* name) { + size_t n = strlen(name); + const char* p = text; + while ((p = strstr(p, "#define ")) != NULL) { + const char* q = p + 8; + if (p == text || p[-1] == '\n' || p[-1] == '\r') { + if (strncmp(q, name, n) == 0 && (q[n] == ' ' || q[n] == '\t')) return q + n; + } + p = q; + } + return NULL; +} + +// A number with an optional C suffix (u, U, ull, ULL, l, L). *end is left after the suffix. +static int pf_number(const char* s, uint64_t* out, const char** end) { + char* e = NULL; + unsigned long long v; + if (!(isdigit((unsigned char)s[0]))) return 0; + v = strtoull(s, &e, 0); + if (e == s) return 0; + while (*e == 'u' || *e == 'U' || *e == 'l' || *e == 'L') ++e; + *out = (uint64_t)v; + *end = e; + return 1; +} + +static int pf_define_u32(const char* text, const char* name, uint32_t* out) { + const char* p = pf_find_define(text, name); + uint64_t v; + const char* e; + if (!p) return 0; + while (*p == ' ' || *p == '\t') ++p; + if (!pf_number(p, &v, &e)) return 0; + *out = (uint32_t)v; + return 1; +} + +static int pf_define_str(const char* text, const char* name, char* out, size_t cap) { + const char* p = pf_find_define(text, name); + const char* q; + size_t n; + if (!p) return 0; + while (*p == ' ' || *p == '\t') ++p; + if (*p != '"') return 0; + ++p; + q = strchr(p, '"'); + if (!q) return 0; + n = (size_t)(q - p); + if (n + 1 > cap) n = cap - 1; + memcpy(out, p, n); out[n] = 0; + return 1; +} + +// Numbers inside the braces that follow `start`, comments and commas skipped, nested braces flattened. +// Returns the count read (at most max). `start` points at or before the opening brace. +static int pf_brace_numbers(const char* start, uint64_t* out, int max) { + const char* p = start; + int depth = 0, n = 0; + while (*p && *p != '{') { if (*p == ';') return 0; ++p; } + if (*p != '{') return 0; + for (;;) { + uint64_t v; + const char* e; + if (!*p) break; + if (*p == '/' && p[1] == '/') { while (*p && *p != '\n') ++p; continue; } + if (*p == '/' && p[1] == '*') { const char* c = strstr(p + 2, "*/"); if (!c) break; p = c + 2; continue; } + if (*p == '{') { ++depth; ++p; continue; } + if (*p == '}') { --depth; ++p; if (depth == 0) break; continue; } + if (pf_number(p, &v, &e)) { if (n < max) out[n] = v; ++n; p = e; continue; } + ++p; + } + return n < max ? n : max; +} + +static int pf_define_words(const char* text, const char* name, uint32_t* out, int max) { + const char* p = pf_find_define(text, name); + uint64_t tmp[PF_MAX_SAMPLES]; + int n, i; + if (!p) return 0; + if (max > PF_MAX_SAMPLES) max = PF_MAX_SAMPLES; + n = pf_brace_numbers(p, tmp, max); + for (i = 0; i < n; ++i) out[i] = (uint32_t)tmp[i]; + return n; +} + +// "NAME" as a C symbol: `... NAME[...] = { ... };` or `... NAME = value;`. Returns the numbers found (at most max). +static int pf_symbol_numbers(const char* text, const char* name, uint64_t* out, int max) { + size_t n = strlen(name); + const char* p = text; + while ((p = strstr(p, name)) != NULL) { + const char* q = p + n; + int before = (p == text) || !(isalnum((unsigned char)p[-1]) || p[-1] == '_'); + p = q; + if (!before) continue; + if (isalnum((unsigned char)*q) || *q == '_') continue; + while (*q == '[' ) { const char* c = strchr(q, ']'); if (!c) return 0; q = c + 1; } + while (*q == ' ' || *q == '\t') ++q; + if (*q != '=') continue; + ++q; + while (*q == ' ' || *q == '\t') ++q; + if (*q == '{') return pf_brace_numbers(q, out, max); + { uint64_t v; const char* e; if (pf_number(q, &v, &e)) { if (max > 0) out[0] = v; return max > 0 ? 1 : 0; } } + return 0; + } + return 0; +} + +static int pf_symbol_u32s(const char* text, const char* name, uint32_t* out, int max) { + uint64_t tmp[PF_MAX_SAMPLES]; + int n, i; + if (max > PF_MAX_SAMPLES) max = PF_MAX_SAMPLES; + n = pf_symbol_numbers(text, name, tmp, max); + for (i = 0; i < n; ++i) out[i] = (uint32_t)tmp[i]; + return n; +} + +static int pf_unhex(const char* s, uint8_t* out, size_t cap, size_t* len) { + size_t n = strlen(s), i; + if (n % 2 || n / 2 > cap) return 0; + for (i = 0; i < n; i += 2) { + unsigned v = 0; + char two[3]; + two[0] = s[i]; two[1] = s[i + 1]; two[2] = 0; + if (!isxdigit((unsigned char)two[0]) || !isxdigit((unsigned char)two[1])) return 0; + if (sscanf(two, "%2x", &v) != 1) return 0; + out[i / 2] = (uint8_t)v; + } + *len = n / 2; + return 1; +} + +// seed_words_from_bytes of igneum-pow/src/seed.rs: FNV-1a 64 with four salts, each finalised. +static void pf_seed_words_from_bytes(const uint8_t* b, size_t n, uint32_t out[8]) { + uint64_t salt; + for (salt = 0; salt < 4; ++salt) { + uint64_t h = 0xcbf29ce484222325ull ^ (salt * 0x9E3779B97F4A7C15ull); + size_t i; + for (i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; } + h ^= h >> 33; h *= 0xff51afd7ed558ccdull; h ^= h >> 33; + out[2 * salt] = (uint32_t)h; + out[2 * salt + 1] = (uint32_t)(h >> 32); + } +} + +static uint64_t pf_fnv1a64(const void* p, size_t n) { + const uint8_t* b = (const uint8_t*)p; + uint64_t h = 0xcbf29ce484222325ull; + size_t i; + for (i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; } + return h; +} + +// One "key value" line of seeds.txt. +static int pf_seeds_line(const char* text, const char* key, char* out, size_t cap) { + size_t n = strlen(key); + const char* p = text; + while ((p = strstr(p, key)) != NULL) { + if ((p == text || p[-1] == '\n') && p[n] == ' ') { + const char* q = p + n + 1; + size_t len = 0; + while (q[len] && q[len] != '\r' && q[len] != '\n') ++len; + if (len + 1 > cap) return 0; + memcpy(out, q, len); out[len] = 0; + return 1; + } + p += n; + } + return 0; +} + +static int pf_fail(char* err, size_t cap, const char* msg) { if (err && cap) { strncpy(err, msg, cap - 1); err[cap - 1] = 0; } return 0; } + +// Reads program.h, seeds.txt (optional) and vectors.h (optional) of a pack directory. Returns 1 on success. +static int pf_load(const char* dir, PfPack* pk, char* err, size_t cap) { + char* prog; + char* seeds; + char* vec; + char ehex[65] = {0}, dhex[PF_HEX_CAP] = {0}; + uint8_t bytes[256]; + size_t blen; + memset(pk, 0, sizeof(*pk)); + prog = pf_read_pack_file(dir, "program.h", NULL); + if (!prog) { char m[480]; snprintf(m, sizeof(m), "cannot read %.400s/program.h", dir); return pf_fail(err, cap, m); } + if (!pf_define_u32(prog, "IGNEUM_DATASET_LOG2", &pk->datasetLog2)) { free(prog); return pf_fail(err, cap, "program.h has no IGNEUM_DATASET_LOG2"); } + if (!pf_define_u32(prog, "IGNEUM_DATASET_MODE", &pk->datasetMode)) pk->datasetMode = 0; + if (pk->datasetMode != 1) { free(prog); return pf_fail(err, cap, "the pack is not memory-hard (IGNEUM_DATASET_MODE 1); the one-click workers serve memory-hard packs only"); } + if (!pf_define_u32(prog, "IGNEUM_CACHE_LOG2_WORDS", &pk->cacheLog2Words)) { free(prog); return pf_fail(err, cap, "program.h has no IGNEUM_CACHE_LOG2_WORDS"); } + if (!pf_define_u32(prog, "IGNEUM_CACHE_SEGMENTS", &pk->cacheSegments)) { free(prog); return pf_fail(err, cap, "program.h has no IGNEUM_CACHE_SEGMENTS"); } + if (!pf_define_u32(prog, "IGNEUM_GENERATOR", &pk->generator)) pk->generator = 1; + if (pf_define_words(prog, "IGNEUM_SEEDW_INIT", pk->seedw, 8) != 8) { free(prog); return pf_fail(err, cap, "program.h has no IGNEUM_SEEDW_INIT with 8 words"); } + if (pf_define_words(prog, "IGNEUM_KEY_INIT", pk->keyw, 8) != 8) { free(prog); return pf_fail(err, cap, "program.h has no IGNEUM_KEY_INIT with 8 words"); } + if (!pf_define_str(prog, "IGNEUM_SEED_STRING", pk->seedString, sizeof(pk->seedString))) strncpy(pk->seedString, "(no IGNEUM_SEED_STRING)", sizeof(pk->seedString) - 1); + pf_define_str(prog, "IGNEUM_SEED_BYTES_HEX", ehex, sizeof(ehex)); + pf_define_str(prog, "IGNEUM_DAY_BYTES_HEX", dhex, sizeof(dhex)); + if (pk->datasetLog2 < 20 || pk->datasetLog2 > 31 || pk->cacheLog2Words < 16 || pk->cacheLog2Words > 30 || pk->cacheSegments == 0) { free(prog); return pf_fail(err, cap, "program.h sizes out of range"); } + free(prog); + + // seeds.txt wins when present; program.h's bytes must agree with it + seeds = pf_read_pack_file(dir, "seeds.txt", NULL); + if (seeds) { + char e2[65] = {0}, d2[PF_HEX_CAP] = {0}; + int he = pf_seeds_line(seeds, "epoch_seed_hex", e2, sizeof(e2)); + int hd = pf_seeds_line(seeds, "day_seed_hex", d2, sizeof(d2)); + free(seeds); + if (!he || !hd) return pf_fail(err, cap, "seeds.txt has no epoch_seed_hex / day_seed_hex line"); + if (ehex[0] && strcmp(ehex, e2) != 0) return pf_fail(err, cap, "seeds.txt epoch_seed_hex differs from program.h IGNEUM_SEED_BYTES_HEX"); + if (dhex[0] && strcmp(dhex, d2) != 0) return pf_fail(err, cap, "seeds.txt day_seed_hex differs from program.h IGNEUM_DAY_BYTES_HEX"); + strcpy(ehex, e2); strcpy(dhex, d2); + } + if (!ehex[0] || !dhex[0]) return pf_fail(err, cap, "no seeds: neither seeds.txt nor IGNEUM_SEED_BYTES_HEX / IGNEUM_DAY_BYTES_HEX in program.h (a pack from igneum-pow export --seed has no byte seeds)"); + if (strlen(ehex) != 64) return pf_fail(err, cap, "epoch seed is not 64 hex characters"); + strcpy(pk->epochHex, ehex); strcpy(pk->dayHex, dhex); + // The seed words derived from the bytes must be the pack's own words: otherwise the pack and the seeds disagree + { + uint32_t w[8]; + if (!pf_unhex(ehex, bytes, 32, &blen) || blen != 32) return pf_fail(err, cap, "epoch seed hex is malformed"); + pf_seed_words_from_bytes(bytes, 32, w); + if (memcmp(w, pk->seedw, 32) != 0) return pf_fail(err, cap, "the epoch seed bytes do not give the pack's IGNEUM_SEEDW_INIT (wrong seeds.txt for this pack?)"); + if (!pf_unhex(dhex, bytes, sizeof(bytes), &blen)) return pf_fail(err, cap, "day seed hex is malformed"); + pf_seed_words_from_bytes(bytes, blen, w); + if (memcmp(w, pk->keyw, 32) != 0) return pf_fail(err, cap, "the day seed bytes do not give the pack's IGNEUM_KEY_INIT (wrong seeds.txt for this pack?)"); + } + + vec = pf_read_pack_file(dir, "vectors.h", NULL); + if (vec) { + uint64_t tmp[PF_MAX_WARPS * 32]; + uint32_t warps = 0; + int n, w, l; + if (pf_define_u32(vec, "IGNEUM_VEC_WARPS", &warps) && warps >= 1 && warps <= PF_MAX_WARPS) { + pk->vecWarps = (int)warps; + n = pf_symbol_u32s(vec, "IGNEUM_VEC_BASE", pk->vecBase, PF_MAX_WARPS); + if (n != pk->vecWarps) pk->vecWarps = 0; + n = pf_symbol_numbers(vec, "IGNEUM_VEC_OUT", tmp, PF_MAX_WARPS * 32); + if (n != pk->vecWarps * 32) pk->vecWarps = 0; + for (w = 0; w < pk->vecWarps; ++w) for (l = 0; l < 32; ++l) pk->vecOut[w][l] = tmp[w * 32 + l]; + } + if (pf_symbol_u32s(vec, "IGNEUM_DS_HEAD", pk->dsHead, 16) == 16 && + pf_symbol_u32s(vec, "IGNEUM_DS_LAST_INDEX", &pk->dsLastIndex, 1) == 1 && + pf_symbol_u32s(vec, "IGNEUM_DS_LAST", &pk->dsLast, 1) == 1 && + pf_symbol_u32s(vec, "IGNEUM_CACHE_HEAD", pk->cacheHead, 16) == 16 && + pf_symbol_u32s(vec, "IGNEUM_CACHE_LAST", pk->cacheLast, 16) == 16 && + pf_symbol_numbers(vec, "IGNEUM_CACHE_FNV64", &pk->cacheFnv, 1) == 1 && pk->vecWarps > 0) { + uint32_t ns = 0; + pk->haveVectors = 1; + if (pf_define_u32(vec, "IGNEUM_DS_SAMPLES", &ns) && ns > 0 && ns <= PF_MAX_SAMPLES) { + int a = pf_symbol_u32s(vec, "IGNEUM_DS_SAMPLE_INDEX", pk->sampleIdx, (int)ns); + int b = pf_symbol_u32s(vec, "IGNEUM_DS_SAMPLE_VALUE", pk->sampleVal, (int)ns); + pk->nSamples = (a == (int)ns && b == (int)ns) ? (int)ns : 0; + } + } + free(vec); + } + return 1; +} + +// The self-test verdict from values the host read back from the device. `vec` holds vecWarps x 32 outputs of the +// bound kernel run with the pack's own seed words as init words (that is igneum_hash of kernel.cu). Writes one line. +static int pf_selftest(const PfPack* pk, const uint32_t* cacheHead, const uint32_t* cacheLast, uint64_t cacheFnv, + const uint32_t* dsHead, uint32_t dsLast, const uint32_t* sampleVals, const uint64_t* vec, + char* out, size_t cap) { + int okCH = memcmp(cacheHead, pk->cacheHead, 64) == 0, okCL = memcmp(cacheLast, pk->cacheLast, 64) == 0; + int okFnv = (cacheFnv == pk->cacheFnv); + int okDH = memcmp(dsHead, pk->dsHead, 64) == 0, okDL = (dsLast == pk->dsLast); + int badS = 0, badV = 0, i, w, l, firstBadWarp = -1, firstBadLane = -1; + for (i = 0; i < pk->nSamples; ++i) if (sampleVals[i] != pk->sampleVal[i]) ++badS; + for (w = 0; w < pk->vecWarps; ++w) for (l = 0; l < 32; ++l) if (vec[w * 32 + l] != pk->vecOut[w][l]) { if (firstBadWarp < 0) { firstBadWarp = w; firstBadLane = l; } ++badV; } + if (okCH && okCL && okFnv && okDH && okDL && badS == 0 && badV == 0) { + snprintf(out, cap, "self-test PASS (cache head, last line and FNV-1a 64 %016llx; dataset head, word [%u] and %d samples; %d of %d vector lanes)", + (unsigned long long)cacheFnv, pk->dsLastIndex, pk->nSamples, pk->vecWarps * 32, pk->vecWarps * 32); + return 1; + } + snprintf(out, cap, "self-test FAIL (cache head %s, cache last %s, cache FNV %016llx vs pack %016llx %s, dataset head %s, dataset last %s, samples %d bad of %d, vector lanes %d bad of %d%s)", + okCH ? "ok" : "BAD", okCL ? "ok" : "BAD", (unsigned long long)cacheFnv, (unsigned long long)pk->cacheFnv, okFnv ? "ok" : "BAD", + okDH ? "ok" : "BAD", okDL ? "ok" : "BAD", badS, pk->nSamples, badV, pk->vecWarps * 32, + firstBadWarp >= 0 ? " (first bad lane in the warp at base nonce" : ""); + if (firstBadWarp >= 0) { + size_t n = strlen(out); + snprintf(out + n, cap > n ? cap - n : 0, " %u lane %d: device %016llx expected %016llx)", pk->vecBase[firstBadWarp], firstBadLane, + (unsigned long long)vec[firstBadWarp * 32 + firstBadLane], (unsigned long long)pk->vecOut[firstBadWarp][firstBadLane]); + } + return 0; +} + +#endif diff --git a/proto-cuda/nvrtc/worker.cpp b/proto-cuda/nvrtc/worker.cpp new file mode 100644 index 00000000..08dab887 --- /dev/null +++ b/proto-cuda/nvrtc/worker.cpp @@ -0,0 +1,791 @@ +// 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 and , 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 +// prepare compile that pack in the background, build its cache +// and dataset, self-test it; a job on it then switches +// quit +// stdout: ready cuda pack dataset-log2 N batch B regs R prepare 1 path nvrtc ... +// found +// done +// error +// prepared ... | prepare-failed +// info ... +// The first pack comes from --pack (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 [--device D] [--batch-log2 22] [--block-warps 1] [--arch sm_120|auto] +// igneum-worker-cuda --check --pack [--device D] compile, build, self-test, print timings, exit 0/1 + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "cuda_api.h" +#include "packfile.h" + +#ifdef _WIN32 +#define WIN32_LEAN_AND_MEAN +#include +#else +#include +#include +#endif + +static const char* WORKER_VERSION = "1.0 (4 October 2026)"; + +// --------------------------------------------------------------------------------------------- +// Helpers + +static double wallMs() { + using namespace std::chrono; + return duration(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_0_0.dll next to the exe (any major), then a toolkit on PATH. IGNEUM_NVRTC_DLL overrides. +static std::vector nvrtcCandidates() { + std::vector 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 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 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 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 image; + std::vector 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& 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; + const char* opts[2] = { archOpt.c_str(), "--std=c++17" }; + r = c.rtc.compileProgram(prog, 2, opts); + { + size_t logSize = 0; + if (c.rtc.getProgramLogSize(prog, &logSize) == NVRTC_SUCCESS && logSize > 1) { + std::vector 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 samples((size_t)(pk.nSamples > 0 ? pk.nSamples : 1), 0u); + std::vector vec((size_t)pk.vecWarps * 32u, 0ull); + std::vector 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 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 GPU worker for igneum-miner --worker: jobs on stdin, found/done lines on stdout\n" + " --check --pack 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 is required (igneum-miner export-pack 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 split(const std::string& line) { + std::vector 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 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 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 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 )", 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; +}