Add `mdspan` reference implementation #299

youyu3 · 2022-08-09T21:58:20Z

Pulls the mdspan reference implementation from branch "stable" of the kokkos repo, https://github.com/kokkos/mdspan, up to PR 172.
Uglified internal identifiers and made some naming convention updates.

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/config.hpp

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/default_accessor.hpp

include/cuda/std/experimental/mdspan

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/extents.hpp

include/cuda/std/experimental/mdspan

include/cuda/std/detail/libcxx/include/experimental/mdspan

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/trait_backports.hpp

mhoemmen · 2022-08-24T14:40:52Z

.upstream-tests/test/std/containers/views/mdspan/mdspan.ctor/ctad_extents.pass.cpp

+int main(int, char**)
+{
+    // workaround to avoid compilation errors
+    cuda::std::experimental::detail::__make_dynamic_extent<int>();


@youyu3 Could you please add details of the compiler version and flags here? This feels very much like a compiler bug. __make_dynamic_extent<T>() literally just returns T(-1). Its definition should be available, so it shouldn't need to be explicitly instantiated.

I suspect this is due to a compiler bug too. I'll create a reproducer and update here.

@youyu3 If __make_dynamic_extent doesn't work, we could try this:

template<class... Integrals> requires((is_convertible_v<Integrals, size_t> && ...)) explicit extents(Integrals... indices) -> extents<size_t, size_t(indices, -1)...>;

What matters here is on the right side of the ->. The (C++17) fold expression size_t(indices, -1)... expands to something like size_t(-1), size_t(-1), ..., size_t(-1), where the number of -1s is sizeof...(indices).

Revision:

template <class... IndexTypes> extents(IndexTypes...) -> extents<std::size_t, std::size_t((IndexTypes(), -1))...>;

The comma expression needed an extra (), else it was being interpreted as a cast with two arguments (not a valid cast).

Thanks! Pushed the workaround.

mhoemmen · 2022-08-31T21:16:03Z

...de/cuda/std/detail/libcxx/include/experimental/__p0009_bits/standard_layout_static_array.hpp

-#include <cuda/std/array>
+#ifndef __cuda_std__
+#include <array>
+#endif
 #ifdef __cpp_lib_span
 #include <span>


Do we also need to protect #include  with #ifndef __cuda_std__ ... #endif ?

We don't have a libcudacxx version of span (or I didn't find one).

