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()