[None][chroe] Mass integration of release/1.2 - 4th

cpp/kernels/xqa/mha_sm90.cu

@@ -1138,6 +1138,10 @@ CUBIN_EXPORT __global__
             auto& xBar = smem.xBar[idxXBuf];
             auto& vBar = smem.vBar[idxVBuf];
             auto const& vBuf = smem.vBuf(idxVBuf);
+#if !SWAP_AB
+            CtaBarrierPair& vtBar = smem.vtBar[idxVBuf];
+            auto& vtBuf = smem.vtBuf(idxVBuf);
+#endif
             xBar.produced.arrive_and_wait();
 #if SKIP_SOFTMAX_ATTN
             bool shouldSkipSoftmaxAttn = smem.skipSoftmaxVotesGemm0ToGemm1[idxXBuf]; // guarded by xBar
@@ -1153,8 +1157,6 @@ CUBIN_EXPORT __global__
             {
                 arrive_tx_and_wait(vBar.produced, exactDiv(sizeof(SharedMem::VBuffer), gemm1NbThrds));
 #if !SWAP_AB
-                CtaBarrierPair& vtBar = smem.vtBar[idxVBuf];
-                auto& vtBuf = smem.vtBuf(idxVBuf);
                 vtBar.consumed.arrive_and_wait();
                 transposeVTile(warpRank, laneId(), vtBuf, vBuf);
                 vBar.consumed.arrive();

cpp/tensorrt_llm/common/attentionOp.cpp

@@ -1271,14 +1271,6 @@ int AttentionOp::mlaGeneration(
                 mXqaDispatcher->run(xqaParams, kv_cache_buffer, kv_scale_cache_buffer);
                 return 0;
             }
-            else if (mIsSpecDecodingEnabled && mUseSpecDecoding)
-            {
-                TLLM_CHECK_WITH_INFO(false, "No available XQA kernels are found for speculative decoding mode.");
-            }
-            else if (mFuseFp4Quant)
-            {
-                TLLM_CHECK_WITH_INFO(false, "No available kernels are found for FP4 output.");
-            }
         }
 
         // Use FMHA otherwise.

docs/source/deployment-guide/configuring-cpu-affinity.md

