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 naive optical flow program
29 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE (
30 const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE;
31
32 __kernel void optical_flow (
33 read_only
34 image2d_t current_image,
35 read_only image2d_t previous_image,
36 write_only image2d_t optical_flow,
37 const float scale,
38 const float offset,
39 const float lambda,
40 const float threshold )
41 {
42 int2 coords = (int2)(get_global_id(0), get_global_id(1));
43 float4 current_pixel = read_imagef(current_image,
44 sampler,
45 coords);
46 float4 previous_pixel = read_imagef(previous_image,
47 sampler,
48 coords);
49 int2 x1 = (int2)(offset, 0.f);
50 int2 y1 = (int2)(0.f, offset);
51
52 //get the difference
53 float4 curdif = previous_pixel - current_pixel;
54
55 //calculate the gradient
56 //Image 2 first
57 float4 gradx = read_imagef(previous_image,
58 sampler,
59 coords+x1) -
60 read_imagef(previous_image,
61 sampler,
62 coords-x1);
63 //Image 1
64 gradx += read_imagef(current_image,
65 sampler,
66 coords+x1) -
67 read_imagef(current_image,
68 sampler,
69 coords-x1);
70 //Image 2 first
71 float4 grady = read_imagef(previous_image,
72 sampler,
73 coords+y1) -
74 read_imagef(previous_image,
75 sampler,
76 coords-y1);
77 //Image 1
78 grady += read_imagef(current_image,
79 sampler,
80 coords+y1) -
81 read_imagef(current_image,
82 sampler,
83 coords-y1);
84
85 float4 sqr = (gradx*gradx) + (grady*grady) +
86 (float4)(lambda,lambda, lambda, lambda);
87 float4 gradmag = sqrt(sqr);
88
89 ///////////////////////////////////////////////////
90 float4 vx = curdif * (gradx / gradmag);
91 float vxd = vx.x;//assumes greyscale
92
93 //format output for flowrepos, out(-x,+x,-y,+y)
94 float2 xout = (float2)(fmax(vxd,0.f),fabs(fmin(vxd,0.f)));
95 xout *= scale;
96 ///////////////////////////////////////////////////
97 float4 vy = curdif*(grady/gradmag);
98 float vyd = vy.x;//assumes greyscale
99
100 //format output for flowrepos, out(-x,+x,-y,+y)
101 float2 yout = (float2)(fmax(vyd,0.f),fabs(fmin(vyd,0.f)));
102 yout *= scale;
103 ///////////////////////////////////////////////////
104
105 float4 out = (float4)(xout, yout);
106 float cond = (float)isgreaterequal(length(out), threshold);
107 out *= cond;
108
109 write_imagef(optical_flow, coords, out);
110 }
111 );
112
113 // This example shows how to read two images or use camera
114 // with OpenCV, transfer the frames to the GPU,
115 // and apply a naive optical flow algorithm
116 // written in OpenCL
main(int argc,char * argv[])117 int main(int argc, char *argv[])
118 {
119 // setup the command line arguments
120 po::options_description desc;
121 desc.add_options()
122 ("help", "show available options")
123 ("camera", po::value<int>()->default_value(-1),
124 "if not default camera, specify a camera id")
125 ("image1", po::value<std::string>(), "path to image file 1")
126 ("image2", po::value<std::string>(), "path to image file 2");
127
128 // Parse the command lines
129 po::variables_map vm;
130 po::store(po::parse_command_line(argc, argv, desc), vm);
131 po::notify(vm);
132
133 //check the command line arguments
134 if(vm.count("help"))
135 {
136 std::cout << desc << std::endl;
137 return 0;
138 }
139
140 //OpenCV variables
141 cv::Mat previous_cv_image;
142 cv::Mat current_cv_image;
143 cv::VideoCapture cap; //OpenCV camera handle
144
145 //check for image paths
146 if(vm.count("image1") && vm.count("image2"))
147 {
148 // Read image 1 with OpenCV
149 previous_cv_image = cv::imread(vm["image1"].as<std::string>(),
150 CV_LOAD_IMAGE_COLOR);
151 if(!previous_cv_image.data){
152 std::cerr << "Failed to load image" << std::endl;
153 return -1;
154 }
155
156 // Read image 2 with opencv
157 current_cv_image = cv::imread(vm["image2"].as<std::string>(),
158 CV_LOAD_IMAGE_COLOR);
159 if(!current_cv_image.data){
160 std::cerr << "Failed to load image" << std::endl;
161 return -1;
162 }
163 }
164 else //by default use camera
165 {
166 //open camera
167 cap.open(vm["camera"].as<int>());
168 // read first frame
169 cap >> previous_cv_image;
170 if(!previous_cv_image.data){
171 std::cerr << "failed to capture frame" << std::endl;
172 return -1;
173 }
174
175 // read second frame
176 cap >> current_cv_image;
177 if(!current_cv_image.data){
178 std::cerr << "failed to capture frame" << std::endl;
179 return -1;
180 }
181
182 }
183
184 // Get default device and setup context
185 compute::device gpu = compute::system::default_device();
186 compute::context context(gpu);
187 compute::command_queue queue(context, gpu);
188
189 // Convert image to BGRA (OpenCL requires 16-byte aligned data)
190 cv::cvtColor(previous_cv_image, previous_cv_image, CV_BGR2BGRA);
191 cv::cvtColor(current_cv_image, current_cv_image, CV_BGR2BGRA);
192
193 // Transfer image to gpu
194 compute::image2d dev_previous_image =
195 compute::opencv_create_image2d_with_mat(
196 previous_cv_image, compute::image2d::read_write, queue
197 );
198 // Transfer image to gpu
199 compute::image2d dev_current_image =
200 compute::opencv_create_image2d_with_mat(
201 current_cv_image, compute::image2d::read_write, queue
202 );
203
204 // Create output image
205 compute::image2d dev_output_image(
206 context,
207 dev_previous_image.width(),
208 dev_previous_image.height(),
209 dev_previous_image.format(),
210 compute::image2d::write_only
211 );
212
213 compute::program optical_program =
214 compute::program::create_with_source(source, context);
215 optical_program.build();
216
217 // create flip kernel and set arguments
218 compute::kernel optical_kernel(optical_program, "optical_flow");
219 float scale = 10;
220 float offset = 1;
221 float lambda = 0.0025;
222 float threshold = 1.0;
223
224 optical_kernel.set_arg(0, dev_previous_image);
225 optical_kernel.set_arg(1, dev_current_image);
226 optical_kernel.set_arg(2, dev_output_image);
227 optical_kernel.set_arg(3, scale);
228 optical_kernel.set_arg(4, offset);
229 optical_kernel.set_arg(5, lambda);
230 optical_kernel.set_arg(6, threshold);
231
232 // run flip kernel
233 size_t origin[2] = { 0, 0 };
234 size_t region[2] = { dev_previous_image.width(),
235 dev_previous_image.height() };
236 queue.enqueue_nd_range_kernel(optical_kernel, 2, origin, region, 0);
237
238 //check for image paths
239 if(vm.count("image1") && vm.count("image2"))
240 {
241 // show host image
242 cv::imshow("Previous Frame", previous_cv_image);
243 cv::imshow("Current Frame", current_cv_image);
244
245 // show gpu image
246 compute::opencv_imshow("filtered image", dev_output_image, queue);
247
248 // wait and return
249 cv::waitKey(0);
250 }
251 else
252 {
253 char key = '\0';
254 while(key != 27) //check for escape key
255 {
256 cap >> current_cv_image;
257
258 // Convert image to BGRA (OpenCL requires 16-byte aligned data)
259 cv::cvtColor(current_cv_image, current_cv_image, CV_BGR2BGRA);
260
261 // Update the device image memory with current frame data
262 compute::opencv_copy_mat_to_image(previous_cv_image,
263 dev_previous_image,
264 queue);
265 compute::opencv_copy_mat_to_image(current_cv_image,
266 dev_current_image,
267 queue);
268
269 // Run the kernel on the device
270 queue.enqueue_nd_range_kernel(optical_kernel, 2, origin, region, 0);
271
272 // Show host image
273 cv::imshow("Previous Frame", previous_cv_image);
274 cv::imshow("Current Frame", current_cv_image);
275
276 // Show GPU image
277 compute::opencv_imshow("filtered image", dev_output_image, queue);
278
279 // Copy current frame container to previous frame container
280 current_cv_image.copyTo(previous_cv_image);
281
282 // wait
283 key = cv::waitKey(10);
284 }
285
286 }
287 return 0;
288 }
289
290