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 #include <iostream>
12 #include <string>
13
14 #include <opencv2/core/core.hpp>
15 #include <opencv2/highgui/highgui.hpp>
16 #include <opencv2/imgproc/imgproc.hpp>
17
18 #include <boost/compute/system.hpp>
19 #include <boost/compute/interop/opencv/core.hpp>
20 #include <boost/compute/interop/opencv/highgui.hpp>
21 #include <boost/compute/utility/source.hpp>
22
23 #include <boost/program_options.hpp>
24
25 namespace compute = boost::compute;
26 namespace po = boost::program_options;
27
28 // Create convolution program
29 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE (
30 __kernel void convolution(__read_only image2d_t sourceImage,
31 __write_only image2d_t outputImage,
32 __constant float* filter,
33 int filterWidth)
34 {
35 const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
36 CLK_ADDRESS_CLAMP_TO_EDGE |
37 CLK_FILTER_NEAREST;
38
39 // Store each work-item's unique row and column
40 int x = get_global_id(0);
41 int y = get_global_id(1);
42
43 // Half the width of the filter is needed for indexing
44 // memory later
45 int halfWidth = (int)(filterWidth/2);
46
47 // All accesses to images return data as four-element vector
48 // (i.e., float4).
49 float4 sum = {0.0f, 0.0f, 0.0f, 0.0f};
50
51 // Iterator for the filter
52 int filterIdx = 0;
53
54 // Each work-item iterates around its local area based on the
55 // size of the filter
56 int2 coords; // Coordinates for accessing the image
57
58 // Iterate the filter rows
59 for(int i = -halfWidth; i <= halfWidth; i++)
60 {
61 coords.y = y + i;
62
63 // Iterate over the filter columns
64 for(int j = -halfWidth; j <= halfWidth; j++)
65 {
66 coords.x = x + j;
67
68 float4 pixel;
69
70 // Read a pixel from the image.
71 // Work on a channel
72 pixel = read_imagef(sourceImage, sampler, coords);
73 sum.x += pixel.x * filter[filterIdx++];
74 //sum.y += pixel.y * filter[filterIdx++];
75 //sum.z += pixel.z * filter[filterIdx++];
76 }
77 }
78
79 barrier(CLK_GLOBAL_MEM_FENCE);
80 // Copy the data to the output image if the
81 // work-item is in bounds
82 if(y < get_image_height(sourceImage) &&
83 x < get_image_width(sourceImage))
84 {
85 coords.x = x;
86 coords.y = y;
87
88 //Same channel is copied in all three channels
89 //write_imagef(outputImage, coords,
90 // (float4)(sum.x,sum.x,sum.x,1.0f));
91
92 write_imagef(outputImage, coords, sum);
93 }
94 }
95 );
96
97 // This example shows how to read two images or use camera
98 // with OpenCV, transfer the frames to the GPU,
99 // and apply a convolution written in OpenCL
main(int argc,char * argv[])100 int main(int argc, char *argv[])
101 {
102 ///////////////////////////////////////////////////////////////////////////
103
104 // setup the command line arguments
105 po::options_description desc;
106 desc.add_options()
107 ("help", "show available options")
108 ("camera", po::value<int>()->default_value(-1),
109 "if not default camera, specify a camera id")
110 ("image", po::value<std::string>(), "path to image file");
111
112 // Parse the command lines
113 po::variables_map vm;
114 po::store(po::parse_command_line(argc, argv, desc), vm);
115 po::notify(vm);
116
117 //check the command line arguments
118 if(vm.count("help"))
119 {
120 std::cout << desc << std::endl;
121 return 0;
122 }
123
124 ///////////////////////////////////////////////////////////////////////////
125
126 //OpenCV variables
127 cv::Mat cv_mat;
128 cv::VideoCapture cap; //OpenCV camera handle.
129
130 //Filter Variables
131 float filter[9] = {
132 -1.0, 0.0, 1.0,
133 -2.0, 0.0, 2.0,
134 -1.0, 0.0, 1.0,
135 };
136
137 // The convolution filter is 3x3
138 int filterWidth = 3;
139
140 //OpenCL variables
141 // Get default device and setup context
142 compute::device gpu = compute::system::default_device();
143 compute::context context(gpu);
144 compute::command_queue queue(context, gpu);
145 compute::buffer dev_filter(context, sizeof(filter),
146 compute::memory_object::read_only |
147 compute::memory_object::copy_host_ptr,
148 filter);
149
150 compute::program filter_program =
151 compute::program::create_with_source(source, context);
152
153 try
154 {
155 filter_program.build();
156 }
157 catch(compute::opencl_error e)
158 {
159 std::cout<<"Build Error: "<<std::endl
160 <<filter_program.build_log();
161 return -1;
162 }
163
164 // create fliter kernel and set arguments
165 compute::kernel filter_kernel(filter_program, "convolution");
166
167 ///////////////////////////////////////////////////////////////////////////
168
169 //check for image paths
170 if(vm.count("image"))
171 {
172 // Read image with OpenCV
173 cv_mat = cv::imread(vm["image"].as<std::string>(),
174 CV_LOAD_IMAGE_COLOR);
175 if(!cv_mat.data){
176 std::cerr << "Failed to load image" << std::endl;
177 return -1;
178 }
179 }
180 else //by default use camera
181 {
182 //open camera
183 cap.open(vm["camera"].as<int>());
184 // read first frame
185 cap >> cv_mat;
186 if(!cv_mat.data){
187 std::cerr << "failed to capture frame" << std::endl;
188 return -1;
189 }
190 }
191
192 // Convert image to BGRA (OpenCL requires 16-byte aligned data)
193 cv::cvtColor(cv_mat, cv_mat, CV_BGR2BGRA);
194
195 // Transfer image/frame data to gpu
196 compute::image2d dev_input_image =
197 compute::opencv_create_image2d_with_mat(
198 cv_mat, compute::image2d::read_write, queue
199 );
200
201 // Create output image
202 // Be sure what will be your ouput image/frame size
203 compute::image2d dev_output_image(
204 context,
205 dev_input_image.width(),
206 dev_input_image.height(),
207 dev_input_image.format(),
208 compute::image2d::write_only
209 );
210
211 filter_kernel.set_arg(0, dev_input_image);
212 filter_kernel.set_arg(1, dev_output_image);
213 filter_kernel.set_arg(2, dev_filter);
214 filter_kernel.set_arg(3, filterWidth);
215
216 // run flip kernel
217 size_t origin[2] = { 0, 0 };
218 size_t region[2] = { dev_input_image.width(),
219 dev_input_image.height() };
220
221 ///////////////////////////////////////////////////////////////////////////
222
223 queue.enqueue_nd_range_kernel(filter_kernel, 2, origin, region, 0);
224
225 //check for image paths
226 if(vm.count("image"))
227 {
228 // show host image
229 cv::imshow("Original Image", cv_mat);
230
231 // show gpu image
232 compute::opencv_imshow("Convoluted Image", dev_output_image, queue);
233
234 // wait and return
235 cv::waitKey(0);
236 }
237 else
238 {
239 char key = '\0';
240 while(key != 27) //check for escape key
241 {
242 cap >> cv_mat;
243
244 // Convert image to BGRA (OpenCL requires 16-byte aligned data)
245 cv::cvtColor(cv_mat, cv_mat, CV_BGR2BGRA);
246
247 // Update the device image memory with current frame data
248 compute::opencv_copy_mat_to_image(cv_mat,
249 dev_input_image,queue);
250
251 // Run the kernel on the device
252 queue.enqueue_nd_range_kernel(filter_kernel, 2, origin, region, 0);
253
254 // Show host image
255 cv::imshow("Camera Frame", cv_mat);
256
257 // Show GPU image
258 compute::opencv_imshow("Convoluted Frame", dev_output_image, queue);
259
260 // wait
261 key = cv::waitKey(10);
262 }
263 }
264 return 0;
265 }
266