Move `cudax::copy(mdspan)` to libcu++ by fbusato · Pull Request #11321 · NVIDIA/cccl

fbusato · 2026-09-09T22:48:49Z

Description

Move the implementations, tests, benchmarks from cudax::copy(mdspan) to libcu++ following Copy and fill for mdspan P3242R4.

The PR mostly moves the files from cuda/experimental to libcudacxx/cuda/. There are a few new files under libcudacxx/include/cuda/std.

The PR also introduces the PSTL entry point in <cuda/std/mdspan>.

template<class ExecutionPolicy, class Src, class Dst>
void copy(ExecutionPolicy&& policy, const Src& src, const Dst& dst);

while the asynchronous version is in <cuda/mdspan>

template <class TpIn,
          class ExtentsIn,
          class LayoutPolicyIn,
          class AccessorPolicyIn,
          class TpOut,
          class ExtentsOut,
          class LayoutPolicyOut,
          class AccessorPolicyOut>
void copy(
        cuda::device_mdspan<TpIn,  ExtentsIn,  LayoutPolicyIn,  AccessorPolicyIn>  src,
        cuda::device_mdspan<TpOut, ExtentsOut, LayoutPolicyOut, AccessorPolicyOut> dst,
        cuda::stream_ref                                                           stream);

Directly call from cub::DeviceCopy will be in a follow-up PR

coderabbitai · 2026-09-09T22:59:23Z

📝 Summary

Summary by CodeRabbit

New Features
- Added asynchronous cuda::copy support for copying between device mdspan objects on a specified CUDA stream.
- Added execution-policy support for cuda::std::copy with device mdspan objects.
- Added public documentation, examples, and feature detection for mdspan copying.
API Changes
- Device mdspan copying now uses cuda::copy instead of the experimental namespace.
- Copying supports destination references assignable from source references.
Removals
- Removed the experimental copy header and related benchmark and test registrations.

Walkthrough

Changes

Mdspan copy migration

Layer / File(s)	Summary
Primary mdspan copy implementation `libcudacxx/include/cuda/__mdspan/__copy/*`	Moved mdspan copy helpers into the primary `cuda` namespace and added `cuda::__raw_tensor`. The device copy dispatch preserves existing copy paths and changes the element constraint to destination-reference assignability.
Experimental helper bridge and removal `cudax/include/cuda/experimental/__copy_bytes/`, `cudax/include/cuda/experimental/__fill_bytes/`, `cudax/include/cuda/experimental/copy.cuh`	Updated experimental helpers to use the new mdspan copy types and namespace utilities. Removed the obsolete experimental header and CMake registrations.
CUDA PSTL mdspan copy `libcudacxx/include/cuda/std/__pstl/*`, `libcudacxx/include/cuda/mdspan`, `libcudacxx/include/cuda/std/mdspan`, `libcudacxx/include/cuda/std/version`	Added execution-policy mdspan copying, CUDA backend dispatch, stream synchronization, aggregate includes, overflow utility use, and a feature-test macro.
Tests and API documentation `libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/`, `docs/libcudacxx/extended_api/mdspan`	Updated tests to call `cuda::copy`, added CUDA PSTL stream and layout coverage, removed shared-memory tiling tests, and documented asynchronous device-mdspan copying.
Build and benchmark wiring `cudax/benchmarks/bench/*`, `cudax/test/CMakeLists.txt`, `libcudacxx/benchmarks/bench/copy/copy_bench.cu`, `thrust/thrust/system/cuda/detail/core/util.h`	Removed obsolete benchmark and test wiring. Updated the libcudacxx benchmark and simplified the local plan type alias.

Priority: ➖ Normal

Merge Risk: 🔵 Low · up to 0ae8e

Device-only consumers can detect mdspan copy support and then fail to find the advertised execution-policy overload. Gate the feature macro with the hosted availability condition before merge.

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

coderabbitai

Actionable comments posted: 1

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)

libcudacxx/include/cuda/__mdspan/__copy/mdspan_d2d.h (1)
146-146: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

important: Match the mdspan copy constraint to assignability. The PSTL overload requires assignability between the source and destination device_mdspan::reference types, then calls cuda::copy, whose element-type convertibility assertion can reject that valid pair. Use the reference types in the direct API and add compile-only coverage.
-  static_assert(::cuda::std::is_convertible_v<_TpIn, _TpOut>, "TpIn must be convertible to TpOut");
+  using __src_reference = typename decltype(__src)::reference;
+  using __dst_reference = typename decltype(__dst)::reference;
+  static_assert(::cuda::std::is_assignable_v<__dst_reference, __src_reference>,
+                "destination mdspan reference must be assignable from source mdspan reference");

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: ee25ae36-9e66-4c43-9e55-d5098a15e73f

