OpenCL with OpenCV (OpenCV-CL)を使ってみた．-RGB編-

前回は取得したフレームをcvtColorでグレー画像に変換してからOpenCL側へ渡す，というまどろっこしいことをしていましたが，実際カラー画像を扱えないとお話にならないので何とか使えないかと試行錯誤．

CUDAの場合はuchar3は構造体として実装されているようで[要出典]，uchar（＝グレー画像）の場合とほとんど同様に扱うことができるが，OpenCLのuchar3は少し違った仕様になっているようで，ucharの場合をそのまま拡張しただけではうまくいかない．（CUDAのコードも適当に整理して上げたいですね）．

ucharなどをスカラ型と呼ぶのに対し，uchar3などはベクタ型と呼ばれますが，このベクタ型の配列にアクセスする場合はvloadnやvstorenを使う．

参考サイト
How vector pointers work in openCL
Vector Data Load and Store Functions

それを考慮して実際にRGB画像を扱って各要素にアクセスした例が以下．

	__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);
	}

view raw kernel.cl hosted with

by GitHub

	#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;
	}

view raw main.cpp hosted with

by GitHub

やっていることは至極単純なグレースケール化のみ．
今回もよくわかっていない点として，widthがいつでもUMatのpitchと等しくなるのか，という点．
例えば，CUDAの場合はx+y * widthとやるとうまくいかなくて，x+y * pitchのようにしなければいけない．つまり要素数はwidth分だが，各行当たりのメモリはもう少し余裕をもって確保されている．このpitchの値はcudaMallocPitchでメモリをアロケートする際に一緒にもらえる．
今回のOpenCV-CLの場合はOpenCV側の提供するUMatを使っているため問題ないのだろうか．
一応UMatもプロパティとしてstepやoffsetを持っているため，これを考慮せずにindexingしても大丈夫なんだろうかいやでも今のところうまくいってるしなあ，という感じ．
つまるところ，まだまだ検証の足りないコードなのでindexing周りで不具合が生じる可能性がある．

ちなみにindexingまわりで不具合があると出力画像はこんな風になりがち．