Merge remote-tracking branch 'origin/stable' into sycl_tests

.gitignore

@@ -1,5 +1,6 @@
 .vs
 CMakeSettings.json
+CMakeUserPresets.json
 out
 build*
 Makefile

CMakeLists.txt

@@ -19,6 +19,7 @@ option(MDSPAN_ENABLE_BENCHMARKS "Enable benchmarks." Off)
 option(MDSPAN_ENABLE_COMP_BENCH "Enable compilation benchmarks." Off)
 option(MDSPAN_ENABLE_CUDA "Enable Cuda tests/benchmarks/examples if tests/benchmarks/examples are enabled." Off)
 option(MDSPAN_ENABLE_SYCL "Enable SYCL tests/benchmarks/examples if tests/benchmarks/examples are enabled." Off)
+option(MDSPAN_ENABLE_HIP "Enable HIP tests/benchmarks/examples if tests/benchmarks/examples are enabled." Off)
 option(MDSPAN_ENABLE_OPENMP "Enable OpenMP benchmarks if benchmarks are enabled." On)
 option(MDSPAN_USE_SYSTEM_GTEST "Use system-installed GoogleTest library for tests." Off)
 
@@ -116,6 +117,22 @@ if(MDSPAN_ENABLE_CUDA)
   endif()
 endif()
 
+if(MDSPAN_ENABLE_HIP)
+  include(CheckLanguage)
+  check_language(HIP)
+  if(CMAKE_HIP_COMPILER)
+    message(STATUS "Using ${CMAKE_CXX_STANDARD} as CMAKE_HIP_STANDARD")
+    set(CMAKE_HIP_STANDARD ${CMAKE_CXX_STANDARD})
+    set(CMAKE_HIP_STANDARD_REQUIRED ON)
+    enable_language(HIP)
+  else()
+    message(FATAL_ERROR "Requested HIP support, but no CMAKE_HIP_COMPILER available")
+  endif()
+  if(MDSPAN_ENABLE_TESTS)
+    set(MDSPAN_TEST_LANGUAGE HIP)
+  endif()
+endif()
+
 ################################################################################
 
 add_library(mdspan INTERFACE)

include/experimental/__p0009_bits/config.hpp

@@ -121,11 +121,15 @@ static_assert(_MDSPAN_CPLUSPLUS >= MDSPAN_CXX_STD_14, "mdspan requires C++14 or
 #  define _MDSPAN_NO_UNIQUE_ADDRESS
 #endif
 
+// AMDs HIP compiler seems to have issues with concepts
+// it pretends concepts exist, but doesn't ship <concept>
+#ifndef __HIPCC__
 #ifndef _MDSPAN_USE_CONCEPTS
 #  if defined(__cpp_concepts) && __cpp_concepts >= 201507L
 #    define _MDSPAN_USE_CONCEPTS 1
 #  endif
 #endif
+#endif
 
 #ifndef _MDSPAN_USE_FOLD_EXPRESSIONS
 #  if (defined(__cpp_fold_expressions) && __cpp_fold_expressions >= 201603L) \

include/experimental/__p0009_bits/macros.hpp

@@ -40,6 +40,16 @@
 #  define MDSPAN_INLINE_FUNCTION inline _MDSPAN_HOST_DEVICE
 #endif
 
+#ifndef MDSPAN_FUNCTION
+#  define MDSPAN_FUNCTION _MDSPAN_HOST_DEVICE
+#endif
+
+#ifdef _MDSPAN_HAS_HIP
+#  define MDSPAN_DEDUCTION_GUIDE _MDSPAN_HOST_DEVICE
+#else
+#  define MDSPAN_DEDUCTION_GUIDE
+#endif
+
 // In CUDA defaulted functions do not need host device markup
 #ifndef MDSPAN_INLINE_FUNCTION_DEFAULTED
 #  define MDSPAN_INLINE_FUNCTION_DEFAULTED

include/experimental/__p0009_bits/mdspan.hpp

@@ -321,9 +321,16 @@ class mdspan
 
   MDSPAN_INLINE_FUNCTION
   friend constexpr void swap(mdspan& x, mdspan& y) noexcept {
+    // can't call the std::swap inside on HIP
+    #ifndef _MDSPAN_HAS_HIP
     swap(x.__ptr_ref(), y.__ptr_ref());
     swap(x.__mapping_ref(), y.__mapping_ref());
     swap(x.__accessor_ref(), y.__accessor_ref());
+    #else
+    mdspan tmp = y;
+    y = x;
+    x = tmp;
+    #endif
   }
 
   //--------------------------------------------------------------------------------
@@ -369,44 +376,44 @@ MDSPAN_TEMPLATE_REQUIRES(
   /* requires */ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(is_integral, SizeTypes) /* && ... */) &&
   (sizeof...(SizeTypes) > 0)
 )
