1 /******************************************************************************
2 * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Redistribution and use in source and binary forms, with or without
5 * modification, are permitted provided that the following conditions are met:
6 * * Redistributions of source code must retain the above copyright
7 * notice, this list of conditions and the following disclaimer.
8 * * Redistributions in binary form must reproduce the above copyright
9 * notice, this list of conditions and the following disclaimer in the
10 * documentation and/or other materials provided with the distribution.
11 * * Neither the name of the NVIDIA CORPORATION nor the
12 * names of its contributors may be used to endorse or promote products
13 * derived from this software without specific prior written permission.
14 *
15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
16 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
17 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
18 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
19 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
20 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
21 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
22 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
23 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
24 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
25 *
26 ******************************************************************************/
27
28 #include <cub/device/device_histogram.cuh>
29
30 using namespace cub;
31
32 template <
33 int NUM_CHANNELS,
34 int ACTIVE_CHANNELS,
35 int NUM_BINS,
36 typename PixelType>
run_cub_histogram(PixelType * d_image,int width,int height,unsigned int * d_hist,bool is_warmup)37 double run_cub_histogram(
38 PixelType *d_image,
39 int width,
40 int height,
41 unsigned int *d_hist,
42 bool is_warmup)
43 {
44 enum {
45 is_float = Equals<PixelType, float4>::VALUE,
46 };
47
48 typedef typename If<is_float, float, unsigned char>::Type SampleT; // Sample type
49 typedef typename If<is_float, float, unsigned int>::Type LevelT; // Level type (uint32 for uchar)
50
51 // Setup data structures
52 unsigned int* d_histogram[ACTIVE_CHANNELS];
53 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 channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
54 LevelT lower_level[ACTIVE_CHANNELS]; ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
55 LevelT upper_level[ACTIVE_CHANNELS]; ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
56
57 for (int CHANNEL = 0; CHANNEL < ACTIVE_CHANNELS; ++CHANNEL)
58 {
59 d_histogram[CHANNEL] = d_hist + (CHANNEL * NUM_BINS);
60 num_levels[CHANNEL] = NUM_BINS + 1;
61 lower_level[CHANNEL] = 0;
62 upper_level[CHANNEL] = (is_float) ? 1 : 256;
63 }
64
65 // Allocate temporary storage
66 size_t temp_storage_bytes = 0;
67 void *d_temp_storage = NULL;
68
69 SampleT* d_image_samples = (SampleT*) d_image;
70
71 // Get amount of temporary storage needed
72 DeviceHistogram::MultiHistogramEven<NUM_CHANNELS, ACTIVE_CHANNELS>(
73 d_temp_storage,
74 temp_storage_bytes,
75 d_image_samples,
76 d_histogram,
77 num_levels,
78 lower_level,
79 upper_level,
80 width * height,
81 (cudaStream_t) 0,
82 is_warmup);
83
84 cudaMalloc(&d_temp_storage, temp_storage_bytes);
85
86 GpuTimer gpu_timer;
87 gpu_timer.Start();
88
89 // Compute histogram
90 DeviceHistogram::MultiHistogramEven<NUM_CHANNELS, ACTIVE_CHANNELS>(
91 d_temp_storage,
92 temp_storage_bytes,
93 d_image_samples,
94 d_histogram,
95 num_levels,
96 lower_level,
97 upper_level,
98 width * height,
99 (cudaStream_t) 0,
100 is_warmup);
101
102 gpu_timer.Stop();
103 float elapsed_millis = gpu_timer.ElapsedMillis();
104
105 cudaFree(d_temp_storage);
106
107 return elapsed_millis;
108 }
109
110