{

// check command line

if(argc < 2){

std::cerr << "usage: " << argv[0] << " FILENAME" << std::endl;

return -1;

}

// read image with opencv

cv::Mat cv_image = cv::imread(argv[1], CV_LOAD_IMAGE_COLOR);

if(!cv_image.data){

std::cerr << "failed to load image" << std::endl;

return -1;

}

// get default device and setup context

compute::device gpu = compute::system::default_device();

compute::context context(gpu);

compute::command_queue queue(context, gpu);

// convert image to BGRA (OpenCL requires 16-byte aligned data)

cv::cvtColor(cv_image, cv_image, CV_BGR2BGRA);

// transfer image to gpu

compute::image2d input_image =

compute::opencv_create_image2d_with_mat(

cv_image, compute::image2d::read_write, queue

);

// create output image

compute::image2d output_image(

context,

input_image.width(),

input_image.height(),

input_image.format(),

compute::image2d::write_only

);

// create flip program

const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(

__kernel void flip_kernel(__read_only image2d_t input,

__write_only image2d_t output)

{

const sampler_t sampler = CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;

int height = get_image_height(input);

int2 input_coord = { get_global_id(0), get_global_id(1) };

int2 output_coord = { input_coord.x, height - input_coord.y - 1 };

float4 value = read_imagef(input, sampler, input_coord);

write_imagef(output, output_coord, value);

}

);

compute::program flip_program =

compute::program::create_with_source(source, context);

flip_program.build();

// create flip kernel and set arguments

compute::kernel flip_kernel(flip_program, "flip_kernel");

flip_kernel.set_arg(0, input_image);

flip_kernel.set_arg(1, output_image);

// run flip kernel

size_t origin[2] = { 0, 0 };

size_t region[2] = { input_image.width(), input_image.height() };

queue.enqueue_nd_range_kernel(flip_kernel, 2, origin, region, 0);

// show host image

cv::imshow("opencv image", cv_image);

// show gpu image

compute::opencv_imshow("filtered image", output_image, queue);

// wait and return

cv::waitKey(0);

return 0;