From bc9a88652dc78ec59f80c0a0404a6bc8b0e9c33d Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 28 Feb 2025 21:08:49 +0100 Subject: [PATCH 01/10] Add tests for raw_reference_cast --- thrust/testing/raw_reference_cast.cu | 48 ++++++++++++++++++++++++++++ 1 file changed, 48 insertions(+) create mode 100644 thrust/testing/raw_reference_cast.cu diff --git a/thrust/testing/raw_reference_cast.cu b/thrust/testing/raw_reference_cast.cu new file mode 100644 index 00000000000..b5a4400d2ae --- /dev/null +++ b/thrust/testing/raw_reference_cast.cu @@ -0,0 +1,48 @@ +#include +#include +#include + +#include + +void TestRawReferenceCast() +{ + using ::cuda::std::is_same_v; + + { + [[maybe_unused]] int i = 42; + [[maybe_unused]] const int ci = 42; + static_assert(is_same_v); + static_assert(is_same_v); + } + { + [[maybe_unused]] thrust::host_vector vec(1); + static_assert(is_same_v); + static_assert(is_same_v); + + [[maybe_unused]] auto zip = thrust::make_zip_iterator(vec.begin(), vec.begin()); + static_assert( + is_same_v>); + + [[maybe_unused]] auto zip2 = thrust::make_zip_iterator(zip, zip); + static_assert( + is_same_v, + thrust::detail::tuple_of_iterator_references>>); + } + { + [[maybe_unused]] thrust::device_vector vec(1); + static_assert(is_same_v); + static_assert(is_same_v); + + [[maybe_unused]] auto zip = thrust::make_zip_iterator(vec.begin(), vec.begin()); + static_assert( + is_same_v>); + + [[maybe_unused]] auto zip2 = thrust::make_zip_iterator(zip, zip); + static_assert( + is_same_v, + thrust::detail::tuple_of_iterator_references>>); + } +} +DECLARE_UNITTEST(TestRawReferenceCast); From 53910e23ef020183aa523c3e1a8bef1ac1578efd Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 28 Feb 2025 21:08:52 +0100 Subject: [PATCH 02/10] Fix raw_reference_cast for tuple_of_iterator_references --- thrust/thrust/detail/raw_reference_cast.h | 41 ++++++++++------------- 1 file changed, 18 insertions(+), 23 deletions(-) diff --git a/thrust/thrust/detail/raw_reference_cast.h b/thrust/thrust/detail/raw_reference_cast.h index d49585a30e8..ce958335e26 100644 --- a/thrust/thrust/detail/raw_reference_cast.h +++ b/thrust/thrust/detail/raw_reference_cast.h @@ -64,10 +64,6 @@ struct is_unwrappable> : ::cuda::std::disjunction...> {}; -template -struct enable_if_unwrappable : ::cuda::std::enable_if::value, Result> -{}; - namespace raw_reference_detail { @@ -141,13 +137,9 @@ struct raw_reference> }; template -struct raw_reference> +struct raw_reference> { -private: - using tuple_type = detail::tuple_of_iterator_references; - -public: - using type = typename raw_reference_detail::raw_reference_tuple_helper::type; + using type = typename raw_reference_detail::raw_reference_tuple_helper>::type; }; } // namespace detail @@ -160,10 +152,8 @@ template _CCCL_HOST_DEVICE typename detail::raw_reference::type raw_reference_cast(const T& ref); template -_CCCL_HOST_DEVICE typename detail::enable_if_unwrappable< - thrust::detail::tuple_of_iterator_references, - typename detail::raw_reference>::type>::type -raw_reference_cast(thrust::detail::tuple_of_iterator_references t); +_CCCL_HOST_DEVICE auto raw_reference_cast(detail::tuple_of_iterator_references t) -> + typename detail::raw_reference>::type; namespace detail { @@ -206,16 +196,21 @@ _CCCL_HOST_DEVICE typename detail::raw_reference::type raw_reference_ca } // end raw_reference_cast template -_CCCL_HOST_DEVICE typename detail::enable_if_unwrappable< - thrust::detail::tuple_of_iterator_references, - typename detail::raw_reference>::type>::type -raw_reference_cast(thrust::detail::tuple_of_iterator_references t) +_CCCL_HOST_DEVICE auto raw_reference_cast(detail::tuple_of_iterator_references t) -> + typename detail::raw_reference>::type { - thrust::detail::raw_reference_caster f; + if constexpr (detail::is_unwrappable>::value) + { + thrust::detail::raw_reference_caster f; - // note that we pass raw_reference_tuple_helper, not raw_reference as the unary metafunction - // the different way that raw_reference_tuple_helper unwraps tuples is important - return thrust::detail::tuple_host_device_transform(t, f); -} // end raw_reference_cast + // note that we pass raw_reference_tuple_helper, not raw_reference as the unary metafunction + // the different way that raw_reference_tuple_helper unwraps tuples is important + return thrust::detail::tuple_host_device_transform(t, f); + } + else + { + return t; + } +} THRUST_NAMESPACE_END From 8901ea2607f4c9199567b72f453140e69af2adfa Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 28 Feb 2025 22:19:40 +0100 Subject: [PATCH 03/10] Simplify thrust::generate --- thrust/thrust/detail/internal_functional.h | 73 ------------------- .../thrust/system/detail/generic/generate.inl | 66 +++++------------ 2 files changed, 19 insertions(+), 120 deletions(-) diff --git a/thrust/thrust/detail/internal_functional.h b/thrust/thrust/detail/internal_functional.h index 346f4a6cd05..fda0fa6838d 100644 --- a/thrust/thrust/detail/internal_functional.h +++ b/thrust/thrust/detail/internal_functional.h @@ -112,79 +112,6 @@ struct tuple_not_binary_predicate mutable Predicate pred; }; -template -struct host_generate_functor -{ - using result_type = void; - - _CCCL_EXEC_CHECK_DISABLE - _CCCL_HOST_DEVICE host_generate_functor(Generator g) - : gen(g) - {} - - // operator() does not take an lvalue reference because some iterators - // produce temporary proxy references when dereferenced. for example, - // consider the temporary tuple of references produced by zip_iterator. - // such temporaries cannot bind to an lvalue reference. - // - // to WAR this, accept a const reference (which is bindable to a temporary), - // and const_cast in the implementation. - // - // XXX change to an rvalue reference upon c++0x (which either a named variable - // or temporary can bind to) - template - _CCCL_HOST void operator()(const T& x) - { - // we have to be naughty and const_cast this to get it to work - T& lvalue = const_cast(x); - - // this assigns correctly whether x is a true reference or proxy - lvalue = gen(); - } - - Generator gen; -}; - -template -struct device_generate_functor -{ - using result_type = void; - - _CCCL_EXEC_CHECK_DISABLE - _CCCL_HOST_DEVICE device_generate_functor(Generator g) - : gen(g) - {} - - // operator() does not take an lvalue reference because some iterators - // produce temporary proxy references when dereferenced. for example, - // consider the temporary tuple of references produced by zip_iterator. - // such temporaries cannot bind to an lvalue reference. - // - // to WAR this, accept a const reference (which is bindable to a temporary), - // and const_cast in the implementation. - // - // XXX change to an rvalue reference upon c++0x (which either a named variable - // or temporary can bind to) - template - _CCCL_HOST_DEVICE void operator()(const T& x) - { - // we have to be naughty and const_cast this to get it to work - T& lvalue = const_cast(x); - - // this assigns correctly whether x is a true reference or proxy - lvalue = gen(); - } - - Generator gen; -}; - -template -struct generate_functor - : thrust::detail::eval_if<::cuda::std::is_convertible::value, - thrust::detail::identity_>, - thrust::detail::identity_>> -{}; - template struct is_non_const_reference : ::cuda::std::_And>, diff --git a/thrust/thrust/system/detail/generic/generate.inl b/thrust/thrust/system/detail/generic/generate.inl index 7705980b2e3..23d0700d810 100644 --- a/thrust/thrust/system/detail/generic/generate.inl +++ b/thrust/thrust/system/detail/generic/generate.inl @@ -25,64 +25,36 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header -#include -#include + #include -#include -#include THRUST_NAMESPACE_BEGIN -namespace system -{ -namespace detail +namespace system::detail::generic { -namespace generic +template +struct generate_functor { + template + _CCCL_HOST_DEVICE void operator()(T&& x) + { + ::cuda::std::forward(x) = gen(); + } + + Generator gen; +}; template _CCCL_HOST_DEVICE void -generate(thrust::execution_policy& exec, ForwardIterator first, ForwardIterator last, Generator gen) +generate(execution_policy& exec, ForwardIterator first, ForwardIterator last, Generator gen) { - // this static assert is necessary due to a workaround in generate_functor - // it takes a const reference to accept temporaries from proxy iterators - // and then const_casts the constness away - // - // this had the weird side effect of allowing generate (and fill, and whatever - // else is implemented in terms of generate) to fill through const iterators. - // this might become unnecessary once Thrust is C++11-and-above only, since the - // other solution is to take an rvalue reference in a second overload of - // operator() of the function object, but until we support pre-11, this is a - // nice solution that validates the const_cast and doesn't take away any - // functionality. - THRUST_STATIC_ASSERT_MSG( - !::cuda::std::is_const<::cuda::std::remove_reference_t>>::value, - "generating to `const` iterators is not allowed"); - thrust::for_each(exec, first, last, typename thrust::detail::generate_functor::type(gen)); -} // end generate() + thrust::for_each(exec, first, last, generate_functor{::cuda::std::move(gen)}); +} template _CCCL_HOST_DEVICE OutputIterator -generate_n(thrust::execution_policy& exec, OutputIterator first, Size n, Generator gen) +generate_n(execution_policy& exec, OutputIterator first, Size n, Generator gen) { - // this static assert is necessary due to a workaround in generate_functor - // it takes a const reference to accept temporaries from proxy iterators - // and then const_casts the constness away - // - // this had the weird side effect of allowing generate (and fill, and whatever - // else is implemented in terms of generate) to fill through const iterators. - // this might become unnecessary once Thrust is C++11-and-above only, since the - // other solution is to take an rvalue reference in a second overload of - // operator() of the function object, but until we support pre-11, this is a - // nice solution that validates the const_cast and doesn't take away any - // functionality. - THRUST_STATIC_ASSERT_MSG( - !::cuda::std::is_const<::cuda::std::remove_reference_t>>::value, - "generating to `const` iterators is not allowed"); - return thrust::for_each_n( - exec, first, n, typename thrust::detail::generate_functor::type(gen)); -} // end generate() - -} // end namespace generic -} // end namespace detail -} // end namespace system + return thrust::for_each_n(exec, first, n, generate_functor{::cuda::std::move(gen)}); +} +} // namespace system::detail::generic THRUST_NAMESPACE_END From 79e8713977479203e747c672f5754c7a1a6237b8 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 28 Feb 2025 23:11:10 +0100 Subject: [PATCH 04/10] Missing header --- thrust/thrust/system/detail/generic/generate.inl | 2 ++ 1 file changed, 2 insertions(+) diff --git a/thrust/thrust/system/detail/generic/generate.inl b/thrust/thrust/system/detail/generic/generate.inl index 23d0700d810..1364a8b7d53 100644 --- a/thrust/thrust/system/detail/generic/generate.inl +++ b/thrust/thrust/system/detail/generic/generate.inl @@ -28,6 +28,8 @@ #include +#include + THRUST_NAMESPACE_BEGIN namespace system::detail::generic { From 9ffde2da7068ee7bcd1042cc3c1488fafa205aef Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 28 Feb 2025 23:11:37 +0100 Subject: [PATCH 05/10] Drop obsolete test --- thrust/testing/generate_const_iterators.cu | 30 ---------------------- 1 file changed, 30 deletions(-) delete mode 100644 thrust/testing/generate_const_iterators.cu diff --git a/thrust/testing/generate_const_iterators.cu b/thrust/testing/generate_const_iterators.cu deleted file mode 100644 index b26af0f38a8..00000000000 --- a/thrust/testing/generate_const_iterators.cu +++ /dev/null @@ -1,30 +0,0 @@ - -#include -#include - -// The runtime_static_assert header needs to come first as we are overwriting thrusts internal static assert -#include - -struct generator -{ - _CCCL_HOST_DEVICE int operator()() const - { - return 1; - } -}; - -void TestGenerateConstIteratorCompilationError() -{ - thrust::host_vector test1(10); - - ASSERT_STATIC_ASSERT(thrust::generate(test1.cbegin(), test1.cend(), generator())); - ASSERT_STATIC_ASSERT(thrust::generate_n(test1.cbegin(), 10, generator())); -} -DECLARE_UNITTEST(TestGenerateConstIteratorCompilationError); - -void TestFillConstIteratorCompilationError() -{ - thrust::host_vector test1(10); - ASSERT_STATIC_ASSERT(thrust::fill(test1.cbegin(), test1.cend(), 1)); -} -DECLARE_UNITTEST(TestFillConstIteratorCompilationError); From 4ff382410c2e4adbb6384dd716a6acfdaadba84b Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 28 Feb 2025 23:12:38 +0100 Subject: [PATCH 06/10] Exec check --- thrust/thrust/system/detail/generic/generate.inl | 1 + 1 file changed, 1 insertion(+) diff --git a/thrust/thrust/system/detail/generic/generate.inl b/thrust/thrust/system/detail/generic/generate.inl index 1364a8b7d53..e888619b419 100644 --- a/thrust/thrust/system/detail/generic/generate.inl +++ b/thrust/thrust/system/detail/generic/generate.inl @@ -36,6 +36,7 @@ namespace system::detail::generic template struct generate_functor { + _CCCL_EXEC_CHECK_DISABLE template _CCCL_HOST_DEVICE void operator()(T&& x) { From 98fc048ee743df9f4f716e8b0253a832986543af Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 28 Feb 2025 23:13:35 +0100 Subject: [PATCH 07/10] Drop runtime_static_assert It's no longer needed --- .../testing/unittest/runtime_static_assert.h | 115 ------------------ thrust/testing/unittest_static_assert.cmake | 10 -- thrust/testing/unittest_static_assert.cu | 33 ----- 3 files changed, 158 deletions(-) delete mode 100644 thrust/testing/unittest/runtime_static_assert.h delete mode 100644 thrust/testing/unittest_static_assert.cmake delete mode 100644 thrust/testing/unittest_static_assert.cu diff --git a/thrust/testing/unittest/runtime_static_assert.h b/thrust/testing/unittest/runtime_static_assert.h deleted file mode 100644 index 73cb1d235b2..00000000000 --- a/thrust/testing/unittest/runtime_static_assert.h +++ /dev/null @@ -1,115 +0,0 @@ -#pragma once - -#include - -#include -#undef THRUST_STATIC_ASSERT -#undef THRUST_STATIC_ASSERT_MSG - -#define THRUST_STATIC_ASSERT(B) unittest::assert_static((B), __FILE__, __LINE__); -#define THRUST_STATIC_ASSERT_MSG(B, msg) unittest::assert_static((B), __FILE__, __LINE__); - -namespace unittest -{ -_CCCL_HOST_DEVICE void assert_static(bool condition, const char* filename, int lineno); -} - -#include -#include - -#include - -#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA - -# define ASSERT_STATIC_ASSERT(X) \ - { \ - bool triggered = false; \ - using ex_t = unittest::static_assert_exception; \ - thrust::device_ptr device_ptr = thrust::device_new(); \ - ex_t* raw_ptr = thrust::raw_pointer_cast(device_ptr); \ - ::cudaMemcpyToSymbol(unittest::detail::device_exception, &raw_ptr, sizeof(ex_t*)); \ - try \ - { \ - X; \ - } \ - catch (ex_t) \ - { \ - triggered = true; \ - } \ - if (!triggered) \ - { \ - triggered = static_cast(*device_ptr).triggered; \ - } \ - thrust::device_free(device_ptr); \ - raw_ptr = nullptr; \ - ::cudaMemcpyToSymbol(unittest::detail::device_exception, &raw_ptr, sizeof(ex_t*)); \ - if (!triggered) \ - { \ - unittest::UnitTestFailure f; \ - f << "[" << __FILE__ << ":" << __LINE__ << "] did not trigger a THRUST_STATIC_ASSERT"; \ - throw f; \ - } \ - } - -#else - -# define ASSERT_STATIC_ASSERT(X) \ - { \ - bool triggered = false; \ - using ex_t = unittest::static_assert_exception; \ - try \ - { \ - X; \ - } \ - catch (ex_t) \ - { \ - triggered = true; \ - } \ - if (!triggered) \ - { \ - unittest::UnitTestFailure f; \ - f << "[" << __FILE__ << ":" << __LINE__ << "] did not trigger a THRUST_STATIC_ASSERT"; \ - throw f; \ - } \ - } - -#endif - -namespace unittest -{ -class static_assert_exception -{ -public: - _CCCL_HOST_DEVICE static_assert_exception() - : triggered(false) - {} - - _CCCL_HOST_DEVICE static_assert_exception(const char* filename, int lineno) - : triggered(true) - , filename(filename) - , lineno(lineno) - {} - - bool triggered; - const char* filename; - int lineno; -}; - -namespace detail -{ -#if _CCCL_COMPILER(GCC) || _CCCL_COMPILER(CLANG) -__attribute__((used)) -#endif -_CCCL_DEVICE static static_assert_exception* device_exception = nullptr; -} // namespace detail - -_CCCL_HOST_DEVICE void assert_static(bool condition, const char* filename, int lineno) -{ - if (!condition) - { - static_assert_exception ex(filename, lineno); - - NV_IF_TARGET(NV_IS_DEVICE, (*detail::device_exception = ex;), (throw ex;)); - } -} -} // namespace unittest diff --git a/thrust/testing/unittest_static_assert.cmake b/thrust/testing/unittest_static_assert.cmake deleted file mode 100644 index 9f65a656dd2..00000000000 --- a/thrust/testing/unittest_static_assert.cmake +++ /dev/null @@ -1,10 +0,0 @@ -# Disable unreachable code warnings. -# This test unconditionally throws in some places, the compiler will detect that -# control flow will never reach some instructions. This is intentional. -target_link_libraries(${test_target} PRIVATE cccl.silence_unreachable_code_warnings) - -# The machinery behind this test is not compatible with NVC++. -# See https://github.com/NVIDIA/thrust/issues/1397 -if ("NVHPC" STREQUAL "${CMAKE_CUDA_COMPILER_ID}") - set_tests_properties(${test_target} PROPERTIES DISABLED True) -endif() diff --git a/thrust/testing/unittest_static_assert.cu b/thrust/testing/unittest_static_assert.cu deleted file mode 100644 index 6c45a80bbfc..00000000000 --- a/thrust/testing/unittest_static_assert.cu +++ /dev/null @@ -1,33 +0,0 @@ -#include -#include - -// The runtime_static_assert header needs to come first as we are overwriting thrusts internal static assert -#include - -template -struct dependent_false -{ - enum - { - value = false - }; -}; - -template -struct static_assertion -{ - _CCCL_HOST_DEVICE T operator()() const - { - THRUST_STATIC_ASSERT(dependent_false::value); - return 0; - } -}; - -template -void TestStaticAssertAssert() -{ - using value_type = typename V::value_type; - V test(10); - ASSERT_STATIC_ASSERT(thrust::generate(test.begin(), test.end(), static_assertion())); -} -DECLARE_VECTOR_UNITTEST(TestStaticAssertAssert); From 4ad587d0b0bda35b51bd9e9e0fdf21906d7eaeed Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Sat, 1 Mar 2025 00:31:39 +0100 Subject: [PATCH 08/10] Fix --- thrust/thrust/mr/disjoint_pool.h | 1 + 1 file changed, 1 insertion(+) diff --git a/thrust/thrust/mr/disjoint_pool.h b/thrust/thrust/mr/disjoint_pool.h index b2441574492..a75070c212c 100644 --- a/thrust/thrust/mr/disjoint_pool.h +++ b/thrust/thrust/mr/disjoint_pool.h @@ -242,6 +242,7 @@ class disjoint_unsynchronized_pool_resource final , previous_allocated_count(other.previous_allocated_count) {} + _CCCL_EXEC_CHECK_DISABLE pool& operator=(const pool&) = default; _CCCL_HOST ~pool() {} From e50d8c33d27912a16f2829c1970c0de2e9ef155e Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Sat, 1 Mar 2025 02:37:58 +0100 Subject: [PATCH 09/10] Fix dangling ref --- thrust/thrust/detail/raw_reference_cast.h | 81 ++++++----------------- thrust/thrust/detail/tuple_transform.h | 64 ------------------ 2 files changed, 21 insertions(+), 124 deletions(-) delete mode 100644 thrust/thrust/detail/tuple_transform.h diff --git a/thrust/thrust/detail/raw_reference_cast.h b/thrust/thrust/detail/raw_reference_cast.h index ce958335e26..3a22bacf069 100644 --- a/thrust/thrust/detail/raw_reference_cast.h +++ b/thrust/thrust/detail/raw_reference_cast.h @@ -26,9 +26,9 @@ # pragma system_header #endif // no system header #include -#include #include #include +#include // the order of declarations and definitions in this file is totally goofy // this header defines raw_reference_cast, which has a few overloads towards the bottom of the file @@ -54,28 +54,26 @@ struct is_unwrappable : is_wrapped_reference // specialize is_unwrappable // a tuple is_unwrappable if any of its elements is_unwrappable template -struct is_unwrappable> : ::cuda::std::disjunction...> +struct is_unwrappable> : ::cuda::std::disjunction...> {}; // specialize is_unwrappable // a tuple_of_iterator_references is_unwrappable if any of its elements is_unwrappable template -struct is_unwrappable> - : ::cuda::std::disjunction...> +struct is_unwrappable> : ::cuda::std::disjunction...> {}; namespace raw_reference_detail { -template +template >::value> struct raw_reference_impl : ::cuda::std::add_lvalue_reference {}; template -struct raw_reference_impl>::value>> -{ - using type = ::cuda::std::add_lvalue_reference_t::type>; -}; +struct raw_reference_impl + : ::cuda::std::add_lvalue_reference::type> +{}; } // namespace raw_reference_detail @@ -106,15 +104,15 @@ struct raw_reference_tuple_helper // recurse on tuples template -struct raw_reference_tuple_helper> +struct raw_reference_tuple_helper> { - using type = thrust::tuple::type...>; + using type = tuple::type...>; }; template -struct raw_reference_tuple_helper> +struct raw_reference_tuple_helper> { - using type = thrust::detail::tuple_of_iterator_references::type...>; + using type = tuple_of_iterator_references::type...>; }; } // namespace raw_reference_detail @@ -125,10 +123,10 @@ struct raw_reference_tuple_helper -struct raw_reference> +struct raw_reference> { private: - using tuple_type = thrust::tuple; + using tuple_type = tuple; public: using type = typename eval_if::value, @@ -145,55 +143,17 @@ struct raw_reference> } // namespace detail // provide declarations of raw_reference_cast's overloads for raw_reference_caster below -template -_CCCL_HOST_DEVICE typename detail::raw_reference::type raw_reference_cast(T& ref); - -template -_CCCL_HOST_DEVICE typename detail::raw_reference::type raw_reference_cast(const T& ref); - -template -_CCCL_HOST_DEVICE auto raw_reference_cast(detail::tuple_of_iterator_references t) -> - typename detail::raw_reference>::type; - -namespace detail -{ - -struct raw_reference_caster -{ - template - _CCCL_HOST_DEVICE typename detail::raw_reference::type operator()(T& ref) - { - return thrust::raw_reference_cast(ref); - } - - template - _CCCL_HOST_DEVICE typename detail::raw_reference::type operator()(const T& ref) - { - return thrust::raw_reference_cast(ref); - } - - template - _CCCL_HOST_DEVICE typename detail::raw_reference>::type - operator()(thrust::detail::tuple_of_iterator_references t, - ::cuda::std::enable_if_t>::value>* = 0) - { - return thrust::raw_reference_cast(t); - } -}; // end raw_reference_caster - -} // namespace detail - template _CCCL_HOST_DEVICE typename detail::raw_reference::type raw_reference_cast(T& ref) { return *thrust::raw_pointer_cast(&ref); -} // end raw_reference_cast +} template _CCCL_HOST_DEVICE typename detail::raw_reference::type raw_reference_cast(const T& ref) { return *thrust::raw_pointer_cast(&ref); -} // end raw_reference_cast +} template _CCCL_HOST_DEVICE auto raw_reference_cast(detail::tuple_of_iterator_references t) -> @@ -201,11 +161,12 @@ _CCCL_HOST_DEVICE auto raw_reference_cast(detail::tuple_of_iterator_references>::value) { - thrust::detail::raw_reference_caster f; - - // note that we pass raw_reference_tuple_helper, not raw_reference as the unary metafunction - // the different way that raw_reference_tuple_helper unwraps tuples is important - return thrust::detail::tuple_host_device_transform(t, f); + using ResultTuple = tuple::type...>; + return ::cuda::std::apply( + [](auto&&... refs) { + return ResultTuple{raw_reference_cast(::cuda::std::forward(refs))...}; + }, + static_cast&>(t)); } else { diff --git a/thrust/thrust/detail/tuple_transform.h b/thrust/thrust/detail/tuple_transform.h deleted file mode 100644 index 1eb835a66b5..00000000000 --- a/thrust/thrust/detail/tuple_transform.h +++ /dev/null @@ -1,64 +0,0 @@ -/* - * Copyright 2008-2013 NVIDIA Corporation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include - -#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) -# pragma GCC system_header -#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) -# pragma clang system_header -#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) -# pragma system_header -#endif // no system header - -#include -#include - -THRUST_NAMESPACE_BEGIN - -namespace detail -{ - -template class UnaryMetaFunction, - typename UnaryFunction, - typename IndexSequence = thrust::make_index_sequence::value>> -struct tuple_transform_functor; - -template class UnaryMetaFunction, typename UnaryFunction, size_t... Is> -struct tuple_transform_functor> -{ - static _CCCL_HOST_DEVICE typename tuple_meta_transform::type - do_it_on_the_host_or_device(const Tuple& t, UnaryFunction f) - { - using XfrmTuple = typename tuple_meta_transform::type; - - return XfrmTuple(f(thrust::get(t))...); - } -}; - -template