Skip to content

Switch the disabled shared_memory_barrier tests to c2h - #12150

Merged
pciolkosz merged 6 commits into
NVIDIA:mainfrom
pciolkosz:shared_memory_barrier_c2h_tests
Oct 24, 2026
Merged

pciolkosz merged 6 commits into
NVIDIA:mainfrom
pciolkosz:shared_memory_barrier_c2h_tests

Conversation

@pciolkosz

Copy link
Copy Markdown
Contributor

shared_memory_barrier lit test is currently disabled because of eager loading in lit and older drivers on some CI machines leading to failure before the main even runs. This PR switches the test to c2h so we can check the driver version and skip the test

@pciolkosz
pciolkosz requested a review from a team as a code owner October 21, 2026 23:17
@pciolkosz
pciolkosz requested a review from wmaxey October 21, 2026 23:17
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Oct 21, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Oct 21, 2026
@coderabbitai

coderabbitai Bot commented Oct 21, 2026 •

Copy link
Copy Markdown
Contributor

Review in Change Stack →

Important

Review skipped

Review was skipped as selected files did not have any reviewable changes.

⚙️ Run configuration
  • Source mode: Auto
  • Repository: NVIDIA/cccl
  • Branch: main
  • File: .coderabbit.yaml
  • Commit: 053d3c7
  • Layers: Repository YAML
  • Review profile: CHILL
  • Plan: Enterprise
  • Run ID: df4474c0-595e-4277-9d5c-3af3ee5804bc
📥 Commits

Reviewing files that changed from the base of the PR and between 4c441a1 and 168b489.

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

  • 🔍 Trigger review

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: d37d78b
  • Layers: Repository YAML
  • Review profile: CHILL
  • Plan: Enterprise
  • Run ID: 6c659651-dbaa-48c9-8979-4b696473d0f7
📥 Commits

Reviewing files that changed from the base of the PR and between 8b9de1e and 3a86210.

📒 Files selected for processing: 1
  • libcudacxx/test/libcudacxx/cuda/barrier/shared_memory_barrier.cu

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

  • Tests
    • Updated shared-memory barrier tests for CUDA 13.4 and newer.
    • Added separate coverage for block- and cluster-scope execution using a two-thread block.
    • Improved validation of barrier status and wait operations.
    • Tests skip execution when the driver is older than CUDA 13.4 or the device does not meet the required compute capability.

The shared-memory barrier test now uses C2H to launch block- and cluster-scope kernels. Compilation is limited to CUDA 13.4 or newer outside NVRTC. The test checks driver version and compute capability before launching, and uses REQUIRE checks in place of assertions.

Changes: Shared-memory barrier tests

Layer / File(s) Summary
Device test checks and CUDA compilation
libcudacxx/test/libcudacxx/cuda/barrier/shared_memory_barrier.cu
The test adds a CUDA 13.4+ non-NVRTC compilation guard and updated includes. Assertions become REQUIRE checks, and concurrent lambdas omit explicit __device__ annotations.
Kernel and C2H test harness
libcudacxx/test/libcudacxx/cuda/barrier/shared_memory_barrier.cu
A templated kernel runs the barrier tests. The C2H test returns without launching when the driver is below 13.4 or the device is below compute capability 9. Otherwise, it launches block- and cluster-scope kernels.

Possibly related issues

Priority: ⬇️ Low

Change: Other

Merge Risk: 🟡 Moderate · up to 3a862

The test can report success without exercising the shared-memory barrier. Ensure the build supplies an SM90-or-newer kernel image before merging.

✨ Finishing Touches: 💡 1
🛠️ Fix failing CI checks: 💡
  • Commit to this branch
  • Create a new PR
  • Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

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

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1


ℹ️ Review info
⚙️ Run configuration
  • Configuration used: Repository: NVIDIA/cccl/.coderabbit.yaml
  • Review profile: CHILL
  • Plan: Enterprise
  • Run ID: 8d433a83-ef87-44a3-bc32-a8ebdf9b7036
📥 Commits

Reviewing files that changed from the base of the PR and between 035f1a6 and 8b9de1e.

📒 Files selected for processing: 1
  • libcudacxx/test/libcudacxx/cuda/barrier/shared_memory_barrier.cu

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

