1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013-2014 Mageswaran.D <mageswaran1989@gmail.com>
3 //
4 // Distributed under the Boost Software License, Version 1.0
5 // See accompanying file LICENSE_1_0.txt or copy at
6 // http://www.boost.org/LICENSE_1_0.txt
7 //
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
10
11 //Code sample for calculating histogram using OpenCL and
12 //displaying image histogram in OpenCV.
13
14 #include <iostream>
15 #include <string>
16
17 #include <opencv2/imgproc/imgproc.hpp>
18 #include <opencv2/highgui/highgui.hpp>
19
20 #include <boost/compute/source.hpp>
21 #include <boost/compute/system.hpp>
22 #include <boost/compute/container/vector.hpp>
23 #include <boost/compute/interop/opencv/core.hpp>
24 #include <boost/compute/interop/opencv/highgui.hpp>
25 #include <boost/program_options.hpp>
26
27 namespace compute = boost::compute;
28 namespace po = boost::program_options;
29
30 // number of bins
31 int histSize = 256;
32
33 // Set the ranges ( for B,G,R) )
34 // TryOut: consider the range in kernel calculation
35 float range[] = { 0, 256 } ;
36 const float* histRange = { range };
37
38 // Create naive histogram program
39 // Needs "cl_khr_local_int32_base_atomics" extension
40 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE (
41 __kernel void histogram(read_only image2d_t src_image,
42 __global int* b_hist,
43 __global int* g_hist,
44 __global int* r_hist)
45 {
46 sampler_t sampler =( CLK_NORMALIZED_COORDS_FALSE |
47 CLK_FILTER_NEAREST |
48 CLK_ADDRESS_CLAMP_TO_EDGE);
49
50 int image_width = get_image_width(src_image);
51 int image_height = get_image_height(src_image);
52
53 int2 coords = (int2)(get_global_id(0), get_global_id(1));
54 float4 pixel = read_imagef(src_image,sampler, coords);
55
56 //boundary condition
57 if ((coords.x < image_width) && (coords.y < image_height))
58 {
59 uchar indx_x, indx_y, indx_z;
60 indx_x = convert_uchar_sat(pixel.x * 255.0f);
61 indx_y = convert_uchar_sat(pixel.y * 255.0f);
62 indx_z = convert_uchar_sat(pixel.z * 255.0f);
63
64 atomic_inc(&b_hist[(uint)indx_z]);
65 atomic_inc(&g_hist[(uint)indx_y]);
66 atomic_inc(&r_hist[(uint)indx_x]);
67 }
68 }
69 );
70
showHistogramWindow(cv::Mat & b_hist,cv::Mat & g_hist,cv::Mat & r_hist,std::string window_name)71 inline void showHistogramWindow(cv::Mat &b_hist, cv::Mat &g_hist, cv::Mat &r_hist,
72 std::string window_name)
73 {
74 // Draw the histograms for B, G and R
75 int hist_w = 1024;
76 int hist_h = 768;
77 int bin_w = cvRound((double)hist_w/histSize);
78
79 cv::Mat histImage(hist_h, hist_w, CV_8UC3, cv::Scalar(0,0,0));
80
81 // Normalize the result to [ 0, histImage.rows ]
82 cv::normalize(b_hist, b_hist, 0, histImage.rows, cv::NORM_MINMAX, -1, cv::Mat());
83 cv::normalize(g_hist, g_hist, 0, histImage.rows, cv::NORM_MINMAX, -1, cv::Mat());
84 cv::normalize(r_hist, r_hist, 0, histImage.rows, cv::NORM_MINMAX, -1, cv::Mat());
85
86 // Draw for each channel
87 for (int i = 1; i < histSize; i++ )
88 {
89 cv::line(histImage,
90 cv::Point(bin_w*(i-1), hist_h - cvRound(b_hist.at<float>(i-1))),
91 cv::Point(bin_w*(i), hist_h - cvRound(b_hist.at<float>(i))),
92 cv::Scalar(255, 0, 0),
93 2,
94 8,
95 0);
96
97 cv::line(histImage,
98 cv::Point(bin_w*(i-1), hist_h - cvRound(g_hist.at<float>(i-1))),
99 cv::Point(bin_w*(i), hist_h - cvRound(g_hist.at<float>(i))),
100 cv::Scalar(0, 255, 0),
101 2,
102 8,
103 0);
104
105 cv::line(histImage,
106 cv::Point( bin_w*(i-1), hist_h - cvRound(r_hist.at<float>(i-1))),
107 cv::Point( bin_w*(i), hist_h - cvRound(r_hist.at<float>(i)) ),
108 cv::Scalar( 0, 0, 255),
109 2,
110 8,
111 0);
112 }
113
114 // Display
115 cv::namedWindow(window_name, CV_WINDOW_AUTOSIZE );
116 cv::imshow(window_name, histImage );
117 }
118
119 //Get the device context
120 //Create GPU array/vector
121 //Copy the image & set up the kernel
122 //Execute the kernel
123 //Copy GPU data back to CPU cv::Mat data pointer
124 //OpenCV conversion for convienient display
calculateHistogramUsingCL(cv::Mat src,compute::command_queue & queue)125 void calculateHistogramUsingCL(cv::Mat src, compute::command_queue &queue)
126 {
127 compute::context context = queue.get_context();
128
129 // Convert image to BGRA (OpenCL requires 16-byte aligned data)
130 cv::cvtColor(src, src, CV_BGR2BGRA);
131
132 //3 channels & 256 bins : alpha channel is ignored
133 compute::vector<int> gpu_b_hist(histSize, context);
134 compute::vector<int> gpu_g_hist(histSize, context);
135 compute::vector<int> gpu_r_hist(histSize, context);
136
137 // Transfer image to gpu
138 compute::image2d gpu_src =
139 compute::opencv_create_image2d_with_mat(
140 src, compute::image2d::read_only,
141 queue
142 );
143
144 compute::program histogram_program =
145 compute::program::create_with_source(source, context);
146 histogram_program.build();
147
148 // create histogram kernel and set arguments
149 compute::kernel histogram_kernel(histogram_program, "histogram");
150 histogram_kernel.set_arg(0, gpu_src);
151 histogram_kernel.set_arg(1, gpu_b_hist.get_buffer());
152 histogram_kernel.set_arg(2, gpu_g_hist.get_buffer());
153 histogram_kernel.set_arg(3, gpu_r_hist.get_buffer());
154
155 // run histogram kernel
156 // each kernel thread updating red, green & blue bins
157 size_t origin[2] = { 0, 0 };
158 size_t region[2] = { gpu_src.width(),
159 gpu_src.height() };
160
161 queue.enqueue_nd_range_kernel(histogram_kernel, 2, origin, region, 0);
162
163 //Make sure kernel get executed and data copied back
164 queue.finish();
165
166 //create Mat and copy GPU bins to CPU memory
167 cv::Mat b_hist(256, 1, CV_32SC1);
168 compute::copy(gpu_b_hist.begin(), gpu_b_hist.end(), b_hist.data, queue);
169 cv::Mat g_hist(256, 1, CV_32SC1);
170 compute::copy(gpu_g_hist.begin(), gpu_g_hist.end(), g_hist.data, queue);
171 cv::Mat r_hist(256, 1, CV_32SC1);
172 compute::copy(gpu_r_hist.begin(), gpu_r_hist.end(), r_hist.data, queue);
173
174 b_hist.convertTo(b_hist, CV_32FC1); //converted for displaying
175 g_hist.convertTo(g_hist, CV_32FC1);
176 r_hist.convertTo(r_hist, CV_32FC1);
177
178 showHistogramWindow(b_hist, g_hist, r_hist, "Histogram");
179 }
180
main(int argc,char ** argv)181 int main( int argc, char** argv )
182 {
183 // Get default device and setup context
184 compute::device gpu = compute::system::default_device();
185 compute::context context(gpu);
186 compute::command_queue queue(context, gpu);
187
188 cv::Mat src;
189
190 // setup the command line arguments
191 po::options_description desc;
192 desc.add_options()
193 ("help", "show available options")
194 ("image", po::value<std::string>(), "path to image file");
195
196 // Parse the command lines
197 po::variables_map vm;
198 po::store(po::parse_command_line(argc, argv, desc), vm);
199 po::notify(vm);
200
201 //check the command line arguments
202 if(vm.count("help"))
203 {
204 std::cout << desc << std::endl;
205 return 0;
206 }
207
208 //check for image paths
209 if(vm.count("image"))
210 {
211 // Read image with OpenCV
212 src = cv::imread(vm["image"].as<std::string>(),
213 CV_LOAD_IMAGE_COLOR);
214 if(!src.data){
215 std::cerr << "Failed to load image" << std::endl;
216 return -1;
217 }
218 calculateHistogramUsingCL(src, queue);
219 cv::imshow("Image", src);
220 cv::waitKey(0);
221 }
222 else
223 {
224 std::cout << desc << std::endl;
225 return 0;
226 }
227 return 0;
228 }
229