Alien-XGBoost
view release on metacpan or search on metacpan
xgboost/cub/cub/device/dispatch/dispatch_radix_sort.cuh view on Meta::CPAN
// Downsweep policies
typedef AgentRadixSortDownsweepPolicy <256, CUB_MAX(1, 16 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, PRIMARY_RADIX_BITS> DownsweepPolicy;
typedef AgentRadixSortDownsweepPolicy <256, CUB_MAX(1, 16 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, ALT_RADIX_BITS> AltDownsweepPolicy;
// Upsweep policies
typedef DownsweepPolicy UpsweepPolicy;
typedef AltDownsweepPolicy AltUpsweepPolicy;
// Single-tile policy
typedef AgentRadixSortDownsweepPolicy <256, CUB_MAX(1, 19 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, PRIMARY_RADIX_BITS> SingleTilePolicy;
// Segmented policies
typedef DownsweepPolicy SegmentedPolicy;
typedef AltDownsweepPolicy AltSegmentedPolicy;
};
/// MaxPolicy
typedef Policy620 MaxPolicy;
};
/******************************************************************************
* Single-problem dispatch
******************************************************************************/
/**
* Utility class for dispatching the appropriately-tuned kernels for device-wide radix sort
*/
template <
bool IS_DESCENDING, ///< Whether or not the sorted-order is high-to-low
typename KeyT, ///< Key type
typename ValueT, ///< Value type
typename OffsetT> ///< Signed integer type for global offsets
struct DispatchRadixSort :
DeviceRadixSortPolicy<KeyT, ValueT, OffsetT>
{
//------------------------------------------------------------------------------
// Constants
//------------------------------------------------------------------------------
enum
{
// Whether this is a keys-only (or key-value) sort
KEYS_ONLY = (Equals<ValueT, NullType>::VALUE),
};
//------------------------------------------------------------------------------
// Problem state
//------------------------------------------------------------------------------
void *d_temp_storage; ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done.
size_t &temp_storage_bytes; ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation
DoubleBuffer<KeyT> &d_keys; ///< [in,out] Double-buffer whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys
DoubleBuffer<ValueT> &d_values; ///< [in,out] Double-buffer whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values
OffsetT num_items; ///< [in] Number of items to sort
int begin_bit; ///< [in] The beginning (least-significant) bit index needed for key comparison
int end_bit; ///< [in] The past-the-end (most-significant) bit index needed for key comparison
cudaStream_t stream; ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous; ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p false.
int ptx_version; ///< [in] PTX version
bool is_overwrite_okay; ///< [in] Whether is okay to overwrite source buffers
//------------------------------------------------------------------------------
// Constructor
//------------------------------------------------------------------------------
/// Constructor
CUB_RUNTIME_FUNCTION __forceinline__
DispatchRadixSort(
void* d_temp_storage,
size_t &temp_storage_bytes,
DoubleBuffer<KeyT> &d_keys,
DoubleBuffer<ValueT> &d_values,
OffsetT num_items,
int begin_bit,
int end_bit,
bool is_overwrite_okay,
cudaStream_t stream,
bool debug_synchronous,
int ptx_version)
:
d_temp_storage(d_temp_storage),
temp_storage_bytes(temp_storage_bytes),
d_keys(d_keys),
d_values(d_values),
num_items(num_items),
begin_bit(begin_bit),
end_bit(end_bit),
stream(stream),
debug_synchronous(debug_synchronous),
ptx_version(ptx_version),
is_overwrite_okay(is_overwrite_okay)
{}
//------------------------------------------------------------------------------
// Small-problem (single tile) invocation
//------------------------------------------------------------------------------
/// Invoke a single block to sort in-core
template <
typename ActivePolicyT, ///< Umbrella policy active for the target device
typename SingleTileKernelT> ///< Function type of cub::DeviceRadixSortSingleTileKernel
CUB_RUNTIME_FUNCTION __forceinline__
cudaError_t InvokeSingleTile(
SingleTileKernelT single_tile_kernel) ///< [in] Kernel function pointer to parameterization of cub::DeviceRadixSortSingleTileKernel
{
#ifndef CUB_RUNTIME_ENABLED
(void)single_tile_kernel;
// Kernel launch not supported from this device
return CubDebug(cudaErrorNotSupported );
#else
cudaError error = cudaSuccess;
do
{
// Return if the caller is simply requesting the size of the storage allocation
if (d_temp_storage == NULL)
{
temp_storage_bytes = 1;
break;
}
// Return if empty problem
if (num_items == 0)
break;
// Log single_tile_kernel configuration
if (debug_synchronous)
_CubLog("Invoking single_tile_kernel<<<%d, %d, 0, %lld>>>(), %d items per thread, %d SM occupancy, current bit %d, bit_grain %d\n",
1, ActivePolicyT::SingleTilePolicy::BLOCK_THREADS, (long long) stream,
ActivePolicyT::SingleTilePolicy::ITEMS_PER_THREAD, 1, begin_bit, ActivePolicyT::SingleTilePolicy::RADIX_BITS);
// Invoke upsweep_kernel with same grid size as downsweep_kernel
single_tile_kernel<<<1, ActivePolicyT::SingleTilePolicy::BLOCK_THREADS, 0, stream>>>(
d_keys.Current(),
d_keys.Alternate(),
d_values.Current(),
d_values.Alternate(),
num_items,
begin_bit,
end_bit);
// Check for failure to launch
if (CubDebug(error = cudaPeekAtLastError())) break;
// Sync the stream if specified to flush runtime errors
if (debug_synchronous && (CubDebug(error = SyncStream(stream)))) break;
// Update selector
d_keys.selector ^= 1;
d_values.selector ^= 1;
}
while (0);
return error;
#endif // CUB_RUNTIME_ENABLED
}
//------------------------------------------------------------------------------
// Normal problem size invocation
//------------------------------------------------------------------------------
/**
* Invoke a three-kernel sorting pass at the current bit.
*/
template <typename PassConfigT>
CUB_RUNTIME_FUNCTION __forceinline__
cudaError_t InvokePass(
const KeyT *d_keys_in,
KeyT *d_keys_out,
const ValueT *d_values_in,
ValueT *d_values_out,
OffsetT *d_spine,
int spine_length,
int ¤t_bit,
PassConfigT &pass_config)
{
cudaError error = cudaSuccess;
do
{
int pass_bits = CUB_MIN(pass_config.radix_bits, (end_bit - current_bit));
// Log upsweep_kernel configuration
if (debug_synchronous)
_CubLog("Invoking upsweep_kernel<<<%d, %d, 0, %lld>>>(), %d items per thread, %d SM occupancy, current bit %d, bit_grain %d\n",
pass_config.even_share.grid_size, pass_config.upsweep_config.block_threads, (long long) stream,
pass_config.upsweep_config.items_per_thread, pass_config.upsweep_config.sm_occupancy, current_bit, pass_bits);
// Invoke upsweep_kernel with same grid size as downsweep_kernel
pass_config.upsweep_kernel<<<pass_config.even_share.grid_size, pass_config.upsweep_config.block_threads, 0, stream>>>(
d_keys_in,
d_spine,
num_items,
current_bit,
pass_bits,
pass_config.even_share);
// Check for failure to launch
if (CubDebug(error = cudaPeekAtLastError())) break;
// Sync the stream if specified to flush runtime errors
if (debug_synchronous && (CubDebug(error = SyncStream(stream)))) break;
// Log scan_kernel configuration
if (debug_synchronous) _CubLog("Invoking scan_kernel<<<%d, %d, 0, %lld>>>(), %d items per thread\n",
1, pass_config.scan_config.block_threads, (long long) stream, pass_config.scan_config.items_per_thread);
// Invoke scan_kernel
pass_config.scan_kernel<<<1, pass_config.scan_config.block_threads, 0, stream>>>(
d_spine,
spine_length);
// Check for failure to launch
if (CubDebug(error = cudaPeekAtLastError())) break;
// Sync the stream if specified to flush runtime errors
if (debug_synchronous && (CubDebug(error = SyncStream(stream)))) break;
// Log downsweep_kernel configuration
if (debug_synchronous) _CubLog("Invoking downsweep_kernel<<<%d, %d, 0, %lld>>>(), %d items per thread, %d SM occupancy\n",
pass_config.even_share.grid_size, pass_config.downsweep_config.block_threads, (long long) stream,
pass_config.downsweep_config.items_per_thread, pass_config.downsweep_config.sm_occupancy);
// Invoke downsweep_kernel
pass_config.downsweep_kernel<<<pass_config.even_share.grid_size, pass_config.downsweep_config.block_threads, 0, stream>>>(
d_keys_in,
d_keys_out,
d_values_in,
d_values_out,
d_spine,
num_items,
current_bit,
pass_bits,
pass_config.even_share);
// Check for failure to launch
if (CubDebug(error = cudaPeekAtLastError())) break;
// Sync the stream if specified to flush runtime errors
if (debug_synchronous && (CubDebug(error = SyncStream(stream)))) break;
// Update current bit
current_bit += pass_bits;
}
while (0);
return error;
}
/// Pass configuration structure
template <
typename UpsweepKernelT,
typename ScanKernelT,
typename DownsweepKernelT>
struct PassConfig
{
UpsweepKernelT upsweep_kernel;
KernelConfig upsweep_config;
ScanKernelT scan_kernel;
KernelConfig scan_config;
DownsweepKernelT downsweep_kernel;
KernelConfig downsweep_config;
int radix_bits;
int radix_digits;
int max_downsweep_grid_size;
GridEvenShare<OffsetT> even_share;
/// Initialize pass configuration
template <
typename UpsweepPolicyT,
typename ScanPolicyT,
typename DownsweepPolicyT>
CUB_RUNTIME_FUNCTION __forceinline__
cudaError_t InitPassConfig(
UpsweepKernelT upsweep_kernel,
ScanKernelT scan_kernel,
DownsweepKernelT downsweep_kernel,
int ptx_version,
int sm_count,
int num_items)
{
cudaError error = cudaSuccess;
do
{
this->upsweep_kernel = upsweep_kernel;
this->scan_kernel = scan_kernel;
this->downsweep_kernel = downsweep_kernel;
radix_bits = DownsweepPolicyT::RADIX_BITS;
radix_digits = 1 << radix_bits;
if (CubDebug(error = upsweep_config.Init<UpsweepPolicyT>(upsweep_kernel))) break;
if (CubDebug(error = scan_config.Init<ScanPolicyT>(scan_kernel))) break;
if (CubDebug(error = downsweep_config.Init<DownsweepPolicyT>(downsweep_kernel))) break;
max_downsweep_grid_size = (downsweep_config.sm_occupancy * sm_count) * CUB_SUBSCRIPTION_FACTOR(ptx_version);
even_share = GridEvenShare<OffsetT>(
num_items,
xgboost/cub/cub/device/dispatch/dispatch_radix_sort.cuh view on Meta::CPAN
d_values.selector = (d_values.selector + num_passes) & 1;
}
while (0);
return error;
#endif // CUB_RUNTIME_ENABLED
}
//------------------------------------------------------------------------------
// Chained policy invocation
//------------------------------------------------------------------------------
/// Invocation
template <typename ActivePolicyT>
CUB_RUNTIME_FUNCTION __forceinline__
cudaError_t Invoke()
{
typedef typename DispatchRadixSort::MaxPolicy MaxPolicyT;
typedef typename ActivePolicyT::SingleTilePolicy SingleTilePolicyT;
// Force kernel code-generation in all compiler passes
if (num_items <= (SingleTilePolicyT::BLOCK_THREADS * SingleTilePolicyT::ITEMS_PER_THREAD))
{
// Small, single tile size
return InvokeSingleTile<ActivePolicyT>(
DeviceRadixSortSingleTileKernel<MaxPolicyT, IS_DESCENDING, KeyT, ValueT, OffsetT>);
}
else
{
// Regular size
return InvokePasses<ActivePolicyT>(
DeviceRadixSortUpsweepKernel< MaxPolicyT, false, IS_DESCENDING, KeyT, OffsetT>,
DeviceRadixSortUpsweepKernel< MaxPolicyT, true, IS_DESCENDING, KeyT, OffsetT>,
RadixSortScanBinsKernel< MaxPolicyT, OffsetT>,
DeviceRadixSortDownsweepKernel< MaxPolicyT, false, IS_DESCENDING, KeyT, ValueT, OffsetT>,
DeviceRadixSortDownsweepKernel< MaxPolicyT, true, IS_DESCENDING, KeyT, ValueT, OffsetT>);
}
}
//------------------------------------------------------------------------------
// Dispatch entrypoints
//------------------------------------------------------------------------------
/**
* Internal dispatch routine
*/
CUB_RUNTIME_FUNCTION __forceinline__
static cudaError_t Dispatch(
void* d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done.
size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation
DoubleBuffer<KeyT> &d_keys, ///< [in,out] Double-buffer whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys
DoubleBuffer<ValueT> &d_values, ///< [in,out] Double-buffer whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values
OffsetT num_items, ///< [in] Number of items to sort
int begin_bit, ///< [in] The beginning (least-significant) bit index needed for key comparison
int end_bit, ///< [in] The past-the-end (most-significant) bit index needed for key comparison
bool is_overwrite_okay, ///< [in] Whether is okay to overwrite source buffers
cudaStream_t stream, ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous) ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p false.
{
typedef typename DispatchRadixSort::MaxPolicy MaxPolicyT;
cudaError_t error;
do {
// Get PTX version
int ptx_version;
if (CubDebug(error = PtxVersion(ptx_version))) break;
// Create dispatch functor
DispatchRadixSort dispatch(
d_temp_storage, temp_storage_bytes,
d_keys, d_values,
num_items, begin_bit, end_bit, is_overwrite_okay,
stream, debug_synchronous, ptx_version);
// Dispatch to chained policy
if (CubDebug(error = MaxPolicyT::Invoke(ptx_version, dispatch))) break;
} while (0);
return error;
}
};
/******************************************************************************
* Segmented dispatch
******************************************************************************/
/**
* Utility class for dispatching the appropriately-tuned kernels for segmented device-wide radix sort
*/
template <
bool IS_DESCENDING, ///< Whether or not the sorted-order is high-to-low
typename KeyT, ///< Key type
typename ValueT, ///< Value type
typename OffsetT> ///< Signed integer type for global offsets
struct DispatchSegmentedRadixSort :
DeviceRadixSortPolicy<KeyT, ValueT, OffsetT>
{
//------------------------------------------------------------------------------
// Constants
//------------------------------------------------------------------------------
enum
{
// Whether this is a keys-only (or key-value) sort
KEYS_ONLY = (Equals<ValueT, NullType>::VALUE),
};
//------------------------------------------------------------------------------
// Parameter members
//------------------------------------------------------------------------------
void *d_temp_storage; ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done.
size_t &temp_storage_bytes; ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation
DoubleBuffer<KeyT> &d_keys; ///< [in,out] Double-buffer whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys
DoubleBuffer<ValueT> &d_values; ///< [in,out] Double-buffer whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values
OffsetT num_items; ///< [in] Number of items to sort
OffsetT num_segments; ///< [in] The number of segments that comprise the sorting data
const OffsetT *d_begin_offsets; ///< [in] %Device-accessible pointer to the sequence of beginning offsets of length \p num_segments, such that <tt>d_begin_offsets[i]</tt> is the first element of the <em>i</em><sup>th</sup> data se...
const OffsetT *d_end_offsets; ///< [in] %Device-accessible pointer to the sequence of ending offsets of length \p num_segments, such that <tt>d_end_offsets[i]-1</tt> is the last element of the <em>i</em><sup>th</sup> data segmen...
int begin_bit; ///< [in] The beginning (least-significant) bit index needed for key comparison
int end_bit; ///< [in] The past-the-end (most-significant) bit index needed for key comparison
cudaStream_t stream; ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous; ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p false.
int ptx_version; ///< [in] PTX version
bool is_overwrite_okay; ///< [in] Whether is okay to overwrite source buffers
//------------------------------------------------------------------------------
// Constructors
//------------------------------------------------------------------------------
/// Constructor
CUB_RUNTIME_FUNCTION __forceinline__
DispatchSegmentedRadixSort(
void* d_temp_storage,
size_t &temp_storage_bytes,
DoubleBuffer<KeyT> &d_keys,
DoubleBuffer<ValueT> &d_values,
OffsetT num_items,
OffsetT num_segments,
const OffsetT *d_begin_offsets,
const OffsetT *d_end_offsets,
int begin_bit,
int end_bit,
bool is_overwrite_okay,
cudaStream_t stream,
bool debug_synchronous,
int ptx_version)
:
d_temp_storage(d_temp_storage),
temp_storage_bytes(temp_storage_bytes),
d_keys(d_keys),
d_values(d_values),
num_items(num_items),
num_segments(num_segments),
d_begin_offsets(d_begin_offsets),
d_end_offsets(d_end_offsets),
begin_bit(begin_bit),
end_bit(end_bit),
is_overwrite_okay(is_overwrite_okay),
stream(stream),
debug_synchronous(debug_synchronous),
ptx_version(ptx_version)
{}
//------------------------------------------------------------------------------
// Multi-segment invocation
//------------------------------------------------------------------------------
/// Invoke a three-kernel sorting pass at the current bit.
template <typename PassConfigT>
CUB_RUNTIME_FUNCTION __forceinline__
cudaError_t InvokePass(
const KeyT *d_keys_in,
KeyT *d_keys_out,
const ValueT *d_values_in,
ValueT *d_values_out,
int ¤t_bit,
PassConfigT &pass_config)
{
cudaError error = cudaSuccess;
do
{
int pass_bits = CUB_MIN(pass_config.radix_bits, (end_bit - current_bit));
// Log kernel configuration
if (debug_synchronous)
_CubLog("Invoking segmented_kernels<<<%d, %d, 0, %lld>>>(), %d items per thread, %d SM occupancy, current bit %d, bit_grain %d\n",
num_segments, pass_config.segmented_config.block_threads, (long long) stream,
pass_config.segmented_config.items_per_thread, pass_config.segmented_config.sm_occupancy, current_bit, pass_bits);
pass_config.segmented_kernel<<<num_segments, pass_config.segmented_config.block_threads, 0, stream>>>(
d_keys_in, d_keys_out,
d_values_in, d_values_out,
d_begin_offsets, d_end_offsets, num_segments,
current_bit, pass_bits);
// Check for failure to launch
if (CubDebug(error = cudaPeekAtLastError())) break;
// Sync the stream if specified to flush runtime errors
if (debug_synchronous && (CubDebug(error = SyncStream(stream)))) break;
// Update current bit
current_bit += pass_bits;
}
while (0);
return error;
}
/// PassConfig data structure
template <typename SegmentedKernelT>
struct PassConfig
{
SegmentedKernelT segmented_kernel;
KernelConfig segmented_config;
int radix_bits;
int radix_digits;
/// Initialize pass configuration
template <typename SegmentedPolicyT>
CUB_RUNTIME_FUNCTION __forceinline__
cudaError_t InitPassConfig(SegmentedKernelT segmented_kernel)
{
this->segmented_kernel = segmented_kernel;
this->radix_bits = SegmentedPolicyT::RADIX_BITS;
this->radix_digits = 1 << radix_bits;
return CubDebug(segmented_config.Init<SegmentedPolicyT>(segmented_kernel));
}
};
/// Invocation (run multiple digit passes)
template <
typename ActivePolicyT, ///< Umbrella policy active for the target device
typename SegmentedKernelT> ///< Function type of cub::DeviceSegmentedRadixSortKernel
CUB_RUNTIME_FUNCTION __forceinline__
cudaError_t InvokePasses(
SegmentedKernelT segmented_kernel, ///< [in] Kernel function pointer to parameterization of cub::DeviceSegmentedRadixSortKernel
SegmentedKernelT alt_segmented_kernel) ///< [in] Alternate kernel function pointer to parameterization of cub::DeviceSegmentedRadixSortKernel
{
#ifndef CUB_RUNTIME_ENABLED
(void)segmented_kernel;
(void)alt_segmented_kernel;
// Kernel launch not supported from this device
return CubDebug(cudaErrorNotSupported );
#else
cudaError error = cudaSuccess;
do
{
// Init regular and alternate kernel configurations
PassConfig<SegmentedKernelT> pass_config, alt_pass_config;
if ((error = pass_config.template InitPassConfig<typename ActivePolicyT::SegmentedPolicy>(segmented_kernel))) break;
if ((error = alt_pass_config.template InitPassConfig<typename ActivePolicyT::AltSegmentedPolicy>(alt_segmented_kernel))) break;
// Temporary storage allocation requirements
void* allocations[2];
xgboost/cub/cub/device/dispatch/dispatch_radix_sort.cuh view on Meta::CPAN
// Invert selectors and update current bit
d_keys_remaining_passes.selector ^= 1;
d_values_remaining_passes.selector ^= 1;
}
// Update selector
if (!is_overwrite_okay) {
num_passes = 1; // Sorted data always ends up in the other vector
}
d_keys.selector = (d_keys.selector + num_passes) & 1;
d_values.selector = (d_values.selector + num_passes) & 1;
}
while (0);
return error;
#endif // CUB_RUNTIME_ENABLED
}
//------------------------------------------------------------------------------
// Chained policy invocation
//------------------------------------------------------------------------------
/// Invocation
template <typename ActivePolicyT>
CUB_RUNTIME_FUNCTION __forceinline__
cudaError_t Invoke()
{
typedef typename DispatchSegmentedRadixSort::MaxPolicy MaxPolicyT;
// Force kernel code-generation in all compiler passes
return InvokePasses<ActivePolicyT>(
DeviceSegmentedRadixSortKernel<MaxPolicyT, false, IS_DESCENDING, KeyT, ValueT, OffsetT>,
DeviceSegmentedRadixSortKernel<MaxPolicyT, true, IS_DESCENDING, KeyT, ValueT, OffsetT>);
}
//------------------------------------------------------------------------------
// Dispatch entrypoints
//------------------------------------------------------------------------------
/// Internal dispatch routine
CUB_RUNTIME_FUNCTION __forceinline__
static cudaError_t Dispatch(
void* d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done.
size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation
DoubleBuffer<KeyT> &d_keys, ///< [in,out] Double-buffer whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys
DoubleBuffer<ValueT> &d_values, ///< [in,out] Double-buffer whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values
int num_items, ///< [in] Number of items to sort
int num_segments, ///< [in] The number of segments that comprise the sorting data
const int *d_begin_offsets, ///< [in] %Device-accessible pointer to the sequence of beginning offsets of length \p num_segments, such that <tt>d_begin_offsets[i]</tt> is the first element of the <em>i</em><sup>th</sup> dat...
const int *d_end_offsets, ///< [in] %Device-accessible pointer to the sequence of ending offsets of length \p num_segments, such that <tt>d_end_offsets[i]-1</tt> is the last element of the <em>i</em><sup>th</sup> data se...
int begin_bit, ///< [in] The beginning (least-significant) bit index needed for key comparison
int end_bit, ///< [in] The past-the-end (most-significant) bit index needed for key comparison
bool is_overwrite_okay, ///< [in] Whether is okay to overwrite source buffers
cudaStream_t stream, ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous) ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p false.
{
typedef typename DispatchSegmentedRadixSort::MaxPolicy MaxPolicyT;
cudaError_t error;
do {
// Get PTX version
int ptx_version;
if (CubDebug(error = PtxVersion(ptx_version))) break;
// Create dispatch functor
DispatchSegmentedRadixSort dispatch(
d_temp_storage, temp_storage_bytes,
d_keys, d_values,
num_items, num_segments, d_begin_offsets, d_end_offsets,
begin_bit, end_bit, is_overwrite_okay,
stream, debug_synchronous, ptx_version);
// Dispatch to chained policy
if (CubDebug(error = MaxPolicyT::Invoke(ptx_version, dispatch))) break;
} while (0);
return error;
}
};
} // CUB namespace
CUB_NS_POSTFIX // Optional outer namespace(s)
( run in 1.809 second using v1.01-cache-2.11-cpan-39bf76dae61 )