| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | import functools |
| 4 | from collections.abc import Iterable |
| 5 | from math import prod |
| 6 | from typing import TYPE_CHECKING |
| 7 | |
| 8 | import torch |
| 9 | import torch.nn.functional as F |
| 10 | from torch import nn |
| 11 | |
| 12 | import vllm.envs as envs |
| 13 | from vllm import _custom_ops as ops |
| 14 | from vllm._aiter_ops import rocm_aiter_ops |
| 15 | from vllm.logger import init_logger |
| 16 | from vllm.model_executor.layers.quantization.utils.config_utils import ( |
| 17 | is_shared_expert_quant_fse_compatible, |
| 18 | ) |
| 19 | from vllm.model_executor.layers.quantization.utils.fp8_utils import ( |
| 20 | per_token_group_quant_fp8, |
| 21 | ) |
| 22 | from vllm.model_executor.layers.quantization.utils.int8_utils import ( |
| 23 | per_token_group_quant_int8, |
| 24 | per_token_quant_int8, |
| 25 | ) |
| 26 | from vllm.model_executor.layers.quantization.utils.mxfp4_utils import ( |
| 27 | quant_dequant_mxfp4, |
| 28 | xpu_mxfp4_quantize, |
| 29 | ) |
| 30 | from vllm.model_executor.layers.quantization.utils.mxfp6_utils import ( |
| 31 | quant_dequant_mxfp6, |
| 32 | ) |
| 33 | from vllm.model_executor.layers.quantization.utils.mxfp8_utils import ( |
| 34 | mxfp8_e4m3_quantize, |
| 35 | xpu_mxfp8_quantize, |
| 36 | ) |
| 37 | from vllm.model_executor.layers.quantization.utils.nvfp4_emulation_utils import ( |
| 38 | ref_nvfp4_quant_dequant, |
| 39 | ) |
| 40 | from vllm.model_executor.layers.quantization.utils.w8a8_utils import ( |
| 41 | per_tensor_dequantize, |
| 42 | ) |
| 43 | from vllm.model_executor.models.utils import PPMissingLayer |
| 44 | from vllm.platforms import current_platform |
| 45 | from vllm.triton_utils import tl, triton |
| 46 | from vllm.utils.math_utils import cdiv |
| 47 | |
| 48 | if TYPE_CHECKING: |
| 49 | from vllm.model_executor.layers.fused_moe.config import FusedMoEConfig |
| 50 | from vllm.model_executor.layers.quantization import QuantizationConfig |
| 51 | |
| 52 | logger = init_logger(__name__) |
| 53 | |
| 54 | |
| 55 | def resolve_layer_fused_shared_expert( |
| 56 | quant_config: "QuantizationConfig | None", |
| 57 | prefix: str, |
| 58 | shared_expert_name: str = "shared_experts", |
| 59 | ) -> bool: |
| 60 | """Resolve whether AITER fused shared-expert execution is enabled. |
| 61 | |
| 62 | Args: |
| 63 | quant_config: Model quantization configuration. |
| 64 | prefix: MoE module prefix. |
| 65 | shared_expert_name: Shared-expert module name under ``prefix``. |
| 66 | |
| 67 | Returns: |
| 68 | Whether AITER fused shared experts are enabled. |
| 69 | |
| 70 | Raises: |
| 71 | ValueError: If requested shared-expert fusion is quantization-incompatible. |
| 72 | |
| 73 | """ |
| 74 | # NOTE: is_fusion_moe_shared_experts_enabled is decorated with @if_aiter_supported |
| 75 | # that returns None if AITER is not available. |
| 76 | fse_requested = bool(rocm_aiter_ops.is_fusion_moe_shared_experts_enabled()) |
| 77 | fse_compatible, fse_reason = ( |
| 78 | is_shared_expert_quant_fse_compatible( |
| 79 | quant_config, |
| 80 | f"{prefix}.experts", |
| 81 | f"{prefix}.{shared_expert_name}", |
| 82 | ) |
| 83 | if fse_requested |
| 84 | else (True, None) |
| 85 | ) |
| 86 | is_fused_shared_expert_enabled = fse_requested and fse_compatible |
| 87 | if fse_requested and not is_fused_shared_expert_enabled: |
| 88 | logger.warning( |
| 89 | "VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS is enabled but " |
| 90 | "cannot be enabled - skipping for this layer: %s.", |
| 91 | fse_reason, |
| 92 | ) |
| 93 | return is_fused_shared_expert_enabled |
| 94 | |
| 95 | |
| 96 | def is_model_fused_shared_expert_compatible( |
| 97 | layers: nn.ModuleList | Iterable[nn.Module], |
| 98 | moe_cls: type[nn.Module], |
| 99 | moe_name: str, |
| 100 | ) -> bool: |
| 101 | """Resolve one fused-shared-expert state for a model's MoE layers.""" |
| 102 | |
| 103 | def get_moe_layer(layer: nn.Module) -> nn.Module | None: |
| 104 | for name in moe_name.split("."): |
| 105 | layer = getattr(layer, name, None) |
| 106 | if layer is None: |
| 107 | return None |
| 108 | return layer |
| 109 | |
| 110 | moe_layers = ( |
| 111 | moe_layer |
| 112 | for layer in layers |
| 113 | if not isinstance(layer, PPMissingLayer) |
| 114 | and (moe_layer := get_moe_layer(layer)) is not None |
| 115 | and isinstance(moe_layer, moe_cls) |
| 116 | ) |
| 117 | |
| 118 | enabled = [ |
| 119 | getattr(layer, "is_fused_shared_expert_enabled", False) for layer in moe_layers |
| 120 | ] |
| 121 | enabled_count = sum(enabled) |
| 122 | disabled_count = len(enabled) - enabled_count |
| 123 | if enabled_count > 0 and disabled_count > 0: |
| 124 | raise NotImplementedError( |
| 125 | "Fused shared experts must be enabled for all MoE layers; found " |
| 126 | f"{enabled_count} enabled and {disabled_count} disabled layers. " |
| 127 | "Per-layer fused shared experts is not yet supported. Please open " |
| 128 | "an issue." |
| 129 | ) |
| 130 | return enabled_count > 0 and disabled_count == 0 |
| 131 | |
| 132 | |
| 133 | @triton.jit |
| 134 | def _count_expert_num_tokens( |
| 135 | topk_ids_ptr, |
| 136 | expert_num_tokens_ptr, |
| 137 | num_experts, |
| 138 | topk_numel, |
| 139 | expert_map, |
| 140 | HAS_EXPERT_MAP: tl.constexpr, |
| 141 | BLOCK_SIZE: tl.constexpr, |
| 142 | ): |
| 143 | curr_expert = tl.program_id(0) |
| 144 | |
| 145 | offsets = tl.arange(0, BLOCK_SIZE) |
| 146 | topk_ids_ptrs = topk_ids_ptr + offsets |
| 147 | |
| 148 | acc = tl.zeros((BLOCK_SIZE,), dtype=tl.int32) |
| 149 | for x in range(tl.cdiv(topk_numel, BLOCK_SIZE)): |
| 150 | mask = offsets < (topk_numel - x * BLOCK_SIZE) |
| 151 | expert_ids = tl.load(topk_ids_ptrs, mask=mask, other=-1) |
| 152 | if HAS_EXPERT_MAP: |
| 153 | expert_map_ptrs = expert_map + expert_ids |
| 154 | expert_map_mask = expert_ids >= 0 |
| 155 | expert_ids = tl.load(expert_map_ptrs, mask=expert_map_mask, other=-1) |
| 156 | |
| 157 | has_curr_expert = tl.where(expert_ids == curr_expert, 1, 0) |
| 158 | acc = acc + has_curr_expert |
| 159 | topk_ids_ptrs += BLOCK_SIZE |
| 160 | |
| 161 | if curr_expert < num_experts: |
| 162 | tl.store(expert_num_tokens_ptr + curr_expert, tl.sum(acc)) |
| 163 | |
| 164 | |
| 165 | def count_expert_num_tokens( |
| 166 | topk_ids: torch.Tensor, num_local_experts: int, expert_map: torch.Tensor | None |
| 167 | ) -> torch.Tensor: |
| 168 | """Count the number to tokens assigned to each expert. |
| 169 | |
| 170 | Args: |
| 171 | topk_ids (torch.Tensor): Tensor mapping each token to its list of experts. |
| 172 | num_local_experts (int): Number of experts in this rank. |
| 173 | expert_map (Optional[torch.Tensor]): A tensor mapping expert indices |
| 174 | from the global expert space to the local expert space of the expert |
| 175 | parallel shard. |
| 176 | |
| 177 | Returns: |
| 178 | A tensor of size num_local_experts, where tensor[i] holds the number |
| 179 | of tokens assigned to the ith expert. |
| 180 | |
| 181 | """ |
| 182 | assert topk_ids.dtype.is_signed, "The kernel uses -1 to represent invalid topk_ids" |
| 183 | expert_num_tokens = torch.empty( |
| 184 | (num_local_experts), device=topk_ids.device, dtype=torch.int32 |
| 185 | ) |
| 186 | |
| 187 | grid = num_local_experts |
| 188 | BLOCK_SIZE = min(topk_ids.numel(), 1024) |
| 189 | BLOCK_SIZE = triton.next_power_of_2(BLOCK_SIZE) |
| 190 | |
| 191 | _count_expert_num_tokens[(grid,)]( |
| 192 | topk_ids, |
| 193 | expert_num_tokens, |
| 194 | num_local_experts, |
| 195 | topk_ids.numel(), |
| 196 | expert_map, |
| 197 | HAS_EXPERT_MAP=expert_map is not None, |
| 198 | BLOCK_SIZE=BLOCK_SIZE, |
| 199 | ) |
| 200 | |
| 201 | return expert_num_tokens |
| 202 | |
| 203 | |
| 204 | def _resize_cache(x: torch.Tensor, v: tuple[int, ...]) -> torch.Tensor: |
| 205 | """Shrink the given tensor and apply the given view to it. This is |
| 206 | used to resize the intermediate fused_moe caches. |
| 207 | """ |
| 208 | assert prod(v) <= x.numel(), ( |
| 209 | f"{v} ({prod(v)}) <= {x.shape} ({x.numel()})" |
| 210 | ) # CUDAGRAPH unfriendly? |
| 211 | return x.flatten()[: prod(v)].view(*v) |
| 212 | |
| 213 | |
| 214 | def _nvfp4_quantize( |
| 215 | A: torch.Tensor, |
| 216 | A_scale: torch.Tensor | None, |
| 217 | is_sf_swizzled_layout: bool, |
| 218 | ) -> tuple[torch.Tensor, torch.Tensor]: |
| 219 | return ops.scaled_fp4_quant(A, A_scale, is_sf_swizzled_layout=is_sf_swizzled_layout) |
| 220 | |
| 221 | |
| 222 | def _fp8_quantize( |
| 223 | A: torch.Tensor, |
| 224 | A_scale: torch.Tensor | None, |
| 225 | per_act_token: bool, |
| 226 | block_shape: list[int] | None = None, |
| 227 | ) -> tuple[torch.Tensor, torch.Tensor]: |
| 228 | """Perform fp8 quantization on the inputs. If a block_shape |
| 229 | is provided, the output will be blocked. |
| 230 | """ |
| 231 | if block_shape is None: |
| 232 | # TODO(luka): use QuantFP8 custom op |
| 233 | # https://github.com/vllm-project/vllm/issues/20711 |
| 234 | A, A_scale = ops.scaled_fp8_quant( |
| 235 | A, A_scale, use_per_token_if_dynamic=per_act_token |
| 236 | ) |
| 237 | else: |
| 238 | assert not per_act_token |
| 239 | assert len(block_shape) == 2 |
| 240 | _, block_k = block_shape[0], block_shape[1] |
| 241 | A, A_scale = per_token_group_quant_fp8(A, block_k) |
| 242 | assert cdiv(A.size(-1), block_k) == A_scale.size(-1) |
| 243 | |
| 244 | return A, A_scale |
| 245 | |
| 246 | |
| 247 | def _int8_quantize( |
| 248 | A: torch.Tensor, |
| 249 | A_scale: torch.Tensor | None, |
| 250 | per_act_token: bool, |
| 251 | block_shape: list[int] | None = None, |
| 252 | ) -> tuple[torch.Tensor, torch.Tensor]: |
| 253 | """Perform int8 quantization on the inputs. If a block_shape |
| 254 | is provided, the output will be blocked. |
| 255 | """ |
| 256 | # If weights are per-channel (per_channel_quant=True), then |
| 257 | # activations apply per-token quantization. Otherwise, assume |
| 258 | # activation tensor-wise fp8/int8 quantization, dynamic or static |
| 259 | if block_shape is None: |
| 260 | if per_act_token: |
| 261 | A, A_scale = per_token_quant_int8(A) |
| 262 | elif A_scale is not None: |
| 263 | # Static per-tensor: use the optimized CUDA kernel |
| 264 | A, A_scale, _ = ops.scaled_int8_quant(A, scale=A_scale) |
| 265 | elif A_scale is None: |
| 266 | # Dynamic per-tensor: compute scale then quantize via kernel |
| 267 | A_scale = torch.clamp(A.abs().max() / 127.0, min=1e-10) |
| 268 | A, A_scale, _ = ops.scaled_int8_quant(A, scale=A_scale) |
| 269 | else: |
| 270 | assert not per_act_token |
| 271 | assert len(block_shape) == 2 |
| 272 | _, block_k = block_shape[0], block_shape[1] |
| 273 | A, A_scale = per_token_group_quant_int8(A, block_k) |
| 274 | assert cdiv(A.size(-1), block_k) == A_scale.size(-1) |
| 275 | |
| 276 | return A, A_scale |
| 277 | |
| 278 | |
| 279 | def _mxfp4_quantize( |
| 280 | A: torch.Tensor, |
| 281 | A_scale: torch.Tensor | None, |
| 282 | per_act_token_quant: bool, |
| 283 | block_shape: list[int] | None = None, |
| 284 | ) -> tuple[torch.Tensor, None]: |
| 285 | if current_platform.is_xpu(): |
| 286 | return xpu_mxfp4_quantize(A) |
| 287 | assert block_shape is None |
| 288 | # TODO: native mxfp4 is currently not integrated in vllm, |
| 289 | # so simulating even on devices supporting this data type natively. |
| 290 | # Once integrated, `current_platform.supports_mx()` should be used to |
| 291 | # control quantize+dequantize, or simply quantize here down to mxfp4. |
| 292 | A = quant_dequant_mxfp4(A) |
| 293 | |
| 294 | return A, None |
| 295 | |
| 296 | |
| 297 | def _fp8_quantize_dequantize( |
| 298 | A: torch.Tensor, |
| 299 | A_scale: torch.Tensor, |
| 300 | ): |
| 301 | qA, qA_scale = ops.scaled_fp8_quant(A, A_scale, use_per_token_if_dynamic=False) |
| 302 | A = per_tensor_dequantize(qA, qA_scale).to(A.dtype) |
| 303 | |
| 304 | return A, None |
| 305 | |
| 306 | |
| 307 | def _mxfp8_e4m3_quantize( |
| 308 | A: torch.Tensor, |
| 309 | A_scale: torch.Tensor | None, |
| 310 | per_act_token_quant: bool, |
| 311 | block_shape: list[int] | None = None, |
| 312 | is_sf_swizzled_layout: bool = False, |
| 313 | mx_alignment: int = 0, |
| 314 | ) -> tuple[torch.Tensor, torch.Tensor]: |
| 315 | if current_platform.is_xpu(): |
| 316 | return xpu_mxfp8_quantize(A) |
| 317 | assert A_scale is None |
| 318 | assert not per_act_token_quant |
| 319 | assert block_shape is None or block_shape == [1, 32] |
| 320 | return mxfp8_e4m3_quantize(A, is_sf_swizzled_layout, mx_alignment) |
| 321 | |
| 322 | |
| 323 | def _mxfp6_e3m2_quantize( |
| 324 | A: torch.Tensor, |
| 325 | A_scale: torch.Tensor | None, |
| 326 | per_act_token_quant: bool, |
| 327 | block_shape: list[int] | None = None, |
| 328 | ) -> tuple[torch.Tensor, None]: |
| 329 | assert block_shape is None |
| 330 | |
| 331 | # TODO: native mxfp6 is currently not integrated in vllm, |
| 332 | # so simulating even on devices supporting this data type natively. |
| 333 | # Eventually, there should be a check based on |
| 334 | # `current_platform.supports_mx()` here. |
| 335 | A = quant_dequant_mxfp6(A, quant_dtype="fp6_e3m2") |
| 336 | |
| 337 | return A, None |
| 338 | |
| 339 | |
| 340 | def _mxfp6_e2m3_quantize( |
| 341 | A: torch.Tensor, |
| 342 | A_scale: torch.Tensor | None, |
| 343 | per_act_token_quant: bool, |
| 344 | block_shape: list[int] | None = None, |
| 345 | ) -> tuple[torch.Tensor, None]: |
| 346 | assert block_shape is None |
| 347 | |
| 348 | # TODO: native mxfp6 is currently not integrated in vllm, |
| 349 | # so simulating even on devices supporting this data type natively. |
| 350 | # Eventually, there should be a check based on |
| 351 | # `current_platform.supports_mx()` here. |
| 352 | A = quant_dequant_mxfp6(A, quant_dtype="fp6_e2m3") |
| 353 | |
| 354 | return A, None |
| 355 | |
| 356 | |
| 357 | def moe_kernel_quantize_input( |
| 358 | A: torch.Tensor, |
| 359 | A_scale: torch.Tensor | None, |
| 360 | quant_dtype: None | torch.dtype | str, |
| 361 | per_act_token_quant: bool, |
| 362 | block_shape: list[int] | None = None, |
| 363 | is_scale_swizzled: bool = True, |
| 364 | quantization_emulation: bool = False, |
| 365 | mx_alignment: int = 0, |
| 366 | ) -> tuple[torch.Tensor, torch.Tensor | None]: |
| 367 | if quant_dtype == current_platform.fp8_dtype(): |
| 368 | if quantization_emulation: |
| 369 | return _fp8_quantize_dequantize(A, A_scale) |
| 370 | else: |
| 371 | return _fp8_quantize(A, A_scale, per_act_token_quant, block_shape) |
| 372 | elif quant_dtype == torch.int8: |
| 373 | if quantization_emulation: |
| 374 | raise NotImplementedError( |
| 375 | "moe_kernel_quantize_input does not support quant_dtype=torch.int8" |
| 376 | " MOE quantization emulation. Please open an issue." |
| 377 | ) |
| 378 | return _int8_quantize(A, A_scale, per_act_token_quant, block_shape) |
| 379 | elif quant_dtype == "nvfp4": |
| 380 | if not quantization_emulation: |
| 381 | return _nvfp4_quantize(A, A_scale, is_sf_swizzled_layout=is_scale_swizzled) |
| 382 | else: |
| 383 | assert A_scale is not None |
| 384 | A = ref_nvfp4_quant_dequant(A, A_scale, block_size=16) |
| 385 | return A, None |
| 386 | elif quant_dtype == "mxfp4": |
| 387 | if not current_platform.is_xpu() and not quantization_emulation: |
| 388 | raise NotImplementedError( |
| 389 | "moe_kernel_quantize_input should not be used for native" |
| 390 | " quant_dtype='mxfp4' MOE. Please open an issue." |
| 391 | ) |
| 392 | return _mxfp4_quantize(A, A_scale, per_act_token_quant, block_shape) |
| 393 | elif quant_dtype == "mxfp8": |
| 394 | # TODO: `quant_dtype == "mxfp8"` is ambiguous, |
| 395 | # should be fp8_e4m3. OCP MX also defines `fp8_e5m2`. |
| 396 | if not current_platform.is_xpu() and quantization_emulation: |
| 397 | raise NotImplementedError( |
| 398 | "moe_kernel_quantize_input does not support quant_dtype='mxfp8' MOE " |
| 399 | "quantization emulation. Please open an issue." |
| 400 | ) |
| 401 | # Non-swizzled (M, K/32) uint8 UE8M0 scales; deepgemm_moe_permute packs |
| 402 | # them for DeepGEMM, TRTLLM takes them as-is. |
| 403 | return _mxfp8_e4m3_quantize( |
| 404 | A, |
| 405 | A_scale, |
| 406 | per_act_token_quant, |
| 407 | block_shape, |
| 408 | is_sf_swizzled_layout=is_scale_swizzled, |
| 409 | mx_alignment=mx_alignment, |
| 410 | ) |
| 411 | elif quant_dtype == "mxfp6_e3m2": |
| 412 | if not quantization_emulation: |
| 413 | raise NotImplementedError( |
| 414 | "moe_kernel_quantize_input should not be used for native " |
| 415 | " quant_dtype='mxfp6_e3m2'MOE. Please open an issue." |
| 416 | ) |
| 417 | |
| 418 | return _mxfp6_e3m2_quantize(A, A_scale, per_act_token_quant, block_shape) |
| 419 | elif quant_dtype == "mxfp6_e2m3": |
| 420 | if not quantization_emulation: |
| 421 | raise NotImplementedError( |
| 422 | "moe_kernel_quantize_input should not be used for native" |
| 423 | " quant_dtype='mxfp6_e2m3' MOE. Please open an issue." |
| 424 | ) |
| 425 | |
| 426 | return _mxfp6_e2m3_quantize(A, A_scale, per_act_token_quant, block_shape) |
| 427 | else: |
| 428 | return A, A_scale |
| 429 | |
| 430 | |
| 431 | def normalize_scales_shape(scales: torch.Tensor | None) -> torch.Tensor | None: |
| 432 | if scales is not None: |
| 433 | if scales.numel() == 1: |
| 434 | scales = scales.view(1, 1) |
| 435 | else: |
| 436 | scales = scales.view(-1, scales.size(-1)) |
| 437 | return scales |
| 438 | |
| 439 | |
| 440 | def normalize_batched_scales_shape( |
| 441 | scales: torch.Tensor | None, |
| 442 | num_experts: int, |
| 443 | ) -> torch.Tensor | None: |
| 444 | if scales is not None and scales.ndim < 3: |
| 445 | if scales.numel() == 1: |
| 446 | scales = scales.view(1) |
| 447 | scales = torch.repeat_interleave(scales, num_experts, dim=0).view( |
| 448 | num_experts, 1, 1 |
| 449 | ) |
| 450 | else: |
| 451 | scales = scales.view(num_experts, -1, scales.size(-1)) |
| 452 | |
| 453 | return scales |
| 454 | |
| 455 | |
| 456 | def fi_moe_largest_bucket(moe_config: "FusedMoEConfig") -> int: |
| 457 | """Estimate FlashInfer's MoE autotuning maximum token count. |
| 458 | |
| 459 | All DP ranks may contribute `max_num_tokens` to one invocation. |
| 460 | Keep FlashInfer's default moe `tune_max_num_tokens=8192` |
| 461 | floor to avoid over-underestimation. |
| 462 | DeepEP, SP, or PCP may make this underestimate, however overestimation |
| 463 | may be dangerous, increasing tuning- cost and memory use. |
| 464 | |
| 465 | NOTE: The DP factor applies even when EP is disabled: |
| 466 | > Without `--enable-expert-parallel`, MoE layers would use tensor parallelism. |
| 467 | |
| 468 | For a detailed explanation, see: `docs/serving/data_parallel_deployment.md` |
| 469 | """ |
| 470 | dp_size = moe_config.elastic_ep_max_dp_size or moe_config.dp_size |
| 471 | return max(moe_config.max_num_tokens * dp_size, 8192) |
| 472 | |
| 473 | |
| 474 | @torch.compile(dynamic=True, backend=current_platform.simple_compile_backend) |
| 475 | def _swiglu_limit_torch( |
| 476 | output: torch.Tensor, |
| 477 | input: torch.Tensor, # first half is gate, second half is up |
| 478 | swiglu_limit: float = 0.0, |
| 479 | ) -> None: |
| 480 | d = input.shape[1] // 2 |
| 481 | gate = input[:, :d] |
| 482 | up = input[:, d:] |
| 483 | |
| 484 | if swiglu_limit > 0: |
| 485 | gate = torch.clamp(gate, max=swiglu_limit) |
| 486 | up = torch.clamp(up, min=-swiglu_limit, max=swiglu_limit) |
| 487 | |
| 488 | output.copy_(F.silu(gate) * up) |
| 489 | |
| 490 | |
| 491 | @triton.jit |
| 492 | def _swiglu_limit_pad_aware_kernel( |
| 493 | input_ptr, # [num_tokens, 2 * hidden_size] |
| 494 | output_ptr, # [num_tokens, hidden_size] |
| 495 | topk_ids_ptr, # [num_tokens, num_topk] |
| 496 | expert_map_ptr, # global -> local expert id, or -1 if non-local |
| 497 | hidden_size, |
| 498 | input_row_stride, |
| 499 | num_tokens, |
| 500 | swiglu_limit, |
| 501 | HAS_LIMIT: tl.constexpr, |
| 502 | HAS_EXPERT_MAP: tl.constexpr, |
| 503 | BLOCK_SIZE: tl.constexpr, |
| 504 | ): |
| 505 | # Persistent over rows: each CTA owns one column tile and processes a |
| 506 | # strided set of token assignments. |
| 507 | pid = tl.program_id(0) |
| 508 | row_stride = tl.num_programs(0) |
| 509 | column_tile = tl.program_id(1) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) |
| 510 | mask = column_tile < hidden_size |
| 511 | |
| 512 | for row in tl.range(pid, num_tokens, row_stride): |
| 513 | expert_id = tl.load(topk_ids_ptr + row) |
| 514 | should_compute = expert_id != -1 |
| 515 | if HAS_EXPERT_MAP: |
| 516 | local_expert_id = tl.load( |
| 517 | expert_map_ptr + expert_id, |
| 518 | mask=expert_id >= 0, |
| 519 | other=-1, |
| 520 | ) |
| 521 | should_compute = should_compute & (local_expert_id != -1) |
| 522 | |
| 523 | if should_compute: |
| 524 | gate_offsets = row.to(tl.int64) * input_row_stride + column_tile |
| 525 | up_offsets = gate_offsets + hidden_size |
| 526 | |
| 527 | gate = tl.load(input_ptr + gate_offsets, mask=mask, other=0.0).to( |
| 528 | tl.float32 |
| 529 | ) |
| 530 | |
| 531 | up = tl.load(input_ptr + up_offsets, mask=mask, other=0.0).to(tl.float32) |
| 532 | |
| 533 | if HAS_LIMIT: |
| 534 | gate = tl.minimum(gate, swiglu_limit) |
| 535 | up = tl.maximum(up, -swiglu_limit) |
| 536 | up = tl.minimum(up, swiglu_limit) |
| 537 | |
| 538 | silu_gate = gate / (1.0 + tl.exp(-gate)) |
| 539 | result = silu_gate * up |
| 540 | tl.store( |
| 541 | output_ptr + row.to(tl.int64) * hidden_size + column_tile, |
| 542 | result.to(output_ptr.dtype.element_ty), |
| 543 | mask=mask, |
| 544 | ) |
| 545 | |
| 546 | |
| 547 | def _swiglu_limit_pad_aware( |
| 548 | output: torch.Tensor, |
| 549 | input: torch.Tensor, |
| 550 | topk_ids: torch.Tensor, |
| 551 | swiglu_limit: float, |
| 552 | expert_map: torch.Tensor | None = None, |
| 553 | ) -> None: |
| 554 | num_tokens, gate_up_size = input.shape |
| 555 | hidden_size = gate_up_size // 2 |
| 556 | if num_tokens == 0: |
| 557 | return |
| 558 | |
| 559 | BLOCK_SIZE = 1024 |
| 560 | grid = (min(num_tokens, 256), triton.cdiv(hidden_size, BLOCK_SIZE)) |
| 561 | _swiglu_limit_pad_aware_kernel[grid]( |
| 562 | input, |
| 563 | output, |
| 564 | topk_ids, |
| 565 | expert_map, |
| 566 | hidden_size, |
| 567 | gate_up_size, |
| 568 | num_tokens, |
| 569 | swiglu_limit, |
| 570 | HAS_LIMIT=swiglu_limit > 0, |
| 571 | HAS_EXPERT_MAP=expert_map is not None, |
| 572 | BLOCK_SIZE=BLOCK_SIZE, |
| 573 | num_warps=4, |
| 574 | ) |
| 575 | |
| 576 | |
| 577 | def swiglu_limit_func( |
| 578 | output: torch.Tensor, |
| 579 | input: torch.Tensor, # first half is gate, second half is up |
| 580 | swiglu_limit: float = 0.0, |
| 581 | topk_ids: torch.Tensor | None = None, |
| 582 | expert_map: torch.Tensor | None = None, |
| 583 | ) -> None: |
| 584 | # The pad-aware Triton kernel skips unrouted token slots (topk_ids == -1) |
| 585 | # and, when expert_map is given, slots routed to non-local experts, so it |
| 586 | # requires topk_ids. Fall back to the torch implementation otherwise. |
| 587 | if topk_ids is not None: |
| 588 | _swiglu_limit_pad_aware(output, input, topk_ids, swiglu_limit, expert_map) |
| 589 | elif current_platform.is_cuda(): |
| 590 | torch.ops._C.silu_and_mul_with_clamp(output, input, swiglu_limit, 1.0, 0.0) |
| 591 | else: |
| 592 | _swiglu_limit_torch(output, input, swiglu_limit) |
| 593 | |
| 594 | |
| 595 | @functools.lru_cache |
| 596 | def enable_swap_ab(BLOCK_SIZE_M: int, BLOCK_SIZE_N: int) -> bool: |
| 597 | return ( |
| 598 | current_platform.is_device_capability(90) |
| 599 | and BLOCK_SIZE_M < 64 |
| 600 | and BLOCK_SIZE_N >= 64 |
| 601 | ) |
| 602 | |
| 603 | |
| 604 | def moe_use_td_hw_supported() -> bool: |
| 605 | """Whether the current device can run the TD (gather) path of |
| 606 | ``fused_moe_kernel`` (ignores the ``VLLM_TRITON_USE_TD`` override). |
| 607 | |
| 608 | The A-load uses ``tensor_descriptor.gather``, which lowers to the PTX |
| 609 | ``tile::gather4`` instruction. That instruction is part of the |
| 610 | ``tcgen05``/Tensor Memory (TMEM) family introduced with Blackwell and has |
| 611 | no Hopper (sm90) equivalent -- ptxas rejects it there ("Feature |
| 612 | '.tile::gather4 ...' requires .target sm_100 or higher"). Unlike |
| 613 | ``scatter4``, ``gather4`` is supported across the whole sm100+ range |
| 614 | including consumer Blackwell (sm120/sm121): see triton-lang/triton#8498, |
| 615 | which enables ``gather4`` on sm120/sm121 while leaving ``scatter4`` |
| 616 | unsupported there. So this gates on a blanket ``has_device_capability(100)`` |
| 617 | rather than the sm100 *family* check used for the scatter store path. |
| 618 | """ |
| 619 | if current_platform.is_xpu(): |
| 620 | return True |
| 621 | if current_platform.is_cuda(): |
| 622 | return current_platform.has_device_capability(100) |
| 623 | return False |
| 624 | |
| 625 | |
| 626 | def resolve_moe_use_td() -> bool: |
| 627 | """Tri-state resolver for ``VLLM_TRITON_USE_TD``. |
| 628 | |
| 629 | Unset auto-selects the TD path on XPU only, mirroring the attention |
| 630 | dispatcher in ``triton_attn.py``. ``1``/``0`` force it on/off regardless |
| 631 | of hardware; forcing ``1`` where it cannot compile (see |
| 632 | ``moe_use_td_hw_supported``) fails at ptxas. Blackwell CUDA (sm100+) can |
| 633 | compile it but is opt-in only, pending validation. |
| 634 | """ |
| 635 | override = envs.VLLM_TRITON_USE_TD |
| 636 | if override is None: |
| 637 | return current_platform.is_xpu() |
| 638 | return override |
| 639 | |
| 640 | |
| 641 | _warned_moe_use_td_ineffective = False |
| 642 | |
| 643 | |
| 644 | def warn_if_moe_use_td_ineffective( |
| 645 | active_backend: str, is_quantized: bool = False |
| 646 | ) -> None: |
| 647 | """One-shot warning when ``VLLM_TRITON_USE_TD`` is set but ignored. |
| 648 | |
| 649 | Fires when the user set the env explicitly and either (a) the active |
| 650 | MoE backend is not the fused Triton kernel, or (b) the model is |
| 651 | quantized (the TD path falls back to the pointer path under any |
| 652 | quantization). |
| 653 | """ |
| 654 | global _warned_moe_use_td_ineffective |
| 655 | if _warned_moe_use_td_ineffective: |
| 656 | return |
| 657 | if envs.VLLM_TRITON_USE_TD is None: |
| 658 | return |
| 659 | is_triton = active_backend.upper() == "TRITON" |
| 660 | if is_triton and not is_quantized: |
| 661 | return |
| 662 | if not is_triton: |
| 663 | reason = ( |
| 664 | f"the active MoE backend is {active_backend!r}; pass " |
| 665 | "`--moe-backend triton` to enable the tensor-descriptor path" |
| 666 | ) |
| 667 | else: |
| 668 | reason = ( |
| 669 | "the model uses quantized MoE weights; the TD path is " |
| 670 | "currently restricted to non-quantized weights and falls " |
| 671 | "back to the pointer path" |
| 672 | ) |
| 673 | logger.warning( |
| 674 | "VLLM_TRITON_USE_TD is set to %s but %s.", |
| 675 | envs.VLLM_TRITON_USE_TD, |
| 676 | reason, |
| 677 | ) |
| 678 | _warned_moe_use_td_ineffective = True |