[https://nvbugs/5655885][fix] fix invalid instruction error in 2shot ar kernel on Ampere

[yilin-void]

https://github.com/NVIDIA/TensorRT-LLM/issues/9086 .acquire and .release qualifiers for fence instruction require sm_90 or higher

Summary by CodeRabbit

• Performance
  • Enhanced barrier synchronization efficiency for newer GPU architectures
  • Maintained backward compatibility with older GPU models

_{✏️ Tip: You can customize this high-level summary in your review settings.}

Description

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.

• [x] Please check this after reviewing the above items as appropriate for this PR.

GitHub Bot Help

/bot [-h] ['run', 'kill', 'skip', 'reuse-pipeline'] ...

Provide a user friendly way for developers to interact with a Jenkins server.

Run /bot [-h|--help] to print this help message.

See details below for each supported subcommand.

run [--reuse-test (optional)pipeline-id --disable-fail-fast --skip-test --stage-list "A10-PyTorch-1, xxx" --gpu-type "A30, H100_PCIe" --test-backend "pytorch, cpp" --add-multi-gpu-test --only-multi-gpu-test --disable-multi-gpu-test --post-merge --extra-stage "H100_PCIe-TensorRT-Post-Merge-1, xxx" --detailed-log --debug(experimental)]

Launch build/test pipelines. All previously running jobs will be killed.

--reuse-test (optional)pipeline-id (OPTIONAL) : Allow the new pipeline to reuse build artifacts and skip successful test stages from a specified pipeline or the last pipeline if no pipeline-id is indicated. If the Git commit ID has changed, this option will be always ignored. The DEFAULT behavior of the bot is to reuse build artifacts and successful test results from the last pipeline.

--disable-reuse-test (OPTIONAL) : Explicitly prevent the pipeline from reusing build artifacts and skipping successful test stages from a previous pipeline. Ensure that all builds and tests are run regardless of previous successes.

--disable-fail-fast (OPTIONAL) : Disable fail fast on build/tests/infra failures.

--skip-test (OPTIONAL) : Skip all test stages, but still run build stages, package stages and sanity check stages. Note: Does NOT update GitHub check status.

--stage-list "A10-PyTorch-1, xxx" (OPTIONAL) : Only run the specified test stages. Examples: "A10-PyTorch-1, xxx". Note: Does NOT update GitHub check status.

--gpu-type "A30, H100_PCIe" (OPTIONAL) : Only run the test stages on the specified GPU types. Examples: "A30, H100_PCIe". Note: Does NOT update GitHub check status.

--test-backend "pytorch, cpp" (OPTIONAL) : Skip test stages which don't match the specified backends. Only support [pytorch, cpp, tensorrt, triton]. Examples: "pytorch, cpp" (does not run test stages with tensorrt or triton backend). Note: Does NOT update GitHub pipeline status.

--only-multi-gpu-test (OPTIONAL) : Only run the multi-GPU tests. Note: Does NOT update GitHub check status.

--disable-multi-gpu-test (OPTIONAL) : Disable the multi-GPU tests. Note: Does NOT update GitHub check status.

--add-multi-gpu-test (OPTIONAL) : Force run the multi-GPU tests in addition to running L0 pre-merge pipeline.

--post-merge (OPTIONAL) : Run the L0 post-merge pipeline instead of the ordinary L0 pre-merge pipeline.

--extra-stage "H100_PCIe-TensorRT-Post-Merge-1, xxx" (OPTIONAL) : Run the ordinary L0 pre-merge pipeline and specified test stages. Examples: --extra-stage "H100_PCIe-TensorRT-Post-Merge-1, xxx".

--detailed-log (OPTIONAL) : Enable flushing out all logs to the Jenkins console. This will significantly increase the log volume and may slow down the job.

--debug (OPTIONAL) : Experimental feature. Enable access to the CI container for debugging purpose. Note: Specify exactly one stage in the stage-list parameter to access the appropriate container environment. Note: Does NOT update GitHub check status.

For guidance on mapping tests to stage names, see docs/source/reference/ci-overview.md and the scripts/test_to_stage_mapping.py helper.

kill

kill

Kill all running builds associated with pull request.

skip

skip --comment COMMENT

Skip testing for latest commit on pull request. --comment "Reason for skipping build/test" is required. IMPORTANT NOTE: This is dangerous since lack of user care and validation can cause top of tree to break.

reuse-pipeline

reuse-pipeline

Reuse a previous pipeline to validate current commit. This action will also kill all currently running builds associated with the pull request. IMPORTANT NOTE: This is dangerous since lack of user care and validation can cause top of tree to break.

[yilin-void]

/bot run

[coderabbitai[bot]]

📝 Walkthrough

Walkthrough

Enhances barrier synchronization in Barrier::sync by introducing arch-specific memory operations for CUDA compute capability 9.0 and newer, replacing previous store operations with inline-assembly st.global.relaxed.sys and adding fence.release.sys while preserving legacy code path for older architectures.

Changes

Cohort / File(s)	Change Summary
Barrier synchronization optimization cpp/tensorrt_llm/kernels/communicationKernels/allReduceFusionKernels.cu	Added architecture-specific conditional compilation for CUDA >= 9.0 to use inline-assembly store (st.global.relaxed.sys) and release fence instead of previous st_flag path. Legacy code path retained for older architectures. ABA-avoidance loop now arch-gated.

Estimated code review effort

🎯 3 (Moderate) | ⏱️ ~20–30 minutes

• Inline assembly correctness: Verify st.global.relaxed.sys syntax and memory ordering semantics match CUDA PTX documentation
• Architecture checks: Confirm __CUDA_ARCH__ conditional logic correctly routes older vs. newer paths
• Memory fence semantics: Validate that fence.release.sys provides required synchronization for newer architectures
• Backward compatibility: Ensure legacy st_flag path remains functional for compute capabilities < 9.0

Pre-merge checks and finishing touches

❌ Failed checks (2 warnings)

Check name	Status	Explanation	Resolution
Description check	⚠️ Warning	The PR description contains only the template structure with no actual content provided. Required sections like Description and Test Coverage are empty, and the PR checklist is unchecked except for one box.	Fill in the Description section explaining the issue and solution, add Test Coverage details, and ensure the PR checklist items are properly reviewed and checked.
Docstring Coverage	⚠️ Warning	Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%.	You can run @coderabbitai generate docstrings to improve docstring coverage.

✅ Passed checks (1 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title follows the required format with ticket ID and type, and clearly identifies the fix for an invalid instruction error in the 2shot AR kernel on Ampere architecture.

✨ Finishing touches

• [ ] 📝 Generate docstrings

🧪 Generate unit tests (beta)

• [ ] Create PR with unit tests
• [ ] Post copyable unit tests in a comment

[!TIP]

📝 Customizable high-level summaries are now available in beta!

You can now customize how CodeRabbit generates the high-level summary in your pull requests — including its content, structure, tone, and formatting.

• Provide your own instructions using the high_level_summary_instructions setting.
• Format the summary however you like (bullet lists, tables, multi-section layouts, contributor stats, etc.).
• Use high_level_summary_in_walkthrough to move the summary from the description to the walkthrough section.

Example instruction:

"Divide the high-level summary into five sections:

1. 📝 Description — Summarize the main change in 50–60 words, explaining what was done.
2. 📓 References — List relevant issues, discussions, documentation, or related PRs.
3. 📦 Dependencies & Requirements — Mention any new/updated dependencies, environment variable changes, or configuration updates.
4. 📊 Contributor Summary — Include a Markdown table showing contributions: | Contributor | Lines Added | Lines Removed | Files Changed |
5. ✔️ Additional Notes — Add any extra reviewer context. Keep each section concise (under 200 words) and use bullet or numbered lists for clarity."

Note: This feature is currently in beta for Pro-tier users, and pricing will be announced later.

---

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

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

[tensorrt-cicd]

PR_Github #25514 [ run ] triggered by Bot. Commit: 25c6bb8

[tensorrt-cicd]

PR_Github #25514 [ run ] completed with state SUCCESS. Commit: 25c6bb8 /LLM/main/L0_MergeRequest_PR pipeline #19320 completed with status: 'FAILURE'

[Funatiq]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #25553 [ run ] triggered by Bot. Commit: 25c6bb8

[tensorrt-cicd]

PR_Github #25553 [ run ] completed with state SUCCESS. Commit: 25c6bb8 /LLM/main/L0_MergeRequest_PR pipeline #19352 completed with status: 'FAILURE'

[yilin-void]

/bot run

[tensorrt-cicd]

PR_Github #25630 [ run ] triggered by Bot. Commit: fb356c1

[tensorrt-cicd]

PR_Github #25630 [ run ] completed with state FAILURE. Commit: fb356c1 LLM/main/L0_MergeRequest_PR #19418 (Blue Ocean) completed with status: ABORTED

[yilin-void]

/bot run

[tensorrt-cicd]

PR_Github #25808 [ run ] triggered by Bot. Commit: e008fc7

[tensorrt-cicd]

PR_Github #25808 [ run ] completed with state SUCCESS. Commit: e008fc7 /LLM/main/L0_MergeRequest_PR pipeline #19575 completed with status: 'FAILURE'

[yilin-void]

/bot run

[yilin-void]

/bot run

[tensorrt-cicd]

PR_Github #25976 [ run ] triggered by Bot. Commit: 0d095c6

[tensorrt-cicd]

PR_Github #25976 [ run ] completed with state SUCCESS. Commit: 0d095c6 /LLM/main/L0_MergeRequest_PR pipeline #19700 completed with status: 'FAILURE'

[Funatiq]

/bot run

[tensorrt-cicd]

PR_Github #26040 [ run ] triggered by Bot. Commit: 0d095c6

[tensorrt-cicd]

PR_Github #26040 [ run ] completed with state SUCCESS. Commit: 0d095c6 /LLM/main/L0_MergeRequest_PR pipeline #19764 completed with status: 'FAILURE'

[yilin-void]

/bot run

[tensorrt-cicd]

PR_Github #26096 [ run ] triggered by Bot. Commit: 0d095c6

[tensorrt-cicd]

PR_Github #26096 [ run ] completed with state FAILURE. Commit: 0d095c6 /LLM/main/L0_MergeRequest_PR pipeline #19815 completed with status: 'FAILURE'

[yilin-void]

/bot run

[tensorrt-cicd]

PR_Github #26124 [ run ] triggered by Bot. Commit: 3d01255

[tensorrt-cicd]

PR_Github #26124 [ run ] completed with state SUCCESS. Commit: 3d01255 /LLM/main/L0_MergeRequest_PR pipeline #19840 completed with status: 'FAILURE'

[ZhanruiSunCh]

unittest/_torch/modules/tests_lora_modules/test_lora_attention_pytorch_flow_vs_trt.py::TestLoraAttentionPytorchFlowVsTRT::test_lora_attention SKIP (https://nvbugs/5701421) The test failed in last pipeline is waived in main branch TOT, not need rebase. We will auto use the waive list of main branch TOT, I rerun the /bot run without rebase to reuse the passed stage of your last pipeline.

[ZhanruiSunCh]

/bot run

[tensorrt-cicd]

PR_Github #26385 [ run ] triggered by Bot. Commit: 3d01255

[tensorrt-cicd]

PR_Github #26385 [ run ] completed with state SUCCESS. Commit: 3d01255 /LLM/main/L0_MergeRequest_PR pipeline #20044 completed with status: 'FAILURE'