opencl worker: the 9070 XT on the eGPU measured (2.5 G random reads/s is the card's ceiling), duplicate AMD platform folded, GPU-side select read-back, kernel report and memprobe; app: --list parser with tests, one pack export at a time

Measured on PC 1 (docs/bench-log.md, "the 9070 XT on the eGPU"): the hash runs at 18.0 MH/s against a dependent random-read ceiling of 2.42 to 2.68 G loads/s at 1 GiB (--memprobe), 128 loads per hash, so 92% of what the card does for this access pattern; group-warps 1/2/4/8 and the old platform all give 18.0; wave32, no spills; the stream probe reads 635 GB/s of the rated 640, random 64 B lines run at the same count as random 4 B loads. Not the eGPU link: the link only carried the 16 MiB per-job read-back (7.3 ms per 2^21-nonce job), which the select pass removes (124.2 to 117.3 ms per job, 16.88 to 17.87 MH/s inside jobs, kernel 116.0 ms on both). host.c: the same card on an older platform of the same vendor is marked and hidden from --list (PC 1 ran two workers on one 9070 XT at 8.9 + 9.4 MH/s); the default pick skips it; printKernelInfo on every path (preferred multiple, private memory, sub-group size); --readback select|full with the transfer and time budget in the stats line; --memprobe (chase, indep x8, 64 B lines, stream, ALU) with a fresh seed per repetition; test_host.c covers the fold. detect.rs: parse_opencl_list with three tests (the PC 1 listing, a CPU device, a repeated index). engine.rs: EXPORT_LOCK serialises igneum-miner export-pack across the per-card threads (PC 1's packs\devnet held one epoch's program.h with another's seeds.txt after an epoch change and every OpenCL worker start refused it). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
2026-10-05 19:15:52 +00:00 · 2026-10-05 19:15:52 +00:00 · 23810dfe58
commit 23810dfe58
parent 4c7a1d0f1a
8 changed files with 767 additions and 43 deletions
--- a/app/igneum-app/src/detect.rs
+++ b/app/igneum-app/src/detect.rs
@ -240,27 +240,12 @@ pub fn detect(bins: &Bins, notes: &mut Vec<String>) -> Vec<CardState> {
    // OpenCL: the worker's own device list (AMD, Intel; NVIDIA shows there too and is skipped)
    if let Some(cl) = bins.opencl.as_ref() {
        if let Some(out) = run_timeout(Command::new(cl).arg("--list"), None, Duration::from_secs(15)) {
-            let lines: Vec<&str> = out.lines().collect();
+            let (listed, list_notes) = parse_opencl_list(&out);
-            for (i, line) in lines.iter().enumerate() {
+            notes.extend(list_notes);
-                let t = line.trim_start_matches(|c| c == ' ' || c == '*').trim();
+            for l in listed {
-                if !t.starts_with('[') {
+                let mut c = card(cards.len(), &l.name, l.vendor, "OpenCL", &l.units, &l.index);
-                    continue;
+                c.device = l.index.clone();
-                }
+                c.kind = if looks_integrated(&l.name) { "integrated".into() } else { "discrete".into() };
                let Some(close) = t.find(']') else { continue };
                let idx = &t[1..close];
                let rest = &t[close + 1..];
                let name = rest.split("  |").next().unwrap_or("").trim();
                let info = lines.get(i + 1).map(|l| l.trim()).unwrap_or("");
                let is_gpu = info.starts_with("GPU");
                let vendor_s = info.split("vendor ").nth(1).unwrap_or("").split(", driver").next().unwrap_or("").trim();
                if !is_gpu || name.contains("NVIDIA") || vendor_s.contains("NVIDIA") {
                    continue;
                }
                let vendor = if vendor_s.contains("Advanced Micro") || name.contains("Radeon") || name.contains("AMD") { "amd" } else { "other" };
                let units = info.split(", ").find(|p| p.contains("compute units")).unwrap_or("").to_string();
                let mut c = card(cards.len(), name, vendor, "OpenCL", &units, idx);
                c.device = idx.to_string();
                c.kind = if looks_integrated(name) { "integrated".into() } else { "discrete".into() };
                c.path = "prebuilt".into();
                apply_defaults(&mut c);
                mark_sweep_support(&mut c);
@ -286,6 +271,106 @@ pub fn detect(bins: &Bins, notes: &mut Vec<String>) -> Vec<CardState> {
    cards
 }
 /// One GPU the OpenCL worker's `--list` shows (AMD, Intel). NVIDIA entries are left out: nvidia-smi lists those.
 #[derive(Debug, Clone, PartialEq)]
 pub struct OpenClListed {
    pub index: String,
    pub name: String,
    pub vendor: &'static str,
    pub units: String,
 }
 /// The worker's `--list` text to cards. A device line is `*[1] name  |  platform (version)` (the star marks the
 /// worker's default) followed by an info line `GPU, vendor ..., driver ..., N compute units, ...`; anything that is
 /// not a GPU, or is NVIDIA, is skipped. A line ` dup [3] name  |  ...` (host.c, 5 October 2026) is the same card
 /// listed again by an older driver's OpenCL platform that is still registered after a driver update: on PC 1 the
 /// 9070 XT appeared as [1] and [3], the app ran two workers on it, and each got half (8.9 and 9.4 MH/s). Such a
 /// line is never a card; it is reported once in `notes`.
 pub fn parse_opencl_list(out: &str) -> (Vec<OpenClListed>, Vec<String>) {
    let lines: Vec<&str> = out.lines().collect();
    let mut listed = Vec::new();
    let mut notes = Vec::new();
    let mut seen = std::collections::HashSet::new();
    for (i, line) in lines.iter().enumerate() {
        let raw = line.trim();
        if let Some(rest) = raw.strip_prefix("dup [") {
            let idx = rest.split(']').next().unwrap_or("").trim();
            let name = rest.split(']').nth(1).unwrap_or("").split("  |").next().unwrap_or("").trim();
            notes.push(format!("OpenCL [{idx}] {name} is the same card on an older driver's platform; hidden (the older driver's OpenCL registration is still present)"));
            continue;
        }
        let t = raw.trim_start_matches(|c| c == ' ' || c == '*').trim();
        if !t.starts_with('[') {
            continue;
        }
        let Some(close) = t.find(']') else { continue };
        let idx = &t[1..close];
        if idx.is_empty() || !idx.chars().all(|c| c.is_ascii_digit()) || !seen.insert(idx.to_string()) {
            continue;
        }
        let rest = &t[close + 1..];
        let name = rest.split("  |").next().unwrap_or("").trim();
        let info = lines.get(i + 1).map(|l| l.trim()).unwrap_or("");
        let is_gpu = info.starts_with("GPU");
        let vendor_s = info.split("vendor ").nth(1).unwrap_or("").split(", driver").next().unwrap_or("").trim();
        if !is_gpu || name.contains("NVIDIA") || vendor_s.contains("NVIDIA") {
            continue;
        }
        let vendor = if vendor_s.contains("Advanced Micro") || name.contains("Radeon") || name.contains("AMD") { "amd" } else { "other" };
        let units = info.split(", ").find(|p| p.contains("compute units")).unwrap_or("").to_string();
        listed.push(OpenClListed { index: idx.to_string(), name: name.to_string(), vendor, units });
    }
    (listed, notes)
 }
 #[cfg(test)]
 mod tests {
    use super::*;
    // PC 1 on 5 October 2026 after the Adrenalin 26.9.2 install, as the 5 October worker prints it: the new platform
    // (3683.0) first, the old one (3652.0) folded into dup lines, the NVIDIA card last.
    const PC1_LIST: &str = "igneum-bench-cl  pack \"igneum-devnet-v4-epoch0\"  (test harness: no pool, no network, no wallet)
 OpenCL devices (5):
 [0] gfx1036  |  AMD Accelerated Parallel Processing (OpenCL 2.1 AMD-APP (3683.0))
      GPU, vendor Advanced Micro Devices, Inc., driver 3683.0 (PAL,LC), OpenCL C 2.0 , 1 compute units, 2200 MHz
      global 59589 MiB, max alloc 48909 MiB, local 32 KiB, max work-group 256, sub-group extension: cl_khr_subgroups (no shuffle extension), AMD wavefront width 32
 [1] gfx1201  |  AMD Accelerated Parallel Processing (OpenCL 2.1 AMD-APP (3683.0))
      GPU, vendor Advanced Micro Devices, Inc., driver 3683.0 (PAL,LC), OpenCL C 2.0 , 32 compute units, 2460 MHz
      global 16304 MiB, max alloc 13858 MiB, local 32 KiB, max work-group 256, sub-group extension: cl_khr_subgroups (no shuffle extension), AMD wavefront width 32
 dup [2] gfx1036  |  AMD Accelerated Parallel Processing (OpenCL 2.1 AMD-APP (3652.0)): the same card as [0] on an older platform (driver 3652.0 (PAL,LC)); hidden, use [0]
 dup [3] gfx1201  |  AMD Accelerated Parallel Processing (OpenCL 2.1 AMD-APP (3652.0)): the same card as [1] on an older platform (driver 3652.0 (PAL,LC)); hidden, use [1]
 [4] NVIDIA GeForce RTX 5090  |  NVIDIA CUDA (OpenCL 3.0 CUDA 13.1.1)
      GPU, vendor NVIDIA Corporation, driver 617.14, OpenCL C 1.2 , 170 compute units, 2407 MHz
      global 32606 MiB, max alloc 8151 MiB, local 48 KiB, max work-group 1024, sub-group extension: none, NVIDIA warp size 32
 platforms: 2 device(s) hidden as the same card on an older platform of the same vendor (an old driver's OpenCL registration is still present)
 ";
    #[test]
    fn pc1_list_gives_two_amd_cards_and_two_notes() {
        let (cards, notes) = parse_opencl_list(PC1_LIST);
        assert_eq!(cards.len(), 2, "{cards:?}");
        assert_eq!(cards[0], OpenClListed { index: "0".into(), name: "gfx1036".into(), vendor: "amd", units: "1 compute units".into() });
        assert_eq!(cards[1], OpenClListed { index: "1".into(), name: "gfx1201".into(), vendor: "amd", units: "32 compute units".into() });
        assert_eq!(notes.len(), 2, "{notes:?}");
        assert!(notes[1].contains("[3] gfx1201") && notes[1].contains("older driver"), "{}", notes[1]);
    }
    #[test]
    fn the_worker_default_star_and_a_cpu_device_are_handled() {
        let text = "OpenCL devices (3):\n*[0] gfx1201  |  AMD Accelerated Parallel Processing (OpenCL 2.1 AMD-APP (3683.0))\n      GPU, vendor Advanced Micro Devices, Inc., driver 3683.0 (PAL,LC), OpenCL C 2.0 , 32 compute units, 2460 MHz\n [1] AMD Ryzen 7 9800X3D  |  AMD Accelerated Parallel Processing (OpenCL 2.1 AMD-APP (3683.0))\n      CPU, vendor Advanced Micro Devices, Inc., driver 3683.0, OpenCL C 2.0 , 16 compute units, 4700 MHz\n [2] Intel(R) Arc(TM) A770 Graphics  |  Intel(R) OpenCL Graphics (OpenCL 3.0)\n      GPU, vendor Intel(R) Corporation, driver 32.0.101, OpenCL C 3.0 , 512 compute units, 2400 MHz\n";
        let (cards, notes) = parse_opencl_list(text);
        assert_eq!(cards.iter().map(|c| (c.index.as_str(), c.vendor)).collect::<Vec<_>>(), vec![("0", "amd"), ("2", "other")]);
        assert!(notes.is_empty());
    }
    #[test]
    fn a_repeated_index_is_taken_once() {
        let text = " [1] gfx1201  |  P (v1)\n      GPU, vendor Advanced Micro Devices, Inc., driver 1, OpenCL C 2.0 , 32 compute units, 1 MHz\n [1] gfx1201  |  P (v2)\n      GPU, vendor Advanced Micro Devices, Inc., driver 2, OpenCL C 2.0 , 32 compute units, 1 MHz\n";
        let (cards, _) = parse_opencl_list(text);
        assert_eq!(cards.len(), 1);
    }
 }
 /// Finds the binaries next to the engine (Windows, a plain folder) or in Contents/Resources/bin (macOS bundle).
 pub fn find_bins() -> Result<Bins, String> {
    let exe = std::env::current_exe().map_err(|e| e.to_string())?;
--- a/app/igneum-app/src/engine.rs
+++ b/app/igneum-app/src/engine.rs
@ -3192,9 +3192,17 @@ fn civil_from_days(z: i64) -> (i64, u32, u32) {
    (if m <= 2 { y + 1 } else { y }, m, d)
 }
 /// One pack export at a time (5 October 2026). prepare_worker runs on a thread per card, so two cards starting
 /// together ran two `igneum-miner export-pack` processes into the same folder; across an epoch change they
 /// interleaved and PC 1's packs\devnet was left with one epoch's program.h and the other's seeds.txt, which every
 /// OpenCL worker start then refused ("the epoch seed bytes do not give the pack's IGNEUM_SEEDW_INIT") until the next
 /// export. The second export of a pair rewrites the same pack, which is harmless.
 static EXPORT_LOCK: std::sync::Mutex<()> = std::sync::Mutex::new(());
 /// Exports this hour's program pack from the node to <app data>\packs\devnet (the prebuilt workers read it with --pack).
 #[allow(unused_variables)]
 fn export_pack(shared: &Arc<Shared>, bins: &Bins) -> Result<(), String> {
    let _one_at_a_time = EXPORT_LOCK.lock().unwrap_or_else(|e| e.into_inner());
    let pack = shared.runtime.app_dir.join("packs").join("devnet");
    let _ = std::fs::create_dir_all(&pack);
    let out = crate::detect::run_timeout(std::process::Command::new(&bins.miner).args(["export-pack", &shared.runtime.rpc_url(), &pack.display().to_string()]), None, Duration::from_secs(120)).unwrap_or_default();
@ -3213,6 +3221,7 @@ fn build_worker_from_source(shared: &Arc<Shared>, bins: &Bins, vendor: &str) ->
    #[cfg(windows)]
    {
        use std::process::Command;
        let _one_at_a_time = EXPORT_LOCK.lock().unwrap_or_else(|e| e.into_inner());
        let pack = shared.runtime.app_dir.join("packs").join("devnet");
        let _ = std::fs::create_dir_all(&pack);
        let out = crate::detect::run_timeout(Command::new(&bins.miner).args(["export-pack", &shared.runtime.rpc_url(), &pack.display().to_string()]), None, Duration::from_secs(120)).unwrap_or_default();
--- a/docs/bench-log.md
+++ b/docs/bench-log.md
@ -1524,3 +1524,76 @@ What is measured: one BLS12-381 aggregate signature over 16 summed G1 keys plus
 | on, split 90 s | v3 | 0 / 2 | none / 3 | 278 / 265 | apart | none | 3 on n0 | 2 (n0 reconnected 6 s after the heal, A's chain at about 58 DAA, inside the table) |
 Reading (the NEW finding, ledger C4). With the module off GHOSTDAG alone converges on the heavier chain and the losing side's records re-determine (F24 works when the chain moves). With the module on the overlay holds during the split (A, with 30% of the frozen table, locks nothing; B locks 7 and 8) and then fails at the heal in the shipped node: B's certificates for blocks off n0's chain are "kept pending until the chain decides (no lock at this index)", n0's chain never decides because GHOSTDAG keeps its heavier tip and nothing turns the certificate into a fork-choice constraint, and once n0's last lock (index 7, DAA 209) is one window old (DAA 329) the frozen table stops applying on A's chain ("no frozen table (no lock on this chain inside the window)"), A's two keys are 100% of A's own window (B's post-cut blocks are red there) and n0 locks 10, 11, 12 alone; B's certificates for 10 and 11 then log CONFLICTING on n0 (n0 log, 17:27:04 to 17:29:54 BST). A finality fork from a 96-s honest partition, no attacker, table intact at the heal; the 150-s run and the v2 control end the same way. The spec's fork choice ("GHOSTDAG among tips through all certified checkpoints", 3.5) is therefore implemented only for certificates over blocks already on the node's chain. Fix named in the ledger entry: verify an off-chain certificate against the table at its own block and let it constrain fork choice (a certificate-driven reorg), then re-determine. Raw: `scratchpad fud-a/c4-results-*.md`, node logs `c4-on90-tmp/`, `c4-v2-control-tmp/`.
 ## 5 October 2026 (evening), the 9070 XT on the eGPU: why 17.9 MH/s, and what moved
 PC 1 (ae432dc7, Windows 11, Ryzen 7 9800X3D with its gfx1036, RTX 5090 on CUDA), an AMD Radeon RX 9070 XT (gfx1201, RDNA 4) in a Sonnet Breakaway Box 850T5 over USB4, Adrenalin 26.9.2 (OpenCL driver string `3683.0 (PAL,LC)`, platform `OpenCL 2.1 AMD-APP (3683.0)`). Branch `opencl-rdna4`. the project lead: "the hashrate is low" (17.9 MH/s with one worker; two workers on the card earlier gave 8.9 and 9.4).
 **Before, from PC 1's own app log** (`node tools/logs.mjs win-ae432dc7-20261005-181046`, the miner's STATUS line for the card `amd:1:gfx1201`, 2^21-nonce jobs): `hash=17.82 MH/s wall (17.83 MH/s inside jobs) ... idle=0.3%`. Wall equals inside, so the host loop (template fetch, job line, read-back, scan) costs nothing measurable; the dispatch itself is slow. The worker's `ready` line: `exchange 0` (local memory: AMD lists `cl_khr_subgroups` and no shuffle extension), `batch 4194304`, `dataset-log2 28` (1 GiB), device `[1] gfx1201` on the 3683.0 platform, `AMD wavefront width 32`. The same card was listed again as `[3] gfx1201` on the older platform `3652.0` (the 32.0.21042 driver's OpenCL registration is still present after the update): that is the two-worker run.
 **Hypotheses, each with its number** (the measurement job `rdna4-bench-1`, 18:39:25 to 18:41:17 UTC, the card switched off in the app through `POST /api/cards` for key `amd:1:gfx1201` only, the 5090 untouched; worker exe sha256 `53c7e8c9…5403e10` built from this branch by `proto-cuda/nvrtc/build-windows.sh`; read back with `node tools/jobs.mjs rdna4-bench-1`):
 | # | Hypothesis | Measured | Verdict |
 |---|---|---|---|
 | 1 | The dataset or program is re-sent over the eGPU link per job | Nothing is re-sent: the dataset (1 GiB) and cache (256 MiB) are built on the device once per pair (`info first pack ... cache 11 dataset 51 ms` on the Mac check); per 2^21-nonce job the old path sent 32 B up and read 16 MiB down; the serve A/B below puts a number on that read-back | Not the cause |
 | 2 | Work-group, occupancy, wave width, the exchange | `clGetKernelSubGroupInfoKHR`: sub-group 32 for a 32-item work-group (wave32), private memory 0 (no spills), preferred multiple 32; `--group-warps 1, 2, 4, 8` = 18.024, 18.063, 18.070, 18.039 MH/s (`--batches 3`, 2^24, device event time); `--batch-log2 21` (the app's job size) = 18.108 | Not the cause: the shape does not move the number |
 | 3 | The wrong AMD platform | The app's worker runs on `[1]`, the 3683.0 platform (ready line). The old platform's `[3]` gives 18.049 MH/s: the same. The duplicate listing is real and is the two-worker halving | Not the cause of 17.9; fixed anyway (below) |
 | 4 | The card's own random-read rate | `--memprobe`: dependent random 4-byte loads over 1024 MiB top out at 2.42 to 2.68 G loads/s from 4,096 lanes up (table below); 128 loads per hash gives a ceiling of 18.9 to 20.9 MH/s; the hash runs at 18.0 to 18.1 | THE CAUSE: the hash is at 87 to 95% of what this card does for this access pattern |
 **The memprobe on the 9070 XT** (`igneum-worker-opencl.exe --device 1 --memprobe`, device event time, best of 3, 256 dependent steps per lane; `chase` = one dependent random 4-byte load per step, `indep x8` = eight independent chains per lane):
 | Buffer | Work-group | Lanes in flight | chase G loads/s | ns per dependent load | indep x8 G loads/s |
 |---|---|---|---|---|---|
 | 4 MiB (inside the 8 MB L2, approximate size) | 256 | 4,096 | 34.95 | 117 | |
 | 4 MiB | 256 | 262,144 | 64.63 | 4,056 | 63.8 (262k lanes) |
 | 64 MiB (the 64 MB Infinity Cache, approximate size) | 256 | 4,096 | 9.17 | 447 | |
 | 64 MiB | 256 | 262,144 | 9.18 | 28,561 | 8.8 (262k lanes) |
 | 1024 MiB (GDDR6) | 32 | 4,096 | 2.64 | 1,552 | |
 | 1024 MiB | 32 | 65,536 | 2.60 | 25,181 | |
 | 1024 MiB | 32 | 4,194,304 | 2.43 | 1,729,136 | |
 | 1024 MiB | 256 | 4,096 | 2.64 | 1,552 | |
 | 1024 MiB | 256 | 262,144 | 2.45 | 106,831 | 2.46 (262k lanes) |
 | 1024 MiB | 256 | 4,194,304 | 2.42 | 1,732,023 | 2.42 (4M lanes) |
 | ALU chain, 1,048,576 lanes x 4,096 steps | 256 | | 6,219 G int ops/s (5 ops per step counted, approximate) | | |
 Reading: at the dataset size the card delivers about 2.5 G random 4-byte reads per second whatever the parallelism (4,096 lanes already saturate it; more lanes only queue, the ns column is Little's law on a fixed throughput). Eight independent loads per lane give the same 2.4 G/s, so it is not a latency-hiding problem in the kernel. Inside the Infinity Cache the same chain runs 3.7x faster and inside L2 26x faster, so the cap is the path to GDDR6 for random reads. The ALU chain says the shader clock is not parked (approximate: 6.2 T int ops/s is of the order of 64 CUs x 64 lanes x 2.46 GHz with quarter-rate multiplies).
 **Against the other two cards** (same probe; the 5090 through NVIDIA's OpenCL `[4]` WHILE its CUDA worker was mining, so a lower bound; the Mac through Apple OpenCL, wall time, a Mac at high load, approximate):
 | Card | 1024 MiB chase at 4,096 lanes | 1024 MiB chase ceiling | indep x8 ceiling | ceiling / 128 = hash ceiling | measured hash rate |
 |---|---|---|---|---|---|
 | RX 9070 XT, eGPU over USB4 | 2.64 G/s, 1,552 ns | 2.42 to 2.68 G/s | 2.42 G/s | 18.9 to 20.9 MH/s | 18.0 to 18.1 MH/s (bench), 17.8 (app) |
 | RTX 5090, PCIe 5 x16, contended | 9.09 G/s, 451 ns | 16.4 to 18.0 G/s | 16.2 to 16.7 G/s | 128 to 141 MH/s | 127 MH/s (app, the project lead), 139.7 alone (M11) |
 | Apple M5 Max, Apple OpenCL | 2.10 G/s, 1,949 ns | 3.41 to 3.49 G/s | 3.45 to 3.47 G/s | 26.6 to 27.3 MH/s | 27.9 Mhash/s (README, Apple OpenCL) |
 Reading: on all three cards the hash runs within a few percent of 1/128 of the card's dependent random-read ceiling, which is what a 128-load program should do; the probe is a good model of the hash. The 5090 does 6.6x the random reads of the 9070 XT for 2.8x the rated bandwidth (1,792 against 640 GB/s, vendor figures): the rest is access granularity and DRAM behaviour on random 4-byte reads, which the kernel cannot change.
 **Is it the eGPU link?** No. 2.42 G loads/s x 64 B lines = 155 GB/s of DRAM traffic, forty times what a USB4 PCIe tunnel carries (about 4 GB/s, approximate); the 1 GiB buffer sits in the card's own memory (the 4 and 64 MiB cases show the card's caches at work above it, and a buffer in host memory would run below 0.1 G/s). A PCIe slot would move the per-job read-back (16 MiB per 2^21-nonce job on the old path, now gone) and nothing else; the random-read ceiling is the card's. What a PCIe slot would give: the same 18 MH/s.
 **What changed on `opencl-rdna4`** (`proto-opencl/host.c`, `app/igneum-app/src/detect.rs`):
 | Change | Before | After |
 |---|---|---|
 | Duplicate platform | `--list` showed the card twice ([1] 3683.0 and [3] 3652.0); the app made two cards and ran two workers (8.9 + 9.4 MH/s) | the older platform's entry prints as ` dup [3] ... hidden, use [1]`, the default pick skips it, the app's parser (`parse_opencl_list`, 3 tests) never makes a card of it; `--device 3` still works for comparison. Verified on PC 1: `platforms: 2 device(s) hidden ...`, cards `amd:0:gfx1036` and `amd:1:gfx1201` only |
 | Kernel report | work-group and local memory | plus preferred multiple, private memory (spills), sub-group size on every exchange path (`info kernel:` in serve mode) |
 | Read-back per dispatch | 8 B per nonce (16 MiB per job) and a host scan of 2^21 words | a GPU select pass: the hits (index, hash) behind an atomic counter plus 34 sentinel words; 276 B per chunk plus 16 B per hit; found lines in nonce order; `--readback full` / `IGNEUM_READBACK=full` keeps the old path; a chunk with over 256 hits falls back to the full read |
 | Transfer accounting | none | bytes up and down per chunk and the mean device time of kernel, select, read-back and scan in the stats line every 200 jobs and at quit |
 | `--memprobe` | none | the tables above, no pack needed |
 Correctness: `proto-opencl/test-generic.sh` on the Mac (Apple OpenCL) PASS on both paths: "15 sampled hashes (both packs, both sides of the 32-bit nonce boundary) equal igneum-pow hash-bound"; select path transfers `5 chunks, up 180 B, down 4452 B`, full path `up 160 B, down 1536 B` (the check's jobs are 32 to 64 nonces with every nonce a hit). The bench on the 9070 XT: cache check PASS, dataset self-test PASS, 6 of 6 vector warps PASS, batch fingerprint `3cc4fbf90fa6366c` at 2^24 for the devnet pack (the Apple OpenCL value in the README), at every `--group-warps`.
 **The serve-mode A/B on the card** (job `rdna4-serve-4`, 19:11 UTC, card off in the app, worker exe sha256 `324a6d9b…2bfdfff`; 200 real `job` lines of 2,097,152 nonces each, the app's `--job-nonces`, against the emulator test pack `pack-a` (epoch `edc4fa84…`, self-test PASS, 96 of 96 vector lanes), target `0000100000000000` so that 408 hits fall in 200 jobs on both paths; `done` ms over jobs 11 to 200; `node tools/jobs.mjs rdna4-serve-4`):
 | Read-back | Bytes down per job | Kernel (device, mean) | Select pass | Read-back (wall) | Host scan | Mean job | Inside-job rate |
 |---|---|---|---|---|---|---|---|
 | full (before) | 16,777,216 | 116.12 ms | 0 | 7.28 ms | 0.55 ms | 124.22 ms | 16.88 MH/s |
 | select (after) | 309 | 116.00 ms | 0.039 ms | 0.78 ms | 0.00 ms | 117.38 ms | 17.87 MH/s |
 | select (repeat) | 309 | 115.96 ms | 0.038 ms | 0.76 ms | 0.00 ms | 117.33 ms | 17.87 MH/s |
 Reading: the kernel is the same 116.0 ms on both paths (18.08 MH/s pure kernel, the bench's number). The old path paid 7.8 ms per job for 16 MiB over the eGPU link (2.3 GB/s, the USB4 tunnel's rate; a PCIe slot would read it in about 1 ms, approximate) and the host scan. The select pass removes it: +5.9% per job on this link, nothing on the kernel. Both paths found the same 408 hits. The `--group-warps` and exchange levers were already shown flat above, so this is the whole host-side gain available on the 9070 XT.
 **Probes with a fresh seed per repetition** (the first probe round replayed the same addresses on repeats, so its low-lane rows were cache hits; fixed in `probeLaunch`, job `rdna4-serve-4`): 1024 MiB chase at 256 lanes 276 ns per dependent load, at 1,024 lanes 422 ns, at 4,096 lanes 1,560 ns (2.63 G/s, the cap). Random 64-byte lines (four `uint4` loads per step) at 1024 MiB: 2.46 to 2.88 G lines/s = 158 to 184 GB/s in lines, the same count per second as the 4-byte chase: every random 4-byte read costs this card a 64-byte line fetch. Coalesced stream over the whole 1024 MiB: 635.2 GB/s against the vendor's 640 GB/s, so the memory clock is in its full state and the card is not parked. Inside the 64 MiB buffer the line probe reaches 8.3 to 14.0 G lines/s (533 to 894 GB/s in lines: the Infinity Cache, approximate).
 **A second defect found on the way: the pack export race.** PC 1's app log since its 19:02 UTC restart (`node tools/logs.mjs win-ae432dc7-20261005-190232`): `worker error: error 0 pack packs\devnet: the epoch seed bytes do not give the pack's IGNEUM_SEEDW_INIT` at 19:07:03, 19:07:19 and 19:08:07, so the 9070 XT was not mining at all in the app while this entry was written (my job `rdna4-serve-1` at 18:43 hit the same folder in the same state). Cause, from `app/igneum-app/src/engine.rs` `prepare_worker`: one thread per card, each running `igneum-miner export-pack` into the one folder `packs\devnet`; across an epoch change the two exports interleave and the folder keeps one epoch's `program.h` with the other's `seeds.txt` until the next export. Fix on this branch: a process-wide mutex around both export sites (`EXPORT_LOCK`); the second export rewrites the same pack. Not measured in the app yet: it ships with the branch.
 **Answer to the project lead.** The 9070 XT does 2.5 G random 4-byte reads per second from its memory for this access pattern, and the hash needs 128 of them, so about 19 MH/s is this card's ceiling for the current program class, on any slot; it was running at 92% of that. The eGPU link cost 6% per job through the read-back, now removed (17.87 against 16.88 MH/s inside jobs standalone). The duplicate platform that halved it to 8.9 + 9.4 is folded away. The pack race that stopped it is serialised. Nothing else in the worker's control moves the number: the next step for this card is the program class itself (fewer, wider loads per hash would favour AMD's 64-byte lines), which is a consensus question, not a worker one.
--- a/proto-opencl/README.md
+++ b/proto-opencl/README.md
@ -32,6 +32,25 @@ the driver) and the Khronos headers fetched by `fetch-redist.sh`. `test-generic.
 OpenCL with the two packs `proto-cuda/nvrtc/emu/test.sh` writes (PASS on 4 October 2026: 192 found lines, prepare and
 swap, 15 sampled hashes equal to `igneum-pow hash-bound`). The bench and the compiled-in serve mode are unchanged.
 ## 5 October 2026: the RX 9070 XT (gfx1201, RDNA 4) on PC 1
 Measured in `docs/bench-log.md` ("the 9070 XT on the eGPU"). What changed in `host.c`:
 - `--list` folds the same card listed by two platforms of one vendor (an old driver's OpenCL registration left behind
  after an update: PC 1 had 3652.0 and 3683.0) into one entry; the older one prints as ` dup [N] ...` and the app's
  parser (`detect.rs`) never makes a card of it. Before, the app ran two workers on one 9070 XT at half rate each.
  Indices stay flat, so `--device N` still reaches the hidden entry for a comparison.
 - Every path prints the kernel as the driver compiled it: `kernel: ... preferred multiple, local and private memory,
  sub-group size`. Private memory above 0 means spilled registers. The sub-group size is queried on the local-memory
  path too (the gfx1201 answers 32: wave32).
 - `--serve` reads back the hits and 34 sentinel words through a GPU-side select pass instead of 8 bytes per nonce
  (`--readback full` or `IGNEUM_READBACK=full` keeps the old path). The stats line every 200 jobs carries the bytes up
  and down per chunk and the mean device time of the hash kernel, the select pass, the read-back and the host scan.
 - `--memprobe [--probe-mib N]`: no pack. Dependent random 4-byte loads against lanes in flight (latency and the
  random-read ceiling), eight independent loads per lane, random 64-byte lines, a coalesced stream and an integer
  chain, at 4, 64 and 1024 MiB. The hash is 128 dependent random 4-byte loads, so the 1024 MiB chase ceiling divided
  by 128 is the card's hash-rate ceiling for this program class.
 ## Layout
 ```
@ -39,6 +58,7 @@ proto-opencl/
  host.c          C99 host: device list, runtime kernel build, cache + dataset fill, self-tests, vectors, bench, sweep, --serve, --pack
  cl_dynamic.h    Windows one-click build: OpenCL.dll loaded at run time (IGNEUM_CL_DYNAMIC)
  test-generic.sh the --pack mode checked here through Apple OpenCL (needs proto-cuda/nvrtc/emu/test.sh's packs)
  test_host.c     device-free unit tests of host.c's rules (the duplicate-platform fold); run with test-host.sh
  build.sh        macOS (-framework OpenCL, or the Khronos ICD loader) and Linux (-lOpenCL)
  build.bat       Windows (MSVC cl.exe + OpenCL.lib)
  WAVEFRONT.md    wave32 vs wave64 on AMD, and why the kernel cannot tell the difference
--- a/proto-opencl/WAVEFRONT.md
+++ b/proto-opencl/WAVEFRONT.md
@ -86,6 +86,15 @@ exchange alone would survive a wave64 sub-group (point 1 above); host.c still re
 device because of point 2 and because of `sub_group_broadcast`. The conservative rule costs nothing in correctness
 and, on a wave64 card, the local-memory path is what runs.
 ## RDNA 4, measured (5 October 2026)
 An RX 9070 XT (gfx1201, Adrenalin 26.9.2, OpenCL driver 3683.0 PAL,LC) on PC 1 reports `AMD wavefront width 32`,
 lists `cl_khr_subgroups` without a shuffle extension, and `clGetKernelSubGroupInfoKHR` on `igneum_hash` answers a
 sub-group of 32 for a 32-item work-group: wave32, so a work-group of 32 is one full wave and path 0 (local memory)
 runs with its barriers elided by the compiler. `--group-warps 1, 2, 4, 8` give 18.02, 18.06, 18.07 and 18.04 MH/s, the
 same number: the exchange path and the work-group shape cost nothing there. The rate is set by the card's random-read
 throughput at the dataset size (`docs/bench-log.md`, "the 9070 XT on the eGPU").
 ## What is not proven here
 - No AMD compiler has compiled `kernel.cl`, and no AMD device has run it. The emulator is clang; pocl is LLVM on a
--- a/proto-opencl/host.c
+++ b/proto-opencl/host.c
@ -202,6 +202,11 @@ typedef struct {
    const char* vendor;  // --vendor S: pick the first GPU whose vendor string contains S (default: first GPU of any vendor)
    const char* packDir; // --pack D (serve only): generic mode, the pack is read from D at run time; the compiled-in pack is
                         // then only the build-time placeholder of the prebuilt exe (4 October 2026)
    int readback;        // --readback: 0 select (a GPU-side pass reads back only the hits and 34 sentinel words), 1 full
                         // (every output word comes back, 8 bytes per nonce, the path before 5 October 2026)
    int memprobe;        // --memprobe: dependent-load latency and throughput, independent-load throughput and an ALU
                         // chain on the chosen device, no pack needed (5 October 2026, the 9070 XT on the eGPU)
    int probeMib;        // --probe-mib N: --memprobe at that one buffer size only (default 0 = 4, 64 and 1024 MiB)
 } Options;
 static int packMib(void) { return (int)(((1ull << IGNEUM_DATASET_LOG2) * 4ull) >> 20); }
@ -228,7 +233,11 @@ static void usage(void) {
        "  --no-prepare      with --serve: no prepare support (the miner then falls back to exit 42 at a seed change).\n"
        "                    Builds the pack's kernel_bound.cl (next to the compiled-in kernel.cl) unless --kernel says otherwise.\n"
        "  --pack D          with --serve: serve the pack in directory D (program.h, seeds.txt, vectors.h, kernel_bound.cl), whatever\n"
-        "                    pack this exe was built against; it is self-tested against its vectors.h first (the one-click worker)\n", packMib(), IGNEUM_KERNEL_PATH);
+        "                    pack this exe was built against; it is self-tested against its vectors.h first (the one-click worker)\n"
        "  --readback M      with --serve: select (default) reads back only the hits and 34 sentinel words of each dispatch through a\n"
        "                    GPU-side pass; full reads back every output (8 bytes per nonce). IGNEUM_READBACK=full does the same.\n"
        "  --memprobe        no pack: dependent random loads (latency and throughput against lanes in flight), independent random\n"
        "                    loads and an ALU chain on the chosen device, at 4, 64 and 1024 MiB (--probe-mib N for one size)\n", packMib(), IGNEUM_KERNEL_PATH);
 }
 static int isPow2(long long v) { return v > 0 && (v & (v - 1)) == 0; }
@ -239,6 +248,7 @@ static Options parseArgs(int argc, char** argv) {
    int i;
    o.datasetMib = 1024; o.batchLog2 = 24; o.batches = 5; o.groupWarps = 1; o.sweep = 0; o.device = -1;
    o.exchange = 0; o.list = 0; o.timeWall = -1; o.kernelPath = IGNEUM_KERNEL_PATH; o.extraOpts = ""; o.serve = 0; o.noPrepare = 0; o.kernelGiven = 0; o.vendor = NULL; o.packDir = NULL;
    o.readback = (getenv("IGNEUM_READBACK") && strcmp(getenv("IGNEUM_READBACK"), "full") == 0) ? 1 : 0; o.memprobe = 0; o.probeMib = 0;
    for (i = 1; i < argc; ++i) {
        const char* a = argv[i];
        int needs = (strcmp(a, "--dataset-mib") == 0 || strcmp(a, "--batch-log2") == 0 || strcmp(a, "--batches") == 0 ||
@ -255,6 +265,16 @@ static Options parseArgs(int argc, char** argv) {
        else if (strcmp(a, "--no-prepare") == 0) o.noPrepare = 1;
        else if (strcmp(a, "--vendor") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.vendor = argv[++i]; }
        else if (strcmp(a, "--pack") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.packDir = argv[++i]; }
        else if (strcmp(a, "--readback") == 0) {
            const char* m;
            if (i + 1 >= argc) { usage(); exit(2); }
            m = argv[++i];
            if (strcmp(m, "select") == 0) o.readback = 0;
            else if (strcmp(m, "full") == 0) o.readback = 1;
            else { printf("--readback must be select or full\n"); exit(2); }
        }
        else if (strcmp(a, "--memprobe") == 0) o.memprobe = 1;
        else if (strcmp(a, "--probe-mib") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.probeMib = atoi(argv[++i]); }
        else if (strcmp(a, "--build-opts") == 0) o.extraOpts = argv[++i];
        else if (strcmp(a, "--time") == 0) {
            const char* m = argv[++i];
@ -297,8 +317,49 @@ typedef struct {
    int cMajor, cMinor;          // OpenCL C version
    int dMajor, dMinor;          // device (platform profile) version
    cl_uint amdWavefront, nvWarp; // 0 if not reported
    int dupOf;                   // index of the same card on a newer platform of the same vendor, -1 if none (5 October 2026)
 } DeviceInfo;
 /* The driver version as a number for ordering ("3683.0 (PAL,LC)" -> 3683.0; "617.14" -> 617.14; 0 when unreadable). */
 static double driverNumber(const char* driver) {
    const char* p = driver;
    while (*p && (*p < '0' || *p > '9')) ++p;
    return *p ? strtod(p, NULL) : 0.0;
 }
 /* Two AMD platforms are registered after a driver update on Windows (PC 1, 5 October 2026: 32.0.21042 and 32.0.32015,
 * OpenCL driver strings 3652.0 and 3683.0): every card is listed twice, the app ran two workers on one 9070 XT, and
 * each got half. The same card on the same-named platform with a different platform version is the one card; the
 * entry whose driver number is lower is the duplicate. Two real cards of one model sit on the SAME platform and are
 * never folded. Indices stay flat (the app passes them back as --device). Returns how many duplicates were marked. */
 static int markDuplicates(DeviceInfo* list, int n) {
    int i, j, marked = 0;
    for (i = 0; i < n; ++i) list[i].dupOf = -1;
    for (i = 0; i < n; ++i) {
        if (list[i].dupOf >= 0) continue;
        for (j = i + 1; j < n; ++j) {
            int older;
            if (list[j].dupOf >= 0) continue;
            if (strcmp(list[i].platformName, list[j].platformName) != 0) continue;
            if (strcmp(list[i].platformVersion, list[j].platformVersion) == 0) continue;
            if (strcmp(list[i].name, list[j].name) != 0 || strcmp(list[i].vendor, list[j].vendor) != 0) continue;
            if (list[i].globalMem != list[j].globalMem || list[i].computeUnits != list[j].computeUnits) continue;
            older = driverNumber(list[j].driver) < driverNumber(list[i].driver) ? j : i;
            if (older == j) { list[j].dupOf = i; }
            else { list[i].dupOf = j; }
            ++marked;
            if (older == i) break;   /* i itself is the duplicate; j stays the real one */
        }
    }
    /* three registrations of one card: every duplicate points at the one that stays, not at another duplicate */
    for (i = 0; i < n; ++i) {
        int k = list[i].dupOf, hops = 0;
        while (k >= 0 && list[k].dupOf >= 0 && hops++ < n) k = list[k].dupOf;
        if (list[i].dupOf >= 0) list[i].dupOf = k;
    }
    return marked;
 }
 static void devStr(cl_device_id d, cl_device_info what, char* out, size_t n) {
    out[0] = 0;
    clGetDeviceInfo(d, what, n - 1, out, NULL);
@ -357,6 +418,7 @@ static int enumerateDevices(DeviceInfo** outList) {
            list[n++] = di;
        }
    }
    if (n) markDuplicates(list, n);
    *outList = list;
    return n;
 }
@ -365,6 +427,13 @@ static void printDevice(int idx, const DeviceInfo* d, int chosen) {
    const char* subExt = strstr(d->extensions, "cl_khr_subgroup_shuffle") ? "cl_khr_subgroup_shuffle" :
                         strstr(d->extensions, "cl_intel_subgroups") ? "cl_intel_subgroups" :
                         strstr(d->extensions, "cl_khr_subgroups") ? "cl_khr_subgroups (no shuffle extension)" : "none";
    if (d->dupOf >= 0) {
        /* Hidden from the app's card list: its parser takes only lines that start with "[" (detect.rs). The index
         * is still valid for --device, so a run on the older platform stays possible for a comparison. */
        printf(" dup [%d] %s  |  %s (%s): the same card as [%d] on an older platform (driver %s); hidden, use [%d]\n",
               idx, d->name, d->platformName, d->platformVersion, d->dupOf, d->driver, d->dupOf);
        return;
    }
    printf("%s[%d] %s  |  %s (%s)\n", chosen ? "*" : " ", idx, d->name, d->platformName, d->platformVersion);
    printf("      %s, vendor %s, driver %s, %s, %u compute units, %u MHz\n", typeName(d->type), d->vendor, d->driver, d->cVersion, d->computeUnits, d->clockMHz);
    printf("      global %llu MiB, max alloc %llu MiB, local %llu KiB, max work-group %llu, sub-group extension: %s",
@ -457,19 +526,45 @@ static void releaseProgram(Device* dv) {
    dv->kHash = dv->kCacheFill = dv->kBuild = dv->kFill = dv->kHashBound = NULL; dv->prog = NULL;
 }
-// Sub-group size of igneum_hash for a work-group of `local` items. 0 if the query is unavailable (reason in *why).
+// Sub-group size of kernel k for a work-group of `local` items. 0 if the query is unavailable (reason in *why).
-static size_t querySubGroupSize(const Device* dv, const DeviceInfo* di, size_t local, char* why, size_t whyLen) {
+static size_t querySubGroupSizeOf(cl_kernel k, const DeviceInfo* di, size_t local, char* why, size_t whyLen) {
    ig_pfn_subgroup_info fn = (ig_pfn_subgroup_info)clGetExtensionFunctionAddressForPlatform(di->platform, "clGetKernelSubGroupInfoKHR");
    const char* via = "clGetKernelSubGroupInfoKHR";
    size_t sg = 0;
    cl_int e;
    if (!fn) { fn = (ig_pfn_subgroup_info)loadSym("clGetKernelSubGroupInfo"); via = "clGetKernelSubGroupInfo (OpenCL 2.1 core, through the loader)"; }
    if (!fn) { snprintf(why, whyLen, "the sub-group size could not be queried (neither clGetKernelSubGroupInfoKHR nor clGetKernelSubGroupInfo is available)"); return 0; }
-    e = fn(dv->kHash, di->device, IG_CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE, sizeof(local), &local, sizeof(sg), &sg, NULL);
+    e = fn(k, di->device, IG_CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE, sizeof(local), &local, sizeof(sg), &sg, NULL);
    if (e != CL_SUCCESS) { snprintf(why, whyLen, "the sub-group size query failed: %s returned %s (%d)", via, clErrName(e), (int)e); return 0; }
    snprintf(why, whyLen, "queried through %s", via);
    return sg;
 }
 static size_t querySubGroupSize(const Device* dv, const DeviceInfo* di, size_t local, char* why, size_t whyLen) {
    return querySubGroupSizeOf(dv->kHash, di, local, why, whyLen);
 }
 /* The compiled kernel as the driver sees it, on every exchange path (5 October 2026, the 9070 XT question): the
 * work-group limit, the preferred multiple (the wave width the compiler chose), local memory, private memory (scratch:
 * anything above 0 means spilled registers, which on AMD costs a memory round trip per spill), and the sub-group size
 * for the built work-group size (wave32 or wave64 on RDNA; the local-memory path never queried it before). */
 #define IG_CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 0x11B3
 #define IG_CL_KERNEL_PRIVATE_MEM_SIZE 0x11B4
 static void printKernelInfo(const DeviceInfo* di, cl_kernel k, const char* kernelName, int groupSize, const char* prefix) {
    size_t wg = 0, mult = 0;
    cl_ulong lmem = 0, pmem = 0;
    char why[256];
    size_t sg;
    clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL);
    clGetKernelWorkGroupInfo(k, di->device, IG_CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(mult), &mult, NULL);
    clGetKernelWorkGroupInfo(k, di->device, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(lmem), &lmem, NULL);
    clGetKernelWorkGroupInfo(k, di->device, IG_CL_KERNEL_PRIVATE_MEM_SIZE, sizeof(pmem), &pmem, NULL);
    sg = querySubGroupSizeOf(k, di, (size_t)groupSize, why, sizeof(why));
    printf("%skernel: %s max work-group %llu, preferred multiple %llu, local memory %llu bytes, private memory %llu bytes%s, work-group %d, sub-group size %llu (%s)%s\n",
           prefix, kernelName, (unsigned long long)wg, (unsigned long long)mult, (unsigned long long)lmem, (unsigned long long)pmem,
           pmem ? " (SPILLED: registers in scratch memory)" : "", groupSize, (unsigned long long)sg, why,
           (sg > (size_t)groupSize) ? " (the work-group fills only part of a wave: see --group-warps)" : "");
    fflush(stdout);
 }
 // Decide the exchange implementation and build. See WAVEFRONT.md for the rule.
 static void setupProgram(Device* dv, const DeviceInfo* di, const Options* o, const char* src, size_t srcLen) {
@ -1166,6 +1261,61 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
    hOut = (uint64_t*)malloc((size_t)batch * sizeof(uint64_t));
    strncpy(devName, di->name, 255); devName[255] = 0;
    for (k = 0; devName[k]; ++k) if (devName[k] == ' ') devName[k] = '_';
    printKernelInfo(di, cur->kHashBound, "igneum_hash_bound", (int)groupSize, "info ");
    /* The select pass (5 October 2026). Before it every dispatch read back 8 bytes per nonce (16 MiB for a 2^21-nonce
     * job) over the bus and scanned them on the host; on an eGPU over USB4 that is a measurable part of every job.
     * Now a tiny kernel built here (no pack involved) writes the hits (index, hash) behind an atomic counter plus 34
     * sentinel words (the first 32 outputs, the middle and the last), and the host reads back a few hundred bytes.
     * The fault detectors read the sentinels; the found lines are printed in nonce order from the sorted hits. If a
     * chunk has more hits than the table holds (a target that loose is a test, not a block), the chunk falls back to the
     * full read. --readback full or IGNEUM_READBACK=full keeps the old path for a comparison. */
 #define IG_MAX_HITS 256u
    cl_program selProg = NULL;
    cl_kernel kSelect = NULL;
    cl_mem dCount = NULL, dHits = NULL, dSentinel = NULL;
    size_t selLocal = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
    uint32_t hCount = 0;
    uint64_t hHits[IG_MAX_HITS * 2];
    uint64_t hSentinel[34];
    unsigned long long bytesUp = 0, bytesDown = 0;
    double kernelMsSum = 0, selectMsSum = 0, readMsSum = 0, scanMsSum = 0;
    unsigned long fullFallbacks = 0;
    if (!o->readback) {
        static const char* SELECT_SRC =
            "__kernel void igneum_select(__global const ulong* out, uint n, ulong target, volatile __global uint* count,\n"
            "                            __global ulong* hits, uint maxHits, __global ulong* sentinel) {\n"
            "    uint i = (uint)get_global_id(0);\n"
            "    if (i < n) {\n"
            "        ulong h = out[i];\n"
            "        if (h <= target) { uint k = atomic_inc(count); if (k < maxHits) { hits[2u * k] = (ulong)i; hits[2u * k + 1u] = h; } }\n"
            "        if (i < 32u) sentinel[i] = h;\n"
            "        if (i == 0u) { sentinel[32] = out[n / 2u]; sentinel[33] = out[n - 1u]; }\n"
            "    }\n"
            "}\n";
        size_t selLen = strlen(SELECT_SRC);
        selProg = clCreateProgramWithSource(dv->ctx, 1, &SELECT_SRC, &selLen, &err); CL_CHECK_ERR(err, "clCreateProgramWithSource select");
        err = clBuildProgram(selProg, 1, &di->device, "-cl-std=CL1.2", NULL, NULL);
        if (err != CL_SUCCESS) {
            size_t logLen = 0; char* log;
            clGetProgramBuildInfo(selProg, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
            log = (char*)calloc(logLen + 1, 1);
            if (logLen) clGetProgramBuildInfo(selProg, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
            printf("info the select pass did not build (%s): %.300s; using the full read-back\n", clErrName(err), log);
            free(log); clReleaseProgram(selProg); selProg = NULL; ((Options*)o)->readback = 1;
        } else {
            kSelect = clCreateKernel(selProg, "igneum_select", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_select");
            dCount = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, 4, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer count");
            dHits = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, IG_MAX_HITS * 2 * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer hits");
            dSentinel = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, 34 * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer sentinel");
            gMemCreated += 3;
            {
                size_t wg = 0;
                if (clGetKernelWorkGroupInfo(kSelect, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL) == CL_SUCCESS && wg && wg < selLocal) selLocal = wg;
            }
        }
    }
    printf("info readback %s (per dispatch of %u nonces: %s)\n", o->readback ? "full" : "select",
           batch, o->readback ? "8 bytes per nonce come back and the host scans them" : "the hits and 34 sentinel words come back; the GPU scans");
    printf("ready opencl %s platform %s pack %s dataset-log2 %d batch %u exchange %d prepare %d path %s\n", devName, di->platformName,
           gGeneric ? gPack.seedString : IGNEUM_SEED_STRING, log2u32(words), batch, dv->exchange, o->noPrepare ? 0 : 1, gGeneric ? "prebuilt-generic" : "compiled-in");
    fflush(stdout);
@ -1178,6 +1328,7 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
    long gFaultTestChunk = getenv("IGNEUM_FAULT_TEST") ? atol(getenv("IGNEUM_FAULT_TEST")) : -1;
    double meanNsPerNonce = 0;      /* running mean of wall ns per nonce over the chunks so far */
    unsigned long chunksSeen = 0, jobsSeen = 0;
    unsigned long long noncesSeen = 0;      /* for the mean chunk in the stats line */
    uint64_t prevSig = 0; int havePrevSig = 0;
    double jobMsSum = 0;
 #define SERVE_FATAL(jobIdStr, fmt, ...) do { printf("error %s worker fault: " fmt "; exiting 3 so the miner restarts the worker\n", jobIdStr, __VA_ARGS__); fflush(stdout); exit(3); } while (0)
@ -1278,11 +1429,15 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
            b[45] = (uint8_t)hi; b[46] = (uint8_t)(hi >> 8); b[47] = (uint8_t)(hi >> 16); b[48] = (uint8_t)(hi >> 24);
            seedWordsFromBytes(b, 49, iw);
            {
-            double c0 = wallMs(), cms;
+            double c0 = wallMs(), cms, kernelMs = -1.0, selectMs = 0.0, readMs = 0.0, scanMs = 0.0, r0;
            cl_int status = 0;
            size_t g = ((chunk + groupSize - 1) / groupSize) * groupSize;
            uint64_t sig;
            int useSelect = (kSelect != NULL);
            uint32_t nHits = 0;
            SERVE_CHECK(jobId, clEnqueueWriteBuffer(dv->q, dInit, CL_TRUE, 0, 32, iw, 0, NULL, NULL));
            bytesUp += 32;
            if (useSelect) { hCount = 0; SERVE_CHECK(jobId, clEnqueueWriteBuffer(dv->q, dCount, CL_TRUE, 0, 4, &hCount, 0, NULL, NULL)); bytesUp += 4; }
            baseNonce = lo;
            SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 0, sizeof(cl_mem), &cur->ds));
            SERVE_CHECK(jobId, clSetKernelArg(cur->kHashBound, 1, sizeof(cl_mem), &dOut));
@ -1302,11 +1457,50 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
            err = clWaitForEvents(1, &ev);
            if (err != CL_SUCCESS) { countRelease(ev); SERVE_FATAL(jobId, "clWaitForEvents on the dispatch returned %s (%d)", clErrName(err), (int)err); }
            err = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(status), &status, NULL);
            kernelMs = eventMs(ev);
            countRelease(ev);
            if (err != CL_SUCCESS) SERVE_FATAL(jobId, "clGetEventInfo on the dispatch returned %s (%d)", clErrName(err), (int)err);
            if (status != CL_COMPLETE) SERVE_FATAL(jobId, "the dispatch event ended with status %d, not CL_COMPLETE (a device reset or a lost context)", (int)status);
            if (useSelect) {
                cl_event evs = NULL;
                cl_uint nArg = chunk, maxArg = IG_MAX_HITS;
                cl_ulong tArg = (cl_ulong)target;
                size_t gs = ((chunk + selLocal - 1) / selLocal) * selLocal;
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 0, sizeof(cl_mem), &dOut));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 1, sizeof(cl_uint), &nArg));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 2, sizeof(cl_ulong), &tArg));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 3, sizeof(cl_mem), &dCount));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 4, sizeof(cl_mem), &dHits));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 5, sizeof(cl_uint), &maxArg));
                SERVE_CHECK(jobId, clSetKernelArg(kSelect, 6, sizeof(cl_mem), &dSentinel));
                SERVE_CHECK(jobId, clEnqueueNDRangeKernel(dv->q, kSelect, 1, NULL, &gs, &selLocal, 0, NULL, &evs));
                ++gEvCreated;
                err = clWaitForEvents(1, &evs);
                if (err != CL_SUCCESS) { countRelease(evs); SERVE_FATAL(jobId, "clWaitForEvents on the select pass returned %s (%d)", clErrName(err), (int)err); }
                selectMs = eventMs(evs);
                countRelease(evs);
            }
-            SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)chunk * sizeof(uint64_t), hOut, 0, NULL, NULL));
+            }
            r0 = wallMs();
            if (useSelect) {
                SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dCount, CL_TRUE, 0, 4, &hCount, 0, NULL, NULL));
                bytesDown += 4;
                nHits = hCount;
                if (nHits > IG_MAX_HITS) {
                    /* more hits than the table holds: this chunk takes the full path (the test target case) */
                    ++fullFallbacks;
                    useSelect = 0;
                } else {
                    if (nHits) { SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dHits, CL_TRUE, 0, (size_t)nHits * 2 * sizeof(uint64_t), hHits, 0, NULL, NULL)); bytesDown += (unsigned long long)nHits * 16; }
                    SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dSentinel, CL_TRUE, 0, 34 * sizeof(uint64_t), hSentinel, 0, NULL, NULL));
                    bytesDown += 34 * 8;
                }
            }
            if (!useSelect) {
                SERVE_CHECK(jobId, clEnqueueReadBuffer(dv->q, dOut, CL_TRUE, 0, (size_t)chunk * sizeof(uint64_t), hOut, 0, NULL, NULL));
                bytesDown += (unsigned long long)chunk * 8;
            }
            readMs = wallMs() - r0;
            cms = wallMs() - c0;
            /* Plausibility: wall time per nonce against the running mean (the first chunk sets it; a chunk is a full
             * batch except at the 32-bit boundary, so per nonce is the comparable unit). 20x faster = the kernel did not run. */
@ -1315,18 +1509,39 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
                if (chunksSeen >= 4 && meanNsPerNonce > 0 && ns * 20.0 < meanNsPerNonce)
                    SERVE_FATAL(jobId, "%u nonces reported complete in %.3f ms, %.0fx faster than the running mean of %.2f ms per million (the runtime is not running the kernel)", chunk, cms, meanNsPerNonce / ns, meanNsPerNonce / 1e3);
                meanNsPerNonce = chunksSeen == 0 ? ns : meanNsPerNonce + (ns - meanNsPerNonce) / (double)(chunksSeen + 1 < 64 ? chunksSeen + 1 : 64);
-                ++chunksSeen;
+                ++chunksSeen; noncesSeen += chunk;
            }
-            /* Stale output: the first, middle and last words plus an FNV of the first 32. Two chunks never share it. */
+            /* Stale output: the first, middle and last words plus an FNV of the first 32. Two chunks never share it.
-            sig = fnv1a64(hOut, 32 * sizeof(uint64_t)) ^ hOut[chunk / 2] ^ hOut[chunk - 1];
+             * On the select path the same 34 words come from the sentinel buffer the select pass wrote. */
            if (useSelect) sig = fnv1a64(hSentinel, 32 * sizeof(uint64_t)) ^ hSentinel[32] ^ hSentinel[33];
            else sig = fnv1a64(hOut, 32 * sizeof(uint64_t)) ^ hOut[chunk / 2] ^ hOut[chunk - 1];
            if (chunk >= 64) {
                if (havePrevSig && sig == prevSig) SERVE_FATAL(jobId, "the output buffer is unchanged since the previous dispatch (signature %016llx): the kernel did not run", (unsigned long long)sig);
                prevSig = sig; havePrevSig = 1;
            }
            r0 = wallMs();
            if (useSelect) {
                /* nonce order, as the full scan printed them (the miner submits the first found of a job) */
                uint32_t a, b;
                for (a = 1; a < nHits; ++a) {
                    uint64_t ki = hHits[2 * a], kh = hHits[2 * a + 1];
                    for (b = a; b > 0 && hHits[2 * (b - 1)] > ki; --b) { hHits[2 * b] = hHits[2 * (b - 1)]; hHits[2 * b + 1] = hHits[2 * (b - 1) + 1]; }
                    hHits[2 * b] = ki; hHits[2 * b + 1] = kh;
                }
                for (a = 0; a < nHits; ++a) {
                    uint32_t idx = (uint32_t)hHits[2 * a];
                    unsigned long long nonce = ((unsigned long long)hi << 32) | (unsigned long long)(uint32_t)(lo + idx);
                    printf("found %s %llu %016llx\n", jobId, nonce, (unsigned long long)hHits[2 * a + 1]);
                }
            } else {
                for (i = 0; i < chunk; ++i) if (hOut[i] <= target) {
                    unsigned long long nonce = ((unsigned long long)hi << 32) | (unsigned long long)(uint32_t)(lo + i);
                    printf("found %s %llu %016llx\n", jobId, nonce, (unsigned long long)hOut[i]);
                }
            }
-            for (i = 0; i < chunk; ++i) if (hOut[i] <= target) {
+            scanMs = wallMs() - r0;
-                unsigned long long nonce = ((unsigned long long)hi << 32) | (unsigned long long)(uint32_t)(lo + i);
+            if (kernelMs > 0) kernelMsSum += kernelMs;
-                printf("found %s %llu %016llx\n", jobId, nonce, (unsigned long long)hOut[i]);
+            selectMsSum += selectMs; readMsSum += readMs; scanMsSum += scanMs;
            }
            fflush(stdout);
            hashes += chunk;
@ -1340,6 +1555,12 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
        if (jobsSeen % 200 == 0) {
            printf("info stats jobs %lu chunks %lu mean job %.1f ms mean %.2f ms per million nonces; events created %lu released %lu live %lu; buffers created %lu released %lu live %lu\n",
                   jobsSeen, chunksSeen, jobMsSum / (double)jobsSeen, meanNsPerNonce / 1e3, gEvCreated, gEvReleased, gEvCreated - gEvReleased, gMemCreated, gMemReleased, gMemCreated - gMemReleased);
            /* The transfer and time budget per chunk (5 October 2026): bytes up (init words, the count reset), bytes
             * down (hits and sentinels, or 8 bytes per nonce on the full path), and the mean device time of the hash
             * kernel (event profiling), the select pass, the blocking read-back and the host scan. */
            printf("info transfers per chunk: up %.0f B, down %.0f B (%s, %lu full fallbacks); mean per chunk: kernel %.2f ms (device), select %.3f ms, read-back %.2f ms, scan %.2f ms, chunk wall %.2f ms\n",
                   (double)bytesUp / (double)chunksSeen, (double)bytesDown / (double)chunksSeen, kSelect ? "select" : "full", fullFallbacks,
                   kernelMsSum / (double)chunksSeen, selectMsSum / (double)chunksSeen, readMsSum / (double)chunksSeen, scanMsSum / (double)chunksSeen, meanNsPerNonce * ((double)noncesSeen / (double)chunksSeen) / 1e6);
            if (gEvCreated - gEvReleased > 16 || gMemCreated - gMemReleased > 12) SERVE_FATAL(jobId, "object leak: %lu events and %lu buffers live after %lu jobs", gEvCreated - gEvReleased, gMemCreated - gMemReleased, jobsSeen);
        }
        if (switched && old) { releasePair(old); old = NULL; printf("info dropped the previous pair (its program, cache and dataset)\n"); }
@ -1349,6 +1570,9 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
    clReleaseMemObject(dInit);
    clReleaseMemObject(dOut);
    gMemReleased += 2;
    if (kSelect) { clReleaseKernel(kSelect); clReleaseMemObject(dCount); clReleaseMemObject(dHits); clReleaseMemObject(dSentinel); gMemReleased += 3; }
    if (selProg) clReleaseProgram(selProg);
    if (chunksSeen) printf("info transfers total: %lu chunks, up %llu B, down %llu B (%s, %lu full fallbacks)\n", chunksSeen, bytesUp, bytesDown, kSelect ? "select" : "full", fullFallbacks);
    if (old) releasePair(old);
    if (prepared) releasePair(prepared);
    releasePair(cur);
@ -1356,6 +1580,224 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) {
 #endif
 }
 /* ---------------------------------------------------------------------------------------------
 * --memprobe (5 October 2026): what the device itself can do with the access pattern of the hash, with no pack and no
 * program. Three kernels built from the text below:
 *   chase  one dependent random 4-byte load per step per lane (the next address comes from the loaded word), so the
 *          time per step at a small lane count is the loaded-latency of one random read, and the loads per second at a
 *          large lane count is the device's random-read throughput for a dependent chain (the hash is 128 of these)
 *   indep  eight independent chains per lane: the throughput when latency is hidden inside one lane
 *   alu    an integer multiply-add-rotate chain with no memory: the achieved integer rate, which moves with the clock
 * Sizes 4 MiB (cache-resident), 64 MiB (the last-level cache on RDNA 3 and 4 is 64 MiB or more, approximate) and
 * 1024 MiB (the dataset size: DRAM). Times from event profiling (wall on Apple). Every number is printed with the
 * configuration that produced it. Rates are G loads/s = 1e9 loads per second; ns per load = time / steps.
 */
 static const char* PROBE_SRC =
    "static inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
    "__kernel void probe_fill(__global uint* ds, uint n) { uint i = (uint)get_global_id(0); if (i < n) ds[i] = pm_mix(i ^ 0x9E3779B9u); }\n"
    "__kernel void probe_chase(__global const uint* ds, uint mask, uint steps, uint seed, __global uint* out) {\n"
    "    uint x = pm_mix((uint)get_global_id(0) ^ seed);\n"
    "    for (uint s = 0u; s < steps; ++s) x = ds[x & mask] ^ (x * 0x9E3779B1u + s);\n"
    "    out[get_global_id(0)] = x;\n"
    "}\n"
    "__kernel void probe_indep(__global const uint* ds, uint mask, uint steps, uint seed, __global uint* out) {\n"
    "    uint g = (uint)get_global_id(0);\n"
    "    uint x0 = pm_mix(g * 8u ^ seed), x1 = pm_mix((g * 8u + 1u) ^ seed), x2 = pm_mix((g * 8u + 2u) ^ seed), x3 = pm_mix((g * 8u + 3u) ^ seed);\n"
    "    uint x4 = pm_mix((g * 8u + 4u) ^ seed), x5 = pm_mix((g * 8u + 5u) ^ seed), x6 = pm_mix((g * 8u + 6u) ^ seed), x7 = pm_mix((g * 8u + 7u) ^ seed);\n"
    "    for (uint s = 0u; s < steps; ++s) {\n"
    "        x0 = ds[x0 & mask] ^ (x0 * 0x9E3779B1u + s); x1 = ds[x1 & mask] ^ (x1 * 0x9E3779B1u + s);\n"
    "        x2 = ds[x2 & mask] ^ (x2 * 0x9E3779B1u + s); x3 = ds[x3 & mask] ^ (x3 * 0x9E3779B1u + s);\n"
    "        x4 = ds[x4 & mask] ^ (x4 * 0x9E3779B1u + s); x5 = ds[x5 & mask] ^ (x5 * 0x9E3779B1u + s);\n"
    "        x6 = ds[x6 & mask] ^ (x6 * 0x9E3779B1u + s); x7 = ds[x7 & mask] ^ (x7 * 0x9E3779B1u + s);\n"
    "    }\n"
    "    out[g] = x0 ^ x1 ^ x2 ^ x3 ^ x4 ^ x5 ^ x6 ^ x7;\n"
    "}\n"
    "__kernel void probe_line(__global const uint4* ds, uint lineMask, uint steps, uint seed, __global uint* out) {\n"
    "    uint x = pm_mix((uint)get_global_id(0) ^ seed);\n"
    "    for (uint s = 0u; s < steps; ++s) {\n"
    "        uint l = (x & lineMask) * 4u;\n"
    "        uint4 a = ds[l], b = ds[l + 1u], c = ds[l + 2u], d = ds[l + 3u];\n"
    "        x = (a.x ^ b.y ^ c.z ^ d.w) ^ (x * 0x9E3779B1u + s);\n"
    "    }\n"
    "    out[get_global_id(0)] = x;\n"
    "}\n"
    "__kernel void probe_stream(__global const uint4* ds, uint perLane, __global uint* out) {\n"
    "    uint g = (uint)get_global_id(0), n = (uint)get_global_size(0);\n"
    "    uint4 acc = (uint4)(0u, 0u, 0u, 0u);\n"
    "    for (uint s = 0u; s < perLane; ++s) acc ^= ds[s * n + g];\n"
    "    out[g] = acc.x ^ acc.y ^ acc.z ^ acc.w;\n"
    "}\n"
    "__kernel void probe_alu(uint steps, uint seed, __global uint* out) {\n"
    "    uint g = (uint)get_global_id(0); uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n"
    "    for (uint s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + rotate(y, 7u); y = (y ^ x) + s; }\n"
    "    out[g] = x ^ y;\n"
    "}\n";
 /* One timed launch, best of `reps`, in ms (event time, or wall when the platform's events are unusable). */
 static double probeLaunch(Device* dv, const Options* o, cl_kernel k, size_t global, size_t local, int reps, int seedArg, cl_uint seed) {
    double best = -1.0;
    int r;
    for (r = 0; r < reps; ++r) {
        cl_event ev = NULL;
        double w0 = wallMs(), ms;
        /* a fresh seed per repetition: a replay of the same addresses would be served from the last-level cache
         * (65,536 reads x 64 B = 4 MiB fits any of them) and read as DRAM latency (seen on the 9070 XT, 5 October) */
        if (seedArg >= 0) { cl_uint s = seed + (cl_uint)r * 0x9E3779B9u; CL_CHECK(clSetKernelArg(k, (cl_uint)seedArg, sizeof(cl_uint), &s)); }
        CL_CHECK(clEnqueueNDRangeKernel(dv->q, k, 1, NULL, &global, &local, 0, NULL, &ev));
        CL_CHECK(clWaitForEvents(1, &ev));
        ms = o->timeWall ? wallMs() - w0 : eventMs(ev);
        if (ms < 0) ms = wallMs() - w0;
        clReleaseEvent(ev);
        if (best < 0 || ms < best) best = ms;
    }
    return best;
 }
 static int runMemprobe(Device* dv, const DeviceInfo* di, const Options* o) {
    cl_int err = 0;
    cl_program prog;
    cl_kernel kFill, kChase, kIndep, kAlu, kLine, kStream;
    size_t srcLen = strlen(PROBE_SRC);
    int sizes[3] = { 4, 64, 1024 }, nSizes = 3, si;
    size_t lanesList[8] = { 256, 1024, 1u << 12, 1u << 14, 1u << 16, 1u << 18, 1u << 20, 1u << 22 };
    const int nLanes = 8;
    size_t groups[2] = { 32, 256 };
    const cl_uint STEPS = 256u, ALU_STEPS = 4096u;
    const size_t maxLanes = 1u << 22;
    cl_mem dOut;
    if (o->probeMib > 0) { sizes[0] = o->probeMib; nSizes = 1; }
    prog = clCreateProgramWithSource(dv->ctx, 1, &PROBE_SRC, &srcLen, &err); CL_CHECK_ERR(err, "clCreateProgramWithSource probe");
    err = clBuildProgram(prog, 1, &di->device, "-cl-std=CL1.2", NULL, NULL);
    if (err != CL_SUCCESS) {
        size_t logLen = 0; char* log;
        clGetProgramBuildInfo(prog, di->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
        log = (char*)calloc(logLen + 1, 1);
        if (logLen) clGetProgramBuildInfo(prog, di->device, CL_PROGRAM_BUILD_LOG, logLen, log, NULL);
        printf("memprobe: build FAILED (%s)\n--- build log ---\n%s\n--- end of build log ---\n", clErrName(err), log);
        free(log);
        return 2;
    }
    kFill = clCreateKernel(prog, "probe_fill", &err); CL_CHECK_ERR(err, "probe_fill");
    kChase = clCreateKernel(prog, "probe_chase", &err); CL_CHECK_ERR(err, "probe_chase");
    kIndep = clCreateKernel(prog, "probe_indep", &err); CL_CHECK_ERR(err, "probe_indep");
    kAlu = clCreateKernel(prog, "probe_alu", &err); CL_CHECK_ERR(err, "probe_alu");
    kLine = clCreateKernel(prog, "probe_line", &err); CL_CHECK_ERR(err, "probe_line");
    kStream = clCreateKernel(prog, "probe_stream", &err); CL_CHECK_ERR(err, "probe_stream");
    dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, maxLanes * 4u, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer probe out");
    printf("memprobe on [%s] %s, driver %s, %u compute units, %u MHz, %s time\n", di->platformName, di->name, di->driver, di->computeUnits, di->clockMHz, o->timeWall ? "wall" : "device event");
    printKernelInfo(di, kChase, "probe_chase", 32, "memprobe ");
    printKernelInfo(di, kChase, "probe_chase", 256, "memprobe ");
    printf("| probe | MiB | work-group | lanes in flight | steps per lane | best ms | G loads/s | ns per dependent load |\n|---|---|---|---|---|---|---|---|\n");
    for (si = 0; si < nSizes; ++si) {
        int mib = sizes[si];
        uint64_t bytes = (uint64_t)mib << 20;
        cl_uint words = (cl_uint)(bytes / 4ull), mask = words - 1u, n = words;
        cl_mem dDs;
        size_t gi, li;
        size_t fillLocal = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
        size_t fillGlobal = ((size_t)words + fillLocal - 1) / fillLocal * fillLocal;
        if ((uint64_t)di->maxAlloc < bytes) { printf("| chase | %d | skipped: max alloc %llu MiB | | | | | |\n", mib, (unsigned long long)(di->maxAlloc >> 20)); continue; }
        dDs = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)bytes, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer probe dataset");
        CL_CHECK(clSetKernelArg(kFill, 0, sizeof(cl_mem), &dDs));
        CL_CHECK(clSetKernelArg(kFill, 1, sizeof(cl_uint), &n));
        probeLaunch(dv, o, kFill, fillGlobal, fillLocal, 1, -1, 0u);
        for (gi = 0; gi < 2; ++gi) {
            size_t local = groups[gi];
            if (local > di->maxWorkGroup) continue;
            for (li = 0; li < (size_t)nLanes; ++li) {
                size_t lanes = lanesList[li];
                cl_uint seed = (cl_uint)(0x1234567u + (cl_uint)li * 977u);
                double ms;
                if (lanes < local) continue;
                CL_CHECK(clSetKernelArg(kChase, 0, sizeof(cl_mem), &dDs));
                CL_CHECK(clSetKernelArg(kChase, 1, sizeof(cl_uint), &mask));
                CL_CHECK(clSetKernelArg(kChase, 2, sizeof(cl_uint), &STEPS));
                CL_CHECK(clSetKernelArg(kChase, 3, sizeof(cl_uint), &seed));
                CL_CHECK(clSetKernelArg(kChase, 4, sizeof(cl_mem), &dOut));
                ms = probeLaunch(dv, o, kChase, lanes, local, 3, 3, seed);
                printf("| chase | %d | %llu | %llu | %u | %.3f | %.3f | %.0f |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
                       (double)lanes * (double)STEPS / (ms / 1000.0) / 1e9, ms * 1e6 / (double)STEPS);
                fflush(stdout);
            }
        }
        {
            size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
            size_t lanes;
            for (lanes = 1u << 16; lanes <= maxLanes; lanes <<= 2) {
                cl_uint seed = 0x7654321u;
                double ms;
                CL_CHECK(clSetKernelArg(kIndep, 0, sizeof(cl_mem), &dDs));
                CL_CHECK(clSetKernelArg(kIndep, 1, sizeof(cl_uint), &mask));
                CL_CHECK(clSetKernelArg(kIndep, 2, sizeof(cl_uint), &STEPS));
                CL_CHECK(clSetKernelArg(kIndep, 3, sizeof(cl_uint), &seed));
                CL_CHECK(clSetKernelArg(kIndep, 4, sizeof(cl_mem), &dOut));
                ms = probeLaunch(dv, o, kIndep, lanes, local, 3, 3, seed);
                printf("| indep x8 | %d | %llu | %llu | %u | %.3f | %.3f | (8 loads in flight per lane) |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
                       (double)lanes * 8.0 * (double)STEPS / (ms / 1000.0) / 1e9);
                fflush(stdout);
            }
        }
        {
            /* Random 64-byte lines (16 words, four uint4 loads) in a dependent chain: lines per second against the
             * 4-byte chase above says what one random 4-byte read costs the memory system. If the two rates are
             * equal, every 4-byte read fetches a whole line; if lines/s is a quarter of loads/s, reads cost a 16-byte
             * sector. */
            size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
            size_t lanes;
            cl_uint lineMask = (words / 16u) - 1u;
            for (lanes = 1u << 14; lanes <= maxLanes; lanes <<= 2) {
                cl_uint seed = 0x3141592u;
                double ms;
                CL_CHECK(clSetKernelArg(kLine, 0, sizeof(cl_mem), &dDs));
                CL_CHECK(clSetKernelArg(kLine, 1, sizeof(cl_uint), &lineMask));
                CL_CHECK(clSetKernelArg(kLine, 2, sizeof(cl_uint), &STEPS));
                CL_CHECK(clSetKernelArg(kLine, 3, sizeof(cl_uint), &seed));
                CL_CHECK(clSetKernelArg(kLine, 4, sizeof(cl_mem), &dOut));
                ms = probeLaunch(dv, o, kLine, lanes, local, 3, 3, seed);
                printf("| line 64 B | %d | %llu | %llu | %u | %.3f | %.3f G lines/s | %.1f GB/s in lines |\n", mib, (unsigned long long)local, (unsigned long long)lanes, STEPS, ms,
                       (double)lanes * (double)STEPS / (ms / 1000.0) / 1e9, (double)lanes * (double)STEPS * 64.0 / (ms / 1000.0) / 1e9);
                fflush(stdout);
            }
        }
        {
            /* Coalesced read of the whole buffer (uint4 per lane per step, consecutive lanes consecutive addresses):
             * the sequential bandwidth. Against the card's rated figure this says whether the memory clock is in its
             * full state; a card parked in a middle memory state shows about half (approximate). */
            size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
            size_t lanes = 1u << 20;
            cl_uint perLane = (cl_uint)((uint64_t)words / 4ull / (uint64_t)lanes);
            double ms, bytes = (double)perLane * (double)lanes * 16.0;
            if (perLane == 0) { perLane = 1; lanes = (size_t)words / 4u; bytes = (double)lanes * 16.0; }
            CL_CHECK(clSetKernelArg(kStream, 0, sizeof(cl_mem), &dDs));
            CL_CHECK(clSetKernelArg(kStream, 1, sizeof(cl_uint), &perLane));
            CL_CHECK(clSetKernelArg(kStream, 2, sizeof(cl_mem), &dOut));
            ms = probeLaunch(dv, o, kStream, lanes, local, 3, -1, 0u);
            printf("| stream | %d | %llu | %llu | %u | %.3f | %.1f GB/s coalesced | (%.0f MiB read once) |\n", mib, (unsigned long long)local, (unsigned long long)lanes, perLane, ms, bytes / (ms / 1000.0) / 1e9, bytes / 1048576.0);
            fflush(stdout);
        }
        clReleaseMemObject(dDs);
    }
    {
        size_t local = di->maxWorkGroup < 256 ? di->maxWorkGroup : 256;
        size_t lanes = 1u << 20;
        cl_uint seed = 0x2468aceu;
        double ms, ops;
        CL_CHECK(clSetKernelArg(kAlu, 0, sizeof(cl_uint), &ALU_STEPS));
        CL_CHECK(clSetKernelArg(kAlu, 1, sizeof(cl_uint), &seed));
        CL_CHECK(clSetKernelArg(kAlu, 2, sizeof(cl_mem), &dOut));
        ms = probeLaunch(dv, o, kAlu, lanes, local, 3, 1, seed);
        ops = (double)lanes * (double)ALU_STEPS * 5.0;   /* mul, add, rotate, xor, add per step */
        printf("| alu | 0 | %llu | %llu | %u | %.3f | %.1f G int ops/s | %.3f G steps/s per compute unit (approximate: 5 ops per step counted) |\n",
               (unsigned long long)local, (unsigned long long)lanes, ALU_STEPS, ms, ops / (ms / 1000.0) / 1e9,
               (double)lanes * (double)ALU_STEPS / (ms / 1000.0) / 1e9 / (double)(di->computeUnits ? di->computeUnits : 1));
    }
    clReleaseMemObject(dOut);
    clReleaseKernel(kFill); clReleaseKernel(kChase); clReleaseKernel(kIndep); clReleaseKernel(kAlu); clReleaseKernel(kLine); clReleaseKernel(kStream);
    clReleaseProgram(prog);
    printf("memprobe: done\n");
    return 0;
 }
 int main(int argc, char** argv) {
    Options o = parseArgs(argc, argv);
    DeviceInfo* devs = NULL;
@ -1397,7 +1839,7 @@ int main(int argc, char** argv) {
        if (o.device >= nDev) { printf("FAIL: --device %d out of range (%d devices)\n", o.device, nDev); return 2; }
        chosen = o.device;
    } else if (o.vendor) {
-        for (i = 0; i < nDev; ++i) if ((devs[i].type & CL_DEVICE_TYPE_GPU) && strstr(devs[i].vendor, o.vendor)) { chosen = i; break; }
+        for (i = 0; i < nDev; ++i) if ((devs[i].type & CL_DEVICE_TYPE_GPU) && devs[i].dupOf < 0 && strstr(devs[i].vendor, o.vendor)) { chosen = i; break; }
        if (chosen < 0) {
            printf("OpenCL devices (%d):\n", nDev);
            for (i = 0; i < nDev; ++i) printDevice(i, &devs[i], 0);
@ -1405,14 +1847,20 @@ int main(int argc, char** argv) {
            return 2;
        }
    } else {
-        for (i = 0; i < nDev; ++i) if (devs[i].type & CL_DEVICE_TYPE_GPU) { chosen = i; break; }
+        for (i = 0; i < nDev; ++i) if ((devs[i].type & CL_DEVICE_TYPE_GPU) && devs[i].dupOf < 0) { chosen = i; break; }
        if (chosen < 0) chosen = 0;
    }
    printf("OpenCL devices (%d):\n", nDev);
    for (i = 0; i < nDev; ++i) printDevice(i, &devs[i], i == chosen && !o.list);
    {
        int dups = 0;
        for (i = 0; i < nDev; ++i) if (devs[i].dupOf >= 0) ++dups;
        if (dups) printf("platforms: %d device(s) hidden as the same card on an older platform of the same vendor (an old driver's OpenCL registration is still present)\n", dups);
    }
    if (o.list) return 0;
    di = &devs[chosen];
    printf("using device [%d] %s\n", chosen, di->name);
    if (di->dupOf >= 0) printf("NOTE: --device %d is the older platform's listing of the card [%d] (driver %s); results on it are for comparison only\n", chosen, di->dupOf, di->driver);
    if (o.timeWall < 0) o.timeWall = (strcmp(di->platformName, "Apple") == 0) ? 1 : 0;
    if (o.timeWall) printf("timing: host wall time (Apple's OpenCL event timestamps are not usable; the rate is still a device rate, see README.md)\n");
    else printf("timing: device event profiling (CL_PROFILING_COMMAND_START/END), like cudaEvent elapsed time\n");
@ -1423,6 +1871,12 @@ int main(int argc, char** argv) {
    dv.q = clCreateCommandQueue(dv.ctx, di->device, CL_QUEUE_PROFILING_ENABLE, &err);
    CL_CHECK_ERR(err, "clCreateCommandQueue");
    if (o.memprobe) {
        int rc = runMemprobe(&dv, di, &o);
        clReleaseCommandQueue(dv.q);
        clReleaseContext(dv.ctx);
        return rc;
    }
    if (o.serve && !o.kernelGiven) {
        /* The bound kernel lives next to the compiled-in kernel.cl as kernel_bound.cl (packs from igneum-pow or igneum-miner export-pack). */
        static char boundPath[1024];
@ -1440,14 +1894,7 @@ int main(int argc, char** argv) {
    printf("build options: %s\n", dv.buildOptions);
    printf("exchange: %s\n", dv.exchangeNote);
    if (o.serve) return runServe(&dv, di, &o);
-    {
+    printKernelInfo(di, dv.kHash, "igneum_hash", dv.groupSize, "");
        size_t wg = 0;
        cl_ulong lmem = 0;
        clGetKernelWorkGroupInfo(dv.kHash, di->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(wg), &wg, NULL);
        clGetKernelWorkGroupInfo(dv.kHash, di->device, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(lmem), &lmem, NULL);
        printf("kernel: igneum_hash max work-group %llu, local memory %llu bytes, work-group %d x 32\n",
               (unsigned long long)wg, (unsigned long long)lmem, o.groupWarps);
    }
    printf("program: %d instructions x %d iterations, loads/hash %d, op mix %s\n",
           IGNEUM_INSTR_COUNT, IGNEUM_ITERATIONS, IGNEUM_LOADS_PER_HASH, IGNEUM_OP_MIX);
    printf("seed words: %08x %08x %08x %08x %08x %08x %08x %08x\n",
--- a/proto-opencl/test-host.sh
+++ b/proto-opencl/test-host.sh
@ -0,0 +1,10 @@
 #!/usr/bin/env bash
 # Builds and runs test_host.c (the device-free rules of host.c) against the placeholder pack. No GPU is touched.
 set -euo pipefail
 cd "$(dirname "$0")"
 P="../proto-cuda/packs/igneum-devnet-v4-epoch0"
 case "$(uname -s)" in
  Darwin) cc -std=c99 -O1 -Wall -Wextra -Wno-deprecated-declarations -Wno-unused-function -I "$P" -DIGNEUM_KERNEL_PATH='"kernel_bound.cl"' -o test_host test_host.c -framework OpenCL ;;
  *) cc -std=c99 -O1 -Wall -Wextra -Wno-unused-function -I "$P" -DIGNEUM_KERNEL_PATH='"kernel_bound.cl"' -o test_host test_host.c -lOpenCL -ldl ;;
 esac
 ./test_host
--- a/proto-opencl/test_host.c
+++ b/proto-opencl/test_host.c
@ -0,0 +1,71 @@
 /* Unit tests for the host-side rules of host.c that need no device (5 October 2026): the duplicate-platform fold
 * (markDuplicates, driverNumber). host.c is included with its main renamed. Build and run: ./test-host.sh */
 #define main host_main
 #include "host.c"
 #undef main
 static int failures = 0;
 #define CHECK(cond, what) do { if (!(cond)) { printf("FAIL: %s (line %d)\n", what, __LINE__); ++failures; } else printf("ok: %s\n", what); } while (0)
 static DeviceInfo dev(const char* platform, const char* version, const char* name, const char* vendor, const char* driver, cl_ulong mem, cl_uint cus) {
    DeviceInfo d;
    memset(&d, 0, sizeof(d));
    strncpy(d.platformName, platform, 255); strncpy(d.platformVersion, version, 255);
    strncpy(d.name, name, 255); strncpy(d.vendor, vendor, 255); strncpy(d.driver, driver, 255);
    d.globalMem = mem; d.computeUnits = cus; d.type = CL_DEVICE_TYPE_GPU; d.dupOf = -1;
    return d;
 }
 int main(void) {
    const char* AMD = "AMD Accelerated Parallel Processing";
    const char* V_NEW = "OpenCL 2.1 AMD-APP (3683.0)", *V_OLD = "OpenCL 2.1 AMD-APP (3652.0)", *V_OLDER = "OpenCL 2.1 AMD-APP (3600.0)";
    const char* AMDV = "Advanced Micro Devices, Inc.";
    CHECK(driverNumber("3683.0 (PAL,LC)") == 3683.0, "driverNumber reads the AMD string");
    CHECK(driverNumber("617.14") > 617.1 && driverNumber("617.14") < 617.2, "driverNumber reads the NVIDIA string");
    CHECK(driverNumber("") == 0.0, "driverNumber of an empty string is 0");
    {
        /* PC 1 on 5 October 2026: the new platform first, the old one second, NVIDIA last */
        DeviceInfo l[5];
        l[0] = dev(AMD, V_NEW, "gfx1036", AMDV, "3683.0 (PAL,LC)", 59589ull << 20, 1);
        l[1] = dev(AMD, V_NEW, "gfx1201", AMDV, "3683.0 (PAL,LC)", 16304ull << 20, 32);
        l[2] = dev(AMD, V_OLD, "gfx1036", AMDV, "3652.0 (PAL,LC)", 59589ull << 20, 1);
        l[3] = dev(AMD, V_OLD, "gfx1201", AMDV, "3652.0 (PAL,LC)", 16304ull << 20, 32);
        l[4] = dev("NVIDIA CUDA", "OpenCL 3.0 CUDA 13.4.96", "NVIDIA GeForce RTX 5090", "NVIDIA Corporation", "617.14", 32579ull << 20, 170);
        CHECK(markDuplicates(l, 5) == 2, "PC 1: two duplicates marked");
        CHECK(l[0].dupOf == -1 && l[1].dupOf == -1 && l[4].dupOf == -1, "PC 1: the new platform's cards and the NVIDIA card stay");
        CHECK(l[2].dupOf == 0 && l[3].dupOf == 1, "PC 1: the old platform's entries point at the new ones");
    }
    {
        /* the old platform enumerated first: the newer entry must still be the one kept */
        DeviceInfo l[2];
        l[0] = dev(AMD, V_OLD, "gfx1201", AMDV, "3652.0 (PAL,LC)", 16304ull << 20, 32);
        l[1] = dev(AMD, V_NEW, "gfx1201", AMDV, "3683.0 (PAL,LC)", 16304ull << 20, 32);
        CHECK(markDuplicates(l, 2) == 1, "old first: one duplicate");
        CHECK(l[0].dupOf == 1 && l[1].dupOf == -1, "old first: the older entry is the duplicate");
    }
    {
        /* two real cards of one model on ONE platform: never folded */
        DeviceInfo l[2];
        l[0] = dev(AMD, V_NEW, "gfx1201", AMDV, "3683.0 (PAL,LC)", 16304ull << 20, 32);
        l[1] = dev(AMD, V_NEW, "gfx1201", AMDV, "3683.0 (PAL,LC)", 16304ull << 20, 32);
        CHECK(markDuplicates(l, 2) == 0 && l[0].dupOf == -1 && l[1].dupOf == -1, "two identical cards on one platform stay two cards");
    }
    {
        /* three registrations of one card: one stays */
        DeviceInfo l[3];
        l[0] = dev(AMD, V_OLD, "gfx1201", AMDV, "3652.0 (PAL,LC)", 16304ull << 20, 32);
        l[1] = dev(AMD, V_OLDER, "gfx1201", AMDV, "3600.0 (PAL,LC)", 16304ull << 20, 32);
        l[2] = dev(AMD, V_NEW, "gfx1201", AMDV, "3683.0 (PAL,LC)", 16304ull << 20, 32);
        CHECK(markDuplicates(l, 3) == 2, "three platforms: two duplicates");
        CHECK(l[2].dupOf == -1 && l[0].dupOf == 2 && l[1].dupOf == 2, "three platforms: the newest stays and both duplicates point at it");
    }
    {
        /* a different card with the same name but other memory (an 8 GB and a 16 GB model) on two platforms: not folded */
        DeviceInfo l[2];
        l[0] = dev(AMD, V_NEW, "gfx1201", AMDV, "3683.0 (PAL,LC)", 16304ull << 20, 32);
        l[1] = dev(AMD, V_OLD, "gfx1201", AMDV, "3652.0 (PAL,LC)", 8000ull << 20, 32);
        CHECK(markDuplicates(l, 2) == 0, "different memory sizes are different cards");
    }
    printf("%s: %d failure(s)\n", failures ? "FAIL" : "PASS", failures);
    return failures ? 1 : 0;
 }