From 301438b1e7ac79a25761526a8bc11f75738185f4 Mon Sep 17 00:00:00 2001 From: Nicolas Morales Date: Thu, 26 Jan 2023 09:16:37 -0800 Subject: [PATCH 01/13] add some more common file patterns to .gitignore, including CMakeUserPresets.json --- .gitignore | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/.gitignore b/.gitignore index ee7b817e..3d826281 100644 --- a/.gitignore +++ b/.gitignore @@ -1,5 +1,9 @@ .vs CMakeSettings.json +CMakeUserPresets.json out build* Makefile +cmake-build-*/ +*.user +.idea/ From fe6df2aa1bfbf2656c7d5ffb4cec318481ba4a5d Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 17:46:03 -0700 Subject: [PATCH 02/13] Add CMake logic to support HIP test build --- CMakeLists.txt | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index 69abff9f..6086cd04 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -18,6 +18,7 @@ option(MDSPAN_ENABLE_EXAMPLES "Build examples." Off) 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_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) @@ -115,6 +116,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) From ab31507a8b683083d687c3e16adc5b4f4dac2776 Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 17:46:22 -0700 Subject: [PATCH 03/13] Add patch file for gtest so it works with HIP compiler in C++20 mode See report here: https://github.com/google/googletest/issues/3659 --- scripts/snl/gtest-hip-patch | 102 ++++++++++++++++++++++++++++++++++++ 1 file changed, 102 insertions(+) create mode 100644 scripts/snl/gtest-hip-patch diff --git a/scripts/snl/gtest-hip-patch b/scripts/snl/gtest-hip-patch new file mode 100644 index 00000000..7fe5644b --- /dev/null +++ b/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_(c), os); + } +-#ifdef __cpp_char8_t ++#ifdef __cpp_lib_char8_t + inline void PrintTo(char8_t c, ::std::ostream* os) { + PrintTo(ImplicitCast_(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_(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 : public UniversalTersePrinter { + }; + +-#ifdef __cpp_char8_t ++#ifdef __cpp_lib_char8_t + template <> + class UniversalTersePrinter { + 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(ch)) != 0; + } +-#ifdef __cpp_char8_t ++#ifdef __cpp_lib_char8_t + inline bool IsXDigit(char8_t ch) { + return isxdigit(static_cast(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); + } From ff4f6a7f2177d21311dc3ea5610479b838da443a Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 17:48:51 -0700 Subject: [PATCH 04/13] Add macro logic for HIP support Note we need the CTAD stuff because clang wants to see device markup on deduction guides. --- include/experimental/__p0009_bits/config.hpp | 4 ++++ include/experimental/__p0009_bits/macros.hpp | 10 ++++++++++ 2 files changed, 14 insertions(+) diff --git a/include/experimental/__p0009_bits/config.hpp b/include/experimental/__p0009_bits/config.hpp index 727a565d..8a47b3c2 100644 --- a/include/experimental/__p0009_bits/config.hpp +++ b/include/experimental/__p0009_bits/config.hpp @@ -115,11 +115,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 +#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) \ diff --git a/include/experimental/__p0009_bits/macros.hpp b/include/experimental/__p0009_bits/macros.hpp index d6496916..469778e5 100644 --- a/include/experimental/__p0009_bits/macros.hpp +++ b/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_CTAD _MDSPAN_HOST_DEVICE +#else +# define MDSPAN_CTAD +#endif + // In CUDA defaulted functions do not need host device markup #ifndef MDSPAN_INLINE_FUNCTION_DEFAULTED # define MDSPAN_INLINE_FUNCTION_DEFAULTED From 8363859195aa751f0d941f4bc21b54de048ffd96 Mon Sep 17 00:00:00 2001 From: Nicolas Morales Date: Thu, 26 Jan 2023 17:05:43 -0800 Subject: [PATCH 05/13] remove some IDE-specific gitignores --- .gitignore | 3 --- 1 file changed, 3 deletions(-) diff --git a/.gitignore b/.gitignore index 3d826281..4e3e747f 100644 --- a/.gitignore +++ b/.gitignore @@ -4,6 +4,3 @@ CMakeUserPresets.json out build* Makefile -cmake-build-*/ -*.user -.idea/ From 1704ced66acdda89377c6f421022854617c1a431 Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 18:54:02 -0700 Subject: [PATCH 06/13] Extent test offload utils to support HIP --- tests/offload_utils.hpp | 31 +++++++++++++++++++++++++++---- 1 file changed, 27 insertions(+), 4 deletions(-) diff --git a/tests/offload_utils.hpp b/tests/offload_utils.hpp index 18f5b52a..971c0039 100644 --- a/tests/offload_utils.hpp +++ b/tests/offload_utils.hpp @@ -14,6 +14,13 @@ // //@HEADER +#ifdef _MDSPAN_HAS_HIP +#include +#include +#endif + +#include + namespace { bool dispatch_host = true; @@ -23,7 +30,23 @@ if (!(LHS == RHS)) { \ errors[0]++; \ } -#ifdef _MDSPAN_HAS_CUDA +#if defined(_MDSPAN_HAS_CUDA) || defined(_MDSPAN_HAS_HIP) + +#if defined(_MDSPAN_HAS_CUDA) +void deviceSynchronize() { (void) cudaDeviceSynchronize(); } +template +void mallocManaged(T** ptr, size_t size) { (void) cudaMallocManaged(ptr, size); } +template +void freeManaged(T* ptr) { (void) cudaFree(ptr); } +#endif + +#if defined(_MDSPAN_HAS_HIP) +void deviceSynchronize() { (void) hipDeviceSynchronize(); } +template +void mallocManaged(T** ptr, size_t size) { (void) hipMallocManaged(ptr, size); } +template +void freeManaged(T* ptr) { (void) hipFree(ptr); } +#endif template __global__ void dispatch_kernel(const LAMBDA f) { @@ -36,7 +59,7 @@ void dispatch(LAMBDA&& f) { static_cast(f)(); } else { dispatch_kernel<<<1,1>>>(static_cast(f)); - cudaDeviceSynchronize(); + deviceSynchronize(); } } @@ -46,7 +69,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; } @@ -55,7 +78,7 @@ void free_array(T* ptr) { if(dispatch_host == true) delete [] ptr; else - cudaFree(ptr); + freeManaged(ptr); } #define __MDSPAN_TESTS_RUN_TEST(A) \ From 022a295a81edc5a950e6383cbcadaeab665792c8 Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 18:54:34 -0700 Subject: [PATCH 07/13] Add CTAD host/device markup to mdspan --- include/experimental/__p0009_bits/mdspan.hpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/include/experimental/__p0009_bits/mdspan.hpp b/include/experimental/__p0009_bits/mdspan.hpp index 23bffe27..9a0808d2 100644 --- a/include/experimental/__p0009_bits/mdspan.hpp +++ b/include/experimental/__p0009_bits/mdspan.hpp @@ -369,28 +369,28 @@ MDSPAN_TEMPLATE_REQUIRES( /* requires */ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(is_integral, SizeTypes) /* && ... */) && (sizeof...(SizeTypes) > 0) ) -explicit mdspan(ElementType*, SizeTypes...) +MDSPAN_CTAD explicit mdspan(ElementType*, SizeTypes...) -> mdspan>; MDSPAN_TEMPLATE_REQUIRES( class Pointer, (_MDSPAN_TRAIT(is_pointer, std::remove_reference_t)) ) -mdspan(Pointer&&) -> mdspan>, extents>; +MDSPAN_CTAD mdspan(Pointer&&) -> mdspan>, extents>; MDSPAN_TEMPLATE_REQUIRES( class CArray, (_MDSPAN_TRAIT(is_array, CArray) && (rank_v == 1)) ) -mdspan(CArray&) -> mdspan, extents>>; +MDSPAN_CTAD mdspan(CArray&) -> mdspan, extents>>; template -mdspan(ElementType*, const ::std::array&) +MDSPAN_CTAD mdspan(ElementType*, const ::std::array&) -> mdspan>; #ifdef __cpp_lib_span template -mdspan(ElementType*, ::std::span) +MDSPAN_CTAD mdspan(ElementType*, ::std::span) -> mdspan>; #endif @@ -398,15 +398,15 @@ mdspan(ElementType*, ::std::span) // `ElementType*`s, and `data_handle_type` is taken from `accessor_type::data_handle_type`, which // seems to throw off automatic deduction guides. template -mdspan(ElementType*, const extents&) +MDSPAN_CTAD mdspan(ElementType*, const extents&) -> mdspan>; template -mdspan(ElementType*, const MappingType&) +MDSPAN_CTAD mdspan(ElementType*, const MappingType&) -> mdspan; template -mdspan(const typename AccessorType::data_handle_type, const MappingType&, const AccessorType&) +MDSPAN_CTAD mdspan(const typename AccessorType::data_handle_type, const MappingType&, const AccessorType&) -> mdspan; #endif From b2ba803bc4904bd95501405c1b1ffb697799389f Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 19:09:38 -0700 Subject: [PATCH 08/13] Fallback to manual implementation of mdspan swap for HIP Compiler errors out when calling std::swap. Note since it was constexpr function it was apparently automatically host device, but the tuple one isn't. --- include/experimental/__p0009_bits/mdspan.hpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/include/experimental/__p0009_bits/mdspan.hpp b/include/experimental/__p0009_bits/mdspan.hpp index 9a0808d2..c6449252 100644 --- a/include/experimental/__p0009_bits/mdspan.hpp +++ b/include/experimental/__p0009_bits/mdspan.hpp @@ -321,9 +321,22 @@ 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 + auto tmp_ptr = y.__ptr_ref(); + auto tmp_mapping = y.__mapping_ref(); + auto tmp_accessor = y.__accessor_ref(); + y.__ptr_ref() = x.__ptr_ref(); + y.__mapping_ref() = x.__mapping_ref(); + y.__accessor_ref() = x.__accessor_ref(); + x.__ptr_ref() = tmp_ptr; + x.__mapping_ref() = tmp_mapping; + x.__accessor_ref() = tmp_accessor; + #endif } //-------------------------------------------------------------------------------- From 95de8b82204fd6b41a4a99f3a9d154a9b62b049a Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 19:10:34 -0700 Subject: [PATCH 09/13] Work around HIP deduction guide markup issue for tuple --- .../experimental/__p2630_bits/submdspan_mapping.hpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/include/experimental/__p2630_bits/submdspan_mapping.hpp b/include/experimental/__p2630_bits/submdspan_mapping.hpp index 5565bfb9..fe930e07 100644 --- a/include/experimental/__p2630_bits/submdspan_mapping.hpp +++ b/include/experimental/__p2630_bits/submdspan_mapping.hpp @@ -125,7 +125,12 @@ submdspan_mapping(const layout_left::mapping &src_mapping, return mapping_offset{ 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{detail::stride_of(slices)...})), + #else tuple{detail::stride_of(slices)...})), + #endif static_cast(src_mapping(detail::first_of(slices)...))}; } #if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__) @@ -226,7 +231,12 @@ submdspan_mapping(const layout_right::mapping &src_mapping, return mapping_offset{ 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{detail::stride_of(slices)...})), + #else tuple{detail::stride_of(slices)...})), + #endif static_cast(src_mapping(detail::first_of(slices)...))}; } #if defined(__NVCC__) && !defined(__CUDA_ARCH__) && defined(__GNUC__) From 30bc07939a17bec58c666470a6293b94be6678f1 Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 19:21:42 -0700 Subject: [PATCH 10/13] Disable for now mdarray test for CUDA and HIP This test doesn't fully work due to using std::vector --- tests/CMakeLists.txt | 3 +++ 1 file changed, 3 insertions(+) diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index f7240b5b..afd2f3b1 100644 --- a/tests/CMakeLists.txt +++ b/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() From f7da6d62016f039e751d2b220397e6a9817c488d Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Thu, 26 Jan 2023 19:22:22 -0700 Subject: [PATCH 11/13] Simplify swap fallback impl and disable foo swap test with HIP --- include/experimental/__p0009_bits/mdspan.hpp | 12 +++--------- tests/test_mdspan_swap.cpp | 5 ++++- 2 files changed, 7 insertions(+), 10 deletions(-) diff --git a/include/experimental/__p0009_bits/mdspan.hpp b/include/experimental/__p0009_bits/mdspan.hpp index c6449252..b6f33d0f 100644 --- a/include/experimental/__p0009_bits/mdspan.hpp +++ b/include/experimental/__p0009_bits/mdspan.hpp @@ -327,15 +327,9 @@ class mdspan swap(x.__mapping_ref(), y.__mapping_ref()); swap(x.__accessor_ref(), y.__accessor_ref()); #else - auto tmp_ptr = y.__ptr_ref(); - auto tmp_mapping = y.__mapping_ref(); - auto tmp_accessor = y.__accessor_ref(); - y.__ptr_ref() = x.__ptr_ref(); - y.__mapping_ref() = x.__mapping_ref(); - y.__accessor_ref() = x.__accessor_ref(); - x.__ptr_ref() = tmp_ptr; - x.__mapping_ref() = tmp_mapping; - x.__accessor_ref() = tmp_accessor; + mdspan tmp = y; + y = x; + x = tmp; #endif } diff --git a/tests/test_mdspan_swap.cpp b/tests/test_mdspan_swap.cpp index 835a956e..4b4dfeb2 100644 --- a/tests/test_mdspan_swap.cpp +++ b/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 sideeffects will fail +#ifndef _MDSPAN_HAS_HIP void test_mdspan_foo_swap_dynamic_extents() { size_t* errors = allocate_array(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 From ec7a97037e8c95445fcb4a9ceb8c7ed94e0202c0 Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Mon, 30 Jan 2023 10:06:24 -0700 Subject: [PATCH 12/13] Rename macro for CTAD as suggested by review --- include/experimental/__p0009_bits/macros.hpp | 4 ++-- include/experimental/__p0009_bits/mdspan.hpp | 16 ++++++++-------- 2 files changed, 10 insertions(+), 10 deletions(-) diff --git a/include/experimental/__p0009_bits/macros.hpp b/include/experimental/__p0009_bits/macros.hpp index 469778e5..1f237fc7 100644 --- a/include/experimental/__p0009_bits/macros.hpp +++ b/include/experimental/__p0009_bits/macros.hpp @@ -45,9 +45,9 @@ #endif #ifdef _MDSPAN_HAS_HIP -# define MDSPAN_CTAD _MDSPAN_HOST_DEVICE +# define MDSPAN_DEDUCTION_GUIDE _MDSPAN_HOST_DEVICE #else -# define MDSPAN_CTAD +# define MDSPAN_DEDUCTION_GUIDE #endif // In CUDA defaulted functions do not need host device markup diff --git a/include/experimental/__p0009_bits/mdspan.hpp b/include/experimental/__p0009_bits/mdspan.hpp index b6f33d0f..928ef631 100644 --- a/include/experimental/__p0009_bits/mdspan.hpp +++ b/include/experimental/__p0009_bits/mdspan.hpp @@ -376,28 +376,28 @@ MDSPAN_TEMPLATE_REQUIRES( /* requires */ _MDSPAN_FOLD_AND(_MDSPAN_TRAIT(is_integral, SizeTypes) /* && ... */) && (sizeof...(SizeTypes) > 0) ) -MDSPAN_CTAD explicit mdspan(ElementType*, SizeTypes...) +MDSPAN_DEDUCTION_GUIDE explicit mdspan(ElementType*, SizeTypes...) -> mdspan>; MDSPAN_TEMPLATE_REQUIRES( class Pointer, (_MDSPAN_TRAIT(is_pointer, std::remove_reference_t)) ) -MDSPAN_CTAD mdspan(Pointer&&) -> mdspan>, extents>; +MDSPAN_DEDUCTION_GUIDE mdspan(Pointer&&) -> mdspan>, extents>; MDSPAN_TEMPLATE_REQUIRES( class CArray, (_MDSPAN_TRAIT(is_array, CArray) && (rank_v == 1)) ) -MDSPAN_CTAD mdspan(CArray&) -> mdspan, extents>>; +MDSPAN_DEDUCTION_GUIDE mdspan(CArray&) -> mdspan, extents>>; template -MDSPAN_CTAD mdspan(ElementType*, const ::std::array&) +MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const ::std::array&) -> mdspan>; #ifdef __cpp_lib_span template -MDSPAN_CTAD mdspan(ElementType*, ::std::span) +MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, ::std::span) -> mdspan>; #endif @@ -405,15 +405,15 @@ MDSPAN_CTAD mdspan(ElementType*, ::std::span) // `ElementType*`s, and `data_handle_type` is taken from `accessor_type::data_handle_type`, which // seems to throw off automatic deduction guides. template -MDSPAN_CTAD mdspan(ElementType*, const extents&) +MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const extents&) -> mdspan>; template -MDSPAN_CTAD mdspan(ElementType*, const MappingType&) +MDSPAN_DEDUCTION_GUIDE mdspan(ElementType*, const MappingType&) -> mdspan; template -MDSPAN_CTAD 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; #endif From f82e9efdf44a200634c2f0b6c1f2dae347260e38 Mon Sep 17 00:00:00 2001 From: Christian Trott Date: Mon, 30 Jan 2023 10:12:42 -0700 Subject: [PATCH 13/13] Fix typo --- tests/test_mdspan_swap.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/test_mdspan_swap.cpp b/tests/test_mdspan_swap.cpp index 4b4dfeb2..31538fea 100644 --- a/tests/test_mdspan_swap.cpp +++ b/tests/test_mdspan_swap.cpp @@ -103,7 +103,7 @@ TEST(TestMDSpanSwap, std_swap_dynamic_extents) { } // On HIP we actually don't call through to swap via ADL -// so the foo swap test which has sideeffects will fail +// 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(1);