[Bugfix] Fix triton import with local TritonPlaceholder

Signed-off-by: Mengqing Cao <[email protected]>

Author: MengqingCao

benchmarks/kernels/benchmark_moe.py

@@ -10,12 +10,12 @@
 
 import ray
 import torch
-import triton
 from ray.experimental.tqdm_ray import tqdm
 from transformers import AutoConfig
 
 from vllm.model_executor.layers.fused_moe.fused_moe import *
 from vllm.platforms import current_platform
+from vllm.triton_utils import triton
 from vllm.utils import FlexibleArgumentParser
 
 FP8_DTYPE = current_platform.fp8_dtype()

benchmarks/kernels/benchmark_rmsnorm.py

@@ -4,11 +4,11 @@
 from typing import Optional, Union
 
 import torch
-import triton
 from flashinfer.norm import fused_add_rmsnorm, rmsnorm
 from torch import nn
 
 from vllm import _custom_ops as vllm_ops
+from vllm.triton_utils import triton
 
 
 class HuggingFaceRMSNorm(nn.Module):

benchmarks/kernels/deepgemm/benchmark_fp8_block_dense_gemm.py

@@ -6,13 +6,13 @@
 # Import DeepGEMM functions
 import deep_gemm
 import torch
-import triton
 from deep_gemm import calc_diff, ceil_div, get_col_major_tma_aligned_tensor
 
 # Import vLLM functions
 from vllm import _custom_ops as ops
 from vllm.model_executor.layers.quantization.utils.fp8_utils import (
     per_token_group_quant_fp8, w8a8_block_fp8_matmul)
+from vllm.triton_utils import triton
 
 
 # Copied from

tests/kernels/attention/test_flashmla.py

@@ -5,11 +5,11 @@
 
 import pytest
 import torch
-import triton
 
 from vllm.attention.ops.flashmla import (flash_mla_with_kvcache,
                                          get_mla_metadata,
                                          is_flashmla_supported)
+from vllm.triton_utils import triton
 
 
 def cal_diff(x: torch.Tensor, y: torch.Tensor, name: str) -> None:

vllm/attention/ops/blocksparse_attention/blocksparse_attention_kernel.py

