CUDA + OpenCVでWebカメラから取得した画像を変換

grayscale.cu

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <helper_cuda.h>

#include <iostream>
#include <opencv2\opencv.hpp>

__global__ void grayscale(
	uchar3* input,
	uchar3* result,
	int width,
	int height
) {
	const int x = blockIdx.x*blockDim.x + threadIdx.x;
	const int y = blockIdx.y*blockDim.y + threadIdx.y;

	int id = x + y * width;
	uchar3 col = input[id];
	uchar gs = (col.x + col.y + col.z) / 3;
	result[id] = { gs, gs, gs };
}

int main() {

	cv::VideoCapture capture(0);

	if (!capture.isOpened()) {
		std::cerr << "Cannot open camera device. " << std::endl;
		exit(EXIT_FAILURE);
	}

	cv::Mat frame, result;

	// ウィンドウ用意
	cv::String windowName = "Test Window";
	cv::namedWindow(windowName, CV_WINDOW_AUTOSIZE);

	// サイズ取得
	capture >> frame;
	cv::Size size = frame.size();

	// 初期化
	result = frame.clone();

	// デバイスメモリの確保(入力)
	uchar3* d_frame; // (256 = 8bit = sizeof uchar) * 3 = uchar3
	size_t d_frame_pitch;
	checkCudaErrors(cudaMallocPitch(&d_frame, &d_frame_pitch, size.width * 3, size.height));

	// デバイスメモリの確保(出力)
	uchar3* d_result;
	size_t d_result_pitch;
	checkCudaErrors(cudaMallocPitch(&d_result, &d_result_pitch, size.width * 3, size.height));

	// サイズ設定
	dim3 dimBlock(32, 32, 1);
	dim3 dimGrid(size.width / dimBlock.x, size.height / dimBlock.y, 1);


	while (cv::waitKey(1) != 113) {
		capture >> frame;

		checkCudaErrors(cudaMemcpy2D(d_frame, d_frame_pitch, frame.data, frame.step, 
			size.width * 3, size.height, cudaMemcpyDefault));

		grayscale<<<dimGrid, dimBlock, 0>>>(d_frame, d_result, size.width, size.height);

		checkCudaErrors(cudaMemcpy2D(result.data, result.step, d_result, d_result_pitch,
			size.width * 3, size.height, cudaMemcpyDefault));

		cv::imshow(windowName, result);
	}

	cv::destroyAllWindows();
	cudaFree(d_frame);

	return 0;
}