📥 Commits

Reviewing files that changed from the base of the PR and between 315453e and 73ebf3e.

📒 Files selected for processing (44)

cudax/benchmarks/bench/CMakeLists.txt
cudax/benchmarks/bench/copy/CMakeLists.txt
cudax/include/cuda/experimental/__copy_bytes/mdspan_d2h_h2d.cuh
cudax/include/cuda/experimental/__copy_bytes/memcpy_batch_tiles.cuh
cudax/include/cuda/experimental/__copy_bytes/print_raw_tensor.cuh
cudax/include/cuda/experimental/__copy_bytes/types.cuh
cudax/include/cuda/experimental/__fill_bytes/fill_bytes_mdspan.cuh
cudax/include/cuda/experimental/__fill_bytes/fill_bytes_mdspan_utils.cuh
cudax/include/cuda/experimental/copy.cuh
cudax/test/CMakeLists.txt
docs/libcudacxx/extended_api/mdspan.rst
docs/libcudacxx/extended_api/mdspan/copy.rst
libcudacxx/benchmarks/bench/copy/copy_bench.cu
libcudacxx/include/cuda/__mdspan/__copy/abs_integer.h
libcudacxx/include/cuda/__mdspan/__copy/copy_contiguous.h
libcudacxx/include/cuda/__mdspan/__copy/copy_dst_contiguous.h
libcudacxx/include/cuda/__mdspan/__copy/copy_optimized.h
libcudacxx/include/cuda/__mdspan/__copy/copy_shared_memory.h
libcudacxx/include/cuda/__mdspan/__copy/copy_shared_memory_utils.h
libcudacxx/include/cuda/__mdspan/__copy/dispatch_by_vector.h
libcudacxx/include/cuda/__mdspan/__copy/mdspan_d2d.h
libcudacxx/include/cuda/__mdspan/__copy/mdspan_to_raw_tensor.h
libcudacxx/include/cuda/__mdspan/__copy/simplify_paired.h
libcudacxx/include/cuda/__mdspan/__copy/tensor_copy_utils.h
libcudacxx/include/cuda/__mdspan/__copy/tensor_iterator.h
libcudacxx/include/cuda/__mdspan/__copy/tensor_query.h
libcudacxx/include/cuda/__mdspan/__copy/types.h
libcudacxx/include/cuda/__mdspan/__copy/vector_access.h
libcudacxx/include/cuda/mdspan
libcudacxx/include/cuda/std/__pstl/cuda/mdspan/copy.h
libcudacxx/include/cuda/std/__pstl/dispatch.h
libcudacxx/include/cuda/std/__pstl/mdspan/copy.h
libcudacxx/include/cuda/std/mdspan
libcudacxx/include/cuda/std/version
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy.cu
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_common.cuh
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_edge_cases.cu
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_llm.cu
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_nvmath.cu
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_nvmath_transpose.cu
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_pstl.cu
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_shared_memory.cu
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_vectorize.cu
libcudacxx/test/libcudacxx/cuda/containers/views/mdspan/copy/copy_vectorize_5d.cu

💤 Files with no reviewable changes (5)

cudax/benchmarks/bench/CMakeLists.txt
cudax/include/cuda/experimental/copy.cuh
cudax/benchmarks/bench/copy/CMakeLists.txt
cudax/test/CMakeLists.txt
cudax/include/cuda/experimental/__copy_bytes/types.cuh

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

miscco · 2026-09-10T06:39:52Z

+  ::cuda::std::array<_ExtentT, _MaxRank> __extents;
+  ::cuda::std::array<_StrideT, _MaxRank> __strides;


Question: Could those also be

Suggested change

::cuda::std::array<_ExtentT, _MaxRank> __extents;

::cuda::std::array<_StrideT, _MaxRank> __strides;

_ExtentT __extents[_MaxRank];

_StrideT __strides[_MaxRank];

unfortunately no, we need array for passing by-value

davebayer

Do we really need 2 versions of the same function? For parallel algorithms, I believe we also allow the asynchronous version, when executed with gpu execution policy with a stream attached even though it violates the standard. CC @miscco

miscco · 2026-09-10T10:57:04Z

We currently do not support asynchronous PSTL algorithms. Even if we provide a stream we always synchronize at the end

bernhardmgruber

Thrust change is fine

coderabbitai

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)

libcudacxx/include/cuda/std/version (1)

89-89: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

important: Gate __cccl_lib_mdspan_copy with _CCCL_HOSTED().

<cuda/std/version> defines the macro in device compilation, but <cuda/std/__pstl/mdspan/copy.h> declares the execution-policy cuda::std::copy overload only inside _CCCL_HOSTED(). A device-only consumer can use the macro to select that overload and then fail during overload lookup.

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 475f5d91-29d4-454c-9342-ddb1464666d9

