| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | """Benchmark the Kimi K3 latent-MoE tail and its up-projection kernels. |
| 4 | |
| 5 | The ``up-projection`` subcommand isolates the TP-local dynamic and static-M |
| 6 | skinny GEMMs. It rotates weights through a working set larger than L2 to model |
| 7 | successive model layers. |
| 8 | |
| 9 | The ``whole-tail`` subcommand measures the distributed operator. Its reference |
| 10 | path includes two AllReduces, RMSNorm, the replicated up-projection, and the |
| 11 | final add. CUDA-event samples report the slowest rank so cross-rank skew is |
| 12 | included. |
| 13 | |
| 14 | Examples: |
| 15 | |
| 16 | .. code-block:: console |
| 17 | |
| 18 | .venv/bin/python \ |
| 19 | benchmarks/kernels/benchmark_kimi_k3_latent_moe_tail.py up-projection |
| 20 | |
| 21 | torchrun --nproc-per-node=8 \ |
| 22 | benchmarks/kernels/benchmark_kimi_k3_latent_moe_tail.py whole-tail |
| 23 | |
| 24 | torchrun --nproc-per-node=8 \ |
| 25 | benchmarks/kernels/benchmark_kimi_k3_latent_moe_tail.py whole-tail \ |
| 26 | --backend fused --routed-input both |
| 27 | |
| 28 | For multi-node runs, launch one ``torchrun`` agent per node and use a shared |
| 29 | rendezvous endpoint. |
| 30 | |
| 31 | """ |
| 32 | |
| 33 | from __future__ import annotations |
| 34 | |
| 35 | import argparse |
| 36 | import json |
| 37 | import math |
| 38 | import os |
| 39 | import statistics |
| 40 | from collections.abc import Callable, Sequence |
| 41 | from dataclasses import asdict |
| 42 | from pathlib import Path |
| 43 | from typing import Any |
| 44 | |
| 45 | import cutlass |
| 46 | import cutlass.utils as utils |
| 47 | import torch |
| 48 | import torch.distributed as dist |
| 49 | import torch.nn.functional as F |
| 50 | from cuda.bindings import driver as cuda |
| 51 | |
| 52 | from vllm.distributed import get_tp_group |
| 53 | from vllm.distributed.parallel_state import ( |
| 54 | init_distributed_environment, |
| 55 | initialize_model_parallel, |
| 56 | set_custom_all_reduce, |
| 57 | ) |
| 58 | from vllm.model_executor.layers.fused_moe.moe_output import UnfinalizedMoEOutput |
| 59 | from vllm.model_executor.warmup.cutedsl_warmup import cutedsl_warmup |
| 60 | from vllm.models.kimi_k3.nvidia.ops import latent_moe_tail |
| 61 | from vllm.models.kimi_k3.nvidia.ops.cute_dsl.latent_moe_tail import ( |
| 62 | fused_add_multicast_gemm, |
| 63 | fused_add_multicast_skinny_gemm, |
| 64 | ) |
| 65 | |
| 66 | HIDDEN_SIZE = 7168 |
| 67 | LATENT_SIZE = 3584 |
| 68 | RMS_EPS = 0.1 |
| 69 | MAX_NUM_TOKENS = 16 |
| 70 | MMA_TILER_MN = (64, 32) |
| 71 | CLUSTER_SHAPE_MN = (1, 8) |
| 72 | B_PRIME_STAGES = 2 |
| 73 | |
| 74 | |
| 75 | def parse_up_projection_config( |
| 76 | value: str, |
| 77 | ) -> fused_add_multicast_skinny_gemm.SkinnyConfig: |
| 78 | try: |
| 79 | values = [int(part) for part in value.split(",")] |
| 80 | except ValueError as error: |
| 81 | raise argparse.ArgumentTypeError( |
| 82 | "config must be BLOCK,OUTPUTS,K_UNROLL[,VECTOR_WIDTH[,PREFETCH_B]]" |
| 83 | ) from error |
| 84 | if len(values) in (3, 4): |
| 85 | return fused_add_multicast_skinny_gemm.SkinnyConfig(*values) |
| 86 | if len(values) == 5 and values[4] in (0, 1): |
| 87 | return fused_add_multicast_skinny_gemm.SkinnyConfig( |
| 88 | *values[:4], |
| 89 | prefetch_b_before_pdl=bool(values[4]), |
| 90 | ) |
| 91 | raise argparse.ArgumentTypeError( |
| 92 | "config must be BLOCK,OUTPUTS,K_UNROLL" |
| 93 | "[,VECTOR_WIDTH[,PREFETCH_B]], where PREFETCH_B is 0 or 1" |
| 94 | ) |
| 95 | |
| 96 | |
| 97 | def parse_tail_skinny_config( |
| 98 | value: str, |
| 99 | ) -> tuple[int, fused_add_multicast_skinny_gemm.SkinnyConfig]: |
| 100 | try: |
| 101 | values = [int(part) for part in value.split(",")] |
| 102 | except ValueError as error: |
| 103 | raise argparse.ArgumentTypeError( |
| 104 | "config must be M,BLOCK,OUTPUTS,K_UNROLL[,VECTOR_WIDTH[,PREFETCH_B]]" |
| 105 | ) from error |
| 106 | if len(values) == 4: |
| 107 | num_tokens, *config = values |
| 108 | return num_tokens, fused_add_multicast_skinny_gemm.SkinnyConfig(*config) |
| 109 | if len(values) == 5: |
| 110 | num_tokens, *config = values |
| 111 | return num_tokens, fused_add_multicast_skinny_gemm.SkinnyConfig(*config) |
| 112 | if len(values) == 6 and values[5] in (0, 1): |
| 113 | num_tokens, block, outputs, unroll, vector_width, prefetch = values |
| 114 | return num_tokens, fused_add_multicast_skinny_gemm.SkinnyConfig( |
| 115 | block, |
| 116 | outputs, |
| 117 | unroll, |
| 118 | vector_width, |
| 119 | bool(prefetch), |
| 120 | ) |
| 121 | raise argparse.ArgumentTypeError( |
| 122 | "config must be M,BLOCK,OUTPUTS,K_UNROLL" |
| 123 | "[,VECTOR_WIDTH[,PREFETCH_B]], where PREFETCH_B is 0 or 1" |
| 124 | ) |
| 125 | |
| 126 | |
| 127 | def parse_args() -> argparse.Namespace: |
| 128 | parser = argparse.ArgumentParser(description=__doc__) |
| 129 | subparsers = parser.add_subparsers(dest="scope", required=True) |
| 130 | |
| 131 | up_projection = subparsers.add_parser( |
| 132 | "up-projection", |
| 133 | help="Benchmark the isolated TP-local up-projection kernels.", |
| 134 | ) |
| 135 | up_projection.add_argument( |
| 136 | "--backend", |
| 137 | choices=("dynamic", "skinny", "both"), |
| 138 | default="both", |
| 139 | ) |
| 140 | up_projection.add_argument("--tp-size", type=int, default=16) |
| 141 | up_projection.add_argument( |
| 142 | "--num-tokens", |
| 143 | type=int, |
| 144 | nargs="+", |
| 145 | default=[*range(1, 9), 16], |
| 146 | ) |
| 147 | up_projection.add_argument( |
| 148 | "--skinny-config", |
| 149 | type=parse_up_projection_config, |
| 150 | action="append", |
| 151 | help="Benchmark a static-M config for every selected token count.", |
| 152 | ) |
| 153 | up_projection.add_argument("--cache-multiplier", type=float, default=2.0) |
| 154 | up_projection.add_argument("--max-weights", type=int, default=64) |
| 155 | up_projection.add_argument("--warmup-replays", type=int, default=10) |
| 156 | up_projection.add_argument("--samples", type=int, default=31) |
| 157 | up_projection.add_argument("--output", type=Path) |
| 158 | |
| 159 | whole_tail = subparsers.add_parser( |
| 160 | "whole-tail", |
| 161 | help="Benchmark the distributed latent-MoE tail operator.", |
| 162 | ) |
| 163 | whole_tail.add_argument( |
| 164 | "--backend", |
| 165 | choices=("reference", "fused", "both"), |
| 166 | default="both", |
| 167 | ) |
| 168 | whole_tail.add_argument( |
| 169 | "--num-tokens", |
| 170 | type=int, |
| 171 | nargs="+", |
| 172 | default=[1, 5, 8, 16], |
| 173 | ) |
| 174 | whole_tail.add_argument( |
| 175 | "--routed-input", |
| 176 | choices=("finalized", "deferred", "both"), |
| 177 | default="finalized", |
| 178 | help="Select whether the fused tail performs the top-k finalize step.", |
| 179 | ) |
| 180 | whole_tail.add_argument( |
| 181 | "--top-k", |
| 182 | type=int, |
| 183 | default=8, |
| 184 | help="Expert count used by the synthetic deferred-finalize input.", |
| 185 | ) |
| 186 | whole_tail.add_argument("--warmup-replays", type=int, default=20) |
| 187 | whole_tail.add_argument("--samples", type=int, default=51) |
| 188 | whole_tail.add_argument( |
| 189 | "--skinny-max-num-tokens", |
| 190 | type=int, |
| 191 | nargs="+", |
| 192 | help="Override the fused operator's static-M cutoff; use 0 for dynamic-only.", |
| 193 | ) |
| 194 | whole_tail.add_argument( |
| 195 | "--skinny-config", |
| 196 | type=parse_tail_skinny_config, |
| 197 | action="append", |
| 198 | help="Override one static-M config for tuning.", |
| 199 | ) |
| 200 | whole_tail.add_argument("--output", type=Path) |
| 201 | return parser.parse_args() |
| 202 | |
| 203 | |
| 204 | def percentile(samples: Sequence[float], fraction: float) -> float: |
| 205 | ordered = sorted(samples) |
| 206 | position = fraction * (len(ordered) - 1) |
| 207 | lower = math.floor(position) |
| 208 | upper = math.ceil(position) |
| 209 | if lower == upper: |
| 210 | return ordered[lower] |
| 211 | upper_weight = position - lower |
| 212 | return ordered[lower] * (1.0 - upper_weight) + ordered[upper] * upper_weight |
| 213 | |
| 214 | |
| 215 | def summarize(samples_us: Sequence[float]) -> dict[str, Any]: |
| 216 | mean_us = statistics.mean(samples_us) |
| 217 | return { |
| 218 | "median_us": statistics.median(samples_us), |
| 219 | "p10_us": percentile(samples_us, 0.1), |
| 220 | "p90_us": percentile(samples_us, 0.9), |
| 221 | "mean_us": mean_us, |
| 222 | "cv_pct": statistics.pstdev(samples_us) / mean_us * 100.0, |
| 223 | "samples_us": list(samples_us), |
| 224 | } |
| 225 | |
| 226 | |
| 227 | def rotating_weight_count( |
| 228 | shard_size: int, |
| 229 | cache_multiplier: float, |
| 230 | limit: int, |
| 231 | ) -> int: |
| 232 | properties = torch.cuda.get_device_properties( |
| 233 | torch.accelerator.current_device_index() |
| 234 | ) |
| 235 | weight_bytes = shard_size * LATENT_SIZE * 2 |
| 236 | target_bytes = math.ceil(properties.L2_cache_size * cache_multiplier) |
| 237 | return max(2, min(limit, math.ceil(target_bytes / weight_bytes))) |
| 238 | |
| 239 | |
| 240 | def capture_up_projection_graph( |
| 241 | launches: Sequence[Callable[[], None]], |
| 242 | ) -> torch.cuda.CUDAGraph: |
| 243 | for launch in launches: |
| 244 | launch() |
| 245 | torch.accelerator.synchronize() |
| 246 | |
| 247 | graph = torch.cuda.CUDAGraph() |
| 248 | with torch.cuda.graph(graph): |
| 249 | for launch in launches: |
| 250 | launch() |
| 251 | torch.accelerator.synchronize() |
| 252 | return graph |
| 253 | |
| 254 | |
| 255 | def benchmark_up_projection_graph( |
| 256 | graph: torch.cuda.CUDAGraph, |
| 257 | *, |
| 258 | operations_per_replay: int, |
| 259 | warmup_replays: int, |
| 260 | samples: int, |
| 261 | ) -> dict[str, Any]: |
| 262 | for _ in range(warmup_replays): |
| 263 | graph.replay() |
| 264 | torch.accelerator.synchronize() |
| 265 | |
| 266 | samples_us = [] |
| 267 | start = torch.cuda.Event(enable_timing=True) |
| 268 | end = torch.cuda.Event(enable_timing=True) |
| 269 | for _ in range(samples): |
| 270 | start.record() |
| 271 | graph.replay() |
| 272 | end.record() |
| 273 | end.synchronize() |
| 274 | samples_us.append(start.elapsed_time(end) * 1000.0 / operations_per_replay) |
| 275 | return summarize(samples_us) |
| 276 | |
| 277 | |
| 278 | class DynamicKernel: |
| 279 | def __init__( |
| 280 | self, |
| 281 | shard_size: int, |
| 282 | mailbox: torch.Tensor, |
| 283 | shared_shard: torch.Tensor, |
| 284 | ) -> None: |
| 285 | self.shard_size = shard_size |
| 286 | self.mailbox = mailbox |
| 287 | self.mailbox_c = fused_add_multicast_gemm._as_cute(mailbox) |
| 288 | compile_latent = torch.empty( |
| 289 | (1, MAX_NUM_TOKENS, LATENT_SIZE), |
| 290 | dtype=torch.bfloat16, |
| 291 | device=mailbox.device, |
| 292 | ) |
| 293 | compile_weight = torch.empty( |
| 294 | (1, shard_size, LATENT_SIZE), |
| 295 | dtype=torch.bfloat16, |
| 296 | device=mailbox.device, |
| 297 | ) |
| 298 | cluster_size = math.prod(CLUSTER_SHAPE_MN) |
| 299 | max_active_clusters = utils.HardwareInfo().get_max_active_clusters(cluster_size) |
| 300 | self.compiled = fused_add_multicast_gemm.compile_kernel( |
| 301 | (MAX_NUM_TOKENS, shard_size, LATENT_SIZE, 1), |
| 302 | fused_add_multicast_gemm._as_cute( |
| 303 | compile_latent, |
| 304 | dynamic_m=True, |
| 305 | ), |
| 306 | fused_add_multicast_gemm._as_cute(compile_weight), |
| 307 | self.mailbox_c, |
| 308 | fused_add_multicast_gemm._as_cute(shared_shard), |
| 309 | HIDDEN_SIZE, |
| 310 | shard_size, |
| 311 | MMA_TILER_MN, |
| 312 | CLUSTER_SHAPE_MN, |
| 313 | max_active_clusters, |
| 314 | B_PRIME_STAGES, |
| 315 | ) |
| 316 | |
| 317 | def launch( |
| 318 | self, |
| 319 | latent: torch.Tensor, |
| 320 | weight: torch.Tensor, |
| 321 | shared_shard: torch.Tensor, |
| 322 | ) -> None: |
| 323 | stream = cuda.CUstream(torch.cuda.current_stream().cuda_stream) |
| 324 | self.compiled( |
| 325 | fused_add_multicast_gemm._as_cute( |
| 326 | latent.unsqueeze(0), |
| 327 | dynamic_m=True, |
| 328 | ), |
| 329 | fused_add_multicast_gemm._as_cute(weight.unsqueeze(0)), |
| 330 | self.mailbox_c, |
| 331 | fused_add_multicast_gemm._as_cute(shared_shard), |
| 332 | cutlass.Int64(latent.shape[0]), |
| 333 | cutlass.Int64(self.mailbox.data_ptr()), |
| 334 | stream, |
| 335 | ) |
| 336 | |
| 337 | |
| 338 | class SkinnyKernel: |
| 339 | def __init__( |
| 340 | self, |
| 341 | num_tokens: int, |
| 342 | shard_size: int, |
| 343 | config: fused_add_multicast_skinny_gemm.SkinnyConfig, |
| 344 | ) -> None: |
| 345 | self.compiled = fused_add_multicast_skinny_gemm.compile_kernel( |
| 346 | num_rows=num_tokens, |
| 347 | latent_dim=LATENT_SIZE, |
| 348 | hidden_dim=HIDDEN_SIZE, |
| 349 | shard_dim=shard_size, |
| 350 | config=config, |
| 351 | ) |
| 352 | |
| 353 | def launch( |
| 354 | self, |
| 355 | latent: torch.Tensor, |
| 356 | weight: torch.Tensor, |
| 357 | shared_shard: torch.Tensor, |
| 358 | mailbox: torch.Tensor, |
| 359 | ) -> None: |
| 360 | self.compiled( |
| 361 | fused_add_multicast_skinny_gemm._as_cute(latent), |
| 362 | fused_add_multicast_skinny_gemm._as_cute(weight), |
| 363 | fused_add_multicast_skinny_gemm._as_cute(shared_shard), |
| 364 | cutlass.Int64(mailbox.data_ptr()), |
| 365 | cuda.CUstream(torch.cuda.current_stream().cuda_stream), |
| 366 | ) |
| 367 | |
| 368 | |
| 369 | def check_up_projection_output( |
| 370 | actual: torch.Tensor, |
| 371 | latent: torch.Tensor, |
| 372 | weight: torch.Tensor, |
| 373 | shared_shard: torch.Tensor, |
| 374 | ) -> None: |
| 375 | gemm = F.linear(latent.float(), weight.float()).to(torch.bfloat16) |
| 376 | expected = (gemm.float() + shared_shard.float()).to(torch.bfloat16) |
| 377 | torch.testing.assert_close(actual, expected, atol=8e-2, rtol=3e-2) |
| 378 | |
| 379 | |
| 380 | def make_up_projection_launches( |
| 381 | launch: Callable[[torch.Tensor, torch.Tensor, torch.Tensor], None], |
| 382 | latent: torch.Tensor, |
| 383 | weights: Sequence[torch.Tensor], |
| 384 | shared_shard: torch.Tensor, |
| 385 | ) -> list[Callable[[], None]]: |
| 386 | return [ |
| 387 | lambda weight=weight: launch(latent, weight, shared_shard) for weight in weights |
| 388 | ] |
| 389 | |
| 390 | |
| 391 | def benchmark_up_projection(args: argparse.Namespace) -> None: |
| 392 | if args.tp_size <= 0 or HIDDEN_SIZE % args.tp_size: |
| 393 | raise ValueError("TP size must be positive and divide the hidden size") |
| 394 | if any(not 1 <= num_tokens <= MAX_NUM_TOKENS for num_tokens in args.num_tokens): |
| 395 | raise ValueError("--num-tokens values must be in [1, 16]") |
| 396 | if args.cache_multiplier <= 0 or args.max_weights <= 0: |
| 397 | raise ValueError("cache multiplier and max weights must be positive") |
| 398 | if args.warmup_replays < 0 or args.samples <= 0: |
| 399 | raise ValueError("warmup replays must be nonnegative and samples positive") |
| 400 | |
| 401 | torch.accelerator.set_device_index(0) |
| 402 | device = torch.device("cuda", 0) |
| 403 | if torch.cuda.get_device_capability(device)[0] != 10: |
| 404 | raise RuntimeError("Kimi K3 latent-MoE tail requires SM100") |
| 405 | |
| 406 | shard_size = HIDDEN_SIZE // args.tp_size |
| 407 | weight_count = rotating_weight_count( |
| 408 | shard_size, |
| 409 | args.cache_multiplier, |
| 410 | args.max_weights, |
| 411 | ) |
| 412 | torch.manual_seed(20260726) |
| 413 | weights = [ |
| 414 | torch.randn( |
| 415 | (shard_size, LATENT_SIZE), |
| 416 | dtype=torch.bfloat16, |
| 417 | device=device, |
| 418 | ) |
| 419 | / LATENT_SIZE**0.5 |
| 420 | for _ in range(weight_count) |
| 421 | ] |
| 422 | mailbox = torch.empty( |
| 423 | (1, MAX_NUM_TOKENS, HIDDEN_SIZE), |
| 424 | dtype=torch.bfloat16, |
| 425 | device=device, |
| 426 | ) |
| 427 | shared = torch.randn( |
| 428 | (MAX_NUM_TOKENS, HIDDEN_SIZE), |
| 429 | dtype=torch.bfloat16, |
| 430 | device=device, |
| 431 | ) |
| 432 | shared_shard = shared[:, :shard_size] |
| 433 | use_dynamic = args.backend in ("dynamic", "both") |
| 434 | use_skinny = args.backend in ("skinny", "both") |
| 435 | dynamic_kernel = ( |
| 436 | DynamicKernel(shard_size, mailbox, shared_shard) if use_dynamic else None |
| 437 | ) |
| 438 | |
| 439 | results = [] |
| 440 | for num_tokens in args.num_tokens: |
| 441 | latent = torch.randn( |
| 442 | (num_tokens, LATENT_SIZE), |
| 443 | dtype=torch.bfloat16, |
| 444 | device=device, |
| 445 | ) |
| 446 | result: dict[str, Any] = {"num_tokens": num_tokens} |
| 447 | if dynamic_kernel is not None: |
| 448 | launches = make_up_projection_launches( |
| 449 | dynamic_kernel.launch, |
| 450 | latent, |
| 451 | weights, |
| 452 | shared_shard, |
| 453 | ) |
| 454 | graph = capture_up_projection_graph(launches) |
| 455 | result["dynamic"] = benchmark_up_projection_graph( |
| 456 | graph, |
| 457 | operations_per_replay=len(launches), |
| 458 | warmup_replays=args.warmup_replays, |
| 459 | samples=args.samples, |
| 460 | ) |
| 461 | check_up_projection_output( |
| 462 | mailbox[0, :num_tokens, :shard_size], |
| 463 | latent, |
| 464 | weights[-1], |
| 465 | shared_shard[:num_tokens], |
| 466 | ) |
| 467 | if use_skinny: |
| 468 | configs = args.skinny_config or [ |
| 469 | fused_add_multicast_skinny_gemm.config_for_m( |
| 470 | num_tokens, |
| 471 | shard_size, |
| 472 | ) |
| 473 | ] |
| 474 | skinny_results = [] |
| 475 | for config in configs: |
| 476 | skinny_kernel = SkinnyKernel(num_tokens, shard_size, config) |
| 477 | |
| 478 | def launch_skinny( |
| 479 | latent: torch.Tensor, |
| 480 | weight: torch.Tensor, |
| 481 | shared_shard: torch.Tensor, |
| 482 | *, |
| 483 | skinny_kernel: SkinnyKernel = skinny_kernel, |
| 484 | num_tokens: int = num_tokens, |
| 485 | ) -> None: |
| 486 | skinny_kernel.launch( |
| 487 | latent, |
| 488 | weight, |
| 489 | shared_shard[:num_tokens], |
| 490 | mailbox, |
| 491 | ) |
| 492 | |
| 493 | launches = make_up_projection_launches( |
| 494 | launch_skinny, |
| 495 | latent, |
| 496 | weights, |
| 497 | shared_shard, |
| 498 | ) |
| 499 | graph = capture_up_projection_graph(launches) |
| 500 | timing = benchmark_up_projection_graph( |
| 501 | graph, |
| 502 | operations_per_replay=len(launches), |
| 503 | warmup_replays=args.warmup_replays, |
| 504 | samples=args.samples, |
| 505 | ) |
| 506 | check_up_projection_output( |
| 507 | mailbox[0, :num_tokens, :shard_size], |
| 508 | latent, |
| 509 | weights[-1], |
| 510 | shared_shard[:num_tokens], |
| 511 | ) |
| 512 | skinny_results.append( |
| 513 | { |
| 514 | "config": asdict(config), |
| 515 | **timing, |
| 516 | } |
| 517 | ) |
| 518 | result["skinny"] = skinny_results |
| 519 | results.append(result) |
| 520 | |
| 521 | properties = torch.cuda.get_device_properties(device) |
| 522 | report = { |
| 523 | "scope": "up-projection", |
| 524 | "device": properties.name, |
| 525 | "compute_capability": list(torch.cuda.get_device_capability(device)), |
| 526 | "tp_size": args.tp_size, |
| 527 | "shard_size": shard_size, |
| 528 | "weight_count": weight_count, |
| 529 | "cache_multiplier": args.cache_multiplier, |
| 530 | "warmup_replays": args.warmup_replays, |
| 531 | "samples": args.samples, |
| 532 | "results": results, |
| 533 | } |
| 534 | rendered = json.dumps(report, indent=2) |
| 535 | print(rendered, flush=True) |
| 536 | if args.output is not None: |
| 537 | args.output.parent.mkdir(parents=True, exist_ok=True) |
| 538 | args.output.write_text(rendered + "\n", encoding="utf-8") |
| 539 | |
| 540 | |
| 541 | def capture_tail_graph( |
| 542 | operation: Callable[[], torch.Tensor], |
| 543 | cpu_group: dist.ProcessGroup, |
| 544 | ) -> tuple[torch.cuda.CUDAGraph, torch.Tensor]: |
| 545 | for _ in range(3): |
| 546 | dist.barrier(group=cpu_group) |
| 547 | output = operation() |
| 548 | torch.accelerator.synchronize() |
| 549 | |
| 550 | dist.barrier(group=cpu_group) |
| 551 | graph = torch.cuda.CUDAGraph() |
| 552 | with torch.cuda.graph(graph): |
| 553 | output = operation() |
| 554 | torch.accelerator.synchronize() |
| 555 | return graph, output |
| 556 | |
| 557 | |
| 558 | def benchmark_tail_graph( |
| 559 | graph: torch.cuda.CUDAGraph, |
| 560 | *, |
| 561 | warmup_replays: int, |
| 562 | samples: int, |
| 563 | device_group: dist.ProcessGroup, |
| 564 | cpu_group: dist.ProcessGroup, |
| 565 | ) -> dict[str, Any]: |
| 566 | for _ in range(warmup_replays): |
| 567 | graph.replay() |
| 568 | torch.accelerator.synchronize() |
| 569 | |
| 570 | dist.barrier(group=cpu_group) |
| 571 | starts = [torch.cuda.Event(enable_timing=True) for _ in range(samples + 1)] |
| 572 | ends = [torch.cuda.Event(enable_timing=True) for _ in range(samples + 1)] |
| 573 | for start, end in zip(starts, ends): |
| 574 | start.record() |
| 575 | graph.replay() |
| 576 | end.record() |
| 577 | torch.accelerator.synchronize() |
| 578 | |
| 579 | samples_us = torch.tensor( |
| 580 | [start.elapsed_time(end) * 1000.0 for start, end in zip(starts, ends)], |
| 581 | dtype=torch.float64, |
| 582 | device=torch.accelerator.current_device_index(), |
| 583 | ) |
| 584 | dist.all_reduce(samples_us, op=dist.ReduceOp.MAX, group=device_group) |
| 585 | return summarize(samples_us[1:].tolist()) |
| 586 | |
| 587 | |
| 588 | def make_inputs( |
| 589 | num_tokens: int, |
| 590 | rank: int, |
| 591 | device: torch.device, |
| 592 | ) -> tuple[torch.Tensor, torch.Tensor]: |
| 593 | torch.manual_seed(20260726 + 100 * num_tokens + rank) |
| 594 | routed = torch.randn( |
| 595 | (num_tokens, LATENT_SIZE), |
| 596 | dtype=torch.bfloat16, |
| 597 | device=device, |
| 598 | ).mul_(0.01) |
| 599 | shared = torch.randn( |
| 600 | (num_tokens, HIDDEN_SIZE), |
| 601 | dtype=torch.bfloat16, |
| 602 | device=device, |
| 603 | ) |
| 604 | return routed, shared |
| 605 | |
| 606 | |
| 607 | def make_deferred_input( |
| 608 | routed: torch.Tensor, |
| 609 | top_k: int, |
| 610 | ) -> UnfinalizedMoEOutput: |
| 611 | num_tokens = routed.shape[0] |
| 612 | gemm2_permuted = ( |
| 613 | routed[:, None, :] |
| 614 | .expand(num_tokens, top_k, LATENT_SIZE) |
| 615 | .contiguous() |
| 616 | .view(num_tokens * top_k, LATENT_SIZE) |
| 617 | ) |
| 618 | expert_weights = torch.zeros( |
| 619 | num_tokens, |
| 620 | top_k, |
| 621 | dtype=routed.dtype, |
| 622 | device=routed.device, |
| 623 | ) |
| 624 | expert_weights[:, 0] = 1 |
| 625 | expanded_idx = torch.arange( |
| 626 | num_tokens * top_k, |
| 627 | dtype=torch.int32, |
| 628 | device=routed.device, |
| 629 | ).view(num_tokens, top_k) |
| 630 | return UnfinalizedMoEOutput( |
| 631 | gemm2_permuted=gemm2_permuted, |
| 632 | expert_weights=expert_weights, |
| 633 | expanded_idx_to_permuted_idx=expanded_idx, |
| 634 | ) |
| 635 | |
| 636 | |
| 637 | def make_reference( |
| 638 | routed: torch.Tensor, |
| 639 | shared: torch.Tensor, |
| 640 | rms_weight: torch.Tensor, |
| 641 | up_weight: torch.Tensor, |
| 642 | device_group: dist.ProcessGroup, |
| 643 | ) -> Callable[[], torch.Tensor]: |
| 644 | routed_workspace = torch.empty_like(routed) |
| 645 | shared_workspace = torch.empty_like(shared) |
| 646 | |
| 647 | def reference() -> torch.Tensor: |
| 648 | routed_workspace.copy_(routed) |
| 649 | dist.all_reduce(routed_workspace, group=device_group) |
| 650 | normalized = F.rms_norm( |
| 651 | routed_workspace, |
| 652 | (LATENT_SIZE,), |
| 653 | rms_weight, |
| 654 | RMS_EPS, |
| 655 | ) |
| 656 | projected = F.linear(normalized, up_weight) |
| 657 | shared_workspace.copy_(shared) |
| 658 | dist.all_reduce(shared_workspace, group=device_group) |
| 659 | return projected.add(shared_workspace) |
| 660 | |
| 661 | return reference |
| 662 | |
| 663 | |
| 664 | def check_fused_output( |
| 665 | fused_output: torch.Tensor, |
| 666 | reference: Callable[[], torch.Tensor], |
| 667 | cpu_group: dist.ProcessGroup, |
| 668 | ) -> None: |
| 669 | dist.barrier(group=cpu_group) |
| 670 | expected = reference() |
| 671 | torch.testing.assert_close(fused_output, expected, atol=8e-2, rtol=3e-2) |
| 672 | |
| 673 | |
| 674 | def benchmark_whole_tail(args: argparse.Namespace) -> None: |
| 675 | if any(not 1 <= num_tokens <= 16 for num_tokens in args.num_tokens): |
| 676 | raise ValueError("--num-tokens values must be in [1, 16]") |
| 677 | if args.warmup_replays < 0 or args.samples <= 0: |
| 678 | raise ValueError("warmup replays must be nonnegative and samples positive") |
| 679 | if args.top_k <= 0: |
| 680 | raise ValueError("--top-k must be positive") |
| 681 | if args.skinny_max_num_tokens is not None and any( |
| 682 | not 0 <= cutoff <= 8 for cutoff in args.skinny_max_num_tokens |
| 683 | ): |
| 684 | raise ValueError("--skinny-max-num-tokens must be in [0, 8]") |
| 685 | skinny_configs = dict(args.skinny_config or ()) |
| 686 | if len(skinny_configs) != len(args.skinny_config or ()): |
| 687 | raise ValueError("--skinny-config must not repeat an M value") |
| 688 | if any(not 1 <= num_tokens <= 8 for num_tokens in skinny_configs): |
| 689 | raise ValueError("--skinny-config M values must be in [1, 8]") |
| 690 | if not {"RANK", "WORLD_SIZE", "LOCAL_RANK"} <= os.environ.keys(): |
| 691 | raise RuntimeError("launch this benchmark with torchrun") |
| 692 | |
| 693 | rank = int(os.environ["RANK"]) |
| 694 | world_size = int(os.environ["WORLD_SIZE"]) |
| 695 | local_rank = int(os.environ["LOCAL_RANK"]) |
| 696 | device = torch.device("cuda", local_rank) |
| 697 | torch.accelerator.set_device_index(device) |
| 698 | init_distributed_environment() |
| 699 | if world_size > 8: |
| 700 | set_custom_all_reduce(False) |
| 701 | initialize_model_parallel(tensor_model_parallel_size=world_size) |
| 702 | device_group = get_tp_group().device_group |
| 703 | cpu_group = dist.new_group(backend="gloo") |
| 704 | |
| 705 | if torch.cuda.get_device_capability(device)[0] != 10: |
| 706 | raise RuntimeError("Kimi K3 latent-MoE tail requires SM100") |
| 707 | |
| 708 | torch.manual_seed(20260726) |
| 709 | rms_weight = 1 + 0.1 * torch.randn( |
| 710 | LATENT_SIZE, |
| 711 | dtype=torch.bfloat16, |
| 712 | device=device, |
| 713 | ) |
| 714 | up_weight = ( |
| 715 | torch.randn( |
| 716 | (HIDDEN_SIZE, LATENT_SIZE), |
| 717 | dtype=torch.bfloat16, |
| 718 | device=device, |
| 719 | ) |
| 720 | / LATENT_SIZE**0.5 |
| 721 | ) |
| 722 | |
| 723 | use_reference = args.backend in ("reference", "both") |
| 724 | use_fused = args.backend in ("fused", "both") |
| 725 | routed_input_modes = ( |
| 726 | ("finalized", "deferred") |
| 727 | if args.routed_input == "both" |
| 728 | else (args.routed_input,) |
| 729 | ) |
| 730 | fused_ops = [] |
| 731 | if use_fused: |
| 732 | production_config_for_m = fused_add_multicast_skinny_gemm.config_for_m |
| 733 | |
| 734 | def config_for_m( |
| 735 | num_rows: int, |
| 736 | shard_dim: int = 896, |
| 737 | ) -> fused_add_multicast_skinny_gemm.SkinnyConfig: |
| 738 | config = skinny_configs.get(num_rows) |
| 739 | if config is not None: |
| 740 | return config |
| 741 | return production_config_for_m(num_rows, shard_dim) |
| 742 | |
| 743 | fused_add_multicast_skinny_gemm.config_for_m = config_for_m |
| 744 | cutoffs = args.skinny_max_num_tokens or [latent_moe_tail._SKINNY_MAX_NUM_TOKENS] |
| 745 | for cutoff in cutoffs: |
| 746 | latent_moe_tail._SKINNY_MAX_NUM_TOKENS = cutoff |
| 747 | latent_moe_tail.KimiK3LatentMoETailOp._instances.clear() |
| 748 | for routed_input_mode in routed_input_modes: |
| 749 | fused_ops.append( |
| 750 | ( |
| 751 | cutoff, |
| 752 | routed_input_mode, |
| 753 | latent_moe_tail.KimiK3LatentMoETailOp.initialize( |
| 754 | hidden_size=HIDDEN_SIZE, |
| 755 | latent_size=LATENT_SIZE, |
| 756 | dtype=torch.bfloat16, |
| 757 | device=device, |
| 758 | rms_eps=RMS_EPS, |
| 759 | experts_per_token=( |
| 760 | args.top_k if routed_input_mode == "deferred" else 0 |
| 761 | ), |
| 762 | ), |
| 763 | ) |
| 764 | ) |
| 765 | cutedsl_warmup() |
| 766 | |
| 767 | results = [] |
| 768 | for num_tokens in args.num_tokens: |
| 769 | routed, shared = make_inputs(num_tokens, rank, device) |
| 770 | routed_inputs: dict[str, torch.Tensor | UnfinalizedMoEOutput] = { |
| 771 | "finalized": routed |
| 772 | } |
| 773 | if "deferred" in routed_input_modes: |
| 774 | routed_inputs["deferred"] = make_deferred_input(routed, args.top_k) |
| 775 | reference = make_reference( |
| 776 | routed, |
| 777 | shared, |
| 778 | rms_weight, |
| 779 | up_weight, |
| 780 | device_group, |
| 781 | ) |
| 782 | result: dict[str, Any] = {"num_tokens": num_tokens} |
| 783 | if use_reference: |
| 784 | reference_graph, _ = capture_tail_graph(reference, cpu_group) |
| 785 | result["reference"] = benchmark_tail_graph( |
| 786 | reference_graph, |
| 787 | warmup_replays=args.warmup_replays, |
| 788 | samples=args.samples, |
| 789 | device_group=device_group, |
| 790 | cpu_group=cpu_group, |
| 791 | ) |
| 792 | for cutoff, routed_input_mode, fused_op in fused_ops: |
| 793 | routed_input = routed_inputs[routed_input_mode] |
| 794 | |
| 795 | def fused( |
| 796 | routed_input: torch.Tensor | UnfinalizedMoEOutput = routed_input, |
| 797 | shared: torch.Tensor = shared, |
| 798 | fused_op: latent_moe_tail.KimiK3LatentMoETailOp = fused_op, |
| 799 | ) -> torch.Tensor: |
| 800 | return fused_op(routed_input, shared, rms_weight, up_weight) |
| 801 | |
| 802 | fused_graph, fused_output = capture_tail_graph(fused, cpu_group) |
| 803 | fused_key_parts = ["fused"] |
| 804 | if routed_input_mode != "finalized" or len(routed_input_modes) > 1: |
| 805 | fused_key_parts.append(routed_input_mode) |
| 806 | if len(cutoffs) > 1: |
| 807 | fused_key_parts.extend(("skinny", "max", str(cutoff))) |
| 808 | fused_key = "_".join(fused_key_parts) |
| 809 | result[fused_key] = benchmark_tail_graph( |
| 810 | fused_graph, |
| 811 | warmup_replays=args.warmup_replays, |
| 812 | samples=args.samples, |
| 813 | device_group=device_group, |
| 814 | cpu_group=cpu_group, |
| 815 | ) |
| 816 | check_fused_output(fused_output, reference, cpu_group) |
| 817 | if "reference" in result: |
| 818 | speedup = ( |
| 819 | result["reference"]["median_us"] / result[fused_key]["median_us"] |
| 820 | ) |
| 821 | if fused_key == "fused": |
| 822 | result["speedup"] = speedup |
| 823 | else: |
| 824 | result[f"{fused_key}_speedup"] = speedup |
| 825 | results.append(result) |
| 826 | |
| 827 | properties = torch.cuda.get_device_properties(device) |
| 828 | report = { |
| 829 | "scope": "whole-tail", |
| 830 | "device": properties.name, |
| 831 | "compute_capability": list(torch.cuda.get_device_capability(device)), |
| 832 | "world_size": world_size, |
| 833 | "torch_version": torch.__version__, |
| 834 | "cuda_version": torch.version.cuda, |
| 835 | "warmup_replays": args.warmup_replays, |
| 836 | "samples": args.samples, |
| 837 | "routed_input": args.routed_input, |
| 838 | "top_k": args.top_k, |
| 839 | "skinny_max_num_tokens": sorted({cutoff for cutoff, _, _ in fused_ops}), |
| 840 | "skinny_configs": { |
| 841 | str(num_tokens): asdict(config) |
| 842 | for num_tokens, config in skinny_configs.items() |
| 843 | }, |
| 844 | "timing_scope": { |
| 845 | "reference": ( |
| 846 | "two input copies, two AllReduces, RMSNorm, full replicated " |
| 847 | "up-projection GEMM, and final add" |
| 848 | ), |
| 849 | "fused": ( |
| 850 | "routed AllReduce/RMSNorm plus shared ReduceScatter, sharded " |
| 851 | "up-projection/multicast, and Lamport copy" |
| 852 | ), |
| 853 | }, |
| 854 | "results": results, |
| 855 | } |
| 856 | if rank == 0: |
| 857 | rendered = json.dumps(report, indent=2) |
| 858 | print(rendered, flush=True) |
| 859 | if args.output is not None: |
| 860 | args.output.parent.mkdir(parents=True, exist_ok=True) |
| 861 | args.output.write_text(rendered + "\n", encoding="utf-8") |
| 862 | |
| 863 | dist.barrier(group=cpu_group) |
| 864 | |
| 865 | |
| 866 | def main() -> None: |
| 867 | args = parse_args() |
| 868 | if args.scope == "up-projection": |
| 869 | benchmark_up_projection(args) |
| 870 | return |
| 871 | |
| 872 | from vllm.config import VllmConfig, set_current_vllm_config |
| 873 | |
| 874 | with set_current_vllm_config(VllmConfig()): |
| 875 | benchmark_whole_tail(args) |
| 876 | |
| 877 | |
| 878 | if __name__ == "__main__": |
| 879 | main() |