Operations Performance

Measured performance

What this page covers. Every reproducible throughput number for the CMP 170HX: compute rates per datatype before and after the unlock, HBM bandwidth, PCIe bandwidth at each link configuration that anyone has actually reached, and how all of it compares against a real A100. Test conditions are stated on every row. Tuning levers (clock offsets, power limits, profiles) live on tuning.md; inference-specific numbers live on llm-inference.md.

The headline result. The compute unlock is a register write, not a clock change. Writing SS0 0x0082381c = 0x88888888 and SS1 0x00823820 = 0x00000008 (after opening the FEAT PLM at 0x00823804) takes FP32 from 0.30-0.41 TFLOPS to 12.2-12.8 TFLOPS, a 26-32x gain, at unchanged clocks. Nothing in the shipping unlocker touches core clock, memory clock, power limit or PCIe link speed, so every figure below is a stock-clock figure unless it says otherwise. See compute-throttle.md for the mechanism.

The control that proves the mechanism. In the same-card A/B where FP32 moves 30.7x, memory bandwidth moves from 1592 GB/s to 1599 GB/s, a ratio of 1.0x. NVIDIA restricted the instruction issue rate, not the memory subsystem.

---

How to read these numbers

• Stock clocks throughout. Sustained SM clock is 1410 MHz (1470 MHz with nvidia-smi -pl 300), base 1140 MHz. The clocks.max.sm = 1935 MHz field that nvidia-smi reports post-unlock is a reported field, not an achievable clock; it is a single report, never re-checked, and the VBIOS table maximum is 1695 MHz. Treat 1935 MHz as low confidence.
• The throttle is per-instruction-class. It is not a global multiplier. On a clean CMP 90HX the ratios were FP64 1/64, FP32 1/32, FP16 1x (untouched), INT32 1/2, INT8 1/16. The 170HX follows the same pattern with different divisors, which is why one datatype can look normal while another is 30x down.
• Two INT8 numbers are both real. The tensor MMA path and the library/OpenCL path differ by 7.6x. Do not average them.
• Compression is off. Shipping patch 0005-ce-scrub-workarounds.patch forces *pteKind = NV_MMU_PTE_KIND_GENERIC_MEMORY for 0x20C2 and 0x2082, where stock returns NV_MMU_PTE_KIND_GENERIC_MEMORY_COMPRESSIBLE_DISABLE_PLC. Any bandwidth comparison against a real A100 must note this.
• Two measurement gotchas that produce garbage. CUBLAS_COMPUTE_16F with a float alpha/beta pointer returns instantly and reports an absurd 10748 TFLOPS: that is a no-op, not a result, so use the fp32-accumulate rows. And CUDA 13 removed cudaDeviceProp::clockRate and memoryClockRate, so query cudaDeviceGetAttribute with cudaDevAttrClockRate / cudaDevAttrMemoryClockRate instead.

Warning

Theoretical peaks masquerading as measurements

Three separate figures in circulation are tool-computed device properties, not run results. 1769.47 GB/sec is exactly 864 MHz x 4 x 4096 bits / 8; 12633.60 GFlops is exactly 4480 x 2 x 1410 MHz. Both are derived from the clock and bus-width fields printed two lines above them in the same mixbench dump. No card has ever measured 1769 GB/s. The same 12633.6 GFLOPS reappears in a second dump as "4480 cores, 12.634 TFLOPs/s", again as a device property.

---

Compute: locked versus unlocked

All rows measured on compute-unlocked cards at stock clocks unless noted. The locked column is the same silicon before the SS0/SS1 writes.