📥 Commits

Reviewing files that changed from the base of the PR and between e13c50e and 0ae8e0f.

📒 Files selected for processing (4)

libcudacxx/include/cuda/__mdspan/__copy/tensor_query.h
libcudacxx/include/cuda/std/__algorithm/inplace_merge.h
libcudacxx/include/cuda/std/__algorithm/stable_sort.h
libcudacxx/include/cuda/std/__pstl/mdspan/copy.h

🚧 Files skipped from review as they are similar to previous changes (1)

libcudacxx/include/cuda/__mdspan/__copy/tensor_query.h

Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review.

github-actions · 2026-09-11T01:56:43Z

🥳 CI Workflow Results

🟩 Finished in 3h 19m: Pass: 100%/290 | Total: 10d 18h | Max: 3h 18m | Hits: 44%/1591894

See results here.

move cudax::copy(mdspan) to libcudacxx

73ebf3e

fbusato self-assigned this Sep 9, 2026

fbusato requested review from a team as code owners September 9, 2026 22:48

fbusato added this to CCCL Sep 9, 2026

fbusato requested review from a team as code owners September 9, 2026 22:48

fbusato requested review from andralex, shwina and wmaxey September 9, 2026 22:48

github-project-automation Bot moved this to Todo in CCCL Sep 9, 2026

fbusato moved this from Todo to In Progress in CCCL Sep 9, 2026

cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Sep 9, 2026

coderabbitai Bot reviewed Sep 9, 2026

View reviewed changes

Comment thread libcudacxx/include/cuda/std/version

This comment has been minimized.

Sign in to view

miscco approved these changes Sep 10, 2026

View reviewed changes

davebayer reviewed Sep 10, 2026

View reviewed changes

Jacobfaib reviewed Sep 10, 2026

View reviewed changes

Comment thread libcudacxx/include/cuda/std/__pstl/cuda/mdspan/copy.h Outdated

Comment thread libcudacxx/include/cuda/__mdspan/__copy/types.h Outdated

fbusato added 2 commits September 10, 2026 13:13

add _CCCL_NODEBUG to aliases

bf84edf

fix gcc/clang compatibility

138170b

Jacobfaib approved these changes Sep 10, 2026

View reviewed changes

fbusato added 5 commits September 10, 2026 13:20

try to fix clang-tidy issue

aec077b

fix GCC7 warning

3526b41

move copy constraints to a concept

91084b6

fix stream name in tests

213f2b2

remove redundant alias

e13c50e

fbusato requested a review from a team as a code owner September 10, 2026 20:56

fbusato requested a review from elstehle September 10, 2026 20:56

bernhardmgruber approved these changes Sep 10, 2026

View reviewed changes

fbusato and others added 3 commits September 10, 2026 15:23

fix concept

8ba3b76

fix clang-tidy warnings

fd815c1

Merge branch 'main' into parallel-mdpan-copy

0ae8e0f

coderabbitai Bot reviewed Sep 10, 2026

View reviewed changes

Navigation Menu

Uh oh!

Move `cudax::copy(mdspan)` to libcu++ - #11321

Move `cudax::copy(mdspan)` to libcu++#11321
fbusato wants to merge 11 commits into
NVIDIA:mainfrom
fbusato:parallel-mdpan-copy

fbusato commented Sep 9, 2026

Uh oh!

coderabbitai Bot commented Sep 9, 2026 •

edited

Loading

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Uh oh!

Uh oh!

This comment has been minimized.

miscco Sep 10, 2026

Uh oh!

fbusato Sep 10, 2026

Uh oh!

Uh oh!

davebayer left a comment

Uh oh!

miscco commented Sep 10, 2026

Uh oh!

Uh oh!

Uh oh!

bernhardmgruber left a comment

Uh oh!

coderabbitai Bot left a comment

Uh oh!

github-actions Bot commented Sep 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

		::cuda::std::array<_ExtentT, _MaxRank> __extents;
		::cuda::std::array<_StrideT, _MaxRank> __strides;

Uh oh!

Conversation

fbusato commented Sep 9, 2026

Description

Uh oh!

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

Uh oh!

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

This comment has been minimized.

miscco Sep 10, 2026

Choose a reason for hiding this comment

Uh oh!

fbusato Sep 10, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

davebayer left a comment

Choose a reason for hiding this comment

Uh oh!

miscco commented Sep 10, 2026

Uh oh!

Uh oh!

Uh oh!

bernhardmgruber left a comment

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Sep 11, 2026

🥳 CI Workflow Results

🟩 Finished in 3h 19m: Pass: 100%/290 | Total: 10d 18h | Max: 3h 18m | Hits: 44%/1591894

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

coderabbitai Bot commented Sep 9, 2026 •

edited

Loading