From 430a2e841957ee65494e76d3aa6257a09b1c5709 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Sun, 9 Nov 2025 03:01:04 -0600 Subject: [PATCH 01/12] fix omp scan bug --- thrust/thrust/system/omp/detail/scan.h | 54 ++++++++++++++++++++++++-- 1 file changed, 50 insertions(+), 4 deletions(-) diff --git a/thrust/thrust/system/omp/detail/scan.h b/thrust/thrust/system/omp/detail/scan.h index d4e5c9db81f7..4604e43dc0af 100644 --- a/thrust/thrust/system/omp/detail/scan.h +++ b/thrust/thrust/system/omp/detail/scan.h @@ -106,7 +106,25 @@ OutputIterator scan_impl( if (start < n) { - block_sums[tid] = ::cuda::std::reduce(first + start, first + end, accum_t{}, wrapped_binary_op); + if constexpr (has_init) + { + if (tid == 0) + { + block_sums[tid] = ::cuda::std::reduce(first + start, first + end, init, wrapped_binary_op); + } + else + { + block_sums[tid] = ::cuda::std::reduce(first + start, first + end, accum_t{}, wrapped_binary_op); + } + } + else + { + // No init value: reduce uses first element as init + if (end > start) + { + block_sums[tid] = ::cuda::std::reduce(first + start + 1, first + end, *(first + start), wrapped_binary_op); + } + } } } @@ -117,7 +135,18 @@ OutputIterator scan_impl( } else { - ::cuda::std::exclusive_scan(block_sums.begin(), block_sums.end(), block_sums.begin(), accum_t{}, wrapped_binary_op); + // For no init inclusive scan, use inclusive_scan on block_sums, then shift + if (num_threads > 1) + { + temporary_array block_sums_copy(exec, num_threads); + for (int i = 0; i < num_threads; ++i) + { + block_sums_copy[i] = block_sums[i]; + } + ::cuda::std::inclusive_scan( + block_sums_copy.begin(), block_sums_copy.begin() + num_threads - 1, block_sums.begin() + 1, wrapped_binary_op); + // block_sums[0] stays as-is (no prefix for first block) + } } // Step 3: Scan each block with offset (N reads/writes) @@ -130,13 +159,30 @@ OutputIterator scan_impl( if (start < n) { - const accum_t prefix = block_sums[tid]; if constexpr (IsInclusive) { - ::cuda::std::inclusive_scan(first + start, first + end, result + start, wrapped_binary_op, prefix); + if constexpr (has_init) + { + const accum_t prefix = block_sums[tid]; + ::cuda::std::inclusive_scan(first + start, first + end, result + start, wrapped_binary_op, prefix); + } + else + { + // For no init: thread 0 has no prefix, others use block_sums + if (tid == 0) + { + ::cuda::std::inclusive_scan(first + start, first + end, result + start, wrapped_binary_op); + } + else + { + const accum_t prefix = block_sums[tid]; + ::cuda::std::inclusive_scan(first + start, first + end, result + start, wrapped_binary_op, prefix); + } + } } else { + const accum_t prefix = block_sums[tid]; ::cuda::std::exclusive_scan(first + start, first + end, result + start, prefix, wrapped_binary_op); } } From e400373cdba85d73492baaddc5a9d1f85708774d Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 10 Nov 2025 10:39:19 +0100 Subject: [PATCH 02/12] Add more scan tests --- thrust/testing/scan.cu | 66 ++++++++++++++++++++++++++++++++++-------- 1 file changed, 54 insertions(+), 12 deletions(-) diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index 3602f6214306..568a30dce24b 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -11,6 +11,8 @@ #include #include +#include + #include template @@ -663,24 +665,40 @@ struct Int _CCCL_HOST_DEVICE explicit Int(int num) : i(num) {} - _CCCL_HOST_DEVICE Int() - : i{} - {} - _CCCL_HOST_DEVICE Int operator+(Int const& o) const + + // TODO(bgruber): I think we should not need this constructors, but CUB fails to compile without it + Int() = default; + + _CCCL_HOST_DEVICE friend Int operator+(Int const& a, Int const& b) { - return Int{this->i + o.i}; + return Int{a.i + b.i}; } -}; -void TestInclusiveScanWithUserDefinedType() -{ - thrust::device_vector vec(5, Int{1}); + _CCCL_HOST_DEVICE friend bool operator==(Int const& a, Int const& b) + { + return a.i == b.i; + } - thrust::inclusive_scan(thrust::device, vec.cbegin(), vec.cend(), vec.begin()); + friend std::ostream& operator<<(std::ostream& os, Int const& a) + { + return os << "Int{" << a.i << "}"; + } +}; - ASSERT_EQUAL(static_cast(vec.back()).i, 5); +void TestScanWithUserDefinedTypeAndInit() +{ + { + thrust::device_vector vec(5, Int{1}); + thrust::exclusive_scan(thrust::device, vec.cbegin(), vec.cend(), vec.begin(), Int{100}, ::cuda::std::plus()); + ASSERT_EQUAL(vec, (thrust::device_vector{Int{100}, Int{101}, Int{102}, Int{103}, Int{104}})); + } + { + thrust::device_vector vec(5, Int{1}); + thrust::inclusive_scan(thrust::device, vec.cbegin(), vec.cend(), vec.begin(), Int{100}, ::cuda::std::plus()); + ASSERT_EQUAL(vec, (thrust::device_vector{Int{101}, Int{102}, Int{103}, Int{104}, Int{105}})); + } } -DECLARE_UNITTEST(TestInclusiveScanWithUserDefinedType); +DECLARE_UNITTEST(TestScanWithUserDefinedTypeAndInit); // Represents a permutation as a tuple of integers, see also: https://en.wikipedia.org/wiki/Permutation // We need a distinct type (instead of an alias) for operator<< to be found via ADL @@ -768,3 +786,27 @@ void TestInclusiveScanWithNonCommutativeOp() {4, 0, 2, 1, 3}})); } DECLARE_UNITTEST(TestInclusiveScanWithNonCommutativeOp); + +// Adapted from issue: https://github.com/NVIDIA/cccl/issues/6317 +void TestScanBug6317() +{ + using T = unsigned int; + const int n = 1729; + { + thrust::device_vector s = unittest::random_integers(n); + thrust::device_vector d(n); + const auto r = thrust::inclusive_scan(s.cbegin(), s.cend(), d.begin(), ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d.end() == r), true); + std::partial_sum(s.cbegin(), s.cend(), s.begin(), std::multiplies<>{}); + ASSERT_EQUAL(s, d); + } + { + thrust::device_vector s = unittest::random_integers(n); + thrust::device_vector d(n); + const auto r = thrust::exclusive_scan(s.cbegin(), s.cend(), d.begin(), 42, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d.end() == r), true); + std::exclusive_scan(s.cbegin(), s.cend(), s.begin(), 42, std::multiplies<>{}); + ASSERT_EQUAL(s, d); + } +} +DECLARE_UNITTEST(TestScanBug6317); From 656f02462d584615dd53a908272d75c280b0d1f8 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Wed, 12 Nov 2025 03:08:07 -0600 Subject: [PATCH 03/12] add scan tests and few changes --- thrust/testing/scan.cu | 131 +++++++++++++++++++++ thrust/testing/transform_scan.cu | 151 +++++++++++++++++++++++++ thrust/thrust/system/omp/detail/scan.h | 34 +----- 3 files changed, 288 insertions(+), 28 deletions(-) diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index 568a30dce24b..539b61af720d 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -810,3 +810,134 @@ void TestScanBug6317() } } DECLARE_UNITTEST(TestScanBug6317); + +// Test edge cases for parallel scan with non-additive operations +void TestScanEdgeCases() +{ + // Test 1: Large array with inclusive_scan, multiplies, and init + // Tests Step 1 fix with many threads + { + const int n = 10000; + thrust::device_vector d_input(n); + for (int i = 0; i < n; ++i) + { + d_input[i] = (i % 5) + 1; + } + + thrust::device_vector d_output(n); + auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 2, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + // Verify with host scan + thrust::host_vector h_input = d_input; + thrust::host_vector h_output(n); + std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 2); + + ASSERT_EQUAL(d_output, h_output); + } + + // Test 2: Boundary exactly at parallel_scan_threshold (1024) + { + const int n = 1024; + thrust::device_vector d_input(n); + for (int i = 0; i < n; ++i) + { + d_input[i] = i + 1; + } + + thrust::device_vector d_output(n); + auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 10, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::host_vector h_input = d_input; + thrust::host_vector h_output(n); + std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 10); + + ASSERT_EQUAL(d_output, h_output); + } + + // Test 3: Below threshold (1023) should use serial path + { + const int n = 1023; + thrust::device_vector d_input(n); + for (int i = 0; i < n; ++i) + { + d_input[i] = (i % 3) + 1; + } + + thrust::device_vector d_output(n); + auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 5, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::host_vector h_input = d_input; + thrust::host_vector h_output(n); + std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 5); + + ASSERT_EQUAL(d_output, h_output); + } + + // Test 4: Very small array (edge case for block distribution) + { + thrust::device_vector d_input = {3, 7}; + thrust::device_vector d_output(2); + + auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 2, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::device_vector expected = {6, 42}; + ASSERT_EQUAL(d_output, expected); + } + + // Test 5: exclusive_scan with large array and multiplies + { + const int n = 10000; + thrust::device_vector d_input(n); + for (int i = 0; i < n; ++i) + { + d_input[i] = (i % 3) + 1; + } + + thrust::device_vector d_output(n); + auto r = thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 5, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::host_vector h_input = d_input; + thrust::host_vector h_output(n); + std::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 5, ::cuda::std::multiplies<>{}); + + ASSERT_EQUAL(d_output, h_output); + } + + // Test 6: exclusive_scan at boundary (1024 elements) + { + const int n = 1024; + thrust::device_vector d_input(n); + for (int i = 0; i < n; ++i) + { + d_input[i] = (i % 5) + 1; + } + + thrust::device_vector d_output(n); + auto r = thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 3, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::host_vector h_input = d_input; + thrust::host_vector h_output(n); + std::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 3, ::cuda::std::multiplies<>{}); + + ASSERT_EQUAL(d_output, h_output); + } + + // Test 7: exclusive_scan with very small array + { + thrust::device_vector d_input = {2, 4}; + thrust::device_vector d_output(2); + + auto r = thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 3, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::device_vector expected = {3, 6}; + ASSERT_EQUAL(d_output, expected); + } +} +DECLARE_UNITTEST(TestScanEdgeCases); diff --git a/thrust/testing/transform_scan.cu b/thrust/testing/transform_scan.cu index c70cc5acfed2..2236fade5256 100644 --- a/thrust/testing/transform_scan.cu +++ b/thrust/testing/transform_scan.cu @@ -5,6 +5,9 @@ #include #include +#include +#include + #include template @@ -400,3 +403,151 @@ void TestValueCategoryDeduction() ASSERT_EQUAL((thrust::device_vector{0, 5, 5, 5, 8, 8, 8, 8, 8, 8}), vec); } DECLARE_GENERIC_UNITTEST(TestValueCategoryDeduction); + +// User-defined Int type with poison default constructor +struct TransformInt +{ + int value; + + _CCCL_HOST_DEVICE TransformInt() + : value(999) + {} + _CCCL_HOST_DEVICE TransformInt(int v) + : value(v) + {} + + _CCCL_HOST_DEVICE friend TransformInt operator+(TransformInt a, TransformInt b) + { + return TransformInt(a.value + b.value); + } + _CCCL_HOST_DEVICE friend TransformInt operator*(TransformInt a, TransformInt b) + { + return TransformInt(a.value * b.value); + } + _CCCL_HOST_DEVICE friend bool operator==(TransformInt a, TransformInt b) + { + return a.value == b.value; + } + friend std::ostream& operator<<(std::ostream& os, TransformInt i) + { + return os << i.value; + } +}; + +// Test edge cases for parallel transform_scan with non-additive operations +void TestTransformScanEdgeCases() +{ + // Test 1: Large array with transform_inclusive_scan, multiplies, and init + { + const int n = 10000; + thrust::device_vector d_input(n); + for (int i = 0; i < n; ++i) + { + d_input[i] = (i % 5) + 1; + } + + thrust::device_vector d_output(n); + auto r = thrust::transform_inclusive_scan( + d_input.begin(), d_input.end(), d_output.begin(), ::cuda::std::negate(), 2, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + // Verify with host scan + thrust::host_vector h_input = d_input; + thrust::host_vector h_output(n); + std::transform(h_input.begin(), h_input.end(), h_input.begin(), [](int x) { + return -x; + }); + std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 2); + + ASSERT_EQUAL(d_output, h_output); + } + + // Test 2: transform_exclusive_scan with multiplies on large array + { + const int n = 10000; + thrust::device_vector d_input(n); + for (int i = 0; i < n; ++i) + { + d_input[i] = (i % 3) + 1; + } + + thrust::device_vector d_output(n); + auto r = thrust::transform_exclusive_scan( + d_input.begin(), d_input.end(), d_output.begin(), ::cuda::std::negate(), 5, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::host_vector h_input = d_input; + thrust::host_vector h_output(n); + std::transform(h_input.begin(), h_input.end(), h_input.begin(), [](int x) { + return -x; + }); + std::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 5, ::cuda::std::multiplies<>{}); + + ASSERT_EQUAL(d_output, h_output); + } + + // Test 3: User-defined type with poison default constructor + { + thrust::device_vector vec(5, TransformInt{1}); + thrust::transform_exclusive_scan( + vec.begin(), + vec.end(), + vec.begin(), + ::cuda::std::identity{}, + TransformInt{100}, + ::cuda::std::plus()); + + ASSERT_EQUAL(vec, + (thrust::device_vector{ + TransformInt{100}, TransformInt{101}, TransformInt{102}, TransformInt{103}, TransformInt{104}})); + } + + // Test 4: In-place transform_exclusive_scan with multiplies + { + thrust::device_vector vec = {2, 3, 4, 5}; + auto r = thrust::transform_exclusive_scan( + vec.begin(), vec.end(), vec.begin(), ::cuda::std::negate(), 10, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((vec.end() == r), true); + + thrust::device_vector expected = {10, -20, 60, -240}; + ASSERT_EQUAL(vec, expected); + } + + // Test 5: Boundary case at threshold (1024 elements) + { + const int n = 1024; + thrust::device_vector d_input(n); + for (int i = 0; i < n; ++i) + { + d_input[i] = i + 1; + } + + thrust::device_vector d_output(n); + auto r = thrust::transform_inclusive_scan( + d_input.begin(), d_input.end(), d_output.begin(), thrust::square(), 1, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::host_vector h_input = d_input; + thrust::host_vector h_output(n); + std::transform(h_input.begin(), h_input.end(), h_input.begin(), [](int x) { + return x * x; + }); + std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 1); + + ASSERT_EQUAL(d_output, h_output); + } + + // Test 6: Very small array (2 elements) + { + thrust::device_vector d_input = {3, 7}; + thrust::device_vector d_output(2); + + auto r = thrust::transform_inclusive_scan( + d_input.begin(), d_input.end(), d_output.begin(), ::cuda::std::negate(), 2, ::cuda::std::multiplies<>{}); + ASSERT_EQUAL((d_output.end() == r), true); + + thrust::device_vector expected = {-6, 42}; + ASSERT_EQUAL(d_output, expected); + } +} +DECLARE_UNITTEST(TestTransformScanEdgeCases); diff --git a/thrust/thrust/system/omp/detail/scan.h b/thrust/thrust/system/omp/detail/scan.h index 4604e43dc0af..29ec5b9719f5 100644 --- a/thrust/thrust/system/omp/detail/scan.h +++ b/thrust/thrust/system/omp/detail/scan.h @@ -106,25 +106,9 @@ OutputIterator scan_impl( if (start < n) { - if constexpr (has_init) - { - if (tid == 0) - { - block_sums[tid] = ::cuda::std::reduce(first + start, first + end, init, wrapped_binary_op); - } - else - { - block_sums[tid] = ::cuda::std::reduce(first + start, first + end, accum_t{}, wrapped_binary_op); - } - } - else - { - // No init value: reduce uses first element as init - if (end > start) - { - block_sums[tid] = ::cuda::std::reduce(first + start + 1, first + end, *(first + start), wrapped_binary_op); - } - } + // For both has_init and no-init cases: reduce each block using first element as init + accum_t first_elem = *(first + start); + block_sums[tid] = ::cuda::std::reduce(first + start + 1, first + end, first_elem, wrapped_binary_op); } } @@ -135,17 +119,11 @@ OutputIterator scan_impl( } else { - // For no init inclusive scan, use inclusive_scan on block_sums, then shift + // For no init inclusive scan: exclusive_scan starting at second element with first element as init if (num_threads > 1) { - temporary_array block_sums_copy(exec, num_threads); - for (int i = 0; i < num_threads; ++i) - { - block_sums_copy[i] = block_sums[i]; - } - ::cuda::std::inclusive_scan( - block_sums_copy.begin(), block_sums_copy.begin() + num_threads - 1, block_sums.begin() + 1, wrapped_binary_op); - // block_sums[0] stays as-is (no prefix for first block) + ::cuda::std::exclusive_scan( + block_sums.begin() + 1, block_sums.end(), block_sums.begin() + 1, block_sums[0], wrapped_binary_op); } } From f549d50ea62a4fcb3cc77a5a480b57f09c03778c Mon Sep 17 00:00:00 2001 From: charan-003 Date: Wed, 12 Nov 2025 12:56:00 -0600 Subject: [PATCH 04/12] use value type for init accum_t --- thrust/thrust/system/omp/detail/scan.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/thrust/thrust/system/omp/detail/scan.h b/thrust/thrust/system/omp/detail/scan.h index 29ec5b9719f5..45939f294f7a 100644 --- a/thrust/thrust/system/omp/detail/scan.h +++ b/thrust/thrust/system/omp/detail/scan.h @@ -122,8 +122,9 @@ OutputIterator scan_impl( // For no init inclusive scan: exclusive_scan starting at second element with first element as init if (num_threads > 1) { + accum_t first_block_sum = block_sums[0]; ::cuda::std::exclusive_scan( - block_sums.begin() + 1, block_sums.end(), block_sums.begin() + 1, block_sums[0], wrapped_binary_op); + block_sums.begin() + 1, block_sums.end(), block_sums.begin() + 1, first_block_sum, wrapped_binary_op); } } From 68f500c3ad287284ad4d6c7c0c1b6071d52cbd01 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Wed, 12 Nov 2025 15:20:49 -0600 Subject: [PATCH 05/12] use thrust::inclusive/exclusive --- thrust/testing/scan.cu | 10 +++++----- thrust/testing/transform_scan.cu | 18 ++++++------------ 2 files changed, 11 insertions(+), 17 deletions(-) diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index 539b61af720d..dcba96a4f221 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -831,7 +831,7 @@ void TestScanEdgeCases() // Verify with host scan thrust::host_vector h_input = d_input; thrust::host_vector h_output(n); - std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 2); + thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 2, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(d_output, h_output); } @@ -851,7 +851,7 @@ void TestScanEdgeCases() thrust::host_vector h_input = d_input; thrust::host_vector h_output(n); - std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 10); + thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 10, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(d_output, h_output); } @@ -871,7 +871,7 @@ void TestScanEdgeCases() thrust::host_vector h_input = d_input; thrust::host_vector h_output(n); - std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 5); + thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 5, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(d_output, h_output); } @@ -903,7 +903,7 @@ void TestScanEdgeCases() thrust::host_vector h_input = d_input; thrust::host_vector h_output(n); - std::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 5, ::cuda::std::multiplies<>{}); + thrust::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 5, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(d_output, h_output); } @@ -923,7 +923,7 @@ void TestScanEdgeCases() thrust::host_vector h_input = d_input; thrust::host_vector h_output(n); - std::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 3, ::cuda::std::multiplies<>{}); + thrust::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 3, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(d_output, h_output); } diff --git a/thrust/testing/transform_scan.cu b/thrust/testing/transform_scan.cu index 2236fade5256..77b425f151e6 100644 --- a/thrust/testing/transform_scan.cu +++ b/thrust/testing/transform_scan.cu @@ -454,10 +454,8 @@ void TestTransformScanEdgeCases() // Verify with host scan thrust::host_vector h_input = d_input; thrust::host_vector h_output(n); - std::transform(h_input.begin(), h_input.end(), h_input.begin(), [](int x) { - return -x; - }); - std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 2); + thrust::transform_inclusive_scan( + h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::negate(), 2, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(d_output, h_output); } @@ -478,10 +476,8 @@ void TestTransformScanEdgeCases() thrust::host_vector h_input = d_input; thrust::host_vector h_output(n); - std::transform(h_input.begin(), h_input.end(), h_input.begin(), [](int x) { - return -x; - }); - std::exclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 5, ::cuda::std::multiplies<>{}); + thrust::transform_exclusive_scan( + h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::negate(), 5, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(d_output, h_output); } @@ -529,10 +525,8 @@ void TestTransformScanEdgeCases() thrust::host_vector h_input = d_input; thrust::host_vector h_output(n); - std::transform(h_input.begin(), h_input.end(), h_input.begin(), [](int x) { - return x * x; - }); - std::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::multiplies<>{}, 1); + thrust::transform_inclusive_scan( + h_input.begin(), h_input.end(), h_output.begin(), thrust::square(), 1, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(d_output, h_output); } From 9ea6902bd1165c4fb011b985a8bcebe17c6e7531 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Thu, 13 Nov 2025 08:07:08 -0600 Subject: [PATCH 06/12] fix --- thrust/testing/scan.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index dcba96a4f221..92d2427403a9 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -805,7 +805,7 @@ void TestScanBug6317() thrust::device_vector d(n); const auto r = thrust::exclusive_scan(s.cbegin(), s.cend(), d.begin(), 42, ::cuda::std::multiplies<>{}); ASSERT_EQUAL((d.end() == r), true); - std::exclusive_scan(s.cbegin(), s.cend(), s.begin(), 42, std::multiplies<>{}); + thrust::exclusive_scan(s.cbegin(), s.cend(), s.begin(), 42, ::cuda::std::multiplies<>{}); ASSERT_EQUAL(s, d); } } From 47060f2bb1e82512586e712ca75fadb1f7266e30 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Tue, 18 Nov 2025 14:56:14 -0600 Subject: [PATCH 07/12] minor changes --- thrust/testing/scan.cu | 35 +++++++++++--------------- thrust/testing/transform_scan.cu | 21 +++++++--------- thrust/thrust/system/omp/detail/scan.h | 4 ++- 3 files changed, 27 insertions(+), 33 deletions(-) diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index 92d2427403a9..6990b8c65ce6 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -819,10 +819,9 @@ void TestScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - for (int i = 0; i < n; ++i) - { - d_input[i] = (i % 5) + 1; - } + thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + return (i % 5) + 1; + }); thrust::device_vector d_output(n); auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 2, ::cuda::std::multiplies<>{}); @@ -840,10 +839,9 @@ void TestScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - for (int i = 0; i < n; ++i) - { - d_input[i] = i + 1; - } + thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + return i + 1; + }); thrust::device_vector d_output(n); auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 10, ::cuda::std::multiplies<>{}); @@ -860,10 +858,9 @@ void TestScanEdgeCases() { const int n = 1023; thrust::device_vector d_input(n); - for (int i = 0; i < n; ++i) - { - d_input[i] = (i % 3) + 1; - } + thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + return (i % 3) + 1; + }); thrust::device_vector d_output(n); auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 5, ::cuda::std::multiplies<>{}); @@ -892,10 +889,9 @@ void TestScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - for (int i = 0; i < n; ++i) - { - d_input[i] = (i % 3) + 1; - } + thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + return (i % 3) + 1; + }); thrust::device_vector d_output(n); auto r = thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 5, ::cuda::std::multiplies<>{}); @@ -912,10 +908,9 @@ void TestScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - for (int i = 0; i < n; ++i) - { - d_input[i] = (i % 5) + 1; - } + thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + return (i % 5) + 1; + }); thrust::device_vector d_output(n); auto r = thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 3, ::cuda::std::multiplies<>{}); diff --git a/thrust/testing/transform_scan.cu b/thrust/testing/transform_scan.cu index 77b425f151e6..bc26bbdf12b9 100644 --- a/thrust/testing/transform_scan.cu +++ b/thrust/testing/transform_scan.cu @@ -441,10 +441,9 @@ void TestTransformScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - for (int i = 0; i < n; ++i) - { - d_input[i] = (i % 5) + 1; - } + thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + return (i % 5) + 1; + }); thrust::device_vector d_output(n); auto r = thrust::transform_inclusive_scan( @@ -464,10 +463,9 @@ void TestTransformScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - for (int i = 0; i < n; ++i) - { - d_input[i] = (i % 3) + 1; - } + thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + return (i % 3) + 1; + }); thrust::device_vector d_output(n); auto r = thrust::transform_exclusive_scan( @@ -513,10 +511,9 @@ void TestTransformScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - for (int i = 0; i < n; ++i) - { - d_input[i] = i + 1; - } + thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + return i + 1; + }); thrust::device_vector d_output(n); auto r = thrust::transform_inclusive_scan( diff --git a/thrust/thrust/system/omp/detail/scan.h b/thrust/thrust/system/omp/detail/scan.h index 45939f294f7a..c4c08a5357f5 100644 --- a/thrust/thrust/system/omp/detail/scan.h +++ b/thrust/thrust/system/omp/detail/scan.h @@ -75,7 +75,7 @@ OutputIterator scan_impl( const int num_threads = omp_get_max_threads(); // Use serial scan for small arrays where parallel overhead dominates - if (static_cast(n) < parallel_scan_threshold || num_threads <= 1) + if (static_cast(n) < ::cuda::std::max(parallel_scan_threshold, num_threads) || num_threads <= 1) { if constexpr (IsInclusive) { @@ -94,6 +94,8 @@ OutputIterator scan_impl( } } + _CCCL_ASSERT(num_threads > 1, "Parallel scan requires multiple threads"); + temporary_array block_sums(exec, num_threads); // Step 1: Reduce each block (N reads) From 660b6f2874270dbd42b416c2f7da1317a5256619 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Tue, 18 Nov 2025 17:46:35 -0600 Subject: [PATCH 08/12] fix --- thrust/thrust/system/omp/detail/scan.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/thrust/thrust/system/omp/detail/scan.h b/thrust/thrust/system/omp/detail/scan.h index c4c08a5357f5..98ade1b255c2 100644 --- a/thrust/thrust/system/omp/detail/scan.h +++ b/thrust/thrust/system/omp/detail/scan.h @@ -75,7 +75,8 @@ OutputIterator scan_impl( const int num_threads = omp_get_max_threads(); // Use serial scan for small arrays where parallel overhead dominates - if (static_cast(n) < ::cuda::std::max(parallel_scan_threshold, num_threads) || num_threads <= 1) + if (static_cast(n) < ::cuda::std::max(parallel_scan_threshold, static_cast(num_threads)) + || num_threads <= 1) { if constexpr (IsInclusive) { From 8a56f7d3bbeb219bc52ce807babf578c0764998d Mon Sep 17 00:00:00 2001 From: charan-003 Date: Wed, 19 Nov 2025 02:45:19 -0600 Subject: [PATCH 09/12] include tabulate --- thrust/testing/scan.cu | 1 + thrust/testing/transform_scan.cu | 1 + 2 files changed, 2 insertions(+) diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index 6990b8c65ce6..768d990f3a26 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -7,6 +7,7 @@ #include #include #include +#include #include #include diff --git a/thrust/testing/transform_scan.cu b/thrust/testing/transform_scan.cu index bc26bbdf12b9..f94284f5942f 100644 --- a/thrust/testing/transform_scan.cu +++ b/thrust/testing/transform_scan.cu @@ -3,6 +3,7 @@ #include #include #include +#include #include #include From a8df5c023c557ca0e49f55f20682b0507de77084 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Wed, 19 Nov 2025 11:46:07 -0600 Subject: [PATCH 10/12] minor fix --- thrust/testing/scan.cu | 10 +++++----- thrust/testing/transform_scan.cu | 6 +++--- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index 768d990f3a26..eca130d9ebe9 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -820,7 +820,7 @@ void TestScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { return (i % 5) + 1; }); @@ -840,7 +840,7 @@ void TestScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { return i + 1; }); @@ -859,7 +859,7 @@ void TestScanEdgeCases() { const int n = 1023; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { return (i % 3) + 1; }); @@ -890,7 +890,7 @@ void TestScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { return (i % 3) + 1; }); @@ -909,7 +909,7 @@ void TestScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { return (i % 5) + 1; }); diff --git a/thrust/testing/transform_scan.cu b/thrust/testing/transform_scan.cu index f94284f5942f..feb7a637e55e 100644 --- a/thrust/testing/transform_scan.cu +++ b/thrust/testing/transform_scan.cu @@ -442,7 +442,7 @@ void TestTransformScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { return (i % 5) + 1; }); @@ -464,7 +464,7 @@ void TestTransformScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { return (i % 3) + 1; }); @@ -512,7 +512,7 @@ void TestTransformScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [] __device__(int i) { + thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { return i + 1; }); From 9a7675d4284881335951810d4fd5e79661b541e2 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Wed, 19 Nov 2025 12:10:42 -0600 Subject: [PATCH 11/12] use thrust::placeholders --- thrust/testing/scan.cu | 25 ++++++++++--------------- thrust/testing/transform_scan.cu | 15 ++++++--------- 2 files changed, 16 insertions(+), 24 deletions(-) diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu index eca130d9ebe9..f692e41149d8 100644 --- a/thrust/testing/scan.cu +++ b/thrust/testing/scan.cu @@ -820,9 +820,8 @@ void TestScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { - return (i % 5) + 1; - }); + using thrust::placeholders::_1; + thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 5) + 1); thrust::device_vector d_output(n); auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 2, ::cuda::std::multiplies<>{}); @@ -840,9 +839,8 @@ void TestScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { - return i + 1; - }); + using thrust::placeholders::_1; + thrust::tabulate(d_input.begin(), d_input.end(), _1 + 1); thrust::device_vector d_output(n); auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 10, ::cuda::std::multiplies<>{}); @@ -859,9 +857,8 @@ void TestScanEdgeCases() { const int n = 1023; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { - return (i % 3) + 1; - }); + using thrust::placeholders::_1; + thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 3) + 1); thrust::device_vector d_output(n); auto r = thrust::inclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 5, ::cuda::std::multiplies<>{}); @@ -890,9 +887,8 @@ void TestScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { - return (i % 3) + 1; - }); + using thrust::placeholders::_1; + thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 3) + 1); thrust::device_vector d_output(n); auto r = thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 5, ::cuda::std::multiplies<>{}); @@ -909,9 +905,8 @@ void TestScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { - return (i % 5) + 1; - }); + using thrust::placeholders::_1; + thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 5) + 1); thrust::device_vector d_output(n); auto r = thrust::exclusive_scan(d_input.begin(), d_input.end(), d_output.begin(), 3, ::cuda::std::multiplies<>{}); diff --git a/thrust/testing/transform_scan.cu b/thrust/testing/transform_scan.cu index feb7a637e55e..0d0d1c56be8d 100644 --- a/thrust/testing/transform_scan.cu +++ b/thrust/testing/transform_scan.cu @@ -442,9 +442,8 @@ void TestTransformScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { - return (i % 5) + 1; - }); + using thrust::placeholders::_1; + thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 5) + 1); thrust::device_vector d_output(n); auto r = thrust::transform_inclusive_scan( @@ -464,9 +463,8 @@ void TestTransformScanEdgeCases() { const int n = 10000; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { - return (i % 3) + 1; - }); + using thrust::placeholders::_1; + thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 3) + 1); thrust::device_vector d_output(n); auto r = thrust::transform_exclusive_scan( @@ -512,9 +510,8 @@ void TestTransformScanEdgeCases() { const int n = 1024; thrust::device_vector d_input(n); - thrust::tabulate(d_input.begin(), d_input.end(), [](int i) { - return i + 1; - }); + using thrust::placeholders::_1; + thrust::tabulate(d_input.begin(), d_input.end(), _1 + 1); thrust::device_vector d_output(n); auto r = thrust::transform_inclusive_scan( From 0b1e611e8f12441733488c7331d173f18bc7e289 Mon Sep 17 00:00:00 2001 From: charan-003 Date: Wed, 19 Nov 2025 18:01:37 -0600 Subject: [PATCH 12/12] remove unnecessary check --- thrust/thrust/system/omp/detail/scan.h | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/thrust/thrust/system/omp/detail/scan.h b/thrust/thrust/system/omp/detail/scan.h index 98ade1b255c2..c3b9f9421416 100644 --- a/thrust/thrust/system/omp/detail/scan.h +++ b/thrust/thrust/system/omp/detail/scan.h @@ -123,12 +123,9 @@ OutputIterator scan_impl( else { // For no init inclusive scan: exclusive_scan starting at second element with first element as init - if (num_threads > 1) - { - accum_t first_block_sum = block_sums[0]; - ::cuda::std::exclusive_scan( - block_sums.begin() + 1, block_sums.end(), block_sums.begin() + 1, first_block_sum, wrapped_binary_op); - } + accum_t first_block_sum = block_sums[0]; + ::cuda::std::exclusive_scan( + block_sums.begin() + 1, block_sums.end(), block_sums.begin() + 1, first_block_sum, wrapped_binary_op); } // Step 3: Scan each block with offset (N reads/writes)