123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265 |
- //---------------------------------------------------------------------------//
- // Copyright (c) 2013-2014 Mageswaran.D <mageswaran1989@gmail.com>
- //
- // Distributed under the Boost Software License, Version 1.0
- // See accompanying file LICENSE_1_0.txt or copy at
- // http://www.boost.org/LICENSE_1_0.txt
- //
- // See http://boostorg.github.com/compute for more information.
- //---------------------------------------------------------------------------//
- #include <iostream>
- #include <string>
- #include <opencv2/core/core.hpp>
- #include <opencv2/highgui/highgui.hpp>
- #include <opencv2/imgproc/imgproc.hpp>
- #include <boost/compute/system.hpp>
- #include <boost/compute/interop/opencv/core.hpp>
- #include <boost/compute/interop/opencv/highgui.hpp>
- #include <boost/compute/utility/source.hpp>
- #include <boost/program_options.hpp>
- namespace compute = boost::compute;
- namespace po = boost::program_options;
- // Create convolution program
- const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE (
- __kernel void convolution(__read_only image2d_t sourceImage,
- __write_only image2d_t outputImage,
- __constant float* filter,
- int filterWidth)
- {
- const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
- CLK_ADDRESS_CLAMP_TO_EDGE |
- CLK_FILTER_NEAREST;
- // Store each work-item's unique row and column
- int x = get_global_id(0);
- int y = get_global_id(1);
- // Half the width of the filter is needed for indexing
- // memory later
- int halfWidth = (int)(filterWidth/2);
- // All accesses to images return data as four-element vector
- // (i.e., float4).
- float4 sum = {0.0f, 0.0f, 0.0f, 0.0f};
- // Iterator for the filter
- int filterIdx = 0;
- // Each work-item iterates around its local area based on the
- // size of the filter
- int2 coords; // Coordinates for accessing the image
- // Iterate the filter rows
- for(int i = -halfWidth; i <= halfWidth; i++)
- {
- coords.y = y + i;
- // Iterate over the filter columns
- for(int j = -halfWidth; j <= halfWidth; j++)
- {
- coords.x = x + j;
- float4 pixel;
- // Read a pixel from the image.
- // Work on a channel
- pixel = read_imagef(sourceImage, sampler, coords);
- sum.x += pixel.x * filter[filterIdx++];
- //sum.y += pixel.y * filter[filterIdx++];
- //sum.z += pixel.z * filter[filterIdx++];
- }
- }
- barrier(CLK_GLOBAL_MEM_FENCE);
- // Copy the data to the output image if the
- // work-item is in bounds
- if(y < get_image_height(sourceImage) &&
- x < get_image_width(sourceImage))
- {
- coords.x = x;
- coords.y = y;
- //Same channel is copied in all three channels
- //write_imagef(outputImage, coords,
- // (float4)(sum.x,sum.x,sum.x,1.0f));
- write_imagef(outputImage, coords, sum);
- }
- }
- );
- // This example shows how to read two images or use camera
- // with OpenCV, transfer the frames to the GPU,
- // and apply a convolution written in OpenCL
- int main(int argc, char *argv[])
- {
- ///////////////////////////////////////////////////////////////////////////
- // setup the command line arguments
- po::options_description desc;
- desc.add_options()
- ("help", "show available options")
- ("camera", po::value<int>()->default_value(-1),
- "if not default camera, specify a camera id")
- ("image", po::value<std::string>(), "path to image file");
- // Parse the command lines
- po::variables_map vm;
- po::store(po::parse_command_line(argc, argv, desc), vm);
- po::notify(vm);
- //check the command line arguments
- if(vm.count("help"))
- {
- std::cout << desc << std::endl;
- return 0;
- }
- ///////////////////////////////////////////////////////////////////////////
- //OpenCV variables
- cv::Mat cv_mat;
- cv::VideoCapture cap; //OpenCV camera handle.
- //Filter Variables
- float filter[9] = {
- -1.0, 0.0, 1.0,
- -2.0, 0.0, 2.0,
- -1.0, 0.0, 1.0,
- };
- // The convolution filter is 3x3
- int filterWidth = 3;
- //OpenCL variables
- // Get default device and setup context
- compute::device gpu = compute::system::default_device();
- compute::context context(gpu);
- compute::command_queue queue(context, gpu);
- compute::buffer dev_filter(context, sizeof(filter),
- compute::memory_object::read_only |
- compute::memory_object::copy_host_ptr,
- filter);
- compute::program filter_program =
- compute::program::create_with_source(source, context);
- try
- {
- filter_program.build();
- }
- catch(compute::opencl_error e)
- {
- std::cout<<"Build Error: "<<std::endl
- <<filter_program.build_log();
- return -1;
- }
- // create fliter kernel and set arguments
- compute::kernel filter_kernel(filter_program, "convolution");
- ///////////////////////////////////////////////////////////////////////////
- //check for image paths
- if(vm.count("image"))
- {
- // Read image with OpenCV
- cv_mat = cv::imread(vm["image"].as<std::string>(),
- CV_LOAD_IMAGE_COLOR);
- if(!cv_mat.data){
- std::cerr << "Failed to load image" << std::endl;
- return -1;
- }
- }
- else //by default use camera
- {
- //open camera
- cap.open(vm["camera"].as<int>());
- // read first frame
- cap >> cv_mat;
- if(!cv_mat.data){
- std::cerr << "failed to capture frame" << std::endl;
- return -1;
- }
- }
- // Convert image to BGRA (OpenCL requires 16-byte aligned data)
- cv::cvtColor(cv_mat, cv_mat, CV_BGR2BGRA);
- // Transfer image/frame data to gpu
- compute::image2d dev_input_image =
- compute::opencv_create_image2d_with_mat(
- cv_mat, compute::image2d::read_write, queue
- );
- // Create output image
- // Be sure what will be your ouput image/frame size
- compute::image2d dev_output_image(
- context,
- dev_input_image.width(),
- dev_input_image.height(),
- dev_input_image.format(),
- compute::image2d::write_only
- );
- filter_kernel.set_arg(0, dev_input_image);
- filter_kernel.set_arg(1, dev_output_image);
- filter_kernel.set_arg(2, dev_filter);
- filter_kernel.set_arg(3, filterWidth);
- // run flip kernel
- size_t origin[2] = { 0, 0 };
- size_t region[2] = { dev_input_image.width(),
- dev_input_image.height() };
- ///////////////////////////////////////////////////////////////////////////
- queue.enqueue_nd_range_kernel(filter_kernel, 2, origin, region, 0);
- //check for image paths
- if(vm.count("image"))
- {
- // show host image
- cv::imshow("Original Image", cv_mat);
- // show gpu image
- compute::opencv_imshow("Convoluted Image", dev_output_image, queue);
- // wait and return
- cv::waitKey(0);
- }
- else
- {
- char key = '\0';
- while(key != 27) //check for escape key
- {
- cap >> cv_mat;
- // Convert image to BGRA (OpenCL requires 16-byte aligned data)
- cv::cvtColor(cv_mat, cv_mat, CV_BGR2BGRA);
- // Update the device image memory with current frame data
- compute::opencv_copy_mat_to_image(cv_mat,
- dev_input_image,queue);
- // Run the kernel on the device
- queue.enqueue_nd_range_kernel(filter_kernel, 2, origin, region, 0);
- // Show host image
- cv::imshow("Camera Frame", cv_mat);
- // Show GPU image
- compute::opencv_imshow("Convoluted Frame", dev_output_image, queue);
- // wait
- key = cv::waitKey(10);
- }
- }
- return 0;
- }
|