diff --git a/cpp/kernels/xqa/mha_sm90.cu b/cpp/kernels/xqa/mha_sm90.cu index dc21872ac696..7f7164c9a24e 100644 --- a/cpp/kernels/xqa/mha_sm90.cu +++ b/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(); diff --git a/cpp/tensorrt_llm/common/attentionOp.cpp b/cpp/tensorrt_llm/common/attentionOp.cpp index a31750ce6ea7..33a92a8c033d 100644 --- a/cpp/tensorrt_llm/common/attentionOp.cpp +++ b/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. diff --git a/docs/source/deployment-guide/configuring-cpu-affinity.md b/docs/source/deployment-guide/configuring-cpu-affinity.md new file mode 100644 index 000000000000..daccaf4e65ff --- /dev/null +++ b/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 | +|---------------------------------|------------------------------------------------------------------------------------------------------------------------------| +| | 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 \ + +``` + +### 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 < +``` + +See the official [OpenMPI Documentation](https://www.open-mpi.org/doc/) for +more details on mapping and binding of MPI ranks. diff --git a/docs/source/index.rst b/docs/source/index.rst index d74ff651e33e..6c0458fafbda 100644 --- a/docs/source/index.rst +++ b/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 diff --git a/docs/source/overview.md b/docs/source/overview.md index e7aeed43a237..0eee6b6f9fd7 100644 --- a/docs/source/overview.md +++ b/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`. diff --git a/docs/source/supported-hardware.md b/docs/source/supported-hardware.md new file mode 100644 index 000000000000..a8c96e99eeb2 --- /dev/null +++ b/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 diff --git a/tensorrt_llm/_torch/models/modeling_qwen3vl.py b/tensorrt_llm/_torch/models/modeling_qwen3vl.py index 312e5ea74583..8a5087731eb6 100644 --- a/tensorrt_llm/_torch/models/modeling_qwen3vl.py +++ b/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): diff --git a/tensorrt_llm/_torch/models/modeling_qwen3vl_moe.py b/tensorrt_llm/_torch/models/modeling_qwen3vl_moe.py index 74cd16ec69c1..66da3c153e2a 100644 --- a/tensorrt_llm/_torch/models/modeling_qwen3vl_moe.py +++ b/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): diff --git a/tensorrt_llm/_torch/pyexecutor/model_engine.py b/tensorrt_llm/_torch/pyexecutor/model_engine.py index 866bc4cf0541..cb9971b9ca4b 100644 --- a/tensorrt_llm/_torch/pyexecutor/model_engine.py +++ b/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, diff --git a/tensorrt_llm/_torch/speculative/interface.py b/tensorrt_llm/_torch/speculative/interface.py index ea623bc44237..523a78129f0b 100644 --- a/tensorrt_llm/_torch/speculative/interface.py +++ b/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 ( diff --git a/tensorrt_llm/llmapi/llm_args.py b/tensorrt_llm/llmapi/llm_args.py index f5c16b9d9cb0..0b4d4e60f272 100644 --- a/tensorrt_llm/llmapi/llm_args.py +++ b/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) diff --git a/tests/integration/defs/disaggregated/test_auto_scaling.py b/tests/integration/defs/disaggregated/test_auto_scaling.py index 038049109ddd..4821ee847638 100644 --- a/tests/integration/defs/disaggregated/test_auto_scaling.py +++ b/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) diff --git a/tests/integration/defs/perf/disagg/execution/executor.py b/tests/integration/defs/perf/disagg/execution/executor.py index b5e2d9a83f6b..68ac6166c787 100644 --- a/tests/integration/defs/perf/disagg/execution/executor.py +++ b/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, diff --git a/tests/integration/defs/perf/disagg/utils/common.py b/tests/integration/defs/perf/disagg/utils/common.py index ec09e225d2fa..f1191020528d 100644 --- a/tests/integration/defs/perf/disagg/utils/common.py +++ b/tests/integration/defs/perf/disagg/utils/common.py @@ -2,49 +2,42 @@ import os -# GPU resource configuration -# Centralized configuration for all GPU-specific parameters +# GPU resource configuration - centralized config for all GPU-specific parameters GPU_RESOURCE_CONFIG = { - # OCI GB200 - "GB200": { - "slurm_extra_args": "--gres=gpu:4", # SLURM extra arguments (empty string if not required) + "GB200": { # OCI GB200 + "slurm_extra_args": "--gres=gpu:4", "set_segment": True, - "lock_freq_graphics_mhz": 2062, # GPU graphics clock lock frequency (MHz) - "lock_freq_memory_mhz": 3996, # GPU memory clock lock frequency (MHz) + "lock_freq_graphics_mhz": 2062, + "lock_freq_memory_mhz": 3996, }, - # Lyris GB200 - "GB200_LYRIS": { - "slurm_extra_args": "", # GB300 does not require extra args + "GB200_LYRIS": { # Lyris GB200 + "slurm_extra_args": "", "set_segment": True, - "lock_freq_graphics_mhz": None, # TODO: Set GB300 lock frequency + "lock_freq_graphics_mhz": None, "lock_freq_memory_mhz": None, }, - # Lyris GB300 - "GB300": { - "slurm_extra_args": "", # GB300 does not require extra args + "GB300": { # Lyris GB300 + "slurm_extra_args": "", "set_segment": True, - "lock_freq_graphics_mhz": None, # TODO: Set GB300 lock frequency + "lock_freq_graphics_mhz": None, "lock_freq_memory_mhz": None, }, - # H100 "H100": { - "slurm_extra_args": "", # H100 does not require extra args + "slurm_extra_args": "", "set_segment": False, - "lock_freq_graphics_mhz": None, # TODO: Set H100 lock frequency + "lock_freq_graphics_mhz": None, "lock_freq_memory_mhz": None, }, - # B200 - "B200": { + "B200": { # OCI B200 "slurm_extra_args": "--gres=gpu:4", "set_segment": False, - "lock_freq_graphics_mhz": None, # TODO: Set B200 lock frequency + "lock_freq_graphics_mhz": None, "lock_freq_memory_mhz": None, }, - # B300 - "B300": { + "B300": { # OCI B300 "slurm_extra_args": "--gres=gpu:4", "set_segment": False, - "lock_freq_graphics_mhz": None, # TODO: Set B300 lock frequency + "lock_freq_graphics_mhz": None, "lock_freq_memory_mhz": None, }, } @@ -67,7 +60,18 @@ def get_slurm_account() -> str: @staticmethod def get_slurm_job_name() -> str: - return os.getenv("SLURM_JOB_NAME", "unified-benchmark") + """Get SLURM job name: {SLURM_ACCOUNT}-{base}. + + Example: myaccount-unified.benchmark + Customize base via SLURM_JOB_BASE_NAME env var (default: unified.benchmark) + """ + account = EnvManager.get_slurm_account() + base = os.getenv("SLURM_JOB_BASE_NAME", "unified.benchmark") + + # Only use account as prefix if it's set and not a placeholder + if account and not account.startswith("<"): + return f"{account}-{base}" + return base @staticmethod def get_slurm_set_segment() -> bool: diff --git a/tests/integration/defs/stress_test/stress_test.py b/tests/integration/defs/stress_test/stress_test.py index f5fed7b0ebbb..2de8a3880199 100644 --- a/tests/integration/defs/stress_test/stress_test.py +++ b/tests/integration/defs/stress_test/stress_test.py @@ -378,7 +378,7 @@ def is_port_available(port: int, ["GUARANTEED_NO_EVICT", "MAX_UTILIZATION"], ids=lambda x: x) @pytest.mark.parametrize("stress_time_timeout", [(180, 300), (300, 450), - (600, 900), (3600, 5400)], + (600, 900), (3600, 10800)], ids=lambda x: f"stress_time_{x[0]}s_timeout_{x[1]}s") @pytest.mark.parametrize( "config", @@ -393,10 +393,22 @@ def is_port_available(port: int, memory_requirement=12), # Configuration for DeepSeek-V3 model ModelConfig(model_dir="DeepSeek-V3", tp_size=8, memory_requirement=96), - # Configuration for DeepSeek-R1 model + # Configuration for DeepSeek-R1 model with FP8 checkpoints (8 GPU setup) ModelConfig(model_dir="DeepSeek-R1/DeepSeek-R1", tp_size=8, memory_requirement=96), + # Configuration for DeepSeek-R1 model with FP8 checkpoints (4 GPU setup, requires GB300 288GB) + ModelConfig(model_dir="DeepSeek-R1/DeepSeek-R1", + tp_size=4, + memory_requirement=256), + # Configuration for DeepSeek-R1 model with NVFP4 checkpoints (8 GPU setup) + ModelConfig(model_dir="DeepSeek-R1/DeepSeek-R1-0528-FP4", + tp_size=8, + memory_requirement=96), + # Configuration for DeepSeek-R1 model with NVFP4 checkpoints (4 GPU setup) + ModelConfig(model_dir="DeepSeek-R1/DeepSeek-R1-0528-FP4", + tp_size=4, + memory_requirement=168), ], ids=lambda x: f"{os.path.basename(x.model_dir)}_tp{x.tp_size}") def test_run_stress_test(config, stress_time_timeout, backend, @@ -506,19 +518,20 @@ def stress_test(config, 36000 # 10 hours for DeepSeek-V3 or DeepSeek-R1, change this value if needed ) - # For DeepSeek-V3 specific server parameters + # For DeepSeek-V3 or DeepSeek-R1 specific server parameters if "DeepSeek-V3" in config.model_dir or "DeepSeek-R1" in config.model_dir: test_server_config = ServerConfig( port=test_server_config.port, host=test_server_config.host, pp_size=test_server_config.pp_size, - ep_size=8, # DeepSeek-V3 or DeepSeek-R1 specific ep_size + ep_size=config. + tp_size, # ep_size matches tp_size for DeepSeek models max_batch_size= 2048, # DeepSeek-V3 or DeepSeek-R1 specific max_batch_size max_num_tokens= - 2048, # DeepSeek-V3 or DeepSeek-R1 specific max_num_tokens + 8192, # DeepSeek-V3 or DeepSeek-R1 specific max_num_tokens kv_cache_free_gpu_memory_fraction= - 0.7, # DeepSeek-V3 or DeepSeek-R1 specific kv_cache fraction + 0.85, # DeepSeek-V3 or DeepSeek-R1 specific kv_cache fraction capacity_scheduler_policy=test_server_config. capacity_scheduler_policy, wait_interval=test_server_config.wait_interval, @@ -583,6 +596,35 @@ def stress_test(config, extra_llm_options["enable_attention_dp"] = True + # Set MOE backend based on GPU architecture and checkpoint type + # B200/GB200 (Blackwell, SM100+) with FP8 checkpoints: use DEEPGEMM backend + # B200/GB200 (Blackwell, SM100+) with NVFP4 checkpoints: use CUTEDSL backend + # H100/H200 (Hopper, SM90) with FP8 checkpoints: use CUTLASS backend (default) + try: + import torch + if torch.cuda.is_available(): + device_capability = torch.cuda.get_device_capability(0) + is_blackwell = device_capability[0] >= 10 + is_nvfp4 = "FP4" in config.model_dir.upper() + + if is_blackwell: + if is_nvfp4: + moe_backend = "CUTEDSL" + else: + moe_backend = "DEEPGEMM" + + extra_llm_options["moe_config"] = { + "backend": moe_backend, + } + checkpoint_type = "NVFP4" if is_nvfp4 else "FP8" + print_info( + f"Detected GPU architecture is SM{device_capability[0]}{device_capability[1]} (Blackwell), " + f"using {moe_backend} MOE backend for DeepSeek-R1/DeepSeek-V3 with {checkpoint_type} checkpoints" + ) + except Exception as e: + print_warning(f"Failed to detect GPU architecture: {e}. " + "Using default MOE backend (CUTLASS).") + if config.backend == "pytorch": extra_llm_options.update({ "cuda_graph_config": { @@ -1047,6 +1089,18 @@ def stress_stage(model_name, request_count = int(stress_request_rate * stress_time) test_timeout = stress_config.stress_timeout + # Cap request count for large MoE models (DeepSeek-V3/R1) to prevent timeout + if "DeepSeek-V3" in model_path or "DeepSeek-R1" in model_path: + max_sustainable_rate = 3.0 # req/s - conservative estimate + max_request_count = int(max_sustainable_rate * + stress_config.stress_time) + if request_count > max_request_count: + print_info( + f"Capping request_count from {request_count} to {max_request_count} " + f"for DeepSeek V3/R1 model (sustainable rate: {max_sustainable_rate} req/s)" + ) + request_count = max_request_count + print_info( f"Running stress test with concurrency={stress_concurrency}, request_count={request_count}" ) diff --git a/tests/integration/defs/test_e2e.py b/tests/integration/defs/test_e2e.py index 918f135c2166..6c9f9240669f 100644 --- a/tests/integration/defs/test_e2e.py +++ b/tests/integration/defs/test_e2e.py @@ -666,11 +666,8 @@ def test_trtllm_bench_llmapi_launch(llm_root, llm_venv, model_name, ids=["llama3_1-8b"]) @pytest.mark.parametrize("model_subdir", ["llama-3.1-model/Meta-Llama-3.1-8B"], ids=["llama_v3_1"]) -@pytest.mark.parametrize("use_pytorch_backend", [False], ids=["trt_backend"]) -def test_trtllm_bench_mig_launch(llm_root, llm_venv, model_name, model_subdir, - use_pytorch_backend): +def test_trtllm_bench_mig_launch(llm_root, llm_venv, model_name, model_subdir): "run bench mark in MIG mode, check if the throughput is increasing by concurrency" - skip_engine_build = False results = {} concurrency_list = [1, 32, 64, 128] @@ -681,19 +678,17 @@ def test_trtllm_bench_mig_launch(llm_root, llm_venv, model_name, model_subdir, model_name=model_name, model_subdir=model_subdir, streaming=False, - use_pytorch_backend=use_pytorch_backend, + use_pytorch_backend=True, use_mpirun=False, tp_size=1, concurrency=concurrency, - num_requests=num_requests, - skip_engine_build=skip_engine_build) + num_requests=num_requests) output = runner() results[concurrency] = output print(f"\n=== Benchmark Results Comparison ===") print(f"Model: {model_name}") - print(f"Backend: {'PyTorch' if use_pytorch_backend else 'TensorRT'}") print( f"{'Concurrency':<15} {'Throughput':<15} {'Latency':<15} {'Num Requests':<15}" ) diff --git a/tests/integration/test_lists/qa/llm_function_stress.txt b/tests/integration/test_lists/qa/llm_function_stress.txt index cabba3def0fd..baee4172c80d 100644 --- a/tests/integration/test_lists/qa/llm_function_stress.txt +++ b/tests/integration/test_lists/qa/llm_function_stress.txt @@ -1,6 +1,10 @@ -stress_test/stress_test.py::test_run_stress_test[DeepSeek-V3_tp8-stress_time_3600s_timeout_5400s-GUARANTEED_NO_EVICT-pytorch-stress-test-with-accuracy] -stress_test/stress_test.py::test_run_stress_test[DeepSeek-V3_tp8-stress_time_3600s_timeout_5400s-MAX_UTILIZATION-pytorch-stress-test-with-accuracy] -stress_test/stress_test.py::test_run_stress_test[DeepSeek-R1_tp8-stress_time_3600s_timeout_5400s-MAX_UTILIZATION-pytorch-stress-test-with-accuracy] +stress_test/stress_test.py::test_run_stress_test[DeepSeek-V3_tp8-stress_time_3600s_timeout_10800s-GUARANTEED_NO_EVICT-pytorch-stress-test-with-accuracy] +stress_test/stress_test.py::test_run_stress_test[DeepSeek-V3_tp8-stress_time_3600s_timeout_10800s-MAX_UTILIZATION-pytorch-stress-test-with-accuracy] +stress_test/stress_test.py::test_run_stress_test[DeepSeek-R1_tp8-stress_time_3600s_timeout_10800s-MAX_UTILIZATION-pytorch-stress-test-with-accuracy] +stress_test/stress_test.py::test_run_stress_test[DeepSeek-R1_tp4-stress_time_3600s_timeout_10800s-GUARANTEED_NO_EVICT-pytorch-stress-test-with-accuracy] +stress_test/stress_test.py::test_run_stress_test[DeepSeek-R1_tp4-stress_time_3600s_timeout_10800s-MAX_UTILIZATION-pytorch-stress-test-with-accuracy] +stress_test/stress_test.py::test_run_stress_test[DeepSeek-R1-0528-FP4_tp4-stress_time_3600s_timeout_10800s-GUARANTEED_NO_EVICT-pytorch-stress-test-with-accuracy] +stress_test/stress_test.py::test_run_stress_test[DeepSeek-R1-0528-FP4_tp4-stress_time_3600s_timeout_10800s-MAX_UTILIZATION-pytorch-stress-test-with-accuracy] disaggregated/test_disaggregated.py::test_disaggregated_stress_test[input8k-output1k-conc512-deepseek_r1_v2_fp4_stress] disaggregated/test_disaggregated.py::test_disaggregated_stress_test[input8k-output1k-conc512-gpt_oss_120b_stress] accuracy/test_llm_api_pytorch.py::TestDeepSeekR1LongBenchV2::test_fp8_8gpus diff --git a/tests/integration/test_lists/waives.txt b/tests/integration/test_lists/waives.txt index 4eb031afe60e..d6b3486b15dc 100644 --- a/tests/integration/test_lists/waives.txt +++ b/tests/integration/test_lists/waives.txt @@ -246,7 +246,6 @@ accuracy/test_llm_api_pytorch.py::TestGPTOSS::test_eagle3_vswa_reuse_4gpus[one_m accuracy/test_llm_api_pytorch.py::TestGPTOSS::test_eagle3_vswa_reuse_4gpus[two_model] SKIP (https://nvbugs/5756028) accuracy/test_llm_api_pytorch.py::TestLlama3_2_1B::test_fp8_prequantized SKIP (https://nvbugs/5785465) accuracy/test_llm_api_pytorch.py::TestMinistral8BInstruct::test_fp8 SKIP (https://nvbugs/5785485) -accuracy/test_llm_api_pytorch.py::TestLlama3_3_70BInstruct::test_fp8_eagle3_tp8[eagle3_one_model=True-torch_compile=False] SKIP (https://nvbugs/5787892) accuracy/test_llm_api_pytorch.py::TestLlama3_3_70BInstruct::test_fp8_eagle3_tp8[eagle3_one_model=False-torch_compile=False] SKIP (https://nvbugs/5787892) accuracy/test_llm_api_pytorch.py::TestLlama4MaverickInstruct::test_fp8_chunked_prefill[tp8ep8-cuda_graph=False] SKIP (https://nvbugs/5795918) accuracy/test_disaggregated_serving.py::TestDeepSeekV3Lite::test_auto_dtype[mtp_nextn=2-overlap_scheduler=True] SKIP (https://nvbugs/5800591) @@ -301,7 +300,6 @@ accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_auto_dtype_4gpus[4-4 accuracy/test_llm_api_pytorch.py::TestGPTOSS::test_w4_2gpus[dp2-cutlass-auto] SKIP (https://nvbugs/5836828) accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_auto_dtype_4gpus[4-4-True-True-True] SKIP (https://nvbugs/5823587) accuracy/test_llm_api_pytorch.py::TestQwen3_4B::test_eagle3 SKIP (https://nvbugs/5836592) -test_e2e.py::test_openai_chat_harmony SKIP (https://nvbugs/5819444) accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[ep4-mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=False] SKIP (https://nvbugs/5800646) accuracy/test_disaggregated_serving.py::TestQwen3_30B_A3B::test_mixed_ctx_gen_model[ctxpp2gentp2] SKIP (https://nvbugs/5748664) examples/test_llama.py::test_llama_3_x_with_bf16_lora_torch[llama-3.2-1b-instruct] SKIP (https://nvbugs/5838178) @@ -344,6 +342,8 @@ accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_auto_dtype_4gpus[4-4 unittest/_torch/thop/serial/test_moe.py::TestMoeFp4::test_gptoss_style_nvfp4[limitinf-beta0-alpha0.1-RoutingGPTOSS-512-512-1] SKIP (https://nvbugs/5819042) disaggregated/test_disaggregated.py::test_disaggregated_benchmark_on_diff_backends[llama-v3-8b-hf] SKIP (https://nvbugs/5839137) accuracy/test_disaggregated_serving.py::TestLlama3_1_8BInstruct::test_ctx_pp_gen_tp_asymmetric[MMLU-gen_tp=1-ctx_pp=4] SKIP (https://nvbugs/5845943) +accuracy/test_cli_flow.py::TestGpt2::test_cuda_graph SKIP (https://nvbugs/5860520) +unittest/_torch/thop/serial/test_moe.py::TestMoeFp4::test_online_eplb288_topk_input[RoutingDSv3-1024-1024-256] SKIP (https://nvbugs/5859881) test_e2e.py::test_openai_chat_guided_decoding[openai/gpt-oss-120b] SKIP (https://nvbugs/5884677) accuracy/test_disaggregated_serving.py::TestGPTOSS::test_auto_dtype[True] SKIP (https://nvbugs/5849648) accuracy/test_disaggregated_serving.py::TestGPTOSS::test_auto_dtype[False] SKIP (https://nvbugs/5849648) diff --git a/tests/unittest/llmapi/apps/_test_openai_chat_harmony.py b/tests/unittest/llmapi/apps/_test_openai_chat_harmony.py index 5eadaf88ae79..e247cb21d9aa 100644 --- a/tests/unittest/llmapi/apps/_test_openai_chat_harmony.py +++ b/tests/unittest/llmapi/apps/_test_openai_chat_harmony.py @@ -160,6 +160,7 @@ async def test_streaming(client: openai.AsyncOpenAI, model: str): "content": "Explain the theory of relativity in brief." }], stream=True, + extra_body={"top_k": 1}, ) collected_messages = [] first_iteration = True