CUDAカーネルをC++クラスのメンバ関数に実装

C++のCUDAコードは、オブジェクト指向で書きづらいため、難解なコードが多い。

そこで、以下のコードを作ってみた。
https://github.com/rueiwoqpqpwoeiru/image-process-CUDA/blob/main/Basic.h

Basic.h の12行目のgeneric_kernelは、可変引数のテンプレート関数で、その中で、
(obj->kernel)(args...)　により、54行目や97行目にある「kernel」という名前の
メンバ関数を呼ぶ。

Basic.h の12行目

template<typename T, class... Args> __global__ void generic_kernel(T* obj, Args... args) {
	(obj->kernel)(args...);
}

Basic.h の54行目（メンバ関数）

__device__ __forceinline__ void kernel(const D_data<T_in> in, D_data<T_out> out) {
	unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
	unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
	if (in.w > x && in.h > y) {
		for (unsigned int k = 0; k < in.ch; k++) {
			out.val(x, y, k) = static_cast<T_out>(in.tex(x, y, k));
		}
	}
}

generic_kernelは、51行目や94行目の通り、自身のクラスの型を
テンプレート引数として与えて、クラス内から呼び出す。このとき
一つ目の引数は「this」で、その後の引数は可変。

Basic.h の51行目（メンバ関数）

void do_proc(const D_mem<T_in>& in, const dim3& block = dim3(32, 8)) {
	Proc_unit p(in.d_data.w, in.d_data.h, block);
	Stopwatch sw;
	K_func::generic_kernel<Cast_img> << < p.grid, p.block >> > (this, in.d_data, _out.d_data); // 51行目
	sw.print_time(__FUNCTION__);
}

この方法だと、CUDAカーネルを実装するメンバ関数の名前が「kernel」に限定され、
１つのクラスから１つのカーネルしか呼べない。しかし、クラスを入れ子にして
階層構造を作ることで、不便なくコードを作れた。むしろ、クラスが複雑にならない
メリットがあった。また、速度面も問題無かった。ただし「kernel」関数は、引数を
参照型にしたら動作しなかったので注意。

以上説明した方法によれば、CUDAカーネルをクラス内に閉じ込めてブラックボックス化できる
ため、クラスを作ってしまえば、部品を組み合わせる感覚で、容易にGPUを利用できる。

なお、ここで紹介した方法を使ったCUDAコードについては、以下で詳細に説明。
https://qiita.com/rueiwoqpqpwoeiru/items/69a931e4042a4f9d4434