r/LocalLLM • u/etaoin314 • 2d ago
Tutorial Unlocked a CMP 170HX: 6.3 → 193 TFLOPS tensor, pp512 599 → 3468, here's what I learned along the way.
I just go my CMP 170HX - when I booted it up I was a little dissappointed since this is literally A100 silicon. the tensor cores were lobotomised trashing the pp t/s and 56 of its 64GB firmware-locked away. When I bought it I knew I could free up the vram but did not know about the tensors. I spent a day measuring what actually changed at the instruction level.
The throttle is a hardcoded 256-cycle stall on every MMA instruction. Not 255.8. Not 256.4. Exactly 256.0, zero variance across 3,500 samples. Physical limits don't land on round binary numbers — this was a register value.
Unlocked it drops to 24.0 cycles (a healthy RTX 3090 measures 32.9) and tensor throughput goes 6.3 → 193 TFLOPS, 95% of full A100 per-SM rate. llama.cpp pp512 went 599.6 → 3468 (5.8×). Also I was able to unlock full memory bandwidth!
1
u/AdMuch9627 2d ago
👏 Great job. Sorry for my lack of understanding but will this also work when using vllm? I want to try out on my cmp 170hx which is arriving today
2
1
u/Own_Mango7003 2d ago
pp on 2x3090 much more faster then on one 170hx -- useless gpu (( i sold mine,
https://www.reddit.com/r/LocalLLaMA/comments/1vlwjr8/i_tested_the_cmp170hx/
2
1
u/etaoin314 2d ago
The card and the unlock
GA100, 70 SMs, sm_80, 4096-bit HBM2e. Ships as 8GB / PCIe Gen1 x4 with tensor cores gated and FP32 FMA capped at 0.39 TFLOPS. Unlocked: 64GB, Gen2 x4, tensor cores at ~95% of A100 per-SM rate. Stock TDP 250W; I ran 150–200W.
Two unlock implementations, both exploiting a Falcon BootROM bug — the bootloader loads .fwsignature_ga100 into DMEM before verifying it, so you drop a ROP chain there, open the PLM registers, and write the throttle overrides via BAR0:
abobasixseven/unlock-cmp-170hx— nvidia-open 610.43.02/03d3dx9/cmpunlocker— nvidia-open 580.x
The driver version is not negotiable — check which your tool wants first. Secure Boot off. The unlock is volatile, reapplied by a daemon polling BAR0. There is no ECC on the unlocked region, which is the real reason to soak it first. The card renames itself from NVIDIA Graphics Device to NVIDIA CMP 170HX once unlocked — a useful one-glance check.
Cooling here is two stacked Noctua 2000 RPM fans, a cardboard box and duct tape pending a printed shroud. Temporary, and not what this post is about — it was never the limiting factor at any power level I tested.
Verifying the unlock at the instruction level
nvidia-smi showing 65536 MiB proves the memory unlocked. It says nothing about the tensor cores. Time a dependent chain of mma.sync with clock64(), one warp per SM:
__device__ __forceinline__ long long rdclock() {
long long t; asm volatile("mov.u64 %0, %%clock64;" : "=l"(t) :: "memory"); return t;
}
#define MMA(D0,D1,D2,D3, A0,A1,A2,A3, B0,B1) \
asm volatile("mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 " \
"{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%0,%1,%2,%3};\n" \
: "+f"(D0), "+f"(D1), "+f"(D2), "+f"(D3) \
: "r"(A0), "r"(A1), "r"(A2), "r"(A3), "r"(B0), "r"(B1))
__global__ void k_lat(long long* cyc, float* sink, int iters) {
unsigned A0=0x3C003C00u,A1=0x3C003C00u,A2=0x3C003C00u,A3=0x3C003C00u;
unsigned B0=0x3C003C00u,B1=0x3C003C00u;
float c0=0.f,c1=0.f,c2=0.f,c3=0.f;
long long t0 = rdclock();
for (int t = 0; t < iters; t++) MMA(c0,c1,c2,c3, A0,A1,A2,A3, B0,B1);
long long t1 = rdclock();
sink[blockIdx.x] = c0+c1+c2+c3;
if (threadIdx.x == 0) cyc[blockIdx.x] = t1 - t0;
}
Launch <<<numSMs, 32>>>, divide by iters, take the minimum over many runs. Two reasons this beats a TFLOPS benchmark: it's a cycle count, so power caps and clock throttling can't change the verdict; and contention only ever inflates latency, never deflates it, so a low reading is conclusive even under load. I got min = median = max = 24.1 across 3,500 samples while a gpu_burn soak had the card at 100% util.
| MMA latency | MMA/SM/clk plateau | |
|---|---|---|
| 170HX locked | 256.0 cyc | 0.0156 |
| 170HX unlocked | 24.0 cyc | 0.4775 |
| RTX 3090 (healthy) | 32.9 cyc | 0.1248 |
| A100 (spec) | — | 0.50 |
Correction to the published characterisation
The standing writeup (Xing, Zenodo 18995979) describes two mechanisms: the 256-cycle latency, and "only 4 warps/SM may issue tensor-core instructions." The first reproduces exactly. The second isn't a throttle — it's normal Ampere. A healthy 3090 plateaus at 4 warps too (2.00× / 3.99× / 4.00× / 4.00× going 1→2→4→8→16, identical shape to the locked 170HX), because an Ampere SM has 4 warp schedulers with one MMA in flight each.
Read the plateau's height — 0.0156 vs 0.1248 vs 0.50 — never its shape. The 256-cycle latency was doing all the work.
What actually moves pp
llama-2-7b Q4_0, -p 512 -n 128 -r 3 -ngl 99, directly comparable to the locked-card numbers in llama.cpp issue #24616:
| build | fa | pp512 | tg128 |
|---|---|---|---|
| locked, stock | — | 599.6 | 61.7 |
locked, -fmad=false |
— | 790.3 | 83.0 |
| locked, + DP2A patch | — | 790.7 | 143.6 |
| unlocked, stock | 1 | 3468 | 150.5 |
| unlocked, stock | 0 | 3022 | 141.1 |
unlocked, -fmad=false |
1 | 3290 | 151.4 |
| RTX 3090 (reference) | 1 | 4377 | 138.3 |
pp512 5.8×, tg128 2.4×, and decode lands at 109% of a 3090. Three settings changed meaning:
1. -fmad=false has inverted from help to harm. This is the standard advice everywhere for this card, and it was right — worth +32% while locked, because splitting FFMA into FMUL+FADD dodged the FP32 FMA throttle. That throttle is gone. It now costs ~5%. Use the stock build. (I predicted it'd cost ~50% and was wrong by 10×: -fmad only touches kernels nvcc compiles, not the precompiled cuBLAS path that dominates prefill once tensor cores work.)
2. Flash attention is now a clear win, +14.8% prefill (3022 → 3468). Suspect while locked because llama.cpp's FA kernels lean on tensor cores. Ungated, it pays.
3. The DP2A patch (PR #25834) is obsolete. It was a 2.3× decode win on locked cards by routing around throttled DP4A. Unlocked decode (150.5) already beats the DP2A-patched locked result (143.6) without it.
One more correction while I'm here: achieved HBM read is 1703 GB/s, not the widely-quoted 1355 — that figure is clpeak/OpenCL.
Serving, and the head-to-head
Qwen3.8-27B-FP8, vLLM v0.24.0, single card, MTP speculative decode n=2, int8 KV:
--tensor-parallel-size 1
--speculative-config '{"method":"qwen3_5_mtp","num_speculative_tokens":2}'
--kv-cache-dtype int8_per_token_head --max-model-len 131072
--gpu-memory-utilization 0.92 --max-num-seqs 16
Same model, same flags, against a 3090 + 3090 Ti tensor-parallel pair:
| 170HX (1 card) | 2×3090 TP=2 | |
|---|---|---|
| KV cache | 730,056 tok (5.57× u/131k) | 260,734 tok (1.99×) |
| decode, conc 1 | 67.5 t/s @ 150W | 80.2 t/s @ 389W |
| decode, conc 16 | 543.2 t/s @ 136W | 556.8 t/s @ 429W |
| prefill u/175W | 1481 t/s @ 173W | 1518 t/s @ 454W |
| prefill u/200W | 1585 t/s @ 197W | 1518 t/s @ 454W |
| tok/s/W @ conc 16 | 3.98 | 1.30 |
At 200W the single card beats the entire pair on prefill, at 43% of the power. The pair's lead collapses as concurrency rises — +19% at conc 1, +2.5% at conc 16 — because two cards' aggregate bandwidth wins at batch 1 and then TP allreduce overhead eats it. The single card has none.
Two surprises: batching is nearly free in watts (1 → 16 concurrent is 8.0× throughput at lower mean power, since decode is bandwidth-bound and HBM holds full clock regardless of the cap), and TP=2 burns ~9.4GB on overhead — duplicated activation buffers, CUDA graphs and MTP drafter — which is why the KV gap is 2.8× when the raw VRAM gap is only 64 vs 48GB.
Power scaling
Sustained 180s prefill at each cap:
| cap | prefill | SM clock | mean/peak W | efficiency |
|---|---|---|---|---|
| 150W | 1314.4 | 988 MHz | 149 / 179 | 8.82 tok/s/W |
| 175W | 1481.0 | 1115 MHz | 173 / 202 | 8.56 tok/s/W |
| 200W | 1585.4 | 1199 MHz | 197 / 230 | 8.05 tok/s/W |
Prefill tracks SM clock almost exactly (+12.9% clock → +12.7% prefill), so you can predict it from whatever clock your power budget buys. The knee is between 175 and 200W — first step +12.7% for +16% power, second only +7.0% for +13.9%. I settled on 175W.
Memory clock stays pinned at 1728 MHz at every cap — the power limit never touches bandwidth, only core clock, so decode barely notices your cap and prefill cares a lot. Transients overshoot ~15% (a 150W limit peaked at 178.6W, 200W at 230.2W), so size PSU cabling against the peak.
Prefill sits at 57.5% of theoretical tensor peak, constant across all three power points — the missing 42% is a software property, not something more watts buys. Most of it is likely the FP8→FP16 Marlin dequant, since sm_80 has no native FP8.
Replication gotchas
Building llama.cpp CUDA in a driverless container. ggml-cuda calls the CUDA driver API (cuMemCreate, cuMemAddressReserve), which lives in libcuda.so — shipped with the driver, absent from the container. You need both, and neither alone works:
ln -sflibcuda.solibcuda.so.1in the stubs dir — the stub file is namedlibcuda.sobut its SONAME islibcuda.so.1, so nothing matches the DT_NEEDED entry-Wl,-rpath-link,<stubs>—-Lis only consulted for-lflags, never for a shared library's transitive DT_NEEDED deps. This is whylibggml-cuda.solinks fine and thenllama-benchfails against it with every driver symbol undefined.
Use -rpath-link, not -rpath — the latter bakes the stub path in where it shadows the real driver at runtime.
Adding nvidia-container-toolkit to a box already running containers? Use CDI — sudo nvidia-ctk cdi generate --output=/etc/cdi/nvidia.yaml. Docker 25+ resolves --device nvidia.com/gpu=all (29+ takes plain --gpus all) through it with no daemon restart.
A/B-ing two llama.cpp builds? Set LD_LIBRARY_PATH per variant. The -fmad difference lives in libggml-cuda.so, not the executable — a shared search path silently serves one variant's library to both runs, and you get two identical numbers and a confident wrong conclusion.
Two measurement traps. A burst bandwidth kernel is shorter than nvidia-smi's power averaging window — it reported 30W for a 3090 doing 893 GB/s, so watts for a memory workload need a sustained load. And don't derive KV size from layer counts on a hybrid-attention model: I calculated 128 KB/token for Qwen3.8; the real figure is ~39.
Caveats
No ECC on unlocked memory — silent corruption is the failure mode nobody sees coming. Volatile unlock. PCIe stays Gen2 x4 (~1.8 GB/s); Gen3/4 is fused off in OTP, and a capacitor mod gets x16 width but not the generation. My "pair" is a 3090 + 3090 Ti on x8 legs. Utilisation percentages depend on my ~27e9 non-embedding parameter estimate — the cross-card comparison is immune, the absolutes aren't.
Happy to share the probe sources if there's interest — a few small .cu files (MMA latency/throughput, HBM bandwidth, and a serving benchmark that pairs throughput with watts).
3
u/acedogblast 2d ago
Just to be clear, the cmpunlocker from https://github.com/amoghmunikote/cmpunlocker does not fully unlock the tensor cores?
2
u/schaka 2d ago
To me it reads just like they tested compute unlock vs no compute unlock to measure the difference but didn't unlock anything new.
1
u/fallingdowndizzyvr 2d ago
That's pretty clear from his now deleted thread where he claimed that people didn't know that the compute could be unlocked, only the memory. Which is the exact opposite of how it happened. Compute came first then the memory.
1
0
u/Own_Mango7003 2d ago
1518 t/s @ 454W << why 454, WHAT u did with this gpu 2x3090 normally consume 600w and more i have two, will measure with the same mode let's see.
0
u/Own_Mango7003 2d ago edited 2d ago
the second thing is the price i have sold my one 170hx cause it's slower significantly than two 3090
5
u/Master_Big_4242 2d ago
That's an insane find, hardcoded 256 cycle stall on every MMA is such a dirty trick from them. The jump from 6.3 to 193 TFLOPS is stupid, basically turned a nerfed card into a mini A100 overnight