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); 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); 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); 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/detail/raw_reference_cast.h b/thrust/thrust/detail/raw_reference_cast.h index d49585a30e8..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,32 +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...> -{}; - -template -struct enable_if_unwrappable : ::cuda::std::enable_if::value, Result> +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 @@ -110,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 @@ -129,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, @@ -141,81 +135,43 @@ 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 // 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 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); - -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 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; - - // 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 + if constexpr (detail::is_unwrappable>::value) + { + using ResultTuple = tuple::type...>; + return ::cuda::std::apply( + [](auto&&... refs) { + return ResultTuple{raw_reference_cast(::cuda::std::forward(refs))...}; + }, + static_cast&>(t)); + } + else + { + return t; + } +} THRUST_NAMESPACE_END 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