Comment on lines +543 to +545
NV_IF_TARGET(NV_PROVIDES_SM_90,
(test_shared_memory_barrier_choreography<Scope>(); test_shared_memory_barrier_common_extensions<Scope>();
test_shared_memory_barrier_sm90_extensions<Scope>();))

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy lift

🔎 Supported by static analysis

🏁 Script executed:

#!/bin/bash
# Inspect architecture settings for this test; also check its generated
# compilation command for an SM90-or-newer image.
rg -n -C 3 'CMAKE_CUDA_ARCHITECTURES|CUDA_ARCHITECTURES|gencode|compute_90|sm_90' libcudacxx/test

Repository: NVIDIA/cccl

Length of output: 23631


🏁 Script executed:

set -u
printf '%s\n' '--- test references and registration ---'
rg -n -F --glob 'CMakeLists.txt' -e 'shared_memory_barrier' -e 'LIBCUDACXX_TEST' libcudacxx/test libcudacxx || test "$?" -eq 1
printf '%s\n' '--- libcudacxx test CMake context ---'
sed -n '1,180p' libcudacxx/test/libcudacxx/CMakeLists.txt
printf '%s\n' '--- relevant test source entrypoints and guards ---'
nl -ba libcudacxx/test/libcudacxx/cuda/barrier/shared_memory_barrier.cu | sed -n '1,90p;520,600p'
printf '%s\n' '--- architecture propagation and defaults ---'
rg -n -C 5 -e 'CMAKE_CUDA_ARCHITECTURES' -e 'LIBCUDACXX.*ARCH' -e 'compute_archs' libcudacxx/test cmake CMakeLists.txt || test "$?" -eq 1

Repository: NVIDIA/cccl

Length of output: 34173


🏁 Script executed:

set -u
printf '%s\n' '--- cccl_add_executable definition and callers ---'
rg -n -C 8 -F -- 'function(cccl_add_executable' . || test "$?" -eq 1
rg -n -C 4 -F -- 'cccl_add_executable(' libcudacxx/test cmake CMakeLists.txt | head -240
printf '%s\n' '--- CUDA architecture defaults and target properties ---'
rg -n -C 6 -e 'CMAKE_CUDA_ARCHITECTURES' -e 'CUDA_ARCHITECTURES' cmake CMakeLists.txt libcudacxx/CMakeLists.txt libcudacxx/test/CMakeLists.txt | head -320

Repository: NVIDIA/cccl

Length of output: 8895


🏁 Script executed:

set -u
printf '%s\n' '--- complete cccl_add_executable implementation ---'
nl -ba cmake/CCCLAddExecutable.cmake | sed -n '1,180p'
printf '%s\n' '--- supported architecture invocations and test build configuration ---'
rg -n -C 5 -e 'all-major-cccl' -e 'all-cccl' -e 'CMAKE_CUDA_ARCHITECTURES=' -e 'CUDA_ARCHITECTURES' .github ci cmake docs libcudacxx 2>/dev/null | head -400

Repository: NVIDIA/cccl

Length of output: 31258


🏁 Script executed:

set -u
rg -n -F -- 'function(cccl_configure_target' cmake
file=$(rg -l -F -- 'function(cccl_configure_target' cmake | head -1)
if [ -n "$file" ]; then
  nl -ba "$file" | sed -n '1,240p'
fi

Repository: NVIDIA/cccl

Length of output: 2491


important: Ensure the shared-memory barrier target contains SM90 device code.

test_shared_memory_barrier_kernel removes all test calls when compiled below SM90. The C2H test checks only the host device capability. A lower-architecture PTX image can therefore run on an SM90 device while executing an empty kernel, and both sections can pass without testing the barrier.

Require an SM90-or-newer compilation target for this test, or skip configurations that cannot provide one. Do not rely on the host device check to select the device code.

Source: Path instructions

@github-actions

This comment has been minimized.

assert(status.complete());
assert(!status.has_report());
REQUIRE_DEVICE(status.complete());
REQUIRE_DEVICE(!status.has_report());

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just REQUIRE().

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 1h 10m: Pass: 100%/147 | Total: 1d 13h | Max: 1h 10m | Hits: 99%/557909

See results here.

@pciolkosz
pciolkosz merged commit 3d72951 into NVIDIA:main Oct 24, 2026
175 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Oct 24, 2026
@github-actions

Copy link
Copy Markdown
Contributor

Successfully created backport PR for branch/3.6.x:

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

2 participants