//---------------------------------------------------------------------------// // Copyright (c) 2013-2014 Mageswaran.D // // 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 #include #include #include #include #include #include #include #include #include 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(i-1))), cv::Point(bin_w*(i), hist_h - cvRound(b_hist.at(i))), cv::Scalar(255, 0, 0), 2, 8, 0); cv::line(histImage, cv::Point(bin_w*(i-1), hist_h - cvRound(g_hist.at(i-1))), cv::Point(bin_w*(i), hist_h - cvRound(g_hist.at(i))), cv::Scalar(0, 255, 0), 2, 8, 0); cv::line(histImage, cv::Point( bin_w*(i-1), hist_h - cvRound(r_hist.at(i-1))), cv::Point( bin_w*(i), hist_h - cvRound(r_hist.at(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 gpu_b_hist(histSize, context); compute::vector gpu_g_hist(histSize, context); compute::vector 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(), "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(), 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; }