vllm - 💡(How to fix) Fix [Bug]: [ROCm][FLA] chunk_gated_delta_rule Triton compilation fails on MI210/gfx90a with num_stages=4

Q: Expected behavior

The vLLM FLA `chunk_gated_delta_rule` kernel should compile successfully on ROCm/gfx90a, or the autotuner should avoid configs known to fail compilation on this backend.

vllm2026-06-09 07:15:32

ON THIS PAGE

Recommended Tools

×6

Utilities matched from this issue’s tags and category — try them while you read without losing context.

GitHub issue graph ai analysis

Paste a GitHub issue URL. We fetch that issue, discover linked issues from bodies/comments/timeline, collect linked pull requests, and produce a structured English report.

The report is written in English Markdown for sharing and archival.

GitHub issue URL

Helpful · Quick feedback

Error Message

The failure happens in chunk_delta_h.py:

vllm/model_executor/layers/fla/ops/chunk_delta_h.py:177:22:
error: 'tt.load' op operation destroyed but still has uses
        b_v = tl.load(p_v, boundary_check=(0, 1)) - b_v
                     ^
LLVM ERROR: operation destroyed but still has uses

The generated MLIR reproducer pipeline includes:

tritonamdgpu-schedule-loops{num_stages=4}

The Python stack ends with:

RuntimeError: PassManager::run failed

Root Cause

I am opening this issue because I could not find an exact vLLM issue for this FLA kernel on MI210/gfx90a, and because filtering out num_stages=4 works around the failure locally.

Fix Action

Workaround

Filtering out num_stages=4 configs for chunk_gated_delta_rule_fwd_kernel_h_blockdim64 avoids the Triton AMD compiler failure locally:

kernel = chunk_delta_h.chunk_gated_delta_rule_fwd_kernel_h_blockdim64
autotuner = kernel.fn
autotuner.configs = [
    config for config in autotuner.configs
    if config.num_stages != 4
]
autotuner.cache.clear()

The benchmark prints:

patched vLLM ROCm autotune configs: 12 -> 8 (disabled num_stages=4 for chunk_delta_h)

The first benchmark case then passes:

output shapes: (1, 64, 4, 32) (1, 64, 4, 32)
state  shapes: (1, 4, 32, 32) (1, 4, 32, 32)
output max abs err: 0.0010410994291305542
state  max abs err: 0.004581227898597717
vLLM ms: {'mean': 0.4414961338043213, 'median': 0.438401997089386, 'min': 0.4179210066795349, 'max': 0.664322018623352}

Code Example

OS                           : Ubuntu 24.04.4 LTS (x86_64)
Clang version                : 22.0.0git (ROCm roc-7.2.2)
PyTorch version              : 2.10.0+rocm7.2.2.git40d237bf
ROCM used to build PyTorch   : 7.2.53211
GPU models and configuration : AMD Instinct MI210 (gfx90a:sramecc+:xnack-)
HIP runtime version          : 7.2.53211
MIOpen runtime version       : 3.5.1
vLLM Version                 : 0.16.0rc2.dev119+g31d992d21 (git sha: 31d992d21)
PYTORCH_ROCM_ARCH            : gfx90a
VLLM_TARGET_DEVICE           : rocm

### 🐛 Describe the bug

`vllm.model_executor.layers.fla.ops.chunk_delta_h.chunk_gated_delta_rule_fwd_kernel_h_blockdim64` fails during Triton AMD compilation on ROCm 7.2.2 / Triton 3.6.0 / AMD Instinct MI210 (`gfx90a`) when the autotuner tries a `num_stages=4` configuration.

The bug was initially observed when running my local benchmark script:

---

The benchmark calls vLLM's `chunk_gated_delta_rule` baseline, and the underlying failure happens in:

---

The error is:

---

This seems related to a known/similar Triton AMD backend issue around `num_stages=4`, especially:

* triton-lang/triton#9815

I am opening this issue because I could not find an exact vLLM issue for this FLA kernel on MI210/gfx90a, and because filtering out `num_stages=4` works around the failure locally.

### Reproduction

The original command that triggered the issue was:

---

The failing benchmark case is:

---

The benchmark calls:

---

with:

---

### Failing autotune configs

Before applying the workaround, the kernel has 12 autotune configs:

---

### Error

The failure happens in `chunk_delta_h.py`:

