Machine Learning on NVIDIA CMP

There exist NVIDIA CMP (Cryptocurrency Mining Processor) card series -- a specialised series of GPUs designed for Ethash mining. Inside, it they have full Ampere generation dice but throttled through VBIOS.

Since Etherium switched to PoS algorithm, they are nearly useless in mining -- even the major card, CMP170HX, does not bring profit above the electricity costs. This is why they are sold on near-metal prices. I bought mine 90HX (crippled 3080) for 60$ and CMP170HX (A100 dice!) for 160$.

These devices will not work out of the box for any ML application (see details below). This post summarises my current findings oh when thy can be applied and how to achieve that.

Microbenchmarking code can be found here. Llamacpp for is in progress.

TL;DR:

What is known up to now:

• Memory subsystem: throughput/latencies remain completely untouched; L1 cache size is the same per SM as for non-mining GPU, L2 is restricted for CMP90HX to 2560KB(?) from 8MB, 170HX has full 32MB L2 cache. In-block and in-SM concurrency remains untouched.
• TensorCores are slowed down ~16x and the scheduler appears restricted to ~1 warp per SM (throughput == latency, no overlap), making them useless for ML. Every mma.sync shape/precision (f16, bf16, tf32, int8) is pinned to a single fixed rate — 182 ns/op on 170HX, 285 ns/op on 90HX — i.e. a hard issue-rate cap, not a per-FLOP slowdown. TF32 matmul is throttled identically to f16, so it is not an escape hatch. See Tensor cores.
• Fused Multiply Accumulate (FMA) is thottled 8x, but can be replaced with FADD+FMUL on CMP170HX. On CMP90HX, any fp32 instructions are throttled 8x.
• Dot Product Accumulate (DP4A) is throttled ~16x on both cards. Since the MMVQ (quantized mat×vec) kernel leans on it heavily, this is the main quantized-decode bottleneck. It can be replaced by 2× PRMT + 2× DP2A (DP2A is unthrottled) with no change to the weight memory layout, which recovers ~2× decode throughput in llama.cpp. See the DP2A patch below.

What may also work

• For 90HX, every fp32 accumulate will be slow -> no matmuls at all. TF32 is not a way out either: TF32 matmul runs on the tensor cores, which are throttled identically to f16 (see Tensor cores below). One possible usage is to use Product Quantisation which effectively replaces matmuls with table lookups.

llama.cpp results

Both mining cards still severely underutilize their memory bandwidth — i.e. decode is compute-bound, not bandwidth-bound. For a 4.5 GB 9B-Q4 model one would expect ~1500 GB/s / 4.5 GB ≈ 330 tps on CMP170HX and ~170 tps on CMP90HX if decode were bandwidth-bound; the throttled DP4A path is what holds it back.

All numbers below are llama-bench -p512 -n128 -r3, in tokens/s (tg128 = decode, pp512 = prefill). master = stock __dp4a; add+mul = the portable 4×mul + 3×add fallback; dp2a = this patch (2× PRMT + 2× DP2A).

CMP 170HX (A100 die, sm_80) — Llama-2-7B-Q4_0

variant	fmad	pp512	tg128	tg vs master
master	on	599.6	61.7	1.00x
add+mul	on	599.7	97.0	1.57x
dp2a	on	599.9	124.6	2.02x
master	off	790.3	83.0	1.34x
add+mul	off	790.8	118.6	1.92x
dp2a	off	790.7	143.6	2.33x

On the 170HX, fmad is throttled but fmul/fadd are not, so compiling with --fmad=false is an extra win for both prefill and decode. Best config: dp2a + --fmad=false.

CMP 90HX (3080 die, sm_86) — Llama-2-7B-Q4_0

variant	fmad	pp512	tg128	tg vs master
master	on	374.5	57.5	1.00x
add+mul	on	378.8	92.2	1.60x
dp2a	on	378.8	114.3	1.99x
master	off	351.1	50.3	0.88x
add+mul	off	355.0	75.5	1.31x
dp2a	off	355.1	92.8	1.62x

On the 90HX, all fp32 is throttled, so --fmad=false (which doubles fp32 op count) hurts. Best config: dp2a + fmad on.

Note that prefill (pp512) barely moves across the three variants: MMQ (prefill) runs on the throttled tensor cores (IMMA), not DP4A, so the patch only helps decode. The only prefill lever here is --fmad.

