1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@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
13 #include <opencv2/core/core.hpp>
14 #include <opencv2/highgui/highgui.hpp>
15 #include <opencv2/imgproc/imgproc.hpp>
16
17 #include <boost/compute/system.hpp>
18 #include <boost/compute/interop/opencv/core.hpp>
19 #include <boost/compute/interop/opencv/highgui.hpp>
20 #include <boost/compute/utility/source.hpp>
21
22 namespace compute = boost::compute;
23
24 // this example shows how to read an image with OpenCV, transfer the
25 // image to the GPU, and apply a simple flip filter written in OpenCL
main(int argc,char * argv[])26 int main(int argc, char *argv[])
27 {
28 // check command line
29 if(argc < 2){
30 std::cerr << "usage: " << argv[0] << " FILENAME" << std::endl;
31 return -1;
32 }
33
34 // read image with opencv
35 cv::Mat cv_image = cv::imread(argv[1], CV_LOAD_IMAGE_COLOR);
36 if(!cv_image.data){
37 std::cerr << "failed to load image" << std::endl;
38 return -1;
39 }
40
41 // get default device and setup context
42 compute::device gpu = compute::system::default_device();
43 compute::context context(gpu);
44 compute::command_queue queue(context, gpu);
45
46 // convert image to BGRA (OpenCL requires 16-byte aligned data)
47 cv::cvtColor(cv_image, cv_image, CV_BGR2BGRA);
48
49 // transfer image to gpu
50 compute::image2d input_image =
51 compute::opencv_create_image2d_with_mat(
52 cv_image, compute::image2d::read_write, queue
53 );
54
55 // create output image
56 compute::image2d output_image(
57 context,
58 input_image.width(),
59 input_image.height(),
60 input_image.format(),
61 compute::image2d::write_only
62 );
63
64 // create flip program
65 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
66 __kernel void flip_kernel(__read_only image2d_t input,
67 __write_only image2d_t output)
68 {
69 const sampler_t sampler = CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
70 int height = get_image_height(input);
71 int2 input_coord = { get_global_id(0), get_global_id(1) };
72 int2 output_coord = { input_coord.x, height - input_coord.y - 1 };
73 float4 value = read_imagef(input, sampler, input_coord);
74 write_imagef(output, output_coord, value);
75 }
76 );
77
78 compute::program flip_program =
79 compute::program::create_with_source(source, context);
80 flip_program.build();
81
82 // create flip kernel and set arguments
83 compute::kernel flip_kernel(flip_program, "flip_kernel");
84 flip_kernel.set_arg(0, input_image);
85 flip_kernel.set_arg(1, output_image);
86
87 // run flip kernel
88 size_t origin[2] = { 0, 0 };
89 size_t region[2] = { input_image.width(), input_image.height() };
90 queue.enqueue_nd_range_kernel(flip_kernel, 2, origin, region, 0);
91
92 // show host image
93 cv::imshow("opencv image", cv_image);
94
95 // show gpu image
96 compute::opencv_imshow("filtered image", output_image, queue);
97
98 // wait and return
99 cv::waitKey(0);
100 return 0;
101 }
102