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
Tasks
Reproduction
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...
System Info
System Information
ptxas: CUDA 13.0, V13.0.88(8, 9)Qwen/Qwen2.5-7B-Instruct(BF16)Description
TRTLLM_ENABLE_PDLdefaults to1, and the sharedtensorrt_llm._torch.flashinfer_utils.get_env_enable_pdl()helper forwardsthis 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_89while also emitting the Hopper-only PDL instructionsgriddepcontrol.waitandgriddepcontrol.launch_dependents.ptxasthencorrectly rejects the generated PTX and TensorRT-LLM executor initialization
fails.
Who can help?
No response
Information
Tasks
examplesfolder (such as GLUE/SQuAD, ...)Reproduction
unset TRTLLM_ENABLE_PDLtrtllm-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=0Proposed 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...