Generalization across models / quants (tg128, base → dp2a best-per-card)

model / quant	CMP170HX	CMP90HX
Llama-2-7B Q4_0	61.7 → 143.6 (2.33x)	57.5 → 114.3 (1.99x)
Qwen3-8B Q4_K_M	30.7 → 76.6 (2.49x)	31.4 → 62.2 (1.98x)
Qwen3-8B Q6_K	34.1 → 64.4 (1.89x)	39.6 → 56.4 (1.42x)
Qwen3-8B IQ4_XS	57.1 → 87.4 (1.53x)	52.4 → 80.1 (1.53x)
Llama-3.2-3B Q4_K_M	66.2 → 146.6 (2.21x)	67.0 → 120.6 (1.80x)

All quant families (legacy / K / IQ) funnel through ggml_cuda_dp4a in MMVQ, so the patch is universal for quantized decode. FP16 models see no benefit (no DP4A; the FP16 path is already unthrottled).

The DP2A patch

The decode bottleneck is DP4A (__dp4a): a single instruction computing a 4-wide int8 dot product with int32 accumulation, c += A[0]·B[0] + A[1]·B[1] + A[2]·B[2] + A[3]·B[3], where A and B are 4×int8 packed into 32-bit words. It is the workhorse of the MMVQ kernel and is throttled ~16x on the 90HX and ~20x on the 170HX (4.45 / 5.68 ns vs 0.28 ns on a 2080S — see the instruction table below).

The fix uses DP2A, which is not throttled (0.31 ns). DP2A is asymmetric: it reads A as 2×int16 and B as 4×int8, and computes either dp2a.lo: c += A[0]·B[0] + A[1]·B[1], or dp2a.hi: c += A[0]·B[2] + A[1]·B[3]. So one DP4A = dp2a.lo + dp2a.hi, provided A's four bytes are first widened to int16. Crucially, only one operand needs widening — the weights in B stay 4×int8, so the memory layout and traffic are unchanged. The widening is two prmt (byte-permute with sign replication), also unthrottled:

// replace one throttled DP4A with 2× PRMT + 2× DP2A (all unthrottled)
int a_lo, a_hi;
asm("prmt.b32 %0, %1, 0, 0x9180;" : "=r"(a_lo) : "r"(a));  // {(s16)a0, (s16)a1}
asm("prmt.b32 %0, %1, 0, 0xB3A2;" : "=r"(a_hi) : "r"(a));  // {(s16)a2, (s16)a3}
int r = c;
asm("dp2a.lo.s32.s32 %0, %1, %2, %0;" : "+r"(r) : "r"(a_lo), "r"(b));
asm("dp2a.hi.s32.s32 %0, %1, %2, %0;" : "+r"(r) : "r"(a_hi), "r"(b));
return r;

Cost: 4 instructions (2 PRMT + 2 DP2A) instead of 1, all on unthrottled pipes. At the instruction level that is 4× the count of a healthy DP4A but ~3.5× faster than the throttled one (4 × 0.31 ns ≈ 1.25 ns vs 4.45 ns); end-to-end it works out to ~2× decode (tables above). The math is bit-exact integer arithmetic, so model output is identical.

It lives in ggml_cuda_dp4a (ggml/src/ggml-cuda/common.cuh) behind a -DDISABLE_DP4A compile flag, off by default; a similar inline-asm path already exists for AMD cards in the same function.

A portable C version (add+mul, the int8 4×mul + 3×add fallback) also beats master, but by less — and unreliably: nvcc often idiom-recognizes it and folds it straight back into a throttled __dp4a. The inline asm is immune to that, which is why the patch uses it.

Memory subsystem

ca - cache all cn - cache none cg - cache global (L1 cleared, L2 used)

Instructions microbenchmarking

