uses unsigned offset types in thrust's scan dispatch

NVIDIA · Jan 20, 2025 · 2eae8f1 · 2eae8f1
1 parent 8630af6
commit 2eae8f1
Show file tree

Hide file tree

Showing 3 changed files with 10 additions and 10 deletions.
diff --git a/thrust/thrust/system/cuda/detail/async/exclusive_scan.h b/thrust/thrust/system/cuda/detail/async/exclusive_scan.h
@@ -73,8 +73,8 @@ unique_eager_event async_exclusive_scan_n(
   execution_policy<DerivedPolicy>& policy, ForwardIt first, Size n, OutputIt out, InitialValueType init, BinaryOp op)
 {
   using InputValueT = cub::detail::InputValue<InitialValueType>;
-  using Dispatch32  = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, InputValueT, std::int32_t, InitialValueType>;
-  using Dispatch64  = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, InputValueT, std::int64_t, InitialValueType>;
+  using Dispatch32  = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, InputValueT, std::uint32_t, InitialValueType>;
+  using Dispatch64  = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, InputValueT, std::uint64_t, InitialValueType>;
 
   InputValueT init_value(init);
 
@@ -85,7 +85,7 @@ unique_eager_event async_exclusive_scan_n(
   cudaError_t status;
   size_t tmp_size = 0;
   {
-    THRUST_INDEX_TYPE_DISPATCH2(
+    THRUST_UNSIGNED_INDEX_TYPE_DISPATCH2(
       status,
       Dispatch32::Dispatch,
       Dispatch64::Dispatch,

diff --git a/thrust/thrust/system/cuda/detail/async/inclusive_scan.h b/thrust/thrust/system/cuda/detail/async/inclusive_scan.h
@@ -68,8 +68,8 @@ unique_eager_event
 async_inclusive_scan_n(execution_policy<DerivedPolicy>& policy, ForwardIt first, Size n, OutputIt out, BinaryOp op)
 {
   using AccumT     = typename thrust::iterator_traits<ForwardIt>::value_type;
-  using Dispatch32 = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, cub::NullType, std::int32_t, AccumT>;
-  using Dispatch64 = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, cub::NullType, std::int64_t, AccumT>;
+  using Dispatch32 = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, cub::NullType, std::uint32_t, AccumT>;
+  using Dispatch64 = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, cub::NullType, std::uint64_t, AccumT>;
 
   auto const device_alloc = get_async_device_allocator(policy);
   unique_eager_event ev;
@@ -78,7 +78,7 @@ async_inclusive_scan_n(execution_policy<DerivedPolicy>& policy, ForwardIt first,
   cudaError_t status;
   size_t tmp_size = 0;
   {
-    THRUST_INDEX_TYPE_DISPATCH2(
+    THRUST_UNSIGNED_INDEX_TYPE_DISPATCH2(
       status,
       Dispatch32::Dispatch,
       Dispatch64::Dispatch,

diff --git a/thrust/thrust/system/cuda/detail/scan.h b/thrust/thrust/system/cuda/detail/scan.h
@@ -128,7 +128,7 @@ _CCCL_HOST_DEVICE OutputIt inclusive_scan_n_impl(
                       OutputIt,
                       ScanOp,
                       InputValueT,
-                      std::int32_t,
+                      std::uint32_t,
                       AccumT,
                       cub::detail::scan::policy_hub<AccumT, ScanOp>,
                       ForceInclusive>;
@@ -137,7 +137,7 @@ _CCCL_HOST_DEVICE OutputIt inclusive_scan_n_impl(
                       OutputIt,
                       ScanOp,
                       InputValueT,
-                      std::int64_t,
+                      std::uint64_t,
                       AccumT,
                       cub::detail::scan::policy_hub<AccumT, ScanOp>,
                       ForceInclusive>;
@@ -154,7 +154,7 @@ _CCCL_HOST_DEVICE OutputIt inclusive_scan_n_impl(
   // Determine temporary storage requirements:
   size_t tmp_size = 0;
   {
-    THRUST_INDEX_TYPE_DISPATCH2(
+    THRUST_UNSIGNED_INDEX_TYPE_DISPATCH2(
       status,
       Dispatch32::Dispatch,
       Dispatch64::Dispatch,
@@ -170,7 +170,7 @@ _CCCL_HOST_DEVICE OutputIt inclusive_scan_n_impl(
   {
     // Allocate temporary storage:
     thrust::detail::temporary_array<std::uint8_t, Derived> tmp{policy, tmp_size};
-    THRUST_INDEX_TYPE_DISPATCH2(
+    THRUST_UNSIGNED_INDEX_TYPE_DISPATCH2(
       status,
       Dispatch32::Dispatch,
       Dispatch64::Dispatch,