CUDAのSurfaceを使ってみる

CUDAのTextureはreadonlyです. (唐突)
「なんでwriteできないんだ！HLSLにはRWTexture2Dがあるのに！」と思うかもしれませんが, Textureの特色はグローバルメモリからデータをフェッチしてくる際に利用されるキャッシュにあるため, そもそも書き込みではその恩恵を受けることができず, よって書き込みはできる必要がないと言えます. (書き込みの際には普通のメモリを使う).
SurfaceはCC2.0以上でしか利用できませんが, Textureと異なり書き込みも行うことができます. じゃあSurfaceには書き込みにもうま味があるんか?というと, 特にそういった記述はProgramming Guide中で見つけられませんでした. TextureとSurfaceは同列に語られているようなので, 単に書き込みの対象にも指定できるようになっただけなのかもしれません.
処理が一段だけであれば単純にTextureを使えばよいのですが, 処理が何段階もあり, 二つのTextureの間を行き来するようにして処理をしていく場合には少し不便なので(ほんまか?)Surfaceを利用してみました.

今回の知見ですが,

CUDAはバージョンによって結構仕様が変わっているっぽいのでちゃんと自分が使用しているバージョンの Programming Guide を読まないといけない(それはそう)
Textureの場合はバージョンの違いに加えて, Low-Level APIとHigh-Level APIの２種類のAPIが存在するため, 一方でうまくいかない場合は他方を試してみると良い

といったところでしょうか.

以下にSurfaceのサンプルコードを示します. cudaMallocPitchで確保した普通のバッファに入れてある, Webカメラから取得した色情報をSurfaceに移して(この時uchar->floatの変換とグレースケール化を行う), また戻すだけです. surfRef2は今後使う用で今は使っていないです. ガウス窓も今後使う用で今は使っていないです.

	//=============================================================================
	// Includes
	//=============================================================================
	#include <cuda_runtime.h>
	#include <device_launch_parameters.h>
	#include <helper_functions.h>
	#include <helper_cuda.h>

	#include <iostream>
	#define _USE_MATH_DEFINES
	#include <cmath>

	#include <opencv2\opencv.hpp>
	//=============================================================================
	// Defines
	//=============================================================================
	#define GAUSS_WINDOW_WIDTH 3
	#define GAUSS_WINDOW_SIZE ((2 * GAUSS_WINDOW_WIDTH + 1) * (2 * GAUSS_WINDOW_WIDTH + 1))


	//=============================================================================
	// Buffers
	//=============================================================================
	__constant__ float gaussWindow[GAUSS_WINDOW_SIZE];

	surface<void, cudaSurfaceType2D> surfRef1;
	cudaArray *d_array1; // surfRef1の実体
	surface<void, cudaSurfaceType2D> surfRef2;
	cudaArray *d_array2; // surfRef2の実体

	//=============================================================================
	// Device Functions
	//=============================================================================
	__global__ void UCharToFloat(
	uchar3* input,
	int width
	) {
	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];
	float gs = (float)(col.x + col.y + col.z) / 3 / 255;
	float4 re;
	re.x = gs;
	re.y = gs;
	re.z = gs;
	re.w = 0;

	surf2Dwrite(re, surfRef1, x * sizeof(float4), y);
	}

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

	float4 gs;
	surf2Dread(&gs, surfRef1, x * sizeof(float4), y);
	uchar3 col;

	col.x = (uchar)(gs.x * 255);
	col.y = (uchar)(gs.y * 255);
	col.z = (uchar)(gs.z * 255);

	result[id] = col;
	}

	//=============================================================================
	// Main Function
	//=============================================================================

	int main() {
	std::cout << "version: " << CV_VERSION << std::endl;

	//=========================================================================
	// ビデオ関係初期化
	//=========================================================================

	cv::VideoCapture capture(0);

	if (!capture.isOpened()) {
	std::cerr << "Error : 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();

	//=========================================================================
	// GPUデバイスメモリ関係
	//=========================================================================

	// スレッド・ブロックサイズ設定
	dim3 dimBlock(32, 32, 1);
	dim3 dimGrid(size.width / dimBlock.x, size.height / dimBlock.y, 1);

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

	// Array用のメモリを確保
	cudaChannelFormatDesc cdesc = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat);
	checkCudaErrors(cudaMallocArray(&d_array1, &cdesc, size.width, size.height, cudaArraySurfaceLoadStore));
	checkCudaErrors(cudaMallocArray(&d_array2, &cdesc, size.width, size.height, cudaArraySurfaceLoadStore));

	// サーフェースにバインド
	size_t offset;
	checkCudaErrors(cudaBindSurfaceToArray(surfRef1, d_array1));
	checkCudaErrors(cudaBindSurfaceToArray(surfRef2, d_array2));

	// ガウス窓の設定
	float *tmpGaussWindow = new float[GAUSS_WINDOW_SIZE];
	const float gauss_sigma = (float)GAUSS_WINDOW_WIDTH / 3;
	for (int i = -GAUSS_WINDOW_WIDTH; i <= GAUSS_WINDOW_WIDTH; i++) {
	for (int j = -GAUSS_WINDOW_WIDTH; j <= GAUSS_WINDOW_WIDTH; j++) {
	tmpGaussWindow[i + (2 * GAUSS_WINDOW_WIDTH + 1) * j] =
	expf((i * i + j * j) / (-2 * gauss_sigma * gauss_sigma)) / (2 * M_PI * gauss_sigma * gauss_sigma);
	}
	}
	checkCudaErrors(cudaMemcpyToSymbol(gaussWindow, tmpGaussWindow, sizeof(float) * GAUSS_WINDOW_SIZE));
	// delete[] tmpGaussWindow;

	//=========================================================================
	// メインループ
	//=========================================================================

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

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

	UCharToFloat << <dimGrid, dimBlock, 0 >> > (d_frame, size.width);
	FloatToUChar << <dimGrid, dimBlock, 0 >> > (d_result, size.width);

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

	cv::imshow(windowName, result);
	}

	cv::destroyAllWindows();
	cudaFreeArray(d_array1);
	cudaFreeArray(d_array2);

	return 0;
	}