Fix device/host call issues in CUDA (when not using relaxed constexpr workaround) #360
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change | ||||||||
|---|---|---|---|---|---|---|---|---|---|---|
|
|
@@ -2,6 +2,8 @@ | |||||||||
|
|
||||||||||
| #include <cstddef> | ||||||||||
| #include <type_traits> | ||||||||||
| #include <array> | ||||||||||
| #include <utility> | ||||||||||
|
|
||||||||||
| namespace MDSPAN_IMPL_STANDARD_NAMESPACE { | ||||||||||
| namespace detail { | ||||||||||
|
|
@@ -64,6 +66,112 @@ constexpr struct | |||||||||
| } | ||||||||||
| } stride; | ||||||||||
|
|
||||||||||
| template<class T> | ||||||||||
| MDSPAN_INLINE_FUNCTION | ||||||||||
| constexpr void maybe_unused_variable(const T&) {} | ||||||||||
|
There was a problem hiding this comment. This function seems unused.
Suggested change
|
||||||||||
|
|
||||||||||
|
|
||||||||||
| // same as std::integral_constant but with __host__ __device__ annotations on | ||||||||||
| // the implicit conversion function and the call operator | ||||||||||
|
There was a problem hiding this comment. Why do we need/want conversion from to the std:: counterpart? There was a problem hiding this comment. 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 |
||||||||||
| template <class T, T v> | ||||||||||
| struct integral_constant { | ||||||||||
| using value_type = T; | ||||||||||
| using type = integral_constant<T, v>; | ||||||||||
|
|
||||||||||
| MDSPAN_INLINE_FUNCTION_DEFAULTED | ||||||||||
| constexpr integral_constant() = default; | ||||||||||
|
|
||||||||||
| MDSPAN_INLINE_FUNCTION_DEFAULTED | ||||||||||
| constexpr integral_constant(std::integral_constant<T,v>) {}; | ||||||||||
|
There was a problem hiding this comment.
Suggested change
There was a problem hiding this comment. Can't be host device in the 1st place There was a problem hiding this comment. 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. There was a problem hiding this comment. yeah clear bug (funnily enough one of the primary bugs we had elsewhere) |
||||||||||
|
|
||||||||||
| static constexpr T value = v; | ||||||||||
| MDSPAN_INLINE_FUNCTION constexpr operator value_type() const noexcept { | ||||||||||
|
There was a problem hiding this comment. Prefer |
||||||||||
| return value; | ||||||||||
| } | ||||||||||
| MDSPAN_INLINE_FUNCTION constexpr value_type operator()() const noexcept { | ||||||||||
| return value; | ||||||||||
| } | ||||||||||
| MDSPAN_INLINE_FUNCTION constexpr operator std::integral_constant<T,v>() const noexcept { | ||||||||||
|
There was a problem hiding this comment. No way that can be annotated host device There was a problem hiding this comment. 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. |
||||||||||
| return std::integral_constant<T,v>{}; | ||||||||||
| } | ||||||||||
| }; | ||||||||||
|
|
||||||||||
| // The tuple implementation only comes in play when using capabilities | ||||||||||
| // such as submdspan which require C++17 anyway | ||||||||||
| #if MDSPAN_HAS_CXX_17 | ||||||||||
| template<class T, size_t Idx> | ||||||||||
| struct tuple_member { | ||||||||||
| using type = T; | ||||||||||
| static constexpr size_t idx = Idx; | ||||||||||
| T val; | ||||||||||
| MDSPAN_FUNCTION constexpr T& get() { return val; } | ||||||||||
| MDSPAN_FUNCTION constexpr const T& get() const { return val; } | ||||||||||
| }; | ||||||||||
|
|
||||||||||
| template<size_t SearchIdx, size_t Idx, class T> | ||||||||||
| struct tuple_idx_matcher { | ||||||||||
| using type = tuple_member<T, Idx>; | ||||||||||
| template<class Other> | ||||||||||
| MDSPAN_FUNCTION | ||||||||||
| constexpr auto operator + (Other v) const { | ||||||||||
|
There was a problem hiding this comment. Using There was a problem hiding this comment. we just need something which expands. |
||||||||||
| if constexpr (Idx == SearchIdx) { return *this; } | ||||||||||
| else { return v; } | ||||||||||
| } | ||||||||||
| }; | ||||||||||
|
|
||||||||||
| template<class IdxSeq, class ... Elements> | ||||||||||
| struct tuple_impl; | ||||||||||
|
|
||||||||||
| template<size_t ... Idx, class ... Elements> | ||||||||||
| struct tuple_impl<std::index_sequence<Idx...>, Elements...>: public tuple_member<Elements, Idx> ... { | ||||||||||
|
|
||||||||||
| MDSPAN_FUNCTION | ||||||||||
| constexpr tuple_impl(Elements ... vals):tuple_member<Elements, Idx>{vals}... {} | ||||||||||
|
|
||||||||||
| template<size_t N> | ||||||||||
| MDSPAN_FUNCTION | ||||||||||
| constexpr auto& get() { | ||||||||||
| using base_t = decltype((tuple_idx_matcher<N, Idx, Elements>() + ...) ); | ||||||||||
| return base_t::type::get(); | ||||||||||
| } | ||||||||||
| template<size_t N> | ||||||||||
| MDSPAN_FUNCTION | ||||||||||
| constexpr const auto& get() const { | ||||||||||
| using base_t = decltype((tuple_idx_matcher<N, Idx, Elements>() + ...) ); | ||||||||||
| return base_t::type::get(); | ||||||||||
| } | ||||||||||
| }; | ||||||||||
|
|
||||||||||
| // A simple tuple-like class for representing slices internally and is compatible with device code | ||||||||||
| // This doesn't support type access since we don't need it | ||||||||||
| // This is not meant as an external API | ||||||||||
| template<class ... Elements> | ||||||||||
| struct tuple: public tuple_impl<decltype(std::make_index_sequence<sizeof...(Elements)>()), Elements...> { | ||||||||||
| MDSPAN_FUNCTION | ||||||||||
| constexpr tuple(Elements ... vals):tuple_impl<decltype(std::make_index_sequence<sizeof...(Elements)>()), Elements ...>(vals ...) {} | ||||||||||
| }; | ||||||||||
|
|
||||||||||
| template<size_t Idx, class ... Args> | ||||||||||
| MDSPAN_FUNCTION | ||||||||||
| constexpr auto& get(tuple<Args...>& vals) { return vals.template get<Idx>(); } | ||||||||||
|
|
||||||||||
| template<size_t Idx, class ... Args> | ||||||||||
| MDSPAN_FUNCTION | ||||||||||
| constexpr const auto& get(const tuple<Args...>& vals) { return vals.template get<Idx>(); } | ||||||||||
|
|
||||||||||
| template<class ... Elements> | ||||||||||
| tuple(Elements ...) -> tuple<Elements...>; | ||||||||||
| #endif | ||||||||||
|
|
||||||||||
| template<class T, size_t ... Idx> | ||||||||||
| constexpr auto c_array_to_std(std::index_sequence<Idx...>, const T(&values)[sizeof...(Idx)]) { | ||||||||||
|
There was a problem hiding this comment. I would prefer calling it There was a problem hiding this comment. Also thinking you could should write a for loop instead of a fold expression There was a problem hiding this comment. got rid of the function for now, not really needed for what we do. |
||||||||||
| return std::array<T, sizeof...(Idx)>{values[Idx]...}; | ||||||||||
| } | ||||||||||
| template<class T, size_t N> | ||||||||||
| constexpr auto c_array_to_std(const T(&values)[N]) { | ||||||||||
| return c_array_to_std(std::make_index_sequence<N>(), values); | ||||||||||
| } | ||||||||||
|
|
||||||||||
| } // namespace detail | ||||||||||
|
|
||||||||||
| constexpr struct mdspan_non_standard_tag { | ||||||||||
|
|
||||||||||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This part of the comment is no longer needed.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
There are users of mdspan which are not Kokkos users and rely on C++14.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
but yeah good observation I just pushed something to remove the comments