[libcu++] Skip fabric status decoding when a shared_memory_barrier wait has no report by SUKUNA-AI · Pull Request #12131 · NVIDIA/cccl

SUKUNA-AI · 2026-10-21T16:39:44Z

Description

operation_status::get_error_count() and for_each_error() called the device-side fabric status decoder even when the completed wait produced no report. On a GeForce RTX 5070 Ti (sm_120) with driver 617.42 (CUDA 13.4), cudaFabricOpErrorStatusCount() returns cudaErrorNotSupported, so __verify_fabric_status() trapped although there was nothing to decode.

With this change, __error_count() returns 0 without calling the decoder when there is no report. When a report is present nothing changes: classify() (which requires a report) and the decoding of real error reports still go through the driver.

Testing (local RTX 5070 Ti, CUDA 13.4, WSL2, gcc 16):

With shared_memory_barrier.pass.cpp un-waived locally, the test fails before this change and passes after it.
All 28 cuda/barrier lit tests pass.
compute-sanitizer racecheck and memcheck are clean on the test.
compute-sanitizer synccheck reports Missing wait on this test both before and after the change. A plain inline-PTX kernel gets the same report as soon as its mbarrier is initialized with layout::v1 (clean with layout v0), so this looks like a compute-sanitizer 2026.3 issue that is unrelated to this change.

The test stays waived in this PR: as described in #11935, older CI drivers cannot load its module.

Checklist

