diff --git a/thrust/system/cuda/detail/async/exclusive_scan.h b/thrust/system/cuda/detail/async/exclusive_scan.h index 4ecbf43b2..0b120a434 100644 --- a/thrust/system/cuda/detail/async/exclusive_scan.h +++ b/thrust/system/cuda/detail/async/exclusive_scan.h @@ -79,12 +79,14 @@ async_exclusive_scan_n(execution_policy& policy, OutputIt, BinaryOp, InputValueT, - thrust::detail::int32_t>; + thrust::detail::int32_t, + InitialValueType>; using Dispatch64 = cub::DispatchScan; + thrust::detail::int64_t, + InitialValueType>; InputValueT init_value(init); diff --git a/thrust/system/cuda/detail/async/inclusive_scan.h b/thrust/system/cuda/detail/async/inclusive_scan.h index ab8e4e97b..363347c35 100644 --- a/thrust/system/cuda/detail/async/inclusive_scan.h +++ b/thrust/system/cuda/detail/async/inclusive_scan.h @@ -72,16 +72,19 @@ async_inclusive_scan_n(execution_policy& policy, OutputIt out, BinaryOp op) { + using AccumT = typename thrust::iterator_traits::value_type; using Dispatch32 = cub::DispatchScan; + thrust::detail::int32_t, + AccumT>; using Dispatch64 = cub::DispatchScan; + thrust::detail::int64_t, + AccumT>; auto const device_alloc = get_async_device_allocator(policy); unique_eager_event ev; diff --git a/thrust/system/cuda/detail/scan.h b/thrust/system/cuda/detail/scan.h index 68434f7e3..fdab8df84 100644 --- a/thrust/system/cuda/detail/scan.h +++ b/thrust/system/cuda/detail/scan.h @@ -60,16 +60,19 @@ OutputIt inclusive_scan_n_impl(thrust::cuda_cub::execution_policy &poli OutputIt result, ScanOp scan_op) { + using AccumT = typename thrust::iterator_traits::value_type; using Dispatch32 = cub::DispatchScan; + thrust::detail::int32_t, + AccumT>; using Dispatch64 = cub::DispatchScan; + thrust::detail::int64_t, + AccumT>; cudaStream_t stream = thrust::cuda_cub::stream(policy); cudaError_t status; @@ -141,12 +144,14 @@ OutputIt exclusive_scan_n_impl(thrust::cuda_cub::execution_policy &poli OutputIt, ScanOp, InputValueT, - thrust::detail::int32_t>; + thrust::detail::int32_t, + InitValueT>; using Dispatch64 = cub::DispatchScan; + thrust::detail::int64_t, + InitValueT>; cudaStream_t stream = thrust::cuda_cub::stream(policy); cudaError_t status; diff --git a/thrust/system/cuda/detail/scan_by_key.h b/thrust/system/cuda/detail/scan_by_key.h index 70077f343..0407779c6 100644 --- a/thrust/system/cuda/detail/scan_by_key.h +++ b/thrust/system/cuda/detail/scan_by_key.h @@ -87,6 +87,7 @@ ValuesOutIt inclusive_scan_by_key_n( thrust::detail::try_unwrap_contiguous_iterator_return_t; using ValuesOutUnwrapIt = thrust::detail::try_unwrap_contiguous_iterator_return_t; + using AccumT = typename thrust::iterator_traits::value_type; auto keys_unwrap = thrust::detail::try_unwrap_contiguous_iterator(keys); auto values_unwrap = thrust::detail::try_unwrap_contiguous_iterator(values); @@ -98,14 +99,16 @@ ValuesOutIt inclusive_scan_by_key_n( EqualityOpT, ScanOpT, cub::NullType, - thrust::detail::int32_t>; + thrust::detail::int32_t, + AccumT>; using Dispatch64 = cub::DispatchScanByKey; + thrust::detail::int64_t, + AccumT>; cudaStream_t stream = thrust::cuda_cub::stream(policy); cudaError_t status{}; @@ -209,14 +212,16 @@ ValuesOutIt exclusive_scan_by_key_n( EqualityOpT, ScanOpT, InitValueT, - thrust::detail::int32_t>; + thrust::detail::int32_t, + InitValueT>; using Dispatch64 = cub::DispatchScanByKey; + thrust::detail::int64_t, + InitValueT>; cudaStream_t stream = thrust::cuda_cub::stream(policy); cudaError_t status{};