# Native FP8 WMMA Support for AMD RDNA4 (RX 9070 XT / R9700) in vLLM

**URL:** https://discuss.vllm.ai/t/native-fp8-wmma-support-for-amd-rdna4-rx-9070-xt-r9700-in-vllm/1900
**Category:** General
**Created:** [November 13, 2025, 12:51pm UTC](https://discuss.vllm.ai/t/native-fp8-wmma-support-for-amd-rdna4-rx-9070-xt-r9700-in-vllm/1900 "2025-11-13T12:51:16Z")
**Posts on this page:** 6
**Page:** 1

<div class="post-metadata">

### Author: ![claviger](https://avatars.discourse-cdn.com/v4/letter/c/dc4da7/32.png) [@claviger](https://discuss.vllm.ai/u/claviger)
#### Post date: [November 13, 2025, 12:51pm UTC](https://discuss.vllm.ai/t/native-fp8-wmma-support-for-amd-rdna4-rx-9070-xt-r9700-in-vllm/1900/1 "2025-11-13T12:51:16Z")

</div>

### 🚀 The feature, motivation and pitch

# Disclosure:

Yes: I used an LLM to prepare this post, obviously, who has time to type all this up…  
Yes: I used an LLM to find and patch in what was necessary.  
Yes: It works and provides significant performance uplift.  
Yes: The work is incomplete to provide full FP8 across all models, but this is a start.  
Yes: I’m bitter AF about having to take the time to do this.  
No: I will not jump through all the hoops to upstream this as a PR myself, I’ll continue to patch locally as required, I don’t have time to deal with PR/CR process.

BLUF: LLM go faster with small code changes, benchmark results:

 ![Image](https://canada1.discourse-cdn.com/flex036/uploads/vllm/original/1X/93892beac174d803a832e74432244576c02592ee.png)

 ![Image](https://canada1.discourse-cdn.com/flex036/uploads/vllm/original/1X/d6c9717c0f9a4b8442590a59c47329ade72cf3ab.png)

# Native FP8 WMMA Support for AMD RDNA4 (RX 9070 XT / R9700) in vLLM

## Summary

Successfully enabled native FP8 WMMA operations on AMD RDNA4 GPUs in vLLM, achieving significant performance improvements by utilizing the hardware’s 128 AI accelerators instead of dequantizing FP8 weights to FP32. This has been in production deployment for 4 days with zero observed aberrations as a result.

## Performance Results

Testing with FP8-quantized Qwen3 models on AMD Radeon R9700:

| Model | Before (TPS) | After (TPS) | Improvement |
| --- | --- | --- | --- |
| **Qwen3-0.6B** | ~160 decode | ~200 decode | **25% faster** |
| **Qwen3-30B-2507** | ~52 decode | ~85 decode | **63% faster** |

Additionally I found nearly **doubling of prefill performance** in certain scenarios with these rough kernel configs at prompt token counts up to 10,000 tokens with tapering gains as token count grows resulting from memory pressure. Significant additional performance remains untapped with further kernel tuning.

With further improvements this will scale well, raising memory speed to 1375 from 1258 shows a further 5% uplift in performance, indicating there is room for memory transfer optimization to make more efficient use of existing bandwidth for further gains.

* * *

## System Environment

- **Hardware** : AMD Radeon AI Pro R9700 (RDNA4 gfx1201)
  - 128 AI Accelerators
  - Native FP8 E4M3FN support
  - 16x16 WMMA instruction tile size

- **Software Stack** :
  - **vLLM Version** : 0.11.1rc6.dev223+g404d7a9d1
  - **ROCm Version** : 7.0.0
  - **AMD Driver** : 6.16.6
  - **Base Image** : rocm/vllm-dev:nightly

* * *

## The Problem

By default, vLLM on RDNA4 was **dequantizing FP8 weights to FP32** for all operations, completely wasting the hardware’s 128 AI accelerators. The execution path fell back to `torch_channelwise_w8a8_scaled_mm()` which explicitly upcasts to FP32:

```python
# Default broken path
output = torch._scaled_mm(
    qinput, weight,
    scale_a=TORCH_DEVICE_IDENTITY,
    scale_b=TORCH_DEVICE_IDENTITY,
    out_dtype=torch.float32 # ← Upcast to FP32!
)

```

This meant zero performance benefit from FP8 quantization on RDNA4.

* * *

## The Solution

Follow the **MI350X Triton kernel path** by:

1. Adding RDNA4 (gfx1201) to platform detection
2. Patching AITER’s architecture mapping to recognize gfx1201
3. Adding RDNA4-specific matrix sizes to kernel tuning
4. Providing optimized kernel configurations

**Important Note** : AITER’s C++/ASM kernels do not work on RDNA4 and must be disabled (`VLLM_ROCM_USE_AITER=0`). However, vLLM’s FP8 code path imports **AITER’s Triton kernels** (`aiter.ops.triton.gemm_a8w8_blockscale`). These Triton kernels check AITER’s architecture mapping and will crash with a `KeyError` if `gfx1201` is not recognized. Therefore, we must **patch AITER’s architecture detection** before vLLM starts, then Triton automatically compiles these kernels down to native WMMA instructions when it detects FP8 data types on gfx1201 hardware.

This routes FP8 operations through native WMMA instructions instead of dequantization.

* * *

## Implementation Details

### 1. vLLM Code Modifications

Two files need to be modified in the vLLM source code:

#### File: `vllm/platforms/rocm.py`

**Original:**

```python
def on_mi3xx() -> bool:
    GPU_ARCH = torch.cuda.get_device_properties("cuda").gcnArchName
    return any(arch in GPU_ARCH for arch in ["gfx942", "gfx950"])

```

**Modified:**

```python
def on_mi3xx() -> bool:
    GPU_ARCH = torch.cuda.get_device_properties("cuda").gcnArchName
    # Added gfx1201 (RDNA4) to enable FP8 Triton kernel path
    return any(arch in GPU_ARCH for arch in ["gfx942", "gfx950", "gfx1201"])

```

**Change** : Added `"gfx1201"` to enable RDNA4 to use the MI350X Triton kernel code path.

* * *

#### File: `vllm/model_executor/layers/quantization/utils/fp8_utils.py`

**Original:**

```python
def is_aiter_triton_kernel_tuned(n, k):
    return (n, k) in [
        (1024, 8192),
        (2112, 7168),
        (3072, 1536),
        (32768, 8192),
        (4096, 7168),
        (4608, 7168),
        (512, 7168),
        (7168, 2048),
        (7168, 256),
        (8192, 1024),
        (8192, 32768),
    ]

```

**Modified:**

```python
def is_aiter_triton_kernel_tuned(n, k):
    # MI350 tuned sizes
    mi350_sizes = [
        (1024, 8192),
        (2112, 7168),
        (3072, 1536),
        (32768, 8192),
        (4096, 7168),
        (4608, 7168),
        (512, 7168),
        (7168, 2048),
        (7168, 256),
        (8192, 1024),
        (8192, 32768),
    ]

    # RDNA4 (gfx1201) specific sizes verified to work
    rdna4_sizes = [
        (1024, 1024), # K, V projections
        (2048, 1024), # Q projection
        (3072, 1024), # Gate, Up projections
        (1024, 3072), # Down projection
        (1024, 2048), # O projection (transposed)
        (512, 512), # Small models
        (1024, 512), # Asymmetric
        (512, 1024), # Asymmetric reverse
        (2048, 2048), # Medium models
        (4096, 4096), # 7B class models
        (8192, 8192), # 70B class models
    ]

    # Check architecture to include RDNA4 sizes
    import torch
    arch_name = torch.cuda.get_device_properties(0).gcnArchName

    if "gfx12" in arch_name:
        # For RDNA4, include both RDNA4 and MI350 sizes
        return (n, k) in (rdna4_sizes + mi350_sizes)
    else:
        # Other architectures use original MI350 sizes only
        return (n, k) in mi350_sizes

```

**Changes** :

- Refactored original list into `mi350_sizes` variable
- Added 11 RDNA4-specific matrix dimensions in `rdna4_sizes`
- Added architecture detection that returns combined list for RDNA4
- The sizes used were tested and found to work correctly

* * *

### 2. AITER Architecture Patch (Required)

**Why This Is Needed:**

While AITER’s C++/ASM kernels don’t work on RDNA4 (hence `VLLM_ROCM_USE_AITER=0`), vLLM still **imports AITER’s Triton kernels** :

```python
from aiter.ops.triton.gemm_a8w8_blockscale import gemm_a8w8_blockscale

```

When these Triton kernels execute, they internally call `arch_info.get_device()` which looks up the GPU architecture in AITER’s `_ARCH_TO_DEVICE` dictionary. Since `gfx1201` is not in this dictionary by default, it throws a `KeyError` and crashes.

**The solution:** Patch AITER’s architecture mapping before vLLM starts.

> **Click to expand: AITER Patch Wrapper Script**
>
> #### `rdna4_aiter_wrapper.sh`
> 
> ```bash
> #!/bin/bash
> # RDNA4 FP8 vLLM Wrapper
> # Patches AITER architecture detection and disables AITER C++/ASM kernels
> 
> echo "=========================================="
> echo "RDNA4 FP8 Startup"
> echo "=========================================="
> 
> # Disable AITER's C++/ASM implementations (they don't work on RDNA4)
> export VLLM_ROCM_USE_AITER=0
> 
> echo "Environment: VLLM_ROCM_USE_AITER=0"
> 
> # Patch AITER's architecture mapping BEFORE vLLM imports it
> echo ""
> echo "Patching AITER architecture detection for gfx1201..."
> python3 -c "
> import aiter.ops.triton.utils.arch_info as arch_info
> if 'gfx1201' not in arch_info._ARCH_TO_DEVICE:
> arch_info._ARCH_TO_DEVICE['gfx1201'] = 'MI350X'
> print('[AITER Patch] ✓ Added gfx1201 -> MI350X mapping')
> else:
> print('[AITER Patch] gfx1201 already mapped')
> "
> 
> if [$? -eq 0]; then
> echo "✓ AITER patch applied"
> else
> echo "✗ AITER patch failed - vLLM will crash"
> exit 1
> fi
> 
> echo ""
> echo "Launching vLLM..."
> exec vllm serve "$@"
> 
> ```
> 
> **What This Does:**
> 
> 1. **Disables AITER C++/ASM kernels** via `VLLM_ROCM_USE_AITER=0`
> 2. **Patches AITER’s Triton code** by adding `gfx1201 -> MI350X` to the architecture mapping
> 3. **Launches vLLM** which can now successfully import and use AITER’s Triton kernels
> 
> **Docker Integration:**
> 
> Set this script as your Docker ENTRYPOINT:
> 
> ```dockerfile
> COPY rdna4_aiter_wrapper.sh /workspace/
> RUN chmod +x /workspace/rdna4_aiter_wrapper.sh
> ENTRYPOINT ["/workspace/rdna4_aiter_wrapper.sh"]
> 
> ```
> 
> Or mount and use it via systemd:
> 
> ```systemd
> -v /path/to/rdna4_aiter_wrapper.sh:/workspace/wrapper.sh \
> --entrypoint /workspace/wrapper.sh \
> 
> ```

**Alternative: Inline Patch**

If you don’t want a wrapper script, add this to your Docker ENTRYPOINT or startup command:

```bash
python3 -c "import aiter.ops.triton.utils.arch_info as arch_info; arch_info._ARCH_TO_DEVICE['gfx1201'] = 'MI350X'" && \
export VLLM_ROCM_USE_AITER=0 && \
vllm serve "$@"

```

* * *

### 3. Kernel Configuration Files

16 JSON configuration files are required to optimize FP8 operations for RDNA4. These files tell vLLM’s Triton compiler how to tile and execute FP8 matrix multiplications. These an educated GUESS but they do result in massive performance uplift, there is significant room for improving throughput by correctly tuning these values.

**File Locations** :

- Linear layers: `vllm/model_executor/layers/quantization/utils/configs/`
- MoE layers: `vllm/model_executor/layers/fused_moe/configs/`

**File Naming** : `N={n},K={k},device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`

- `device_name=0x7551` = AMD Radeon RX 9070 XT device ID
- `dtype=fp8_w8a8` = FP8 weights and activations
- `block_shape=[128,128]` = Quantization block size

* * *

## Kernel Configuration Files

### Linear Layer Configs (15 files)

All linear configs use this structure with batch size keys (“16”, “32”, “64”):

> **Click to expand all 15 linear config files**
>
> #### `N=1024,K=1024,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=1024,K=1536,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=1024,K=2048,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=1024,K=3072,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=2048,K=1024,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=2048,K=2048,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=2048,K=4096,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=2048,K=768,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=2560,K=2048,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=256,K=2048,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=3072,K=1024,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=384,K=2048,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=4096,K=2048,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=512,K=2048,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```
> 
> #### `N=768,K=2048,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`
> 
> ```json
> {
> "16": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "32": {
> "BLOCK_SIZE_K": 32,
> "BLOCK_SIZE_M": 32,
> "BLOCK_SIZE_N": 32,
> "GROUP_SIZE_M": 16,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> },
> "64": {
> "BLOCK_SIZE_K": 64,
> "BLOCK_SIZE_M": 64,
> "BLOCK_SIZE_N": 64,
> "GROUP_SIZE_M": 8,
> "kpack": 1,
> "matrix_instr_nonkdim": 16,
> "num_warps": 4
> }
> }
> 
> ```

* * *

### MoE Layer Config (1 file)

MoE configs use `num_stages` instead of `kpack` and `matrix_instr_nonkdim`:

#### `E=128,N=384,device_name=0x7551,dtype=fp8_w8a8,block_shape=[128,128].json`

```json
{
    "16": {
        "BLOCK_SIZE_M": 32,
        "BLOCK_SIZE_N": 32,
        "BLOCK_SIZE_K": 32,
        "GROUP_SIZE_M": 8,
        "num_warps": 4,
        "num_stages": 2
    },
    "32": {
        "BLOCK_SIZE_M": 32,
        "BLOCK_SIZE_N": 32,
        "BLOCK_SIZE_K": 32,
        "GROUP_SIZE_M": 16,
        "num_warps": 4,
        "num_stages": 2
    },
    "64": {
        "BLOCK_SIZE_M": 64,
        "BLOCK_SIZE_N": 64,
        "BLOCK_SIZE_K": 64,
        "GROUP_SIZE_M": 8,
        "num_warps": 4,
        "num_stages": 2
    }
}

```

* * *

## Configuration Parameters Explained

| Parameter | Description | RDNA4 Values |
| --- | --- | --- |
| `BLOCK_SIZE_M/N/K` | Tile size for GEMM operations | 32, 64 (multiples of 16) |
| `GROUP_SIZE_M` | Batch grouping for processing | 8, 16 |
| `num_warps` | Warps per thread block | 4 (aligns with CU SIMDs) |
| `matrix_instr_nonkdim` | WMMA instruction size | 16 (fixed for RDNA4) |
| `kpack` | K-dimension packing factor | 1 (conservative) |
| `num_stages` | Pipeline stages (MoE only) | 2 |

**Notes** :

- All dimensions must be multiples of 16 (WMMA tile size)
- Configs are conservative; significant tuning potential remains
- `num_warps=4` aligns well with RDNA4’s compute unit structure
- Larger block sizes (128x128) could be explored for batch operations but caused performance instability/inconsistency during prefill so were removed to fall back to default values

* * *

## Key Technical Details

### How Triton Automatically Compiles to WMMA

The magic happens because **Triton automatically detects and uses WMMA instructions** when all these conditions are met:

1. **GPU Architecture** : gfx1201 (RDNA4) is detected
2. **Data Types** : FP8 (float8\_e4m3fn) input tensors
3. **Operation** : Matrix multiplication via `tl.dot()`
4. **Tile Sizes** : Dimensions are multiples of 16 (WMMA tile size)

When vLLM imports `aiter.ops.triton.gemm_a8w8_blockscale`, it’s just importing Triton code. The Triton compiler examines the target architecture, sees FP8 operations on gfx1201, and automatically generates assembly code using RDNA4’s 16x16x16 WMMA instructions (`v_wmma_f32_16x16x16_fp8_fp8`).

**No runtime patching or AITER library support is needed** - it’s pure Triton compilation doing the heavy lifting.

### Execution Path

```auto
FP8 Model Load
    ↓
on_mi3xx() returns True (gfx1201 recognized)
    ↓
rocm_aiter_gemm_w8a8_blockscale_impl() [fp8_utils.py:71]
    ↓
is_aiter_triton_kernel_tuned(n, k) returns True
    ↓
Imports aiter.ops.triton.gemm_a8w8_blockscale (Triton code)
    ↓
Triton compiler detects gfx1201 + FP8 types
    ↓
Auto-generates WMMA instructions (v_wmma_f32_16x16x16_fp8_fp8)
    ↓
✓ Native FP8 WMMA execution!

```

* * *

## Verification

### Check Platform Detection

```python
from vllm.platforms import rocm
import torch
print(f"GPU: {torch.cuda.get_device_properties(0).gcnArchName}") # Should be gfx1201
print(f"Detected as MI3xx: {rocm.on_mi3xx()}") # Should be True

```

### Check AITER is Disabled

```bash
echo $VLLM_ROCM_USE_AITER # Should output: 0

```

### Look for Log Messages

```auto
✅ GOOD: "Using configuration from .../N=2048,K=1024,device_name=0x7551,dtype=fp8_w8a8..."
❌ BAD: "Using default W8A8 Block FP8 kernel config"

```

* * *

## Known Limitations

1. **Minimum Batch Size** : RDNA4 requires M ≥ 16 for FP8 operations. Single-token generation may need padding, it showed failures in testing.
2. **AITER Runtime Patch Required** : AITER’s C++/ASM kernels don’t work on RDNA4, but we still need AITER’s Triton kernels. This requires patching AITER’s architecture mapping at runtime before vLLM starts, otherwise you’ll get a `KeyError: 'gfx1201'` crash.
3. **Kernel Configs** : Current configs are conservative. Significant performance improvements possible with tuning.
4. **Model Coverage** : Configs cover common LLM dimensions. New models may need additional configs.

* * *

## Future Optimization Opportunities

1. **Kernel Tuning** :

2. **Additional Matrix Sizes** :

3. **Batch Size Handling** :

* * *

## Contributing

This work enables native FP8 support on RDNA4 for the first time in vLLM. If you use these modifications:

1. **Test other models** : Try different FP8 models and report results
2. **Tune configs** : Experiment with kernel parameters and share improvements
3. **Add matrix sizes** : Contribute configs for models not yet covered
4. **Upstream** : Help prepare patches for vLLM mainline

* * *

## Acknowledgments

This implementation follows the MI350X Triton kernel path and builds upon AMD’s AITER library and vLLM’s FP8 quantization framework.

**Key insight** : RDNA4 uses the same standard FP8 E4M3FN format as MI350X (not FNUZ like MI300), making it compatible with MI350X’s Triton kernels after proper platform detection.

* * *

**Status** : Production-ready with ongoing optimization opportunities.

**Last Updated** : 2025-11-13

---

<div class="post-metadata">

### Author: ![RunLLM](https://yyz2.discourse-cdn.com/flex036/user_avatar/discuss.vllm.ai/runllm/32/233_2.png) [@RunLLM](https://discuss.vllm.ai/u/RunLLM)
#### Post date: [November 13, 2025, 12:52pm UTC](https://discuss.vllm.ai/t/native-fp8-wmma-support-for-amd-rdna4-rx-9070-xt-r9700-in-vllm/1900/2 "2025-11-13T12:52:07Z")

</div>

Your summary and implementation approach for enabling native FP8 WMMA support on AMD RDNA4 (gfx1201) in vLLM are correct and well-aligned with vLLM’s quantization and kernel selection logic. The key points—modifying platform detection to include gfx1201, disabling AITER, and providing tuned Triton kernel configs—are all supported by vLLM’s documented architecture and recent kernel integration practices. Triton’s auto-compilation to WMMA for FP8 on RDNA4 is also accurate, as it leverages the hardware’s native instructions when the correct data types and tile sizes are used. Your kernel config structure and parameter choices (BLOCK\_SIZE\_M/N/K, GROUP\_SIZE\_M, num\_warps, etc.) are consistent with vLLM’s tuning methodology for new hardware backends, and the need for further tuning is expected and encouraged by the vLLM community.

The only caveats are: (1) AITER must remain disabled on RDNA4, as it is not compatible, and (2) minimum batch size (M ≥ 16) is a hardware constraint for WMMA, so single-token or very small batch inference may require padding or fallback. Your approach to config file naming and placement matches vLLM’s config lookup logic, and your method for extending is\_aiter\_triton\_kernel\_tuned() is in line with how vLLM supports new architectures. For further optimization, you are encouraged to contribute your configs and findings upstream, as vLLM’s maintainers actively seek hardware-specific tuning contributions. For more details, see [vLLM Quantization Docs](https://docs.vllm.ai/en/latest/features/quantization/fp8/), [Kernel Tuning Example](https://github.com/vllm-project/vllm/pull/25656), and [Platform Detection](https://github.com/vllm-project/vllm/issues/26700).

Would you like a step-by-step breakdown of the config file structure, kernel selection logic, or further guidance on contributing these changes upstream?

Sources:

- [vLLM Quantization Docs](https://docs.vllm.ai/en/latest/features/quantization/fp8/)
- [Kernel Tuning Example](https://github.com/vllm-project/vllm/pull/25656)
- [Platform Detection](https://github.com/vllm-project/vllm/issues/26700)

Hint: Mention @RunLLM in the post for followups.

---

<div class="post-metadata">

### Author: ![jmander11](https://yyz2.discourse-cdn.com/flex036/user_avatar/discuss.vllm.ai/jmander11/32/1146_2.png) [@jmander11](https://discuss.vllm.ai/u/jmander11)
#### Post date: [February 12, 2026, 5:10am UTC](https://discuss.vllm.ai/t/native-fp8-wmma-support-for-amd-rdna4-rx-9070-xt-r9700-in-vllm/1900/3 "2026-02-12T05:10:03Z")

</div>

@claviger The code is looking different now in fp8\_utils.py. I am trying to get fp8 working on AMD R9700 and wondering how I should proceed. In particular, I need [[ROCm][Quantization] GPT\_OSS in amd-quark format model loading and emulations by xuebwang-amd · Pull Request #29008 · vllm-project/vllm](https://github.com/vllm-project/vllm/pull/29008) and [[Kernel][MoE] fix computation order of MoE weight multiplication and improve flow by xuebwang-amd · Pull Request #31962 · vllm-project/vllm](https://github.com/vllm-project/vllm/pull/31962) if I want to run amd/gpt-oss-20b fp8 version of the model and merging the first pull request updates my vllm to 0.15. I realize this is a pretty niche question, but this model appears to be the sweet spot for an fp8 model with a large context window still that could fit the R9700. What fp8 models have you run?

@RunLLM Is RDNA4 fp8 support tracked in a pull request or already merged into vllm?

---

<div class="post-metadata">

### Author: ![RunLLM](https://yyz2.discourse-cdn.com/flex036/user_avatar/discuss.vllm.ai/runllm/32/233_2.png) [@RunLLM](https://discuss.vllm.ai/u/RunLLM)
#### Post date: [February 12, 2026, 5:10am UTC](https://discuss.vllm.ai/t/native-fp8-wmma-support-for-amd-rdna4-rx-9070-xt-r9700-in-vllm/1900/4 "2026-02-12T05:10:49Z")

</div>

RDNA4 (gfx1201) native FP8 support is not yet fully merged into vLLM mainline; it requires local patches and kernel config files as described in community discussions. The key PRs for running amd/gpt-oss-20b in FP8 on AMD (including R9700) are [#29008](https://github.com/vllm-project/vllm/pull/29008) (amd-quark format/model loading and emulation) and [#31962](https://github.com/vllm-project/vllm/pull/31962) (MoE computation order fix)—both are merged, but you may need to use vLLM 0.15+ and ensure your local repo includes these changes. For RDNA4, you must also add kernel configs and ensure AITER is disabled (VLLM\_ROCM\_USE\_AITER=0), as upstream support for RDNA4-specific WMMA is not yet in main.

FP8 models successfully run by the community on AMD include Qwen3, Mixtral, and GPT-OSS variants, with Qwen3-30B and GPT-OSS-20B cited as working well on R9700 after applying the above patches and configs. There is no single PR tracking full RDNA4 FP8 support; progress is documented in issues like [#28649](https://github.com/vllm-project/vllm/issues/28649) and related discussions. For step-by-step details, see the implementation notes and config requirements in those issues and PRs.

Would you like a detailed step-by-step guide for patching and running FP8 models on R9700?

Sources:

- [#29008 ROCm Quantization GPT\_OSS in amd-quark format](https://github.com/vllm-project/vllm/pull/29008)
- [#31962 Kernel MoE computation order fix](https://github.com/vllm-project/vllm/pull/31962)
- [#28649 RDNA4 FP8 Patch Issue](https://github.com/vllm-project/vllm/issues/28649)

Hint: Mention @RunLLM in the post for followups.

---

<div class="post-metadata">

### Author: ![jmander11](https://yyz2.discourse-cdn.com/flex036/user_avatar/discuss.vllm.ai/jmander11/32/1146_2.png) [@jmander11](https://discuss.vllm.ai/u/jmander11)
#### Post date: [February 12, 2026, 5:21am UTC](https://discuss.vllm.ai/t/native-fp8-wmma-support-for-amd-rdna4-rx-9070-xt-r9700-in-vllm/1900/5 "2026-02-12T05:21:28Z")

</div>

@RunLLM A detailed guide for patching vLLM to run FP8 models on R9700 would be awesome. The latest AMD docker image is rocm/vllm-dev:rocm7.2\_navi\_ubuntu24.04\_py3.12\_pytorch\_2.9\_vllm\_0.14.0rc0, but will a more recent version of vLLM with the PRs 29008 and 32962 have steps to add FP8 support or are FP8 support steps only available for a vLLM version more in line with the date of the original post here?

---

<div class="post-metadata">

### Author: ![RunLLM](https://yyz2.discourse-cdn.com/flex036/user_avatar/discuss.vllm.ai/runllm/32/233_2.png) [@RunLLM](https://discuss.vllm.ai/u/RunLLM)
#### Post date: [February 12, 2026, 5:22am UTC](https://discuss.vllm.ai/t/native-fp8-wmma-support-for-amd-rdna4-rx-9070-xt-r9700-in-vllm/1900/6 "2026-02-12T05:22:46Z")

</div>

FP8 support for AMD RDNA4 (R9700) is not fully upstreamed in vLLM mainline as of v0.14.0rc0 or even v0.15.0; it requires both the merged PRs ([#29008](https://github.com/vllm-project/vllm/pull/29008), [#31962](https://github.com/vllm-project/vllm/pull/31962)) and additional manual steps. The official AMD Docker images do not yet include all the necessary RDNA4-specific kernel configs or platform detection changes for native FP8 WMMA. You must patch your vLLM source and add kernel config files as described in [#28649](https://github.com/vllm-project/vllm/issues/28649).

Here’s a step-by-step guide:

1. **Clone vLLM and Checkout/Cherry-pick PRs**

2. **Patch Platform Detection for RDNA4**

3. **Add/Modify Kernel Config Files**

4. **Disable AITER**

5. **(Optional) Patch Kernel Tuning Logic**

6. **Build and Run**

7. **Verify**

**Note:**

- These steps are required regardless of the vLLM version, as full RDNA4 FP8 support is not yet in mainline or official Docker images.
- You may need to tune or add kernel configs for new models or batch sizes.

Would you like a sample Dockerfile or example kernel config file?

Sources:

- [#29008 ROCm Quantization GPT\_OSS in amd-quark format](https://github.com/vllm-project/vllm/pull/29008)
- [#28649 RDNA4 FP8 Patch Issue](https://github.com/vllm-project/vllm/issues/28649)

Hint: Mention @RunLLM in the post for followups.