| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | """Dispatch DSV4.1 mHC operations and overlap coefficient generation.""" |
| 4 | |
| 5 | from functools import partial |
| 6 | from typing import TYPE_CHECKING, cast |
| 7 | |
| 8 | import torch |
| 9 | |
| 10 | from vllm.config import VllmConfig |
| 11 | from vllm.distributed import get_tp_group |
| 12 | from vllm.logger import init_logger |
| 13 | from vllm.model_executor.kernels.mhc.tilelang import ( |
| 14 | mhc_fused_post_pre_delayed_tilelang, |
| 15 | mhc_post_tilelang, |
| 16 | ) |
| 17 | from vllm.model_executor.layers.fused_moe.moe_output import ( |
| 18 | MoEOutput, |
| 19 | UnfinalizedMoEOutput, |
| 20 | ) |
| 21 | from vllm.platforms import current_platform |
| 22 | from vllm.utils.deep_gemm import is_deep_gemm_supported |
| 23 | |
| 24 | from .mega_mhc import ( |
| 25 | can_use_mega_mhc, |
| 26 | mhc_shifted_post_pre_deep_gemm, |
| 27 | ) |
| 28 | |
| 29 | if TYPE_CHECKING: |
| 30 | from vllm.distributed.device_communicators.cuda_communicator import CudaCommunicator |
| 31 | |
| 32 | from .cute_dsl import AllReduceMHC |
| 33 | |
| 34 | logger = init_logger(__name__) |
| 35 | |
| 36 | # GB200 TP4/FlashInfer improves through 16 tokens; larger screens tie or regress. |
| 37 | MHC_OVERLAP_MAX_TOKENS = 16 |
| 38 | |
| 39 | # The fused all-reduce + mHC kernel of the overlap path, built once by |
| 40 | # init_mhc_all_reduce. |
| 41 | _all_reduce_mhc: "AllReduceMHC | None" = None |
| 42 | |
| 43 | |
| 44 | def supports_mhc_overlap(vllm_config: VllmConfig) -> bool: |
| 45 | """Check kernel requirements and safety of sharing the coefficient stream.""" |
| 46 | config = vllm_config.model_config.hf_config |
| 47 | # DeepGEMM's prenorm kernel requires K % 64 == 0, N % 8 == 0, and N <= 32. |
| 48 | mix_size = config.hc_mult * (config.hc_mult + 2) |
| 49 | return ( |
| 50 | current_platform.is_device_capability_family(100) |
| 51 | and is_deep_gemm_supported() |
| 52 | and config.hidden_size % 64 == 0 |
| 53 | and 0 < mix_size <= 32 |
| 54 | and mix_size % 8 == 0 |
| 55 | and not vllm_config.parallel_config.use_ubatching |
| 56 | ) |
| 57 | |
| 58 | |
| 59 | def supports_mhc_all_reduce(vllm_config: VllmConfig) -> bool: |
| 60 | parallel = vllm_config.parallel_config |
| 61 | config = vllm_config.model_config.hf_config |
| 62 | if ( |
| 63 | parallel.tensor_parallel_size != 4 |
| 64 | or parallel.enable_expert_parallel |
| 65 | or config.hidden_size != 5120 |
| 66 | or config.hc_mult != 4 |
| 67 | ): |
| 68 | return False |
| 69 | # The custom all-reduce's MNNVL buffers succeed exactly when NVLink |
| 70 | # multicast does, which the fused kernel's own mailbox needs too. |
| 71 | comm = cast("CudaCommunicator", get_tp_group().device_communicator).ca_comm |
| 72 | return comm is not None and bool(comm.mnnvl_lamport_ag_multicast_ptr) |
| 73 | |
| 74 | |
| 75 | def init_mhc_all_reduce(vllm_config: VllmConfig) -> None: |
| 76 | """Build the fused all-reduce + mHC kernel the overlap path reduces with. |
| 77 | |
| 78 | Collective over the TP group, so every rank must call it, and only when |
| 79 | ``supports_mhc_all_reduce`` holds. |
| 80 | """ |
| 81 | global _all_reduce_mhc |
| 82 | from .cute_dsl import AllReduceMHC |
| 83 | |
| 84 | config = vllm_config.model_config.hf_config |
| 85 | _all_reduce_mhc = AllReduceMHC( |
| 86 | hidden_size=config.hidden_size, |
| 87 | hc_mult=config.hc_mult, |
| 88 | max_num_tokens=MHC_OVERLAP_MAX_TOKENS, |
| 89 | top_k=config.num_experts_per_tok, |
| 90 | device=current_platform.current_device(), |
| 91 | ) |
| 92 | logger.info_once( |
| 93 | "DSV4.1 mHC: CuTe DSL Lamport all-reduce fused with mHC for up to %d tokens.", |
| 94 | MHC_OVERLAP_MAX_TOKENS, |
| 95 | ) |
| 96 | |
| 97 | |
| 98 | def mhc_pre_delayed_overlap( |
| 99 | residual: torch.Tensor, |
| 100 | fn: torch.Tensor, |
| 101 | hc_scale: torch.Tensor, |
| 102 | hc_base: torch.Tensor, |
| 103 | rms_eps: float, |
| 104 | hc_pre_eps: float, |
| 105 | hc_sinkhorn_eps: float, |
| 106 | hc_post_mult_value: float, |
| 107 | sinkhorn_repeat: int, |
| 108 | pre_mix: torch.Tensor | None = None, |
| 109 | x: torch.Tensor | None = None, |
| 110 | norm_weight: torch.Tensor | None = None, |
| 111 | norm_eps: float = 1e-6, |
| 112 | *, |
| 113 | stream: torch.cuda.Stream, |
| 114 | layer_input: torch.Tensor | None = None, |
| 115 | ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]: |
| 116 | """Prepare the input on the caller stream and coefficients on another stream. |
| 117 | |
| 118 | Returns post mix, residual mix, normalized input, and next pre-mix. Only the |
| 119 | input is ready on the caller stream; join stream before using coefficients. |
| 120 | The caller must retain the weights and coefficient outputs until the join. |
| 121 | """ |
| 122 | from vllm.model_executor.kernels.mhc.warmup import ( |
| 123 | MHC_PRE_NORM_KERNEL, |
| 124 | compute_mhc_pre_num_splits, |
| 125 | ) |
| 126 | from vllm.utils.deep_gemm import tf32_hc_prenorm_gemm |
| 127 | |
| 128 | n, hc, hidden = residual.shape |
| 129 | assert residual.is_contiguous() and residual.dtype == torch.bfloat16 |
| 130 | assert norm_weight is not None |
| 131 | if x is None: |
| 132 | x = residual.view(n, hc * hidden) |
| 133 | assert x.is_contiguous() |
| 134 | post = torch.empty((n, hc), device=residual.device, dtype=torch.float32) |
| 135 | comb = torch.empty((n, hc * hc), device=residual.device, dtype=torch.float32) |
| 136 | next_pre = torch.empty_like(post) |
| 137 | compute_input = layer_input is None |
| 138 | if layer_input is None: |
| 139 | layer_input = torch.empty( |
| 140 | (n, hidden), device=residual.device, dtype=torch.bfloat16 |
| 141 | ) |
| 142 | outputs = post.unsqueeze(-1), comb.view(n, hc, hc), layer_input, next_pre |
| 143 | if n == 0: |
| 144 | return outputs |
| 145 | splits = compute_mhc_pre_num_splits(x.shape[1], n) |
| 146 | mix = torch.empty( |
| 147 | (splits, n, hc * (hc + 2)), device=residual.device, dtype=torch.float32 |
| 148 | ) |
| 149 | sqr = torch.empty((splits, n), device=residual.device, dtype=torch.float32) |
| 150 | epilogue = partial( |
| 151 | MHC_PRE_NORM_KERNEL, |
| 152 | mix, |
| 153 | sqr, |
| 154 | hc_scale, |
| 155 | hc_base, |
| 156 | residual, |
| 157 | post, |
| 158 | comb, |
| 159 | layer_input, |
| 160 | norm_weight, |
| 161 | pre_mix if pre_mix is not None else post, |
| 162 | next_pre, |
| 163 | layer_input, # Unused aux output in split modes. |
| 164 | hidden_size=hidden, |
| 165 | rms_eps=rms_eps, |
| 166 | hc_pre_eps=hc_pre_eps, |
| 167 | hc_sinkhorn_eps=hc_sinkhorn_eps, |
| 168 | hc_post_mult_value=hc_post_mult_value, |
| 169 | sinkhorn_repeat=sinkhorn_repeat, |
| 170 | norm_eps=norm_eps, |
| 171 | hc_mult=hc, |
| 172 | use_pre_mix_in=pre_mix is not None, |
| 173 | save_pre_mix=True, |
| 174 | rms_numel=x.shape[1], |
| 175 | ) |
| 176 | main = torch.cuda.current_stream() |
| 177 | # Prioritize input readiness for small batches before releasing statistics. |
| 178 | if compute_input and n <= 8: |
| 179 | epilogue(split_mode="input") |
| 180 | stream.wait_stream(main) |
| 181 | with torch.cuda.stream(stream): |
| 182 | tf32_hc_prenorm_gemm(x, fn, mix, sqr, splits) |
| 183 | epilogue(split_mode="stats") |
| 184 | for tensor in (x, mix, sqr): |
| 185 | tensor.record_stream(stream) |
| 186 | if compute_input and n > 8: |
| 187 | epilogue(split_mode="input") |
| 188 | return outputs |
| 189 | |
| 190 | |
| 191 | def mhc_shifted_post_pre( |
| 192 | x: torch.Tensor | MoEOutput, |
| 193 | residual: torch.Tensor, |
| 194 | post_layer_mix: torch.Tensor, |
| 195 | comb_res_mix: torch.Tensor, |
| 196 | fn: torch.Tensor, |
| 197 | hc_scale: torch.Tensor, |
| 198 | hc_base: torch.Tensor, |
| 199 | rms_eps: float, |
| 200 | hc_pre_eps: float, |
| 201 | hc_sinkhorn_eps: float, |
| 202 | hc_post_mult_value: float, |
| 203 | sinkhorn_repeat: int, |
| 204 | pre_mix: torch.Tensor | None = None, |
| 205 | norm_weight: torch.Tensor | None = None, |
| 206 | norm_eps: float = 1e-6, |
| 207 | capture_aux: bool = False, |
| 208 | *, |
| 209 | stream: torch.cuda.Stream | None = None, |
| 210 | reduce_results: bool = False, |
| 211 | ) -> tuple[ |
| 212 | torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor |
| 213 | ]: |
| 214 | """Dispatch shifted post/pre to overlap, Mega-mHC, or fused TileLang. |
| 215 | |
| 216 | A MoE output left unfinalized is finalized inside the fused all-reduce. |
| 217 | When stream is supplied, join it before consuming the returned coefficients. |
| 218 | """ |
| 219 | layer_input = None |
| 220 | if isinstance(x, MoEOutput): |
| 221 | assert _all_reduce_mhc is not None, "init_mhc_all_reduce was not called" |
| 222 | assert pre_mix is not None and norm_weight is not None |
| 223 | routed = x.routed |
| 224 | assert isinstance(routed, UnfinalizedMoEOutput) |
| 225 | assert x.shared_output is not None |
| 226 | residual, layer_input = _all_reduce_mhc.finalize( |
| 227 | routed.gemm2_permuted, |
| 228 | routed.expert_weights, |
| 229 | routed.expanded_idx_to_permuted_idx, |
| 230 | x.shared_output, |
| 231 | residual, |
| 232 | post_layer_mix, |
| 233 | comb_res_mix, |
| 234 | pre_mix, |
| 235 | norm_weight, |
| 236 | norm_eps, |
| 237 | ) |
| 238 | if stream is None: |
| 239 | # Outside FULL graphs, compute the next mixes inline. |
| 240 | stream = torch.cuda.current_stream() |
| 241 | elif reduce_results: |
| 242 | if stream is not None and 0 < x.shape[0] <= MHC_OVERLAP_MAX_TOKENS: |
| 243 | assert _all_reduce_mhc is not None, "init_mhc_all_reduce was not called" |
| 244 | assert pre_mix is not None and norm_weight is not None |
| 245 | residual, layer_input = _all_reduce_mhc( |
| 246 | x, |
| 247 | residual, |
| 248 | post_layer_mix, |
| 249 | comb_res_mix, |
| 250 | pre_mix, |
| 251 | norm_weight, |
| 252 | norm_eps, |
| 253 | ) |
| 254 | else: |
| 255 | x = get_tp_group().all_reduce(x) |
| 256 | if stream is not None: |
| 257 | if layer_input is None: |
| 258 | assert isinstance(x, torch.Tensor) |
| 259 | residual = mhc_post_tilelang(x, residual, post_layer_mix, comb_res_mix) |
| 260 | aux = ( |
| 261 | residual.mean(dim=1) |
| 262 | if capture_aux |
| 263 | else residual.new_empty(0, residual.shape[-1]) |
| 264 | ) |
| 265 | pre_outputs = mhc_pre_delayed_overlap( |
| 266 | residual, |
| 267 | fn, |
| 268 | hc_scale, |
| 269 | hc_base, |
| 270 | rms_eps, |
| 271 | hc_pre_eps, |
| 272 | hc_sinkhorn_eps, |
| 273 | hc_post_mult_value, |
| 274 | sinkhorn_repeat, |
| 275 | pre_mix=pre_mix, |
| 276 | norm_weight=norm_weight, |
| 277 | norm_eps=norm_eps, |
| 278 | stream=stream, |
| 279 | layer_input=layer_input, |
| 280 | ) |
| 281 | return residual, *pre_outputs, aux |
| 282 | |
| 283 | assert isinstance(x, torch.Tensor) |
| 284 | if can_use_mega_mhc(x, residual, pre_mix, norm_weight, capture_aux): |
| 285 | assert pre_mix is not None and norm_weight is not None |
| 286 | outputs = mhc_shifted_post_pre_deep_gemm( |
| 287 | x, |
| 288 | residual, |
| 289 | pre_mix, |
| 290 | post_layer_mix, |
| 291 | comb_res_mix, |
| 292 | fn, |
| 293 | hc_scale, |
| 294 | hc_base, |
| 295 | rms_eps, |
| 296 | hc_pre_eps, |
| 297 | hc_post_mult_value, |
| 298 | hc_sinkhorn_eps, |
| 299 | sinkhorn_repeat, |
| 300 | norm_weight, |
| 301 | norm_eps, |
| 302 | ) |
| 303 | return *outputs, x.new_empty(0, x.shape[1]) |
| 304 | |
| 305 | return mhc_fused_post_pre_delayed_tilelang( |
| 306 | x, |
| 307 | residual, |
| 308 | post_layer_mix, |
| 309 | comb_res_mix, |
| 310 | fn, |
| 311 | hc_scale, |
| 312 | hc_base, |
| 313 | rms_eps, |
| 314 | hc_pre_eps, |
| 315 | hc_sinkhorn_eps, |
| 316 | hc_post_mult_value, |
| 317 | sinkhorn_repeat, |
| 318 | pre_mix=pre_mix, |
| 319 | norm_weight=norm_weight, |
| 320 | norm_eps=norm_eps, |
| 321 | capture_aux=capture_aux, |
| 322 | ) |