From 2ab3dc331363eaf15cd09cf8897d198ba6ad0973 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 15 Jan 2025 10:54:19 +0100 Subject: [PATCH] NVRTC fixes --- thrust/thrust/detail/alignment.h | 7 +- .../detail/allocator_aware_execution_policy.h | 17 +- .../detail/execute_with_allocator_fwd.h | 8 +- thrust/thrust/detail/type_deduction.h | 22 +- thrust/thrust/iterator/counting_iterator.h | 2 +- .../thrust/iterator/detail/any_system_tag.h | 2 +- .../iterator/detail/discard_iterator_base.h | 4 +- .../detail/iterator_facade_category.h | 54 +++-- thrust/thrust/iterator/discard_iterator.h | 2 +- thrust/thrust/iterator/iterator_adaptor.h | 4 +- thrust/thrust/iterator/iterator_categories.h | 216 +++++++++++------- thrust/thrust/iterator/iterator_facade.h | 2 +- thrust/thrust/iterator/iterator_traits.h | 13 +- .../system/cpp/detail/execution_policy.h | 2 +- .../system/cuda/detail/execution_policy.h | 9 +- .../detail/sequential/execution_policy.h | 2 +- 16 files changed, 223 insertions(+), 143 deletions(-) diff --git a/thrust/thrust/detail/alignment.h b/thrust/thrust/detail/alignment.h index 9a43eeec226..56e10edd215 100644 --- a/thrust/thrust/detail/alignment.h +++ b/thrust/thrust/detail/alignment.h @@ -30,10 +30,9 @@ #endif // no system header #include +#include #include -#include // For `std::size_t` and `std::max_align_t`. - THRUST_NAMESPACE_BEGIN namespace detail { @@ -49,7 +48,7 @@ using alignment_of = ::cuda::std::alignment_of; /// type whose alignment requirement is a divisor of `Align`. /// /// The behavior is undefined if `Align` is not a power of 2. -template +template <::cuda::std::size_t Align> struct aligned_type { struct alignas(Align) type @@ -74,7 +73,7 @@ _CCCL_HOST_DEVICE T aligned_reinterpret_cast(U u) return reinterpret_cast(reinterpret_cast(u)); } -_CCCL_HOST_DEVICE inline std::size_t aligned_storage_size(std::size_t n, std::size_t align) +_CCCL_HOST_DEVICE inline ::cuda::std::size_t aligned_storage_size(::cuda::std::size_t n, ::cuda::std::size_t align) { return ::cuda::ceil_div(n, align) * align; } diff --git a/thrust/thrust/detail/allocator_aware_execution_policy.h b/thrust/thrust/detail/allocator_aware_execution_policy.h index c6589314709..0ca0c33dbff 100644 --- a/thrust/thrust/detail/allocator_aware_execution_policy.h +++ b/thrust/thrust/detail/allocator_aware_execution_policy.h @@ -28,7 +28,7 @@ #include #include -#include +#include THRUST_NAMESPACE_BEGIN @@ -61,19 +61,20 @@ struct allocator_aware_execution_policy }; template - typename execute_with_memory_resource_type::type operator()(MemoryResource* mem_res) const + _CCCL_HOST_DEVICE typename execute_with_memory_resource_type::type + operator()(MemoryResource* mem_res) const { return typename execute_with_memory_resource_type::type(mem_res); } template - typename execute_with_allocator_type::type operator()(Allocator& alloc) const + _CCCL_HOST_DEVICE typename execute_with_allocator_type::type operator()(Allocator& alloc) const { return typename execute_with_allocator_type::type(alloc); } template - typename execute_with_allocator_type::type operator()(const Allocator& alloc) const + _CCCL_HOST_DEVICE typename execute_with_allocator_type::type operator()(const Allocator& alloc) const { return typename execute_with_allocator_type::type(alloc); } @@ -81,10 +82,12 @@ struct allocator_aware_execution_policy // just the rvalue overload // perfect forwarding doesn't help, because a const reference has to be turned // into a value by copying for the purpose of storing it in execute_with_allocator - template ::value>::type* = nullptr> - typename execute_with_allocator_type::type operator()(Allocator&& alloc) const + _CCCL_EXEC_CHECK_DISABLE + template ::value>::type* = nullptr> + _CCCL_HOST_DEVICE typename execute_with_allocator_type::type operator()(Allocator&& alloc) const { - return typename execute_with_allocator_type::type(std::move(alloc)); + return typename execute_with_allocator_type::type(::cuda::std::move(alloc)); } }; diff --git a/thrust/thrust/detail/execute_with_allocator_fwd.h b/thrust/thrust/detail/execute_with_allocator_fwd.h index 54fc6e55fa8..9b98b886188 100644 --- a/thrust/thrust/detail/execute_with_allocator_fwd.h +++ b/thrust/thrust/detail/execute_with_allocator_fwd.h @@ -26,7 +26,9 @@ # pragma system_header #endif // no system header -#include +#if !_CCCL_COMPILER(NVRTC) +# include +#endif // !_CCCL_COMPILER(NVRTC) #include THRUST_NAMESPACE_BEGIN @@ -53,11 +55,12 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH // because of execute_with_allocator_and_dependen : alloc(alloc_) {} - ::cuda::std::remove_reference_t& get_allocator() + _CCCL_HOST_DEVICE ::cuda::std::remove_reference_t& get_allocator() { return alloc; } +#if !_CCCL_COMPILER(NVRTC) template CCCL_DEPRECATED _CCCL_HOST execute_with_allocator_and_dependencies after(Dependencies&&... dependencies) const @@ -97,6 +100,7 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH // because of execute_with_allocator_and_dependen { return {alloc, capture_as_dependency(std::move(dependencies))}; } +#endif // !_CCCL_COMPILER(NVRTC) }; _CCCL_SUPPRESS_DEPRECATED_POP diff --git a/thrust/thrust/detail/type_deduction.h b/thrust/thrust/detail/type_deduction.h index 717edc366ee..8127940f777 100644 --- a/thrust/thrust/detail/type_deduction.h +++ b/thrust/thrust/detail/type_deduction.h @@ -18,20 +18,20 @@ #endif // no system header #include -#include -#include +#include +#include /////////////////////////////////////////////////////////////////////////////// /// \def THRUST_FWD(x) /// \brief Performs universal forwarding of a universal reference. /// -#define THRUST_FWD(x) ::std::forward(x) +#define THRUST_FWD(x) ::cuda::std::forward(x) /// \def THRUST_MVCAP(x) /// \brief Capture `x` into a lambda by moving. /// deprecated [Since 2.8] -#define THRUST_MVCAP(x) x = ::std::move(x) +#define THRUST_MVCAP(x) x = ::cuda::std::move(x) /// \def THRUST_RETOF(invocable, ...) /// \brief Expands to the type returned by invoking an instance of the invocable @@ -40,9 +40,9 @@ /// deprecated [Since 2.8] #define THRUST_RETOF(...) THRUST_PP_DISPATCH(THRUST_RETOF, __VA_ARGS__) /// deprecated [Since 2.8] -#define THRUST_RETOF1(C) decltype(::std::declval()()) +#define THRUST_RETOF1(C) decltype(::cuda::std::declval()()) /// deprecated [Since 2.8] -#define THRUST_RETOF2(C, V) decltype(::std::declval()(::std::declval())) +#define THRUST_RETOF2(C, V) decltype(::cuda::std::declval()(::cuda::std::declval())) /// \def THRUST_RETURNS(...) /// \brief Expands to a function definition that returns the expression @@ -91,11 +91,11 @@ /**/ #else /// deprecated [Since 2.8] -# define THRUST_DECLTYPE_RETURNS_WITH_SFINAE_CONDITION(condition, ...) \ - noexcept(noexcept(__VA_ARGS__))->typename std::enable_if::type \ - { \ - return (__VA_ARGS__); \ - } \ +# define THRUST_DECLTYPE_RETURNS_WITH_SFINAE_CONDITION(condition, ...) \ + noexcept(noexcept(__VA_ARGS__))->::cuda::std::enable_if_t \ + { \ + return (__VA_ARGS__); \ + } \ /**/ #endif diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h index bb98bc3b4d5..fe4ae3061dc 100644 --- a/thrust/thrust/iterator/counting_iterator.h +++ b/thrust/thrust/iterator/counting_iterator.h @@ -142,7 +142,7 @@ class _CCCL_DECLSPEC_EMPTY_BASES counting_iterator */ using super_t = typename detail::counting_iterator_base::type; - friend class thrust::iterator_core_access; + friend class iterator_core_access; public: using reference = typename super_t::reference; diff --git a/thrust/thrust/iterator/detail/any_system_tag.h b/thrust/thrust/iterator/detail/any_system_tag.h index 7b3172bfef5..9aa9827e40c 100644 --- a/thrust/thrust/iterator/detail/any_system_tag.h +++ b/thrust/thrust/iterator/detail/any_system_tag.h @@ -34,7 +34,7 @@ struct any_system_tag : thrust::execution_policy // allow any_system_tag to convert to any type at all // XXX make this safer using enable_if> upon c++11 template - operator T() const + _CCCL_HOST_DEVICE operator T() const { return T(); } diff --git a/thrust/thrust/iterator/detail/discard_iterator_base.h b/thrust/thrust/iterator/detail/discard_iterator_base.h index 7f4e637bbc4..8191ae1ba8d 100644 --- a/thrust/thrust/iterator/detail/discard_iterator_base.h +++ b/thrust/thrust/iterator/detail/discard_iterator_base.h @@ -29,7 +29,7 @@ #include #include -#include // for std::ptrdiff_t +#include THRUST_NAMESPACE_BEGIN @@ -47,7 +47,7 @@ struct discard_iterator_base // but this interferes with zip_iterator using value_type = any_assign; using reference = any_assign&; - using incrementable = std::ptrdiff_t; + using incrementable = ::cuda::std::ptrdiff_t; using base_iterator = typename thrust::counting_iterator; diff --git a/thrust/thrust/iterator/detail/iterator_facade_category.h b/thrust/thrust/iterator/detail/iterator_facade_category.h index 02dd93e914e..797e8ae6dcd 100644 --- a/thrust/thrust/iterator/detail/iterator_facade_category.h +++ b/thrust/thrust/iterator/detail/iterator_facade_category.h @@ -94,6 +94,7 @@ struct iterator_facade_default_category; // Instead, it simply assumes that if is_convertible, // then the category is input_iterator_tag +#if !_CCCL_COMPILER(NVRTC) // this is the function for standard system iterators template struct iterator_facade_default_category_std @@ -110,6 +111,7 @@ struct iterator_facade_default_category_std thrust::detail::identity_, thrust::detail::identity_>> {}; // end iterator_facade_default_category_std +#endif // !_CCCL_COMPILER(NVRTC) // this is the function for host system iterators template @@ -127,7 +129,6 @@ struct iterator_facade_default_category_host thrust::detail::identity_, thrust::detail::identity_>> {}; // end iterator_facade_default_category_host - // this is the function for device system iterators template struct iterator_facade_default_category_device @@ -146,6 +147,7 @@ struct iterator_facade_default_category_device thrust::detail::identity_>> {}; // end iterator_facade_default_category_device +#if !_CCCL_COMPILER(NVRTC) // this is the function for any system iterators template struct iterator_facade_default_category_any @@ -155,29 +157,41 @@ struct iterator_facade_default_category_any thrust::any_system_tag, Traversal>; }; // end iterator_facade_default_category_any +#endif // !_CCCL_COMPILER(NVRTC) template struct iterator_facade_default_category + : +#if !_CCCL_COMPILER(NVRTC) // check for any system - : thrust::detail::eval_if< - ::cuda::std::is_convertible::value, - iterator_facade_default_category_any, - - // check for host system - thrust::detail::eval_if< - ::cuda::std::is_convertible::value, - iterator_facade_default_category_host, - - // check for device system - thrust::detail::eval_if<::cuda::std::is_convertible::value, - iterator_facade_default_category_device, - - // if we don't recognize the system, get a standard iterator category - // and combine it with System & Traversal - thrust::detail::identity_::type, - System, - Traversal>>>>> + thrust::detail::eval_if< + ::cuda::std::is_convertible::value, + iterator_facade_default_category_any, +#endif // !_CCCL_COMPILER(NVRTC) + // check for host system + thrust::detail::eval_if< + ::cuda::std::is_convertible::value, + iterator_facade_default_category_host, + + // check for device system + thrust::detail::eval_if<::cuda::std::is_convertible::value, + iterator_facade_default_category_device, + + // if we don't recognize the system, get a standard iterator category + // and combine it with System & Traversal (for NVRTC fall back to host categories) + thrust::detail::identity_::type, + System, + Traversal>>>> +#if !_CCCL_COMPILER(NVRTC) + > +#endif // !_CCCL_COMPILER(NVRTC) {}; template diff --git a/thrust/thrust/iterator/discard_iterator.h b/thrust/thrust/iterator/discard_iterator.h index 5b64aabd5a3..e7f2754bba8 100644 --- a/thrust/thrust/iterator/discard_iterator.h +++ b/thrust/thrust/iterator/discard_iterator.h @@ -102,7 +102,7 @@ class discard_iterator : public detail::discard_iterator_base::type { /*! \cond */ - friend class thrust::iterator_core_access; + friend class iterator_core_access; using super_t = typename detail::discard_iterator_base::type; using incrementable = typename detail::discard_iterator_base::incrementable; using base_iterator = typename detail::discard_iterator_base::base_iterator; diff --git a/thrust/thrust/iterator/iterator_adaptor.h b/thrust/thrust/iterator/iterator_adaptor.h index a427b628e11..6987c197caf 100644 --- a/thrust/thrust/iterator/iterator_adaptor.h +++ b/thrust/thrust/iterator/iterator_adaptor.h @@ -55,6 +55,8 @@ THRUST_NAMESPACE_BEGIN * \{ */ +class iterator_core_access; + /*! \p iterator_adaptor is an iterator which adapts an existing type of iterator to create a new type of * iterator. Most of Thrust's fancy iterators are defined via inheritance from \p iterator_adaptor. * While composition of these existing Thrust iterators is often sufficient for expressing the desired @@ -132,7 +134,7 @@ class _CCCL_DECLSPEC_EMPTY_BASES iterator_adaptor /*! \cond */ - friend class thrust::iterator_core_access; + friend class iterator_core_access; protected: using super_t = diff --git a/thrust/thrust/iterator/iterator_categories.h b/thrust/thrust/iterator/iterator_categories.h index 5fbe90ab2c3..1749879f82e 100644 --- a/thrust/thrust/iterator/iterator_categories.h +++ b/thrust/thrust/iterator/iterator_categories.h @@ -43,7 +43,11 @@ #include // #include this for stl's iterator tags -#include +#if _CCCL_COMPILER(NVRTC) +# include +#else // _CCCL_COMPILER(NVRTC) +# include +#endif // _CCCL_COMPILER(NVRTC) THRUST_NAMESPACE_BEGIN @@ -55,6 +59,96 @@ THRUST_NAMESPACE_BEGIN * \{ */ +/*! \p input_host_iterator_tag is an empty class: it has no member + * functions, member variables, or nested types. It is used solely as a "tag": a + * representation of the Input Host Iterator concept within the C++ + * type system. + * + * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags + * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, + * forward_device_iterator_tag, bidirectional_device_iterator_tag, + * random_access_device_iterator_tag, + * output_host_iterator_tag, forward_host_iterator_tag, + * bidirectional_host_iterator_tag, random_access_host_iterator_tag + */ +using input_host_iterator_tag = +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::input_iterator_tag; + +/*! \p output_host_iterator_tag is an empty class: it has no member + * functions, member variables, or nested types. It is used solely as a "tag": a + * representation of the Output Host Iterator concept within the C++ + * type system. + * + * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags + * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, + * forward_device_iterator_tag, bidirectional_device_iterator_tag, + * random_access_device_iterator_tag, + * input_host_iterator_tag, forward_host_iterator_tag, + * bidirectional_host_iterator_tag, random_access_host_iterator_tag + */ +using output_host_iterator_tag = +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::output_iterator_tag; + +/*! \p forward_host_iterator_tag is an empty class: it has no member + * functions, member variables, or nested types. It is used solely as a "tag": a + * representation of the Forward Host Iterator concept within the C++ + * type system. + * + * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags + * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, + * forward_device_iterator_tag, bidirectional_device_iterator_tag, + * random_access_device_iterator_tag, + * input_host_iterator_tag, output_host_iterator_tag, + * bidirectional_host_iterator_tag, random_access_host_iterator_tag + */ +using forward_host_iterator_tag = +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::forward_iterator_tag; + +/*! \p bidirectional_host_iterator_tag is an empty class: it has no member + * functions, member variables, or nested types. It is used solely as a "tag": a + * representation of the Forward Host Iterator concept within the C++ + * type system. + * + * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags + * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, + * forward_device_iterator_tag, bidirectional_device_iterator_tag, + * random_access_device_iterator_tag, + * input_host_iterator_tag, output_host_iterator_tag, + * forward_host_iterator_tag, random_access_host_iterator_tag + */ +using bidirectional_host_iterator_tag = +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::bidirectional_iterator_tag; + +/*! \p random_access_host_iterator_tag is an empty class: it has no member + * functions, member variables, or nested types. It is used solely as a "tag": a + * representation of the Forward Host Iterator concept within the C++ + * type system. + * + * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags + * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, + * forward_device_iterator_tag, bidirectional_device_iterator_tag, + * random_access_device_iterator_tag, + * input_host_iterator_tag, output_host_iterator_tag, + * forward_host_iterator_tag, bidirectional_host_iterator_tag + */ +using random_access_host_iterator_tag = +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::random_access_iterator_tag; + /*! \p input_device_iterator_tag is an empty class: it has no member functions, * member variables, or nested types. It is used solely as a "tag": a * representation of the Input Device Iterator concept within the C++ type @@ -67,9 +161,13 @@ THRUST_NAMESPACE_BEGIN * bidirectional_host_iterator_tag, random_access_host_iterator_tag */ struct input_device_iterator_tag - : thrust::detail::iterator_category_with_system_and_traversal + : thrust::detail::iterator_category_with_system_and_traversal< +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::input_iterator_tag, + thrust::device_system_tag, + thrust::single_pass_traversal_tag> {}; /*! \p output_device_iterator_tag is an empty class: it has no member functions, @@ -84,9 +182,13 @@ struct input_device_iterator_tag * bidirectional_host_iterator_tag, random_access_host_iterator_tag */ struct output_device_iterator_tag - : thrust::detail::iterator_category_with_system_and_traversal + : thrust::detail::iterator_category_with_system_and_traversal< +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::output_iterator_tag, + thrust::device_system_tag, + thrust::single_pass_traversal_tag> {}; /*! \p forward_device_iterator_tag is an empty class: it has no member functions, @@ -101,9 +203,13 @@ struct output_device_iterator_tag * bidirectional_host_iterator_tag, random_access_host_iterator_tag */ struct forward_device_iterator_tag - : thrust::detail::iterator_category_with_system_and_traversal + : thrust::detail::iterator_category_with_system_and_traversal< +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::forward_iterator_tag, + thrust::device_system_tag, + thrust::forward_traversal_tag> {}; /*! \p bidirectional_device_iterator_tag is an empty class: it has no member @@ -118,9 +224,13 @@ struct forward_device_iterator_tag * bidirectional_host_iterator_tag, random_access_host_iterator_tag */ struct bidirectional_device_iterator_tag - : thrust::detail::iterator_category_with_system_and_traversal + : thrust::detail::iterator_category_with_system_and_traversal< +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::bidirectional_iterator_tag, + thrust::device_system_tag, + thrust::bidirectional_traversal_tag> {}; /*! \p random_access_device_iterator_tag is an empty class: it has no member @@ -135,81 +245,15 @@ struct bidirectional_device_iterator_tag * bidirectional_host_iterator_tag, random_access_host_iterator_tag */ struct random_access_device_iterator_tag - : thrust::detail::iterator_category_with_system_and_traversal + : thrust::detail::iterator_category_with_system_and_traversal< +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::random_access_iterator_tag, + thrust::device_system_tag, + thrust::random_access_traversal_tag> {}; -/*! \p input_host_iterator_tag is an empty class: it has no member - * functions, member variables, or nested types. It is used solely as a "tag": a - * representation of the Input Host Iterator concept within the C++ - * type system. - * - * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags - * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, - * forward_device_iterator_tag, bidirectional_device_iterator_tag, - * random_access_device_iterator_tag, - * output_host_iterator_tag, forward_host_iterator_tag, - * bidirectional_host_iterator_tag, random_access_host_iterator_tag - */ -using input_host_iterator_tag = std::input_iterator_tag; - -/*! \p output_host_iterator_tag is an empty class: it has no member - * functions, member variables, or nested types. It is used solely as a "tag": a - * representation of the Output Host Iterator concept within the C++ - * type system. - * - * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags - * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, - * forward_device_iterator_tag, bidirectional_device_iterator_tag, - * random_access_device_iterator_tag, - * input_host_iterator_tag, forward_host_iterator_tag, - * bidirectional_host_iterator_tag, random_access_host_iterator_tag - */ -using output_host_iterator_tag = std::output_iterator_tag; - -/*! \p forward_host_iterator_tag is an empty class: it has no member - * functions, member variables, or nested types. It is used solely as a "tag": a - * representation of the Forward Host Iterator concept within the C++ - * type system. - * - * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags - * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, - * forward_device_iterator_tag, bidirectional_device_iterator_tag, - * random_access_device_iterator_tag, - * input_host_iterator_tag, output_host_iterator_tag, - * bidirectional_host_iterator_tag, random_access_host_iterator_tag - */ -using forward_host_iterator_tag = std::forward_iterator_tag; - -/*! \p bidirectional_host_iterator_tag is an empty class: it has no member - * functions, member variables, or nested types. It is used solely as a "tag": a - * representation of the Forward Host Iterator concept within the C++ - * type system. - * - * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags - * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, - * forward_device_iterator_tag, bidirectional_device_iterator_tag, - * random_access_device_iterator_tag, - * input_host_iterator_tag, output_host_iterator_tag, - * forward_host_iterator_tag, random_access_host_iterator_tag - */ -using bidirectional_host_iterator_tag = std::bidirectional_iterator_tag; - -/*! \p random_access_host_iterator_tag is an empty class: it has no member - * functions, member variables, or nested types. It is used solely as a "tag": a - * representation of the Forward Host Iterator concept within the C++ - * type system. - * - * \see https://en.cppreference.com/w/cpp/iterator/iterator_tags - * iterator_traits, input_device_iterator_tag, output_device_iterator_tag, - * forward_device_iterator_tag, bidirectional_device_iterator_tag, - * random_access_device_iterator_tag, - * input_host_iterator_tag, output_host_iterator_tag, - * forward_host_iterator_tag, bidirectional_host_iterator_tag - */ -using random_access_host_iterator_tag = std::random_access_iterator_tag; - /*! \} // end iterator_tag_classes */ diff --git a/thrust/thrust/iterator/iterator_facade.h b/thrust/thrust/iterator/iterator_facade.h index 7ce354a5b3a..5d26c8eaeb3 100644 --- a/thrust/thrust/iterator/iterator_facade.h +++ b/thrust/thrust/iterator/iterator_facade.h @@ -306,7 +306,7 @@ template + typename Difference = ::cuda::std::ptrdiff_t> class iterator_facade { private: diff --git a/thrust/thrust/iterator/iterator_traits.h b/thrust/thrust/iterator/iterator_traits.h index 4001410613b..e087982496d 100644 --- a/thrust/thrust/iterator/iterator_traits.h +++ b/thrust/thrust/iterator/iterator_traits.h @@ -39,7 +39,11 @@ # pragma system_header #endif // no system header -#include +#if _CCCL_COMPILER(NVRTC) +# include +#else // _CCCL_COMPILER(NVRTC) +# include +#endif // _CCCL_COMPILER(NVRTC) THRUST_NAMESPACE_BEGIN @@ -47,7 +51,12 @@ THRUST_NAMESPACE_BEGIN * interface for querying the properties of iterators at compile-time. */ template -struct iterator_traits : std::iterator_traits +struct iterator_traits + : +#if _CCCL_COMPILER(NVRTC) + ::cuda:: +#endif // _CCCL_COMPILER(NVRTC) + std::iterator_traits {}; template diff --git a/thrust/thrust/system/cpp/detail/execution_policy.h b/thrust/thrust/system/cpp/detail/execution_policy.h index e41517c7137..f8304faddcc 100644 --- a/thrust/thrust/system/cpp/detail/execution_policy.h +++ b/thrust/thrust/system/cpp/detail/execution_policy.h @@ -64,7 +64,7 @@ template struct execution_policy : thrust::system::detail::sequential::execution_policy { using tag_type = tag; - operator tag() const + _CCCL_HOST_DEVICE operator tag() const { return tag(); } diff --git a/thrust/thrust/system/cuda/detail/execution_policy.h b/thrust/thrust/system/cuda/detail/execution_policy.h index cd82a4dceb8..8674f509179 100644 --- a/thrust/thrust/system/cuda/detail/execution_policy.h +++ b/thrust/thrust/system/cuda/detail/execution_policy.h @@ -40,7 +40,10 @@ #include #include -#include +#if !_CCCL_COMPILER(NVRTC) +# include +#endif // !_CCCL_COMPILER(NVRTC) + #include #include #include @@ -65,7 +68,9 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH struct tag : execution_policy , thrust::detail::allocator_aware_execution_policy +#if !_CCCL_COMPILER(NVRTC) , thrust::detail::dependencies_aware_execution_policy +#endif // !_CCCL_COMPILER(NVRTC) {}; _CCCL_SUPPRESS_DEPRECATED_POP @@ -73,7 +78,7 @@ template struct execution_policy : thrust::execution_policy { using tag_type = tag; - operator tag() const + _CCCL_HOST_DEVICE operator tag() const { return tag(); } diff --git a/thrust/thrust/system/detail/sequential/execution_policy.h b/thrust/thrust/system/detail/sequential/execution_policy.h index d4834ebc89b..ce7dba69d2f 100644 --- a/thrust/thrust/system/detail/sequential/execution_policy.h +++ b/thrust/thrust/system/detail/sequential/execution_policy.h @@ -64,7 +64,7 @@ template struct execution_policy : thrust::execution_policy { // allow conversion to tag - inline operator tag() const + _CCCL_HOST_DEVICE inline operator tag() const { return tag(); }