New or existing tests cover these changes. (Existing shared_memory_barrier.pass.cpp, still waived in CI because of Add CUDA Forward Compatibility package to our CI runners to allow testing new CUDA features #11935.)
The documentation is up to date with these changes. (No user-visible API change.)

copy-pr-bot · 2026-10-21T16:46:01Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

coderabbitai · 2026-10-21T16:48:55Z

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Source mode: Auto
Repository: NVIDIA/cccl
Branch: main
File: .coderabbit.yaml
Commit: 9f6ed13
Layers: Repository YAML
Review profile: CHILL
Plan: Enterprise
Run ID: dd2eed6b-f234-4d13-a9de-ada454c2aa0b

📥 Commits

Reviewing files that changed from the base of the PR and between d95ea5e and 8936c00.

📒 Files selected for processing: 1

libcudacxx/include/cuda/__barrier/shared_memory_barrier.h

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.

📝 Summary

Summary by CodeRabbit

Bug Fixes
- Barrier status now reports zero errors when a completed wait has no report.
- When a report is present and includes fabric error information, the error count is decoded as before.

operation_status::__error_count now returns zero without invoking the fabric status decoder when no report is present. When a report is present and its source encodes fabric errors, the method continues to decode the error count.

Changes: Shared memory barrier status

Layer / File(s)	Summary
Guard fabric status decoding `libcudacxx/include/cuda/__barrier/shared_memory_barrier.h`	`operation_status::__error_count` returns zero without decoding fabric status when no report is present.

Assessment against linked issues

Objective	Addressed	Explanation
[`#12130`] No-report waits return zero errors without invoking the fabric-status decoder.	✅
[`#12130`] `for_each_error()` does not invoke the callback when no report is present.	❓	The supplied change summary describes `__error_count`, but does not describe the callback path.

Possibly related issues

NVIDIA/cccl#12130: The change skips fabric-status decoding when no report is present, addressing the reported decoder failure.

Suggested reviewers: fbusato

Priority: ➖ Normal

Change: Bug fix · Severity of issue fixed: Medium

Merge Risk: ⚪ Minimal · up to 8936c

No concrete merge-blocking risk remains in the reviewed change.

- [ ] Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

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

…report operation_status::get_error_count() and for_each_error() called the device-side fabric status decoder even when the completed wait produced no report. On a GeForce RTX 5070 Ti (sm_120) with driver 617.42 (CUDA 13.4), cudaFabricOpErrorStatusCount() returns cudaErrorNotSupported, so __verify_fabric_status() trapped although there was nothing to decode. Return 0 without calling the decoder when there is no report.

pciolkosz · 2026-10-21T20:48:48Z

/ok to test c9b0e2a

copy-pr-bot · 2026-10-21T20:48:52Z

/ok to test c9b0e2a

@pciolkosz, there was an error processing your request: E2

See the following link for more information: https://docs.gha-runners.nvidia.com/cpr/e/2/

pciolkosz

Good catch! I think I assumed someone would call has_report() first, but I think doesn't hurt to check it again in the status checks

pciolkosz · 2026-10-21T20:50:18Z

/ok to test c9b0e2a

SUKUNA-AI · 2026-10-22T06:35:41Z

The only failing job here, cub.test.device.radix_sort_runs_on.lid_0 in the sm86-only build, is unrelated to this PR: on main that test aborts in any single-architecture build because it passes RunsOn a compute capability that was not compiled for. Details in #12154, fix in #12155.

pciolkosz · 2026-10-22T20:23:56Z

/ok to test d95ea5e

SUKUNA-AI · 2026-10-23T05:46:42Z

Thanks for the review! The only failure in the last run was clang-tidy in cudax/test/coop/group/make_group.cu (unrelated to this PR), which #12184 fixed on main after the previous merge. I merged main again (8936c00), so a new /ok to test should come back green.

pciolkosz · 2026-10-23T16:22:32Z

/ok to test 8936c00

github-actions · 2026-10-23T18:19:14Z

🥳 CI Workflow Results

🟩 Finished in 1h 52m: Pass: 100%/228 | Total: 3d 05h | Max: 1h 51m | Hits: 99%/629481

See results here.

github-actions · 2026-10-23T18:20:43Z

Successfully created backport PR for branch/3.6.x:

[Backport branch/3.6.x] [libcu++] Skip fabric status decoding when a shared_memory_barrier wait has no report #12228

…report (#12131) (#12228) operation_status::get_error_count() and for_each_error() called the device-side fabric status decoder even when the completed wait produced no report. On a GeForce RTX 5070 Ti (sm_120) with driver 617.42 (CUDA 13.4), cudaFabricOpErrorStatusCount() returns cudaErrorNotSupported, so __verify_fabric_status() trapped although there was nothing to decode. Return 0 without calling the decoder when there is no report. (cherry picked from commit 4811a1c) Co-authored-by: Ilya Koryaev <ilyakoryaev@yandex.ru> Co-authored-by: SUKUNA-AI <180781973+SUKUNA-AI@users.noreply.github.com> Co-authored-by: pciolkosz <pciolkosz@nvidia.com>

SUKUNA-AI requested a review from a team as a code owner October 21, 2026 16:39

SUKUNA-AI requested a review from pciolkosz October 21, 2026 16:39

github-project-automation Bot added this to CCCL Oct 21, 2026

github-project-automation Bot moved this to Todo in CCCL Oct 21, 2026

SUKUNA-AI mentioned this pull request Oct 21, 2026

Add CUDA Forward Compatibility package to our CI runners to allow testing new CUDA features #11935

Open

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Oct 21, 2026

SUKUNA-AI force-pushed the fix/shared-memory-barrier-no-report-decode branch from a2fd7ec to c9b0e2a Compare October 21, 2026 18:31

pciolkosz added the backport branch/3.6.x label Oct 21, 2026

pciolkosz approved these changes Oct 21, 2026

View reviewed changes

This comment has been minimized.

Sign in to view

This was referenced Oct 22, 2026

[BUG]: cub.test.device.radix_sort_runs_on.lid_0 aborts in single-architecture builds #12154

Closed

[CUB] Only use compiled compute capabilities in the radix_sort_runs_on tests #12155

Closed

Merge branch 'main' into fix/shared-memory-barrier-no-report-decode

d95ea5e

This comment has been minimized.

Sign in to view

Merge branch 'main' into fix/shared-memory-barrier-no-report-decode

8936c00

pciolkosz merged commit 4811a1c into NVIDIA:main Oct 23, 2026
256 checks passed

github-project-automation Bot moved this from In Review to Done in CCCL Oct 23, 2026

github-actions Bot mentioned this pull request Oct 23, 2026

[Backport branch/3.6.x] [libcu++] Skip fabric status decoding when a shared_memory_barrier wait has no report #12228

Merged

coderabbitai Bot mentioned this pull request Oct 26, 2026

[libcu++] Shared memory barrier wait with a time hint #12264

Open

Conversation

SUKUNA-AI commented Oct 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Checklist

Uh oh!

copy-pr-bot Bot commented Oct 21, 2026

Uh oh!

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

Uh oh!

Summary by CodeRabbit

Uh oh!

pciolkosz commented Oct 21, 2026

Uh oh!

copy-pr-bot Bot commented Oct 21, 2026

Uh oh!

pciolkosz left a comment

Choose a reason for hiding this comment

Uh oh!

pciolkosz commented Oct 21, 2026

Uh oh!

This comment has been minimized.

SUKUNA-AI commented Oct 22, 2026

Uh oh!

pciolkosz commented Oct 22, 2026

Uh oh!

This comment has been minimized.

SUKUNA-AI commented Oct 23, 2026

Uh oh!

pciolkosz commented Oct 23, 2026

Uh oh!

github-actions Bot commented Oct 23, 2026

🥳 CI Workflow Results

🟩 Finished in 1h 52m: Pass: 100%/228 | Total: 3d 05h | Max: 1h 51m | Hits: 99%/629481

Uh oh!

Uh oh!

github-actions Bot commented Oct 23, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

SUKUNA-AI commented Oct 21, 2026 •

edited

Loading

coderabbitai Bot commented Oct 21, 2026 •

edited

Loading