Datatype / path	Locked	Unlocked	Ratio	Conditions
FP32 non-tensor	0.30 / 0.40 / 0.41 TFLOPS	12.58 TFLOPS	28.4x-30.7x	1024³ / 4096³ / 8192³ modded sweep, same card
FP32 SGEMM	393 Gflop/s	12,233-12,256 Gflop/s	31x	gpu-burn, peak 67 C
FP32 (rented cohort)	0.39 TFLOPS	12.6 TFLOPS	32.3x	mean of 8 cards, driver 610.43.02
FP64 non-tensor	0.20 TFLOPS	6.223-6.31 TFLOPS	~31x	1/2 of FP32, the full unrestricted GA100 rate
FP64 tensor	197 / 191 Gflop/s (paper)	11.6-11.96 TFLOPS	59x-62x (paper)	11,668 / 11,786 Gflop/s DGEMM / FP64 tensor
TF32 tensor	2.96-3.21 TFLOPS	79-94 TFLOPS	15x-28x	widest spread of any datatype
FP16 tensor	6.01 TFLOPS	158.7-190 TFLOPS	~27x	fp32 accumulate at the low end
BF16 tensor	6.41 TFLOPS	164.4-192.7 TFLOPS	~28x	ceiling 202.1 TFLOPS
INT8 tensor MMA	1.60 TOPS	335.0-335.6 TOPS	n/a	per-instruction microbenchmark
INT8 library / OpenCL	1.60 TOPS	43.33-47.894 TOPS	27.0x	torch, cuBLAS, OpenCL-Benchmark
INT4 tensor MMA	n/a	320.2 TOPS	n/a	mma_s4s4s32_8_8_32
FP4 / FP6 / FP8 MMA	not supported	not supported	n/a	expected for sm_80

FP32, the load-bearing number

Seven tools on at least a dozen distinct cards land inside a 12.2-12.8 TFLOPS band.

Value	Tool / conditions
12.72 TFLOPS	torch GEMM 8192²
12.76 TFLOPS	gemm_probe.cu n=8192, 30 iterations
12.58 TFLOPS	8192³ modded sweep
12,565.14 GFLOPS	clpeak, driver 13.0 / CUDA 13.3
12.493 TFLOPs/s	OpenCL-Benchmark, 10 GB to 40 GB card
12.6 TFLOPS	mean of eight rented cards
12,233-12,256 Gflop/s	paper Table 2, gpu-burn
12,229-12,254 Gflop/s	sustained burn-in, 268435456 B buffers, 24 iterations
11.1 TFLOPS	per-instruction scalar fma_fp32 microbenchmark (lower bound)

The theoretical peak for a 70-SM GA100 at 1410 MHz is 12,633.6 GFLOPS, so the card achieves roughly 99% of arithmetic peak. Locked, independent measurements sit at 0.3159 TFLOPS FFMA / 0.32 / 0.39 TFLOPS / 393 Gflop/s SGEMM / 367 GFLOPS clpeak, all inside the 1/32 issue-rate model.

FP64 has two rates and they must not be conflated

Non-tensor FP64 runs at exactly 1/2 of FP32 (6.223 TFLOPs/s OpenCL, 6308.65 GFLOPS clpeak, ~6,200 GFLOPS DGEMM, 5.6 TFLOPS for a pure scalar fma_fp64 microbenchmark). FP64 tensor runs at roughly 2x that (11.65 TFLOPS, 11.96 TFLOPS clpeak WMMA fp64 8x8x4, 11.6 TFLOPS across eight rented cards, 11,668-11,786 Gflop/s in the paper). The 1/2 ratio is the full unrestricted GA100 rate: FP64 is genuinely restored, not partially.

Tensor throughput detail

