[Bug] Default PDL enablement fails FlashInfer RMSNorm JIT on SM89 GPUs

System Info

System Information

• TensorRT-LLM: 1.2.0
• FlashInfer: 0.6.7.post3
• NVIDIA CUTLASS DSL: 4.3.4
• PyTorch: 2.9.1+cu128
• PyTorch CUDA runtime: 12.8
• CUDA ptxas: CUDA 13.0, V13.0.88
• GPU: NVIDIA L20
• Compute capability: SM89 (8, 9)
• NVIDIA driver: 580.105.08
• Model: Qwen/Qwen2.5-7B-Instruct (BF16)

Description

TRTLLM_ENABLE_PDL defaults to 1, and the shared
tensorrt_llm._torch.flashinfer_utils.get_env_enable_pdl() helper forwards
this value to FlashInfer operators without checking the GPU architecture.

On an SM89 GPU, this causes FlashInfer's CuTe DSL RMSNorm path to compile a
module targeting sm_89 while also emitting the Hopper-only PDL instructions
griddepcontrol.wait and griddepcontrol.launch_dependents. ptxas then
correctly rejects the generated PTX and TensorRT-LLM executor initialization
fails.

Who can help?

No response

Information

• The official example scripts
• My own modified scripts

Tasks

• An officially supported task in the examples folder (such as GLUE/SQuAD, ...)
• My own task or dataset (give details below)

Reproduction

unset TRTLLM_ENABLE_PDL

trtllm-llmapi-launch trtllm-serve Qwen/Qwen2.5-7B-Instruct
--host 0.0.0.0
--port 8000
--backend pytorch

Expected behavior

PDL should only be enabled on SM90+ GPUs. On SM89, TensorRT-LLM should
automatically disable PDL and start successfully when the environment variable
is unset or set to 1.

actual behavior

cutlass.base_dsl.common.DSLRuntimeError: DSLRuntimeError: 🧊🧊🧊 ICE 🧊🧊🧊
Caused exception: Failure while executing pass pipeline:
error: unknown: NVPTX compiler invocation failed, error log: ptxas application ptx input, line 38; error : Modifier '.wait' requires .target sm_90 or higher
ptxas application ptx input, line 38; error : Instruction 'griddepcontrol' requires .target sm_90 or higher
ptxas application ptx input, line 709; error : Modifier '.launch_dependents' requires .target sm_90 or higher
ptxas application ptx input, line 709; error : Instruction 'griddepcontrol' requires .target sm_90 or higher
ptxas fatal : Ptx assembly aborted due to errors
...
error: "@cute.jit"(".../flashinfer/norm/kernels/rmsnorm.py":203:5):
An error happened while serializing the module.

CuTe DSL debug artifacts confirm that the module target is sm_89 and contains:
llvm.inline_asm ... "griddepcontrol.wait;"
llvm.inline_asm ... "griddepcontrol.launch_dependents;"
The FlashInfer RMSNorm JIT cache key also records enable_pdl=True.

additional notes

Workaround

The following workaround allows the server to start:

export TRTLLM_ENABLE_PDL=0

Proposed Fix

Gate PDL in TensorRT-LLM's shared FlashInfer helper:
requested = os.environ.get("TRTLLM_ENABLE_PDL", "1") == "1"
enabled = requested and get_sm_version() >= 90
This protects all TensorRT-LLM paths that reuse the helper, including
FlashInfer RMSNorm, fused RMSNorm, activation, sampling, attention, and FP4
quantization.
FlashInfer may additionally benefit from defensive validation for
enable_pdl=True on SM<90, but TensorRT-LLM should not forward a
hardware-unsupported PDL request to downstream dependencies.

Before submitting a new issue...

• Make sure you already searched for relevant issues, and checked the documentation and examples for answers to frequently asked questions.

Search for the shared get_env_enable_pdl() helper and inspect how it obtains the SM version before forwarding the environment setting to FlashInfer. Reproduce with TRTLLM_ENABLE_PDL unset on an SM89 GPU, then verify that PDL is disabled for SM<90 while remaining enabled when supported and that executor initialization succeeds.