1884 lines
129 KiB
C++
1884 lines
129 KiB
C++
// igneum-worker-cuda: the one-click NVIDIA worker for igneum-miner --worker. 4 October 2026.
|
|
//
|
|
// Nothing to install but the NVIDIA driver. The pack's kernels (kernel.cu: cache fill and dataset build;
|
|
// kernel_bound.cu: the header-bound hash) are compiled at run time by NVRTC, the toolkit's runtime compiler, which
|
|
// ships next to this exe as nvrtc64_120_0.dll plus nvrtc-builtins64_128.dll (NVIDIA's redistributable, see
|
|
// THIRD-PARTY.md). The GPU is driven through the driver API in nvcuda.dll, which every NVIDIA driver installs. Both
|
|
// libraries are loaded with LoadLibrary/GetProcAddress (cuda_api.h), so no import library is linked and the exe is
|
|
// cross-compiled on the Mac with mingw (build-windows.sh). Plain C++17 otherwise.
|
|
//
|
|
// The source handed to NVRTC is the pack's own text: kernel.cu and kernel_bound.cu up to the host-side launch
|
|
// wrappers (which nvcc compiles for the host and NVRTC has no use for), with the pack's program.h and memhard.h as
|
|
// named headers, byte for byte. Two stub headers stand in for <cuda_runtime.h> and <cstdint>, which nvcc takes from
|
|
// the toolkit. emu/test.sh checks the equality on the Mac. The memory-hard core is therefore the same text host.cu
|
|
// compiles, and the miner's CPU re-check of every found nonce covers the rest.
|
|
//
|
|
// Worker protocol (the same lines as proto-cuda/host.cu --serve and proto-opencl/host.c --serve):
|
|
// stdin: job <job_id> <header_prehash_hex 64> <target_hex 16> <nonce_start u64> <nonce_count u64> <epoch_seed_hex 64> <day_seed_hex>
|
|
// prepare <epoch_seed_hex 64> <day_seed_hex> <pack_dir> compile that pack in the background, build its cache
|
|
// and dataset, self-test it; a job on it then switches
|
|
// quit
|
|
// stdout: ready cuda <device> pack <seed string> dataset-log2 N batch B regs R prepare 1 path nvrtc ...
|
|
// found <job_id> <nonce u64> <hash_hex 16>
|
|
// done <job_id> <hashes> <ms>
|
|
// error <job_id> <text>
|
|
// need <epoch_seed_hex> <day_seed_hex> before the mismatch error: the pair this worker lacks (the miner prepares it)
|
|
// prepared <epoch_seed_hex> <day_seed_hex> <ms> ... | prepare-failed <epoch_seed_hex> <day_seed_hex> <text>
|
|
// info ...
|
|
// The first pack comes from --pack <dir> (igneum-miner export-pack writes it; the launcher passes it). Every pack is
|
|
// self-tested before it serves a job: cache head, last line and FNV-1a 64, dataset head, last word and 64 samples,
|
|
// and the three vector warps of vectors.h through the bound kernel with the pack's own seed words. A pack that fails
|
|
// is refused.
|
|
//
|
|
// Variant racing (4 October 2026, evening; docs/design/miner-tuning.md): every pack's bound kernel is compiled in
|
|
// several variants (loop unrolling, the dataset load path: plain, __ldg, __ldcg, __ldcs; a register budget through
|
|
// -maxrregcount or __launch_bounds__; threads per block), each self-tested against the pack's vectors (bit-exact or
|
|
// discarded) and run for about two seconds on the card; the fastest serves the hour. A race runs inside the prepare
|
|
// (the hourly compile-ahead, one lead before the boundary) and never delays the swap: it has a time budget, "base"
|
|
// (the pack's text as shipped) is always the first entry, and a prepare that runs out of budget keeps the best so far.
|
|
// While a variant is timed the job loop pauses (one mutex): the numbers are exclusive, mining resumes between
|
|
// variants. One line per race: `race <epoch16> device <name> ... variants N a=MH/s b=MH/s ... winner <name> <MH/s>
|
|
// gain <pct> ...`. A tuning file (--tuning, or IGNEUM_TUNING_FILE from the app) may pin a variant for this card
|
|
// model or order the candidates; the app's over-the-air manifest carries it (fleet learning). Under IGNEUM_EMU the
|
|
// race is off (the stand-in checks that the handed-over text equals the pack's).
|
|
//
|
|
// Usage: igneum-worker-cuda --serve --pack <dir> [--device D] [--batch-log2 22] [--block-warps 1] [--arch sm_120|auto]
|
|
// [--race on|off|<name,name,...>] [--race-bench-ms 2000] [--race-budget-s 120] [--race-rounds 1]
|
|
// [--variant <name>] [--tuning <file>]
|
|
// igneum-worker-cuda --check --pack <dir> [--device D] compile, build, self-test, print timings, exit 0/1
|
|
// igneum-worker-cuda --race --pack <dir> [--device D] [--race-rounds 3] the race alone: one line per variant, exit 0/1
|
|
|
|
#include <cstdint>
|
|
#include <cstdarg>
|
|
#include <cstdio>
|
|
#include <cstdlib>
|
|
#include <cstring>
|
|
#include <ctime>
|
|
#include <chrono>
|
|
#include <string>
|
|
#include <vector>
|
|
#include <thread>
|
|
#include <atomic>
|
|
#include <mutex>
|
|
#include <algorithm>
|
|
#include <iostream>
|
|
|
|
#include "cuda_api.h"
|
|
#include "packfile.h"
|
|
|
|
#ifdef _WIN32
|
|
#define WIN32_LEAN_AND_MEAN
|
|
#include <windows.h>
|
|
#else
|
|
#include <dlfcn.h>
|
|
#include <dirent.h>
|
|
#endif
|
|
|
|
static const char* WORKER_VERSION = "1.0 (4 October 2026)";
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Helpers
|
|
|
|
static double wallMs() {
|
|
using namespace std::chrono;
|
|
return duration<double, std::milli>(steady_clock::now().time_since_epoch()).count();
|
|
}
|
|
|
|
static void emit(const std::string& s) { std::fputs(s.c_str(), stdout); std::fputc('\n', stdout); std::fflush(stdout); }
|
|
static void info(const std::string& s) { emit("info " + s); }
|
|
|
|
static std::string fmt(const char* f, ...) {
|
|
char buf[2048];
|
|
va_list ap;
|
|
va_start(ap, f);
|
|
vsnprintf(buf, sizeof(buf), f, ap);
|
|
va_end(ap);
|
|
return buf;
|
|
}
|
|
|
|
static std::string readText(const std::string& path, bool& ok) {
|
|
size_t n = 0;
|
|
char* b = pf_read_file(path.c_str(), &n);
|
|
if (!b) { ok = false; return ""; }
|
|
std::string s(b, n);
|
|
free(b);
|
|
ok = true;
|
|
return s;
|
|
}
|
|
|
|
static std::string exeDir() {
|
|
#ifdef _WIN32
|
|
char buf[MAX_PATH];
|
|
DWORD n = GetModuleFileNameA(nullptr, buf, MAX_PATH);
|
|
std::string p(buf, n);
|
|
size_t i = p.find_last_of("\\/");
|
|
return i == std::string::npos ? "." : p.substr(0, i);
|
|
#else
|
|
return ".";
|
|
#endif
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Loading the two libraries
|
|
|
|
static void* libOpen(const std::string& name) {
|
|
#ifdef _WIN32
|
|
return (void*)LoadLibraryA(name.c_str());
|
|
#else
|
|
return dlopen(name.c_str(), RTLD_NOW);
|
|
#endif
|
|
}
|
|
static void* libSym(void* lib, const char* name) {
|
|
#ifdef _WIN32
|
|
return (void*)GetProcAddress((HMODULE)lib, name);
|
|
#else
|
|
return dlsym(lib, name);
|
|
#endif
|
|
}
|
|
|
|
#define LOAD_SYM(table, field, name) do { table.field = (decltype(table.field))libSym(lib, name); if (!table.field) { missing += std::string(missing.empty() ? "" : ", ") + name; } } while (0)
|
|
|
|
static bool loadDriver(Drv& d, std::string& err, std::string& libName) {
|
|
#ifdef IGNEUM_EMU
|
|
emu_fill_driver(d); libName = "emulation (host threads, no GPU)"; (void)err; return true;
|
|
#else
|
|
#ifdef _WIN32
|
|
const char* names[] = { "nvcuda.dll" };
|
|
#else
|
|
const char* names[] = { "libcuda.so.1", "libcuda.so" };
|
|
#endif
|
|
void* lib = nullptr;
|
|
for (const char* n : names) { lib = libOpen(n); if (lib) { libName = n; break; } }
|
|
if (!lib) { err = "the CUDA driver library (nvcuda.dll) is not installed: install or update the NVIDIA driver"; return false; }
|
|
std::string missing;
|
|
LOAD_SYM(d, init, "cuInit");
|
|
LOAD_SYM(d, driverGetVersion, "cuDriverGetVersion");
|
|
LOAD_SYM(d, deviceGetCount, "cuDeviceGetCount");
|
|
LOAD_SYM(d, deviceGet, "cuDeviceGet");
|
|
LOAD_SYM(d, deviceGetName, "cuDeviceGetName");
|
|
LOAD_SYM(d, deviceGetAttribute, "cuDeviceGetAttribute");
|
|
LOAD_SYM(d, deviceTotalMem, "cuDeviceTotalMem_v2");
|
|
LOAD_SYM(d, primaryCtxSetFlags, "cuDevicePrimaryCtxSetFlags_v2");
|
|
LOAD_SYM(d, primaryCtxRetain, "cuDevicePrimaryCtxRetain");
|
|
LOAD_SYM(d, primaryCtxRelease, "cuDevicePrimaryCtxRelease_v2");
|
|
LOAD_SYM(d, ctxSetCurrent, "cuCtxSetCurrent");
|
|
LOAD_SYM(d, ctxSynchronize, "cuCtxSynchronize");
|
|
LOAD_SYM(d, memGetInfo, "cuMemGetInfo_v2");
|
|
LOAD_SYM(d, memAlloc, "cuMemAlloc_v2");
|
|
LOAD_SYM(d, memFree, "cuMemFree_v2");
|
|
LOAD_SYM(d, memcpyDtoH, "cuMemcpyDtoH_v2");
|
|
LOAD_SYM(d, moduleLoadData, "cuModuleLoadData");
|
|
LOAD_SYM(d, moduleUnload, "cuModuleUnload");
|
|
LOAD_SYM(d, moduleGetFunction, "cuModuleGetFunction");
|
|
LOAD_SYM(d, launchKernel, "cuLaunchKernel");
|
|
LOAD_SYM(d, streamCreate, "cuStreamCreate");
|
|
LOAD_SYM(d, streamSynchronize, "cuStreamSynchronize");
|
|
LOAD_SYM(d, streamDestroy, "cuStreamDestroy_v2");
|
|
LOAD_SYM(d, funcGetAttribute, "cuFuncGetAttribute");
|
|
LOAD_SYM(d, occupancy, "cuOccupancyMaxActiveBlocksPerMultiprocessor");
|
|
LOAD_SYM(d, getErrorString, "cuGetErrorString");
|
|
LOAD_SYM(d, getErrorName, "cuGetErrorName");
|
|
d.texObjectCreate = (decltype(d.texObjectCreate))libSym(lib, "cuTexObjectCreate"); // optional, --microbench only
|
|
d.texObjectDestroy = (decltype(d.texObjectDestroy))libSym(lib, "cuTexObjectDestroy");
|
|
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("libnvrtc.so.12"); v.push_back("/usr/local/cuda-12.8/lib64/libnvrtc.so.12"); v.push_back("/usr/local/cuda/lib64/libnvrtc.so.12"); // the box's card-free --list-race check
|
|
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;
|
|
// read-width experiment, variant 5 (5 October 2026): persistent warps and their scratch; --warps caps the launch
|
|
int warps = 0; // 0 = the resident capacity from the occupancy query, rounded down to a power of two
|
|
uint32_t salt = 1; // the running per-unit tag salt (+= units per launch)
|
|
int batches = 5; // --bench: timed dispatches
|
|
CUdevice dev = 0;
|
|
CUcontext ctx = nullptr;
|
|
std::string name;
|
|
int major = 0, minor = 0, sms = 0, driverVersion = 0, rtcMajor = 0, rtcMinor = 0;
|
|
std::vector<int> rtcArchs; // what NVRTC can target (empty when the query is unavailable)
|
|
std::string archOpt; // "sm_120" or "compute_120": what the packs are compiled for
|
|
bool ptx = false; // true when archOpt is compute_XY (PTX, driver JIT)
|
|
std::string why; // how archOpt was chosen
|
|
int blockWarps = 1;
|
|
// variant racing (see the header): which variants, how long each is timed, the budget of a race, rounds
|
|
std::string race = "on"; // on | off | comma list of variant names
|
|
int raceBenchMs = 2000, raceBudgetS = 120, raceRounds = 1, batchLog2 = 22;
|
|
std::string pinned; // --variant: use this variant, no race
|
|
std::string tuning; // the tuning file's text ("" = none)
|
|
std::string err(CUresult r) { const char* s = nullptr; if (drv.getErrorString) drv.getErrorString(r, &s); return s ? s : "CUDA driver error"; }
|
|
};
|
|
|
|
#define DRV_CHECK(c, call, what) do { CUresult r_ = (call); if (r_ != CUDA_SUCCESS) { err = std::string(what) + ": " + (c).err(r_); return false; } } while (0)
|
|
|
|
static bool openDevice(Ctx& c, int device, const std::string& archArg, std::string& err) {
|
|
DRV_CHECK(c, c.drv.init(0), "cuInit");
|
|
int count = 0;
|
|
DRV_CHECK(c, c.drv.deviceGetCount(&count), "cuDeviceGetCount");
|
|
if (count == 0) { err = "no CUDA device"; return false; }
|
|
if (device < 0 || device >= count) { err = fmt("device %d out of range (%d devices)", device, count); return false; }
|
|
DRV_CHECK(c, c.drv.deviceGet(&c.dev, device), "cuDeviceGet");
|
|
char name[256] = {0};
|
|
DRV_CHECK(c, c.drv.deviceGetName(name, 255, c.dev), "cuDeviceGetName");
|
|
c.name = name;
|
|
for (char& ch : c.name) if (ch == ' ') ch = '_';
|
|
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, c.dev), "compute capability major");
|
|
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, c.dev), "compute capability minor");
|
|
DRV_CHECK(c, c.drv.deviceGetAttribute(&c.sms, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, c.dev), "multiprocessor count");
|
|
c.drv.driverGetVersion(&c.driverVersion);
|
|
// Blocking sync, set before the context exists: the host thread sleeps in cuStreamSynchronize instead of spinning
|
|
// (one full core per worker at the default spin schedule, measured on the RTX 5090 with eight workers, 3 Oct 2026).
|
|
c.drv.primaryCtxSetFlags(c.dev, CU_CTX_SCHED_BLOCKING_SYNC);
|
|
DRV_CHECK(c, c.drv.primaryCtxRetain(&c.ctx, c.dev), "cuDevicePrimaryCtxRetain");
|
|
DRV_CHECK(c, c.drv.ctxSetCurrent(c.ctx), "cuCtxSetCurrent");
|
|
c.rtc.version(&c.rtcMajor, &c.rtcMinor);
|
|
if (c.rtc.getNumSupportedArchs && c.rtc.getSupportedArchs) {
|
|
int n = 0;
|
|
if (c.rtc.getNumSupportedArchs(&n) == NVRTC_SUCCESS && n > 0 && n < 256) { c.rtcArchs.assign((size_t)n, 0); if (c.rtc.getSupportedArchs(c.rtcArchs.data()) != NVRTC_SUCCESS) c.rtcArchs.clear(); }
|
|
}
|
|
// The target: the device's own SASS (sm_XY) when this NVRTC knows the architecture, else PTX for the newest
|
|
// architecture it knows below the device's, which the driver JIT-compiles forward. --arch overrides.
|
|
int cc = c.major * 10 + c.minor;
|
|
if (archArg != "auto" && !archArg.empty()) {
|
|
c.archOpt = archArg; c.ptx = archArg.rfind("compute_", 0) == 0; c.why = "--arch";
|
|
} else if (c.rtcArchs.empty()) {
|
|
c.archOpt = fmt("sm_%d", cc); c.why = "the device's architecture (NVRTC did not list its targets)";
|
|
} else {
|
|
bool known = false; int best = 0;
|
|
for (int a : c.rtcArchs) { if (a == cc) known = true; if (a <= cc && a > best) best = a; }
|
|
if (known) { c.archOpt = fmt("sm_%d", cc); c.why = "the device's architecture, listed by NVRTC"; }
|
|
else if (best > 0) { c.archOpt = fmt("compute_%d", best); c.ptx = true; c.why = fmt("this NVRTC does not know sm_%d; PTX for compute_%d, JIT-compiled by the driver", cc, best); }
|
|
else { c.archOpt = fmt("compute_%d", c.rtcArchs.front()); c.ptx = true; c.why = fmt("this NVRTC knows nothing at or below sm_%d; PTX for its oldest target", cc); }
|
|
}
|
|
return true;
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// NVRTC: compile one of the pack's kernel files
|
|
|
|
// The pack's kernel files end in host-side launch wrappers (cudaError_t igneum_launch_* with <<< >>> launches) that
|
|
// nvcc compiles for the host. NVRTC compiles device code only, so the text is cut there. The cut is checked: nothing
|
|
// device-side may follow it.
|
|
static bool deviceOnly(const std::string& text, std::string& out, std::string& err) {
|
|
size_t cut = text.find("\n// Host-side launch wrappers");
|
|
if (cut == std::string::npos) cut = text.find("\ncudaError_t ");
|
|
if (cut == std::string::npos) { out = text; return true; }
|
|
std::string tail = text.substr(cut + 1);
|
|
if (tail.find("__global__") != std::string::npos || tail.find("__device__") != std::string::npos) { err = "device code after the host launch wrappers; the pack layout is not the one this worker knows"; return false; }
|
|
out = text.substr(0, cut + 1);
|
|
return true;
|
|
}
|
|
|
|
static const char* STUB_CUDA_RUNTIME =
|
|
"// igneum-worker-cuda: stand-in for <cuda_runtime.h> under NVRTC, which has the device built-ins already\n"
|
|
"#pragma once\n"
|
|
"#ifndef __CUDACC_RTC__\n#error \"this stub is for NVRTC only\"\n#endif\n"
|
|
"#ifdef __SIZE_TYPE__\ntypedef __SIZE_TYPE__ size_t;\n#elif defined(__LP64__) || defined(_LP64)\ntypedef unsigned long size_t;\n#else\ntypedef unsigned long long size_t;\n#endif\n";
|
|
static const char* STUB_CSTDINT =
|
|
"// igneum-worker-cuda: stand-in for <cstdint> under NVRTC (the fixed-width types the packs use)\n"
|
|
"#pragma once\n"
|
|
"typedef signed char int8_t; typedef unsigned char uint8_t; typedef short int16_t; typedef unsigned short uint16_t;\n"
|
|
"typedef int int32_t; typedef unsigned int uint32_t;\n"
|
|
"#if defined(__LP64__) || defined(_LP64)\ntypedef long int64_t; typedef unsigned long uint64_t;\n"
|
|
"#else\ntypedef long long int64_t; typedef unsigned long long uint64_t;\n#endif\n";
|
|
|
|
struct Compiled {
|
|
std::vector<char> image;
|
|
std::vector<std::string> lowered;
|
|
double ms = 0;
|
|
std::string log;
|
|
};
|
|
|
|
static bool rtcCompile(Ctx& c, const std::string& src, const char* name, const std::string& programH, const std::string& memhardH,
|
|
const std::vector<std::string>& nameExprs, Compiled& out, std::string& err, const std::vector<std::string>& extraOpts = {}) {
|
|
double t0 = wallMs();
|
|
const char* headers[4] = { STUB_CUDA_RUNTIME, STUB_CSTDINT, programH.c_str(), memhardH.c_str() };
|
|
const char* names[4] = { "cuda_runtime.h", "cstdint", "program.h", "memhard.h" };
|
|
nvrtcProgram prog = nullptr;
|
|
nvrtcResult r = c.rtc.createProgram(&prog, src.c_str(), name, 4, headers, names);
|
|
if (r != NVRTC_SUCCESS) { err = std::string("nvrtcCreateProgram: ") + c.rtc.getErrorString(r); return false; }
|
|
for (const std::string& e : nameExprs) {
|
|
r = c.rtc.addNameExpression(prog, e.c_str());
|
|
if (r != NVRTC_SUCCESS) { err = "nvrtcAddNameExpression " + e + ": " + c.rtc.getErrorString(r); c.rtc.destroyProgram(&prog); return false; }
|
|
}
|
|
std::string archOpt = "--gpu-architecture=" + c.archOpt;
|
|
// -default-device: NVRTC rejects unannotated functions as host code (nvcc treats them as host and discards them);
|
|
// the pack headers (program.h, memhard.h) carry plain inline helpers, so every unannotated function is device code here.
|
|
std::vector<const char*> opts = { archOpt.c_str(), "--std=c++17", "-default-device" };
|
|
for (const std::string& o : extraOpts) opts.push_back(o.c_str());
|
|
r = c.rtc.compileProgram(prog, (int)opts.size(), opts.data());
|
|
{
|
|
size_t logSize = 0;
|
|
if (c.rtc.getProgramLogSize(prog, &logSize) == NVRTC_SUCCESS && logSize > 1) {
|
|
std::vector<char> log(logSize);
|
|
c.rtc.getProgramLog(prog, log.data());
|
|
out.log.assign(log.data(), logSize - 1);
|
|
}
|
|
}
|
|
if (r != NVRTC_SUCCESS) {
|
|
std::string one;
|
|
for (char ch : out.log) { if (ch == '\n' || ch == '\r') { if (one.size() && one.back() != '|') one += " | "; } else one += ch; if (one.size() > 600) break; }
|
|
err = std::string("nvrtcCompileProgram ") + name + " for " + c.archOpt + ": " + c.rtc.getErrorString(r) + ": " + one;
|
|
c.rtc.destroyProgram(&prog);
|
|
return false;
|
|
}
|
|
for (const std::string& e : nameExprs) {
|
|
const char* lowered = nullptr;
|
|
r = c.rtc.getLoweredName(prog, e.c_str(), &lowered);
|
|
if (r != NVRTC_SUCCESS || !lowered) { err = "nvrtcGetLoweredName " + e + ": " + c.rtc.getErrorString(r); c.rtc.destroyProgram(&prog); return false; }
|
|
out.lowered.push_back(lowered);
|
|
}
|
|
size_t n = 0;
|
|
if (c.ptx) {
|
|
r = c.rtc.getPTXSize(prog, &n);
|
|
if (r == NVRTC_SUCCESS) { out.image.resize(n); r = c.rtc.getPTX(prog, out.image.data()); }
|
|
} else {
|
|
r = c.rtc.getCUBINSize(prog, &n);
|
|
if (r == NVRTC_SUCCESS) { out.image.resize(n); r = c.rtc.getCUBIN(prog, out.image.data()); }
|
|
}
|
|
c.rtc.destroyProgram(&prog);
|
|
if (r != NVRTC_SUCCESS || n == 0) { err = std::string(c.ptx ? "nvrtcGetPTX" : "nvrtcGetCUBIN") + ": " + c.rtc.getErrorString(r); return false; }
|
|
out.ms = wallMs() - t0;
|
|
return true;
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// A resident pair: one pack compiled, its cache and dataset on the device, self-tested
|
|
|
|
static bool hexEq(const std::string& a, const std::string& b) {
|
|
if (a.size() != b.size()) return false;
|
|
for (size_t i = 0; i < a.size(); ++i) if (std::tolower((unsigned char)a[i]) != std::tolower((unsigned char)b[i])) return false;
|
|
return true;
|
|
}
|
|
|
|
struct Pair;
|
|
// A pair is the pair of a job when the job's seeds (the hex the node sent) are the pair's seeds. The derived seed
|
|
// words are no identity: a retried program's words are its attempt's words, not the bare seed's (packfile.h,
|
|
// 5 October 2026), so comparing words refused every job of a retried program.
|
|
static bool pairIs(const Pair* p, const std::string& epochHex, const std::string& dayHex);
|
|
|
|
struct Pair {
|
|
std::string dir, epochHex, dayHex, seedString;
|
|
uint32_t sw[8] = {0}, kw[8] = {0};
|
|
uint32_t datasetLog2 = 0, words = 0, cacheWords = 0, cacheSegments = 0;
|
|
CUmodule modKernel = nullptr, modBound = nullptr;
|
|
CUfunction fCacheFill = nullptr, fBuild = nullptr, fHashBound = nullptr;
|
|
CUdeviceptr cache = 0, ds = 0;
|
|
double compileMs = 0, cacheMs = 0, dsMs = 0, checkMs = 0;
|
|
std::string check;
|
|
bool checkPass = false, checked = false;
|
|
int regs = 0, blocksPerSM = 0;
|
|
int blockWarps = 1; // threads per block = 32 x this (the winning variant's, else the worker's default)
|
|
std::string variant = "base"; // the bound kernel in service: a variant name (see allVariants)
|
|
int sparseBlocks = 0; // the SM-sparse variants (sp<N>): the persistent grid in blocks, 0 = the plain grid
|
|
std::string raceLine; // the race's one-line report, emitted by the main thread with "prepared"
|
|
double raceMs = 0;
|
|
// read-width experiment (5 October 2026): the pack's load class and, for variant 5, the persistent-warp scratch
|
|
std::string loadClass = "v2";
|
|
std::string programClass = "v2", eraHex; // Counter ASIC 2.0: the pack's class and era seed (packfile.h)
|
|
uint32_t loadsPerHash = 128, bytesPerHash = 512, scratchOps = 0;
|
|
bool persistent = false;
|
|
CUdeviceptr scratch = 0;
|
|
int warps = 0; // persistent warps launched (the arena holds this many)
|
|
int residentWarps = 0; // the occupancy query's capacity: blocks/SM x warps/block x SMs
|
|
size_t scratchBytes = 0;
|
|
// hot-table experiment (5 October 2026, docs/plans/hot-table.md): the epoch's hot table, filled on the device by the
|
|
// pack's igneum_hot_fill, the argument after the init words
|
|
uint32_t hotMb = 0, hotWords = 0, hotSegments = 0, hotSlots = 0;
|
|
CUfunction fHotFill = nullptr;
|
|
CUdeviceptr hot = 0;
|
|
double hotMs = 0;
|
|
};
|
|
|
|
static bool pairIs(const Pair* p, const std::string& epochHex, const std::string& dayHex) {
|
|
return p && hexEq(p->epochHex, epochHex) && hexEq(p->dayHex, dayHex);
|
|
}
|
|
|
|
// Counter ASIC 2.0: a job that names a class (and an era) belongs to a pair of that class (and era) only, so a pack
|
|
// of the old class for the same seeds is not this job's pair and the prepared pack of the right class wins.
|
|
static bool pairIsClass(const Pair* p, const std::string& epochHex, const std::string& dayHex, const std::string& cls, const std::string& era) {
|
|
char why[256];
|
|
return pairIs(p, epochHex, dayHex) && pf_pack_class_ok(p->programClass.c_str(), p->eraHex.c_str(), cls.c_str(), era.c_str(), why, sizeof(why));
|
|
}
|
|
|
|
// The trailing `class=` and `era=` tokens of a job or prepare line (absent on every class v2 line), removed from `f`.
|
|
static void takeClassTokens(std::vector<std::string>& f, std::string& cls, std::string& era) {
|
|
while (!f.empty()) {
|
|
char c[8] = {0}, e[65] = {0};
|
|
if (!pf_class_token(f.back().c_str(), c, sizeof(c), e, sizeof(e))) break;
|
|
if (c[0]) cls = c;
|
|
if (e[0]) era = e;
|
|
f.pop_back();
|
|
}
|
|
}
|
|
|
|
// The job loop and a race take turns on the card: a variant is timed with no job running (exclusive numbers), and
|
|
// mining resumes between variants. Held per chunk by the job loop, per variant by the race.
|
|
static std::mutex gpuMutex;
|
|
|
|
static void releasePair(Ctx& c, Pair* p) {
|
|
if (!p) return;
|
|
if (p->ds) c.drv.memFree(p->ds);
|
|
if (p->cache) c.drv.memFree(p->cache);
|
|
if (p->scratch) c.drv.memFree(p->scratch);
|
|
if (p->hot) c.drv.memFree(p->hot);
|
|
if (p->modBound) c.drv.moduleUnload(p->modBound);
|
|
if (p->modKernel) c.drv.moduleUnload(p->modKernel);
|
|
delete p;
|
|
}
|
|
|
|
struct IgneumInitWordsArg { uint32_t w[8]; };
|
|
|
|
// The launch shape of a non-persistent dispatch: grid blocks for `nonces` lanes at `block` threads per block, or the
|
|
// SM-sparse grid of `sparseBlocks` persistent blocks (Counter ASIC 4.0 research). Pure, so the emulation test can
|
|
// read it (emu/variant-test.cpp): the base shape is nonces / block, the sparse shape is the block count itself.
|
|
static unsigned launchGridBlocks(int sparseBlocks, uint32_t nonces, uint32_t block) {
|
|
return sparseBlocks > 0 ? (unsigned)sparseBlocks : (unsigned)(nonces / block);
|
|
}
|
|
|
|
// Whether a --bench/--memprobe/--microbench run keeps the race off: yes unless --variant names a kernel to serve
|
|
// (8 October 2026: the SM-sparse run of 7 October served base on 48 rows because this read true with a --variant).
|
|
static bool benchRaceOff(bool bench, const std::string& pinned) { return bench && pinned.empty(); }
|
|
|
|
// `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);
|
|
if (p->persistent) {
|
|
// Variant 5: N persistent warps over nonces / 32 units; the arena was sized for p->warps warps in buildPair.
|
|
uint32_t units = nonces / 32u, warps = (uint32_t)p->warps;
|
|
if (warps > units) warps = units;
|
|
while (warps > 1u && units % warps != 0u) warps >>= 1;
|
|
if (block != 32u) { err = "a variant-5 pack runs one warp per block (--block-warps 1)"; return false; }
|
|
uint32_t salt = c.salt; c.salt += units;
|
|
// the hot table (when the pack has one) sits between the init words and the scratch triple
|
|
void* args[9] = { &p->ds, &out, &baseNonce, &mask, &a, &p->scratch, &units, &salt, nullptr };
|
|
if (p->hot) { args[5] = &p->hot; args[6] = &p->scratch; args[7] = &units; args[8] = &salt; }
|
|
DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, warps, 1, 1, 32, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound (persistent)");
|
|
return true;
|
|
}
|
|
if (p->sparseBlocks > 0) {
|
|
// The SM-sparse wrapper: a grid of sparseBlocks blocks, the trailing argument the dispatch's nonce count (after
|
|
// the hot table when the pack has one); a dispatch smaller than the grid leaves the surplus blocks idle.
|
|
if (nonces % block != 0u) { err = "nonces not a multiple of the block"; return false; }
|
|
void* args[7] = { &p->ds, &out, &baseNonce, &mask, &a, &nonces, nullptr };
|
|
if (p->hot) { args[5] = &p->hot; args[6] = &nonces; }
|
|
DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, launchGridBlocks(p->sparseBlocks, nonces, block), 1, 1, block, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound (SM-sparse)");
|
|
return true;
|
|
}
|
|
void* args[6] = { &p->ds, &out, &baseNonce, &mask, &a, &p->hot };
|
|
DRV_CHECK(c, c.drv.launchKernel(p->fHashBound, launchGridBlocks(0, nonces, block), 1, 1, block, 1, 1, 0, s, args, nullptr), "cuLaunchKernel igneum_hash_bound");
|
|
return true;
|
|
}
|
|
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Variant racing
|
|
|
|
struct Variant {
|
|
std::string name;
|
|
int unroll = 0; // 0: the iteration loop as emitted; N: "#pragma unroll N" before it (8 = fully unrolled)
|
|
int load = 0; // 0: plain ds[i]; 1: __ldg (read-only data path); 2: __ldcg (L2 only, no L1); 3: __ldcs (streaming)
|
|
int maxrreg = 0; // 0: none; N: --maxrregcount=N (registers per thread, occupancy against spills)
|
|
int blockWarps = 0; // 0: the worker's --block-warps; N: 32 x N threads per block
|
|
int minBlocks = 0; // N > 0: __launch_bounds__(32 x blockWarps, N) (the compiler fits N blocks per SM)
|
|
int sparseBlocks = 0; // N > 0: the SM-sparse shape (Counter ASIC 4.0 research, 7 October 2026): a persistent grid of N
|
|
// blocks, every thread looping over the dispatch's nonces with stride gridDim.x * blockDim.x, so the
|
|
// hash runs on about N SMs at 32 warps per block and the other SMs idle; the kernel gains a
|
|
// trailing `uint32_t nonces` argument. Opt-in by name only (sp<N> or sp<N>-w<W>), never in a race by default
|
|
};
|
|
|
|
// The catalogue. Names are stable: the tuning file and the fleet records use them. "base" is the pack's text as
|
|
// shipped with the worker's default block and is always the first entry of a race.
|
|
static std::vector<Variant> allVariants() {
|
|
std::vector<Variant> v;
|
|
auto add = [&](const char* n, int unroll, int load, int maxrreg, int bw, int minBlocks) { Variant x; x.name = n; x.unroll = unroll; x.load = load; x.maxrreg = maxrreg; x.blockWarps = bw; x.minBlocks = minBlocks; v.push_back(x); };
|
|
add("base", 0, 0, 0, 0, 0);
|
|
add("w2", 0, 0, 0, 2, 0);
|
|
add("w4", 0, 0, 0, 4, 0);
|
|
add("w8", 0, 0, 0, 8, 0);
|
|
add("u2", 2, 0, 0, 0, 0);
|
|
add("u8", 8, 0, 0, 0, 0);
|
|
add("ldg", 0, 1, 0, 0, 0);
|
|
add("ldcg", 0, 2, 0, 0, 0);
|
|
add("ldcs", 0, 3, 0, 0, 0);
|
|
add("r32", 0, 0, 32, 0, 0);
|
|
add("r64", 0, 0, 64, 0, 0);
|
|
add("lb4-w4", 0, 0, 0, 4, 4);
|
|
add("lb8-w2", 0, 0, 0, 2, 8);
|
|
add("u2-ldg", 2, 1, 0, 0, 0);
|
|
add("u2-w4", 2, 0, 0, 4, 0);
|
|
add("ldg-w4", 0, 1, 0, 4, 0);
|
|
add("ldcg-w4", 0, 2, 0, 4, 0);
|
|
return v;
|
|
}
|
|
|
|
// sp<N> and sp<N>-w<W>: the SM-sparse variants, made on demand (not in the catalogue, so a default race never runs them):
|
|
// N persistent blocks (1 to 4096) of W warps (default 32, the largest block, one block per SM at the hash's register count).
|
|
static bool parseSparseVariant(const std::string& name, Variant& out) {
|
|
if (name.rfind("sp", 0) != 0) return false;
|
|
size_t i = 2, n = 0; int blocks = 0, warps = 32;
|
|
while (i < name.size() && name[i] >= '0' && name[i] <= '9') { blocks = blocks * 10 + (name[i] - '0'); ++i; ++n; }
|
|
if (n == 0 || blocks < 1 || blocks > 4096) return false;
|
|
if (i < name.size()) {
|
|
if (name.compare(i, 2, "-w") != 0) return false;
|
|
i += 2; n = 0; warps = 0;
|
|
while (i < name.size() && name[i] >= '0' && name[i] <= '9') { warps = warps * 10 + (name[i] - '0'); ++i; ++n; }
|
|
if (n == 0 || i != name.size() || warps < 1 || warps > 32) return false;
|
|
}
|
|
out = Variant(); out.name = name; out.blockWarps = warps; out.sparseBlocks = blocks;
|
|
return true;
|
|
}
|
|
|
|
static const Variant* findVariant(const std::vector<Variant>& all, const std::string& name) {
|
|
for (const Variant& v : all) if (v.name == name) return &v;
|
|
static std::vector<Variant> made; // the on-demand sparse variants, kept so the pointer stays valid
|
|
Variant sp;
|
|
if (parseSparseVariant(name, sp)) {
|
|
for (const Variant& v : made) if (v.name == name) return &v;
|
|
made.reserve(64);
|
|
if (made.size() < 64) { made.push_back(sp); return &made.back(); }
|
|
}
|
|
return nullptr;
|
|
}
|
|
|
|
|
|
// The variant's source: the pack's bound-kernel text with the variant's rewrites. Every rewrite has an exact anchor
|
|
// in the text igneum-pow emits; a text without the anchor refuses the variant (why), it is never guessed.
|
|
static bool variantSource(const std::string& base, const Variant& v, int blockWarps, std::string& out, std::string& why) {
|
|
out = base;
|
|
if (v.unroll > 0) {
|
|
const char* anchor = "\n for (uint32_t it = 0u; it < ";
|
|
size_t p = out.find(anchor);
|
|
if (p == std::string::npos) { why = "no iteration loop in the bound kernel text"; return false; }
|
|
out.insert(p + 1, fmt("#pragma unroll %d\n", v.unroll));
|
|
}
|
|
if (v.load > 0) {
|
|
const char* fn = v.load == 1 ? "__ldg" : v.load == 2 ? "__ldcg" : "__ldcs";
|
|
size_t body = out.find("igneum_hash_bound(");
|
|
if (body == std::string::npos) { why = "no igneum_hash_bound in the text"; return false; }
|
|
size_t p = body; int n = 0;
|
|
while ((p = out.find(" ^ ds[", p)) != std::string::npos) {
|
|
size_t close = out.find(']', p);
|
|
if (close == std::string::npos) { why = "an unterminated dataset load"; return false; }
|
|
out.insert(close + 1, ")"); // " ^ ds[idx]" -> " ^ __ldg(&ds[idx])"
|
|
out.insert(p + 3, std::string(fn) + "(&");
|
|
p += 6; ++n;
|
|
}
|
|
if (n == 0) { why = "no dataset loads in the bound kernel"; return false; }
|
|
}
|
|
if (v.minBlocks > 0) {
|
|
const char* a = "__global__ void igneum_hash_bound(";
|
|
size_t p = out.find(a);
|
|
if (p == std::string::npos) { why = "no kernel declaration anchor"; return false; }
|
|
out.replace(p, std::strlen(a), fmt("__global__ void __launch_bounds__(%d, %d) igneum_hash_bound(", 32 * blockWarps, v.minBlocks));
|
|
}
|
|
if (v.sparseBlocks > 0) {
|
|
// The SM-sparse shape: the emitted kernel becomes a per-thread unit function taking its gid, and a persistent
|
|
// wrapper of the kernel's name loops the unit over the dispatch's nonces. Anchors: the kernel declaration as
|
|
// igneum-pow emits it (no __launch_bounds__ on it: the two rewrites are not combined) and its first line.
|
|
if (v.minBlocks > 0) { why = "sp and __launch_bounds__ minBlocks are not combined"; return false; }
|
|
// line-ending-agnostic: a pack copied through Windows may carry CRLF; the anchors below are LF, so the text is
|
|
// normalised first (the rewritten kernel then has the same bytes from LF and CRLF input)
|
|
out.erase(std::remove(out.begin(), out.end(), '\r'), out.end());
|
|
const char* a = "__global__ void igneum_hash_bound(";
|
|
size_t p = out.find(a);
|
|
if (p == std::string::npos) { why = "no kernel declaration anchor"; return false; }
|
|
size_t close = out.find(") {", p);
|
|
if (close == std::string::npos) { why = "no parameter list close"; return false; }
|
|
std::string params = out.substr(p + std::strlen(a), close - (p + std::strlen(a)));
|
|
if (params.find("scratch") != std::string::npos || params.find("units") != std::string::npos) { why = "a persistent (variant-5) pack has its own loop"; return false; }
|
|
// the argument names: the last token of each parameter
|
|
std::string args; size_t from = 0;
|
|
while (from <= params.size()) {
|
|
size_t comma = params.find(',', from);
|
|
std::string one = params.substr(from, comma == std::string::npos ? std::string::npos : comma - from);
|
|
size_t e = one.find_last_not_of(" \t");
|
|
if (e == std::string::npos) { why = "an empty parameter"; return false; }
|
|
size_t b = one.find_last_of(" \t*&", e);
|
|
size_t start = b == std::string::npos ? 0 : b + 1;
|
|
std::string nm = one.substr(start, e - start + 1); // e is the last character's index: the length is e - start + 1
|
|
// (the 03:23Z card run read `d, ou, baseNonc, mas, i`: the length was written e - start and dropped every name's last character)
|
|
if (nm.empty()) { why = "a parameter without a name"; return false; }
|
|
args += (args.empty() ? "" : ", ") + nm;
|
|
if (comma == std::string::npos) break;
|
|
from = comma + 1;
|
|
}
|
|
const char* gidLine = "\n uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x;\n";
|
|
size_t g = out.find(gidLine, close);
|
|
if (g == std::string::npos) { why = "no gid line after the declaration"; return false; }
|
|
out.replace(g, std::strlen(gidLine), "\n");
|
|
out.replace(p, close - p, "__device__ __forceinline__ void igneum_hash_bound_unit(" + params + ", uint32_t gid");
|
|
out += fmt("\n// SM-sparse wrapper (variant %s): %d persistent blocks of %d threads over the dispatch's nonces\n", v.name.c_str(), v.sparseBlocks, 32 * blockWarps);
|
|
out += fmt("__global__ void __launch_bounds__(%d) igneum_hash_bound(", 32 * blockWarps) + params + ", uint32_t nonces) {\n";
|
|
out += " for (uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x; gid < nonces; gid += gridDim.x * blockDim.x)\n";
|
|
out += " igneum_hash_bound_unit(" + args + ", gid);\n}\n";
|
|
}
|
|
return true;
|
|
}
|
|
|
|
// The tuning file: {"cards": {"<device name as this worker prints it>": {"variant": "u2-ldg", "race": false,
|
|
// "candidates": ["u2-ldg", "ldg", "base"]}}, ...}. Read with plain string scanning (no JSON library in this exe);
|
|
// a file that does not parse means no tuning. Keys and names are [A-Za-z0-9_.-].
|
|
struct Tuning {
|
|
bool found = false;
|
|
std::string variant; // pinned variant ("" = none)
|
|
bool race = true; // false: use the pinned variant without a race
|
|
std::vector<std::string> candidates;
|
|
};
|
|
|
|
static std::string jsonStringAfter(const std::string& t, size_t from, const char* key, size_t limit) {
|
|
size_t k = t.find(std::string("\"") + key + "\"", from);
|
|
if (k == std::string::npos || k > limit) return "";
|
|
size_t q = t.find('"', t.find(':', k) + 1);
|
|
if (q == std::string::npos) return "";
|
|
size_t e = t.find('"', q + 1);
|
|
return e == std::string::npos ? "" : t.substr(q + 1, e - q - 1);
|
|
}
|
|
|
|
// The race's order: base first, then the pinned or tuned candidates, then the rest (or the --race list only). Pure, so
|
|
// the emulation test reads it (emu/variant-test.cpp). Membership in `order` is by NAME: findVariant answers for the
|
|
// on-demand sp<N> names whatever list it is asked about, which is the fault of the 8 October 02:52Z run (the pinned
|
|
// sparse variant was looked up in `order`, found by the on-demand path, and never pushed: "variants 1 base only").
|
|
static std::vector<Variant> raceOrder(const std::vector<Variant>& all, const std::string& pinned, const Tuning& tu, const std::string& race) {
|
|
std::vector<Variant> order;
|
|
auto has = [&](const std::string& n) { for (const Variant& v : order) if (v.name == n) return true; return false; };
|
|
auto push = [&](const std::string& n) { const Variant* v = findVariant(all, n); if (v && !has(n)) order.push_back(*v); };
|
|
push("base");
|
|
if (!pinned.empty()) push(pinned);
|
|
else {
|
|
for (const std::string& n : tu.candidates) push(n);
|
|
if (race != "on" && race != "off") { std::string rest = race; size_t i = 0; while (i <= rest.size()) { size_t j = rest.find(',', i); if (j == std::string::npos) j = rest.size(); if (j > i) push(rest.substr(i, j - i)); i = j + 1; } }
|
|
else if (race == "on") for (const Variant& v : all) push(v.name);
|
|
}
|
|
return order;
|
|
}
|
|
|
|
static Tuning readTuning(const std::string& text, const std::string& device) {
|
|
Tuning tu;
|
|
if (text.empty()) return tu;
|
|
size_t cards = text.find("\"cards\"");
|
|
if (cards == std::string::npos) return tu;
|
|
size_t k = text.find("\"" + device + "\"", cards);
|
|
if (k == std::string::npos) return tu;
|
|
size_t open = text.find('{', k);
|
|
if (open == std::string::npos) return tu;
|
|
size_t close = open; int depth = 0;
|
|
for (; close < text.size(); ++close) { if (text[close] == '{') ++depth; else if (text[close] == '}' && --depth == 0) break; }
|
|
if (close >= text.size()) return tu;
|
|
tu.found = true;
|
|
tu.variant = jsonStringAfter(text, open, "variant", close);
|
|
size_t r = text.find("\"race\"", open);
|
|
if (r != std::string::npos && r < close) { size_t c = text.find(':', r); tu.race = text.compare(text.find_first_not_of(" \t\r\n", c + 1), 5, "false") != 0; }
|
|
size_t cand = text.find("\"candidates\"", open);
|
|
if (cand != std::string::npos && cand < close) {
|
|
size_t a = text.find('[', cand), b = text.find(']', a == std::string::npos ? cand : a);
|
|
if (a != std::string::npos && b != std::string::npos && b < close) {
|
|
size_t i = a;
|
|
while ((i = text.find('"', i + 1)) != std::string::npos && i < b) { size_t e = text.find('"', i + 1); if (e == std::string::npos || e > b) break; tu.candidates.push_back(text.substr(i + 1, e - i - 1)); i = e; }
|
|
}
|
|
}
|
|
return tu;
|
|
}
|
|
|
|
struct RaceEntry {
|
|
Variant v;
|
|
int blockWarps = 1; // the block this entry runs with
|
|
Compiled cb;
|
|
CUmodule mod = nullptr;
|
|
CUfunction fn = nullptr;
|
|
int regs = 0, blocksPerSM = 0;
|
|
double mhs = 0; // best round
|
|
bool ok = false; // compiled, loaded, self-tested
|
|
std::string note; // why not, or a detail
|
|
std::string src;
|
|
};
|
|
|
|
// One timed window on the card for an entry: the pack's vector warps (bit-exact or the entry is out), then
|
|
// launches of `batch` nonces until benchMs elapsed (the first launch warms up and is not counted). Holds gpuMutex.
|
|
static bool raceTime(Ctx& c, Pair* p, const PfPack& pk, RaceEntry& e, CUdeviceptr dOut, uint32_t batch, int benchMs, CUstream s, bool selfTest) {
|
|
std::lock_guard<std::mutex> hold(gpuMutex);
|
|
CUfunction keep = p->fHashBound;
|
|
int keepSparse = p->sparseBlocks;
|
|
p->fHashBound = e.fn;
|
|
p->sparseBlocks = e.v.sparseBlocks;
|
|
std::string err;
|
|
uint32_t block = 32u * (uint32_t)e.blockWarps;
|
|
bool ok = true;
|
|
if (selfTest && pk.haveVectors) {
|
|
std::vector<uint64_t> vec(32);
|
|
for (int w = 0; w < pk.vecWarps && ok; ++w) {
|
|
if (!launchHash(c, p, dOut, pk.vecBase[w], p->sw, block, block, s, err)) { e.note = "launch: " + err; ok = false; break; }
|
|
CUresult r = c.drv.streamSynchronize(s);
|
|
if (r == CUDA_SUCCESS) r = c.drv.memcpyDtoH(vec.data(), dOut, 32u * 8u);
|
|
if (r != CUDA_SUCCESS) { e.note = "vector warp: " + c.err(r); ok = false; break; }
|
|
for (int l = 0; l < 32; ++l) if (vec[(size_t)l] != pk.vecOut[w][l]) { e.note = fmt("vector warp %d lane %d: device %016llx expected %016llx (discarded)", w, l, (unsigned long long)vec[(size_t)l], (unsigned long long)pk.vecOut[w][l]); ok = false; break; }
|
|
}
|
|
}
|
|
if (ok) {
|
|
uint32_t iw[8]; std::memcpy(iw, p->sw, 32);
|
|
uint32_t n = batch - (batch % block);
|
|
if (n == 0) n = block;
|
|
double t0 = 0; uint64_t hashes = 0; int launches = 0;
|
|
while (true) {
|
|
if (!launchHash(c, p, dOut, 0x10000000u + (uint32_t)launches * n, iw, n, block, s, err)) { e.note = "launch: " + err; ok = false; break; }
|
|
CUresult r = c.drv.streamSynchronize(s);
|
|
if (r != CUDA_SUCCESS) { e.note = "bench: " + c.err(r); ok = false; break; }
|
|
double now = wallMs();
|
|
if (launches == 0) t0 = now; else hashes += n;
|
|
++launches;
|
|
if (launches >= 3 && now - t0 >= benchMs) { double mhs = (double)hashes / (now - t0) / 1000.0; if (mhs > e.mhs) e.mhs = mhs; break; }
|
|
}
|
|
}
|
|
p->fHashBound = keep;
|
|
p->sparseBlocks = keepSparse;
|
|
return ok;
|
|
}
|
|
|
|
// Races the bound kernel of `p` (its cache and dataset are built, its base kernel self-tested) and installs the
|
|
// winner: p->modBound, fHashBound, regs, blockWarps, variant. The pair keeps serving its base kernel if every other
|
|
// entry fails. `boundDev`, `programH`, `memhardH` are the texts the base was compiled from. Sets p->raceLine.
|
|
static void racePair(Ctx& c, Pair* p, const PfPack& pk, const std::string& boundDev, const std::string& programH, const std::string& memhardH, CUstream s) {
|
|
double t0 = wallMs();
|
|
std::vector<Variant> all = allVariants();
|
|
Tuning tu = readTuning(c.tuning, c.name);
|
|
std::string pinned = !c.pinned.empty() ? c.pinned : (tu.found && !tu.race ? tu.variant : "");
|
|
std::vector<Variant> order = raceOrder(all, pinned, tu, c.race);
|
|
#ifdef IGNEUM_EMU
|
|
order.resize(1); // the stand-in checks that the handed-over text is the pack's; no rewrites under emulation
|
|
#endif
|
|
if (order.size() < 2) {
|
|
// nothing to race against base (emulation, or --race with no known name): no timing, the base kernel serves
|
|
p->blockWarps = c.blockWarps; p->variant = "base"; p->raceMs = wallMs() - t0;
|
|
p->raceLine = fmt("race %.16s device %s variants 1 base only, no race (%s)", p->epochHex.c_str(), c.name.c_str(),
|
|
#ifdef IGNEUM_EMU
|
|
"emulation");
|
|
#else
|
|
"no other variant named");
|
|
#endif
|
|
return;
|
|
}
|
|
const bool pinnedOnly = !pinned.empty() && order.size() == 2;
|
|
uint32_t batch = 1u << c.batchLog2;
|
|
int benchMs = c.raceBenchMs;
|
|
double deadline = t0 + c.raceBudgetS * 1000.0;
|
|
std::vector<RaceEntry> entries;
|
|
for (const Variant& v : order) {
|
|
RaceEntry e; e.v = v; e.blockWarps = v.blockWarps > 0 ? v.blockWarps : c.blockWarps;
|
|
if (v.name == "base") { e.mod = p->modBound; e.fn = p->fHashBound; e.regs = p->regs; e.blocksPerSM = p->blocksPerSM; e.ok = true; }
|
|
else if (!variantSource(boundDev, v, e.blockWarps, e.src, e.note)) e.ok = false;
|
|
else e.ok = true; // compiled below
|
|
entries.push_back(std::move(e));
|
|
}
|
|
// Compile the variants, up to four at a time (NVRTC is thread-safe; the compile is CPU work)
|
|
{
|
|
std::vector<size_t> todo;
|
|
for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok) todo.push_back(i);
|
|
size_t next = 0;
|
|
std::mutex m;
|
|
auto work = [&]() {
|
|
while (true) {
|
|
size_t i;
|
|
{ std::lock_guard<std::mutex> g(m); if (next >= todo.size() || wallMs() > deadline - benchMs) return; i = todo[next++]; }
|
|
RaceEntry& e = entries[i];
|
|
std::vector<std::string> extra;
|
|
if (e.v.maxrreg > 0) extra.push_back(fmt("--maxrregcount=%d", e.v.maxrreg));
|
|
std::string err;
|
|
if (!rtcCompile(c, e.src, "kernel_bound.cu", programH, memhardH, { "igneum_hash_bound" }, e.cb, err, extra)) { e.ok = false; e.note = "compile: " + err.substr(0, 200); }
|
|
}
|
|
};
|
|
int threads = (int)std::min<size_t>(4, std::max<size_t>(1, todo.size()));
|
|
std::vector<std::thread> ts;
|
|
for (int t = 0; t < threads; ++t) ts.emplace_back(work);
|
|
for (std::thread& t : ts) t.join();
|
|
for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok && entries[i].cb.image.empty()) { entries[i].ok = false; entries[i].note = "not compiled: the race budget ran out"; }
|
|
}
|
|
// Load the modules (the context is current on this thread)
|
|
for (size_t i = 1; i < entries.size(); ++i) {
|
|
RaceEntry& e = entries[i];
|
|
if (!e.ok) continue;
|
|
CUresult r = c.drv.moduleLoadData(&e.mod, e.cb.image.data());
|
|
if (r != CUDA_SUCCESS) { e.ok = false; e.note = "cuModuleLoadData: " + c.err(r); e.mod = nullptr; continue; }
|
|
if (c.drv.moduleGetFunction(&e.fn, e.mod, e.cb.lowered[0].c_str()) != CUDA_SUCCESS) { e.ok = false; e.note = "function not in the module"; continue; }
|
|
c.drv.funcGetAttribute(&e.regs, CU_FUNC_ATTRIBUTE_NUM_REGS, e.fn);
|
|
c.drv.occupancy(&e.blocksPerSM, e.fn, 32 * e.blockWarps, 0);
|
|
}
|
|
double compileMs = wallMs() - t0;
|
|
// Time them: rounds over the entries, interleaved, best per entry. A pinned variant is only self-tested.
|
|
CUdeviceptr dOut = 0;
|
|
std::string benchErr;
|
|
if (c.drv.memAlloc(&dOut, (size_t)batch * 8u) != CUDA_SUCCESS) { benchErr = "cuMemAlloc for the race"; for (RaceEntry& e : entries) if (e.v.name != "base") e.ok = false; }
|
|
int rounds = pinnedOnly ? 1 : std::max(1, c.raceRounds);
|
|
for (int round = 0; round < rounds && benchErr.empty(); ++round) {
|
|
for (size_t i = 0; i < entries.size(); ++i) {
|
|
RaceEntry& e = entries[i];
|
|
if (!e.ok) continue;
|
|
if (i > 0 && round == 0 && wallMs() > deadline) { e.ok = false; e.note = "not timed: the race budget ran out"; continue; }
|
|
if (pinnedOnly && i == 0) continue;
|
|
if (!raceTime(c, p, pk, e, dOut, batch, pinnedOnly ? 0 : benchMs, s, round == 0)) e.ok = false;
|
|
// the mutex is not fair: give the job loop the card between windows (measured on the Mac, 4 October
|
|
// 2026: without this a queued job waited the whole race, 36 s)
|
|
std::this_thread::sleep_for(std::chrono::milliseconds(150));
|
|
}
|
|
}
|
|
if (dOut) c.drv.memFree(dOut);
|
|
// The winner: the fastest entry; base keeps its place unless a variant is at least 0.5% faster (noise guard).
|
|
size_t win = 0;
|
|
if (pinnedOnly && entries.size() == 2 && entries[1].ok) win = 1;
|
|
else for (size_t i = 1; i < entries.size(); ++i) if (entries[i].ok && entries[i].mhs > entries[win].mhs * (win == 0 ? 1.005 : 1.0)) win = i;
|
|
double baseMhs = entries[0].mhs, winMhs = entries[win].mhs;
|
|
if (win != 0) {
|
|
RaceEntry& w = entries[win];
|
|
c.drv.moduleUnload(p->modBound);
|
|
p->modBound = w.mod; p->fHashBound = w.fn; p->regs = w.regs; p->blocksPerSM = w.blocksPerSM; p->blockWarps = w.blockWarps; p->variant = w.v.name;
|
|
p->sparseBlocks = w.v.sparseBlocks;
|
|
w.mod = nullptr;
|
|
} else {
|
|
p->blockWarps = c.blockWarps; p->variant = "base"; p->sparseBlocks = 0;
|
|
}
|
|
for (size_t i = 1; i < entries.size(); ++i) if (entries[i].mod) c.drv.moduleUnload(entries[i].mod);
|
|
p->raceMs = wallMs() - t0;
|
|
// The one line. Variants in race order: name=MH/s (regs), or name=- (why).
|
|
uint32_t loads = 0, wide = 0;
|
|
pf_define_u32(programH.c_str(), "IGNEUM_LOADS_PER_HASH", &loads);
|
|
pf_define_u32(programH.c_str(), "IGNEUM_WIDE_LOADS_PER_HASH", &wide);
|
|
std::string line = fmt("race %.16s device %s driver %d.%d arch %s loads %u wide %u variants %zu", p->epochHex.c_str(), c.name.c_str(), c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.archOpt.c_str(), loads, wide, entries.size());
|
|
for (const RaceEntry& e : entries) {
|
|
if (e.ok && (e.mhs > 0 || pinnedOnly)) line += fmt(" %s=%.3f/%dr/%dw", e.v.name.c_str(), e.mhs, e.regs, e.blockWarps);
|
|
else line += fmt(" %s=-", e.v.name.c_str());
|
|
}
|
|
line += fmt(" winner %s %.3f base %.3f gain %+.2f%% compile %.0f bench %.0f total %.0f ms%s%s", p->variant.c_str(), winMhs, baseMhs, baseMhs > 0 ? (winMhs / baseMhs - 1.0) * 100.0 : 0.0, compileMs, p->raceMs - compileMs, p->raceMs,
|
|
pinnedOnly ? " pinned by tuning" : (tu.found ? " tuned order" : ""), benchErr.empty() ? "" : (" " + benchErr).c_str());
|
|
for (const RaceEntry& e : entries) if (!e.ok && !e.note.empty()) line += " | " + e.v.name + ": " + e.note;
|
|
p->raceLine = line;
|
|
}
|
|
|
|
// Compiles the pack in `dir`, builds its cache and dataset on stream `s`, runs the self-test, races the variants
|
|
// (`race`). Returns the pair or null with `err` set. Runs on the main thread for --pack and --check, on the prepare
|
|
// thread for `prepare`.
|
|
static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string& err, bool race) {
|
|
PfPack pk;
|
|
char perr[512];
|
|
if (!pf_load(dir.c_str(), &pk, perr, sizeof(perr))) { err = std::string("pack ") + dir + ": " + perr; return nullptr; }
|
|
bool ok1, ok2, ok3, ok4;
|
|
std::string kernelCu = readText(dir + "/kernel.cu", ok1), boundCu = readText(dir + "/kernel_bound.cu", ok2);
|
|
std::string programH = readText(dir + "/program.h", ok3), memhardH = readText(dir + "/memhard.h", ok4);
|
|
if (!ok1 || !ok2 || !ok3 || !ok4) { err = "pack " + dir + " lacks kernel.cu, kernel_bound.cu, program.h or memhard.h"; return nullptr; }
|
|
std::string kernelDev, boundDev;
|
|
if (!deviceOnly(kernelCu, kernelDev, err) || !deviceOnly(boundCu, boundDev, err)) { err = "pack " + dir + ": " + err; return nullptr; }
|
|
Pair* p = new Pair();
|
|
p->dir = dir; p->epochHex = pk.epochHex; p->dayHex = pk.dayHex; p->seedString = pk.seedString;
|
|
std::memcpy(p->sw, pk.seedw, 32); std::memcpy(p->kw, pk.keyw, 32);
|
|
p->datasetLog2 = pk.datasetLog2; p->words = 1u << pk.datasetLog2; p->cacheWords = 1u << pk.cacheLog2Words; p->cacheSegments = pk.cacheSegments;
|
|
// Compile
|
|
Compiled ck, cb;
|
|
std::vector<std::string> kernelNames = { "igneum_cache_fill", "igneum_build" };
|
|
if (pk.hotMb) kernelNames.push_back("igneum_hot_fill");
|
|
if (!rtcCompile(c, kernelDev, "kernel.cu", programH, memhardH, kernelNames, 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; }
|
|
if (pk.hotMb && c.drv.moduleGetFunction(&p->fHotFill, p->modKernel, ck.lowered[2].c_str()) != CUDA_SUCCESS) { err = "igneum_hot_fill (" + ck.lowered[2] + ") 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);
|
|
}
|
|
p->loadClass = pk.loadClass; p->loadsPerHash = pk.loadsPerHash; p->bytesPerHash = pk.bytesPerHash; p->scratchOps = pk.scratchOps;
|
|
p->programClass = pk.programClass; p->eraHex = pk.eraHex;
|
|
p->persistent = pk.persistent != 0;
|
|
p->hotMb = pk.hotMb; p->hotWords = pk.hotWords; p->hotSegments = pk.hotSegments; p->hotSlots = pk.hotSlots;
|
|
p->residentWarps = p->blocksPerSM * c.blockWarps * c.sms;
|
|
size_t scratchBytes = 0, hotBytes = (size_t)pk.hotWords * 4u;
|
|
if (p->persistent) {
|
|
// Variant 5: one arena per launched warp. The launch is the resident capacity (the occupancy query), rounded
|
|
// down to a power of two so it divides every batch, or --warps; the allocation cannot change the occupancy
|
|
// (registers and shared memory decide it), and the number is re-queried after the allocation below to show it.
|
|
if (c.blockWarps != 1) { err = "a variant-5 pack runs one warp per block: use --block-warps 1"; releasePair(c, p); return nullptr; }
|
|
int w = c.warps > 0 ? c.warps : p->residentWarps;
|
|
int pw = 1; while (pw * 2 <= w) pw *= 2;
|
|
p->warps = pw;
|
|
scratchBytes = (size_t)p->warps * 32u * (size_t)pk.scratchWordsPerLane * 4u;
|
|
}
|
|
// 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 + scratchBytes + hotBytes + (64u << 20)) {
|
|
err = fmt("%llu MiB free on the device, this pack needs %llu MiB (cache %llu + dataset %llu + scratch %llu + hot %llu)", (unsigned long long)(freeB >> 20), (unsigned long long)((cacheBytes + dsBytes + scratchBytes + hotBytes) >> 20), (unsigned long long)(cacheBytes >> 20), (unsigned long long)(dsBytes >> 20), (unsigned long long)(scratchBytes >> 20), (unsigned long long)(hotBytes >> 20));
|
|
releasePair(c, p); return nullptr;
|
|
}
|
|
}
|
|
if (p->persistent) {
|
|
CUresult r = c.drv.memAlloc(&p->scratch, scratchBytes);
|
|
if (r != CUDA_SUCCESS) { err = "cuMemAlloc scratch: " + c.err(r); p->scratch = 0; releasePair(c, p); return nullptr; }
|
|
p->scratchBytes = scratchBytes;
|
|
int after = 0;
|
|
c.drv.occupancy(&after, p->fHashBound, 32 * c.blockWarps, 0);
|
|
info(fmt("variant 5: %d persistent warps (resident capacity %d = %d blocks/SM x %d warps/block x %d SMs; occupancy query after the allocation %d blocks/SM), scratch %llu MiB (%u KiB per warp)",
|
|
p->warps, p->residentWarps, p->blocksPerSM, c.blockWarps, c.sms, after, (unsigned long long)(scratchBytes >> 20), pk.scratchWordsPerLane * 4u * 32u / 1024u));
|
|
}
|
|
{
|
|
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;
|
|
// Hot table (hot-table experiment): filled from the epoch seed by the pack's own kernel, never shipped
|
|
if (pk.hotMb) {
|
|
t0 = wallMs();
|
|
CUresult r = c.drv.memAlloc(&p->hot, hotBytes);
|
|
if (r != CUDA_SUCCESS) { err = "cuMemAlloc hot table: " + c.err(r); p->hot = 0; releasePair(c, p); return nullptr; }
|
|
uint32_t nSeg = pk.hotSegments, block = 256u, grid = (nSeg + block - 1u) / block;
|
|
void* args[2] = { &p->hot, &nSeg };
|
|
r = c.drv.launchKernel(p->fHotFill, grid, 1, 1, block, 1, 1, 0, s, args, nullptr);
|
|
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(s);
|
|
if (r != CUDA_SUCCESS) { err = "hot table fill: " + c.err(r); releasePair(c, p); return nullptr; }
|
|
p->hotMs = 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);
|
|
uint32_t hotHead[16] = {0}, hotLast[16] = {0};
|
|
uint64_t hotFnv = 0;
|
|
if (p->hot) {
|
|
std::vector<uint32_t> hw(p->hotWords);
|
|
if ((r = c.drv.memcpyDtoH(hw.data(), p->hot, hotBytes)) != CUDA_SUCCESS) { err = "cuMemcpyDtoH hot table: " + c.err(r); releasePair(c, p); return nullptr; }
|
|
std::memcpy(hotHead, hw.data(), 64);
|
|
std::memcpy(hotLast, hw.data() + p->hotWords - 16u, 64);
|
|
hotFnv = pf_fnv1a64(hw.data(), hotBytes);
|
|
}
|
|
char line[1024];
|
|
p->checkPass = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples.data(), vec.data(), p->hot ? hotHead : nullptr, p->hot ? hotLast : nullptr, hotFnv, line, sizeof(line)) != 0;
|
|
p->checked = true;
|
|
p->check = line;
|
|
}
|
|
p->checkMs = wallMs() - t0;
|
|
if (!p->checkPass) { err = p->check; releasePair(c, p); return nullptr; }
|
|
p->blockWarps = c.blockWarps;
|
|
if (race && c.race != "off") racePair(c, p, pk, boundDev, programH, memhardH, s);
|
|
return p;
|
|
}
|
|
|
|
static std::string pairSummary(const Pair* p) {
|
|
return fmt("nvrtc %.0f cache %.0f dataset %.0f hot %.0f check %.0f race %.0f ms variant %s; %s", p->compileMs, p->cacheMs, p->dsMs, p->hotMs, p->checkMs, p->raceMs, p->variant.c_str(), p->check.c_str());
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Prepare, on its own thread
|
|
|
|
struct PrepareTask {
|
|
std::string epochHex, dayHex, dir, error;
|
|
std::string wantClass, wantEra; // the class and era the prepare line named (empty: any)
|
|
std::atomic<bool> done{false};
|
|
Pair* result = nullptr;
|
|
double t0 = 0;
|
|
std::thread thread;
|
|
};
|
|
|
|
static void prepareRun(Ctx* c, PrepareTask* t) {
|
|
CUstream s = nullptr;
|
|
std::string err;
|
|
if (c->drv.ctxSetCurrent(c->ctx) != CUDA_SUCCESS) { t->error = "cuCtxSetCurrent on the prepare thread"; t->done = true; return; }
|
|
if (c->drv.streamCreate(&s, CU_STREAM_NON_BLOCKING) != CUDA_SUCCESS) { t->error = "cuStreamCreate on the prepare thread"; t->done = true; return; }
|
|
Pair* p = buildPair(*c, t->dir, s, err, true);
|
|
c->drv.streamDestroy(s);
|
|
if (p && (p->epochHex != t->epochHex || p->dayHex != t->dayHex)) {
|
|
err = "the pack in " + t->dir + " is for epoch " + p->epochHex.substr(0, 16) + " day " + p->dayHex + ", not the prepared seeds";
|
|
releasePair(*c, p); p = nullptr;
|
|
}
|
|
char why[256];
|
|
if (p && !pf_pack_class_ok(p->programClass.c_str(), p->eraHex.c_str(), t->wantClass.c_str(), t->wantEra.c_str(), why, sizeof(why))) {
|
|
err = "pack " + t->dir + ": " + why;
|
|
releasePair(*c, p); p = nullptr;
|
|
}
|
|
t->result = p;
|
|
t->error = err;
|
|
t->done = true;
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Serve
|
|
|
|
struct Options {
|
|
bool serve = false, check = false, raceOnly = false;
|
|
bool bench = false, memprobe = false; // read-width experiment (5 October 2026)
|
|
bool microbench = false; // Counter ASIC 4.0 research (7 October 2026): one kernel per GPU hardware block, sustained
|
|
int mbSeconds = 60; // --mb-seconds: the sustained window per probe (the power sampler reads it)
|
|
std::string mbOnly; // --mb-only a,b,c: these probes only
|
|
bool listRace = false; // --list-race: print the race order --bench would run (card-free), then exit
|
|
int batches = 5, warps = 0, probeMib = 0;
|
|
int device = 0, batchLog2 = 22, blockWarps = 1;
|
|
std::string pack, arch = "auto";
|
|
std::string race = "on", pinned, tuningPath;
|
|
int raceBenchMs = 2000, raceBudgetS = 120, raceRounds = 0; // rounds 0 = 1 in --serve, 3 in --race
|
|
};
|
|
|
|
static void usage() {
|
|
std::printf("igneum-worker-cuda %s\n"
|
|
" --serve --pack <dir> GPU worker for igneum-miner --worker: jobs on stdin, found/done lines on stdout\n"
|
|
" --check --pack <dir> compile the pack, build its cache and dataset, self-test, print timings, exit 0 or 1\n"
|
|
" --device D CUDA device index (default 0)\n"
|
|
" --batch-log2 B nonces per dispatch = 2^B (default 22)\n"
|
|
" --block-warps W warps per thread block (default 1)\n"
|
|
" --arch sm_XY|compute_XY|auto NVRTC target (default auto: the device's architecture)\n"
|
|
" --bench --pack <dir> read-width experiment: build and self-test the pack, time --batches dispatches of 2^B nonces,\n"
|
|
" print the 2^B fingerprint at base nonce 0 (one RESULT line); a variant-5 pack runs --warps persistent warps\n"
|
|
" --memprobe [--probe-mib N] no pack: dependent random 4, 16 and 64-byte reads, independent reads, a coalesced stream and an\n"
|
|
" integer chain at 4, 64 and 1024 MiB (the same table as igneum-worker-opencl --memprobe)\n"
|
|
" --list-race [--pack <dir>] [--variant <name>] [--race ...] card-free: print the race order the run would build (variants N, names) and,\n"
|
|
" with a pack, whether the named variant's kernel rewrite applies to its bound kernel; then exit 0, or 1 when a named variant is missing\n"
|
|
" --microbench [--mb-seconds 60] [--mb-only a,b] no pack: one sustained kernel per GPU hardware block (ALU families, shuffle,\n"
|
|
" byte permute, FP32 and FP16 FMA, tensor tiles per precision, L2-resident chases, texture fetches, a DRAM\n"
|
|
" chase and a sleeping-SM floor), each for --mb-seconds with UTC start and end stamps for a power sampler\n"
|
|
" --batches N --bench: timed dispatches (default 5)\n"
|
|
" --warps N --bench on a variant-5 pack: persistent warps (default: the occupancy capacity, rounded down to a power of two)\n"
|
|
" --race --pack <dir> the variant race alone (3 rounds): one line per variant, the race line, exit 0 or 1\n"
|
|
" --race on|off|a,b,c in --serve: race every variant (default), none, or these names\n"
|
|
" --race-bench-ms N timed window per variant (default 2000)\n"
|
|
" --race-budget-s N a race stops compiling and timing after this (default 120; base is kept)\n"
|
|
" --race-rounds N interleaved rounds, best per variant (default 1 in --serve, 3 in --race)\n"
|
|
" --variant <name> use this variant without a race (also from the tuning file); sp<N> or sp<N>-w<W> is the\n"
|
|
" SM-sparse shape (N persistent blocks of W warps, default 32; Counter ASIC 4.0 research)\n"
|
|
" --tuning <file> the per-card tuning file (default: IGNEUM_TUNING_FILE from the environment)\n", WORKER_VERSION);
|
|
}
|
|
|
|
static Options parseArgs(int argc, char** argv) {
|
|
Options o;
|
|
for (int i = 1; i < argc; ++i) {
|
|
std::string a = argv[i];
|
|
auto next = [&]() -> std::string { if (i + 1 >= argc) { usage(); std::exit(2); } return argv[++i]; };
|
|
if (a == "--serve") o.serve = true;
|
|
else if (a == "--check") o.check = true;
|
|
else if (a == "--bench") o.bench = true;
|
|
else if (a == "--memprobe") o.memprobe = true;
|
|
else if (a == "--microbench") o.microbench = true;
|
|
else if (a == "--mb-seconds") o.mbSeconds = std::atoi(next().c_str());
|
|
else if (a == "--mb-only") o.mbOnly = next();
|
|
else if (a == "--list-race") o.listRace = true;
|
|
else if (a == "--batches") o.batches = std::atoi(next().c_str());
|
|
else if (a == "--warps") o.warps = std::atoi(next().c_str());
|
|
else if (a == "--probe-mib") o.probeMib = std::atoi(next().c_str());
|
|
else if (a == "--race" && (i + 1 >= argc || std::string(argv[i + 1]).rfind("--", 0) == 0)) o.raceOnly = true;
|
|
else if (a == "--race") o.race = next();
|
|
else if (a == "--race-bench-ms") o.raceBenchMs = std::atoi(next().c_str());
|
|
else if (a == "--race-budget-s") o.raceBudgetS = std::atoi(next().c_str());
|
|
else if (a == "--race-rounds") o.raceRounds = std::atoi(next().c_str());
|
|
else if (a == "--variant") o.pinned = next();
|
|
else if (a == "--tuning") o.tuningPath = next();
|
|
else if (a == "--pack") o.pack = next();
|
|
else if (a == "--device") o.device = std::atoi(next().c_str());
|
|
else if (a == "--batch-log2") o.batchLog2 = std::atoi(next().c_str());
|
|
else if (a == "--block-warps") o.blockWarps = std::atoi(next().c_str());
|
|
else if (a == "--arch") o.arch = next();
|
|
else if (a == "--no-prepare") { /* accepted for symmetry with the other workers; prepare is always on here */ }
|
|
else if (a == "-h" || a == "--help") { usage(); std::exit(0); }
|
|
else { std::printf("unknown argument %s\n", argv[i]); usage(); std::exit(2); }
|
|
}
|
|
if (o.batchLog2 < 10 || o.batchLog2 > 28) { std::printf("--batch-log2 must be between 10 and 28\n"); std::exit(2); }
|
|
if (o.blockWarps < 1 || o.blockWarps > 32) { std::printf("--block-warps must be between 1 and 32\n"); std::exit(2); }
|
|
if (!o.serve && !o.check && !o.raceOnly && !o.bench && !o.memprobe && !o.microbench && !o.listRace) { usage(); std::exit(2); }
|
|
if (o.mbSeconds < 5 || o.mbSeconds > 600) { std::printf("--mb-seconds must be between 5 and 600\n"); std::exit(2); }
|
|
if (o.raceBenchMs < 200 || o.raceBenchMs > 20000) { std::printf("--race-bench-ms must be between 200 and 20000\n"); std::exit(2); }
|
|
if (o.raceBudgetS < 5 || o.raceBudgetS > 540) { std::printf("--race-budget-s must be between 5 and 540 (the prepare lead is 600 DAA)\n"); std::exit(2); }
|
|
if (o.raceRounds == 0) o.raceRounds = o.raceOnly ? 3 : 1;
|
|
if (o.tuningPath.empty()) if (const char* t = std::getenv("IGNEUM_TUNING_FILE")) o.tuningPath = t;
|
|
if (o.pack.empty() && !o.memprobe && !o.microbench && !o.listRace) { std::printf("--pack <dir> is required (igneum-miner export-pack <node> <dir> writes one)\n"); std::exit(2); }
|
|
while (o.pack.size() > 1 && (o.pack.back() == '/' || o.pack.back() == '\\')) o.pack.pop_back();
|
|
return o;
|
|
}
|
|
|
|
static std::vector<std::string> split(const std::string& line) {
|
|
std::vector<std::string> f;
|
|
size_t i = 0;
|
|
while (i < line.size()) {
|
|
while (i < line.size() && (line[i] == ' ' || line[i] == '\t' || line[i] == '\r')) ++i;
|
|
size_t j = i;
|
|
while (j < line.size() && line[j] != ' ' && line[j] != '\t' && line[j] != '\r') ++j;
|
|
if (j > i) f.push_back(line.substr(i, j - i));
|
|
i = j;
|
|
}
|
|
return f;
|
|
}
|
|
|
|
// A pack directory for the given seeds under `root` (one subdirectory per pack, each with seeds.txt), or "".
|
|
static std::string findPackFor(const std::string& root, const std::string& epochHex, const std::string& dayHex) {
|
|
if (root.empty()) return "";
|
|
std::vector<std::string> names;
|
|
#ifdef _WIN32
|
|
WIN32_FIND_DATAA fd;
|
|
HANDLE h = FindFirstFileA((root + "\\*").c_str(), &fd);
|
|
if (h == INVALID_HANDLE_VALUE) return "";
|
|
do { if ((fd.dwFileAttributes & FILE_ATTRIBUTE_DIRECTORY) && fd.cFileName[0] != '.') names.push_back(root + "\\" + fd.cFileName); } while (FindNextFileA(h, &fd));
|
|
FindClose(h);
|
|
#else
|
|
DIR* d = opendir(root.c_str());
|
|
if (!d) return "";
|
|
while (dirent* e = readdir(d)) if (e->d_name[0] != '.') names.push_back(root + "/" + e->d_name);
|
|
closedir(d);
|
|
#endif
|
|
for (const std::string& dir : names) {
|
|
bool ok = false;
|
|
std::string seeds = readText(dir + "/seeds.txt", ok);
|
|
if (!ok) continue;
|
|
char e[65] = {0}, dd[PF_HEX_CAP] = {0};
|
|
if (pf_seeds_line(seeds.c_str(), "epoch_seed_hex", e, sizeof(e)) && pf_seeds_line(seeds.c_str(), "day_seed_hex", dd, sizeof(dd)) && epochHex == e && dayHex == dd) return dir;
|
|
}
|
|
return "";
|
|
}
|
|
|
|
static std::string parentDir(const std::string& p) { size_t i = p.find_last_of("/\\"); return i == std::string::npos ? "." : p.substr(0, i); }
|
|
|
|
static int runServe(Ctx& c, const Options& o, Pair* cur) {
|
|
const uint32_t batch = 1u << o.batchLog2;
|
|
CUdeviceptr dOut = 0;
|
|
std::string err;
|
|
if (c.drv.memAlloc(&dOut, (size_t)batch * 8u) != CUDA_SUCCESS) { emit("error 0 cuMemAlloc out buffer"); return 2; }
|
|
std::vector<uint64_t> hOut(batch);
|
|
Pair* prepared = nullptr;
|
|
Pair* old = nullptr;
|
|
PrepareTask* task = nullptr;
|
|
std::string prepareRoot; // the parent of the last prepare's pack directory: where the miner writes its packs
|
|
emit(fmt("ready cuda %s pack %s dataset-log2 %u batch %u regs %d prepare 1 path nvrtc %d.%d driver %d.%d arch %s variant %s race %s worker %s",
|
|
c.name.c_str(), cur->seedString.c_str(), cur->datasetLog2, batch, cur->regs, c.rtcMajor, c.rtcMinor, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.archOpt.c_str(), cur->variant.c_str(), c.race.c_str(), WORKER_VERSION));
|
|
info(fmt("first pack %s: %s", cur->dir.c_str(), pairSummary(cur).c_str()));
|
|
if (!cur->raceLine.empty()) emit(cur->raceLine);
|
|
std::string line;
|
|
while (std::getline(std::cin, line)) {
|
|
if (line == "quit") break;
|
|
std::vector<std::string> f = split(line);
|
|
if (f.empty()) continue;
|
|
// A finished prepare is reported here, between lines
|
|
if (task && task->done) {
|
|
task->thread.join();
|
|
if (task->result) {
|
|
if (prepared) releasePair(c, prepared);
|
|
prepared = task->result;
|
|
if (!prepared->raceLine.empty()) emit(prepared->raceLine);
|
|
emit(fmt("prepared %s %s %.1f %s resident 2 programs 2 datasets", prepared->epochHex.c_str(), prepared->dayHex.c_str(), wallMs() - task->t0, pairSummary(prepared).c_str()));
|
|
} else {
|
|
emit(fmt("prepare-failed %s %s %s", task->epochHex.c_str(), task->dayHex.c_str(), task->error.c_str()));
|
|
}
|
|
delete task; task = nullptr;
|
|
}
|
|
std::string wantClass, wantEra;
|
|
takeClassTokens(f, wantClass, wantEra);
|
|
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->wantClass = wantClass; task->wantEra = wantEra;
|
|
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 (!pairIsClass(cur, f[6], f[7], wantClass, wantEra) && !task && !pairIsClass(prepared, f[6], f[7], wantClass, wantEra)) {
|
|
// Self-heal: a job on seeds this worker has no pair for and no prepare in flight (a prepare failed, or
|
|
// the miner never sent one). The miner writes a pack per pair under its --prepare-packs root; find it by
|
|
// seeds.txt and build it now, in the foreground. The miner only re-sends prepare for the pair after this one.
|
|
std::string dir = findPackFor(prepareRoot, f[6], f[7]);
|
|
if (dir.empty()) dir = findPackFor(parentDir(o.pack) + "/prepare", f[6], f[7]);
|
|
if (dir.empty()) dir = findPackFor(parentDir(o.pack), f[6], f[7]);
|
|
if (!dir.empty()) {
|
|
info(fmt("job %s is for epoch %.16s day %s, which is not resident; building its pack %s now (foreground)", jobId.c_str(), f[6].c_str(), f[7].c_str(), dir.c_str()));
|
|
std::string berr;
|
|
Pair* p = buildPair(c, dir, nullptr, berr, true);
|
|
if (p && (p->epochHex != f[6] || p->dayHex != f[7])) { berr = "the pack in " + dir + " is for other seeds"; releasePair(c, p); p = nullptr; }
|
|
if (p) { if (prepared) releasePair(c, prepared); prepared = p; info(fmt("built %s: %s", dir.c_str(), pairSummary(p).c_str())); if (!p->raceLine.empty()) emit(p->raceLine); }
|
|
else emit("error " + jobId + " could not build " + dir + ": " + berr);
|
|
}
|
|
}
|
|
if (!pairIsClass(cur, f[6], f[7], wantClass, wantEra)) {
|
|
char why[256];
|
|
if (pairIsClass(prepared, f[6], f[7], wantClass, wantEra)) {
|
|
if (old) releasePair(c, old);
|
|
old = cur; cur = prepared; prepared = nullptr; switched = true;
|
|
info(fmt("switched to the prepared pair epoch %.16s day %s (class %s) in %.2f ms", cur->epochHex.c_str(), cur->dayHex.c_str(), cur->programClass.c_str(), wallMs() - t0));
|
|
} else if (pairIs(cur, f[6], f[7]) && !pf_pack_class_ok(cur->programClass.c_str(), cur->eraHex.c_str(), wantClass.c_str(), wantEra.c_str(), why, sizeof(why))) {
|
|
// Counter ASIC 2.0: right seeds, wrong class or era. The miner prepares the pair again from a pack of
|
|
// the class the chain is on (the `need` line), and the pack of the other class is never mined.
|
|
emit(fmt("need %s %s", f[6].c_str(), f[7].c_str()));
|
|
emit(fmt("error %s pack %s: %s", jobId.c_str(), cur->dir.c_str(), why));
|
|
continue;
|
|
} else if (!hexEq(cur->epochHex, f[6])) {
|
|
emit(fmt("need %s %s", f[6].c_str(), f[7].c_str())); // the miner prepares this pair (4 October 2026)
|
|
emit(fmt("error %s epoch seed mismatch: this worker holds epoch %.16s (program words %08x %08x ...)%s, the job is for epoch %.16s (bare seed words %08x %08x ...); send prepare with a pack directory",
|
|
jobId.c_str(), cur->epochHex.c_str(), cur->sw[0], cur->sw[1], prepared ? " plus one prepared pair" : "", f[6].c_str(), sw[0], sw[1]));
|
|
continue;
|
|
} else {
|
|
emit(fmt("need %s %s", f[6].c_str(), f[7].c_str()));
|
|
emit(fmt("error %s day seed mismatch: this worker's cache is for key %08x %08x ..., the job's day seed %s gives %08x %08x ...; send prepare with a pack directory",
|
|
jobId.c_str(), cur->kw[0], cur->kw[1], f[7].c_str(), kw[0], kw[1]));
|
|
continue;
|
|
}
|
|
}
|
|
uint64_t remaining = nonceCount, hashes = 0;
|
|
uint32_t hi = (uint32_t)(nonceStart >> 32), lo = (uint32_t)nonceStart;
|
|
bool failed = false;
|
|
while (remaining > 0) {
|
|
uint64_t room = (uint64_t)(0xffffffffu - lo) + 1ull;
|
|
uint64_t chunk64 = remaining < batch ? remaining : batch;
|
|
if (chunk64 > room) chunk64 = room;
|
|
uint32_t chunk = (uint32_t)chunk64;
|
|
uint32_t iw[8];
|
|
{
|
|
uint8_t b[49];
|
|
std::memcpy(b, "igneum-block/", 13);
|
|
std::memcpy(b + 13, prehash, 32);
|
|
b[45] = (uint8_t)hi; b[46] = (uint8_t)(hi >> 8); b[47] = (uint8_t)(hi >> 16); b[48] = (uint8_t)(hi >> 24);
|
|
pf_seed_words_from_bytes(b, 49, iw);
|
|
}
|
|
// A chunk that is not a multiple of the block is finished one 32-lane block at a time. The block is the
|
|
// pair's (its winning variant's). The mutex gives a race its exclusive windows between chunks.
|
|
std::lock_guard<std::mutex> hold(gpuMutex);
|
|
uint32_t block = 32u * (uint32_t)cur->blockWarps;
|
|
uint32_t main = chunk - (chunk % block);
|
|
CUresult r = CUDA_SUCCESS;
|
|
if (main > 0 && !launchHash(c, cur, dOut, lo, iw, main, block, nullptr, err)) { emit("error " + jobId + " dispatch failed: " + err); failed = true; break; }
|
|
if (main < chunk) {
|
|
for (uint32_t off = main; off < chunk && !failed; off += 32u) if (!launchHash(c, cur, dOut + (CUdeviceptr)off * 8u, lo + off, iw, 32u, 32u, nullptr, err)) { emit("error " + jobId + " dispatch failed: " + err); failed = true; }
|
|
if (failed) break;
|
|
}
|
|
r = c.drv.streamSynchronize(nullptr);
|
|
if (r == CUDA_SUCCESS) r = c.drv.memcpyDtoH(hOut.data(), dOut, (size_t)chunk * 8u);
|
|
if (r != CUDA_SUCCESS) { emit("error " + jobId + " dispatch failed: " + c.err(r)); failed = true; break; }
|
|
for (uint32_t i = 0; i < chunk; ++i) if (hOut[i] <= target) {
|
|
uint64_t nonce = ((uint64_t)hi << 32) | (uint64_t)(uint32_t)(lo + i);
|
|
std::printf("found %s %llu %016llx\n", jobId.c_str(), (unsigned long long)nonce, (unsigned long long)hOut[i]);
|
|
}
|
|
std::fflush(stdout);
|
|
hashes += chunk;
|
|
remaining -= chunk;
|
|
if (chunk64 == room) { hi += 1u; lo = 0u; } else lo += chunk;
|
|
}
|
|
if (failed) continue;
|
|
emit(fmt("done %s %llu %.2f", jobId.c_str(), (unsigned long long)hashes, wallMs() - t0));
|
|
if (switched && old) { releasePair(c, old); old = nullptr; info("dropped the previous pair (its program, cache and dataset)"); }
|
|
}
|
|
if (task) { task->thread.join(); if (task->result) releasePair(c, task->result); delete task; }
|
|
c.drv.memFree(dOut);
|
|
if (old) releasePair(c, old);
|
|
if (prepared) releasePair(c, prepared);
|
|
releasePair(c, cur);
|
|
return 0;
|
|
}
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Main
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Read-width experiment (5 October 2026, docs/plans/read-width.md): --bench and --memprobe
|
|
|
|
static uint64_t fnv1a64Bytes(const void* p, size_t n) {
|
|
const uint8_t* b = (const uint8_t*)p;
|
|
uint64_t h = 0xcbf29ce484222325ull;
|
|
for (size_t i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
|
|
return h;
|
|
}
|
|
|
|
// --bench: the pair is built and self-tested (vectors through the bound kernel with the seed words); then a warm-up
|
|
// dispatch at base nonce 0 (fingerprinted) and --batches timed dispatches of 2^B nonces, wall time around
|
|
// cuStreamSynchronize (the driver API path loads no event symbols; a 2^24 dispatch is 60 to 900 ms on the cards here,
|
|
// so the launch overhead is under 1 percent).
|
|
static int runBench(Ctx& c, const Options& o, Pair* p) {
|
|
uint32_t nonces = 1u << o.batchLog2, block = 32u * (uint32_t)p->blockWarps;
|
|
if (p->persistent) { uint32_t unit = 32u * (uint32_t)p->warps; nonces = (nonces / unit) * unit; if (nonces == 0) nonces = unit; }
|
|
CUdeviceptr dOut = 0;
|
|
std::string err;
|
|
if (c.drv.memAlloc(&dOut, (size_t)nonces * 8u) != CUDA_SUCCESS) { std::printf("FAIL: cuMemAlloc out\n"); return 2; }
|
|
std::vector<uint64_t> hOut(nonces);
|
|
double sum = 0, warm = 0;
|
|
uint64_t fp = 0;
|
|
for (int b = -1; b < o.batches; ++b) {
|
|
double t0 = wallMs();
|
|
if (!launchHash(c, p, dOut, (uint32_t)(b + 1) * nonces, p->sw, nonces, block, nullptr, err)) { std::printf("FAIL: %s\n", err.c_str()); return 2; }
|
|
CUresult r = c.drv.streamSynchronize(nullptr);
|
|
if (r != CUDA_SUCCESS) { std::printf("FAIL: dispatch %d: %s\n", b, c.err(r).c_str()); return 2; }
|
|
double ms = wallMs() - t0;
|
|
if (b < 0) {
|
|
warm = ms;
|
|
if (c.drv.memcpyDtoH(hOut.data(), dOut, (size_t)nonces * 8u) != CUDA_SUCCESS) { std::printf("FAIL: read-back\n"); return 2; }
|
|
fp = fnv1a64Bytes(hOut.data(), (size_t)nonces * 8u);
|
|
} else sum += ms;
|
|
}
|
|
c.drv.memFree(dOut);
|
|
std::string dev = c.name; for (char& ch : dev) if (ch == ' ') ch = '_';
|
|
std::printf("warm-up dispatch (base 0): %.2f ms; %d timed dispatches of %u nonces: mean %.2f ms\n", warm, o.batches, nonces, sum / o.batches);
|
|
std::printf("RESULT pack=%s class=%s device=%s arch=%s regs=%d blocks_per_sm=%d warps=%d resident=%d arena_mib=%llu hot_mib=%u hot_slots=%u hot_fill_ms=%.2f nonces=%u batches=%d check=%s fingerprint=%016llx mhs=%.3f loads=%u bytes=%u scratch_ops=%u variant=%s sparse_blocks=%d block_warps=%d time=wall\n",
|
|
p->dir.c_str(), p->loadClass.c_str(), dev.c_str(), c.archOpt.c_str(), p->regs, p->blocksPerSM, p->warps, p->residentWarps, (unsigned long long)(p->scratchBytes >> 20), p->hotMb, p->hotSlots, p->hotMs, nonces, o.batches,
|
|
p->checked ? (p->checkPass ? "PASS" : "FAIL") : "skipped", (unsigned long long)fp, (double)nonces * (double)o.batches / (sum / 1000.0) / 1e6,
|
|
p->loadsPerHash, p->bytesPerHash, p->scratchOps * 8u, p->variant.c_str(), p->sparseBlocks, p->blockWarps);
|
|
if (!o.pinned.empty() && p->variant != o.pinned) std::printf("RESULT variant_not_installed requested=%s served=%s race=\"%s\"\n", o.pinned.c_str(), p->variant.c_str(), p->raceLine.c_str());
|
|
return 0;
|
|
}
|
|
|
|
// --memprobe: the OpenCL worker's table (proto-opencl/host.c, 5 October 2026) in CUDA C through NVRTC, so the two
|
|
// vendors are probed with the same access patterns: a dependent chain of random 4-byte reads (the hash's pattern),
|
|
// eight independent chains per lane, dependent random 16-byte and 64-byte reads, a coalesced stream and an integer
|
|
// chain, at 4, 64 and 1024 MiB. Wall time around cuStreamSynchronize, best of 3, a fresh seed per repetition.
|
|
static const char* PROBE_CUDA =
|
|
"#include <cstdint>\n"
|
|
"__device__ __forceinline__ uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
|
|
"extern \"C\" __global__ void probe_fill(uint32_t* ds, uint32_t n) { uint32_t i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) ds[i] = pm_mix(i ^ 0x9E3779B9u); }\n"
|
|
"extern \"C\" __global__ void probe_chase(const uint32_t* ds, uint32_t mask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
|
|
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) x = ds[x & mask] ^ (x * 0x9E3779B1u + s);\n"
|
|
" out[g] = x;\n"
|
|
"}\n"
|
|
"extern \"C\" __global__ void probe_indep(const uint32_t* ds, uint32_t mask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
|
|
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;\n"
|
|
" uint32_t x0 = pm_mix(g * 8u ^ seed), x1 = pm_mix((g * 8u + 1u) ^ seed), x2 = pm_mix((g * 8u + 2u) ^ seed), x3 = pm_mix((g * 8u + 3u) ^ seed);\n"
|
|
" uint32_t x4 = pm_mix((g * 8u + 4u) ^ seed), x5 = pm_mix((g * 8u + 5u) ^ seed), x6 = pm_mix((g * 8u + 6u) ^ seed), x7 = pm_mix((g * 8u + 7u) ^ seed);\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) {\n"
|
|
" x0 = ds[x0 & mask] ^ (x0 * 0x9E3779B1u + s); x1 = ds[x1 & mask] ^ (x1 * 0x9E3779B1u + s);\n"
|
|
" x2 = ds[x2 & mask] ^ (x2 * 0x9E3779B1u + s); x3 = ds[x3 & mask] ^ (x3 * 0x9E3779B1u + s);\n"
|
|
" x4 = ds[x4 & mask] ^ (x4 * 0x9E3779B1u + s); x5 = ds[x5 & mask] ^ (x5 * 0x9E3779B1u + s);\n"
|
|
" x6 = ds[x6 & mask] ^ (x6 * 0x9E3779B1u + s); x7 = ds[x7 & mask] ^ (x7 * 0x9E3779B1u + s);\n"
|
|
" }\n"
|
|
" out[g] = x0 ^ x1 ^ x2 ^ x3 ^ x4 ^ x5 ^ x6 ^ x7;\n"
|
|
"}\n"
|
|
"extern \"C\" __global__ void probe_line16(const uint4* ds, uint32_t vecMask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
|
|
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) { uint4 a = ds[x & vecMask]; x = (a.x ^ a.y ^ a.z ^ a.w) ^ (x * 0x9E3779B1u + s); }\n"
|
|
" out[g] = x;\n"
|
|
"}\n"
|
|
"extern \"C\" __global__ void probe_line(const uint4* ds, uint32_t lineMask, uint32_t steps, uint32_t seed, uint32_t* out) {\n"
|
|
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed);\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) { uint32_t l = (x & lineMask) * 4u; uint4 a = ds[l], b = ds[l + 1u], c = ds[l + 2u], d = ds[l + 3u]; x = (a.x ^ b.y ^ c.z ^ d.w) ^ (x * 0x9E3779B1u + s); }\n"
|
|
" out[g] = x;\n"
|
|
"}\n"
|
|
"extern \"C\" __global__ void probe_stream(const uint4* ds, uint32_t perLane, uint32_t* out) {\n"
|
|
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x, n = gridDim.x * blockDim.x; uint4 acc = make_uint4(0u, 0u, 0u, 0u);\n"
|
|
" for (uint32_t s = 0u; s < perLane; ++s) { uint4 v = ds[s * n + g]; acc.x ^= v.x; acc.y ^= v.y; acc.z ^= v.z; acc.w ^= v.w; }\n"
|
|
" out[g] = acc.x ^ acc.y ^ acc.z ^ acc.w;\n"
|
|
"}\n"
|
|
"extern \"C\" __global__ void probe_alu(uint32_t steps, uint32_t seed, uint32_t* out) {\n"
|
|
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + ((y << 7u) | (y >> 25u)); y = (y ^ x) + s; }\n"
|
|
" out[g] = x ^ y;\n"
|
|
"}\n";
|
|
|
|
static double probeLaunch(Ctx& c, CUfunction f, size_t lanes, size_t local, int reps, int seedArg, uint32_t seed, void** args) {
|
|
double best = -1;
|
|
for (int r = 0; r < reps; ++r) {
|
|
uint32_t s = seed + (uint32_t)r * 0x9E3779B9u;
|
|
if (seedArg >= 0) args[seedArg] = &s;
|
|
double t0 = wallMs();
|
|
if (c.drv.launchKernel(f, (unsigned)(lanes / local), 1, 1, (unsigned)local, 1, 1, 0, nullptr, args, nullptr) != CUDA_SUCCESS) return -1;
|
|
if (c.drv.streamSynchronize(nullptr) != CUDA_SUCCESS) return -1;
|
|
double ms = wallMs() - t0;
|
|
if (best < 0 || ms < best) best = ms;
|
|
}
|
|
return best;
|
|
}
|
|
|
|
static int runMemprobe(Ctx& c, const Options& o) {
|
|
Compiled cp;
|
|
std::string err;
|
|
if (!rtcCompile(c, PROBE_CUDA, "probe.cu", "", "", {}, cp, err)) { std::printf("memprobe: build FAILED: %s\n", err.c_str()); return 2; }
|
|
CUmodule mod = nullptr;
|
|
if (c.drv.moduleLoadData(&mod, cp.image.data()) != CUDA_SUCCESS) { std::printf("memprobe: cuModuleLoadData failed\n"); return 2; }
|
|
CUfunction kFill, kChase, kIndep, kLine16, kLine, kStream, kAlu;
|
|
const char* names[7] = { "probe_fill", "probe_chase", "probe_indep", "probe_line16", "probe_line", "probe_stream", "probe_alu" };
|
|
CUfunction* fns[7] = { &kFill, &kChase, &kIndep, &kLine16, &kLine, &kStream, &kAlu };
|
|
for (int i = 0; i < 7; ++i) if (c.drv.moduleGetFunction(fns[i], mod, names[i]) != CUDA_SUCCESS) { std::printf("memprobe: %s not in the module\n", names[i]); return 2; }
|
|
int sizes[3] = { 4, 64, 1024 }, nSizes = 3;
|
|
if (o.probeMib > 0) { sizes[0] = o.probeMib; nSizes = 1; }
|
|
const size_t lanesList[8] = { 256, 1024, 1u << 12, 1u << 14, 1u << 16, 1u << 18, 1u << 20, 1u << 22 };
|
|
const size_t groups[2] = { 32, 256 };
|
|
const uint32_t STEPS = 256u, ALU_STEPS = 4096u;
|
|
const size_t maxLanes = 1u << 22;
|
|
CUdeviceptr dOut = 0;
|
|
if (c.drv.memAlloc(&dOut, maxLanes * 4u) != CUDA_SUCCESS) { std::printf("memprobe: cuMemAlloc out\n"); return 2; }
|
|
std::printf("memprobe on %s (sm_%d%d, %d SMs, driver %d.%d, NVRTC %d.%d), wall time around cuStreamSynchronize\n", c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor);
|
|
std::printf("| probe | MiB | block | lanes in flight | steps per lane | best ms | G loads/s | ns per dependent load |\n|---|---|---|---|---|---|---|---|\n");
|
|
for (int si = 0; si < nSizes; ++si) {
|
|
int mib = sizes[si];
|
|
uint64_t bytes = (uint64_t)mib << 20;
|
|
uint32_t words = (uint32_t)(bytes / 4ull), mask = words - 1u, n = words;
|
|
CUdeviceptr dDs = 0;
|
|
if (c.drv.memAlloc(&dDs, (size_t)bytes) != CUDA_SUCCESS) { std::printf("| chase | %d | skipped: cuMemAlloc failed | | | | | |\n", mib); continue; }
|
|
{ void* a[2] = { &dDs, &n }; probeLaunch(c, kFill, ((size_t)words + 255) / 256 * 256, 256, 1, -1, 0, a); }
|
|
for (int gi = 0; gi < 2; ++gi) {
|
|
size_t local = groups[gi];
|
|
for (int li = 0; li < 8; ++li) {
|
|
size_t lanes = lanesList[li];
|
|
if (lanes < local) continue;
|
|
uint32_t seed = 0x1234567u + (uint32_t)li * 977u, steps = STEPS;
|
|
void* a[5] = { &dDs, &mask, &steps, &seed, &dOut };
|
|
double ms = probeLaunch(c, kChase, lanes, local, 3, 3, seed, a);
|
|
std::printf("| chase | %d | %zu | %zu | %u | %.3f | %.3f | %.0f |\n", mib, local, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, ms * 1e6 / STEPS);
|
|
std::fflush(stdout);
|
|
}
|
|
}
|
|
for (size_t lanes = 1u << 16; lanes <= maxLanes; lanes <<= 2) {
|
|
uint32_t seed = 0x7654321u, steps = STEPS;
|
|
void* a[5] = { &dDs, &mask, &steps, &seed, &dOut };
|
|
double ms = probeLaunch(c, kIndep, lanes, 256, 3, 3, seed, a);
|
|
std::printf("| indep x8 | %d | 256 | %zu | %u | %.3f | %.3f | (8 loads in flight per lane) |\n", mib, lanes, STEPS, ms, (double)lanes * 8.0 * STEPS / (ms / 1000.0) / 1e9);
|
|
}
|
|
for (size_t lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
|
|
uint32_t vecMask = (words / 4u) - 1u, seed = 0x2718281u, steps = STEPS;
|
|
void* a[5] = { &dDs, &vecMask, &steps, &seed, &dOut };
|
|
double ms = probeLaunch(c, kLine16, lanes, 256, 3, 3, seed, a);
|
|
std::printf("| line 16 B | %d | 256 | %zu | %u | %.3f | %.3f G reads/s | %.1f GB/s in 16 B reads |\n", mib, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, (double)lanes * STEPS * 16.0 / (ms / 1000.0) / 1e9);
|
|
}
|
|
for (size_t lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
|
|
uint32_t lineMask = (words / 16u) - 1u, seed = 0x3141592u, steps = STEPS;
|
|
void* a[5] = { &dDs, &lineMask, &steps, &seed, &dOut };
|
|
double ms = probeLaunch(c, kLine, lanes, 256, 3, 3, seed, a);
|
|
std::printf("| line 64 B | %d | 256 | %zu | %u | %.3f | %.3f G lines/s | %.1f GB/s in lines |\n", mib, lanes, STEPS, ms, (double)lanes * STEPS / (ms / 1000.0) / 1e9, (double)lanes * STEPS * 64.0 / (ms / 1000.0) / 1e9);
|
|
}
|
|
{
|
|
size_t lanes = 1u << 20;
|
|
uint32_t perLane = (uint32_t)((uint64_t)words / 4ull / (uint64_t)lanes);
|
|
if (perLane == 0) { perLane = 1; lanes = (size_t)words / 4u; }
|
|
double bytesRead = (double)perLane * (double)lanes * 16.0;
|
|
void* a[3] = { &dDs, &perLane, &dOut };
|
|
double ms = probeLaunch(c, kStream, lanes, 256, 3, -1, 0, a);
|
|
std::printf("| stream | %d | 256 | %zu | %u | %.3f | %.1f GB/s coalesced | (%.0f MiB read once) |\n", mib, lanes, perLane, ms, bytesRead / (ms / 1000.0) / 1e9, bytesRead / 1048576.0);
|
|
}
|
|
c.drv.memFree(dDs);
|
|
std::fflush(stdout);
|
|
}
|
|
{
|
|
size_t lanes = 1u << 20;
|
|
uint32_t seed = 0x2468aceu, steps = ALU_STEPS;
|
|
void* a[3] = { &steps, &seed, &dOut };
|
|
double ms = probeLaunch(c, kAlu, lanes, 256, 3, 1, seed, a);
|
|
double ops = (double)lanes * ALU_STEPS * 5.0;
|
|
std::printf("| alu | 0 | 256 | %zu | %u | %.3f | %.1f G int ops/s | %.3f G steps/s per SM (approximate: 5 ops per step counted) |\n", lanes, ALU_STEPS, ms, ops / (ms / 1000.0) / 1e9, (double)lanes * ALU_STEPS / (ms / 1000.0) / 1e9 / (c.sms ? c.sms : 1));
|
|
}
|
|
c.drv.memFree(dOut);
|
|
c.drv.moduleUnload(mod);
|
|
std::printf("memprobe: done\n");
|
|
return 0;
|
|
}
|
|
|
|
|
|
// ---------------------------------------------------------------------------------------------
|
|
// Counter ASIC 4.0 research, 7 October 2026: --microbench. One kernel per GPU hardware block that gaming and AI
|
|
// already paid for, each run sustained for --mb-seconds at full residency so a 1 Hz power sampler reads its watts, with
|
|
// the counted operations per second beside it, so the energy per operation on this card is (watts minus the sleeping
|
|
// floor) over operations per second. Every probe prints one RESULT line with UTC start and end stamps and a checksum of
|
|
// its outputs. Probes compile one at a time, so a form the card or NVRTC refuses (an FP8 tile on a card under sm_89,
|
|
// a texture the driver cannot bind) drops that probe alone with its error on the row. Nothing here touches a pack.
|
|
|
|
struct MicroProbe {
|
|
const char* name; // the row's name
|
|
const char* unit; // what one counted operation is
|
|
double opsPerStep; // counted operations per lane per step
|
|
uint32_t steps; // steps per launch (sized so one launch is 50 ms to 2 s on a 5090, approximate)
|
|
int tableMib; // a table of uint32 the kernel reads (0 = none); filled by kb_fill mode 0 (or 1 = floats in [0,1))
|
|
int tableMode;
|
|
int tex; // 0 none; 1 a point-sampled u32 texture over the table; 2 a linearly filtered float texture over it
|
|
const char* body; // the kernel body: has steps, seed, out, tbl, mask, tex, g (the lane), x0..x3 (mixed seeds)
|
|
};
|
|
|
|
static const char* MICRO_PRELUDE =
|
|
"#include <cstdint>\n"
|
|
"__device__ __forceinline__ uint32_t kb_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
|
|
"__device__ __forceinline__ uint32_t kb_rotl(uint32_t x, uint32_t r) { return (x << r) | (x >> (32u - r)); }\n"
|
|
"extern \"C\" __global__ void kb_fill(uint32_t* t, uint32_t n, uint32_t mode) { uint32_t i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= n) return;\n"
|
|
" uint32_t v = kb_mix(i ^ 0x9E3779B9u); if (mode == 1u) { float f = (float)(v >> 8) * (1.0f / 16777216.0f); t[i] = __float_as_uint(f); } else t[i] = v; }\n"
|
|
"#define KB_KERNEL(NAME) extern \"C\" __global__ void NAME(uint32_t steps, uint32_t seed, uint32_t* out, const uint32_t* tbl, uint32_t mask, unsigned long long tex) {\\\n"
|
|
" uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t x0 = kb_mix(g * 4u ^ seed), x1 = kb_mix((g * 4u + 1u) ^ seed), x2 = kb_mix((g * 4u + 2u) ^ seed), x3 = kb_mix((g * 4u + 3u) ^ seed);\n";
|
|
|
|
// Tensor tiles: the dependency carries the accumulator back into the A fragment so the chain is real; two tiles per step for issue overlap.
|
|
#define KB_MMA2(INSTR, NA, NB, NC, ATY) \
|
|
" uint32_t a0 = x0, a1 = x1, a2 = x2, a3 = x3, b0 = x1 ^ 0x5bd1e995u, b1 = x2 * 0x27d4eb2fu; " ATY " c0 = 0, c1 = 0, c2 = 0, c3 = 0, e0 = 0, e1 = 0, e2 = 0, e3 = 0;\n" \
|
|
" for (uint32_t s = 0u; s < steps; ++s) {\n" \
|
|
" asm volatile(\"" INSTR "\" : \"=r\"(c0), \"=r\"(c1), \"=r\"(c2), \"=r\"(c3) : \"r\"(a0), \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(b0), \"r\"(b1), \"r\"(c0), \"r\"(c1), \"r\"(c2), \"r\"(c3));\n" \
|
|
" asm volatile(\"" INSTR "\" : \"=r\"(e0), \"=r\"(e1), \"=r\"(e2), \"=r\"(e3) : \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(a0), \"r\"(b1), \"r\"(b0), \"r\"(e0), \"r\"(e1), \"r\"(e2), \"r\"(e3));\n" \
|
|
" a0 ^= (uint32_t)c0 + s; a1 ^= (uint32_t)e1; a2 += (uint32_t)c2; a3 ^= (uint32_t)e3 * 0x9E3779B1u;\n" \
|
|
" }\n" \
|
|
" out[g] = (uint32_t)c0 ^ (uint32_t)c1 ^ (uint32_t)c2 ^ (uint32_t)c3 ^ (uint32_t)e0 ^ (uint32_t)e1 ^ (uint32_t)e2 ^ (uint32_t)e3;\n}\n"
|
|
|
|
static const MicroProbe MICRO_PROBES[] = {
|
|
{ "sleep", "none (the SM-resident floor: full occupancy, __nanosleep, no issue)", 0, 2000, 0, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { __nanosleep(1000); x0 += s; }\n out[g] = x0;\n}\n" },
|
|
{ "int_arx", "int32 add, xor or rotate (4 independent chains, 3 ops each per step)", 12, 16384, 0, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { x0 = kb_rotl(x0 + x1, 7u) ^ s; x1 = kb_rotl(x1 + x2, 13u) ^ x0; x2 = kb_rotl(x2 + x3, 17u) ^ x1; x3 = kb_rotl(x3 + x0, 23u) ^ x2; }\n out[g] = x0 ^ x1 ^ x2 ^ x3;\n}\n" },
|
|
{ "int_mul", "int32 multiply-add (4 independent chains)", 4, 16384, 0, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { x0 = x0 * 0x9E3779B1u + s; x1 = x1 * 0x85EBCA77u + x0; x2 = x2 * 0xC2B2AE3Du + x1; x3 = x3 * 0x27D4EB2Fu + x2; }\n out[g] = x0 ^ x1 ^ x2 ^ x3;\n}\n" },
|
|
{ "int_mulhi", "int32 high multiply (4 independent chains)", 4, 16384, 0, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { x0 = __umulhi(x0, 0x9E3779B1u) + s; x1 = __umulhi(x1, x0 | 1u) ^ x1; x2 = __umulhi(x2, 0xC2B2AE3Du) + x1; x3 = __umulhi(x3, x2 | 1u) ^ x3; }\n out[g] = x0 ^ x1 ^ x2 ^ x3;\n}\n" },
|
|
{ "prmt", "byte permute, __byte_perm (4 independent chains)", 4, 16384, 0, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { x0 = __byte_perm(x0, x1, 0x2103u ^ (s & 7u)); x1 = __byte_perm(x1, x2, 0x3012u); x2 = __byte_perm(x2, x3, 0x1230u); x3 = __byte_perm(x3, x0, 0x0321u) + s; }\n out[g] = x0 ^ x1 ^ x2 ^ x3;\n}\n" },
|
|
{ "lop3", "three-input logic (and, or, xor fused to one LOP3; 4 independent chains)", 4, 16384, 0, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { x0 = (x0 & x1) ^ (x2 | s); x1 = (x1 & x2) ^ (x3 | x0); x2 = (x2 & x3) ^ (x0 | x1); x3 = (x3 & x0) ^ (x1 | x2); }\n out[g] = x0 ^ x1 ^ x2 ^ x3;\n}\n" },
|
|
{ "shfl", "warp shuffle, __shfl_xor_sync (4 independent chains, one xor each beside it)", 4, 8192, 0, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { x0 = __shfl_xor_sync(0xffffffffu, x0, (int)((s & 31u) | 1u)) ^ x1; x1 = __shfl_xor_sync(0xffffffffu, x1, 2) ^ x2; x2 = __shfl_xor_sync(0xffffffffu, x2, 4) ^ x3; x3 = __shfl_xor_sync(0xffffffffu, x3, 8) ^ x0; }\n out[g] = x0 ^ x1 ^ x2 ^ x3;\n}\n" },
|
|
{ "fp32_fma", "FP32 fused multiply-add (4 independent chains)", 4, 16384, 0, 0, 0,
|
|
" float f0 = __uint_as_float((x0 & 0x007fffffu) | 0x3f800000u), f1 = __uint_as_float((x1 & 0x007fffffu) | 0x3f800000u), f2 = __uint_as_float((x2 & 0x007fffffu) | 0x3f800000u), f3 = __uint_as_float((x3 & 0x007fffffu) | 0x3f800000u);\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) { f0 = fmaf(f0, 0.999f, 0.5f); f1 = fmaf(f1, 1.001f, -0.5f); f2 = fmaf(f2, 0.998f, 0.25f); f3 = fmaf(f3, 1.002f, -0.25f); }\n out[g] = __float_as_uint(f0) ^ __float_as_uint(f1) ^ __float_as_uint(f2) ^ __float_as_uint(f3);\n}\n" },
|
|
{ "fp16x2_fma", "FP16 fused multiply-add, two per fma.rn.f16x2 (4 independent chains)", 8, 16384, 0, 0, 0,
|
|
" uint32_t h0 = (x0 & 0x03ff03ffu) | 0x3c003c00u, h1 = (x1 & 0x03ff03ffu) | 0x3c003c00u, h2 = (x2 & 0x03ff03ffu) | 0x3c003c00u, h3 = (x3 & 0x03ff03ffu) | 0x3c003c00u;\n"
|
|
" const uint32_t m = 0x3bff3bffu, a = 0x38003800u;\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) { asm volatile(\"fma.rn.f16x2 %0, %0, %1, %2;\" : \"+r\"(h0) : \"r\"(m), \"r\"(a)); asm volatile(\"fma.rn.f16x2 %0, %0, %1, %2;\" : \"+r\"(h1) : \"r\"(m), \"r\"(a)); asm volatile(\"fma.rn.f16x2 %0, %0, %1, %2;\" : \"+r\"(h2) : \"r\"(m), \"r\"(a)); asm volatile(\"fma.rn.f16x2 %0, %0, %1, %2;\" : \"+r\"(h3) : \"r\"(m), \"r\"(a)); }\n"
|
|
" out[g] = h0 ^ h1 ^ h2 ^ h3;\n}\n" },
|
|
{ "mma_u8_m8n8k16", "int8 multiply-add inside mma.sync m8n8k16 u8 (1,024 per tile per warp, 32 per lane; two tiles per step)", 64, 8192, 0, 0, 0,
|
|
" uint32_t a0 = x0, b0 = x1, c0 = 0, c1 = 0, a1 = x2, b1 = x3, e0 = 0, e1 = 0;\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) {\n"
|
|
" asm volatile(\"mma.sync.aligned.m8n8k16.row.col.s32.u8.u8.s32 {%0,%1}, {%2}, {%3}, {%4,%5};\" : \"=r\"(c0), \"=r\"(c1) : \"r\"(a0), \"r\"(b0), \"r\"(c0), \"r\"(c1));\n"
|
|
" asm volatile(\"mma.sync.aligned.m8n8k16.row.col.s32.u8.u8.s32 {%0,%1}, {%2}, {%3}, {%4,%5};\" : \"=r\"(e0), \"=r\"(e1) : \"r\"(a1), \"r\"(b1), \"r\"(e0), \"r\"(e1));\n"
|
|
" a0 ^= c0 + s; b0 = kb_rotl(b0, 7u) ^ c1; a1 ^= e1; b1 = kb_rotl(b1, 5u) ^ e0;\n }\n out[g] = c0 ^ c1 ^ e0 ^ e1;\n}\n" },
|
|
{ "mma_s8_m16n8k32", "int8 multiply-add inside mma.sync m16n8k32 s8 (4,096 per tile per warp, 128 per lane; two tiles per step)", 256, 4096, 0, 0, 0,
|
|
KB_MMA2("mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};", 4, 2, 4, "int32_t") },
|
|
{ "mma_f16_m16n8k16", "FP16 multiply-add with FP32 accumulate inside mma.sync m16n8k16 (2,048 per tile per warp, 64 per lane; two tiles per step)", 128, 4096, 0, 0, 0,
|
|
" uint32_t a0 = (x0 & 0x03ff03ffu) | 0x3c003c00u, a1 = (x1 & 0x03ff03ffu) | 0x3c003c00u, a2 = (x2 & 0x03ff03ffu) | 0x3c003c00u, a3 = (x3 & 0x03ff03ffu) | 0x3c003c00u, b0 = 0x3c003c00u ^ (x1 & 0x00ff00ffu), b1 = 0x3c003c00u ^ (x2 & 0x00ff00ffu);\n"
|
|
" float c0 = 0.f, c1 = 0.f, c2 = 0.f, c3 = 0.f, e0 = 0.f, e1 = 0.f, e2 = 0.f, e3 = 0.f;\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) {\n"
|
|
" asm volatile(\"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\" : \"=f\"(c0), \"=f\"(c1), \"=f\"(c2), \"=f\"(c3) : \"r\"(a0), \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(b0), \"r\"(b1), \"f\"(c0), \"f\"(c1), \"f\"(c2), \"f\"(c3));\n"
|
|
" asm volatile(\"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\" : \"=f\"(e0), \"=f\"(e1), \"=f\"(e2), \"=f\"(e3) : \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(a0), \"r\"(b1), \"r\"(b0), \"f\"(e0), \"f\"(e1), \"f\"(e2), \"f\"(e3));\n"
|
|
" a0 = ((a0 ^ __float_as_uint(c0)) & 0x03ff03ffu) | 0x3c003c00u; a2 = ((a2 ^ __float_as_uint(e2)) & 0x03ff03ffu) | 0x3c003c00u; c0 *= 0.5f; c1 *= 0.5f; c2 *= 0.5f; c3 *= 0.5f; e0 *= 0.5f; e1 *= 0.5f; e2 *= 0.5f; e3 *= 0.5f;\n }\n"
|
|
" out[g] = __float_as_uint(c0) ^ __float_as_uint(c1) ^ __float_as_uint(c2) ^ __float_as_uint(c3) ^ __float_as_uint(e0) ^ __float_as_uint(e1) ^ __float_as_uint(e2) ^ __float_as_uint(e3);\n}\n" },
|
|
{ "mma_bf16_m16n8k16", "BF16 multiply-add with FP32 accumulate inside mma.sync m16n8k16 (64 per lane; two tiles per step)", 128, 4096, 0, 0, 0,
|
|
" uint32_t a0 = (x0 & 0x007f007fu) | 0x3f803f80u, a1 = (x1 & 0x007f007fu) | 0x3f803f80u, a2 = (x2 & 0x007f007fu) | 0x3f803f80u, a3 = (x3 & 0x007f007fu) | 0x3f803f80u, b0 = 0x3f803f80u ^ (x1 & 0x003f003fu), b1 = 0x3f803f80u ^ (x2 & 0x003f003fu);\n"
|
|
" float c0 = 0.f, c1 = 0.f, c2 = 0.f, c3 = 0.f, e0 = 0.f, e1 = 0.f, e2 = 0.f, e3 = 0.f;\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) {\n"
|
|
" asm volatile(\"mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\" : \"=f\"(c0), \"=f\"(c1), \"=f\"(c2), \"=f\"(c3) : \"r\"(a0), \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(b0), \"r\"(b1), \"f\"(c0), \"f\"(c1), \"f\"(c2), \"f\"(c3));\n"
|
|
" asm volatile(\"mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\" : \"=f\"(e0), \"=f\"(e1), \"=f\"(e2), \"=f\"(e3) : \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(a0), \"r\"(b1), \"r\"(b0), \"f\"(e0), \"f\"(e1), \"f\"(e2), \"f\"(e3));\n"
|
|
" a0 = ((a0 ^ __float_as_uint(c0)) & 0x007f007fu) | 0x3f803f80u; a2 = ((a2 ^ __float_as_uint(e2)) & 0x007f007fu) | 0x3f803f80u; c0 *= 0.5f; c1 *= 0.5f; c2 *= 0.5f; c3 *= 0.5f; e0 *= 0.5f; e1 *= 0.5f; e2 *= 0.5f; e3 *= 0.5f;\n }\n"
|
|
" out[g] = __float_as_uint(c0) ^ __float_as_uint(c1) ^ __float_as_uint(c2) ^ __float_as_uint(c3) ^ __float_as_uint(e0) ^ __float_as_uint(e1) ^ __float_as_uint(e2) ^ __float_as_uint(e3);\n}\n" },
|
|
{ "mma_e4m3_m16n8k32", "FP8 e4m3 multiply-add with FP32 accumulate inside mma.sync m16n8k32 (4,096 per tile per warp, 128 per lane; two tiles per step; sm_89 and later)", 256, 4096, 0, 0, 0,
|
|
" uint32_t a0 = x0 & 0x7f7f7f7fu, a1 = x1 & 0x7f7f7f7fu, a2 = x2 & 0x7f7f7f7fu, a3 = x3 & 0x7f7f7f7fu, b0 = (x1 >> 1) & 0x3f3f3f3fu, b1 = (x2 >> 1) & 0x3f3f3f3fu;\n"
|
|
" float c0 = 0.f, c1 = 0.f, c2 = 0.f, c3 = 0.f, e0 = 0.f, e1 = 0.f, e2 = 0.f, e3 = 0.f;\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) {\n"
|
|
" asm volatile(\"mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\" : \"=f\"(c0), \"=f\"(c1), \"=f\"(c2), \"=f\"(c3) : \"r\"(a0), \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(b0), \"r\"(b1), \"f\"(c0), \"f\"(c1), \"f\"(c2), \"f\"(c3));\n"
|
|
" asm volatile(\"mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\" : \"=f\"(e0), \"=f\"(e1), \"=f\"(e2), \"=f\"(e3) : \"r\"(a1), \"r\"(a2), \"r\"(a3), \"r\"(a0), \"r\"(b1), \"r\"(b0), \"f\"(e0), \"f\"(e1), \"f\"(e2), \"f\"(e3));\n"
|
|
" a0 = (a0 ^ __float_as_uint(c0)) & 0x7f7f7f7fu; a2 = (a2 ^ __float_as_uint(e2)) & 0x7f7f7f7fu; c0 *= 0.5f; c1 *= 0.5f; c2 *= 0.5f; c3 *= 0.5f; e0 *= 0.5f; e1 *= 0.5f; e2 *= 0.5f; e3 *= 0.5f;\n }\n"
|
|
" out[g] = __float_as_uint(c0) ^ __float_as_uint(c1) ^ __float_as_uint(c2) ^ __float_as_uint(c3) ^ __float_as_uint(e0) ^ __float_as_uint(e1) ^ __float_as_uint(e2) ^ __float_as_uint(e3);\n}\n" },
|
|
{ "l2_chase_32m", "dependent random 4-byte read in a 32 MiB table (L2-resident on a 5090's 96 MB; one chain per lane)", 1, 2048, 32, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) x0 = tbl[x0 & mask] ^ (x0 * 0x9E3779B1u + s);\n out[g] = x0;\n}\n" },
|
|
{ "l2_chase_64m", "dependent random 4-byte read in a 64 MiB table (inside a 5090's L2, outside a 4070's 36 MB)", 1, 2048, 64, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) x0 = tbl[x0 & mask] ^ (x0 * 0x9E3779B1u + s);\n out[g] = x0;\n}\n" },
|
|
{ "l2_indep4_32m", "random 4-byte read in a 32 MiB table, 4 independent chains per lane (L2 throughput)", 4, 2048, 32, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { x0 = tbl[x0 & mask] ^ (x0 * 0x9E3779B1u + s); x1 = tbl[x1 & mask] ^ (x1 * 0x9E3779B1u + s); x2 = tbl[x2 & mask] ^ (x2 * 0x9E3779B1u + s); x3 = tbl[x3 & mask] ^ (x3 * 0x9E3779B1u + s); }\n out[g] = x0 ^ x1 ^ x2 ^ x3;\n}\n" },
|
|
{ "dram_chase_1g", "dependent random 4-byte read in a 1 GiB table (the hash's own pattern, the control)", 1, 512, 1024, 0, 0,
|
|
" for (uint32_t s = 0u; s < steps; ++s) x0 = tbl[x0 & mask] ^ (x0 * 0x9E3779B1u + s);\n out[g] = x0;\n}\n" },
|
|
{ "tex_point_u32_32m", "texture fetch, point sampled, tex.1d.v4.u32.s32 over a 32 MiB u32 table (the sampler's address path; dependent chain)", 1, 2048, 32, 0, 1,
|
|
" for (uint32_t s = 0u; s < steps; ++s) { int cx = (int)(x0 & mask); uint32_t t0, t1, t2, t3; asm volatile(\"tex.1d.v4.u32.s32 {%0,%1,%2,%3}, [%4, {%5}];\" : \"=r\"(t0), \"=r\"(t1), \"=r\"(t2), \"=r\"(t3) : \"l\"(tex), \"r\"(cx)); x0 = t0 ^ (x0 * 0x9E3779B1u + s); }\n out[g] = x0;\n}\n" },
|
|
{ "tex_linear_f32_256k", "texture fetch, linearly filtered, tex.1d.v4.f32.f32 over a 1 MiB float table (the sampler's interpolation; 4 independent chains)", 4, 4096, 1, 1, 2,
|
|
" float p0 = (float)(x0 & 0xffffu), p1 = (float)(x1 & 0xffffu), p2 = (float)(x2 & 0xffffu), p3 = (float)(x3 & 0xffffu);\n"
|
|
" for (uint32_t s = 0u; s < steps; ++s) { float t0, t1, t2, t3, u0, u1, u2, u3, v0, v1, v2, v3, w0, w1, w2, w3;\n"
|
|
" asm volatile(\"tex.1d.v4.f32.f32 {%0,%1,%2,%3}, [%4, {%5}];\" : \"=f\"(t0), \"=f\"(t1), \"=f\"(t2), \"=f\"(t3) : \"l\"(tex), \"f\"(p0));\n"
|
|
" asm volatile(\"tex.1d.v4.f32.f32 {%0,%1,%2,%3}, [%4, {%5}];\" : \"=f\"(u0), \"=f\"(u1), \"=f\"(u2), \"=f\"(u3) : \"l\"(tex), \"f\"(p1));\n"
|
|
" asm volatile(\"tex.1d.v4.f32.f32 {%0,%1,%2,%3}, [%4, {%5}];\" : \"=f\"(v0), \"=f\"(v1), \"=f\"(v2), \"=f\"(v3) : \"l\"(tex), \"f\"(p2));\n"
|
|
" asm volatile(\"tex.1d.v4.f32.f32 {%0,%1,%2,%3}, [%4, {%5}];\" : \"=f\"(w0), \"=f\"(w1), \"=f\"(w2), \"=f\"(w3) : \"l\"(tex), \"f\"(p3));\n"
|
|
" p0 = fmaf(t0, 65535.0f, 0.37f); p1 = fmaf(u0, 65535.0f, 0.61f); p2 = fmaf(v0, 65535.0f, 0.13f); p3 = fmaf(w0, 65535.0f, 0.89f); }\n"
|
|
" out[g] = __float_as_uint(p0) ^ __float_as_uint(p1) ^ __float_as_uint(p2) ^ __float_as_uint(p3);\n}\n" },
|
|
};
|
|
|
|
static std::string utcNow() {
|
|
std::time_t t = std::time(nullptr); char b[32]; std::strftime(b, sizeof b, "%Y-%m-%dT%H:%M:%SZ", std::gmtime(&t)); return b;
|
|
}
|
|
|
|
static int runMicrobench(Ctx& c, const Options& o) {
|
|
std::printf("microbench on %s (sm_%d%d, %d SMs, driver %d.%d, NVRTC %d.%d): %d s per probe, full residency, wall time around cuStreamSynchronize; a 1 Hz power sampler aligns on start_utc and end_utc\n",
|
|
c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor, o.mbSeconds);
|
|
const size_t local = 256;
|
|
CUdeviceptr dOut = 0;
|
|
const size_t maxLanes = (size_t)(c.sms > 0 ? c.sms : 1) * 2048u;
|
|
if (c.drv.memAlloc(&dOut, maxLanes * 4u) != CUDA_SUCCESS) { std::printf("microbench: cuMemAlloc out\n"); return 2; }
|
|
std::vector<uint32_t> hOut(maxLanes);
|
|
int ran = 0, failed = 0;
|
|
for (const MicroProbe& pr : MICRO_PROBES) {
|
|
if (!o.mbOnly.empty() && ("," + o.mbOnly + ",").find(std::string(",") + pr.name + ",") == std::string::npos) continue;
|
|
std::string src = std::string(MICRO_PRELUDE) + "KB_KERNEL(kb_probe)\n" + pr.body;
|
|
Compiled cp; std::string err;
|
|
if (!rtcCompile(c, src, (std::string(pr.name) + ".cu").c_str(), "", "", {}, cp, err)) {
|
|
std::string e = err; for (char& ch : e) if (ch == '\n') ch = ' ';
|
|
std::printf("RESULT microbench probe=%s status=compile_failed error=\"%.300s\"\n", pr.name, e.c_str()); ++failed; std::fflush(stdout); continue;
|
|
}
|
|
CUmodule mod = nullptr; CUfunction kFill = nullptr, kProbe = nullptr;
|
|
if (c.drv.moduleLoadData(&mod, cp.image.data()) != CUDA_SUCCESS || c.drv.moduleGetFunction(&kFill, mod, "kb_fill") != CUDA_SUCCESS || c.drv.moduleGetFunction(&kProbe, mod, "kb_probe") != CUDA_SUCCESS) {
|
|
std::printf("RESULT microbench probe=%s status=load_failed\n", pr.name); ++failed; std::fflush(stdout); if (mod) c.drv.moduleUnload(mod); continue;
|
|
}
|
|
int blocksPerSM = 0; c.drv.occupancy(&blocksPerSM, kProbe, (int)local, 0);
|
|
if (blocksPerSM < 1) blocksPerSM = 1;
|
|
size_t lanes = (size_t)blocksPerSM * local * (size_t)(c.sms > 0 ? c.sms : 1);
|
|
if (lanes > maxLanes) lanes = maxLanes;
|
|
int regs = 0; c.drv.funcGetAttribute(®s, CU_FUNC_ATTRIBUTE_NUM_REGS, kProbe);
|
|
CUdeviceptr dTbl = 0; uint32_t mask = 0; unsigned long long tex = 0; CUtexObject texObj = 0; bool ok = true; std::string why;
|
|
if (pr.tableMib > 0) {
|
|
uint64_t bytes = (uint64_t)pr.tableMib << 20; uint32_t words = (uint32_t)(bytes / 4ull); mask = words - 1u;
|
|
if (c.drv.memAlloc(&dTbl, (size_t)bytes) != CUDA_SUCCESS) { ok = false; why = "cuMemAlloc table"; }
|
|
else { uint32_t n = words, mode = (uint32_t)pr.tableMode; void* a[3] = { &dTbl, &n, &mode };
|
|
if (c.drv.launchKernel(kFill, (unsigned)((words + 255u) / 256u), 1, 1, 256, 1, 1, 0, nullptr, a, nullptr) != CUDA_SUCCESS || c.drv.streamSynchronize(nullptr) != CUDA_SUCCESS) { ok = false; why = "fill"; } }
|
|
}
|
|
if (ok && pr.tex > 0) {
|
|
if (!c.drv.texObjectCreate || !c.drv.texObjectDestroy) { ok = false; why = "the driver has no cuTexObjectCreate"; }
|
|
else {
|
|
CUDA_RESOURCE_DESC rd; std::memset(&rd, 0, sizeof rd); rd.resType = CU_RESOURCE_TYPE_LINEAR;
|
|
rd.res.linear.devPtr = dTbl; rd.res.linear.format = pr.tex == 1 ? CU_AD_FORMAT_UNSIGNED_INT32 : CU_AD_FORMAT_FLOAT; rd.res.linear.numChannels = 1; rd.res.linear.sizeInBytes = (size_t)pr.tableMib << 20;
|
|
CUDA_TEXTURE_DESC td; std::memset(&td, 0, sizeof td);
|
|
td.addressMode[0] = CU_TR_ADDRESS_MODE_CLAMP; td.addressMode[1] = CU_TR_ADDRESS_MODE_CLAMP; td.addressMode[2] = CU_TR_ADDRESS_MODE_CLAMP;
|
|
td.filterMode = pr.tex == 1 ? CU_TR_FILTER_MODE_POINT : CU_TR_FILTER_MODE_LINEAR; td.flags = pr.tex == 1 ? CU_TRSF_READ_AS_INTEGER : 0;
|
|
CUresult r = c.drv.texObjectCreate(&texObj, &rd, &td, nullptr);
|
|
if (r != CUDA_SUCCESS) { ok = false; why = "cuTexObjectCreate: " + c.err(r); } else tex = (unsigned long long)texObj;
|
|
}
|
|
}
|
|
if (ok) {
|
|
uint32_t steps = pr.steps, seed = 0x51ed270bu; void* a[6] = { &steps, &seed, &dOut, &dTbl, &mask, &tex };
|
|
// warm-up (also the checksum's launch)
|
|
CUresult r = c.drv.launchKernel(kProbe, (unsigned)(lanes / local), 1, 1, (unsigned)local, 1, 1, 0, nullptr, a, nullptr);
|
|
if (r == CUDA_SUCCESS) r = c.drv.streamSynchronize(nullptr);
|
|
if (r != CUDA_SUCCESS) { ok = false; why = "warm-up: " + c.err(r); }
|
|
else {
|
|
c.drv.memcpyDtoH(hOut.data(), dOut, lanes * 4u);
|
|
uint64_t fp = fnv1a64Bytes(hOut.data(), lanes * 4u);
|
|
std::string start = utcNow(); double t0 = wallMs(); uint64_t launches = 0; double last = 0;
|
|
while (true) {
|
|
double l0 = wallMs();
|
|
seed += 0x9E3779B9u;
|
|
if (c.drv.launchKernel(kProbe, (unsigned)(lanes / local), 1, 1, (unsigned)local, 1, 1, 0, nullptr, a, nullptr) != CUDA_SUCCESS || c.drv.streamSynchronize(nullptr) != CUDA_SUCCESS) { ok = false; why = "launch"; break; }
|
|
last = wallMs() - l0; ++launches;
|
|
if (wallMs() - t0 >= o.mbSeconds * 1000.0) break;
|
|
}
|
|
double secs = (wallMs() - t0) / 1000.0; std::string end = utcNow();
|
|
if (ok) {
|
|
double stepsPerS = (double)lanes * (double)pr.steps * (double)launches / secs;
|
|
std::printf("RESULT microbench probe=%s status=ok unit=\"%s\" ops_per_step=%g lanes=%zu regs=%d blocks_per_sm=%d steps=%u launches=%llu launch_ms=%.1f seconds=%.1f start_utc=%s end_utc=%s G_steps_s=%.3f G_ops_s=%.3f checksum=%016llx\n",
|
|
pr.name, pr.unit, pr.opsPerStep, lanes, regs, blocksPerSM, pr.steps, (unsigned long long)launches, last, secs, start.c_str(), end.c_str(), stepsPerS / 1e9, stepsPerS * pr.opsPerStep / 1e9, (unsigned long long)fp);
|
|
++ran;
|
|
}
|
|
}
|
|
}
|
|
if (!ok) { std::printf("RESULT microbench probe=%s status=skipped why=\"%s\"\n", pr.name, why.c_str()); ++failed; }
|
|
std::fflush(stdout);
|
|
if (texObj) c.drv.texObjectDestroy(texObj);
|
|
if (dTbl) c.drv.memFree(dTbl);
|
|
c.drv.moduleUnload(mod);
|
|
}
|
|
c.drv.memFree(dOut);
|
|
std::printf("microbench: done, %d probes ran, %d skipped or failed\n", ran, failed);
|
|
return failed > 0 && ran == 0 ? 2 : 0;
|
|
}
|
|
|
|
int main(int argc, char** argv) {
|
|
Options o = parseArgs(argc, argv);
|
|
Ctx c;
|
|
c.blockWarps = o.blockWarps;
|
|
c.race = o.race; c.raceBenchMs = o.raceBenchMs; c.raceBudgetS = o.raceBudgetS; c.raceRounds = o.raceRounds; c.batchLog2 = o.batchLog2; c.pinned = o.pinned;
|
|
c.warps = o.warps; c.batches = o.batches;
|
|
// --bench runs the pack's base kernel with no race; with --variant <name> (Counter ASIC 4.0 research, 8 October 2026:
|
|
// the SM-sparse job of 7 October ran 48 rows on base because --bench turned the race off and buildPair never raced) it
|
|
// runs the pinned race instead: base and the named variant only, no timing, the variant installed whatever its speed,
|
|
// so the bench reads that kernel; the race line names it
|
|
if (benchRaceOff(o.bench || o.memprobe || o.microbench, o.pinned)) c.race = "off";
|
|
if (!o.tuningPath.empty()) { bool ok = false; c.tuning = readText(o.tuningPath, ok); if (!ok) c.tuning.clear(); }
|
|
if (o.listRace) {
|
|
// Counter ASIC 4.0 research (8 October 2026, the third fix of the --bench --variant fault): the race order as the
|
|
// bench's own option handling builds it, with no device: "variants 2" with the named variant is the pass line,
|
|
// "variants 1 base only" the 02:52Z fault; with a pack the rewrite is applied to the bound kernel and reported
|
|
bool raceOff = benchRaceOff(o.bench || o.memprobe || o.microbench, o.pinned);
|
|
std::vector<Variant> all = allVariants();
|
|
Tuning tu = readTuning(c.tuning, "");
|
|
std::string pinned = !o.pinned.empty() ? o.pinned : (tu.found && !tu.race ? tu.variant : "");
|
|
std::vector<Variant> order = raceOff ? std::vector<Variant>{ *findVariant(all, "base") } : raceOrder(all, pinned, tu, c.race);
|
|
std::string names;
|
|
for (const Variant& v : order) names += (names.empty() ? "" : ",") + v.name;
|
|
std::printf("RESULT list-race bench=%d race_off=%d pinned=%s variants=%zu names=%s\n", o.bench ? 1 : 0, raceOff ? 1 : 0, pinned.c_str(), order.size(), names.c_str());
|
|
bool ok = pinned.empty() || (order.size() >= 2 && order.back().name == pinned);
|
|
if (!o.pack.empty() && !pinned.empty()) {
|
|
bool rok = true; std::string text = readText(o.pack + "/kernel_bound.cu", rok), dev, err2, src, why;
|
|
const Variant* v = findVariant(all, pinned);
|
|
bool rewritten = rok && v && deviceOnly(text, dev, err2) && variantSource(dev, *v, v->blockWarps > 0 ? v->blockWarps : o.blockWarps, src, why);
|
|
std::printf("RESULT list-race pack=%s variant=%s sparse_blocks=%d block_warps=%d rewrite=%s bytes=%zu nonces_arg=%d unit_fn=%d why=\"%s\"\n", o.pack.c_str(), pinned.c_str(), v ? v->sparseBlocks : 0, v ? (v->blockWarps > 0 ? v->blockWarps : o.blockWarps) : 0,
|
|
rewritten ? "applied" : "none", src.size(), rewritten && src.find(", uint32_t nonces) {") != std::string::npos ? 1 : 0, rewritten && src.find("igneum_hash_bound_unit(") != std::string::npos ? 1 : 0, why.empty() ? err2.c_str() : why.c_str());
|
|
ok = ok && rewritten;
|
|
if (rewritten && v && v->sparseBlocks > 0) {
|
|
// the wrapper's call against the unit function's parameters, name by name (the 03:23Z fault: every
|
|
// argument one character short), then the whole rewritten text through NVRTC when its library is here
|
|
size_t u = src.find("void igneum_hash_bound_unit("), uc = u == std::string::npos ? u : src.find(")", u);
|
|
size_t cl = src.rfind("igneum_hash_bound_unit("), cc = cl == std::string::npos ? cl : src.find(")", cl);
|
|
std::string params = u == std::string::npos ? "" : src.substr(u + 27, uc - (u + 27)), call = cl == std::string::npos ? "" : src.substr(cl + 23, cc - (cl + 23));
|
|
auto names = [](const std::string& list, bool lastToken) { std::vector<std::string> out; size_t i = 0; while (i <= list.size()) { size_t j = list.find(',', i); if (j == std::string::npos) j = list.size(); std::string one = list.substr(i, j - i); size_t e = one.find_last_not_of(" \t"); if (e != std::string::npos) { size_t b = lastToken ? one.find_last_of(" \t*&", e) : std::string::npos; size_t st = b == std::string::npos ? one.find_first_not_of(" \t") : b + 1; out.push_back(one.substr(st, e - st + 1)); } i = j + 1; } return out; };
|
|
std::vector<std::string> pn = names(params, true), cn = names(call, false);
|
|
bool match = pn.size() == cn.size() && !pn.empty();
|
|
for (size_t i = 0; match && i < pn.size(); ++i) if (pn[i] != cn[i]) match = false;
|
|
std::printf("RESULT list-race call=\"igneum_hash_bound_unit(%s)\" params=%zu args=%zu names_match=%d\n", call.c_str(), pn.size(), cn.size(), match ? 1 : 0);
|
|
ok = ok && match;
|
|
std::string rerr, rlib;
|
|
if (loadNvrtc(c.rtc, rerr, rlib)) {
|
|
c.archOpt = "sm_120"; c.rtcMajor = 12; c.rtcMinor = 8;
|
|
Compiled cb; std::string cerr;
|
|
bool pr = true; std::string programH = readText(o.pack + "/program.h", pr), memhardH = readText(o.pack + "/memhard.h", pr);
|
|
bool compiled = pr && rtcCompile(c, src, "kernel_bound.cu", programH, memhardH, { "igneum_hash_bound" }, cb, cerr);
|
|
for (char& ch : cerr) if (ch == '\n') ch = ' ';
|
|
std::printf("RESULT list-race nvrtc=%s arch=sm_120 compiled=%d image_bytes=%zu error=\"%.400s\"\n", rlib.c_str(), compiled ? 1 : 0, cb.image.size(), cerr.c_str());
|
|
ok = ok && compiled;
|
|
} else {
|
|
std::printf("RESULT list-race nvrtc=none (%s): the rewritten text was not compiled here\n", rerr.c_str());
|
|
}
|
|
}
|
|
}
|
|
return ok ? 0 : 1;
|
|
}
|
|
std::string err, drvLib, rtcLib;
|
|
if (!loadDriver(c.drv, err, drvLib)) { emit("error 0 " + err); return 2; }
|
|
if (!loadNvrtc(c.rtc, err, rtcLib)) { emit("error 0 " + err); return 2; }
|
|
if (!openDevice(c, o.device, o.arch, err)) { emit("error 0 " + err); return 2; }
|
|
info(fmt("igneum-worker-cuda %s: device %d %s (sm_%d%d, %d SMs), driver %d.%d from %s, NVRTC %d.%d from %s, target %s (%s)",
|
|
WORKER_VERSION, o.device, c.name.c_str(), c.major, c.minor, c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, drvLib.c_str(), c.rtcMajor, c.rtcMinor, rtcLib.c_str(), c.archOpt.c_str(), c.why.c_str()));
|
|
if (!c.tuning.empty()) info(fmt("tuning file %s (%zu bytes): %s", o.tuningPath.c_str(), c.tuning.size(), readTuning(c.tuning, c.name).found ? "has an entry for this card" : "no entry for this card"));
|
|
if (o.memprobe) { int rc = runMemprobe(c, o); c.drv.primaryCtxRelease(c.dev); return rc; }
|
|
if (o.microbench) { int rc = runMicrobench(c, o); c.drv.primaryCtxRelease(c.dev); return rc; }
|
|
double t0 = wallMs();
|
|
Pair* cur = buildPair(c, o.pack, nullptr, err, !o.check && !benchRaceOff(o.bench, o.pinned));
|
|
if (!cur) { emit("error 0 " + err); return 1; }
|
|
if (o.bench) {
|
|
std::printf("pack %s on %s: %s\n", o.pack.c_str(), c.name.c_str(), pairSummary(cur).c_str());
|
|
int rc = runBench(c, o, cur);
|
|
releasePair(c, cur);
|
|
c.drv.primaryCtxRelease(c.dev);
|
|
return rc;
|
|
}
|
|
if (o.raceOnly) {
|
|
std::printf("race %s on %s (%s, %d SMs, driver %d.%d, NVRTC %d.%d, %s): %s\n", o.pack.c_str(), c.name.c_str(), c.archOpt.c_str(), c.sms, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.rtcMajor, c.rtcMinor, c.why.c_str(), pairSummary(cur).c_str());
|
|
std::printf("%s\n", cur->raceLine.c_str());
|
|
std::printf("winner %s: %d registers, %d blocks/SM at %d warp(s)/block\n", cur->variant.c_str(), cur->regs, cur->blocksPerSM, cur->blockWarps);
|
|
releasePair(c, cur);
|
|
return 0;
|
|
}
|
|
if (o.check) {
|
|
std::printf("check PASS %s in %.0f ms: %s\n", o.pack.c_str(), wallMs() - t0, pairSummary(cur).c_str());
|
|
std::printf(" epoch %s day %s, dataset 2^%u words, cache 2^%u words in %u segments, %d registers, %d blocks/SM at %d warp(s)/block, target %s\n",
|
|
cur->epochHex.c_str(), cur->dayHex.c_str(), cur->datasetLog2, (unsigned)__builtin_ctz(cur->cacheWords), cur->cacheSegments, cur->regs, cur->blocksPerSM, c.blockWarps, c.archOpt.c_str());
|
|
releasePair(c, cur);
|
|
return 0;
|
|
}
|
|
int rc = runServe(c, o, cur);
|
|
c.drv.primaryCtxRelease(c.dev);
|
|
return rc;
|
|
}
|