Skip to content
Merged
Show file tree
Hide file tree
Changes from 29 commits
Commits
Show all changes
33 commits
Select commit Hold shift + click to select a range
a864404
Add the option to turn on hipBLASLt online tuning
hanlin12-AMD Mar 16, 2026
96cf2d1
move hip_online_tuning option into serve.py
hanlin12-AMD Mar 17, 2026
d5753b3
use environment variable instead of CLI
hanlin12-AMD Mar 17, 2026
8136fe7
Merge branch 'main' into hip_online_tuning
hanlin12-AMD Mar 17, 2026
7abf916
fixup suffix of environment variable
hanlin12-AMD Mar 17, 2026
9c8fd55
Merge branch 'main' into hip_online_tuning
hanlin12-AMD Mar 18, 2026
838438d
Merge branch 'main' into hip_online_tuning
hanlin12-AMD Mar 18, 2026
f6c46d1
Merge branch 'main' into hip_online_tuning
hanlin12-AMD Mar 18, 2026
6443ef9
Merge branch 'vllm-project:main' into hip_online_tuning
hanlin12-AMD Apr 7, 2026
47744cf
add unit test for AITER hipBLASLt online tuning
hanlin12-AMD Apr 7, 2026
eccbced
fix typos in comment
hanlin12-AMD Apr 8, 2026
2e0a6a5
Merge branch 'vllm-project:main' into hip_online_tuning
hanlin12-AMD Apr 13, 2026
f08dd93
Merge branch 'vllm-project:main' into hip_online_tuning
hanlin12-AMD Apr 20, 2026
f33bfe5
Add AITER hipBLASLt GEMM kernel in vLLM
hanlin12-AMD Apr 21, 2026
879ccbe
Merge branch 'vllm-project:main' into hip_online_tuning
hanlin12-AMD Apr 21, 2026
7b56a01
Merge branch 'main' into hip_online_tuning
hanlin12-AMD Apr 21, 2026
ea9cef5
Update vllm/model_executor/kernels/linear/scaled_mm/aiter.py
hanlin12-AMD Apr 21, 2026
d44e5b8
Update vllm/_aiter_ops.py
hanlin12-AMD Apr 21, 2026
226e030
Remove the contiguous() after preshuffle
hanlin12-AMD Apr 21, 2026
50a8a1d
fix the env name and logic of aiter hipblaslt gemm and online tuning
hanlin12-AMD Apr 28, 2026
c4ff925
ensure VLLM_ROCM_USE_AITER_LINEAR_HIPBMM working
hanlin12-AMD Apr 28, 2026
2fdc8d0
Merge branch 'vllm-project:main' into aiter_hipbmm_online_tuning
hanlin12-AMD May 8, 2026
1ebc71e
Change the conditions of hipblaslt online tuning
hanlin12-AMD May 8, 2026
fde5878
Resolve the condition of hipBLASLt online tuning
hanlin12-AMD May 9, 2026
46bfb7b
Merge branch 'vllm-project:main' into aiter_hipbmm_online_tuning
hanlin12-AMD May 13, 2026
d001385
fix some variable name
hanlin12-AMD May 13, 2026
bdb9d33
fix missing line in aiter_ops
hanlin12-AMD May 15, 2026
98ed39b
Merge branch 'main' into aiter_hipbmm_online_tuning
tjtanaa May 27, 2026
30599ba
Merge branch 'main' into aiter_hipbmm_online_tuning
tjtanaa May 28, 2026
6b52d97
fix pre-commit
hanlin12-AMD May 28, 2026
6c0d85e
Add accuracy unit-test of Aiter hipBlaslt
hanlin12-AMD May 29, 2026
f6aed71
fix pre-commit
hanlin12-AMD May 29, 2026
df814ac
Merge branch 'main' into aiter_hipbmm_online_tuning
tjtanaa Jun 4, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
73 changes: 73 additions & 0 deletions vllm/_aiter_ops.py
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,16 @@
# on ROCm the fp8_dtype always calls is_fp8_fnuz
# which is a host op, so we cache it once here.
FP8_DTYPE = current_platform.fp8_dtype()
_HIPB_MM_INITIALIZED_DEVICES: set[int] = set()


def _ensure_hipb_mm_extension_initialized() -> None:
import aiter

device = torch.accelerator.current_device_index()
if device not in _HIPB_MM_INITIALIZED_DEVICES:
aiter.hipb_create_extension()
_HIPB_MM_INITIALIZED_DEVICES.add(device)


