# TurboQuant on AMD Instinct MI300X & MI355X (ROCm/HIP)

## Summary

TurboQuant KV cache compression (turbo2/turbo3/turbo4) builds and runs correctly on AMD Instinct MI300X (gfx942) and MI355X (gfx950). MI300X requires zero code changes. MI355X requires adding CDNA4 arch defines to the HIP vendor header.

## Test Environment

| Component | MI300X | MI355X |
|-----------|--------|--------|
| GPU | MI300X (gfx942), 192 GB HBM3 | MI355X (gfx950), 288 GB HBM3e |
| ROCm | 7.0.2 | 7.0.1 |
| Wave Size | 64 | 64 |
| Build | `-DAMDGPU_TARGETS="gfx942"` | `-DAMDGPU_TARGETS="gfx950"` |
| Model | Qwen2.5-1.5B Q4_K_M (1.04 GiB) | same |

## WHT Kernel Correctness

Standalone roundtrip test (forward WHT → inverse WHT) confirms the Walsh-Hadamard Transform kernel works correctly on HIP with 64-wide wavefronts:

```
=== TurboQuant WHT Roundtrip Test (HIP/gfx942) ===
Total elements: 512 (4 heads x 128 dim)
Forward WHT zeros: 0 / 512
Roundtrip max error: 2.980232e-07
Roundtrip RMSE:      6.816018e-08
Result: PASS ✅
```

The kernel uses shared memory + `__syncthreads()` (no warp shuffles), so it works correctly with GCN's 64-thread wavefronts without modification.

## Performance Results

### MI300X (single GPU, Qwen2.5-1.5B Q4_K_M)

| KV Cache | pp512 (tok/s) | tg128 (tok/s) | Prefill vs f16 | Decode vs f16 |
|----------|--------------|--------------|----------------|---------------|
| f16 | 24,453 ± 230 | 181.2 ± 2.0 | baseline | baseline |
| turbo3 | ~25,200 | ~160 | **+3%** | 88% |
| turbo4 | 25,427 ± 17 | 161.1 ± 0.2 | **+4%** | 89% |

### MI355X (single GPU, Qwen2.5-1.5B Q4_K_M)

| KV Cache | pp512 (tok/s) | tg128 (tok/s) | Prefill vs f16 | Decode vs f16 |
|----------|--------------|--------------|----------------|---------------|
| f16+FA | 40,013 ± 902 | 254.5 ± 1.0 | baseline | baseline |
| turbo3 | 39,140 ± 475 | 162.3 ± 0.1 | 98% | 64% |
| turbo4 | 39,232 ± 508 | 214.1 ± 0.7 | 98% | **84%** |

### Key Observations

1. **MI300X prefill is faster with TurboQuant** (+3-4%) — less KV cache data to write to HBM.
2. **MI300X decode at 88-89% of f16** — consistent with Apple Silicon community results.
3. **MI355X turbo4 decode at 84%** — turbo4 outperforms turbo3 in decode due to simpler 4-bit dequant.
4. **MI355X turbo3 decode at 64%** — the 3-bit codebook + sign extraction is more expensive on gfx950.
5. **MI355X non-FA MMQ path crashes** (xf32 MFMA issue) — turbo types force FA and work correctly.

## Build Instructions

```bash
git clone https://github.com/TheTom/llama-cpp-turboquant.git
cd llama-cpp-turboquant
git checkout feature/turboquant-kv-cache

# MI300X (gfx942) — works without code changes
cmake -B build -DGGML_HIP=ON -DCMAKE_BUILD_TYPE=Release -DAMDGPU_TARGETS="gfx942"
cmake --build build --config Release -j

# MI355X (gfx950) — requires CDNA4 define patch (see commit)
cmake -B build -DGGML_HIP=ON -DCMAKE_BUILD_TYPE=Release -DAMDGPU_TARGETS="gfx950"
cmake --build build --config Release -j

# Test
HIP_VISIBLE_DEVICES=0 ./build/bin/llama-bench \
  -m model.gguf -ctk turbo3 -ctv turbo3 -ngl 99 -r 3 -p 512 -n 128
```

## Code Changes for gfx950 (MI355X)

Three files modified to add CDNA4 (gfx950) architecture support:

1. **`ggml/src/ggml-cuda/vendors/hip.h`** — Add `CDNA4` define for `__gfx950__`, include in `CDNA` family
2. **`ggml/src/ggml-cuda/common.cuh`** — Add `GGML_CUDA_CC_CDNA4` constant and `GGML_CUDA_CC_IS_CDNA4` macro
3. **`ggml/src/ggml-cuda/mma.cuh`** — Route CDNA4 to compatible MFMA instructions (bf16_1k, i32x16x32_i8, f32x16x4f32 — NOT xf32 which doesn't exist on gfx950)

## Known Limitations

- **MI355X non-FA MMQ crashes**: The default (non-flash-attention) matrix multiply path crashes on gfx950 due to the xf32 MFMA instruction (`mfma_f32_16x16x8_xf32`) not being available. TurboQuant types force flash attention and work correctly. Standard f16/q8_0 KV cache types need `-fa 1` flag on MI355X.
- **llama-cli text output**: Interactive mode produces empty tokens on ROCm (display issue), but `llama-bench` confirms computation is correct.

## Tested By

Andy Luo (@andyluo7)
- AMD Instinct MI300X (gfx942), ROCm 7.0.2 — April 2026
- AMD Instinct MI355X (gfx950), ROCm 7.0.1 — April 2026