Datatype	Value	Tool / shape
TF32	79.0 TFLOPS	8 rented cards, torch GEMM
TF32	80.59 / 84.75 / 51.53 TFLOPS	8192³ / 4096³ / 1024³ modded sweep
TF32	81.35 TFLOPS	torch GEMM 8192²
TF32	83.2 TFLOPS	mma_tf32tf32f32_16_16_8
TF32	88.9-91.9 TFLOPS	gemm_probe.cu n=8192
TF32	89.69 TFLOPS	clpeak mma.sync m16n8k8
TF32	94,103 Gflop/s	paper, gpu-burn, 64 C
FP16 (fp32 acc)	158.7-160.0 TFLOPS	gemm_probe.cu n=8192
FP16 (fp32 acc)	162.7 TFLOPS	8 rented cards
FP16	174.11 TFLOPS	4096³ modded sweep
FP16	175.79 TFLOPS	torch GEMM 4096²
FP16 (fp32 acc)	179.1 TFLOPS	mma_f16f16f32, both tile shapes
FP16 (fp16 acc)	180.2 / 180.3 TFLOPS	mma_f16f16f16_16_16_16 / _32_8_16
FP16 (fp16 acc)	189.66 TFLOPS	clpeak mma.sync m16n8k16
BF16	164.4 TFLOPS	mma_bf16bf16f32, both tile shapes
BF16	171.4 TFLOPS	8 rented cards
BF16	180.09 TFLOPS	torch GEMM 4096²
BF16	183.75 TFLOPS	4096³ modded sweep
BF16	188.1-192.7 TFLOPS	gemm_probe.cu n=8192, fp32 accumulate
BF16 ceiling	202.1 TFLOPS	arithmetic: 2048 x 70 SM x 1410 MHz, verified exact

Two spreads are unexplained and are recorded as such. TF32 varies 19% across seven tools while FP16 and BF16 stay tight. FP16 with fp16 accumulate reads consistently above FP16 with fp32 accumulate, where on A100 the two should be the same rate; the likeliest explanation is that mmapeak-style microbenchmarks keep operands in shared memory and therefore flatter the card.

Note

Open problem: INT4 measures below INT8

mma_s4s4s32_8_8_32 returns 320.2 TOPS against INT8's 335.0/335.6 TOPS. On Ampere INT4 tensor throughput should be roughly 2x INT8. Nobody re-ran it. The INT8 side is sound (two tile shapes agree); the INT4 side is one run. Re-running the INT4 shape with a longer target time and varied tiles is the cheapest open lead in this domain.

Note

Open problem: the INT8 library path is 7.6x below the INT8 tensor path

Direct MMA gives 335 TOPS; torch, cuBLAS and OpenCL all land at 43-48 TOPS. Either those libraries are not issuing IMMA on this device, or the unlock leaves an INT8 issue-rate restriction partially in place. The suggested test is an explicit CUBLAS_COMPUTE_32I GEMM with INT8 inputs against the raw MMA figure.

Silicon-to-silicon reproducibility

Eight unlocked 64 GB cards from one rental, benchmarked in a single session on driver 610.43.02 at PCIe Gen1 x4, showed under 2.5% per-card spread and 8/8 passes on a full byte-compare VRAM integrity test. Eight cards, one report. Stock cards, no capacitor mod, no Gen2.

FP16	BF16	TF32	FP32 (locked)	FP64	INT8	HBM
162.7 TFLOPS	171.4 TFLOPS	79.0 TFLOPS	12.6 TFLOPS (0.39)	11.6 TFLOPS	44.1 TOPS	1600 GB/s

This is the strongest evidence in the archive that the unlock lands identically across cards.

Failed-unlock signature

If BF16 comes back around 12 TFLOPS instead of ~185, the payload did not land. One user reporting mixbench 12 BF16 TFLOPS, clpeak 367 GFLOPS and a custom GEMM at 6.25 TFLOPS was diagnosed as seeing "the performance of tf32 in locked mode" against the known 202 TFLOPS ceiling, accepted the diagnosis and retried. An order-of-magnitude shortfall against 202 TFLOPS is the tell. See verify.md.

Note also that gaming frame rates are not a valid verification method: one tester measured identical FPS with and without the compute unlock while LLM and CUTLASS throughput clearly moved. The unlock targets SM issue rate, not the graphics path.

---

Memory bandwidth

