histogram_cub.h
4.27 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
/******************************************************************************
* 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 <cub/device/device_histogram.cuh>
using namespace cub;
template <
int NUM_CHANNELS,
int ACTIVE_CHANNELS,
int NUM_BINS,
typename PixelType>
double run_cub_histogram(
PixelType *d_image,
int width,
int height,
unsigned int *d_hist,
bool is_warmup)
{
enum {
is_float = Equals<PixelType, float4>::VALUE,
};
typedef typename If<is_float, float, unsigned char>::Type SampleT; // Sample type
typedef typename If<is_float, float, unsigned int>::Type LevelT; // 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 channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 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<NUM_CHANNELS, ACTIVE_CHANNELS>(
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<NUM_CHANNELS, ACTIVE_CHANNELS>(
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;
}