diff --git a/projects/rocthrust/.clang-format b/projects/rocthrust/.clang-format index 676f9822da5f..991d86cb827f 100644 --- a/projects/rocthrust/.clang-format +++ b/projects/rocthrust/.clang-format @@ -28,12 +28,12 @@ AttributeMacros: [ 'THRUST_FORCEINLINE', 'THRUST_HOST_DEVICE', 'THRUST_HOST', - 'THRUST_NODISCARD', 'THRUST_DEPRECATED', 'THRUST_GLOBAL_CONSTANT', # thrust '_CCCL_ALIGNAS_TYPE', '_CCCL_ALIGNAS', + '_CCCL_API', '_CCCL_CONSTEXPR_CXX14', '_CCCL_CONSTEXPR_CXX17', '_CCCL_CONSTEXPR_CXX20', @@ -45,7 +45,6 @@ AttributeMacros: [ '_CCCL_HOST', '_CCCL_NO_UNIQUE_ADDRESS', '_CCCL_NODISCARD_FRIEND', - '_CCCL_NODISCARD', '_CCCL_NORETURN', '_CCCL_VISIBILITY_HIDDEN', 'CUB_RUNTIME_FUNCTION', @@ -66,8 +65,6 @@ AttributeMacros: [ '_LIBCUDACXX_EXPORTED_FROM_ABI', '_LIBCUDACXX_EXTERN_TEMPLATE_TYPE_VIS', '_LIBCUDACXX_HIDDEN', - '_LIBCUDACXX_HIDE_FROM_ABI_AFTER_V1', - '_LIBCUDACXX_HIDE_FROM_ABI', '_LIBCUDACXX_INLINE_VISIBILITY', '_LIBCUDACXX_INTERNAL_LINKAGE', '_LIBCUDACXX_METHOD_TEMPLATE_IMPLICIT_INSTANTIATION_VIS', diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index 31cc541827bc..c46a6f9c66da 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -3,8 +3,31 @@ 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) + +### Added + +* Add `THRUST_DISABLE_BROKEN_GCC_VECTORIZER` +* Add `thrust::reduce_into`. +* Add `thrust::shuffle_iterator`. +* Add `thrust::transform_n` and `thrust::transform_if_n`. + +### Changed + +* Replaced bad_alloc, cassert, cfloat, cmath, complex, cstddef, cstdint, cstring, initializer_list, limits, move and type_traits operators with `hip::std` operators. +* Replaced `hip::std::__identity` with `hip::std::identity`. +* Replaced ``_LIBCUDACXX_HIDE_FROM_ABI`` with ``_CCCL_API``. +* Defined `THRUST_COMP_DEPR_IMPL` with `_CCCL_WARNING`. +* Replaced `thrust::constant_iterator`, `thrust::make_discard_iterator`, `thrust::permutation_iterator`, `thrust::strided_iterator`, `thrust::tabulate_output_iterator` and `thrust::transform_output_iterator` with `hip::constant_iterator`, `hip::make_discard_iterator`, `hip::permutation_iterator`, `hip::strided_iterator`, `hip::tabulate_output_iterator` and `hip::transform_output_iterator` instead. + +### Removed + +* Deprecated `thrust::advance`, `thrust::bit_and`, `thrust::bit_or`, `thrust::bit_xor`, `thrust::distance`, `thrust::equal_to`, `thrust::greater`, `thrust::greater_equal`, `thrust::less`, `thrust::less_equal`, `thrust::logical_and`, `thrust::logical_not`, `thrust::logical_or`, `thrust::maximum`, `thrust::minimum`, `thrust::modulus`, `thrust::multiplies`, `thrust::negate`, `thrust::next`, `thrust::not_equal_to`, `thrust::not_fn`, `thrust::operator_name`, and `thrust::plus`. Use `hip::std::advance`, `hip::std::bit_and`, `hip::std::bit_or`, `hip::std::bit_xor`, `hip::std::distance`, `hip::std::equal_to`, `hip::std::greater`, `hip::std::greater_equal`, `hip::std::less`, `hip::std::less_equal`, `hip::std::logical_and`, `hip::std::logical_not`, `hip::std::logical_or`, `hip::maximum`, `hip::minimum`, `hip::modulus`, `hip::std::multiplies`, `hip::std::negate`, `hip::std::next`, `hip::std::not_equal_to`, `hip::std::not_fn`, `hip::std::operator_name` and `hip::std::plus` instead. +* Deprecated `thrust::detail::::identity_`. Use `hip::std::type_identity` instead. +* 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/benchmark/bench/copy/if.cu b/projects/rocthrust/benchmark/bench/copy/if.cu index c2407f76985d..b8e394a4f60d 100644 --- a/projects/rocthrust/benchmark/bench/copy/if.cu +++ b/projects/rocthrust/benchmark/bench/copy/if.cu @@ -38,17 +38,6 @@ PRIMBENCH_REGISTER_TYPE(bench_utils::large_data, "large_data") -template -struct less_then_t -{ - T m_val; - - __host__ __device__ bool operator()(const T& val) const - { - return val < m_val; - } -}; - template struct copy_benchmark : public primbench::benchmark_interface { @@ -69,8 +58,8 @@ struct copy_benchmark : public primbench::benchmark_interface void run(primbench::state& state) override { - T val = bench_utils::value_from_entropy(bench_utils::get_entropy_percentage(entropy)); - less_then_t select_op{val}; + const T val = bench_utils::lerp_min_max(bench_utils::get_entropy_percentage(entropy)); + bench_utils::less_then_t select_op{val}; bench_utils::caching_allocator_t alloc{}; thrust::detail::device_t policy{}; diff --git a/projects/rocthrust/benchmark/bench/partition/basic.cu b/projects/rocthrust/benchmark/bench/partition/basic.cu index e645c7e918e8..3226acb93c4c 100644 --- a/projects/rocthrust/benchmark/bench/partition/basic.cu +++ b/projects/rocthrust/benchmark/bench/partition/basic.cu @@ -34,17 +34,6 @@ #include "bench_utils.hpp" -template -struct less_then_t -{ - T m_val; - - __host__ __device__ bool operator()(const T& val) const - { - return val < m_val; - } -}; - template struct partition_benchmark : public primbench::benchmark_interface { @@ -68,9 +57,9 @@ struct partition_benchmark : public primbench::benchmark_interface bench_utils::caching_allocator_t alloc{}; thrust::detail::device_t policy{}; - T val = bench_utils::value_from_entropy(bench_utils::get_entropy_percentage(entropy_reduction)); + const T val = bench_utils::lerp_min_max(bench_utils::get_entropy_percentage(entropy_reduction)); - less_then_t select_op{val}; + bench_utils::less_then_t select_op{val}; thrust::device_vector in = bench_utils::generate(m_items, state.seed); thrust::device_vector out(m_items); diff --git a/projects/rocthrust/benchmark/bench/reduce/by_key.cu b/projects/rocthrust/benchmark/bench/reduce/by_key.cu index b7be50e18e20..5e44e40dd6d5 100644 --- a/projects/rocthrust/benchmark/bench/reduce/by_key.cu +++ b/projects/rocthrust/benchmark/bench/reduce/by_key.cu @@ -71,7 +71,8 @@ struct reduce_benchmark : public primbench::benchmark_interface out_keys = in_keys; in_vals = thrust::device_vector(m_items); - const size_t unique_keys = thrust::distance(out_keys.begin(), thrust::unique(out_keys.begin(), out_keys.end())); + const size_t unique_keys = + _THRUST_STD::distance(out_keys.begin(), thrust::unique(out_keys.begin(), out_keys.end())); thrust::device_vector out_vals; out_vals = thrust::device_vector(unique_keys); diff --git a/projects/rocthrust/benchmark/bench/set_operations/base.hpp b/projects/rocthrust/benchmark/bench/set_operations/base.hpp index 4b410bf1a4d1..8c8214e0c3a6 100644 --- a/projects/rocthrust/benchmark/bench/set_operations/base.hpp +++ b/projects/rocthrust/benchmark/bench/set_operations/base.hpp @@ -83,7 +83,7 @@ struct base_set_benchmark : public primbench::benchmark_interface in.cend(), out.begin()); - const size_t items_in_AB = thrust::distance(out.begin(), result_ends); + const size_t items_in_AB = _THRUST_STD::distance(out.begin(), result_ends); state.set_items(m_items); state.add_reads(m_items); diff --git a/projects/rocthrust/benchmark/bench/set_operations/by_key.hpp b/projects/rocthrust/benchmark/bench/set_operations/by_key.hpp index c8b3e845888e..8b8e5f062c77 100644 --- a/projects/rocthrust/benchmark/bench/set_operations/by_key.hpp +++ b/projects/rocthrust/benchmark/bench/set_operations/by_key.hpp @@ -90,7 +90,7 @@ struct by_key_set_benchmark : public primbench::benchmark_interface out_keys.begin(), out_vals.begin()); - const size_t item_in_AB = thrust::distance(out_keys.begin(), result_ends.first); + const size_t item_in_AB = _THRUST_STD::distance(out_keys.begin(), result_ends.first); const size_t values_read = OpT::read_all_values ? m_items : items_in_A; state.set_items(m_items); diff --git a/projects/rocthrust/benchmark/bench/transform/basic.cu b/projects/rocthrust/benchmark/bench/transform/basic.cu index 5ba9cd10899f..e96824daf69d 100644 --- a/projects/rocthrust/benchmark/bench/transform/basic.cu +++ b/projects/rocthrust/benchmark/bench/transform/basic.cu @@ -33,7 +33,6 @@ #include #include #include -#include #include #include #include @@ -174,7 +173,7 @@ struct mul { const T scalar = startScalar; thrust::transform( - policy, c.begin(), c.end(), b.begin(), ::internal::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ci) { + policy, c.begin(), c.end(), b.begin(), _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ci) { return ci * scalar; })); } @@ -195,7 +194,7 @@ struct add a.end(), b.begin(), c.begin(), - ::internal::proclaim_copyable_arguments([] THRUST_DEVICE(const T& ai, const T& bi) -> T { + _THRUST_LIBCXX::proclaim_copyable_arguments([] THRUST_DEVICE(const T& ai, const T& bi) -> T { return ai + bi; })); } @@ -217,7 +216,7 @@ struct triad b.end(), c.begin(), a.begin(), - ::internal::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& bi, const T& ci) { + _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& bi, const T& ci) { return bi + scalar * ci; })); } @@ -239,7 +238,7 @@ struct nstream thrust::make_zip_iterator(a.end(), b.end(), c.end()), a.begin(), thrust::make_zip_function( - ::internal::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ai, const T& bi, const T& ci) { + _THRUST_LIBCXX::proclaim_copyable_arguments([=] THRUST_DEVICE(const T& ai, const T& bi, const T& ci) { return ai + bi + scalar * ci; }))); } diff --git a/projects/rocthrust/benchmark/bench/transform_reduce/sum.cu b/projects/rocthrust/benchmark/bench/transform_reduce/sum.cu index d34b861ca448..3197072164b9 100644 --- a/projects/rocthrust/benchmark/bench/transform_reduce/sum.cu +++ b/projects/rocthrust/benchmark/bench/transform_reduce/sum.cu @@ -72,7 +72,7 @@ struct transform_reduce_benchmark : public primbench::benchmark_interface state.run([&] { thrust::transform_reduce( - policy(alloc).on(state.stream), in.begin(), in.end(), square_t{}, T{}, thrust::plus{}); + policy(alloc).on(state.stream), in.begin(), in.end(), square_t{}, T{}, _THRUST_STD::plus{}); }); } diff --git a/projects/rocthrust/benchmark/bench/unique/basic.cu b/projects/rocthrust/benchmark/bench/unique/basic.cu index a6fcf20af7b8..f8d3b119c6c3 100644 --- a/projects/rocthrust/benchmark/bench/unique/basic.cu +++ b/projects/rocthrust/benchmark/bench/unique/basic.cu @@ -65,7 +65,7 @@ struct unique_benchmark : public primbench::benchmark_interface // not a warm-up run, we need to run once to determine the size of the output const auto new_end = thrust::unique_copy(policy(alloc).on(state.stream), in.cbegin(), in.cend(), out.begin()); - const size_t unique_items = thrust::distance(out.begin(), new_end); + const size_t unique_items = _THRUST_STD::distance(out.begin(), new_end); state.set_items(m_items); state.add_reads(m_items); diff --git a/projects/rocthrust/benchmark/bench/unique/by_key.cu b/projects/rocthrust/benchmark/bench/unique/by_key.cu index ce762a7053b2..bd96a5b390fe 100644 --- a/projects/rocthrust/benchmark/bench/unique/by_key.cu +++ b/projects/rocthrust/benchmark/bench/unique/by_key.cu @@ -93,7 +93,7 @@ struct unique_benchmark : public primbench::benchmark_interface out_keys.begin(), out_vals.begin()); - const size_t unique_elements = thrust::distance(out_keys.begin(), new_key_end); + const size_t unique_elements = _THRUST_STD::distance(out_keys.begin(), new_key_end); state.set_items(m_items); state.add_reads(m_items); diff --git a/projects/rocthrust/benchmark/generation_utils.hpp b/projects/rocthrust/benchmark/generation_utils.hpp index 13350a832821..df946136e5c8 100644 --- a/projects/rocthrust/benchmark/generation_utils.hpp +++ b/projects/rocthrust/benchmark/generation_utils.hpp @@ -45,6 +45,10 @@ // rocRAND #include +// LIBHIPCXX +#include _THRUST_STD_INCLUDE(cmath) +#include _THRUST_STD_INCLUDE(functional) + // STL #include #include @@ -85,25 +89,29 @@ float get_entropy_percentage(int entropy_reduction) } } +// Creates an interpolated value of type T between min (at = 0.0) and max (at = 1.0). template -T value_from_entropy(double percentage) +[[nodiscard]] T lerp_min_max(double at) noexcept { - if (percentage == 100.0) + if (at == 1.0) { - return std::numeric_limits::max(); + return _THRUST_STD::numeric_limits::max(); } - - percentage /= 100; // convert percentage to per one - - // Select value from the line between the lowest and the highest representable - // values of type T based on the entropy value. - const auto max_val = static_cast(std::numeric_limits::max()); - const auto min_val = static_cast(std::numeric_limits::lowest()); - const auto result = min_val + percentage * (max_val - min_val); - return static_cast(result); + const auto min_val = static_cast(_THRUST_STD::numeric_limits::lowest()); + const auto max_val = static_cast(_THRUST_STD::numeric_limits::max()); + return static_cast(_THRUST_STD::lerp(min_val, max_val, at)); } -const int entropy_reductions[] = {0, 2, 4, 6}; +template +struct less_then_t +{ + T m_val; + + [[nodiscard]] __device__ bool operator()(const T& val) const noexcept + { + return val < m_val; + } +}; namespace detail { @@ -353,7 +361,7 @@ std::size_t gen_uniform_offsets( policy, segment_offsets.data(), segment_offsets.data() + segment_offsets.size(), - segment_offsets.data() /*, thrust::plus<>{}*/); + segment_offsets.data() /*, _THRUST_STD::plus<>{}*/); // Find first sum of offsets greater than 'elements', we are sure that there is // going to be one because we added elements + 1 at the end of the segment_offsets. @@ -361,7 +369,7 @@ std::size_t gen_uniform_offsets( policy, segment_offsets.data(), segment_offsets.data() + segment_offsets.size(), geq_t{elements}); // Compute the element's index. - auto dist = thrust::distance(segment_offsets.data(), iter); + auto dist = _THRUST_STD::distance(segment_offsets.data(), iter); // Fill next item with 'elements'. thrust::fill_n(policy, segment_offsets.data() + dist, 1, elements); // Return next item's index. diff --git a/projects/rocthrust/doc/CHANGELOG.md b/projects/rocthrust/doc/CHANGELOG.md index 470051873f19..b242506bfebd 100644 --- a/projects/rocthrust/doc/CHANGELOG.md +++ b/projects/rocthrust/doc/CHANGELOG.md @@ -1605,8 +1605,8 @@ Finally, a new TBB reduce_by_key implementation provides 80% faster - `thrust::stable_partition_copy with stencil` - `thrust::tabulate` - Memory Allocation: - - `thrust::malloc` - - `thrust::free` + - `thrust::malloc` + - `thrust::free` - `thrust::get_temporary_buffer` - `thrust::return_temporary_buffer` 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/examples/constant_iterator.cu b/projects/rocthrust/examples/constant_iterator.cu index 11486c6d5a74..5e90f358eb26 100644 --- a/projects/rocthrust/examples/constant_iterator.cu +++ b/projects/rocthrust/examples/constant_iterator.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -36,7 +36,8 @@ int main(void) data[3] = 5; // add 10 to all values in data - thrust::transform(data.begin(), data.end(), thrust::constant_iterator(10), data.begin(), thrust::plus()); + thrust::transform( + data.begin(), data.end(), thrust::constant_iterator(10), data.begin(), _THRUST_STD::plus()); // data is now [13, 17, 12, 15] diff --git a/projects/rocthrust/examples/counting_iterator.cu b/projects/rocthrust/examples/counting_iterator.cu index 24a4e546895e..83d8c7007ecb 100644 --- a/projects/rocthrust/examples/counting_iterator.cu +++ b/projects/rocthrust/examples/counting_iterator.cu @@ -19,7 +19,6 @@ // SOFTWARE. #include -#include #include #include #include @@ -52,7 +51,7 @@ int main() // compute indices of nonzero elements using IndexIterator = thrust::device_vector::iterator; - IndexIterator indices_end = thrust::copy_if(first, last, stencil.begin(), indices.begin(), ::internal::identity{}); + IndexIterator indices_end = thrust::copy_if(first, last, stencil.begin(), indices.begin(), _THRUST_STD::identity{}); // indices now contains [1,2,5,7] // print result diff --git a/projects/rocthrust/examples/cuda/async_reduce.cu b/projects/rocthrust/examples/cuda/async_reduce.cu index 4e1120155acb..4c92471e7c33 100644 --- a/projects/rocthrust/examples/cuda/async_reduce.cu +++ b/projects/rocthrust/examples/cuda/async_reduce.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -59,7 +59,7 @@ int main() cudaStreamCreate(&s); // launch a CUDA kernel with only 1 thread on our stream - reduce_kernel<<<1, 1, 0, s>>>(data.begin(), data.end(), 0, thrust::plus(), result.data()); + reduce_kernel<<<1, 1, 0, s>>>(data.begin(), data.end(), 0, _THRUST_STD::plus(), result.data()); // wait for the stream to finish cudaStreamSynchronize(s); @@ -77,7 +77,7 @@ int main() auto begin = data.begin(); auto end = data.end(); unsigned int init = 0; - auto binary_op = thrust::plus(); + auto binary_op = _THRUST_STD::plus(); // std::async captures the algorithm parameters by value // use std::launch::async to ensure the creation of a new thread diff --git a/projects/rocthrust/examples/cuda/range_view.cu b/projects/rocthrust/examples/cuda/range_view.cu index b32c82bc5f0b..ce747e095869 100644 --- a/projects/rocthrust/examples/cuda/range_view.cu +++ b/projects/rocthrust/examples/cuda/range_view.cu @@ -60,7 +60,7 @@ public: __host__ __device__ difference_type size() const { - return thrust::distance(first, last); + return _THRUST_STD::distance(first, last); } __host__ __device__ reference operator[](difference_type n) @@ -216,7 +216,7 @@ int main() make_range_view(thrust::make_transform_iterator(X.cbegin(), f1()), thrust::make_transform_iterator(X.cend(), f1())), // range view of normal_iterators - make_range_view(Y.begin(), thrust::distance(Y.begin(), Y.end())), + make_range_view(Y.begin(), _THRUST_STD::distance(Y.begin(), Y.end())), // range view of naked pointers make_range_view(Z.data().get(), 4)); diff --git a/projects/rocthrust/examples/expand.cu b/projects/rocthrust/examples/expand.cu index 974a58fd444b..33d86e93e29b 100644 --- a/projects/rocthrust/examples/expand.cu +++ b/projects/rocthrust/examples/expand.cu @@ -44,7 +44,7 @@ OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator // otherwise it falls back to ::std. using difference_type = typename _THRUST_STD::iterator_traits::difference_type; - difference_type input_size = thrust::distance(first1, last1); + difference_type input_size = _THRUST_STD::distance(first1, last1); difference_type output_size = thrust::reduce(first1, last1); // scan the counts to obtain output offsets for each input element @@ -62,13 +62,13 @@ OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator // compute max-scan over the output indices, filling in the holes thrust::inclusive_scan( - output_indices.begin(), output_indices.end(), output_indices.begin(), thrust::maximum()); + output_indices.begin(), output_indices.end(), output_indices.begin(), _THRUST_LIBCXX::maximum()); // gather input values according to index array (output = first2[output_indices]) thrust::gather(output_indices.begin(), output_indices.end(), first2, output); // return output + output_size - thrust::advance(output, output_size); + _THRUST_STD::advance(output, output_size); return output; } diff --git a/projects/rocthrust/examples/histogram.cu b/projects/rocthrust/examples/histogram.cu index 768c7b7677f3..be8d50b710d2 100644 --- a/projects/rocthrust/examples/histogram.cu +++ b/projects/rocthrust/examples/histogram.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -139,8 +139,8 @@ void sparse_histogram(const Vector1& input, Vector2& histogram_values, Vector3& data.end() - 1, data.begin() + 1, IndexType(1), - thrust::plus(), - thrust::not_equal_to()); + _THRUST_STD::plus(), + _THRUST_STD::not_equal_to()); // resize histogram storage histogram_values.resize(num_bins); diff --git a/projects/rocthrust/examples/max_abs_diff.cu b/projects/rocthrust/examples/max_abs_diff.cu index f0d83b262260..441ce7f10991 100644 --- a/projects/rocthrust/examples/max_abs_diff.cu +++ b/projects/rocthrust/examples/max_abs_diff.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -55,7 +55,7 @@ int main(void) float init = 0; // binary operations - thrust::maximum binary_op1; + _THRUST_LIBCXX::maximum binary_op1; abs_diff binary_op2; float max_abs_diff = thrust::inner_product(d_a.begin(), d_a.end(), d_b.begin(), init, binary_op1, binary_op2); diff --git a/projects/rocthrust/examples/minimal_custom_backend.cu b/projects/rocthrust/examples/minimal_custom_backend.cu index 4d9213cf3ee8..2c294d13ebdf 100644 --- a/projects/rocthrust/examples/minimal_custom_backend.cu +++ b/projects/rocthrust/examples/minimal_custom_backend.cu @@ -29,8 +29,6 @@ # include #endif -#include - // This example demonstrates how to build a minimal custom // Thrust backend by intercepting for_each's dispatch. @@ -77,7 +75,7 @@ int main() // cause our version of for_each to be invoked when we pass an instance of my_system as the first parameter. // Even though we did not define a special version of transform, Thrust dispatches the version it knows // for thrust::device_execution_policy, which my_system inherits. - thrust::transform(sys, vec.begin(), vec.end(), vec.begin(), ::internal::identity{}); + thrust::transform(sys, vec.begin(), vec.end(), vec.begin(), _THRUST_STD::identity{}); // Invocations without my_system are handled normally. // Note: _THRUST_STD refers to ::hip::std if libhipcxx is available; diff --git a/projects/rocthrust/examples/monte_carlo.cu b/projects/rocthrust/examples/monte_carlo.cu index 4971a89b04fc..634552eb7df3 100644 --- a/projects/rocthrust/examples/monte_carlo.cu +++ b/projects/rocthrust/examples/monte_carlo.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -88,7 +88,11 @@ int main(void) int M = 30000; float estimate = thrust::transform_reduce( - thrust::counting_iterator(0), thrust::counting_iterator(M), estimate_pi(), 0.0f, thrust::plus()); + thrust::counting_iterator(0), + thrust::counting_iterator(M), + estimate_pi(), + 0.0f, + _THRUST_STD::plus()); estimate /= M; std::cout << std::setprecision(3); diff --git a/projects/rocthrust/examples/monte_carlo_disjoint_sequences.cu b/projects/rocthrust/examples/monte_carlo_disjoint_sequences.cu index c7062ed2b95e..39f342fd4cb1 100644 --- a/projects/rocthrust/examples/monte_carlo_disjoint_sequences.cu +++ b/projects/rocthrust/examples/monte_carlo_disjoint_sequences.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -95,7 +95,11 @@ int main(void) int M = 30000; float estimate = thrust::transform_reduce( - thrust::counting_iterator(0), thrust::counting_iterator(M), estimate_pi(), 0.0f, thrust::plus()); + thrust::counting_iterator(0), + thrust::counting_iterator(M), + estimate_pi(), + 0.0f, + _THRUST_STD::plus()); estimate /= M; std::cout << "pi is around " << estimate << std::endl; diff --git a/projects/rocthrust/examples/norm.cu b/projects/rocthrust/examples/norm.cu index 53857133e888..548a21f1960d 100644 --- a/projects/rocthrust/examples/norm.cu +++ b/projects/rocthrust/examples/norm.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -59,7 +59,7 @@ int main() // setup arguments square unary_op; - thrust::plus binary_op; + _THRUST_STD::plus binary_op; float init = 0; // compute norm diff --git a/projects/rocthrust/examples/padded_grid_reduction.cu b/projects/rocthrust/examples/padded_grid_reduction.cu index b536661362e4..1dab677812fa 100644 --- a/projects/rocthrust/examples/padded_grid_reduction.cu +++ b/projects/rocthrust/examples/padded_grid_reduction.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -131,8 +131,8 @@ int main(void) reduce_tuple binary_op; // reduction operator result_type result = thrust::transform_reduce( - thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator(0), data.begin())), - thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator(0), data.begin())) + data.size(), + thrust::make_zip_iterator(thrust::counting_iterator(0), data.begin()), + thrust::make_zip_iterator(thrust::counting_iterator(0), data.begin()) + data.size(), unary_op, init, binary_op); diff --git a/projects/rocthrust/examples/remove_points2d.cu b/projects/rocthrust/examples/remove_points2d.cu index a2f9b380b716..264897bc7c6b 100644 --- a/projects/rocthrust/examples/remove_points2d.cu +++ b/projects/rocthrust/examples/remove_points2d.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -78,10 +78,10 @@ int main(void) // remove points where x^2 + y^2 > 1 and determine new array sizes size_t new_size = - thrust::remove_if(thrust::make_zip_iterator(thrust::make_tuple(x.begin(), y.begin())), - thrust::make_zip_iterator(thrust::make_tuple(x.end(), y.end())), + thrust::remove_if(thrust::make_zip_iterator(x.begin(), y.begin()), + thrust::make_zip_iterator(x.end(), y.end()), is_outside_circle()) - - thrust::make_zip_iterator(thrust::make_tuple(x.begin(), y.begin())); + - thrust::make_zip_iterator(x.begin(), y.begin()); // resize the vectors (note: this does not free any memory) x.resize(new_size); diff --git a/projects/rocthrust/examples/saxpy.cu b/projects/rocthrust/examples/saxpy.cu index 41a7cb5e1aae..1ab0f34905be 100644 --- a/projects/rocthrust/examples/saxpy.cu +++ b/projects/rocthrust/examples/saxpy.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -66,10 +66,10 @@ void saxpy_slow(float A, thrust::device_vector& X, thrust::device_vector< thrust::fill(temp.begin(), temp.end(), A); // temp <- A * X - thrust::transform(X.begin(), X.end(), temp.begin(), temp.begin(), thrust::multiplies()); + thrust::transform(X.begin(), X.end(), temp.begin(), temp.begin(), _THRUST_STD::multiplies()); // Y <- A * X + Y - thrust::transform(temp.begin(), temp.end(), Y.begin(), Y.begin(), thrust::plus()); + thrust::transform(temp.begin(), temp.end(), Y.begin(), Y.begin(), _THRUST_STD::plus()); } int main(void) diff --git a/projects/rocthrust/examples/scan_by_key.cu b/projects/rocthrust/examples/scan_by_key.cu index 806b5e094de9..4aa70be3d2f6 100644 --- a/projects/rocthrust/examples/scan_by_key.cu +++ b/projects/rocthrust/examples/scan_by_key.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -27,7 +27,7 @@ #include "include/host_device.h" // BinaryPredicate for the head flag segment representation -// equivalent to thrust::not_fn(thrust::project2nd())); +// equivalent to _THRUST_STD::not_fn(thrust::project2nd())); template struct head_flag_predicate { diff --git a/projects/rocthrust/examples/sort.cu b/projects/rocthrust/examples/sort.cu index b4aa29cd169b..6c52bdcf14c0 100644 --- a/projects/rocthrust/examples/sort.cu +++ b/projects/rocthrust/examples/sort.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -148,7 +148,7 @@ int main(void) thrust::device_vector keys(N); initialize(keys); print(keys); - thrust::sort(keys.begin(), keys.end(), thrust::greater()); + thrust::sort(keys.begin(), keys.end(), _THRUST_STD::greater()); print(keys); } @@ -195,7 +195,7 @@ int main(void) thrust::device_vector values(N); initialize(keys, values); print(keys, values); - thrust::sort_by_key(keys.begin(), keys.end(), values.begin(), thrust::greater()); + thrust::sort_by_key(keys.begin(), keys.end(), values.begin(), _THRUST_STD::greater()); print(keys, values); } diff --git a/projects/rocthrust/examples/sparse_vector.cu b/projects/rocthrust/examples/sparse_vector.cu index 5e7ec67d99c3..652ce4cdca17 100644 --- a/projects/rocthrust/examples/sparse_vector.cu +++ b/projects/rocthrust/examples/sparse_vector.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -86,8 +86,8 @@ void sum_sparse_vectors( temp_index.end() - 1, temp_index.begin() + 1, size_t(0), - thrust::plus(), - thrust::not_equal_to()) + _THRUST_STD::plus(), + _THRUST_STD::not_equal_to()) + 1; // allocate space for output @@ -101,8 +101,8 @@ void sum_sparse_vectors( temp_value.begin(), C_index.begin(), C_value.begin(), - thrust::equal_to(), - thrust::plus()); + _THRUST_STD::equal_to(), + _THRUST_STD::plus()); } int main(void) diff --git a/projects/rocthrust/examples/sum.cu b/projects/rocthrust/examples/sum.cu index 19911c47d767..9df6da929a8e 100644 --- a/projects/rocthrust/examples/sum.cu +++ b/projects/rocthrust/examples/sum.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -45,7 +45,7 @@ int main(void) int init = 0; // binary operation used to reduce values - thrust::plus binary_op; + _THRUST_STD::plus binary_op; // compute sum on the device int sum = thrust::reduce(d_vec.begin(), d_vec.end(), init, binary_op); diff --git a/projects/rocthrust/examples/sum_rows.cu b/projects/rocthrust/examples/sum_rows.cu index 3838d2b6d97b..3db98a232a01 100644 --- a/projects/rocthrust/examples/sum_rows.cu +++ b/projects/rocthrust/examples/sum_rows.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -70,8 +70,8 @@ int main(void) array.begin(), row_indices.begin(), row_sums.begin(), - thrust::equal_to(), - thrust::plus()); + _THRUST_STD::equal_to(), + _THRUST_STD::plus()); // print data for (int i = 0; i < R; i++) diff --git a/projects/rocthrust/examples/transform_input_output_iterator.cu b/projects/rocthrust/examples/transform_input_output_iterator.cu index 436ffe863de9..f1cca6ceb594 100644 --- a/projects/rocthrust/examples/transform_input_output_iterator.cu +++ b/projects/rocthrust/examples/transform_input_output_iterator.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -112,7 +112,7 @@ int main(void) C.begin(), ValueToScaledInteger{C_scale}, ScaledIntegerToValue{C_scale}); // Sum A and B as ScaledIntegers, storing the scaled result in C - thrust::transform(A_begin, A_end, B_begin, C_begin, thrust::plus{}); + thrust::transform(A_begin, A_end, B_begin, C_begin, _THRUST_STD::plus{}); thrust::host_vector A_h(A); thrust::host_vector B_h(B); diff --git a/projects/rocthrust/examples/transform_iterator.cu b/projects/rocthrust/examples/transform_iterator.cu index 5d17f566afff..0b122aa35e93 100644 --- a/projects/rocthrust/examples/transform_iterator.cu +++ b/projects/rocthrust/examples/transform_iterator.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -132,10 +132,10 @@ int main(void) //// // combine transform_iterator with another transform_iterator - using NegatedClampedCountingIterator = thrust::transform_iterator, ClampedCountingIterator>; + using NegatedClampedCountingIterator = thrust::transform_iterator<_THRUST_STD::negate, ClampedCountingIterator>; - NegatedClampedCountingIterator ncs_begin = thrust::make_transform_iterator(cs_begin, thrust::negate()); - NegatedClampedCountingIterator ncs_end = thrust::make_transform_iterator(cs_end, thrust::negate()); + NegatedClampedCountingIterator ncs_begin = thrust::make_transform_iterator(cs_begin, _THRUST_STD::negate()); + NegatedClampedCountingIterator ncs_end = thrust::make_transform_iterator(cs_end, _THRUST_STD::negate()); print_range("negated sequence ", ncs_begin, ncs_end); diff --git a/projects/rocthrust/examples/transform_output_iterator.cu b/projects/rocthrust/examples/transform_output_iterator.cu index 02321eb0aa7d..5967ba328c22 100644 --- a/projects/rocthrust/examples/transform_output_iterator.cu +++ b/projects/rocthrust/examples/transform_output_iterator.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -52,7 +52,7 @@ int main(void) // gather multiple elements and apply a function before writing result in memory thrust::gather(IDX.begin(), IDX.end(), - thrust::make_zip_iterator(thrust::make_tuple(U.begin(), V.begin())), + thrust::make_zip_iterator(U.begin(), V.begin()), thrust::make_transform_output_iterator(W.begin(), Functor())); std::cout << "result= [ "; diff --git a/projects/rocthrust/examples/word_count.cu b/projects/rocthrust/examples/word_count.cu index ec9bceb625f9..f74c27d14e56 100644 --- a/projects/rocthrust/examples/word_count.cu +++ b/projects/rocthrust/examples/word_count.cu @@ -1,4 +1,4 @@ -// Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2020-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 @@ -62,7 +62,7 @@ int word_count(const thrust::device_vector& input) input.end() - 1, // sequence of left characters input.begin() + 1, // sequence of right characters 0, // initialize sum to 0 - thrust::plus(), // sum values together + _THRUST_STD::plus(), // sum values together is_word_start()); // how to compare the left and right characters // if the first character is alphabetical, then it also begins a word 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/CMakeLists.txt b/projects/rocthrust/test/CMakeLists.txt index 8bf2f4ffee9e..cd2095ae383f 100644 --- a/projects/rocthrust/test/CMakeLists.txt +++ b/projects/rocthrust/test/CMakeLists.txt @@ -205,6 +205,7 @@ if(BUILD_TEST) add_rocthrust_test("is_partitioned") add_rocthrust_test("is_sorted") add_rocthrust_test("is_sorted_until") + add_rocthrust_test("libcxx_iterators") add_rocthrust_test("logical") add_rocthrust_test("max_element") add_rocthrust_test("memory") @@ -214,6 +215,8 @@ if(BUILD_TEST) add_rocthrust_test("metaprogramming") add_rocthrust_test("min_and_max") add_rocthrust_test("min_element") + add_rocthrust_test("minimum_system") + add_rocthrust_test("minimum_type") add_rocthrust_test("minmax_element") add_rocthrust_test("mismatch") add_rocthrust_test("mr_disjoint_pool") @@ -238,6 +241,7 @@ if(BUILD_TEST) add_rocthrust_test("raw_reference_cast") add_rocthrust_test("reduce") add_rocthrust_test("reduce_by_key") + add_rocthrust_test("reduce_into") add_rocthrust_test("reduce_large") add_rocthrust_test("remove") # Temporarily disable this test for gfx950 when CODE_COVERAGE is ON due to a compiler issue @@ -283,6 +287,7 @@ if(BUILD_TEST) add_rocthrust_test("stable_sort_by_key_large_keys_and_values") add_rocthrust_test("stable_sort_by_key_large_values") add_rocthrust_test("stable_sort_large") + add_rocthrust_test("strided_iterator") add_rocthrust_test("swap_ranges") add_rocthrust_test("tabulate") add_rocthrust_test("tabulate_output_iterator") @@ -379,6 +384,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 diff --git a/projects/rocthrust/test/bitwise_repro/bwr_utils.hpp b/projects/rocthrust/test/bitwise_repro/bwr_utils.hpp index e3860e98f8e8..9631471e95f0 100644 --- a/projects/rocthrust/test/bitwise_repro/bwr_utils.hpp +++ b/projects/rocthrust/test/bitwise_repro/bwr_utils.hpp @@ -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 @@ -281,8 +281,8 @@ std::string get_iterator_token(const std::string& vec_hash, const std::string& d /*! \brief Returns a "token" string that uniquely identifies a functor (callable object). * * \tparam T The data type that the functor operates on. - * \param functor_type String representing the functor type, without the datatype (eg. "thrust::plus" for functor - * thrust::plus). \return Strin gof the form "function(function_type)" + * \param functor_type String representing the functor type, without the datatype (eg. "_THRUST_STD::plus" for + * functor _THRUST_STD::plus). \return Strin gof the form "function(function_type)" */ template std::string get_functor_token(const std::string& functor_type) @@ -343,14 +343,14 @@ std::string build_output_token(const std::vector& list) * \code * // Suppose we want to record this call: * std::vector input = {...}; - * thrust::inclusive_scan(policy, d_input.begin(), d_input.end(), d_output.begin(), thrust::plus); + * thrust::inclusive_scan(policy, d_input.begin(), d_input.end(), d_output.begin(), _THRUST_STD::plus); * * TokenHelper helper; * token_helper.build_input_token( * "thrust::inclusive_scan", * d_input.begin(), * d_input.end(), - * {bwr_utils::get_functor_token("thrust::plus")} + * {bwr_utils::get_functor_token("_THRUST_STD::plus")} * ); * * token_helper.build_output_token(d_output.begin(), d_output.size()); diff --git a/projects/rocthrust/test/generate_resource_spec.cpp b/projects/rocthrust/test/generate_resource_spec.cpp index caad9e32a083..ac146274aa75 100644 --- a/projects/rocthrust/test/generate_resource_spec.cpp +++ b/projects/rocthrust/test/generate_resource_spec.cpp @@ -49,7 +49,7 @@ // ./generate_resource_spec // // Sample output: -// { +// { // "version": { // "major": 1, // "minor": 0 diff --git a/projects/rocthrust/test/hipstdpar/CMakeLists.txt b/projects/rocthrust/test/hipstdpar/CMakeLists.txt index 0bfee8a44e44..c2e9fe99f8d9 100644 --- a/projects/rocthrust/test/hipstdpar/CMakeLists.txt +++ b/projects/rocthrust/test/hipstdpar/CMakeLists.txt @@ -1,5 +1,5 @@ # ######################################################################## -# Copyright 2024-2025 Advanced Micro Devices, Inc. +# Copyright 2024-2026 Advanced Micro Devices, Inc. # ######################################################################## function(add_hipstdpar_test TEST TEST_TYPE INTERPOSE_ALLOC) diff --git a/projects/rocthrust/test/test_address_stability.cpp b/projects/rocthrust/test/test_address_stability.cpp index f7f4a619ae48..3400b87087a8 100644 --- a/projects/rocthrust/test/test_address_stability.cpp +++ b/projects/rocthrust/test/test_address_stability.cpp @@ -15,12 +15,14 @@ * limitations under the License. */ -#include - #include "test_param_fixtures.hpp" #include "test_utils.hpp" -#if !_THRUST_HAS_DEVICE_SYSTEM_STD +#if _THRUST_HAS_DEVICE_SYSTEM_STD +# include _THRUST_LIBCXX_INCLUDE(functional) +#else +# include + # include #endif @@ -38,8 +40,8 @@ TEST(AddressStabilityTests, TestAddressStabilityLibcuxx) { SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); - using ::internal::proclaim_copyable_arguments; - using ::internal::proclaims_copyable_arguments; + using _THRUST_LIBCXX::proclaim_copyable_arguments; + using _THRUST_LIBCXX::proclaims_copyable_arguments; // libcu++ function objects with known types static_assert(proclaims_copyable_arguments<_THRUST_STD::plus>::value, ""); @@ -59,21 +61,21 @@ TEST(AddressStabilityTests, TestAddressStabilityThrust) { SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); - using ::internal::proclaim_copyable_arguments; - using ::internal::proclaims_copyable_arguments; + using _THRUST_LIBCXX::proclaim_copyable_arguments; + using _THRUST_LIBCXX::proclaims_copyable_arguments; // thrust function objects with known types - static_assert(proclaims_copyable_arguments>::value, ""); - static_assert(!proclaims_copyable_arguments>::value, ""); + static_assert(proclaims_copyable_arguments<_THRUST_STD::plus>::value, ""); + static_assert(!proclaims_copyable_arguments<_THRUST_STD::plus<>>::value, ""); // thrust function objects with unknown types - static_assert(!proclaims_copyable_arguments>::value, ""); - static_assert(!proclaims_copyable_arguments>::value, ""); + static_assert(!proclaims_copyable_arguments<_THRUST_STD::plus>::value, ""); + static_assert(!proclaims_copyable_arguments<_THRUST_STD::plus<>>::value, ""); // thrust function objects with unknown types and opt-in - static_assert(proclaims_copyable_arguments{}))>::value, - ""); - static_assert(proclaims_copyable_arguments{}))>::value, ""); + static_assert( + proclaims_copyable_arguments{}))>::value, ""); + static_assert(proclaims_copyable_arguments{}))>::value, ""); } template @@ -89,8 +91,8 @@ TEST(AddressStabilityTests, TestAddressStabilityUserDefinedFunctionObject) { SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); - using ::internal::proclaim_copyable_arguments; - using ::internal::proclaims_copyable_arguments; + using _THRUST_LIBCXX::proclaim_copyable_arguments; + using _THRUST_LIBCXX::proclaims_copyable_arguments; // by-value overload static_assert(!proclaims_copyable_arguments>::value, ""); @@ -115,8 +117,8 @@ TEST(AddressStabilityTests, TestAddressStabilityLambda) { SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); - using ::internal::proclaim_copyable_arguments; - using ::internal::proclaims_copyable_arguments; + using _THRUST_LIBCXX::proclaim_copyable_arguments; + using _THRUST_LIBCXX::proclaims_copyable_arguments; { auto l = [](const int& i) { @@ -133,8 +135,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_adjacent_difference.cpp b/projects/rocthrust/test/test_adjacent_difference.cpp index 730df13460da..2f25a38d02a9 100644 --- a/projects/rocthrust/test/test_adjacent_difference.cpp +++ b/projects/rocthrust/test/test_adjacent_difference.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. @@ -45,7 +45,7 @@ TYPED_TEST(AdjacentDifferenceTests, TestAdjacentDifferenceSimple) Vector ref{1, 3, 2, 1}; ASSERT_EQ(output, ref); - result = thrust::adjacent_difference(input.begin(), input.end(), output.begin(), thrust::plus()); + result = thrust::adjacent_difference(input.begin(), input.end(), output.begin(), _THRUST_STD::plus()); ASSERT_EQ(result - output.begin(), 4); ref = {1, 5, 10, 13}; @@ -90,16 +90,16 @@ TYPED_TEST(AdjacentDifferenceVariableTests, TestAdjacentDifference) ASSERT_EQ(std::size_t(d_result - d_output.begin()), size); ASSERT_EQ(h_output, d_output); - h_result = thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin(), thrust::plus()); - d_result = thrust::adjacent_difference(d_input.begin(), d_input.end(), d_output.begin(), thrust::plus()); + h_result = thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin(), _THRUST_STD::plus()); + d_result = thrust::adjacent_difference(d_input.begin(), d_input.end(), d_output.begin(), _THRUST_STD::plus()); ASSERT_EQ(std::size_t(h_result - h_output.begin()), size); ASSERT_EQ(std::size_t(d_result - d_output.begin()), size); ASSERT_EQ(h_output, d_output); // in-place operation - h_result = thrust::adjacent_difference(h_input.begin(), h_input.end(), h_input.begin(), thrust::plus()); - d_result = thrust::adjacent_difference(d_input.begin(), d_input.end(), d_input.begin(), thrust::plus()); + h_result = thrust::adjacent_difference(h_input.begin(), h_input.end(), h_input.begin(), _THRUST_STD::plus()); + d_result = thrust::adjacent_difference(d_input.begin(), d_input.end(), d_input.begin(), _THRUST_STD::plus()); ASSERT_EQ(std::size_t(h_result - h_input.begin()), size); ASSERT_EQ(std::size_t(d_result - d_input.begin()), size); @@ -133,12 +133,12 @@ TYPED_TEST(AdjacentDifferenceVariableTests, TestAdjacentDifferenceInPlaceWithRel typename thrust::host_vector::iterator h_result; typename thrust::device_vector::iterator d_result; - h_result = thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin(), thrust::plus()); - d_result = thrust::adjacent_difference(d_input.begin(), d_input.end(), d_output.begin(), thrust::plus()); + h_result = thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin(), _THRUST_STD::plus()); + d_result = thrust::adjacent_difference(d_input.begin(), d_input.end(), d_output.begin(), _THRUST_STD::plus()); // in-place operation with different iterator types - h_result = thrust::adjacent_difference(h_input.cbegin(), h_input.cend(), h_input.begin(), thrust::plus()); - d_result = thrust::adjacent_difference(d_input.cbegin(), d_input.cend(), d_input.begin(), thrust::plus()); + h_result = thrust::adjacent_difference(h_input.cbegin(), h_input.cend(), h_input.begin(), _THRUST_STD::plus()); + d_result = thrust::adjacent_difference(d_input.cbegin(), d_input.cend(), d_input.begin(), _THRUST_STD::plus()); ASSERT_EQ(std::size_t(h_result - h_input.begin()), size); ASSERT_EQ(std::size_t(d_result - d_input.begin()), size); @@ -224,7 +224,7 @@ __global__ THRUST_HIP_LAUNCH_BOUNDS_DEFAULT void AdjacentDifferenceKernel(int co { thrust::device_ptr begin(array); thrust::device_ptr end(array + N); - // thrust::advance(begin,2); + // _THRUST_STD::advance(begin,2); thrust::adjacent_difference(thrust::hip::par, begin, end, begin); } } diff --git a/projects/rocthrust/test/test_advance.cpp b/projects/rocthrust/test/test_advance.cpp index 8f362e9bbf92..0b651f7bc17b 100644 --- a/projects/rocthrust/test/test_advance.cpp +++ b/projects/rocthrust/test/test_advance.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. @@ -39,15 +39,15 @@ TYPED_TEST(AdvanceVectorTests, TestAdvance) Iterator i = v.begin(); - thrust::advance(i, 1); + _THRUST_STD::advance(i, 1); ASSERT_EQ(*i, T(1)); - thrust::advance(i, 8); + _THRUST_STD::advance(i, 8); ASSERT_EQ(*i, T(9)); - thrust::advance(i, -4); + _THRUST_STD::advance(i, -4); ASSERT_EQ(*i, T(5)); @@ -56,15 +56,15 @@ TYPED_TEST(AdvanceVectorTests, TestAdvance) i = v.begin(); - thrust::advance(i, 7); + _THRUST_STD::advance(i, 7); ASSERT_EQ(*i, T(7)); - thrust::advance(i, 13); + _THRUST_STD::advance(i, 13); ASSERT_EQ(*i, T(20)); - thrust::advance(i, -10); + _THRUST_STD::advance(i, -10); ASSERT_EQ(*i, T(10)); } @@ -82,18 +82,18 @@ TYPED_TEST(AdvanceVectorTests, TestNext) Iterator const i0 = v.begin(); - Iterator const i1 = thrust::next(i0); + Iterator const i1 = _THRUST_STD::next(i0); ASSERT_EQ(*i0, T(0)); ASSERT_EQ(*i1, T(1)); - Iterator const i2 = thrust::next(i1, 8); + Iterator const i2 = _THRUST_STD::next(i1, 8); ASSERT_EQ(*i0, T(0)); ASSERT_EQ(*i1, T(1)); ASSERT_EQ(*i2, T(9)); - Iterator const i3 = thrust::next(i2, -4); + Iterator const i3 = _THRUST_STD::next(i2, -4); ASSERT_EQ(*i0, T(0)); ASSERT_EQ(*i1, T(1)); @@ -114,18 +114,18 @@ TYPED_TEST(AdvanceVectorTests, TestPrev) Iterator const i0 = v.end(); - Iterator const i1 = thrust::prev(i0); + Iterator const i1 = _THRUST_STD::prev(i0); ASSERT_EQ_QUIET(i0, v.end()); ASSERT_EQ(*i1, T(9)); - Iterator const i2 = thrust::prev(i1, 8); + Iterator const i2 = _THRUST_STD::prev(i1, 8); ASSERT_EQ_QUIET(i0, v.end()); ASSERT_EQ(*i1, T(9)); ASSERT_EQ(*i2, T(1)); - Iterator const i3 = thrust::prev(i2, -4); + Iterator const i3 = _THRUST_STD::prev(i2, -4); ASSERT_EQ_QUIET(i0, v.end()); ASSERT_EQ(*i1, T(9)); diff --git a/projects/rocthrust/test/test_allocator.cpp b/projects/rocthrust/test/test_allocator.cpp index 28dde4c6897b..7ae434b47b3f 100644 --- a/projects/rocthrust/test/test_allocator.cpp +++ b/projects/rocthrust/test/test_allocator.cpp @@ -156,7 +156,7 @@ struct my_allocator_with_custom_destroy return !(*this == other); } - using is_always_equal = thrust::detail::true_type; + using is_always_equal = _THRUST_STD::true_type; // use composition rather than inheritance // to avoid inheriting std::allocator's member diff --git a/projects/rocthrust/test/test_binary_search.cpp b/projects/rocthrust/test/test_binary_search.cpp index 71a243d41aa3..f931ed6b8f56 100644 --- a/projects/rocthrust/test/test_binary_search.cpp +++ b/projects/rocthrust/test/test_binary_search.cpp @@ -1011,21 +1011,21 @@ void TestBoundsWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(1); thrust::counting_iterator end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); _THRUST_STD::intmax_t distance_low_value = - thrust::distance(begin, thrust::lower_bound(thrust::device, begin, end, 17)); + _THRUST_STD::distance(begin, thrust::lower_bound(thrust::device, begin, end, 17)); _THRUST_STD::intmax_t distance_high_value = - thrust::distance(begin, thrust::lower_bound(thrust::device, begin, end, (1ll << magnitude) - 17)); + _THRUST_STD::distance(begin, thrust::lower_bound(thrust::device, begin, end, (1ll << magnitude) - 17)); ASSERT_EQ(distance_low_value, 16); ASSERT_EQ(distance_high_value, (1ll << magnitude) - 18); - distance_low_value = thrust::distance(begin, thrust::upper_bound(thrust::device, begin, end, 17)); + distance_low_value = _THRUST_STD::distance(begin, thrust::upper_bound(thrust::device, begin, end, 17)); distance_high_value = - thrust::distance(begin, thrust::upper_bound(thrust::device, begin, end, (1ll << magnitude) - 17)); + _THRUST_STD::distance(begin, thrust::upper_bound(thrust::device, begin, end, (1ll << magnitude) - 17)); ASSERT_EQ(distance_low_value, 17); ASSERT_EQ(distance_high_value, (1ll << magnitude) - 17); diff --git a/projects/rocthrust/test/test_binary_search_descending.cpp b/projects/rocthrust/test/test_binary_search_descending.cpp index c8db5f2c88c6..a76c4b67c515 100644 --- a/projects/rocthrust/test/test_binary_search_descending.cpp +++ b/projects/rocthrust/test/test_binary_search_descending.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. @@ -39,16 +39,16 @@ TYPED_TEST(BinarySearchDescendingTests, TestScalarLowerBoundDescendingSimple) Vector vec{8, 7, 5, 2, 0}; - ASSERT_EQ_QUIET(vec.begin() + 4, thrust::lower_bound(vec.begin(), vec.end(), T{0}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 4, thrust::lower_bound(vec.begin(), vec.end(), T{1}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), T{2}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), T{3}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), T{4}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 2, thrust::lower_bound(vec.begin(), vec.end(), T{5}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 2, thrust::lower_bound(vec.begin(), vec.end(), T{6}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 1, thrust::lower_bound(vec.begin(), vec.end(), T{7}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 0, thrust::lower_bound(vec.begin(), vec.end(), T{8}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 0, thrust::lower_bound(vec.begin(), vec.end(), T{9}, thrust::greater())); + ASSERT_EQ_QUIET(vec.begin() + 4, thrust::lower_bound(vec.begin(), vec.end(), T{0}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 4, thrust::lower_bound(vec.begin(), vec.end(), T{1}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), T{2}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), T{3}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), T{4}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 2, thrust::lower_bound(vec.begin(), vec.end(), T{5}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 2, thrust::lower_bound(vec.begin(), vec.end(), T{6}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 1, thrust::lower_bound(vec.begin(), vec.end(), T{7}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 0, thrust::lower_bound(vec.begin(), vec.end(), T{8}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 0, thrust::lower_bound(vec.begin(), vec.end(), T{9}, _THRUST_STD::greater())); } TYPED_TEST(BinarySearchDescendingTests, TestScalarUpperBoundDescendingSimple) @@ -60,16 +60,16 @@ TYPED_TEST(BinarySearchDescendingTests, TestScalarUpperBoundDescendingSimple) Vector vec{8, 7, 5, 2, 0}; - ASSERT_EQ_QUIET(vec.begin() + 5, thrust::upper_bound(vec.begin(), vec.end(), T{0}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 4, thrust::upper_bound(vec.begin(), vec.end(), T{1}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 4, thrust::upper_bound(vec.begin(), vec.end(), T{2}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), T{3}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), T{4}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), T{5}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 2, thrust::upper_bound(vec.begin(), vec.end(), T{6}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 2, thrust::upper_bound(vec.begin(), vec.end(), T{7}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 1, thrust::upper_bound(vec.begin(), vec.end(), T{8}, thrust::greater())); - ASSERT_EQ_QUIET(vec.begin() + 0, thrust::upper_bound(vec.begin(), vec.end(), T{9}, thrust::greater())); + ASSERT_EQ_QUIET(vec.begin() + 5, thrust::upper_bound(vec.begin(), vec.end(), T{0}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 4, thrust::upper_bound(vec.begin(), vec.end(), T{1}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 4, thrust::upper_bound(vec.begin(), vec.end(), T{2}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), T{3}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), T{4}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), T{5}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 2, thrust::upper_bound(vec.begin(), vec.end(), T{6}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 2, thrust::upper_bound(vec.begin(), vec.end(), T{7}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 1, thrust::upper_bound(vec.begin(), vec.end(), T{8}, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(vec.begin() + 0, thrust::upper_bound(vec.begin(), vec.end(), T{9}, _THRUST_STD::greater())); } TYPED_TEST(BinarySearchDescendingTests, TestScalarBinarySearchDescendingSimple) @@ -81,16 +81,16 @@ TYPED_TEST(BinarySearchDescendingTests, TestScalarBinarySearchDescendingSimple) Vector vec{8, 7, 5, 2, 0}; - ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{0}, thrust::greater())); - ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{1}, thrust::greater())); - ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{2}, thrust::greater())); - ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{3}, thrust::greater())); - ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{4}, thrust::greater())); - ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{5}, thrust::greater())); - ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{6}, thrust::greater())); - ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{7}, thrust::greater())); - ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{8}, thrust::greater())); - ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{9}, thrust::greater())); + ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{0}, _THRUST_STD::greater())); + ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{1}, _THRUST_STD::greater())); + ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{2}, _THRUST_STD::greater())); + ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{3}, _THRUST_STD::greater())); + ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{4}, _THRUST_STD::greater())); + ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{5}, _THRUST_STD::greater())); + ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{6}, _THRUST_STD::greater())); + ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{7}, _THRUST_STD::greater())); + ASSERT_EQ(true, thrust::binary_search(vec.begin(), vec.end(), T{8}, _THRUST_STD::greater())); + ASSERT_EQ(false, thrust::binary_search(vec.begin(), vec.end(), T{9}, _THRUST_STD::greater())); } TYPED_TEST(BinarySearchDescendingTests, TestScalarEqualRangeDescendingSimple) @@ -102,25 +102,25 @@ TYPED_TEST(BinarySearchDescendingTests, TestScalarEqualRangeDescendingSimple) Vector vec{8, 7, 5, 2, 0}; - ASSERT_EQ_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), T{0}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), T{1}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{2}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{3}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{4}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), T{5}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), T{6}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 1, thrust::equal_range(vec.begin(), vec.end(), T{7}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), T{8}, thrust::greater()).first); - ASSERT_EQ_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), T{9}, thrust::greater()).first); - - ASSERT_EQ_QUIET(vec.begin() + 5, thrust::equal_range(vec.begin(), vec.end(), T{0}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), T{1}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), T{2}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{3}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{4}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{5}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), T{6}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), T{7}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 1, thrust::equal_range(vec.begin(), vec.end(), T{8}, thrust::greater()).second); - ASSERT_EQ_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), T{9}, thrust::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), T{0}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), T{1}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{2}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{3}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{4}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), T{5}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), T{6}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 1, thrust::equal_range(vec.begin(), vec.end(), T{7}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), T{8}, _THRUST_STD::greater()).first); + ASSERT_EQ_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), T{9}, _THRUST_STD::greater()).first); + + ASSERT_EQ_QUIET(vec.begin() + 5, thrust::equal_range(vec.begin(), vec.end(), T{0}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), T{1}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), T{2}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{3}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{4}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), T{5}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), T{6}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), T{7}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 1, thrust::equal_range(vec.begin(), vec.end(), T{8}, _THRUST_STD::greater()).second); + ASSERT_EQ_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), T{9}, _THRUST_STD::greater()).second); } diff --git a/projects/rocthrust/test/test_binary_search_vector_descending.cpp b/projects/rocthrust/test/test_binary_search_vector_descending.cpp index cdc9a8b435b5..19461b5ae5d5 100644 --- a/projects/rocthrust/test/test_binary_search_vector_descending.cpp +++ b/projects/rocthrust/test/test_binary_search_vector_descending.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. @@ -60,7 +60,7 @@ TYPED_TEST(BinarySearchVectorDescendingTests, TestVectorLowerBoundDescendingSimp // test with integral output type IntVector integral_output(10); typename IntVector::iterator output_end = thrust::lower_bound( - vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), thrust::greater()); + vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), _THRUST_STD::greater()); ASSERT_EQ_QUIET(integral_output.end(), output_end); @@ -86,7 +86,7 @@ TYPED_TEST(BinarySearchVectorDescendingTests, TestVectorUpperBoundDescendingSimp // test with integral output type IntVector integral_output(10); typename IntVector::iterator output_end = thrust::upper_bound( - vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), thrust::greater()); + vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), _THRUST_STD::greater()); ASSERT_EQ_QUIET(output_end, integral_output.end()); @@ -113,7 +113,7 @@ TYPED_TEST(BinarySearchVectorDescendingTests, TestVectorBinarySearchDescendingSi // test with boolean output type BoolVector bool_output(10); typename BoolVector::iterator bool_output_end = thrust::binary_search( - vec.begin(), vec.end(), input.begin(), input.end(), bool_output.begin(), thrust::greater()); + vec.begin(), vec.end(), input.begin(), input.end(), bool_output.begin(), _THRUST_STD::greater()); ASSERT_EQ_QUIET(bool_output_end, bool_output.end()); @@ -123,7 +123,7 @@ TYPED_TEST(BinarySearchVectorDescendingTests, TestVectorBinarySearchDescendingSi // test with integral output type IntVector integral_output(10, 2); typename IntVector::iterator int_output_end = thrust::binary_search( - vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), thrust::greater()); + vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), _THRUST_STD::greater()); ASSERT_EQ_QUIET(int_output_end, integral_output.end()); @@ -148,7 +148,7 @@ TYPED_TEST(BinarySearchVectorDescendingIntegerTests, TestVectorLowerBoundDescend thrust::host_vector h_vec = get_random_data(size, get_default_limits::min(), get_default_limits::max(), seed); - thrust::sort(h_vec.begin(), h_vec.end(), thrust::greater()); + thrust::sort(h_vec.begin(), h_vec.end(), _THRUST_STD::greater()); thrust::device_vector d_vec = h_vec; thrust::host_vector h_input = get_random_data( @@ -160,9 +160,9 @@ TYPED_TEST(BinarySearchVectorDescendingIntegerTests, TestVectorLowerBoundDescend thrust::device_vector d_output(2 * size); thrust::lower_bound( - h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), thrust::greater()); + h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), _THRUST_STD::greater()); thrust::lower_bound( - d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), thrust::greater()); + d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), _THRUST_STD::greater()); ASSERT_EQ(h_output, d_output); } @@ -185,7 +185,7 @@ TYPED_TEST(BinarySearchVectorDescendingIntegerTests, TestVectorUpperBoundDescend thrust::host_vector h_vec = get_random_data(size, get_default_limits::min(), get_default_limits::max(), seed); - thrust::sort(h_vec.begin(), h_vec.end(), thrust::greater()); + thrust::sort(h_vec.begin(), h_vec.end(), _THRUST_STD::greater()); thrust::device_vector d_vec = h_vec; thrust::host_vector h_input = get_random_data( @@ -197,9 +197,9 @@ TYPED_TEST(BinarySearchVectorDescendingIntegerTests, TestVectorUpperBoundDescend thrust::device_vector d_output(2 * size); thrust::upper_bound( - h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), thrust::greater()); + h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), _THRUST_STD::greater()); thrust::upper_bound( - d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), thrust::greater()); + d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), _THRUST_STD::greater()); ASSERT_EQ(h_output, d_output); } @@ -222,7 +222,7 @@ TYPED_TEST(BinarySearchVectorDescendingIntegerTests, TestVectorBinarySearchDesce thrust::host_vector h_vec = get_random_data(size, get_default_limits::min(), get_default_limits::max(), seed); - thrust::sort(h_vec.begin(), h_vec.end(), thrust::greater()); + thrust::sort(h_vec.begin(), h_vec.end(), _THRUST_STD::greater()); thrust::device_vector d_vec = h_vec; thrust::host_vector h_input = get_random_data( @@ -234,9 +234,9 @@ TYPED_TEST(BinarySearchVectorDescendingIntegerTests, TestVectorBinarySearchDesce thrust::device_vector d_output(2 * size); thrust::binary_search( - h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), thrust::greater()); + h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), _THRUST_STD::greater()); thrust::binary_search( - d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), thrust::greater()); + d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), _THRUST_STD::greater()); ASSERT_EQ(h_output, d_output); } 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_complex_various.cpp b/projects/rocthrust/test/test_complex_various.cpp index fe79b1aaa6ab..8862520f365d 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. @@ -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) { @@ -153,15 +155,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) diff --git a/projects/rocthrust/test/test_constant_iterator.cpp b/projects/rocthrust/test/test_constant_iterator.cpp index 4ef7e1f477bd..5cb5e7e9419e 100644 --- a/projects/rocthrust/test/test_constant_iterator.cpp +++ b/projects/rocthrust/test/test_constant_iterator.cpp @@ -59,7 +59,7 @@ TEST(ConstantIteratorTests, TestConstantIteratorTraits) static_assert(_THRUST_STD::is_same_v, thrust::random_access_traversal_tag>); - static_assert(::internal::is_cpp17_random_access_iterator::value); + static_assert(_THRUST_STD::__has_random_access_traversal); #if _THRUST_HAS_DEVICE_SYSTEM_STD || THRUST_STD_VER >= 2020 static_assert(!_THRUST_STD::output_iterator); @@ -126,7 +126,7 @@ TEST(ConstantIteratorTests, TestConstantIteratorIncrementBig) thrust::constant_iterator begin(1); thrust::constant_iterator end = begin + n; - ASSERT_EQ(thrust::distance(begin, end), n); + ASSERT_EQ(_THRUST_STD::distance(begin, end), n); } TEST(ConstantIteratorTests, TestConstantIteratorComparison) @@ -174,6 +174,11 @@ TEST(ConstantIteratorTests, TestMakeConstantIterator) ASSERT_EQ(13, *iter1); ASSERT_EQ(7, iter1 - iter0); + + // ensure CTAD words + constant_iterator deduced_iter{42}; + static_assert(_THRUST_STD::is_same_v>); + ASSERT_EQ(42, *deduced_iter); } TYPED_TEST(ConstantIteratorTests, TestConstantIteratorCopy) @@ -212,12 +217,12 @@ TYPED_TEST(ConstantIteratorTests, TestConstantIteratorTransform) ConstIter last1 = first1 + result.size(); ConstIter first2 = make_constant_iterator(3); - thrust::transform(first1, last1, result.begin(), thrust::negate()); + thrust::transform(first1, last1, result.begin(), _THRUST_STD::negate()); Vector ref(4, -7); ASSERT_EQ(ref, result); - thrust::transform(first1, last1, first2, result.begin(), thrust::plus()); + thrust::transform(first1, last1, first2, result.begin(), _THRUST_STD::plus()); ref = Vector(4, 10); ASSERT_EQ(ref, result); diff --git a/projects/rocthrust/test/test_copy.cpp b/projects/rocthrust/test/test_copy.cpp index 40d492d8b5de..b2c7e18f0fb2 100644 --- a/projects/rocthrust/test/test_copy.cpp +++ b/projects/rocthrust/test/test_copy.cpp @@ -41,6 +41,11 @@ # include #endif +#if _CCCL_COMPILER(GCC, >=, 11) +# define THRUST_DISABLE_BROKEN_GCC_VECTORIZER __attribute__((optimize("no-tree-vectorize"))) +#else +# define THRUST_DISABLE_BROKEN_GCC_VECTORIZER +#endif using IntegralVariableParams = ::testing::Types, Params, @@ -197,15 +202,15 @@ TEST(CopyTests, TestCopyToDiscardIteratorZipped) // copy from host_vector ZipIterator1 h_result = thrust::copy( - thrust::make_zip_iterator(thrust::make_tuple(h_input.begin(), h_input.begin())), - thrust::make_zip_iterator(thrust::make_tuple(h_input.end(), h_input.end())), - thrust::make_zip_iterator(thrust::make_tuple(thrust::make_discard_iterator(), h_output.begin()))); + thrust::make_zip_iterator(h_input.begin(), h_input.begin()), + thrust::make_zip_iterator(h_input.end(), h_input.end()), + thrust::make_zip_iterator(thrust::make_discard_iterator(), h_output.begin())); // copy from device_vector ZipIterator2 d_result = thrust::copy( - thrust::make_zip_iterator(thrust::make_tuple(d_input.begin(), d_input.begin())), - thrust::make_zip_iterator(thrust::make_tuple(d_input.end(), d_input.end())), - thrust::make_zip_iterator(thrust::make_tuple(thrust::make_discard_iterator(), d_output.begin()))); + thrust::make_zip_iterator(d_input.begin(), d_input.begin()), + thrust::make_zip_iterator(d_input.end(), d_input.end()), + thrust::make_zip_iterator(thrust::make_discard_iterator(), d_output.begin())); ASSERT_EQ(h_output, h_input); ASSERT_EQ(d_output, d_input); @@ -685,7 +690,7 @@ struct always_true THRUST_HOST_DEVICE bool operator()(const object_with_non_trivial_ctor&) { return true; - }; + } }; } // namespace @@ -753,16 +758,10 @@ TYPED_TEST(CopyTests, TestCopyCountingIterator) ASSERT_EQ(vec[3], 4); } -TYPED_TEST(CopyTests, TestCopyZipIterator) +template +THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestCopyZipIteratorImpl() { - using Vector = typename TestFixture::input_type; - using T = typename Vector::value_type; - - SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); - - // initializer list doesn't work with GCC when - // Vector = thrust::host_vector - // Vector v1{1, 2, 3}; + using T = typename Vector::value_type; Vector v1(3); v1[0] = 1; @@ -775,14 +774,23 @@ TYPED_TEST(CopyTests, TestCopyZipIterator) Vector v3(3, T(0)); Vector v4(3, T(0)); - thrust::copy(thrust::make_zip_iterator(thrust::make_tuple(v1.begin(), v2.begin())), - thrust::make_zip_iterator(thrust::make_tuple(v1.end(), v2.end())), - thrust::make_zip_iterator(thrust::make_tuple(v3.begin(), v4.begin()))); + thrust::copy(thrust::make_zip_iterator(v1.begin(), v2.begin()), + thrust::make_zip_iterator(v1.end(), v2.end()), + thrust::make_zip_iterator(v3.begin(), v4.begin())); ASSERT_EQ(v1, v3); ASSERT_EQ(v2, v4); } +TYPED_TEST(CopyTests, TestCopyZipIterator) +{ + using Vector = typename TestFixture::input_type; + + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); + + TestCopyZipIteratorImpl(); +} + TYPED_TEST(CopyTests, TestCopyConstantIteratorToZipIterator) { using Vector = typename TestFixture::input_type; @@ -795,7 +803,7 @@ TYPED_TEST(CopyTests, TestCopyConstantIteratorToZipIterator) thrust::copy(thrust::make_constant_iterator(thrust::tuple(4, 7)), thrust::make_constant_iterator(thrust::tuple(4, 7)) + v1.size(), - thrust::make_zip_iterator(thrust::make_tuple(v1.begin(), v2.begin()))); + thrust::make_zip_iterator(v1.begin(), v2.begin())); Vector ref1{4, 4, 4}; Vector ref2{7, 7, 7}; @@ -967,8 +975,7 @@ namespace detail // We need this type to pass as a non-const ref for unary_transform_functor // to compile: template <> -struct is_non_const_reference : thrust::true_type -{}; +inline constexpr bool is_non_const_reference_v = true; } // end namespace detail THRUST_NAMESPACE_END @@ -1001,7 +1008,7 @@ void TestCopyWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(0); thrust::counting_iterator end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); thrust::device_ptr has_executed = thrust::device_malloc(1); *has_executed = false; diff --git a/projects/rocthrust/test/test_copy_n.cpp b/projects/rocthrust/test/test_copy_n.cpp index e6244a18277a..b669935385af 100644 --- a/projects/rocthrust/test/test_copy_n.cpp +++ b/projects/rocthrust/test/test_copy_n.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. @@ -221,9 +221,8 @@ TYPED_TEST(CopyNTests, TestCopyNZipIterator) Vector v3(4, T(0)); Vector v4(4, T(0)); - thrust::copy_n(thrust::make_zip_iterator(thrust::make_tuple(v1.begin(), v2.begin())), - 4, - thrust::make_zip_iterator(thrust::make_tuple(v3.begin(), v4.begin()))); + thrust::copy_n( + thrust::make_zip_iterator(v1.begin(), v2.begin()), 4, thrust::make_zip_iterator(v3.begin(), v4.begin())); ASSERT_EQ(v1, v3); ASSERT_EQ(v2, v4); @@ -241,7 +240,7 @@ TYPED_TEST(CopyNTests, TestCopyNConstantIteratorToZipIterator) thrust::copy_n(thrust::make_constant_iterator(thrust::tuple(4, 7)), v1.size(), - thrust::make_zip_iterator(thrust::make_tuple(v1.begin(), v2.begin()))); + thrust::make_zip_iterator(v1.begin(), v2.begin())); Vector ref1(4, 4); Vector ref2(4, 7); diff --git a/projects/rocthrust/test/test_count.cpp b/projects/rocthrust/test/test_count.cpp index 8c6dc3ea01ad..365049994748 100644 --- a/projects/rocthrust/test/test_count.cpp +++ b/projects/rocthrust/test/test_count.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. @@ -180,7 +180,7 @@ void TestCountWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(1); thrust::counting_iterator end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); long long result = thrust::count(thrust::device, begin, end, (1ll << magnitude) - 17); diff --git a/projects/rocthrust/test/test_counting_iterator.cpp b/projects/rocthrust/test/test_counting_iterator.cpp index bc09f769d456..2a12cfa7861f 100644 --- a/projects/rocthrust/test/test_counting_iterator.cpp +++ b/projects/rocthrust/test/test_counting_iterator.cpp @@ -21,15 +21,18 @@ #include #include +#include _THRUST_STD_INCLUDE(__algorithm_) +#include _THRUST_STD_INCLUDE(iterator) +#include _THRUST_STD_INCLUDE(type_traits) + +#include #include +#include #include "test_param_fixtures.hpp" #include "test_real_assertions.hpp" #include "test_utils.hpp" -#include _THRUST_STD_INCLUDE(iterator) -#include _THRUST_STD_INCLUDE(type_traits) - TESTS_DEFINE(CountingIteratorTests, NumericalTestsParams); template @@ -83,7 +86,7 @@ TEST(CountingIteratorTests, TestCountingIteratorTraits) static_assert(_THRUST_STD::is_same_v, thrust::random_access_traversal_tag>); - static_assert(::internal::is_cpp17_random_access_iterator::value); + static_assert(_THRUST_STD::__has_random_access_traversal); #if _THRUST_HAS_DEVICE_SYSTEM_STD || THRUST_STD_VER >= 2020 static_assert(!_THRUST_STD::output_iterator); @@ -250,19 +253,19 @@ TEST(CountingIteratorTests, TestCountingIteratorDistance) thrust::counting_iterator iter1(0); thrust::counting_iterator iter2(5); - ASSERT_EQ(thrust::distance(iter1, iter2), 5); + ASSERT_EQ(_THRUST_STD::distance(iter1, iter2), 5); iter1++; - ASSERT_EQ(thrust::distance(iter1, iter2), 4); + ASSERT_EQ(_THRUST_STD::distance(iter1, iter2), 4); iter2 += 100; - ASSERT_EQ(thrust::distance(iter1, iter2), 104); + ASSERT_EQ(_THRUST_STD::distance(iter1, iter2), 104); iter2 += 1000; - ASSERT_EQ(thrust::distance(iter1, iter2), 1104); + ASSERT_EQ(_THRUST_STD::distance(iter1, iter2), 1104); } TEST(CountingIteratorTests, TestCountingIteratorUnsignedType) @@ -330,4 +333,62 @@ TEST(CountingIteratorTests, TestCountingIteratorDifference) ASSERT_EQ(diff, last - first); } +TEST(CountingIteratorTests, TestCountingIteratorDynamicStride) +{ + auto iter = thrust::make_counting_iterator(0, 2); + static_assert(sizeof(iter) == 2 * sizeof(int)); + + ASSERT_EQ(*iter, 0); + iter++; + ASSERT_EQ(*iter, 2); + iter++; + iter++; + ASSERT_EQ(*iter, 6); + iter += 5; + ASSERT_EQ(*iter, 16); + iter -= 10; + ASSERT_EQ(*iter, -4); +} + +TEST(CountingIteratorTests, TestCountingIteratorStaticStride) +{ + auto iter = thrust::make_counting_iterator<2>(0); + static_assert(sizeof(decltype(iter)) == sizeof(int)); + + ASSERT_EQ(*iter, 0); + iter++; + ASSERT_EQ(*iter, 2); + iter++; + iter++; + ASSERT_EQ(*iter, 6); + iter += 5; + ASSERT_EQ(*iter, 16); + iter -= 10; + ASSERT_EQ(*iter, -4); +} + +TEST(CountingIteratorTests, TestCountingIteratorPointer) +{ + int arr[11]; + std::iota(arr, arr + 11, 0); + + auto iter = thrust::make_counting_iterator(&arr[2]); + + ASSERT_EQ(*iter, &arr[2]); + ASSERT_EQ(**iter, 2); + iter++; + ASSERT_EQ(*iter, &arr[3]); + ASSERT_EQ(**iter, 3); + iter++; + iter++; + ASSERT_EQ(*iter, &arr[5]); + ASSERT_EQ(**iter, 5); + iter += 5; + ASSERT_EQ(*iter, &arr[10]); + ASSERT_EQ(**iter, 10); + iter -= 10; + ASSERT_EQ(*iter, &arr[0]); + ASSERT_EQ(**iter, 0); +} + THRUST_DIAG_POP diff --git a/projects/rocthrust/test/test_dereference.cpp b/projects/rocthrust/test/test_dereference.cpp index 566de2dc67d7..26f7af626e6e 100644 --- a/projects/rocthrust/test/test_dereference.cpp +++ b/projects/rocthrust/test/test_dereference.cpp @@ -15,7 +15,6 @@ * limitations under the License. */ -#include #include #include #include @@ -97,8 +96,8 @@ TEST(DereferenceTests, TestDeviceDereferenceTransformIterator) get_random_data(100, get_default_limits::min(), get_default_limits::max(), seed); thrust::device_vector output(input.size(), 0); - simple_copy(thrust::make_transform_iterator(input.begin(), ::internal::identity{}), - thrust::make_transform_iterator(input.end(), ::internal::identity{}), + simple_copy(thrust::make_transform_iterator(input.begin(), _THRUST_STD::identity{}), + thrust::make_transform_iterator(input.end(), _THRUST_STD::identity{}), output.begin()); ASSERT_EQ(input, output); @@ -112,8 +111,8 @@ TEST(DereferenceTests, TestDeviceDereferenceTransformIteratorInputConversion) thrust::device_vector input = unittest::random_integers(100); thrust::device_vector output(input.size(), 0); - simple_copy(thrust::make_transform_iterator(input.begin(), ::internal::identity{}), - thrust::make_transform_iterator(input.end(), ::internal::identity{}), + simple_copy(thrust::make_transform_iterator(input.begin(), _THRUST_STD::identity{}), + thrust::make_transform_iterator(input.end(), _THRUST_STD::identity{}), output.begin()); ASSERT_EQ(input == output, true); @@ -126,8 +125,8 @@ TEST(DereferenceTests, TestDeviceDereferenceTransformIteratorOutputConversion) thrust::device_vector input = unittest::random_integers(100); thrust::device_vector output(input.size(), 0); - simple_copy(thrust::make_transform_iterator(input.begin(), ::internal::identity{}), - thrust::make_transform_iterator(input.end(), ::internal::identity{}), + simple_copy(thrust::make_transform_iterator(input.begin(), _THRUST_STD::identity{}), + thrust::make_transform_iterator(input.end(), _THRUST_STD::identity{}), output.begin()); ASSERT_EQ(input == output, true); @@ -157,8 +156,8 @@ TEST(DereferenceTests, TestDeviceDereferenceTransformedCountingIterator) thrust::device_vector output(5); - simple_copy(thrust::make_transform_iterator(first, thrust::negate()), - thrust::make_transform_iterator(last, thrust::negate()), + simple_copy(thrust::make_transform_iterator(first, _THRUST_STD::negate()), + thrust::make_transform_iterator(last, _THRUST_STD::negate()), output.begin()); thrust::device_vector ref{-1, -2, -3, -4, -5}; diff --git a/projects/rocthrust/test/test_device_ptr.cpp b/projects/rocthrust/test/test_device_ptr.cpp index cf4cbfa549c3..b3a74a4b2adf 100644 --- a/projects/rocthrust/test/test_device_ptr.cpp +++ b/projects/rocthrust/test/test_device_ptr.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. @@ -18,11 +18,19 @@ #include #include +#include _THRUST_STD_INCLUDE(iterator) + #include +#include #include "test_param_fixtures.hpp" #include "test_utils.hpp" +#ifdef __cpp_lib_concepts +static_assert(std::indirectly_writable, uint8_t>); +#endif // __cpp_lib_concepts +static_assert(_THRUST_STD::indirectly_writable, uint8_t>); + #define HIP_CHECK_HMM(condition) \ { \ hipError_t error = condition; \ diff --git a/projects/rocthrust/test/test_discard_iterator.cpp b/projects/rocthrust/test/test_discard_iterator.cpp index 3f493c90b5ae..4b88b439f41e 100644 --- a/projects/rocthrust/test/test_discard_iterator.cpp +++ b/projects/rocthrust/test/test_discard_iterator.cpp @@ -46,7 +46,7 @@ TEST(DiscardIteratorTests, TestDiscardIteratorTraits) static_assert(_THRUST_STD::is_same_v, thrust::random_access_traversal_tag>); - static_assert(::internal::is_cpp17_random_access_iterator::value); + static_assert(_THRUST_STD::__has_random_access_traversal); #if _THRUST_HAS_DEVICE_SYSTEM_STD || THRUST_STD_VER >= 2020 static_assert(_THRUST_STD::output_iterator); @@ -152,7 +152,7 @@ TEST(DiscardIteratorTests, TestZippedDiscardIterator) using IteratorTuple2 = tuple>; using ZipIterator2 = zip_iterator; - ZipIterator2 z_iter_first = thrust::make_zip_iterator(thrust::make_tuple((int*) 0, thrust::make_discard_iterator())); + ZipIterator2 z_iter_first = thrust::make_zip_iterator((int*) 0, thrust::make_discard_iterator()); ZipIterator2 z_iter_last = z_iter_first + 10; for (; z_iter_first != z_iter_last; ++z_iter_first) diff --git a/projects/rocthrust/test/test_distance.cpp b/projects/rocthrust/test/test_distance.cpp index 4079a983ac95..aa7fc6580a4f 100644 --- a/projects/rocthrust/test/test_distance.cpp +++ b/projects/rocthrust/test/test_distance.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. @@ -35,17 +35,17 @@ TYPED_TEST(DistanceTests, TestDistance) Iterator i = v.begin(); - ASSERT_EQ(thrust::distance(i, v.end()), 100); + ASSERT_EQ(_THRUST_STD::distance(i, v.end()), 100); i++; - ASSERT_EQ(thrust::distance(i, v.end()), 99); + ASSERT_EQ(_THRUST_STD::distance(i, v.end()), 99); i += 49; - ASSERT_EQ(thrust::distance(i, v.end()), 50); + ASSERT_EQ(_THRUST_STD::distance(i, v.end()), 50); - ASSERT_EQ(thrust::distance(i, i), 0); + ASSERT_EQ(_THRUST_STD::distance(i, i), 0); } TYPED_TEST(DistanceTests, TestDistanceLarge) @@ -59,19 +59,19 @@ TYPED_TEST(DistanceTests, TestDistanceLarge) Iterator i = v.begin(); - ASSERT_EQ(thrust::distance(i, v.end()), 1000); + ASSERT_EQ(_THRUST_STD::distance(i, v.end()), 1000); i++; - ASSERT_EQ(thrust::distance(i, v.end()), 999); + ASSERT_EQ(_THRUST_STD::distance(i, v.end()), 999); i += 49; - ASSERT_EQ(thrust::distance(i, v.end()), 950); + ASSERT_EQ(_THRUST_STD::distance(i, v.end()), 950); i += 950; - ASSERT_EQ(thrust::distance(i, v.end()), 0); + ASSERT_EQ(_THRUST_STD::distance(i, v.end()), 0); - ASSERT_EQ(thrust::distance(i, i), 0); + ASSERT_EQ(_THRUST_STD::distance(i, i), 0); } diff --git a/projects/rocthrust/test/test_equal.cpp b/projects/rocthrust/test/test_equal.cpp index 0ca57c2d908a..46935fccbcd7 100644 --- a/projects/rocthrust/test/test_equal.cpp +++ b/projects/rocthrust/test/test_equal.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. @@ -46,13 +46,14 @@ TYPED_TEST(EqualTests, TestEqualSimple) ASSERT_EQ(thrust::equal(v1.begin(), v1.begin() + 3, v2.begin()), true); ASSERT_EQ(thrust::equal(v1.begin(), v1.begin() + 4, v2.begin()), false); - ASSERT_EQ(thrust::equal(v1.begin(), v1.end(), v2.begin(), thrust::less_equal()), true); - ASSERT_EQ(thrust::equal(v1.begin(), v1.end(), v2.begin(), thrust::greater()), false); + ASSERT_EQ(thrust::equal(v1.begin(), v1.end(), v2.begin(), _THRUST_STD::less_equal()), true); + ASSERT_EQ(thrust::equal(v1.begin(), v1.end(), v2.begin(), _THRUST_STD::greater()), false); } -TYPED_TEST(EqualsPrimitiveTests, TestEqual) +TYPED_TEST(EqualTests, TestEqual) { - using T = typename TestFixture::input_type; + using Vector = typename TestFixture::input_type; + using T = typename Vector::value_type; SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); @@ -91,10 +92,12 @@ TYPED_TEST(EqualsPrimitiveTests, TestEqual) ASSERT_EQ(thrust::equal(d_data1.begin(), d_data1.end(), d_data2.begin()), false); // different predicates - ASSERT_EQ(thrust::equal(h_data1.begin(), h_data1.begin() + 1, h_data2.begin(), thrust::less()), true); - ASSERT_EQ(thrust::equal(d_data1.begin(), d_data1.begin() + 1, d_data2.begin(), thrust::less()), true); - ASSERT_EQ(thrust::equal(h_data1.begin(), h_data1.begin() + 1, h_data2.begin(), thrust::greater()), false); - ASSERT_EQ(thrust::equal(d_data1.begin(), d_data1.begin() + 1, d_data2.begin(), thrust::greater()), false); + ASSERT_EQ(thrust::equal(h_data1.begin(), h_data1.begin() + 1, h_data2.begin(), _THRUST_STD::less()), true); + ASSERT_EQ(thrust::equal(d_data1.begin(), d_data1.begin() + 1, d_data2.begin(), _THRUST_STD::less()), true); + ASSERT_EQ(thrust::equal(h_data1.begin(), h_data1.begin() + 1, h_data2.begin(), _THRUST_STD::greater()), + false); + ASSERT_EQ(thrust::equal(d_data1.begin(), d_data1.begin() + 1, d_data2.begin(), _THRUST_STD::greater()), + false); } } } @@ -158,7 +161,7 @@ void TestEqualWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(1); thrust::counting_iterator end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); thrust::device_ptr has_executed = thrust::device_malloc(1); *has_executed = false; diff --git a/projects/rocthrust/test/test_fill.cpp b/projects/rocthrust/test/test_fill.cpp index 99b570df8048..15de63c9c7de 100644 --- a/projects/rocthrust/test/test_fill.cpp +++ b/projects/rocthrust/test/test_fill.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. @@ -275,8 +275,8 @@ TYPED_TEST(FillTests, TestFillZipIterator) Vector v2(3, T(0)); Vector v3(3, T(0)); - thrust::fill(thrust::make_zip_iterator(thrust::make_tuple(v1.begin(), v2.begin(), v3.begin())), - thrust::make_zip_iterator(thrust::make_tuple(v1.end(), v2.end(), v3.end())), + thrust::fill(thrust::make_zip_iterator(v1.begin(), v2.begin(), v3.begin()), + thrust::make_zip_iterator(v1.end(), v2.end(), v3.end()), thrust::tuple(4, 7, 13)); Vector ref1{4, 4, 4}; diff --git a/projects/rocthrust/test/test_find.cpp b/projects/rocthrust/test/test_find.cpp index 0426cd4c8858..32bc47fa5722 100644 --- a/projects/rocthrust/test/test_find.cpp +++ b/projects/rocthrust/test/test_find.cpp @@ -15,7 +15,6 @@ * limitations under the License. */ -#include #include #include #include @@ -158,7 +157,7 @@ TEST(FindTests, TestFindIfDispatchExplicit) thrust::device_vector vec(1); my_system sys(0); - thrust::find_if(sys, vec.begin(), vec.end(), ::internal::identity{}); + thrust::find_if(sys, vec.begin(), vec.end(), _THRUST_STD::identity{}); ASSERT_EQ(true, sys.is_valid()); } @@ -176,7 +175,7 @@ TEST(FindTests, TestFindIfDispatchImplicit) thrust::device_vector vec(1); - thrust::find_if(thrust::retag(vec.begin()), thrust::retag(vec.end()), ::internal::identity{}); + thrust::find_if(thrust::retag(vec.begin()), thrust::retag(vec.end()), _THRUST_STD::identity{}); ASSERT_EQ(13, vec.front()); } @@ -212,7 +211,7 @@ TEST(FindTests, TestFindIfNotDispatchExplicit) thrust::device_vector vec(1); my_system sys(0); - thrust::find_if_not(sys, vec.begin(), vec.end(), ::internal::identity{}); + thrust::find_if_not(sys, vec.begin(), vec.end(), _THRUST_STD::identity{}); ASSERT_EQ(true, sys.is_valid()); } @@ -230,7 +229,7 @@ TEST(FindTests, TestFindIfNotDispatchImplicit) thrust::device_vector vec(1); - thrust::find_if_not(thrust::retag(vec.begin()), thrust::retag(vec.end()), ::internal::identity{}); + thrust::find_if_not(thrust::retag(vec.begin()), thrust::retag(vec.end()), _THRUST_STD::identity{}); ASSERT_EQ(13, vec.front()); } @@ -347,12 +346,12 @@ void TestFindWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(1); thrust::counting_iterator end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); - _THRUST_STD::intmax_t distance_low_value = thrust::distance(begin, thrust::find(thrust::device, begin, end, 17)); + _THRUST_STD::intmax_t distance_low_value = _THRUST_STD::distance(begin, thrust::find(thrust::device, begin, end, 17)); _THRUST_STD::intmax_t distance_high_value = - thrust::distance(begin, thrust::find(thrust::device, begin, end, (1ll << magnitude) - 17)); + _THRUST_STD::distance(begin, thrust::find(thrust::device, begin, end, (1ll << magnitude) - 17)); ASSERT_EQ(distance_low_value, 16); ASSERT_EQ(distance_high_value, (1ll << magnitude) - 18); diff --git a/projects/rocthrust/test/test_for_each.cpp b/projects/rocthrust/test/test_for_each.cpp index 213177f11497..5407263eedca 100644 --- a/projects/rocthrust/test/test_for_each.cpp +++ b/projects/rocthrust/test/test_for_each.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. @@ -519,7 +519,7 @@ void TestForEachWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(0); thrust::counting_iterator end = begin + (1ull << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); thrust::device_ptr has_executed = thrust::device_malloc(1); *has_executed = false; diff --git a/projects/rocthrust/test/test_functional.cpp b/projects/rocthrust/test/test_functional.cpp index c3ffa62b6be2..dc710d10a80d 100644 --- a/projects/rocthrust/test/test_functional.cpp +++ b/projects/rocthrust/test/test_functional.cpp @@ -15,7 +15,6 @@ * limitations under the License. */ -#include #include #include #include @@ -137,63 +136,63 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestBinaryFunctional() } // op(T) -> T -#define DECLARE_UNARY_ARITHMETIC_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ - TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ - { \ - using Vector = typename TestFixture::input_type; \ - using data_type = typename Vector::value_type; \ - \ - SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ - \ - TestUnaryFunctional, std::operator_name>(); \ +#define DECLARE_UNARY_ARITHMETIC_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ + TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ + { \ + using Vector = typename TestFixture::input_type; \ + using data_type = typename Vector::value_type; \ + \ + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ + \ + TestUnaryFunctional, ::std::operator_name>(); \ } // op(T) -> bool -#define DECLARE_UNARY_LOGICAL_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ - TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ - { \ - using Vector = typename TestFixture::input_type; \ - using data_type = typename Vector::value_type; \ - \ - SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ - \ - TestUnaryFunctional, std::operator_name>(); \ +#define DECLARE_UNARY_LOGICAL_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ + TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ + { \ + using Vector = typename TestFixture::input_type; \ + using data_type = typename Vector::value_type; \ + \ + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ + \ + TestUnaryFunctional, ::std::operator_name>(); \ } // op(T,T) -> T -#define DECLARE_BINARY_ARITHMETIC_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ - TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ - { \ - using Vector = typename TestFixture::input_type; \ - using data_type = typename Vector::value_type; \ - \ - SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ - \ - TestBinaryFunctional, std::operator_name>(); \ +#define DECLARE_BINARY_ARITHMETIC_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ + TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ + { \ + using Vector = typename TestFixture::input_type; \ + using data_type = typename Vector::value_type; \ + \ + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ + \ + TestBinaryFunctional, ::std::operator_name>(); \ } // op(T,T) -> T (for integer T only) -#define DECLARE_BINARY_INTEGER_ARITHMETIC_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ - TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ - { \ - using Vector = typename TestFixture::input_type; \ - using data_type = typename Vector::value_type; \ - \ - SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ - \ - TestBinaryFunctional, std::operator_name>(); \ +#define DECLARE_BINARY_INTEGER_ARITHMETIC_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ + TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ + { \ + using Vector = typename TestFixture::input_type; \ + using data_type = typename Vector::value_type; \ + \ + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ + \ + TestBinaryFunctional, ::std::operator_name>(); \ } // op(T,T) -> bool -#define DECLARE_BINARY_LOGICAL_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ - TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ - { \ - using Vector = typename TestFixture::input_type; \ - using data_type = typename Vector::value_type; \ - \ - SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ - \ - TestBinaryFunctional, std::operator_name>(); \ +#define DECLARE_BINARY_LOGICAL_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ + TYPED_TEST(AllTypesTests, Test##OperatorName##Functional) \ + { \ + using Vector = typename TestFixture::input_type; \ + using data_type = typename Vector::value_type; \ + \ + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); \ + \ + TestBinaryFunctional, ::std::operator_name>(); \ } THRUST_DIAG_PUSH @@ -219,7 +218,7 @@ TYPED_TEST(VectorTests, TestIdentityFunctionalVector) THRUST_DISABLE_BROKEN_GCC_ using Vector = typename TestFixture::input_type; Vector input{0, 1, 2, 3}; Vector output(4); - thrust::transform(input.begin(), input.end(), output.begin(), ::internal::identity{}); + thrust::transform(input.begin(), input.end(), output.begin(), _THRUST_STD::identity{}); ASSERT_EQ(input, output); } @@ -269,7 +268,7 @@ TYPED_TEST(VectorTests, TestMaximumFunctional) THRUST_DISABLE_BROKEN_GCC_VECTORI Vector output(4); - thrust::transform(input1.begin(), input1.end(), input2.begin(), output.begin(), thrust::maximum()); + thrust::transform(input1.begin(), input1.end(), input2.begin(), output.begin(), _THRUST_LIBCXX::maximum()); Vector ref{8, 6, 9, 7}; ASSERT_EQ(output, ref); @@ -287,7 +286,7 @@ TYPED_TEST(VectorTests, TestMinimumFunctional) THRUST_DISABLE_BROKEN_GCC_VECTORI Vector output(4); - thrust::transform(input1.begin(), input1.end(), input2.begin(), output.begin(), thrust::minimum()); + thrust::transform(input1.begin(), input1.end(), input2.begin(), output.begin(), _THRUST_LIBCXX::minimum()); Vector ref{5, 3, 7, 3}; ASSERT_EQ(output, ref); @@ -303,7 +302,7 @@ TYPED_TEST(IntegralVectorTests, TestNot1) THRUST_DISABLE_BROKEN_GCC_VECTORIZER Vector output(5); - thrust::transform(input.begin(), input.end(), output.begin(), thrust::not_fn(::internal::identity{})); + thrust::transform(input.begin(), input.end(), output.begin(), _THRUST_STD::not_fn(_THRUST_STD::identity{})); Vector ref{0, 1, 0, 0, 1}; ASSERT_EQ(output, ref); @@ -321,7 +320,8 @@ TYPED_TEST(VectorTests, TestNot2) THRUST_DISABLE_BROKEN_GCC_VECTORIZER Vector output(5); - thrust::transform(input1.begin(), input1.end(), input2.begin(), output.begin(), thrust::not_fn(thrust::equal_to())); + thrust::transform( + input1.begin(), input1.end(), input2.begin(), output.begin(), _THRUST_STD::not_fn(_THRUST_STD::equal_to())); Vector ref{0, 1, 1, 0, 1}; ASSERT_EQ(output, ref); diff --git a/projects/rocthrust/test/test_functional_arithmetic.cpp b/projects/rocthrust/test/test_functional_arithmetic.cpp index 801c0b33e831..561ea0f59217 100644 --- a/projects/rocthrust/test/test_functional_arithmetic.cpp +++ b/projects/rocthrust/test/test_functional_arithmetic.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2025-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. @@ -68,8 +68,8 @@ void TestBinaryFunctional() #define INSTANTIATE_BINARY_ARITHMETIC_FUNCTIONAL_TEST(vector_type, operator_name, data_type) \ TestBinaryFunctional, \ thrust::vector_type, \ - thrust::operator_name, \ - std::operator_name>(); + _THRUST_STD::operator_name, \ + ::std::operator_name>(); // op(T,T) -> T #define DECLARE_BINARY_ARITHMETIC_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ TEST(FunctionalArithmeticTests, Test##OperatorName##FunctionalHost) \ diff --git a/projects/rocthrust/test/test_functional_bitwise.cpp b/projects/rocthrust/test/test_functional_bitwise.cpp index a881f7445378..0272c8beac41 100644 --- a/projects/rocthrust/test/test_functional_bitwise.cpp +++ b/projects/rocthrust/test/test_functional_bitwise.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2025-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,39 +27,6 @@ const size_t NUM_SAMPLES = 10000; -// STL doesn't necessarily have these available -namespace ref -{ - -template -struct bit_and -{ - T operator()(const T& lhs, const T& rhs) const - { - return lhs & rhs; - } -}; - -template -struct bit_or -{ - T operator()(const T& lhs, const T& rhs) const - { - return lhs | rhs; - } -}; - -template -struct bit_xor -{ - T operator()(const T& lhs, const T& rhs) const - { - return lhs ^ rhs; - } -}; - -} // namespace ref - template void TestBinaryFunctional() { @@ -107,8 +74,8 @@ Macro(vector_type, operator_name, uint64_t) #define INSTANTIATE_BINARY_BITWISE_FUNCTIONAL_TEST(vector_type, operator_name, data_type) \ TestBinaryFunctional, \ thrust::vector_type, \ - thrust::operator_name, \ - ref::operator_name>(); + _THRUST_STD::operator_name, \ + std::operator_name>(); // op(T,T) -> T (for bitwise op and integer T only) #define DECLARE_BINARY_BITWISE_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ diff --git a/projects/rocthrust/test/test_functional_logical.cpp b/projects/rocthrust/test/test_functional_logical.cpp index 46a3e64b8940..24e0d821b117 100644 --- a/projects/rocthrust/test/test_functional_logical.cpp +++ b/projects/rocthrust/test/test_functional_logical.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2025-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. @@ -72,8 +72,8 @@ Macro(vector_type, operator_name, uint64_t) #define INSTANTIATE_BINARY_LOGICAL_FUNCTIONAL_TEST(vector_type, operator_name, data_type) \ TestBinaryFunctional, \ thrust::vector_type, \ - thrust::operator_name, \ - std::operator_name>(); + _THRUST_STD::operator_name, \ + ::std::operator_name>(); // op(T,T) -> bool #define DECLARE_BINARY_LOGICAL_FUNCTIONAL_UNITTEST(operator_name, OperatorName) \ diff --git a/projects/rocthrust/test/test_functional_placeholders_arithmetic.cpp b/projects/rocthrust/test/test_functional_placeholders_arithmetic.cpp index 1f09194c812f..19cdf659d888 100644 --- a/projects/rocthrust/test/test_functional_placeholders_arithmetic.cpp +++ b/projects/rocthrust/test/test_functional_placeholders_arithmetic.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2025-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. @@ -100,11 +100,11 @@ TESTS_DEFINE(UnaryTests, UnaryTestsParams); ASSERT_EQ(reference, result); \ } -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Plus, +, thrust::plus, BinaryArithmeticTests); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Minus, -, thrust::minus, BinaryArithmeticTests); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Multiplies, *, thrust::multiplies, BinaryArithmeticTests); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Divides, /, thrust::divides, BinaryArithmeticTests); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Modulus, %, thrust::modulus, ModulusTests); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Plus, +, _THRUST_STD::plus, BinaryArithmeticTests); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Minus, -, _THRUST_STD::minus, BinaryArithmeticTests); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Multiplies, *, _THRUST_STD::multiplies, BinaryArithmeticTests); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Divides, /, _THRUST_STD::divides, BinaryArithmeticTests); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Modulus, %, _THRUST_STD::modulus, ModulusTests); #define UNARY_FUNCTIONAL_PLACEHOLDERS_TEST(name, reference_operator, functor) \ TYPED_TEST(UnaryTests, TestFunctionalPlaceholders##name) \ @@ -137,6 +137,6 @@ struct unary_plus_reference }; UNARY_FUNCTIONAL_PLACEHOLDERS_TEST(UnaryPlus, +, unary_plus_reference); -UNARY_FUNCTIONAL_PLACEHOLDERS_TEST(Negate, -, thrust::negate); +UNARY_FUNCTIONAL_PLACEHOLDERS_TEST(Negate, -, _THRUST_STD::negate); THRUST_DIAG_POP diff --git a/projects/rocthrust/test/test_functional_placeholders_bitwise.cpp b/projects/rocthrust/test/test_functional_placeholders_bitwise.cpp index aba5cdc933cf..70aae02632aa 100644 --- a/projects/rocthrust/test/test_functional_placeholders_bitwise.cpp +++ b/projects/rocthrust/test/test_functional_placeholders_bitwise.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2025-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. @@ -117,9 +117,9 @@ struct rebind_vector, U> THRUST_DIAG_PUSH THRUST_DIAG_SUPPRESS_MSVC(4244) // warning C4244: '=': conversion from 'int' to '_Ty', possible loss of data -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(BitAnd, &, thrust::bit_and); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(BitOr, |, thrust::bit_or); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(BitXor, ^, thrust::bit_xor); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(BitAnd, &, _THRUST_STD::bit_and); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(BitOr, |, _THRUST_STD::bit_or); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(BitXor, ^, _THRUST_STD::bit_xor); template struct bit_negate_reference diff --git a/projects/rocthrust/test/test_functional_placeholders_logical.cpp b/projects/rocthrust/test/test_functional_placeholders_logical.cpp index 82af1b76477a..72eb0d47db47 100644 --- a/projects/rocthrust/test/test_functional_placeholders_logical.cpp +++ b/projects/rocthrust/test/test_functional_placeholders_logical.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2025-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. @@ -96,8 +96,8 @@ struct rebind_vector, U> ASSERT_EQ(reference, result); \ } -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(LogicalAnd, &&, thrust::logical_and); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(LogicalOr, ||, thrust::logical_or); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(LogicalAnd, &&, _THRUST_STD::logical_and); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(LogicalOr, ||, _THRUST_STD::logical_or); TYPED_TEST(FunctionalPlaceholdersLogicalNotTest, TestFunctionalPlaceholdersLogicalNot) { @@ -115,7 +115,7 @@ TYPED_TEST(FunctionalPlaceholdersLogicalNotTest, TestFunctionalPlaceholdersLogic } // end if bool_vector reference(input.size()); - thrust::transform(input.begin(), input.end(), reference.begin(), thrust::logical_not()); + thrust::transform(input.begin(), input.end(), reference.begin(), _THRUST_STD::logical_not()); using namespace thrust::placeholders; bool_vector result(input.size()); diff --git a/projects/rocthrust/test/test_functional_placeholders_relational.cpp b/projects/rocthrust/test/test_functional_placeholders_relational.cpp index 8bb39e4d19aa..cf2025854b16 100644 --- a/projects/rocthrust/test/test_functional_placeholders_relational.cpp +++ b/projects/rocthrust/test/test_functional_placeholders_relational.cpp @@ -1,6 +1,6 @@ /* * Copyright 2008-2013 NVIDIA Corporation - * Modifications Copyright© 2025 Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright© 2025-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. @@ -85,9 +85,9 @@ struct rebind_vector, U> ASSERT_EQ(reference, result); \ } -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(EqualTo, ==, thrust::equal_to); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(NotEqualTo, !=, thrust::not_equal_to); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Greater, >, thrust::greater); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Less, <, thrust::less); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(GreaterEqual, >=, thrust::greater_equal); -BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(LessEqual, <=, thrust::less_equal); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(EqualTo, ==, _THRUST_STD::equal_to); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(NotEqualTo, !=, _THRUST_STD::not_equal_to); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Greater, >, _THRUST_STD::greater); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(Less, <, _THRUST_STD::less); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(GreaterEqual, >=, _THRUST_STD::greater_equal); +BINARY_FUNCTIONAL_PLACEHOLDERS_TEST(LessEqual, <=, _THRUST_STD::less_equal); diff --git a/projects/rocthrust/test/test_generate.cpp b/projects/rocthrust/test/test_generate.cpp index 4dba986a6b7e..89c5727c540b 100644 --- a/projects/rocthrust/test/test_generate.cpp +++ b/projects/rocthrust/test/test_generate.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. @@ -242,8 +242,8 @@ TYPED_TEST(GenerateVectorTests, TestGenerateZipIterator) Vector v1(3, T(0)); Vector v2(3, T(0)); - thrust::generate(thrust::make_zip_iterator(thrust::make_tuple(v1.begin(), v2.begin())), - thrust::make_zip_iterator(thrust::make_tuple(v1.end(), v2.end())), + thrust::generate(thrust::make_zip_iterator(v1.begin(), v2.begin()), + thrust::make_zip_iterator(v1.end(), v2.end()), return_value>(thrust::tuple(4, 7))); Vector ref1(3, 4); diff --git a/projects/rocthrust/test/test_inner_product.cpp b/projects/rocthrust/test/test_inner_product.cpp index 6cf7d0850ec8..3d41ea0ae897 100644 --- a/projects/rocthrust/test/test_inner_product.cpp +++ b/projects/rocthrust/test/test_inner_product.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. @@ -131,7 +131,8 @@ TYPED_TEST(VectorUnitInnerProductTests, TestInnerProductWithOperator) // compute (v1 - v2) and perform a multiplies reduction T init = 3; - T result = thrust::inner_product(v1.begin(), v1.end(), v2.begin(), init, thrust::multiplies(), thrust::minus()); + T result = thrust::inner_product( + v1.begin(), v1.end(), v2.begin(), init, _THRUST_STD::multiplies(), _THRUST_STD::minus()); ASSERT_EQ(result, 90); } @@ -180,14 +181,14 @@ void TestInnerProductWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(1); thrust::counting_iterator end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); thrust::device_ptr has_executed = thrust::device_malloc(1); *has_executed = false; only_set_when_both_expected fn = {(1ll << magnitude) - 1, thrust::raw_pointer_cast(has_executed)}; - ASSERT_EQ(thrust::inner_product(thrust::device, begin, end, begin, 0ll, thrust::plus(), fn), + ASSERT_EQ(thrust::inner_product(thrust::device, begin, end, begin, 0ll, _THRUST_STD::plus(), fn), (1ll << magnitude)); bool has_executed_h = *has_executed; @@ -217,7 +218,8 @@ TEST(InnerProductTests, TestInnerProductPlaceholders) thrust::device_vector v1(100, 1.f); thrust::device_vector v2(100, 1.f); - auto result = thrust::inner_product(v1.begin(), v1.end(), v2.begin(), 0.0f, thrust::plus{}, _1 * _2 + 1.0f); + auto result = + thrust::inner_product(v1.begin(), v1.end(), v2.begin(), 0.0f, _THRUST_STD::plus{}, _1 * _2 + 1.0f); auto error_margin = 1e-4 * (::std::abs(result) + 200.f) * 1e-4; ASSERT_NEAR(static_cast(result), static_cast(200.f), static_cast(error_margin)); diff --git a/projects/rocthrust/test/test_is_contiguous_iterator.cpp b/projects/rocthrust/test/test_is_contiguous_iterator.cpp index 9444a0a161dd..729f670bf6be 100644 --- a/projects/rocthrust/test/test_is_contiguous_iterator.cpp +++ b/projects/rocthrust/test/test_is_contiguous_iterator.cpp @@ -44,11 +44,11 @@ TESTS_DEFINE(IsContiguousIteratorTests, FullTestsParams); -static_assert(thrust::is_contiguous_iterator::value); -static_assert(thrust::is_contiguous_iterator::value); -static_assert(thrust::is_contiguous_iterator::value); -static_assert(thrust::is_contiguous_iterator::value); -static_assert(!thrust::is_contiguous_iterator::iterator>::value); +static_assert(thrust::is_contiguous_iterator_v); +static_assert(thrust::is_contiguous_iterator_v); +static_assert(thrust::is_contiguous_iterator_v); +static_assert(thrust::is_contiguous_iterator_v); +static_assert(!thrust::is_contiguous_iterator_v::iterator>); TYPED_TEST(IsContiguousIteratorTests, test_is_contiguous_iterator) { @@ -57,24 +57,44 @@ TYPED_TEST(IsContiguousIteratorTests, test_is_contiguous_iterator) SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); - static_assert(thrust::is_contiguous_iterator::value); - static_assert(thrust::is_contiguous_iterator::value); - static_assert(thrust::is_contiguous_iterator>::value); - static_assert(thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::reverse_iterator>::value); - static_assert(thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator::iterator>::value); - static_assert(!thrust::is_contiguous_iterator>::value); - static_assert(!thrust::is_contiguous_iterator>::value); + static_assert(thrust::is_contiguous_iterator_v); + static_assert(thrust::is_contiguous_iterator_v); + static_assert(thrust::is_contiguous_iterator_v>); + static_assert(thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::reverse_iterator>); + static_assert(thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v::iterator>); + static_assert(!thrust::is_contiguous_iterator_v>); + static_assert(!thrust::is_contiguous_iterator_v>); +} + +TYPED_TEST(IsContiguousIteratorTests, test_is_contiguous_iterator_cvref) +{ + using Vector = typename TestFixture::input_type; + using T = typename Vector::value_type; + + SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); + + static_assert(thrust::is_contiguous_iterator_v); + static_assert(thrust::is_contiguous_iterator_v); + static_assert(thrust::is_contiguous_iterator_v); + static_assert(thrust::is_contiguous_iterator_v); + static_assert(thrust::is_contiguous_iterator_v); + + static_assert(!thrust::is_contiguous_iterator_v::iterator const>); + static_assert(!thrust::is_contiguous_iterator_v::iterator volatile>); + static_assert(!thrust::is_contiguous_iterator_v::iterator&>); + static_assert(!thrust::is_contiguous_iterator_v::iterator const&>); + static_assert(!thrust::is_contiguous_iterator_v::iterator volatile&>); } TYPED_TEST(IsContiguousIteratorTests, test_is_contiguous_iterator_vectors) @@ -83,7 +103,7 @@ TYPED_TEST(IsContiguousIteratorTests, test_is_contiguous_iterator_vectors) SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest()); - static_assert(thrust::is_contiguous_iterator::value); + static_assert(thrust::is_contiguous_iterator_v); } struct expect_pointer diff --git a/projects/rocthrust/test/test_is_operator_function_object.cpp b/projects/rocthrust/test/test_is_operator_function_object.cpp index 9bb9fae8dc87..1af512e69f93 100644 --- a/projects/rocthrust/test/test_is_operator_function_object.cpp +++ b/projects/rocthrust/test/test_is_operator_function_object.cpp @@ -45,10 +45,10 @@ TYPED_TEST(IsOperatorFunctionObjectTest, test_is_operator_less_function_object) using T = typename TestFixture::input_type; - static_assert(thrust::is_operator_less_function_object>::value); - static_assert(!thrust::is_operator_less_function_object>::value); - static_assert(!thrust::is_operator_less_function_object>::value); - static_assert(!thrust::is_operator_less_function_object>::value); + static_assert(thrust::is_operator_less_function_object<_THRUST_STD::less>::value); + static_assert(!thrust::is_operator_less_function_object<_THRUST_STD::greater>::value); + static_assert(!thrust::is_operator_less_function_object<_THRUST_STD::less_equal>::value); + static_assert(!thrust::is_operator_less_function_object<_THRUST_STD::less_equal>::value); static_assert(thrust::is_operator_less_function_object>::value); static_assert(!thrust::is_operator_less_function_object>::value); static_assert(!thrust::is_operator_less_function_object>::value); @@ -62,10 +62,10 @@ TYPED_TEST(IsOperatorFunctionObjectTest, test_is_operator_greater_function_objec using T = typename TestFixture::input_type; - static_assert(!thrust::is_operator_greater_function_object>::value); - static_assert(thrust::is_operator_greater_function_object>::value); - static_assert(!thrust::is_operator_greater_function_object>::value); - static_assert(!thrust::is_operator_greater_function_object>::value); + static_assert(!thrust::is_operator_greater_function_object<_THRUST_STD::less>::value); + static_assert(thrust::is_operator_greater_function_object<_THRUST_STD::greater>::value); + static_assert(!thrust::is_operator_greater_function_object<_THRUST_STD::less_equal>::value); + static_assert(!thrust::is_operator_greater_function_object<_THRUST_STD::less_equal>::value); static_assert(!thrust::is_operator_greater_function_object>::value); static_assert(thrust::is_operator_greater_function_object>::value); static_assert(!thrust::is_operator_greater_function_object>::value); @@ -79,10 +79,10 @@ TYPED_TEST(IsOperatorFunctionObjectTest, test_is_operator_less_or_greater_functi using T = typename TestFixture::input_type; - static_assert(thrust::is_operator_less_or_greater_function_object>::value); - static_assert(thrust::is_operator_less_or_greater_function_object>::value); - static_assert(!thrust::is_operator_less_or_greater_function_object>::value); - static_assert(!thrust::is_operator_less_or_greater_function_object>::value); + static_assert(thrust::is_operator_less_or_greater_function_object<_THRUST_STD::less>::value); + static_assert(thrust::is_operator_less_or_greater_function_object<_THRUST_STD::greater>::value); + static_assert(!thrust::is_operator_less_or_greater_function_object<_THRUST_STD::less_equal>::value); + static_assert(!thrust::is_operator_less_or_greater_function_object<_THRUST_STD::less_equal>::value); static_assert(thrust::is_operator_less_or_greater_function_object>::value); static_assert(thrust::is_operator_less_or_greater_function_object>::value); static_assert(!thrust::is_operator_less_or_greater_function_object>::value); @@ -96,10 +96,10 @@ TYPED_TEST(IsOperatorFunctionObjectTest, test_is_operator_plus_function_object) using T = typename TestFixture::input_type; - static_assert(thrust::is_operator_plus_function_object>::value); - static_assert(!thrust::is_operator_plus_function_object>::value); - static_assert(!thrust::is_operator_plus_function_object>::value); - static_assert(!thrust::is_operator_plus_function_object>::value); + static_assert(thrust::is_operator_plus_function_object<_THRUST_STD::plus>::value); + static_assert(!thrust::is_operator_plus_function_object<_THRUST_STD::minus>::value); + static_assert(!thrust::is_operator_plus_function_object<_THRUST_STD::less>::value); + static_assert(!thrust::is_operator_plus_function_object<_THRUST_STD::greater>::value); static_assert(thrust::is_operator_plus_function_object>::value); static_assert(!thrust::is_operator_plus_function_object>::value); static_assert(!thrust::is_operator_plus_function_object>::value); diff --git a/projects/rocthrust/test/test_is_partitioned.cpp b/projects/rocthrust/test/test_is_partitioned.cpp index 500360df6c10..1a721137661f 100644 --- a/projects/rocthrust/test/test_is_partitioned.cpp +++ b/projects/rocthrust/test/test_is_partitioned.cpp @@ -15,7 +15,6 @@ * limitations under the License. */ -#include #include #include #include @@ -46,24 +45,24 @@ TYPED_TEST(IsPartitionedVectorTests, TestIsPartitionedSimple) Vector v{1, 1, 1, 0}; // empty partition - ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin(), v.begin(), ::internal::identity{})); + ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin(), v.begin(), _THRUST_STD::identity{})); // one element true partition - ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin(), v.begin() + 1, ::internal::identity{})); + ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin(), v.begin() + 1, _THRUST_STD::identity{})); // just true partition - ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin(), v.begin() + 2, ::internal::identity{})); + ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin(), v.begin() + 2, _THRUST_STD::identity{})); // both true & false partitions - ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin(), v.end(), ::internal::identity{})); + ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin(), v.end(), _THRUST_STD::identity{})); // one element false partition - ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin() + 3, v.end(), ::internal::identity{})); + ASSERT_EQ_QUIET(true, thrust::is_partitioned(v.begin() + 3, v.end(), _THRUST_STD::identity{})); v = {1, 0, 1, 1}; // not partitioned - ASSERT_EQ_QUIET(false, thrust::is_partitioned(v.begin(), v.end(), ::internal::identity{})); + ASSERT_EQ_QUIET(false, thrust::is_partitioned(v.begin(), v.end(), _THRUST_STD::identity{})); } TYPED_TEST(IsPartitionedVectorTests, TestIsPartitioned) diff --git a/projects/rocthrust/test/test_is_sorted.cpp b/projects/rocthrust/test/test_is_sorted.cpp index 98b71149ca1a..7982053d0317 100644 --- a/projects/rocthrust/test/test_is_sorted.cpp +++ b/projects/rocthrust/test/test_is_sorted.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. @@ -47,10 +47,10 @@ TYPED_TEST(IsSortedVectorTests, TestIsSortedSimple) ASSERT_EQ(thrust::is_sorted(v.begin(), v.begin() + 3), true); ASSERT_EQ(thrust::is_sorted(v.begin(), v.begin() + 4), false); - ASSERT_EQ(thrust::is_sorted(v.begin(), v.begin() + 3, thrust::less()), true); + ASSERT_EQ(thrust::is_sorted(v.begin(), v.begin() + 3, _THRUST_STD::less()), true); - ASSERT_EQ(thrust::is_sorted(v.begin(), v.begin() + 1, thrust::greater()), true); - ASSERT_EQ(thrust::is_sorted(v.begin(), v.begin() + 4, thrust::greater()), false); + ASSERT_EQ(thrust::is_sorted(v.begin(), v.begin() + 1, _THRUST_STD::greater()), true); + ASSERT_EQ(thrust::is_sorted(v.begin(), v.begin() + 4, _THRUST_STD::greater()), false); ASSERT_EQ(thrust::is_sorted(v.begin(), v.end()), false); } diff --git a/projects/rocthrust/test/test_is_sorted_until.cpp b/projects/rocthrust/test/test_is_sorted_until.cpp index fa4ab3da6a91..480f4391a42d 100644 --- a/projects/rocthrust/test/test_is_sorted_until.cpp +++ b/projects/rocthrust/test/test_is_sorted_until.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. @@ -59,24 +59,24 @@ TYPED_TEST(IsSortedUntilVectorTests, TestIsSortedUntilSimple) last = v.begin() + 3; ref = v.begin() + 3; - ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, thrust::less())); + ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, _THRUST_STD::less())); last = v.begin() + 4; ref = v.begin() + 3; - ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, thrust::less())); + ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, _THRUST_STD::less())); last = v.begin() + 1; ref = v.begin() + 1; - ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, thrust::greater())); + ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, _THRUST_STD::greater())); last = v.begin() + 4; ref = v.begin() + 1; - ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, thrust::greater())); + ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, _THRUST_STD::greater())); first = v.begin() + 2; last = v.begin() + 4; ref = v.begin() + 4; - ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, thrust::greater())); + ASSERT_EQ_QUIET(ref, thrust::is_sorted_until(first, last, _THRUST_STD::greater())); } TYPED_TEST(IsSortedUntilVectorTests, TestIsSortedUntilRepeatedElements) diff --git a/projects/rocthrust/test/test_libcxx_iterators.cpp b/projects/rocthrust/test/test_libcxx_iterators.cpp new file mode 100644 index 000000000000..3a1afd8407d3 --- /dev/null +++ b/projects/rocthrust/test/test_libcxx_iterators.cpp @@ -0,0 +1,226 @@ +/* + * Copyright 2026 NVIDIA Corporation + * Modifications Copyright© 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. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#if _THRUST_HAS_DEVICE_SYSTEM_STD + +# include _THRUST_LIBCXX_INCLUDE(iterator) + +# include + +# include "test_utils.hpp" + +TEST(LibcxxIteratorTests, TestLibcxxDiscardIterator) +{ + auto discard = _THRUST_LIBCXX::discard_iterator{}; + { // device system + thrust::device_vector vec{1, 2, 3, 4}; + thrust::copy(vec.begin(), vec.end(), discard); + } + + { // host system + thrust::host_vector vec{1, 2, 3, 4}; + thrust::copy(vec.begin(), vec.end(), discard); + } + + { // plain std::vector + std::vector vec{1, 2, 3, 4}; + thrust::copy(vec.begin(), vec.end(), discard); + } +} + +TEST(LibcxxIteratorTests, TestLibcxxConstantIterator) +{ + { // device system + thrust::device_vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::constant_iterator{42, 0}, _THRUST_LIBCXX::constant_iterator{42, 4}, vec.begin()); + } + + { // host system + thrust::host_vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::constant_iterator{42, 0}, _THRUST_LIBCXX::constant_iterator{42, 4}, vec.begin()); + } + + { // plain std::vector + std::vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::constant_iterator{42, 0}, _THRUST_LIBCXX::constant_iterator{42, 4}, vec.begin()); + } +} + +TEST(LibcxxIteratorTests, TestLibcxxCountingIterator) +{ + { // device system + thrust::device_vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::counting_iterator{0}, _THRUST_LIBCXX::counting_iterator{4}, vec.begin()); + } + + { // host system + thrust::host_vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::counting_iterator{0}, _THRUST_LIBCXX::counting_iterator{4}, vec.begin()); + } + + { // plain std::vector + std::vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::counting_iterator{0}, _THRUST_LIBCXX::counting_iterator{4}, vec.begin()); + } +} + +TEST(LibcxxIteratorTests, TestLibcxxPermutationIterator) +{ + { // device system + thrust::device_vector vec{1, 2, 3, 4, 5, 6, 7, 8, 9}; + thrust::device_vector off{5, 2, 7, 0}; + thrust::device_vector res{-1, -1, -1, -1, -1}; + thrust::copy(_THRUST_LIBCXX::permutation_iterator{vec.begin(), off.begin()}, + _THRUST_LIBCXX::permutation_iterator{vec.begin(), off.end()}, + res.begin()); + ASSERT_TRUE(res == thrust::device_vector{6, 3, 8, 1, -1}); + } + + { // host system + thrust::host_vector vec{1, 2, 3, 4, 5, 6, 7, 8, 9}; + thrust::host_vector off{5, 2, 7, 0}; + thrust::host_vector res{-1, -1, -1, -1, -1}; + thrust::copy(_THRUST_LIBCXX::permutation_iterator{vec.begin(), off.begin()}, + _THRUST_LIBCXX::permutation_iterator{vec.begin(), off.end()}, + res.begin()); + ASSERT_TRUE(res == thrust::host_vector{6, 3, 8, 1, -1}); + } + + { // plain std::vector + std::vector vec{1, 2, 3, 4, 5, 6, 7, 8, 9}; + std::vector off{5, 2, 7, 0}; + std::vector res{-1, -1, -1, -1, -1}; + thrust::copy(_THRUST_LIBCXX::permutation_iterator{vec.begin(), off.begin()}, + _THRUST_LIBCXX::permutation_iterator{vec.begin(), off.end()}, + res.begin()); + ASSERT_TRUE(res == std::vector{6, 3, 8, 1, -1}); + } +} + +TEST(LibcxxIteratorTests, TestLibcxxStridedIterator) +{ + auto discard = _THRUST_LIBCXX::discard_iterator{}; + { // device system + thrust::device_vector vec{1, 2, 3, 4, 5, 6}; + thrust::copy( + _THRUST_LIBCXX::strided_iterator{vec.begin(), 2}, _THRUST_LIBCXX::strided_iterator{vec.end(), 2}, discard); + } + + { // host system + thrust::host_vector vec{1, 2, 3, 4, 5, 6}; + thrust::copy( + _THRUST_LIBCXX::strided_iterator{vec.begin(), 2}, _THRUST_LIBCXX::strided_iterator{vec.end(), 2}, discard); + } + + { // plain std::vector + std::vector vec{1, 2, 3, 4, 5, 6}; + thrust::copy( + _THRUST_LIBCXX::strided_iterator{vec.begin(), 2}, _THRUST_LIBCXX::strided_iterator{vec.end(), 2}, discard); + } +} + +struct is_equal_index +{ + THRUST_HOST_DEVICE constexpr void + operator()([[maybe_unused]] const int index, [[maybe_unused]] const int expected) const noexcept + { + _CCCL_VERIFY(index == expected, "should have right value"); + } +}; + +TEST(LibcxxIteratorTests, TestLibcxxTabulateOutputIterator) +{ + { // device system + thrust::device_vector vec{5, 6, 7, 8, 9}; + thrust::copy(vec.begin(), vec.end(), _THRUST_LIBCXX::make_tabulate_output_iterator(is_equal_index{}, 5)); + } + + { // host system + thrust::host_vector vec{5, 6, 7, 8, 9}; + thrust::copy(vec.begin(), vec.end(), _THRUST_LIBCXX::make_tabulate_output_iterator(is_equal_index{}, 5)); + } + + { // plain std::vector + std::vector vec{5, 6, 7, 8, 9}; + thrust::copy(vec.begin(), vec.end(), _THRUST_LIBCXX::make_tabulate_output_iterator(is_equal_index{}, 5)); + } +} + +struct plus_one +{ + [[nodiscard]] THRUST_HOST_DEVICE constexpr int operator()(const int val) const noexcept + { + return val + 1; + } +}; + +TEST(LibcxxIteratorTests, TestLibcxxTransformOutputIterator) +{ + { // device system + thrust::device_vector vec{-1, -1, -1, -1, -1}; + thrust::copy(_THRUST_LIBCXX::counting_iterator{0}, + _THRUST_LIBCXX::counting_iterator{5}, + _THRUST_LIBCXX::make_transform_output_iterator(vec.begin(), plus_one{})); + thrust::device_vector expected{1, 2, 3, 4, 5}; + ASSERT_TRUE(thrust::equal(vec.begin(), vec.end(), expected.begin())); + } + + { // host system + thrust::host_vector vec{-1, -1, -1, -1, -1}; + thrust::copy(_THRUST_LIBCXX::counting_iterator{0}, + _THRUST_LIBCXX::counting_iterator{5}, + _THRUST_LIBCXX::make_transform_output_iterator(vec.begin(), plus_one{})); + thrust::host_vector expected{1, 2, 3, 4, 5}; + ASSERT_TRUE(thrust::equal(vec.begin(), vec.end(), expected.begin())); + } + + { // plain std::vector + std::vector vec{-1, -1, -1, -1, -1}; + thrust::copy(_THRUST_LIBCXX::counting_iterator{0}, + _THRUST_LIBCXX::counting_iterator{5}, + _THRUST_LIBCXX::make_transform_output_iterator(vec.begin(), plus_one{})); + std::vector expected{1, 2, 3, 4, 5}; + ASSERT_TRUE(thrust::equal(vec.begin(), vec.end(), expected.begin())); + } +} + +TEST(LibcxxIteratorTests, TestLibcxxTransformIterator) +{ + auto discard = _THRUST_LIBCXX::discard_iterator{}; + { // device system + thrust::device_vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::make_transform_iterator(vec.begin(), plus_one{}), + _THRUST_LIBCXX::make_transform_iterator(vec.end(), plus_one{}), + discard); + } + + { // host system + thrust::host_vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::make_transform_iterator(vec.begin(), plus_one{}), + _THRUST_LIBCXX::make_transform_iterator(vec.end(), plus_one{}), + discard); + } + + { // plain std::vector + std::vector vec{1, 2, 3, 4}; + thrust::copy(_THRUST_LIBCXX::make_transform_iterator(vec.begin(), plus_one{}), + _THRUST_LIBCXX::make_transform_iterator(vec.end(), plus_one{}), + discard); + } +} + +#endif diff --git a/projects/rocthrust/test/test_logical.cpp b/projects/rocthrust/test/test_logical.cpp index a52ad713c051..045bf46d0def 100644 --- a/projects/rocthrust/test/test_logical.cpp +++ b/projects/rocthrust/test/test_logical.cpp @@ -15,7 +15,6 @@ * limitations under the License. */ -#include #include #include #include @@ -35,16 +34,16 @@ TYPED_TEST(LogicalTests, TestAllOf) Vector v(3, T{1}); - ASSERT_EQ(thrust::all_of(v.begin(), v.end(), ::internal::identity{}), true); + ASSERT_EQ(thrust::all_of(v.begin(), v.end(), _THRUST_STD::identity{}), true); v[1] = T{0}; - ASSERT_EQ(thrust::all_of(v.begin(), v.end(), ::internal::identity{}), false); + ASSERT_EQ(thrust::all_of(v.begin(), v.end(), _THRUST_STD::identity{}), false); - ASSERT_EQ(thrust::all_of(v.begin() + 0, v.begin() + 0, ::internal::identity{}), true); - ASSERT_EQ(thrust::all_of(v.begin() + 0, v.begin() + 1, ::internal::identity{}), true); - ASSERT_EQ(thrust::all_of(v.begin() + 0, v.begin() + 2, ::internal::identity{}), false); - ASSERT_EQ(thrust::all_of(v.begin() + 1, v.begin() + 2, ::internal::identity{}), false); + ASSERT_EQ(thrust::all_of(v.begin() + 0, v.begin() + 0, _THRUST_STD::identity{}), true); + ASSERT_EQ(thrust::all_of(v.begin() + 0, v.begin() + 1, _THRUST_STD::identity{}), true); + ASSERT_EQ(thrust::all_of(v.begin() + 0, v.begin() + 2, _THRUST_STD::identity{}), false); + ASSERT_EQ(thrust::all_of(v.begin() + 1, v.begin() + 2, _THRUST_STD::identity{}), false); } template @@ -93,16 +92,16 @@ TYPED_TEST(LogicalTests, TestAnyOf) Vector v(3, T{1}); - ASSERT_EQ(thrust::any_of(v.begin(), v.end(), ::internal::identity{}), true); + ASSERT_EQ(thrust::any_of(v.begin(), v.end(), _THRUST_STD::identity{}), true); v[1] = 0; - ASSERT_EQ(thrust::any_of(v.begin(), v.end(), ::internal::identity{}), true); + ASSERT_EQ(thrust::any_of(v.begin(), v.end(), _THRUST_STD::identity{}), true); - ASSERT_EQ(thrust::any_of(v.begin() + 0, v.begin() + 0, ::internal::identity{}), false); - ASSERT_EQ(thrust::any_of(v.begin() + 0, v.begin() + 1, ::internal::identity{}), true); - ASSERT_EQ(thrust::any_of(v.begin() + 0, v.begin() + 2, ::internal::identity{}), true); - ASSERT_EQ(thrust::any_of(v.begin() + 1, v.begin() + 2, ::internal::identity{}), false); + ASSERT_EQ(thrust::any_of(v.begin() + 0, v.begin() + 0, _THRUST_STD::identity{}), false); + ASSERT_EQ(thrust::any_of(v.begin() + 0, v.begin() + 1, _THRUST_STD::identity{}), true); + ASSERT_EQ(thrust::any_of(v.begin() + 0, v.begin() + 2, _THRUST_STD::identity{}), true); + ASSERT_EQ(thrust::any_of(v.begin() + 1, v.begin() + 2, _THRUST_STD::identity{}), false); } template @@ -151,16 +150,16 @@ TYPED_TEST(LogicalTests, TestNoneOf) Vector v(3, T{1}); - ASSERT_EQ(thrust::none_of(v.begin(), v.end(), ::internal::identity{}), false); + ASSERT_EQ(thrust::none_of(v.begin(), v.end(), _THRUST_STD::identity{}), false); v[1] = 0; - ASSERT_EQ(thrust::none_of(v.begin(), v.end(), ::internal::identity{}), false); + ASSERT_EQ(thrust::none_of(v.begin(), v.end(), _THRUST_STD::identity{}), false); - ASSERT_EQ(thrust::none_of(v.begin() + 0, v.begin() + 0, ::internal::identity{}), true); - ASSERT_EQ(thrust::none_of(v.begin() + 0, v.begin() + 1, ::internal::identity{}), false); - ASSERT_EQ(thrust::none_of(v.begin() + 0, v.begin() + 2, ::internal::identity{}), false); - ASSERT_EQ(thrust::none_of(v.begin() + 1, v.begin() + 2, ::internal::identity{}), true); + ASSERT_EQ(thrust::none_of(v.begin() + 0, v.begin() + 0, _THRUST_STD::identity{}), true); + ASSERT_EQ(thrust::none_of(v.begin() + 0, v.begin() + 1, _THRUST_STD::identity{}), false); + ASSERT_EQ(thrust::none_of(v.begin() + 0, v.begin() + 2, _THRUST_STD::identity{}), false); + ASSERT_EQ(thrust::none_of(v.begin() + 1, v.begin() + 2, _THRUST_STD::identity{}), true); } template diff --git a/projects/rocthrust/test/test_max_element.cpp b/projects/rocthrust/test/test_max_element.cpp index 37e1a2e459a2..4907ddb9fb70 100644 --- a/projects/rocthrust/test/test_max_element.cpp +++ b/projects/rocthrust/test/test_max_element.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. @@ -54,8 +54,8 @@ TYPED_TEST(MaxElementTests, TestMaxElementSimple) ASSERT_EQ(*thrust::max_element(data.begin(), data.end()), 5); ASSERT_EQ(thrust::max_element(data.begin(), data.end()) - data.begin(), 1); - ASSERT_EQ(*thrust::max_element(data.begin(), data.end(), thrust::greater()), 1); - ASSERT_EQ(thrust::max_element(data.begin(), data.end(), thrust::greater()) - data.begin(), 2); + ASSERT_EQ(*thrust::max_element(data.begin(), data.end(), _THRUST_STD::greater()), 1); + ASSERT_EQ(thrust::max_element(data.begin(), data.end(), _THRUST_STD::greater()) - data.begin(), 2); } TYPED_TEST(MaxElementVectorUnitTests, TestMaxElementWithTransform) @@ -67,12 +67,12 @@ TYPED_TEST(MaxElementVectorUnitTests, TestMaxElementWithTransform) Vector data{3, 5, 1, 2, 5, 1}; - ASSERT_EQ(*thrust::max_element(thrust::make_transform_iterator(data.begin(), thrust::negate()), - thrust::make_transform_iterator(data.end(), thrust::negate())), + ASSERT_EQ(*thrust::max_element(thrust::make_transform_iterator(data.begin(), _THRUST_STD::negate()), + thrust::make_transform_iterator(data.end(), _THRUST_STD::negate())), -1); - ASSERT_EQ(*thrust::max_element(thrust::make_transform_iterator(data.begin(), thrust::negate()), - thrust::make_transform_iterator(data.end(), thrust::negate()), - thrust::greater()), + ASSERT_EQ(*thrust::max_element(thrust::make_transform_iterator(data.begin(), _THRUST_STD::negate()), + thrust::make_transform_iterator(data.end(), _THRUST_STD::negate()), + _THRUST_STD::greater()), -5); } @@ -100,9 +100,9 @@ TYPED_TEST(MaxElementPrimitiveTests, TestMaxElement) ASSERT_EQ(h_max - h_data.begin(), d_max - d_data.begin()); typename thrust::host_vector::iterator h_min = - thrust::max_element(h_data.begin(), h_data.end(), thrust::greater()); + thrust::max_element(h_data.begin(), h_data.end(), _THRUST_STD::greater()); typename thrust::device_vector::iterator d_min = - thrust::max_element(d_data.begin(), d_data.end(), thrust::greater()); + thrust::max_element(d_data.begin(), d_data.end(), _THRUST_STD::greater()); ASSERT_EQ(h_min - h_data.begin(), d_min - d_data.begin()); } @@ -150,7 +150,7 @@ void TestMaxElementWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(1); thrust::counting_iterator end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); ASSERT_EQ(*thrust::max_element(thrust::device, begin, end), (1ll << magnitude)); } diff --git a/projects/rocthrust/test/test_merge.cpp b/projects/rocthrust/test/test_merge.cpp index ed856e63822a..f1bf34593ff2 100644 --- a/projects/rocthrust/test/test_merge.cpp +++ b/projects/rocthrust/test/test_merge.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. @@ -205,8 +205,8 @@ TYPED_TEST(PrimitiveMergeTests, MergeDescending) thrust::host_vector h_b = get_random_data( size, get_default_limits::min(), get_default_limits::max(), seed + seed_value_addition); - thrust::stable_sort(h_a.begin(), h_a.end(), thrust::greater()); - thrust::stable_sort(h_b.begin(), h_b.end(), thrust::greater()); + thrust::stable_sort(h_a.begin(), h_a.end(), _THRUST_STD::greater()); + thrust::stable_sort(h_b.begin(), h_b.end(), _THRUST_STD::greater()); const thrust::device_vector d_a = h_a; const thrust::device_vector d_b = h_b; @@ -215,10 +215,10 @@ TYPED_TEST(PrimitiveMergeTests, MergeDescending) thrust::device_vector d_result(d_a.size() + d_b.size()); const auto h_end = - thrust::merge(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), h_result.begin(), thrust::greater()); + thrust::merge(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), h_result.begin(), _THRUST_STD::greater()); const auto d_end = - thrust::merge(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), d_result.begin(), thrust::greater()); + thrust::merge(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), d_result.begin(), _THRUST_STD::greater()); ASSERT_EQ(h_result, d_result); ASSERT_TRUE(h_end == h_result.end()); @@ -261,8 +261,8 @@ TEST(PrimitiveMergeTests, TestMergeDevice) thrust::host_vector h_b = get_random_data( size, get_default_limits::min(), get_default_limits::max(), seed + seed_value_addition); - thrust::stable_sort(h_a.begin(), h_a.end(), thrust::greater()); - thrust::stable_sort(h_b.begin(), h_b.end(), thrust::greater()); + thrust::stable_sort(h_a.begin(), h_a.end(), _THRUST_STD::greater()); + thrust::stable_sort(h_b.begin(), h_b.end(), _THRUST_STD::greater()); const thrust::device_vector d_a = h_a; const thrust::device_vector d_b = h_b; diff --git a/projects/rocthrust/test/test_merge_by_key.cpp b/projects/rocthrust/test/test_merge_by_key.cpp index e35c9bb5c5d0..bb198d3992f7 100644 --- a/projects/rocthrust/test/test_merge_by_key.cpp +++ b/projects/rocthrust/test/test_merge_by_key.cpp @@ -32,7 +32,7 @@ # include #endif -template > +template > struct ParamsMerge { using input_type = Input; @@ -48,22 +48,22 @@ class MergeByKeyTestsClass : public ::testing::Test }; using MergeByKeyTestsParams = ::testing::Types< - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>, - ParamsMerge>>; + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>, + ParamsMerge>>; TYPED_TEST_SUITE(MergeByKeyTestsClass, MergeByKeyTestsParams); @@ -188,7 +188,7 @@ auto call_merge_by_key(Args&&... args) -> decltype(thrust::merge_by_key(std::for else { // TODO(bgruber): remove next line in C++17 and pass CompareOp{} directly to stable_sort - using C = _THRUST_STD::conditional_t<_THRUST_STD::is_void::value, thrust::less, CompareOp>; + using C = _THRUST_STD::conditional_t<_THRUST_STD::is_void::value, _THRUST_STD::less, CompareOp>; return thrust::merge_by_key(std::forward(args)..., C{}); } __builtin_unreachable(); @@ -234,8 +234,8 @@ TYPED_TEST(MergeByKeyTestsClass, TestMergeByKey) else { // TODO(bgruber): remove next line in C++17 and pass compare_function{} directly to stable_sort - using C = - _THRUST_STD::conditional_t<_THRUST_STD::is_void::value, thrust::less, compare_function>; + using C = _THRUST_STD:: + conditional_t<_THRUST_STD::is_void::value, _THRUST_STD::less, compare_function>; thrust::stable_sort(h_a_keys.begin(), h_a_keys.end(), C{}); thrust::stable_sort(h_b_keys.begin(), h_b_keys.end(), C{}); } @@ -468,8 +468,8 @@ TEST(MergeByKeyTests, TestMergeByKeyDevice) thrust::host_vector h_values_b = get_random_data( size, get_default_limits::min(), get_default_limits::max(), seed + seed_value_addition); - thrust::stable_sort(h_keys_a.begin(), h_keys_a.end(), thrust::greater()); - thrust::stable_sort(h_keys_b.begin(), h_keys_b.end(), thrust::greater()); + thrust::stable_sort(h_keys_a.begin(), h_keys_a.end(), _THRUST_STD::greater()); + thrust::stable_sort(h_keys_b.begin(), h_keys_b.end(), _THRUST_STD::greater()); thrust::device_vector d_keys_a = h_keys_a; thrust::device_vector d_keys_b = h_keys_b; diff --git a/projects/rocthrust/test/test_merge_key_value.cpp b/projects/rocthrust/test/test_merge_key_value.cpp index b6577ec85311..2b8ea1325106 100644 --- a/projects/rocthrust/test/test_merge_key_value.cpp +++ b/projects/rocthrust/test/test_merge_key_value.cpp @@ -28,7 +28,7 @@ # include #endif -template >> +template >> struct ParamsMerge { using input_type = Input; @@ -44,22 +44,22 @@ class MergeKeyValueTestsClass : public ::testing::Test }; using MergeKeyValueTestsParams = ::testing::Types< - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>, - ParamsMerge>>>; + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>, + ParamsMerge>>>; TYPED_TEST_SUITE(MergeKeyValueTestsClass, MergeKeyValueTestsParams); @@ -75,7 +75,7 @@ auto call_merge(Args&&... args) -> decltype(thrust::merge(std::forward(arg else { // TODO(bgruber): remove next line in C++17 and pass CompareOp{} directly to stable_sort - using C = _THRUST_STD::conditional_t<_THRUST_STD::is_void::value, thrust::less, CompareOp>; + using C = _THRUST_STD::conditional_t<_THRUST_STD::is_void::value, _THRUST_STD::less, CompareOp>; return thrust::merge(std::forward(args)..., C{}); } __builtin_unreachable(); @@ -121,8 +121,8 @@ TYPED_TEST(MergeKeyValueTestsClass, TestMergeKeyValue) else { // TODO(bgruber): remove next line in C++17 and pass compare_function{} directly to stable_sort - using C = - _THRUST_STD::conditional_t<_THRUST_STD::is_void::value, thrust::less, compare_function>; + using C = _THRUST_STD:: + conditional_t<_THRUST_STD::is_void::value, _THRUST_STD::less, compare_function>; thrust::stable_sort(h_a.begin(), h_a.end(), C{}); thrust::stable_sort(h_b.begin(), h_b.end(), C{}); } diff --git a/projects/rocthrust/test/test_min_and_max.cpp b/projects/rocthrust/test/test_min_and_max.cpp index 431934df3418..149d8bb289cf 100644 --- a/projects/rocthrust/test/test_min_and_max.cpp +++ b/projects/rocthrust/test/test_min_and_max.cpp @@ -47,13 +47,13 @@ TYPED_TEST(MinAndMaxTests, TestMin) // 2 < 3 T two(2), three(3); ASSERT_EQ(two, _THRUST_STD::min(two, three)); - ASSERT_EQ(two, _THRUST_STD::min(two, three, thrust::less())); + ASSERT_EQ(two, _THRUST_STD::min(two, three, _THRUST_STD::less())); ASSERT_EQ(two, _THRUST_STD::min(three, two)); - ASSERT_EQ(two, _THRUST_STD::min(three, two, thrust::less())); + ASSERT_EQ(two, _THRUST_STD::min(three, two, _THRUST_STD::less())); - ASSERT_EQ(three, _THRUST_STD::min(two, three, thrust::greater())); - ASSERT_EQ(three, _THRUST_STD::min(three, two, thrust::greater())); + ASSERT_EQ(three, _THRUST_STD::min(two, three, _THRUST_STD::greater())); + ASSERT_EQ(three, _THRUST_STD::min(three, two, _THRUST_STD::greater())); using KV = key_value; KV two_and_two(two, two); @@ -63,11 +63,11 @@ TYPED_TEST(MinAndMaxTests, TestMin) ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::min(two_and_two, two_and_three)); ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::min(two_and_three, two_and_two)); - ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::min(two_and_two, two_and_three, thrust::less())); - ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::min(two_and_three, two_and_two, thrust::less())); + ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::min(two_and_two, two_and_three, _THRUST_STD::less())); + ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::min(two_and_three, two_and_two, _THRUST_STD::less())); - ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::min(two_and_two, two_and_three, thrust::greater())); - ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::min(two_and_three, two_and_two, thrust::greater())); + ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::min(two_and_two, two_and_three, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::min(two_and_three, two_and_two, _THRUST_STD::greater())); } TYPED_TEST(MinAndMaxTests, TestMax) @@ -79,13 +79,13 @@ TYPED_TEST(MinAndMaxTests, TestMax) // 2 < 3 T two(2), three(3); ASSERT_EQ(three, _THRUST_STD::max(two, three)); - ASSERT_EQ(three, _THRUST_STD::max(two, three, thrust::less())); + ASSERT_EQ(three, _THRUST_STD::max(two, three, _THRUST_STD::less())); ASSERT_EQ(three, _THRUST_STD::max(three, two)); - ASSERT_EQ(three, _THRUST_STD::max(three, two, thrust::less())); + ASSERT_EQ(three, _THRUST_STD::max(three, two, _THRUST_STD::less())); - ASSERT_EQ(two, _THRUST_STD::max(two, three, thrust::greater())); - ASSERT_EQ(two, _THRUST_STD::max(three, two, thrust::greater())); + ASSERT_EQ(two, _THRUST_STD::max(two, three, _THRUST_STD::greater())); + ASSERT_EQ(two, _THRUST_STD::max(three, two, _THRUST_STD::greater())); using KV = key_value; KV two_and_two(two, two); @@ -95,9 +95,9 @@ TYPED_TEST(MinAndMaxTests, TestMax) ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::max(two_and_two, two_and_three)); ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::max(two_and_three, two_and_two)); - ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::max(two_and_two, two_and_three, thrust::less())); - ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::max(two_and_three, two_and_two, thrust::less())); + ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::max(two_and_two, two_and_three, _THRUST_STD::less())); + ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::max(two_and_three, two_and_two, _THRUST_STD::less())); - ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::max(two_and_two, two_and_three, thrust::greater())); - ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::max(two_and_three, two_and_two, thrust::greater())); + ASSERT_EQ_QUIET(two_and_two, _THRUST_STD::max(two_and_two, two_and_three, _THRUST_STD::greater())); + ASSERT_EQ_QUIET(two_and_three, _THRUST_STD::max(two_and_three, two_and_two, _THRUST_STD::greater())); } diff --git a/projects/rocthrust/test/test_min_element.cpp b/projects/rocthrust/test/test_min_element.cpp index 693e950ba085..17e4c7ec275b 100644 --- a/projects/rocthrust/test/test_min_element.cpp +++ b/projects/rocthrust/test/test_min_element.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,8 +52,8 @@ TYPED_TEST(MinElementTests, TestMinElementSimple) ASSERT_EQ(*thrust::min_element(data.begin(), data.end()), 1); ASSERT_EQ(thrust::min_element(data.begin(), data.end()) - data.begin(), 2); - ASSERT_EQ(*thrust::min_element(data.begin(), data.end(), thrust::greater()), 5); - ASSERT_EQ(thrust::min_element(data.begin(), data.end(), thrust::greater()) - data.begin(), 1); + ASSERT_EQ(*thrust::min_element(data.begin(), data.end(), _THRUST_STD::greater()), 5); + ASSERT_EQ(thrust::min_element(data.begin(), data.end(), _THRUST_STD::greater()) - data.begin(), 1); } TYPED_TEST(MinElementVectorUnitTests, TestMinElementWithTransform) @@ -65,12 +65,12 @@ TYPED_TEST(MinElementVectorUnitTests, TestMinElementWithTransform) Vector data{3, 5, 1, 2, 5, 1}; - ASSERT_EQ(*thrust::min_element(thrust::make_transform_iterator(data.begin(), thrust::negate()), - thrust::make_transform_iterator(data.end(), thrust::negate())), + ASSERT_EQ(*thrust::min_element(thrust::make_transform_iterator(data.begin(), _THRUST_STD::negate()), + thrust::make_transform_iterator(data.end(), _THRUST_STD::negate())), -5); - ASSERT_EQ(*thrust::min_element(thrust::make_transform_iterator(data.begin(), thrust::negate()), - thrust::make_transform_iterator(data.end(), thrust::negate()), - thrust::greater()), + ASSERT_EQ(*thrust::min_element(thrust::make_transform_iterator(data.begin(), _THRUST_STD::negate()), + thrust::make_transform_iterator(data.end(), _THRUST_STD::negate()), + _THRUST_STD::greater()), -1); } @@ -98,9 +98,9 @@ TYPED_TEST(MinElementPrimitiveTests, TestMinElement) ASSERT_EQ(h_min - h_data.begin(), d_min - d_data.begin()); typename thrust::host_vector::iterator h_max = - thrust::min_element(h_data.begin(), h_data.end(), thrust::greater()); + thrust::min_element(h_data.begin(), h_data.end(), _THRUST_STD::greater()); typename thrust::device_vector::iterator d_max = - thrust::min_element(d_data.begin(), d_data.end(), thrust::greater()); + thrust::min_element(d_data.begin(), d_data.end(), _THRUST_STD::greater()); ASSERT_EQ(h_max - h_data.begin(), d_max - d_data.begin()); } @@ -148,9 +148,9 @@ void TestMinElementWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(1); thrust::counting_iterator end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); - ASSERT_EQ(*thrust::min_element(thrust::device, begin, end, thrust::greater()), (1ll << magnitude)); + ASSERT_EQ(*thrust::min_element(thrust::device, begin, end, _THRUST_STD::greater()), (1ll << magnitude)); } TEST(MinElementTests, TestMinElementWithBigIndexes) diff --git a/projects/rocthrust/test/test_minimum_system.cpp b/projects/rocthrust/test/test_minimum_system.cpp new file mode 100644 index 000000000000..087a395f0433 --- /dev/null +++ b/projects/rocthrust/test/test_minimum_system.cpp @@ -0,0 +1,49 @@ +/* + * Copyright 2026 NVIDIA Corporation + * Modifications Copyright© 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. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include _THRUST_STD_INCLUDE(type_traits) + +#include "test_param_fixtures.hpp" +#include "test_utils.hpp" + +template +void check() +{ + EXPECT_TRUE((_THRUST_STD::is_convertible_v) ); + EXPECT_TRUE((!_THRUST_STD::is_convertible_v) ); + EXPECT_TRUE((_THRUST_STD::is_same_v, SeqSystem>) ); + EXPECT_TRUE((_THRUST_STD::is_same_v, SeqSystem>) ); +} + +TEST(MinimumSystemTests, MinimumSystem) +{ + using seq = decltype(thrust::seq); + using seq_tag = decltype(thrust::seq)::tag_type; + using dev = decltype(thrust::device); + using dev_tag = thrust::device_system_tag; + using host = decltype(thrust::host); + using host_tag = thrust::host_system_tag; + + check(); + check(); + check(); + check(); + check(); + check(); +} diff --git a/projects/rocthrust/test/test_minimum_type.cpp b/projects/rocthrust/test/test_minimum_type.cpp new file mode 100644 index 000000000000..10ace9256c4a --- /dev/null +++ b/projects/rocthrust/test/test_minimum_type.cpp @@ -0,0 +1,78 @@ +/* + * Copyright 2026 NVIDIA Corporation + * Modifications Copyright© 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. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include + +#include _THRUST_STD_INCLUDE(type_traits) + +#include "test_param_fixtures.hpp" +#include "test_utils.hpp" + +template +using mt = thrust::detail::minimum_type; + +template +inline constexpr bool mt_fails_impl = true; + +template +inline constexpr bool mt_fails_impl<_THRUST_STD::void_t>, Ts...> = false; + +template +inline constexpr bool mt_fails = mt_fails_impl; + +struct A +{}; +struct B : A +{}; +struct C : B +{}; +struct C2 : B +{}; + +TEST(MinimumTypeTests, MinimumType) +{ + using _THRUST_STD::is_same_v; + + EXPECT_TRUE((is_same_v, int>) ); + EXPECT_TRUE((is_same_v, int>) ); + EXPECT_TRUE((is_same_v, int>) ); + + EXPECT_TRUE((is_same_v, char>) ); + EXPECT_TRUE((is_same_v, int>) ); + + EXPECT_TRUE((is_same_v, A>) ); + EXPECT_TRUE((is_same_v, A>) ); + + EXPECT_TRUE((is_same_v, A>) ); + EXPECT_TRUE((is_same_v, A>) ); + EXPECT_TRUE((is_same_v, A>) ); + + EXPECT_TRUE( + (is_same_v< + mt<_THRUST_STD::random_access_iterator_tag, _THRUST_STD::input_iterator_tag, _THRUST_STD::forward_iterator_tag>, + _THRUST_STD::input_iterator_tag>) ); + EXPECT_TRUE((is_same_v, + _THRUST_STD::random_access_iterator_tag>) ); + + EXPECT_TRUE((mt_fails) ); + EXPECT_TRUE((mt_fails) ); + EXPECT_TRUE((mt_fails) ); + EXPECT_TRUE((mt_fails) ); +} diff --git a/projects/rocthrust/test/test_minmax_element.cpp b/projects/rocthrust/test/test_minmax_element.cpp index f2af397c3628..89e79aa91a08 100644 --- a/projects/rocthrust/test/test_minmax_element.cpp +++ b/projects/rocthrust/test/test_minmax_element.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. @@ -63,12 +63,12 @@ TYPED_TEST(MinMaxElementVectorUnitTests, TestMinMaxElementWithTransform) Vector data{3, 5, 1, 2, 5, 1}; - ASSERT_EQ(*thrust::minmax_element(thrust::make_transform_iterator(data.begin(), thrust::negate()), - thrust::make_transform_iterator(data.end(), thrust::negate())) + ASSERT_EQ(*thrust::minmax_element(thrust::make_transform_iterator(data.begin(), _THRUST_STD::negate()), + thrust::make_transform_iterator(data.end(), _THRUST_STD::negate())) .first, -5); - ASSERT_EQ(*thrust::minmax_element(thrust::make_transform_iterator(data.begin(), thrust::negate()), - thrust::make_transform_iterator(data.end(), thrust::negate())) + ASSERT_EQ(*thrust::minmax_element(thrust::make_transform_iterator(data.begin(), _THRUST_STD::negate()), + thrust::make_transform_iterator(data.end(), _THRUST_STD::negate())) .second, -1); } @@ -104,10 +104,10 @@ TYPED_TEST(MinMaxElementPrimitiveTests, TestMinMaxElement) ASSERT_EQ(h_min - h_data.begin(), d_min - d_data.begin()); ASSERT_EQ(h_max - h_data.begin(), d_max - d_data.begin()); - h_max = thrust::minmax_element(h_data.begin(), h_data.end(), thrust::greater()).first; - d_max = thrust::minmax_element(d_data.begin(), d_data.end(), thrust::greater()).first; - h_min = thrust::minmax_element(h_data.begin(), h_data.end(), thrust::greater()).second; - d_min = thrust::minmax_element(d_data.begin(), d_data.end(), thrust::greater()).second; + h_max = thrust::minmax_element(h_data.begin(), h_data.end(), _THRUST_STD::greater()).first; + d_max = thrust::minmax_element(d_data.begin(), d_data.end(), _THRUST_STD::greater()).first; + h_min = thrust::minmax_element(h_data.begin(), h_data.end(), _THRUST_STD::greater()).second; + d_min = thrust::minmax_element(d_data.begin(), d_data.end(), _THRUST_STD::greater()).second; ASSERT_EQ(h_min - h_data.begin(), d_min - d_data.begin()); ASSERT_EQ(h_max - h_data.begin(), d_max - d_data.begin()); @@ -157,13 +157,13 @@ void TestMinMaxElementWithBigIndexesHelper(int magnitude) using Iter = thrust::counting_iterator; Iter begin(1); Iter end = begin + (1ll << magnitude); - ASSERT_EQ(thrust::distance(begin, end), 1ll << magnitude); + ASSERT_EQ(_THRUST_STD::distance(begin, end), 1ll << magnitude); thrust::pair result = thrust::minmax_element(thrust::device, begin, end); ASSERT_EQ(*result.first, 1); ASSERT_EQ(*result.second, (1ll << magnitude)); - result = thrust::minmax_element(thrust::device, begin, end, thrust::greater()); + result = thrust::minmax_element(thrust::device, begin, end, _THRUST_STD::greater()); ASSERT_EQ(*result.second, 1); ASSERT_EQ(*result.first, (1ll << magnitude)); } 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