Quantity	Value	Conditions
Theoretical peak	1555.2 GB/s = 1448.4 GiB/s	1215 MHz DDR x 5120-bit; the two figures are the same number in different units
Peak line of a real sweep	1448 GiB/s	79.3 GiB total / 79.0 free
Measured, 8 rented 64 GB cards	1600 GB/s	identical across all eight
Stock versus modded, same card	1592 to 1599 GB/s (1.0x)	256 MB working set, the unlock control row
OpenCL coalesced read / write	1305.86 / 1521.62 GB/s	10 GB to 40 GB card, driver 580.159.03
OpenCL misaligned read / write	789.82 / 161.76 GB/s	same card
One OpenCL test	1333 GB/s	stock VBIOS
2023 external review ceiling	1355 GB/s	used for its roofline ridge points

Note

There is no single canonical HBM number

Measured HBM bandwidth spans 1305.86 GB/s to 1600 GB/s and no methodology reconciles the range. Partial reconciliation: the 8 GB card carries 4 stacks of faster HBM2e (4096-bit) while the 10 GB card carries 5 stacks of slower HBM2 (5120-bit), and read/write/misaligned patterns differ by nearly 10x within a single tool. Quote the range, not a point estimate. Separately, "1493 GB/s datasheet" and "1555 GB/s theoretical at full boost" are different quantities (the first is the A100 40 GB PCIe datasheet figure, the second is what 1215 MHz x 5120-bit computes to) and are frequently used interchangeably in circulating documents.

The 79% plateau above an 8 GiB offset

Measured on a 10 GB card fired to the abandoned 80 GB geometry, with a 1 GB memset sweep:

Offset	Bandwidth	Percent of peak	Time
0 GB	1416 GiB/s	98%	0.70-0.71 ms
1 GB	1422 GiB/s	98%	0.70-0.71 ms
2 GB	1416 GiB/s	98%	0.70-0.71 ms
4 GB	1419 GiB/s	98%	0.70-0.71 ms
8 GB through 76 GB	1147-1151 GiB/s (flat)	79%	~0.87 ms

The gap closes entirely with a large enough chunk: at a 32 GB chunk the sweep gives 1452 GiB/s (100%) at offset 0 against 1443 GiB/s (100%) at offset 40 GB, versus 1419 vs 1149 GiB/s for a 1 GB chunk.

Two shipping-code facts bear on this without settling it:

• 0004-bar0-pramin-clamp.patch pins the BAR0 window to (0x2000ULL << 20) - DRF_SIZE(NV_PRAMIN) whenever devId is 0x20C2 or 0x2082 and fbAddrSpaceSizeMb > 0x2000. That is a genuine 8 GiB discontinuity in shipping code at exactly the offset of the observed step, but PRAMIN is a CPU-side BAR0 aperture and the sweep was a device-side memset, so causality is not established.
• 0001-sec2-postbl-plm-ss-cfg.patch extends the last FB region to targetFbBytes - 1 with supportCompressed = NV_TRUE, supportISO = NV_TRUE and performance = 20. The resource manager therefore models the whole unlocked range as uniform-performance memory and has no way to prefer the fast region.

Note

Open problem: does the plateau exist on shipping geometries?

The sweep was only ever run on the abandoned 80 GB configuration. Repeating the identical offset and chunk sweeps on a shipping 8 GB to 64 GB card would settle whether this is a property of any unlocked geometry or an artefact of the over-fire. See 80gb.md.

---

PCIe bandwidth by link configuration

Two independent mechanisms, never to be conflated. Link speed (Gen1 to Gen2) is a driver/firmware change and exists only on unreleased branches. Link width (x4 to x16) is a physical board modification: 12 of the 16 lanes ship with their AC-coupling capacitors depopulated, and restoring them means hand-soldering 24 0402 parts. See pcie-gen2.md and physical-mods.md.