-explicit mdspan(ElementType*, SizeTypes...)
+MDSPAN_DEDUCTION_GUIDE explicit mdspan(ElementType*, SizeTypes...)
   -> mdspan<ElementType, ::std::experimental::dextents<size_t, sizeof...(SizeTypes)>>;
 
 MDSPAN_TEMPLATE_REQUIRES(
   class Pointer,
   (_MDSPAN_TRAIT(is_pointer, std::remove_reference_t<Pointer>))
 )
-mdspan(Pointer&&) -> mdspan<std::remove_pointer_t<std::remove_reference_t<Pointer>>, extents<size_t>>;
+MDSPAN_DEDUCTION_GUIDE mdspan(Pointer&&) -> mdspan<std::remove_pointer_t<std::remove_reference_t<Pointer>>, extents<size_t>>;
 
 MDSPAN_TEMPLATE_REQUIRES(
   class CArray,
   (_MDSPAN_TRAIT(is_array, CArray) && (rank_v<CArray> == 1))
 )
-mdspan(CArray&) -> mdspan<std::remove_all_extents_t<CArray>, extents<size_t, ::std::extent_v<CArray,0>>>;
+MDSPAN_DEDUCTION_GUIDE mdspan(CArray&) -> mdspan<std::remove_all_extents_t<CArray>, extents<size_t, ::std::extent_v<CArray,0>>>;
 
 template <class ElementType, class SizeType, size_t N>
-mdspan(ElementType*, const ::std::array<SizeType, N>&)
+MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const ::std::array<SizeType, N>&)
   -> mdspan<ElementType, ::std::experimental::dextents<size_t, N>>;
 
 #ifdef __cpp_lib_span
 template <class ElementType, class SizeType, size_t N>
-mdspan(ElementType*, ::std::span<SizeType, N>)
+MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, ::std::span<SizeType, N>)
   -> mdspan<ElementType, ::std::experimental::dextents<size_t, N>>;
 #endif
 
 // This one is necessary because all the constructors take `data_handle_type`s, not
 // `ElementType*`s, and `data_handle_type` is taken from `accessor_type::data_handle_type`, which
 // seems to throw off automatic deduction guides.
 template <class ElementType, class SizeType, size_t... ExtentsPack>
-mdspan(ElementType*, const extents<SizeType, ExtentsPack...>&)
+MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const extents<SizeType, ExtentsPack...>&)
   -> mdspan<ElementType, ::std::experimental::extents<SizeType, ExtentsPack...>>;
 
 template <class ElementType, class MappingType>
-mdspan(ElementType*, const MappingType&)
+MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const MappingType&)
   -> mdspan<ElementType, typename MappingType::extents_type, typename MappingType::layout_type>;
 
 template <class MappingType, class AccessorType>
-mdspan(const typename AccessorType::data_handle_type, const MappingType&, const AccessorType&)
+MDSPAN_DEDUCTION_GUIDE mdspan(const typename AccessorType::data_handle_type, const MappingType&, const AccessorType&)
   -> mdspan<typename AccessorType::element_type, typename MappingType::extents_type, typename MappingType::layout_type, AccessorType>;
 #endif
 

include/experimental/__p2630_bits/submdspan_mapping.hpp

@@ -125,7 +125,12 @@ submdspan_mapping(const layout_left::mapping<Extents> &src_mapping,
     return mapping_offset<dst_mapping_t>{
         dst_mapping_t(dst_ext, detail::construct_sub_strides(
                                    src_mapping, inv_map,
+    // HIP needs deduction guides to have markups so we need to be explicit
+    #ifdef _MDSPAN_HAS_HIP
+                                   tuple<decltype(detail::stride_of(slices))...>{detail::stride_of(slices)...})),
+    #else
                                    tuple{detail::stride_of(slices)...})),
+    #endif
         static_cast<size_t>(src_mapping(detail::first_of(slices)...))};
   }
 #if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)