@@ -0,0 +1,208 @@
+# CPU Affinity configuration in TensorRT-LLM
+
+## NUMA-aware affinity in TensorRT-LLM
+
+TensorRT-LLM is frequently deployed on
+[NUMA](https://en.wikipedia.org/wiki/Non-uniform_memory_access) systems. In
+order to ensure consistent and optimal performance on these systems, it is
+critical to set the CPU affinity of the workers/tasks launched as part of a
+particular TRT-LLM instance so as to minimize latency and maximize bandwidth of
+CPU↔GPU and CPU↔DRAM communication.
+
+Because TensorRT-LLM does the work of allocating GPU/CUDA devices to ranks, it
+is logically the ideal place for the CPU affinity to be determined and set. For
+this reason, TensorRT-LLM provides a mechanism to automatically set CPU
+affinity according to NUMA topology. In some situations/deployments, the user
+may wish to configure CPU affinity manually (i.e. using
+[numactl](https://github.com/numactl/numactl), [wrappers around the
+same](https://github.com/NVIDIA/mlperf-common/blob/main/client/bindpcie), or
+mpirun). For this reason, this feature is only activated if it is explicitly
+enabled or if CPU affinity is not already constrained by the user or
+environment. It is controlled by the TLLM_NUMA_AWARE_WORKER_AFFINITY
+environment variable as follows:
+
+
+| TLLM_NUMA_AWARE_WORKER_AFFINITY | Behavior                                                                                                                     |
+|---------------------------------|------------------------------------------------------------------------------------------------------------------------------|
+| <unset>                         | Affinity is auto-configured if it is unconstrained, and cleared if it is constrained by the user and/or environment          |
+| 1                               | Affinity is unconditionally auto-configured.                                                                                 |
+| 0 or any other value            | Affinity remains as configured by the user and/or environment                                                                |
+
+
+## Other environmental considerations
+
+Whether or not the user chooses to manually configure CPU affinity or have
+TensorRT-LLM configure it automatically, the environment can also constrain the
+CPU affinity in a way that subverts the user's intent. Both OpenMPI and Slurm
+may configure CPU affinity, so the following additional configuration is
+recommended to avoid this.
+
+### OpenMPI
+
+By default, OpenMPI chooses a rank-wise CPU affinity that is not sensitized to
+the NUMA-topology of the system. Because it does not know which GPU a
+particular rank will be communicating with (this is determined by TRT-LLM at
+runtime), it cannot set the CPU affinity accordingly. For this reason, it is
+recommended that OpenMPI's default binding policy be disabled as follows:
+
+```bash
+export OMPI_MCA_hwloc_base_binding_policy=none
+export OMPI_MCA_rmaps_base_inherit=1
+```
+
+The first environment variable ensures that OpenMPI will not attempt to bind or
+set the affinity of the ranks that are created at launch.
+
+The second ensures that OpenMPI's binding policy will propagate to MPI workers
+that are spawned by `mpi4py`'s `MPIPoolExecutor` class within TensorRT-LLM
+(when using mpirun).
+
+### Slurm
+
+If Slurm is configured to use a affinity or cgroup task plugin, then Slurm may
+also configure CPU affinity by default in a way that is not sensitized to NUMA
+topology. To prevent this, Slurm jobs should be launched accordingly:
+
+#### srun
+
+The srun parameters should include `--cpu-bind=none` and exclude `--exclusive`:
+
+```bash
+srun --cpu-bind=none ...
+```
+
+#### sbatch
+
+The sbatch script should set `SLURM_CPU_BIND` environment variable to "none":
+
+```bash
+export SLURM_CPU_BIND=none
+```
+
+Note: if this environment variable is set, it is not necessary to supply the
+`--cpu-bind=none` to each job step (srun invocation)
+
+## CPU affinity configuration examples
+
+### Using NUMA-aware autoconfiguration
+
+To explicitly enable the NUMA-aware autoconfiguration feature in TensorRT-LLM,
+simply set `TLLM_NUMA_AWARE_WORKER_AFFINITY` in the launch script (prior to
+`trtllm-bench` or `trtllm-serve`) as follows:
+
+```bash
+export TLLM_NUMA_AWARE_WORKER_AFFINITY=1
+```
+
+Because autoconfiguration happens within TensorRT-LLM itself, it will override
+any CPU affinity or binding that has been previously set by OpenMPI or Slurm.
+
+### NUMA-aware CPU affinity using [bindpcie](https://github.com/NVIDIA/mlperf-common/blob/main/client/bindpcie)
+
+The bindpcie script is designed to set a per-rank CPU affinity that is ideal
+for NUMA topology. While setting `TLLM_NUMA_AWARE_WORKER_AFFINITY=1` usually
+achieves the same result in terms of the CPU affinity that is set, this
+approach has the distinct advantage that the optimal CPU affinity gets set
+_upon launching_ TensorRT-LLM, guaranteeing that each worker/rank executes on
+the optimal NUMA node from inception. The NUMA-aware CPU affinity
+autoconfiguration mechanism in TensorRT-LLM, on the other hand, is triggered by
+each worker/rank upon its own PID _after_ it has already launched. If the
+worker/rank executes on a NUMA node other than the optimal NUMA node at some
+point between the launch of the process and the NUMA-aware autoconfiguration,
+it is possible that some CPU memory may have been allocated/touched on what
+will become a remote NUMA node after the point of autoconfiguration,
+potentially negatively impacting performance. In practice, this effect has been
+observed to have minimal performance impact, but some degradation of
+performance due to remote NUMA node access is still theoretically possible.
+
+The `bindpcie` script can only be applied to deployments that make use of
+`trtllm-llmapi-launch` within an sbatch script. One example of how to apply
+bindpcie to `trtllm-serve` in an sbatch script is as follows:
+
+```bash
+# Prevent TensorRT-LLM from autoconfiguring or clearing CPU affinity
+export TLLM_NUMA_AWARE_WORKER_AFFINITY=0
+
+# Prevent OpenMPI from overriding affinity set by bindpcie
+export OMPI_MCA_hwloc_base_binding_policy=none
+
+# Ensure that MPI binding policy propagates to any MPI workers dynamically
+# spawned by MPIPoolExecutor
+export OMPI_MCA_rmaps_base_inherit=1
+
+# Prevent Slurm from assigning a default CPU affinity
+export SLURM_CPU_BIND=none
+
+srun -l \
+    --container-image=${CONTAINER_IMAGE} \
+    --container-mounts=${MOUNT_DIR}:${MOUNT_DEST} \
+    --container-workdir=${WORKDIR} \
+    --export=ALL,PYTHONPATH=${SOURCE_ROOT} \
+    --mpi=pmix \
+    bash -c "
+        set -ex
+        $PROLOGUE
+        export PATH=$PATH:~/.local/bin
+
+        bindpcie trtllm-llmapi-launch \
+         trtllm-serve $LOCAL_MODEL \
+            ${ADDITIONAL_OPTIONS}
+```
+
+> [!NOTE]
+> This is not a complete or exhaustive example of an sbatch script to launch
+> trtllm-serve and is only intended to highlight the application of bindpcie
+> within an existing sbatch script.
+
+### Using [numactl](https://github.com/numactl/numactl)
+
+```bash
+# Prevent TensorRT-LLM from autoconfiguring or clearing CPU affinity
+export TLLM_NUMA_AWARE_WORKER_AFFINITY=0
+
+# Prevent OpenMPI from overriding affinity set by numactl
+export OMPI_MCA_hwloc_base_binding_policy=none
+
+# Ensure that MPI binding policy propagates to any MPI workers dynamically
+# spawned by MPIPoolExecutor
+export OMPI_MCA_rmaps_base_inherit=1
+
+# Use numactl to specify CPU and memory binding for all ranks (not per-rank)
+numactl --physcpubind=0,1,16,17 --membind=0 mpirun --report-bindings --oversubscribe --allow-run-as-root \
+<trtllm-serve | trtllm-bench> <arguments>
+```
+
+### Using mpirun
+
+If a manually-specified per-rank CPU affinity is desired when running on a
+single node with mpirun, this can be achieved most easily using an OpenMPI
+rankfile. The following is an example of how a rankfile can be used to
+arbitrarily map each of 4 MPI ranks to a distinct set of 4 cores:
+
+```bash
+# Prevent TensorRT-LLM from autoconfiguring or clearing CPU affinity
+export TLLM_NUMA_AWARE_WORKER_AFFINITY=0
+
+# Not strictly needed here, since we are overriding with explicit bindings from
+# a rankfile
+# export OMPI_MCA_hwloc_base_binding_policy=none
+
+# Ensure that MPI binding policy propagates to any MPI workers dynamically
+# spawned by MPIPoolExecutor
+export OMPI_MCA_rmaps_base_inherit=1
+
+# Create a rankfile to enumerate a set of 4 cores to which each rank is bound
+cat > ./rankfile <<EOF
+rank 0=localhost slot=0,1,2,3
+rank 1=localhost slot=4,5,6,7
+rank 2=localhost slot=8,9,10,11
+rank 3=localhost slot=12,13,14,15
+EOF
+
+# Run with (and report to verify) the bindings from the rankfile
+mpirun --rankfile ./rankfile -n 1 --report-bindings --oversubscribe --allow-run-as-root \
+<trtllm-serve | trtllm-bench> <arguments>
+```
+
+See the official [OpenMPI Documentation](https://www.open-mpi.org/doc/) for
+more details on mapping and binding of MPI ranks.

docs/source/index.rst

@@ -14,6 +14,7 @@ Welcome to TensorRT LLM's Documentation!
    overview.md
    quick-start-guide.md
    installation/index.rst
+   supported-hardware.md
 
 
 .. toctree::
@@ -25,6 +26,7 @@ Welcome to TensorRT LLM's Documentation!
    examples/trtllm_serve_examples
    examples/dynamo_k8s_example.rst
    deployment-guide/index.rst
+   deployment-guide/configuring-cpu-affinity.md
 
 .. toctree::
    :maxdepth: 2

docs/source/overview.md

@@ -51,13 +51,6 @@ TensorRT LLM strives to support the most popular models on **Day 0**.
 - **[Guided Decoding](./features/guided-decoding.md)**: Advanced sampling with stop words, bad words, and custom constraints
 - **[Disaggregated Serving (Beta)](./features/disagg-serving.md)**: Separate context and generation phases across different GPUs for optimal resource utilization
 
-### 🔧 **Latest GPU Architecture Support**
-
-TensorRT LLM supports the full spectrum of NVIDIA GPU architectures:
-- **NVIDIA Blackwell**: B200, GB200, B300, GB300, and RTX Pro 6000 SE with FP4 optimization
-- **NVIDIA Ada Lovelace**: L40/L40S, RTX 40 series with FP8 acceleration
-- **NVIDIA Ampere**: A100, RTX 30 series for production workloads
-
 ## What Can You Do With TensorRT LLM?
 
 Whether you're building the next generation of AI applications, optimizing existing LLM deployments, or exploring the frontiers of large language model technology, TensorRT LLM provides the tools, performance, and flexibility you need to succeed in the era of generative AI.To get started, refer to the {ref}`quick-start-guide`.

docs/source/supported-hardware.md

@@ -0,0 +1,7 @@
+# Supported Hardware
+
+TensorRT LLM supports the full spectrum of NVIDIA GPU architectures:
+- **NVIDIA Blackwell**: B200, GB200, B300, GB300, DGX Spark
+- **NVIDIA Hopper**: H100, H200, GH200
+- **NVIDIA Ada Lovelace**: L20, L40/L40S
+- **NVIDIA Ampere**: A100

tensorrt_llm/_torch/models/modeling_qwen3vl.py

@@ -1138,6 +1138,7 @@ def _get_requests_with_mm_data(self, multimodal_params):
             "video": "<|vision_start|><|video_pad|><|vision_end|>",
         },
         placeholder_placement=MultimodalPlaceholderPlacement.BEFORE_TEXT,
+        placeholders_separator="",
     ),
 )
 class Qwen3VLModel(Qwen3VLModelBase):