Link configuration	How reached	Measured bandwidth	Tool / conditions	Confidence
Gen1 x4 (stock, shipping unlock)	nothing required	send 0.80 / receive 0.84 / bidirectional 0.81 GB/s	One OpenCL-Benchmark screenshot, relayed from an outside hardware group, on a 10 GB-to-40 GB card; the tool printed the link as "Gen1 x16"	medium
Gen1 x4	nothing required	~0.85 GB/s	clpeak	high
Gen1 x4 under inference load	nothing required	~1.0 GB/s, did not ramp	8-card rig, device max reported Gen2 x16	high
Gen1 x16	capacitor mod only	2.88 GB/s (nominal ~4 GB/s)	one report, one modded card, tool unnamed; the measurer expected 3.2 GB/s and was told the gap is Gen1 signal overhead	medium
Gen2 x4	Gen2-family branch only	send 1.68 / receive 1.71 GB/s	OpenCL-Benchmark, one archived screenshot, unmodded card; the setup script independently predicts "~0.85 to ~1.7 GB/s, exactly 2x"	medium
Gen1 → Gen2 on a cap-modded card that negotiated x8	branch and a partial capacitor mod	1.67 → 3.24 GB/s	one A/B, single card, Asus Prime Z370 / i3-8100 / 8 GB RAM. Not a Gen2 x4 figure: 3.24 GB/s is above the ~2.0 GB/s Gen2 x4 ceiling because the link was x8	medium
Gen2 x4 (vendor claim)	distributed package README	"2 GB/s"	not independently logged, not achievable on shipping master	low
Gen2 x16	branch and full 24-capacitor mod	6.63-6.67 GB/s (~83% of the 8 GB/s line rate); nvtop TX 7.061 GiB/s	ocl_pcie_bw	medium

Note

Gen2 x16 bandwidth figures

Two rigs have published captures: 6.63 to 6.67 GB/s on one card on 2026-07-26, and 5.97 GB/s on each of four cards with zero AER errors over 90 minutes. Nobody has published a long burn-in. Gen2 itself ships in master, so these figures are what a capacitor-modded card with the unlocker installed should be expected to reach.

Also relevant: 12 of 24 capacitors populated yields x8, because width negotiation falls back to the next legal width (16, 8, 4, 1). An x8 result after a mod means incomplete or bridged solder work, not a distinct hardware limit.

How much PCIe actually costs you

PCIe sensitivity is engine- and topology-dependent, and the two headline results are not in conflict.

Case	Change	Result
Single card, llama.cpp	x4 to x16 (same generation)	pp 439 to 448, tg 81.91 to 85.75 (about +2%)
Three cards, llama.cpp	x4 to x16	pp 441 to 461, tg 86 to 89
Three cards, GLM-5.2 with almost all of the model on CPU (one layer plus buffers on the GPUs)	Gen1 x4 run to a later "gen2 x4 attempt"	pp2048 33.44 ± 0.37 to 48.22 ± 1.36 t/s (+44.2%), tg512 5.90 ± 0.03 to 6.39 ± 0.09, time-to-first-response 61,253.57 ± 675.94 ms to 42,510.25 ± 1,217.69 ms (-30.6%). Both percentages are computed here from two separately posted runs, not stated by the tester
Single card, Qwen3.6-27B-MTP-UD-Q8_K_XL on ik_llama with MTP, model fully VRAM-resident	Gen1 x4 to Gen2 x4, "all other factors unchanged"	pp2048 328.81 to 449.41 t/s, pp8192 363.25 to 493.86, tg128 38.15 to 41.52, tg512 37.69 to 40.12. Same tester as the row above, runs five days apart
Single card, model load	Gen1 x4	~30 s, a one-time cost
Graphics (BeamNG.drive)	Gen1 x16 versus x4, cap mod done	15 fps versus 5 fps, "still awful" either way

