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

kernel.cl

__kernel void negaposi(
	__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 from = mad24(y, input_step, x + input_offset);
	int to = mad24(y, result_step, x + result_offset);

	result[to] = (uchar)fmax(0, fmin(255, 255 - input[from]));
}

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, gray;
	cap >> frame; // aquire frame size
	cv::cvtColor(frame, gray, cv::COLOR_BGR2GRAY);

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

	// kernel settings
	cl::Kernel kernel("negaposi", 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
	//===========================================
	while (cv::waitKey(10) != 'q') {
		cap >> frame;
		cv::cvtColor(frame, gray, cv::COLOR_BGR2GRAY);
		if (!kernel.run(3, dimThreads, NULL, true)) {
			std::cerr
				<< "Failed to execute kernel. "
				<< std::endl;
			continue;
		}
		cv::imshow(windowName, d_result);
	}

	return 0;
}