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 <noreply@anthropic.com>
This commit is contained in:
igneum-labs
parent
4c3f8420ce
commit
f535abb59d
10 changed files with 1691 additions and 0 deletions
8
.gitignore
vendored
8
.gitignore
vendored
|
|
@ -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
|
||||
|
|
|
|||
62
proto-cuda/nvrtc/THIRD-PARTY.md
Normal file
62
proto-cuda/nvrtc/THIRD-PARTY.md
Normal file
|
|
@ -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.
|
||||
31
proto-cuda/nvrtc/build-windows.sh
Executable file
31
proto-cuda/nvrtc/build-windows.sh
Executable file
|
|
@ -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
|
||||
61
proto-cuda/nvrtc/cuda_api.h
Normal file
61
proto-cuda/nvrtc/cuda_api.h
Normal file
|
|
@ -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 <cuda.h>
|
||||
#include <nvrtc.h>
|
||||
|
||||
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
|
||||
213
proto-cuda/nvrtc/emu/emu_backend.cpp
Normal file
213
proto-cuda/nvrtc/emu/emu_backend.cpp
Normal file
|
|
@ -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 <cstdint>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstring>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <list>
|
||||
#include <map>
|
||||
|
||||
#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<std::string, std::string> headers;
|
||||
std::list<std::string> 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 <class T> static T arg(void** params, int i) { return *(T*)params[i]; }
|
||||
template <class T> 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<uint32_t>(params, 0), arg<uint32_t>(params, 1)); return CUDA_SUCCESS;
|
||||
case 2: emu_launch(emu_pack_a::igneum_build, gx, bx, dptr<uint32_t>(params, 0), dptr<const uint32_t>(params, 1), arg<uint32_t>(params, 2)); return CUDA_SUCCESS;
|
||||
case 3: emu_launch(emu_pack_a::igneum_hash_bound, gx, bx, dptr<const uint32_t>(params, 0), dptr<uint64_t>(params, 1), arg<uint32_t>(params, 2), arg<uint32_t>(params, 3), arg<emu_pack_a::IgneumInitWords>(params, 4)); return CUDA_SUCCESS;
|
||||
#ifdef IGNEUM_EMU_TWO_PACKS
|
||||
case 4: emu_launch(emu_pack_b::igneum_cache_fill, gx, bx, dptr<uint32_t>(params, 0), arg<uint32_t>(params, 1)); return CUDA_SUCCESS;
|
||||
case 5: emu_launch(emu_pack_b::igneum_build, gx, bx, dptr<uint32_t>(params, 0), dptr<const uint32_t>(params, 1), arg<uint32_t>(params, 2)); return CUDA_SUCCESS;
|
||||
case 6: emu_launch(emu_pack_b::igneum_hash_bound, gx, bx, dptr<const uint32_t>(params, 0), dptr<uint64_t>(params, 1), arg<uint32_t>(params, 2), arg<uint32_t>(params, 3), arg<emu_pack_b::IgneumInitWords>(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;
|
||||
}
|
||||
62
proto-cuda/nvrtc/emu/serve-check.sh
Executable file
62
proto-cuda/nvrtc/emu/serve-check.sh
Executable file
|
|
@ -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 <pack A dir> <pack B dir> <log file> <worker command...>
|
||||
# The worker command is run as given plus the stdin script; it must already carry --serve --pack <pack A> 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"
|
||||
61
proto-cuda/nvrtc/emu/test.sh
Executable file
61
proto-cuda/nvrtc/emu/test.sh
Executable file
|
|
@ -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 <cuda_runtime.h>'; echo '#include <cstdint>'; 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"
|
||||
52
proto-cuda/nvrtc/fetch-redist.sh
Executable file
52
proto-cuda/nvrtc/fetch-redist.sh
Executable file
|
|
@ -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)
|
||||
350
proto-cuda/nvrtc/packfile.h
Normal file
350
proto-cuda/nvrtc/packfile.h
Normal file
|
|
@ -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 <stdint.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <ctype.h>
|
||||
|
||||
#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 <name> 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
|
||||
791
proto-cuda/nvrtc/worker.cpp
Normal file
791
proto-cuda/nvrtc/worker.cpp
Normal file
|
|
@ -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 <cuda_runtime.h> and <cstdint>, which nvcc takes from
|
||||
// the toolkit. emu/test.sh checks the equality on the Mac. The memory-hard core is therefore the same text host.cu
|
||||
// compiles, and the miner's CPU re-check of every found nonce covers the rest.
|
||||
//
|
||||
// Worker protocol (the same lines as proto-cuda/host.cu --serve and proto-opencl/host.c --serve):
|
||||
// stdin: job <job_id> <header_prehash_hex 64> <target_hex 16> <nonce_start u64> <nonce_count u64> <epoch_seed_hex 64> <day_seed_hex>
|
||||
// prepare <epoch_seed_hex 64> <day_seed_hex> <pack_dir> compile that pack in the background, build its cache
|
||||
// and dataset, self-test it; a job on it then switches
|
||||
// quit
|
||||
// stdout: ready cuda <device> pack <seed string> dataset-log2 N batch B regs R prepare 1 path nvrtc ...
|
||||
// found <job_id> <nonce u64> <hash_hex 16>
|
||||
// done <job_id> <hashes> <ms>
|
||||
// error <job_id> <text>
|
||||
// prepared <epoch_seed_hex> <day_seed_hex> <ms> ... | prepare-failed <epoch_seed_hex> <day_seed_hex> <text>
|
||||
// info ...
|
||||
// The first pack comes from --pack <dir> (igneum-miner export-pack writes it; the launcher passes it). Every pack is
|
||||
// self-tested before it serves a job: cache head, last line and FNV-1a 64, dataset head, last word and 64 samples,
|
||||
// and the three vector warps of vectors.h through the bound kernel with the pack's own seed words. A pack that fails
|
||||
// is refused.
|
||||
//
|
||||
// Usage: igneum-worker-cuda --serve --pack <dir> [--device D] [--batch-log2 22] [--block-warps 1] [--arch sm_120|auto]
|
||||
// igneum-worker-cuda --check --pack <dir> [--device D] compile, build, self-test, print timings, exit 0/1
|
||||
|
||||
#include <cstdint>
|
||||
#include <cstdarg>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstring>
|
||||
#include <chrono>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <thread>
|
||||
#include <atomic>
|
||||
#include <iostream>
|
||||
|
||||
#include "cuda_api.h"
|
||||
#include "packfile.h"
|
||||
|
||||
#ifdef _WIN32
|
||||
#define WIN32_LEAN_AND_MEAN
|
||||
#include <windows.h>
|
||||
#else
|
||||
#include <dlfcn.h>
|
||||
#include <dirent.h>
|
||||
#endif
|
||||
|
||||
static const char* WORKER_VERSION = "1.0 (4 October 2026)";
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// Helpers
|
||||
|
||||
static double wallMs() {
|
||||
using namespace std::chrono;
|
||||
return duration<double, std::milli>(steady_clock::now().time_since_epoch()).count();
|
||||
}
|
||||
|
||||
static void emit(const std::string& s) { std::fputs(s.c_str(), stdout); std::fputc('\n', stdout); std::fflush(stdout); }
|
||||
static void info(const std::string& s) { emit("info " + s); }
|
||||
|
||||
static std::string fmt(const char* f, ...) {
|
||||
char buf[2048];
|
||||
va_list ap;
|
||||
va_start(ap, f);
|
||||
vsnprintf(buf, sizeof(buf), f, ap);
|
||||
va_end(ap);
|
||||
return buf;
|
||||
}
|
||||
|
||||
static std::string readText(const std::string& path, bool& ok) {
|
||||
size_t n = 0;
|
||||
char* b = pf_read_file(path.c_str(), &n);
|
||||
if (!b) { ok = false; return ""; }
|
||||
std::string s(b, n);
|
||||
free(b);
|
||||
ok = true;
|
||||
return s;
|
||||
}
|
||||
|
||||
static std::string exeDir() {
|
||||
#ifdef _WIN32
|
||||
char buf[MAX_PATH];
|
||||
DWORD n = GetModuleFileNameA(nullptr, buf, MAX_PATH);
|
||||
std::string p(buf, n);
|
||||
size_t i = p.find_last_of("\\/");
|
||||
return i == std::string::npos ? "." : p.substr(0, i);
|
||||
#else
|
||||
return ".";
|
||||
#endif
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// Loading the two libraries
|
||||
|
||||
static void* libOpen(const std::string& name) {
|
||||
#ifdef _WIN32
|
||||
return (void*)LoadLibraryA(name.c_str());
|
||||
#else
|
||||
return dlopen(name.c_str(), RTLD_NOW);
|
||||
#endif
|
||||
}
|
||||
static void* libSym(void* lib, const char* name) {
|
||||
#ifdef _WIN32
|
||||
return (void*)GetProcAddress((HMODULE)lib, name);
|
||||
#else
|
||||
return dlsym(lib, name);
|
||||
#endif
|
||||
}
|
||||
|
||||
#define LOAD_SYM(table, field, name) do { table.field = (decltype(table.field))libSym(lib, name); if (!table.field) { missing += std::string(missing.empty() ? "" : ", ") + name; } } while (0)
|
||||
|
||||
static bool loadDriver(Drv& d, std::string& err, std::string& libName) {
|
||||
#ifdef IGNEUM_EMU
|
||||
emu_fill_driver(d); libName = "emulation (host threads, no GPU)"; (void)err; return true;
|
||||
#else
|
||||
#ifdef _WIN32
|
||||
const char* names[] = { "nvcuda.dll" };
|
||||
#else
|
||||
const char* names[] = { "libcuda.so.1", "libcuda.so" };
|
||||
#endif
|
||||
void* lib = nullptr;
|
||||
for (const char* n : names) { lib = libOpen(n); if (lib) { libName = n; break; } }
|
||||
if (!lib) { err = "the CUDA driver library (nvcuda.dll) is not installed: install or update the NVIDIA driver"; return false; }
|
||||
std::string missing;
|
||||
LOAD_SYM(d, init, "cuInit");
|
||||
LOAD_SYM(d, driverGetVersion, "cuDriverGetVersion");
|
||||
LOAD_SYM(d, deviceGetCount, "cuDeviceGetCount");
|
||||
LOAD_SYM(d, deviceGet, "cuDeviceGet");
|
||||
LOAD_SYM(d, deviceGetName, "cuDeviceGetName");
|
||||
LOAD_SYM(d, deviceGetAttribute, "cuDeviceGetAttribute");
|
||||
LOAD_SYM(d, deviceTotalMem, "cuDeviceTotalMem_v2");
|
||||
LOAD_SYM(d, primaryCtxSetFlags, "cuDevicePrimaryCtxSetFlags_v2");
|
||||
LOAD_SYM(d, primaryCtxRetain, "cuDevicePrimaryCtxRetain");
|
||||
LOAD_SYM(d, primaryCtxRelease, "cuDevicePrimaryCtxRelease_v2");
|
||||
LOAD_SYM(d, ctxSetCurrent, "cuCtxSetCurrent");
|
||||
LOAD_SYM(d, ctxSynchronize, "cuCtxSynchronize");
|
||||
LOAD_SYM(d, memGetInfo, "cuMemGetInfo_v2");
|
||||
LOAD_SYM(d, memAlloc, "cuMemAlloc_v2");
|
||||
LOAD_SYM(d, memFree, "cuMemFree_v2");
|
||||
LOAD_SYM(d, memcpyDtoH, "cuMemcpyDtoH_v2");
|
||||
LOAD_SYM(d, moduleLoadData, "cuModuleLoadData");
|
||||
LOAD_SYM(d, moduleUnload, "cuModuleUnload");
|
||||
LOAD_SYM(d, moduleGetFunction, "cuModuleGetFunction");
|
||||
LOAD_SYM(d, launchKernel, "cuLaunchKernel");
|
||||
LOAD_SYM(d, streamCreate, "cuStreamCreate");
|
||||
LOAD_SYM(d, streamSynchronize, "cuStreamSynchronize");
|
||||
LOAD_SYM(d, streamDestroy, "cuStreamDestroy_v2");
|
||||
LOAD_SYM(d, funcGetAttribute, "cuFuncGetAttribute");
|
||||
LOAD_SYM(d, occupancy, "cuOccupancyMaxActiveBlocksPerMultiprocessor");
|
||||
LOAD_SYM(d, getErrorString, "cuGetErrorString");
|
||||
LOAD_SYM(d, getErrorName, "cuGetErrorName");
|
||||
if (!missing.empty()) { err = "the driver library lacks " + missing + " (driver too old; CUDA 11 or newer is needed)"; return false; }
|
||||
return true;
|
||||
#endif
|
||||
}
|
||||
|
||||
#ifdef _WIN32
|
||||
// nvrtc64_<major>0_0.dll next to the exe (any major), then a toolkit on PATH. IGNEUM_NVRTC_DLL overrides.
|
||||
static std::vector<std::string> nvrtcCandidates() {
|
||||
std::vector<std::string> v;
|
||||
if (const char* o = std::getenv("IGNEUM_NVRTC_DLL")) v.push_back(o);
|
||||
std::string dir = exeDir();
|
||||
WIN32_FIND_DATAA fd;
|
||||
HANDLE h = FindFirstFileA((dir + "\\nvrtc64_*_0.dll").c_str(), &fd);
|
||||
if (h != INVALID_HANDLE_VALUE) {
|
||||
do { std::string n = fd.cFileName; if (n.find(".alt.") == std::string::npos) v.push_back(dir + "\\" + n); } while (FindNextFileA(h, &fd));
|
||||
FindClose(h);
|
||||
}
|
||||
v.push_back("nvrtc64_120_0.dll");
|
||||
v.push_back("nvrtc64_130_0.dll");
|
||||
if (const char* cp = std::getenv("CUDA_PATH")) { v.push_back(std::string(cp) + "\\bin\\nvrtc64_120_0.dll"); v.push_back(std::string(cp) + "\\bin\\nvrtc64_130_0.dll"); }
|
||||
return v;
|
||||
}
|
||||
#endif
|
||||
|
||||
static bool loadNvrtc(Rtc& r, std::string& err, std::string& libName) {
|
||||
#ifdef IGNEUM_EMU
|
||||
emu_fill_nvrtc(r); libName = "emulation (source recorded and checked, nothing compiled)"; (void)err; return true;
|
||||
#else
|
||||
void* lib = nullptr;
|
||||
#ifdef _WIN32
|
||||
for (const std::string& n : nvrtcCandidates()) { lib = libOpen(n); if (lib) { libName = n; break; } }
|
||||
if (!lib) { err = "nvrtc64_120_0.dll (and nvrtc-builtins64_128.dll) must sit next to " + exeDir() + "\\igneum-worker-cuda.exe; they are in the package"; return false; }
|
||||
#else
|
||||
const char* names[] = { "libnvrtc.so.12", "libnvrtc.so" };
|
||||
for (const char* n : names) { lib = libOpen(n); if (lib) { libName = n; break; } }
|
||||
if (!lib) { err = "libnvrtc.so.12 not found"; return false; }
|
||||
#endif
|
||||
std::string missing;
|
||||
LOAD_SYM(r, version, "nvrtcVersion");
|
||||
LOAD_SYM(r, createProgram, "nvrtcCreateProgram");
|
||||
LOAD_SYM(r, destroyProgram, "nvrtcDestroyProgram");
|
||||
LOAD_SYM(r, compileProgram, "nvrtcCompileProgram");
|
||||
LOAD_SYM(r, getProgramLogSize, "nvrtcGetProgramLogSize");
|
||||
LOAD_SYM(r, getProgramLog, "nvrtcGetProgramLog");
|
||||
LOAD_SYM(r, getPTXSize, "nvrtcGetPTXSize");
|
||||
LOAD_SYM(r, getPTX, "nvrtcGetPTX");
|
||||
LOAD_SYM(r, getCUBINSize, "nvrtcGetCUBINSize");
|
||||
LOAD_SYM(r, getCUBIN, "nvrtcGetCUBIN");
|
||||
LOAD_SYM(r, addNameExpression, "nvrtcAddNameExpression");
|
||||
LOAD_SYM(r, getLoweredName, "nvrtcGetLoweredName");
|
||||
LOAD_SYM(r, getErrorString, "nvrtcGetErrorString");
|
||||
if (!missing.empty()) { err = "the NVRTC library lacks " + missing; return false; }
|
||||
r.getNumSupportedArchs = (decltype(r.getNumSupportedArchs))libSym(lib, "nvrtcGetNumSupportedArchs");
|
||||
r.getSupportedArchs = (decltype(r.getSupportedArchs))libSym(lib, "nvrtcGetSupportedArchs");
|
||||
return true;
|
||||
#endif
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// The device context
|
||||
|
||||
struct Ctx {
|
||||
Drv drv;
|
||||
Rtc rtc;
|
||||
CUdevice dev = 0;
|
||||
CUcontext ctx = nullptr;
|
||||
std::string name;
|
||||
int major = 0, minor = 0, sms = 0, driverVersion = 0, rtcMajor = 0, rtcMinor = 0;
|
||||
std::vector<int> rtcArchs; // what NVRTC can target (empty when the query is unavailable)
|
||||
std::string archOpt; // "sm_120" or "compute_120": what the packs are compiled for
|
||||
bool ptx = false; // true when archOpt is compute_XY (PTX, driver JIT)
|
||||
std::string why; // how archOpt was chosen
|
||||
int blockWarps = 1;
|
||||
std::string err(CUresult r) { const char* s = nullptr; if (drv.getErrorString) drv.getErrorString(r, &s); return s ? s : "CUDA driver error"; }
|
||||
};
|
||||
|
||||
#define DRV_CHECK(c, call, what) do { CUresult r_ = (call); if (r_ != CUDA_SUCCESS) { err = std::string(what) + ": " + (c).err(r_); return false; } } while (0)
|
||||
|
||||
static bool openDevice(Ctx& c, int device, const std::string& archArg, std::string& err) {
|
||||
DRV_CHECK(c, c.drv.init(0), "cuInit");
|
||||
int count = 0;
|
||||
DRV_CHECK(c, c.drv.deviceGetCount(&count), "cuDeviceGetCount");
|
||||
if (count == 0) { err = "no CUDA device"; return false; }
|
||||
if (device < 0 || device >= count) { err = fmt("device %d out of range (%d devices)", device, count); return false; }
|
||||
DRV_CHECK(c, c.drv.deviceGet(&c.dev, device), "cuDeviceGet");
|
||||
char name[256] = {0};
|
||||
DRV_CHECK(c, c.drv.deviceGetName(name, 255, c.dev), "cuDeviceGetName");
|
||||
c.name = name;
|
||||
for (char& ch : c.name) if (ch == ' ') ch = '_';
|
||||
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, c.dev), "compute capability major");
|
||||
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, c.dev), "compute capability minor");
|
||||
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.sms, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, c.dev), "multiprocessor count");
|
||||
c.drv.driverGetVersion(&c.driverVersion);
|
||||
// Blocking sync, set before the context exists: the host thread sleeps in cuStreamSynchronize instead of spinning
|
||||
// (one full core per worker at the default spin schedule, measured on the RTX 5090 with eight workers, 3 Oct 2026).
|
||||
c.drv.primaryCtxSetFlags(c.dev, CU_CTX_SCHED_BLOCKING_SYNC);
|
||||
DRV_CHECK(c, c.drv.primaryCtxRetain(&c.ctx, c.dev), "cuDevicePrimaryCtxRetain");
|
||||
DRV_CHECK(c, c.drv.ctxSetCurrent(c.ctx), "cuCtxSetCurrent");
|
||||
c.rtc.version(&c.rtcMajor, &c.rtcMinor);
|
||||
if (c.rtc.getNumSupportedArchs && c.rtc.getSupportedArchs) {
|
||||
int n = 0;
|
||||
if (c.rtc.getNumSupportedArchs(&n) == NVRTC_SUCCESS && n > 0 && n < 256) { c.rtcArchs.assign((size_t)n, 0); if (c.rtc.getSupportedArchs(c.rtcArchs.data()) != NVRTC_SUCCESS) c.rtcArchs.clear(); }
|
||||
}
|
||||
// The target: the device's own SASS (sm_XY) when this NVRTC knows the architecture, else PTX for the newest
|
||||
// architecture it knows below the device's, which the driver JIT-compiles forward. --arch overrides.
|
||||
int cc = c.major * 10 + c.minor;
|
||||
if (archArg != "auto" && !archArg.empty()) {
|
||||
c.archOpt = archArg; c.ptx = archArg.rfind("compute_", 0) == 0; c.why = "--arch";
|
||||
} else if (c.rtcArchs.empty()) {
|
||||
c.archOpt = fmt("sm_%d", cc); c.why = "the device's architecture (NVRTC did not list its targets)";
|
||||
} else {
|
||||
bool known = false; int best = 0;
|
||||
for (int a : c.rtcArchs) { if (a == cc) known = true; if (a <= cc && a > best) best = a; }
|
||||
if (known) { c.archOpt = fmt("sm_%d", cc); c.why = "the device's architecture, listed by NVRTC"; }
|
||||
else if (best > 0) { c.archOpt = fmt("compute_%d", best); c.ptx = true; c.why = fmt("this NVRTC does not know sm_%d; PTX for compute_%d, JIT-compiled by the driver", cc, best); }
|
||||
else { c.archOpt = fmt("compute_%d", c.rtcArchs.front()); c.ptx = true; c.why = fmt("this NVRTC knows nothing at or below sm_%d; PTX for its oldest target", cc); }
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// NVRTC: compile one of the pack's kernel files
|
||||
|
||||
// The pack's kernel files end in host-side launch wrappers (cudaError_t igneum_launch_* with <<< >>> launches) that
|
||||
// nvcc compiles for the host. NVRTC compiles device code only, so the text is cut there. The cut is checked: nothing
|
||||
// device-side may follow it.
|
||||
static bool deviceOnly(const std::string& text, std::string& out, std::string& err) {
|
||||
size_t cut = text.find("\n// Host-side launch wrappers");
|
||||
if (cut == std::string::npos) cut = text.find("\ncudaError_t ");
|
||||
if (cut == std::string::npos) { out = text; return true; }
|
||||
std::string tail = text.substr(cut + 1);
|
||||
if (tail.find("__global__") != std::string::npos || tail.find("__device__") != std::string::npos) { err = "device code after the host launch wrappers; the pack layout is not the one this worker knows"; return false; }
|
||||
out = text.substr(0, cut + 1);
|
||||
return true;
|
||||
}
|
||||
|
||||
static const char* STUB_CUDA_RUNTIME =
|
||||
"// igneum-worker-cuda: stand-in for <cuda_runtime.h> under NVRTC, which has the device built-ins already\n"
|
||||
"#pragma once\n"
|
||||
"#ifndef __CUDACC_RTC__\n#error \"this stub is for NVRTC only\"\n#endif\n"
|
||||
"#ifdef __SIZE_TYPE__\ntypedef __SIZE_TYPE__ size_t;\n#elif defined(__LP64__) || defined(_LP64)\ntypedef unsigned long size_t;\n#else\ntypedef unsigned long long size_t;\n#endif\n";
|
||||
static const char* STUB_CSTDINT =
|
||||
"// igneum-worker-cuda: stand-in for <cstdint> under NVRTC (the fixed-width types the packs use)\n"
|
||||
"#pragma once\n"
|
||||
"typedef signed char int8_t; typedef unsigned char uint8_t; typedef short int16_t; typedef unsigned short uint16_t;\n"
|
||||
"typedef int int32_t; typedef unsigned int uint32_t;\n"
|
||||
"#if defined(__LP64__) || defined(_LP64)\ntypedef long int64_t; typedef unsigned long uint64_t;\n"
|
||||
"#else\ntypedef long long int64_t; typedef unsigned long long uint64_t;\n#endif\n";
|
||||
|
||||
struct Compiled {
|
||||
std::vector<char> image;
|
||||
std::vector<std::string> lowered;
|
||||
double ms = 0;
|
||||
std::string log;
|
||||
};
|
||||
|
||||
static bool rtcCompile(Ctx& c, const std::string& src, const char* name, const std::string& programH, const std::string& memhardH,
|
||||
const std::vector<std::string>& nameExprs, Compiled& out, std::string& err) {
|
||||
double t0 = wallMs();
|
||||
const char* headers[4] = { STUB_CUDA_RUNTIME, STUB_CSTDINT, programH.c_str(), memhardH.c_str() };
|
||||
const char* names[4] = { "cuda_runtime.h", "cstdint", "program.h", "memhard.h" };
|
||||
nvrtcProgram prog = nullptr;
|
||||
nvrtcResult r = c.rtc.createProgram(&prog, src.c_str(), name, 4, headers, names);
|
||||
if (r != NVRTC_SUCCESS) { err = std::string("nvrtcCreateProgram: ") + c.rtc.getErrorString(r); return false; }
|
||||
for (const std::string& e : nameExprs) {
|
||||
r = c.rtc.addNameExpression(prog, e.c_str());
|
||||
if (r != NVRTC_SUCCESS) { err = "nvrtcAddNameExpression " + e + ": " + c.rtc.getErrorString(r); c.rtc.destroyProgram(&prog); return false; }
|
||||
}
|
||||
std::string archOpt = "--gpu-architecture=" + c.archOpt;
|
||||
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<char> log(logSize);
|
||||
c.rtc.getProgramLog(prog, log.data());
|
||||
out.log.assign(log.data(), logSize - 1);
|
||||
}
|
||||
}
|
||||
if (r != NVRTC_SUCCESS) {
|
||||
std::string one;
|
||||
for (char ch : out.log) { if (ch == '\n' || ch == '\r') { if (one.size() && one.back() != '|') one += " | "; } else one += ch; if (one.size() > 600) break; }
|
||||
err = std::string("nvrtcCompileProgram ") + name + " for " + c.archOpt + ": " + c.rtc.getErrorString(r) + ": " + one;
|
||||
c.rtc.destroyProgram(&prog);
|
||||
return false;
|
||||
}
|
||||
for (const std::string& e : nameExprs) {
|
||||
const char* lowered = nullptr;
|
||||
r = c.rtc.getLoweredName(prog, e.c_str(), &lowered);
|
||||
if (r != NVRTC_SUCCESS || !lowered) { err = "nvrtcGetLoweredName " + e + ": " + c.rtc.getErrorString(r); c.rtc.destroyProgram(&prog); return false; }
|
||||
out.lowered.push_back(lowered);
|
||||
}
|
||||
size_t n = 0;
|
||||
if (c.ptx) {
|
||||
r = c.rtc.getPTXSize(prog, &n);
|
||||
if (r == NVRTC_SUCCESS) { out.image.resize(n); r = c.rtc.getPTX(prog, out.image.data()); }
|
||||
} else {
|
||||
r = c.rtc.getCUBINSize(prog, &n);
|
||||
if (r == NVRTC_SUCCESS) { out.image.resize(n); r = c.rtc.getCUBIN(prog, out.image.data()); }
|
||||
}
|
||||
c.rtc.destroyProgram(&prog);
|
||||
if (r != NVRTC_SUCCESS || n == 0) { err = std::string(c.ptx ? "nvrtcGetPTX" : "nvrtcGetCUBIN") + ": " + c.rtc.getErrorString(r); return false; }
|
||||
out.ms = wallMs() - t0;
|
||||
return true;
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// A resident pair: one pack compiled, its cache and dataset on the device, self-tested
|
||||
|
||||
struct Pair {
|
||||
std::string dir, epochHex, dayHex, seedString;
|
||||
uint32_t sw[8] = {0}, kw[8] = {0};
|
||||
uint32_t datasetLog2 = 0, words = 0, cacheWords = 0, cacheSegments = 0;
|
||||
CUmodule modKernel = nullptr, modBound = nullptr;
|
||||
CUfunction fCacheFill = nullptr, fBuild = nullptr, fHashBound = nullptr;
|
||||
CUdeviceptr cache = 0, ds = 0;
|
||||
double compileMs = 0, cacheMs = 0, dsMs = 0, checkMs = 0;
|
||||
std::string check;
|
||||
bool checkPass = false, checked = false;
|
||||
int regs = 0, blocksPerSM = 0;
|
||||
};
|
||||
|
||||
static void releasePair(Ctx& c, Pair* p) {
|
||||
if (!p) return;
|
||||
if (p->ds) c.drv.memFree(p->ds);
|
||||
if (p->cache) c.drv.memFree(p->cache);
|
||||
if (p->modBound) c.drv.moduleUnload(p->modBound);
|
||||
if (p->modKernel) c.drv.moduleUnload(p->modKernel);
|
||||
delete p;
|
||||
}
|
||||
|
||||
struct IgneumInitWordsArg { uint32_t w[8]; };
|
||||
|
||||
// `block` threads per block (32 x warps); `nonces` must be a multiple of it.
|
||||
static bool launchHash(Ctx& c, Pair* p, CUdeviceptr out, uint32_t baseNonce, const uint32_t iw[8], uint32_t nonces, uint32_t block, CUstream s, std::string& err) {
|
||||
uint32_t mask = p->words - 1u;
|
||||
IgneumInitWordsArg a; std::memcpy(a.w, iw, 32);
|
||||
void* args[5] = { &p->ds, &out, &baseNonce, &mask, &a };
|
||||
DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, nonces / block, 1, 1, block, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound");
|
||||
return true;
|
||||
}
|
||||
|
||||
// Compiles the pack in `dir`, builds its cache and dataset on stream `s`, runs the self-test. Returns the pair or
|
||||
// null with `err` set. Runs on the main thread for --pack and --check, on the prepare thread for `prepare`.
|
||||
static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string& err) {
|
||||
PfPack pk;
|
||||
char perr[512];
|
||||
if (!pf_load(dir.c_str(), &pk, perr, sizeof(perr))) { err = std::string("pack ") + dir + ": " + perr; return nullptr; }
|
||||
bool ok1, ok2, ok3, ok4;
|
||||
std::string kernelCu = readText(dir + "/kernel.cu", ok1), boundCu = readText(dir + "/kernel_bound.cu", ok2);
|
||||
std::string programH = readText(dir + "/program.h", ok3), memhardH = readText(dir + "/memhard.h", ok4);
|
||||
if (!ok1 || !ok2 || !ok3 || !ok4) { err = "pack " + dir + " lacks kernel.cu, kernel_bound.cu, program.h or memhard.h"; return nullptr; }
|
||||
std::string kernelDev, boundDev;
|
||||
if (!deviceOnly(kernelCu, kernelDev, err) || !deviceOnly(boundCu, boundDev, err)) { err = "pack " + dir + ": " + err; return nullptr; }
|
||||
Pair* p = new Pair();
|
||||
p->dir = dir; p->epochHex = pk.epochHex; p->dayHex = pk.dayHex; p->seedString = pk.seedString;
|
||||
std::memcpy(p->sw, pk.seedw, 32); std::memcpy(p->kw, pk.keyw, 32);
|
||||
p->datasetLog2 = pk.datasetLog2; p->words = 1u << pk.datasetLog2; p->cacheWords = 1u << pk.cacheLog2Words; p->cacheSegments = pk.cacheSegments;
|
||||
// Compile
|
||||
Compiled ck, cb;
|
||||
if (!rtcCompile(c, kernelDev, "kernel.cu", programH, memhardH, { "igneum_cache_fill", "igneum_build" }, ck, err)) { releasePair(c, p); return nullptr; }
|
||||
if (!rtcCompile(c, boundDev, "kernel_bound.cu", programH, memhardH, { "igneum_hash_bound" }, cb, err)) { releasePair(c, p); return nullptr; }
|
||||
p->compileMs = ck.ms + cb.ms;
|
||||
// Load
|
||||
{
|
||||
CUresult r = c.drv.moduleLoadData(&p->modKernel, ck.image.data());
|
||||
if (r != CUDA_SUCCESS) { err = "cuModuleLoadData kernel.cu (" + c.archOpt + "): " + c.err(r); releasePair(c, p); return nullptr; }
|
||||
r = c.drv.moduleLoadData(&p->modBound, cb.image.data());
|
||||
if (r != CUDA_SUCCESS) { err = "cuModuleLoadData kernel_bound.cu (" + c.archOpt + "): " + c.err(r); releasePair(c, p); return nullptr; }
|
||||
if (c.drv.moduleGetFunction(&p->fCacheFill, p->modKernel, ck.lowered[0].c_str()) != CUDA_SUCCESS) { err = "igneum_cache_fill (" + ck.lowered[0] + ") not in the module"; releasePair(c, p); return nullptr; }
|
||||
if (c.drv.moduleGetFunction(&p->fBuild, p->modKernel, ck.lowered[1].c_str()) != CUDA_SUCCESS) { err = "igneum_build (" + ck.lowered[1] + ") not in the module"; releasePair(c, p); return nullptr; }
|
||||
if (c.drv.moduleGetFunction(&p->fHashBound, p->modBound, cb.lowered[0].c_str()) != CUDA_SUCCESS) { err = "igneum_hash_bound (" + cb.lowered[0] + ") not in the module"; releasePair(c, p); return nullptr; }
|
||||
c.drv.funcGetAttribute(&p->regs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->fHashBound);
|
||||
c.drv.occupancy(&p->blocksPerSM, p->fHashBound, 32 * c.blockWarps, 0);
|
||||
}
|
||||
// Cache
|
||||
double t0 = wallMs();
|
||||
size_t cacheBytes = (size_t)p->cacheWords * 4u, dsBytes = (size_t)p->words * 4u;
|
||||
{
|
||||
size_t freeB = 0, totalB = 0;
|
||||
if (c.drv.memGetInfo(&freeB, &totalB) == CUDA_SUCCESS && freeB < cacheBytes + dsBytes + (64u << 20)) {
|
||||
err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20));
|
||||
releasePair(c, p); return nullptr;
|
||||
}
|
||||
}
|
||||
{
|
||||
CUresult r = c.drv.memAlloc(&p->cache, cacheBytes);
|
||||
if (r != CUDA_SUCCESS) { err = "cuMemAlloc cache: " + c.err(r); p->cache = 0; releasePair(c, p); return nullptr; }
|
||||
uint32_t nSeg = p->cacheSegments, block = 256u, grid = (nSeg + block - 1u) / block;
|
||||
void* args[2] = { &p->cache, &nSeg };
|
||||
r = c.drv.launchKernel(p->fCacheFill, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
|
||||
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
|
||||
if (r != CUDA_SUCCESS) { err = "cache fill: " + c.err(r); releasePair(c, p); return nullptr; }
|
||||
}
|
||||
p->cacheMs = wallMs() - t0;
|
||||
// Dataset
|
||||
t0 = wallMs();
|
||||
{
|
||||
CUresult r = c.drv.memAlloc(&p->ds, dsBytes);
|
||||
if (r != CUDA_SUCCESS) { err = "cuMemAlloc dataset: " + c.err(r); p->ds = 0; releasePair(c, p); return nullptr; }
|
||||
uint32_t nItems = p->words / 16u, block = 256u, grid = (nItems + block - 1u) / block;
|
||||
void* args[3] = { &p->ds, &p->cache, &nItems };
|
||||
r = c.drv.launchKernel(p->fBuild, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
|
||||
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
|
||||
if (r != CUDA_SUCCESS) { err = "dataset build: " + c.err(r); releasePair(c, p); return nullptr; }
|
||||
}
|
||||
p->dsMs = wallMs() - t0;
|
||||
// Self-test against vectors.h
|
||||
t0 = wallMs();
|
||||
if (!pk.haveVectors) {
|
||||
p->checked = false; p->checkPass = true;
|
||||
p->check = "self-test skipped (no vectors.h in the pack); the miner's CPU re-check covers every found nonce";
|
||||
} else {
|
||||
uint32_t cacheHead[16], cacheLast[16], dsHead[16], dsLast = 0;
|
||||
std::vector<uint32_t> samples((size_t)(pk.nSamples > 0 ? pk.nSamples : 1), 0u);
|
||||
std::vector<uint64_t> vec((size_t)pk.vecWarps * 32u, 0ull);
|
||||
std::vector<uint32_t> whole(p->cacheWords);
|
||||
CUresult r = c.drv.memcpyDtoH(whole.data(), p->cache, cacheBytes);
|
||||
if (r != CUDA_SUCCESS) { err = "cuMemcpyDtoH cache: " + c.err(r); releasePair(c, p); return nullptr; }
|
||||
std::memcpy(cacheHead, whole.data(), 64);
|
||||
std::memcpy(cacheLast, whole.data() + p->cacheWords - 16u, 64);
|
||||
uint64_t fnv = pf_fnv1a64(whole.data(), cacheBytes);
|
||||
whole.clear(); whole.shrink_to_fit();
|
||||
if ((r = c.drv.memcpyDtoH(dsHead, p->ds, 64)) != CUDA_SUCCESS) { err = "cuMemcpyDtoH dataset head: " + c.err(r); releasePair(c, p); return nullptr; }
|
||||
if (pk.dsLastIndex < p->words) r = c.drv.memcpyDtoH(&dsLast, p->ds + (CUdeviceptr)pk.dsLastIndex * 4u, 4);
|
||||
for (int i = 0; i < pk.nSamples && r == CUDA_SUCCESS; ++i) if (pk.sampleIdx[i] < p->words) r = c.drv.memcpyDtoH(&samples[(size_t)i], p->ds + (CUdeviceptr)pk.sampleIdx[i] * 4u, 4);
|
||||
if (r != CUDA_SUCCESS) { err = "cuMemcpyDtoH dataset words: " + c.err(r); releasePair(c, p); return nullptr; }
|
||||
CUdeviceptr out = 0;
|
||||
if ((r = c.drv.memAlloc(&out, 32u * (size_t)c.blockWarps * 8u)) != CUDA_SUCCESS) { err = "cuMemAlloc vector out: " + c.err(r); releasePair(c, p); return nullptr; }
|
||||
for (int w = 0; w < pk.vecWarps; ++w) {
|
||||
// One block of 32 x block-warps lanes; the vector warp is its first 32 lanes (lane nonce = base + gid)
|
||||
if (!launchHash(c, p, out, pk.vecBase[w], p->sw, 32u * (uint32_t)c.blockWarps, 32u * (uint32_t)c.blockWarps, s, err)) { c.drv.memFree(out); releasePair(c, p); return nullptr; }
|
||||
if ((r = c.drv.streamSynchronize(s)) != CUDA_SUCCESS || (r = c.drv.memcpyDtoH(&vec[(size_t)w * 32u], out, 32u * 8u)) != CUDA_SUCCESS) { err = "vector warp: " + c.err(r); c.drv.memFree(out); releasePair(c, p); return nullptr; }
|
||||
}
|
||||
c.drv.memFree(out);
|
||||
char line[1024];
|
||||
p->checkPass = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples.data(), vec.data(), line, sizeof(line)) != 0;
|
||||
p->checked = true;
|
||||
p->check = line;
|
||||
}
|
||||
p->checkMs = wallMs() - t0;
|
||||
if (!p->checkPass) { err = p->check; releasePair(c, p); return nullptr; }
|
||||
return p;
|
||||
}
|
||||
|
||||
static std::string pairSummary(const Pair* p) {
|
||||
return fmt("nvrtc %.0f cache %.0f dataset %.0f check %.0f ms; %s", p->compileMs, p->cacheMs, p->dsMs, p->checkMs, p->check.c_str());
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// Prepare, on its own thread
|
||||
|
||||
struct PrepareTask {
|
||||
std::string epochHex, dayHex, dir, error;
|
||||
std::atomic<bool> done{false};
|
||||
Pair* result = nullptr;
|
||||
double t0 = 0;
|
||||
std::thread thread;
|
||||
};
|
||||
|
||||
static void prepareRun(Ctx* c, PrepareTask* t) {
|
||||
CUstream s = nullptr;
|
||||
std::string err;
|
||||
if (c->drv.ctxSetCurrent(c->ctx) != CUDA_SUCCESS) { t->error = "cuCtxSetCurrent on the prepare thread"; t->done = true; return; }
|
||||
if (c->drv.streamCreate(&s, CU_STREAM_NON_BLOCKING) != CUDA_SUCCESS) { t->error = "cuStreamCreate on the prepare thread"; t->done = true; return; }
|
||||
Pair* p = buildPair(*c, t->dir, s, err);
|
||||
c->drv.streamDestroy(s);
|
||||
if (p && (p->epochHex != t->epochHex || p->dayHex != t->dayHex)) {
|
||||
err = "the pack in " + t->dir + " is for epoch " + p->epochHex.substr(0, 16) + " day " + p->dayHex + ", not the prepared seeds";
|
||||
releasePair(*c, p); p = nullptr;
|
||||
}
|
||||
t->result = p;
|
||||
t->error = err;
|
||||
t->done = true;
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// Serve
|
||||
|
||||
struct Options {
|
||||
bool serve = false, check = false;
|
||||
int device = 0, batchLog2 = 22, blockWarps = 1;
|
||||
std::string pack, arch = "auto";
|
||||
};
|
||||
|
||||
static void usage() {
|
||||
std::printf("igneum-worker-cuda %s\n"
|
||||
" --serve --pack <dir> GPU worker for igneum-miner --worker: jobs on stdin, found/done lines on stdout\n"
|
||||
" --check --pack <dir> compile the pack, build its cache and dataset, self-test, print timings, exit 0 or 1\n"
|
||||
" --device D CUDA device index (default 0)\n"
|
||||
" --batch-log2 B nonces per dispatch = 2^B (default 22)\n"
|
||||
" --block-warps W warps per thread block (default 1)\n"
|
||||
" --arch sm_XY|compute_XY|auto NVRTC target (default auto: the device's architecture)\n", WORKER_VERSION);
|
||||
}
|
||||
|
||||
static Options parseArgs(int argc, char** argv) {
|
||||
Options o;
|
||||
for (int i = 1; i < argc; ++i) {
|
||||
std::string a = argv[i];
|
||||
auto next = [&]() -> std::string { if (i + 1 >= argc) { usage(); std::exit(2); } return argv[++i]; };
|
||||
if (a == "--serve") o.serve = true;
|
||||
else if (a == "--check") o.check = true;
|
||||
else if (a == "--pack") o.pack = next();
|
||||
else if (a == "--device") o.device = std::atoi(next().c_str());
|
||||
else if (a == "--batch-log2") o.batchLog2 = std::atoi(next().c_str());
|
||||
else if (a == "--block-warps") o.blockWarps = std::atoi(next().c_str());
|
||||
else if (a == "--arch") o.arch = next();
|
||||
else if (a == "--no-prepare") { /* accepted for symmetry with the other workers; prepare is always on here */ }
|
||||
else if (a == "-h" || a == "--help") { usage(); std::exit(0); }
|
||||
else { std::printf("unknown argument %s\n", argv[i]); usage(); std::exit(2); }
|
||||
}
|
||||
if (o.batchLog2 < 10 || o.batchLog2 > 28) { std::printf("--batch-log2 must be between 10 and 28\n"); std::exit(2); }
|
||||
if (o.blockWarps < 1 || o.blockWarps > 32) { std::printf("--block-warps must be between 1 and 32\n"); std::exit(2); }
|
||||
if (!o.serve && !o.check) { usage(); std::exit(2); }
|
||||
if (o.pack.empty()) { std::printf("--pack <dir> is required (igneum-miner export-pack <node> <dir> writes one)\n"); std::exit(2); }
|
||||
while (o.pack.size() > 1 && (o.pack.back() == '/' || o.pack.back() == '\\')) o.pack.pop_back();
|
||||
return o;
|
||||
}
|
||||
|
||||
static std::vector<std::string> split(const std::string& line) {
|
||||
std::vector<std::string> f;
|
||||
size_t i = 0;
|
||||
while (i < line.size()) {
|
||||
while (i < line.size() && (line[i] == ' ' || line[i] == '\t' || line[i] == '\r')) ++i;
|
||||
size_t j = i;
|
||||
while (j < line.size() && line[j] != ' ' && line[j] != '\t' && line[j] != '\r') ++j;
|
||||
if (j > i) f.push_back(line.substr(i, j - i));
|
||||
i = j;
|
||||
}
|
||||
return f;
|
||||
}
|
||||
|
||||
// A pack directory for the given seeds under `root` (one subdirectory per pack, each with seeds.txt), or "".
|
||||
static std::string findPackFor(const std::string& root, const std::string& epochHex, const std::string& dayHex) {
|
||||
if (root.empty()) return "";
|
||||
std::vector<std::string> names;
|
||||
#ifdef _WIN32
|
||||
WIN32_FIND_DATAA fd;
|
||||
HANDLE h = FindFirstFileA((root + "\\*").c_str(), &fd);
|
||||
if (h == INVALID_HANDLE_VALUE) return "";
|
||||
do { if ((fd.dwFileAttributes & FILE_ATTRIBUTE_DIRECTORY) && fd.cFileName[0] != '.') names.push_back(root + "\\" + fd.cFileName); } while (FindNextFileA(h, &fd));
|
||||
FindClose(h);
|
||||
#else
|
||||
DIR* d = opendir(root.c_str());
|
||||
if (!d) return "";
|
||||
while (dirent* e = readdir(d)) if (e->d_name[0] != '.') names.push_back(root + "/" + e->d_name);
|
||||
closedir(d);
|
||||
#endif
|
||||
for (const std::string& dir : names) {
|
||||
bool ok = false;
|
||||
std::string seeds = readText(dir + "/seeds.txt", ok);
|
||||
if (!ok) continue;
|
||||
char e[65] = {0}, dd[PF_HEX_CAP] = {0};
|
||||
if (pf_seeds_line(seeds.c_str(), "epoch_seed_hex", e, sizeof(e)) && pf_seeds_line(seeds.c_str(), "day_seed_hex", dd, sizeof(dd)) && epochHex == e && dayHex == dd) return dir;
|
||||
}
|
||||
return "";
|
||||
}
|
||||
|
||||
static std::string parentDir(const std::string& p) { size_t i = p.find_last_of("/\\"); return i == std::string::npos ? "." : p.substr(0, i); }
|
||||
|
||||
static int runServe(Ctx& c, const Options& o, Pair* cur) {
|
||||
const uint32_t batch = 1u << o.batchLog2;
|
||||
CUdeviceptr dOut = 0;
|
||||
std::string err;
|
||||
if (c.drv.memAlloc(&dOut, (size_t)batch * 8u) != CUDA_SUCCESS) { emit("error 0 cuMemAlloc out buffer"); return 2; }
|
||||
std::vector<uint64_t> hOut(batch);
|
||||
Pair* prepared = nullptr;
|
||||
Pair* old = nullptr;
|
||||
PrepareTask* task = nullptr;
|
||||
std::string prepareRoot; // the parent of the last prepare's pack directory: where the miner writes its packs
|
||||
emit(fmt("ready cuda %s pack %s dataset-log2 %u batch %u regs %d prepare 1 path nvrtc %d.%d driver %d.%d arch %s worker %s",
|
||||
c.name.c_str(), cur->seedString.c_str(), cur->datasetLog2, batch, cur->regs, c.rtcMajor, c.rtcMinor, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.archOpt.c_str(), WORKER_VERSION));
|
||||
info(fmt("first pack %s: %s", cur->dir.c_str(), pairSummary(cur).c_str()));
|
||||
std::string line;
|
||||
while (std::getline(std::cin, line)) {
|
||||
if (line == "quit") break;
|
||||
std::vector<std::string> f = split(line);
|
||||
if (f.empty()) continue;
|
||||
// A finished prepare is reported here, between lines
|
||||
if (task && task->done) {
|
||||
task->thread.join();
|
||||
if (task->result) {
|
||||
if (prepared) releasePair(c, prepared);
|
||||
prepared = task->result;
|
||||
emit(fmt("prepared %s %s %.1f %s resident 2 programs 2 datasets", prepared->epochHex.c_str(), prepared->dayHex.c_str(), wallMs() - task->t0, pairSummary(prepared).c_str()));
|
||||
} else {
|
||||
emit(fmt("prepare-failed %s %s %s", task->epochHex.c_str(), task->dayHex.c_str(), task->error.c_str()));
|
||||
}
|
||||
delete task; task = nullptr;
|
||||
}
|
||||
if (f[0] == "prepare") {
|
||||
if (f.size() < 4) { emit(fmt("prepare-failed %s %s a pack directory is needed as the third field (igneum-miner --prepare-packs <dir>)", f.size() > 1 ? f[1].c_str() : "0", f.size() > 2 ? f[2].c_str() : "0")); continue; }
|
||||
if (f[1].size() != 64) { emit(fmt("prepare-failed %s %s bad field (epoch_seed 64 hex, day_seed hex)", f[1].c_str(), f[2].c_str())); continue; }
|
||||
if (task) { emit(fmt("prepare-failed %s %s a prepare is still running", f[1].c_str(), f[2].c_str())); continue; }
|
||||
if (prepared && prepared->epochHex == f[1] && prepared->dayHex == f[2]) { emit(fmt("prepared %s %s 0 (already resident)", f[1].c_str(), f[2].c_str())); continue; }
|
||||
if (cur->epochHex == f[1] && cur->dayHex == f[2]) { emit(fmt("prepared %s %s 0 (already the current pair)", f[1].c_str(), f[2].c_str())); continue; }
|
||||
std::string dir = f[3];
|
||||
for (size_t i = 4; i < f.size(); ++i) dir += " " + f[i]; // a directory with spaces arrives as several fields
|
||||
prepareRoot = parentDir(dir);
|
||||
task = new PrepareTask();
|
||||
task->epochHex = f[1]; task->dayHex = f[2]; task->dir = dir; task->t0 = wallMs();
|
||||
task->thread = std::thread(prepareRun, &c, task);
|
||||
info(fmt("prepare started for epoch %.16s day %s from %s (NVRTC %s in the background)", f[1].c_str(), f[2].c_str(), dir.c_str(), c.archOpt.c_str()));
|
||||
continue;
|
||||
}
|
||||
if (f[0] != "job") { info("ignored: " + line); continue; }
|
||||
std::string jobId = f.size() > 1 ? f[1] : "0";
|
||||
if (f.size() < 8) { emit("error " + jobId + " malformed job line (need 7 fields after job)"); continue; }
|
||||
uint8_t prehash[32], epochSeed[32], daySeed[256];
|
||||
size_t pl = 0, el = 0, dl = 0;
|
||||
unsigned long long target = 0, nonceStart = 0, nonceCount = 0;
|
||||
if (!pf_unhex(f[2].c_str(), prehash, 32, &pl) || pl != 32 || std::sscanf(f[3].c_str(), "%llx", &target) != 1 ||
|
||||
std::sscanf(f[4].c_str(), "%llu", &nonceStart) != 1 || std::sscanf(f[5].c_str(), "%llu", &nonceCount) != 1 ||
|
||||
!pf_unhex(f[6].c_str(), epochSeed, 32, &el) || el != 32 || !pf_unhex(f[7].c_str(), daySeed, sizeof(daySeed), &dl)) {
|
||||
emit("error " + jobId + " bad field (prehash 64 hex, target 16 hex, nonce_start u64, nonce_count u64, epoch_seed 64 hex, day_seed hex)"); continue;
|
||||
}
|
||||
if (nonceCount == 0 || nonceCount % 32 != 0 || (nonceStart & 31) != 0) { emit("error " + jobId + " nonce_start must be 32-aligned and nonce_count a non-zero multiple of 32"); continue; }
|
||||
uint32_t sw[8], kw[8];
|
||||
pf_seed_words_from_bytes(epochSeed, 32, sw);
|
||||
pf_seed_words_from_bytes(daySeed, dl, kw);
|
||||
double t0 = wallMs();
|
||||
bool switched = false;
|
||||
if ((std::memcmp(sw, cur->sw, 32) != 0 || std::memcmp(kw, cur->kw, 32) != 0) && !task &&
|
||||
!(prepared && std::memcmp(sw, prepared->sw, 32) == 0 && std::memcmp(kw, prepared->kw, 32) == 0)) {
|
||||
// Self-heal: a job on seeds this worker has no pair for and no prepare in flight (a prepare failed, or
|
||||
// the miner never sent one). The miner writes a pack per pair under its --prepare-packs root; find it by
|
||||
// seeds.txt and build it now, in the foreground. The miner only re-sends prepare for the pair after this one.
|
||||
std::string dir = findPackFor(prepareRoot, f[6], f[7]);
|
||||
if (dir.empty()) dir = findPackFor(parentDir(o.pack) + "/prepare", f[6], f[7]);
|
||||
if (dir.empty()) dir = findPackFor(parentDir(o.pack), f[6], f[7]);
|
||||
if (!dir.empty()) {
|
||||
info(fmt("job %s is for epoch %.16s day %s, which is not resident; building its pack %s now (foreground)", jobId.c_str(), f[6].c_str(), f[7].c_str(), dir.c_str()));
|
||||
std::string berr;
|
||||
Pair* p = buildPair(c, dir, nullptr, berr);
|
||||
if (p && (p->epochHex != f[6] || p->dayHex != f[7])) { berr = "the pack in " + dir + " is for other seeds"; releasePair(c, p); p = nullptr; }
|
||||
if (p) { if (prepared) releasePair(c, prepared); prepared = p; info(fmt("built %s: %s", dir.c_str(), pairSummary(p).c_str())); }
|
||||
else emit("error " + jobId + " could not build " + dir + ": " + berr);
|
||||
}
|
||||
}
|
||||
if (std::memcmp(sw, cur->sw, 32) != 0 || std::memcmp(kw, cur->kw, 32) != 0) {
|
||||
if (prepared && std::memcmp(sw, prepared->sw, 32) == 0 && std::memcmp(kw, prepared->kw, 32) == 0) {
|
||||
if (old) releasePair(c, old);
|
||||
old = cur; cur = prepared; prepared = nullptr; switched = true;
|
||||
info(fmt("switched to the prepared pair epoch %.16s day %s in %.2f ms", cur->epochHex.c_str(), cur->dayHex.c_str(), wallMs() - t0));
|
||||
} else if (std::memcmp(sw, cur->sw, 32) != 0) {
|
||||
emit(fmt("error %s epoch seed mismatch: this worker holds epoch %.16s (seed words %08x %08x ...)%s, the job's epoch seed %.16s gives %08x %08x ...; send prepare with a pack directory",
|
||||
jobId.c_str(), cur->epochHex.c_str(), cur->sw[0], cur->sw[1], prepared ? " plus one prepared pair" : "", f[6].c_str(), sw[0], sw[1]));
|
||||
continue;
|
||||
} else {
|
||||
emit(fmt("error %s day seed mismatch: this worker's cache is for key %08x %08x ..., the job's day seed %s gives %08x %08x ...; send prepare with a pack directory",
|
||||
jobId.c_str(), cur->kw[0], cur->kw[1], f[7].c_str(), kw[0], kw[1]));
|
||||
continue;
|
||||
}
|
||||
}
|
||||
uint64_t remaining = nonceCount, hashes = 0;
|
||||
uint32_t hi = (uint32_t)(nonceStart >> 32), lo = (uint32_t)nonceStart;
|
||||
bool failed = false;
|
||||
while (remaining > 0) {
|
||||
uint64_t room = (uint64_t)(0xffffffffu - lo) + 1ull;
|
||||
uint64_t chunk64 = remaining < batch ? remaining : batch;
|
||||
if (chunk64 > room) chunk64 = room;
|
||||
uint32_t chunk = (uint32_t)chunk64;
|
||||
uint32_t iw[8];
|
||||
{
|
||||
uint8_t b[49];
|
||||
std::memcpy(b, "igneum-block/", 13);
|
||||
std::memcpy(b + 13, prehash, 32);
|
||||
b[45] = (uint8_t)hi; b[46] = (uint8_t)(hi >> 8); b[47] = (uint8_t)(hi >> 16); b[48] = (uint8_t)(hi >> 24);
|
||||
pf_seed_words_from_bytes(b, 49, iw);
|
||||
}
|
||||
// A chunk that is not a multiple of the block is finished one 32-lane block at a time
|
||||
uint32_t block = 32u * (uint32_t)o.blockWarps;
|
||||
uint32_t main = chunk - (chunk % block);
|
||||
CUresult r = CUDA_SUCCESS;
|
||||
if (main > 0 && !launchHash(c, cur, dOut, lo, iw, main, block, nullptr, err)) { emit("error " + jobId + " dispatch failed: " + err); failed = true; break; }
|
||||
if (main < chunk) {
|
||||
for (uint32_t off = main; off < chunk && !failed; off += 32u) if (!launchHash(c, cur, dOut + (CUdeviceptr)off * 8u, lo + off, iw, 32u, 32u, nullptr, err)) { emit("error " + jobId + " dispatch failed: " + err); failed = true; }
|
||||
if (failed) break;
|
||||
}
|
||||
r = c.drv.streamSynchronize(nullptr);
|
||||
if (r == CUDA_SUCCESS) r = c.drv.memcpyDtoH(hOut.data(), dOut, (size_t)chunk * 8u);
|
||||
if (r != CUDA_SUCCESS) { emit("error " + jobId + " dispatch failed: " + c.err(r)); failed = true; break; }
|
||||
for (uint32_t i = 0; i < chunk; ++i) if (hOut[i] <= target) {
|
||||
uint64_t nonce = ((uint64_t)hi << 32) | (uint64_t)(uint32_t)(lo + i);
|
||||
std::printf("found %s %llu %016llx\n", jobId.c_str(), (unsigned long long)nonce, (unsigned long long)hOut[i]);
|
||||
}
|
||||
std::fflush(stdout);
|
||||
hashes += chunk;
|
||||
remaining -= chunk;
|
||||
if (chunk64 == room) { hi += 1u; lo = 0u; } else lo += chunk;
|
||||
}
|
||||
if (failed) continue;
|
||||
emit(fmt("done %s %llu %.2f", jobId.c_str(), (unsigned long long)hashes, wallMs() - t0));
|
||||
if (switched && old) { releasePair(c, old); old = nullptr; info("dropped the previous pair (its program, cache and dataset)"); }
|
||||
}
|
||||
if (task) { task->thread.join(); if (task->result) releasePair(c, task->result); delete task; }
|
||||
c.drv.memFree(dOut);
|
||||
if (old) releasePair(c, old);
|
||||
if (prepared) releasePair(c, prepared);
|
||||
releasePair(c, cur);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
// Main
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
Options o = parseArgs(argc, argv);
|
||||
Ctx c;
|
||||
c.blockWarps = o.blockWarps;
|
||||
std::string err, drvLib, rtcLib;
|
||||
if (!loadDriver(c.drv, err, drvLib)) { emit("error 0 " + err); return 2; }
|
||||
if (!loadNvrtc(c.rtc, err, rtcLib)) { emit("error 0 " + err); return 2; }
|
||||
if (!openDevice(c, o.device, o.arch, err)) { emit("error 0 " + err); return 2; }
|
||||
info(fmt("igneum-worker-cuda %s: device %d %s (sm_%d%d, %d SMs), driver %d.%d from %s, NVRTC %d.%d from %s, target %s (%s)",
|
||||
WORKER_VERSION, o.device, c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, drvLib.c_str(), c.rtcMajor, c.rtcMinor, rtcLib.c_str(), c.archOpt.c_str(), c.why.c_str()));
|
||||
double t0 = wallMs();
|
||||
Pair* cur = buildPair(c, o.pack, nullptr, err);
|
||||
if (!cur) { emit("error 0 " + err); return 1; }
|
||||
if (o.check) {
|
||||
std::printf("check PASS %s in %.0f ms: %s\n", o.pack.c_str(), wallMs() - t0, pairSummary(cur).c_str());
|
||||
std::printf(" epoch %s day %s, dataset 2^%u words, cache 2^%u words in %u segments, %d registers, %d blocks/SM at %d warp(s)/block, target %s\n",
|
||||
cur->epochHex.c_str(), cur->dayHex.c_str(), cur->datasetLog2, (unsigned)__builtin_ctz(cur->cacheWords), cur->cacheSegments, cur->regs, cur->blocksPerSM, c.blockWarps, c.archOpt.c_str());
|
||||
releasePair(c, cur);
|
||||
return 0;
|
||||
}
|
||||
int rc = runServe(c, o, cur);
|
||||
c.drv.primaryCtxRelease(c.dev);
|
||||
return rc;
|
||||
}
|
||||
Loading…
Reference in a new issue