def is_aiter_found() -> bool:
Expand Down Expand Up @@ -625,6 +635,43 @@ def _rocm_aiter_preshuffled_per_token_w8a8_gemm_fake(
return torch.empty(m, n, dtype=output_dtype, device=A.device)


def _rocm_aiter_hipb_mm_fp8_impl(
A: torch.Tensor,
B: torch.Tensor,
As: torch.Tensor,
Bs: torch.Tensor,
bias: torch.Tensor | None = None,
output_dtype: torch.dtype = torch.bfloat16,
) -> torch.Tensor:
from aiter import hipb_mm

_ensure_hipb_mm_extension_initialized()
return hipb_mm(
A,
B,
solution_index=-1,
bias=bias,
out_dtype=output_dtype,
scaleA=As,
scaleB=Bs,
scaleOut=None,
bpreshuffle=True,
)


def _rocm_aiter_hipb_mm_fp8_fake(
A: torch.Tensor,
B: torch.Tensor,
As: torch.Tensor,
Bs: torch.Tensor,
bias: torch.Tensor | None = None,
output_dtype: torch.dtype = torch.bfloat16,
) -> torch.Tensor:
m = A.shape[0]
n = B.shape[1]
return torch.empty(m, n, dtype=output_dtype, device=A.device)


def _rocm_aiter_triton_gemm_a8w8_blockscale_impl(
A: torch.Tensor,
B: torch.Tensor,
Expand Down Expand Up @@ -1308,6 +1355,7 @@ def get_moe_dispatch_policy(cls) -> int:
# TODO: Consolidate under _LINEAR_ENABLED
_FP8BMM_ENABLED = envs.VLLM_ROCM_USE_AITER_FP8BMM
_FP4BMM_ENABLED = envs.VLLM_ROCM_USE_AITER_FP4BMM
_LINEAR_HIPBMM_ENABLED = envs.VLLM_ROCM_USE_AITER_LINEAR_HIPBMM
# TODO: Consolidate under _LINEAR_ENABLED
_FP4_GEMM_DYNAMIC_QUANT_ASM = envs.VLLM_ROCM_USE_AITER_FP4_ASM_GEMM
# TODO: Consolidate under VLLM_ROCM_USE_AITER_ROPE
Expand Down Expand Up @@ -1340,6 +1388,7 @@ def refresh_env_variables(cls):
cls._TRITON_UNIFIED_ATTN_ENABLED = envs.VLLM_ROCM_USE_AITER_UNIFIED_ATTENTION
cls._FP8BMM_ENABLED = envs.VLLM_ROCM_USE_AITER_FP8BMM
cls._FP4BMM_ENABLED = envs.VLLM_ROCM_USE_AITER_FP4BMM
cls._LINEAR_HIPBMM_ENABLED = envs.VLLM_ROCM_USE_AITER_LINEAR_HIPBMM
cls._FP4_GEMM_DYNAMIC_QUANT_ASM = envs.VLLM_ROCM_USE_AITER_FP4_ASM_GEMM
cls._TRITON_ROTARY_EMBED = envs.VLLM_ROCM_USE_AITER_TRITON_ROPE
cls._MOE_SHARED_EXPERTS_ENABLED = envs.VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS
Expand Down Expand Up @@ -1512,6 +1561,13 @@ def is_fp4bmm_enabled(cls) -> bool:

return cls._AITER_ENABLED and cls._FP4BMM_ENABLED and on_gfx950()

@classmethod
@if_aiter_supported
def is_linear_hipbmm_enabled(cls) -> bool:
from vllm.platforms.rocm import on_mi3xx

return cls.is_linear_enabled() and on_mi3xx() and cls._LINEAR_HIPBMM_ENABLED

@classmethod
@if_aiter_supported
def is_asm_fp4_gemm_dynamic_quant_enabled(cls) -> bool:
Expand Down Expand Up @@ -1668,6 +1724,12 @@ def register_ops_once() -> None:
fake_impl=_rocm_aiter_preshuffled_per_token_w8a8_gemm_fake,
)

direct_register_custom_op(
op_name="rocm_aiter_hipb_mm_fp8",
op_func=_rocm_aiter_hipb_mm_fp8_impl,
fake_impl=_rocm_aiter_hipb_mm_fp8_fake,
)

direct_register_custom_op(
op_name="rocm_aiter_triton_gemm_a8w8_blockscale",
op_func=_rocm_aiter_triton_gemm_a8w8_blockscale_impl,
Expand Down Expand Up @@ -1858,6 +1920,17 @@ def preshuffled_per_token_w8a8_gemm(
A, B, As, Bs, bias, output_dtype
)

@staticmethod
def hipb_mm_fp8(
A: torch.Tensor,
B: torch.Tensor,
As: torch.Tensor,
Bs: torch.Tensor,
bias: torch.Tensor | None = None,
output_dtype: torch.dtype = torch.bfloat16,
) -> torch.Tensor:
return torch.ops.vllm.rocm_aiter_hipb_mm_fp8(A, B, As, Bs, bias, output_dtype)

@staticmethod
def triton_gemm_a8w8_blockscale(
A: torch.Tensor,
Expand Down
5 changes: 5 additions & 0 deletions vllm/envs.py
Original file line number Diff line number Diff line change
Expand Up @@ -115,6 +115,7 @@
VLLM_ROCM_USE_AITER: bool = False
VLLM_ROCM_USE_AITER_PAGED_ATTN: bool = False
VLLM_ROCM_USE_AITER_LINEAR: bool = True
VLLM_ROCM_USE_AITER_LINEAR_HIPBMM: bool = False

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please add example documentation of where this env var can prove useful. Use cases of models or set ups that exhibit perf boost or some kind of advantage. People are not going to know how to use this env var.

@tjtanaa tjtanaa May 28, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Currently we don't have a page on vLLM documentation page that we logged down all of the aiter flags. (Let me code them up this week).
This new kernel will be kept as experimental for now as it is not enabled by default. However, it does have one good benefit over AITER's CK PTPC kernel is that this kernel can be tuned on the fly with vllm serve. The AITER CK kernels are extremely not friendly as we need to perform offline tuning and make sure to upstream to aiter before we can consume in vLLM.

VLLM_ROCM_USE_AITER_MOE: bool = True
VLLM_ROCM_AITER_MOE_DISPATCH_POLICY: int = 0
VLLM_ROCM_USE_AITER_RMSNORM: bool = True
Expand Down Expand Up @@ -1099,6 +1100,10 @@ def _resolve_rust_frontend_path() -> str | None:
"VLLM_ROCM_USE_AITER_LINEAR": lambda: (
os.getenv("VLLM_ROCM_USE_AITER_LINEAR", "True").lower() in ("true", "1")
),
"VLLM_ROCM_USE_AITER_LINEAR_HIPBMM": lambda: (
os.getenv("VLLM_ROCM_USE_AITER_LINEAR_HIPBMM", "False").lower()
in ("true", "1")
),
# Whether to use aiter moe ops.
# By default is enabled.
"VLLM_ROCM_USE_AITER_MOE": lambda: (
Expand Down
3 changes: 3 additions & 0 deletions vllm/model_executor/kernels/linear/__init__.py
Original file line number Diff line number Diff line change
Expand Up @@ -119,6 +119,7 @@
)
from vllm.model_executor.kernels.linear.scaled_mm.aiter import (
AiterFp8BlockScaledMMKernel,
AiterHipbMMPerTokenFp8ScaledMMLinearKernel,
AiterInt8ScaledMMLinearKernel,
AiterPerTokenFp8ScaledMMLinearKernel,
AiterPreshuffledPerTokenFp8ScaledMMLinearKernel,
Expand Down Expand Up @@ -272,6 +273,7 @@ def _filter_kernels_by_backend(
ChannelWiseTorchFP8ScaledMMLinearKernel,
],
PlatformEnum.ROCM: [
AiterHipbMMPerTokenFp8ScaledMMLinearKernel,
AiterPreshuffledPerTokenFp8ScaledMMLinearKernel,
AiterPerTokenFp8ScaledMMLinearKernel,
ROCmFP8ScaledMMLinearKernel,
Expand Down Expand Up @@ -1005,6 +1007,7 @@ def register_linear_kernel(
"FP8ScaledMMLinearLayerConfig",
"Int8ScaledMMLinearLayerConfig",
"ScaledMMLinearLayerConfig",
"AiterHipbMMPerTokenFp8ScaledMMLinearKernel",
"AiterPreshuffledPerTokenFp8ScaledMMLinearKernel",
"AiterPerTokenFp8ScaledMMLinearKernel",
"NvFp4LinearKernel",
Expand Down
94 changes: 94 additions & 0 deletions vllm/model_executor/kernels/linear/scaled_mm/aiter.py
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@

import torch

import vllm.envs as envs
from vllm import _custom_ops as ops
from vllm._aiter_ops import (
rocm_aiter_ops,
Expand Down Expand Up @@ -212,6 +213,99 @@
)


class AiterHipbMMPerTokenFp8ScaledMMLinearKernel(FP8ScaledMMLinearKernel):
Comment thread
hanlin12-AMD marked this conversation as resolved.
@classmethod
def is_supported(
cls, compute_capability: int | None = None
) -> tuple[bool, str | None]:
if not current_platform.is_rocm():
Comment thread
hanlin12-AMD marked this conversation as resolved.
return False, "requires ROCm."

if not rocm_aiter_ops.is_linear_hipbmm_enabled():
return (
False,
"requires setting `VLLM_ROCM_USE_AITER=1` "
"and `VLLM_ROCM_USE_AITER_LINEAR=1` "
"and `VLLM_ROCM_USE_AITER_LINEAR_HIPBMM=1`. ",
)
try:
import aiter # noqa: F401
except Exception:
return False, "requires aiter library to be installed."

if not hasattr(aiter, "hipb_mm"):
return False, "requires aiter hipb_mm support."

return True, None

@classmethod
def can_implement(cls, c: FP8ScaledMMLinearLayerConfig) -> tuple[bool, str | None]:
is_ptpc = (
c.activation_quant_key.scale.group_shape.is_per_token()
and c.weight_quant_key.scale.group_shape.is_per_channel()
)
if c.weight_shape is None:
return False, "weight_shape is required for Aiter kernels"
N, K = c.weight_shape

Check failure on line 249 in vllm/model_executor/kernels/linear/scaled_mm/aiter.py

View workflow job for this annotation

GitHub Actions / pre-commit

Ruff (F841)

vllm/model_executor/kernels/linear/scaled_mm/aiter.py:249:9: F841 Local variable `fp8_dtype` is assigned to but never used
fp8_dtype = current_platform.fp8_dtype()

if c.out_dtype is not torch.bfloat16:
return False, "requires bfloat16 output dtype."

if not is_ptpc:
return (
False,
"requires per token activation scales and per channel weight scales.",
)

if not (N >= 16 and N % 16 == 0 and K % 16 == 0):
return (
False,

Check failure on line 263 in vllm/model_executor/kernels/linear/scaled_mm/aiter.py

View workflow job for this annotation

GitHub Actions / pre-commit

Ruff (E501)

vllm/model_executor/kernels/linear/scaled_mm/aiter.py:263:89: E501 Line too long (96 > 88)
f"requires N >= 16 and both N and K divisible by 16, received N={N} and K={K}.",
)

return True, None

def process_weights_after_loading(self, layer: torch.nn.Module) -> None:
w_name, w_s_name, *_ = self.layer_param_names
w, w_s, *_ = self._get_layer_params(layer)

# Pre-apply the transposes that used to live in
# _rocm_aiter_hipb_mm_fp8_impl so the kernel can consume B/Bs directly.
# The `.t()` on the shuffled weight is kept as a non-contiguous view —
# materializing it with `.contiguous()` would re-arrange the bytes and
# break the `bpreshuffle` layout.
shuffled_w = rocm_aiter_ops.shuffle_weight(w.t().contiguous())
replace_parameter(
layer,
w_name,
torch.nn.Parameter(shuffled_w.t(), requires_grad=False),
)

if w_s.ndim > 1:
replace_parameter(
layer,
w_s_name,
torch.nn.Parameter(w_s.t().contiguous(), requires_grad=False),
)

def apply_scaled_mm(
self,
*,
A: torch.Tensor,
B: torch.Tensor,
out_dtype: torch.dtype,
As: torch.Tensor,
Bs: torch.Tensor,
bias: torch.Tensor | None,
output_shape: list,
) -> torch.Tensor:
output_shape[-1] = B.shape[1]
return rocm_aiter_ops.hipb_mm_fp8(A, B, As, Bs, bias, out_dtype).view(
*output_shape
)


class AiterPerTokenFp8ScaledMMLinearKernel(FP8ScaledMMLinearKernel):
@classmethod
def is_supported(
Expand Down
8 changes: 7 additions & 1 deletion vllm/platforms/rocm.py
Original file line number Diff line number Diff line change
Expand Up @@ -134,6 +134,8 @@
# Sync at import time - catches misconfigurations from process start.
_sync_hip_cuda_env_vars()



# AMDSMI utils
# Note that NVML is not affected by `{CUDA/HIP}_VISIBLE_DEVICES`,
# all the related functions work on real physical device ids.
Expand Down Expand Up @@ -296,7 +298,11 @@
def on_gfx950() -> bool:
return _ON_GFX950


# Enable HIP online tuning early, before hipBLASLt initializes.
# Turn on hipBLASLt online tuning if use AITER hipBLASLt GEMM.
if envs.VLLM_ROCM_USE_AITER and envs.VLLM_ROCM_USE_AITER_LINEAR and envs.VLLM_ROCM_USE_AITER_LINEAR_HIPBMM and on_mi3xx():

Check failure on line 303 in vllm/platforms/rocm.py

View workflow job for this annotation

GitHub Actions / pre-commit

Ruff (E501)

vllm/platforms/rocm.py:303:89: E501 Line too long (122 > 88)
os.environ["HIP_ONLINE_TUNING"] = "1"

@cache
def use_rocm_custom_paged_attention(
qtype: torch.dtype,
Expand Down
Loading