@youyu3 Oh, oof : - ( Should we file an issue to add one? We would need it for mdspan just for a few things that take span as input.

Now we have and it worked for my initial tests. More tests to be added.

mhoemmen · 2022-09-02T14:43:22Z

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/layout_left.hpp

@@ -73,21 +73,21 @@ class layout_left::mapping {
    struct __rank_count {};

    template <::std::size_t _r, ::std::size_t _Rank, class _Ip, class... _Indices>
-    __MDSPAN_INLINE_FUNCTION
+    __MDSPAN_HOST_DEVICE


@youyu3 I'm a bit confused -- doesn't MDSPAN_INLINE_FUNCTION include MDSPAN_HOST_DEVICE?

That's right. But I'm only adding __host__ __device__ while matching the declaration before the change, where the function wasn't declared as inline.

mhoemmen · 2022-09-08T20:25:46Z

.upstream-tests/test/std/containers/views/mdspan/mdspan.layout.right.ctor/list_init.pass.cpp

+
+        static_assert( m.is_exhaustive() == true );
+
+        assert( m.extents().rank()         == 2  );


[OPTIONAL] rank() or rank_dynamic() are always static constexpr and thus can be checked with static_assert.

Switched to static_assert.

mhoemmen · 2022-09-08T20:26:18Z

.upstream-tests/test/std/containers/views/mdspan/mdspan.layout.right.ctor/list_init.pass.cpp

+
+        stdex::layout_right::mapping<stdex::extents<size_t,dyn, dyn>> m{stdex::dextents<size_t,2>{16, 32}};
+
+        static_assert( m.is_exhaustive() == true );


is_exhaustive() happens to be static constexpr for layout_right, but please note that this is not generally true for all layout mappings.

Thanks. I'll change this to assert then.

It's OK here -- I'm just pointing this out in case you want to write generic tests. Thanks!

miscco · 2022-09-20T09:35:05Z

Generally speaking I am extremely concerned about the implementation complexity here.

I only looked at extends in detail, because I did implement it previously.

Especially compile times are going to be incredibly difficult here.

Also what I am generally wondering is why we are putting this in namespace stdex. It was merged into the standard after all

mhoemmen · 2022-09-20T15:58:49Z

Hi @miscco ! We talked a bit about extents offline. An excellent first step would be to run the benchmarks both with the current version, and with the current version with your extents implementation substituted in. It would also be interesting to see if your implementation performs well at lower optimization levels. (This would cover the important use case of debug builds not being horribly slow, which is important for game development.)

Also what I am generally wondering is why we are putting this in namespace stdex. It was merged into the standard after all

This is a user library (not a vendor's implementation of the Standard Library). Thus, it's not allowed to put things in namespace std.

mdspan was merged into C++23. The reference mdspan implementation works with C++14, 17, and 20 as well. For those back-porting use cases, it's important that we don't break applications' builds, in case those applications adopt a later compiler version that provides an mdspan implementation. Thus, we don't want to collide with std::mdspan.

mhoemmen · 2022-09-20T16:46:35Z

@miscco I wonder if some of the complexity of the reference mdspan implementation comes from an earlier requirement to back-port to C++11 (with its much more restrictive requirements for constexpr functions). That requirement has since been dropped. The reference mdspan implementation currently requires at least C++14. I'll file an issue in the reference mdspan implementation to investigate this.

crtrott · 2022-09-20T17:26:46Z

I think the comment about extents could be simpler are fair. Part of the problem was trying to work around a number of issues with empty base optimization in MSVC and Intel compiler. MSVC did get much better though in the last couple years. That said: looking at some assembly it does look like the reference implementation compiles cleaner:

https://godbolt.org/z/Gvsfrn31o

crtrott · 2022-09-20T17:29:25Z

Thinking more about this: there was a lot of work on extents, to essentially make the .extent(i) function optimize really well. Without the complication I think you end up with something which goes through some loop structures to figure out which element to access. I.e. you kinda rebuild the mapping for index to "which element from the dynamic extents array do I need to access" every time or so? But am not 100% sure.

miscco · 2022-09-20T20:13:35Z

Yeah I believe my _Dynamic_index function is the issue, I will try to memoize that in a second static array, so there is no runtime cost

crtrott · 2022-09-20T20:24:23Z

Note: with -O3 and GCC 11 it does compile into the same instruction sequence. I just told Mark, I am totally open to entertaining a whole sale replacement on our side. I think we should underpin this with some data down the line, for example we are working on using mdspan in Kokkos which then would make a whole lot of applications use it. If we collect data from that side regarding compile time/binary size/performance and from your applications we should have a good basis to make decisions.

jrhemstad · 2022-09-21T13:39:52Z

Can we defer any significant refactoring discussion to future work/PR? I don't want to delay @youyu3's PR any more than necessary.

mhoemmen · 2022-09-21T13:46:56Z

@jrhemstad Don't worry! These investigations should not affect the current release. Discussions continue on the reference mdspan implementation issue kokkos/mdspan#190 .

miscco · 2022-09-21T15:30:18Z

include/cuda/std/detail/libcxx/include/experimental/mdspan

+// 1. Redistributions of source code must retain the above copyright
+// notice, this list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright


Just to be sure, is that license compatible with the usual one.

Do we have SWIPAT clearance for adding. Binary distribution clauses scare me a lot

mhoemmen · 2022-10-02T22:25:25Z

libcxx/test/std/containers/views/mdspan/mdspan.extents.cons/convertible_to_size_t.pass.cpp

+#include <experimental/mdspan>
+#include <cassert>
+#include <array>
+#include <span>


Suggested change

#include 

#ifdef __cpp_lib_span

#include 

#endif

I don't think you can even include  unless that feature test macro is defined.

Thanks Mark, you are right. However, I don't think an inclusion of  is even needed here, since it's already included by <mdspan>. Tests ran fine without this line.
Actually, #include <array> is not needed here either.

@youyu3 That's reasonable, though please note that mdspan only includes it if __cpp_lib_span is defined. For example: https://github.com/kokkos/mdspan/blob/9e2b68a3be59c867587cdfbb641153199c813ff5/include/experimental/__p0009_bits/standard_layout_static_array.hpp#L56

Yeah I'm aware of that. And any test that uses span should be guarded by #ifdef __cpp_lib_span as well (which has been fixed).

jrhemstad · 2022-10-03T20:05:03Z

Hey @youyu3 sorry for the last minute change, but based on some internal conversation, I think we can drop the experimental namespace for mdspan.

youyu3 · 2022-10-03T22:13:21Z

Hey @youyu3 sorry for the last minute change, but based on some internal conversation, I think we can drop the experimental namespace for mdspan.

Okay. Will do. Then I suppose that we should move the files out of the experimental sub-directory as well?

youyu3 · 2022-10-04T22:44:04Z

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/dynamic_extent.hpp

@@ -52,9 +52,6 @@
 #endif

 _LIBCUDACXX_BEGIN_NAMESPACE_STD
-namespace experimental {
-
-__MDSPAN_INLINE_VARIABLE constexpr auto dynamic_extent = ::std::numeric_limits<::std::size_t>::max();


This is needed when span is not available. Fixed in a later commit...

mhoemmen · 2022-11-07T18:52:32Z

.upstream-tests/test/std/containers/views/mdspan/mdspan.extents.cons/array.pass.cpp

@@ -0,0 +1,64 @@
+//===----------------------------------------------------------------------===//


Excellent test!

mhoemmen · 2022-11-07T22:11:08Z

It's far too much for me to review in detail, but I'm really impressed and happy with all the new unit tests! Thank you also for addressing those code comments we talked about earlier today!

miscco

I am not fully through with the product code.

That said, I want to reiterate, that I do not believe this is code is maintainable for us. I accept that we currently do not have capacity to write it from scratch and working code is definitely better than no code.

Nevertheless, I would want a commitment that we actually get the time to clean this up once we get the big features we are working on right now out of the window.

Before we can merge it we would need to remove all remnants of std:: and make sure we are using the proper namespaces throughout.

include/cuda/std/detail/libcxx/include/__mdspan/macros.hpp

miscco · 2022-11-08T07:55:50Z

include/cuda/std/detail/libcxx/include/__mdspan/compressed_pair.hpp

+  }
+
+  __MDSPAN_INLINE_FUNCTION_DEFAULTED
+  constexpr __compressed_pair() noexcept = default;


Conventionally, we do not decorate defaulted or delete functions, so this would be
__compressed_pair() = default;

That said this is definitely something we can do afterwards on our own, so no change requested

@miscco wrote:

Conventionally, we do not decorate defaulted or delete functions....

Does this mean that =defaulted functions do not need the decoration? Historically, we've had issues with that.

include/cuda/std/detail/libcxx/include/__mdspan/type_list.hpp

include/cuda/std/detail/libcxx/include/__mdspan/trait_backports.hpp

miscco · 2022-11-08T08:33:28Z

include/cuda/std/detail/libcxx/include/__mdspan/standard_layout_static_array.hpp

+ * @tparam _ValsSeq
+ * @tparam __sentinal
+ */
+template <class _Tag, class _Tp, class _static_t, class _ValsSeq, _static_t __sentinal = static_cast<_static_t>(dynamic_extent),


This whole file is pain

@miscco I can never get a certain someone to spell "sentinel" right ; - ) .

wmaxey

Commit history appears to be broken. I'll try to resolve this locally and possibly issue a correction, but AFAICT there seems to be a replay of several commits.

youyu3 · 2022-11-15T22:44:00Z

I pulled commits from the span PR. There are further changes in those files I believe.

Commit history appears to be broken. I'll try to resolve this locally and possibly issue a correction, but AFAICT there seems to be a replay of several commits.

wmaxey · 2022-11-17T08:35:26Z

...tream-tests/test/std/containers/views/mdspan/mdspan.accessor.default.members/access.pass.cpp

+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+


All tests will likely need //UNSUPPORTED: c++11 below the license and above the includes.

This letds lit know that the tests are not designed for pre-C++14.

https://blossom.nvidia.com/sw-cuda-cccl-instance/blue/organizations/jenkins/cluster_launch_test/detail/cluster_launch_test/386/pipeline/1129

Results are crunchy, but it looks like there is maybe only a few failure points.

I made a PR on gitlab that fixes the history. It drops the  PR and gets rid of some strange re-merge, but I think the content is identical.

Thanks Wesley. I can add the //UNSUPPORTED: c++11 to the github PR. Would you be able to merge that to the gitlab PR for CI?

Also, what CUDA versions do our CI tests cover?

mhoemmen · 2022-11-28T01:23:42Z

include/cuda/std/detail/libcxx/include/__mdspan/standard_layout_static_array.hpp

@@ -534,17 +534,9 @@ struct __standard_layout_psa<_Tag, _Tp, _static_t, _CUDA_VSTD::integer_sequence<
  //--------------------------------------------------------------------------

  __MDSPAN_INLINE_FUNCTION_DEFAULTED
-  constexpr __standard_layout_psa() noexcept
-#if defined(__clang__) || defined(__MDSPAN_DEFAULTED_CONSTRUCTORS_INHERITANCE_WORKAROUND)


@youyu3
Did the above change fix this issue, making this Clang work-around no longer needed?

In my tests with clang, yes the work-around is no longer needed.

mhoemmen · 2022-11-28T19:13:35Z

libcxx/test/std/containers/views/mdspan/mdspan.extents.cons/span.pass.cpp

@@ -19,7 +21,7 @@ void test_span_con()
    using TestFixture = TestExtents<T>;
    TestFixture t;

-    auto s = std::span( t.dyn_sizes );
+    auto s = std::span<const size_t, t.dyn_sizes.size()>( t.dyn_sizes );


@youyu3 If you're using std::span, you'll need C++20. Would you consider using cuda::std::span (if it exists) instead?

This test is for testing libcudacxx in the libcxx mode. std::span is implemented in this mode.

Ah excellent, thank you for explaning!

mhoemmen · 2022-11-30T22:27:56Z

.upstream-tests/test/std/containers/views/mdspan/mdspan.extents.cons/array.pass.cpp

@@ -49,18 +49,18 @@ int main(int, char**)
    test_array_con< cuda::std::tuple_element_t< 4, extents_test_types > >();
    test_array_con< cuda::std::tuple_element_t< 5, extents_test_types > >();

-    static_assert( is_array_cons_avail_v< cuda::std::dextents<int,2>, int   , 2 > == true  );
+    static_assert( is_array_cons_avail_v< cuda::std::dextents<int,2>, int   , 2 > == true , "" );


Just FYI, these are changes to support C++14 -- C++17 implements one-argument static_assert.

Also make it available in C++14 onwards

…_VAR`

…`constexpr`

mhoemmen reviewed Aug 16, 2022

View reviewed changes

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/config.hpp Outdated Show resolved Hide resolved

wmaxey reviewed Aug 17, 2022

View reviewed changes

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/default_accessor.hpp Outdated Show resolved Hide resolved

wmaxey reviewed Aug 17, 2022

View reviewed changes

include/cuda/std/experimental/mdspan Outdated Show resolved Hide resolved

wmaxey reviewed Aug 17, 2022

View reviewed changes

include/cuda/std/experimental/mdspan Outdated Show resolved Hide resolved

mhoemmen reviewed Aug 19, 2022

View reviewed changes

wmaxey reviewed Aug 19, 2022

View reviewed changes

include/cuda/std/experimental/mdspan Outdated Show resolved Hide resolved

wmaxey reviewed Aug 19, 2022

View reviewed changes

include/cuda/std/detail/libcxx/include/experimental/mdspan Outdated Show resolved Hide resolved

wmaxey reviewed Aug 19, 2022

View reviewed changes

include/cuda/std/detail/libcxx/include/experimental/__p0009_bits/trait_backports.hpp Outdated Show resolved Hide resolved

mhoemmen reviewed Aug 24, 2022

View reviewed changes

mhoemmen reviewed Aug 31, 2022

View reviewed changes

mhoemmen reviewed Sep 8, 2022

View reviewed changes

mhoemmen mentioned this pull request Sep 20, 2022

Investigate making the implementation simpler, and benchmark the effects kokkos/mdspan#190

Open

miscco reviewed Sep 21, 2022

View reviewed changes

youyu3 changed the title ~~Adding mdspan reference implementation~~ Add mdspan reference implementation Sep 30, 2022

mhoemmen suggested changes Oct 2, 2022

View reviewed changes

youyu3 commented Oct 4, 2022

View reviewed changes

miscco force-pushed the mdspan branch from edfcb7d to 7f097a4 Compare October 11, 2022 08:01

jrhemstad requested review from miscco and mhoemmen November 7, 2022 18:06

jrhemstad requested review from mhoemmen and wmaxey and removed request for mhoemmen November 7, 2022 18:06

mhoemmen approved these changes Nov 7, 2022

View reviewed changes

wmaxey added the testing: internal ci in progress Currently testing on internal NVIDIA CI (DVS). label Nov 8, 2022

miscco reviewed Nov 8, 2022

View reviewed changes

wmaxey suggested changes Nov 15, 2022

View reviewed changes

wmaxey reviewed Nov 17, 2022

View reviewed changes

mhoemmen reviewed Nov 28, 2022

View reviewed changes

mhoemmen reviewed Nov 30, 2022

View reviewed changes

miscco and others added 7 commits February 9, 2023 14:10

Move as_const into its own file

8beb664

Also make it available in C++14 onwards

Implement std::mdspan

41d52db

Make the header conform to our internal architecture

c834cf7

Replace usages of __MDSPAN_INLINE_VARIABLE with `_LIBCUDACXX_INLINE…

5cb6a2a

…_VAR`

We are always at least C++14 so replace __MDSPAN_CONSTEXPR_14 with …

4f45d79

…`constexpr`

Add some necessary __uglification

1e25eb1

Bring back all the tests

24702f0

youyu3 force-pushed the mdspan branch from a4bcd00 to 24702f0 Compare February 13, 2023 19:47

miscco approved these changes Feb 13, 2023

View reviewed changes

wmaxey approved these changes Feb 13, 2023

View reviewed changes

miscco merged commit 639bcbc into NVIDIA:main Feb 14, 2023


		static_assert( m.is_exhaustive() == true );

		assert( m.extents().rank() == 2 );


		stdex::layout_right::mapping<stdex::extents<size_t,dyn, dyn>> m{stdex::dextents<size_t,2>{16, 32}};

		static_assert( m.is_exhaustive() == true );

		@@ -0,0 +1,64 @@
		//===----------------------------------------------------------------------===//

Add mdspan reference implementation #299

Add mdspan reference implementation #299

Conversation

youyu3 commented Aug 9, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

mhoemmen Aug 24, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

miscco commented Sep 20, 2022

mhoemmen commented Sep 20, 2022

mhoemmen commented Sep 20, 2022 • edited Loading

crtrott commented Sep 20, 2022

crtrott commented Sep 20, 2022

miscco commented Sep 20, 2022 • edited Loading

crtrott commented Sep 20, 2022

jrhemstad commented Sep 21, 2022

mhoemmen commented Sep 21, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

jrhemstad commented Oct 3, 2022

youyu3 commented Oct 3, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

mhoemmen commented Nov 7, 2022 • edited Loading

miscco left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

wmaxey left a comment • edited Loading

Choose a reason for hiding this comment

youyu3 commented Nov 15, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Add `mdspan` reference implementation #299

Add `mdspan` reference implementation #299

mhoemmen Aug 24, 2022 •

edited

Loading

mhoemmen commented Sep 20, 2022 •

edited

Loading

miscco commented Sep 20, 2022 •

edited

Loading

youyu3 commented Oct 3, 2022 •

edited

Loading

mhoemmen commented Nov 7, 2022 •

edited

Loading

wmaxey left a comment •

edited

Loading

youyu3 commented Nov 15, 2022 •

edited

Loading