The reconciliation: single-card work is largely bandwidth-local (weights are resident, only the model load crosses the link), while multi-card and CPU-offload prefill is link-bound. Caveat carried from the source for the +44.2% row: the post-Gen2 run used a Q4-labelled quant where the pre-Gen2 run was unlabelled, so the prompt-processing delta may be overstated. A Q2_K_XL quant on the same rig at Gen2 x4 gave pp2048 49.00 ± 1.08 and tg512 6.81 ± 0.06. The single-card row complicates the picture rather than settling it: that model was entirely resident in VRAM, and neither the tester nor anyone in the channel could explain why the link speed moved it at all. Both rows are the same tester; there is no independent Gen2 x4 inference measurement. Full tables and conditions on llm-inference.md.

Note

Open problem: nobody has measured Gen2 x16 on a multi-card LLM rig

The Gen2 x16 bandwidth figure and the Gen1-to-Gen2 inference runs were measured on different systems and never combined. The 2x2 matrix of {Gen1, Gen2} x {x4, x16} on one rig with one model would settle both this and the width-versus-generation dispute.

---

Thermals and power, in brief

Full treatment is on thermals.md, cooling.md and power-and-psu.md; what matters for benchmarking is here.

Observation	Value
Sustained full-rate GEMM burn-in	12,229-12,254 Gflop/s flat while the die went 62 to 64 to 69 to 71 to 73 C over ~30-40 s, zero errors
Peak load temperatures (paper)	67 C FP32, 64 C FP64 tensor and TF32 tensor; a full-capability part throttles only above ~85 C
Default gpu_burn	~70 C ± 2
Idle draw / stock cap	~42 W / 250 W on the controlled rig
Power under hashcat versus FP32 burn	160+ W versus 60-75 W (2023, locked card)

Practical rule: do not validate stability or cooling with a conventional FP32 burn-in. This card is hard to load. Integer and memory benchmarks pull far more power than FP32 tools do.

The accepted post-unlock validation recipe:

# github.com/wilicc/gpu-burn
make COMPUTE=80
./gpu_burn -tc -m 90% 1200          # 20 minutes, tensor cores, 90% of VRAM
# variant used by one distributed package, expecting 0 memory errors:
./gpu_burn -m 63500 -d 30

Reported clean runs: 30 minutes on a tuned single card; 2 hours on each of four 8 GB to 64 GB cards with no instability; a 5-minute pass on a 10 GB to 40 GB card. Ensure adequate cooling first.

Caution

Never validate the 80 GB geometry as if it worked

Firing a 10 GB card to 80 GB produces gpu-burn errors, independently reproduced. 10 GB cards ship at 40 GB for this reason. See 80gb.md.

Memory overclocking was reported to buy about +2.5% (gpu_burn 12,180 Gflop/s average at default versus 12,472-12,485 Gflop/s sustained) at a cost of about 5 C (70 C ± 2 versus 75-77 C). Note that no branch named mem_overclock exists in the archived branch set, and no archived branch contains any clock, boost or p-state code at all: a grep across all thirteen trees for overclock/memclk/pstate/boost returns nothing. The result stands as a tester report that cannot be checked against code.

---

Comparison with the A100

The 170HX carries a complete GA100 die (826 mm², PMC_BOOT_0 = 0x170000a1, identical to all three A100 SKUs and the Drive A100), floorswept to 70 of the die's SMs. So the honest comparison is "same architecture, fewer SMs, worse I/O, no ECC, no NVLink".

