From bb364d608b2f2f7c511f715df00b9cac38c55a8d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Tue, 2 Jun 2026 12:44:14 +0000 Subject: [PATCH 01/86] Add offload-compress to generate_resource_spec --- projects/rocthrust/test/CMakeLists.txt | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/rocthrust/test/CMakeLists.txt b/projects/rocthrust/test/CMakeLists.txt index 8bf2f4ffee9e..898fec65cd34 100644 --- a/projects/rocthrust/test/CMakeLists.txt +++ b/projects/rocthrust/test/CMakeLists.txt @@ -379,6 +379,9 @@ rocm_install( add_executable(generate_resource_spec ${CMAKE_SOURCE_DIR}/test/generate_resource_spec.cpp) set_target_properties(generate_resource_spec PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}") target_link_libraries(generate_resource_spec PRIVATE hip::host) +# This test may still get the offload-compress flag passed. Since it does not include any kernel +# code it will give an unused-command-line-argument warning. +target_compile_options(generate_resource_spec PRIVATE "-Wno-unused-command-line-argument") # Install the executable using the standard install command instead of rocm_install # since we want to put it in a custom DESTINATION (rocm_install does not support From 6f1a065bf280ac93417a54d3ffadb694789fb768 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Tue, 2 Jun 2026 11:08:11 +0000 Subject: [PATCH 02/86] Remove ``_CCCL_INLINE_VAR`` replacement recommendation, since ``_CCCL_INLINE_VAR`` is removed from CCCL --- projects/rocthrust/thrust/detail/raw_reference_cast.h | 3 --- .../thrust/detail/type_traits/iterator/is_output_iterator.h | 2 -- .../thrust/iterator/detail/iterator_facade_category.h | 3 --- projects/rocthrust/thrust/iterator/detail/retag.h | 1 - 4 files changed, 9 deletions(-) diff --git a/projects/rocthrust/thrust/detail/raw_reference_cast.h b/projects/rocthrust/thrust/detail/raw_reference_cast.h index 774323bf4be8..a52a3e3209d6 100644 --- a/projects/rocthrust/thrust/detail/raw_reference_cast.h +++ b/projects/rocthrust/thrust/detail/raw_reference_cast.h @@ -54,19 +54,16 @@ __THRUST_DEFINE_HAS_NESTED_TYPE(is_wrapped_reference, wrapped_reference_hint) // wrapped reference-like things which aren't strictly wrapped references // (e.g. tuples of wrapped references) are considered unwrappable template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool can_unwrap = is_wrapped_reference::value; // specialize is_unwrappable // a tuple is_unwrappable if any of its elements is_unwrappable template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool can_unwrap> = (can_unwrap || ...); // specialize is_unwrappable // a tuple_of_iterator_references is_unwrappable if any of its elements is_unwrappable template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool can_unwrap> = (can_unwrap || ...); namespace raw_reference_detail diff --git a/projects/rocthrust/thrust/detail/type_traits/iterator/is_output_iterator.h b/projects/rocthrust/thrust/detail/type_traits/iterator/is_output_iterator.h index f30b7bf4df9e..6606b074adb9 100644 --- a/projects/rocthrust/thrust/detail/type_traits/iterator/is_output_iterator.h +++ b/projects/rocthrust/thrust/detail/type_traits/iterator/is_output_iterator.h @@ -36,11 +36,9 @@ THRUST_NAMESPACE_BEGIN namespace detail { template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool is_output_iterator = true; template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool is_output_iterator>> = _THRUST_STD::is_void_v> || _THRUST_STD::is_same_v, any_assign>; } // namespace detail diff --git a/projects/rocthrust/thrust/iterator/detail/iterator_facade_category.h b/projects/rocthrust/thrust/iterator/detail/iterator_facade_category.h index 7a8a7c6b9162..da4a536adcd8 100644 --- a/projects/rocthrust/thrust/iterator/detail/iterator_facade_category.h +++ b/projects/rocthrust/thrust/iterator/detail/iterator_facade_category.h @@ -47,19 +47,16 @@ THRUST_NAMESPACE_BEGIN namespace detail { template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool is_host_iterator_category = _THRUST_STD::is_convertible_v || _THRUST_STD::is_convertible_v; template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool is_device_iterator_category = _THRUST_STD::is_convertible_v || _THRUST_STD::is_convertible_v; template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool is_iterator_category = is_host_iterator_category || is_device_iterator_category; // adapted from http://www.boost.org/doc/libs/1_37_0/libs/iterator/doc/iterator_facade.html#iterator-category diff --git a/projects/rocthrust/thrust/iterator/detail/retag.h b/projects/rocthrust/thrust/iterator/detail/retag.h index b56d2879ee8c..387a951de504 100644 --- a/projects/rocthrust/thrust/iterator/detail/retag.h +++ b/projects/rocthrust/thrust/iterator/detail/retag.h @@ -38,7 +38,6 @@ namespace detail { // we can retag an iterator if FromTag converts to ToTag or vice versa template -// TODO(libhipcxx): replace inline with _CCCL_INLINE_VAR once libhipcxx gets ready inline constexpr bool is_retaggable = _THRUST_STD::is_convertible_v || _THRUST_STD::is_convertible_v; } // namespace detail From 19ba7d4e439ae4f40f6dbfc86f91ac09e934b2b0 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Tue, 2 Jun 2026 11:12:49 +0000 Subject: [PATCH 03/86] Allow cuda::par*.on() to take cuda::stream_ref --- .../rocthrust/thrust/system/cuda/detail/par.h | 22 +++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/projects/rocthrust/thrust/system/cuda/detail/par.h b/projects/rocthrust/thrust/system/cuda/detail/par.h index b11d266f226d..48e78f66c454 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/par.h +++ b/projects/rocthrust/thrust/system/cuda/detail/par.h @@ -39,6 +39,8 @@ #include #include +#include + THRUST_NAMESPACE_BEGIN namespace cuda_cub { @@ -62,6 +64,11 @@ struct execute_on_stream_base : execution_policy return result; } + Derived on(::cuda::stream_ref const& s) const + { + return on(s.get()); + } + private: friend _CCCL_HOST_DEVICE cudaStream_t get_stream(const execute_on_stream_base& exec) { @@ -87,6 +94,11 @@ struct execute_on_stream_nosync_base : execution_policy return result; } + Derived on(::cuda::stream_ref const& s) const + { + return on(s.get()); + } + private: friend _CCCL_HOST_DEVICE cudaStream_t get_stream(const execute_on_stream_nosync_base& exec) { @@ -135,6 +147,11 @@ struct par_t { return execute_on_stream(stream); } + + stream_attachment_type on(::cuda::stream_ref const& s) const + { + return on(s.get()); + } }; struct par_nosync_t @@ -154,6 +171,11 @@ struct par_nosync_t return execute_on_stream_nosync(stream); } + stream_attachment_type on(::cuda::stream_ref const& s) const + { + return on(s.get()); + } + private: // this function is defined to allow non-blocking calls on the default_stream() with thrust::cuda::par_nosync // without explicitly using thrust::cuda::par_nosync.on(default_stream()) From 06913c40ecc5ea0075b3794762e0432fac58c51f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Tue, 2 Jun 2026 11:49:16 +0000 Subject: [PATCH 04/86] Simplify Thrust assign_value --- .../thrust/system/cuda/detail/assign_value.h | 88 ++++++++++++++----- .../cuda/detail/internal/copy_cross_system.h | 3 + 2 files changed, 68 insertions(+), 23 deletions(-) diff --git a/projects/rocthrust/thrust/system/cuda/detail/assign_value.h b/projects/rocthrust/thrust/system/cuda/detail/assign_value.h index 6965db5c4f54..982b73c423d7 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/assign_value.h +++ b/projects/rocthrust/thrust/system/cuda/detail/assign_value.h @@ -27,57 +27,99 @@ #endif // no system header #if _CCCL_HAS_CUDA_COMPILER() -# include - # include -# include +# include # include +# include # include THRUST_NAMESPACE_BEGIN namespace cuda_cub { + +namespace detail +{ +template +CCCL_DETAIL_KERNEL_ATTRIBUTES void assign_value_kernel(T* dst, const U* src) +{ + *dst = *src; +} +} // namespace detail + template _CCCL_HOST_DEVICE void assign_value(execution_policy& exec, Pointer1 dst, Pointer2 src) { + using V1 = thrust::detail::it_value_t; + using V2 = thrust::detail::it_value_t; // Because of https://docs.nvidia.com/cuda/cuda-c-programming-guide/#cuda-arch point 2., if a call from a __host__ // __device__ function leads to the template instantiation of a __global__ function, then this instantiation needs to // happen regardless of whether __CUDA_ARCH__ is defined. Therefore, we make the host path visible outside the // _THRUST_IF_TARGET switch. See also NVBug 881631. struct HostPath { - _CCCL_HOST auto operator()(execution_policy& exec, Pointer1 dst, Pointer2 src) + execution_policy& exec; + Pointer1 dst; + Pointer2 src; + + _CCCL_HOST void operator()() const { - cuda_cub::copy(exec, src, src + 1, dst); + if constexpr (::cuda::std::is_same_v) + { + const cudaError status = trivial_copy_device_to_device(exec, raw_pointer_cast(dst), raw_pointer_cast(src), 1); + throw_on_error(status, "__copy:: D->D: failed"); + } + else + { + const cudaError status = + cuda_cub::detail::triple_chevron(1, 1, 0, stream(exec)) + .doit(detail::assign_value_kernel, raw_pointer_cast(dst), raw_pointer_cast(src)); + throw_on_error(status, "__copy:: D->D with different data types: kernel failed"); + throw_on_error(synchronize_optional(exec), "__copy:: D->D with different data types: sync failed"); + } } }; _THRUST_IF_TARGET( - _THRUST_IS_HOST, (HostPath{}(exec, dst, src);), *thrust::raw_pointer_cast(dst) = *thrust::raw_pointer_cast(src);); + _THRUST_IS_HOST, + // on host, perform device -> device memcpy + (HostPath{exec, dst, src}();), + // on device, simply assign + *thrust::raw_pointer_cast(dst) = *thrust::raw_pointer_cast(src);); } +namespace detail +{ +struct cross_system_assign_host_path +{ + // device -> host copy executed from host + template + _CCCL_HOST void operator()(System1&, execution_policy& system2, Pointer1 dst, Pointer2 src) + { + thrust::detail::it_value_t copy_dst; + const cudaError status = trivial_copy_from_device(©_dst, raw_pointer_cast(src), 1, stream(system2)); + *dst = copy_dst; // may convert type + throw_on_error(status, "__copy:: D->H: failed"); + } + + // host -> device copy executed from host + template + _CCCL_HOST void operator()(execution_policy& system1, System2&, Pointer1 dst, Pointer2 src) + { + thrust::detail::it_value_t copy_src = *src; // may convert type + const cudaError status = trivial_copy_to_device(raw_pointer_cast(dst), ©_src, 1, stream(system1)); + throw_on_error(status, "__copy:: H->D: failed"); + } +}; +} // namespace detail + template _CCCL_HOST_DEVICE void assign_value(cross_system& systems, Pointer1 dst, Pointer2 src) { - // Because of https://docs.nvidia.com/cuda/cuda-c-programming-guide/#cuda-arch point 2., if a call from a __host__ - // __device__ function leads to the template instantiation of a __global__ function, then this instantiation needs to - // happen regardless of whether __CUDA_ARCH__ is defined. Therefore, we make the host path visible outside the - // _THRUST_IF_TARGET switch. See also NVBug 881631. - struct HostPath - { - _CCCL_HOST auto operator()(cross_system& systems, Pointer1 dst, Pointer2 src) - { - // rotate the systems so that they are ordered the same as (src, dst) for the call to thrust::copy - cross_system rotated_systems = systems.rotate(); - cuda_cub::copy(rotated_systems, src, src + 1, dst); - } - }; _THRUST_IF_TARGET( _THRUST_IS_HOST, - (HostPath{}(systems, dst, src);), - ( - // XXX forward the true cuda::execution_policy inside systems here instead of materializing a tag - cuda::tag cuda_tag; cuda_cub::assign_value(cuda_tag, dst, src);)); + (detail::cross_system_assign_host_path{}( + thrust::detail::derived_cast(systems.sys1), thrust::detail::derived_cast(systems.sys2), dst, src);), + (*thrust::raw_pointer_cast(dst) = *thrust::raw_pointer_cast(src);)); } } // namespace cuda_cub THRUST_NAMESPACE_END diff --git a/projects/rocthrust/thrust/system/cuda/detail/internal/copy_cross_system.h b/projects/rocthrust/thrust/system/cuda/detail/internal/copy_cross_system.h index 614ed1fd5629..f871e4c875d1 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/internal/copy_cross_system.h +++ b/projects/rocthrust/thrust/system/cuda/detail/internal/copy_cross_system.h @@ -185,6 +185,9 @@ OutputIt _CCCL_HOST cross_system_copy_n(cross_system systems, begin, n, result, + // FIXME(bgruber): I think this is a pessimization. We should only check if the iterator is contiguous and the value + // types are the same, and not whether value_t is trivially copyable, since we memcpy the content + // regardless in the non-trivial path, but pay for a temporary storage allocation. typename is_indirectly_trivially_relocatable_to::type()); } From fc7230c73e8a17ee595cc3547cd84445d8dee411 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Tue, 2 Jun 2026 12:38:10 +0000 Subject: [PATCH 05/86] Use ``[[maybe_unused]]`` --- .../rocthrust/test/test_address_stability.cpp | 3 +- projects/rocthrust/test/test_complex.cpp | 3 +- .../test/test_transform_iterator.cpp | 30 +++++----------- .../rocthrust/testing/address_stability.cu | 3 +- projects/rocthrust/testing/complex.cu | 3 +- projects/rocthrust/testing/cuda/merge_sort.cu | 8 ++--- .../rocthrust/testing/transform_iterator.cu | 36 +++++++------------ .../rocthrust/thrust/detail/complex/catrig.h | 3 +- .../thrust/detail/contiguous_storage.h | 8 ++--- .../thrust/device_malloc_allocator.h | 5 +-- .../rocthrust/thrust/device_new_allocator.h | 3 +- projects/rocthrust/thrust/mr/new.h | 6 ++-- .../linear_congruential_engine_discard.h | 13 +++---- .../system/cuda/detail/core/agent_launcher.h | 4 +-- .../cuda/detail/core/triple_chevron_launch.h | 4 +-- .../thrust/system/cuda/memory_resource.h | 9 ++--- .../thrust/system/hip/memory_resource.h | 9 ++--- 17 files changed, 47 insertions(+), 103 deletions(-) diff --git a/projects/rocthrust/test/test_address_stability.cpp b/projects/rocthrust/test/test_address_stability.cpp index f7f4a619ae48..99491b6c288b 100644 --- a/projects/rocthrust/test/test_address_stability.cpp +++ b/projects/rocthrust/test/test_address_stability.cpp @@ -133,8 +133,7 @@ TEST(AddressStabilityTests, TestAddressStabilityLambda) return i + 2; }; static_assert(!proclaims_copyable_arguments::value, ""); - auto pr_device_l = proclaim_copyable_arguments(l); - (void) &pr_device_l; + [[maybe_unused]] auto pr_device_l = proclaim_copyable_arguments(l); static_assert(proclaims_copyable_arguments::value, ""); } diff --git a/projects/rocthrust/test/test_complex.cpp b/projects/rocthrust/test/test_complex.cpp index 5b0ea8583312..8bfa087b22b3 100644 --- a/projects/rocthrust/test/test_complex.cpp +++ b/projects/rocthrust/test/test_complex.cpp @@ -832,8 +832,7 @@ TYPED_TEST(ComplexTests, TestComplexExplicitConstruction) SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); const thrust::complex input(42.0, 1337.0); - const user_complex result = thrust::exp(input); - (void) result; + [[maybe_unused]] const user_complex result = thrust::exp(input); } TESTS_PAIRS_DEFINE(ComplexPairsTests, PairsTestsParams) diff --git a/projects/rocthrust/test/test_transform_iterator.cpp b/projects/rocthrust/test/test_transform_iterator.cpp index d18fce851622..c14f3a78d096 100644 --- a/projects/rocthrust/test/test_transform_iterator.cpp +++ b/projects/rocthrust/test/test_transform_iterator.cpp @@ -225,26 +225,22 @@ TEST(TransformIteratorTests, TestTransformIteratorReferenceAndValueType) static_assert(is_same::value, ""); // ordinary reference static_assert(is_same::value, ""); - auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); + [[maybe_unused]] auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_val; - auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); + [[maybe_unused]] auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_ref; - auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); + [[maybe_unused]] auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_fwd; auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); static_assert(is_same::value, ""); // special handling by // transform_iterator_reference static_assert(is_same::value, ""); - (void) it_tr_cid; } { @@ -254,26 +250,22 @@ TEST(TransformIteratorTests, TestTransformIteratorReferenceAndValueType) static_assert(is_same>::value, ""); // proxy reference static_assert(is_same::value, ""); - auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); + [[maybe_unused]] auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_val; - auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); + [[maybe_unused]] auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_ref; - auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); + [[maybe_unused]] auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); static_assert(is_same::value, ""); // wrapped reference is decayed static_assert(is_same::value, ""); - (void) it_tr_fwd; auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); static_assert(is_same::value, ""); // special handling by // transform_iterator_reference static_assert(is_same::value, ""); - (void) it_tr_cid; } { @@ -283,26 +275,22 @@ TEST(TransformIteratorTests, TestTransformIteratorReferenceAndValueType) static_assert(is_same::reference>::value, ""); // proxy reference static_assert(is_same::value, ""); - auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); + [[maybe_unused]] auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_val; - auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); + [[maybe_unused]] auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_ref; - auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); + [[maybe_unused]] auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); static_assert(is_same::value, ""); // proxy reference is decayed static_assert(is_same::value, ""); - (void) it_tr_fwd; auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); static_assert(is_same::value, ""); // special handling by // transform_iterator_reference static_assert(is_same::value, ""); - (void) it_tr_cid; } } diff --git a/projects/rocthrust/testing/address_stability.cu b/projects/rocthrust/testing/address_stability.cu index 2dabcf1dc041..8d0e6451c490 100644 --- a/projects/rocthrust/testing/address_stability.cu +++ b/projects/rocthrust/testing/address_stability.cu @@ -125,8 +125,7 @@ void TestAddressStabilityLambda() return i + 2; }; static_assert(!proclaims_copyable_arguments::value, ""); - auto pr_device_l = proclaim_copyable_arguments(l); - (void) &pr_device_l; + [[maybe_unused]] auto pr_device_l = proclaim_copyable_arguments(l); static_assert(proclaims_copyable_arguments::value, ""); } diff --git a/projects/rocthrust/testing/complex.cu b/projects/rocthrust/testing/complex.cu index bc0214ce91cc..fe003bac10f9 100644 --- a/projects/rocthrust/testing/complex.cu +++ b/projects/rocthrust/testing/complex.cu @@ -730,8 +730,7 @@ struct TestComplexExplicitConstruction void operator()() { const thrust::complex input(42.0, 1337.0); - const user_complex result = thrust::exp(input); - (void) result; + [[maybe_unused]] const user_complex result = thrust::exp(input); } }; SimpleUnitTest TestComplexExplicitConstructionInstance; diff --git a/projects/rocthrust/testing/cuda/merge_sort.cu b/projects/rocthrust/testing/cuda/merge_sort.cu index a1be652e873c..a48554886999 100644 --- a/projects/rocthrust/testing/cuda/merge_sort.cu +++ b/projects/rocthrust/testing/cuda/merge_sort.cu @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -191,7 +191,7 @@ void TestMergeSortDescendingKey(void) DECLARE_UNITTEST(TestMergeSortDescendingKey); template -void TestMergeSortAscendingKeyValue(const size_t n) +void TestMergeSortAscendingKeyValue([[maybe_unused]] const size_t n) { #if 0 thrust::host_vector h_keys = unittest::random_integers(n); @@ -208,7 +208,6 @@ void TestMergeSortAscendingKeyValue(const size_t n) ASSERT_EQUAL(h_keys, d_keys); ASSERT_EQUAL(h_values, d_values); #else - (void) n; KNOWN_FAILURE; #endif } @@ -239,7 +238,7 @@ void TestMergeSortDescendingKeyValue(void) DECLARE_UNITTEST(TestMergeSortDescendingKeyValue); template -void TestMergeSortKeyValue(size_t n) +void TestMergeSortKeyValue([[maybe_unused]] size_t n) { #if 0 using T = key_value; @@ -261,7 +260,6 @@ void TestMergeSortKeyValue(size_t n) ASSERT_EQUAL_QUIET(h_data, d_data); #else - (void) n; KNOWN_FAILURE; #endif } diff --git a/projects/rocthrust/testing/transform_iterator.cu b/projects/rocthrust/testing/transform_iterator.cu index d01fc3fb2938..ff8f7d5c7522 100644 --- a/projects/rocthrust/testing/transform_iterator.cu +++ b/projects/rocthrust/testing/transform_iterator.cu @@ -199,26 +199,22 @@ void TestTransformIteratorReferenceAndValueType() static_assert(is_same::value, ""); // ordinary reference static_assert(is_same::value, ""); - auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); + [[maybe_unused]] auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_val; - auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); + [[maybe_unused]] auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_ref; - auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); + [[maybe_unused]] auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_fwd; - auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); + [[maybe_unused]] auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); static_assert(is_same::value, ""); // special handling by // transform_iterator_reference static_assert(is_same::value, ""); - (void) it_tr_cid; } { @@ -228,26 +224,22 @@ void TestTransformIteratorReferenceAndValueType() static_assert(is_same>::value, ""); // proxy reference static_assert(is_same::value, ""); - auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); + [[maybe_unused]] auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_val; - auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); + [[maybe_unused]] auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_ref; - auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); + [[maybe_unused]] auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); static_assert(is_same::value, ""); // wrapped reference is decayed static_assert(is_same::value, ""); - (void) it_tr_fwd; - auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); + [[maybe_unused]] auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); static_assert(is_same::value, ""); // special handling by // transform_iterator_reference static_assert(is_same::value, ""); - (void) it_tr_cid; } { @@ -257,26 +249,22 @@ void TestTransformIteratorReferenceAndValueType() static_assert(is_same::reference>::value, ""); // proxy reference static_assert(is_same::value, ""); - auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); + [[maybe_unused]] auto it_tr_val = thrust::make_transform_iterator(it, flip_value{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_val; - auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); + [[maybe_unused]] auto it_tr_ref = thrust::make_transform_iterator(it, pass_ref{}); static_assert(is_same::value, ""); static_assert(is_same::value, ""); - (void) it_tr_ref; - auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); + [[maybe_unused]] auto it_tr_fwd = thrust::make_transform_iterator(it, forward{}); static_assert(is_same::value, ""); // proxy reference is decayed static_assert(is_same::value, ""); - (void) it_tr_fwd; - auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); + [[maybe_unused]] auto it_tr_cid = thrust::make_transform_iterator(it, ::internal::identity{}); static_assert(is_same::value, ""); // special handling by // transform_iterator_reference static_assert(is_same::value, ""); - (void) it_tr_cid; } } DECLARE_UNITTEST(TestTransformIteratorReferenceAndValueType); diff --git a/projects/rocthrust/thrust/detail/complex/catrig.h b/projects/rocthrust/thrust/detail/complex/catrig.h index 16077f4155c3..8cc984c22f19 100644 --- a/projects/rocthrust/thrust/detail/complex/catrig.h +++ b/projects/rocthrust/thrust/detail/complex/catrig.h @@ -68,8 +68,7 @@ THRUST_HOST_DEVICE inline void raise_inexact() { const volatile float tiny = 7.888609052210118054117286e-31; /* 0x1p-100; */ // needs the volatile to prevent compiler from ignoring it - volatile float junk = 1 + tiny; - (void) junk; + [[maybe_unused]] volatile float junk = 1 + tiny; } THRUST_HOST_DEVICE inline complex clog_for_large_values(complex z); diff --git a/projects/rocthrust/thrust/detail/contiguous_storage.h b/projects/rocthrust/thrust/detail/contiguous_storage.h index 3e319cf5ff23..33b6f3a2b31d 100644 --- a/projects/rocthrust/thrust/detail/contiguous_storage.h +++ b/projects/rocthrust/thrust/detail/contiguous_storage.h @@ -177,8 +177,8 @@ class contiguous_storage } THRUST_EXEC_CHECK_DISABLE - THRUST_HOST_DEVICE void - destroy_on_allocator_mismatch(const contiguous_storage& other, iterator first, iterator last) noexcept + THRUST_HOST_DEVICE void destroy_on_allocator_mismatch( + const contiguous_storage& other, [[maybe_unused]] iterator first, [[maybe_unused]] iterator last) noexcept { if constexpr (allocator_traits::propagate_on_container_copy_assignment::value) { @@ -187,10 +187,6 @@ class contiguous_storage destroy(first, last); } } -#if THRUST_COMPILER(GCC, <, 10) - (void) first; - (void) last; -#endif } THRUST_HOST_DEVICE void set_allocator(const allocator_type& alloc); diff --git a/projects/rocthrust/thrust/device_malloc_allocator.h b/projects/rocthrust/thrust/device_malloc_allocator.h index 1faaca360341..1098702dbf25 100644 --- a/projects/rocthrust/thrust/device_malloc_allocator.h +++ b/projects/rocthrust/thrust/device_malloc_allocator.h @@ -155,11 +155,8 @@ class device_malloc_allocator * \note Memory deallocated by this function must previously have been * allocated with \p allocate. */ - THRUST_HOST inline void deallocate(pointer p, size_type cnt) noexcept + THRUST_HOST inline void deallocate(pointer p, [[maybe_unused]] size_type cnt) noexcept { - // silence unused parameter warning while still leaving the parameter name for Doxygen - (void) (cnt); - device_free(p); } // end deallocate() diff --git a/projects/rocthrust/thrust/device_new_allocator.h b/projects/rocthrust/thrust/device_new_allocator.h index b91901e9f8e5..742c246a132d 100644 --- a/projects/rocthrust/thrust/device_new_allocator.h +++ b/projects/rocthrust/thrust/device_new_allocator.h @@ -150,10 +150,9 @@ class device_new_allocator * \note Memory deallocated by this function must previously have been * allocated with \p allocate. */ - THRUST_HOST inline void deallocate(pointer p, size_type cnt) noexcept + THRUST_HOST inline void deallocate(pointer p, [[maybe_unused]] size_type cnt) noexcept { // use "::operator delete" rather than keyword delete - (void) cnt; device_delete(p); } // end deallocate() diff --git a/projects/rocthrust/thrust/mr/new.h b/projects/rocthrust/thrust/mr/new.h index 3021f45578bb..f1949fb43321 100644 --- a/projects/rocthrust/thrust/mr/new.h +++ b/projects/rocthrust/thrust/mr/new.h @@ -67,17 +67,17 @@ class new_delete_resource_base : public memory_resource<> # endif } - void do_deallocate(void* p, std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override + void do_deallocate(void* p, + [[maybe_unused]] std::size_t bytes, + [[maybe_unused]] std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override { # if defined(__cpp_aligned_new) # if defined(__cpp_sized_deallocation) ::operator delete(p, bytes, std::align_val_t(alignment)); # else - (void) bytes; ::operator delete(p, std::align_val_t(alignment)); # endif # else - (void) alignment; char* ptr = static_cast(p); // calculate where the offset is stored std::size_t* offset = reinterpret_cast(ptr + bytes); diff --git a/projects/rocthrust/thrust/random/detail/linear_congruential_engine_discard.h b/projects/rocthrust/thrust/random/detail/linear_congruential_engine_discard.h index 8fbcc67eb0c0..37de89c5b49c 100644 --- a/projects/rocthrust/thrust/random/detail/linear_congruential_engine_discard.h +++ b/projects/rocthrust/thrust/random/detail/linear_congruential_engine_discard.h @@ -88,15 +88,10 @@ struct linear_congruential_engine_discard template THRUST_HOST_DEVICE static void discard(LinearCongruentialEngine& lcg, unsigned long long z) { - using result_type = typename LinearCongruentialEngine::result_type; - const result_type c = LinearCongruentialEngine::increment; - const result_type a = LinearCongruentialEngine::multiplier; - const result_type m = LinearCongruentialEngine::modulus; - - // XXX WAR unused variable warnings - (void) c; - (void) a; - (void) m; + using result_type = typename LinearCongruentialEngine::result_type; + [[maybe_unused]] const result_type c = LinearCongruentialEngine::increment; + [[maybe_unused]] const result_type a = LinearCongruentialEngine::multiplier; + [[maybe_unused]] const result_type m = LinearCongruentialEngine::modulus; linear_congruential_engine_discard_implementation::discard(lcg.m_x, z); } diff --git a/projects/rocthrust/thrust/system/cuda/detail/core/agent_launcher.h b/projects/rocthrust/thrust/system/cuda/detail/core/agent_launcher.h index 914a6caca2cc..ac58cda94db9 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/core/agent_launcher.h +++ b/projects/rocthrust/thrust/system/cuda/detail/core/agent_launcher.h @@ -199,7 +199,7 @@ struct AgentLauncher : Agent } template - THRUST_RUNTIME_FUNCTION void print_info(K k) const + THRUST_RUNTIME_FUNCTION void print_info([[maybe_unused]] K k) const { # if THRUST_DEBUG_SYNC_FLAG cuda_optional occ = max_sm_occupancy(k); @@ -234,8 +234,6 @@ struct AgentLauncher : Agent (!has_shmem ? (int) plan.shared_memory_size : 0), (int) ptx_version); } -# else - (void) k; # endif } diff --git a/projects/rocthrust/thrust/system/cuda/detail/core/triple_chevron_launch.h b/projects/rocthrust/thrust/system/cuda/detail/core/triple_chevron_launch.h index 7c87d677c5e3..e31345b56569 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/core/triple_chevron_launch.h +++ b/projects/rocthrust/thrust/system/cuda/detail/core/triple_chevron_launch.h @@ -129,7 +129,7 @@ struct _CCCL_VISIBILITY_HIDDEN triple_chevron size_t _CCCL_DEVICE argument_pack_size(size_t size, Args const&...) const { // TODO(bgruber): replace by fold over comma in C++17 (make sure order of evaluation is left to right!) - int dummy[] = {(size = align_up(size) + sizeof(Args), 0)...}; + [[maybe_unused]] int dummy[] = {(size = align_up(size) + sizeof(Args), 0)...}; static_cast(dummy); return size; } @@ -150,7 +150,7 @@ struct _CCCL_VISIBILITY_HIDDEN triple_chevron _CCCL_DEVICE void fill_arguments(char* buffer, size_t offset, Args const&... args) const { // TODO(bgruber): replace by fold over comma in C++17 (make sure order of evaluation is left to right!) - int dummy[] = {(copy_arg(buffer, offset, args), 0)...}; + [[maybe_unused]] int dummy[] = {(copy_arg(buffer, offset, args), 0)...}; static_cast(dummy); } diff --git a/projects/rocthrust/thrust/system/cuda/memory_resource.h b/projects/rocthrust/thrust/system/cuda/memory_resource.h index 23373e6b2f0b..595f318d2402 100644 --- a/projects/rocthrust/thrust/system/cuda/memory_resource.h +++ b/projects/rocthrust/thrust/system/cuda/memory_resource.h @@ -55,10 +55,8 @@ template class cuda_memory_resource final : public mr::memory_resource { public: - Pointer do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override + Pointer do_allocate(std::size_t bytes, [[maybe_unused]] std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override { - (void) alignment; - void* ret; cudaError_t status = Alloc(&ret, bytes); @@ -71,11 +69,8 @@ class cuda_memory_resource final : public mr::memory_resource return Pointer(ret); } - void do_deallocate(Pointer p, std::size_t bytes, std::size_t alignment) override + void do_deallocate(Pointer p, [[maybe_unused]] std::size_t bytes, [[maybe_unused]] std::size_t alignment) override { - (void) bytes; - (void) alignment; - cudaError_t status = Dealloc(thrust::detail::pointer_traits::get(p)); if (status != cudaSuccess) diff --git a/projects/rocthrust/thrust/system/hip/memory_resource.h b/projects/rocthrust/thrust/system/hip/memory_resource.h index 41d9cbed5cf1..7a6db285057b 100644 --- a/projects/rocthrust/thrust/system/hip/memory_resource.h +++ b/projects/rocthrust/thrust/system/hip/memory_resource.h @@ -56,10 +56,8 @@ template class hip_memory_resource final : public mr::memory_resource { public: - Pointer do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override + Pointer do_allocate(std::size_t bytes, [[maybe_unused]] std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override { - (void) alignment; - void* ret; hipError_t status = Alloc(&ret, bytes); @@ -72,11 +70,8 @@ class hip_memory_resource final : public mr::memory_resource return Pointer(ret); } - void do_deallocate(Pointer p, std::size_t bytes, std::size_t alignment) override + void do_deallocate(Pointer p, [[maybe_unused]] std::size_t bytes, [[maybe_unused]] std::size_t alignment) override { - (void) bytes; - (void) alignment; - hipError_t status = Dealloc(thrust::detail::pointer_traits::get(p)); if (status != hipSuccess) From 0f6a6dd125a8547635429fa2f0019eb9b9a2c426 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Tue, 2 Jun 2026 15:04:57 +0000 Subject: [PATCH 06/86] Make ``cuda::stream_ref`` constructible on device --- .../rocthrust/thrust/system/cuda/detail/par.h | 37 ++++--------------- 1 file changed, 8 insertions(+), 29 deletions(-) diff --git a/projects/rocthrust/thrust/system/cuda/detail/par.h b/projects/rocthrust/thrust/system/cuda/detail/par.h index 48e78f66c454..afb6b5a82d7f 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/par.h +++ b/projects/rocthrust/thrust/system/cuda/detail/par.h @@ -52,23 +52,17 @@ struct execute_on_stream_base : execution_policy cudaStream_t stream; public: - _CCCL_EXEC_CHECK_DISABLE _CCCL_HOST_DEVICE execute_on_stream_base(cudaStream_t stream_ = default_stream()) : stream(stream_) {} - THRUST_RUNTIME_FUNCTION Derived on(cudaStream_t const& s) const + _CCCL_HOST_DEVICE Derived on(::cuda::stream_ref s) const { Derived result = derived_cast(*this); - result.stream = s; + result.stream = s.get(); return result; } - Derived on(::cuda::stream_ref const& s) const - { - return on(s.get()); - } - private: friend _CCCL_HOST_DEVICE cudaStream_t get_stream(const execute_on_stream_base& exec) { @@ -87,18 +81,13 @@ struct execute_on_stream_nosync_base : execution_policy : stream(stream_) {} - THRUST_RUNTIME_FUNCTION Derived on(cudaStream_t const& s) const + _CCCL_HOST_DEVICE Derived on(::cuda::stream_ref s) const { Derived result = derived_cast(*this); - result.stream = s; + result.stream = s.get(); return result; } - Derived on(::cuda::stream_ref const& s) const - { - return on(s.get()); - } - private: friend _CCCL_HOST_DEVICE cudaStream_t get_stream(const execute_on_stream_nosync_base& exec) { @@ -143,14 +132,9 @@ struct par_t using stream_attachment_type = execute_on_stream; - THRUST_RUNTIME_FUNCTION stream_attachment_type on(cudaStream_t const& stream) const + _CCCL_HOST_DEVICE stream_attachment_type on(::cuda::stream_ref s) const { - return execute_on_stream(stream); - } - - stream_attachment_type on(::cuda::stream_ref const& s) const - { - return on(s.get()); + return execute_on_stream(s.get()); } }; @@ -166,14 +150,9 @@ struct par_nosync_t using stream_attachment_type = execute_on_stream_nosync; - THRUST_RUNTIME_FUNCTION stream_attachment_type on(cudaStream_t const& stream) const - { - return execute_on_stream_nosync(stream); - } - - stream_attachment_type on(::cuda::stream_ref const& s) const + _CCCL_HOST_DEVICE stream_attachment_type on(::cuda::stream_ref s) const { - return on(s.get()); + return execute_on_stream_nosync(s.get()); } private: From a918bbe0b6588f761a2522a518f4cf975c49f128 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Tue, 2 Jun 2026 15:05:40 +0000 Subject: [PATCH 07/86] Drop ``THRUST_NODISCARD`` --- projects/rocthrust/.clang-format | 2 -- projects/rocthrust/CHANGELOG.md | 8 +++++++- projects/rocthrust/docs/doxygen/Doxyfile | 1 - .../libcxx_wrapper/__functional/address_stability.h | 2 +- .../detail/libcxx_wrapper/__functional/maximum.h | 4 ++-- .../detail/libcxx_wrapper/__functional/minimum.h | 4 ++-- .../detail/libcxx_wrapper/std/__cccl/attributes.h | 13 +++---------- .../libcxx_wrapper/std/__functional/identity.h | 2 +- .../detail/libcxx_wrapper/std/__functional/not_fn.h | 2 +- projects/rocthrust/thrust/device_allocator.h | 2 +- projects/rocthrust/thrust/mr/allocator.h | 2 +- projects/rocthrust/thrust/mr/disjoint_pool.h | 2 +- projects/rocthrust/thrust/mr/disjoint_sync_pool.h | 2 +- .../rocthrust/thrust/mr/fancy_pointer_resource.h | 2 +- projects/rocthrust/thrust/mr/memory_resource.h | 4 ++-- projects/rocthrust/thrust/mr/pool.h | 2 +- projects/rocthrust/thrust/mr/sync_pool.h | 2 +- 17 files changed, 26 insertions(+), 30 deletions(-) diff --git a/projects/rocthrust/.clang-format b/projects/rocthrust/.clang-format index 676f9822da5f..840e62ec97fb 100644 --- a/projects/rocthrust/.clang-format +++ b/projects/rocthrust/.clang-format @@ -28,7 +28,6 @@ AttributeMacros: [ 'THRUST_FORCEINLINE', 'THRUST_HOST_DEVICE', 'THRUST_HOST', - 'THRUST_NODISCARD', 'THRUST_DEPRECATED', 'THRUST_GLOBAL_CONSTANT', # thrust @@ -45,7 +44,6 @@ AttributeMacros: [ '_CCCL_HOST', '_CCCL_NO_UNIQUE_ADDRESS', '_CCCL_NODISCARD_FRIEND', - '_CCCL_NODISCARD', '_CCCL_NORETURN', '_CCCL_VISIBILITY_HIDDEN', 'CUB_RUNTIME_FUNCTION', diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index 31cc541827bc..6ae585ce303e 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -3,8 +3,14 @@ Documentation for rocThrust available at [https://rocm.docs.amd.com/projects/rocThrust/en/latest/](https://rocm.docs.amd.com/projects/rocThrust/en/latest/). +## rocThrust *.*.* for ROCm *.*.* (CCCL 3.1 parity) + +### Removed + +* Removed `THRUST_NODISCARD`. + ## rocThrust 5.0.0 for ROCm 10.2.0 - + ### Added * Largely in feature parity with CCCL/thrust v3.0.3. diff --git a/projects/rocthrust/docs/doxygen/Doxyfile b/projects/rocthrust/docs/doxygen/Doxyfile index 10f8f7447d8c..89956d6ebf06 100644 --- a/projects/rocthrust/docs/doxygen/Doxyfile +++ b/projects/rocthrust/docs/doxygen/Doxyfile @@ -2396,7 +2396,6 @@ PREDEFINED = THRUST_DOXYGEN_INVOKED \ THRUST_DEVICE_SYSTEM=5 \ THRUST_STD_VER=2017 \ THRUST_GLOBAL_CONSTANT="inline constexpr" \ - THRUST_NODISCARD=[[nodiscard]] \ THRUST_MR_DEFAULT_ALIGNMENT=alignof(std::max_align_t) \ __cpp_lib_remove_cvref=201711 \ THRUST_PREVENT_MACRO_SUBSTITUTION= \ diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/address_stability.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/address_stability.h index c20bbec71dba..565df1eff35e 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/address_stability.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/address_stability.h @@ -59,7 +59,7 @@ struct proclaims_copyable_arguments -THRUST_NODISCARD inline THRUST_HOST_DEVICE constexpr auto +[[nodiscard]] inline THRUST_HOST_DEVICE constexpr auto proclaim_copyable_arguments(F&& f) -> detail::callable_permitting_copied_arguments<::internal::decay_t> { return {::std::forward(f)}; diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/maximum.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/maximum.h index c31032031586..b196841320bc 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/maximum.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/maximum.h @@ -38,7 +38,7 @@ template struct maximum { THRUST_EXEC_CHECK_DISABLE - THRUST_NODISCARD inline THRUST_HOST_DEVICE constexpr Tp operator()(const Tp& lhs, const Tp& rhs) const + [[nodiscard]] inline THRUST_HOST_DEVICE constexpr Tp operator()(const Tp& lhs, const Tp& rhs) const noexcept(noexcept((lhs < rhs) ? rhs : lhs)) { return (lhs < rhs) ? rhs : lhs; @@ -51,7 +51,7 @@ struct maximum { THRUST_EXEC_CHECK_DISABLE template - THRUST_NODISCARD inline THRUST_HOST_DEVICE constexpr ::std::common_type_t + [[nodiscard]] inline THRUST_HOST_DEVICE constexpr ::std::common_type_t operator()(const T1& lhs, const T2& rhs) const noexcept(noexcept((lhs < rhs) ? rhs : lhs)) { return (lhs < rhs) ? rhs : lhs; diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/minimum.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/minimum.h index 5f5f1baa1ef6..5e3e9f031e43 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/minimum.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/__functional/minimum.h @@ -38,7 +38,7 @@ template struct minimum { THRUST_EXEC_CHECK_DISABLE - THRUST_NODISCARD inline THRUST_HOST_DEVICE constexpr Tp operator()(const Tp& lhs, const Tp& rhs) const + [[nodiscard]] inline THRUST_HOST_DEVICE constexpr Tp operator()(const Tp& lhs, const Tp& rhs) const noexcept(noexcept((lhs < rhs) ? lhs : rhs)) { return (lhs < rhs) ? lhs : rhs; @@ -51,7 +51,7 @@ struct minimum { THRUST_EXEC_CHECK_DISABLE template - THRUST_NODISCARD inline THRUST_HOST_DEVICE constexpr ::std::common_type_t + [[nodiscard]] inline THRUST_HOST_DEVICE constexpr ::std::common_type_t operator()(const T1& lhs, const T2& rhs) const noexcept(noexcept((lhs < rhs) ? lhs : rhs)) { return (lhs < rhs) ? lhs : rhs; diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__cccl/attributes.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__cccl/attributes.h index be6d1168730b..5038b36b8e50 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__cccl/attributes.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__cccl/attributes.h @@ -12,8 +12,8 @@ #ifndef LIBCXX_WRAPPER_STD__CCCL_ATTRIBUTES_H #define LIBCXX_WRAPPER_STD__CCCL_ATTRIBUTES_H -// TODO(libhipcxx): remove this file and replace THRUST_DECLSPEC_EMPTY_BASES, THRUST_NODISCARD* and -// THRUST_ALIAS_ATTRIBUTE with _CCCL_DECLSPEC_EMPTY_BASES, _CCCL_NODISCARD* and _CCCL_ALIAS_ATTRIBUTE in rocThrust +// TODO(libhipcxx): remove this file and replace THRUST_DECLSPEC_EMPTY_BASES, THRUST_NODISCARD_FRIEND and +// THRUST_ALIAS_ATTRIBUTE with _CCCL_DECLSPEC_EMPTY_BASES, _CCCL_NODISCARD_FRIEND and _CCCL_ALIAS_ATTRIBUTE in rocThrust // once libhipcxx gets ready #include @@ -26,7 +26,6 @@ # define THRUST_HAS_ATTRIBUTE(__x) _CCCL_HAS_ATTRIBUTE(__x) # define THRUST_HAS_CPP_ATTRIBUTE(__x) _CCCL_HAS_CPP_ATTRIBUTE(__x) # define THRUST_DECLSPEC_EMPTY_BASES _CCCL_DECLSPEC_EMPTY_BASES -# define THRUST_NODISCARD _CCCL_NODISCARD # define THRUST_NODISCARD_FRIEND _CCCL_NODISCARD_FRIEND # define THRUST_NORETURN _CCCL_NORETURN @@ -56,16 +55,10 @@ # define THRUST_DECLSPEC_EMPTY_BASES # endif // !THRUST_COMPILER(MSVC) -# if THRUST_HAS_CPP_ATTRIBUTE(nodiscard) || THRUST_COMPILER(MSVC) -# define THRUST_NODISCARD [[nodiscard]] -# else // ^^^ has nodiscard ^^^ / vvv no nodiscard vvv -# define THRUST_NODISCARD -# endif // no nodiscard - # if THRUST_COMPILER(CLANG) # define THRUST_NODISCARD_FRIEND friend # else -# define THRUST_NODISCARD_FRIEND THRUST_NODISCARD friend +# define THRUST_NODISCARD_FRIEND [[nodiscard]] friend # endif # if THRUST_COMPILER(MSVC) diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/identity.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/identity.h index 207f966f53dc..20b7c16a783e 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/identity.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/identity.h @@ -37,7 +37,7 @@ using identity = _THRUST_STD::__identity; struct identity { template - THRUST_NODISCARD inline THRUST_HOST_DEVICE constexpr Tp&& operator()(Tp&& t) const noexcept + [[nodiscard]] inline THRUST_HOST_DEVICE constexpr Tp&& operator()(Tp&& t) const noexcept { return ::std::forward(t); } diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/not_fn.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/not_fn.h index 74e4624495be..7d17f9ae4489 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/not_fn.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/not_fn.h @@ -133,7 +133,7 @@ template , Fn> && cccl_true, int> = 0, ::std::enable_if_t<::std::is_move_constructible_v> && cccl_true, int> = 0> -THRUST_NODISCARD inline THRUST_HOST_DEVICE constexpr auto not_fn(Fn&& f) +[[nodiscard]] inline THRUST_HOST_DEVICE constexpr auto not_fn(Fn&& f) { return not_fn_t>(::std::forward(f)); } diff --git a/projects/rocthrust/thrust/device_allocator.h b/projects/rocthrust/thrust/device_allocator.h index c3ad71a570cf..dbb486c28138 100644 --- a/projects/rocthrust/thrust/device_allocator.h +++ b/projects/rocthrust/thrust/device_allocator.h @@ -75,7 +75,7 @@ class device_ptr_memory_resource final : public thrust::mr::memory_resourcedo_allocate(bytes, alignment).get()); diff --git a/projects/rocthrust/thrust/mr/allocator.h b/projects/rocthrust/thrust/mr/allocator.h index 863f2555b6f4..1e2ab779289a 100644 --- a/projects/rocthrust/thrust/mr/allocator.h +++ b/projects/rocthrust/thrust/mr/allocator.h @@ -127,7 +127,7 @@ class allocator : private validator * \param n number of elements to allocate * \return a pointer to the newly allocated storage. */ - THRUST_NODISCARD THRUST_HOST pointer allocate(size_type n) + [[nodiscard]] THRUST_HOST pointer allocate(size_type n) { return static_cast(mem_res->do_allocate(n * sizeof(T), alignof(T))); } diff --git a/projects/rocthrust/thrust/mr/disjoint_pool.h b/projects/rocthrust/thrust/mr/disjoint_pool.h index 85a4c7184fa4..bc639dcc80d6 100644 --- a/projects/rocthrust/thrust/mr/disjoint_pool.h +++ b/projects/rocthrust/thrust/mr/disjoint_pool.h @@ -297,7 +297,7 @@ class disjoint_unsynchronized_pool_resource final m_cached_oversized.clear(); } - THRUST_NODISCARD virtual void_ptr + [[nodiscard]] virtual void_ptr do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override { bytes = (std::max)(bytes, m_options.smallest_block_size); diff --git a/projects/rocthrust/thrust/mr/disjoint_sync_pool.h b/projects/rocthrust/thrust/mr/disjoint_sync_pool.h index edd21a654117..7dcae3660caf 100644 --- a/projects/rocthrust/thrust/mr/disjoint_sync_pool.h +++ b/projects/rocthrust/thrust/mr/disjoint_sync_pool.h @@ -95,7 +95,7 @@ struct disjoint_synchronized_pool_resource : public memory_resource(m_upstream->do_allocate(bytes, alignment)); diff --git a/projects/rocthrust/thrust/mr/memory_resource.h b/projects/rocthrust/thrust/mr/memory_resource.h index 6893f0932589..65429fbd5e0b 100644 --- a/projects/rocthrust/thrust/mr/memory_resource.h +++ b/projects/rocthrust/thrust/mr/memory_resource.h @@ -73,7 +73,7 @@ class memory_resource * \throws thrust::bad_alloc when no memory with requested size and alignment can be allocated. * \return A pointer to void to the newly allocated memory. */ - THRUST_NODISCARD pointer allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) + [[nodiscard]] pointer allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) { return do_allocate(bytes, alignment); } @@ -145,7 +145,7 @@ class memory_resource virtual ~memory_resource() = default; - THRUST_NODISCARD pointer allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) + [[nodiscard]] pointer allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) { return do_allocate(bytes, alignment); } diff --git a/projects/rocthrust/thrust/mr/pool.h b/projects/rocthrust/thrust/mr/pool.h index 519f63b0f706..4813633d338b 100644 --- a/projects/rocthrust/thrust/mr/pool.h +++ b/projects/rocthrust/thrust/mr/pool.h @@ -257,7 +257,7 @@ class unsynchronized_pool_resource final m_cached_oversized = oversized_block_descriptor_ptr(); } - THRUST_NODISCARD virtual void_ptr + [[nodiscard]] virtual void_ptr do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override { bytes = (std::max)(bytes, m_options.smallest_block_size); diff --git a/projects/rocthrust/thrust/mr/sync_pool.h b/projects/rocthrust/thrust/mr/sync_pool.h index 94891c8c6dc1..85bd0b8cce33 100644 --- a/projects/rocthrust/thrust/mr/sync_pool.h +++ b/projects/rocthrust/thrust/mr/sync_pool.h @@ -91,7 +91,7 @@ struct synchronized_pool_resource : public memory_resource Date: Tue, 2 Jun 2026 15:20:23 +0000 Subject: [PATCH 08/86] Use libhipcxx includes - std::move --- projects/rocthrust/CHANGELOG.md | 4 ++++ .../rocthrust/test/test_unique_ptr_general.cpp | 12 ++++++------ projects/rocthrust/test/test_vector.cpp | 6 +++--- .../rocthrust/test/test_vector_allocators.cpp | 8 ++++---- projects/rocthrust/testing/vector.cu | 6 +++--- projects/rocthrust/testing/vector_allocators.cu | 8 ++++---- projects/rocthrust/thrust/allocate_unique.h | 16 ++++++++-------- .../thrust/detail/contiguous_storage.inl | 6 ++++-- .../libcxx_wrapper/std/__functional/not_fn.h | 14 +++++++------- projects/rocthrust/thrust/detail/reference.h | 6 +++--- projects/rocthrust/thrust/detail/vector_base.inl | 6 +++--- projects/rocthrust/thrust/device_vector.h | 9 +++++---- projects/rocthrust/thrust/host_vector.h | 9 +++++---- .../thrust/system/cpp/detail/vector.inl | 6 +++--- .../thrust/system/detail/generic/sequence.inl | 2 +- projects/rocthrust/thrust/unique_ptr.h | 6 +++--- projects/rocthrust/thrust/zip_function.h | 2 +- 17 files changed, 67 insertions(+), 59 deletions(-) diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index 6ae585ce303e..b122e72cf19c 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -5,6 +5,10 @@ Documentation for rocThrust available at ## rocThrust *.*.* for ROCm *.*.* (CCCL 3.1 parity) +### Changed + +* Replaced `std::move` operators to `hip::std` operators. + ### Removed * Removed `THRUST_NODISCARD`. diff --git a/projects/rocthrust/test/test_unique_ptr_general.cpp b/projects/rocthrust/test/test_unique_ptr_general.cpp index 746b7fb0f3f8..b5f6b18c8355 100644 --- a/projects/rocthrust/test/test_unique_ptr_general.cpp +++ b/projects/rocthrust/test/test_unique_ptr_general.cpp @@ -136,7 +136,7 @@ TEST(UniquePtrGeneralTests, TestUniquePtrAsgnMove) ASSERT_NE(p1.get_raw(), nullptr); ASSERT_NE(p2.get_raw(), nullptr); - p2 = std::move(p1); + p2 = _THRUST_STD::move(p1); ASSERT_EQ(p2.get_raw(), raw_p1); ASSERT_EQ(p1.get_raw(), nullptr); @@ -168,7 +168,7 @@ TEST(UniquePtrGeneralTests, TestUniquePtrAsgnMoveArray) ASSERT_NE(p1.get_raw(), nullptr); ASSERT_NE(p2.get_raw(), nullptr); - p2 = std::move(p1); + p2 = _THRUST_STD::move(p1); ASSERT_EQ(p2.get_raw(), raw_p1); ASSERT_EQ(p1.get_raw(), nullptr); @@ -192,7 +192,7 @@ TEST(UniquePtrGeneralTests, TestUniquePtrAsgnSelfMove) int* raw_p = p.get_raw(); THRUST_DIAG_PUSH THRUST_DIAG_SUPPRESS_CLANG("-Wself-move") - p = std::move(p); + p = _THRUST_STD::move(p); ASSERT_EQ(p.get_raw(), raw_p); } } @@ -208,7 +208,7 @@ TEST(UniquePtrGeneralTests, TestUnqiuePtrAsgnSelfMoveArray) int* raw_p = p.get_raw(); THRUST_DIAG_PUSH THRUST_DIAG_SUPPRESS_CLANG("-Wself-move") - p = std::move(p); + p = _THRUST_STD::move(p); ASSERT_EQ(p.get_raw(), raw_p); } } @@ -280,7 +280,7 @@ TEST(UniquePtrGeneralTests, TestUniquePtrCtorMove) { thrust::unique_ptr p1 = thrust::make_unique(42); int* raw_p1 = p1.get_raw(); - thrust::unique_ptr p2(std::move(p1)); + thrust::unique_ptr p2(_THRUST_STD::move(p1)); ASSERT_EQ(p2.get_raw(), raw_p1); ASSERT_EQ(p1, nullptr); } @@ -333,7 +333,7 @@ TEST(UniquePtrGeneralTests, TestUniquePtrCtorMoveArray) { thrust::unique_ptr p1 = thrust::make_unique(3); int* raw_p1 = p1.get_raw(); - thrust::unique_ptr p2(std::move(p1)); + thrust::unique_ptr p2(_THRUST_STD::move(p1)); ASSERT_EQ(p2.get_raw(), raw_p1); ASSERT_EQ(p1, nullptr); } diff --git a/projects/rocthrust/test/test_vector.cpp b/projects/rocthrust/test/test_vector.cpp index ca45d04f1b72..fa203fea8437 100644 --- a/projects/rocthrust/test/test_vector.cpp +++ b/projects/rocthrust/test/test_vector.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -863,7 +863,7 @@ TYPED_TEST(VectorTests, TestVectorMove) const auto ptr1 = v1.data(); const auto size1 = v1.size(); - Vector v2(std::move(v1)); + Vector v2(_THRUST_STD::move(v1)); const auto ptr2 = v2.data(); const auto size2 = v2.size(); @@ -884,7 +884,7 @@ TYPED_TEST(VectorTests, TestVectorMove) const auto ptr3 = v3.data(); const auto size3 = v3.size(); - v2 = std::move(v3); + v2 = _THRUST_STD::move(v3); const auto ptr4 = v2.data(); const auto size4 = v2.size(); diff --git a/projects/rocthrust/test/test_vector_allocators.cpp b/projects/rocthrust/test/test_vector_allocators.cpp index 0e6068433afc..4ee4075d4d2e 100644 --- a/projects/rocthrust/test/test_vector_allocators.cpp +++ b/projects/rocthrust/test/test_vector_allocators.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -52,7 +52,7 @@ class stateful_allocator : public BaseAlloc } stateful_allocator(stateful_allocator&& other) - : BaseAlloc(std::move(other)) + : BaseAlloc(_THRUST_STD::move(other)) , state(other.state) { other.state = 0; @@ -172,7 +172,7 @@ void TestVectorAllocatorConstructors() // FIXME: uncomment this after the vector_base(vector_base&&, const Alloc&) // is fixed and implemented - // Vector v5(std::move(v3), alloc2); + // Vector v5(_THRUST_STD::move(v3), alloc2); // ASSERT_EQ((v4 == v5), true); // ASSERT_EQ(v5.get_allocator(), alloc2); // ASSERT_EQ(Alloc::last_allocated, 1); @@ -249,7 +249,7 @@ void TestVectorAllocatorPropagateOnMoveAssignment() Vector v1(10, alloc1); Vector v2(15, alloc2); - v2 = std::move(v1); + v2 = _THRUST_STD::move(v1); ASSERT_EQ(v2.get_allocator(), alloc1); ASSERT_EQ(Alloc::last_allocated, 2); ASSERT_EQ(Alloc::last_deallocated, 2); diff --git a/projects/rocthrust/testing/vector.cu b/projects/rocthrust/testing/vector.cu index 673e4c9ac863..70071ba09761 100644 --- a/projects/rocthrust/testing/vector.cu +++ b/projects/rocthrust/testing/vector.cu @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -812,7 +812,7 @@ void TestVectorMove() const auto ptr1 = v1.data(); const auto size1 = v1.size(); - Vector v2(std::move(v1)); + Vector v2(_THRUST_STD::move(v1)); const auto ptr2 = v2.data(); const auto size2 = v2.size(); @@ -833,7 +833,7 @@ void TestVectorMove() const auto ptr3 = v3.data(); const auto size3 = v3.size(); - v2 = std::move(v3); + v2 = _THRUST_STD::move(v3); const auto ptr4 = v2.data(); const auto size4 = v2.size(); diff --git a/projects/rocthrust/testing/vector_allocators.cu b/projects/rocthrust/testing/vector_allocators.cu index c77d46cc0f3c..99067e32dea3 100644 --- a/projects/rocthrust/testing/vector_allocators.cu +++ b/projects/rocthrust/testing/vector_allocators.cu @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -50,7 +50,7 @@ public: } stateful_allocator(stateful_allocator&& other) - : BaseAlloc(std::move(other)) + : BaseAlloc(_THRUST_STD::move(other)) , state(other.state) { other.state = 0; @@ -170,7 +170,7 @@ void TestVectorAllocatorConstructors() // FIXME: uncomment this after the vector_base(vector_base&&, const Alloc&) // is fixed and implemented - // Vector v5(std::move(v3), alloc2); + // Vector v5(_THRUST_STD::move(v3), alloc2); // ASSERT_EQUAL((v4 == v5), true); // ASSERT_EQUAL(v5.get_allocator(), alloc2); // ASSERT_EQUAL(Alloc::last_allocated, 1); @@ -243,7 +243,7 @@ void TestVectorAllocatorPropagateOnMoveAssignment() Vector v1(10, alloc1); Vector v2(15, alloc2); - v2 = std::move(v1); + v2 = _THRUST_STD::move(v1); ASSERT_EQUAL(v2.get_allocator(), alloc1); ASSERT_EQUAL(Alloc::last_allocated, 2); ASSERT_EQUAL(Alloc::last_deallocated, 2); diff --git a/projects/rocthrust/thrust/allocate_unique.h b/projects/rocthrust/thrust/allocate_unique.h index 049ea54eb939..5448a7f51f08 100644 --- a/projects/rocthrust/thrust/allocate_unique.h +++ b/projects/rocthrust/thrust/allocate_unique.h @@ -47,7 +47,7 @@ struct allocator_delete final {} template allocator_delete(allocator_delete&& other) noexcept - : alloc_(std::move(other.get_allocator())) + : alloc_(_THRUST_STD::move(other.get_allocator())) {} template @@ -59,7 +59,7 @@ struct allocator_delete final template allocator_delete& operator=(allocator_delete&& other) noexcept { - alloc_ = std::move(other.get_allocator()); + alloc_ = _THRUST_STD::move(other.get_allocator()); return *this; } @@ -120,7 +120,7 @@ struct array_allocator_delete final {} template array_allocator_delete(array_allocator_delete&& other) noexcept - : alloc_(std::move(other.get_allocator())) + : alloc_(_THRUST_STD::move(other.get_allocator())) , count_(other.count_) {} @@ -134,7 +134,7 @@ struct array_allocator_delete final template array_allocator_delete& operator=(array_allocator_delete&& other) noexcept { - alloc_ = std::move(other.get_allocator()); + alloc_ = _THRUST_STD::move(other.get_allocator()); count_ = other.count_; return *this; } @@ -219,7 +219,7 @@ allocate_unique(Allocator const& alloc, Args&&... args) traits::construct(alloc_T, thrust::raw_pointer_cast(hold.get()), THRUST_FWD(args)...); auto deleter = allocator_delete(alloc); - return std::unique_ptr(hold.release(), std::move(deleter)); + return std::unique_ptr(hold.release(), _THRUST_STD::move(deleter)); } //! Creates a \p std::unique_ptr holding storage for a new object of type \p T without constructing it, using \p alloc @@ -243,7 +243,7 @@ uninitialized_allocate_unique(Allocator const& alloc) auto hold = hold_t(traits::allocate(alloc_T, 1), hold_deleter); auto deleter = uninitialized_allocator_delete(alloc_T); - return std::unique_ptr(hold.release(), std::move(deleter)); + return std::unique_ptr(hold.release(), _THRUST_STD::move(deleter)); } //! Creates a \p std::unique_ptr holding an array of objects of type \p T, each one constructed with \p args, using \p @@ -268,7 +268,7 @@ allocate_unique_n(Allocator const& alloc, Size n, Args&&... args) uninitialized_construct_n_with_allocator(alloc_T, hold.get(), n, THRUST_FWD(args)...); auto deleter = array_allocator_delete(alloc_T, n); - return std::unique_ptr(hold.release(), std::move(deleter)); + return std::unique_ptr(hold.release(), _THRUST_STD::move(deleter)); } //! Creates a \p std::unique_ptr holding storage for an array of objects of type \p T without constructing them, using @@ -292,7 +292,7 @@ uninitialized_allocate_unique_n(Allocator const& alloc, Size n) auto hold = hold_t(traits::allocate(alloc_T, n), hold_deleter); auto deleter = uninitialized_array_allocator_delete(alloc_T, n); - return std::unique_ptr(hold.release(), std::move(deleter)); + return std::unique_ptr(hold.release(), _THRUST_STD::move(deleter)); } /////////////////////////////////////////////////////////////////////////////// diff --git a/projects/rocthrust/thrust/detail/contiguous_storage.inl b/projects/rocthrust/thrust/detail/contiguous_storage.inl index 1040b4dac00b..e52805d345f2 100644 --- a/projects/rocthrust/thrust/detail/contiguous_storage.inl +++ b/projects/rocthrust/thrust/detail/contiguous_storage.inl @@ -35,6 +35,8 @@ #include // for std::runtime_error +#include _THRUST_STD_INCLUDE(utility) + THRUST_NAMESPACE_BEGIN namespace detail @@ -256,8 +258,8 @@ THRUST_HOST_DEVICE contiguous_storage& contiguous_storage::o m_allocator = _THRUST_STD::move(other.m_allocator); } - m_begin = std::move(other.m_begin); - m_size = std::move(other.m_size); + m_begin = _THRUST_STD::move(other.m_begin); + m_size = _THRUST_STD::move(other.m_size); other.m_begin = pointer(static_cast(0)); other.m_size = 0; diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/not_fn.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/not_fn.h index 7d17f9ae4489..f851a06fefd8 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/not_fn.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/not_fn.h @@ -101,10 +101,10 @@ struct not_fn_t bool cccl_true = true, ::std::enable_if_t && cccl_true, int> = 0> inline THRUST_HOST_DEVICE constexpr auto - operator()(Args&&... args) && noexcept(noexcept(!detail::invoke(::std::move(f), ::std::forward(args)...))) - -> decltype(!detail::invoke(::std::move(f), ::std::forward(args)...)) + operator()(Args&&... args) && noexcept(noexcept(!detail::invoke(_THRUST_STD::move(f), ::std::forward(args)...))) + -> decltype(!detail::invoke(_THRUST_STD::move(f), ::std::forward(args)...)) { - return !detail::invoke(::std::move(f), ::std::forward(args)...); + return !detail::invoke(_THRUST_STD::move(f), ::std::forward(args)...); } template && cccl_true, int> = 0> - inline THRUST_HOST_DEVICE constexpr auto - operator()(Args&&... args) const&& noexcept(noexcept(!detail::invoke(::std::move(f), ::std::forward(args)...))) - -> decltype(!detail::invoke(::std::move(f), ::std::forward(args)...)) + inline THRUST_HOST_DEVICE constexpr auto operator()(Args&&... args) const&& noexcept( + noexcept(!detail::invoke(_THRUST_STD::move(f), ::std::forward(args)...))) + -> decltype(!detail::invoke(_THRUST_STD::move(f), ::std::forward(args)...)) { - return !detail::invoke(::std::move(f), ::std::forward(args)...); + return !detail::invoke(_THRUST_STD::move(f), ::std::forward(args)...); } template ::vector_base(vector_base&& v) : m_storage(copy_allocator_t(), v.m_storage) , m_size(0) { - *this = std::move(v); + *this = _THRUST_STD::move(v); } // end vector_base::vector_base() template @@ -145,8 +145,8 @@ template vector_base& vector_base::operator=(vector_base&& v) { m_storage.destroy(begin(), end()); - m_storage = std::move(v.m_storage); - m_size = std::move(v.m_size); + m_storage = _THRUST_STD::move(v.m_storage); + m_size = _THRUST_STD::move(v.m_size); v.m_storage = contiguous_storage(copy_allocator_t(), m_storage); v.m_size = 0; diff --git a/projects/rocthrust/thrust/device_vector.h b/projects/rocthrust/thrust/device_vector.h index d9d5916a244b..a73b685e24b6 100644 --- a/projects/rocthrust/thrust/device_vector.h +++ b/projects/rocthrust/thrust/device_vector.h @@ -33,8 +33,9 @@ #include #include +#include _THRUST_STD_INCLUDE(utility) + #include -#include #include THRUST_NAMESPACE_BEGIN @@ -143,7 +144,7 @@ class device_vector : public detail::vector_base * \param v The device_vector to move. */ device_vector(device_vector&& v) - : Parent(std::move(v)) + : Parent(_THRUST_STD::move(v)) {} /*! Move constructor moves from another \p device_vector. @@ -151,7 +152,7 @@ class device_vector : public detail::vector_base * \param alloc The allocator to use by this device_vector. */ device_vector(device_vector&& v, const Alloc& alloc) - : Parent(std::move(v), alloc) + : Parent(_THRUST_STD::move(v), alloc) {} /*! Copy assign operator copies another \p device_vector with the same type. @@ -168,7 +169,7 @@ class device_vector : public detail::vector_base */ device_vector& operator=(device_vector&& v) { - Parent::operator=(std::move(v)); + Parent::operator=(_THRUST_STD::move(v)); return *this; } diff --git a/projects/rocthrust/thrust/host_vector.h b/projects/rocthrust/thrust/host_vector.h index 5b25792825a7..634fd7c21d55 100644 --- a/projects/rocthrust/thrust/host_vector.h +++ b/projects/rocthrust/thrust/host_vector.h @@ -33,8 +33,9 @@ #include #include +#include _THRUST_STD_INCLUDE(utility) + #include -#include #include THRUST_NAMESPACE_BEGIN @@ -144,7 +145,7 @@ class host_vector : public detail::vector_base * \param v The host_vector to move. */ THRUST_HOST host_vector(host_vector&& v) - : Parent(std::move(v)) + : Parent(_THRUST_STD::move(v)) {} /*! Move constructor moves from another host_vector. @@ -152,7 +153,7 @@ class host_vector : public detail::vector_base * \param alloc The allocator to use by this host_vector. */ THRUST_HOST host_vector(host_vector&& v, const Alloc& alloc) - : Parent(std::move(v), alloc) + : Parent(_THRUST_STD::move(v), alloc) {} /*! Assign operator copies from an exemplar \p host_vector. @@ -169,7 +170,7 @@ class host_vector : public detail::vector_base */ THRUST_HOST host_vector& operator=(host_vector&& v) { - Parent::operator=(std::move(v)); + Parent::operator=(_THRUST_STD::move(v)); return *this; } diff --git a/projects/rocthrust/thrust/system/cpp/detail/vector.inl b/projects/rocthrust/thrust/system/cpp/detail/vector.inl index 0e3c2626ec0b..c1fedeefe200 100644 --- a/projects/rocthrust/thrust/system/cpp/detail/vector.inl +++ b/projects/rocthrust/thrust/system/cpp/detail/vector.inl @@ -27,7 +27,7 @@ #endif // no system header #include -#include +#include _THRUST_STD_INCLUDE(utility) THRUST_NAMESPACE_BEGIN namespace system @@ -57,7 +57,7 @@ vector::vector(const vector& x) template vector::vector(vector&& x) - : super_t(std::move(x)) + : super_t(_THRUST_STD::move(x)) {} template @@ -88,7 +88,7 @@ vector& vector::operator=(const vector& x) template vector& vector::operator=(vector&& x) { - super_t::operator=(std::move(x)); + super_t::operator=(_THRUST_STD::move(x)); return *this; } diff --git a/projects/rocthrust/thrust/system/detail/generic/sequence.inl b/projects/rocthrust/thrust/system/detail/generic/sequence.inl index 0cbe55c0bbcd..cfdb88fb62ba 100644 --- a/projects/rocthrust/thrust/system/detail/generic/sequence.inl +++ b/projects/rocthrust/thrust/system/detail/generic/sequence.inl @@ -85,7 +85,7 @@ template THRUST_HOST_DEVICE void sequence(thrust::execution_policy& exec, ForwardIterator first, ForwardIterator last, T init, T step) { - thrust::tabulate(exec, first, last, detail::compute_sequence_value{std::move(init), std::move(step)}); + thrust::tabulate(exec, first, last, detail::compute_sequence_value{_THRUST_STD::move(init), _THRUST_STD::move(step)}); } // end sequence() } // end namespace generic diff --git a/projects/rocthrust/thrust/unique_ptr.h b/projects/rocthrust/thrust/unique_ptr.h index 4444c2b7ec73..a5959ce799a1 100644 --- a/projects/rocthrust/thrust/unique_ptr.h +++ b/projects/rocthrust/thrust/unique_ptr.h @@ -471,7 +471,7 @@ class unique_ptr template >> THRUST_HOST THRUST_CONSTEXPR_CXX23 unique_ptr(pointer p, GoodRValRefType d) noexcept : m_ptr(p) - , m_deleter(std::move(d)) + , m_deleter(_THRUST_STD::move(d)) { static_assert(!std::is_reference_v, "rvalue deleter bound to reference"); } @@ -911,7 +911,7 @@ class unique_ptr class = EnableIfPointerConvertible> THRUST_HOST THRUST_CONSTEXPR_CXX23 unique_ptr(Pp p, GoodRValRefType deleter) noexcept : m_ptr(p) - , m_deleter(std::move(deleter)) + , m_deleter(_THRUST_STD::move(deleter)) { static_assert(!std::is_reference_v, "rvalue deleter bound to reference"); } @@ -923,7 +923,7 @@ class unique_ptr template >> THRUST_HOST THRUST_CONSTEXPR_CXX23 unique_ptr(std::nullptr_t, GoodRValRefType deleter) noexcept : m_ptr(nullptr) - , m_deleter(std::move(deleter)) + , m_deleter(_THRUST_STD::move(deleter)) { static_assert(!std::is_reference_v, "rvalue deleter bound to reference"); } diff --git a/projects/rocthrust/thrust/zip_function.h b/projects/rocthrust/thrust/zip_function.h index 6f5b9f513ef8..4c369b6413ef 100644 --- a/projects/rocthrust/thrust/zip_function.h +++ b/projects/rocthrust/thrust/zip_function.h @@ -150,7 +150,7 @@ class zip_function /*! Constructs a \p zip_function with the provided function object \p func. */ THRUST_HOST_DEVICE zip_function(Function func) - : func(std::move(func)) + : func(_THRUST_STD::move(func)) {} /*! Applies the N-ary function object to elements of the tuple \p args. */ From e043e5522171998d0676b2fd5d5a012d98362ab0 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Fri, 12 Jun 2026 11:34:16 +0200 Subject: [PATCH 09/86] Use libhipcxx includes - limits, cmath, cfloat, cstdint, cstddef, cstring --- projects/rocthrust/CHANGELOG.md | 2 +- .../detail/allocator/tagged_allocator.inl | 4 +- .../thrust/detail/complex/arithmetic.h | 4 +- .../rocthrust/thrust/detail/complex/c99math.h | 206 ++---------------- .../rocthrust/thrust/detail/complex/catrig.h | 7 +- .../rocthrust/thrust/detail/complex/catrigf.h | 7 +- .../rocthrust/thrust/detail/complex/cproj.h | 11 +- .../rocthrust/thrust/detail/complex/csinh.h | 4 +- .../rocthrust/thrust/detail/complex/csinhf.h | 5 +- .../rocthrust/thrust/detail/complex/csqrt.h | 5 +- .../rocthrust/thrust/detail/complex/csqrtf.h | 5 +- .../rocthrust/thrust/detail/complex/ctanh.h | 2 +- .../rocthrust/thrust/detail/complex/ctanhf.h | 2 +- .../thrust/detail/complex/math_private.h | 2 +- .../thrust/detail/config/memory_resource.h | 2 +- .../rocthrust/thrust/detail/integer_math.h | 3 +- projects/rocthrust/thrust/detail/pointer.h | 2 +- projects/rocthrust/thrust/detail/util/align.h | 2 +- projects/rocthrust/thrust/device_allocator.h | 3 +- projects/rocthrust/thrust/device_malloc.h | 2 +- .../thrust/device_malloc_allocator.h | 5 +- projects/rocthrust/thrust/device_new.h | 2 +- projects/rocthrust/thrust/mr/allocator.h | 4 +- projects/rocthrust/thrust/mr/pool_options.h | 2 +- projects/rocthrust/thrust/random.h | 2 +- .../rocthrust/thrust/random/detail/erfcinv.h | 10 +- .../linear_congruential_engine_discard.h | 2 +- .../random/detail/normal_distribution.inl | 3 +- .../random/detail/normal_distribution_base.h | 12 +- .../random/detail/xor_combine_engine_max.h | 7 +- .../thrust/random/discard_block_engine.h | 3 +- .../random/linear_congruential_engine.h | 3 +- .../random/linear_feedback_shift_engine.h | 3 +- .../random/subtract_with_carry_engine.h | 5 +- .../thrust/random/xor_combine_engine.h | 2 +- .../thrust/system/cpp/detail/memory.inl | 2 +- .../system/cuda/detail/adjacent_difference.h | 3 +- .../thrust/system/cuda/detail/copy_if.h | 2 +- .../thrust/system/cuda/detail/dispatch.h | 2 +- .../thrust/system/cuda/detail/extrema.h | 3 +- .../thrust/system/cuda/detail/memory.inl | 2 +- .../thrust/system/cuda/detail/merge.h | 2 +- .../thrust/system/cuda/detail/partition.h | 2 +- .../thrust/system/cuda/detail/reduce.h | 2 +- .../thrust/system/cuda/detail/reduce_by_key.h | 3 +- .../thrust/system/cuda/detail/scan.h | 2 +- .../thrust/system/cuda/detail/scan_by_key.h | 2 +- .../system/cuda/detail/set_operations.h | 3 +- .../thrust/system/cuda/detail/sort.h | 3 +- .../thrust/system/cuda/detail/transform.h | 3 +- .../system/cuda/detail/transform_reduce.h | 3 +- .../thrust/system/cuda/detail/unique.h | 2 +- .../thrust/system/cuda/detail/unique_by_key.h | 2 +- .../system/detail/generic/reduce_by_key.inl | 2 +- .../system/detail/generic/scan_by_key.inl | 2 +- .../thrust/system/detail/generic/shuffle.inl | 2 +- .../detail/sequential/malloc_and_free.h | 6 +- .../detail/sequential/stable_radix_sort.inl | 6 +- .../system/detail/sequential/trivial_copy.h | 14 +- .../thrust/system/hip/detail/binary_search.h | 2 +- .../thrust/system/hip/detail/copy_if.h | 2 +- .../thrust/system/hip/detail/dispatch.h | 2 +- .../thrust/system/hip/detail/extrema.h | 2 +- .../system/hip/detail/general/temp_storage.h | 4 +- .../thrust/system/hip/detail/memory.inl | 2 +- .../thrust/system/hip/detail/merge.h | 2 +- .../thrust/system/hip/detail/partition.h | 2 +- .../thrust/system/hip/detail/reduce.h | 2 +- .../thrust/system/hip/detail/reduce_by_key.h | 2 +- .../rocthrust/thrust/system/hip/detail/scan.h | 2 +- .../thrust/system/hip/detail/scan_by_key.h | 2 +- .../thrust/system/hip/detail/set_operations.h | 2 +- .../rocthrust/thrust/system/hip/detail/sort.h | 3 +- .../thrust/system/hip/detail/transform.h | 2 +- .../system/hip/detail/transform_reduce.h | 3 +- .../thrust/system/hip/detail/unique.h | 2 +- .../thrust/system/hip/detail/unique_by_key.h | 2 +- .../thrust/system/omp/detail/memory.inl | 2 +- .../thrust/system/tbb/detail/memory.inl | 2 +- 79 files changed, 151 insertions(+), 313 deletions(-) diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index b122e72cf19c..5296df49cb2a 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -7,7 +7,7 @@ Documentation for rocThrust available at ### Changed -* Replaced `std::move` operators to `hip::std` operators. +* Replaced cmath, cfloat, cstdint, cstddef, cstring, move, limits operators to `hip::std` operators. ### Removed diff --git a/projects/rocthrust/thrust/detail/allocator/tagged_allocator.inl b/projects/rocthrust/thrust/detail/allocator/tagged_allocator.inl index 601be80aa2cb..dee43ce08445 100644 --- a/projects/rocthrust/thrust/detail/allocator/tagged_allocator.inl +++ b/projects/rocthrust/thrust/detail/allocator/tagged_allocator.inl @@ -27,7 +27,7 @@ #endif // no system header #include -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace detail @@ -67,7 +67,7 @@ tagged_allocator::address(const_reference x) const template typename tagged_allocator::size_type tagged_allocator::max_size() const { - return (std::numeric_limits::max)() / sizeof(T); + return (_THRUST_STD::numeric_limits::max)() / sizeof(T); } template diff --git a/projects/rocthrust/thrust/detail/complex/arithmetic.h b/projects/rocthrust/thrust/detail/complex/arithmetic.h index 78451401d132..ae1b209b521b 100644 --- a/projects/rocthrust/thrust/detail/complex/arithmetic.h +++ b/projects/rocthrust/thrust/detail/complex/arithmetic.h @@ -24,8 +24,8 @@ #include #include -#include -#include +#include _THRUST_STD_INCLUDE(cfloat) +#include _THRUST_STD_INCLUDE(cmath) THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/detail/complex/c99math.h b/projects/rocthrust/thrust/detail/complex/c99math.h index 37b8bc820331..4e72976dfe22 100644 --- a/projects/rocthrust/thrust/detail/complex/c99math.h +++ b/projects/rocthrust/thrust/detail/complex/c99math.h @@ -21,7 +21,7 @@ #include -#include +#include _THRUST_STD_INCLUDE(cmath) #include @@ -38,191 +38,25 @@ namespace complex // Some platforms define these as macros, others as free functions. // Avoid using the std:: form of these as nvcc may treat std::foo() as __host__ functions. -using ::acos; -using ::asin; -using ::atan; -using ::cos; -using ::cosh; -using ::exp; -using ::log; -using ::sin; -using ::sinh; -using ::sqrt; -using ::tan; - -template -inline THRUST_HOST_DEVICE T infinity(); - -template <> -inline THRUST_HOST_DEVICE float infinity() -{ - float res; - set_float_word(res, 0x7f800000); - return res; -} - -template <> -inline THRUST_HOST_DEVICE double infinity() -{ - double res; - insert_words(res, 0x7ff00000, 0); - return res; -} - -#if defined _MSC_VER -THRUST_HOST_DEVICE inline int isinf(float x) -{ - return std::abs(x) == infinity(); -} - -THRUST_HOST_DEVICE inline int isinf(double x) -{ - return std::abs(x) == infinity(); -} - -THRUST_HOST_DEVICE inline int isnan(float x) -{ - return x != x; -} - -THRUST_HOST_DEVICE inline int isnan(double x) -{ - return x != x; -} - -THRUST_HOST_DEVICE inline int signbit(float x) -{ - return ((*((uint32_t*) &x)) & 0x80000000) != 0 ? 1 : 0; -} - -THRUST_HOST_DEVICE inline int signbit(double x) -{ - return ((*((uint64_t*) &x)) & 0x8000000000000000) != 0ull ? 1 : 0; -} - -THRUST_HOST_DEVICE inline int isfinite(float x) -{ - return !isnan(x) && !isinf(x); -} - -THRUST_HOST_DEVICE inline int isfinite(double x) -{ - return !isnan(x) && !isinf(x); -} - -#else - -# if (defined(__CUDACC__) && !(defined(__CUDA__) && defined(__clang__)) && !defined(_NVHPC_CUDA)) || defined(__HIP__) -// NVCC implements at least some signature of these as functions not macros. -using ::isfinite; -using ::isinf; -using ::isnan; -using ::signbit; -# else -// Some compilers do not provide these in the global scope, because they are -// supposed to be macros. The versions in `std` are supposed to be functions. -// Since we're not compiling with nvcc, it's safe to use the functions in std:: -using std::isfinite; -using std::isinf; -using std::isnan; -using std::signbit; -# endif // __CUDACC__|| __HIP__ -#endif // _MSC_VER - -using ::atanh; - -#if defined _MSC_VER - -THRUST_HOST_DEVICE inline double copysign(double x, double y) -{ - uint32_t hx, hy; - get_high_word(hx, x); - get_high_word(hy, y); - set_high_word(x, (hx & 0x7fffffff) | (hy & 0x80000000)); - return x; -} - -THRUST_HOST_DEVICE inline float copysignf(float x, float y) -{ - uint32_t ix, iy; - get_float_word(ix, x); - get_float_word(iy, y); - set_float_word(x, (ix & 0x7fffffff) | (iy & 0x80000000)); - return x; -} - -# if !defined(__CUDACC__) && !defined(_NVHPC_CUDA) - -// Simple approximation to log1p as Visual Studio is lacking one -# if THRUST_HAS_HIP_COMPILER() && defined(_MSC_VER) // fix HIP on Windows error -THRUST_HOST_DEVICE -# endif -inline double -log1p(double x) -{ - double u = 1.0 + x; - if (u == 1.0) - { - return x; - } - else - { - if (u > 2.0) - { - // Use normal log for large arguments - return log(u); - } - else - { - return log(u) * (x / (u - 1.0)); - } - } -} - -# if THRUST_HAS_HIP_COMPILER() && defined(_MSC_VER) // fix HIP on Windows error -THRUST_HOST_DEVICE -# endif -inline float -log1pf(float x) -{ - float u = 1.0f + x; - if (u == 1.0f) - { - return x; - } - else - { - if (u > 2.0f) - { - // Use normal log for large arguments - return logf(u); - } - else - { - return logf(u) * (x / (u - 1.0f)); - } - } -} - -// add !THRUST_HAS_HIP_COMPILER() to fix HIP on Windows error -# if _MSV_VER <= 1500 && !THRUST_HAS_HIP_COMPILER() -# include - -inline float hypotf(float x, float y) -{ - return abs(std::complex(x, y)); -} - -inline double hypot(double x, double y) -{ - return _hypot(x, y); -} - -# endif // _MSC_VER <= 1500 && !THRUST_HAS_HIP_COMPILER() - -# endif // __CUDACC__ - -#endif // _MSC_VER +using _THRUST_STD::acos; +using _THRUST_STD::asin; +using _THRUST_STD::atan; +using _THRUST_STD::atanh; +using _THRUST_STD::copysign; +using _THRUST_STD::cos; +using _THRUST_STD::cosh; +using _THRUST_STD::exp; +using _THRUST_STD::hypot; +using _THRUST_STD::isfinite; +using _THRUST_STD::isinf; +using _THRUST_STD::isnan; +using _THRUST_STD::log; +using _THRUST_STD::log1p; +using _THRUST_STD::signbit; +using _THRUST_STD::sin; +using _THRUST_STD::sinh; +using _THRUST_STD::sqrt; +using _THRUST_STD::tan; } // namespace complex diff --git a/projects/rocthrust/thrust/detail/complex/catrig.h b/projects/rocthrust/thrust/detail/complex/catrig.h index 8cc984c22f19..08e32bf92112 100644 --- a/projects/rocthrust/thrust/detail/complex/catrig.h +++ b/projects/rocthrust/thrust/detail/complex/catrig.h @@ -53,8 +53,9 @@ #include #include -#include -#include +#include _THRUST_STD_INCLUDE(cfloat) +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace detail @@ -430,7 +431,7 @@ THRUST_HOST_DEVICE inline complex cacos(complex z) /* cacos(+-Inf + I*NaN) = NaN + I*opt(-)Inf */ if (isinf(x)) { - return (complex(y + y, -infinity())); + return (complex(y + y, -_THRUST_STD::numeric_limits::infinity())); } /* cacos(NaN + I*+-Inf) = NaN + I*-+Inf */ if (isinf(y)) diff --git a/projects/rocthrust/thrust/detail/complex/catrigf.h b/projects/rocthrust/thrust/detail/complex/catrigf.h index d847e6dcc1ce..91c8cb29daa2 100644 --- a/projects/rocthrust/thrust/detail/complex/catrigf.h +++ b/projects/rocthrust/thrust/detail/complex/catrigf.h @@ -53,8 +53,9 @@ #include #include -#include -#include +#include _THRUST_STD_INCLUDE(cfloat) +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace detail @@ -278,7 +279,7 @@ THRUST_HOST_DEVICE inline complex cacosf(complex z) { if (isinf(x)) { - return (complex(y + y, -infinity())); + return (complex(y + y, -_THRUST_STD::numeric_limits::infinity())); } if (isinf(y)) { diff --git a/projects/rocthrust/thrust/detail/complex/cproj.h b/projects/rocthrust/thrust/detail/complex/cproj.h index 7cda5804d12a..41f6a7933347 100644 --- a/projects/rocthrust/thrust/detail/complex/cproj.h +++ b/projects/rocthrust/thrust/detail/complex/cproj.h @@ -22,7 +22,8 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace detail @@ -37,8 +38,8 @@ THRUST_HOST_DEVICE inline complex cprojf(const complex& z) } else { - // std::numeric_limits::infinity() doesn't run on the GPU - return complex(infinity(), copysignf(0.0, z.imag())); + // _THRUST_STD::numeric_limits::infinity() doesn't run on the GPU + return complex(_THRUST_STD::numeric_limits::infinity(), copysignf(0.0, z.imag())); } } @@ -50,8 +51,8 @@ THRUST_HOST_DEVICE inline complex cproj(const complex& z) } else { - // std::numeric_limits::infinity() doesn't run on the GPU - return complex(infinity(), copysign(0.0, z.imag())); + // _THRUST_STD::numeric_limits::infinity() doesn't run on the GPU + return complex(_THRUST_STD::numeric_limits::infinity(), copysign(0.0, z.imag())); } } diff --git a/projects/rocthrust/thrust/detail/complex/csinh.h b/projects/rocthrust/thrust/detail/complex/csinh.h index dfbfa4e03831..10f9a3f8f9a5 100644 --- a/projects/rocthrust/thrust/detail/complex/csinh.h +++ b/projects/rocthrust/thrust/detail/complex/csinh.h @@ -53,6 +53,8 @@ #include #include +#include _THRUST_STD_INCLUDE(limits) + THRUST_NAMESPACE_BEGIN namespace detail { @@ -167,7 +169,7 @@ THRUST_HOST_DEVICE inline complex csinh(const complex& z) { return (complex(x * x, x * (y - y))); } - return (complex(x * cos(y), infinity() * sin(y))); + return (complex(x * cos(y), _THRUST_STD::numeric_limits::infinity() * sin(y))); } /* diff --git a/projects/rocthrust/thrust/detail/complex/csinhf.h b/projects/rocthrust/thrust/detail/complex/csinhf.h index 62fe0abe41a4..a4e870bdab73 100644 --- a/projects/rocthrust/thrust/detail/complex/csinhf.h +++ b/projects/rocthrust/thrust/detail/complex/csinhf.h @@ -54,6 +54,9 @@ #include #include +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(limits) + THRUST_NAMESPACE_BEGIN namespace detail { @@ -135,7 +138,7 @@ THRUST_HOST_DEVICE inline complex csinhf(const complex& z) { return (complex(x * x, x * (y - y))); } - return (complex(x * cosf(y), infinity() * sinf(y))); + return (complex(x * cosf(y), _THRUST_STD::numeric_limits::infinity() * sinf(y))); } return (complex((x * x) * (y - y), (x + x) * (y - y))); diff --git a/projects/rocthrust/thrust/detail/complex/csqrt.h b/projects/rocthrust/thrust/detail/complex/csqrt.h index fa8c95c57903..731e2f673f48 100644 --- a/projects/rocthrust/thrust/detail/complex/csqrt.h +++ b/projects/rocthrust/thrust/detail/complex/csqrt.h @@ -53,7 +53,8 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace detail @@ -83,7 +84,7 @@ THRUST_HOST_DEVICE inline complex csqrt(const complex& z) } if (isinf(b)) { - return (complex(infinity(), b)); + return (complex(_THRUST_STD::numeric_limits::infinity(), b)); } if (isnan(a)) { diff --git a/projects/rocthrust/thrust/detail/complex/csqrtf.h b/projects/rocthrust/thrust/detail/complex/csqrtf.h index 9dbfd75cafad..fda6183dbfd3 100644 --- a/projects/rocthrust/thrust/detail/complex/csqrtf.h +++ b/projects/rocthrust/thrust/detail/complex/csqrtf.h @@ -53,7 +53,8 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace detail @@ -80,7 +81,7 @@ THRUST_HOST_DEVICE inline complex csqrtf(const complex& z) } if (isinf(b)) { - return (complex(infinity(), b)); + return (complex(_THRUST_STD::numeric_limits::infinity(), b)); } if (isnan(a)) { diff --git a/projects/rocthrust/thrust/detail/complex/ctanh.h b/projects/rocthrust/thrust/detail/complex/ctanh.h index 0798ac895031..948cbd46a1d8 100644 --- a/projects/rocthrust/thrust/detail/complex/ctanh.h +++ b/projects/rocthrust/thrust/detail/complex/ctanh.h @@ -92,7 +92,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cmath) THRUST_NAMESPACE_BEGIN namespace detail diff --git a/projects/rocthrust/thrust/detail/complex/ctanhf.h b/projects/rocthrust/thrust/detail/complex/ctanhf.h index 631f3aa1b5b9..ca397c599f9b 100644 --- a/projects/rocthrust/thrust/detail/complex/ctanhf.h +++ b/projects/rocthrust/thrust/detail/complex/ctanhf.h @@ -57,7 +57,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cmath) THRUST_NAMESPACE_BEGIN namespace detail diff --git a/projects/rocthrust/thrust/detail/complex/math_private.h b/projects/rocthrust/thrust/detail/complex/math_private.h index b65849e7978d..9653428b6dd9 100644 --- a/projects/rocthrust/thrust/detail/complex/math_private.h +++ b/projects/rocthrust/thrust/detail/complex/math_private.h @@ -35,7 +35,7 @@ #include -#include +#include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN namespace detail diff --git a/projects/rocthrust/thrust/detail/config/memory_resource.h b/projects/rocthrust/thrust/detail/config/memory_resource.h index 8b5b7506bdfb..eb3fee562e97 100644 --- a/projects/rocthrust/thrust/detail/config/memory_resource.h +++ b/projects/rocthrust/thrust/detail/config/memory_resource.h @@ -28,7 +28,7 @@ #include -#include +#include _THRUST_STD_INCLUDE(cstddef) #define THRUST_MR_DEFAULT_ALIGNMENT alignof(THRUST_NS_QUALIFIER::detail::max_align_t) diff --git a/projects/rocthrust/thrust/detail/integer_math.h b/projects/rocthrust/thrust/detail/integer_math.h index 06d6aa02d727..d046a7c22703 100644 --- a/projects/rocthrust/thrust/detail/integer_math.h +++ b/projects/rocthrust/thrust/detail/integer_math.h @@ -28,10 +28,9 @@ #endif // no system header #include +#include _THRUST_STD_INCLUDE(limits) #include _THRUST_STD_INCLUDE(type_traits) -#include - THRUST_NAMESPACE_BEGIN namespace detail { diff --git a/projects/rocthrust/thrust/detail/pointer.h b/projects/rocthrust/thrust/detail/pointer.h index 747b4234f3d8..acba58f795c0 100644 --- a/projects/rocthrust/thrust/detail/pointer.h +++ b/projects/rocthrust/thrust/detail/pointer.h @@ -38,9 +38,9 @@ #include #include +#include _THRUST_STD_INCLUDE(cstddef) #include _THRUST_STD_INCLUDE(type_traits) -#include #include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/detail/util/align.h b/projects/rocthrust/thrust/detail/util/align.h index b51fe0c774fd..11adae15daf0 100644 --- a/projects/rocthrust/thrust/detail/util/align.h +++ b/projects/rocthrust/thrust/detail/util/align.h @@ -26,7 +26,7 @@ # pragma system_header #endif // no system header -#include +#include _THRUST_STD_INCLUDE(cstdint) // functions to handle memory alignment diff --git a/projects/rocthrust/thrust/device_allocator.h b/projects/rocthrust/thrust/device_allocator.h index dbb486c28138..e97e8201aa5c 100644 --- a/projects/rocthrust/thrust/device_allocator.h +++ b/projects/rocthrust/thrust/device_allocator.h @@ -34,8 +34,7 @@ #include #include -#include -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/device_malloc.h b/projects/rocthrust/thrust/device_malloc.h index bc2b7c9a829f..eb08ee7d6eb2 100644 --- a/projects/rocthrust/thrust/device_malloc.h +++ b/projects/rocthrust/thrust/device_malloc.h @@ -31,7 +31,7 @@ #endif // no system header #include -#include // for std::size_t +#include _THRUST_STD_INCLUDE(cstddef) // for std::size_t THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/device_malloc_allocator.h b/projects/rocthrust/thrust/device_malloc_allocator.h index 1098702dbf25..70fe54879327 100644 --- a/projects/rocthrust/thrust/device_malloc_allocator.h +++ b/projects/rocthrust/thrust/device_malloc_allocator.h @@ -34,7 +34,8 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(limits) + #include THRUST_NAMESPACE_BEGIN @@ -165,7 +166,7 @@ class device_malloc_allocator */ inline size_type max_size() const { - return (std::numeric_limits::max)() / sizeof(T); + return (_THRUST_STD::numeric_limits::max)() / sizeof(T); } // end max_size() /*! Compares against another \p device_malloc_allocator for equality. diff --git a/projects/rocthrust/thrust/device_new.h b/projects/rocthrust/thrust/device_new.h index dc1dd19054e2..b236511acac4 100644 --- a/projects/rocthrust/thrust/device_new.h +++ b/projects/rocthrust/thrust/device_new.h @@ -33,7 +33,7 @@ // #include this for size_t #include -#include +#include _THRUST_STD_INCLUDE(cstddef) THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/mr/allocator.h b/projects/rocthrust/thrust/mr/allocator.h index 1e2ab779289a..2c62fd18801c 100644 --- a/projects/rocthrust/thrust/mr/allocator.h +++ b/projects/rocthrust/thrust/mr/allocator.h @@ -37,7 +37,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace mr @@ -105,7 +105,7 @@ class allocator : private validator THRUST_EXEC_CHECK_DISABLE THRUST_HOST_DEVICE size_type max_size() const { - return (std::numeric_limits::max)() / sizeof(T); + return (_THRUST_STD::numeric_limits::max)() / sizeof(T); } /*! Constructor. diff --git a/projects/rocthrust/thrust/mr/pool_options.h b/projects/rocthrust/thrust/mr/pool_options.h index ed79fd08650e..3de8f0d29cf4 100644 --- a/projects/rocthrust/thrust/mr/pool_options.h +++ b/projects/rocthrust/thrust/mr/pool_options.h @@ -34,7 +34,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cstddef) THRUST_NAMESPACE_BEGIN namespace mr diff --git a/projects/rocthrust/thrust/random.h b/projects/rocthrust/thrust/random.h index d6c4e3149851..911fd0aac268 100644 --- a/projects/rocthrust/thrust/random.h +++ b/projects/rocthrust/thrust/random.h @@ -29,7 +29,7 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header -#include +#include _THRUST_STD_INCLUDE(cstdint) // RNGs #include diff --git a/projects/rocthrust/thrust/random/detail/erfcinv.h b/projects/rocthrust/thrust/random/detail/erfcinv.h index 1e6a6bceb4e2..250b03dc1e95 100644 --- a/projects/rocthrust/thrust/random/detail/erfcinv.h +++ b/projects/rocthrust/thrust/random/detail/erfcinv.h @@ -1,6 +1,6 @@ /* Copyright © 2003-2013 SciPy Developers. -Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. +Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: @@ -21,8 +21,8 @@ IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #pragma once -#include -#include +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(limits) /* sqrt(2pi) */ static double s2pi = 2.50662827463100050242E0; @@ -196,11 +196,11 @@ static inline double erfcinv(double x) } if (x == 0.0) { - return std::numeric_limits::infinity(); + return _THRUST_STD::numeric_limits::infinity(); } if (x == 2.0) { - return -std::numeric_limits::infinity(); + return -_THRUST_STD::numeric_limits::infinity(); } return -ndtri(0.5 * x) * one_o_sqrt2; } diff --git a/projects/rocthrust/thrust/random/detail/linear_congruential_engine_discard.h b/projects/rocthrust/thrust/random/detail/linear_congruential_engine_discard.h index 37de89c5b49c..5c9aacecb06a 100644 --- a/projects/rocthrust/thrust/random/detail/linear_congruential_engine_discard.h +++ b/projects/rocthrust/thrust/random/detail/linear_congruential_engine_discard.h @@ -28,7 +28,7 @@ #include -#include +#include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/random/detail/normal_distribution.inl b/projects/rocthrust/thrust/random/detail/normal_distribution.inl index 0713eaa8bd3c..eb2cbe28f8d6 100644 --- a/projects/rocthrust/thrust/random/detail/normal_distribution.inl +++ b/projects/rocthrust/thrust/random/detail/normal_distribution.inl @@ -29,11 +29,10 @@ #include #include +#include _THRUST_STD_INCLUDE(cstdint) // Use rocprim::numeric_limits if thrust/detail/type_traits.h uses rocprim::arithmetic #include _THRUST_STD_INCLUDE(limits) -#include - THRUST_NAMESPACE_BEGIN namespace random diff --git a/projects/rocthrust/thrust/random/detail/normal_distribution_base.h b/projects/rocthrust/thrust/random/detail/normal_distribution_base.h index d6662a7ecccc..277b7ab6d2e6 100644 --- a/projects/rocthrust/thrust/random/detail/normal_distribution_base.h +++ b/projects/rocthrust/thrust/random/detail/normal_distribution_base.h @@ -36,8 +36,8 @@ #include #include -#include -#include +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(limits) #if THRUST_HAS_HIP_COMPILER() # include @@ -119,10 +119,10 @@ class normal_distribution_portable { // implementation from Boost // allow for Koenig lookup - using std::cos; - using std::log; - using std::sin; - using std::sqrt; + using _THRUST_STD::cos; + using _THRUST_STD::log; + using _THRUST_STD::sin; + using _THRUST_STD::sqrt; if (!m_valid) { diff --git a/projects/rocthrust/thrust/random/detail/xor_combine_engine_max.h b/projects/rocthrust/thrust/random/detail/xor_combine_engine_max.h index f232d26288a3..60102e716c5e 100644 --- a/projects/rocthrust/thrust/random/detail/xor_combine_engine_max.h +++ b/projects/rocthrust/thrust/random/detail/xor_combine_engine_max.h @@ -29,8 +29,9 @@ #include #include +#include _THRUST_STD_INCLUDE(limits) + #include -#include THRUST_NAMESPACE_BEGIN @@ -69,7 +70,7 @@ struct lshift_w }; template -struct lshift : lshift_w::digits, lhs, rhs> +struct lshift : lshift_w::digits, lhs, rhs> {}; template @@ -193,7 +194,7 @@ struct xor_combine_engine_max_aux : xor_combine_engine_max_aux_1 struct xor_combine_engine_max { - static const size_t w = std::numeric_limits::digits; + static const size_t w = _THRUST_STD::numeric_limits::digits; static const result_type m1 = math:: min::value - 1>::value; diff --git a/projects/rocthrust/thrust/random/discard_block_engine.h b/projects/rocthrust/thrust/random/discard_block_engine.h index 526d09bbf69a..7cfc23807c71 100644 --- a/projects/rocthrust/thrust/random/discard_block_engine.h +++ b/projects/rocthrust/thrust/random/discard_block_engine.h @@ -33,7 +33,8 @@ #include -#include +#include _THRUST_STD_INCLUDE(cstdint) + #include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/random/linear_congruential_engine.h b/projects/rocthrust/thrust/random/linear_congruential_engine.h index b838d7f55d0e..08e7992849f6 100644 --- a/projects/rocthrust/thrust/random/linear_congruential_engine.h +++ b/projects/rocthrust/thrust/random/linear_congruential_engine.h @@ -32,7 +32,8 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cstdint) + #include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/random/linear_feedback_shift_engine.h b/projects/rocthrust/thrust/random/linear_feedback_shift_engine.h index 327236869730..e28349ee9426 100644 --- a/projects/rocthrust/thrust/random/linear_feedback_shift_engine.h +++ b/projects/rocthrust/thrust/random/linear_feedback_shift_engine.h @@ -41,7 +41,8 @@ #include #include -#include // for size_t +#include _THRUST_STD_INCLUDE(cstddef) // for size_t + #include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/random/subtract_with_carry_engine.h b/projects/rocthrust/thrust/random/subtract_with_carry_engine.h index 0a6c9b26bac9..e68a1c6eb07c 100644 --- a/projects/rocthrust/thrust/random/subtract_with_carry_engine.h +++ b/projects/rocthrust/thrust/random/subtract_with_carry_engine.h @@ -33,8 +33,9 @@ #include -#include // for size_t -#include +#include _THRUST_STD_INCLUDE(cstddef) // for size_t +#include _THRUST_STD_INCLUDE(cstdint) + #include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/random/xor_combine_engine.h b/projects/rocthrust/thrust/random/xor_combine_engine.h index 9bf9e568dfca..7e32c6c83173 100644 --- a/projects/rocthrust/thrust/random/xor_combine_engine.h +++ b/projects/rocthrust/thrust/random/xor_combine_engine.h @@ -36,7 +36,7 @@ #include #include -#include // for size_t +#include _THRUST_STD_INCLUDE(cstddef) // for size_t #include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/cpp/detail/memory.inl b/projects/rocthrust/thrust/system/cpp/detail/memory.inl index ab6976b48a7e..e5acb8ff9ff4 100644 --- a/projects/rocthrust/thrust/system/cpp/detail/memory.inl +++ b/projects/rocthrust/thrust/system/cpp/detail/memory.inl @@ -28,7 +28,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace system diff --git a/projects/rocthrust/thrust/system/cuda/detail/adjacent_difference.h b/projects/rocthrust/thrust/system/cuda/detail/adjacent_difference.h index 4830f7fd3f32..e3d8846cc57c 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/adjacent_difference.h +++ b/projects/rocthrust/thrust/system/cuda/detail/adjacent_difference.h @@ -53,10 +53,9 @@ # include # include +# include # include -# include - THRUST_NAMESPACE_BEGIN template diff --git a/projects/rocthrust/thrust/system/cuda/detail/copy_if.h b/projects/rocthrust/thrust/system/cuda/detail/copy_if.h index 2ee5bd8947d7..7152008395f4 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/copy_if.h +++ b/projects/rocthrust/thrust/system/cuda/detail/copy_if.h @@ -56,7 +56,7 @@ # include # include -# include +# include THRUST_NAMESPACE_BEGIN // XXX declare generic copy_if interface diff --git a/projects/rocthrust/thrust/system/cuda/detail/dispatch.h b/projects/rocthrust/thrust/system/cuda/detail/dispatch.h index c161d9c3cf15..274f209d8e58 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/dispatch.h +++ b/projects/rocthrust/thrust/system/cuda/detail/dispatch.h @@ -29,11 +29,11 @@ #include #include +#include #include #include #include -#include #include #if defined(THRUST_FORCE_32_BIT_OFFSET_TYPE) && defined(THRUST_FORCE_64_BIT_OFFSET_TYPE) diff --git a/projects/rocthrust/thrust/system/cuda/detail/extrema.h b/projects/rocthrust/thrust/system/cuda/detail/extrema.h index 5065ad52f925..c4d348311b8a 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/extrema.h +++ b/projects/rocthrust/thrust/system/cuda/detail/extrema.h @@ -51,10 +51,9 @@ # include # include +# include # include -# include - THRUST_NAMESPACE_BEGIN namespace cuda_cub { diff --git a/projects/rocthrust/thrust/system/cuda/detail/memory.inl b/projects/rocthrust/thrust/system/cuda/detail/memory.inl index b99cc5744c94..3aa2b755be89 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/memory.inl +++ b/projects/rocthrust/thrust/system/cuda/detail/memory.inl @@ -28,7 +28,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace cuda_cub diff --git a/projects/rocthrust/thrust/system/cuda/detail/merge.h b/projects/rocthrust/thrust/system/cuda/detail/merge.h index f0b2e21852dd..c81bae9dea1a 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/merge.h +++ b/projects/rocthrust/thrust/system/cuda/detail/merge.h @@ -47,7 +47,7 @@ # include # include -# include +# include THRUST_NAMESPACE_BEGIN namespace cuda_cub diff --git a/projects/rocthrust/thrust/system/cuda/detail/partition.h b/projects/rocthrust/thrust/system/cuda/detail/partition.h index f9117978f27c..512820182949 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/partition.h +++ b/projects/rocthrust/thrust/system/cuda/detail/partition.h @@ -55,7 +55,7 @@ # include # include -# include +# include THRUST_NAMESPACE_BEGIN namespace cuda_cub diff --git a/projects/rocthrust/thrust/system/cuda/detail/reduce.h b/projects/rocthrust/thrust/system/cuda/detail/reduce.h index 552e1c2b7a8e..5c2f44a57369 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/reduce.h +++ b/projects/rocthrust/thrust/system/cuda/detail/reduce.h @@ -56,7 +56,7 @@ # include # include -# include +# include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/cuda/detail/reduce_by_key.h b/projects/rocthrust/thrust/system/cuda/detail/reduce_by_key.h index fe9a090f492d..0a0270be182c 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/reduce_by_key.h +++ b/projects/rocthrust/thrust/system/cuda/detail/reduce_by_key.h @@ -58,10 +58,9 @@ # include # include +# include # include -# include - THRUST_NAMESPACE_BEGIN template # include -# include +# include THRUST_NAMESPACE_BEGIN namespace cuda_cub diff --git a/projects/rocthrust/thrust/system/cuda/detail/scan_by_key.h b/projects/rocthrust/thrust/system/cuda/detail/scan_by_key.h index 042a32ef4893..2a9fce710719 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/scan_by_key.h +++ b/projects/rocthrust/thrust/system/cuda/detail/scan_by_key.h @@ -55,7 +55,7 @@ # include # include -# include +# include THRUST_NAMESPACE_BEGIN namespace cuda_cub diff --git a/projects/rocthrust/thrust/system/cuda/detail/set_operations.h b/projects/rocthrust/thrust/system/cuda/detail/set_operations.h index deb58d488f30..f08864a99ca5 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/set_operations.h +++ b/projects/rocthrust/thrust/system/cuda/detail/set_operations.h @@ -54,8 +54,7 @@ # include # include - -# include +# include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/cuda/detail/sort.h b/projects/rocthrust/thrust/system/cuda/detail/sort.h index 15c673e850aa..645b7b54aac6 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/sort.h +++ b/projects/rocthrust/thrust/system/cuda/detail/sort.h @@ -59,8 +59,7 @@ # include # include - -# include +# include # if _CCCL_HAS_NVFP16() # include diff --git a/projects/rocthrust/thrust/system/cuda/detail/transform.h b/projects/rocthrust/thrust/system/cuda/detail/transform.h index a04b468dab0c..1152219154d7 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/transform.h +++ b/projects/rocthrust/thrust/system/cuda/detail/transform.h @@ -49,8 +49,7 @@ # include # include - -# include +# include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/cuda/detail/transform_reduce.h b/projects/rocthrust/thrust/system/cuda/detail/transform_reduce.h index 16e8c935421f..8d2b4c7ec657 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/transform_reduce.h +++ b/projects/rocthrust/thrust/system/cuda/detail/transform_reduce.h @@ -55,7 +55,8 @@ # include # include -# include +# include + # include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/cuda/detail/unique.h b/projects/rocthrust/thrust/system/cuda/detail/unique.h index 985143cf4be0..983ec958a0ae 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/unique.h +++ b/projects/rocthrust/thrust/system/cuda/detail/unique.h @@ -53,7 +53,7 @@ # include # include -# include +# include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/cuda/detail/unique_by_key.h b/projects/rocthrust/thrust/system/cuda/detail/unique_by_key.h index c3bb8203e20b..4fff42995af2 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/unique_by_key.h +++ b/projects/rocthrust/thrust/system/cuda/detail/unique_by_key.h @@ -55,7 +55,7 @@ # include # include -# include +# include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/detail/generic/reduce_by_key.inl b/projects/rocthrust/thrust/system/detail/generic/reduce_by_key.inl index af5ceaf54ed4..59663f058766 100644 --- a/projects/rocthrust/thrust/system/detail/generic/reduce_by_key.inl +++ b/projects/rocthrust/thrust/system/detail/generic/reduce_by_key.inl @@ -42,7 +42,7 @@ # include _THRUST_STD_INCLUDE(iterator) #endif -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace system diff --git a/projects/rocthrust/thrust/system/detail/generic/scan_by_key.inl b/projects/rocthrust/thrust/system/detail/generic/scan_by_key.inl index 16f0585fa349..b4dc3f51daa7 100644 --- a/projects/rocthrust/thrust/system/detail/generic/scan_by_key.inl +++ b/projects/rocthrust/thrust/system/detail/generic/scan_by_key.inl @@ -35,7 +35,7 @@ #include #include -#include +# include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN namespace system diff --git a/projects/rocthrust/thrust/system/detail/generic/shuffle.inl b/projects/rocthrust/thrust/system/detail/generic/shuffle.inl index 7d258da752cc..3700931f6cd2 100644 --- a/projects/rocthrust/thrust/system/detail/generic/shuffle.inl +++ b/projects/rocthrust/thrust/system/detail/generic/shuffle.inl @@ -24,7 +24,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN namespace system diff --git a/projects/rocthrust/thrust/system/detail/sequential/malloc_and_free.h b/projects/rocthrust/thrust/system/detail/sequential/malloc_and_free.h index d0d89126c063..8ea75d7b1027 100644 --- a/projects/rocthrust/thrust/system/detail/sequential/malloc_and_free.h +++ b/projects/rocthrust/thrust/system/detail/sequential/malloc_and_free.h @@ -28,7 +28,7 @@ #include #include -#include // for malloc & free +#include _THRUST_STD_INCLUDE(cstdlib) // for malloc & free THRUST_NAMESPACE_BEGIN namespace system @@ -41,13 +41,13 @@ namespace sequential template inline THRUST_HOST_DEVICE void* malloc(execution_policy&, std::size_t n) { - return std::malloc(n); + return _THRUST_STD::malloc(n); } // end mallc() template inline THRUST_HOST_DEVICE void free(sequential::execution_policy&, Pointer ptr) { - std::free(thrust::raw_pointer_cast(ptr)); + _THRUST_STD::free(thrust::raw_pointer_cast(ptr)); } // end mallc() } // namespace sequential diff --git a/projects/rocthrust/thrust/system/detail/sequential/stable_radix_sort.inl b/projects/rocthrust/thrust/system/detail/sequential/stable_radix_sort.inl index 486bfe22486e..6ba93279e488 100644 --- a/projects/rocthrust/thrust/system/detail/sequential/stable_radix_sort.inl +++ b/projects/rocthrust/thrust/system/detail/sequential/stable_radix_sort.inl @@ -36,8 +36,8 @@ #include _THRUST_STD_INCLUDE(utility) -#include -#include +#include _THRUST_STD_INCLUDE(cstdint) +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace system @@ -63,7 +63,7 @@ struct RadixEncoder { THRUST_HOST_DEVICE unsigned char operator()(char x) const { - if (std::numeric_limits::is_signed) + if (_THRUST_STD::numeric_limits::is_signed) { return static_cast(x) ^ static_cast(1) << (8 * sizeof(unsigned char) - 1); } diff --git a/projects/rocthrust/thrust/system/detail/sequential/trivial_copy.h b/projects/rocthrust/thrust/system/detail/sequential/trivial_copy.h index 9804132a6fc1..f0508d0604fd 100644 --- a/projects/rocthrust/thrust/system/detail/sequential/trivial_copy.h +++ b/projects/rocthrust/thrust/system/detail/sequential/trivial_copy.h @@ -32,7 +32,7 @@ #endif // no system header #include -#include +#include _THRUST_STD_INCLUDE(cstring) THRUST_NAMESPACE_BEGIN namespace system @@ -48,18 +48,12 @@ THRUST_HOST_DEVICE T* trivial_copy_n(const T* first, std::ptrdiff_t n, T* result if (n == 0) { // If `first` or `result` is an invalid pointer, - // the behavior of `std::memmove` is undefined, even if `n` is zero. + // the behavior of `_THRUST_STD::memmove` is undefined, even if `n` is zero. return result; } - T* return_value = nullptr; - - _THRUST_IF_TARGET(_THRUST_IS_HOST, - (std::memmove(result, first, n * sizeof(T)); return_value = result + n;), - ( // _THRUST_IS_DEVICE: - return_value = thrust::system::detail::sequential::general_copy_n(first, n, result);)); - - return return_value; + _THRUST_STD::memmove(result, first, n * sizeof(T)); + return result + n; } // end trivial_copy_n() } // end namespace sequential diff --git a/projects/rocthrust/thrust/system/hip/detail/binary_search.h b/projects/rocthrust/thrust/system/hip/detail/binary_search.h index f1cd44363cef..8a997570c1d0 100644 --- a/projects/rocthrust/thrust/system/hip/detail/binary_search.h +++ b/projects/rocthrust/thrust/system/hip/detail/binary_search.h @@ -36,7 +36,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN namespace hip_rocprim diff --git a/projects/rocthrust/thrust/system/hip/detail/copy_if.h b/projects/rocthrust/thrust/system/hip/detail/copy_if.h index 21cea6a15f35..8f7c49aa68ff 100644 --- a/projects/rocthrust/thrust/system/hip/detail/copy_if.h +++ b/projects/rocthrust/thrust/system/hip/detail/copy_if.h @@ -56,7 +56,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN // XXX declare generic copy_if interface diff --git a/projects/rocthrust/thrust/system/hip/detail/dispatch.h b/projects/rocthrust/thrust/system/hip/detail/dispatch.h index ad325ad39a57..3f5200f0b9b7 100644 --- a/projects/rocthrust/thrust/system/hip/detail/dispatch.h +++ b/projects/rocthrust/thrust/system/hip/detail/dispatch.h @@ -32,11 +32,11 @@ #include #if _THRUST_HAS_DEVICE_SYSTEM_STD +# include _THRUST_STD_INCLUDE(cstdint) # include _THRUST_STD_INCLUDE(limits) # include _THRUST_STD_INCLUDE(type_traits) #endif -#include #include #if defined(THRUST_FORCE_32_BIT_OFFSET_TYPE) && defined(THRUST_FORCE_64_BIT_OFFSET_TYPE) diff --git a/projects/rocthrust/thrust/system/hip/detail/extrema.h b/projects/rocthrust/thrust/system/hip/detail/extrema.h index a6cd5a532f43..e1ed851a076a 100644 --- a/projects/rocthrust/thrust/system/hip/detail/extrema.h +++ b/projects/rocthrust/thrust/system/hip/detail/extrema.h @@ -56,7 +56,7 @@ # include _THRUST_STD_INCLUDE(iterator) -# include +# include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN namespace hip_rocprim diff --git a/projects/rocthrust/thrust/system/hip/detail/general/temp_storage.h b/projects/rocthrust/thrust/system/hip/detail/general/temp_storage.h index 40c8b677903b..4af6ee4f6bfe 100644 --- a/projects/rocthrust/thrust/system/hip/detail/general/temp_storage.h +++ b/projects/rocthrust/thrust/system/hip/detail/general/temp_storage.h @@ -1,4 +1,4 @@ -// Copyright (c) 2024-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2024-2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -25,7 +25,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cstddef) THRUST_NAMESPACE_BEGIN namespace system diff --git a/projects/rocthrust/thrust/system/hip/detail/memory.inl b/projects/rocthrust/thrust/system/hip/detail/memory.inl index b7817f4ded9e..d287a2afeea0 100644 --- a/projects/rocthrust/thrust/system/hip/detail/memory.inl +++ b/projects/rocthrust/thrust/system/hip/detail/memory.inl @@ -29,7 +29,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace hip_rocprim diff --git a/projects/rocthrust/thrust/system/hip/detail/merge.h b/projects/rocthrust/thrust/system/hip/detail/merge.h index 035c3350fad1..954eaeaa1492 100644 --- a/projects/rocthrust/thrust/system/hip/detail/merge.h +++ b/projects/rocthrust/thrust/system/hip/detail/merge.h @@ -46,7 +46,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) // rocPRIM includes # include diff --git a/projects/rocthrust/thrust/system/hip/detail/partition.h b/projects/rocthrust/thrust/system/hip/detail/partition.h index ac899d350ad0..86733b76a068 100644 --- a/projects/rocthrust/thrust/system/hip/detail/partition.h +++ b/projects/rocthrust/thrust/system/hip/detail/partition.h @@ -52,7 +52,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) // rocprim include # include diff --git a/projects/rocthrust/thrust/system/hip/detail/reduce.h b/projects/rocthrust/thrust/system/hip/detail/reduce.h index df87dc1749ac..6b48b8470a03 100644 --- a/projects/rocthrust/thrust/system/hip/detail/reduce.h +++ b/projects/rocthrust/thrust/system/hip/detail/reduce.h @@ -52,7 +52,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) // rocprim include # include diff --git a/projects/rocthrust/thrust/system/hip/detail/reduce_by_key.h b/projects/rocthrust/thrust/system/hip/detail/reduce_by_key.h index c373197cfbfd..177cfecdb337 100644 --- a/projects/rocthrust/thrust/system/hip/detail/reduce_by_key.h +++ b/projects/rocthrust/thrust/system/hip/detail/reduce_by_key.h @@ -59,7 +59,7 @@ # include _THRUST_STD_INCLUDE(iterator) # endif -# include +# include _THRUST_STD_INCLUDE(cstdint) // rocprim include # include diff --git a/projects/rocthrust/thrust/system/hip/detail/scan.h b/projects/rocthrust/thrust/system/hip/detail/scan.h index f527d7722158..68ca548f3aa3 100644 --- a/projects/rocthrust/thrust/system/hip/detail/scan.h +++ b/projects/rocthrust/thrust/system/hip/detail/scan.h @@ -48,7 +48,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) // rocprim include # include diff --git a/projects/rocthrust/thrust/system/hip/detail/scan_by_key.h b/projects/rocthrust/thrust/system/hip/detail/scan_by_key.h index 7b0084efd87b..2219ed2e738d 100644 --- a/projects/rocthrust/thrust/system/hip/detail/scan_by_key.h +++ b/projects/rocthrust/thrust/system/hip/detail/scan_by_key.h @@ -52,7 +52,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) // rocprim include # include diff --git a/projects/rocthrust/thrust/system/hip/detail/set_operations.h b/projects/rocthrust/thrust/system/hip/detail/set_operations.h index e4f098b2687b..b40ae583b224 100644 --- a/projects/rocthrust/thrust/system/hip/detail/set_operations.h +++ b/projects/rocthrust/thrust/system/hip/detail/set_operations.h @@ -64,7 +64,7 @@ # include -# include +# include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/hip/detail/sort.h b/projects/rocthrust/thrust/system/hip/detail/sort.h index 788ecc331286..7bd71ead4fa4 100644 --- a/projects/rocthrust/thrust/system/hip/detail/sort.h +++ b/projects/rocthrust/thrust/system/hip/detail/sort.h @@ -53,7 +53,8 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) + # include # include diff --git a/projects/rocthrust/thrust/system/hip/detail/transform.h b/projects/rocthrust/thrust/system/hip/detail/transform.h index a1d644a1fd08..12259dd92d42 100644 --- a/projects/rocthrust/thrust/system/hip/detail/transform.h +++ b/projects/rocthrust/thrust/system/hip/detail/transform.h @@ -48,7 +48,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/hip/detail/transform_reduce.h b/projects/rocthrust/thrust/system/hip/detail/transform_reduce.h index 2c21208e3c28..a92012a37425 100644 --- a/projects/rocthrust/thrust/system/hip/detail/transform_reduce.h +++ b/projects/rocthrust/thrust/system/hip/detail/transform_reduce.h @@ -52,7 +52,8 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) + # include THRUST_NAMESPACE_BEGIN diff --git a/projects/rocthrust/thrust/system/hip/detail/unique.h b/projects/rocthrust/thrust/system/hip/detail/unique.h index 1ba2322cd4aa..d24a011ab762 100644 --- a/projects/rocthrust/thrust/system/hip/detail/unique.h +++ b/projects/rocthrust/thrust/system/hip/detail/unique.h @@ -51,7 +51,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) // rocPRIM includes # include diff --git a/projects/rocthrust/thrust/system/hip/detail/unique_by_key.h b/projects/rocthrust/thrust/system/hip/detail/unique_by_key.h index 09cdbf1ac90b..a2f14886a90e 100644 --- a/projects/rocthrust/thrust/system/hip/detail/unique_by_key.h +++ b/projects/rocthrust/thrust/system/hip/detail/unique_by_key.h @@ -52,7 +52,7 @@ # include # include -# include +# include _THRUST_STD_INCLUDE(cstdint) // rocPRIM includes # include diff --git a/projects/rocthrust/thrust/system/omp/detail/memory.inl b/projects/rocthrust/thrust/system/omp/detail/memory.inl index 0057db7a02fb..cf5842589382 100644 --- a/projects/rocthrust/thrust/system/omp/detail/memory.inl +++ b/projects/rocthrust/thrust/system/omp/detail/memory.inl @@ -29,7 +29,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace system diff --git a/projects/rocthrust/thrust/system/tbb/detail/memory.inl b/projects/rocthrust/thrust/system/tbb/detail/memory.inl index 3c0f4c661ff1..adceaa1face8 100644 --- a/projects/rocthrust/thrust/system/tbb/detail/memory.inl +++ b/projects/rocthrust/thrust/system/tbb/detail/memory.inl @@ -29,7 +29,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(limits) THRUST_NAMESPACE_BEGIN namespace system From 518c37a4dd938962c8b4777aaf69d757d1dbdbcc Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Fri, 12 Jun 2026 11:36:22 +0200 Subject: [PATCH 10/86] Use libhipcxx includes - initializer_list --- projects/rocthrust/CHANGELOG.md | 2 +- projects/rocthrust/thrust/detail/vector_base.h | 9 +++++---- projects/rocthrust/thrust/detail/vector_base.inl | 6 +++--- projects/rocthrust/thrust/device_vector.h | 8 ++++---- projects/rocthrust/thrust/host_vector.h | 8 ++++---- 5 files changed, 17 insertions(+), 16 deletions(-) diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index 5296df49cb2a..cfff138335b3 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -7,7 +7,7 @@ Documentation for rocThrust available at ### Changed -* Replaced cmath, cfloat, cstdint, cstddef, cstring, move, limits operators to `hip::std` operators. +* Replaced cfloat, cmath, cstddef, cstdint, cstring, initializer_list, limits and move operators to `hip::std` operators. ### Removed diff --git a/projects/rocthrust/thrust/detail/vector_base.h b/projects/rocthrust/thrust/detail/vector_base.h index a778e28248c0..56fb727bff37 100644 --- a/projects/rocthrust/thrust/detail/vector_base.h +++ b/projects/rocthrust/thrust/detail/vector_base.h @@ -41,12 +41,13 @@ #include _THRUST_STD_INCLUDE(utility) -#include #include #if !_THRUST_HAS_DEVICE_SYSTEM_STD # include #endif +#include _THRUST_STD_INCLUDE(initializer_list) + THRUST_NAMESPACE_BEGIN namespace detail @@ -144,18 +145,18 @@ class vector_base /*! This constructor builds a \p vector_base from an intializer_list. * \param il The intializer_list. */ - vector_base(std::initializer_list il); + vector_base(_THRUST_STD::initializer_list il); /*! This constructor builds a \p vector_base from an intializer_list. * \param il The intializer_list. * \param alloc The allocator to use by this device_vector. */ - vector_base(std::initializer_list il, const Alloc& alloc); + vector_base(_THRUST_STD::initializer_list il, const Alloc& alloc); /*! Assign operator copies from an initializer_list * \param il The initializer_list. */ - vector_base& operator=(std::initializer_list il); + vector_base& operator=(_THRUST_STD::initializer_list il); /*! Copy constructor copies from an exemplar vector_base with different * type. diff --git a/projects/rocthrust/thrust/detail/vector_base.inl b/projects/rocthrust/thrust/detail/vector_base.inl index 9d4e86777194..3f6512086012 100644 --- a/projects/rocthrust/thrust/detail/vector_base.inl +++ b/projects/rocthrust/thrust/detail/vector_base.inl @@ -191,7 +191,7 @@ vector_base& vector_base::operator=(const std::vector -vector_base::vector_base(std::initializer_list il) +vector_base::vector_base(_THRUST_STD::initializer_list il) : m_storage() , m_size(0) { @@ -199,7 +199,7 @@ vector_base::vector_base(std::initializer_list il) } // end vector_base::vector_base() template -vector_base::vector_base(std::initializer_list il, const Alloc& alloc) +vector_base::vector_base(_THRUST_STD::initializer_list il, const Alloc& alloc) : m_storage(alloc) , m_size(0) { @@ -207,7 +207,7 @@ vector_base::vector_base(std::initializer_list il, const Alloc& all } // end vector_base::vector_base() template -vector_base& vector_base::operator=(std::initializer_list il) +vector_base& vector_base::operator=(_THRUST_STD::initializer_list il) { assign(il.begin(), il.end()); diff --git a/projects/rocthrust/thrust/device_vector.h b/projects/rocthrust/thrust/device_vector.h index a73b685e24b6..291a052c908b 100644 --- a/projects/rocthrust/thrust/device_vector.h +++ b/projects/rocthrust/thrust/device_vector.h @@ -33,9 +33,9 @@ #include #include +#include _THRUST_STD_INCLUDE(initializer_list) #include _THRUST_STD_INCLUDE(utility) -#include #include THRUST_NAMESPACE_BEGIN @@ -232,7 +232,7 @@ class device_vector : public detail::vector_base /*! This constructor builds a \p device_vector from an intializer_list. * \param il The intializer_list. */ - device_vector(std::initializer_list il) + device_vector(_THRUST_STD::initializer_list il) : Parent(il) {} @@ -240,14 +240,14 @@ class device_vector : public detail::vector_base * \param il The intializer_list. * \param alloc The allocator to use by this device_vector. */ - device_vector(std::initializer_list il, const Alloc& alloc) + device_vector(_THRUST_STD::initializer_list il, const Alloc& alloc) : Parent(il, alloc) {} /*! Assign an \p intializer_list with a matching element type * \param il The intializer_list. */ - device_vector& operator=(std::initializer_list il) + device_vector& operator=(_THRUST_STD::initializer_list il) { Parent::operator=(il); return *this; diff --git a/projects/rocthrust/thrust/host_vector.h b/projects/rocthrust/thrust/host_vector.h index 634fd7c21d55..e329cfdcfacc 100644 --- a/projects/rocthrust/thrust/host_vector.h +++ b/projects/rocthrust/thrust/host_vector.h @@ -33,9 +33,9 @@ #include #include +#include _THRUST_STD_INCLUDE(initializer_list) #include _THRUST_STD_INCLUDE(utility) -#include #include THRUST_NAMESPACE_BEGIN @@ -234,7 +234,7 @@ class host_vector : public detail::vector_base /*! This constructor builds a \p host_vector from an intializer_list. * \param il The intializer_list. */ - host_vector(std::initializer_list il) + host_vector(_THRUST_STD::initializer_list il) : Parent(il) {} @@ -242,14 +242,14 @@ class host_vector : public detail::vector_base * \param il The intializer_list. * \param alloc The allocator to use by this host_vector. */ - host_vector(std::initializer_list il, const Alloc& alloc) + host_vector(_THRUST_STD::initializer_list il, const Alloc& alloc) : Parent(il, alloc) {} /*! Assign an \p intializer_list with a matching element type * \param il The intializer_list. */ - host_vector& operator=(std::initializer_list il) + host_vector& operator=(_THRUST_STD::initializer_list il) { Parent::operator=(il); return *this; From b94c30c62da2a7674b1ee4fc0e832611c0091c7b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Fri, 12 Jun 2026 11:36:39 +0200 Subject: [PATCH 11/86] Use libhipcxx includes - cassert --- projects/rocthrust/CHANGELOG.md | 2 +- .../rocthrust/thrust/detail/allocator/temporary_allocator.inl | 2 +- projects/rocthrust/thrust/mr/disjoint_pool.h | 3 +-- projects/rocthrust/thrust/mr/pool.h | 3 +-- .../rocthrust/thrust/system/cuda/detail/core/agent_launcher.h | 2 +- projects/rocthrust/thrust/system/tbb/detail/reduce_by_key.inl | 2 +- projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h | 2 +- 7 files changed, 7 insertions(+), 9 deletions(-) diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index cfff138335b3..bf2cfee2a13c 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -7,7 +7,7 @@ Documentation for rocThrust available at ### Changed -* Replaced cfloat, cmath, cstddef, cstdint, cstring, initializer_list, limits and move operators to `hip::std` operators. +* Replaced cassert, cfloat, cmath, cstddef, cstdint, cstring, initializer_list, limits and move operators to `hip::std` operators. ### Removed diff --git a/projects/rocthrust/thrust/detail/allocator/temporary_allocator.inl b/projects/rocthrust/thrust/detail/allocator/temporary_allocator.inl index a53704fa37e0..64f85aa61da1 100644 --- a/projects/rocthrust/thrust/detail/allocator/temporary_allocator.inl +++ b/projects/rocthrust/thrust/detail/allocator/temporary_allocator.inl @@ -30,7 +30,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(cassert) #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA # if (defined(_NVHPC_CUDA) || defined(__CUDA_ARCH__)) diff --git a/projects/rocthrust/thrust/mr/disjoint_pool.h b/projects/rocthrust/thrust/mr/disjoint_pool.h index bc639dcc80d6..09b953e2123e 100644 --- a/projects/rocthrust/thrust/mr/disjoint_pool.h +++ b/projects/rocthrust/thrust/mr/disjoint_pool.h @@ -41,10 +41,9 @@ #include #include +#include _THRUST_STD_INCLUDE(cassert) #include _THRUST_STD_INCLUDE(cstdint) -#include - THRUST_NAMESPACE_BEGIN namespace mr { diff --git a/projects/rocthrust/thrust/mr/pool.h b/projects/rocthrust/thrust/mr/pool.h index 4813633d338b..a7ece2b78252 100644 --- a/projects/rocthrust/thrust/mr/pool.h +++ b/projects/rocthrust/thrust/mr/pool.h @@ -38,10 +38,9 @@ #include #include +#include _THRUST_STD_INCLUDE(cassert) #include _THRUST_STD_INCLUDE(cstdint) -#include - THRUST_NAMESPACE_BEGIN namespace mr { diff --git a/projects/rocthrust/thrust/system/cuda/detail/core/agent_launcher.h b/projects/rocthrust/thrust/system/cuda/detail/core/agent_launcher.h index ac58cda94db9..004530a06f52 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/core/agent_launcher.h +++ b/projects/rocthrust/thrust/system/cuda/detail/core/agent_launcher.h @@ -40,7 +40,7 @@ # include # include -# include +# include # include diff --git a/projects/rocthrust/thrust/system/tbb/detail/reduce_by_key.inl b/projects/rocthrust/thrust/system/tbb/detail/reduce_by_key.inl index 100e8e18bbce..3d38304ac925 100644 --- a/projects/rocthrust/thrust/system/tbb/detail/reduce_by_key.inl +++ b/projects/rocthrust/thrust/system/tbb/detail/reduce_by_key.inl @@ -37,8 +37,8 @@ #include #include #include +#include -#include #include #include diff --git a/projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h b/projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h index e3552aca2fb0..6ca2025b4b37 100644 --- a/projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h +++ b/projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h @@ -32,8 +32,8 @@ #include #include +#include -#include #include #include From a3752d578f4757ddd886af332ffadb5adc0da028 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Fri, 12 Jun 2026 13:30:27 +0200 Subject: [PATCH 12/86] Use libhipcxx includes - make ``complex`` less dependent on host headers --- projects/rocthrust/CHANGELOG.md | 2 +- projects/rocthrust/thrust/complex.h | 41 +++++++++++++------ .../thrust/detail/complex/complex.inl | 24 +++++++---- 3 files changed, 45 insertions(+), 22 deletions(-) diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index bf2cfee2a13c..f0f0c226fb7e 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -7,7 +7,7 @@ Documentation for rocThrust available at ### Changed -* Replaced cassert, cfloat, cmath, cstddef, cstdint, cstring, initializer_list, limits and move operators to `hip::std` operators. +* Replaced cassert, cfloat, cmath, complex, cstddef, cstdint, cstring, initializer_list, limits and move operators to `hip::std` operators. ### Removed diff --git a/projects/rocthrust/thrust/complex.h b/projects/rocthrust/thrust/complex.h index cee2c715dfb2..50c9a8d114bb 100644 --- a/projects/rocthrust/thrust/complex.h +++ b/projects/rocthrust/thrust/complex.h @@ -35,9 +35,12 @@ #include #include -#include -#include -#include +#include _THRUST_STD_INCLUDE(cmath) +#if !THRUST_COMPILER(NVRTC) +# include +# include +#endif // !THRUST_COMPILER(NVRTC) + #if !_THRUST_HAS_DEVICE_SYSTEM_STD # include #endif @@ -117,12 +120,13 @@ struct complex template THRUST_HOST_DEVICE complex(const complex& z); +#if !THRUST_COMPILER(NVRTC) /*! This converting copy constructor copies from a std::complex with * a type that is convertible to this \p complex's \c value_type. * * \param z The \p complex to copy from. */ - THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex(const std::complex& z); + THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex(const ::std::complex& z); /*! This converting copy constructor copies from a std::complex with * a type that is convertible to this \p complex's \c value_type. @@ -132,7 +136,8 @@ struct complex * \tparam U is convertible to \c value_type. */ template - THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex(const std::complex& z); + THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex(const ::std::complex& z); +#endif // !THRUST_COMPILER(NVRTC) /* --- Assignment Operators --- */ @@ -160,12 +165,13 @@ struct complex template THRUST_HOST_DEVICE complex& operator=(const complex& z); +#if !THRUST_COMPILER(NVRTC) /*! Assign `z.real()` and `z.imag()` to the real and imaginary parts of this * \p complex respectively. * * \param z The \p complex to copy from. */ - THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& operator=(const std::complex& z); + THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& operator=(const ::std::complex& z); /*! Assign `z.real()` and `z.imag()` to the real and imaginary parts of this * \p complex respectively. @@ -175,7 +181,8 @@ struct complex * \tparam U is convertible to \c value_type. */ template - THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& operator=(const std::complex& z); + THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& operator=(const ::std::complex& z); +#endif // !THRUST_COMPILER(NVRTC) /* --- Compound Assignment Operators --- */ @@ -333,14 +340,16 @@ struct complex data.y = im; } +#if !THRUST_COMPILER(NVRTC) /* --- Casting functions --- */ /*! Casts this \p complex to a std::complex of the same type. */ - THRUST_HOST operator std::complex() const + THRUST_HOST operator ::std::complex() const { - return std::complex(real(), imag()); + return ::std::complex(real(), imag()); } +#endif // !THRUST_COMPILER(NVRTC) private: struct alignas(sizeof(T) * 2) storage @@ -722,6 +731,7 @@ THRUST_HOST_DEVICE complex asinh(const complex& z); template THRUST_HOST_DEVICE complex atanh(const complex& z); +#if !THRUST_COMPILER(NVRTC) /* --- Stream Operators --- */ /*! Writes to an output stream a \p complex number in the form (real, imaginary). @@ -745,7 +755,8 @@ std::basic_ostream& operator<<(std::basic_ostream& * \param z The \p complex number to set. */ template -THRUST_HOST std::basic_istream& operator>>(std::basic_istream& is, complex& z); +THRUST_HOST ::std::basic_istream& operator>>(std::basic_istream& is, complex& z); +#endif // !THRUST_COMPILER(NVRTC) /* --- Equality Operators --- */ @@ -757,6 +768,7 @@ THRUST_HOST std::basic_istream& operator>>(std::basic_istream THRUST_HOST_DEVICE bool operator==(const complex& x, const complex& y); +#if !THRUST_COMPILER(NVRTC) /*! Returns true if two \p complex numbers are equal and false otherwise. * * \param x The first \p complex. @@ -771,7 +783,8 @@ THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const complex& x, cons * \param y The second \p complex. */ template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const std::complex& x, const complex& y); +THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const ::std::complex& x, const complex& y); +#endif // !THRUST_COMPILER(NVRTC) /*! Returns true if the imaginary part of the \p complex number is zero and * the real part is equal to the scalar. Returns false otherwise. @@ -791,6 +804,7 @@ THRUST_HOST_DEVICE bool operator==(const T0& x, const complex& y); template THRUST_HOST_DEVICE bool operator==(const complex& x, const T1& y); +#if !THRUST_COMPILER(NVRTC) /*! Returns true if two \p complex numbers are different and false otherwise. * * \param x The first \p complex. @@ -805,7 +819,7 @@ THRUST_HOST_DEVICE bool operator!=(const complex& x, const complex& y); * \param y The second \p complex. */ template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const complex& x, const std::complex& y); +THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const complex& x, const ::std::complex& y); /*! Returns true if two \p complex numbers are different and false otherwise. * @@ -813,7 +827,8 @@ THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const complex& x, cons * \param y The second \p complex. */ template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const std::complex& x, const complex& y); +THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const ::std::complex& x, const complex& y); +#endif // !THRUST_COMPILER(NVRTC) /*! Returns true if the imaginary part of the \p complex number is not zero or * the real part is different from the scalar. Returns false otherwise. diff --git a/projects/rocthrust/thrust/detail/complex/complex.inl b/projects/rocthrust/thrust/detail/complex/complex.inl index 1427f9e66d91..e7f64c9da667 100644 --- a/projects/rocthrust/thrust/detail/complex/complex.inl +++ b/projects/rocthrust/thrust/detail/complex/complex.inl @@ -49,8 +49,9 @@ THRUST_HOST_DEVICE complex::complex(const complex& z) : data{T(z.real()), T(z.imag())} {} +#if !THRUST_COMPILER(NVRTC) template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const std::complex& z) +THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const ::std::complex& z) // Initialize the storage in the member initializer list using C++ unicorn // initialization. This allows `complex` to work. : data{THRUST_STD_COMPLEX_REAL(z), THRUST_STD_COMPLEX_IMAG(z)} @@ -58,12 +59,13 @@ THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const std::complex& template template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const std::complex& z) +THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const ::std::complex& z) // Initialize the storage in the member initializer list using C++ unicorn // initialization. This allows `complex` to work. // We do a functional-style cast here to suppress conversion warnings. : data{T(THRUST_STD_COMPLEX_REAL(z)), T(THRUST_STD_COMPLEX_IMAG(z))} {} +#endif // !THRUST_COMPILER(NVRTC) /* --- Assignment Operators --- */ @@ -84,8 +86,9 @@ THRUST_HOST_DEVICE complex& complex::operator=(const complex& z) return *this; } +#if !THRUST_COMPILER(NVRTC) template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& complex::operator=(const std::complex& z) +THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& complex::operator=(const ::std::complex& z) { real(THRUST_STD_COMPLEX_REAL(z)); imag(THRUST_STD_COMPLEX_IMAG(z)); @@ -94,12 +97,13 @@ THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& complex::operator=(const st template template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& complex::operator=(const std::complex& z) +THRUST_HOST THRUST_STD_COMPLEX_DEVICE complex& complex::operator=(const ::std::complex& z) { real(T(THRUST_STD_COMPLEX_REAL(z))); imag(T(THRUST_STD_COMPLEX_IMAG(z))); return *this; } +#endif // !THRUST_COMPILER(NVRTC) /* --- Compound Assignment Operators --- */ @@ -175,17 +179,19 @@ THRUST_HOST_DEVICE bool operator==(const complex& x, const complex& y) return x.real() == y.real() && x.imag() == y.imag(); } +#if !THRUST_COMPILER(NVRTC) template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const complex& x, const std::complex& y) +THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const complex& x, const ::std::complex& y) { return x.real() == THRUST_STD_COMPLEX_REAL(y) && x.imag() == THRUST_STD_COMPLEX_IMAG(y); } template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const std::complex& x, const complex& y) +THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const ::std::complex& x, const complex& y) { return THRUST_STD_COMPLEX_REAL(x) == y.real() && THRUST_STD_COMPLEX_IMAG(x) == y.imag(); } +#endif // !THRUST_COMPILER(NVRTC) template THRUST_HOST_DEVICE bool operator==(const T0& x, const complex& y) @@ -205,17 +211,19 @@ THRUST_HOST_DEVICE bool operator!=(const complex& x, const complex& y) return !(x == y); } +#if !THRUST_COMPILER(NVRTC) template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const complex& x, const std::complex& y) +THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const complex& x, const ::std::complex& y) { return !(x == y); } template -THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const std::complex& x, const complex& y) +THRUST_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const ::std::complex& x, const complex& y) { return !(x == y); } +#endif // !THRUST_COMPILER(NVRTC) template THRUST_HOST_DEVICE bool operator!=(const T0& x, const complex& y) From a077a80c31597412da42a4f015a933b166bc1b1a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Fri, 12 Jun 2026 14:39:36 +0200 Subject: [PATCH 13/86] Use libhipcxx includes - remove infinity test --- projects/rocthrust/test/test_complex_various.cpp | 11 +---------- 1 file changed, 1 insertion(+), 10 deletions(-) diff --git a/projects/rocthrust/test/test_complex_various.cpp b/projects/rocthrust/test/test_complex_various.cpp index fe79b1aaa6ab..f1f13aafdd86 100644 --- a/projects/rocthrust/test/test_complex_various.cpp +++ b/projects/rocthrust/test/test_complex_various.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -153,15 +153,6 @@ void run_trig_tests( // // =================================================================== -TYPED_TEST_SUITE(VariousComplexTest, FloatDouble); -TYPED_TEST(VariousComplexTest, getInf) -{ - using T = typename TestFixture::T; - - T inf = thrust::detail::complex::infinity(); - ASSERT_TRUE(std::isinf(inf)); -} - #if defined _MSC_VER TYPED_TEST(VariousComplexTest, isinf) From dc5ac9676b91b503aa2bc0515519f628803d7729 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Fri, 12 Jun 2026 14:25:59 +0200 Subject: [PATCH 14/86] Use libhipcxx includes - type_traits --- projects/rocthrust/CHANGELOG.md | 2 +- projects/rocthrust/thrust/system/cpp/pointer.h | 4 ++-- projects/rocthrust/thrust/system/cuda/pointer.h | 4 ++-- projects/rocthrust/thrust/system/hip/pointer.h | 4 ++-- projects/rocthrust/thrust/system/omp/pointer.h | 4 ++-- .../rocthrust/thrust/system/tbb/detail/reduce_intervals.h | 5 ++--- projects/rocthrust/thrust/system/tbb/pointer.h | 4 ++-- projects/rocthrust/thrust/zip_function.h | 6 +++--- 8 files changed, 16 insertions(+), 17 deletions(-) diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index f0f0c226fb7e..a78d5c07ea91 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -7,7 +7,7 @@ Documentation for rocThrust available at ### Changed -* Replaced cassert, cfloat, cmath, complex, cstddef, cstdint, cstring, initializer_list, limits and move operators to `hip::std` operators. +* Replaced cassert, cfloat, cmath, complex, cstddef, cstdint, cstring, initializer_list, limits, move and type_traits operators to `hip::std` operators. ### Removed diff --git a/projects/rocthrust/thrust/system/cpp/pointer.h b/projects/rocthrust/thrust/system/cpp/pointer.h index 04963e2cf618..2951782d70a7 100644 --- a/projects/rocthrust/thrust/system/cpp/pointer.h +++ b/projects/rocthrust/thrust/system/cpp/pointer.h @@ -32,7 +32,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(type_traits) THRUST_NAMESPACE_BEGIN namespace system @@ -88,7 +88,7 @@ using pointer = thrust::pointer -using universal_pointer = thrust::pointer::type>; +using universal_pointer = thrust::pointer>; /*! \p reference is a wrapped reference to an object stored in memory available * to the \p cpp system. \p reference is the type of the result of diff --git a/projects/rocthrust/thrust/system/cuda/pointer.h b/projects/rocthrust/thrust/system/cuda/pointer.h index 5ab5109acca2..8a82b90232a3 100644 --- a/projects/rocthrust/thrust/system/cuda/pointer.h +++ b/projects/rocthrust/thrust/system/cuda/pointer.h @@ -33,7 +33,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(type_traits) THRUST_NAMESPACE_BEGIN namespace cuda_cub @@ -87,7 +87,7 @@ using pointer = thrust::pointer -using universal_pointer = thrust::pointer::type>; +using universal_pointer = thrust::pointer>; /*! \p cuda::reference is a wrapped reference to an object stored in memory * accessible by the \p cuda system. \p cuda::reference is the type of the diff --git a/projects/rocthrust/thrust/system/hip/pointer.h b/projects/rocthrust/thrust/system/hip/pointer.h index 8ac6984e0b79..e1e0d7a6a112 100644 --- a/projects/rocthrust/thrust/system/hip/pointer.h +++ b/projects/rocthrust/thrust/system/hip/pointer.h @@ -1,6 +1,6 @@ /* * Copyright 2008-2020 NVIDIA Corporation - * Modifications Copyright (c) 2019-2025, Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright (c) 2019-2026, Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -88,7 +88,7 @@ using pointer = thrust::pointer -using universal_pointer = thrust::pointer::type>; +using universal_pointer = thrust::pointer>; /*! \p hip::reference is a wrapped reference to an object stored in memory * accessible by the \p hip system. \p hip::reference is the type of the diff --git a/projects/rocthrust/thrust/system/omp/pointer.h b/projects/rocthrust/thrust/system/omp/pointer.h index 8071701a5cbc..c691144d5dae 100644 --- a/projects/rocthrust/thrust/system/omp/pointer.h +++ b/projects/rocthrust/thrust/system/omp/pointer.h @@ -33,7 +33,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(type_traits) THRUST_NAMESPACE_BEGIN namespace system @@ -89,7 +89,7 @@ using pointer = thrust::pointer -using universal_pointer = thrust::pointer::type>; +using universal_pointer = thrust::pointer>; /*! \p reference is a wrapped reference to an object stored in memory available * to the \p omp system. \p reference is the type of the result of diff --git a/projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h b/projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h index 6ca2025b4b37..aad3dc913c36 100644 --- a/projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h +++ b/projects/rocthrust/thrust/system/tbb/detail/reduce_intervals.h @@ -33,8 +33,7 @@ #include #include - -#include +#include #include @@ -83,7 +82,7 @@ struct body RandomAccessIterator1 my_last = first + offset_to_last; // carefully pass the init value for the interval with raw_reference_cast - using sum_type = typename std::decay::type; + using sum_type = _THRUST_STD::decay_t::type; result[interval_idx] = thrust::reduce(thrust::seq, my_first + 1, my_last, sum_type(thrust::raw_reference_cast(*my_first)), binary_op); } diff --git a/projects/rocthrust/thrust/system/tbb/pointer.h b/projects/rocthrust/thrust/system/tbb/pointer.h index 910e1ec93af1..c065a43c6753 100644 --- a/projects/rocthrust/thrust/system/tbb/pointer.h +++ b/projects/rocthrust/thrust/system/tbb/pointer.h @@ -33,7 +33,7 @@ #include #include -#include +#include _THRUST_STD_INCLUDE(type_traits) THRUST_NAMESPACE_BEGIN namespace system @@ -89,7 +89,7 @@ using pointer = thrust::pointer -using universal_pointer = thrust::pointer::type>; +using universal_pointer = thrust::pointer>; /*! \p reference is a wrapped reference to an object stored in memory available * to the \p tbb system. \p reference is the type of the result of diff --git a/projects/rocthrust/thrust/zip_function.h b/projects/rocthrust/thrust/zip_function.h index 4c369b6413ef..6409ca44ab47 100644 --- a/projects/rocthrust/thrust/zip_function.h +++ b/projects/rocthrust/thrust/zip_function.h @@ -73,7 +73,7 @@ THRUST_HOST_DEVICE decltype(auto) apply_impl(Function&& func, Tuple&& args, inde template THRUST_HOST_DEVICE decltype(auto) apply(Function&& func, Tuple&& args) { - constexpr auto tuple_size = thrust::tuple_size::type>::value; + constexpr auto tuple_size = thrust::tuple_size<_THRUST_STD::decay_t>::value; return apply_impl(THRUST_FWD(func), THRUST_FWD(args), make_index_sequence{}); } @@ -180,9 +180,9 @@ class zip_function * \see zip_function */ template -THRUST_HOST_DEVICE zip_function::type> make_zip_function(Function&& fun) +THRUST_HOST_DEVICE zip_function<_THRUST_STD::decay_t> make_zip_function(Function&& fun) { - using func_t = typename std::decay::type; + using func_t = _THRUST_STD::decay_t; return zip_function(THRUST_FWD(fun)); } From 8988e08a1a2b3a64d00cd1e1289868f0ce8ccc91 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Fri, 31 Jul 2026 09:51:09 +0200 Subject: [PATCH 15/86] Use libhipcxx includes - bad_alloc --- projects/rocthrust/CHANGELOG.md | 2 +- projects/rocthrust/thrust/device_malloc_allocator.h | 5 ++--- projects/rocthrust/thrust/device_new_allocator.h | 12 +++++------- 3 files changed, 8 insertions(+), 11 deletions(-) diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index a78d5c07ea91..6a9f42fd7282 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -7,7 +7,7 @@ Documentation for rocThrust available at ### Changed -* Replaced cassert, cfloat, cmath, complex, cstddef, cstdint, cstring, initializer_list, limits, move and type_traits operators to `hip::std` operators. +* Replaced bad_alloc, cassert, cfloat, cmath, complex, cstddef, cstdint, cstring, initializer_list, limits, move and type_traits operators to `hip::std` operators. ### Removed diff --git a/projects/rocthrust/thrust/device_malloc_allocator.h b/projects/rocthrust/thrust/device_malloc_allocator.h index 70fe54879327..0fad73a3790c 100644 --- a/projects/rocthrust/thrust/device_malloc_allocator.h +++ b/projects/rocthrust/thrust/device_malloc_allocator.h @@ -34,10 +34,9 @@ #include #include +#include _THRUST_STD_INCLUDE(__new/bad_alloc.h) #include _THRUST_STD_INCLUDE(limits) -#include - THRUST_NAMESPACE_BEGIN // forward declarations to WAR circular #includes @@ -144,7 +143,7 @@ class device_malloc_allocator { if (cnt > this->max_size()) { - throw std::bad_alloc(); + _THRUST_STD::__throw_bad_alloc(); } // end if return pointer(device_malloc(cnt)); diff --git a/projects/rocthrust/thrust/device_new_allocator.h b/projects/rocthrust/thrust/device_new_allocator.h index 742c246a132d..8663f2bd3124 100644 --- a/projects/rocthrust/thrust/device_new_allocator.h +++ b/projects/rocthrust/thrust/device_new_allocator.h @@ -34,18 +34,16 @@ #include #include -// Use rocprim::numeric_limits if thrust/detail/type_traits.h uses rocprim::arithmetic -#include _THRUST_STD_INCLUDE(limits) +#include _THRUST_STD_INCLUDE(__new/bad_alloc.h) #if _THRUST_HAS_DEVICE_SYSTEM_STD # include _THRUST_STD_INCLUDE(cstdint) -#endif - -#include -#if !_THRUST_HAS_DEVICE_SYSTEM_STD +#else # include #endif +#include _THRUST_STD_INCLUDE(limits) + THRUST_NAMESPACE_BEGIN /*! \addtogroup allocators Allocators @@ -137,7 +135,7 @@ class device_new_allocator { if (cnt > this->max_size()) { - throw std::bad_alloc(); + _THRUST_STD::__throw_bad_alloc(); } // end if // use "::operator new" rather than keyword new From df3e6f617a8ba3c5f065be57bc7f74dbdbb6ade7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Fri, 12 Jun 2026 14:45:23 +0200 Subject: [PATCH 16/86] Update version to 3.1.0 --- projects/rocthrust/thrust/version.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rocthrust/thrust/version.h b/projects/rocthrust/thrust/version.h index 857e70b9e51c..05c3dc497645 100644 --- a/projects/rocthrust/thrust/version.h +++ b/projects/rocthrust/thrust/version.h @@ -63,7 +63,7 @@ * THRUST_VERSION / 100 % 1000 is the minor version. * THRUST_VERSION / 100000 is the major version. */ -#define THRUST_VERSION 300003 // macro expansion with ## requires this to be a single value +#define THRUST_VERSION 300100 // macro expansion with ## requires this to be a single value /*! \def THRUST_MAJOR_VERSION * \brief The preprocessor macro \p THRUST_MAJOR_VERSION encodes the From 3b1728c967ecbca7a1c861809eae6c553365c71c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Mon, 15 Jun 2026 15:23:32 +0200 Subject: [PATCH 17/86] Fix complex type test suite --- projects/rocthrust/test/test_complex_various.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/rocthrust/test/test_complex_various.cpp b/projects/rocthrust/test/test_complex_various.cpp index f1f13aafdd86..8862520f365d 100644 --- a/projects/rocthrust/test/test_complex_various.cpp +++ b/projects/rocthrust/test/test_complex_various.cpp @@ -78,6 +78,8 @@ class VariousComplexTest : public ::testing::Test using T = Type; }; +TYPED_TEST_SUITE(VariousComplexTest, FloatDouble); + template void run_rng_test(const StdFunc& sf, const ThrustFunc& tf) { From 2a03e787b7cdd9c4226abd445469919ef046a87a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Mon, 15 Jun 2026 15:25:50 +0200 Subject: [PATCH 18/86] Replace ``__identity`` --- projects/rocthrust/CHANGELOG.md | 1 + .../detail/libcxx_wrapper/std/__functional/identity.h | 2 +- projects/rocthrust/thrust/system/cuda/detail/gather.h | 6 +++--- projects/rocthrust/thrust/system/cuda/detail/mismatch.h | 2 +- projects/rocthrust/thrust/system/cuda/detail/scatter.h | 6 +++--- 5 files changed, 9 insertions(+), 8 deletions(-) diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index 6a9f42fd7282..2d71cea09a63 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -7,6 +7,7 @@ Documentation for rocThrust available at ### Changed +* Replaced `hip::std::__identity` to `hip::std::identity`. * Replaced bad_alloc, cassert, cfloat, cmath, complex, cstddef, cstdint, cstring, initializer_list, limits, move and type_traits operators to `hip::std` operators. ### Removed diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/identity.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/identity.h index 20b7c16a783e..9554da2961fe 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/identity.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__functional/identity.h @@ -30,7 +30,7 @@ namespace internal { #if _THRUST_HAS_DEVICE_SYSTEM_STD -using identity = _THRUST_STD::__identity; +using identity = _THRUST_STD::identity; #else diff --git a/projects/rocthrust/thrust/system/cuda/detail/gather.h b/projects/rocthrust/thrust/system/cuda/detail/gather.h index 08244d115db3..353e2d29e6b2 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/gather.h +++ b/projects/rocthrust/thrust/system/cuda/detail/gather.h @@ -55,7 +55,7 @@ gather(execution_policy& policy, MapIt map_first, MapIt map_last, Items thrust::make_permutation_iterator(items, map_first), thrust::make_permutation_iterator(items, map_last), result, - ::cuda::std::__identity{}); + ::cuda::std::identity{}); } template @@ -74,7 +74,7 @@ ResultIt _CCCL_HOST_DEVICE gather_if( thrust::make_permutation_iterator(items, map_last), stencil, result, - ::cuda::std::__identity{}, + ::cuda::std::identity{}, predicate); } @@ -82,7 +82,7 @@ template & policy, MapIt map_first, MapIt map_last, StencilIt stencil, ItemsIt items, ResultIt result) { - return cuda_cub::gather_if(policy, map_first, map_last, stencil, items, result, ::cuda::std::__identity{}); + return cuda_cub::gather_if(policy, map_first, map_last, stencil, items, result, ::cuda::std::identity{}); } } // namespace cuda_cub diff --git a/projects/rocthrust/thrust/system/cuda/detail/mismatch.h b/projects/rocthrust/thrust/system/cuda/detail/mismatch.h index d026b973c1c4..8661949c2101 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/mismatch.h +++ b/projects/rocthrust/thrust/system/cuda/detail/mismatch.h @@ -187,7 +187,7 @@ mismatch(execution_policy& policy, InputIt1 first1, InputIt1 last1, Inp const auto transform_first = detail::transform_pair_of_input_iterators_t(first1, first2, binary_pred); const auto result = cuda_cub::find_if_not( - policy, transform_first, transform_first + thrust::distance(first1, last1), ::cuda::std::__identity{}); + policy, transform_first, transform_first + thrust::distance(first1, last1), ::cuda::std::identity{}); return thrust::make_pair(first1 + thrust::distance(transform_first, result), first2 + thrust::distance(transform_first, result)); diff --git a/projects/rocthrust/thrust/system/cuda/detail/scatter.h b/projects/rocthrust/thrust/system/cuda/detail/scatter.h index 80f2ea80eaa0..1b1d60485b7a 100644 --- a/projects/rocthrust/thrust/system/cuda/detail/scatter.h +++ b/projects/rocthrust/thrust/system/cuda/detail/scatter.h @@ -50,7 +50,7 @@ template void _CCCL_HOST_DEVICE scatter(execution_policy& policy, ItemsIt first, ItemsIt last, MapIt map, ResultIt result) { - cuda_cub::transform(policy, first, last, thrust::make_permutation_iterator(result, map), ::cuda::std::__identity{}); + cuda_cub::transform(policy, first, last, thrust::make_permutation_iterator(result, map), ::cuda::std::identity{}); } template @@ -64,14 +64,14 @@ void _CCCL_HOST_DEVICE scatter_if( Predicate predicate) { cuda_cub::transform_if( - policy, first, last, stencil, thrust::make_permutation_iterator(result, map), ::cuda::std::__identity{}, predicate); + policy, first, last, stencil, thrust::make_permutation_iterator(result, map), ::cuda::std::identity{}, predicate); } template void _CCCL_HOST_DEVICE scatter_if( execution_policy& policy, ItemsIt first, ItemsIt last, MapIt map, StencilIt stencil, ResultIt result) { - cuda_cub::scatter_if(policy, first, last, map, stencil, result, ::cuda::std::__identity{}); + cuda_cub::scatter_if(policy, first, last, map, stencil, result, ::cuda::std::identity{}); } } // namespace cuda_cub From efa4126555f01b2a7af666e0f46a3a2aeee93333 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Wed, 17 Jun 2026 09:55:33 +0200 Subject: [PATCH 19/86] Replace deprecated C++17 iterators --- .../libcxx_wrapper/std/__iterator/iterator_traits.h | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__iterator/iterator_traits.h b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__iterator/iterator_traits.h index 2d4182b15a3d..03b852e4322f 100644 --- a/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__iterator/iterator_traits.h +++ b/projects/rocthrust/thrust/detail/libcxx_wrapper/std/__iterator/iterator_traits.h @@ -40,9 +40,13 @@ namespace internal #if _THRUST_HAS_DEVICE_SYSTEM_STD template -using is_cpp17_input_iterator = _THRUST_STD::__is_cpp17_input_iterator; +using is_cpp17_input_iterator = + _THRUST_STD::is_base_of<_THRUST_STD::input_iterator_tag, typename _THRUST_STD::iterator_traits::iterator_category>; + template -using is_cpp17_random_access_iterator = _THRUST_STD::__is_cpp17_random_access_iterator; +using is_cpp17_random_access_iterator = + _THRUST_STD::is_base_of<_THRUST_STD::random_access_iterator_tag, + typename _THRUST_STD::iterator_traits::iterator_category>; #else From 8b6a264f0c45f05843332526b67b106071655511 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Wed, 17 Jun 2026 13:37:52 +0200 Subject: [PATCH 20/86] Fix thrust.compiler_interface_cpp dependency --- projects/rocthrust/thrust/detail/config/config.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/rocthrust/thrust/detail/config/config.h b/projects/rocthrust/thrust/detail/config/config.h index 237bf404f218..3af1bb01a671 100644 --- a/projects/rocthrust/thrust/detail/config/config.h +++ b/projects/rocthrust/thrust/detail/config/config.h @@ -42,3 +42,6 @@ #include // IWYU pragma: export #include // IWYU pragma: export #include // IWYU pragma: export + +// Fix thrust.compiler_interface_cpp dependency +#define THRUST_DISABLE_COMPILER_INTERFACE 1 From 47923aa8d7e92144d99adfeadafe42f700cdfeea Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Thu, 18 Jun 2026 16:17:35 +0200 Subject: [PATCH 21/86] Remove extra semicolons --- projects/rocthrust/test/test_copy.cpp | 2 +- projects/rocthrust/testing/copy.cu | 2 +- projects/rocthrust/testing/unittest/testframework.h | 4 ++-- projects/rocthrust/thrust/detail/contiguous_storage.h | 2 +- projects/rocthrust/thrust/detail/range/head_flags.h | 2 +- projects/rocthrust/thrust/detail/range/tail_flags.h | 2 +- projects/rocthrust/thrust/detail/vector_base.h | 4 ++-- projects/rocthrust/thrust/system/detail/bad_alloc.h | 2 +- projects/rocthrust/thrust/system/system_error.h | 2 +- 9 files changed, 11 insertions(+), 11 deletions(-) diff --git a/projects/rocthrust/test/test_copy.cpp b/projects/rocthrust/test/test_copy.cpp index 40d492d8b5de..edf09e2367ac 100644 --- a/projects/rocthrust/test/test_copy.cpp +++ b/projects/rocthrust/test/test_copy.cpp @@ -685,7 +685,7 @@ struct always_true THRUST_HOST_DEVICE bool operator()(const object_with_non_trivial_ctor&) { return true; - }; + } }; } // namespace diff --git a/projects/rocthrust/testing/copy.cu b/projects/rocthrust/testing/copy.cu index 157246fa16e6..e5db1444199f 100644 --- a/projects/rocthrust/testing/copy.cu +++ b/projects/rocthrust/testing/copy.cu @@ -463,7 +463,7 @@ struct always_true THRUST_HOST_DEVICE bool operator()(const object_with_non_trivial_ctor&) { return true; - }; + } }; } // namespace diff --git a/projects/rocthrust/testing/unittest/testframework.h b/projects/rocthrust/testing/unittest/testframework.h index 588f4aa443be..1ec7bb166806 100644 --- a/projects/rocthrust/testing/unittest/testframework.h +++ b/projects/rocthrust/testing/unittest/testframework.h @@ -198,7 +198,7 @@ class custom_numeric DEFINE_OPERATOR(>) DEFINE_OPERATOR(>=) DEFINE_OPERATOR(&&) - DEFINE_OPERATOR(||); + DEFINE_OPERATOR(||) #undef DEFINE_OPERATOR @@ -341,7 +341,7 @@ class UnitTestDriver virtual bool post_test_smoke_check(const UnitTest& test, bool concise); public: - inline virtual ~UnitTestDriver(){}; + inline virtual ~UnitTestDriver() {} void register_test(UnitTest* test); virtual bool run_tests(const ArgumentSet& args, const ArgumentMap& kwargs); diff --git a/projects/rocthrust/thrust/detail/contiguous_storage.h b/projects/rocthrust/thrust/detail/contiguous_storage.h index 33b6f3a2b31d..551bdb14b279 100644 --- a/projects/rocthrust/thrust/detail/contiguous_storage.h +++ b/projects/rocthrust/thrust/detail/contiguous_storage.h @@ -212,7 +212,7 @@ class contiguous_storage // allow move assignment for a sane implementation of allocator propagation THRUST_HOST_DEVICE contiguous_storage& operator=(contiguous_storage&& other); - THRUST_SYNTHESIZE_SEQUENCE_ACCESS(contiguous_storage, const_iterator); + THRUST_SYNTHESIZE_SEQUENCE_ACCESS(contiguous_storage, const_iterator) private: // XXX we could inherit from this to take advantage of empty base class optimization diff --git a/projects/rocthrust/thrust/detail/range/head_flags.h b/projects/rocthrust/thrust/detail/range/head_flags.h index dd4907af32ee..9d737491e4b7 100644 --- a/projects/rocthrust/thrust/detail/range/head_flags.h +++ b/projects/rocthrust/thrust/detail/range/head_flags.h @@ -93,7 +93,7 @@ class head_flags return *(begin() + i); } - THRUST_SYNTHESIZE_SEQUENCE_ACCESS(head_flags, iterator); + THRUST_SYNTHESIZE_SEQUENCE_ACCESS(head_flags, iterator) private: iterator m_begin; diff --git a/projects/rocthrust/thrust/detail/range/tail_flags.h b/projects/rocthrust/thrust/detail/range/tail_flags.h index 88db60726b5c..2dc00acbcb20 100644 --- a/projects/rocthrust/thrust/detail/range/tail_flags.h +++ b/projects/rocthrust/thrust/detail/range/tail_flags.h @@ -107,7 +107,7 @@ class tail_flags return *(begin() + i); } - THRUST_SYNTHESIZE_SEQUENCE_ACCESS(tail_flags, iterator); + THRUST_SYNTHESIZE_SEQUENCE_ACCESS(tail_flags, iterator) private: iterator m_begin, m_end; diff --git a/projects/rocthrust/thrust/detail/vector_base.h b/projects/rocthrust/thrust/detail/vector_base.h index 56fb727bff37..d4a24f17032f 100644 --- a/projects/rocthrust/thrust/detail/vector_base.h +++ b/projects/rocthrust/thrust/detail/vector_base.h @@ -485,8 +485,8 @@ class vector_base */ allocator_type get_allocator() const; - THRUST_SYNTHESIZE_SEQUENCE_ACCESS(vector_base, const_iterator); - THRUST_SYNTHESIZE_SEQUENCE_REVERSE_ACCESS(vector_base, const_reverse_iterator); + THRUST_SYNTHESIZE_SEQUENCE_ACCESS(vector_base, const_iterator) + THRUST_SYNTHESIZE_SEQUENCE_REVERSE_ACCESS(vector_base, const_reverse_iterator) protected: // Our storage diff --git a/projects/rocthrust/thrust/system/detail/bad_alloc.h b/projects/rocthrust/thrust/system/detail/bad_alloc.h index 583d445ebf22..2b95e11671b8 100644 --- a/projects/rocthrust/thrust/system/detail/bad_alloc.h +++ b/projects/rocthrust/thrust/system/detail/bad_alloc.h @@ -48,7 +48,7 @@ class bad_alloc : public std::bad_alloc m_what += w; } // end bad_alloc() - inline virtual ~bad_alloc() noexcept {}; + inline virtual ~bad_alloc() noexcept {} inline virtual const char* what() const noexcept { diff --git a/projects/rocthrust/thrust/system/system_error.h b/projects/rocthrust/thrust/system/system_error.h index 70ed79124953..d4cb12022940 100644 --- a/projects/rocthrust/thrust/system/system_error.h +++ b/projects/rocthrust/thrust/system/system_error.h @@ -146,7 +146,7 @@ class system_error : public std::runtime_error /*! Destructor does not throw. */ - inline virtual ~system_error() noexcept {}; + inline virtual ~system_error() noexcept {} /*! Returns an object encoding the error. * \return ec or error_code(ev, ecat), from the From b3630fb5f3326dda70043f5025413f421182f85a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Thu, 18 Jun 2026 16:43:14 +0200 Subject: [PATCH 22/86] Drop superfluous semicolons --- projects/rocthrust/testing/unittest/assertions.h | 2 +- projects/rocthrust/testing/unittest/exceptions.h | 4 ++-- projects/rocthrust/testing/unittest/random.h | 4 ++-- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/rocthrust/testing/unittest/assertions.h b/projects/rocthrust/testing/unittest/assertions.h index 9ee3bed874ac..ee27cdae252e 100644 --- a/projects/rocthrust/testing/unittest/assertions.h +++ b/projects/rocthrust/testing/unittest/assertions.h @@ -790,4 +790,4 @@ void check_assert_throws( } } -}; // end namespace unittest +} // end namespace unittest diff --git a/projects/rocthrust/testing/unittest/exceptions.h b/projects/rocthrust/testing/unittest/exceptions.h index 50e424f2dc59..08f949536821 100644 --- a/projects/rocthrust/testing/unittest/exceptions.h +++ b/projects/rocthrust/testing/unittest/exceptions.h @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -76,4 +76,4 @@ class UnitTestKnownFailure : public UnitTestException {} }; -}; // end namespace unittest +} // end namespace unittest diff --git a/projects/rocthrust/testing/unittest/random.h b/projects/rocthrust/testing/unittest/random.h index be7e7d217ebd..fcfeafb540ef 100644 --- a/projects/rocthrust/testing/unittest/random.h +++ b/projects/rocthrust/testing/unittest/random.h @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -139,4 +139,4 @@ THRUST_NS_QUALIFIER::host_vector random_samples(const size_t N) return vec; } -}; // end namespace unittest +} // end namespace unittest From daf3d03ca9918374e5409baf210ec92c773d4244 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Thu, 18 Jun 2026 16:36:07 +0200 Subject: [PATCH 23/86] Attempt to recover from upstream OOM in disjoint_pool --- projects/rocthrust/rtest.xml | 2 +- .../rocthrust/test/test_mr_disjoint_pool.cpp | 267 +++++++++++++++++- .../rocthrust/testing/mr_disjoint_pool.cu | 246 +++++++++++++++- projects/rocthrust/thrust/mr/disjoint_pool.h | 104 +++++-- 4 files changed, 579 insertions(+), 40 deletions(-) diff --git a/projects/rocthrust/rtest.xml b/projects/rocthrust/rtest.xml index bf1ff1e58e6a..96effa1cdd0c 100644 --- a/projects/rocthrust/rtest.xml +++ b/projects/rocthrust/rtest.xml @@ -2,7 +2,7 @@ - + diff --git a/projects/rocthrust/test/test_mr_disjoint_pool.cpp b/projects/rocthrust/test/test_mr_disjoint_pool.cpp index 5b6143d37afc..952b8c05858e 100644 --- a/projects/rocthrust/test/test_mr_disjoint_pool.cpp +++ b/projects/rocthrust/test/test_mr_disjoint_pool.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2019-2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2019-2026 Advanced Micro Devices, Inc. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -27,7 +27,7 @@ struct alloc_id std::size_t id; std::size_t size; std::size_t alignment; - std::size_t offset; + std::size_t offset{}; __host__ __device__ bool operator==(const alloc_id& other) const { @@ -64,20 +64,43 @@ struct thrust::detail::pointer_traits class dummy_resource final : public thrust::mr::memory_resource { public: - dummy_resource() - : id_to_allocate(0) - , id_to_deallocate(0) - {} + dummy_resource() = default; ~dummy_resource() { EXPECT_EQ(id_to_allocate, 0u); EXPECT_EQ(id_to_deallocate, 0u); + EXPECT_EQ(used_bytes, 0u); + EXPECT_EQ(allocation_ids.size(), 0u); + } + + void assert_empty_and_reset() + { + ASSERT_EQ(used_bytes, 0u); + ASSERT_EQ(allocation_ids.size(), 0u); + + free_bytes = 1ull << 63; + id_to_allocate = 0; + id_to_deallocate = 0; } virtual alloc_id do_allocate(std::size_t bytes, std::size_t alignment) override { - EXPECT_EQ(static_cast(id_to_allocate), true); + if (bytes > free_bytes) + { + throw thrust::system::detail::bad_alloc("Dummy allocation failed: insufficient free bytes."); + } + + // id_to_allocate must not be zero + EXPECT_NE(id_to_allocate, 0u); + + // id_to_allocate must not already exist in allocation_ids + auto it = std::find(allocation_ids.begin(), allocation_ids.end(), id_to_allocate); + EXPECT_EQ(it, allocation_ids.end()) << "Allocation ID already exists"; + + free_bytes -= bytes; + used_bytes += bytes; + allocation_ids.push_back(id_to_allocate); alloc_id ret; ret.id = id_to_allocate; @@ -93,6 +116,13 @@ class dummy_resource final : public thrust::mr::memory_resource { ASSERT_EQ(p.size, bytes); ASSERT_EQ(p.alignment, alignment); + ASSERT_LE(bytes, used_bytes); + // Check that the id has been previously allocated + ASSERT_NE(find(allocation_ids.begin(), allocation_ids.end(), p.id), allocation_ids.end()); + + free_bytes += bytes; + used_bytes -= bytes; + allocation_ids.erase(find(allocation_ids.begin(), allocation_ids.end(), p.id)); if (id_to_deallocate != 0) { @@ -101,8 +131,12 @@ class dummy_resource final : public thrust::mr::memory_resource } } - std::size_t id_to_allocate; - std::size_t id_to_deallocate; + std::size_t free_bytes{1ull << 63}; + std::size_t used_bytes{0}; + std::vector allocation_ids; + + std::size_t id_to_allocate{}; + std::size_t id_to_deallocate{}; }; template