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 the project lead); layer 7 dp4a probes for Metal, CUDA and OpenCL (standalone, no lottery kernel)
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs
parent
8c898d3f60
commit
32944d12d4
4 changed files with 694 additions and 0 deletions
227
docs/analysis/sram-mirror.md
Normal file
227
docs/analysis/sram-mirror.md
Normal file
|
|
@ -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 the project lead, 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.
|
||||
115
proto-cuda/dot4-probe.cu
Normal file
115
proto-cuda/dot4-probe.cu
Normal file
|
|
@ -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 <cuda_runtime.h>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstring>
|
||||
#include <cstdint>
|
||||
|
||||
__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<<<lanes / 256, 256>>>(steps, seed, d_out);
|
||||
else if (k == 1) probe_dot4i<<<lanes / 256, 256>>>(steps, seed, d_out);
|
||||
else probe_dot4e<<<lanes / 256, 256>>>(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;
|
||||
}
|
||||
165
proto-metal/dot4-probe.swift
Normal file
165
proto-metal/dot4-probe.swift
Normal file
|
|
@ -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 <metal_stdlib>
|
||||
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<char4>(a));
|
||||
int4 vb = int4(as_type<char4>(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<uchar4>(a));
|
||||
uint4 vb = uint4(as_type<uchar4>(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..<steps { x = x &* 0x9E3779B1 &+ rotl(y, 7); y = (y ^ x) &+ s }
|
||||
return x ^ y
|
||||
case "probe_dot4s":
|
||||
var acc = Int32(bitPattern: pmMix(x))
|
||||
for s in 0..<steps { acc = dot4sRef(x, y, acc); x = x &* 0x9E3779B1 &+ UInt32(bitPattern: acc); y = rotl(y, 7) ^ (UInt32(bitPattern: acc) &+ s) }
|
||||
return UInt32(bitPattern: acc) ^ x ^ y
|
||||
default:
|
||||
var acc = pmMix(x)
|
||||
for s in 0..<steps { acc = dot4uRef(x, y, acc); x = x &* 0x9E3779B1 &+ acc; y = rotl(y, 7) ^ (acc &+ s) }
|
||||
return acc ^ x ^ y
|
||||
}
|
||||
}
|
||||
|
||||
var lanes = 1 << 20, steps: UInt32 = 4096, reps = 3
|
||||
var args = Array(CommandLine.arguments.dropFirst())
|
||||
while !args.isEmpty {
|
||||
let a = args.removeFirst()
|
||||
switch a {
|
||||
case "--lanes": lanes = Int(args.removeFirst())!
|
||||
case "--steps": steps = UInt32(args.removeFirst())!
|
||||
case "--reps": reps = Int(args.removeFirst())!
|
||||
default: print("unknown argument \(a)"); exit(2)
|
||||
}
|
||||
}
|
||||
|
||||
guard let dev = MTLCreateSystemDefaultDevice() else { print("no Metal device"); exit(1) }
|
||||
let lib: MTLLibrary
|
||||
do { lib = try dev.makeLibrary(source: source, options: nil) } catch { print("compile failed: \(error)"); exit(1) }
|
||||
let queue = dev.makeCommandQueue()!
|
||||
let outBuf = dev.makeBuffer(length: lanes * 4, options: .storageModeShared)!
|
||||
print("dot4-probe on \(dev.name), lanes \(lanes), steps \(steps), best of \(reps), GPU start-to-end time")
|
||||
print("| kernel | lanes | steps | best ms | G steps/s (= G dot4/s for dot4 rows) | ns per dependent step | lane 0 ok |")
|
||||
print("|---|---|---|---|---|---|---|")
|
||||
for name in ["probe_alu", "probe_dot4s", "probe_dot4u"] {
|
||||
let fn = lib.makeFunction(name: name)!
|
||||
let pso = try! dev.makeComputePipelineState(function: fn)
|
||||
let tg = min(256, pso.maxTotalThreadsPerThreadgroup)
|
||||
var best = Double.infinity
|
||||
var okAll = true
|
||||
for r in 0..<reps {
|
||||
var st = steps
|
||||
var seed = UInt32(0x2468ace) &+ UInt32(r) &* 0x9E3779B9
|
||||
let cb = queue.makeCommandBuffer()!
|
||||
let enc = cb.makeComputeCommandEncoder()!
|
||||
enc.setComputePipelineState(pso)
|
||||
enc.setBytes(&st, length: 4, index: 0)
|
||||
enc.setBytes(&seed, length: 4, index: 1)
|
||||
enc.setBuffer(outBuf, offset: 0, index: 2)
|
||||
enc.dispatchThreads(MTLSize(width: lanes, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: tg, height: 1, depth: 1))
|
||||
enc.endEncoding()
|
||||
cb.commit()
|
||||
cb.waitUntilCompleted()
|
||||
let ms = (cb.gpuEndTime - cb.gpuStartTime) * 1000.0
|
||||
if ms < best { best = ms }
|
||||
// bit-exactness of the kernel against the CPU reference on two lanes
|
||||
let p = outBuf.contents().bindMemory(to: UInt32.self, capacity: lanes)
|
||||
for g in [UInt32(0), UInt32(lanes - 1)] {
|
||||
let want = laneRef(kernel: name, g: g, seed: seed, steps: steps)
|
||||
if p[Int(g)] != want { okAll = false; print("MISMATCH \(name) lane \(g): gpu \(String(p[Int(g)], radix: 16)) cpu \(String(want, radix: 16))") }
|
||||
}
|
||||
}
|
||||
let stepsPerS = Double(lanes) * Double(steps) / (best / 1000.0)
|
||||
print(String(format: "| %@ | %d | %u | %.3f | %.2f | %.3f | %@ |", name, lanes, steps, best, stepsPerS / 1e9, best * 1e6 / Double(steps), okAll ? "yes" : "NO"))
|
||||
}
|
||||
print("dot4-probe: done")
|
||||
187
proto-opencl/dot4-probe.c
Normal file
187
proto-opencl/dot4-probe.c
Normal file
|
|
@ -0,0 +1,187 @@
|
|||
/* dot4-probe (OpenCL): dp4a-class throughput on AMD and NVIDIA through OpenCL, standalone (no pack, no lottery kernel).
|
||||
* Counter ASIC 2.0 layer 7 (docs/analysis/int8-matrix-family.md), 5 October 2026. PC job; the Mac can run it on Apple
|
||||
* OpenCL for the scalar rows only.
|
||||
*
|
||||
* Dependent chains, same shape as --memprobe's ALU chain (host.c, probe_alu) and the Metal and CUDA probes:
|
||||
* 1,048,576 lanes x 4,096 steps, best of 3, device event time.
|
||||
* alu x = x * K + rotate(y, 7); y = (y ^ x) + s the card's integer baseline
|
||||
* dot4e scalar emulation: 4 sign-extended byte products summed into a wrapping int accumulator
|
||||
* dot4_amd __builtin_amdgcn_sudot4(a, 1, b, 1, acc, 0): V_DOT4_I32_IU8 with both operands signed, no clamp (RDNA 3;
|
||||
* the same builtin on RDNA 4 is the thing to check: LLVM gates it on the dot7-insts / dot8-insts features)
|
||||
* dot4_khr cl_khr_integer_dot_product: acc + dot(as_char4(a), as_char4(b)) (OpenCL C 3.0 extension, section
|
||||
* 6.2.2.16 of the OpenCL C specification; listed by a platform or not)
|
||||
* dot4_nv inline PTX "dp4a.s32.s32" (NVIDIA's OpenCL compiler accepts inline PTX asm; PTX ISA 9.7.1.24)
|
||||
* Each variant is built on its own: a variant the platform cannot compile prints one "build failed" row with the first
|
||||
* line of the build log and the run goes on. Every variant that runs is checked bit for bit against the CPU reference
|
||||
* on two lanes, so an intrinsic with different semantics (a saturating accumulate, say) shows as NO in the ok column.
|
||||
*
|
||||
* Build (Mac, Apple OpenCL): cc -std=c99 -O2 -o dot4-probe-cl dot4-probe.c -framework OpenCL
|
||||
* Build (Windows, mingw, no SDK): x86_64-w64-mingw32-gcc -std=c99 -O2 -static -DIGNEUM_CL_DYNAMIC -DCL_TARGET_OPENCL_VERSION=120 \
|
||||
* -I <redist>/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 <OpenCL/cl.h>
|
||||
#else
|
||||
#include <CL/cl.h>
|
||||
#endif
|
||||
#ifdef IGNEUM_CL_DYNAMIC
|
||||
#include "cl_dynamic.h"
|
||||
#endif
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <stdint.h>
|
||||
|
||||
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;
|
||||
}
|
||||
Loading…
Reference in a new issue