| /****************************************************************************** | |
| * 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. | |
| * | |
| ******************************************************************************/ | |
| /** | |
| * @file cub::DeviceHistogram provides device-wide parallel operations for | |
| * constructing histogram(s) from a sequence of samples data residing | |
| * within device-accessible memory. | |
| */ | |
| #pragma once | |
| #include <stdio.h> | |
| #include <iterator> | |
| #include <limits> | |
| #include <cub/config.cuh> | |
| #include <cub/device/dispatch/dispatch_histogram.cuh> | |
| #include <cub/util_deprecated.cuh> | |
| CUB_NAMESPACE_BEGIN | |
| /** | |
| * @brief DeviceHistogram provides device-wide parallel operations for | |
| * constructing histogram(s) from a sequence of samples data residing | |
| * within device-accessible memory.  | |
| * @ingroup SingleModule | |
| * | |
| * @par Overview | |
| * A <a href="http://en.wikipedia.org/wiki/Histogram"><em>histogram</em></a> | |
| * counts the number of observations that fall into each of the disjoint categories (known as <em>bins</em>). | |
| * | |
| * @par Usage Considerations | |
| * @cdp_class{DeviceHistogram} | |
| * | |
| */ | |
| struct DeviceHistogram | |
| { | |
| /******************************************************************//** | |
| * @name Evenly-segmented bin ranges | |
| *********************************************************************/ | |
| //@{ | |
| /** | |
| * @brief Computes an intensity histogram from a sequence of data samples | |
| * using equal-width bins. | |
| * | |
| * @par | |
| * - The number of histogram bins is (`num_levels - 1`) | |
| * - All bins comprise the same width of sample values: | |
| * `(upper_level - lower_level) / (num_levels - 1)`. | |
| * - If the common type of `SampleT` and `LevelT` is of integral type, the bin for a sample is | |
| * computed as `(sample - lower_level) * (num_levels - 1) / (upper_level - lower_level)`, round | |
| * down to the nearest whole number. To protect against potential overflows, if the product | |
| * `(upper_level - lower_level) * (num_levels - 1)` exceeds the number representable by an | |
| * `uint64_t`, the cuda error `cudaErrorInvalidValue` is returned. If the common type is 128 | |
| * bits wide, bin computation will use 128-bit arithmetic and `cudaErrorInvalidValue` will only | |
| * be returned if bin computation would overflow for 128-bit arithmetic. | |
| * - The ranges `[d_samples, d_samples + num_samples)` and | |
| * `[d_histogram, d_histogram + num_levels - 1)` shall not overlap | |
| * in any way. | |
| * - `cuda::std::common_type<LevelT, SampleT>` must be valid, and both LevelT | |
| * and SampleT must be valid arithmetic types. The common type must be | |
| * convertible to `int` and trivially copyable. | |
| * - @devicestorage | |
| * | |
| * @par Snippet | |
| * The code snippet below illustrates the computation of a six-bin histogram | |
| * from a sequence of float samples | |
| * | |
| * @par | |
| * @code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for | |
| * // input samples and output histogram | |
| * int num_samples; // e.g., 10 | |
| * float* d_samples; // e.g., [2.2, 6.1, 7.1, 2.9, 3.5, 0.3, 2.9, 2.1, 6.1, 999.5] | |
| * int* d_histogram; // e.g., [ -, -, -, -, -, -] | |
| * int num_levels; // e.g., 7 (seven level boundaries for six bins) | |
| * float lower_level; // e.g., 0.0 (lower sample value boundary of lowest bin) | |
| * float upper_level; // e.g., 12.0 (upper sample value boundary of upper bin) | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void* d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceHistogram::HistogramEven( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, | |
| * lower_level, upper_level, num_samples); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Compute histograms | |
| * cub::DeviceHistogram::HistogramEven( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, | |
| * lower_level, upper_level, num_samples); | |
| * | |
| * // d_histogram <-- [1, 5, 0, 3, 0, 0]; | |
| * @endcode | |
| * | |
| * @tparam SampleIteratorT | |
| * **[inferred]** Random-access input iterator type for reading input | |
| * samples \iterator | |
| * | |
| * @tparam CounterT | |
| * **[inferred]** Integer type for histogram bin counters | |
| * | |
| * @tparam LevelT | |
| * **[inferred]** Type for specifying boundaries (levels) | |
| * | |
| * @tparam OffsetT | |
| * **[inferred]** Signed integer type for sequence offsets, list lengths, | |
| * pointer differences, etc. \offset_size1 | |
| * | |
| * @param[in] d_temp_storage | |
| * Device-accessible allocation of temporary storage. When `nullptr`, the | |
| * required allocation size is written to `temp_storage_bytes` and no | |
| * work is done. | |
| * | |
| * @param[in,out] temp_storage_bytes | |
| * Reference to size in bytes of `d_temp_storage` allocation | |
| * | |
| * @param[in] d_samples | |
| * The pointer to the input sequence of data samples. | |
| * | |
| * @param[out] d_histogram | |
| * The pointer to the histogram counter output array of length | |
| * `num_levels - 1`. | |
| * | |
| * @param[in] num_levels | |
| * The number of boundaries (levels) for delineating histogram samples. | |
| * Implies that the number of bins is `num_levels - 1`. | |
| * | |
| * @param[in] lower_level | |
| * The lower sample value bound (inclusive) for the lowest histogram bin. | |
| * | |
| * @param[in] upper_level | |
| * The upper sample value bound (exclusive) for the highest histogram bin. | |
| * | |
| * @param[in] num_samples | |
| * The number of input samples (i.e., the length of `d_samples`) | |
| * | |
| * @param[in] stream | |
| * **[optional]** CUDA stream to launch kernels within. | |
| * Default is stream<sub>0</sub>. | |
| */ | |
| template <typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| HistogramEven(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram, | |
| int num_levels, | |
| LevelT lower_level, | |
| LevelT upper_level, | |
| OffsetT num_samples, | |
| cudaStream_t stream = 0) | |
| { | |
| /// The sample value type of the input iterator | |
| using SampleT = cub::detail::value_t<SampleIteratorT>; | |
| CounterT *d_histogram1[1] = {d_histogram}; | |
| int num_levels1[1] = {num_levels}; | |
| LevelT lower_level1[1] = {lower_level}; | |
| LevelT upper_level1[1] = {upper_level}; | |
| return MultiHistogramEven<1, 1>(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram1, | |
| num_levels1, | |
| lower_level1, | |
| upper_level1, | |
| num_samples, | |
| static_cast<OffsetT>(1), | |
| sizeof(SampleT) * num_samples, | |
| stream); | |
| } | |
| template <typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| HistogramEven(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram, | |
| int num_levels, | |
| LevelT lower_level, | |
| LevelT upper_level, | |
| OffsetT num_samples, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return HistogramEven(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| lower_level, | |
| upper_level, | |
| num_samples, | |
| stream); | |
| } | |
| /** | |
| * @brief Computes an intensity histogram from a sequence of data samples | |
| * using equal-width bins. | |
| * | |
| * @par | |
| * - A two-dimensional *region of interest* within `d_samples` can be | |
| * specified using the `num_row_samples`, `num_rows`, and | |
| * `row_stride_bytes` parameters. | |
| * - The row stride must be a whole multiple of the sample data type | |
| * size, i.e., `(row_stride_bytes % sizeof(SampleT)) == 0`. | |
| * - The number of histogram bins is (`num_levels - 1`) | |
| * - All bins comprise the same width of sample values: | |
| * `(upper_level - lower_level) / (num_levels - 1)` | |
| * - If the common type of `SampleT` and `LevelT` is of integral type, the bin for a sample is | |
| * computed as `(sample - lower_level) * (num_levels - 1) / (upper_level - lower_level)`, round | |
| * down to the nearest whole number. To protect against potential overflows, if the product | |
| * `(upper_level - lower_level) * (num_levels - 1)` exceeds the number representable by an | |
| * `uint64_t`, the cuda error `cudaErrorInvalidValue` is returned. If the common type is 128 | |
| * bits wide, bin computation will use 128-bit arithmetic and `cudaErrorInvalidValue` will only | |
| * be returned if bin computation would overflow for 128-bit arithmetic. | |
| * - For a given row `r` in `[0, num_rows)`, let | |
| * `row_begin = d_samples + r * row_stride_bytes / sizeof(SampleT)` and | |
| * `row_end = row_begin + num_row_samples`. The ranges | |
| * `[row_begin, row_end)` and `[d_histogram, d_histogram + num_levels - 1)` | |
| * shall not overlap in any way. | |
| * - `cuda::std::common_type<LevelT, SampleT>` must be valid, and both LevelT | |
| * and SampleT must be valid arithmetic types. The common type must be | |
| * convertible to `int` and trivially copyable. | |
| * - @devicestorage | |
| * | |
| * @par Snippet | |
| * The code snippet below illustrates the computation of a six-bin histogram | |
| * from a 2x5 region of interest within a flattened 2x7 array of float samples. | |
| * | |
| * @par | |
| * @code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for | |
| * // input samples and output histogram | |
| * int num_row_samples; // e.g., 5 | |
| * int num_rows; // e.g., 2; | |
| * size_t row_stride_bytes; // e.g., 7 * sizeof(float) | |
| * float* d_samples; // e.g., [2.2, 6.1, 7.1, 2.9, 3.5, -, -, | |
| * // 0.3, 2.9, 2.1, 6.1, 999.5, -, -] | |
| * int* d_histogram; // e.g., [ -, -, -, -, -, -] | |
| * int num_levels; // e.g., 7 (seven level boundaries for six bins) | |
| * float lower_level; // e.g., 0.0 (lower sample value boundary of lowest bin) | |
| * float upper_level; // e.g., 12.0 (upper sample value boundary of upper bin) | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void* d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceHistogram::HistogramEven( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, lower_level, upper_level, | |
| * num_row_samples, num_rows, row_stride_bytes); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Compute histograms | |
| * cub::DeviceHistogram::HistogramEven( | |
| * d_temp_storage, temp_storage_bytes, d_samples, d_histogram, | |
| * d_samples, d_histogram, num_levels, lower_level, upper_level, | |
| * num_row_samples, num_rows, row_stride_bytes); | |
| * | |
| * // d_histogram <-- [1, 5, 0, 3, 0, 0]; | |
| * @endcode | |
| * | |
| * @tparam SampleIteratorT | |
| * **[inferred]** Random-access input iterator type for reading | |
| * input samples. \iterator | |
| * | |
| * @tparam CounterT | |
| * **[inferred]** Integer type for histogram bin counters | |
| * | |
| * @tparam LevelT | |
| * **[inferred]** Type for specifying boundaries (levels) | |
| * | |
| * @tparam OffsetT | |
| * **[inferred]** Signed integer type for sequence offsets, list lengths, | |
| * pointer differences, etc. \offset_size1 | |
| * @param[in] d_temp_storage | |
| * Device-accessible allocation of temporary storage. When `nullptr`, the | |
| * required allocation size is written to `temp_storage_bytes` and no | |
| * work is done. | |
| * | |
| * @param[in,out] temp_storage_bytes | |
| * Reference to size in bytes of `d_temp_storage` allocation | |
| * | |
| * @param[in] d_samples | |
| * The pointer to the input sequence of data samples. | |
| * | |
| * @param[out] d_histogram | |
| * The pointer to the histogram counter output array of | |
| * length `num_levels - 1`. | |
| * | |
| * @param[in] num_levels | |
| * The number of boundaries (levels) for delineating histogram samples. | |
| * Implies that the number of bins is `num_levels - 1`. | |
| * | |
| * @param[in] lower_level | |
| * The lower sample value bound (inclusive) for the lowest histogram bin. | |
| * | |
| * @param[in] upper_level | |
| * The upper sample value bound (exclusive) for the highest histogram bin. | |
| * | |
| * @param[in] num_row_samples | |
| * The number of data samples per row in the region of interest | |
| * | |
| * @param[in] num_rows | |
| * The number of rows in the region of interest | |
| * | |
| * @param[in] row_stride_bytes | |
| * The number of bytes between starts of consecutive rows in | |
| * the region of interest | |
| * | |
| * @param[in] stream | |
| * **[optional]** CUDA stream to launch kernels within. | |
| * Default is stream<sub>0</sub>. | |
| */ | |
| template <typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| HistogramEven(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram, | |
| int num_levels, | |
| LevelT lower_level, | |
| LevelT upper_level, | |
| OffsetT num_row_samples, | |
| OffsetT num_rows, | |
| size_t row_stride_bytes, | |
| cudaStream_t stream = 0) | |
| { | |
| CounterT *d_histogram1[1] = {d_histogram}; | |
| int num_levels1[1] = {num_levels}; | |
| LevelT lower_level1[1] = {lower_level}; | |
| LevelT upper_level1[1] = {upper_level}; | |
| return MultiHistogramEven<1, 1>(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram1, | |
| num_levels1, | |
| lower_level1, | |
| upper_level1, | |
| num_row_samples, | |
| num_rows, | |
| row_stride_bytes, | |
| stream); | |
| } | |
| template <typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| HistogramEven(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram, | |
| int num_levels, | |
| LevelT lower_level, | |
| LevelT upper_level, | |
| OffsetT num_row_samples, | |
| OffsetT num_rows, | |
| size_t row_stride_bytes, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return HistogramEven(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| lower_level, | |
| upper_level, | |
| num_row_samples, | |
| num_rows, | |
| row_stride_bytes, | |
| stream); | |
| } | |
| /** | |
| * @brief Computes per-channel intensity histograms from a sequence of | |
| * multi-channel "pixel" data samples using equal-width bins. | |
| * | |
| * @par | |
| * - The input is a sequence of *pixel* structures, where each pixel comprises | |
| * a record of `NUM_CHANNELS` consecutive data samples | |
| * (e.g., an *RGBA* pixel). | |
| * - Of the `NUM_CHANNELS` specified, the function will only compute | |
| * histograms for the first `NUM_ACTIVE_CHANNELS` | |
| * (e.g., only *RGB* histograms from *RGBA* pixel samples). | |
| * - The number of histogram bins for channel<sub><em>i</em></sub> is | |
| * `num_levels[i] - 1`. | |
| * - For channel<sub><em>i</em></sub>, the range of values for all histogram bins | |
| * have the same width: | |
| * `(upper_level[i] - lower_level[i]) / (num_levels[i] - 1)` | |
| * - If the common type of sample and level is of integral type, the bin for a sample is | |
| * computed as `(sample - lower_level[i]) * (num_levels - 1) / (upper_level[i] - | |
| * lower_level[i])`, round down to the nearest whole number. To protect against potential | |
| * overflows, if, for any channel `i`, the product `(upper_level[i] - lower_level[i]) * | |
| * (num_levels[i] - 1)` exceeds the number representable by an `uint64_t`, the cuda error | |
| * `cudaErrorInvalidValue` is returned. If the common type is 128 bits wide, bin computation | |
| * will use 128-bit arithmetic and `cudaErrorInvalidValue` will only be returned if bin | |
| * computation would overflow for 128-bit arithmetic. | |
| * - For a given channel `c` in `[0, NUM_ACTIVE_CHANNELS)`, the ranges | |
| * `[d_samples, d_samples + NUM_CHANNELS * num_pixels)` and | |
| * `[d_histogram[c], d_histogram[c] + num_levels[c] - 1)` shall not overlap | |
| * in any way. | |
| * - `cuda::std::common_type<LevelT, SampleT>` must be valid, and both LevelT | |
| * and SampleT must be valid arithmetic types. The common type must be | |
| * convertible to `int` and trivially copyable. | |
| * - @devicestorage | |
| * | |
| * @par Snippet | |
| * The code snippet below illustrates the computation of three 256-bin <em>RGB</em> histograms | |
| * from a quad-channel sequence of <em>RGBA</em> pixels (8 bits per channel per pixel) | |
| * | |
| * @par | |
| * @code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for | |
| * // input samples and output histograms | |
| * int num_pixels; // e.g., 5 | |
| * unsigned char* d_samples; // e.g., [(2, 6, 7, 5), (3, 0, 2, 1), (7, 0, 6, 2), | |
| * // (0, 6, 7, 5), (3, 0, 2, 6)] | |
| * int* d_histogram[3]; // e.g., three device pointers to three device buffers, | |
| * // each allocated with 256 integer counters | |
| * int num_levels[3]; // e.g., {257, 257, 257}; | |
| * unsigned int lower_level[3]; // e.g., {0, 0, 0}; | |
| * unsigned int upper_level[3]; // e.g., {256, 256, 256}; | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void* d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceHistogram::MultiHistogramEven<4, 3>( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, | |
| * lower_level, upper_level, num_pixels); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Compute histograms | |
| * cub::DeviceHistogram::MultiHistogramEven<4, 3>( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, | |
| * lower_level, upper_level, num_pixels); | |
| * | |
| * // d_histogram <-- [ [1, 0, 1, 2, 0, 0, 0, 1, 0, 0, 0, ..., 0], | |
| * // [0, 3, 0, 0, 0, 0, 2, 0, 0, 0, 0, ..., 0], | |
| * // [0, 0, 2, 0, 0, 0, 1, 2, 0, 0, 0, ..., 0] ] | |
| * @endcode | |
| * | |
| * @tparam NUM_CHANNELS | |
| * Number of channels interleaved in the input data (may be greater than | |
| * the number of channels being actively histogrammed) | |
| * | |
| * @tparam NUM_ACTIVE_CHANNELS | |
| * **[inferred]** Number of channels actively being histogrammed | |
| * | |
| * @tparam SampleIteratorT | |
| * **[inferred]** Random-access input iterator type for reading | |
| * input samples. \iterator | |
| * | |
| * @tparam CounterT | |
| * **[inferred]** Integer type for histogram bin counters | |
| * | |
| * @tparam LevelT | |
| * **[inferred]** Type for specifying boundaries (levels) | |
| * | |
| * @tparam OffsetT | |
| * **[inferred]** Signed integer type for sequence offsets, list lengths, | |
| * pointer differences, etc. \offset_size1 | |
| * | |
| * @param[in] d_temp_storage | |
| * Device-accessible allocation of temporary storage. When `nullptr`, the | |
| * required allocation size is written to `temp_storage_bytes` and no | |
| * work is done. | |
| * | |
| * @param[in,out] temp_storage_bytes | |
| * Reference to size in bytes of `d_temp_storage` allocation | |
| * | |
| * @param[in] d_samples | |
| * The pointer to the multi-channel input sequence of data samples. | |
| * The samples from different channels are assumed to be interleaved | |
| * (e.g., an array of 32-bit pixels where each pixel consists of four | |
| * *RGBA* 8-bit samples). | |
| * | |
| * @param[out] d_histogram | |
| * The pointers to the histogram counter output arrays, one for each active | |
| * channel. For channel<sub><em>i</em></sub>, the allocation length of | |
| * `d_histogram[i]` should be `num_levels[i] - 1`. | |
| * | |
| * @param[in] num_levels | |
| * The number of boundaries (levels) for delineating histogram samples in | |
| * each active channel. Implies that the number of bins for | |
| * channel<sub><em>i</em></sub> is `num_levels[i] - 1`. | |
| * | |
| * @param[in] lower_level | |
| * The lower sample value bound (inclusive) for the lowest histogram bin in | |
| * each active channel. | |
| * | |
| * @param[in] upper_level | |
| * The upper sample value bound (exclusive) for the highest histogram bin | |
| * in each active channel. | |
| * | |
| * @param[in] num_pixels | |
| * The number of multi-channel pixels | |
| * (i.e., the length of `d_samples / NUM_CHANNELS`) | |
| * | |
| * @param[in] stream | |
| * **[optional]** CUDA stream to launch kernels within. | |
| * Default is stream<sub>0</sub>. | |
| */ | |
| template <int NUM_CHANNELS, | |
| int NUM_ACTIVE_CHANNELS, | |
| typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| MultiHistogramEven(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram[NUM_ACTIVE_CHANNELS], | |
| int num_levels[NUM_ACTIVE_CHANNELS], | |
| LevelT lower_level[NUM_ACTIVE_CHANNELS], | |
| LevelT upper_level[NUM_ACTIVE_CHANNELS], | |
| OffsetT num_pixels, | |
| cudaStream_t stream = 0) | |
| { | |
| /// The sample value type of the input iterator | |
| using SampleT = cub::detail::value_t<SampleIteratorT>; | |
| return MultiHistogramEven<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>( | |
| d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| lower_level, | |
| upper_level, | |
| num_pixels, | |
| static_cast<OffsetT>(1), | |
| sizeof(SampleT) * NUM_CHANNELS * num_pixels, | |
| stream); | |
| } | |
| template <int NUM_CHANNELS, | |
| int NUM_ACTIVE_CHANNELS, | |
| typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| MultiHistogramEven(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram[NUM_ACTIVE_CHANNELS], | |
| int num_levels[NUM_ACTIVE_CHANNELS], | |
| LevelT lower_level[NUM_ACTIVE_CHANNELS], | |
| LevelT upper_level[NUM_ACTIVE_CHANNELS], | |
| OffsetT num_pixels, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return MultiHistogramEven(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| lower_level, | |
| upper_level, | |
| num_pixels, | |
| stream); | |
| } | |
| /** | |
| * @brief Computes per-channel intensity histograms from a sequence of | |
| * multi-channel "pixel" data samples using equal-width bins. | |
| * | |
| * @par | |
| * - The input is a sequence of *pixel* structures, where each pixel | |
| * comprises a record of `NUM_CHANNELS` consecutive data samples | |
| * (e.g., an *RGBA* pixel). | |
| * - Of the `NUM_CHANNELS` specified, the function will only compute | |
| * histograms for the first `NUM_ACTIVE_CHANNELS` (e.g., only *RGB* | |
| * histograms from *RGBA* pixel samples). | |
| * - A two-dimensional *region of interest* within `d_samples` can be | |
| * specified using the `num_row_samples`, `num_rows`, and | |
| * `row_stride_bytes` parameters. | |
| * - The row stride must be a whole multiple of the sample data type | |
| * size, i.e., `(row_stride_bytes % sizeof(SampleT)) == 0`. | |
| * - The number of histogram bins for channel<sub><em>i</em></sub> is | |
| * `num_levels[i] - 1`. | |
| * - For channel<sub><em>i</em></sub>, the range of values for all histogram | |
| * bins have the same width: | |
| * `(upper_level[i] - lower_level[i]) / (num_levels[i] - 1)` | |
| * - If the common type of sample and level is of integral type, the bin for a sample is | |
| * computed as `(sample - lower_level[i]) * (num_levels - 1) / (upper_level[i] - | |
| * lower_level[i])`, round down to the nearest whole number. To protect against potential | |
| * overflows, if, for any channel `i`, the product `(upper_level[i] - lower_level[i]) * | |
| * (num_levels[i] - 1)` exceeds the number representable by an `uint64_t`, the cuda error | |
| * `cudaErrorInvalidValue` is returned. If the common type is 128 bits wide, bin computation | |
| * will use 128-bit arithmetic and `cudaErrorInvalidValue` will only be returned if bin | |
| * computation would overflow for 128-bit arithmetic. | |
| * - For a given row `r` in `[0, num_rows)`, and sample `s` in | |
| * `[0, num_row_pixels)`, let | |
| * `row_begin = d_samples + r * row_stride_bytes / sizeof(SampleT)`, | |
| * `sample_begin = row_begin + s * NUM_CHANNELS`, and | |
| * `sample_end = sample_begin + NUM_ACTIVE_CHANNELS`. For a given channel | |
| * `c` in `[0, NUM_ACTIVE_CHANNELS)`, the ranges | |
| * `[sample_begin, sample_end)` and | |
| * `[d_histogram[c], d_histogram[c] + num_levels[c] - 1)` shall not overlap | |
| * in any way. | |
| * - `cuda::std::common_type<LevelT, SampleT>` must be valid, and both LevelT | |
| * and SampleT must be valid arithmetic types. The common type must be | |
| * convertible to `int` and trivially copyable. | |
| * - @devicestorage | |
| * | |
| * @par Snippet | |
| * The code snippet below illustrates the computation of three 256-bin | |
| * *RGB* histograms from a 2x3 region of interest of within a flattened 2x4 | |
| * array of quad-channel *RGBA* pixels (8 bits per channel per pixel). | |
| * | |
| * @par | |
| * @code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input | |
| * // samples and output histograms | |
| * int num_row_pixels; // e.g., 3 | |
| * int num_rows; // e.g., 2 | |
| * size_t row_stride_bytes; // e.g., 4 * sizeof(unsigned char) * NUM_CHANNELS | |
| * unsigned char* d_samples; // e.g., [(2, 6, 7, 5), (3, 0, 2, 1), (7, 0, 6, 2), (-, -, -, -), | |
| * // (0, 6, 7, 5), (3, 0, 2, 6), (1, 1, 1, 1), (-, -, -, -)] | |
| * int* d_histogram[3]; // e.g., three device pointers to three device buffers, | |
| * // each allocated with 256 integer counters | |
| * int num_levels[3]; // e.g., {257, 257, 257}; | |
| * unsigned int lower_level[3]; // e.g., {0, 0, 0}; | |
| * unsigned int upper_level[3]; // e.g., {256, 256, 256}; | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void* d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceHistogram::MultiHistogramEven<4, 3>( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, lower_level, upper_level, | |
| * num_row_pixels, num_rows, row_stride_bytes); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Compute histograms | |
| * cub::DeviceHistogram::MultiHistogramEven<4, 3>( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, lower_level, upper_level, | |
| * num_row_pixels, num_rows, row_stride_bytes); | |
| * | |
| * // d_histogram <-- [ [1, 1, 1, 2, 0, 0, 0, 1, 0, 0, 0, ..., 0], | |
| * // [0, 4, 0, 0, 0, 0, 2, 0, 0, 0, 0, ..., 0], | |
| * // [0, 1, 2, 0, 0, 0, 1, 2, 0, 0, 0, ..., 0] ] | |
| * @endcode | |
| * | |
| * @tparam NUM_CHANNELS | |
| * Number of channels interleaved in the input data (may be greater than | |
| * the number of channels being actively histogrammed) | |
| * | |
| * @tparam NUM_ACTIVE_CHANNELS | |
| * **[inferred]** Number of channels actively being histogrammed | |
| * | |
| * @tparam SampleIteratorT | |
| * **[inferred]** Random-access input iterator type for reading input | |
| * samples. \iterator | |
| * | |
| * @tparam CounterT | |
| * **[inferred]** Integer type for histogram bin counters | |
| * | |
| * @tparam LevelT | |
| * **[inferred]** Type for specifying boundaries (levels) | |
| * | |
| * @tparam OffsetT | |
| * **[inferred]** Signed integer type for sequence offsets, list lengths, | |
| * pointer differences, etc. \offset_size1 | |
| * | |
| * @param[in] d_temp_storage | |
| * Device-accessible allocation of temporary storage. When `nullptr`, the | |
| * required allocation size is written to `temp_storage_bytes` and no | |
| * work is done. | |
| * | |
| * @param[in,out] temp_storage_bytes | |
| * Reference to size in bytes of `d_temp_storage` allocation | |
| * | |
| * @param[in] d_samples | |
| * The pointer to the multi-channel input sequence of data samples. The | |
| * samples from different channels are assumed to be interleaved (e.g., | |
| * an array of 32-bit pixels where each pixel consists of four | |
| * *RGBA* 8-bit samples). | |
| * | |
| * @param[out] d_histogram | |
| * The pointers to the histogram counter output arrays, one for each | |
| * active channel. For channel<sub><em>i</em></sub>, the allocation length | |
| * of `d_histogram[i]` should be `num_levels[i] - 1`. | |
| * | |
| * @param[in] num_levels | |
| * The number of boundaries (levels) for delineating histogram samples in | |
| * each active channel. Implies that the number of bins for | |
| * channel<sub><em>i</em></sub> is `num_levels[i] - 1`. | |
| * | |
| * @param[in] lower_level | |
| * The lower sample value bound (inclusive) for the lowest histogram bin in | |
| * each active channel. | |
| * | |
| * @param[in] upper_level | |
| * The upper sample value bound (exclusive) for the highest histogram bin | |
| * in each active channel. | |
| * | |
| * @param[in] num_row_pixels | |
| * The number of multi-channel pixels per row in the region of interest | |
| * | |
| * @param[in] num_rows | |
| * The number of rows in the region of interest | |
| * | |
| * @param[in] row_stride_bytes | |
| * The number of bytes between starts of consecutive rows in the region of | |
| * interest | |
| * | |
| * @param[in] stream | |
| * **[optional]** CUDA stream to launch kernels within. | |
| * Default is stream<sub>0</sub>. | |
| */ | |
| template <int NUM_CHANNELS, | |
| int NUM_ACTIVE_CHANNELS, | |
| typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| MultiHistogramEven(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram[NUM_ACTIVE_CHANNELS], | |
| int num_levels[NUM_ACTIVE_CHANNELS], | |
| LevelT lower_level[NUM_ACTIVE_CHANNELS], | |
| LevelT upper_level[NUM_ACTIVE_CHANNELS], | |
| OffsetT num_row_pixels, | |
| OffsetT num_rows, | |
| size_t row_stride_bytes, | |
| cudaStream_t stream = 0) | |
| { | |
| /// The sample value type of the input iterator | |
| using SampleT = cub::detail::value_t<SampleIteratorT>; | |
| Int2Type<sizeof(SampleT) == 1> is_byte_sample; | |
| if ((sizeof(OffsetT) > sizeof(int)) && | |
| ((unsigned long long)(num_rows * row_stride_bytes) < | |
| (unsigned long long)INT_MAX)) | |
| { | |
| // Down-convert OffsetT data type | |
| return DispatchHistogram<NUM_CHANNELS, | |
| NUM_ACTIVE_CHANNELS, | |
| SampleIteratorT, | |
| CounterT, | |
| LevelT, | |
| int>::DispatchEven(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| lower_level, | |
| upper_level, | |
| (int)num_row_pixels, | |
| (int)num_rows, | |
| (int)(row_stride_bytes / | |
| sizeof(SampleT)), | |
| stream, | |
| is_byte_sample); | |
| } | |
| return DispatchHistogram<NUM_CHANNELS, | |
| NUM_ACTIVE_CHANNELS, | |
| SampleIteratorT, | |
| CounterT, | |
| LevelT, | |
| OffsetT>::DispatchEven(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| lower_level, | |
| upper_level, | |
| num_row_pixels, | |
| num_rows, | |
| (OffsetT)(row_stride_bytes / | |
| sizeof(SampleT)), | |
| stream, | |
| is_byte_sample); | |
| } | |
| template <int NUM_CHANNELS, | |
| int NUM_ACTIVE_CHANNELS, | |
| typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| MultiHistogramEven(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram[NUM_ACTIVE_CHANNELS], | |
| int num_levels[NUM_ACTIVE_CHANNELS], | |
| LevelT lower_level[NUM_ACTIVE_CHANNELS], | |
| LevelT upper_level[NUM_ACTIVE_CHANNELS], | |
| OffsetT num_row_pixels, | |
| OffsetT num_rows, | |
| size_t row_stride_bytes, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return MultiHistogramEven(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| lower_level, | |
| upper_level, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_bytes, | |
| stream); | |
| } | |
| //@} end member group | |
| /******************************************************************//** | |
| * @name Custom bin ranges | |
| *********************************************************************/ | |
| //@{ | |
| /** | |
| * @brief Computes an intensity histogram from a sequence of data samples | |
| * using the specified bin boundary levels. | |
| * | |
| * @par | |
| * - The number of histogram bins is (`num_levels - 1`) | |
| * - The value range for bin<sub><em>i</em></sub> is `[level[i], level[i+1])` | |
| * - The range `[d_histogram, d_histogram + num_levels - 1)` shall not | |
| * overlap `[d_samples, d_samples + num_samples)` nor | |
| * `[d_levels, d_levels + num_levels)` in any way. The ranges | |
| * `[d_levels, d_levels + num_levels)` and | |
| * `[d_samples, d_samples + num_samples)` may overlap. | |
| * - @devicestorage | |
| * | |
| * @par Snippet | |
| * The code snippet below illustrates the computation of an six-bin histogram | |
| * from a sequence of float samples | |
| * | |
| * @par | |
| * @code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input | |
| * // samples and output histogram | |
| * int num_samples; // e.g., 10 | |
| * float* d_samples; // e.g., [2.2, 6.0, 7.1, 2.9, 3.5, 0.3, 2.9, 2.0, 6.1, 999.5] | |
| * int* d_histogram; // e.g., [ -, -, -, -, -, -] | |
| * int num_levels // e.g., 7 (seven level boundaries for six bins) | |
| * float* d_levels; // e.g., [0.0, 2.0, 4.0, 6.0, 8.0, 12.0, 16.0] | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void* d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceHistogram::HistogramRange( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, d_levels, num_samples); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Compute histograms | |
| * cub::DeviceHistogram::HistogramRange( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, d_levels, num_samples); | |
| * | |
| * // d_histogram <-- [1, 5, 0, 3, 0, 0]; | |
| * | |
| * @endcode | |
| * | |
| * @tparam SampleIteratorT | |
| * **[inferred]** Random-access input iterator type for reading | |
| * input samples.\iterator | |
| * | |
| * @tparam CounterT | |
| * **[inferred]** Integer type for histogram bin counters | |
| * | |
| * @tparam LevelT | |
| * **[inferred]** Type for specifying boundaries (levels) | |
| * | |
| * @tparam OffsetT | |
| * **[inferred]** Signed integer type for sequence offsets, list lengths, | |
| * pointer differences, etc. \offset_size1 | |
| * | |
| * @param[in] d_temp_storage | |
| * Device-accessible allocation of temporary storage. When `nullptr`, the | |
| * required allocation size is written to `temp_storage_bytes` and no work | |
| * is done. | |
| * | |
| * @param[in,out] temp_storage_bytes | |
| * Reference to size in bytes of `d_temp_storage` allocation | |
| * | |
| * @param[in] d_samples | |
| * The pointer to the input sequence of data samples. | |
| * | |
| * @param[out] d_histogram | |
| * The pointer to the histogram counter output array of length | |
| * `num_levels - 1`. | |
| * | |
| * @param[in] num_levels | |
| * The number of boundaries (levels) for delineating histogram samples. | |
| * Implies that the number of bins is `num_levels - 1`. | |
| * | |
| * @param[in] d_levels | |
| * The pointer to the array of boundaries (levels). Bin ranges are defined | |
| * by consecutive boundary pairings: lower sample value boundaries are | |
| * inclusive and upper sample value boundaries are exclusive. | |
| * | |
| * @param[in] num_samples | |
| * The number of data samples per row in the region of interest | |
| * | |
| * @param[in] stream | |
| * **[optional]** CUDA stream to launch kernels within. | |
| * Default is stream<sub>0</sub>. | |
| */ | |
| template <typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| HistogramRange(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram, | |
| int num_levels, | |
| LevelT *d_levels, | |
| OffsetT num_samples, | |
| cudaStream_t stream = 0) | |
| { | |
| /// The sample value type of the input iterator | |
| using SampleT = cub::detail::value_t<SampleIteratorT>; | |
| CounterT *d_histogram1[1] = {d_histogram}; | |
| int num_levels1[1] = {num_levels}; | |
| LevelT *d_levels1[1] = {d_levels}; | |
| return MultiHistogramRange<1, 1>(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram1, | |
| num_levels1, | |
| d_levels1, | |
| num_samples, | |
| (OffsetT)1, | |
| (size_t)(sizeof(SampleT) * num_samples), | |
| stream); | |
| } | |
| template <typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| HistogramRange(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram, | |
| int num_levels, | |
| LevelT *d_levels, | |
| OffsetT num_samples, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return HistogramRange(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| d_levels, | |
| num_samples, | |
| stream); | |
| } | |
| /** | |
| * @brief Computes an intensity histogram from a sequence of data samples | |
| * using the specified bin boundary levels. | |
| * | |
| * @par | |
| * - A two-dimensional *region of interest* within `d_samples` can be | |
| * specified using the `num_row_samples`, `num_rows`, and | |
| * `row_stride_bytes` parameters. | |
| * - The row stride must be a whole multiple of the sample data type | |
| * size, i.e., `(row_stride_bytes % sizeof(SampleT)) == 0`. | |
| * - The number of histogram bins is (`num_levels - 1`) | |
| * - The value range for bin<sub><em>i</em></sub> is `[level[i], level[i+1])` | |
| * - For a given row `r` in `[0, num_rows)`, let | |
| * `row_begin = d_samples + r * row_stride_bytes / sizeof(SampleT)` and | |
| * `row_end = row_begin + num_row_samples`. The range | |
| * `[d_histogram, d_histogram + num_levels - 1)` shall not overlap | |
| * `[row_begin, row_end)` nor `[d_levels, d_levels + num_levels)`. | |
| * The ranges `[d_levels, d_levels + num_levels)` and `[row_begin, row_end)` | |
| * may overlap. | |
| * - @devicestorage | |
| * | |
| * @par Snippet | |
| * The code snippet below illustrates the computation of a six-bin histogram | |
| * from a 2x5 region of interest within a flattened 2x7 array of float samples. | |
| * | |
| * @par | |
| * @code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input samples and | |
| * // output histogram | |
| * int num_row_samples; // e.g., 5 | |
| * int num_rows; // e.g., 2; | |
| * int row_stride_bytes; // e.g., 7 * sizeof(float) | |
| * float* d_samples; // e.g., [2.2, 6.0, 7.1, 2.9, 3.5, -, -, | |
| * // 0.3, 2.9, 2.0, 6.1, 999.5, -, -] | |
| * int* d_histogram; // e.g., [ -, -, -, -, -, -] | |
| * int num_levels // e.g., 7 (seven level boundaries for six bins) | |
| * float *d_levels; // e.g., [0.0, 2.0, 4.0, 6.0, 8.0, 12.0, 16.0] | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void* d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceHistogram::HistogramRange( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, d_levels, | |
| * num_row_samples, num_rows, row_stride_bytes); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Compute histograms | |
| * cub::DeviceHistogram::HistogramRange( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, d_levels, | |
| * num_row_samples, num_rows, row_stride_bytes); | |
| * | |
| * // d_histogram <-- [1, 5, 0, 3, 0, 0]; | |
| * @endcode | |
| * | |
| * @tparam SampleIteratorT | |
| * **[inferred]** Random-access input iterator type for reading | |
| * input samples. \iterator | |
| * | |
| * @tparam CounterT | |
| * **[inferred]** Integer type for histogram bin counters | |
| * | |
| * @tparam LevelT | |
| * **[inferred]** Type for specifying boundaries (levels) | |
| * | |
| * @tparam OffsetT | |
| * **[inferred]** Signed integer type for sequence offsets, list lengths, | |
| * pointer differences, etc. \offset_size1 | |
| * | |
| * @param[in] d_temp_storage | |
| * Device-accessible allocation of temporary storage. When `nullptr`, the | |
| * required allocation size is written to `temp_storage_bytes` and no | |
| * work is done. | |
| * | |
| * @param[in,out] temp_storage_bytes | |
| * Reference to size in bytes of `d_temp_storage` allocation | |
| * | |
| * @param[in] d_samples | |
| * The pointer to the input sequence of data samples. | |
| * | |
| * @param[out] d_histogram | |
| * The pointer to the histogram counter output array of length | |
| * `num_levels - 1`. | |
| * | |
| * @param[in] num_levels | |
| * The number of boundaries (levels) for delineating histogram samples. | |
| * Implies that the number of bins is `num_levels - 1`. | |
| * | |
| * @param[in] d_levels | |
| * The pointer to the array of boundaries (levels). Bin ranges are defined | |
| * by consecutive boundary pairings: lower sample value boundaries are | |
| * inclusive and upper sample value boundaries are exclusive. | |
| * | |
| * @param[in] num_row_samples | |
| * The number of data samples per row in the region of interest | |
| * | |
| * @param[in] num_rows | |
| * The number of rows in the region of interest | |
| * | |
| * @param[in] row_stride_bytes | |
| * The number of bytes between starts of consecutive rows in the region | |
| * of interest | |
| * | |
| * @param[in] stream | |
| * **[optional]** CUDA stream to launch kernels within. | |
| * Default is stream<sub>0</sub>. | |
| */ | |
| template <typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| HistogramRange(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram, | |
| int num_levels, | |
| LevelT *d_levels, | |
| OffsetT num_row_samples, | |
| OffsetT num_rows, | |
| size_t row_stride_bytes, | |
| cudaStream_t stream = 0) | |
| { | |
| CounterT *d_histogram1[1] = {d_histogram}; | |
| int num_levels1[1] = {num_levels}; | |
| LevelT *d_levels1[1] = {d_levels}; | |
| return MultiHistogramRange<1, 1>(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram1, | |
| num_levels1, | |
| d_levels1, | |
| num_row_samples, | |
| num_rows, | |
| row_stride_bytes, | |
| stream); | |
| } | |
| template <typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| HistogramRange(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram, | |
| int num_levels, | |
| LevelT *d_levels, | |
| OffsetT num_row_samples, | |
| OffsetT num_rows, | |
| size_t row_stride_bytes, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return HistogramRange(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| d_levels, | |
| num_row_samples, | |
| num_rows, | |
| row_stride_bytes, | |
| stream); | |
| } | |
| /** | |
| * @brief Computes per-channel intensity histograms from a sequence of | |
| * multi-channel "pixel" data samples using the specified bin | |
| * boundary levels. | |
| * | |
| * @par | |
| * - The input is a sequence of *pixel* structures, where each pixel | |
| * comprises a record of `NUM_CHANNELS` consecutive data samples | |
| * (e.g., an *RGBA* pixel). | |
| * - Of the `NUM_CHANNELS` specified, the function will only compute | |
| * histograms for the first `NUM_ACTIVE_CHANNELS` (e.g., *RGB* histograms | |
| * from *RGBA* pixel samples). | |
| * - The number of histogram bins for channel<sub><em>i</em></sub> is | |
| * `num_levels[i] - 1`. | |
| * - For channel<sub><em>i</em></sub>, the range of values for all histogram | |
| * bins have the same width: | |
| * `(upper_level[i] - lower_level[i]) / (num_levels[i] - 1)` | |
| * - For given channels `c1` and `c2` in `[0, NUM_ACTIVE_CHANNELS)`, the | |
| * range `[d_histogram[c1], d_histogram[c1] + num_levels[c1] - 1)` shall | |
| * not overlap `[d_samples, d_samples + NUM_CHANNELS * num_pixels)` nor | |
| * `[d_levels[c2], d_levels[c2] + num_levels[c2])` in any way. | |
| * The ranges `[d_levels[c2], d_levels[c2] + num_levels[c2])` and | |
| * `[d_samples, d_samples + NUM_CHANNELS * num_pixels)` may overlap. | |
| * - @devicestorage | |
| * | |
| * @par Snippet | |
| * The code snippet below illustrates the computation of three 4-bin *RGB* | |
| * histograms from a quad-channel sequence of *RGBA* pixels | |
| * (8 bits per channel per pixel) | |
| * | |
| * @par | |
| * @code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for | |
| * // input samples and output histograms | |
| * int num_pixels; // e.g., 5 | |
| * unsigned char *d_samples; // e.g., [(2, 6, 7, 5),(3, 0, 2, 1),(7, 0, 6, 2), | |
| * // (0, 6, 7, 5),(3, 0, 2, 6)] | |
| * unsigned int *d_histogram[3]; // e.g., [[ -, -, -, -],[ -, -, -, -],[ -, -, -, -]]; | |
| * int num_levels[3]; // e.g., {5, 5, 5}; | |
| * unsigned int *d_levels[3]; // e.g., [ [0, 2, 4, 6, 8], | |
| * // [0, 2, 4, 6, 8], | |
| * // [0, 2, 4, 6, 8] ]; | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void* d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceHistogram::MultiHistogramRange<4, 3>( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, d_levels, num_pixels); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Compute histograms | |
| * cub::DeviceHistogram::MultiHistogramRange<4, 3>( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, d_levels, num_pixels); | |
| * | |
| * // d_histogram <-- [ [1, 3, 0, 1], | |
| * // [3, 0, 0, 2], | |
| * // [0, 2, 0, 3] ] | |
| * | |
| * @endcode | |
| * | |
| * @tparam NUM_CHANNELS | |
| * Number of channels interleaved in the input data (may be greater than | |
| * the number of channels being actively histogrammed) | |
| * | |
| * @tparam NUM_ACTIVE_CHANNELS | |
| * **[inferred]** Number of channels actively being histogrammed | |
| * | |
| * @tparam SampleIteratorT | |
| * **[inferred]** Random-access input iterator type for reading | |
| * input samples. \iterator | |
| * | |
| * @tparam CounterT | |
| * **[inferred]** Integer type for histogram bin counters | |
| * | |
| * @tparam LevelT | |
| * **[inferred]** Type for specifying boundaries (levels) | |
| * | |
| * @tparam OffsetT | |
| * **[inferred]** Signed integer type for sequence offsets, list lengths, | |
| * pointer differences, etc. \offset_size1 | |
| * | |
| * @param[in] d_temp_storage | |
| * Device-accessible allocation of temporary storage. When `nullptr`, the | |
| * required allocation size is written to `temp_storage_bytes` and no | |
| * work is done. | |
| * | |
| * @param[in,out] temp_storage_bytes | |
| * Reference to size in bytes of `d_temp_storage` allocation | |
| * | |
| * @param[in] d_samples | |
| * The pointer to the multi-channel input sequence of data samples. | |
| * The samples from different channels are assumed to be interleaved (e.g., | |
| * an array of 32-bit pixels where each pixel consists of four *RGBA* | |
| * 8-bit samples). | |
| * | |
| * @param[out] d_histogram | |
| * The pointers to the histogram counter output arrays, one for each active | |
| * channel. For channel<sub><em>i</em></sub>, the allocation length of | |
| * `d_histogram[i]` should be `num_levels[i] - 1`. | |
| * | |
| * @param[in] num_levels | |
| * The number of boundaries (levels) for delineating histogram samples in | |
| * each active channel. Implies that the number of bins for | |
| * channel<sub><em>i</em></sub> is `num_levels[i] - 1`. | |
| * | |
| * @param[in] d_levels | |
| * The pointers to the arrays of boundaries (levels), one for each active | |
| * channel. Bin ranges are defined by consecutive boundary pairings: lower | |
| * sample value boundaries are inclusive and upper sample value boundaries | |
| * are exclusive. | |
| * | |
| * @param[in] num_pixels | |
| * The number of multi-channel pixels | |
| * (i.e., the length of `d_samples / NUM_CHANNELS`) | |
| * | |
| * @param[in] stream | |
| * **[optional]** CUDA stream to launch kernels within. | |
| * Default is stream<sub>0</sub>. | |
| */ | |
| template <int NUM_CHANNELS, | |
| int NUM_ACTIVE_CHANNELS, | |
| typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| MultiHistogramRange(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram[NUM_ACTIVE_CHANNELS], | |
| int num_levels[NUM_ACTIVE_CHANNELS], | |
| LevelT *d_levels[NUM_ACTIVE_CHANNELS], | |
| OffsetT num_pixels, | |
| cudaStream_t stream = 0) | |
| { | |
| /// The sample value type of the input iterator | |
| using SampleT = cub::detail::value_t<SampleIteratorT>; | |
| return MultiHistogramRange<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>( | |
| d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| d_levels, | |
| num_pixels, | |
| (OffsetT)1, | |
| (size_t)(sizeof(SampleT) * NUM_CHANNELS * num_pixels), | |
| stream); | |
| } | |
| template <int NUM_CHANNELS, | |
| int NUM_ACTIVE_CHANNELS, | |
| typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| MultiHistogramRange(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram[NUM_ACTIVE_CHANNELS], | |
| int num_levels[NUM_ACTIVE_CHANNELS], | |
| LevelT *d_levels[NUM_ACTIVE_CHANNELS], | |
| OffsetT num_pixels, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return MultiHistogramRange(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| d_levels, | |
| num_pixels, | |
| stream); | |
| } | |
| /** | |
| * @brief Computes per-channel intensity histograms from a sequence of | |
| * multi-channel "pixel" data samples using the specified bin boundary | |
| * levels. | |
| * | |
| * @par | |
| * - The input is a sequence of *pixel* structures, where each pixel comprises | |
| * a record of `NUM_CHANNELS` consecutive data samples | |
| * (e.g., an *RGBA* pixel). | |
| * - Of the `NUM_CHANNELS` specified, the function will only compute | |
| * histograms for the first `NUM_ACTIVE_CHANNELS` (e.g., *RGB* histograms | |
| * from *RGBA* pixel samples). | |
| * - A two-dimensional *region of interest* within `d_samples` can be | |
| * specified using the `num_row_samples`, `num_rows`, and `row_stride_bytes` | |
| * parameters. | |
| * - The row stride must be a whole multiple of the sample data type | |
| * size, i.e., `(row_stride_bytes % sizeof(SampleT)) == 0`. | |
| * - The number of histogram bins for channel<sub><em>i</em></sub> is | |
| * `num_levels[i] - 1`. | |
| * - For channel<sub><em>i</em></sub>, the range of values for all histogram | |
| * bins have the same width: | |
| * `(upper_level[i] - lower_level[i]) / (num_levels[i] - 1)` | |
| * - For a given row `r` in `[0, num_rows)`, and sample `s` in | |
| * `[0, num_row_pixels)`, let | |
| * `row_begin = d_samples + r * row_stride_bytes / sizeof(SampleT)`, | |
| * `sample_begin = row_begin + s * NUM_CHANNELS`, and | |
| * `sample_end = sample_begin + NUM_ACTIVE_CHANNELS`. For given channels | |
| * `c1` and `c2` in `[0, NUM_ACTIVE_CHANNELS)`, the range | |
| * `[d_histogram[c1], d_histogram[c1] + num_levels[c1] - 1)` shall not | |
| * overlap `[sample_begin, sample_end)` nor | |
| * `[d_levels[c2], d_levels[c2] + num_levels[c2])` in any way. The ranges | |
| * `[d_levels[c2], d_levels[c2] + num_levels[c2])` and | |
| * `[sample_begin, sample_end)` may overlap. | |
| * - @devicestorage | |
| * | |
| * @par Snippet | |
| * The code snippet below illustrates the computation of three 4-bin *RGB* | |
| * histograms from a 2x3 region of interest of within a flattened 2x4 array | |
| * of quad-channel *RGBA* pixels (8 bits per channel per pixel). | |
| * | |
| * @par | |
| * @code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input | |
| * // samples and output histograms | |
| * int num_row_pixels; // e.g., 3 | |
| * int num_rows; // e.g., 2 | |
| * size_t row_stride_bytes; // e.g., 4 * sizeof(unsigned char) * NUM_CHANNELS | |
| * unsigned char* d_samples; // e.g., [(2, 6, 7, 5),(3, 0, 2, 1),(1, 1, 1, 1),(-, -, -, -), | |
| * // (7, 0, 6, 2),(0, 6, 7, 5),(3, 0, 2, 6),(-, -, -, -)] | |
| * int* d_histogram[3]; // e.g., [[ -, -, -, -],[ -, -, -, -],[ -, -, -, -]]; | |
| * int num_levels[3]; // e.g., {5, 5, 5}; | |
| * unsigned int* d_levels[3]; // e.g., [ [0, 2, 4, 6, 8], | |
| * // [0, 2, 4, 6, 8], | |
| * // [0, 2, 4, 6, 8] ]; | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void* d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceHistogram::MultiHistogramRange<4, 3>( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, d_levels, | |
| * num_row_pixels, num_rows, row_stride_bytes); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Compute histograms | |
| * cub::DeviceHistogram::MultiHistogramRange<4, 3>( | |
| * d_temp_storage, temp_storage_bytes, | |
| * d_samples, d_histogram, num_levels, | |
| * d_levels, num_row_pixels, num_rows, row_stride_bytes); | |
| * | |
| * // d_histogram <-- [ [2, 3, 0, 1], | |
| * // [3, 0, 0, 2], | |
| * // [1, 2, 0, 3] ] | |
| * | |
| * @endcode | |
| * | |
| * @tparam NUM_CHANNELS | |
| * Number of channels interleaved in the input data (may be greater than | |
| * the number of channels being actively histogrammed) | |
| * | |
| * @tparam NUM_ACTIVE_CHANNELS | |
| * **[inferred]** Number of channels actively being histogrammed | |
| * | |
| * @tparam SampleIteratorT | |
| * **[inferred]** Random-access input iterator type for reading input | |
| * samples. \iterator | |
| * | |
| * @tparam CounterT | |
| * **[inferred]** Integer type for histogram bin counters | |
| * | |
| * @tparam LevelT | |
| * **[inferred]** Type for specifying boundaries (levels) | |
| * | |
| * @tparam OffsetT | |
| * **[inferred]** Signed integer type for sequence offsets, list lengths, | |
| * pointer differences, etc. \offset_size1 | |
| * | |
| * @param[in] d_temp_storage | |
| * Device-accessible allocation of temporary storage. When `nullptr`, the | |
| * required allocation size is written to \p temp_storage_bytes and no work is done. | |
| * | |
| * @param[in,out] temp_storage_bytes | |
| * Reference to size in bytes of `d_temp_storage` allocation | |
| * | |
| * @param[in] d_samples | |
| * The pointer to the multi-channel input sequence of data samples. The | |
| * samples from different channels are assumed to be interleaved (e.g., an | |
| * array of 32-bit pixels where each pixel consists of four | |
| * *RGBA* 8-bit samples). | |
| * | |
| * @param[out] d_histogram | |
| * The pointers to the histogram counter output arrays, one for each active | |
| * channel. For channel<sub><em>i</em></sub>, the allocation length of | |
| * `d_histogram[i]` should be `num_levels[i] - 1`. | |
| * | |
| * @param[in] num_levels | |
| * The number of boundaries (levels) for delineating histogram samples in | |
| * each active channel. Implies that the number of bins for | |
| * channel<sub><em>i</em></sub> is `num_levels[i] - 1`. | |
| * | |
| * @param[in] d_levels | |
| * The pointers to the arrays of boundaries (levels), one for each active | |
| * channel. Bin ranges are defined by consecutive boundary pairings: lower | |
| * sample value boundaries are inclusive and upper sample value boundaries | |
| * are exclusive. | |
| * | |
| * @param[in] num_row_pixels | |
| * The number of multi-channel pixels per row in the region of interest | |
| * | |
| * @param[in] num_rows | |
| * The number of rows in the region of interest | |
| * | |
| * @param[in] row_stride_bytes | |
| * The number of bytes between starts of consecutive rows in the | |
| * region of interest | |
| * | |
| * @param[in] stream | |
| * **[optional]** CUDA stream to launch kernels within. | |
| * Default is stream<sub>0</sub>. | |
| */ | |
| template <int NUM_CHANNELS, | |
| int NUM_ACTIVE_CHANNELS, | |
| typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| MultiHistogramRange(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram[NUM_ACTIVE_CHANNELS], | |
| int num_levels[NUM_ACTIVE_CHANNELS], | |
| LevelT *d_levels[NUM_ACTIVE_CHANNELS], | |
| OffsetT num_row_pixels, | |
| OffsetT num_rows, | |
| size_t row_stride_bytes, | |
| cudaStream_t stream = 0) | |
| { | |
| /// The sample value type of the input iterator | |
| using SampleT = cub::detail::value_t<SampleIteratorT>; | |
| Int2Type<sizeof(SampleT) == 1> is_byte_sample; | |
| if ((sizeof(OffsetT) > sizeof(int)) && | |
| ((unsigned long long)(num_rows * row_stride_bytes) < | |
| (unsigned long long)INT_MAX)) | |
| { | |
| // Down-convert OffsetT data type | |
| return DispatchHistogram<NUM_CHANNELS, | |
| NUM_ACTIVE_CHANNELS, | |
| SampleIteratorT, | |
| CounterT, | |
| LevelT, | |
| int>::DispatchRange(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| d_levels, | |
| (int)num_row_pixels, | |
| (int)num_rows, | |
| (int)(row_stride_bytes / | |
| sizeof(SampleT)), | |
| stream, | |
| is_byte_sample); | |
| } | |
| return DispatchHistogram<NUM_CHANNELS, | |
| NUM_ACTIVE_CHANNELS, | |
| SampleIteratorT, | |
| CounterT, | |
| LevelT, | |
| OffsetT>::DispatchRange(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| d_levels, | |
| num_row_pixels, | |
| num_rows, | |
| (OffsetT)(row_stride_bytes / | |
| sizeof(SampleT)), | |
| stream, | |
| is_byte_sample); | |
| } | |
| template <int NUM_CHANNELS, | |
| int NUM_ACTIVE_CHANNELS, | |
| typename SampleIteratorT, | |
| typename CounterT, | |
| typename LevelT, | |
| typename OffsetT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTION static cudaError_t | |
| MultiHistogramRange(void *d_temp_storage, | |
| size_t &temp_storage_bytes, | |
| SampleIteratorT d_samples, | |
| CounterT *d_histogram[NUM_ACTIVE_CHANNELS], | |
| int num_levels[NUM_ACTIVE_CHANNELS], | |
| LevelT *d_levels[NUM_ACTIVE_CHANNELS], | |
| OffsetT num_row_pixels, | |
| OffsetT num_rows, | |
| size_t row_stride_bytes, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return MultiHistogramRange(d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_histogram, | |
| num_levels, | |
| d_levels, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_bytes, | |
| stream); | |
| } | |
| //@} end member group | |
| }; | |
| CUB_NAMESPACE_END | |