Conversation
Contributor
|
Navigate logical layers of code changes, visualize relationships, and explore their blast radius. No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Repository: NVIDIA/cccl/.coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (13)
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review. 📝 SummarySummary by CodeRabbit
WalkthroughThe changes update mdspan copy extent validation, type-aware copy dispatch, shared-memory staging constraints, and contiguous-copy handling for outer extents beyond the device grid-y limit. Tests cover mismatched empty shapes, single-element conversions, and large copies. ChangesMdspan Copy Behavior
Priority: ➖ Normal Change: Bug fix Merge Risk: ⚪ Minimal · up to The copy changes preserve type conversion and provide valid fallbacks when shared-memory staging is unsupported. No actionable merge-blocking issue was identified; merge after normal build and CUDA test checks. Comment ✨ Finishing Touches 💡 1🛠️ Fix failing CI checks 💡
|
Jacobfaib
reviewed
Sep 29, 2026
| { | ||
| typename _ExtentsIn::rank_type __i = 0; | ||
| typename _ExtentsOut::rank_type __j = 0; | ||
| while (true) |
Contributor
There was a problem hiding this comment.
This could be a for-loop since you increment ++__i and ++__j.
| typename _ExtentsOut::rank_type __j = 0; | ||
| while (true) | ||
| { | ||
| while (__i != _ExtentsIn::rank() && __extents_in.extent(__i) == 1) |
Contributor
There was a problem hiding this comment.
Suggested change
| while (__i != _ExtentsIn::rank() && __extents_in.extent(__i) == 1) | |
| while (__i < _ExtentsIn::rank() && __extents_in.extent(__i) == 1) |
< is safer. I could implement an extents impls that returns negative ranks.
Comment on lines
+59
to
+60
| copy_stream.sync(); | ||
| REQUIRE(d_dst[0] == 42.5); |
Contributor
There was a problem hiding this comment.
Question: does this run after each section, or after both sections? My intuition is that this runs after both sections because that's how the setup code works as well.
| thrust::raw_pointer_cast(d_src.data()), extents_t(0, 3)); | ||
| const cuda::device_mdspan<float, extents_t, layout_right> dst(thrust::raw_pointer_cast(d_dst.data()), extents_t(0, 2)); | ||
|
|
||
| CHECK_THROWS_AS(cuda::copy(src, dst, copy_stream), std::invalid_argument); |
Contributor
There was a problem hiding this comment.
REQUIRES_THROWS_MATCHES() to also match the exception message please.
| constexpr int M = 65537; | ||
| constexpr int N = 128 * 1024; | ||
| constexpr int Ld = N + 128; | ||
| const auto required = size_t{M} * (Ld + N); |
Contributor
There was a problem hiding this comment.
Suggested change
| const auto required = size_t{M} * (Ld + N); | |
| constexpr auto required = size_t{M} * (Ld + N); |
| thrust::device_vector<char> d_src(size_t{M} * Ld, static_cast<char>(0x42)); | ||
| thrust::device_vector<char> d_dst(size_t{M} * N, static_cast<char>(0x00)); | ||
|
|
||
| using cuda::std::layout_stride; |
Contributor
There was a problem hiding this comment.
Do you really need the using decl for just one use :)?
| // references. | ||
| template <typename _TpIn, typename _SrcAccessor, typename _DstAccessor> | ||
| inline constexpr bool __can_stage_in_shared_mem_v = | ||
| ::cuda::is_trivially_copyable_v<::cuda::std::remove_cv_t<_TpIn>> |
Contributor
There was a problem hiding this comment.
remove_cvref_t instead to be safe here?
| template <typename _TpIn, typename _SrcAccessor, typename _DstAccessor> | ||
| inline constexpr bool __can_stage_in_shared_mem_v = | ||
| ::cuda::is_trivially_copyable_v<::cuda::std::remove_cv_t<_TpIn>> | ||
| && ::cuda::std::is_assignable_v<::cuda::std::remove_cv_t<_TpIn>&, typename _SrcAccessor::reference> |
Contributor
There was a problem hiding this comment.
You are looking for the cuda::std::assignable_from concept here.
| "grid y-dimension exceeds the maximum grid size"); | ||
| const auto __grid_dims = ::dim3(static_cast<unsigned>(__num_inner_tiles), static_cast<unsigned>(__outer_size)); | ||
| const auto __config = ::cuda::make_config(::cuda::block_dims<__block_size>(), ::cuda::grid_dims(__grid_dims)); | ||
| const auto __grid_dim_y = ::cuda::std::min(__outer_size, _ExtentT(__arch_limits.max_grid_dim_y)); |
Contributor
There was a problem hiding this comment.
_ExtentT{} to forbid narrowing?
Contributor
😬 CI Workflow Results🟥 Finished in 1h 41m: Pass: 91%/265 | Total: 3d 15h | Max: 1h 38m | Hits: 89%/680881See results here. AI failure analysis1. MDSpan rank-zero comparison triggers MSVC unreachable-code errors · 19 jobsExplanation: The newly added `__same_non_singleton_extents` helper is instantiated with rank-zero extents, allowing MSVC to identify its loop bodies and subsequent comparison path as unreachable. Every MSVC build uses `/WX`, so C4702 becomes fatal C2220 in both libcudacxx and cudax targets across the tested CUDA, MSVC, and C++ versions. Evidence: Copy this prompt into a coding agentJobs:
2. Captured-bool transform test is a strict XPASS with numba-cuda-mlir 0.5.4 · 3 jobsExplanation: All three jobs installed `numba-cuda-mlir` 0.5.4, where the test now passes; its unconditional `strict=True` xfail therefore converts the corrected behavior into a failure. The same signature occurs on Linux, Windows, and the v2/HostJIT backend. Evidence: Copy this prompt into a coding agentJobs: 3. Third-party NVCC builds fail while sccache packages missing compiler outputs · 2 jobsExplanation: Both optional third-party builds use `sccache` 0.17.0-rapids.4 and fail after compilation because an NVCC-generated output disappears before sccache can archive it. The affected MatX and cuDF sources differ, but the cache packaging mechanism and decisive error are equivalent. Evidence: Copy this prompt into a coding agentJobs: 4. libcudacxx buffer LLDB pretty-printer exceeds all timeout retries · 1 jobExplanation: The LLDB session successfully loaded the formatter, printed several buffer cases, and then stopped making progress after continuing from `inspect_before_update`; all three attempts exceeded the per-attempt timeout. Neighboring LLDB tests passed, and the available log does not establish whether this was runner load, CUDA runtime delay, or a deterministic hang. Evidence: Copy this prompt into a coding agentJobs: |
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
Fixes cccl_qa_agent_findings/issues/253
Actual issues:
noexcept