Replace internal Thrust sequential system with libcuc++#9666
Replace internal Thrust sequential system with libcuc++#9666charan-003 wants to merge 4 commits into
Conversation
charan-003
force-pushed
the
replace-internal-thrust-system
branch
from
40cf3c3 to
f007483
Compare
|
@bernhardmgruber I'm assuming we should, since libcu++ already has those algorithms |
Replaced Thrust’s sequential backend algorithm implementations with delegations to libcu++/ Key updates:
WalkthroughSequential Thrust algorithm implementations (find_if, binary_search, extrema, for_each, copy_backward, adjacent_difference, reduce, scan, remove, unique, partition, merge, set_operations) are rewritten to delegate to corresponding ::cuda::std algorithms, wrapping predicates/operators via wrapped_function. A new Catch2 test file validates this behavior. ChangesSequential algorithms delegated to libcu++
Assessment against linked issues
Suggested reviewers: Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 7
🧹 Nitpick comments (2)
thrust/thrust/system/detail/sequential/merge.h (1)
19-19: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion:
thrust/detail/copy.hlooks unused now that the merge loop delegates to::cuda::std::mergeinstead of callingthrust::copyfor trailing elements.merge_by_keynever used it either (hand-written loop). Confirm no other symbol in this file needs it, then drop the include.diff
-#include <thrust/detail/copy.h> `#include` <thrust/detail/function.h>As per coding guidelines, "Remove unneeded headers."
Also applies to: 44-46
Source: Coding guidelines
thrust/thrust/system/detail/sequential/for_each.h (1)
33-36: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: make both wrapped functor locals
const.Neither
wrapped_fis reassigned, and these headers require unmodified locals to be declaredconst. As per coding guidelines, "All variables that are not modified must be declaredconst."Also applies to: 43-45
Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: d33296bf-e63e-4157-95ae-74a03aea673b
📒 Files selected for processing (14)
thrust/testing/catch2_test_wrapped_function.cuthrust/thrust/system/detail/sequential/adjacent_difference.hthrust/thrust/system/detail/sequential/binary_search.hthrust/thrust/system/detail/sequential/copy_backward.hthrust/thrust/system/detail/sequential/extrema.hthrust/thrust/system/detail/sequential/find.hthrust/thrust/system/detail/sequential/for_each.hthrust/thrust/system/detail/sequential/merge.hthrust/thrust/system/detail/sequential/partition.hthrust/thrust/system/detail/sequential/reduce.hthrust/thrust/system/detail/sequential/remove.hthrust/thrust/system/detail/sequential/scan.hthrust/thrust/system/detail/sequential/set_operations.hthrust/thrust/system/detail/sequential/unique.h
| // Test that sequential algorithms correctly handle predicates | ||
| // that require implicit type conversions (e.g., float → const double&). | ||
| // This validates that wrapped_function properly handles the conversion chain | ||
| // when user-provided callables expect a different type than the iterator's value_type. | ||
| // When the sequential backend is invoked from device code (CDP) with device_vector | ||
| // iterators, wrapped_function also unwraps proxy references (device_reference<T> → T&) | ||
| // before the implicit conversion occurs. |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy lift
important: These tests do not exercise proxy references. Every case uses thrust::host_vector<float>, so dereference stays float& and the suite only covers implicit float -> const double& conversion. A regression in wrapped_function/raw_reference_cast for device_reference or another proxy iterator would still pass here, despite the file comments and test names claiming proxy-reference coverage. Add at least one real proxy-reference case, or narrow the comments and names to conversion-only coverage. As per path instructions, thrust/**/* review should focus on iterator/system interactions and regressions. Based on PR objectives, this suite is supposed to validate proxy-reference unwrapping.
Also applies to: 93-379
Source: Path instructions
| // Test that sequential algorithms correctly handle predicates | ||
| // that require implicit type conversions (e.g., float → const double&). | ||
| // This validates that wrapped_function properly handles the conversion chain | ||
| // when user-provided callables expect a different type than the iterator's value_type. | ||
| // When the sequential backend is invoked from device code (CDP) with device_vector | ||
| // iterators, wrapped_function also unwraps proxy references (device_reference<T> → T&) | ||
| // before the implicit conversion occurs. | ||
|
|
||
| #include <thrust/adjacent_difference.h> | ||
| #include <thrust/binary_search.h> | ||
| #include <thrust/count.h> | ||
| #include <thrust/extrema.h> | ||
| #include <thrust/find.h> | ||
| #include <thrust/for_each.h> | ||
| #include <thrust/host_vector.h> | ||
| #include <thrust/logical.h> | ||
| #include <thrust/merge.h> | ||
| #include <thrust/partition.h> | ||
| #include <thrust/reduce.h> | ||
| #include <thrust/remove.h> | ||
| #include <thrust/scan.h> | ||
| #include <thrust/set_operations.h> | ||
| #include <thrust/unique.h> | ||
|
|
||
| #include "catch2_test_helper.h" |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
important: Add the standard SPDX/license header to this new .cu file. CCCL requires the correct license text on source files, and this file currently starts with test comments instead. As per coding guidelines, "All headers, and also source files, must have the correct license text."
Source: Coding guidelines
| TEST_CASE("SequentialInclusiveScanProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> input{1.0f, 2.0f, 3.0f, 4.0f}; | ||
| thrust::host_vector<float> output(4); | ||
|
|
||
| thrust::inclusive_scan(thrust::seq, input.begin(), input.end(), output.begin(), double_plus{}); | ||
| CHECK(output[0] == 1.0f); | ||
| CHECK(output[1] == 3.0f); | ||
| CHECK(output[2] == 6.0f); | ||
| CHECK(output[3] == 10.0f); | ||
| } | ||
|
|
||
| TEST_CASE("SequentialExclusiveScanProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> input{1.0f, 2.0f, 3.0f, 4.0f}; | ||
| thrust::host_vector<float> output(4); | ||
|
|
||
| thrust::exclusive_scan(thrust::seq, input.begin(), input.end(), output.begin(), 0.0f, double_plus{}); | ||
| CHECK(output[0] == 0.0f); | ||
| CHECK(output[1] == 1.0f); | ||
| CHECK(output[2] == 3.0f); | ||
| CHECK(output[3] == 6.0f); | ||
| } | ||
|
|
||
| TEST_CASE("SequentialStablePartitionCopyProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> input{1.0f, 2.0f, 3.0f, 4.0f, 5.0f}; | ||
| thrust::host_vector<float> out_true(5); | ||
| thrust::host_vector<float> out_false(5); | ||
|
|
||
| auto result = thrust::stable_partition_copy( | ||
| thrust::seq, input.begin(), input.end(), out_true.begin(), out_false.begin(), double_greater_than_two{}); | ||
| auto n_true = result.first - out_true.begin(); | ||
| auto n_false = result.second - out_false.begin(); | ||
| CHECK(n_true == 3); | ||
| CHECK(n_false == 2); | ||
| CHECK(out_true[0] == 3.0f); | ||
| CHECK(out_true[1] == 4.0f); | ||
| CHECK(out_true[2] == 5.0f); | ||
| CHECK(out_false[0] == 1.0f); | ||
| CHECK(out_false[1] == 2.0f); | ||
| } | ||
|
|
||
| TEST_CASE("SequentialPartitionProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> vec{1.0f, 3.0f, 2.0f, 5.0f, 4.0f}; | ||
|
|
||
| auto mid = thrust::partition(thrust::seq, vec.begin(), vec.end(), double_greater_than_two{}); | ||
| CHECK(mid - vec.begin() == 3); | ||
| // all elements in [begin, mid) satisfy pred; none in [mid, end) do | ||
| for (auto it = vec.begin(); it != mid; ++it) | ||
| { | ||
| CHECK(*it > 2.0f); | ||
| } | ||
| for (auto it = mid; it != vec.end(); ++it) | ||
| { | ||
| CHECK(*it <= 2.0f); | ||
| } | ||
| } | ||
|
|
||
| TEST_CASE("SequentialInclusiveScanWithInitProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> input{1.0f, 2.0f, 3.0f, 4.0f}; | ||
| thrust::host_vector<float> output(4); | ||
|
|
||
| thrust::inclusive_scan(thrust::seq, input.begin(), input.end(), output.begin(), 10.0f, double_plus{}); | ||
| CHECK(output[0] == 11.0f); // 10 + 1 | ||
| CHECK(output[1] == 13.0f); // 11 + 2 | ||
| CHECK(output[2] == 16.0f); // 13 + 3 | ||
| CHECK(output[3] == 20.0f); // 16 + 4 | ||
| } | ||
|
|
||
| TEST_CASE("SequentialCountIfProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> vec{1.0f, 2.0f, 3.0f, 4.0f, 5.0f}; | ||
|
|
||
| auto n = thrust::count_if(thrust::seq, vec.begin(), vec.end(), double_greater_than_two{}); | ||
| CHECK(n == 3); | ||
| } | ||
|
|
||
| TEST_CASE("SequentialAllOfProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> vec{3.0f, 4.0f, 5.0f}; | ||
|
|
||
| CHECK(thrust::all_of(thrust::seq, vec.begin(), vec.end(), double_greater_than_two{})); | ||
| } | ||
|
|
||
| TEST_CASE("SequentialAnyOfProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> vec{1.0f, 2.0f, 3.0f}; | ||
|
|
||
| CHECK(thrust::any_of(thrust::seq, vec.begin(), vec.end(), double_greater_than_two{})); | ||
| } | ||
|
|
||
| TEST_CASE("SequentialNoneOfProxyReference", "[sequential][proxy_reference]") | ||
| { | ||
| thrust::host_vector<float> vec{1.0f, 2.0f}; | ||
|
|
||
| CHECK(thrust::none_of(thrust::seq, vec.begin(), vec.end(), double_greater_than_two{})); | ||
| } |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
suggestion: Mark immutable locals const throughout this file. result, new_end, mid, n, n_true, n_false, and similar locals are never reassigned, so the new tests miss the repo's immutable-local rule. As per coding guidelines, "All variables that are not modified must be declared const, including cast results, function return values, and loop-invariant computations."
Source: Coding guidelines
| // wrap binary_op to handle proxy references | ||
| thrust::detail::wrapped_function<BinaryFunction, InputType> wrapped_op{binary_op}; | ||
| return ::cuda::std::adjacent_difference(first, last, result, wrapped_op); |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Inspect the changed file and the wrapped_function definition.
git ls-files thrust/thrust/system/detail/sequential/adjacent_difference.h thrust/thrust/detail/function.h
echo "---- adjacent_difference.h ----"
nl -ba thrust/thrust/system/detail/sequential/adjacent_difference.h | sed -n '1,120p'
echo "---- function.h (wrapped_function) ----"
rg -n "wrapped_function|operator\\(" thrust/thrust/detail/function.h
nl -ba thrust/thrust/detail/function.h | sed -n '1,220p'Repository: NVIDIA/cccl
Length of output: 314
🏁 Script executed:
#!/bin/bash
set -euo pipefail
echo "---- adjacent_difference.h ----"
awk 'NR>=1 && NR<=120 { printf "%4d %s\n", NR, $0 }' thrust/thrust/system/detail/sequential/adjacent_difference.h
echo "---- function.h (wrapped_function) ----"
rg -n "wrapped_function|operator\\(" thrust/thrust/detail/function.h
awk 'NR>=1 && NR<=240 { printf "%4d %s\n", NR, $0 }' thrust/thrust/detail/function.hRepository: NVIDIA/cccl
Length of output: 3318
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Inspect libcudacxx adjacent_difference semantics to compare result typing.
rg -n "adjacent_difference" libcudacxx/include/cuda/std -g '*.h' -g '*.cuh' -g '*.hpp'Repository: NVIDIA/cccl
Length of output: 1917
🏁 Script executed:
#!/bin/bash
set -euo pipefail
awk 'NR>=1 && NR<=120 { printf "%4d %s\n", NR, $0 }' libcudacxx/include/cuda/std/__numeric/adjacent_difference.hRepository: NVIDIA/cccl
Length of output: 3111
important: wrapped_function<BinaryFunction, InputType> forces every binary_op result through InputType before it reaches ::cuda::std::adjacent_difference. That changes semantics for valid heterogeneous cases like InputType=int with OutputIterator=double* and a double-returning op, where the fractional part is lost. Keep the wrapper/result type aligned with the callable/output contract instead of InputType.
Source: Path instructions
| // wrap comp | ||
| thrust::detail::wrapped_function<StrictWeakOrdering, bool> wrapped_comp{comp}; | ||
|
|
||
| using difference_type = thrust::detail::it_difference_t<ForwardIterator>; | ||
|
|
||
| difference_type len = ::cuda::std::distance(first, last); | ||
|
|
||
| while (len > 0) | ||
| { | ||
| difference_type half = len >> 1; | ||
| ForwardIterator middle = first; | ||
|
|
||
| ::cuda::std::advance(middle, half); | ||
|
|
||
| if (wrapped_comp(*middle, val)) | ||
| { | ||
| first = middle; | ||
| ++first; | ||
| len = len - half - 1; | ||
| } | ||
| else | ||
| { | ||
| len = half; | ||
| } | ||
| } | ||
|
|
||
| return first; | ||
| return ::cuda::std::lower_bound(first, last, val, wrapped_comp); |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
suggestion: Make the wrapped comparators const in both helpers. Neither wrapped_comp is reassigned, so the current form violates the repo rule for immutable locals. As per coding guidelines, "All variables that are not modified must be declared const, including cast results, function return values, and loop-invariant computations."
Also applies to: 53-55
Source: Coding guidelines
| // wrap comp | ||
| thrust::detail::wrapped_function<BinaryPredicate, bool> wrapped_comp{comp}; | ||
|
|
||
| ForwardIterator imin = first; | ||
|
|
||
| for (; first != last; ++first) | ||
| { | ||
| if (wrapped_comp(*first, *imin)) | ||
| { | ||
| imin = first; | ||
| } | ||
| } | ||
|
|
||
| return imin; | ||
| return ::cuda::std::min_element(first, last, wrapped_comp); |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
suggestion: Make the wrapped comparators const in these helpers. Each wrapped_comp is immutable after construction, so this misses the repo's const-for-immutable-locals rule. As per coding guidelines, "All variables that are not modified must be declared const, including cast results, function return values, and loop-invariant computations."
Also applies to: 45-47, 55-57
Source: Coding guidelines
| // wrap pred | ||
| thrust::detail::wrapped_function<Predicate, bool> wrapped_pred{pred}; | ||
|
|
||
| while (first != last) | ||
| { | ||
| if (wrapped_pred(*first)) | ||
| { | ||
| return first; | ||
| } | ||
|
|
||
| ++first; | ||
| } | ||
|
|
||
| // return first so zip_iterator works correctly | ||
| return first; | ||
| return ::cuda::std::find_if(first, last, wrapped_pred); |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
suggestion: Make wrapped_pred const. It is never mutated in this helper, and the repo style requires immutable locals to be declared const. As per coding guidelines, "All variables that are not modified must be declared const, including cast results, function return values, and loop-invariant computations."
Source: Coding guidelines
1 participant
Replace Thrust sequential backend algorithms with libcu++ equivalents using wrapped_function for proxy reference handling.
Add test to validate implicit type conversions and proxy reference unwrapping.
Closes #5157