Op	CMP90HX ns/op	2080S ns/op	CMP170HX ns/op	CMP90HX ratio	CMP170HX ratio
b32_and	0.123	0.102	0.157	1.2x	1.5x
b32_bfi	0.131	0.114	0.168	1.1x	1.5x
b32_brev	1.186	0.951	1.514	1.2x	1.6x
b32_clz	1.624	1.346	2.074	1.2x	1.5x
b32_cnot	0.122	0.102	0.157	1.2x	1.5x
b32_lop3	0.318	0.280	0.406	1.1x	1.4x
b32_not	0.123	0.102	0.157	1.2x	1.5x
b32_or	0.123	0.103	0.157	1.2x	1.5x
b32_popc	1.186	0.951	1.514	1.2x	1.6x
b32_prmt	0.314	0.280	0.400	1.1x	1.4x
b32_shl	0.453	0.406	0.578	1.1x	1.4x
b32_shr	0.123	0.103	0.157	1.2x	1.5x
b32_xor	0.123	0.103	0.157	1.2x	1.5x
dp2a_hi_ss	0.313	0.280	0.401	1.1x	1.4x
dp2a_lo_ss	0.314	0.280	0.401	1.1x	1.4x
dp4a_ss	4.450	0.280	5.681	15.9x	20.3x
dp4a_su	4.451	0.280	5.689	15.9x	20.3x
dp4a_us	4.451	0.280	5.681	15.9x	20.3x
dp4a_uu	4.451	0.280	5.682	15.9x	20.3x
f16_abs	0.123	0.103	0.157	1.2x	1.5x
f16_add	0.314	0.399	0.400	0.8x	1.0x
f16_ex2	1.186	0.951	1.514	1.2x	1.6x
f16_fma	0.314	0.399	0.400	0.8x	1.0x
f16_max	0.313	—	0.400	-	-
f16_min	0.313	—	0.400	-	-
f16_mul	0.313	0.398	0.401	0.8x	1.0x
f16_neg	0.005	0.004	0.007	1.2x	1.8x
f16_sub	0.313	0.398	0.400	0.8x	1.0x
f16_tanh	1.186	0.951	1.514	1.2x	1.6x
f16x2_abs	0.122	0.102	0.157	1.2x	1.5x
f16x2_add	0.314	0.395	0.722	0.8x	1.8x
f16x2_ex2	1.269	1.081	1.620	1.2x	1.5x
f16x2_fma	0.314	0.398	0.744	0.8x	1.9x
f16x2_max	0.313	—	0.400	-	-
f16x2_min	0.314	—	0.400	-	-
f16x2_mul	0.313	0.399	0.744	0.8x	1.9x
f16x2_neg	0.123	0.102	0.157	1.2x	1.5x
f16x2_sub	0.314	0.395	0.722	0.8x	1.8x
f16x2_tanh	1.269	1.081	1.620	1.2x	1.5x
f32_abs	0.152	0.102	0.157	1.5x	1.5x
f32_add	2.346	0.336	0.401	7.0x	1.2x
f32_cos	2.274	1.320	2.035	1.7x	1.5x
f32_div_approx	0.561	0.102	0.157	5.5x	1.5x
f32_div_rn	90.096	21.451	226.970	4.2x	10.6x
f32_ex2	1.913	1.398	2.135	1.4x	1.5x
f32_fma	2.228	0.278	5.681	8.0x	20.4x
f32_lg2	4.503	2.464	3.598	1.8x	1.5x
f32_max	0.313	0.278	0.400	1.1x	1.4x
f32_min	0.314	0.278	0.401	1.1x	1.4x
f32_mul	2.228	0.278	0.401	8.0x	1.4x
f32_neg	0.005	0.004	0.007	1.2x	1.8x
f32_rcp_approx	4.533	2.542	3.819	1.8x	1.5x
f32_rcp_rn	63.264	28.327	91.251	2.2x	3.2x
f32_rsqrt	4.503	2.464	3.598	1.8x	1.5x
f32_sin	2.274	1.320	2.035	1.7x	1.5x
f32_sqrt_approx	4.507	2.464	3.605	1.8x	1.5x
f32_sqrt_rn	73.479	18.991	90.818	3.9x	4.8x
f32_sub	2.228	0.278	0.401	8.0x	1.4x
f32_tanh	1.186	0.944	1.514	1.3x	1.6x
s32_abs	0.123	0.103	0.157	1.2x	1.5x
s32_add	0.344	0.308	0.462	1.1x	1.5x
s32_add_cc	0.344	0.308	0.462	1.1x	1.5x
s32_addcc_addc	0.821	0.730	1.049	1.1x	1.4x
s32_bfe	0.383	0.343	0.489	1.1x	1.4x
s32_bfind	1.186	0.951	1.514	1.2x	1.6x
s32_div	9.434	8.590	11.776	1.1x	1.4x
s32_mad24_lo	0.730	0.655	0.933	1.1x	1.4x
s32_mad_hi	0.670	0.592	0.866	1.1x	1.5x
s32_mad_lo	0.322	0.280	0.412	1.1x	1.5x
s32_max	0.318	0.280	0.406	1.1x	1.4x
s32_min	0.314	0.280	0.400	1.1x	1.4x
s32_mul24_lo	0.731	0.655	0.932	1.1x	1.4x
s32_mul_hi	0.999	0.908	1.276	1.1x	1.4x
s32_mul_lo	0.123	0.103	0.157	1.2x	1.5x
s32_neg	0.005	0.004	0.007	1.2x	1.8x
s32_rem	9.486	8.515	12.020	1.1x	1.4x
s32_sad	0.314	0.280	0.400	1.1x	1.4x
s32_shr	0.123	0.102	0.157	1.2x	1.5x
s32_sub	0.123	0.103	0.157	1.2x	1.5x
s32_sub_cc	0.122	0.103	0.157	1.2x	1.5x
u32_add	0.344	0.308	0.462	1.1x	1.5x
u32_bfe	0.383	0.343	0.489	1.1x	1.4x
u32_div	3.809	3.402	4.861	1.1x	1.4x
u32_mad_lo	0.322	0.280	0.412	1.1x	1.5x
u32_max	0.318	0.280	0.406	1.1x	1.4x
u32_min	0.314	0.280	0.401	1.1x	1.4x
u32_mul_hi	1.000	0.907	1.276	1.1x	1.4x
u32_mul_lo	0.123	0.102	0.157	1.2x	1.5x
u32_rem	3.704	3.309	4.728	1.1x	1.4x
u32_sub	0.123	0.102	0.157	1.2x	1.5x