tensorrt_llm/_torch/models/modeling_qwen3vl_moe.py

@@ -41,6 +41,7 @@
             "video": "<|vision_start|><|video_pad|><|vision_end|>",
         },
         placeholder_placement=MultimodalPlaceholderPlacement.BEFORE_TEXT,
+        placeholders_separator="",
     ),
 )
 class Qwen3MoeVLModel(Qwen3VLModelBase):

tensorrt_llm/_torch/pyexecutor/model_engine.py

@@ -3445,13 +3445,9 @@ def forward(self,
                                                        no_cache=kv_cache_manager
                                                        is None)
             # attn_metadata now depends on spec_metadata since it determines the shape/content of spec_dec parameter Tensors
-            enable_mla = is_mla(self.model.model_config.pretrained_config)
             is_spec_dec_mode = spec_metadata.spec_dec_mode.attention_need_spec_dec_mode(
-                spec_resource_manager,
-                self.is_draft_model,
-                self.attn_backend,
-                self.model_is_wrapped,
-                is_mla=enable_mla)
+                spec_resource_manager, self.is_draft_model, self.attn_backend,
+                self.model_is_wrapped)
             attn_metadata.update_spec_dec_param(
                 batch_size=scheduled_requests.batch_size,
                 is_spec_decoding_enabled=is_spec_dec_mode,

tensorrt_llm/_torch/speculative/interface.py

@@ -164,12 +164,11 @@ def extend_ctx(self, attention_backend: Type[AttentionBackend]):
                               TrtllmAttention) or not xqa_supported
 
     def attention_need_spec_dec_mode(
-        self,
-        spec_resource_manager: Optional[BaseResourceManager],
-        is_draft_model: bool,
-        attention_backend: Type[AttentionBackend],
-        use_chain_drafter: bool,  # CDL
-        is_mla: bool,
+            self,
+            spec_resource_manager: Optional[BaseResourceManager],
+            is_draft_model: bool,
+            attention_backend: Type[AttentionBackend],
+            use_chain_drafter: bool,  # CDL
     ):
         """
         If true, the attention backend kernel needs to run in spec-dec mode (multi-token query mode).
@@ -182,8 +181,7 @@ def attention_need_spec_dec_mode(
         is_trtllm_attention = issubclass(attention_backend, TrtllmAttention)
 
         # Always use the multi-token query mode for 1-model if the kernels are available.
-        xqa_supported = not is_mla or get_sm_version() < 120
-        use_case_1 = self.use_one_engine() and xqa_supported
+        use_case_1 = self.use_one_engine()
         # For 2-model, we need to enable it when we process multiple tokens at once. This occurs with
         # the target model (verification) or on the first draft for CDL based speculation.
         use_case_2 = not self.use_one_engine() and (

tensorrt_llm/llmapi/llm_args.py

@@ -1183,6 +1183,12 @@ def __init__(self, **kwargs):
             self.max_total_draft_tokens = kwargs[
                 'num_nextn_predict_layers']  # Current MTP only support linear tree
 
+        if not self.mtp_eagle_one_model:
+            logger.warning(
+                "2-model style MTP is deprecated. The mtp_eagle_one_model flag will do nothing "
+                "in release 1.3. After that, the flag will be removed entirely."
+            )
+
     @classmethod
     def from_dict(cls, data: dict):
         out = cls(**data)

tests/integration/defs/disaggregated/test_auto_scaling.py

@@ -2,7 +2,6 @@
 import os
 import shutil
 import subprocess
-import sys
 import tempfile
 import traceback
 import uuid
@@ -187,14 +186,13 @@ def run_disagg_server(disagg_cluster_config, work_dir, port=0, save_log=False):
     cmds = ["trtllm-serve", "disaggregated", "-c", disagg_server_config_path]
     log_file = None
     log_path = None
+    stdout = None
+    stderr = None
     if save_log:
         log_path = os.path.join(work_dir, "disagg_server.log")
         log_file = open(log_path, "w+")
         stdout = log_file
         stderr = log_file
-    else:
-        stdout = sys.stdout
-        stderr = sys.stderr
     p = subprocess.Popen(cmds, stdout=stdout, stderr=stderr)
     return ProcessWrapper(p, log_file=log_file, log_path=log_path, port=port)
 

tests/integration/defs/perf/disagg/execution/executor.py

@@ -219,6 +219,7 @@ def submit_test_job(test_config) -> tuple:
 
             # Write temporary config file with replaced environment variables
             logger.info(f"Creating temporary config: {temp_config_path}")
+
             with open(temp_config_path, "w") as f:
                 yaml.dump(
                     test_config.config_data,