| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | |
| 4 | from __future__ import annotations |
| 5 | |
| 6 | import functools |
| 7 | from collections.abc import Callable |
| 8 | from functools import cache |
| 9 | from typing import TYPE_CHECKING, Any |
| 10 | |
| 11 | from vllm.platforms import current_platform |
| 12 | from vllm.utils.import_utils import has_tilelang |
| 13 | |
| 14 | if TYPE_CHECKING or current_platform.is_cuda(): |
| 15 | if not has_tilelang(): |
| 16 | raise ImportError( |
| 17 | "tilelang is required for mhc but is not installed. Install it with " |
| 18 | "`pip install tilelang`." |
| 19 | ) |
| 20 | import tilelang |
| 21 | import tilelang.language as T |
| 22 | else: |
| 23 | tilelang = None # type: ignore[assignment] |
| 24 | T = None # type: ignore[assignment] |
| 25 | |
| 26 | |
| 27 | def _ensure_tilelang_imported() -> None: |
| 28 | """Bind the `tilelang` and `T` module globals, importing them if needed. |
| 29 | |
| 30 | On ROCm, this runs on the first kernel call instead of at import time. |
| 31 | |
| 32 | Raises: |
| 33 | ImportError: If TileLang is not installed. |
| 34 | |
| 35 | """ |
| 36 | global T, tilelang |
| 37 | |
| 38 | if tilelang is not None: |
| 39 | return |
| 40 | if not has_tilelang(): |
| 41 | raise ImportError( |
| 42 | "tilelang is required for mhc but is not installed. Install it with " |
| 43 | "`pip install tilelang`." |
| 44 | ) |
| 45 | import tilelang as tilelang_module |
| 46 | import tilelang.language as tilelang_language |
| 47 | |
| 48 | tilelang = tilelang_module |
| 49 | T = tilelang_language |
| 50 | |
| 51 | |
| 52 | @cache |
| 53 | def _get_pass_configs() -> dict[Any, Any]: |
| 54 | _ensure_tilelang_imported() |
| 55 | pass_configs: dict[Any, Any] = { |
| 56 | tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True, |
| 57 | tilelang.PassConfigKey.TL_DISABLE_TMA_LOWER: True, |
| 58 | } |
| 59 | if current_platform.is_cuda(): |
| 60 | pass_configs[tilelang.PassConfigKey.TL_PTXAS_REGISTER_USAGE_LEVEL] = 10 |
| 61 | return pass_configs |
| 62 | |
| 63 | |
| 64 | class _DeferredTileLangJitKernel: |
| 65 | """Stand-in for a `tilelang.jit` kernel that decorates on first use. |
| 66 | |
| 67 | Both attribute access and calling apply the decoration. Required for |
| 68 | compile-only JIT warmup on platforms that defer import of tilelang. |
| 69 | """ |
| 70 | |
| 71 | _kernel_function: Callable[..., Any] | None = None |
| 72 | _jit_kernel: Any = None |
| 73 | |
| 74 | def __init__(self, kernel_function: Callable[..., Any]) -> None: |
| 75 | self._kernel_function = kernel_function |
| 76 | functools.update_wrapper(self, kernel_function) |
| 77 | |
| 78 | def _ensure_jit_kernel(self) -> Any: |
| 79 | if self._jit_kernel is None: |
| 80 | _ensure_tilelang_imported() |
| 81 | kernel_function = self._kernel_function |
| 82 | assert kernel_function is not None |
| 83 | kernel_function.__globals__["tilelang"] = tilelang |
| 84 | kernel_function.__globals__["T"] = T |
| 85 | self._jit_kernel = tilelang.jit(pass_configs=_get_pass_configs())( |
| 86 | kernel_function |
| 87 | ) |
| 88 | return self._jit_kernel |
| 89 | |
| 90 | def __call__(self, *args: Any, **kwargs: Any) -> Any: |
| 91 | return self._ensure_jit_kernel()(*args, **kwargs) |
| 92 | |
| 93 | def __getattr__(self, name: str) -> Any: |
| 94 | return getattr(self._ensure_jit_kernel(), name) |
| 95 | |
| 96 | |
| 97 | def tilelang_jit(kernel_function: Callable[..., Any]) -> Callable[..., Any]: |
| 98 | """Apply `tilelang.jit`, deferring until first use on ROCm. |
| 99 | |
| 100 | ROCm defers JIT decoration so importing the caller's module does not |
| 101 | require TileLang immediately. CUDA keeps the eager decoration behavior. |
| 102 | |
| 103 | The kernel body parsed by TileLang references `T` as an unqualified |
| 104 | global, so on the deferred ROCm path this rebinds `T`/`tilelang` in the |
| 105 | decorated function's own module globals once they become available. |
| 106 | """ |
| 107 | if not current_platform.is_rocm(): |
| 108 | _ensure_tilelang_imported() |
| 109 | return tilelang.jit(pass_configs=_get_pass_configs())(kernel_function) |
| 110 | |
| 111 | return _DeferredTileLangJitKernel(kernel_function) |