二値化 (deviceメモリ周りを扱いやすく)

deviceメモリの確保・解放をやるクラスを作った。 カーネルに渡す時はdeviceメモリのポインタと縦・横のバイト数なんかのデータを保持する構造体を渡すようにしてみた。

binarize.cu:

#include "binarize.cuh"
#include "DeviceMemory.cuh"
 
using namespace FM7b5;
 
void
FM7b5::binarize_gpu(ImageGray& out, const ImageGray& in, const uint8_t thres)
{
	if (in.width() != out.width() || in.height() != out.height()) {
		throw std::runtime_error("sizes of input and output images are diferent.");
	}
 
	const size_t width(in.width()), height(in.height()), bpp(in.bpp());
 
	const size_t threads_per_dim(32);
	dim3 threads_per_block(threads_per_dim, threads_per_dim);
	dim3 blocks_per_grid((width + threads_per_block.x - 1)/ threads_per_block.x,
	                     (height + threads_per_block.y - 1)/ threads_per_block.y);
 
	// allocate input/output memories
	memory::LinearPitch<uint8_t> d_in(width, height), d_out(width, height);
 
	// copy an input image to device memory
	d_in.copy_from(in.data(), bpp * width, height, in.stride());
 
	// launch kernel
	binarize<<<blocks_per_grid, threads_per_block>>>(d_out.ref(), d_in.ref(), width, height, thres);
 
	// copy the result back to host memory
	d_out.copy_to(out.data(), bpp * width, height, out.stride());
}

ひょんな事からカーネルはヘッダファイルに移動。

binarize.cuh:

#include "binarize.h"
#include "MDView.cuh"
 
namespace FM7b5
{
	template <class T, class U>
	__global__ void
	binarize(const MDView<T, 2> out, const MDView<U, 2> in, const size_t width, const size_t height, const uint8_t thres)
	{
		const size_t w(blockDim.x * blockIdx.x + threadIdx.x);
		const size_t h(blockDim.y * blockIdx.y + threadIdx.y);
 
		if (w >= width || h >= height) {
			return;
		}
 
		out[h][w] = (in[h][w] < thres) ? 0 : 255;
	}
}

これでメモリ周りが多少すっきりしたかなぁ。2次元データを(x, y)で扱うか[y][x]で扱うか迷ってたり。

あとtemplateを使ったせいかReleaseでコンパイルしないと非効率なPTXが生成されてしまうみたい。デバッグ用の情報を出力する(-G)オプション のせいで途中の一時的なクラスたちが残るのかも。

それにしても1命令でも削ったり効率のいいものに置き換えたりしたいdeviceコードで、こんな悠長なことしててもいいんだろうか。 実は*.cu/*.gpuから生成される*.ptxを眺めると、ちょっとした型の違いとかで出てくるものに違いが生じるのが分かる。 特にビットサイズまわりが怪しくて16bit->32bit->16bitみたいな無意味な変換を平気で挟んでくるので、CUDAのコンパイラをあまり信用していなかったり。

最終的にはPTXからもう一段階コンパイルが行われてGPUに渡されるのでそのときに最適化されていれば問題ないと思うけど、ぱっと見で無駄なコードがあると精神衛生上よくない。 幸い計算結果が合わないとか致命的なヤツにはまだ遭遇していないので、コンパイラの性能向上に期待して待つことにする。