@@ -1,8 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import torch
-import triton
-import triton.language as tl
+
+from vllm.triton_utils import tl, triton
 
 
 def blocksparse_flash_attn_varlen_fwd(

vllm/attention/ops/blocksparse_attention/utils.py

@@ -8,7 +8,8 @@
 
 import numpy as np
 import torch
-import triton
+
+from vllm.triton_utils import triton
 
 
 class csr_matrix:

vllm/attention/ops/chunked_prefill_paged_decode.py

@@ -7,11 +7,10 @@
 #  - Thomas Parnell <[email protected]>
 
 import torch
-import triton
-import triton.language as tl
 
 from vllm import _custom_ops as ops
 from vllm.platforms.rocm import use_rocm_custom_paged_attention
+from vllm.triton_utils import tl, triton
 
 from .prefix_prefill import context_attention_fwd
 

vllm/attention/ops/prefix_prefill.py

@@ -4,10 +4,9 @@
 # https://github.com/ModelTC/lightllm/blob/main/lightllm/models/llama/triton_kernel/context_flashattention_nopad.py
 
 import torch
-import triton
-import triton.language as tl
 
 from vllm.platforms import current_platform
+from vllm.triton_utils import tl, triton
 
 # Static kernels parameters
 BASE_BLOCK = 128 if current_platform.has_device_capability(80) else 64

vllm/attention/ops/triton_decode_attention.py

@@ -30,10 +30,8 @@
 
 import logging
 
-import triton
-import triton.language as tl
-
 from vllm.platforms import current_platform
+from vllm.triton_utils import tl, triton
 
 is_hip_ = current_platform.is_rocm()
 

vllm/attention/ops/triton_flash_attention.py

@@ -25,11 +25,10 @@
 from typing import Optional
 
 import torch
-import triton
-import triton.language as tl
 
 from vllm import _custom_ops as ops
 from vllm.platforms import current_platform
+from vllm.triton_utils import tl, triton
 
 SUPPORTED_LAYOUTS = ['thd', 'bhsd', 'bshd']
 

vllm/attention/ops/triton_merge_attn_states.py

@@ -2,8 +2,8 @@
 from typing import Optional
 
 import torch
-import triton
-import triton.language as tl
+
+from vllm.triton_utils import tl, triton
 
 
 # Implements section 2.2 of https://www.arxiv.org/pdf/2501.01005

vllm/lora/ops/triton_ops/kernel_utils.py

@@ -2,8 +2,7 @@
 """
 Utilities for Punica kernel construction.
 """
-import triton
-import triton.language as tl
+from vllm.triton_utils import tl, triton
 
 
 @triton.jit

vllm/model_executor/layers/fused_moe/fused_moe.py

@@ -6,8 +6,6 @@
 from typing import Any, Callable, Dict, List, Optional, Tuple
 
 import torch
-import triton
-import triton.language as tl
 
 import vllm.envs as envs
 from vllm import _custom_ops as ops
@@ -21,6 +19,7 @@
 from vllm.model_executor.layers.quantization.utils.int8_utils import (
     per_token_group_quant_int8, per_token_quant_int8)
 from vllm.platforms import current_platform
+from vllm.triton_utils import tl, triton
 from vllm.utils import direct_register_custom_op
 
 from .rocm_aiter_fused_moe import is_rocm_aiter_moe_enabled

vllm/model_executor/layers/fused_moe/moe_align_block_size.py

@@ -2,11 +2,10 @@
 from typing import Optional, Tuple
 
 import torch
-import triton
-import triton.language as tl
 
 import vllm.envs as envs
 from vllm import _custom_ops as ops
+from vllm.triton_utils import tl, triton
 from vllm.utils import round_up
 
 

vllm/model_executor/layers/lightning_attn.py

@@ -1,9 +1,9 @@
 # SPDX-License-Identifier: Apache-2.0
 import torch
-import triton
-import triton.language as tl
 from einops import rearrange
 
+from vllm.triton_utils import tl, triton
+
 
 @triton.jit
 def _fwd_diag_kernel(Q, K, V, Out, S, b: tl.constexpr, h: tl.constexpr, n,

vllm/model_executor/layers/mamba/ops/mamba_ssm.py

@@ -4,13 +4,11 @@
 # Adapted from https://github.com/state-spaces/mamba/blob/v2.2.4/mamba_ssm/ops/triton/selective_state_update.py
 
 import torch
-import triton
-import triton.language as tl
 from packaging import version
 
 from vllm import _custom_ops as ops
 from vllm.attention.backends.utils import PAD_SLOT_ID
-from vllm.triton_utils import HAS_TRITON
+from vllm.triton_utils import HAS_TRITON, tl, triton
 
 TRITON3 = HAS_TRITON and (version.parse(triton.__version__)
                           >= version.parse("3.0.0"))

vllm/model_executor/layers/mamba/ops/ssd_bmm.py

@@ -8,8 +8,8 @@
 import math
 
 import torch
-import triton
-import triton.language as tl
+
+from vllm.triton_utils import tl, triton
 
 
 @triton.autotune(

vllm/model_executor/layers/mamba/ops/ssd_chunk_scan.py

@@ -6,10 +6,10 @@
 # ruff: noqa: E501,SIM102
 
 import torch
-import triton
-import triton.language as tl
 from packaging import version
 
+from vllm.triton_utils import tl, triton
+
 TRITON_22 = version.parse(triton.__version__) >= version.parse('2.2.0')
 
 

vllm/model_executor/layers/mamba/ops/ssd_chunk_state.py

@@ -8,8 +8,8 @@
 import math
 
 import torch
-import triton
-import triton.language as tl
+
+from vllm.triton_utils import tl, triton
 
 from .mamba_ssm import softplus
 

vllm/model_executor/layers/mamba/ops/ssd_combined.py

@@ -6,10 +6,11 @@
 # ruff: noqa: E501
 
 import torch
-import triton
 from einops import rearrange
 from packaging import version
 
+from vllm.triton_utils import triton
+
 from .ssd_bmm import _bmm_chunk_fwd
 from .ssd_chunk_scan import _chunk_scan_fwd
 from .ssd_chunk_state import (_chunk_cumsum_fwd, _chunk_state_fwd,

vllm/model_executor/layers/mamba/ops/ssd_state_passing.py

@@ -6,8 +6,8 @@
 # ruff: noqa: E501
 
 import torch
-import triton
-import triton.language as tl
+
+from vllm.triton_utils import tl, triton
 
 
 @triton.autotune(

vllm/model_executor/layers/quantization/awq_triton.py

@@ -1,8 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import torch
-import triton
-import triton.language as tl
+
+from vllm.triton_utils import tl, triton
 
 AWQ_TRITON_SUPPORTED_GROUP_SIZES = [-1, 32, 64, 128]
 

vllm/model_executor/layers/quantization/compressed_tensors/triton_scaled_mm.py

@@ -3,8 +3,8 @@
 from typing import Optional, Type
 
 import torch
-import triton
-import triton.language as tl
+
+from vllm.triton_utils import tl, triton
 
 
 def is_weak_contiguous(x: torch.Tensor):

vllm/model_executor/layers/quantization/utils/fp8_utils.py

@@ -7,8 +7,6 @@
 from typing import Any, Dict, List, Optional, Tuple, Union
 
 import torch
-import triton
-import triton.language as tl
 
 from vllm import _custom_ops as ops
 from vllm.logger import init_logger
@@ -17,6 +15,7 @@
 from vllm.model_executor.layers.quantization.utils.w8a8_utils import (
     CUTLASS_BLOCK_FP8_SUPPORTED)
 from vllm.platforms import current_platform
+from vllm.triton_utils import tl, triton
 from vllm.utils import direct_register_custom_op
 
 logger = init_logger(__name__)

vllm/model_executor/layers/quantization/utils/int8_utils.py

@@ -8,10 +8,9 @@
 from typing import Any, Dict, List, Optional, Tuple
 
 import torch
-import triton
-import triton.language as tl
 
 from vllm.platforms import current_platform
+from vllm.triton_utils import tl, triton
 
 logger = logging.getLogger(__name__)
 

vllm/triton_utils/__init__.py

@@ -1,5 +1,13 @@
 # SPDX-License-Identifier: Apache-2.0
 
-from vllm.triton_utils.importing import HAS_TRITON
+from vllm.triton_utils.importing import (HAS_TRITON, TritonLanguagePlaceholder,
+                                         TritonPlaceholder)
 
-__all__ = ["HAS_TRITON"]
+if HAS_TRITON:
+    import triton
+    import triton.language as tl
+else:
+    triton = TritonPlaceholder()
+    tl = TritonLanguagePlaceholder()
+
+__all__ = ["HAS_TRITON", "triton", "tl"]

vllm/triton_utils/importing.py

@@ -16,32 +16,34 @@
     logger.info("Triton not installed or not compatible; certain GPU-related"
                 " functions will not be available.")
 
-    class TritonPlaceholder(types.ModuleType):
-
-        def __init__(self):
-            super().__init__("triton")
-            self.jit = self._dummy_decorator("jit")
-            self.autotune = self._dummy_decorator("autotune")
-            self.heuristics = self._dummy_decorator("heuristics")
-            self.language = TritonLanguagePlaceholder()
-            logger.warning_once(
-                "Triton is not installed. Using dummy decorators. "
-                "Install it via `pip install triton` to enable kernel"
-                "compilation.")
-
-        def _dummy_decorator(self, name):
-
-            def decorator(func=None, **kwargs):
-                if func is None:
-                    return lambda f: f
-                return func
-
-            return decorator
-
-    class TritonLanguagePlaceholder(types.ModuleType):
-
-        def __init__(self):
-            super().__init__("triton.language")
-            self.constexpr = None
-            self.dtype = None
-            self.int64 = None
+
+class TritonPlaceholder(types.ModuleType):
+
+    def __init__(self):
+        super().__init__("triton")
+        self.jit = self._dummy_decorator("jit")
+        self.autotune = self._dummy_decorator("autotune")
+        self.heuristics = self._dummy_decorator("heuristics")
+        self.language = TritonLanguagePlaceholder()
+        logger.warning_once(
+            "Triton is not installed. Using dummy decorators. "
+            "Install it via `pip install triton` to enable kernel"
+            " compilation.")
+
+    def _dummy_decorator(self, name):
+
+        def decorator(func=None, **kwargs):
+            if func is None:
+                return lambda f: f
+            return func
+
+        return decorator
+
+
+class TritonLanguagePlaceholder(types.ModuleType):
+
+    def __init__(self):
+        super().__init__("triton.language")
+        self.constexpr = None
+        self.dtype = None
+        self.int64 = None