Property	CMP 170HX (unlocked)	A100 40 GB reference	Notes
Die	GA100, 826 mm²	GA100, same die	PMC_BOOT_0 0x170000a1 on both
SMs / CUDA cores	70 / 4480	108 / 6912 (A100 SXM4 40 GB)	5 active GPCs, 35 TPCs
Compute capability	8.0 (sm_80)	8.0	identical ISA, no FP8, no NVFP4
L2 cache	32 MB (32768 KB)	n/a	TechPowerUp's 8 MB for the 170HX is wrong, corroborated by latency-spike measurement
Capacity	64 GB (8 GB SKU) or 40 GB (10 GB SKU)	40 GB	see memory-geometry.md
Bus width	4096-bit (8 GB, 4 stacks) / 5120-bit (10 GB, 5 stacks)	5120-bit	GPU-Z on an A100-PCIE-40GB reports 5120-bit, 1555.2 GB/s, 1215 MHz memory
HBM bandwidth	1305.86-1600 GB/s measured	1493 GB/s datasheet	the eight-card 1600 GB/s figure is above the A100 40 GB datasheet number
Host link	Gen1 x4 stock; Gen2 x4 on a branch; x16 only after soldering	PCIe 4.0 x16	the single largest gap
NVLink / P2P	fused off, no lever found; 0 of 56 GPU pairs report peer access	present	see nvlink.md, p2p.md
ECC	fused off, no telemetry	on	see ecc.md
Memory compression	forced off by the shipping patch	on	affects any bandwidth comparison
MIG	only a 1g.64gb profile exists; standard A100 profiles are rejected	full profile set

Application-level comparisons

The FluidX3D and hashcat rows below were measured in 2023 on a locked card, before the register unlock existed. The FluidX3D rows used the no-FMA source workaround; hashcat was run unmodified, being integer work that the FP32 FMA throttle does not reach. They are the only whole-application comparisons against a named A100 in the archive, and they are a floor, not a ceiling, for what an unlocked card should do.

