kanshan/TensorRT-LLMs
1
0
Fork 0
mirror of https://github.com/NVIDIA/TensorRT-LLM.git synced 2026-01-13 22:18:36 +08:00
Code Issues Actions 1 Packages Projects Releases Wiki Activity
4,607 Commits 25 Branches 59 Tags 2.1 GiB
main
Commit Graph

881 Commits

Author SHA1 Message Date
Chang Liu
e47c787dd7
[TRTLLM-8535][feat] Support DeepSeek V3.2 with FP8 + BF16 KV cache/NVFP4 + BF16 KV cache (#8405) 
Signed-off-by: Fanrong Li <23290157+lfr-0531@users.noreply.github.com>
Signed-off-by: Chang Liu <9713593+chang-l@users.noreply.github.com>
Signed-off-by: Tracin <10434017+Tracin@users.noreply.github.com>
 2025-10-24 13:40:41 -04:00
Chuang Zhu
2420918e5b
[TRTLLM-7078][chore] optimal kvcache transfer for VWSA (#7952) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-10-24 08:58:16 -04:00
Aurelien Chartier
32e1ad68e1
[None][chore] Cleanup GDS code (#8475) 
Signed-off-by: Aurelien Chartier <2567591+achartier@users.noreply.github.com>
 2025-10-23 12:36:31 -07:00
Shijie
928247a3f9
[https://nvbugs/5451205][feat] Add cuBLASLt NVFP4 GEMM backend support (#7943) 
Signed-off-by: Shijie Wang <jaywan@nvidia.com>
 2025-10-23 15:55:10 +08:00
dongxuy04
a7c2c8c212
[None][fix] Allow multi-threaded copy for GDRCopy wrapper (#8535) 
Signed-off-by: Dongxu Yang <78518666+dongxuy04@users.noreply.github.com>
 2025-10-23 10:25:04 +08:00
Anthony Chang
8a3b870e09
[None][feat] Update TRTLLM MoE MxFP4 cubins; autotune tileN (#8156) 
Signed-off-by: Anthony Chang <27950904+rosenrodt@users.noreply.github.com>
 2025-10-23 09:14:18 +08:00
Anish Shanbhag
15de45d782
[TRTLLM-8682][chore] Remove auto_parallel module (#8329) 
Signed-off-by: Anish Shanbhag <ashanbhag@nvidia.com>
 2025-10-22 20:53:08 -04:00
dongxuy04
df689f8fed
[None][fix] Fix EPLB CPU thread NUMA binding (#8579) 
Signed-off-by: Dongxu Yang <78518666+dongxuy04@users.noreply.github.com>
 2025-10-22 10:52:09 -04:00
Patrice Castonguay
879039f6d5
[https://nvbugs/5429636][feat] Kv transfer timeout (#8459) 
Signed-off-by: raayandhar <raayan.dhar@gmail.com>
Signed-off-by: Patrice Castonguay <55748270+pcastonguay@users.noreply.github.com>
Co-authored-by: raayandhar <raayan.dhar@gmail.com>
 2025-10-22 09:29:02 -04:00
qsang-nv
07edac2818
[None][feat] Add vLLM KV Pool support for XQA mla kernel (#8560) 
Signed-off-by: Qidi Sang <200703406+qsang-nv@users.noreply.github.com>
 2025-10-22 14:12:57 +08:00
Shi Xiaowei
a0024f4d34
[None][doc] Facilitates the integration of the transfer agent (#7867) 
Signed-off-by: Shixiaowei02 <39303645+Shixiaowei02@users.noreply.github.com>
 2025-10-21 20:06:24 +08:00
Yuxian Qiu
ec32711b1e
[https://nvbugs/5542862][fix] Upgrade fmha_v2. (#8364) 
Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com>
 2025-10-20 10:20:23 +08:00
ChristinaZ
c8b9998acb
[TRTLLM-8637][feat] Optimize the routing kernel for DeepseekV3 (MoE CUTLASS backend); Add support for KimiK2 and Qwen-next (MoE TRTLLM backend) (#7761) 
Signed-off-by: Christina Zhang <83400082+ChristinaZ@users.noreply.github.com>
 2025-10-20 10:08:31 +08:00
Yueh-Ting (eop) Chen
128a351bdc
[None][fix] Avoid overwrite of kv_cache_config.max_tokens for VSWA scheme for the KVCacheManager (#8219) 
For VSWA scheme, we do not want `kv_cache_cnonfig.max_token` to control
and cap the maximum memory of a block pool because block pool size are
not identical amongst different window sizes. This MR omits the effect
of `kv_cache_config.max_tokens` under `kvCacheManager.cpp` to allow the
setting of block pool size to rely on the window size to share ratio
and the total gpu memory analyzed and fed to the kv cache manager.

Only skipping for VSWA scheme, no extra coverage was added.

Signed-off-by: eopXD <yuehtingc@nvidia.com>
 2025-10-20 10:48:40 +09:00
Bo Deng
dd25595ae8
[TRTLLM-7964][infra] Set nixl to default cache transceiver backend (#7926) 
Signed-off-by: Bo Deng <deemod@nvidia.com>
 2025-10-19 19:24:43 +08:00
jthomson04
852316886e
[None][fix] Fix KV event consumption (#6346) 
Signed-off-by: jthomson04 <jwillthomson19@gmail.com>
 2025-10-18 15:41:26 -07:00
Wanli Jiang
56f697be2e
[None][feat] Add fmha_v2 kernel for head_dim=80 and sm=100 to support VLM (#8392) 
Signed-off-by: Wanli Jiang <35160485+Wanli-Jiang@users.noreply.github.com>
 2025-10-17 19:42:47 +08:00
Perkz Zheng
0722717ec0
[None][fix] trtllm-gen regression in PR 8301 (#8426) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-10-17 03:21:31 -07:00
Iman Tabrizian
22eb1633ae
[None][bug] Set NCCL_GRAPH_REGISTER to false to avoid hang (#8413) 
Signed-off-by: Iman Tabrizian <10105175+tabrizian@users.noreply.github.com>
 2025-10-16 18:59:18 +02:00
Patrice Castonguay
b7602f7bd4 [https://nvbugs/5534837][fix] Fix KV cache split on long context (#8247) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
Signed-off-by: Patrice Castonguay <55748270+pcastonguay@users.noreply.github.com>
Co-authored-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
Signed-off-by: Mike Iovine <6158008+mikeiovine@users.noreply.github.com>
 2025-10-16 22:46:19 +08:00
Min Yu
0a0159fdd8
[https://nvbugs/5378031] [feat] W4A8 AWQ MoE supports Per Expert Pre-quant Scale Factor for PyT backend (#7286) 
Signed-off-by: Min Yu <171526537+yumin066@users.noreply.github.com>
 2025-10-16 11:07:48 +08:00
ChristinaZ
db1c271bc6
[None][feat] Revise the calculation related to TileN in routing of MOE TRTLLM backend (#8148) 
Signed-off-by: Christina Zhang <83400082+ChristinaZ@users.noreply.github.com>
 2025-10-16 09:15:46 +08:00
Chuang Zhu
40d129a415
[None][fix] Fix cache buffer size for window (#8320) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-10-16 09:01:11 +08:00
Fanrong Li
0d20a8fd61
[TRTLLM-8536][feat] Add the sparse attention framework and one use case--RocketKV support (#8086) 
Signed-off-by: Fanrong Li <23290157+lfr-0531@users.noreply.github.com>
Signed-off-by: yuhangh <58161490+heyuhhh@users.noreply.github.com>
Co-authored-by: yuhangh <58161490+heyuhhh@users.noreply.github.com>
 2025-10-14 08:23:16 -07:00
Aurelien Chartier
7291cdc422
[https://nvbugs/5404000][fix] Ensure consistency between firstTokenTime and lastTokenTime (#8294) 
Signed-off-by: Aurelien Chartier <2567591+achartier@users.noreply.github.com>
 2025-10-14 08:15:08 -04:00
Chuang Zhu
8733e830fc
[None][fix] Add lock for request_to_session in sendReadySingal (#8310) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-10-14 04:32:37 -07:00
Yueh-Ting (eop) Chen
4882815fa1
[TLLM-6777][feature] Support SWA KV cache reuse OOW block detach (#7922) 
This MR is a continuation of #6768. In the previous merge request,
OOW (out-of-window) blocks are only detached when reuse is not enabled,
that is, the block movement behavior is identical between SWA and full
attention when reuse is enabled.

This merge request attempts to enable OOW block detach when reuse is
enabled. The required changes are:

- Let KV cache manager keep track of which block is used by which
  sequence
- Remove restriction for the eviction policy to be able to release a
  non-leaf block

Along with the development, bugs inside freeChildren and offload
mechanism under getFreeBlock is resolved because they will affect the
functionality this merge request is trying to achieve.

When a block goes OOW, it is released from the sequence, it will be
available to be reclaimed and the block is held by the eviction policy
for another sequence to acquire upon calling. On the other hand, we
want to potentially store the sequence for reuse. To safely achieve
this, the record of block ownership is done under
WindowBlockManager::getFreeBlock. If the block acquired was originally
owned by another sequence that is live inside the manager, then we
invalidate the sequence for store for reuse.

At the end of a sequence (when removeSequence is called toward it),
the KV cache manager will check if the sequence has all blocks not
reclaimed by another sequence. If so, then the sequence is safe to
be stored for reuse and store for reuse action will be performed.

Signed-off-by: eopXD <yuehtingc@nvidia.com>
 2025-10-13 09:18:12 -07:00
Fanrong Li
1e0fbb776d
[TRTLLM-8536][feat] Update trtllm gen fmha kernels to support block sparse attention (#8301) 
Signed-off-by: Fanrong Li <23290157+lfr-0531@users.noreply.github.com>
 2025-10-13 05:54:48 -07:00
xiweny
5ce9719759
[https://nvbugs/5503138] [fix] Remove compile warnings (#8167) 
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
 2025-10-13 13:24:23 +08:00
Zhenhuan Chen
84d2f12818
[TRTLLM-6748][feat] add PDL support for more kernels (#7977) 
Signed-off-by: Zhenhuan Chen <chenzhh3671@gmail.com>
 2025-10-11 08:32:05 +08:00
Chuang Zhu
85f157f389
[None][fix] Add Lock to protect mReqeustToSession (#8085) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
Co-authored-by: Xianjie Qiao <5410381+qiaoxj07@users.noreply.github.com>
 2025-10-10 21:51:50 +08:00
Jonas Li
76a47c7bef
[None][fix] Enable FP8 ContextMLA on GB300 (#8080) 
Signed-off-by: Jonas Li <6110159+longlee0622@users.noreply.github.com>
 2025-10-10 10:20:46 +08:00
dongfengy
9f2a3ae88c
[None][fix] Restrict tinygemm use to certain SMs (#8182) 
Signed-off-by: Dongfeng Yu <dongfengy@nvidia.com>
Signed-off-by: dongfengy <99041270+dongfengy@users.noreply.github.com>
 2025-10-08 17:55:57 -07:00
xiweny
9298f1bdcc
[None] [test] Add B300 cases to CI (#8056) 
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
 2025-10-06 19:23:31 -07:00
Faraz
27a5091fcb
[None][feat] GPT-OSS Sm120/Sm121 Support (#7937) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
Signed-off-by: list <58580514+farazkh80@users.noreply.github.com>
Signed-off-by: Vincent Huang <vincenth@nvidia.com>
Co-authored-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
Co-authored-by: Vincent Huang <vincenth@nvidia.com>
 2025-10-06 16:59:06 -04:00
Jonas Yang CN
88ea2c4ee9
[TRTLLM-7349][feat] Adding new orchestrator type -- ray (#7520) 
Signed-off-by: Erin Ho <14718778+hchings@users.noreply.github.com>
Co-authored-by: Yuan Tong <13075180+tongyuantongyu@users.noreply.github.com>
Co-authored-by: Erin Ho <14718778+hchings@users.noreply.github.com>
 2025-10-04 08:12:24 +08:00
Robin Kobus
e2f69c5c23
[None] [refactor] Minor cleanup and improvements (#7619) 
Signed-off-by: Robin Kobus <19427718+Funatiq@users.noreply.github.com>
 2025-10-03 11:40:06 +02:00
Nikita Korobov
9b3d7cc3e6
[None][feat] Update TRT-LLM Gen MoE kernels (#7970) 
Signed-off-by: Nikita Korobov <14355239+nekorobov@users.noreply.github.com>
 2025-10-03 09:22:45 +08:00
Yilin Fan
01423ac183
[None][feat] perf_metrics endpoint functionality improvement (#8005) 
Signed-off-by: Yilin Fan <206948969+nv-yilinf@users.noreply.github.com>
Signed-off-by: nv-yilinf <206948969+nv-yilinf@users.noreply.github.com>
 2025-10-02 17:43:25 -07:00
Patrice Castonguay
fefa7d8fa3
[None][feat] Support for cancelling requests with disaggregation (#8114) 
Signed-off-by: Shunkang <182541032+Shunkangz@users.noreply.github.co>
Signed-off-by: Patrice Castonguay <55748270+pcastonguay@users.noreply.github.com>
Co-authored-by: Shunkang <182541032+Shunkangz@users.noreply.github.co>
 2025-10-02 11:04:26 -07:00
dongfengy
6568e565db
[TRTLLM-7775][feat] Integrate tinygemm2 for gpt-oss (#7916) 
Signed-off-by: Dongfeng Yu <dongfengy@nvidia.com>
Signed-off-by: dongfengy <99041270+dongfengy@users.noreply.github.com>
Co-authored-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
 2025-10-02 10:47:04 -07:00
yifeizhang-c
34d158b6da
[TRTLLM-6589][feat] Support CUDA graph for DeepEP (#7514) 
Signed-off-by: Yifei Zhang <219273404+yifeizhang-c@users.noreply.github.com>
 2025-10-02 10:13:24 -07:00
bhsueh_NV
38d6e4e60b
[None][feat] Support Qwen3 next (#7892) 
Signed-off-by: mengw <12670782+wm2012011492@users.noreply.github.com>
Signed-off-by: bhsueh <11360707+byshiue@users.noreply.github.com>
Signed-off-by: nv-guomingz <137257613+nv-guomingz@users.noreply.github.com>
Co-authored-by: nv-guomingz <137257613+nv-guomingz@users.noreply.github.com>
 2025-09-29 21:16:07 +08:00
xiweny
48e779ae8c
[https://nvbugs/5541494] [fix] add back missing sm100f bmm kernels (#8051) 
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
 2025-09-29 05:35:44 -04:00
Iman Tabrizian
33282351a2
[TRTLLM-6106][feat] Add support for KVCache transfer from KVCache reuse path (#6348) 
Signed-off-by: Iman Tabrizian <10105175+tabrizian@users.noreply.github.com>
 2025-09-27 19:29:30 -04:00
Jhao-Ting Chen
c33f43e13a
[https://nvbugs/5518713][fix] Trtllm-gen moe backend for blockwise fp8 ckpt (Qwen3-235B-A22B-FP8) (#7856) 
Signed-off-by: Jhao-Ting Chen <jhaotingc@nvidia.com>
 2025-09-26 14:29:32 -07:00
Yueh-Ting (eop) Chen
2db22fb4e5
[None][feature] Add environment variable to adjust block pool allocation ration under kv cache manager (#7923) 
By default, we allocate equal proportion shares of memory for all
window sizes (see the else case). With TRTLLM_WINDOW_SIZE_SHARES,
we can override this behavior to adjust the memory share of each
window size. For example, if we have window size of [512, 32768],
then setting TRTLLM_WINDOW_SIZE_SHARES=0.4,0.6 will be allocating
40% of the memory to window size 512 and 60% of the memory to window
size 32768.

Signed-off-by: eopXD <yuehtingc@nvidia.com>
 2025-09-26 14:09:01 +08:00
Chuang Zhu
f98fa0cf8b
[None][feat] Optimize kv cache transfer TEP (#7613) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-09-25 20:20:04 -07:00
Yuan Tong
fae83c387b
[#6102][fix] support non-system python installation (#7763) 
Signed-off-by: Yuan Tong <13075180+tongyuantongyu@users.noreply.github.com>
 2025-09-26 10:16:15 +08:00
Matthias Jouanneaux
eda1467061
[TRTLLM-5966][feat] Helix: add alltoall op (#6815) 
Signed-off-by: Matthias Jouanneaux <mjoux@nvidia.com>
 2025-09-25 07:18:29 -07:00
Yueh-Ting (eop) Chen
c5012423f5
[None][chore] Remove developer name in comment (#7981) 
Signed-off-by: eopXD <yuehtingc@nvidia.com>
 2025-09-25 06:43:38 -07:00
Guoming Zhang
202bed4574 [None][chroe] Rename TensorRT-LLM to TensorRT LLM for source code. (#7851) 
Signed-off-by: nv-guomingz <137257613+nv-guomingz@users.noreply.github.com>
Signed-off-by: Wangshanshan <30051912+dominicshanshan@users.noreply.github.com>
 2025-09-25 21:02:35 +08:00
Guoming Zhang
9f0f52249e [None][doc] Rename TensorRT-LLM to TensorRT LLM for homepage and the … (#7850) 
Signed-off-by: nv-guomingz <137257613+nv-guomingz@users.noreply.github.com>
Signed-off-by: Wangshanshan <30051912+dominicshanshan@users.noreply.github.com>
 2025-09-25 21:02:35 +08:00
Jinyang Yuan
b622cde5d5
[None][perf] Fix the tactic sorting in TrtllmGenBatchedGemmRunner::getValidConfigIndices (#7419) 
Signed-off-by: Jinyang Yuan <154768711+jinyangyuan-nvidia@users.noreply.github.com>
 2025-09-25 10:27:57 +02:00
Void
336c2ef540
[None][feat] DeepEP LL fp8 dispatch/combine (#7927) 
Signed-off-by: Yilin Zhang <18275976+yilin-void@users.noreply.github.com>
 2025-09-25 09:20:24 +08:00
sychen52
5a65af24cd
[OMNIML-2336][feat] Add NVFP4 x FP8 moe kernels (#7821) 
Signed-off-by: Shiyang Chen <shiychen@nvidia.com>
 2025-09-24 12:14:35 -07:00
Perkz Zheng
60101eb8a5
[None][fix] trtllm-gen cubins compiled with wrong arch. (#7953) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-09-24 04:13:36 -07:00
HuiGao-NV
29e63d3bc2
[https://nvbugs/5532248][fix] Fix fused_moe OOM (#7931) 
Signed-off-by: Hui Gao <huig@nvidia.com>
 2025-09-24 02:22:38 -07:00
qsang-nv
929ef4c474
[None][chore] remove cubins for ci cases (#7902) 
Signed-off-by: Qidi Sang <200703406+qsang-nv@users.noreply.github.com>
 2025-09-24 14:56:31 +08:00
xiweny
276d83c898
[https://nvbugs/5532225] [fix] MoE use stream-dependent workspace (#7940) 
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
 2025-09-24 14:44:27 +08:00
Yueh-Ting (eop) Chen
cf100933cc
[TRTLLM-6341][feature] Support SWA KV cache reuse (#6768) 
This merge request attempts to support more SWA KV cache functionality
inside the KV cache manager. Before this merge request, the KV cache for
sliding window attention (SWA) only holds "window size" number of blocks
and reuse them in a cyclic manner. We will not be able to utilize more
GPU memory with this design, leading to a limited max batch size
throughput. Additionally, we will not be able to support KV cache reuse
with this design.

In this MR, we change such behavior to let the manager write blocks in
a linear manner. With a linear block writing behavior, as the attention
window moves on, the out-of-window (OOW) blocks will be detached. Right
now for the sake of a correct feature first, we directly offload the
OOW block from the primary block pool (GPU memory) to the secondary
block pool (host memory). We will improve this in the future by
delegating the block movement to the eviction policy.

KV cache reuse for SWA is not developed in this merge request and will
be amended in a follow-up merge request.

Writing the blocks linearly, the maximum number of blocks allocated for
a sequence(`GenerationRequest`) is the "max sequence length" specified.
The `GenerationRequest` that stores the cache block bookkeeping
structure will now keep "max sequence length" tokens of blocks.

Given the above, main changes are (more context in the MR):
- Remove "cyclic" concept under the kv cache manager, such concept
  originally guards the block reuse under kv cache manager.
- Add detach mechanism and have it under `KVCacheManager::addToken`.
  Please note that detach is still guarded off for SWA when reuse
  is enabled. A follow-up merge request will proceed to improve this.
- Enforce "max sequence length" to be a non-optional parameter to
  the `KVCacheManager`/`BlockManager`
- Let all window size resource pool get identical proportion of memory
- Fix free memory calculation under `resource_manager.py`

Signed-off-by: eopXD <yuehtingc@nvidia.com>
Co-authored-by: Tomer Asida <tasida@nvidia.com>
 2025-09-24 14:28:24 +08:00
Jhao-Ting Chen
220dc01372
[None][feat] support JIT mha.cu for SPEC_DEC in runtime (#6078) 
Signed-off-by: Jhao-Ting Chen <jhaotingc@nvidia.com>
 2025-09-23 14:56:17 -07:00
Zheng Duan
e3c1a9409f
[TRTLLM-6549][fix] add kv cache time output back (#7798) 
Signed-off-by: zhengd-nv <200704041+zhengd-nv@users.noreply.github.com>
 2025-09-23 14:12:42 -04:00
Perkz Zheng
bb64e7462c
[None][fix] fix a bug with trtllm-gen kernels + attention sinks (#7919) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-09-23 00:32:04 -07:00
Pengbo Wang
a4b4ed4535
[None][fix] Fix and add test for TRTLLM MoE backend (#7755) 
Signed-off-by: Pengbo Wang <221450789+pengbowang-nv@users.noreply.github.com>
 2025-09-23 11:26:25 +08:00
Pengbo Wang
08cc7a041f
[https://nvbugs/5355128][fix] Add missing wgmma intrinsic for starcoder (#7643) 
Signed-off-by: Pengbo Wang <221450789+pengbowang-nv@users.noreply.github.com>
 2025-09-23 10:38:58 +08:00
Enwei Zhu
8330d5363a
[TRTLLM-8209][feat] Support new structural tag API (upgrade XGrammar to 0.1.25) (#7893) 
Signed-off-by: Enwei Zhu <21126786+syuoni@users.noreply.github.com>
 2025-09-23 09:10:09 +08:00
ChristinaZ
be576a3152
[None] [feat] Enable run_post_quant_allgather for MoE TRTLLM backend (#6794) 
Signed-off-by: Christina Zhang <83400082+ChristinaZ@users.noreply.github.com>
 2025-09-23 08:24:21 +08:00
Bo Deng
8cf95681e6
[TRTLLM-7989][infra] Bundle UCX and NIXL libs in the TRTLLM python package (#7766) 
Signed-off-by: Bo Deng <deemod@nvidia.com>
 2025-09-22 16:43:35 +08:00
brb-nv
8879ec4d35 [https://nvbugs/5501557][fix] Fix out-of-bounds vector access for model with multiple layer types (#7636) 
Signed-off-by: Balaram Buddharaju <169953907+brb-nv@users.noreply.github.com>
Signed-off-by: Wangshanshan <30051912+dominicshanshan@users.noreply.github.com>
 2025-09-22 14:28:38 +08:00
xiweny
822cb0115b
[TRTLLM-6286] [perf] Add NoSmem epilogue schedule and dynamic cluster shape for sm10x group gemm (#7757) 
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
Signed-off-by: djns99 <40156487+djns99@users.noreply.github.com>
Co-authored-by: djns99 <40156487+djns99@users.noreply.github.com>
 2025-09-21 11:38:17 +08:00
brb-nv
e10a027a03
[TRTLLM-7731][feat] KV cache transmission in disagg with CP on gen side (#7624) 
Signed-off-by: Balaram Buddharaju <169953907+brb-nv@users.noreply.github.com>
 2025-09-20 06:15:26 -07:00
Mike Iovine
8030b540ac
[https://nvbugs/5522462][fix] Fix FP8 scout illegal memory access (#7845) 
Signed-off-by: Mike Iovine <6158008+mikeiovine@users.noreply.github.com>
 2025-09-19 10:30:37 -04:00
Matthias Jouanneaux
1be7faef37
[TRTLLM-5966][feat] Helix: add custom position ids to MLA kernels (#6904) 
Signed-off-by: Matthias Jouanneaux <mjoux@nvidia.com>
Co-authored-by: brb-nv <169953907+brb-nv@users.noreply.github.com>
 2025-09-19 20:55:32 +08:00
Chuang Zhu
c98b9468af
[None][fix] get Local IP by connect remote (#7719) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-09-19 10:01:03 +08:00
xiweny
423e5f6a3c
[TRTLLM-6286] [feat] Update CUTLASS to 4.2 and enable SM103 group gemm (#7832) 
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
 2025-09-19 09:50:54 +08:00
Yuxian Qiu
d6ebcf7c4a
[TRTLLM-6994][feat] FP8 Context MLA integration (Cherry-pick https://github.com/NVIDIA/TensorRT-LLM/pull/6059 from release/1.1.0rc2) (#7610) 
Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com>
 2025-09-19 09:40:49 +08:00
QI JUN
7f87b278bc
[None][chore] remove generated fmha_cubin.h from source tree (#7836) 
Signed-off-by: junq <22017000+QiJune@users.noreply.github.com>
 2025-09-18 20:10:04 +08:00
Wanli Jiang
a7ca0fff54
[TRTLLM-6577][feat] Support nano_v2_vlm in pytorch backend (#7207) 
Signed-off-by: Wanli Jiang <35160485+Wanli-Jiang@users.noreply.github.com>
 2025-09-18 16:26:20 +08:00
Matthias Jouanneaux
022d77807d
[TRTLLM-5966][feat] Helix: make softmax stats pointer available to attention gen (#6865) 
Signed-off-by: Matthias Jouanneaux <mjoux@nvidia.com>
Co-authored-by: brb-nv <169953907+brb-nv@users.noreply.github.com>
 2025-09-18 05:01:24 +08:00
Shiyu Li
8bdbb48264
[https://nvbugs/5489015][fix] Support communicator split in MNNVL allreduce and fix the binding issues. (#7387) 
Signed-off-by: Shiyu Li <shili@nvidia.com>
 2025-09-17 07:43:20 +08:00
Iman Tabrizian
6ce0624208
[TRTLLM-8044][refactor] Rename data -> cache for cacheTransceiver (#7659) 2025-09-16 08:43:56 -04:00
xiweny
c076a02b38
[TRTLLM-4629] [feat] Add support of CUDA13 and sm103 devices (#7568) 
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
Signed-off-by: Tian Zheng <29906817+Tom-Zheng@users.noreply.github.com>
Signed-off-by: Daniel Stokes <dastokes@nvidia.com>
Signed-off-by: Zhanrui Sun <zhanruis@nvidia.com>
Signed-off-by: Xiwen Yu <xiweny@nvidia.com>
Signed-off-by: Jiagan Cheng <jiaganc@nvidia.com>
Signed-off-by: Yiqing Yan <yiqingy@nvidia.com>
Signed-off-by: Bo Deng <deemod@nvidia.com>
Signed-off-by: ZhanruiSunCh <184402041+ZhanruiSunCh@users.noreply.github.com>
Signed-off-by: xiweny <13230610+VALLIS-NERIA@users.noreply.github.com>
Co-authored-by: Tian Zheng <29906817+Tom-Zheng@users.noreply.github.com>
Co-authored-by: Daniel Stokes <dastokes@nvidia.com>
Co-authored-by: Zhanrui Sun <zhanruis@nvidia.com>
Co-authored-by: Jiagan Cheng <jiaganc@nvidia.com>
Co-authored-by: Yiqing Yan <yiqingy@nvidia.com>
Co-authored-by: Bo Deng <deemod@nvidia.com>
Co-authored-by: Zhanrui Sun <184402041+ZhanruiSunCh@users.noreply.github.com>
 2025-09-16 09:56:18 +08:00
jmydurant
7deefb3d2b
[TRTLLM-7192][feat] optimize MLA chunked prefill && support fp8 mla chunked prefill (#7477) 
Signed-off-by: Mingyang Jiang <13463932+jmydurant@users.noreply.github.com>
 2025-09-15 21:43:49 +08:00
Zheng Duan
24fc1f9acf
[None][fix] using arrival time in llmapi when creating LlmRequest in pytorch workflow (#7553) 
Signed-off-by: zhengd-nv <200704041+zhengd-nv@users.noreply.github.com>
 2025-09-15 07:26:01 -04:00
Perkz Zheng
1b29c2e731
[None][feat] support gpt-oss with fp8 kv cache (#7612) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-09-15 02:17:37 +08:00
Fan - Yunfan
e3117731b3
[None][fix] Fix the incorrect header file import in dataType.h (#7133) 
Signed-off-by: fanyunfan <2569548856@qq.com>
Co-authored-by: fanyunfan <2569658856@qq.com>
Co-authored-by: Yunfan Fan <46273019+fyf2016@users.noreply.github.com>
Co-authored-by: Kanghwan <861393+karljang@users.noreply.github.com>
 2025-09-11 08:59:04 +08:00
Linda
a4312ba743
[https://nvbugs/5477359][fix] Nanobind: Allow none types for fields in result (#7672) 
Signed-off-by: Linda-Stadter <57756729+Linda-Stadter@users.noreply.github.com>
 2025-09-10 14:13:46 +01:00
Chuang Zhu
f412f5c4b0
[None][fix]UCX zmq ip support ipv6 (#7530) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-09-10 10:24:41 +08:00
Linda
0566df672d
[TRTLLM-6707][fix] nanobind fix for executor exit call (#7565) 
Signed-off-by: Linda-Stadter <57756729+Linda-Stadter@users.noreply.github.com>
 2025-09-09 14:56:04 +01:00
NVJiangShao
cc7593987b
[https://nvbugs/5434424][fix] A quick fix for the wrong output issue of SM89 blocked scaling batched GEMM when the input tensor is non-contiguous. (#7615) 
Signed-off-by: Jiang Shao <91270701+StudyingShao@users.noreply.github.com>
 2025-09-09 08:58:15 -04:00
William Tambellini
a6ed0d17d6
[#6798][fix] fix compilation error in ub_allocator in single device build (#6874) 
Signed-off-by: William Tambellini <wtambellini@sdl.com>
 2025-09-09 07:13:53 -04:00
Perkz Zheng
da6cb541a2
[None][feat] Optimize MLA kernels with separate reduction kernels (#7597) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-09-09 16:58:44 +08:00
William Tambellini
6ba1c8421c
[#6529][feat] CMake option to link statically with cublas/curand (#7178) 
Close #6529.

Signed-off-by: William Tambellini <wtambellini@sdl.com>
 2025-09-09 14:26:45 +08:00
Tomer Shmilovich
ecc0e687c6
[None][feat] Nixl support for GDS (#5488) 
Signed-off-by: Tomer Shmilovich <tshmilovich@nvidia.com>
Signed-off-by: Guy Lev <glev@nvidia.com>
Co-authored-by: Guy Lev <glev@nvidia.com>
 2025-09-09 13:00:38 +08:00
Yiqing Yan
5c616da2fd
[TRTLLM-5877][infra] Add fmha tests and auto trigger rules (#6050) 
Signed-off-by: Yiqing Yan <yiqingy@nvidia.com>
Co-authored-by: Yanchao Lu <yanchaol@nvidia.com>
 2025-09-09 11:33:09 +08:00
Chuang Zhu
77657a1c12
[TRTLLM-7361][feat] KV cache transfer for uneven pp (#7117) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-09-08 13:37:46 -04:00
dominicshanshan
c9dca69e1b
[None][chore] Mass integration of release/1.0 - 3rd (#7519) 
Signed-off-by: Nave Assaf <nassaf@nvidia.com>
Signed-off-by: Wangshanshan <30051912+dominicshanshan@users.noreply.github.com>
Signed-off-by: yechank <161688079+yechank-nvidia@users.noreply.github.com>
Signed-off-by: Balaram Buddharaju <169953907+brb-nv@users.noreply.github.com>
Signed-off-by: Iman Tabrizian <10105175+tabrizian@users.noreply.github.com>
Signed-off-by: qqiao <qqiao@nvidia.com>
Signed-off-by: Superjomn <328693+Superjomn@users.noreply.github.com>
Signed-off-by: Bo Deng <deemod@nvidia.com>
Signed-off-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
Signed-off-by: Yifei Zhang <219273404+yifeizhang-c@users.noreply.github.com>
Signed-off-by: Amit Zuker <203509407+amitz-nv@users.noreply.github.com>
Signed-off-by: Erin Ho <14718778+hchings@users.noreply.github.com>
Signed-off-by: Chenfei Zhang <chenfeiz@nvidia.com>
Signed-off-by: Christina Zhang <83400082+ChristinaZ@users.noreply.github.com>
Signed-off-by: Venky Ganesh <23023424+venkywonka@users.noreply.github.com>
Signed-off-by: Pamela <179191831+pamelap-nvidia@users.noreply.github.com>
Signed-off-by: Hui Gao <huig@nvidia.com>
Signed-off-by: Alexandre Milesi <30204471+milesial@users.noreply.github.com>
Signed-off-by: Shixiaowei02 <39303645+Shixiaowei02@users.noreply.github.com>
Signed-off-by: Michal Guzek <mguzek@nvidia.com>
Signed-off-by: peaceh <103117813+peaceh-nv@users.noreply.github.com>
Signed-off-by: nv-guomingz <137257613+nv-guomingz@users.noreply.github.com>
Signed-off-by: Wanli Jiang <35160485+Wanli-Jiang@users.noreply.github.com>
Signed-off-by: Patrice Castonguay <55748270+pcastonguay@users.noreply.github.com>
Signed-off-by: ruodil <200874449+ruodil@users.noreply.github.com>
Signed-off-by: Linda-Stadter <57756729+Linda-Stadter@users.noreply.github.com>
Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com>
Signed-off-by: Jiagan Cheng <jiaganc@nvidia.com>
Signed-off-by: William Zhang <133824995+2ez4bz@users.noreply.github.com>
Signed-off-by: Dom Brown <3886319+DomBrown@users.noreply.github.com>
Co-authored-by: Nave Assaf <55059536+Naveassaf@users.noreply.github.com>
Co-authored-by: Yechan Kim <161688079+yechank-nvidia@users.noreply.github.com>
Co-authored-by: brb-nv <169953907+brb-nv@users.noreply.github.com>
Co-authored-by: Iman Tabrizian <10105175+Tabrizian@users.noreply.github.com>
Co-authored-by: Emma Qiao <qqiao@nvidia.com>
Co-authored-by: Yan Chunwei <328693+Superjomn@users.noreply.github.com>
Co-authored-by: Bo Deng <deemod@nvidia.com>
Co-authored-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
Co-authored-by: yifeizhang-c <219273404+yifeizhang-c@users.noreply.github.com>
Co-authored-by: amitz-nv <203509407+amitz-nv@users.noreply.github.com>
Co-authored-by: Erin <14718778+hchings@users.noreply.github.com>
Co-authored-by: chenfeiz0326 <chenfeiz@nvidia.com>
Co-authored-by: ChristinaZ <83400082+ChristinaZ@users.noreply.github.com>
Co-authored-by: Venky <23023424+venkywonka@users.noreply.github.com>
Co-authored-by: Pamela Peng <179191831+pamelap-nvidia@users.noreply.github.com>
Co-authored-by: HuiGao-NV <huig@nvidia.com>
Co-authored-by: milesial <milesial@users.noreply.github.com>
Co-authored-by: Shi Xiaowei <39303645+Shixiaowei02@users.noreply.github.com>
Co-authored-by: Michal Guzek <moraxu@users.noreply.github.com>
Co-authored-by: peaceh-nv <103117813+peaceh-nv@users.noreply.github.com>
Co-authored-by: Guoming Zhang <137257613+nv-guomingz@users.noreply.github.com>
Co-authored-by: Wanli Jiang <35160485+Wanli-Jiang@users.noreply.github.com>
Co-authored-by: pcastonguay <55748270+pcastonguay@users.noreply.github.com>
Co-authored-by: ruodil <200874449+ruodil@users.noreply.github.com>
Co-authored-by: Linda <57756729+Linda-Stadter@users.noreply.github.com>
Co-authored-by: Zhanrui Sun <184402041+ZhanruiSunCh@users.noreply.github.com>
Co-authored-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com>
Co-authored-by: Jiagan Cheng <jiaganc@nvidia.com>
Co-authored-by: William Zhang <133824995+2ez4bz@users.noreply.github.com>
Co-authored-by: Larry <197874197+LarryXFly@users.noreply.github.com>
Co-authored-by: Sharan Chetlur <116769508+schetlur-nv@users.noreply.github.com>
Co-authored-by: Dom Brown <3886319+DomBrown@users.noreply.github.com>
 2025-09-08 14:03:04 +08:00
Raayan Dhar
bae9560e62
[https://nvbugs/5448767][fix] sync termination of requests across PP ranks (#7455) 
Signed-off-by: raayandhar <rdhar@nvidia.com>
Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
Co-authored-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
 2025-09-07 08:45:49 -04:00
xiweny
0fdc6c7278
[TRTLLM-4629] [feat] trtllm-gen kernels support sm103 (#7570) 
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
 2025-09-07 10:04:10 +08:00
Chang Liu
23500b55c3
[TRTLLM-7398][feat] Support KV cache salting for secure KV cache reuse (#7106) 
Signed-off-by: Chang Liu (Enterprise Products) <9713593+chang-l@users.noreply.github.com>
Signed-off-by: Chang Liu <9713593+chang-l@users.noreply.github.com>
 2025-09-06 17:58:32 -04:00
Shunkangz
bddf183e15
[None][feat] Add Request specific exception (#6931) 
Signed-off-by: Shunkang <182541032+Shunkangz@users.noreply.github.co>
 2025-09-04 18:43:42 -04:00
sychen52
98a1bffb7c
[OMNIML-2336][feat] Add NVFP4 x FP8 (#6809) 
Signed-off-by: Shiyang Chen <shiychen@nvidia.com>
 2025-09-04 09:03:38 -07:00
Enwei Zhu
1745102e72
[TRTLLM-7027][feat] Fuse d2t to logitsBitmaskKernel and fix a race condition in one-model spec (#7481) 
Signed-off-by: Enwei Zhu <21126786+syuoni@users.noreply.github.com>
Co-authored-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
 2025-09-04 23:30:14 +08:00
Yao Yao
c1aa7f31d9
[None][fix] Fix a numerical stability issue for XQA with spec dec (#7114) 
Signed-off-by: Yao Yao <lowsfer@users.noreply.github.com>
 2025-09-03 20:40:05 -04:00
Enwei Zhu
5ff3a65b23
[TRTLLM-7028][feat] Enable guided decoding with speculative decoding (part 2: one-model engine) (#6948) 
Signed-off-by: Enwei Zhu <21126786+syuoni@users.noreply.github.com>
 2025-09-03 15:16:11 -07:00
Daniel Stokes
109f27265c
[None][perf] Add MOE support for dynamic cluster shapes and custom epilogue schedules (#6126) 
Signed-off-by: djns99 <40156487+djns99@users.noreply.github.com>
 2025-09-02 21:54:43 -04:00
Tian Zheng
1b9c4cc2f7
[None][fix] Fix nanobind failure (#7425) 
Signed-off-by: Tian Zheng <29906817+Tom-Zheng@users.noreply.github.com>
 2025-09-01 17:26:40 -04:00
Yilin Fan
261ffacfa4 [https://nvbugs/5412562][feat] Allocate MoE workspace only when necessary (release/1.0 retargeted) (#6955) 
Signed-off-by: Yilin Fan <206948969+nv-yilinf@users.noreply.github.com>
Signed-off-by: Wangshanshan <30051912+dominicshanshan@users.noreply.github.com>
 2025-09-01 11:02:31 +08:00
Tian Zheng
e257cb3533
[None][feat] Support NVFP4 KV Cache (#6244) 
Signed-off-by: Tian Zheng <29906817+Tom-Zheng@users.noreply.github.com>
 2025-09-01 09:24:52 +08:00
brb-nv
43cb50f788
[None][feat] Update TargetInfo to accommodate CP in disagg (#7224) 
Signed-off-by: Balaram Buddharaju <169953907+brb-nv@users.noreply.github.com>
 2025-08-29 15:56:20 -04:00
Richard Huo
ce580ce4f5
[None][feat] KV Cache Connector API (#7228) 
Signed-off-by: jthomson04 <jwillthomson19@gmail.com>
Signed-off-by: richardhuo-nv <rihuo@nvidia.com>
Co-authored-by: jthomson04 <jwillthomson19@gmail.com>
Co-authored-by: Iman Tabrizian <10105175+Tabrizian@users.noreply.github.com>
Co-authored-by: Sharan Chetlur <116769508+schetlur-nv@users.noreply.github.com>
 2025-08-28 23:09:27 -04:00
Daniel Stokes
e0253ee805
[None][perf] Disable Swap AB when num tokens exceeds N dimension (#7104) 
Signed-off-by: djns99 <40156487+djns99@users.noreply.github.com>
 2025-08-28 21:29:55 -04:00
Yuan Tong
ccb800f909
[TRTLLM-7457][ci] Update unittest parallel config (#7297) 
Signed-off-by: Yuan Tong <13075180+tongyuantongyu@users.noreply.github.com>
 2025-08-29 09:28:04 +08:00
Zongfei Jing
53163bf1df
[TRTLLM-6876][feat] Add low precision all2all for mnnvl (#7155) 
Signed-off-by: Zongfei Jing <20381269+zongfeijing@users.noreply.github.com>
 2025-08-28 18:26:16 +08:00
dongxuy04
abdb2735be
[None][fix] Fix possible hang issue in WideEP and move some tests to pre-merge (#7262) 
Signed-off-by: Dongxu Yang <78518666+dongxuy04@users.noreply.github.com>
 2025-08-27 01:39:24 -04:00
Jin Li
028235404b
[TRTLLM-6633][feat] Padding for piecewise cudagraph (#6750) 
Signed-off-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
 2025-08-26 18:31:33 -04:00
Void
040f4c70d3
[None][perf] Accelerate global scale calculations for deepEP fp4 combine (#7126) 
Signed-off-by: Yilin Zhang <18275976+yilin-void@users.noreply.github.com>
 2025-08-27 00:13:13 +08:00
Zhou Yuxin
f01101f687
[None][feat] Hopper Fp8 context mla (#7116) 
Signed-off-by: Yuxin <yuxinz@nvidia.com>
 2025-08-26 17:10:20 +08:00
qixiang-99
b165f8bc97
fix/improve kvcache allocation in PyTorch runtime (#5933) 
Signed-off-by: qixiang-99 <203170375+qixiang-99@users.noreply.github.com>
 2025-08-26 12:40:22 +08:00
Bo Li
bf1b958f1a
[TRTLLM-7319][perf] Fuse slicing into MoE. (#6728) 
Signed-off-by: Bo Li <22713281+bobboli@users.noreply.github.com>
Signed-off-by: Sergey Klevtsov <sklevtsov@nvidia.com>
Co-authored-by: Sergey Klevtsov <sklevtsov@nvidia.com>
 2025-08-25 16:52:30 -04:00
Robin Kobus
31979aefac
[None] [ci] Reorganize CMake and Python integration test infrastructure for C++ tests (#6754) 
Signed-off-by: Robin Kobus <19427718+Funatiq@users.noreply.github.com>
 2025-08-24 20:53:17 +02:00
dongxuy04
19a0ea363b
[TRTLLM-6743][feat] Optimize and refactor alltoall in WideEP (#6973) 
Signed-off-by: Dongxu Yang <78518666+dongxuy04@users.noreply.github.com>
Signed-off-by: Fred Wei <20514172+WeiHaocheng@users.noreply.github.com>
Signed-off-by: Dongxu Yang <dongxuy@nvidia.com>
Co-authored-by: Fred Wei <20514172+WeiHaocheng@users.noreply.github.com>
 2025-08-24 08:15:29 -04:00
Robin Kobus
37543a9ad7
[None][refactor] Simplify decoder state initialization for speculative decoding (#6869) 
Signed-off-by: Robin Kobus <19427718+Funatiq@users.noreply.github.com>
 2025-08-22 18:44:17 +02:00
Linda
898f37faa0
[None][feat] Enable nanobind as the default binding library (#6608) 
Signed-off-by: Linda-Stadter <57756729+Linda-Stadter@users.noreply.github.com>
 2025-08-22 09:48:41 +02:00
dominicshanshan
6f245ec78b
[None][chore] Mass integration of release/1.0 (#6864) 
Signed-off-by: Stanley Sun <190317771+StanleySun639@users.noreply.github.com>
Signed-off-by: Wangshanshan <30051912+dominicshanshan@users.noreply.github.com>
Signed-off-by: ruodil <200874449+ruodil@users.noreply.github.com>
Signed-off-by: Yiqing Yan <yiqingy@nvidia.com>
Signed-off-by: Yanchao Lu <yanchaol@nvidia.com>
Signed-off-by: Balaram Buddharaju <169953907+brb-nv@users.noreply.github.com>
Signed-off-by: Ivy Zhang <25222398+crazydemo@users.noreply.github.com>
Signed-off-by: Bo Deng <deemod@nvidia.com>
Signed-off-by: Chang Liu <9713593+chang-l@users.noreply.github.com>
Signed-off-by: Stefan Niebler <82932102+stnie@users.noreply.github.com>
Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com>
Signed-off-by: Superjomn <328693+Superjomn@users.noreply.github.com>
Signed-off-by: qqiao <qqiao@nvidia.com>
Signed-off-by: yechank <161688079+yechank-nvidia@users.noreply.github.com>
Signed-off-by: William Zhang <133824995+2ez4bz@users.noreply.github.com>
Signed-off-by: raayandhar <rdhar@nvidia.com>
Co-authored-by: Stanley Sun <190317771+StanleySun639@users.noreply.github.com>
Co-authored-by: ruodil <200874449+ruodil@users.noreply.github.com>
Co-authored-by: Yiqing Yan <yiqingy@nvidia.com>
Co-authored-by: Yanchao Lu <yanchaol@nvidia.com>
Co-authored-by: brb-nv <169953907+brb-nv@users.noreply.github.com>
Co-authored-by: Ivy Zhang <25222398+crazydemo@users.noreply.github.com>
Co-authored-by: Larry <197874197+LarryXFly@users.noreply.github.com>
Co-authored-by: Bo Deng <deemod@nvidia.com>
Co-authored-by: Guoming Zhang <137257613+nv-guomingz@users.noreply.github.com>
Co-authored-by: Stefan Niebler <82932102+stnie@users.noreply.github.com>
Co-authored-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com>
Co-authored-by: Yan Chunwei <328693+Superjomn@users.noreply.github.com>
Co-authored-by: Emma Qiao <qqiao@nvidia.com>
Co-authored-by: Yechan Kim <161688079+yechank-nvidia@users.noreply.github.com>
Co-authored-by: 2ez4bz <133824995+2ez4bz@users.noreply.github.com>
Co-authored-by: Raayan Dhar <58057652+raayandhar@users.noreply.github.com>
Co-authored-by: Zhanrui Sun <184402041+ZhanruiSunCh@users.noreply.github.com>
 2025-08-22 09:25:15 +08:00
Daniel Stokes
f7c597ec40
[None][perf] Make finalize fusion part of the tactic selection logic (#6915) 
Signed-off-by: djns99 <40156487+djns99@users.noreply.github.com>
 2025-08-21 14:08:03 -07:00
brb-nv
9a2b44d0f2
[None][chore] No-op changes to support context parallelism in disaggregated serving later (#7063) 
Signed-off-by: Balaram Buddharaju <169953907+brb-nv@users.noreply.github.com>
 2025-08-21 08:21:27 -07:00
Yuan Tong
90bfc8cc29
[https://nvbugs/5453827][fix] Fix RPATH of th_common shared library to find pip-installed NCCL (#6984) 
Signed-off-by: Yuan Tong <13075180+tongyuantongyu@users.noreply.github.com>
 2025-08-21 17:58:30 +08:00
ChristinaZ
c7269ea93a
[https://nvbugs/5392414] [fix] Add customized default routing method (#6818) 
Signed-off-by: Christina Zhang <83400082+ChristinaZ@users.noreply.github.com>
 2025-08-21 16:58:41 +08:00
Yao Yao
cbcea33279
[fix]: use safeInitRowMax instead of fp32_lowest to avoid NaN (#7087) 
Signed-off-by: Yao Yao <lowsfer@users.noreply.github.com>
 2025-08-20 22:12:21 -07:00
Fan - Yunfan
41ff4901ee
[None][fix] Fix const modifier inconsistency in log function declaration/implementation (#6679) 
Signed-off-by: fanyunfan <2569548856@qq.com>
Co-authored-by: fanyunfan <2569658856@qq.com>
Co-authored-by: Yunfan Fan <46273019+fyf2016@users.noreply.github.com>
 2025-08-21 11:08:11 +08:00
BatshevaBlack
9f51f8d20c
[None][infra] Upgrade UCX to v1.19.x and NIXL to 0.5.0 (#7024) 
Signed-off-by: Batsheva Black <132911331+BatshevaBlack@users.noreply.github.com>
Signed-off-by: Bo Deng <deemod@nvidia.com>
Co-authored-by: Bo Deng <deemod@nvidia.com>
 2025-08-20 22:49:55 -04:00
Dom Brown
92daec1115
[TRTLLM-7348] [feat] Enable Cross-Attention to use XQA kernels for Whisper (#7035) 
Signed-off-by: Dom Brown <3886319+DomBrown@users.noreply.github.com>
 2025-08-20 10:11:25 -04:00
Yuhao Yao
8ac7dec623
[None][fix] Fix W4A8 MoE kernel issue (#7072) 
Signed-off-by: yuhyao <827623970@qq.com>
 2025-08-20 06:52:47 -04:00
Yueh-Ting (eop) Chen
020fed97b6
[TRTLLM-6341][chore] Preliminary refactors on the kv cache manager before supporting swa kv cache reuse (#6767) 
This MR is a preliminary MR for implementing the SWA reuse mechanism for
the kv cache manager. Please be aware that **no functional change is
intended** in this merge request. The purpose of the clean-up is to
decouple and remove existing functions for the up-coming SWA KV cache
reuse change to be more natural and easier to review.

Right now, (1) streamLLM, and (2) beam search with SWA, are broken. We
do not want to complicate the code base by stacking more features upon
something that does not work. This MR prunes out the logic and add
assertions so we can come back and re-support the broken feature and
remove the assertion.

Since streamLLM (sink attention) is broken now, assertion is added
under `KVCacheManager` ctor to guard for the value of
`mSinkBlockTokenLength` and `mSinkBubbleLength`. Compute logics relate
to it are pruned.

The beam search with SWA will still be broke when introducing the SWA
KV cache reuse. We will revisit this problem in the future.

On top of this, we should make an effort to update the [supporting matrix](https://github.com/NVIDIA/TensorRT-LLM/blob/feat/1.0_doc_dev/docs/source/1.0/features/feature-combination-matrix.md)
of the kv cache manager after merging the support of SWA KV cache reuse.

Changes are listed as following:
- Separate `KVCacheManager::updateToken` into `KVCacheManager::addToken`
  and `KVCacheManager::removeToken`. The functionality should be
  decoupled.
- Push utility `cacheSequenceBlockOffsets` and `cacheNewBlockOffset` from
  `KVCacheManager` down to `WindowBlockManager`. `KVCacheManager`-exposed
  functions should be real utilities that users of the structure can
  leverage. Implementation-detailed function calls should not exist at
  this level.
- Simplify "is shared last context block" logic under
  `KVCacheManager::addSequence`.

Since no functional change is intended in this merge request, no test
case is added. Several comments are added for future test coverage
reminder.

For `LlmRequestTest.ParamTest`, `streaming=True` is commented out
because we guard sink attention with assertion now.

In `capacitySchedulerTest`, `addToken` action to `crossKVCacheManager`
is removed because in encoder-decoder model, generation tokens are
added only to the decoder and not to the encoder.

Signed-off-by: eopXD <yuehtingc@nvidia.com>
 2025-08-20 13:57:57 +08:00
zhhuang-nv
7e135d2ea7
[None][feat] Use Separate QKV Input Layout for Context MLA (#6538) 
Signed-off-by: Zhen Huang <145532724+zhhuang-nv@users.noreply.github.com>
 2025-08-19 22:04:48 +08:00
amitz-nv
a54c53652b
[TRTLLM-7263][fix] Prevent recreation of cublas handles in lora_grouped_gemm every call (#6968) 
Signed-off-by: Amit Zuker <203509407+amitz-nv@users.noreply.github.com>
 2025-08-19 15:39:56 +03:00
Zero Zeng
953f4fd69e
[None][fix] acceptance rate calculation fix in benchmark_serving (#6746) 
Signed-off-by: Zero Zeng <38289304+zerollzeng@users.noreply.github.com>
 2025-08-19 17:29:36 +08:00
Martin Marciniszyn Mehringer
425dad01fd
[None][fix] Clean up linking to CUDA stub libraries in build_wheel.py (#6823) 
Signed-off-by: Linda-Stadter <57756729+Linda-Stadter@users.noreply.github.com>
Signed-off-by: Martin Marciniszyn Mehringer <11665257+MartinMarciniszyn@users.noreply.github.com>
Co-authored-by: Linda-Stadter <57756729+Linda-Stadter@users.noreply.github.com>
 2025-08-18 11:20:51 -04:00
ChristinaZ
55f4f2d80c
[None] [fix] Fix the macro name (#6983) 
Signed-off-by: Christina Zhang <83400082+ChristinaZ@users.noreply.github.com>
 2025-08-18 03:08:32 -04:00
ChristinaZ
1e72721e8c
[None][feat] Add single block version renormalized routing kernel (#6756) 
Signed-off-by: Christina Zhang <83400082+ChristinaZ@users.noreply.github.com>
 2025-08-17 13:47:13 +08:00
bhsueh_NV
85cbd0263b
[None][feat] Support Yarn on Qwen3 (#6785) 
Signed-off-by: bhsueh <11360707+byshiue@users.noreply.github.com>
 2025-08-17 07:21:29 +08:00
Fan - Yunfan
22d59a6f61
[None][fix] Using RAII to automatically manage the allocation and release of va_list for potential resource leak (#6758) 
Signed-off-by: fanyunfan <2569548856@qq.com>
Co-authored-by: fanyunfan <2569658856@qq.com>
Co-authored-by: Yunfan Fan <46273019+fyf2016@users.noreply.github.com>
Co-authored-by: Yuan Tong <13075180+tongyuantongyu@users.noreply.github.com>
 2025-08-16 15:19:19 +08:00
Yuening Li
1f8ae2b2db
[TRTLLM-5863][feat] Support MoE INT8 Weight-Only-Quantization in PyTorch Workflow (#6629) 
Signed-off-by: Yuening Li <62227368+yueningl@users.noreply.github.com>
 2025-08-15 17:15:49 -04:00
yifeizhang-c
4127d77678
[https://nvbugs/5394392][fix] Enlarge scheduler capacity under disagg bs == 1 (#6537) 
Signed-off-by: Yifei Zhang <219273404+yifeizhang-c@users.noreply.github.com>
 2025-08-15 09:52:06 -07:00
Perkz Zheng
6037fe3716
[https://nvbugs/5394685][fix] proper fix for the accuracy issue in 2CTA MLA kernels (#6941) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-08-15 23:29:36 +08:00
peaceh-nv
1c1d5d2495
[https://nvbugs/5451373][fix] : Fix the accuracy issue when using FP8 context MLA (#6881) 
Signed-off-by: peaceh <103117813+peaceh-nv@users.noreply.github.com>
 2025-08-15 16:53:56 +08:00
Yanchao Lu
3a987891d8
[TRTLLM-7141][infra] Use repo mirrors to avoid intermittent network failures (#6836) 
Signed-off-by: Yanchao Lu <yanchaol@nvidia.com>
Co-authored-by: Zhanrui Sun <184402041+ZhanruiSunCh@users.noreply.github.com>
 2025-08-15 11:16:07 +08:00
Wanli Jiang
9a133e9b41
[https://nvbugs/5415862][fix] Update cublas as 12.9.1 and cuda memory alignment as 256 (#6501) 
Signed-off-by: Wanli Jiang <35160485+Wanli-Jiang@users.noreply.github.com>
 2025-08-15 11:10:59 +08:00
Yunfan Fan
11d08c33af
[None][fix] Fix responsibility boundary between the assert and tllmException files (#6723) 
Signed-off-by: fanyunfan <2569548856@qq.com>
Co-authored-by: Yuan Tong <13075180+tongyuantongyu@users.noreply.github.com>
 2025-08-15 10:34:49 +08:00
Perkz Zheng
11d89a3732
[https://nvbugs/5394685][fix] using static scheduler 2CTA MLA as WAR for an accuracy issue (#6896) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-08-15 08:51:04 +08:00
jmydurant
4200fa46d1
[None][feat] Add support for Hopper MLA chunked prefill (#6655) 
Signed-off-by: Mingyang Jiang <13463932+jmydurant@users.noreply.github.com>
 2025-08-14 10:39:26 +08:00
Linda
eb4ed18a63
[None][fix] max_num_sequences argument in nanobind (#6862) 
Signed-off-by: Linda-Stadter <57756729+Linda-Stadter@users.noreply.github.com>
 2025-08-13 19:16:17 -04:00
Perkz Zheng
58f7783ea4
[https://nvbugs/5394685][fix] the bug with spec-decoding + SWA && an accuracy issue related to 2CTA MLA (#6834) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-08-13 13:55:56 -07:00
Tin-Yin Lai
6c52bb07ff
[https://nvbugs/5302040][feat] Add whisper support (Bert Attention on SM100 and GPTAttention for cross attention on SM100) (#5527) 
Signed-off-by: tinyinl <tinyinl@nvidia.com>
 2025-08-13 11:19:13 -07:00
Perkz Zheng
0fad6029f7
[TRTLLM-7093][fix] the perf regression to cvt_fp4 kernels (#6851) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
 2025-08-13 19:13:40 +08:00
Void
1d80df0955
[None][feat] DeepEP LL combine FP4 (#6822) 
Signed-off-by: Yilin Zhang <18275976+yilin-void@users.noreply.github.com>
 2025-08-13 04:20:21 -04:00
Zhou Yuxin
50e5e725e9
[https://nvbugs/5412456][fix] Fix an illegal instruction was encountered (#6776) 
Signed-off-by: Zhou Yuxin <yuxinz@nvidia.com>
 2025-08-13 15:45:59 +08:00
Robin Kobus
45c7518032
[None][refactor] Simplify decoder state initialization (#6559) 
Signed-off-by: Robin Kobus <19427718+Funatiq@users.noreply.github.com>
 2025-08-12 21:44:41 +02:00
QI JUN
8845e0f065
[None][fix] fix ci (#6814) 2025-08-12 02:21:50 -07:00
Liao Lanyu
f7c13a4aa7
[TRTLLM-6906][chore] Using pybind to bind functions in thop/attentionOp (#6745) 
Signed-off-by: Lanyu Liao <lancelly@users.noreply.github.com>
 2025-08-12 16:45:16 +08:00
Sergey Klevtsov
27fc35175e
[None][feat] CUTLASS MoE FC2+Finalize fusion (#3294) 
Signed-off-by: Sergey Klevtsov <sklevtsov@nvidia.com>
 2025-08-12 15:56:48 +08:00
bhsueh_NV
83dbc6c75d
[TRTLLM-5532][feat] store the block of context request into kv cache (#6683) 
Signed-off-by: bhsueh <11360707+byshiue@users.noreply.github.com>
 2025-08-11 16:14:52 +08:00
Martin Marciniszyn Mehringer
9a8195ef88
fix: Ensure that Python stub generation works against libnvidia-ml stubs (#6188) 
Signed-off-by: Martin Marciniszyn Mehringer <11665257+MartinMarciniszyn@users.noreply.github.com>
 2025-08-11 09:18:17 +02:00
Chuang Zhu
c566a8d2a2
[None][fix] fix same pp disagg (#6730) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-08-10 22:45:15 -04:00
Yueh-Ting (eop) Chen
199f306984
[None][chore][kv cache manager] Dead code elimination, we no longer record/fetch through WindowBlockManager:: mContextBlocksByHash (#6249) 
No functional change is intended in this MR.

`WindowBlockManager::mCachedBlocksRoot` is now who is responsible
for the bookkeeping of the `KVCacheBlock`, and the `mNextBlocks` is
now the actual hash map that fetches the block.

The `mEnableHashKey` knob and related hashing is removed.

Signed-off-by: eopXD <yuehtingc@nvidia.com>
 2025-08-10 09:10:10 -04:00
Ziyi Xiong
de472828b9
[TRTLLM-6637][feat] Resolve KV cache divergence issue (#6628) 
Signed-off-by: ziyixiong-nv <219238287+ziyixiong-nv@users.noreply.github.com>
 2025-08-09 23:15:04 +08:00
Chuang Zhu
e251f7c00b
[None][fix]revert kvcache transfer (#6709) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-08-08 07:18:53 -04:00
Zheng Duan
ebdc43e69d
[None][feat] move kv cache measure into transfer session (#6633) 
Signed-off-by: zhengd-nv <200704041+zhengd-nv@users.noreply.github.com>
 2025-08-08 17:49:22 +08:00
NVJiangShao
2f2f5cc72c
[TRTLLM-6744][feat] Remove input_sf swizzle for module WideEPMoE (#6231) 
Signed-off-by: Jiang Shao <91270701+StudyingShao@users.noreply.github.com>
 2025-08-08 11:13:42 +08:00
Daniel Cámpora
efca359b66
[TRTLLM-6785][feat] BREAKING CHANGE Enable TRTLLM sampler by default (#6216) 
Signed-off-by: Daniel Campora <961215+dcampora@users.noreply.github.com>
 2025-08-07 22:19:37 -04:00
Iman Tabrizian
82276167e6
[None][feat] Add NCCL Symmetric Integration for All Reduce (#4500) 
Signed-off-by: Iman Tabrizian <10105175+tabrizian@users.noreply.github.com>
 2025-08-07 17:28:14 -07:00
Yuan Tong
db8dc97b7b
[None][fix] Migrate to new cuda binding package name (#6700) 
Signed-off-by: Yuan Tong <13075180+tongyuantongyu@users.noreply.github.com>
 2025-08-07 16:29:55 -04:00
pcastonguay
453a06e6ab
[TRTLLM-6881][feat] Include attention dp rank info with KV cache events (#6563) 
Signed-off-by: Patrice Castonguay <55748270+pcastonguay@users.noreply.github.com>
 2025-08-07 14:17:07 +02:00
Enwei Zhu
1b9781e8e7
[TRTLLM-6409][feat] Enable guided decoding with speculative decoding (part 1: two-model engine) (#6300) 
Signed-off-by: Enwei Zhu <21126786+syuoni@users.noreply.github.com>
 2025-08-07 05:53:48 -04:00
peaceh-nv
8ec3b1de10
[None][feat] : Add FP8 context MLA support for SM120 (#6059) 
Signed-off-by: peaceh <103117813+peaceh-nv@users.noreply.github.com>
 2025-08-07 16:16:34 +08:00
hlu1
8207d5fd39
[None] [feat] Add model gpt-oss (#6645) 
Signed-off-by: Hao Lu <14827759+hlu1@users.noreply.github.com>
 2025-08-07 03:04:18 -04:00
amitz-nv
85af62184b
[TRTLLM-6683][feat] Support LoRA reload CPU cache evicted adapter (#6510) 
Signed-off-by: Amit Zuker <203509407+amitz-nv@users.noreply.github.com>
 2025-08-07 09:05:36 +03:00
Chuang Zhu
ee471df07c
[None][chore] optimize kv cache transfer for context TEP and gen DEP (#6657) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-08-07 11:36:05 +08:00
Zongfei Jing
0ff8df95b7
[https://nvbugs/5433581][fix] DeepGEMM installation on SBSA (#6588) 
Signed-off-by: Zongfei Jing <20381269+zongfeijing@users.noreply.github.com>
 2025-08-06 16:44:21 +08:00
Ransiki
19b7524ff6
[None][feat] Add vLLM KV Pool support for XQA kernel (#6013) 
Signed-off-by: Ransiki Zhang <ransikiz@nvidia.com>
 2025-08-06 09:29:37 +08:00
ixlmar
1ebceb790d
[TRTLLM-5508][feat] check input tokens + improve error handling (#5170) 
Signed-off-by: ixlmar <206748156+ixlmar@users.noreply.github.com>
 2025-08-05 18:27:43 +01:00
Haohang Huang
c9eebcb454
[TRTLLM-6674][feat] (Breaking Change) Hopper SWA non-cyclic kernels + KV reuse + Spec Dec (#6379) 
Signed-off-by: Haohang Huang <31998628+symphonylyh@users.noreply.github.com>
Signed-off-by: symphonylyh <31998628+symphonylyh@users.noreply.github.com>
 2025-08-05 07:47:41 +00:00
Chuang Zhu
4d040b50b7
[None][chore] ucx establish connection with zmq (#6090) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-08-05 02:50:45 -04:00
Olya Kozlova
13cc1c4878
[TRTLLM-5271][feat] best_of/n for pytorch workflow (#5997) 
Signed-off-by: Olya Kozlova <okozlova@nvidia.com>
 2025-08-04 14:08:06 +02:00
Bruce-Lee-LY
8c82ee2803
[fix] xqa precision for fp16/bf16 kv cache (#6573) 
Signed-off-by: Bruce-Lee-LY <yong-li14@tsinghua.org.cn>
Co-authored-by: Bruce-Lee-LY <yong-li14@tsinghua.org.cn>
 2025-08-04 14:34:20 +08:00
Yuan Tong
a2f271c8e0
[TRTLLM-4406][feat] LLM sleep & wakeup Part 1: virtual device memory (#5034) 
Signed-off-by: Yuan Tong <13075180+tongyuantongyu@users.noreply.github.com>
 2025-08-04 13:51:01 +08:00
Perkz Zheng
03430ed379 [https://nvbugspro.nvidia.com/bug/5415268] fix illegal smem access with chunked attention (#6401) 
Signed-off-by: Perkz Zheng <67892460+PerkzZheng@users.noreply.github.com>
Co-authored-by: Sharan Chetlur <116769508+schetlur-nv@users.noreply.github.com>
 2025-08-04 11:19:58 +08:00
Jhao-Ting Chen
6edaa23c1c
[None][feat] Multi-block mode for Hopper spec dec XQA kernel (#4416) 
Signed-off-by: Jhao-Ting Chen <jhaotingc@nvidia.com>
 2025-08-03 14:31:33 -07:00
Chuang Zhu
542f552d0b
use cudaSetDevice to create context ,fix nvbug 5394497 (#6403) 
Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>
 2025-08-03 13:32:55 -04:00
Robin Kobus
918fedf952
[None][refactor] Simplify finish reasons handling in DecoderState (#6524) 
Signed-off-by: Robin Kobus <19427718+Funatiq@users.noreply.github.com>
 2025-08-02 07:17:43 +02:00
yunruis
a20ab5cbdb
[https://nvbugs/5381276][fix] fix warning for fused_a_gemm (#6402) 
Signed-off-by: yunruis <205571022+yunruis@users.noreply.github.com>
 2025-08-01 09:37:21 -04:00
Yang Li
ac23f4a80d
[TRTLLM-4279] fix: Add a protection test for checking trtllm custom ops (#6515) 
Signed-off-by: Yang Li <56944310+yali-arch@users.noreply.github.com>
Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
 2025-08-01 15:59:09 +08:00
Robin Kobus
d3c14682f0
refactor: Remove unused buffers and bindings from sampler (#6484) 
Signed-off-by: Robin Kobus <19427718+Funatiq@users.noreply.github.com>
 2025-08-01 00:43:03 -04:00
Yao Yao
942e080415
[fix] Fix missing fields in xqa kernel cache key (#6282) 
Signed-off-by: Yao Yao <lowsfer@users.noreply.github.com>
 2025-08-01 10:41:26 +08:00
Jaedeok Kim
fbee279909
fix: remove duplicate layer multiplication in KV cache size calculation (#6481) 
Signed-off-by: Jaedeok Kim <jaedeokk@nvidia.com>
 2025-07-31 22:34:34 -04:00
Michal Guzek
b8719fe96d
[nvbug/5374773] chore: Update nanobind with fail_fast_on_attention_window_too_large changes (#6491) 
Signed-off-by: Michal Guzek <mguzek@nvidia.com>
 2025-07-31 20:25:29 +01:00
Enwei Zhu
4b299cb77e
feat: Support structural tag in C++ runtime and upgrade xgrammar to 0.1.21 (#6408) 
Signed-off-by: Enwei Zhu <21126786+syuoni@users.noreply.github.com>
 2025-07-31 09:53:52 +08:00
pcastonguay
e7ae5e2824
feat: Add support for disaggregation with pp with pytorch backend (#6369) 
Signed-off-by: Patrice Castonguay <55748270+pcastonguay@users.noreply.github.com>
Signed-off-by: raayandhar <rdhar@nvidia.com>
Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
Signed-off-by: pcastonguay <55748270+pcastonguay@users.noreply.github.com>
Co-authored-by: raayandhar <rdhar@nvidia.com>
Co-authored-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
 2025-07-30 09:42:13 -04:00