Use CCCL's mdspan implementation #2836

bdice · 2025-10-14T21:45:37Z

This PR adopts CCCL's mdspan implementation and deletes the vendored implementation.

Closes #1616.

Note

This branch was meant to be a stress test of Cursor with GPT-5, and I did not anticipate it becoming a PR worth opening. To my surprise, it got to a state that actually compiles (with a lot of retries). After skimming over the changes, I think it's fairly close to what we might actually want.

copy-pr-bot · 2025-10-14T21:45:41Z

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

bdice · 2025-10-14T21:47:22Z

cpp/include/raft/core/mdspan.hpp


+// Padded layout shims compatible with CCCL mdspan mappings
+template <size_t PaddingStride>
+struct layout_left_padded_impl {


Marking this for closer review.

This feature has been requested in CCCL, in NVIDIA/cccl#2808.

Should this be detail so that it's not in a public namespace?

Yes, that's a good idea.

bdice · 2025-10-14T21:57:37Z

cpp/tests/sparse/preprocess.cu

-                   raft::device_vector_view<T1> rows,
-                   raft::device_vector_view<T1> columns,
-                   raft::device_vector_view<T2> values,
+                   raft::device_vector_view<T1, int64_t> rows,


Need to check if this change is correct.

It is explicit but unnecessary.

bdice · 2025-10-14T21:57:50Z

cpp/tests/util/preprocess_utils.cu

 };

-template <typename T1, typename T2>
+template <typename T1, typename T2, typename DenseVectorView, typename ResultVectorView>


Need to check if this change is correct.

bdice · 2025-10-14T21:59:14Z

python/pylibraft/pylibraft/common/cpp/mdspan.pxd

        namespace "raft":
-    ctypedef layout_right row_major
-    ctypedef layout_left col_major
+    cdef cppclass row_major:


I think this is equivalent, Cython probably compiles the old and new code the same way.

bdice · 2025-10-14T21:59:39Z

/ok to test

mhoemmen · 2025-10-14T22:06:38Z

