/home/runner/work/cccl/cccl/cub/cub/device/device_scan.cuh
File members: /home/runner/work/cccl/cccl/cub/cub/device/device_scan.cuh
/******************************************************************************
* Copyright (c) 2011, Duane Merrill. All rights reserved.
* Copyright (c) 2011-2022, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are met:
* * Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
* * Neither the name of the NVIDIA CORPORATION nor the
* names of its contributors may be used to endorse or promote products
* derived from this software without specific prior written permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
* AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
* ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
#include <cub/config.cuh>
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#include <cub/detail/nvtx.cuh>
#include <cub/device/dispatch/dispatch_scan.cuh>
#include <cub/device/dispatch/dispatch_scan_by_key.cuh>
#include <cub/thread/thread_operators.cuh>
#include <cub/util_deprecated.cuh>
CUB_NAMESPACE_BEGIN
struct DeviceScan
{
template <typename InputIteratorT, typename OutputIteratorT>
CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveSum(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
int num_items,
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::ExclusiveSum");
// Signed integer type for global offsets
using OffsetT = int;
using InitT = cub::detail::value_t<InputIteratorT>;
// Initial value
InitT init_value{};
return DispatchScan<InputIteratorT, OutputIteratorT, Sum, detail::InputValue<InitT>, OffsetT>::Dispatch(
d_temp_storage, temp_storage_bytes, d_in, d_out, Sum(), detail::InputValue<InitT>(init_value), num_items, stream);
}
template <typename InputIteratorT, typename OutputIteratorT>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveSum(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return ExclusiveSum<InputIteratorT, OutputIteratorT>(
d_temp_storage, temp_storage_bytes, d_in, d_out, num_items, stream);
}
template <typename IteratorT>
CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveSum(
void* d_temp_storage, size_t& temp_storage_bytes, IteratorT d_data, int num_items, cudaStream_t stream = 0)
{
return ExclusiveSum(d_temp_storage, temp_storage_bytes, d_data, d_data, num_items, stream);
}
template <typename IteratorT>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveSum(
void* d_temp_storage,
size_t& temp_storage_bytes,
IteratorT d_data,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return ExclusiveSum<IteratorT>(d_temp_storage, temp_storage_bytes, d_data, num_items, stream);
}
template <typename InputIteratorT, typename OutputIteratorT, typename ScanOpT, typename InitValueT>
CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
ScanOpT scan_op,
InitValueT init_value,
int num_items,
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::ExclusiveScan");
// Signed integer type for global offsets
using OffsetT = int;
return DispatchScan<InputIteratorT, OutputIteratorT, ScanOpT, detail::InputValue<InitValueT>, OffsetT>::Dispatch(
d_temp_storage,
temp_storage_bytes,
d_in,
d_out,
scan_op,
detail::InputValue<InitValueT>(init_value),
num_items,
stream);
}
template <typename InputIteratorT, typename OutputIteratorT, typename ScanOpT, typename InitValueT>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
ScanOpT scan_op,
InitValueT init_value,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return ExclusiveScan<InputIteratorT, OutputIteratorT, ScanOpT, InitValueT>(
d_temp_storage, temp_storage_bytes, d_in, d_out, scan_op, init_value, num_items, stream);
}
template <typename IteratorT, typename ScanOpT, typename InitValueT>
CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
IteratorT d_data,
ScanOpT scan_op,
InitValueT init_value,
int num_items,
cudaStream_t stream = 0)
{
return ExclusiveScan(d_temp_storage, temp_storage_bytes, d_data, d_data, scan_op, init_value, num_items, stream);
}
template <typename IteratorT, typename ScanOpT, typename InitValueT>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
IteratorT d_data,
ScanOpT scan_op,
InitValueT init_value,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return ExclusiveScan<IteratorT, ScanOpT, InitValueT>(
d_temp_storage, temp_storage_bytes, d_data, scan_op, init_value, num_items, stream);
}
template <typename InputIteratorT,
typename OutputIteratorT,
typename ScanOpT,
typename InitValueT,
typename InitValueIterT = InitValueT*>
CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
ScanOpT scan_op,
FutureValue<InitValueT, InitValueIterT> init_value,
int num_items,
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::ExclusiveScan");
// Signed integer type for global offsets
using OffsetT = int;
return DispatchScan<InputIteratorT, OutputIteratorT, ScanOpT, detail::InputValue<InitValueT>, OffsetT>::Dispatch(
d_temp_storage,
temp_storage_bytes,
d_in,
d_out,
scan_op,
detail::InputValue<InitValueT>(init_value),
num_items,
stream);
}
template <typename InputIteratorT,
typename OutputIteratorT,
typename ScanOpT,
typename InitValueT,
typename InitValueIterT = InitValueT*>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
ScanOpT scan_op,
FutureValue<InitValueT, InitValueIterT> init_value,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return ExclusiveScan<InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, InitValueIterT>(
d_temp_storage, temp_storage_bytes, d_in, d_out, scan_op, init_value, num_items, stream);
}
template <typename IteratorT, typename ScanOpT, typename InitValueT, typename InitValueIterT = InitValueT*>
CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
IteratorT d_data,
ScanOpT scan_op,
FutureValue<InitValueT, InitValueIterT> init_value,
int num_items,
cudaStream_t stream = 0)
{
return ExclusiveScan(d_temp_storage, temp_storage_bytes, d_data, d_data, scan_op, init_value, num_items, stream);
}
template <typename IteratorT, typename ScanOpT, typename InitValueT, typename InitValueIterT = InitValueT*>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
IteratorT d_data,
ScanOpT scan_op,
FutureValue<InitValueT, InitValueIterT> init_value,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return ExclusiveScan<IteratorT, ScanOpT, InitValueT, InitValueIterT>(
d_temp_storage, temp_storage_bytes, d_data, scan_op, init_value, num_items, stream);
}
template <typename InputIteratorT, typename OutputIteratorT>
CUB_RUNTIME_FUNCTION static cudaError_t InclusiveSum(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
int num_items,
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::InclusiveSum");
// Signed integer type for global offsets
using OffsetT = int;
return DispatchScan<InputIteratorT, OutputIteratorT, Sum, NullType, OffsetT>::Dispatch(
d_temp_storage, temp_storage_bytes, d_in, d_out, Sum(), NullType(), num_items, stream);
}
template <typename InputIteratorT, typename OutputIteratorT>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t InclusiveSum(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return InclusiveSum<InputIteratorT, OutputIteratorT>(
d_temp_storage, temp_storage_bytes, d_in, d_out, num_items, stream);
}
template <typename IteratorT>
CUB_RUNTIME_FUNCTION static cudaError_t InclusiveSum(
void* d_temp_storage, size_t& temp_storage_bytes, IteratorT d_data, int num_items, cudaStream_t stream = 0)
{
return InclusiveSum(d_temp_storage, temp_storage_bytes, d_data, d_data, num_items, stream);
}
template <typename IteratorT>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t InclusiveSum(
void* d_temp_storage,
size_t& temp_storage_bytes,
IteratorT d_data,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return InclusiveSum<IteratorT>(d_temp_storage, temp_storage_bytes, d_data, num_items, stream);
}
template <typename InputIteratorT, typename OutputIteratorT, typename ScanOpT>
CUB_RUNTIME_FUNCTION static cudaError_t InclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
ScanOpT scan_op,
int num_items,
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::InclusiveScan");
// Signed integer type for global offsets
using OffsetT = int;
return DispatchScan<InputIteratorT, OutputIteratorT, ScanOpT, NullType, OffsetT>::Dispatch(
d_temp_storage, temp_storage_bytes, d_in, d_out, scan_op, NullType(), num_items, stream);
}
template <typename InputIteratorT, typename OutputIteratorT, typename ScanOpT>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t InclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
ScanOpT scan_op,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return InclusiveScan<InputIteratorT, OutputIteratorT, ScanOpT>(
d_temp_storage, temp_storage_bytes, d_in, d_out, scan_op, num_items, stream);
}
template <typename IteratorT, typename ScanOpT>
CUB_RUNTIME_FUNCTION static cudaError_t InclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
IteratorT d_data,
ScanOpT scan_op,
int num_items,
cudaStream_t stream = 0)
{
return InclusiveScan(d_temp_storage, temp_storage_bytes, d_data, d_data, scan_op, num_items, stream);
}
template <typename IteratorT, typename ScanOpT>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t InclusiveScan(
void* d_temp_storage,
size_t& temp_storage_bytes,
IteratorT d_data,
ScanOpT scan_op,
int num_items,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return InclusiveScan<IteratorT, ScanOpT>(d_temp_storage, temp_storage_bytes, d_data, scan_op, num_items, stream);
}
template <typename KeysInputIteratorT,
typename ValuesInputIteratorT,
typename ValuesOutputIteratorT,
typename EqualityOpT = Equality>
CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveSumByKey(
void* d_temp_storage,
size_t& temp_storage_bytes,
KeysInputIteratorT d_keys_in,
ValuesInputIteratorT d_values_in,
ValuesOutputIteratorT d_values_out,
int num_items,
EqualityOpT equality_op = EqualityOpT(),
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::ExclusiveSumByKey");
// Signed integer type for global offsets
using OffsetT = int;
using InitT = cub::detail::value_t<ValuesInputIteratorT>;
// Initial value
InitT init_value{};
return DispatchScanByKey<
KeysInputIteratorT,
ValuesInputIteratorT,
ValuesOutputIteratorT,
EqualityOpT,
Sum,
InitT,
OffsetT>::Dispatch(d_temp_storage,
temp_storage_bytes,
d_keys_in,
d_values_in,
d_values_out,
equality_op,
Sum(),
init_value,
num_items,
stream);
}
template <typename KeysInputIteratorT,
typename ValuesInputIteratorT,
typename ValuesOutputIteratorT,
typename EqualityOpT = Equality>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveSumByKey(
void* d_temp_storage,
size_t& temp_storage_bytes,
KeysInputIteratorT d_keys_in,
ValuesInputIteratorT d_values_in,
ValuesOutputIteratorT d_values_out,
int num_items,
EqualityOpT equality_op,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return ExclusiveSumByKey<KeysInputIteratorT, ValuesInputIteratorT, ValuesOutputIteratorT, EqualityOpT>(
d_temp_storage, temp_storage_bytes, d_keys_in, d_values_in, d_values_out, num_items, equality_op, stream);
}
template <typename KeysInputIteratorT,
typename ValuesInputIteratorT,
typename ValuesOutputIteratorT,
typename ScanOpT,
typename InitValueT,
typename EqualityOpT = Equality>
CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScanByKey(
void* d_temp_storage,
size_t& temp_storage_bytes,
KeysInputIteratorT d_keys_in,
ValuesInputIteratorT d_values_in,
ValuesOutputIteratorT d_values_out,
ScanOpT scan_op,
InitValueT init_value,
int num_items,
EqualityOpT equality_op = EqualityOpT(),
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::ExclusiveScanByKey");
// Signed integer type for global offsets
using OffsetT = int;
return DispatchScanByKey<
KeysInputIteratorT,
ValuesInputIteratorT,
ValuesOutputIteratorT,
EqualityOpT,
ScanOpT,
InitValueT,
OffsetT>::Dispatch(d_temp_storage,
temp_storage_bytes,
d_keys_in,
d_values_in,
d_values_out,
equality_op,
scan_op,
init_value,
num_items,
stream);
}
template <typename KeysInputIteratorT,
typename ValuesInputIteratorT,
typename ValuesOutputIteratorT,
typename ScanOpT,
typename InitValueT,
typename EqualityOpT = Equality>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t ExclusiveScanByKey(
void* d_temp_storage,
size_t& temp_storage_bytes,
KeysInputIteratorT d_keys_in,
ValuesInputIteratorT d_values_in,
ValuesOutputIteratorT d_values_out,
ScanOpT scan_op,
InitValueT init_value,
int num_items,
EqualityOpT equality_op,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return ExclusiveScanByKey<KeysInputIteratorT,
ValuesInputIteratorT,
ValuesOutputIteratorT,
ScanOpT,
InitValueT,
EqualityOpT>(
d_temp_storage,
temp_storage_bytes,
d_keys_in,
d_values_in,
d_values_out,
scan_op,
init_value,
num_items,
equality_op,
stream);
}
template <typename KeysInputIteratorT,
typename ValuesInputIteratorT,
typename ValuesOutputIteratorT,
typename EqualityOpT = Equality>
CUB_RUNTIME_FUNCTION static cudaError_t InclusiveSumByKey(
void* d_temp_storage,
size_t& temp_storage_bytes,
KeysInputIteratorT d_keys_in,
ValuesInputIteratorT d_values_in,
ValuesOutputIteratorT d_values_out,
int num_items,
EqualityOpT equality_op = EqualityOpT(),
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::InclusiveSumByKey");
// Signed integer type for global offsets
using OffsetT = int;
return DispatchScanByKey<
KeysInputIteratorT,
ValuesInputIteratorT,
ValuesOutputIteratorT,
EqualityOpT,
Sum,
NullType,
OffsetT>::Dispatch(d_temp_storage,
temp_storage_bytes,
d_keys_in,
d_values_in,
d_values_out,
equality_op,
Sum(),
NullType(),
num_items,
stream);
}
template <typename KeysInputIteratorT,
typename ValuesInputIteratorT,
typename ValuesOutputIteratorT,
typename EqualityOpT = Equality>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t InclusiveSumByKey(
void* d_temp_storage,
size_t& temp_storage_bytes,
KeysInputIteratorT d_keys_in,
ValuesInputIteratorT d_values_in,
ValuesOutputIteratorT d_values_out,
int num_items,
EqualityOpT equality_op,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return InclusiveSumByKey<KeysInputIteratorT, ValuesInputIteratorT, ValuesOutputIteratorT, EqualityOpT>(
d_temp_storage, temp_storage_bytes, d_keys_in, d_values_in, d_values_out, num_items, equality_op, stream);
}
template <typename KeysInputIteratorT,
typename ValuesInputIteratorT,
typename ValuesOutputIteratorT,
typename ScanOpT,
typename EqualityOpT = Equality>
CUB_RUNTIME_FUNCTION static cudaError_t InclusiveScanByKey(
void* d_temp_storage,
size_t& temp_storage_bytes,
KeysInputIteratorT d_keys_in,
ValuesInputIteratorT d_values_in,
ValuesOutputIteratorT d_values_out,
ScanOpT scan_op,
int num_items,
EqualityOpT equality_op = EqualityOpT(),
cudaStream_t stream = 0)
{
CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceScan::InclusiveScanByKey");
// Signed integer type for global offsets
using OffsetT = int;
return DispatchScanByKey<
KeysInputIteratorT,
ValuesInputIteratorT,
ValuesOutputIteratorT,
EqualityOpT,
ScanOpT,
NullType,
OffsetT>::Dispatch(d_temp_storage,
temp_storage_bytes,
d_keys_in,
d_values_in,
d_values_out,
equality_op,
scan_op,
NullType(),
num_items,
stream);
}
template <typename KeysInputIteratorT,
typename ValuesInputIteratorT,
typename ValuesOutputIteratorT,
typename ScanOpT,
typename EqualityOpT = Equality>
CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED CUB_RUNTIME_FUNCTION static cudaError_t InclusiveScanByKey(
void* d_temp_storage,
size_t& temp_storage_bytes,
KeysInputIteratorT d_keys_in,
ValuesInputIteratorT d_values_in,
ValuesOutputIteratorT d_values_out,
ScanOpT scan_op,
int num_items,
EqualityOpT equality_op,
cudaStream_t stream,
bool debug_synchronous)
{
CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG
return InclusiveScanByKey<KeysInputIteratorT, ValuesInputIteratorT, ValuesOutputIteratorT, ScanOpT, EqualityOpT>(
d_temp_storage, temp_storage_bytes, d_keys_in, d_values_in, d_values_out, scan_op, num_items, equality_op, stream);
}
};
CUB_NAMESPACE_END