From 5d5ba15ed738ccf71bb11230a15e853f089e2afa Mon Sep 17 00:00:00 2001 From: igneum-josh <337424239+igneum-josh@users.noreply.github.com> Date: Mon, 5 Oct 2026 20:09:10 +0000 Subject: [PATCH] Counter ASIC 2.0 layer 6: SRAM mirror analysis (cited bit cells N7 to N2 and 18A, area and cost per node, no cache growth rule in the spec, options A to E for Josh); layer 7 dp4a probes for Metal, CUDA and OpenCL (standalone, no lottery kernel) Co-Authored-By: Claude Fable 5.1 --- docs/analysis/sram-mirror.md | 227 +++++++++++++++++++++++++++++++++++ proto-cuda/dot4-probe.cu | 115 ++++++++++++++++++ proto-metal/dot4-probe.swift | 165 +++++++++++++++++++++++++ proto-opencl/dot4-probe.c | 187 +++++++++++++++++++++++++++++ 4 files changed, 694 insertions(+) create mode 100644 docs/analysis/sram-mirror.md create mode 100644 proto-cuda/dot4-probe.cu create mode 100644 proto-metal/dot4-probe.swift create mode 100644 proto-opencl/dot4-probe.c diff --git a/docs/analysis/sram-mirror.md b/docs/analysis/sram-mirror.md new file mode 100644 index 000000000..970239776 --- /dev/null +++ b/docs/analysis/sram-mirror.md @@ -0,0 +1,227 @@ +# Layer 6: the SRAM mirror of the cache against published SRAM density, year 0 to 10 + +5 October 2026 (night), Counter ASIC 2.0 (`docs/plans/counter-asic-2.md`, layer 6), branch `ca2-analysis`. Every figure +below is either cited (paper, vendor document, URL, date) or labelled approximate. Nothing here is a measurement of a +chip. Numbers in this file were computed with the arithmetic shown; the script is in section 9. + +## 1. The question + +The lottery hash derives every dataset item from a 256 MiB cache (spec 01 sections 1.5 and 1.8). A chip that holds the +cache in on-die SRAM can recompute items instead of reading the dataset (ledger M16, the recompute attacker). Layer 6 +asks whether the cache size, as the specification schedules it, keeps that SRAM mirror unaffordable for ten years of +the genesis schedule, and if not what growth rule would. + +Two things also sit in a chip's SRAM budget if it mirrors the full read-only working set: the layer 5 hot table (32, +64 or 96 MB, a class parameter on `readwidth` b970dda, coordinator's note of 5 October) beside the 256 MiB cache. The +per-warp scratch of layer 3 (32 or 128 KB per warp, written, not read-only) is not mirrorable and is left out of the +mirror; it is counted in the 6 GB working-set budget in section 7. + +## 2. What the specification schedules for the cache + +| Quantity | Rule | Where | +|---|---|---| +| Dataset | 2 GiB at genesis plus 0.5 GiB per year (`N_d` grows about 23 KiB per day) | spec 01 section 1.13.3, Designed | +| Cache | 256 MiB, "prototype value, to be fixed at gate 1"; the rule that fixes it: "the cache must exceed the largest on-chip cache of any card that mines, and 96 MiB of L2 on the 5090 is the figure to beat" | spec 01 sections 1.5 and 1.16 | +| Cache growth | None. No section of `docs/spec/` grows the cache (grep of `docs/spec` for cache growth, schedule, doubling: only the dataset rule of 1.13.3 and the README's "growth" word, which refers to it) | this analysis, 5 October 2026 | + +So the plan's layer 6 row ("already in the design; confirm the schedule") is half right: dataset growth is in the +design, cache growth is not. The cache is flat at 256 MiB for every year of the schedule as the spec stands. M16's +closing line names the rule the cache should get ("exceeds what one die can hold, and grows") as a gate 1 decision +that has not been taken. + +## 3. SRAM bit cell per node, cited + +| Node (vendor) | HD 6T bit cell, um^2 | Raw density, Mbit/mm^2 (1/cell) | Year of volume (approximate) | Source | +|---|---|---|---|---| +| N7 (TSMC) | 0.027 | 37.0 | 2018 | WikiChip, "TSMC Details 5 nm" (ISSCC/IEDM disclosures), https://fuse.wikichip.org/news/3398/tsmc-details-5-nm/ | +| N5 (TSMC) | 0.021 | 47.6 | 2020 | same (two N5 cells: HD 0.021, HP 0.025) | +| N3B (TSMC) | 0.0199 | 50.3 | 2022 to 2023 | WikiChip, "IEDM 2022: Did We Just Witness The Death Of SRAM?", https://fuse.wikichip.org/news/7343/iedm-2022-did-we-just-witness-the-death-of-sram/ (TSMC's IEDM 2022 N3 paper) | +| N3E (TSMC) | 0.021 | 47.6 | 2023 | same; Tom's Hardware, "TSMC's 3nm Node: No SRAM Scaling", https://www.tomshardware.com/news/no-sram-scaling-implies-on-more-expensive-cpus-and-gpus | +| N2 (TSMC) | 0.0175 | 57.1 | 2025 to 2026 | TSMC at IEDM 2024, reported by Tom's Hardware, https://www.tomshardware.com/tech-industry/tsmc-shares-deep-dive-details-about-its-cutting-edge-2nm-process-node-at-iedm-2024-35-percent-less-power-or-15-percent-more-performance ; ISSCC 2025 paper "A 38.1Mb/mm2 SRAM in a 2nm-CMOS-Nanosheet Technology", https://research.tsmc.com/page/memory/4.html | +| Intel 18A | 0.021 | 47.6 | 2025 to 2026 | ISSCC 2025 paper 29.2, "A 0.021 um^2 High-Density SRAM in Intel 18A RibbonFET Technology with PowerVia", https://www.researchgate.net/publication/389644177 ; IEEE Spectrum 26 Feb 2025, https://spectrum.ieee.org/sram-intel-tsmc | +| Samsung SF3 / SF2 | not disclosed as a bit cell area in anything found tonight (Samsung's ISSCC papers give assist circuits and macro figures, not the HD cell) | | | search of ISSCC 2021 to 2025 coverage, 5 October 2026; left out of the tables | + +The stall. N3B's cell is 5% smaller than N5's and N3E's is the same size as N5's (0.021 um^2 both): zero SRAM +scaling from N5 to N3E (WikiChip IEDM 2022 article above; Tom's Hardware above; SemiAnalysis "TSMC's 3nm Conundrum", +https://newsletter.semianalysis.com/p/tsmcs-3nm-conundrum-does-it-even). N2's nanosheet cell recovers 17% (0.021 to +0.0175 um^2). So across 2020 to 2026 the HD bit cell shrank once, by 17%. + +Array efficiency (bit cell to macro). The usable density of a macro is below 1/cell because of word-line and +bit-line drivers, sense amplifiers, decoders and redundancy. The factor used here is 0.70, WikiChip's convention +(their 31.8 Mib/mm^2 for the 0.021 um^2 N3E cell is 1/0.021 x 0.70 in Mib). The two ISSCC 2025 macros bracket it: +TSMC N2 38.1 Mb/mm^2 at a 0.0175 um^2 cell is 67%; Intel 18A 38.1 Mb/mm^2 array density and 34.3 Mb/mm^2 for the +volume macro at a 0.021 um^2 cell are 80% and 72% (ISSCC 2025 29.2, above). Both lie within 10% of 0.70. + +GPU on-die SRAM for scale: the RTX 5090 carries 96 MB of L2 (98,304 KB) on a 750 mm^2 TSMC 4N die with 92.2 billion +transistors; the full GB202 has 128 MB; the RTX 4090 had 72 MB and the RTX 3090 6 MB (NVIDIA, "RTX Blackwell GPU +Architecture" whitepaper v1.1, appendix table "L2 Cache Size", https://images.nvidia.com/aem-dam/Solutions/geforce/blackwell/nvidia-rtx-blackwell-gpu-architecture.pdf). +At the N5-class cell and 0.70 that L2 is about 24 mm^2 of the 750 (3%), approximate. The RX 9070 XT carries 64 MB +of Infinity Cache plus 8 MB of L2 (vendor figures, approximate, bench-log "the 9070 XT on the eGPU"). + +Reticle: the EUV field is 26 x 33 mm = 858 mm^2, about 830 mm^2 usable after scribe lanes (SemiAnalysis, "Die Size +And Reticle Conundrum", https://newsletter.semianalysis.com/p/die-size-and-reticle-conundrum-cost ; WikiChip "Mask", +https://en.wikichip.org/wiki/mask). The 5090's 750 mm^2 is 90% of it. + +Wafer prices (approximate; TSMC publishes none, every figure is supply-chain reporting): N7 about $9,500, N5 and N3 +about $20,000 (Silicon Analysts, "Wafer Pricing by Node", September 2026, https://siliconanalysts.com/data/wafer-pricing); +N2 about $30,000 (Tom's Hardware, https://www.tomshardware.com/tech-industry/semiconductors/tsmc-could-charge-up-to-usd45-000-for-1-6nm-wafers-rumors-allege-a-50-percent-increase-in-pricing-over-prior-gen-wafers). + +## 4. Die area to mirror the cache, per node + +Area = bits / (raw density x 0.70). The columns are the 256 MiB cache alone, the cache plus each hot-table size of +layer 5 (32, 64, 96 MB taken as MiB), and the larger caches of the options in section 6. + +| Node | Macro Mbit/mm^2 at 0.70 | 256 MiB | 256 + 32 | 256 + 64 | 256 + 96 | 512 MiB | 1 GiB | 2 GiB | 4 GiB | +|---|---|---|---|---|---|---|---|---|---| +| N7 | 25.9 | 83 mm^2 | 93 | 104 | 114 | 166 | 331 | 663 | 1,325 (2 dies) | +| N5 | 33.3 | 64 | 72 | 81 | 89 | 129 | 258 | 515 | 1,031 (2 dies) | +| N3B | 35.2 | 61 | 69 | 76 | 84 | 122 | 244 | 488 | 977 (2 dies) | +| N3E, Intel 18A | 33.3 | 64 | 72 | 81 | 89 | 129 | 258 | 515 | 1,031 (2 dies) | +| N2 | 40.0 | 54 | 60 | 67 | 74 | 107 | 215 | 429 | 859 (2 dies) | + +One reticle (830 mm^2) holds 2.5 GiB of SRAM at N7, 3.2 GiB at N5, N3E and 18A, 3.9 GiB at N2 (same arithmetic). + +Against the figures the ledger carries: M16's "100 to 300 mm^2" (low end from a 0.02 um^2 cell with overhead, high +end from wafer-scale parts at about 1 MB/mm^2) and the plan's "about 45 mm^2 at a leading node" both bracket the +cited 54 to 64 mm^2; the wafer-scale high end is a different efficiency (Cerebras-class arrays sit beside logic) and +is not the right number for a pure SRAM die. The right figure for the ledger is 54 to 83 mm^2 depending on node, +cited above. + +## 5. Cost per good die + +Dies per 300 mm wafer by the usual approximation pi x 150^2 / A minus the edge term pi x 300 / sqrt(2A); yield by +Poisson exp(-A x D0) with D0 = 0.1 defects per cm^2 (an assumption, approximate; SRAM arrays carry redundancy so +real yield is higher, which lowers these costs). Cost per good die = wafer price / (dies x yield). Packaging, test, +the logic beside the SRAM and the design (masks at N5 and below run into the tens of millions of dollars, +approximate) are not in these numbers; they are per-die silicon only. + +| Node, wafer price | 256 MiB | 256 + 96 MiB | 1 GiB | 4 GiB (2 dies) | +|---|---|---|---|---| +| N7, $9,500 | 83 mm^2, 780 dies, yield 0.92, $13 | $19 | 331 mm^2, 177 dies, 0.72, $75 | $456 | +| N5, $20,000 | 64 mm^2, 1,014 dies, 0.94, $21 | $30 | 258 mm^2, 233 dies, 0.77, $111 | $621 | +| N3E, $20,000 | $21 | $30 | $111 | $621 | +| N2, $30,000 | 54 mm^2, 1,226 dies, 0.95, $26 | $37 | 215 mm^2, 284 dies, 0.81, $131 | $696 | + +Reading. The silicon for a 256 MiB mirror is tens of dollars per die on any node from N7 up. With the hot table it +is still under $40. It was never the SRAM that priced the recompute attacker out; the plan's premise for layer 6 +("the SRAM mirror stays unaffordable") does not hold for the cache as a mirror and did not hold at genesis either. + +## 6. What the mirror buys the attacker, year by year + +From M16 (`docs/analysis/m16-recompute-attacker-2026-10-05.md`): with the cache on die the attacker recomputes 128 +items per hash at about 1,170 integer operations and 8 dependent 64-byte cache reads each, about 150,000 operations +and 1,024 dependent SRAM reads per hash. At a 5090-class integer budget (about 50 T op/s, approximate) that is +0.33 Ghash/s against the honest 141 Mhash/s projected for version 2 programs: 2.4x at equal silicon before any +fixed-function factor, 3x to 6x with one (approximate). The SRAM is 54 to 83 mm^2 of that chip (7 to 11% of a +750 mm^2 die), so the mirror is cheap and the recompute route is bound by integer throughput, not by SRAM. + +The layer 5 hot table changes nothing in that arithmetic: the hot table is read-only and derived from the day key +like the cache, so a chip mirrors it in the same SRAM (another 32 to 96 MB, 7 to 24 mm^2 at N5) and reads it at +SRAM latency, which is exactly what a GPU's L2 does with it. Layer 5 taxes the DRAM-only chip (the one without SRAM); +it does not tax the SRAM chip. + +Dataset growth does not touch the recompute attacker: the attacker never holds the dataset. It taxes the +partial-store attacker (O-1.6, the time-memory curve, not drawn) and the honest card. + +Year by year under the schedule as it stands (flat 256 MiB), the mirror's area at the best node available that +year. Node years are approximate; the density trend from 2018 to 2025 is 37.0 to 57.1 Mbit/mm^2 raw, 1.54x in +7 years, about 6% per year, and it came in one step (N2); the extrapolation past 2026 assumes that average holds +(approximate, and optimistic for the attacker: A16 and A14 have no disclosed SRAM cell yet). + +| Year | Calendar (approximate) | Dataset, GiB | Cache (spec) | Best node, raw Mbit/mm^2 | Mirror of the cache, mm^2 | With a 96 MiB hot table, mm^2 | Mirror as a share of a 750 mm^2 die | +|---|---|---|---|---|---|---|---| +| 0 | 2027 | 2.0 | 256 MiB | N2, 57.1 (cited) | 54 | 74 | 7% | +| 1 | 2028 | 2.5 | 256 MiB | N2 or A16, 57 to 61 | 50 to 54 | 69 to 74 | 7% | +| 2 | 2029 | 3.0 | 256 MiB | about 64 (trend) | 48 | 66 | 6% | +| 3 | 2030 | 3.5 | 256 MiB | about 68 | 45 | 62 | 6% | +| 4 | 2031 | 4.0 | 256 MiB | about 72 | 43 | 59 | 6% | +| 5 | 2032 | 4.5 | 256 MiB | about 76 | 40 | 55 | 5% | +| 6 | 2033 | 5.0 | 256 MiB | about 81 | 38 | 52 | 5% | +| 7 | 2034 | 5.5 | 256 MiB | about 86 | 36 | 49 | 5% | +| 8 | 2035 | 6.0 | 256 MiB | about 91 | 34 | 46 | 5% | +| 9 | 2036 | 6.5 | 256 MiB | about 97 | 32 | 44 | 4% | +| 10 | 2037 | 7.0 | 256 MiB | about 102 | 30 | 41 | 4% | + +Reading. A flat cache's mirror shrinks from 7% to 4% of a large die over the decade, and a 5090-class consumer GPU +already carries 96 MB of L2 on one die with the full GB202 at 128 MB; at the 2020 to 2025 pace of GPU L2 growth +(6 MB, 72 MB, 96 MB on the three NVIDIA flagships in the whitepaper table) a consumer GPU could hold 256 MiB on die +within the decade. The spec's own rule for the cache ("must exceed the largest on-chip cache of any card that +mines") would then be broken by a flat cache. That is the real reason to grow it: not to price a chip out (section +5 shows the SRAM cannot do that) but to keep the cache out of every GPU's own cache, so the honest hash stays +DRAM-latency-bound and the recompute route stays a route only a custom chip can take. + +## 7. Answer to the layer 6 question, and the options + +Does the flat 256 MiB cache keep the SRAM mirror unaffordable through year 10? No. It is affordable at year 0 +(tens of dollars of silicon per die, section 5) and gets cheaper. What keeps the recompute attacker near 1x is +M16's integer arithmetic and the mixer-cost lever (4x the mixer cost puts the equal-silicon gain at 0.36x, bounded +by the CPU verify gate), not the cache size. The cache size does one other job, keeping the cache larger than any +GPU's L2, and that job needs growth. + +Options for the cache rule, with the honest costs each implies. Verifier fill time is 0.2 s per 256 MiB on one core +(spec 1.12: "a 0.2 s CPU cache fill", from the measured 175 to 190 ms of section 1.8.3), scaled linearly; the +verifier holds the whole cache (section 1.11), so its memory is the cache size plus the program and the interpreter. +GPU fill: 0.67 ms per 256 MiB on the 5090 (section 1.8.3), linear. The GPU dataset build (13.4 ms per 1 GiB on the +5090, section 1.8.3) depends on the dataset size, not the cache size; a larger cache spreads the build's 8 dependent +reads per item over more memory, which on a GPU means more of them miss L2 and the build slows by some factor +between 1x and the L2-to-DRAM latency ratio, which is a measurement to take (approximate; owed). Mirror area is at N2 +(cited density), the node of the first years; at the trend's year-10 density divide by about 1.8. + +| Option | Rule | Cache at year 0 / 4 / 10 | Mirror at N2, year 0 / 4 / 10 (mm^2) | Dies at year 10 (830 mm^2 reticle) | Verifier fill, one core, year 0 / 10 | Verifier memory, year 10 | GPU cache fill (5090), year 10 | Keeps the cache above a 96 MB L2 at year 10 | Keeps it above a 256 MB L2 | +|---|---|---|---|---|---|---|---|---|---| +| A, as specified | flat 256 MiB | 256 / 256 / 256 MiB | 54 / 54 / 54 | 1 | 0.2 / 0.2 s | 256 MiB | 0.7 ms | yes, 2.7x | no | +| B | cache = dataset / 8 (today's ratio) | 256 / 512 / 896 MiB | 54 / 107 / 188 | 1 | 0.2 / 0.7 s | 896 MiB | 2.3 ms | yes, 9.3x | yes, 3.5x | +| C | cache doubles when the dataset doubles (the dataset's own clock: year 4, then year 12) | 256 / 512 / 512 MiB | 54 / 107 / 107 | 1 | 0.2 / 0.4 s | 512 MiB | 1.3 ms | yes, 5.3x | yes, 2x | +| D | cache = dataset / 4 | 512 / 1,024 / 1,792 MiB | 107 / 215 / 376 | 1 | 0.4 / 1.4 s | 1.75 GiB | 4.7 ms | yes | yes, 7x | +| E, one reticle | cache sized so the mirror exceeds one reticle at the node of the day: 4 GiB at N2 (section 4), growing with density | 4 GiB / about 4.5 / about 7 GiB | 859 / 860 / 860 (by construction) | 2 | 3.2 / 5.6 s | 7 GiB | 11 / 19 ms | yes | yes | + +Where the working set enters (coordinator's budget: 1 GiB table + hot table + scratch for every resident warp + +buffers under 6 GB on an 8 GB card): the cache is not in the miner's working set at hash time (the dataset is built +from it once a day and the cache can be dropped or kept), so options A to D do not move that budget; the dataset's own +growth does (2 GiB at genesis, 4 GiB at year 4, 7 GiB at year 10, which is past an 8 GB card at about year 8 on its +own). Option E's 4 GiB cache would have to be built on the card and dropped, which is fine for a 16 GB card and tight +on an 8 GB one at build time (4 GiB cache + 2 GiB dataset + hot table). The per-warp scratch at 170 SMs x 64 warps +(approximate, readwidth) is 340 MB at 32 KB and 1.36 GB at 128 KB per warp; with the 1 GiB table, a 96 MB hot table +and buffers that is 1.5 to 2.5 GB at the prototype dataset size, 2.5 to 3.5 GB at the 2 GiB genesis size, inside +6 GB either way. + +Recommendation. Option C (the cache doubles when the dataset doubles) is the one that keeps the spec's own rule true +with the smallest verifier cost: it ties the cache to a clock the spec already has, keeps `AND MASK` (a power of two +every step, which is the 1.13.3 option (b) argument again), costs the verifier 0.4 s and 512 MiB at year 4 and nothing +more until year 12, and keeps the cache 2x above a 256 MB GPU L2 if one appears. It does not price a chip out; nothing +about cache size does (section 5). The lever that does is the mixer cost multiplier of M16, which is the gate 1 +decision to take beside this one. Option B is the same idea in a smooth form and costs the verifier 0.7 s at year 10. +Option E is the only one that makes the mirror a multi-die part and it costs every verifier 3.2 s and 4 GiB at +genesis, which fails the spirit of the 10 ms verify gate (the fill is once a day, but a light node joining pays it on +every day it syncs across). + +Decision for Josh, at gate 1: A, B, C, D or E above, together with M16's mixer multiplier. Nothing here changes a +vector today: the cache size is a prototype value of spec 1.16 and the growth rule would be a new sentence in 1.13.3. + +## 8. What is cited, what is approximate, what is owed + +| Item | Status | +|---|---| +| Bit cells for N7, N5, N3B, N3E, N2, Intel 18A | cited (section 3) | +| Samsung SF2 or SF3 bit cell | not found; left out | +| Array efficiency 0.70 | WikiChip's convention, bracketed by two ISSCC 2025 macros (67 to 80%) | +| Wafer prices | approximate, supply-chain reporting, cited | +| D0 = 0.1 per cm^2, Poisson yield | assumption, stated | +| Node years and the 6% per year density trend past 2026 | approximate, extrapolated from cited 2018 to 2025 points | +| GPU L2 sizes | cited (NVIDIA whitepaper); AMD Infinity Cache approximate | +| Recompute attacker arithmetic | M16, which is itself arithmetic on measured rates, not a chip measurement | +| Dataset-build slowdown at a larger cache on a GPU | owed, a measurement (5090 at a 512 MiB and 1 GiB cache) | +| The on-die emulation of M16 (inline kernel with a 64 MiB cache inside the 5090's L2) | still a PC job (M16) | + +## 9. The arithmetic + +``` +MiB = 2^20; bits = cache_MiB * MiB * 8 +raw_Mbit_per_mm2 = 1 / cell_um2 (1e6 cells per mm2 per um2 of cell) +area_mm2 = bits / (raw * 0.70 * 1e6) +dies_per_wafer = pi * 150^2 / area - pi * 300 / sqrt(2 * area) +yield = exp(-area_mm2 * 0.001) (D0 = 0.1 per cm2) +cost_per_good_die = wafer_price / (dies * yield) +reticle_GiB = 830 * raw * 0.70 * 1e6 / 8 / 2^30 +``` +Run on 5 October 2026 with Python 3 on the M5 Max; the printed tables are the ones above, rounded. diff --git a/proto-cuda/dot4-probe.cu b/proto-cuda/dot4-probe.cu new file mode 100644 index 000000000..05b66fb68 --- /dev/null +++ b/proto-cuda/dot4-probe.cu @@ -0,0 +1,115 @@ +// dot4-probe (CUDA): dp4a-class throughput on NVIDIA, standalone (no pack, no lottery kernel). +// Counter ASIC 2.0 layer 7 (docs/analysis/int8-matrix-family.md), 5 October 2026. PC job; not run on the Mac. +// +// Three dependent chains, same shape as the OpenCL --memprobe ALU chain (proto-opencl/host.c, probe_alu) and the +// Metal probe (proto-metal/dot4-probe.swift): 1,048,576 lanes x 4,096 steps, best of 3, event time. +// alu x = x * K + rotl(y, 7); y = (y ^ x) + s the card's integer baseline, 5 ops per step counted +// dot4i acc = __dp4a(x, y, acc) (PTX dp4a.s32.s32, sm_61+) one hardware dot4 per step per lane +// dot4e the scalar emulation of the same (4 sign-extended byte products summed, wrapping int32) +// The emulation and the intrinsic must agree bit for bit with the CPU reference (checked on two lanes per run). +// +// Build (Windows, CUDA Toolkit): nvcc -O2 -arch=sm_120 -o dot4-probe-cuda.exe dot4-probe.cu +// Build (Linux): nvcc -O2 -arch=sm_120 -o dot4-probe-cuda dot4-probe.cu +// Run: dot4-probe-cuda [--lanes N] [--steps N] [--reps N] [--device N] +#include +#include +#include +#include +#include + +__host__ __device__ inline uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; } +__host__ __device__ inline uint32_t rotl32(uint32_t v, uint32_t n) { return (v << n) | (v >> (32u - n)); } +__host__ __device__ inline int32_t dot4_emul(uint32_t a, uint32_t b, int32_t acc) { + int32_t r = acc; + for (int i = 0; i < 4; ++i) { + int32_t ba = (int32_t)(int8_t)((a >> (8 * i)) & 0xffu); + int32_t bb = (int32_t)(int8_t)((b >> (8 * i)) & 0xffu); + r = (int32_t)((uint32_t)r + (uint32_t)(ba * bb)); + } + return r; +} + +__global__ void probe_alu(uint32_t steps, uint32_t seed, uint32_t* out) { + uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + for (uint32_t s = 0; s < steps; ++s) { x = x * 0x9E3779B1u + rotl32(y, 7u); y = (y ^ x) + s; } + out[g] = x ^ y; +} +__global__ void probe_dot4i(uint32_t steps, uint32_t seed, uint32_t* out) { + uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + int32_t acc = (int32_t)pm_mix(x); + for (uint32_t s = 0; s < steps; ++s) { + acc = __dp4a((int)x, (int)y, acc); + x = x * 0x9E3779B1u + (uint32_t)acc; + y = rotl32(y, 7u) ^ ((uint32_t)acc + s); + } + out[g] = (uint32_t)acc ^ x ^ y; +} +__global__ void probe_dot4e(uint32_t steps, uint32_t seed, uint32_t* out) { + uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + int32_t acc = (int32_t)pm_mix(x); + for (uint32_t s = 0; s < steps; ++s) { + acc = dot4_emul(x, y, acc); + x = x * 0x9E3779B1u + (uint32_t)acc; + y = rotl32(y, 7u) ^ ((uint32_t)acc + s); + } + out[g] = (uint32_t)acc ^ x ^ y; +} + +static uint32_t lane_ref(const char* name, uint32_t g, uint32_t seed, uint32_t steps) { + uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + if (strcmp(name, "alu") == 0) { + for (uint32_t s = 0; s < steps; ++s) { x = x * 0x9E3779B1u + rotl32(y, 7u); y = (y ^ x) + s; } + return x ^ y; + } + int32_t acc = (int32_t)pm_mix(x); + for (uint32_t s = 0; s < steps; ++s) { + acc = dot4_emul(x, y, acc); + x = x * 0x9E3779B1u + (uint32_t)acc; + y = rotl32(y, 7u) ^ ((uint32_t)acc + s); + } + return (uint32_t)acc ^ x ^ y; +} + +#define CK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { printf("CUDA error %s at %s:%d\n", cudaGetErrorString(e), __FILE__, __LINE__); return 1; } } while (0) + +int main(int argc, char** argv) { + uint32_t lanes = 1u << 20, steps = 4096u; int reps = 3, device = 0; + for (int i = 1; i < argc; ++i) { + if (!strcmp(argv[i], "--lanes") && i + 1 < argc) lanes = (uint32_t)strtoul(argv[++i], 0, 10); + else if (!strcmp(argv[i], "--steps") && i + 1 < argc) steps = (uint32_t)strtoul(argv[++i], 0, 10); + else if (!strcmp(argv[i], "--reps") && i + 1 < argc) reps = atoi(argv[++i]); + else if (!strcmp(argv[i], "--device") && i + 1 < argc) device = atoi(argv[++i]); + else { printf("unknown argument %s\n", argv[i]); return 2; } + } + CK(cudaSetDevice(device)); + cudaDeviceProp p; CK(cudaGetDeviceProperties(&p, device)); + printf("dot4-probe (CUDA) on %s, sm_%d%d, %d SMs, %d MHz, lanes %u, steps %u, best of %d, event time\n", p.name, p.major, p.minor, p.multiProcessorCount, p.clockRate / 1000, lanes, steps, reps); + uint32_t* d_out; CK(cudaMalloc(&d_out, (size_t)lanes * 4)); + uint32_t* h_out = (uint32_t*)malloc((size_t)lanes * 4); + cudaEvent_t e0, e1; CK(cudaEventCreate(&e0)); CK(cudaEventCreate(&e1)); + printf("| kernel | lanes | steps | best ms | G steps/s (= G dot4/s for dot4 rows) | ns per dependent step | lanes 0 and last ok |\n|---|---|---|---|---|---|---|\n"); + const char* names[3] = { "alu", "dot4i", "dot4e" }; + for (int k = 0; k < 3; ++k) { + float best = 1e30f; int ok = 1; + for (int r = 0; r < reps; ++r) { + uint32_t seed = 0x2468aceu + (uint32_t)r * 0x9E3779B9u; + CK(cudaEventRecord(e0)); + if (k == 0) probe_alu<<>>(steps, seed, d_out); + else if (k == 1) probe_dot4i<<>>(steps, seed, d_out); + else probe_dot4e<<>>(steps, seed, d_out); + CK(cudaEventRecord(e1)); CK(cudaEventSynchronize(e1)); CK(cudaGetLastError()); + float ms = 0; CK(cudaEventElapsedTime(&ms, e0, e1)); if (ms < best) best = ms; + CK(cudaMemcpy(h_out, d_out, (size_t)lanes * 4, cudaMemcpyDeviceToHost)); + uint32_t gs[2] = { 0u, lanes - 1u }; + for (int j = 0; j < 2; ++j) { uint32_t want = lane_ref(names[k], gs[j], seed, steps); if (h_out[gs[j]] != want) { ok = 0; printf("MISMATCH %s lane %u: gpu %08x cpu %08x\n", names[k], gs[j], h_out[gs[j]], want); } } + } + double sps = (double)lanes * (double)steps / (best / 1000.0); + printf("| %s | %u | %u | %.3f | %.2f | %.3f | %s |\n", names[k], lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ok ? "yes" : "NO"); + printf("RESULT DOT4 vendor=nvidia device=\"%s\" kernel=%s lanes=%u steps=%u best_ms=%.3f gsteps_per_s=%.2f ok=%d\n", p.name, names[k], lanes, steps, best, sps / 1e9, ok); + } + printf("dot4-probe: done\n"); + return 0; +} diff --git a/proto-metal/dot4-probe.swift b/proto-metal/dot4-probe.swift new file mode 100644 index 000000000..2fa1b4024 --- /dev/null +++ b/proto-metal/dot4-probe.swift @@ -0,0 +1,165 @@ +// dot4-probe: dp4a-class throughput on Apple silicon, standalone (no pack, no lottery kernel). +// Counter ASIC 2.0 layer 7 (docs/analysis/int8-matrix-family.md), 5 October 2026. +// +// Metal has no dp4a intrinsic and no integer simdgroup_matrix (MSL 4.1 section 2.4 lists half, bfloat and float only; +// the Metal 4 tensor op matmul2d does carry char x char -> int, MSL 4.1 table 7.3, measured separately when it is). +// So the per-lane dot4 here is the scalar emulation a conforming Apple miner would run: four sign-extended bytes of +// each operand multiplied and summed into a wrapping int32 accumulator, exactly the PTX dp4a semantics +// (PTX ISA 9.4 section 9.7.1.24: d = c; d += Va[i] * Vb[i] for i in 0..3, bytes sign- or zero-extended). +// +// Two kernels, same shape as the OpenCL --memprobe ALU chain (proto-opencl/host.c, probe_alu: 1,048,576 lanes x 4,096 +// steps, best of 3): +// alu x = x * K + rotate(y, 7); y = (y ^ x) + s the card's integer baseline, 5 ops per step counted +// dot4 acc = dot4(x, y, acc); x = x * K + acc; y = rotate(y, 7) ^ acc one dependent dot4 per step per lane +// Rates: G steps/s per lane-step, so G dot4/s for the second kernel. Timing is the command buffer's GPU start to end. +// +// Build: swiftc -O -o dot4-probe dot4-probe.swift -framework Metal +// Run: ./dot4-probe [--lanes N] [--steps N] [--reps N] [--signed|--unsigned] +import Foundation +import Metal + +let source = """ +#include +using namespace metal; + +inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; } + +// dp4a, signed bytes, wrapping int32 accumulate: the exact PTX dp4a.s32.s32 semantics. +inline int dot4_s(uint a, uint b, int acc) { + int4 va = int4(as_type(a)); + int4 vb = int4(as_type(b)); + return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w; +} +// dp4a, unsigned bytes, wrapping uint32 accumulate: dp4a.u32.u32. +inline uint dot4_u(uint a, uint b, uint acc) { + uint4 va = uint4(as_type(a)); + uint4 vb = uint4(as_type(b)); + return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w; +} + +kernel void probe_alu(constant uint& steps [[buffer(0)]], constant uint& seed [[buffer(1)]], + device uint* out [[buffer(2)]], uint g [[thread_position_in_grid]]) { + uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + for (uint s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + rotate(y, 7u); y = (y ^ x) + s; } + out[g] = x ^ y; +} + +kernel void probe_dot4s(constant uint& steps [[buffer(0)]], constant uint& seed [[buffer(1)]], + device uint* out [[buffer(2)]], uint g [[thread_position_in_grid]]) { + uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + int acc = int(pm_mix(x)); + for (uint s = 0u; s < steps; ++s) { + acc = dot4_s(x, y, acc); + x = x * 0x9E3779B1u + uint(acc); + y = rotate(y, 7u) ^ (uint(acc) + s); + } + out[g] = uint(acc) ^ x ^ y; +} + +kernel void probe_dot4u(constant uint& steps [[buffer(0)]], constant uint& seed [[buffer(1)]], + device uint* out [[buffer(2)]], uint g [[thread_position_in_grid]]) { + uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; + uint acc = pm_mix(x); + for (uint s = 0u; s < steps; ++s) { + acc = dot4_u(x, y, acc); + x = x * 0x9E3779B1u + acc; + y = rotate(y, 7u) ^ (acc + s); + } + out[g] = acc ^ x ^ y; +} +""" + +// CPU reference of the dot4 chain for one lane, to check the kernel is the arithmetic it claims (bit-exact). +func pmMix(_ v: UInt32) -> UInt32 { + var x = v + x ^= x >> 16; x = x &* 0x7feb352d; x ^= x >> 15; x = x &* 0x846ca68b; x ^= x >> 16 + return x +} +func dot4sRef(_ a: UInt32, _ b: UInt32, _ acc: Int32) -> Int32 { + var r = acc + for i in 0..<4 { + let ba = Int32(Int8(truncatingIfNeeded: a >> (8 * UInt32(i)))) + let bb = Int32(Int8(truncatingIfNeeded: b >> (8 * UInt32(i)))) + r = r &+ ba &* bb + } + return r +} +func dot4uRef(_ a: UInt32, _ b: UInt32, _ acc: UInt32) -> UInt32 { + var r = acc + for i in 0..<4 { + let ba = UInt32(UInt8(truncatingIfNeeded: a >> (8 * UInt32(i)))) + let bb = UInt32(UInt8(truncatingIfNeeded: b >> (8 * UInt32(i)))) + r = r &+ ba &* bb + } + return r +} +func rotl(_ v: UInt32, _ n: UInt32) -> UInt32 { (v << n) | (v >> (32 - n)) } +func laneRef(kernel: String, g: UInt32, seed: UInt32, steps: UInt32) -> UInt32 { + var x = pmMix(g ^ seed), y = x ^ 0x5bd1e995 + switch kernel { + case "probe_alu": + for s in 0../include -o dot4-probe-cl.exe dot4-probe.c (OpenCL.dll loaded at run time) + * Run: dot4-probe-cl [--list] [--device N] [--lanes N] [--steps N] [--reps N] + */ +#define CL_TARGET_OPENCL_VERSION 120 +#ifdef __APPLE__ +#include +#else +#include +#endif +#ifdef IGNEUM_CL_DYNAMIC +#include "cl_dynamic.h" +#endif +#include +#include +#include +#include + +static const char* COMMON = + "static inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n" + "#define CHAIN_HEAD uint g = (uint)get_global_id(0); uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; int acc = (int)pm_mix(x);\n" + "#define CHAIN_TAIL x = x * 0x9E3779B1u + (uint)acc; y = rotate(y, 7u) ^ ((uint)acc + s);\n" + "#define CHAIN_OUT out[g] = (uint)acc ^ x ^ y;\n"; + +static const char* K_ALU = + "__kernel void probe(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"; +static const char* K_EMUL = + "static inline int dot4e(uint a, uint b, int acc) {\n" + " int4 va = convert_int4(as_char4(a)); int4 vb = convert_int4(as_char4(b));\n" + " return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w;\n" + "}\n" + "__kernel void probe(uint steps, uint seed, __global uint* out) {\n" + " CHAIN_HEAD\n" + " for (uint s = 0u; s < steps; ++s) { acc = dot4e(x, y, acc); CHAIN_TAIL }\n" + " CHAIN_OUT\n" + "}\n"; +static const char* K_AMD = + "__kernel void probe(uint steps, uint seed, __global uint* out) {\n" + " CHAIN_HEAD\n" + " for (uint s = 0u; s < steps; ++s) { acc = __builtin_amdgcn_sudot4(true, (int)x, true, (int)y, acc, false); CHAIN_TAIL }\n" + " CHAIN_OUT\n" + "}\n"; +static const char* K_KHR = + "#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable\n" + "__kernel void probe(uint steps, uint seed, __global uint* out) {\n" + " CHAIN_HEAD\n" + " for (uint s = 0u; s < steps; ++s) { acc = acc + dot(as_char4(x), as_char4(y)); CHAIN_TAIL }\n" + " CHAIN_OUT\n" + "}\n"; +static const char* K_NV = + "static inline int dot4nv(uint a, uint b, int acc) { int d; asm(\"dp4a.s32.s32 %0, %1, %2, %3;\" : \"=r\"(d) : \"r\"(a), \"r\"(b), \"r\"(acc)); return d; }\n" + "__kernel void probe(uint steps, uint seed, __global uint* out) {\n" + " CHAIN_HEAD\n" + " for (uint s = 0u; s < steps; ++s) { acc = dot4nv(x, y, acc); CHAIN_TAIL }\n" + " CHAIN_OUT\n" + "}\n"; + +static uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; } +static uint32_t rotl32(uint32_t v, uint32_t n) { return (v << n) | (v >> (32u - n)); } +static int32_t dot4_ref(uint32_t a, uint32_t b, int32_t acc) { + int32_t r = acc; int i; + for (i = 0; i < 4; ++i) { int32_t ba = (int8_t)((a >> (8 * i)) & 0xffu), bb = (int8_t)((b >> (8 * i)) & 0xffu); r = (int32_t)((uint32_t)r + (uint32_t)(ba * bb)); } + return r; +} +static uint32_t lane_ref(int alu, uint32_t g, uint32_t seed, uint32_t steps) { + uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u, s; int32_t acc; + if (alu) { for (s = 0; s < steps; ++s) { x = x * 0x9E3779B1u + rotl32(y, 7u); y = (y ^ x) + s; } return x ^ y; } + acc = (int32_t)pm_mix(x); + for (s = 0; s < steps; ++s) { acc = dot4_ref(x, y, acc); x = x * 0x9E3779B1u + (uint32_t)acc; y = rotl32(y, 7u) ^ ((uint32_t)acc + s); } + return (uint32_t)acc ^ x ^ y; +} + +typedef struct { cl_platform_id p; cl_device_id d; char pname[128], dname[128], driver[64], ver[64]; } Dev; +static Dev devs[32]; static int ndevs = 0; +static void enumerate(void) { + cl_platform_id ps[8]; cl_uint np = 0, i; + if (clGetPlatformIDs(8, ps, &np) != CL_SUCCESS) return; + for (i = 0; i < np; ++i) { + cl_device_id ds[8]; cl_uint nd = 0, j; + if (clGetDeviceIDs(ps[i], CL_DEVICE_TYPE_GPU, 8, ds, &nd) != CL_SUCCESS) continue; + for (j = 0; j < nd && ndevs < 32; ++j) { + Dev* v = &devs[ndevs++]; v->p = ps[i]; v->d = ds[j]; + clGetPlatformInfo(ps[i], CL_PLATFORM_NAME, sizeof v->pname, v->pname, NULL); + clGetDeviceInfo(ds[j], CL_DEVICE_NAME, sizeof v->dname, v->dname, NULL); + clGetDeviceInfo(ds[j], CL_DRIVER_VERSION, sizeof v->driver, v->driver, NULL); + clGetDeviceInfo(ds[j], CL_DEVICE_VERSION, sizeof v->ver, v->ver, NULL); + } + } +} + +static int run_variant(cl_context ctx, cl_command_queue q, cl_device_id dev, const char* name, const char* body, int alu, cl_uint lanes, cl_uint steps, int reps, cl_mem out, uint32_t* host, const char* vendor, const char* dname) { + const char* srcs[2] = { COMMON, body }; cl_int err; cl_program prog; cl_kernel k; int r, ok = 1; double best = -1; + prog = clCreateProgramWithSource(ctx, 2, srcs, NULL, &err); + if (err != CL_SUCCESS) { printf("| %s | build failed | clCreateProgramWithSource %d | | | | |\n", name, (int)err); return 0; } + err = clBuildProgram(prog, 1, &dev, "-cl-std=CL1.2", NULL, NULL); + if (err != CL_SUCCESS) { + size_t n = 0; char* log; char* nl; + clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &n); log = (char*)calloc(n + 1, 1); + if (n) clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, n, log, NULL); + while (*log == '\n' || *log == '\r') ++log; + nl = strpbrk(log, "\r\n"); if (nl) *nl = 0; + printf("| %s | build failed | %.160s | | | | |\n", name, log); + printf("RESULT DOT4 vendor=%s device=\"%s\" kernel=%s build=failed\n", vendor, dname, name); + clReleaseProgram(prog); return 0; + } + k = clCreateKernel(prog, "probe", &err); + if (err != CL_SUCCESS) { printf("| %s | build failed | clCreateKernel %d | | | | |\n", name, (int)err); clReleaseProgram(prog); return 0; } + for (r = 0; r < reps; ++r) { + cl_uint seed = 0x2468aceu + (cl_uint)r * 0x9E3779B9u; size_t global = lanes, local = 256; cl_event ev; cl_ulong t0, t1; double ms; uint32_t gs[2]; int j; + clSetKernelArg(k, 0, sizeof(cl_uint), &steps); clSetKernelArg(k, 1, sizeof(cl_uint), &seed); clSetKernelArg(k, 2, sizeof(cl_mem), &out); + err = clEnqueueNDRangeKernel(q, k, 1, NULL, &global, &local, 0, NULL, &ev); + if (err != CL_SUCCESS) { printf("| %s | launch failed | %d | | | | |\n", name, (int)err); clReleaseKernel(k); clReleaseProgram(prog); return 0; } + clWaitForEvents(1, &ev); + clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START, sizeof t0, &t0, NULL); clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END, sizeof t1, &t1, NULL); + ms = (double)(t1 - t0) / 1e6; clReleaseEvent(ev); + if (best < 0 || ms < best) best = ms; + clEnqueueReadBuffer(q, out, CL_TRUE, 0, (size_t)lanes * 4, host, 0, NULL, NULL); + gs[0] = 0; gs[1] = lanes - 1; + for (j = 0; j < 2; ++j) { uint32_t want = lane_ref(alu, gs[j], seed, steps); if (host[gs[j]] != want) { ok = 0; printf("MISMATCH %s lane %u: gpu %08x cpu %08x\n", name, gs[j], host[gs[j]], want); } } + } + { + double sps = (double)lanes * (double)steps / (best / 1000.0); + printf("| %s | %u | %u | %.3f | %.2f | %.3f | %s |\n", name, lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ok ? "yes" : "NO"); + printf("RESULT DOT4 vendor=%s device=\"%s\" kernel=%s lanes=%u steps=%u best_ms=%.3f gsteps_per_s=%.2f ok=%d\n", vendor, dname, name, lanes, steps, best, sps / 1e9, ok); + } + clReleaseKernel(k); clReleaseProgram(prog); return 1; +} + +int main(int argc, char** argv) { + cl_uint lanes = 1u << 20, steps = 4096u; int reps = 3, device = 0, list = 0, i; Dev* v; cl_int err; cl_context ctx; cl_command_queue q; cl_mem out; uint32_t* host; const char* vendor; + for (i = 1; i < argc; ++i) { + if (!strcmp(argv[i], "--list")) list = 1; + else if (!strcmp(argv[i], "--device") && i + 1 < argc) device = atoi(argv[++i]); + else if (!strcmp(argv[i], "--lanes") && i + 1 < argc) lanes = (cl_uint)strtoul(argv[++i], 0, 10); + else if (!strcmp(argv[i], "--steps") && i + 1 < argc) steps = (cl_uint)strtoul(argv[++i], 0, 10); + else if (!strcmp(argv[i], "--reps") && i + 1 < argc) reps = atoi(argv[++i]); + else { printf("unknown argument %s\n", argv[i]); return 2; } + } +#ifdef IGNEUM_CL_DYNAMIC + if (!ig_cl_load()) { printf("%s\n", ig_cl_error); return 1; } +#endif + enumerate(); + if (list || ndevs == 0) { for (i = 0; i < ndevs; ++i) printf("[%d] %s | %s | driver %s | %s\n", i, devs[i].dname, devs[i].pname, devs[i].driver, devs[i].ver); if (ndevs == 0) printf("no OpenCL GPU devices\n"); return ndevs ? 0 : 1; } + if (device < 0 || device >= ndevs) { printf("no device %d (have %d)\n", device, ndevs); return 2; } + v = &devs[device]; + vendor = strstr(v->pname, "NVIDIA") ? "nvidia" : (strstr(v->pname, "AMD") ? "amd" : (strstr(v->pname, "Apple") ? "apple" : "other")); + ctx = clCreateContext(NULL, 1, &v->d, NULL, NULL, &err); if (err != CL_SUCCESS) { printf("clCreateContext %d\n", (int)err); return 1; } + q = clCreateCommandQueue(ctx, v->d, CL_QUEUE_PROFILING_ENABLE, &err); if (err != CL_SUCCESS) { printf("clCreateCommandQueue %d\n", (int)err); return 1; } + out = clCreateBuffer(ctx, CL_MEM_READ_WRITE, (size_t)lanes * 4, NULL, &err); if (err != CL_SUCCESS) { printf("clCreateBuffer %d\n", (int)err); return 1; } + host = (uint32_t*)malloc((size_t)lanes * 4); + printf("dot4-probe (OpenCL) on [%d] %s | %s | driver %s | %s, lanes %u, steps %u, best of %d, device event time\n", device, v->dname, v->pname, v->driver, v->ver, lanes, steps, reps); + { + char ext[8192]; ext[0] = 0; clGetDeviceInfo(v->d, CL_DEVICE_EXTENSIONS, sizeof ext, ext, NULL); + printf("cl_khr_integer_dot_product listed: %s\n", strstr(ext, "cl_khr_integer_dot_product") ? "yes" : "no"); + } + printf("| kernel | lanes | steps | best ms | G steps/s (= G dot4/s for dot4 rows) | ns per dependent step | lanes 0 and last ok |\n|---|---|---|---|---|---|---|\n"); + run_variant(ctx, q, v->d, "alu", K_ALU, 1, lanes, steps, reps, out, host, vendor, v->dname); + run_variant(ctx, q, v->d, "dot4e", K_EMUL, 0, lanes, steps, reps, out, host, vendor, v->dname); + run_variant(ctx, q, v->d, "dot4_khr", K_KHR, 0, lanes, steps, reps, out, host, vendor, v->dname); + if (!strcmp(vendor, "amd")) run_variant(ctx, q, v->d, "dot4_amd", K_AMD, 0, lanes, steps, reps, out, host, vendor, v->dname); + if (!strcmp(vendor, "nvidia")) run_variant(ctx, q, v->d, "dot4_nv", K_NV, 0, lanes, steps, reps, out, host, vendor, v->dname); + clReleaseMemObject(out); clReleaseCommandQueue(q); clReleaseContext(ctx); free(host); + printf("dot4-probe: done\n"); + return 0; +}