/****************************************************************************** * Copyright (c) 2011-2018, 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. * ******************************************************************************/ #include using namespace cub; template double run_cub_histogram( PixelType *d_image, int width, int height, unsigned int *d_hist, bool is_warmup) { constexpr bool is_float = std::is_same::value; using SampleT = cub::detail::conditional_t; // Sample type using LevelT = cub::detail::conditional_t; // Level type (uint32 for uchar) // Setup data structures unsigned int* d_histogram[ACTIVE_CHANNELS]; int num_levels[ACTIVE_CHANNELS]; ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channeli is num_levels[i] - 1. LevelT lower_level[ACTIVE_CHANNELS]; ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel. LevelT upper_level[ACTIVE_CHANNELS]; ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel. for (int CHANNEL = 0; CHANNEL < ACTIVE_CHANNELS; ++CHANNEL) { d_histogram[CHANNEL] = d_hist + (CHANNEL * NUM_BINS); num_levels[CHANNEL] = NUM_BINS + 1; lower_level[CHANNEL] = 0; upper_level[CHANNEL] = (is_float) ? 1 : 256; } // Allocate temporary storage size_t temp_storage_bytes = 0; void *d_temp_storage = NULL; SampleT* d_image_samples = (SampleT*) d_image; // Get amount of temporary storage needed DeviceHistogram::MultiHistogramEven( d_temp_storage, temp_storage_bytes, d_image_samples, d_histogram, num_levels, lower_level, upper_level, width * height, (cudaStream_t) 0, is_warmup); cudaMalloc(&d_temp_storage, temp_storage_bytes); GpuTimer gpu_timer; gpu_timer.Start(); // Compute histogram DeviceHistogram::MultiHistogramEven( d_temp_storage, temp_storage_bytes, d_image_samples, d_histogram, num_levels, lower_level, upper_level, width * height, (cudaStream_t) 0, is_warmup); gpu_timer.Stop(); float elapsed_millis = gpu_timer.ElapsedMillis(); cudaFree(d_temp_storage); return elapsed_millis; }