A high-throughput and memory-efficient inference and serving engine for LLMs
Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git
Synced Nov 13, 08:15 UTC
[XPU] Support Triton path for LoRA operations on XPU (#28511)
Signed-off-by: Fanli Lin <fanli.lin@intel.com>
Fanli Lin committed
Nov 13, 2025 at 13:31 UTC
dbbe0c756a41e5a64d6e364c131fd7d12a56b926
3 files changed
+7
-1
vllm/lora/ops/triton_ops/lora_expand_op.py
+1
@@ -48,6 +48,7 @@ def _lora_expand_kernel(
48
SLICE_NUM: tl.constexpr,
49
SAME_STRIDE: tl.constexpr,
50
USE_GDC: tl.constexpr,
51
+ launch_pdl: tl.constexpr,
52
):
53
cta_n_num = tl.cdiv(N, BLOCK_N)
54
cta_m_num = tl.cdiv(M, BLOCK_M)
vllm/lora/ops/triton_ops/lora_shrink_op.py
+1
@@ -46,6 +46,7 @@ def _lora_shrink_kernel(
46
GROUP_SIZE_M: tl.constexpr,
47
SLICE_NUM: tl.constexpr,
48
USE_GDC: tl.constexpr,
49
+ launch_pdl: tl.constexpr,
50
):
51
cta_n_num = tl.cdiv(N, BLOCK_N)
52
cta_m_num = tl.cdiv(M, BLOCK_M)
vllm/platforms/xpu.py
+5
-1
@@ -101,7 +101,11 @@ class XPUPlatform(Platform):
101
102
@classmethod
103
def get_punica_wrapper(cls) -> str:
104
- return "vllm.lora.punica_wrapper.punica_xpu.PunicaWrapperXPU"
104
+ xpu_use_triton_kernel = os.getenv("XPU_USE_TRITON_KERNEL", "0") == "1"
105
+ if not xpu_use_triton_kernel:
106
+ return "vllm.lora.punica_wrapper.punica_xpu.PunicaWrapperXPU"
107
+ else:
108
+ return "vllm.lora.punica_wrapper.punica_gpu.PunicaWrapperGPU"
109
110
@classmethod
111
def get_device_total_memory(cls, device_id: int = 0) -> int: