[Bug]: Tensor parallelism on A40 results in CUDA illegal instruction error

[AlessioNetti]

System Info

Nvidia A40 GPUs, AMD Epyc 7352 CPU, NGC 25.09 PyTorch image (CUDA 13.0, NCCL 2.27.7, PyTorch 2.9.0a0), TensorRT-LLM 1.2.0rc2.

Who can help?

@Funatiq

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

Using the official quickstart_advanced.py script for the LLM API (PyTorch workflow) is enough to trigger the issue. In this case, we are using a Qwen3 1.7B model as an example:

python quickstart_advanced.py --tp_size 2 --pp_size 1 --max_batch_size 64 --use_cuda_graph --model_dir ./Qwen3-1.7B

Expected behavior

The model should load correctly and inference should proceed as normal.

actual behavior

During initialization time (specifically, during CUDA graph capturing) a fatal CUDA error: an illegal instruction was encountered is raised. If CUDA graph capturing is disabled, the model loads correctly, and the error occurs at the first batch of requests with a sufficiently large size. In general, the error only arises if --max_batch_size is large enough (in our experiments, >= 32).

Looking at the traceback (omitting some lines for brevity) reveals that the error happens during execution of a linear module:

  File "/virtualenv/lib/python3.12/site-packages/tensorrt_llm/_torch/modules/attention.py", line 548, in forward
    qkv = self.qkv_proj(hidden_states)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/local/lib/python3.12/dist-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/local/lib/python3.12/dist-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/virtualenv/lib/python3.12/site-packages/tensorrt_llm/_torch/modules/linear.py", line 2041, in forward
    output = self.apply_linear(input, self.bias, lora_params, layer_idx)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/virtualenv/lib/python3.12/site-packages/tensorrt_llm/_torch/modules/linear.py", line 1994, in apply_linear
    output = self.quant_method.apply(self, input, bias)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/virtualenv/lib/python3.12/site-packages/tensorrt_llm/_torch/modules/linear.py", line 312, in apply
    output = F.linear(input, module.weight, bias)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
RuntimeError: CUDA error: CUBLAS_STATUS_EXECUTION_FAILED when calling `cublasGemmEx( handle, opa, opb, m, n, k, &falpha, a, CUDA_R_16BF, lda, b, CUDA_R_16BF, ldb, &fbeta, c, std::is_same_v<C_Dtype, float> ? CUDA_R_32F : CUDA_R_16BF, ldc, compute_type, CUBLAS_GEMM_DEFAULT_TENSOR_OP)`

Exception ignored in: <function PyTorchModelEngine.__del__ at 0x7fee8dfdde40>
Traceback (most recent call last):
  File "/virtualenv/lib/python3.12/site-packages/tensorrt_llm/_torch/pyexecutor/model_engine.py", line 976, in __del__
    release_gc()
  File "/virtualenv/lib/python3.12/site-packages/tensorrt_llm/_utils.py", line 728, in release_gc
    torch.cuda.empty_cache()
  File "/usr/local/lib/python3.12/dist-packages/torch/cuda/memory.py", line 224, in empty_cache
    torch._C._cuda_emptyCache()
torch.AcceleratorError: CUDA error: an illegal instruction was encountered

additional notes

The issue only occurs on A40 (compute capability 86) but not on H100 or GB200. Furthermore, this is strictly connected to use of tensor parallelism. The issue appears to be only present in TensorRT-LLM 1.2.0rc2, and 1.2.0rc0.post1 does not suffer from it.

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.

[Funatiq]

@NVIDIA/trt-llm-torch-attention-devs any idea what's going on here?

[yuxianq]

@AlessioNetti I cannot get A40x2 now, so I validate it on A16x2 with PCIe, which is also sm86. I cannot reproduce this issue with both v1.2.0rc2 and latest main branch (commit f1d637e). I use the same dependency versions with you. Are there NVLink connections in your A40x2 environment?

[AlessioNetti]

Thanks for looking into this, @yuxianq - yes, our A40 hosts actually have 8 GPUs with pairwise NVLink connections.

[yuxianq]

@AlessioNetti I cannot get sm86 + NVLink environment for further investigation. If it is related to NVLink, using NCCL instead of custom allreduce op may workaround it. You can patch trtllm's tensorrt_llm/_torch/distributed/ops.py with self.strategy = AllReduceStrategy.NCCL in

TensorRT-LLM/tensorrt_llm/_torch/distributed/ops.py

Line 602 in eeb56c2

self.strategy = strategy

.

[AlessioNetti]

Hardcoding the AllReduce strategy to AllReduceStrategy.NCCL indeed fixes the issue. Do you happen to know the performance impact of this change and if this is a viable temporary workaround? I'm also wondering if the problem might be caused by something else in our stack (e.g., drivers) or if we should just wait for a fix on the TRT-LLM side.

[yuxianq]

@AlessioNetti It should be slower in the small allreduce message cases (e.g. < 16MB in TP2 cases, see

TensorRT-LLM/cpp/tensorrt_llm/thop/allreduceOp.cpp

Line 927 in eeb56c2

AllReduceStrategyType selectImplementation(size_t seq_len, size_t hidden_size)

for all rules).
This is a temporary workaround, but we cannot fix it if we don't have environment to reproduce it.
You can try whether export NCCL_P2P_DISABLE=1 can disable custom allreduce kernels without patching, since custom allreduce kernels are only enabled when P2P is supported, see

TensorRT-LLM/cpp/tensorrt_llm/thop/allreduceOp.cpp

Line 975 in eeb56c2

if (message_size_bytes > max_workspace_size || !mIsP2PSupported || !mIsNVLINKSupported)

.

[AlessioNetti]

Just exporting NCCL_P2P_DISABLE=1 doesn't seem to work unfortunately, so we'll patch the code for the time being.

[AlessioNetti]

Another small update, relevant for @Funatiq : I rebuilt and tested TRT-LLM 1.2.0rc2 under the 25.10 PyTorch base image, but that did not help either. As mentioned above, we will be patching the code to work around the issue until a fix is implemented. If necessary, we can also help with testing on our A40 + NVLink setup.

[Funatiq]

@yuxianq I was able to reproduce the error with A30x2 with NVLink, so it doesn't seem to be limited to sm86. Could you give that a try?

[yuxianq]

@AlessioNetti @Funatiq Please Add CUDA_LAUNCH_BLOCKING=1 CUDA_LAUNCH_SYNC=1 TLLM_LOG_LEVEL=DEBUG and rerun the command. Find the last log containing "AllReduceOp runtime strategy for rank" to see which allreduce kernel caused it. I will triage this bug to the kernel author. Thanks~

[Funatiq]

Her is the log:

[TensorRT-LLM][DEBUG] All reduce message size is 262144
[TensorRT-LLM][DEBUG] AllReduceOp runtime strategy for rank 1: TWOSHOT
terminate called after throwing an instance of 'tensorrt_llm::common::TllmException'
  what():  [TensorRT-LLM][ERROR] CUDA runtime error in cudaLaunchKernelEx(&cfg, allreduce_fusion_kernel_twoshot_sync<Pattern, DType, NRanks, Fp32Acc>, params, begin_tokens, token_num_per_ranks): an illegal instruction was encountered (../../../tensorrt_llm/kernels/communicationKernels/allReduceFusionKernels.cu:619)