---

The generated MLIR reproducer pipeline includes:

---

The Python stack ends with:

---

### Workaround

Filtering out `num_stages=4` configs for `chunk_gated_delta_rule_fwd_kernel_h_blockdim64` avoids the Triton AMD compiler failure locally:

---

The benchmark prints:

---

The first benchmark case then passes:

---

### Expected behavior

The vLLM FLA `chunk_gated_delta_rule` kernel should compile successfully on ROCm/gfx90a, or the autotuner should avoid configs known to fail compilation on this backend.

### Actual behavior

When the autotuner tries a `num_stages=4` candidate, Triton AMD compilation fails with:

---

### Suggested fix direction

If vLLM wants to handle this locally, one possible direction is to filter `num_stages=4` for this specific FLA kernel on ROCm and affected architectures/versions, rather than removing it globally. This would keep CUDA and other ROCm targets unchanged unless they are also affected.

I am not sure whether the final fix should be in vLLM, Triton, or ROCm. I would appreciate guidance on whether vLLM wants a local workaround for this, or whether this should only be tracked in Triton. I can help test proposed fixes on MI210/gfx90a.

### Existing issues checked

I searched existing vLLM issues with the following queries but could not find an exact match:

RAW_BUFFERClick to expand / collapse

Your current environment

<details> <summary>The output of <code>python collect_env.py</code></summary>

OS                           : Ubuntu 24.04.4 LTS (x86_64)
Clang version                : 22.0.0git (ROCm roc-7.2.2)
PyTorch version              : 2.10.0+rocm7.2.2.git40d237bf
ROCM used to build PyTorch   : 7.2.53211
GPU models and configuration : AMD Instinct MI210 (gfx90a:sramecc+:xnack-)
HIP runtime version          : 7.2.53211
MIOpen runtime version       : 3.5.1
vLLM Version                 : 0.16.0rc2.dev119+g31d992d21 (git sha: 31d992d21)
PYTORCH_ROCM_ARCH            : gfx90a
VLLM_TARGET_DEVICE           : rocm

### 🐛 Describe the bug

`vllm.model_executor.layers.fla.ops.chunk_delta_h.chunk_gated_delta_rule_fwd_kernel_h_blockdim64` fails during Triton AMD compilation on ROCm 7.2.2 / Triton 3.6.0 / AMD Instinct MI210 (`gfx90a`) when the autotuner tries a `num_stages=4` configuration.

The bug was initially observed when running my local benchmark script:

```text
bench_stage1_vllm_chunk_gdn.py

The benchmark calls vLLM's chunk_gated_delta_rule baseline, and the underlying failure happens in:

vllm/model_executor/layers/fla/ops/chunk_delta_h.py

The error is:

LLVM ERROR: operation destroyed but still has uses
RuntimeError: PassManager::run failed

This seems related to a known/similar Triton AMD backend issue around num_stages=4, especially:

triton-lang/triton#9815

I am opening this issue because I could not find an exact vLLM issue for this FLA kernel on MI210/gfx90a, and because filtering out num_stages=4 works around the failure locally.

Reproduction

The original command that triggered the issue was:

cd /path/to/avelang/test/examples/linear_attention/vllm_compare
HIP_VISIBLE_DEVICES=0 python bench_stage1_vllm_chunk_gdn.py

The failing benchmark case is:

B=1
T=64
Hk=2
Hv=4
K=32
V=32
chunk_size=64
dtype=torch.bfloat16

The benchmark calls:

from vllm.model_executor.layers.fla.ops import chunk_gated_delta_rule

with:

output_final_state=True
cu_seqlens=None
head_first=False
use_qk_l2norm_in_kernel=False

Failing autotune configs

Before applying the workaround, the kernel has 12 autotune configs:

autotune configs: 12
config: {'BV': 32} num_warps= 2 num_stages= 2
config: {'BV': 64} num_warps= 2 num_stages= 2
config: {'BV': 32} num_warps= 2 num_stages= 3
config: {'BV': 64} num_warps= 2 num_stages= 3
config: {'BV': 32} num_warps= 2 num_stages= 4
config: {'BV': 64} num_warps= 2 num_stages= 4
config: {'BV': 32} num_warps= 4 num_stages= 2
config: {'BV': 64} num_warps= 4 num_stages= 2
config: {'BV': 32} num_warps= 4 num_stages= 3
config: {'BV': 64} num_warps= 4 num_stages= 3
config: {'BV': 32} num_warps= 4 num_stages= 4
config: {'BV': 64} num_warps= 4 num_stages= 4

Error

The failure happens in chunk_delta_h.py:

vllm/model_executor/layers/fla/ops/chunk_delta_h.py:177:22:
error: 'tt.load' op operation destroyed but still has uses
        b_v = tl.load(p_v, boundary_check=(0, 1)) - b_v
                     ^
LLVM ERROR: operation destroyed but still has uses

The generated MLIR reproducer pipeline includes:

tritonamdgpu-schedule-loops{num_stages=4}

The Python stack ends with:

RuntimeError: PassManager::run failed

Workaround

Filtering out num_stages=4 configs for chunk_gated_delta_rule_fwd_kernel_h_blockdim64 avoids the Triton AMD compiler failure locally:

kernel = chunk_delta_h.chunk_gated_delta_rule_fwd_kernel_h_blockdim64
autotuner = kernel.fn
autotuner.configs = [
    config for config in autotuner.configs
    if config.num_stages != 4
]
autotuner.cache.clear()

The benchmark prints:

patched vLLM ROCm autotune configs: 12 -> 8 (disabled num_stages=4 for chunk_delta_h)

The first benchmark case then passes:

output shapes: (1, 64, 4, 32) (1, 64, 4, 32)
state  shapes: (1, 4, 32, 32) (1, 4, 32, 32)
output max abs err: 0.0010410994291305542
state  max abs err: 0.004581227898597717
vLLM ms: {'mean': 0.4414961338043213, 'median': 0.438401997089386, 'min': 0.4179210066795349, 'max': 0.664322018623352}

Expected behavior

The vLLM FLA chunk_gated_delta_rule kernel should compile successfully on ROCm/gfx90a, or the autotuner should avoid configs known to fail compilation on this backend.

Actual behavior

When the autotuner tries a num_stages=4 candidate, Triton AMD compilation fails with:

LLVM ERROR: operation destroyed but still has uses
RuntimeError: PassManager::run failed

Suggested fix direction

If vLLM wants to handle this locally, one possible direction is to filter num_stages=4 for this specific FLA kernel on ROCm and affected architectures/versions, rather than removing it globally. This would keep CUDA and other ROCm targets unchanged unless they are also affected.

I am not sure whether the final fix should be in vLLM, Triton, or ROCm. I would appreciate guidance on whether vLLM wants a local workaround for this, or whether this should only be tracked in Triton. I can help test proposed fixes on MI210/gfx90a.

Existing issues checked

I searched existing vLLM issues with the following queries but could not find an exact match:

repo:vllm-project/vllm is:issue chunk_delta_h ROCm
repo:vllm-project/vllm is:issue chunk_gated_delta_rule
repo:vllm-project/vllm is:issue "operation destroyed but still has uses"
repo:vllm-project/vllm is:issue num_stages=4 ROCm
repo:vllm-project/vllm is:issue FLA gfx90a

Vote matrix · Quick signals

Works

Did the solution work? Tap to confirm.

Easy Fix

Was it a quick fix?

Time Saver

Did it save you time?

Blocking

Was it severely blocking?

Common Issue

Are others likely hitting this too?

Flaky / Intermittent

Is it intermittent?

Verified / Reproducible

Can you reproduce it reliably?

FAQ

Expected behavior

The vLLM FLA chunk_gated_delta_rule kernel should compile successfully on ROCm/gfx90a, or the autotuner should avoid configs known to fail compilation on this backend.

Still need to ship something?

×6

Another batch ranked right after the header list — different links, same matching logic.

Data

Security

Network

Code

UI/UX

Text

System

Multimedia

Protocol

API

Engineering

vllm - 💡(How to fix) Fix [Bug]: [ROCm][FLA] chunk_gated_delta_rule Triton compilation fails on MI210/gfx90a with num_stages=4

Recommended Tools

GitHub issue graph ai analysis

Error Message

Root Cause

Fix Action

Workaround

Code Example

Your current environment

Reproduction

Failing autotune configs

Error

Workaround

Expected behavior

Actual behavior

Suggested fix direction

Existing issues checked

FAQ

Expected behavior

Still need to ship something?

TRENDING