123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228 |
- //---------------------------------------------------------------------------//
- // 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.
- //---------------------------------------------------------------------------//
- //Code sample for calculating histogram using OpenCL and
- //displaying image histogram in OpenCV.
- #include <iostream>
- #include <string>
- #include <opencv2/imgproc/imgproc.hpp>
- #include <opencv2/highgui/highgui.hpp>
- #include <boost/compute/source.hpp>
- #include <boost/compute/system.hpp>
- #include <boost/compute/container/vector.hpp>
- #include <boost/compute/interop/opencv/core.hpp>
- #include <boost/compute/interop/opencv/highgui.hpp>
- #include <boost/program_options.hpp>
- namespace compute = boost::compute;
- namespace po = boost::program_options;
- // number of bins
- int histSize = 256;
- // Set the ranges ( for B,G,R) )
- // TryOut: consider the range in kernel calculation
- float range[] = { 0, 256 } ;
- const float* histRange = { range };
- // Create naive histogram program
- // Needs "cl_khr_local_int32_base_atomics" extension
- const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE (
- __kernel void histogram(read_only image2d_t src_image,
- __global int* b_hist,
- __global int* g_hist,
- __global int* r_hist)
- {
- sampler_t sampler =( CLK_NORMALIZED_COORDS_FALSE |
- CLK_FILTER_NEAREST |
- CLK_ADDRESS_CLAMP_TO_EDGE);
- int image_width = get_image_width(src_image);
- int image_height = get_image_height(src_image);
- int2 coords = (int2)(get_global_id(0), get_global_id(1));
- float4 pixel = read_imagef(src_image,sampler, coords);
- //boundary condition
- if ((coords.x < image_width) && (coords.y < image_height))
- {
- uchar indx_x, indx_y, indx_z;
- indx_x = convert_uchar_sat(pixel.x * 255.0f);
- indx_y = convert_uchar_sat(pixel.y * 255.0f);
- indx_z = convert_uchar_sat(pixel.z * 255.0f);
- atomic_inc(&b_hist[(uint)indx_z]);
- atomic_inc(&g_hist[(uint)indx_y]);
- atomic_inc(&r_hist[(uint)indx_x]);
- }
- }
- );
- inline void showHistogramWindow(cv::Mat &b_hist, cv::Mat &g_hist, cv::Mat &r_hist,
- std::string window_name)
- {
- // Draw the histograms for B, G and R
- int hist_w = 1024;
- int hist_h = 768;
- int bin_w = cvRound((double)hist_w/histSize);
- cv::Mat histImage(hist_h, hist_w, CV_8UC3, cv::Scalar(0,0,0));
- // Normalize the result to [ 0, histImage.rows ]
- cv::normalize(b_hist, b_hist, 0, histImage.rows, cv::NORM_MINMAX, -1, cv::Mat());
- cv::normalize(g_hist, g_hist, 0, histImage.rows, cv::NORM_MINMAX, -1, cv::Mat());
- cv::normalize(r_hist, r_hist, 0, histImage.rows, cv::NORM_MINMAX, -1, cv::Mat());
- // Draw for each channel
- for (int i = 1; i < histSize; i++ )
- {
- cv::line(histImage,
- cv::Point(bin_w*(i-1), hist_h - cvRound(b_hist.at<float>(i-1))),
- cv::Point(bin_w*(i), hist_h - cvRound(b_hist.at<float>(i))),
- cv::Scalar(255, 0, 0),
- 2,
- 8,
- 0);
- cv::line(histImage,
- cv::Point(bin_w*(i-1), hist_h - cvRound(g_hist.at<float>(i-1))),
- cv::Point(bin_w*(i), hist_h - cvRound(g_hist.at<float>(i))),
- cv::Scalar(0, 255, 0),
- 2,
- 8,
- 0);
- cv::line(histImage,
- cv::Point( bin_w*(i-1), hist_h - cvRound(r_hist.at<float>(i-1))),
- cv::Point( bin_w*(i), hist_h - cvRound(r_hist.at<float>(i)) ),
- cv::Scalar( 0, 0, 255),
- 2,
- 8,
- 0);
- }
- // Display
- cv::namedWindow(window_name, CV_WINDOW_AUTOSIZE );
- cv::imshow(window_name, histImage );
- }
- //Get the device context
- //Create GPU array/vector
- //Copy the image & set up the kernel
- //Execute the kernel
- //Copy GPU data back to CPU cv::Mat data pointer
- //OpenCV conversion for convienient display
- void calculateHistogramUsingCL(cv::Mat src, compute::command_queue &queue)
- {
- compute::context context = queue.get_context();
- // Convert image to BGRA (OpenCL requires 16-byte aligned data)
- cv::cvtColor(src, src, CV_BGR2BGRA);
- //3 channels & 256 bins : alpha channel is ignored
- compute::vector<int> gpu_b_hist(histSize, context);
- compute::vector<int> gpu_g_hist(histSize, context);
- compute::vector<int> gpu_r_hist(histSize, context);
- // Transfer image to gpu
- compute::image2d gpu_src =
- compute::opencv_create_image2d_with_mat(
- src, compute::image2d::read_only,
- queue
- );
- compute::program histogram_program =
- compute::program::create_with_source(source, context);
- histogram_program.build();
- // create histogram kernel and set arguments
- compute::kernel histogram_kernel(histogram_program, "histogram");
- histogram_kernel.set_arg(0, gpu_src);
- histogram_kernel.set_arg(1, gpu_b_hist.get_buffer());
- histogram_kernel.set_arg(2, gpu_g_hist.get_buffer());
- histogram_kernel.set_arg(3, gpu_r_hist.get_buffer());
- // run histogram kernel
- // each kernel thread updating red, green & blue bins
- size_t origin[2] = { 0, 0 };
- size_t region[2] = { gpu_src.width(),
- gpu_src.height() };
- queue.enqueue_nd_range_kernel(histogram_kernel, 2, origin, region, 0);
- //Make sure kernel get executed and data copied back
- queue.finish();
- //create Mat and copy GPU bins to CPU memory
- cv::Mat b_hist(256, 1, CV_32SC1);
- compute::copy(gpu_b_hist.begin(), gpu_b_hist.end(), b_hist.data, queue);
- cv::Mat g_hist(256, 1, CV_32SC1);
- compute::copy(gpu_g_hist.begin(), gpu_g_hist.end(), g_hist.data, queue);
- cv::Mat r_hist(256, 1, CV_32SC1);
- compute::copy(gpu_r_hist.begin(), gpu_r_hist.end(), r_hist.data, queue);
- b_hist.convertTo(b_hist, CV_32FC1); //converted for displaying
- g_hist.convertTo(g_hist, CV_32FC1);
- r_hist.convertTo(r_hist, CV_32FC1);
- showHistogramWindow(b_hist, g_hist, r_hist, "Histogram");
- }
- int main( int argc, char** argv )
- {
- // Get default device and setup context
- compute::device gpu = compute::system::default_device();
- compute::context context(gpu);
- compute::command_queue queue(context, gpu);
- cv::Mat src;
- // setup the command line arguments
- po::options_description desc;
- desc.add_options()
- ("help", "show available options")
- ("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;
- }
- //check for image paths
- if(vm.count("image"))
- {
- // Read image with OpenCV
- src = cv::imread(vm["image"].as<std::string>(),
- CV_LOAD_IMAGE_COLOR);
- if(!src.data){
- std::cerr << "Failed to load image" << std::endl;
- return -1;
- }
- calculateHistogramUsingCL(src, queue);
- cv::imshow("Image", src);
- cv::waitKey(0);
- }
- else
- {
- std::cout << desc << std::endl;
- return 0;
- }
- return 0;
- }
|