From 94039ab0ebb24131de80cbab8517290480b5ca49 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 3 Feb 2025 15:42:47 +0100 Subject: [PATCH] Fix transform_iterator and drop result_of_adaptable_function Thrust's transform_iterator relied on thrust::identity::result_type via result_of_adaptable_function to avoid dangling references. However, thrust::identity, cuda::std::identity, and the placeholder expression _1 have the same issue without a workaround. This change centralizes and adds workarounds for all of these by introducing a new trait thrust::detail::transform_iterator_reference, which decays the return value type of any of the aforementioned function objects. --- thrust/testing/transform_iterator.cu | 24 ++++--- thrust/thrust/detail/functional/actor.h | 7 -- thrust/thrust/detail/functional/operators.h | 1 - .../result_of_adaptable_function.h | 67 ------------------- thrust/thrust/functional.h | 1 - .../iterator/detail/transform_iterator.inl | 44 +++++++++--- 6 files changed, 48 insertions(+), 96 deletions(-) delete mode 100644 thrust/thrust/detail/type_traits/result_of_adaptable_function.h diff --git a/thrust/testing/transform_iterator.cu b/thrust/testing/transform_iterator.cu index b95012286d5..c8e1242fd27 100644 --- a/thrust/testing/transform_iterator.cu +++ b/thrust/testing/transform_iterator.cu @@ -164,8 +164,9 @@ void TestTransformIteratorReferenceAndValueType() static_assert(is_same::value, ""); (void) it_tr_tid; - auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::__identity{}); - static_assert(is_same::value, ""); // inferred, like forward + auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::identity{}); + static_assert(is_same::value, ""); // special handling by + // transform_iterator_reference static_assert(is_same::value, ""); (void) it_tr_cid; } @@ -197,8 +198,9 @@ void TestTransformIteratorReferenceAndValueType() static_assert(is_same::value, ""); (void) it_tr_tid; - auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::__identity{}); - static_assert(is_same::value, ""); // inferred, like forward + auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::identity{}); + static_assert(is_same::value, ""); // special handling by + // transform_iterator_reference static_assert(is_same::value, ""); (void) it_tr_cid; } @@ -235,8 +237,9 @@ void TestTransformIteratorReferenceAndValueType() static_assert(is_same::value, ""); (void) it_tr_tid; - auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::__identity{}); - static_assert(is_same::value, ""); // inferred, like forward + auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::identity{}); + static_assert(is_same::value, ""); // special handling by + // transform_iterator_reference static_assert(is_same::value, ""); (void) it_tr_cid; } @@ -248,11 +251,10 @@ void TestTransformIteratorIdentity() thrust::device_vector v(3, 42); ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), thrust::identity{}), 42); - // FIXME(bgruber): fix transform_iterator to get these tests compiling: - // ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), thrust::identity<>{}), 42); - // ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), cuda::std::identity{}), 42); - // using namespace thrust::placeholders; - // ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), _1), 42); + ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), thrust::identity<>{}), 42); + ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), cuda::std::identity{}), 42); + using namespace thrust::placeholders; + ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), _1), 42); } DECLARE_UNITTEST(TestTransformIteratorIdentity); diff --git a/thrust/thrust/detail/functional/actor.h b/thrust/thrust/detail/functional/actor.h index c727be09306..e43e07176c0 100644 --- a/thrust/thrust/detail/functional/actor.h +++ b/thrust/thrust/detail/functional/actor.h @@ -35,7 +35,6 @@ # pragma system_header #endif // no system header #include -#include #include #include @@ -211,11 +210,5 @@ _CCCL_HOST_DEVICE auto compose(Eval e, const SubExpr1& subexpr1, const SubExpr2& {{::cuda::std::move(e)}, make_actor(subexpr1), make_actor(subexpr2)}}; } } // namespace functional - -template -struct result_of_adaptable_function(Args...)> -{ - using type = decltype(::cuda::std::declval>()(::cuda::std::declval()...)); -}; } // namespace detail THRUST_NAMESPACE_END diff --git a/thrust/thrust/detail/functional/operators.h b/thrust/thrust/detail/functional/operators.h index 57c52b92520..8f87e3250a1 100644 --- a/thrust/thrust/detail/functional/operators.h +++ b/thrust/thrust/detail/functional/operators.h @@ -34,7 +34,6 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header -#include #include #include diff --git a/thrust/thrust/detail/type_traits/result_of_adaptable_function.h b/thrust/thrust/detail/type_traits/result_of_adaptable_function.h deleted file mode 100644 index 44138e5f909..00000000000 --- a/thrust/thrust/detail/type_traits/result_of_adaptable_function.h +++ /dev/null @@ -1,67 +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 - -#include - -THRUST_NAMESPACE_BEGIN -namespace detail -{ -// Sets `type` to the result of the specified Signature invocation. If the callable defines a `result_type` alias -// member, that type is used instead. Use invoke_result / result_of when FuncType::result_type is not defined. -template -struct result_of_adaptable_function -{ -private: - template - struct impl; - - template - struct impl - { - using type = invoke_result_t; - }; - -public: - using type = typename impl::type; -}; - -// TODO(bgruber): remove this specialization eventually -// specialization for invocations which define result_type -_CCCL_SUPPRESS_DEPRECATED_PUSH -template -struct result_of_adaptable_function> -{ - using type = typename Functor::result_type; -}; -_CCCL_SUPPRESS_DEPRECATED_POP - -} // namespace detail -THRUST_NAMESPACE_END diff --git a/thrust/thrust/functional.h b/thrust/thrust/functional.h index a8887d839e8..0b3a74133ea 100644 --- a/thrust/thrust/functional.h +++ b/thrust/thrust/functional.h @@ -176,7 +176,6 @@ using ::cuda::std::bit_xor; template struct identity { - // FIXME(bgruber): we cannot remove this yet, because transform_iterator depends on it using result_type _LIBCUDACXX_DEPRECATED_IN_CXX11 = T; _CCCL_EXEC_CHECK_DISABLE diff --git a/thrust/thrust/iterator/detail/transform_iterator.inl b/thrust/thrust/iterator/detail/transform_iterator.inl index 68145e65c50..de828cec523 100644 --- a/thrust/thrust/iterator/detail/transform_iterator.inl +++ b/thrust/thrust/iterator/detail/transform_iterator.inl @@ -25,10 +25,12 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header -#include +#include +#include #include #include +#include #include THRUST_NAMESPACE_BEGIN @@ -39,19 +41,43 @@ class transform_iterator; namespace detail { +template +struct transform_iterator_reference +{ + // by default, dereferencing the iterator yields the same as the function. + using type = decltype(::cuda::std::declval()(::cuda::std::declval>())); +}; + +// for certain function objects, we need to tweak the reference type. Notably, identity functions must decay to values. +// See the implementation of transform_iterator<...>::dereference() for several comments on why this is necessary. +template +struct transform_iterator_reference, Iterator> +{ + using type = T; +}; +template +struct transform_iterator_reference, Iterator> +{ + using type = iterator_value_t; +}; +template +struct transform_iterator_reference<::cuda::std::identity, Iterator> +{ + using type = iterator_value_t; +}; +template +struct transform_iterator_reference, Iterator> +{ + using type = ::cuda::std::remove_reference_t>()( + ::cuda::std::declval>()))>; +}; + // Type function to compute the iterator_adaptor instantiation to be used for transform_iterator template struct make_transform_iterator_base { private: - // FIXME(bgruber): the next line should be correct, but thrust::identity lies and advertises a ::return_type of T, - // while its operator() returns const T& (which __invoke_of correctly detects), which causes transform_iterator to - // crash (or cause UB) during dereferencing. Check the test `thrust.test.dereference` for the OMP and TBB backends. - // using wrapped_func_ret_t = ::cuda::std::__invoke_of>; - using wrapped_func_ret_t = result_of_adaptable_function)>; - - // By default, dereferencing the iterator yields the same as the function. - using reference = typename ia_dflt_help::type; + using reference = typename ia_dflt_help>::type; using value_type = typename ia_dflt_help>::type; public: