Skip to content

[Backport branch/3.2.x] fix omp scan bug #6704

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
elstehle merged 2 commits into from
Nov 20, 2025
Merged

[Backport branch/3.2.x] fix omp scan bug #6704

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
86c41db
fix omp scan bug (#6560)
charan-003 Nov 20, 2025
664098b
Merge branch 'branch/3.2.x' into backport-6560-to-branch/3.2.x
elstehle Nov 20, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
188 changes: 176 additions & 12 deletions thrust/testing/scan.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,10 +7,13 @@
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/retag.h>
#include <thrust/scan.h>
#include <thrust/tabulate.h>

#include <cuda/functional>
#include <cuda/std/array>

#include <numeric>

#include <unittest/unittest.h>

template <class Vector>
Expand Down Expand Up @@ -663,24 +666,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<Int> 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<Int>(vec.back()).i, 5);
void TestScanWithUserDefinedTypeAndInit()
{
{
thrust::device_vector<Int> vec(5, Int{1});
thrust::exclusive_scan(thrust::device, vec.cbegin(), vec.cend(), vec.begin(), Int{100}, ::cuda::std::plus<Int>());
ASSERT_EQUAL(vec, (thrust::device_vector<Int>{Int{100}, Int{101}, Int{102}, Int{103}, Int{104}}));
}
{
thrust::device_vector<Int> vec(5, Int{1});
thrust::inclusive_scan(thrust::device, vec.cbegin(), vec.cend(), vec.begin(), Int{100}, ::cuda::std::plus<Int>());
ASSERT_EQUAL(vec, (thrust::device_vector<Int>{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
Expand Down Expand Up @@ -768,3 +787,148 @@ 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<T> s = unittest::random_integers<T>(n);
thrust::device_vector<T> 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<T> s = unittest::random_integers<T>(n);
thrust::device_vector<T> d(n);
const auto r = thrust::exclusive_scan(s.cbegin(), s.cend(), d.begin(), 42, ::cuda::std::multiplies<>{});
ASSERT_EQUAL((d.end() == r), true);
thrust::exclusive_scan(s.cbegin(), s.cend(), s.begin(), 42, ::cuda::std::multiplies<>{});
ASSERT_EQUAL(s, d);
}
}
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<int> d_input(n);
using thrust::placeholders::_1;
thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 5) + 1);

thrust::device_vector<int> 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<int> h_input = d_input;
thrust::host_vector<int> h_output(n);
thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 2, ::cuda::std::multiplies<>{});

ASSERT_EQUAL(d_output, h_output);
}

// Test 2: Boundary exactly at parallel_scan_threshold (1024)
{
const int n = 1024;
thrust::device_vector<int> d_input(n);
using thrust::placeholders::_1;
thrust::tabulate(d_input.begin(), d_input.end(), _1 + 1);

thrust::device_vector<int> 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<int> h_input = d_input;
thrust::host_vector<int> h_output(n);
thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 10, ::cuda::std::multiplies<>{});

ASSERT_EQUAL(d_output, h_output);
}

// Test 3: Below threshold (1023) should use serial path
{
const int n = 1023;
thrust::device_vector<int> d_input(n);
using thrust::placeholders::_1;
thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 3) + 1);

thrust::device_vector<int> 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<int> h_input = d_input;
thrust::host_vector<int> h_output(n);
thrust::inclusive_scan(h_input.begin(), h_input.end(), h_output.begin(), 5, ::cuda::std::multiplies<>{});

ASSERT_EQUAL(d_output, h_output);
}

// Test 4: Very small array (edge case for block distribution)
{
thrust::device_vector<int> d_input = {3, 7};
thrust::device_vector<int> 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<int> expected = {6, 42};
ASSERT_EQUAL(d_output, expected);
}

// Test 5: exclusive_scan with large array and multiplies
{
const int n = 10000;
thrust::device_vector<int> d_input(n);
using thrust::placeholders::_1;
thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 3) + 1);

thrust::device_vector<int> 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<int> h_input = d_input;
thrust::host_vector<int> h_output(n);
thrust::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<int> d_input(n);
using thrust::placeholders::_1;
thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 5) + 1);

thrust::device_vector<int> 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<int> h_input = d_input;
thrust::host_vector<int> h_output(n);
thrust::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<int> d_input = {2, 4};
thrust::device_vector<int> 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<int> expected = {3, 6};
ASSERT_EQUAL(d_output, expected);
}
}
DECLARE_UNITTEST(TestScanEdgeCases);
140 changes: 140 additions & 0 deletions thrust/testing/transform_scan.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,8 +3,12 @@
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/iterator/retag.h>
#include <thrust/tabulate.h>
#include <thrust/transform_scan.h>

#include <algorithm>
#include <numeric>

#include <unittest/unittest.h>

template <typename InputIterator, typename OutputIterator, typename UnaryFunction, typename AssociativeOperator>
Expand Down Expand Up @@ -400,3 +404,139 @@ void TestValueCategoryDeduction()
ASSERT_EQUAL((thrust::device_vector<T>{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<int> d_input(n);
using thrust::placeholders::_1;
thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 5) + 1);

thrust::device_vector<int> d_output(n);
auto r = thrust::transform_inclusive_scan(
d_input.begin(), d_input.end(), d_output.begin(), ::cuda::std::negate<int>(), 2, ::cuda::std::multiplies<>{});
ASSERT_EQUAL((d_output.end() == r), true);

// Verify with host scan
thrust::host_vector<int> h_input = d_input;
thrust::host_vector<int> h_output(n);
thrust::transform_inclusive_scan(
h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::negate<int>(), 2, ::cuda::std::multiplies<>{});

ASSERT_EQUAL(d_output, h_output);
}

// Test 2: transform_exclusive_scan with multiplies on large array
{
const int n = 10000;
thrust::device_vector<int> d_input(n);
using thrust::placeholders::_1;
thrust::tabulate(d_input.begin(), d_input.end(), (_1 % 3) + 1);

thrust::device_vector<int> d_output(n);
auto r = thrust::transform_exclusive_scan(
d_input.begin(), d_input.end(), d_output.begin(), ::cuda::std::negate<int>(), 5, ::cuda::std::multiplies<>{});
ASSERT_EQUAL((d_output.end() == r), true);

thrust::host_vector<int> h_input = d_input;
thrust::host_vector<int> h_output(n);
thrust::transform_exclusive_scan(
h_input.begin(), h_input.end(), h_output.begin(), ::cuda::std::negate<int>(), 5, ::cuda::std::multiplies<>{});

ASSERT_EQUAL(d_output, h_output);
}

// Test 3: User-defined type with poison default constructor
{
thrust::device_vector<TransformInt> vec(5, TransformInt{1});
thrust::transform_exclusive_scan(
vec.begin(),
vec.end(),
vec.begin(),
::cuda::std::identity{},
TransformInt{100},
::cuda::std::plus<TransformInt>());

ASSERT_EQUAL(vec,
(thrust::device_vector<TransformInt>{
TransformInt{100}, TransformInt{101}, TransformInt{102}, TransformInt{103}, TransformInt{104}}));
}

// Test 4: In-place transform_exclusive_scan with multiplies
{
thrust::device_vector<int> vec = {2, 3, 4, 5};
auto r = thrust::transform_exclusive_scan(
vec.begin(), vec.end(), vec.begin(), ::cuda::std::negate<int>(), 10, ::cuda::std::multiplies<>{});
ASSERT_EQUAL((vec.end() == r), true);

thrust::device_vector<int> expected = {10, -20, 60, -240};
ASSERT_EQUAL(vec, expected);
}

// Test 5: Boundary case at threshold (1024 elements)
{
const int n = 1024;
thrust::device_vector<int> d_input(n);
using thrust::placeholders::_1;
thrust::tabulate(d_input.begin(), d_input.end(), _1 + 1);

thrust::device_vector<int> d_output(n);
auto r = thrust::transform_inclusive_scan(
d_input.begin(), d_input.end(), d_output.begin(), thrust::square<int>(), 1, ::cuda::std::multiplies<>{});
ASSERT_EQUAL((d_output.end() == r), true);

thrust::host_vector<int> h_input = d_input;
thrust::host_vector<int> h_output(n);
thrust::transform_inclusive_scan(
h_input.begin(), h_input.end(), h_output.begin(), thrust::square<int>(), 1, ::cuda::std::multiplies<>{});

ASSERT_EQUAL(d_output, h_output);
}

// Test 6: Very small array (2 elements)
{
thrust::device_vector<int> d_input = {3, 7};
thrust::device_vector<int> d_output(2);

auto r = thrust::transform_inclusive_scan(
d_input.begin(), d_input.end(), d_output.begin(), ::cuda::std::negate<int>(), 2, ::cuda::std::multiplies<>{});
ASSERT_EQUAL((d_output.end() == r), true);

thrust::device_vector<int> expected = {-6, 42};
ASSERT_EQUAL(d_output, expected);
}
}
DECLARE_UNITTEST(TestTransformScanEdgeCases);
Loading