Fix device/host call issues in CUDA (when not using relaxed constexpr workaround) #360

nmm0 · 2024-10-02T19:27:58Z

This includes code that Christian wrote and will be included into Kokkos in kokkos/kokkos#7396

This addresses warnings for calling constexpr device code inside of Kokkos with BasicView

…code enabled

…_constant

nmm0 · 2024-10-02T21:14:10Z

include/experimental/__p0009_bits/layout_stride.hpp

@@ -199,7 +199,15 @@ struct layout_stride {
      template<class IntegralType>
      MDSPAN_INLINE_FUNCTION
      static constexpr const __strides_storage_t fill_strides(const std::array<IntegralType,extents_type::rank()>& s) {
+        // avoid warning for use of host std::array operator[]
+  #if defined(_MDSPAN_HAS_CUDA) || defined(_MDSPAN_HAS_HIP)
+        const IntegralType* s_ptr = reinterpret_cast<const IntegralType*>(&s);


We need to figure out what to do properly here, since reinterpret_cast can't be used in a constant expression, so we will remove it for now

std::array isn't usable on Cuda devices anyway unless you are using -expt-relaxed-constexpr, right? Wouldn't an interface in terms of C-style arrays be a more natural choice?

Yes and no: we have constructors which have to take std::array because that is the standard. Also a bunch of operations actually do work on device (i.e. construction, destruction, structured binding - because they don't go through member functions, or at least go through defaulted member functions).

That said the above casting approach doesn't work in constexpr so we changed this.

nmm0 · 2024-10-02T21:27:57Z

Move device_support to utility
Make construct_sub_strides use a custom c-array like type instead of std::array
Remove fake-tuple type access
Make std::array constructors for layout_stride non-device

…s to work on device

… about what it's used for

tcclevenger

Just curious, if C++17 is the minimum req for Kokkos core, why do we test C++14 for mdspan. Is this the same for other tpls or containers/algorithms?

tcclevenger · 2024-10-03T14:49:42Z

include/experimental/__p0009_bits/layout_stride.hpp

@@ -206,7 +205,7 @@ struct layout_stride {
        class IntegralType,
        // The is_convertible condition is added to make sfinae valid
        // the extents_type::rank() > 0 is added to avoid use of non-standard zero length c-array


Suggested change

// the extents_type::rank() > 0 is added to avoid use of non-standard zero length c-array

This part of the comment is no longer needed.

There are users of mdspan which are not Kokkos users and rely on C++14.

but yeah good observation I just pushed something to remove the comments

masterleinad · 2024-10-03T15:27:53Z

include/experimental/__p0009_bits/layout_stride.hpp

@@ -206,7 +205,7 @@ struct layout_stride {
        class IntegralType,
        // The is_convertible condition is added to make sfinae valid
        // the extents_type::rank() > 0 is added to avoid use of non-standard zero length c-array
-        (std::is_convertible<IntegralType, typename extents_type::index_type>::value && (extents_type::rank() > 0))
+        (std::is_convertible<IntegralType, typename extents_type::index_type>::value)
      )
      MDSPAN_INLINE_FUNCTION
      // despite the requirement some compilers still complain about zero length array during parsing


This comment seems out-of-date now.

masterleinad · 2024-10-03T15:28:13Z

include/experimental/__p0009_bits/layout_stride.hpp

@@ -206,7 +205,7 @@ struct layout_stride {
        class IntegralType,
        // The is_convertible condition is added to make sfinae valid
        // the extents_type::rank() > 0 is added to avoid use of non-standard zero length c-array


Suggested change

// the extents_type::rank() > 0 is added to avoid use of non-standard zero length c-array

masterleinad · 2024-10-03T15:29:58Z

include/experimental/__p0009_bits/utility.hpp

+template<class T>
+MDSPAN_INLINE_FUNCTION
+constexpr void maybe_unused_variable(const T&) {}


This function seems unused.

Suggested change

template<class T>

MDSPAN_INLINE_FUNCTION

constexpr void maybe_unused_variable(const T&) {}

masterleinad · 2024-10-03T15:36:57Z

include/experimental/__p0009_bits/utility.hpp

+  MDSPAN_INLINE_FUNCTION_DEFAULTED
+  constexpr integral_constant(std::integral_constant<T,v>) {};


DEFAULTED seems inappropriate. Do you need the extra annotations at all here?

Suggested change

MDSPAN_INLINE_FUNCTION_DEFAULTED

constexpr integral_constant(std::integral_constant<T,v>) {};

MDSPAN_INLINE_FUNCTION

constexpr integral_constant(std::integral_constant<T,v>) {};

Can't be host device in the 1st place

yes it can. Nothing in this calls device functions. Constructor of std::integral_constant is defaulted, so it works too. The only thing which doesn't work of std::integral_constant is the conversion operator to its scaler type.

Then what Daniel said

(minus the inline)

yeah clear bug (funnily enough one of the primary bugs we had elsewhere)

masterleinad · 2024-10-03T15:49:21Z

include/experimental/__p0009_bits/utility.hpp

+  using type = tuple_member<T, Idx>;
+  template<class Other>
+  MDSPAN_FUNCTION
+  constexpr auto operator + (Other v) const {


Using + seems a little weird. What | or || be slightly better?

we just need something which expands.

dalg24 · 2024-10-03T15:56:51Z

include/experimental/__p0009_bits/utility.hpp

+  constexpr integral_constant(std::integral_constant<T,v>) {};
+
+  static constexpr T value = v;
+  MDSPAN_INLINE_FUNCTION constexpr operator value_type() const noexcept {


Prefer MDSPAN_FUNCTION here and below

dalg24 · 2024-10-03T15:57:59Z

include/experimental/__p0009_bits/utility.hpp

+
+
+// same as std::integral_constant but with __host__ __device__ annotations on
+// the implicit conversion function and the call operator


Why do we need/want conversion from to the std:: counterpart?

because there are places where we pass in std::integral_constant, and std::integral_constant is the standard approved way to pass certain ways to certain functions. For example submdspan(a, std::integral_constant<int, 1>()) is fine in device code.

dalg24 · 2024-10-03T15:58:29Z

include/experimental/__p0009_bits/utility.hpp

+  MDSPAN_INLINE_FUNCTION constexpr value_type operator()() const noexcept {
+    return value;
+  }
+  MDSPAN_INLINE_FUNCTION constexpr operator std::integral_constant<T,v>() const noexcept {


No way that can be annotated host device

Of course it works:

https://godbolt.org/z/bb8qdT31Y

We can use any standard library stuff as long as we don't call non-defaulted member functions. And we don't here.

dalg24 · 2024-10-03T15:58:44Z

include/experimental/__p0009_bits/utility.hpp

+  MDSPAN_INLINE_FUNCTION_DEFAULTED
+  constexpr integral_constant(std::integral_constant<T,v>) {};


Can't be host device in the 1st place

dalg24 · 2024-10-03T22:48:31Z

include/experimental/__p0009_bits/utility.hpp

+#endif
+
+template<class T, size_t ... Idx>
+constexpr auto c_array_to_std(std::index_sequence<Idx...>, const T(&values)[sizeof...(Idx)]) {


I would prefer calling it to_std_array

Also thinking you could should write a for loop instead of a fold expression

got rid of the function for now, not really needed for what we do.

dalg24 · 2024-10-03T22:52:37Z

include/experimental/__p2630_bits/submdspan_mapping.hpp

+template<class T, size_t N>
+struct sub_strides
+{
+  T values[N];


Suggested change

T values[N];

T values[N > 0 ? N : 1];

Done (and delete the specialization below) unless you prefer the specialization path.

dalg24 · 2024-10-04T16:38:53Z

include/experimental/__p0009_bits/layout_stride.hpp

@@ -476,7 +468,8 @@ struct layout_stride {
    MDSPAN_INLINE_FUNCTION
    constexpr index_type required_span_size() const noexcept {
      index_type span_size = 1;
-      for(unsigned r = 0; r < extents_type::rank(); r++) {
+      // using int here to avoid warning about pointless comparison to 0


I suppose you mean the rank 0 case

dalg24 · 2024-10-04T16:42:13Z

include/experimental/__p2630_bits/submdspan_extents.hpp

+template <class IdxT1, class IdxT2, class IndexType>
+struct index_pair_like<tuple<IdxT1, IdxT2>, IndexType> {


Might be worth a comment "our device-compatible poor man's implementation of tuple" for the future reader that might be surprised to see that specialization interleaved with all the std:: ones.

dalg24 · 2024-10-04T16:43:13Z

include/experimental/__p2630_bits/submdspan_extents.hpp

+template<class Integral, Integral v>
+MDSPAN_INLINE_FUNCTION
+constexpr Integral first_of(const std::integral_constant<Integral, v>&) {
+  return integral_constant<Integral, v>();
+}


This is weird

Why do we need this one

dalg24 · 2024-10-04T16:46:01Z

include/experimental/__p2630_bits/submdspan_extents.hpp

+MDSPAN_TEMPLATE_REQUIRES(
+  class IdxT1, class IdxT2,
+  /* requires */ (index_pair_like<std::tuple<IdxT1, IdxT2>, size_t>::value)
+  )
+constexpr auto first_of(const std::tuple<IdxT1, IdxT2>& i) {
+  return get<0>(i);
+}
+
+MDSPAN_TEMPLATE_REQUIRES(
+  class IdxT1, class IdxT2,
+  /* requires */ (index_pair_like<std::pair<IdxT1, IdxT2>, size_t>::value)
+  )
+MDSPAN_INLINE_FUNCTION
+constexpr auto first_of(const std::pair<IdxT1, IdxT2>& i) {
+  return i.first;


Can't we use get() for all of these

get on std::pair is not host-device capable.

nmm0 added 3 commits October 2, 2024 14:37

initial workarounds to get mdspan copmiling without constexpr device …

a0c6886

…code enabled

make device-compatible tuple standin constexpr

56d8424

Make sure StaticExtentFromRange recognizes device-compatible integral…

00969d8

…_constant

nmm0 added a commit to nmm0/kokkos that referenced this pull request Oct 2, 2024

include kokkos/mdspan#360 changes

2aa3582

nmm0 mentioned this pull request Oct 2, 2024

MDSpan library changes for constexpr device code kokkos/kokkos#7396

Closed

avoid using arraay.data() since in C++14 it's not constexpr

aaa2b87

nmm0 added a commit to nmm0/kokkos that referenced this pull request Oct 2, 2024

include kokkos/mdspan#360 changes

52770ad

nmm0 commented Oct 2, 2024

View reviewed changes

nmm0 and others added 6 commits October 2, 2024 17:31

merge device_support.hpp into utility.hpp

37283cd

make layout_stride std::array ctors host-only

97c1813

make submdspan_mapping create strided mapping without std::array so a…

0fcdad7

…s to work on device

remove type access from our tuple standin for device code and comment…

9f9a555

… about what it's used for

Add missing includes in utility

43d0595

Fix C++14 build

b0c2452

tcclevenger reviewed Oct 3, 2024

View reviewed changes

Fix comments

e818f4d

masterleinad reviewed Oct 3, 2024

View reviewed changes

dalg24 reviewed Oct 3, 2024

View reviewed changes

crtrott added 2 commits October 3, 2024 17:23

More review comment fixes

f86bfe8

Use | operator instead of +

f5017f3

masterleinad approved these changes Oct 4, 2024

View reviewed changes

crtrott approved these changes Oct 4, 2024

View reviewed changes

crtrott changed the title ~~Remove constexpr device code warnings~~ Fix device/host call issues in CUDA (when not using relaxed constexpr workaround) Oct 4, 2024

crtrott merged commit 260f525 into kokkos:stable Oct 4, 2024
15 checks passed

dalg24 reviewed Oct 4, 2024

View reviewed changes

	template<class T>
	MDSPAN_INLINE_FUNCTION
	constexpr void maybe_unused_variable(const T&) {}

		MDSPAN_INLINE_FUNCTION_DEFAULTED
		constexpr integral_constant(std::integral_constant<T,v>) {};



		// same as std::integral_constant but with __host__ __device__ annotations on
		// the implicit conversion function and the call operator

		template <class IdxT1, class IdxT2, class IndexType>
		struct index_pair_like<tuple<IdxT1, IdxT2>, IndexType> {

Fix device/host call issues in CUDA (when not using relaxed constexpr workaround) #360

Fix device/host call issues in CUDA (when not using relaxed constexpr workaround) #360

Conversation

nmm0 commented Oct 2, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

nmm0 commented Oct 2, 2024 • edited Loading

tcclevenger 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

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

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

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

nmm0 commented Oct 2, 2024 •

edited

Loading