[https://nvbugs/5875031][fix] Compile XQA with sm_120f by pamelap-nvidia · Pull Request #12170 · NVIDIA/TensorRT-LLM

pamelap-nvidia · 2026-03-12T20:56:48Z

Summary by CodeRabbit

New Features
- Added support for CUDA architecture 1210.
Refactor
- Streamlined compilation process for improved build efficiency.
Chores
- Removed unnecessary build dependencies.

Description

We did 2 stage compilation for sm120f before to workaround some accuracy issue. Looks like we don't need to do that anymore and can compile with sm_120f e2e.

Test Coverage

PR Checklist

Please review the following before submitting your PR:

PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.
PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.
Test cases are provided for new code paths (see test instructions)
Any new dependencies have been scanned for license and vulnerabilities
CODEOWNERS updated if ownership changes
Documentation updated as needed
Update tava architecture diagram if there is a significant design change in PR.
The reviewers assigned automatically/manually are appropriate for the PR.
Please check this after reviewing the above items as appropriate for this PR.

GitHub Bot Help

To see a list of available CI bot commands, please comment /bot help.

pamelap-nvidia · 2026-03-12T20:59:35Z

/bot run --extra-stage "RTXPro6000D-PyTorch-Post-Merge-1,RTXPro6000D-4_GPUs-PyTorch-Post-Merge-1,RTXPro6000D-4_GPUs-PyTorch-Post-Merge-2,GB10-PyTorch-Post-Merge-1" --disable-fail-fast

tensorrt-cicd · 2026-03-12T21:09:36Z

PR_Github #38779 [ run ] triggered by Bot. Commit: 0a3f3c7 Link to invocation

tensorrt-cicd · 2026-03-13T03:05:38Z

PR_Github #38779 [ run ] completed with state SUCCESS. Commit: 0a3f3c7
/LLM/main/L0_MergeRequest_PR pipeline #30093 completed with status: 'SUCCESS'

CI Report

Link to invocation

coderabbitai · 2026-03-13T15:04:18Z

📝 Walkthrough

Walkthrough

The changes add support for CUDA architecture 1210 in compiler guards, consolidate NVRTC compilation from a two-stage PTX+cubin approach to a single NVRTC-based path, eliminate precompiled cubin data storage, and remove the NVPTX_LIB dependency from the build configuration.

Changes

Cohort / File(s)	Summary
CUDA Architecture 1210 Support `cpp/kernels/xqa/mha.cu`, `cpp/kernels/xqa/utils.cuh`	Extended conditional preprocessor guards to include `__CUDA_ARCH__ == 1210` alongside existing architectures 860, 890, and 1200 for code path selection and memory size definitions.
NVRTC Compilation Refactoring `cpp/tensorrt_llm/kernels/decoderMaskedMultiheadAttention/decoderXQAImplJIT/nvrtcWrapper/src/nvrtcWrapper.cpp`	Removed two-stage NVRTC compilation flow (PTX generation via NVRTC followed by cubin generation via nvPTXCompiler), eliminated precompiled cubin data storage for SM120/121, simplified architecture flag handling to target sm_120f directly, and consolidated to single NVRTC-based compilation path.
Build Configuration `cpp/tensorrt_llm/kernels/decoderMaskedMultiheadAttention/decoderXQAImplJIT/nvrtcWrapper/CMakeLists.txt`	Removed `NVPTX_LIB` from the PUBLIC link libraries list while retaining `NVRTC_LIB`, `NVRTC_BUILTINS_LIB`, `CUDA_DRV_LIB`, and `CUDA_RT_LIB`.

Estimated code review effort

🎯 3 (Moderate) | ⏱️ ~30 minutes

🚥 Pre-merge checks | ✅ 1 | ❌ 2

❌ Failed checks (2 warnings)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.
Description check	⚠️ Warning	The PR description is incomplete. While it explains the change (removing two-stage compilation), it lacks proper PR title format and leaves Test Coverage section empty.	Add a properly formatted PR title (e.g., '[https://nvbugs/5875031][fix] Compile XQA with sm_120f') and document test coverage that validates the sm_120f compilation changes.

✅ Passed checks (1 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly and specifically describes the main change: compiling XQA with sm_120f instead of using a two-stage workaround.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing Touches

🧪 Generate unit tests (beta)

Create PR with unit tests
Post copyable unit tests in a comment

📝 Coding Plan

Generate coding plan for human review comments

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

Signed-off-by: Pamela <179191831+pamelap-nvidia@users.noreply.github.com>

pamelap-nvidia · 2026-03-17T15:12:38Z

/bot run --extra-stage "RTXPro6000D-PyTorch-Post-Merge-1,RTXPro6000D-4_GPUs-PyTorch-Post-Merge-1,RTXPro6000D-4_GPUs-PyTorch-Post-Merge-2,GB10-PyTorch-Post-Merge-1" --disable-fail-fast

tensorrt-cicd · 2026-03-17T15:20:00Z

PR_Github #39280 [ run ] triggered by Bot. Commit: 2dbc7a3 Link to invocation

tensorrt-cicd · 2026-03-17T20:27:47Z

PR_Github #39280 [ run ] completed with state SUCCESS. Commit: 2dbc7a3
/LLM/main/L0_MergeRequest_PR pipeline #30527 completed with status: 'SUCCESS'

CI Report

Link to invocation

jhaotingc

LGTM, thanks for the fix!!

tburt-nv

No license changes, approved

Signed-off-by: Pamela <179191831+pamelap-nvidia@users.noreply.github.com>

## Description Ports the architecture-table part of NVIDIA/TensorRT-LLM#12170 ("Compile XQA with sm_120f") to ORT's XQA decode kernels: `__CUDA_ARCH__ == 1210` (sm_121 — GB10 / DGX Spark and other Blackwell Arm devices) is now listed next to sm_120 in the two compile-time arch tables. sm_121 has the same 99 KB opt-in shared-memory limit as sm_120 and uses the same K/V tile configuration. This is explicit-parity hardening, not a user-visible bug fix: sm_121 already built and ran through the "unknown architecture" `#else` default, which has the same values as sm_120. See Testing for the evidence that the generated code is identical before and after. ## Summary of Changes ### XQA arch tables | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/xqa/mha_impl.cuh` | Add `__CUDA_ARCH__ == 1210` to the `860 / 890 / 1200` branch that selects `preferedKHeadPartBytes = 64` and `cacheVTileSeqLen = 32` (previously reached only via the `#else` default). Mention sm_121 in the related comment and in the shared-memory-limit hint text. | | `onnxruntime/contrib_ops/cuda/bert/xqa/utils.cuh` | Add `__CUDA_ARCH__ == 1210` to the 99 KB `kMAX_SMEM_SIZE` branch. Before, 1210 fell through to the 48 KB default. `kMAX_SMEM_SIZE` is currently only used by a `#if 0` `static_assert`, so this has no runtime effect today; it just keeps the table correct. | ### Not ported from TRT-LLM #12170 (not applicable) The rest of #12170 changes TensorRT-LLM's NVRTC JIT wrapper (`-arch=sm_120f`, dropping the compute_89 PTX to sm_120 two-stage compile through nvPTXCompiler, and the `nvptxcompiler_static` link). ORT compiles XQA ahead of time with nvcc, so there is nothing to port. In TRT-LLM the default arch branch was an `#error` and `kMAX_SMEM_SIZE` was undefined for 1210, which is why sm_121 failed to compile there; ORT's `#else` fallbacks mean it never did. ## Do the shipped libraries include sm_121? Yes, so no packaging or CMake change is needed: - Linux aarch64 CUDA 13 wheel: `87-real;89-real;90-real;100-real;103-real;110-real;120-real;121-real` (`py-cuda13-packaging-pipeline.yml`). - Unified Linux aarch64 wheel (`py-cuda13-unified-packaging-pipeline.yml`): passes no `-a`, so it takes the `build_linux_python_package.sh` default for CUDA 13.0 on aarch64, which is the same list. - CUDA plugin (`plugin-cuda13-pipeline.yml` / `plugin-cuda-packaging-stage.yml`): Linux aarch64 uses the same list; Windows ARM64 uses `120-real;121-real`. - x86_64 lists only carry `120-real` (sm_121 parts are Arm). - CMake normalizes `121-real` to `121a-real`, and XQA sources are compiled for every selected arch >= 80, so XQA gets sm_121a SASS (no PTX). The CUDA plugin's XQA target (`onnxruntime_providers_cuda_plugin_xqa`) uses the same `MIN_SM 80` filter, with no virtual replacement (that option is only used for the plugin's separate LLM target), so it gets the same SASS on Linux and Windows. ## Testing Built the XQA object library of the CUDA provider (`onnxruntime_providers_cuda_xqa`, CUDA 13.0) with `CMAKE_CUDA_ARCHITECTURES='120-real;121-real'`. The configure log shows `CMAKE_CUDA_ARCHITECTURES: 120a-real;121a-real`. - All 44 XQA objects contain both `sm_120a` and `sm_121a` cubins (`cuobjdump -lelf`), with 158 `kernel_mha` kernels each. Per-kernel register / shared / local / constant usage is identical between sm_120a and sm_121a. - A/B against `main` (without this change): SASS is identical for all 44 objects on sm_120a and sm_121a (600,232 instructions compared; all 88 cubins byte-identical), and also identical on sm_90a for 8 representative translation units (loader, paged, spec-dec, fp8/int8/int4, head sizes 64-512). So existing archs are unaffected and sm_121 behaves exactly as before. - The XQA kernels are present in the final `libonnxruntime_providers_cuda.so` (checked on an sm_90a build: 158 `kernel_mha` kernels). - `lintrunner` and `clang-format` are clean. - Not run on real sm_121 hardware (none available), so a smoke test on a DGX Spark would be welcome. To reproduce the cubin check: ``` python tools/ci_build/build.py --build_dir build/xqa_sm121 --config Release --use_cuda \ --cuda_home <cuda13> --cudnn_home <cudnn> --build_shared_lib --skip_tests --update \ --cmake_extra_defines 'CMAKE_CUDA_ARCHITECTURES=120-real;121-real' ninja -C build/xqa_sm121/Release onnxruntime_providers_cuda_xqa cuobjdump -lelf <obj dir>/xqa_loader_fp16_128.cu.o # lists sm_120a and sm_121a cubins ``` ## Motivation and Context Keeps ORT's XQA arch tables in sync with TensorRT-LLM (NVIDIA/TensorRT-LLM#12170) so sm_121 is handled explicitly rather than by the unknown-arch fallback, and so the `kMAX_SMEM_SIZE` table is right if that check is re-enabled. Independent of the XQA race fixes in #33703. Overlap with #33373: its commit "xqa: treat SM121 like SM120 in the per-architecture constants" makes the same two `|| __CUDA_ARCH__ == 1210` edits with identical spelling. This PR carries the XQA part on its own so it does not have to wait for that larger CMake change. The two PRs merge cleanly in either order (checked with `git merge-tree` against the head of #33373), so whichever lands second just has a no-op hunk. The extra edits here (the sm_121 mention in the comment and the shared-memory-limit hint text) are in different hunks. ## Checklist - [ ] Tests added/updated (N/A: arch-table-only change with identical generated code; verified by cubin comparison above) - [x] No breaking changes - [x] Documentation updated (N/A: only an error-message hint and comments changed) Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description - **Output ownership:** Only the final reducing CTA publishes multi-block attention output; other CTAs cannot overwrite it with partial results. - **Synchronization:** All threads latch the final-CTA flag before its shared memory is reused. Four more `cp.async` drains / warp barriers on the same path (after GEMM0, before `mergeAndSaveOutTile`, at the start of `loadBufAsync` and after the `waitGroup<1>()` in the multi-block merge loop) were compiled only for `XQA_PAGED_INT4`; they are now enabled for every variant. Two more `__syncwarp()`s, after the K-tile wait in the GEMM0 loop and after the V-tile wait in `syncVTileLoad`, are ported from [NVIDIA/TensorRT-LLM#18141](NVIDIA/TensorRT-LLM#18141): `cp.async.wait_group` only waits for the copies of the executing thread, but every lane reads tile data that other lanes of the warp copied. - **Regression coverage:** `XqaMultiBlockTest` (fp16 and bf16) runs 24 independent decode problems per launch on 4 concurrent streams, each with its own workspace and output, with K/V in mapped host memory. That is what makes the races observable; a single decode problem on one stream passes on the unpatched kernel. It covers head sizes 64, 128 and 256 (256 has a single GEMM1 warp group, so the CTA output is staged in a different shared-memory buffer), group sizes 1, 2 and 4, batch 1, 2 and 4, a sliding window with per-head attention sinks, 1 and 4 sub-sequences, direct launches and CUDA graph replay. Every variant captures a 4-stream fork/join graph and asserts that the multi-block grid is launched. The test is skipped below sm_80. ### Motivation and Context GQA decoding with aliased past/present KV caches can intermittently corrupt attention and change generated tokens, especially under CUDA graph capture. Eligible shared-cache decoding selects XQA independently of `sdpa_kernel`, so requesting Flash does not exclude the affected path. ### Testing The A100 section below was measured on the first revision of this PR (8 tests, without the two K / V `__syncwarp()`s and the head size 256 tests); the H200 section covers the current revision. #### A100-SXM4-80GB (sm_80), first revision CUDA 13.0, Release, built with `CMAKE_CUDA_ARCHITECTURES=80`, `onnxruntime_BUILD_UNIT_TESTS=ON`, `onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON`, targets `onnxruntime_provider_test` and `onnxruntime_providers_cuda_ut`: ```bash cd build/<config>/Release GTEST_FILTER='XqaMultiBlockTest*' ./onnxruntime_provider_test --gtest_filter='CUDA_EP_Unittest.All' ``` **The test detects the bug.** The previous version of the test passed in 24/24 processes on the unpatched kernel. The current version, 24 independent processes per row with all 8 tests in each: | `mha_impl.cuh` | Result | |---|---| | `main` (both race-fix hunks reverted) | **0/24 pass**: all 8 tests fail in every run, typically within the first few launches per stream | | the two race-fix hunks only | 24/24 pass | | this PR | 24/24 pass (192 test runs, none skipped) | **The four extra sync sites matter.** `compute-sanitizer --tool racecheck` on the fp16 head size 64 test, one repeat per configuration: | | errors | racing instruction pairs | pairs involving `cp.async` | |---|---|---|---| | the two race-fix hunks only | 634 | 347 | 328 | | this PR | 168 | 19 | 0 | The 19 remaining pairs (`LDSM`/`STS`, `LDS`/`STS`) are also present without this change. I believe they are the GEMM0 → GEMM1 tile hand-off, which the `xBar` / `ctaRowMax` mbarriers order and racecheck does not model, but I have not proven that. **Cost of the four sync sites.** CUDA graph replay of 100 back-to-back launches, median of 3 runs on one A100, 12 shapes (head size 64/128, group size 1/4, batch 1/8, length 505–2503): +0.0% to +1.1%, for example 15.57 → 15.68 µs and 25.16 → 25.42 µs. **Regression suites** (no failures; skipped cases need WebGPU, Flash Attention or a PagedAttention EP that this build does not enable, SM89+ for FP8, or SM90+ for cuDNN paged SDPA and H512 MEA): | Suites | Passed | Skipped | |---|---|---| | `GroupQueryAttention{,GrowingCache,Position,ShapeInference}Test.*` | 107 | 25 | | `PagedAttention.*`, `SparsePagedAttention.*` | 45 | 53 | | `*PagedAttentionXqa*`, `*SparsePagedAttentionInt8*` (INT8 XQA) | 12 | 2 (FP8) | | internal `Xqa*` (8 multi-block, 2 H512) | 10 | 0 | #### H200 (sm_90a), current revision CUDA 13.0, Release, built for `sm_90a` with the same options and targets; Flash Attention and memory-efficient attention are not compiled in. ORT runs the same `mha_impl.cuh` kernel on sm_90 (the Hopper `gmma_impl.cuh` is not built). | `mha_impl.cuh` | `XqaMultiBlockTest*` (10 tests) | |---|---| | `main` | **0/10 pass**: every test fails, including both head size 256 tests | | this PR before the K / V `__syncwarp()`s | 10/10 pass | | this revision | 10/10 pass in each of 7 parallel processes (one per GPU), none skipped | The two K / V `__syncwarp()`s are hardening, not a failure I can reproduce: the tests also pass without them on H200, and racecheck does not model intra-warp `cp.async` visibility. They follow from the `cp.async.wait_group` semantics and match [NVIDIA/TensorRT-LLM#18141](NVIDIA/TensorRT-LLM#18141), which includes them in its fix for the same race. The head size 256 tests cover the original race fix for a kernel configuration that was not tested before. **Cost of the two K / V `__syncwarp()`s.** CUDA graph replay of 50 back-to-back launches (min of 9 trials per run), median of 4 alternating A/B rounds, 10 shapes (head size 64/128/256, batch 1–32, length 1024–8192), fp16 and bf16: median change +0.02%, range −0.43% to +0.86%, within the round-to-round noise (about ±1–2%). For example 14.87 → 14.83 µs (head size 128, batch 1, length 8192), 12.01 → 12.06 µs (head size 64) and 20.30 → 20.47 µs (head size 256). **Regression suites on this revision:** | Suites | Passed | Skipped | Failed | |---|---|---|---| | `GroupQueryAttention{,GrowingCache,Position,ShapeInference}Test.*` | 108 | 24 | 0 | | `PagedAttention.*`, `SparsePagedAttention.*` | 47 | 46 | 5 | | `*PagedAttentionXqa*`, `*SparsePagedAttentionInt8*` (INT8 and FP8 XQA) | 14 | 0 | 0 | | internal `Xqa*` (10 multi-block, 2 H512) | 12 | 0 | 0 | The 5 `PagedAttention` failures (`Cuda_SkipInputSanitizerCudaGraphReplayChangingLengths`, `Cuda_SkipInputSanitizerRaggedSmallBatchParity`, `Cuda_XqaNativeFp16CacheMultiTokenBoundFallsBack`, `Cuda_XqaSpecDecInt8CacheHeadSize128NonCausalFallsBack`, `Cuda_XqaSpecDecInt8CacheHeadSize128UnsupportedWidthFallsBack`) are not caused by this PR: they expect a Flash Attention / memory-efficient fallback (`PagedAttention requires FlashAttention ...`), which this build does not compile in, and they fail the same way on `main` and on the PR before this revision. **Not covered:** GPUs other than sm_80 and sm_90a, Windows (the `_WIN32` environment-variable handling in the test was kept from the earlier revision but not compiled here), builds with Flash Attention enabled, and the Python GQA tests. ### Comparison with TensorRT-LLM Compared with [NVIDIA/TensorRT-LLM#18141](NVIDIA/TensorRT-LLM#18141) and every later change to `cpp/kernels/xqa` in TensorRT-LLM (the numbers in the table are TensorRT-LLM pull requests): | TensorRT-LLM change | Status here | |---|---| | [TRT-LLM#18141]: only the final reducing CTA writes the output, `isLastCta` is latched by all threads before shared memory is reused, GEMM0 `cp.async` drain | In this PR | | [TRT-LLM#18141]: `__syncwarp()`s inside the multi-block merge loop | Covered by the barriers at the start of `loadBufAsync` and after `waitGroup<1>()` that this PR enables for every variant | | [TRT-LLM#18141]: `__syncwarp()` after the K-tile and V-tile `cp.async` waits | **Added in this revision** | | [TRT-LLM#18141]: `gemm0ToPostProcBarrier` mbarrier | Not needed: the `__syncthreads()` at the end of GEMM0 pairs with the first `__syncthreads()` of the GEMM1 post-processing (`needMergeGlobal` or `mergeAndSaveOutTile`), which orders the same hand-off; warps that already exited are not waited for | | [TRT-LLM#18141] (sink handled by CTA 0 only) and [TRT-LLM#19220] (sink in the initial row max, zero-denominator guard) | Not applicable: the sink is still added once in the final reduction, with the merged row max. The sliding-window tests check per-head sinks of 4.0–6.0 against the analytic reference with 1 and 4 sub-sequences, and a sink applied twice would move the result by at least 13%, far above the tolerance | | [TRT-LLM#14459] (paged sliding window with unmapped pages) | Not needed: `LaunchSanitizeBlockTable` maps page ids outside `[-1, num_blocks)` to -1, `HeadPtr::operator+` turns them into `nullptr`, and `copyPartialHeadsAsync` zero-fills them | | [TRT-LLM#15739] (speculative-decoding sinks) | Already equivalent here | | [TRT-LLM#17976] (SM120 performance) | `writesOutput` is part of this PR. The V-page cursor rewrite is not needed: ORT only instantiates 128-token pages, and for every head size and `cacheVTileSeqLen`, `nbVItersPerXIter > 1` only occurs with `xIterSeqStride <= tokensPerPage`, where the old and the new cursor stepping are identical | | [TRT-LLM#12170] (`__CUDA_ARCH__ == 1210` in the XQA per-architecture tables) | Not needed for correctness: ORT's unknown-architecture `#else` default already gives sm_121 the sm_120 values. Ported as explicit parity in #33762 | | [TRT-LLM#10335] (performance), [TRT-LLM#11066] and the Hopper part of [TRT-LLM#14459] (`gmma_impl.cuh`, not built by ORT), [TRT-LLM#15312] (host runner) | Not applicable | [TRT-LLM#10335]: NVIDIA/TensorRT-LLM#10335 [TRT-LLM#11066]: NVIDIA/TensorRT-LLM#11066 [TRT-LLM#12170]: NVIDIA/TensorRT-LLM#12170 [TRT-LLM#14459]: NVIDIA/TensorRT-LLM#14459 [TRT-LLM#15312]: NVIDIA/TensorRT-LLM#15312 [TRT-LLM#15739]: NVIDIA/TensorRT-LLM#15739 [TRT-LLM#17976]: NVIDIA/TensorRT-LLM#17976 [TRT-LLM#18141]: NVIDIA/TensorRT-LLM#18141 [TRT-LLM#19220]: NVIDIA/TensorRT-LLM#19220  - Fixes #33562 --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

github-actions Bot assigned pamelap-nvidia Mar 12, 2026

pamelap-nvidia requested a review from pengbowang-nv March 13, 2026 14:57

pamelap-nvidia marked this pull request as ready for review March 13, 2026 14:59

pamelap-nvidia requested a review from jhaotingc March 16, 2026 00:12

pengbowang-nv requested a review from lowsfer March 17, 2026 02:38

pengbowang-nv reviewed Mar 17, 2026

View reviewed changes

Comment thread cpp/kernels/xqa/utils.cuh Resolved

Comment thread ...rt_llm/kernels/decoderMaskedMultiheadAttention/decoderXQAImplJIT/nvrtcWrapper/CMakeLists.txt Resolved

pamelap-nvidia requested a review from a team as a code owner March 17, 2026 05:06

pamelap-nvidia added 2 commits March 17, 2026 11:12

update xqa compilation

adfd3a1

Signed-off-by: Pamela <179191831+pamelap-nvidia@users.noreply.github.com>

cleanup

2dbc7a3

Signed-off-by: Pamela <179191831+pamelap-nvidia@users.noreply.github.com>

pamelap-nvidia force-pushed the sm120_xqa branch from cd65c77 to 2dbc7a3 Compare March 17, 2026 15:12

jhaotingc approved these changes Mar 17, 2026

View reviewed changes

pengbowang-nv approved these changes Mar 18, 2026

View reviewed changes

pamelap-nvidia enabled auto-merge (squash) March 18, 2026 03:32

tburt-nv approved these changes Mar 19, 2026

View reviewed changes

pamelap-nvidia merged commit e37493a into NVIDIA:main Mar 19, 2026
5 checks passed

pamelap-nvidia deleted the sm120_xqa branch March 30, 2026 18:49

longcheng-nv pushed a commit to longcheng-nv/TensorRT-LLM that referenced this pull request Mar 31, 2026

[https://nvbugs/5875031][fix] Compile XQA with sm_120f (NVIDIA#12170)

f0d1f6e

Signed-off-by: Pamela <179191831+pamelap-nvidia@users.noreply.github.com>

ConnorBaker mentioned this pull request Oct 14, 2026

[CUDA] Fix CUDA architecture gating for SM12x, family/arch-specific targets and -virtual entries microsoft/onnxruntime#33373

Closed

This was referenced Oct 30, 2026

Fix XQA reduction races in shared-cache GQA decoding microsoft/onnxruntime#33703

Merged

[CUDA] Add sm_121 to XQA arch-specific branches microsoft/onnxruntime#33762

Merged

Navigation Menu

Repository navigation

[https://nvbugs/5875031][fix] Compile XQA with sm_120f - #12170

[https://nvbugs/5875031][fix] Compile XQA with sm_120f#12170
pamelap-nvidia merged 2 commits into
NVIDIA:mainfrom
pamelap-nvidia:sm120_xqa

pamelap-nvidia commented Mar 12, 2026 •

edited by coderabbitai Bot

Loading

Uh oh!

pamelap-nvidia commented Mar 12, 2026

Uh oh!

tensorrt-cicd commented Mar 12, 2026

Uh oh!

tensorrt-cicd commented Mar 13, 2026

Uh oh!

coderabbitai Bot commented Mar 13, 2026 •

edited

Loading

Walkthrough

Changes

Estimated code review effort

❌ Failed checks (2 warnings)

Uh oh!

Uh oh!

Uh oh!

pamelap-nvidia commented Mar 17, 2026

Uh oh!

tensorrt-cicd commented Mar 17, 2026

Uh oh!

tensorrt-cicd commented Mar 17, 2026

Uh oh!

jhaotingc left a comment

Uh oh!

tburt-nv left a comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

pamelap-nvidia commented Mar 12, 2026 • edited by coderabbitai Bot Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Description

Test Coverage

PR Checklist

GitHub Bot Help

Uh oh!

pamelap-nvidia commented Mar 12, 2026

Uh oh!

tensorrt-cicd commented Mar 12, 2026

Uh oh!

tensorrt-cicd commented Mar 13, 2026

Uh oh!

coderabbitai Bot commented Mar 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Walkthrough

Changes

Estimated code review effort

❌ Failed checks (2 warnings)

Uh oh!

Uh oh!

Uh oh!

pamelap-nvidia commented Mar 17, 2026

Uh oh!

tensorrt-cicd commented Mar 17, 2026

Uh oh!

tensorrt-cicd commented Mar 17, 2026

Uh oh!

jhaotingc left a comment

Choose a reason for hiding this comment

Uh oh!

tburt-nv left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

pamelap-nvidia commented Mar 12, 2026 •

edited by coderabbitai Bot

Loading

coderabbitai Bot commented Mar 13, 2026 •

edited

Loading