cpp/include/raft/core/mdspan.hpp

                "std::experimental::padding ValueType has to be provided without "
                "const or volatile specifiers.");
  static_assert(ByteAlignment % sizeof(ValueType) == 0 || sizeof(ValueType) % ByteAlignment == 0,
                "std::experimental::padding sizeof(ValueType) has to be multiple or "


Not quite sure where this std::experimental::padding comes from, though it was preexisting. There's another instance of this above.

mhoemmen · 2025-10-14T22:09:06Z

cpp/include/raft/core/mdspan_types.hpp

+using cuda::std::dynamic_extent;
+using cuda::std::extents;


If we're using these using declarations anyway, then should we consider just using them consistently for all cuda::std namespace symbols we want to use in namespace raft, so that we don't have to type cuda::std (and so that we could change the namespace in one place, instead of touching all the files)?

I'd prefer not to have these using declarations. I'd like to keep all the types in cuda::std namespace. But we can deal with this in a follow-up issue to keep the diff as minimal as possible. Please do not bring in any more types in raft:: namespace than what are already there.

Actually, I'm re-reading Mark's comment and I think what he is saying is a good idea. We using raft:: namespace outside of this header and bring cuda::std:: symbols in raft:: namespace in this header.

The thing that's happening now is that there are using declarations in the raft namespace, AND code in the raft namespace uses explicit namespace qualification (e.g., ::cuda::std::extents). I don't think it makes sense to do both. On the other hand, fixing that is perhaps out of scope of this PR.

Agreed. Code outside this header should use types from the raft namespace.

mhoemmen · 2025-10-14T22:10:36Z

cpp/include/raft/core/mdspan_types.hpp

+using extent_5d = cuda::std::extents<IndexType,
+                                     dynamic_extent,
+                                     dynamic_extent,
+                                     dynamic_extent,
+                                     dynamic_extent,
+                                     dynamic_extent>;


The more concise way to spell this is

Suggested change

using extent_5d = cuda::std::extents<IndexType,

dynamic_extent,

dynamic_extent,

dynamic_extent,

dynamic_extent,

dynamic_extent>;

using extent_5d = cuda::std::dims<5, IndexType>;

but dims<5, IndexType> is more concise than extent_5d<IndexType> anyway; should we consider just using dims?

mhoemmen · 2025-10-14T22:11:55Z

cpp/tests/core/mdarray.cu

  thrust::sequence(rmm::exec_policy(stream), a.begin(), a.end());
-  stdex::mdspan<float, stdex::extents<int, raft::dynamic_extent, raft::dynamic_extent>> span{
-    a.data(), 4, 4};
+  cuda::std::mdspan<float, cuda::std::extents<int, raft::dynamic_extent, raft::dynamic_extent>>


Suggested change

cuda::std::mdspan<float, cuda::std::extents<int, raft::dynamic_extent, raft::dynamic_extent>>

cuda::std::mdspan<float, cuda::std::dims<2, int>>

The preferred option, though, would be to use CTAD.

cuda::std::mdspan{a.data(), cuda::std::dims<2, int>{4, 4}}

mhoemmen · 2025-10-14T22:13:39Z

cpp/tests/core/mdarray.cu

 void test_mdarray_basic()
 {
-  using matrix_extent = stdex::extents<int, dynamic_extent, dynamic_extent>;
+  using matrix_extent = cuda::std::extents<int, dynamic_extent, dynamic_extent>;


Suggested change

using matrix_extent = cuda::std::extents<int, dynamic_extent, dynamic_extent>;

using matrix_extent = cuda::std::dims<2, int>;

mhoemmen · 2025-10-14T22:14:05Z

cpp/tests/core/mdarray.cu

 {
  raft::resources handle;
-  using matrix_extent = stdex::extents<size_t, dynamic_extent, dynamic_extent>;
+  using matrix_extent = cuda::std::extents<size_t, dynamic_extent, dynamic_extent>;


Suggested change

using matrix_extent = cuda::std::extents<size_t, dynamic_extent, dynamic_extent>;

using matrix_extent = cuda::std::dims<2>;

(Default index type is size_t.)

mhoemmen · 2025-10-14T22:14:16Z

cpp/tests/core/mdarray.cu

 TEST(MDArray, CopyMove)
 {
-  using matrix_extent = stdex::extents<size_t, dynamic_extent, dynamic_extent>;
+  using matrix_extent = cuda::std::extents<size_t, dynamic_extent, dynamic_extent>;


Suggested change

using matrix_extent = cuda::std::extents<size_t, dynamic_extent, dynamic_extent>;

using matrix_extent = cuda::std::dims<2>;

mhoemmen · 2025-10-14T22:14:46Z

cpp/tests/core/mdarray.cu

 void test_mdarray_padding()
 {
-  using extents_type = stdex::extents<size_t, dynamic_extent, dynamic_extent>;
+  using extents_type = cuda::std::extents<size_t, dynamic_extent, dynamic_extent>;


Suggested change

using extents_type = cuda::std::extents<size_t, dynamic_extent, dynamic_extent>;

using extents_type = cuda::std::dims<2>;

mhoemmen · 2025-10-14T22:15:00Z

cpp/tests/core/mdarray.cu

  // use simple 1D array to allocate some space
  auto s          = rmm::cuda_stream_default;
-  using extent_1d = stdex::extents<size_t, dynamic_extent>;
+  using extent_1d = cuda::std::extents<size_t, dynamic_extent>;


Suggested change

using extent_1d = cuda::std::extents<size_t, dynamic_extent>;

using extent_1d = cuda::std::dims<1>;

divyegala · 2025-10-15T00:15:30Z

cpp/include/raft/core/mdspan.hpp


+// Padded layout shims compatible with CCCL mdspan mappings
+template <size_t PaddingStride>
+struct layout_left_padded_impl {


Yes, that's a good idea.

divyegala · 2025-10-15T00:17:01Z

cpp/include/raft/core/mdspan_types.hpp

+using cuda::std::dynamic_extent;
+using cuda::std::extents;


I'd prefer not to have these using declarations. I'd like to keep all the types in cuda::std namespace. But we can deal with this in a follow-up issue to keep the diff as minimal as possible. Please do not bring in any more types in raft:: namespace than what are already there.

divyegala · 2025-10-15T00:21:23Z

cpp/tests/core/mdarray.cu

     */
    using padded_layout_row_major =
-      stdex::layout_right_padded<detail::padding<float, alignment_bytes>::value>;
+      raft::layout_right_padded_impl<detail::padding<float, alignment_bytes>::value>;


Should this use raft::layout_right_padded<float>?

divyegala · 2025-10-15T00:22:19Z

cpp/tests/sparse/preprocess.cu

-                   raft::device_vector_view<T1> rows,
-                   raft::device_vector_view<T1> columns,
-                   raft::device_vector_view<T2> values,
+                   raft::device_vector_view<T1, int64_t> rows,


It is explicit but unnecessary.

Use CCCL's mdspan implementation

4178acf

github-project-automation bot added this to Vector Search, ML, & Data Mining Release Board Oct 14, 2025

github-project-automation bot moved this to Todo in Vector Search, ML, & Data Mining Release Board Oct 14, 2025

Remove extraneous comments

f30e7e5

bdice mentioned this pull request Oct 14, 2025

[FEA]: Implement P2642: Padded layouts NVIDIA/cccl#5752

Closed

3 tasks

Add comment about upstream feature request

7f628af

bdice commented Oct 14, 2025

View reviewed changes

bdice added feature request New feature or request breaking Breaking change labels Oct 14, 2025

mhoemmen reviewed Oct 14, 2025

View reviewed changes

This was referenced Oct 14, 2025

Test RAFT with CCCL mdspan rapidsai/cuvs#1430

Draft

Test RAFT with CCCL mdspan rapidsai/cuml#7342

Draft

Test RAFT with CCCL mdspan rapidsai/cugraph#5313

Draft

divyegala reviewed Oct 15, 2025

View reviewed changes

	cuda::std::mdspan<float, cuda::std::extents<int, raft::dynamic_extent, raft::dynamic_extent>>
	cuda::std::mdspan<float, cuda::std::dims<2, int>>

	using matrix_extent = cuda::std::extents<int, dynamic_extent, dynamic_extent>;
	using matrix_extent = cuda::std::dims<2, int>;

	using matrix_extent = cuda::std::extents<size_t, dynamic_extent, dynamic_extent>;
	using matrix_extent = cuda::std::dims<2>;

	using extents_type = cuda::std::extents<size_t, dynamic_extent, dynamic_extent>;
	using extents_type = cuda::std::dims<2>;

	using extent_1d = cuda::std::extents<size_t, dynamic_extent>;
	using extent_1d = cuda::std::dims<1>;

Use CCCL's mdspan implementation #2836

Are you sure you want to change the base?

Use CCCL's mdspan implementation #2836

Conversation

bdice commented Oct 14, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Note

Uh oh!

copy-pr-bot bot commented Oct 14, 2025

Uh oh!

bdice Oct 14, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bdice commented Oct 14, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

bdice commented Oct 14, 2025 •

edited

Loading

bdice Oct 14, 2025 •

edited

Loading