CUDAで画像処理画像のグレースケール変換ソフト

本日は前回のCUDA C++でGPUプログラミング【CUDA 第1回】に引き続いて、CUDAでのプログラミングをやっていく予定です。

前回は環境構築と、取り合えず動かしてみたというところがメインだったのですが、基本的に画像処理にCUDAを用いたいと思っているので、前回よりは少し画像処理チックなことをやってみます。

各ピクセルで独立に処理をすればよい教材ということで、今回はカラー画像をグレースケールに変換するプログラムをCUDAを利用して書いてみました。

V = 0.299*R + 0.587*G + 0.114*B

グレースケールとカラーの輝度値の変換は、上記の式に従って行っています。

今回の環境

・OS : Windows10(64bit)
・GPU: GeForce GTX 1060
・Visual Studio 2015 インストール済
・CUDA 9.0 インストール済
・OpenCV 3.1 環境構築済

ソースコード

	__global__ void convertToGray(uchar3 color_pixel, unsigned char gray_pixel) {

	int ID = blockIdx.x*blockDim.x + threadIdx.x;

	// カラーからグレースケール変換を実施
	gray_pixel[ID] =
	(unsigned char)(0.299f*color_pixel[ID].x
	+ 0.587f*(float)color_pixel[ID].y
	+ 0.114f*(float)color_pixel[ID].z);
	}

view raw device.cu hosted with

by GitHub

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

	__global__ void convertToGray(uchar3 color_pixel, unsigned char gray_pixel);


	int main(int argc, char *argv) {

	// 画像をカラーで読み込む
	cv::Mat input_img = cv::imread("test.jpg", 1);
	if (input_img.empty() == true) {
	std::cerr << "Error : cannot find input image" << std::endl;
	}

	// 画像のサイズを得る
	int width = input_img.cols;
	int height = input_img.rows;
	std::cout << "Image_size : " << width << "×" << height << std::endl;

	// ホスト(CPU)側配列(RGBを扱えるようunsigned char * 3で扱う)
	uchar3* host_img_array_color = new uchar3[width * height];
	unsigned char* host_img_array_gray = new unsigned char[width * height];

	// 画素値を1次元配列化(OpenCVは画素がBGRの順なのでRGBの順に直して入れる)
	for (int y = 0; y < height; y++) {
	for (int x = 0; x < width; x++) {
	host_img_array_color[x + y * width]
	= make_uchar3(input_img.at<cv::Vec3b>(y, x)[2], input_img.at<cv::Vec3b>(y, x)[1], input_img.at<cv::Vec3b>(y, x)[0]);
	}
	}

	// デバイス(GPU)側のメモリ確保
	uchar3* device_img_array_color;
	unsigned char* device_img_array_gray;
	int datasize_color = sizeof(uchar3) * width * height;
	int datasize_gray = sizeof(unsigned char) * width * height;
	cudaMalloc((void**)&device_img_array_color, datasize_color);
	cudaMalloc((void**)&device_img_array_gray, datasize_gray);

	// データ転送(CPU → GPU)
	cudaMemcpy(device_img_array_color, host_img_array_color, datasize_color, cudaMemcpyHostToDevice);

	// GPU側実行関数
	convertToGray << <width * height, 1 >> > (device_img_array_color, device_img_array_gray);

	// データ転送(GPU → CPU)
	cudaMemcpy(host_img_array_gray, device_img_array_gray, datasize_gray, cudaMemcpyDeviceToHost);

	// 結果画像の出力
	cv::Mat1b output_img(height, width);
	for (int y = 0; y < height; y++) {
	for (int x = 0; x < width; x++) {
	output_img.at<unsigned char>(y, x) = host_img_array_gray[x + y * width];
	}
	}
	cv::imwrite("test_gray.jpg", output_img);

	// 解放
	cudaFree(device_img_array_color);
	cudaFree(device_img_array_gray);
	delete host_img_array_color;
	delete host_img_array_gray;
	}