From 523d74e2ac78867aa35d535835e14f78fb7433af Mon Sep 17 00:00:00 2001 From: Beichen Huang <71347503+BeichenHuang@users.noreply.github.com> Date: Fri, 16 Oct 2026 18:42:30 -0500 Subject: [PATCH] =?UTF-8?q?[Bugfix][LoRA]=20Fix=20garbage/NaN=20lora=5Fexp?= =?UTF-8?q?and=20output=20on=20SM90=20with=20fp32=20i=E2=80=A6=20(#61906)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: BeichenHuang Co-authored-by: Claude Opus 5.5 Co-authored-by: aoshen02 --- vllm/lora/ops/triton_ops/lora_expand_op.py | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/vllm/lora/ops/triton_ops/lora_expand_op.py b/vllm/lora/ops/triton_ops/lora_expand_op.py index 184e58247b..05fea52a71 100644 --- a/vllm/lora/ops/triton_ops/lora_expand_op.py +++ b/vllm/lora/ops/triton_ops/lora_expand_op.py @@ -15,6 +15,7 @@ from vllm.lora.ops.triton_ops.utils import ( get_lora_op_configs, supports_pdl, ) +from vllm.platforms import current_platform from vllm.triton_utils import tl, triton from vllm.utils.torch_utils import direct_register_custom_op @@ -231,6 +232,10 @@ def _lora_expand( torch.bfloat16, ]: CAST_TYPE = True + # Pipelined wgmma with a register A operand is miscompiled by + # ptxas < 13.3 on SM90 (triton-lang/triton#9433). + if current_platform.is_cuda() and current_platform.is_device_capability(90): + NUM_STAGES = 1 # TODO (varun): This grid formulation maximizes parallelization at the # cost of wasteful thread block launch when only a few input tokens require