An example of OpenCV-CL (deal with RGB textures)

kernel.cl

__kernel void grayscale(
	__global uchar* input,
	int input_step, int input_offset,
	__global uchar* result,
	int result_step, int result_offset,
	int height, int width
) {
	int x = get_global_id(0);
	int y = get_global_id(1);

	if (x >= width || y >= height) {
		return;
	}

	int coord = x + y * width;

	uchar3 data = vload3(coord, input);
	uchar gray = (data.x + data.y + data.z) / 3;
	uchar3 res = (uchar3)(gray, gray, gray);
	vstore3(res, coord, result);
}

main.cpp

#include <iostream>
#include <fstream>
#include <string>
#include <iterator>
#include <opencv2/opencv.hpp>
#include <opencv2/core/ocl.hpp>
namespace cl = cv::ocl;
using uint = unsigned int;

const std::string projectName = "OpenCV CL";
const std::string projectPath = "path/to/this/project";

int main() {
	//===========================================
	// Prepare GPU Code
	//===========================================

	// check if opencl is available
	if (!cl::haveOpenCL()) {
		std::cerr 
			<< "OpenCL is not available. " 
			<< std::endl;
		return EXIT_FAILURE;
	}

	// create opencl context of gpu
	cl::Context ctx;
	if (!ctx.create(cl::Device::TYPE_GPU)) {
		std::cerr 
			<< "Failed to create GPU context. " 
			<< std::endl;
		return EXIT_FAILURE;
	}

	// select a device
	cl::Device(ctx.device(0));

	// read the kernel code
	std::ifstream kernel_file("kernel.cl");
	if (kernel_file.fail()) {
		std::cerr 
			<< "Failed to load the kernel code. " 
			<< std::endl;
		return EXIT_FAILURE;
	}
	// interpret as chars
	std::string kernel_code(
		(std::istreambuf_iterator<char>(kernel_file)), 
		std::istreambuf_iterator<char>());
	cl::ProgramSource kernel_source(kernel_code);

	// Compile the kernel code
	cv::String err_msg;
	cv::String option = "";
	cl::Program program = ctx.getProg(
		kernel_source, option, err_msg);
	if (!err_msg.empty()) {
		std::cerr
			<< "Compile Error has occurred. \n"
			<< err_msg
			<< std::endl;
		return EXIT_FAILURE;
	}

	//===========================================
	// Prepare resources
	//===========================================

	// prepare video capture
	cv::VideoCapture cap(0);
	if (!cap.isOpened()) {
		std::cerr 
			<< "Cannot open webcam device. " 
			<< std::endl;
		return EXIT_FAILURE;
	}

	// prepare window
	cv::String windowName = projectName;
	cv::namedWindow(windowName, CV_WINDOW_AUTOSIZE);

	// Texture Buffers
	cv::Mat frame;
	cap >> frame; // aquire frame size

	cv::UMat d_frame = frame.getUMat(
		cv::ACCESS_READ,
		cv::USAGE_ALLOCATE_DEVICE_MEMORY
	);
	cv::UMat d_result(
		frame.size(),
		frame.type(),
		cv::ACCESS_WRITE,
		cv::USAGE_ALLOCATE_DEVICE_MEMORY
	);

	// kernel settings
	cl::Kernel kernel("grayscale", program);
	kernel.args(
		cl::KernelArg::ReadOnlyNoSize(d_frame),
		cl::KernelArg::ReadWrite(d_result)
	);
	size_t dimThreads[3] = {
		frame.size().width,
		frame.size().height,
		1
	};

	//===========================================
	// Main Loop
	//===========================================
	bool is_ok = true;
	while (cv::waitKey(10) != 'q') {
		cap >> frame;
		is_ok = kernel.run(3, dimThreads, NULL, true);
		if (!is_ok) {
			std::cerr
				<< "Failed to execute kernel. "
				<< std::endl;
			continue;
		}
		cv::imshow(windowName, d_result);
	}

	return 0;
}