Tensor cores

Dedicated mma.sync microbenchmark, single SM × 4 warps, so the per-SM tensor-core issue rate is what is measured. Numbers are ns per mma.sync instruction in throughput mode (4 independent chains); latency mode (1 dependent chain) lands within <0.1 ns of these, i.e. throughput ≈ latency.

MMA shape	operands → acc	CMP 170HX (sm_80)	CMP 90HX (sm_86)
m16n8k16	f16 → f32	181.7	284.5
m16n8k16	bf16 → f32	181.7	284.5
m16n8k8	tf32 → f32	181.7	284.5
m16n8k32	s8 → s32 (IMMA)	181.7	284.5
m16n8k32	u8 → s32 (IMMA)	181.7	284.5
m8n8k4	f64 → f64 (DMMA)	726.3	331.1

• Throughput == latency. Adding independent MMA chains buys no speedup, so the scheduler does not overlap tensor-core ops across warps — consistent with the ~1-warp-per-SM restriction.
• One fixed rate for every precision. f16, bf16, tf32 and int8 all collapse onto the same per-card number (182 ns on 170HX, 285 ns on 90HX) regardless of the K dimension or data type. That is the signature of a hard cap on the MMA issue rate (a fixed gate between issues), not a per-FLOP throttle — were it per-FLOP, int8 m16n8k32 (4× the MACs of f16 m16n8k16) could not match f16. For reference a healthy Turing 2080S runs a comparable HMMA in ~16 ns, so this is roughly 11× (170HX) / 18× (90HX) slower.
• TF32 is not an escape hatch. TF32 matmul rides the same throttled MMA path as f16/bf16.
• FP64 (DMMA) is the one inversion. The 170HX (GA100 / A100 die) has real FP64 tensor cores, but they are throttled hard (726 ns); the 90HX (GA102) has only token FP64 hardware yet is actually faster here (331 ns) — the only op where the 90HX beats the 170HX. Neither is usable for FP64 compute.

Net: there is no usable tensor-core path on either card. Quantized and FP16 ML have to stay on the scalar pipelines (FP16x2 + integer DP2A), as in the DP2A patch above.

Roadmap

1. Tensor cores microbenchmarking — done: all mma.sync precisions throttled to one fixed issue rate, ~11–18× slower; tput == lat (see Tensor cores)
2. TF32 arithmetics microbenchmarking — TF32 matmul done (throttled like f16); scalar cvt.rna.tf32.f32 throughput still TODO
3. GGUF re-packing into 2×16 bits from 4×8 bits per 32-bit block — not needed: the PRMT widening keeps the original 4×8 layout
4. Verify my llama.cpp patching — done: DP2A patch, ~2× decode, bit-exact (see above); upstream issue/PR in progress
5. Maybe patch triton to enforce compilation without fma and dp4a
6. Product Quantization