diff --git a/benchmarks/kernels/benchmark_flydsl_moe_w4a16.py b/benchmarks/kernels/benchmark_flydsl_moe_w4a16.py index 9e4f4157a8a..09a01b301be 100644 --- a/benchmarks/kernels/benchmark_flydsl_moe_w4a16.py +++ b/benchmarks/kernels/benchmark_flydsl_moe_w4a16.py @@ -17,7 +17,7 @@ from vllm.model_executor.layers.fused_moe.fused_flydsl_moe import fused_flydsl_m from vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe import ( # noqa: E501 compressed_tensors_moe_w4a16_flydsl, ) -from vllm.platforms import current_platform +from vllm.utils.platform_utils import get_device_name_as_file_name RoutingBuffers = tuple[ torch.Tensor, # sorted_token_ids @@ -259,7 +259,7 @@ def tune_flydsl_moe_w4a16( ) us_best = us tuned_config[str(num_tokens)] = tile_config - device_name = current_platform.get_device_name().replace(" ", "_") + device_name = get_device_name_as_file_name() tuned_config_file_name = ( f"E={num_experts},N={inter_dim},device_name={device_name}," f"dtype=int4_w4a16,backend=flydsl.json" diff --git a/benchmarks/kernels/benchmark_w8a8_block_fp8.py b/benchmarks/kernels/benchmark_w8a8_block_fp8.py index 36dce1b6388..590d4cfdc6d 100644 --- a/benchmarks/kernels/benchmark_w8a8_block_fp8.py +++ b/benchmarks/kernels/benchmark_w8a8_block_fp8.py @@ -19,6 +19,7 @@ from vllm.model_executor.layers.quantization.utils.fp8_utils import ( from vllm.platforms import current_platform from vllm.triton_utils import triton from vllm.utils.argparse_utils import FlexibleArgumentParser +from vllm.utils.platform_utils import get_device_name_as_file_name mp.set_start_method("spawn", force=True) @@ -264,7 +265,7 @@ def save_configs( input_type="fp8", ) -> None: os.makedirs(save_path, exist_ok=True) - device_name = current_platform.get_device_name().replace(" ", "_") + device_name = get_device_name_as_file_name() json_file_name = ( f"N={N},K={K},device_name={device_name},dtype={input_type}_w8a8," f"block_shape=[{block_n},{block_k}].json" diff --git a/vllm/model_executor/layers/fused_moe/fused_flydsl_moe.py b/vllm/model_executor/layers/fused_moe/fused_flydsl_moe.py index cf49e01e628..b1588c4a2a4 100644 --- a/vllm/model_executor/layers/fused_moe/fused_flydsl_moe.py +++ b/vllm/model_executor/layers/fused_moe/fused_flydsl_moe.py @@ -15,7 +15,7 @@ from aiter.ops.flydsl.kernels.moe_gemm_2stage import ( ) from vllm.logger import init_logger -from vllm.platforms import current_platform +from vllm.utils.platform_utils import get_device_name_as_file_name from vllm.utils.torch_utils import direct_register_custom_op logger = init_logger(__name__) @@ -106,7 +106,7 @@ def build_routing_buffers( @functools.lru_cache def try_get_optimal_config(num_experts, inter_dim): - device_name = current_platform.get_device_name().replace(" ", "_") + device_name = get_device_name_as_file_name() json_file_name = ( f"E={num_experts},N={inter_dim},device_name={device_name}," "dtype=int4_w4a16,backend=flydsl.json" diff --git a/vllm/model_executor/layers/fused_moe/fused_moe.py b/vllm/model_executor/layers/fused_moe/fused_moe.py index 269b6e3da0b..62cf12ea24e 100644 --- a/vllm/model_executor/layers/fused_moe/fused_moe.py +++ b/vllm/model_executor/layers/fused_moe/fused_moe.py @@ -31,6 +31,7 @@ from vllm.model_executor.layers.fused_moe.utils import ( ) from vllm.platforms import current_platform from vllm.triton_utils import tl, triton +from vllm.utils.platform_utils import get_device_name_as_file_name from vllm.utils.torch_utils import direct_register_custom_op logger = init_logger(__name__) @@ -1034,7 +1035,7 @@ def zero_experts_compute_triton( def get_config_file_name( E: int, N: int, dtype: str | None, block_shape: list[int] | None = None ) -> str: - device_name = current_platform.get_device_name().replace(" ", "_") + device_name = get_device_name_as_file_name() # Set device_name to H200 if a device from the H200 family is detected if "H200" in device_name.split("_"): device_name = "NVIDIA_H200" diff --git a/vllm/model_executor/layers/mamba/ops/mamba_ssm.py b/vllm/model_executor/layers/mamba/ops/mamba_ssm.py index 8c5a6355803..d348defcc76 100644 --- a/vllm/model_executor/layers/mamba/ops/mamba_ssm.py +++ b/vllm/model_executor/layers/mamba/ops/mamba_ssm.py @@ -19,6 +19,7 @@ from vllm.logger import init_logger from vllm.model_executor.layers.mamba.ops.triton_helpers import fast_exp from vllm.platforms import current_platform from vllm.triton_utils import HAS_TRITON, tl, triton +from vllm.utils.platform_utils import get_device_name_as_file_name from vllm.v1.attention.backends.utils import NULL_BLOCK_ID if current_platform.is_xpu(): @@ -53,7 +54,7 @@ def get_ssm_config_file_name( def get_ssm_device_name() -> str: - return current_platform.get_device_name().replace(" ", "_") + return get_device_name_as_file_name() def _canonical_cache_dtype(cache_dtype: str) -> str: diff --git a/vllm/model_executor/layers/quantization/utils/fp8_utils.py b/vllm/model_executor/layers/quantization/utils/fp8_utils.py index 32a2d86899c..49eadc5a152 100644 --- a/vllm/model_executor/layers/quantization/utils/fp8_utils.py +++ b/vllm/model_executor/layers/quantization/utils/fp8_utils.py @@ -33,6 +33,7 @@ from vllm.utils.deep_gemm import ( is_deep_gemm_e8m0_used, transform_sf_into_required_layout, ) +from vllm.utils.platform_utils import get_device_name_as_file_name from vllm.utils.torch_utils import direct_register_custom_op logger = init_logger(__name__) @@ -864,7 +865,7 @@ def get_w8a8_block_fp8_configs( # First look up if an optimized configuration is available in the configs # directory - device_name = current_platform.get_device_name().replace(" ", "_") + device_name = get_device_name_as_file_name() json_file_name = f"N={N},K={K},device_name={device_name},dtype=fp8_w8a8,block_shape=[{block_n},{block_k}].json" # noqa: E501 config_file_path = os.path.join( diff --git a/vllm/model_executor/layers/quantization/utils/int8_utils.py b/vllm/model_executor/layers/quantization/utils/int8_utils.py index e0db2526948..4f624cf4963 100644 --- a/vllm/model_executor/layers/quantization/utils/int8_utils.py +++ b/vllm/model_executor/layers/quantization/utils/int8_utils.py @@ -12,6 +12,7 @@ import torch from vllm.platforms import current_platform from vllm.triton_utils import tl, triton +from vllm.utils.platform_utils import get_device_name_as_file_name logger = logging.getLogger(__name__) @@ -329,7 +330,7 @@ def get_w8a8_block_int8_configs( # First look up if an optimized configuration is available in the configs # directory - device_name = current_platform.get_device_name().replace(" ", "_") + device_name = get_device_name_as_file_name() json_file_name = f"N={N},K={K},device_name={device_name},dtype=int8_w8a8,block_shape=[{block_n}, {block_k}].json" # noqa: E501 config_file_path = os.path.join( diff --git a/vllm/utils/platform_utils.py b/vllm/utils/platform_utils.py index cc69d9a241c..5d7fed3c990 100644 --- a/vllm/utils/platform_utils.py +++ b/vllm/utils/platform_utils.py @@ -7,6 +7,7 @@ from concurrent.futures.process import ProcessPoolExecutor from functools import cache from typing import Any +import regex as re import torch @@ -62,3 +63,12 @@ def num_compute_units(device_id: int = 0) -> int: from vllm.platforms import current_platform return current_platform.num_compute_units(device_id) + + +@cache +def get_device_name_as_file_name(device_id: int = 0) -> str: + from vllm.platforms import current_platform + + name = current_platform.get_device_name(device_id) + name = re.sub(r"[\s/]+", "_", name) + return name