From 27bcb8149426a949557b2d63f0338070ad74e07f Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Sun, 4 Oct 2026 19:08:45 +0000 Subject: [PATCH] Variant racing: the NVRTC and Metal workers compile several kernel variants at every prepare and keep the fastest for the hour the project lead, 4 Oct 2026 evening: "we need to make our miner better than anything else can be". The compile-ahead pipeline built one kernel per program; it now builds a catalogue (unroll 2 or 8; the dataset load path __ldg, __ldcg, __ldcs on NVIDIA; a register budget by --maxrregcount or __launch_bounds__, max_total_threads_per_threadgroup on Apple; 2, 4, 8 warps per block or 64 to 256 threads per threadgroup; combinations), self-tests each against the pack's vector warps (NVIDIA) or the base kernel over 2^16 nonces (Metal), bit for bit or out, and times each for about two seconds with the job loop paused (one mutex, mining resumes between variants). Base is the pack's text as shipped, always first, never discarded; a race has a budget (default 120 s against the 600-DAA lead) and keeps the best so far when it runs out, so the swap is never delayed. One line per race: variants, MH/s each, winner, gain, time. A tuning file (--tuning, IGNEUM_TUNING_FILE) pins a variant or orders the candidates per card model. NVIDIA: textual rewrites on the pack's own kernel_bound.cu with exact anchors from igneum-pow's emitter, so the pack format, the miner and igneum-pow are untouched and old packs race. --race --pack runs the race alone. Under IGNEUM_EMU the race is off (the stand-in checks the handed-over text is the pack's). Metal: the MSL hooks in generateMSL, a lock-protected kernel slot per program, a pair compiled inline races after its first job, --race-test --seed --day runs the race alone with a table. OpenCL is not raced yet. docs/design/miner-tuning.md. Co-Authored-By: Claude Fable 5.1 --- docs/design/miner-tuning.md | 156 +++++++++++++++ proto-cuda/nvrtc/worker.cpp | 389 ++++++++++++++++++++++++++++++++++-- proto-metal/main.swift | 267 +++++++++++++++++++++++-- 3 files changed, 778 insertions(+), 34 deletions(-) create mode 100644 docs/design/miner-tuning.md diff --git a/docs/design/miner-tuning.md b/docs/design/miner-tuning.md new file mode 100644 index 000000000..9a22f2003 --- /dev/null +++ b/docs/design/miner-tuning.md @@ -0,0 +1,156 @@ +# Miner tuning: variant racing and fleet learning + +4 October 2026, evening. the project lead: "we need to make our miner better than anything else can be". Two levers, both +measured: a race between kernel variants at every hourly swap, and a fleet that remembers which variant each card +model likes. Numbers live in `docs/bench-log.md` ("miner performance: variant racing"); the PC job is in +`docs/plans/miner-perf.md`. Nothing here changes the hash: every variant is the same instruction text in a +different shape for the compiler, and a variant that is not bit-exact is discarded before it is timed. + +## 1. Why a race + +The lottery program changes every hour. The generator draws 64 instructions with 16 loads; the compiler sees a +different straight-line body each time, and what suits one body (full unrolling, a read-only load path, more +threads per block) does not suit the next. A fixed compile is a guess. The compile-ahead pipeline already builds +the next hour's kernel one lead (600 DAA, about 600 s) before the boundary, so there is time to build several +and let the card pick. + +## 2. The variants + +| Knob | NVIDIA (NVRTC worker, `proto-cuda/nvrtc/worker.cpp`) | Apple (Metal worker, `proto-metal/main.swift`) | +|---|---|---| +| Unroll | `#pragma unroll 2` or `8` before the iteration loop (`u2`, `u8`) | the same pragma (`u2`, `u8`) | +| Load path | `ds[i]` as shipped; `__ldg` read-only path (`ldg`); `__ldcg` L2 only (`ldcg`); `__ldcs` streaming (`ldcs`) | none (one address space on Apple silicon) | +| Register budget | `--maxrregcount=32` or `64` (`r32`, `r64`); `__launch_bounds__(128, 4)` (`lb4-w4`), `(64, 8)` (`lb8-w2`) | `[[max_total_threads_per_threadgroup(N)]]` 256, 512, 1024 (`mt256` ...) | +| Threads per block | 2, 4, 8 warps (`w2`, `w4`, `w8`) | 64, 128, 256 threads per threadgroup (`g64`, `g128`, `g256`) | +| Compiler | | `optimizationLevel = .size` (`osize`) | +| Combinations | `u2-ldg`, `u2-w4`, `ldg-w4`, `ldcg-w4` | `u2-g128`, `u8-g128`, `mt256-g128`, `mt512-g256` | + +17 names on NVIDIA, 14 on Metal. `base` is always the pack's text as shipped with the worker's default block: +the kernel every machine ran before this change. Names are stable; the tuning file and the fleet records use them. + +NVIDIA rewrites are textual, on the pack's own `kernel_bound.cu`, with exact anchors from `igneum-pow`'s emitter +(`"\n for (uint32_t it = 0u; it < "`, `" ^ ds["`, `"__global__ void igneum_hash_bound("`); a text without the +anchor refuses the variant instead of guessing. The pack format, the miner and `igneum-pow` are untouched, so a +new worker races old packs. OpenCL (AMD, Intel) is not raced yet: `proto-opencl/host.c` builds through +`clBuildProgram` with `-D IGNEUM_GROUP` already, so the same catalogue (unroll pragma, group size, `-cl-` options) +is the next step; see "open". + +## 3. The race inside the prepare + +``` +prepare (the miner, one lead before the boundary) + compile kernel.cu, build cache + dataset, self-test base (as before) + compile the variants (NVRTC: 4 threads; Metal: in turn) budget: --race-budget-s, default 120 + for each round (1 in --serve): + for each variant: lock the card, self-test, time ~2 s, unlock + winner = fastest; base keeps its place unless beaten by 0.5% + one "race ..." line, then "prepared ..." as before +job on the new pair at the boundary (swaps as before; the winner serves the hour) +``` + +Rules that keep the swap safe: + +| Rule | Where | +|---|---| +| Base is the first entry and is never discarded; a race that runs out of budget keeps the best so far | `racePair`, `raceProgram` | +| Every variant must reproduce the pack's vector warps (NVIDIA) or the base kernel's output over 2^16 nonces (Metal), bit for bit, or it is out | `raceTime`, `raceProgram` | +| The card is exclusive while a variant is timed: one mutex, held per chunk by the job loop and per window by the race. Mining pauses about 2 s per variant and resumes between variants | `gpuMutex`, `gpuLock` | +| `--race-budget-s` is capped at 540 (the lead is 600 DAA); default 120 | option parsing | +| A pair compiled inline (nobody prepared it) races after its first job, in the background | Metal `raceDue`; NVIDIA self-heal path | +| `--race off` restores the old behaviour; `--race a,b,c` limits the catalogue | both workers | +| Under `IGNEUM_EMU` (the Mac's emulation test) the race is off: the stand-in checks that the handed-over text is the pack's | `racePair` | + +Cost per hour: on the 5090 about 17 variants x (2 s window + a self-test) of paused mining, under 1% of the hour, +plus the compiles on the CPU. The expected gain is what the race measures; nothing is claimed for it. + +The line, one per race (the miner logs it as `worker: race ...`): + +``` +race device driver arch loads wide variants base=/r/w u2=... ldg=- + winner base gain <+pct>% compile bench total ms [pinned by tuning|tuned order] + [| : ] +``` + +`--race --pack ` (NVIDIA) and `--race-test --seed --day ` (Metal) run the race alone, three rounds, +and print a table; that is what the bench log and the PC job use. + +## 4. Fleet learning + +### 4.1 The record + +The app's engine reads the race line (`engine.rs` `race_line`) and writes one `TUNING {json}` line to the app +log, which the existing intake receives with every upload (Neon `miner_logs`, the same table `tools/logs.mjs` +reads). Fields: + +| Field | From | +|---|---| +| `ts`, `machine` (id8), `app` (version) | the engine | +| `card` (the worker's device name, spaces as underscores: the key everything else uses), `vendor`, `worker` (CUDA, Metal, OpenCL) | the race line, the card state | +| `driver`, `arch` (sm_120, metal) | the race line | +| `epoch` (16 hex), `loads`, `wide` (the program class features the generator fixes: loads per hash and wide loads per hash) | the race line | +| `variants` {name: MH/s or null} | the race line | +| `winner`, `mhs`, `base_mhs`, `gain_pct`, `total_ms`, `pinned`, `tuned`, `notes` | the race line | +| `power_limit_w`, `power_w`, `power_pct`, `mh_per_w` (= mhs / power_w, 0 when the card reports no draw) | the card state (nvidia-smi telemetry; Apple reports none) | + +The card state also carries `variant`, `race_mhs`, `race_gain_pct`, `race_variants` for the dashboard, and one +event per race ("RTX 5090 kernel race: u2-ldg at 118.3 MH/s (+2.1% over base, 17 variants, 41 s)"). + +### 4.2 The aggregation + +`tools/tuning.mjs` on the Mac (or a small job): every app-log upload of the window (default 7 days) with a +TUNING line, de-duplicated on (machine, card, epoch) because the log is re-sent every minute, then per card +model and variant the sample count, the median MH/s and the median MH per watt. The winner is the best median +with at least `--min-samples` (3) samples; `--by mhw` ranks by MH per watt instead. `--write tuning.json` writes: + +```json +{"updated": "2026-10-04T21:00:00Z", "window_days": 7, + "cards": {"NVIDIA_GeForce_RTX_5090": {"variant": "u2-ldg", "race": true, "candidates": ["u2-ldg", "ldg", "base"], + "samples": 41, "races": 41, "machines": 2, "mhs": 118.3, "base_mhs": 115.9, + "gain_pct": 2.07, "mh_per_w": 0.254, "worker": "CUDA", "by": "mhs"}}} +``` + +A program class split (by `loads`, `wide`) is in the record and not yet in the aggregation: the generator fixes +16 loads per instruction block, so every program has 128 loads per hash today; the split starts to matter when the +era draw changes the mix. + +### 4.3 The way back: the manifest + +`packaging/ota/publish-manifest.sh --tuning tuning.json` puts the object under `tuning` in the signed +`igneum-app-latest.json` (carried over from the current manifest when not given; `--no-tuning` drops it; the +same script now also takes `--override '{json}'` for `consensus.override` and carries that over too). +`manifest.rs` parses `tuning` (an object with a `cards` object, else the manifest is refused). The updater writes +it as is to `/app/tuning.json` (`ota.rs` `write_tuning`, next to `override.json`), logs one event, and +the engine starts every miner with `IGNEUM_TUNING_FILE=` (`procs::spawn` gained an environment +parameter); the miner's child, the worker, reads it at every prepare. No restart for a change: a worker that has +the path reads the file again at its next prepare; a miner started before the file existed is restarted by the +usual hourly path. + +What a worker does with its entry (`readTuning`, `readMetalTuning`): + +| Entry | Behaviour | +|---|---| +| none for this card model | the full race | +| `race: true` with `candidates` | the candidates are raced first, then the rest as the budget allows (the fleet keeps learning; the card starts from the known best) | +| `race: false` with `variant` | the variant is compiled and self-tested, no timing (about 1 s); a failed self-test falls back to the full race | +| `--variant ` on the worker | the same as a pinned entry, for tests | + +### 4.4 What is implemented and what is not + +| Piece | State | +|---|---| +| NVRTC worker race, `--race` mode, tuning file | code, syntax-checked for mingw and the emulation build; emulation suite; **not yet run on a GPU** (the PC job does that) | +| Metal worker race, `--race-test`, deferred race, tuning file | code and the Mac measurement (bench log) | +| OpenCL worker race | not implemented (open) | +| Engine: race line to TUNING record, card state, event, `IGNEUM_TUNING_FILE` | code, unit tests of the manifest parse | +| `publish-manifest.sh --tuning`, `--override`, carry-over | code (dry run against a scratch folder in the plan) | +| `tools/tuning.mjs` | code; needs records, so no table yet | +| Aggregation as a job on a PC or the observer | not needed yet: the Mac script reads the intake | +| Dashboard field for the variant | state only; the UI does not show it yet | + +## 5. Open + +- OpenCL: the same catalogue through `clBuildProgram` options and the pragma; `--group-warps` exists already. +- The program class split in the aggregation once eras change the instruction mix. +- A per-card cap on how long a race may pause mining (today the 2 s windows plus the budget); and whether to race + only every N hours once a card's winner is stable (the pinned entry does that by hand). +- Power: the record has MH per watt, the race does not touch the power cap; a race across caps is a later lever. diff --git a/proto-cuda/nvrtc/worker.cpp b/proto-cuda/nvrtc/worker.cpp index 3e884faf1..aee50baf0 100644 --- a/proto-cuda/nvrtc/worker.cpp +++ b/proto-cuda/nvrtc/worker.cpp @@ -30,8 +30,23 @@ // 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 device ... variants N a=MH/s b=MH/s ... winner +// gain ...`. 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 [--device D] [--batch-log2 22] [--block-warps 1] [--arch sm_120|auto] +// [--race on|off|] [--race-bench-ms 2000] [--race-budget-s 120] [--race-rounds 1] +// [--variant ] [--tuning ] // igneum-worker-cuda --check --pack [--device D] compile, build, self-test, print timings, exit 0/1 +// igneum-worker-cuda --race --pack [--device D] [--race-rounds 3] the race alone: one line per variant, exit 0/1 #include #include @@ -43,6 +58,8 @@ #include #include #include +#include +#include #include #include "cuda_api.h" @@ -233,6 +250,11 @@ struct Ctx { 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"; } }; @@ -317,7 +339,7 @@ struct Compiled { }; static bool rtcCompile(Ctx& c, const std::string& src, const char* name, const std::string& programH, const std::string& memhardH, - const std::vector& nameExprs, Compiled& out, std::string& err) { + const std::vector& nameExprs, Compiled& out, std::string& err, const std::vector& 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" }; @@ -331,8 +353,9 @@ static bool rtcCompile(Ctx& c, const std::string& src, const char* name, const s std::string archOpt = "--gpu-architecture=" + c.archOpt; // -default-device: NVRTC rejects unannotated functions as host code (nvcc treats them as host and discards them); // the pack headers (program.h, memhard.h) carry plain inline helpers, so every unannotated function is device code here. - const char* opts[3] = { archOpt.c_str(), "--std=c++17", "-default-device" }; - r = c.rtc.compileProgram(prog, 3, opts); + std::vector 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) { @@ -382,8 +405,16 @@ struct Pair { 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) + std::string raceLine; // the race's one-line report, emitted by the main thread with "prepared" + double raceMs = 0; }; +// 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); @@ -404,9 +435,295 @@ static bool launchHash(Ctx& c, Pair* p, CUdeviceptr out, uint32_t baseNonce, con return true; } -// Compiles the pack in `dir`, builds its cache and dataset on stream `s`, runs the self-test. Returns the pair or -// null with `err` set. Runs on the main thread for --pack and --check, on the prepare thread for `prepare`. -static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string& err) { + +// --------------------------------------------------------------------------------------------- +// 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) +}; + +// 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 allVariants() { + std::vector 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; +} + +static const Variant* findVariant(const std::vector& all, const std::string& name) { + for (const Variant& v : all) if (v.name == name) return &v; + 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)); + } + return true; +} + +// The tuning file: {"cards": {"": {"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 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); +} + +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 hold(gpuMutex); + CUfunction keep = p->fHashBound; + p->fHashBound = e.fn; + std::string err; + uint32_t block = 32u * (uint32_t)e.blockWarps; + bool ok = true; + if (selfTest && pk.haveVectors) { + std::vector 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; + 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 all = allVariants(); + Tuning tu = readTuning(c.tuning, c.name); + std::string pinned = !c.pinned.empty() ? c.pinned : (tu.found && !tu.race ? tu.variant : ""); + // the order: base first, then the pinned or tuned candidates, then the rest (or the --race list only) + std::vector order; + auto push = [&](const std::string& n) { const Variant* v = findVariant(all, n); if (v && !findVariant(order, n)) order.push_back(*v); }; + push("base"); + if (!pinned.empty()) push(pinned); + else { + for (const std::string& n : tu.candidates) push(n); + if (c.race != "on" && c.race != "off") { std::string rest = c.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 (c.race == "on") for (const Variant& v : all) push(v.name); + } +#ifdef IGNEUM_EMU + order.resize(1); // the stand-in checks that the handed-over text is the pack's; no rewrites under emulation +#endif + 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 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 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 g(m); if (next >= todo.size() || wallMs() > deadline - benchMs) return; i = todo[next++]; } + RaceEntry& e = entries[i]; + std::vector 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(4, std::max(1, todo.size())); + std::vector 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; + } + } + 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; + w.mod = nullptr; + } else { + p->blockWarps = c.blockWarps; p->variant = "base"; + } + 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; } @@ -504,11 +821,13 @@ static Pair* buildPair(Ctx& c, const std::string& dir, CUstream s, std::string& } 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 check %.0f ms; %s", p->compileMs, p->cacheMs, p->dsMs, p->checkMs, p->check.c_str()); + return fmt("nvrtc %.0f cache %.0f dataset %.0f check %.0f race %.0f ms variant %s; %s", p->compileMs, p->cacheMs, p->dsMs, p->checkMs, p->raceMs, p->variant.c_str(), p->check.c_str()); } // --------------------------------------------------------------------------------------------- @@ -527,7 +846,7 @@ static void prepareRun(Ctx* c, PrepareTask* t) { 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); + 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"; @@ -542,9 +861,11 @@ static void prepareRun(Ctx* c, PrepareTask* t) { // Serve struct Options { - bool serve = false, check = false; + bool serve = false, check = false, raceOnly = false; 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() { @@ -554,7 +875,14 @@ static void usage() { " --device D CUDA device index (default 0)\n" " --batch-log2 B nonces per dispatch = 2^B (default 22)\n" " --block-warps W warps per thread block (default 1)\n" - " --arch sm_XY|compute_XY|auto NVRTC target (default auto: the device's architecture)\n", WORKER_VERSION); + " --arch sm_XY|compute_XY|auto NVRTC target (default auto: the device's architecture)\n" + " --race --pack 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 use this variant without a race (also from the tuning file)\n" + " --tuning the per-card tuning file (default: IGNEUM_TUNING_FILE from the environment)\n", WORKER_VERSION); } static Options parseArgs(int argc, char** argv) { @@ -564,6 +892,13 @@ static Options parseArgs(int argc, char** argv) { 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 == "--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()); @@ -575,7 +910,11 @@ static Options parseArgs(int argc, char** argv) { } if (o.batchLog2 < 10 || o.batchLog2 > 28) { std::printf("--batch-log2 must be between 10 and 28\n"); std::exit(2); } if (o.blockWarps < 1 || o.blockWarps > 32) { std::printf("--block-warps must be between 1 and 32\n"); std::exit(2); } - if (!o.serve && !o.check) { usage(); std::exit(2); } + if (!o.serve && !o.check && !o.raceOnly) { usage(); 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()) { std::printf("--pack is required (igneum-miner export-pack writes one)\n"); std::exit(2); } while (o.pack.size() > 1 && (o.pack.back() == '/' || o.pack.back() == '\\')) o.pack.pop_back(); return o; @@ -632,9 +971,10 @@ static int runServe(Ctx& c, const Options& o, Pair* cur) { Pair* old = nullptr; PrepareTask* task = nullptr; std::string prepareRoot; // the parent of the last prepare's pack directory: where the miner writes its packs - emit(fmt("ready cuda %s pack %s dataset-log2 %u batch %u regs %d prepare 1 path nvrtc %d.%d driver %d.%d arch %s worker %s", - c.name.c_str(), cur->seedString.c_str(), cur->datasetLog2, batch, cur->regs, c.rtcMajor, c.rtcMinor, c.driverVersion / 1000, (c.driverVersion % 100) / 10, c.archOpt.c_str(), WORKER_VERSION)); + 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; @@ -646,6 +986,7 @@ static int runServe(Ctx& c, const Options& o, Pair* cur) { 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())); @@ -695,9 +1036,9 @@ static int runServe(Ctx& c, const Options& o, Pair* cur) { 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); + 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) { 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); } } @@ -734,8 +1075,10 @@ static int runServe(Ctx& c, const Options& o, Pair* cur) { b[45] = (uint8_t)hi; b[46] = (uint8_t)(hi >> 8); b[47] = (uint8_t)(hi >> 16); b[48] = (uint8_t)(hi >> 24); pf_seed_words_from_bytes(b, 49, iw); } - // A chunk that is not a multiple of the block is finished one 32-lane block at a time - uint32_t block = 32u * (uint32_t)o.blockWarps; + // 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 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; } @@ -774,15 +1117,25 @@ 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; + if (!o.tuningPath.empty()) { bool ok = false; c.tuning = readText(o.tuningPath, ok); if (!ok) c.tuning.clear(); } 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")); double t0 = wallMs(); - Pair* cur = buildPair(c, o.pack, nullptr, err); + Pair* cur = buildPair(c, o.pack, nullptr, err, !o.check); if (!cur) { emit("error 0 " + err); return 1; } + 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", diff --git a/proto-metal/main.swift b/proto-metal/main.swift index 7287bff2b..508b77328 100644 --- a/proto-metal/main.swift +++ b/proto-metal/main.swift @@ -35,6 +35,16 @@ struct Options { // Generator levers (MEMHARD.md section 7). Defaults reproduce the original generator exactly. var loadWeight = 25 // --load-weight W: percent weight of the load op (default 25) var wideFrac = 0 // --wide-frac P: percent of load instructions emitted as warp-coalesced wide loads + // Variant racing (4 October 2026, evening; docs/design/miner-tuning.md): in --serve every prepared program is + // compiled in several variants, each checked bit for bit against the base kernel and timed for about two + // seconds with the job loop paused; the fastest serves the hour. --race-test runs the race alone and prints it. + var race = "on" // --race on|off|a,b,c + var raceBenchMs = 2000 // --race-bench-ms + var raceBudgetS = 120 // --race-budget-s (the prepare lead is 600 DAA; base is kept when the budget runs out) + var raceRounds = 0 // --race-rounds (0 = 1 in --serve, 3 in --race-test) + var pinnedVariant: String? = nil // --variant: this one, no race + var tuningPath: String? = nil // --tuning (default IGNEUM_TUNING_FILE) + var raceTest = false // --race-test: the race for --seed on --day, rounds, table, exit var anyTest: Bool { fuzz != nil || edge || stats || determinism || memcheck } } @@ -66,6 +76,13 @@ func parseArgs() -> Options { case "--closed-form": o.closedForm = true case "--load-weight": o.loadWeight = Int(take()) ?? o.loadWeight case "--wide-frac": o.wideFrac = Int(take()) ?? o.wideFrac + case "--race": o.race = take() + case "--race-bench-ms": o.raceBenchMs = Int(take()) ?? o.raceBenchMs + case "--race-budget-s": o.raceBudgetS = Int(take()) ?? o.raceBudgetS + case "--race-rounds": o.raceRounds = Int(take()) ?? o.raceRounds + case "--variant": o.pinnedVariant = take() + case "--tuning": o.tuningPath = take() + case "--race-test": o.raceTest = true case "-h", "--help": print(""" igneum-bench [--seed ] [--hours N] [--batch-log2 22] [--batches 4] @@ -84,6 +101,10 @@ func parseArgs() -> Options { [--memcheck] static dataset-index mask check, 4 MiB run with wrapping nonces shortcut measurement: [--inline-dataset] bench variant: every load computes ds_elem(index) inline, no memory read + variant racing (4 October 2026; in --serve every prepared program is raced, the fastest variant serves the hour): + [--race on|off|a,b,c] [--race-bench-ms 2000] [--race-budget-s 120] [--race-rounds N] + [--variant ] use this variant, no race [--tuning ] per-card tuning (IGNEUM_TUNING_FILE) + [--race-test] the race alone for --seed on --day (3 rounds): a table per variant, exit 0/1 """) exit(0) default: @@ -782,7 +803,27 @@ enum LoadSource { case inlineMemhard(MixParams) // memory-hard: mh_word(cache, index), 8 dependent cache reads per word } -func generateMSL(_ p: Program, datasetLog2: Int, source: LoadSource = .stored, bound: Bool = false) -> String { +// A kernel variant for the race (4 October 2026): the same instruction text, a different shape for the compiler. +// Names are stable: the tuning file and the fleet records use them. "base" is the kernel as it has always shipped. +struct MetalVariant { + let name: String + var unroll = 0 // 0: the iteration loop as emitted; N: "#pragma unroll N" before it (8 = fully unrolled) + var maxThreads = 0 // N > 0: [[max_total_threads_per_threadgroup(N)]] (fewer threads per group, more registers per thread) + var groupWidth = 32 // threads per threadgroup at dispatch (a multiple of 32; SIMD groups stay 32 wide) + var sizeOpt = false // MTLCompileOptions.optimizationLevel = .size +} + +func metalVariants() -> [MetalVariant] { + [MetalVariant(name: "base"), + MetalVariant(name: "g64", groupWidth: 64), MetalVariant(name: "g128", groupWidth: 128), MetalVariant(name: "g256", groupWidth: 256), + MetalVariant(name: "u2", unroll: 2), MetalVariant(name: "u8", unroll: 8), + MetalVariant(name: "mt256", maxThreads: 256), MetalVariant(name: "mt512", maxThreads: 512), MetalVariant(name: "mt1024", maxThreads: 1024), + MetalVariant(name: "osize", sizeOpt: true), + MetalVariant(name: "u2-g128", unroll: 2, groupWidth: 128), MetalVariant(name: "u8-g128", unroll: 8, groupWidth: 128), + MetalVariant(name: "mt256-g128", maxThreads: 256, groupWidth: 128), MetalVariant(name: "mt512-g256", maxThreads: 512, groupWidth: 256)] +} + +func generateMSL(_ p: Program, datasetLog2: Int, source: LoadSource = .stored, bound: Bool = false, variant: MetalVariant? = nil) -> String { let mask = UInt32((1 << datasetLog2) - 1) var s = """ #include @@ -820,6 +861,7 @@ func generateMSL(_ p: Program, datasetLog2: Int, source: LoadSource = .stored, b let kernelName = bound ? "igneum_hash_bound" : "igneum_hash" let initArg = bound ? " constant uint* initw [[buffer(3)]],\n" : "" let iw = bound ? "initw" : "SEEDW" + if let v = variant, v.maxThreads > 0 { s += "[[max_total_threads_per_threadgroup(\(v.maxThreads))]]\n" } s += """ kernel void \(kernelName)(\(buffer0), device ulong* out [[buffer(1)]], @@ -833,6 +875,7 @@ func generateMSL(_ p: Program, datasetLog2: Int, source: LoadSource = .stored, b for i in 0..<8 { s += " { uint x = nonce ^ \(iw)[\(i)]; x += 0x9e3779b9u * \(i + 1)u; x = splitmix32(x); r\(i) = x ^ \(iw)[\((i + 1) & 7)]; }\n" } + if let v = variant, v.unroll > 0 { s += "\n#pragma unroll \(v.unroll)" } s += "\n for (uint it = 0u; it < \(Program.iterations)u; ++it) {\n uint sel = r0;\n" // The word index expression for a load: plain = a & MASK; wide = lane 0's a, aligned to 32 words, plus lane. func wordIndex(_ a: String, wide: Bool) -> String { wide ? "(simd_broadcast(\(a), 0) & WMASK) + lane" : "\(a) & MASK" } @@ -2007,6 +2050,8 @@ struct CompiledHash { let pipeline: MTLComputePipelineState let libraryMs: Double let pipelineMs: Double + var groupWidth = 32 // threads per threadgroup at dispatch (the variant's) + var variant = "base" var totalMs: Double { libraryMs + pipelineMs } } @@ -2737,9 +2782,164 @@ func blockInitWords(prehash: [UInt8], nonceHi: UInt32) -> [UInt32] { final class ServeProgram { let seedHex: String let program: Program - let compiled: CompiledHash - init(seedHex: String, program: Program, compiled: CompiledHash) { self.seedHex = seedHex; self.program = program; self.compiled = compiled } + private var slot: CompiledHash + private let lock = NSLock() + /// true until a race ran for this program (an inline compile races after the first job on it) + var raceDue = true + var raceLine = "" + init(seedHex: String, program: Program, compiled: CompiledHash) { self.seedHex = seedHex; self.program = program; self.slot = compiled } + var compiled: CompiledHash { lock.lock(); defer { lock.unlock() }; return slot } + func install(_ c: CompiledHash) { lock.lock(); slot = c; lock.unlock() } } + +// 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 window by the race. +let gpuLock = NSLock() + +// The tuning file: {"cards": {"": {"variant": "u2", "race": false, "candidates": [...]}}} +struct MetalTuning { + var found = false + var variant = "" + var race = true + var candidates = [String]() +} +func readMetalTuning(path: String?, device: String) -> MetalTuning { + var t = MetalTuning() + guard let path = path, let data = FileManager.default.contents(atPath: path), + let root = try? JSONSerialization.jsonObject(with: data) as? [String: Any], + let cards = root["cards"] as? [String: Any], let e = cards[device] as? [String: Any] else { return t } + t.found = true + t.variant = e["variant"] as? String ?? "" + t.race = e["race"] as? Bool ?? true + t.candidates = e["candidates"] as? [String] ?? [] + return t +} + +// Dispatches `count` lane nonces (a multiple of 32) from `base` with the kernel's group width; a tail under the +// width goes in 32-wide groups (same pipeline, a threadgroup size is a dispatch parameter in Metal). +func encodeHash(_ enc: MTLComputeCommandEncoder, _ k: CompiledHash, dataset: MTLBuffer, out: MTLBuffer, base: UInt32, initw: inout [UInt32], count: Int) { + enc.setComputePipelineState(k.pipeline) + enc.setBuffer(dataset, offset: 0, index: 0) + enc.setBuffer(out, offset: 0, index: 1) + var b = base + enc.setBytes(&b, length: 4, index: 2) + enc.setBytes(&initw, length: 32, index: 3) + let w = k.groupWidth + let main = count - count % w + if main > 0 { enc.dispatchThreadgroups(MTLSize(width: main / w, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: w, height: 1, depth: 1)) } + if main < count { + var b2 = base &+ UInt32(main) + enc.setBytes(&b2, length: 4, index: 2) + enc.setBuffer(out, offset: main * 8, index: 1) + enc.dispatchThreadgroups(MTLSize(width: (count - main) / 32, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: 32, height: 1, depth: 1)) + } +} + +struct RaceEntry { + let v: MetalVariant + var compiled: CompiledHash? = nil + var mhs = 0.0 // best round + var note = "" // why it is out, or a detail + var ok: Bool { compiled != nil && note.isEmpty } +} + +// Races the variants of `program` on `dataset` and returns the race line; the winner is installed in `program`. +// Every variant must equal the base kernel bit for bit over 2^16 nonces (the base is the kernel the miner's CPU +// re-check has always covered); a window is `benchMs` of launches of `batch` nonces, the first launch warming up. +// `rounds` interleaved rounds, best per variant. Stops compiling and timing after `budgetS` (base is always kept). +func raceProgram(_ gpu: GPU, _ program: ServeProgram, dataset: MTLBuffer, datasetLog2: Int, opts: Options, rounds: Int, table: Bool = false) -> String { + let t0 = nowNs() + let device = gpu.device.name.replacingOccurrences(of: " ", with: "_") + let tuning = readMetalTuning(path: opts.tuningPath ?? ProcessInfo.processInfo.environment["IGNEUM_TUNING_FILE"], device: device) + let all = metalVariants() + let pinned = opts.pinnedVariant ?? (tuning.found && !tuning.race ? tuning.variant : "") + var order = [MetalVariant]() + func push(_ n: String) { if let v = all.first(where: { $0.name == n }), !order.contains(where: { $0.name == n }) { order.append(v) } } + push("base") + if !pinned.isEmpty { push(pinned) } + else { + tuning.candidates.forEach(push) + if opts.race == "on" { all.forEach { push($0.name) } } + else if opts.race != "off" { opts.race.split(separator: ",").forEach { push(String($0)) } } + } + let pinnedOnly = !pinned.isEmpty && order.count == 2 + let batch = 1 << opts.batchLog2 + let checkN = 1 << 16 + let deadline = t0 + UInt64(opts.raceBudgetS) * 1_000_000_000 + guard let outBuf = gpu.device.makeBuffer(length: batch * 8, options: .storageModeShared), + let refBuf = gpu.device.makeBuffer(length: checkN * 8, options: .storageModeShared) else { return "race \(program.seedHex.prefix(16)) failed: no buffers" } + var entries = order.map { RaceEntry(v: $0) } + entries[0].compiled = program.compiled + // Compile (the base is already compiled) + for i in 1.. deadline { entries[i].note = "not compiled: the race budget ran out"; continue } + do { entries[i].compiled = try compileBound(gpu, msl: generateMSL(program.program, datasetLog2: datasetLog2, source: .stored, bound: true, variant: entries[i].v), variant: entries[i].v) } + catch { entries[i].note = "compile: \(error)".prefix(200).description } + } + let compileMs = ms(t0, nowNs()) + var initw = blockInitWords(prehash: [UInt8](repeating: 0x5a, count: 32), nonceHi: 7) + func run(_ k: CompiledHash, base: UInt32, count: Int, into: MTLBuffer) -> Bool { + let cb = gpu.queue.makeCommandBuffer()! + let enc = cb.makeComputeCommandEncoder()! + encodeHash(enc, k, dataset: dataset, out: into, base: base, initw: &initw, count: count) + enc.endEncoding() + cb.commit(); cb.waitUntilCompleted() + return cb.error == nil + } + // Reference output of the base kernel (exclusive window) + gpuLock.lock() + let refOk = run(entries[0].compiled!, base: 0x1000_0000, count: checkN, into: refBuf) + gpuLock.unlock() + if !refOk { return "race \(program.seedHex.prefix(16)) failed: the base kernel did not run" } + let ref = refBuf.contents().bindMemory(to: UInt64.self, capacity: checkN) + let out = outBuf.contents().bindMemory(to: UInt64.self, capacity: batch) + for round in 0.. 0 && nowNs() > deadline { entries[i].note = "not timed: the race budget ran out"; continue } + gpuLock.lock() + defer { gpuLock.unlock() } + if round == 0 && i > 0 { + if !run(k, base: 0x1000_0000, count: checkN, into: outBuf) { entries[i].note = "did not run"; continue } + var bad = -1 + for j in 0..= 0 { entries[i].note = String(format: "lane %d: %016llx, base %016llx (discarded)", bad, out[bad], ref[bad]); continue } + } + if pinnedOnly { continue } + var launches = 0, hashes = 0, tStart: UInt64 = 0 + while true { + if !run(k, base: 0x2000_0000 &+ UInt32(launches * batch), count: batch, into: outBuf) { entries[i].note = "bench: did not run"; break } + let now = nowNs() + if launches == 0 { tStart = now } else { hashes += batch } + launches += 1 + if launches >= 3 && ms(tStart, now) >= Double(opts.raceBenchMs) { entries[i].mhs = max(entries[i].mhs, Double(hashes) / ms(tStart, now) / 1000.0); break } + } + } + } + var win = 0 + if pinnedOnly && entries.count == 2 && entries[1].ok { win = 1 } + else { for i in 1.. entries[win].mhs * (win == 0 ? 1.005 : 1.0) { win = i } } + let baseMhs = entries[0].mhs, winMhs = entries[win].mhs + if win != 0, let k = entries[win].compiled { program.install(k) } + let total = ms(t0, nowNs()) + let os = ProcessInfo.processInfo.operatingSystemVersion + var line = "race \(program.seedHex.prefix(16)) device \(device) driver macos-\(os.majorVersion).\(os.minorVersion).\(os.patchVersion) arch metal loads \(program.program.loadsPerHash) wide \(program.program.wideLoadsPerHash) variants \(entries.count)" + for e in entries { line += e.ok && (e.mhs > 0 || pinnedOnly) ? " \(e.v.name)=\(fmt(e.mhs, 3))/\(e.compiled!.pipeline.maxTotalThreadsPerThreadgroup)t/\(e.v.groupWidth)w" : " \(e.v.name)=-" } + let gain = baseMhs > 0 ? (winMhs / baseMhs - 1) * 100 : 0 + line += " winner \(entries[win].v.name) \(fmt(winMhs, 3)) base \(fmt(baseMhs, 3)) gain \(gain >= 0 ? "+" : "")\(fmt(gain, 2))% compile \(fmt(compileMs, 0)) bench \(fmt(total - compileMs, 0)) total \(fmt(total, 0)) ms" + if pinnedOnly { line += " pinned by tuning" } else if tuning.found { line += " tuned order" } + for e in entries where !e.note.isEmpty { line += " | \(e.v.name): \(e.note)" } + program.raceDue = false + program.raceLine = line + if table { + print("| variant | threads/group | max threads | MH/s | vs base | note |") + print("|---|---|---|---|---|---|") + for e in entries { print("| \(e.v.name) | \(e.v.groupWidth) | \(e.compiled.map { String($0.pipeline.maxTotalThreadsPerThreadgroup) } ?? "-") | \(e.ok ? fmt(e.mhs, 3) : "-") | \(e.ok && baseMhs > 0 ? (e.mhs / baseMhs - 1 >= 0 ? "+" : "") + fmt((e.mhs / baseMhs - 1) * 100, 2) + "%" : "-") | \(e.note) |") } + } + return line +} + final class ServeDataset { let dayHex: String let ctx: DatasetContext @@ -2782,14 +2982,19 @@ final class ServeStore { } } -func compileBound(_ gpu: GPU, msl: String) throws -> CompiledHash { +func compileBound(_ gpu: GPU, msl: String, variant: MetalVariant? = nil) throws -> CompiledHash { let t0 = nowNs() - let lib = try gpu.device.makeLibrary(source: msl, options: MTLCompileOptions()) + let copts = MTLCompileOptions() + if let v = variant, v.sizeOpt { if #available(macOS 13.0, *) { copts.optimizationLevel = .size } } + let lib = try gpu.device.makeLibrary(source: msl, options: copts) let t1 = nowNs() guard let fn = lib.makeFunction(name: "igneum_hash_bound") else { throw IgneumError("no igneum_hash_bound function in library") } let pipe = try gpu.device.makeComputePipelineState(function: fn) let t2 = nowNs() - return CompiledHash(pipeline: pipe, libraryMs: ms(t0, t1), pipelineMs: ms(t1, t2)) + var c = CompiledHash(pipeline: pipe, libraryMs: ms(t0, t1), pipelineMs: ms(t1, t2)) + if let v = variant { c.groupWidth = v.groupWidth; c.variant = v.name } + if c.groupWidth > pipe.maxTotalThreadsPerThreadgroup { throw IgneumError("variant \(c.variant): group width \(c.groupWidth) is over the pipeline's maxTotalThreadsPerThreadgroup \(pipe.maxTotalThreadsPerThreadgroup)") } + return c } // Builds the program for an epoch seed (hex) unless resident. Returns (program, compile ms) or throws. @@ -2823,7 +3028,7 @@ func runServe(_ opts: Options) -> Never { let store = ServeStore() let prepareQueue = DispatchQueue(label: "igneum.prepare") // one prepare at a time, off the job loop guard let outBuf = gpu.device.makeBuffer(length: batch * 8, options: .storageModeShared) else { emit("error 0 cannot allocate the output buffer"); exit(1) } - emit("ready metal \(gpu.device.name.replacingOccurrences(of: " ", with: "_")) dataset-log2 \(datasetLog2) batch \(batch) prepare \(opts.noPrepare ? 0 : 1)") + emit("ready metal \(gpu.device.name.replacingOccurrences(of: " ", with: "_")) dataset-log2 \(datasetLog2) batch \(batch) prepare \(opts.noPrepare ? 0 : 1) race \(opts.race)") var lastPair: (String, String)? = nil while let line = readLine(strippingNewline: true) { let f = line.split(separator: " ").map(String.init) @@ -2843,8 +3048,15 @@ func runServe(_ opts: Options) -> Never { do { let (sp, progMs) = try serveProgram(gpu, store, seedHex: epochHex, seed: epochSeed, datasetLog2: datasetLog2) let (sd, dsMs) = serveDataset(gpu, store, dayHex: dayHex, day: daySeed, datasetLog2: datasetLog2) + // The race: the prepared program against its own dataset, exclusive windows between jobs + var raceMs = 0.0 + if sp.raceDue && opts.race != "off" { + let r0 = nowNs() + emit(raceProgram(gpu, sp, dataset: sd.buffer, datasetLog2: datasetLog2, opts: opts, rounds: opts.raceRounds)) + raceMs = ms(r0, nowNs()) + } let (np, nd) = store.counts() - emit("prepared \(epochHex) \(dayHex) \(fmt(ms(t0, nowNs()), 1)) program \(fmt(progMs, 1)) dataset \(fmt(dsMs, 1)) loads/hash \(sp.program.loadsPerHash) cache-fill \(fmt(sd.ctx.cacheFillGPUms, 1)) resident \(np) programs \(nd) datasets") + emit("prepared \(epochHex) \(dayHex) \(fmt(ms(t0, nowNs()), 1)) program \(fmt(progMs, 1)) dataset \(fmt(dsMs, 1)) race \(fmt(raceMs, 1)) variant \(sp.compiled.variant) loads/hash \(sp.program.loadsPerHash) cache-fill \(fmt(sd.ctx.cacheFillGPUms, 1)) resident \(np) programs \(nd) datasets") } catch { emit("prepare-failed \(epochHex) \(dayHex) Metal compile failed: \(error)") } } continue @@ -2885,21 +3097,18 @@ func runServe(_ opts: Options) -> Never { var found = 0 var failed = false let outPtr = outBuf.contents().bindMemory(to: UInt64.self, capacity: batch) + let kernel = program.compiled // the pair's kernel for this job (a race may swap it for the next) while remaining > 0 { let room = UInt64(UInt32.max - lo) + 1 // lane nonces left before the high word steps let chunk = Int(min(min(remaining, UInt64(batch)), room)) var initw = blockInitWords(prehash: prehash, nonceHi: hi) + gpuLock.lock() // a race's exclusive windows fall between chunks let cb = gpu.queue.makeCommandBuffer()! let enc = cb.makeComputeCommandEncoder()! - enc.setComputePipelineState(program.compiled.pipeline) - enc.setBuffer(dataset.buffer, offset: 0, index: 0) - enc.setBuffer(outBuf, offset: 0, index: 1) - var b = lo - enc.setBytes(&b, length: 4, index: 2) - enc.setBytes(&initw, length: 32, index: 3) - enc.dispatchThreadgroups(MTLSize(width: chunk / 32, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: 32, height: 1, depth: 1)) + encodeHash(enc, kernel, dataset: dataset.buffer, out: outBuf, base: lo, initw: &initw, count: chunk) enc.endEncoding() cb.commit(); cb.waitUntilCompleted() + gpuLock.unlock() if let e = cb.error { emit("error \(jobId) dispatch failed: \(e)"); failed = true; break } for i in 0.. Never { let (dp, dd) = store.prune(to: pair) if dp + dd > 0 { emit("info dropped \(dp) program(s) and \(dd) dataset(s) of the previous pair") } } + if program.raceDue && opts.race != "off" { + // A pair compiled inline (nobody prepared it) races now, in the background, exclusive windows between chunks + program.raceDue = false + prepareQueue.async { emit(raceProgram(gpu, program, dataset: dataset.buffer, datasetLog2: datasetLog2, opts: opts, rounds: opts.raceRounds)) } + } lastPair = pair } exit(0) @@ -2926,9 +3140,30 @@ func runServe(_ opts: Options) -> Never { // MARK: - Main -let opts = parseArgs() +// The race alone (the Mac measurement, docs/bench-log.md "miner performance: variant racing"): the memory-hard +// dataset for --day, the version-2 program for --seed, every variant timed for --race-rounds rounds, a table. +func runRaceTest(_ opts: Options) -> Never { + let gpu = GPU() + print("igneum-bench --race-test on \(gpu.device.name): seed \"\(opts.seed)\", day \"\(opts.day)\", dataset 2^\(opts.datasetLog2) words, batch 2^\(opts.batchLog2), \(opts.raceRounds) rounds, \(opts.raceBenchMs) ms per window") + let ctx = DatasetContext(gpu: gpu, closedForm: false, dayString: opts.day) + let t0 = nowNs() + let dataset = ctx.makeDataset(log2: opts.datasetLog2) + print("dataset built in \(fmt(ms(t0, nowNs()), 0)) ms (cache fill \(fmt(ctx.cacheFillGPUms, 0)) ms GPU, build \(fmt(ctx.lastBuildGPUms, 0)) ms GPU)") + let p = generateProgramV2(seedString: opts.seed, bytes: Array(opts.seed.utf8)) + print("program: \(describeProgram(p))") + let base: CompiledHash + do { base = try compileBound(gpu, msl: generateMSL(p, datasetLog2: opts.datasetLog2, source: .stored, bound: true)) } catch { print("FAIL: \(error)"); exit(1) } + let sp = ServeProgram(seedHex: String(repeating: "0", count: 64), program: p, compiled: base) + let line = raceProgram(gpu, sp, dataset: dataset, datasetLog2: opts.datasetLog2, opts: opts, rounds: opts.raceRounds, table: true) + print(line) + exit(line.contains("winner ") ? 0 : 1) +} + +var opts = parseArgs() +if opts.raceRounds == 0 { opts.raceRounds = opts.raceTest ? 3 : 1 } generatorConfig = GeneratorConfig(loadWeight: opts.loadWeight, wideFrac: opts.wideFrac) if opts.exportPack != nil { exportPack(opts) } +if opts.raceTest { runRaceTest(opts) } if opts.serve { runServe(opts) } if opts.anyTest { runTests(opts) } let gpu = GPU()