Workload	CMP 170HX	A100	Ratio
FluidX3D FP32/FP32, no-FMA	7681 MLUPs/s at 1175 GB/s (458 steps/s)	8526 MLUPs/s (A100 40 GB PCIe)	90.1%
FluidX3D FP32/FP16S, no-FMA	12386 MLUPs/s at 954 GB/s (738 steps/s)	16035 MLUPs/s	77.2% (and +11.7% over an RTX 4090's 11091)
hashcat MD5	43930.0 MH/s (53.01 ms) @ Accel:64 Loops:512 Thr:1024 Vec:1	~64900 MH/s	67.7% (also slower than an RTX 3080's 54000.1 MH/s)
GLM-5.2 decode, 8-way pipeline parallel	30.2 t/s on 8 unlocked 64 GB cards	the circulated reference recipe targeted 8x A100 80 GB at an expected ~40 t/s	see llm-inference.md

Note on the FluidX3D rows: with FMA enabled on the locked card the same kernel is compute-bound at 2276 MLUPs/s and only 348 GB/s. Dropping to FP32/FP16S halves memory traffic to 173 GB/s but leaves throughput at 2250 MLUPs/s. Flat throughput across halved bandwidth is the diagnostic signature of the throttle. With FMA removed the kernel becomes memory-bound again, achieving 1175 GB/s, which is 87% of that review's 1355 GB/s ceiling.

A single-A100 GLM-5.2 figure of 55 tok/s circulates as a comparison baseline. It is second-hand with no configuration attached and is rated low confidence.

Roofline selection rule

From the 2023 external review, still the best guidance for deciding whether a kernel suits this card in its locked state: useful below 0.3 FLOPs/byte of arithmetic intensity with stock FMA, or below 4.6 FLOPs/byte after disabling FMA. Those ridge points follow from 394 GFLOPS and 6250 GFLOPS over a 1355 GB/s measured ceiling (0.291 and 4.61 exactly). After the compute unlock the FP32 ridge moves out by roughly the same 30x factor, so the rule stops binding: an unlocked card is a normal memory-bound GA100 for most kernels.

Two worked examples of the rule in action:

• A SYCL FDTD kernel at 0.25 FLOPs/byte needed no workaround at all on a locked card: 10110 MC/s and 1156992 MiB/s (16777216 cells x 1000 timesteps in 1.66 s), about 1.5x a Radeon VII / Instinct MI50's ~6000 MC/s, using unmodified source. Its achieved 181.98 GFLOPS never approaches the 394 GFLOPS throttled ceiling. Note 1156992 MiB/s is base-2 and so looks smaller than the base-10 GB/s figures elsewhere.
• FluidX3D at 1.7 FLOPs/byte (261 FP32 + 102 INT32 ops over 153 B of traffic, 363 ops per cell update) sits above the locked ridge and is exactly where the FMA workaround pays.

---

Tools and the evidentiary standard

The community standard for a claimed unlock is a screenshot of ProjectPhysX/OpenCL-Benchmark. AI-written summaries were explicitly rejected as proof.

Tool	Use	Caveat
ProjectPhysX/OpenCL-Benchmark	proof-of-unlock artefact, full-device dump	does not measure tensor cores at all; reading only its output leads to "the card does 12.5 TFLOPS" and misses the ~190 TFLOPS tensor path
mixbench (CUDA)	compute sweep	prints theoretical peaks next to measurements; see the warning above
cuBLAS tensor tests / torch GEMM sweeps	representative tensor numbers	torch GEMM reads from HBM, so it is more representative than MMA microbenchmarks
ReinForce-II/mmapeak	per-instruction MMA sweep	optimistic: operands stay in shared memory
gemm_probe.cu	highest published FP32 and BF16 numbers	its TF32 (88.9-91.9) is below the paper's 94.1
clpeak	full-device dump	explicitly unsuitable for measuring the FMA/DP4A patches
gpu-burn	stability and sustained flops	build with make COMPUTE=80
ocl_pcie_bw	PCIe bandwidth	the tool behind the Gen2 x16 figure

---

Non-LLM workload results

Workload	Result	Conditions
SDXL 1024², 30 steps	4.73 s (6.35 it/s, 10.5 GB) versus 7.59 s (3.95 it/s) on an RTX 3090 = 1.60x	identical script
Wan2.1-T2V, 81 frames at 480p	73.4 s (0.91 s/frame, 18.5 GB) versus 132.8 s = 1.81x	identical script
LTX-Video, 81 frames	11.0 s (0.14 s/frame, 15.9 GB) versus 20.0 s = 1.82x	identical script
Wan2.1, 129 frames at 720p	1,485 s using 33.3 GB versus OOM on the 3090's 24 GB	identical script
pearlhash mining	3 TH to 147 TH after the unlock (~49x, one tester); 140-170 "th" at 200 W after a wildrig update	the Pearl network's aggregate hashrate doubled after the unlocker was released
Gravity bench, 50k asteroids	18 FPS	best reported; unlock state not stated
FurMark	56 fps	pre-memory-unlock

Diffusion is the standout non-LLM fit: the workload is compute-bound and fits entirely in VRAM, so the Gen1 x4 link never bites.

Note

Open problem: the mining unit is uninterpretable

"140-170 th @ 200 W" is reported with no unit expansion and no second per-card source. The network-level doubling is solid; the per-card figure is not usable as written.

---

Open questions in this domain

1. Why INT4 measures below INT8.
2. Why the INT8 library path sits 7.6x below the INT8 tensor path.
3. Whether the 79%-of-peak bandwidth plateau above 8 GiB applies to shipping 64 GB and 40 GB geometries.
4. Whether the n_ubatch scheduling cliff seen on a CMP 100-210 (pp512 353.59 to 977.20 with flash attention off, 380.96 to 1159.39 with it on, a 3.04x gain from one flag) also exists on a 70-SM 170HX. This is the single cheapest untested lead in the archive: sweep n_ubatch from 48 to 80 with llama-bench.
5. Why TF32 spreads 79-94 TFLOPS when BF16 and FP16 are tight.
6. Whether the Gen2 x16 bandwidth result translates into multi-card LLM throughput.
7. What practical P2P bandwidth between two 170HX cards is, if any.

See open-questions.md and dead-ends.md for the full register.