@@ -226,7 +231,12 @@ submdspan_mapping(const layout_right::mapping<Extents> &src_mapping,
     return mapping_offset<dst_mapping_t>{
         dst_mapping_t(dst_ext, detail::construct_sub_strides(
                                    src_mapping, inv_map,
+    // HIP needs deduction guides to have markups so we need to be explicit
+    #ifdef _MDSPAN_HAS_HIP
+                                   tuple<decltype(detail::stride_of(slices))...>{detail::stride_of(slices)...})),
+    #else
                                    tuple{detail::stride_of(slices)...})),
+    #endif
         static_cast<size_t>(src_mapping(detail::first_of(slices)...))};
   }
 #if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__)

scripts/snl/gtest-hip-patch

@@ -0,0 +1,102 @@
+diff --git a/googletest/include/gtest/gtest-printers.h b/googletest/include/gtest/gtest-printers.h
+index 7d7e77c1..025ed642 100644
+--- a/googletest/include/gtest/gtest-printers.h
++++ b/googletest/include/gtest/gtest-printers.h
+@@ -473,7 +473,7 @@ GTEST_API_ void PrintTo(char32_t c, ::std::ostream* os);
+ inline void PrintTo(char16_t c, ::std::ostream* os) {
+   PrintTo(ImplicitCast_<char32_t>(c), os);
+ }
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ inline void PrintTo(char8_t c, ::std::ostream* os) {
+   PrintTo(ImplicitCast_<char32_t>(c), os);
+ }
+@@ -586,7 +586,7 @@ inline void PrintTo(const unsigned char* s, ::std::ostream* os) {
+ inline void PrintTo(unsigned char* s, ::std::ostream* os) {
+   PrintTo(ImplicitCast_<const void*>(s), os);
+ }
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ // Overloads for u8 strings.
+ GTEST_API_ void PrintTo(const char8_t* s, ::std::ostream* os);
+ inline void PrintTo(char8_t* s, ::std::ostream* os) {
+@@ -906,7 +906,7 @@ void UniversalPrintArray(const T* begin, size_t len, ::std::ostream* os) {
+ GTEST_API_ void UniversalPrintArray(const char* begin, size_t len,
+                                     ::std::ostream* os);
+
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ // This overload prints a (const) char8_t array compactly.
+ GTEST_API_ void UniversalPrintArray(const char8_t* begin, size_t len,
+                                     ::std::ostream* os);
+@@ -1002,7 +1002,7 @@ template <>
+ class UniversalTersePrinter<char*> : public UniversalTersePrinter<const char*> {
+ };
+
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ template <>
+ class UniversalTersePrinter<const char8_t*> {
+  public:
+diff --git a/googletest/include/gtest/internal/gtest-port.h b/googletest/include/gtest/internal/gtest-port.h
+index b4fa3f07..594a038c 100644
+--- a/googletest/include/gtest/internal/gtest-port.h
++++ b/googletest/include/gtest/internal/gtest-port.h
+@@ -1939,7 +1939,7 @@ inline bool IsUpper(char ch) {
+ inline bool IsXDigit(char ch) {
+   return isxdigit(static_cast<unsigned char>(ch)) != 0;
+ }
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ inline bool IsXDigit(char8_t ch) {
+   return isxdigit(static_cast<unsigned char>(ch)) != 0;
+ }
+diff --git a/googletest/src/gtest-printers.cc b/googletest/src/gtest-printers.cc
+index d475ad36..7ccd181d 100644
+--- a/googletest/src/gtest-printers.cc
++++ b/googletest/src/gtest-printers.cc
+@@ -214,7 +214,7 @@ static const char* GetCharWidthPrefix(signed char) { return ""; }
+
+ static const char* GetCharWidthPrefix(unsigned char) { return ""; }
+
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ static const char* GetCharWidthPrefix(char8_t) { return "u8"; }
+ #endif
+
+@@ -230,7 +230,7 @@ static CharFormat PrintAsStringLiteralTo(char c, ostream* os) {
+   return PrintAsStringLiteralTo(ToChar32(c), os);
+ }
+
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ static CharFormat PrintAsStringLiteralTo(char8_t c, ostream* os) {
+   return PrintAsStringLiteralTo(ToChar32(c), os);
+ }
+@@ -393,7 +393,7 @@ void UniversalPrintArray(const char* begin, size_t len, ostream* os) {
+   UniversalPrintCharArray(begin, len, os);
+ }
+
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ // Prints a (const) char8_t array of 'len' elements, starting at address
+ // 'begin'.
+ void UniversalPrintArray(const char8_t* begin, size_t len, ostream* os) {
+@@ -436,7 +436,7 @@ void PrintCStringTo(const Char* s, ostream* os) {
+
+ void PrintTo(const char* s, ostream* os) { PrintCStringTo(s, os); }
+
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ void PrintTo(const char8_t* s, ostream* os) { PrintCStringTo(s, os); }
+ #endif
+
+@@ -528,7 +528,7 @@ void PrintStringTo(const ::std::string& s, ostream* os) {
+   }
+ }
+
+-#ifdef __cpp_char8_t
++#ifdef __cpp_lib_char8_t
+ void PrintU8StringTo(const ::std::u8string& s, ostream* os) {
+   PrintCharsAsStringTo(s.data(), s.size(), os);
+ }

tests/CMakeLists.txt

@@ -61,4 +61,7 @@ if(NOT CMAKE_CXX_STANDARD STREQUAL "14")
 mdspan_add_test(test_submdspan)
 mdspan_add_test(test_submdspan_static_slice)
 endif()
+# both of those don't work yet since its using vector
+if(NOT MDSPAN_ENABLE_CUDA AND NOT MDSPAN_ENABLE_HIP)
 mdspan_add_test(test_mdarray_ctors)
+endif()

tests/offload_utils.hpp

@@ -18,6 +18,13 @@
 #include <sycl/sycl.hpp>
 #endif
 
+#ifdef _MDSPAN_HAS_HIP
+#include <hip/hip_runtime.h>
+#include <hip/hip_runtime_api.h>
+#endif
+
+#include<cstdio>
+
 namespace {
 bool dispatch_host = true;
 
@@ -35,7 +42,23 @@ if (!(LHS == RHS)) { \
 }
 #endif
 
-#ifdef _MDSPAN_HAS_CUDA
+#if defined(_MDSPAN_HAS_CUDA) || defined(_MDSPAN_HAS_HIP)
+
+#if defined(_MDSPAN_HAS_CUDA)
+void deviceSynchronize() { (void) cudaDeviceSynchronize(); }
+template<class T>
+void mallocManaged(T** ptr, size_t size) { (void) cudaMallocManaged(ptr, size); }
+template<class T>
+void freeManaged(T* ptr) { (void) cudaFree(ptr); }
+#endif
+
+#if defined(_MDSPAN_HAS_HIP)
+void deviceSynchronize() { (void) hipDeviceSynchronize(); }
+template<class T>
+void mallocManaged(T** ptr, size_t size) { (void) hipMallocManaged(ptr, size); }
+template<class T>
+void freeManaged(T* ptr) { (void) hipFree(ptr); }
+#endif
 
 template<class LAMBDA>
 __global__ void dispatch_kernel(const LAMBDA f) {
@@ -48,7 +71,7 @@ void dispatch(LAMBDA&& f) {
     static_cast<LAMBDA&&>(f)();
   } else {
     dispatch_kernel<<<1,1>>>(static_cast<LAMBDA&&>(f));
-    cudaDeviceSynchronize();
+    deviceSynchronize();
   }
 }
 
@@ -58,7 +81,7 @@ T* allocate_array(size_t size) {
   if(dispatch_host == true)
     ptr = new T[size];
   else
-    cudaMallocManaged(&ptr, sizeof(T)*size);
+    mallocManaged(&ptr, sizeof(T)*size);
   return ptr;
 }
 
@@ -67,7 +90,7 @@ void free_array(T* ptr) {
   if(dispatch_host == true)
     delete [] ptr;
   else
-    cudaFree(ptr);
+    freeManaged(ptr);
 }
 
 #define __MDSPAN_TESTS_RUN_TEST(A) \

tests/test_mdspan_swap.cpp

@@ -102,7 +102,9 @@ TEST(TestMDSpanSwap, std_swap_dynamic_extents) {
   __MDSPAN_TESTS_RUN_TEST(test_mdspan_std_swap_dynamic_extents())
 }
 
-
+// On HIP we actually don't call through to swap via ADL
+// so the foo swap test which has side effects will fail
+#ifndef _MDSPAN_HAS_HIP
 void test_mdspan_foo_swap_dynamic_extents() {
   size_t* errors = allocate_array<size_t>(1);
   errors[0] = 0;
@@ -158,3 +160,4 @@ void test_mdspan_foo_swap_dynamic_extents() {
 TEST(TestMDSpanSwap, foo_swap_dynamic_extents) {
   __MDSPAN_TESTS_RUN_TEST(test_mdspan_foo_swap_dynamic_extents())
 }
+#endif