【CUDA】テクスチャメモリの使用方法　その１

テクスチャメモリ(TextureMemory) とは、グローバルメモリと同様に扱えるメモリです。

使用するメリットですが、スレッド内で連続したメモリ領域にアクセスした場合、キャッシュが効くようなのでグローバルメモリよりも高速化できる可能性があります。

ちょっとしたコードの追加で使えるようになるので、試しに使ってみましょう。

API
コード例
速度比較
補足

API

テクスチャメモリをバインドする際に使用するAPIです。以下は2次元で読み出すときに使います。
１次元や３次元で使用できるものもありますが、今回は省略します。

API名	解説
cudaBindTexture2D	デバイスメモリにテクスチャメモリをバインドします。
cudaUnbindTexture	テクスチャメモリのバインドを解除します。
tex2D	カーネル内で使用します。テクスチャメモリからデータを取得します。

コード例

コード例を示します。
こちらで示したコードを、テクスチャメモリ化しています。

【CUDA】ブロック・スレッドを2次元で使う

CUDAのスレッドは２次元で持つことができるので、その使い方の解説をします。

カーネルは以下の通り。


//テクスチャのオブジェクト <int型,2次元>　グローバルで宣言
texture<int, cudaTextureType2D> tex;

//テクスチャメモリで平均を取得するカーネル
__global__ void kernel_mean_texture(int *pResult, int sizex, int sizey)
{
	int idx = blockDim.x * blockIdx.x + threadIdx.x; //自分のスレッドxのindex
	int idy = blockDim.y * blockIdx.y + threadIdx.y; //自分のスレッドyのindex

	if (idx < 1 || idy < 1 || idx >= sizex - 1 || idy >= sizey - 1) {
		return; //配列の左右上下端は無視
	}

	unsigned int sum = 0;

	for (int y = 0; y < 3; y++)
	{
		for (int x = 0; x < 3; x++) {
			sum += tex2D(tex, idx + x - 1, idy + y - 1); //２次元のテクスチャメモリにアクセス
		}
	}

	pResult[idy * sizex + idx] = sum / 9; //結果を書き込み　書き込みにはテクスチャメモリを使えない

	return;
}

//テクスチャのオブジェクト <int型,2次元>　グローバルで宣言

texture<int, cudaTextureType2D> tex;

//テクスチャメモリで平均を取得するカーネル

__global__ void kernel_mean_texture(int *pResult, int sizex, int sizey)

{

int idx = blockDim.x * blockIdx.x + threadIdx.x; //自分のスレッドxのindex

int idy = blockDim.y * blockIdx.y + threadIdx.y; //自分のスレッドyのindex

if (idx < 1 || idy < 1 || idx >= sizex - 1 || idy >= sizey - 1) {

return; //配列の左右上下端は無視

}

unsigned int sum = 0;

for (int y = 0; y < 3; y++)

{

for (int x = 0; x < 3; x++) {

sum += tex2D(tex, idx + x - 1, idy + y - 1); //２次元のテクスチャメモリにアクセス

}

pResult[idy * sizex + idx] = sum / 9; //結果を書き込み　書き込みにはテクスチャメモリを使えない

return;

}

処理関数は以下の通り。


//処理関数
void Process()
{
	int sizex = 8192;
	int sizey = 8192;
	int length = sizex * sizey;
	size_t size = sizeof(int) * length;


	//ホストメモリのポインタ
	int* pHostSrc;		//数値
	int* pHostResult;	//結果
	//デバイスメモリのポインタ
	int* pDevSrc;		//数値
	int* pDevResult;	//結果


	//ホストメモリの確保
	cudaMallocHost(&pHostSrc, size);
	cudaMallocHost(&pHostResult, size);
	//デバイスメモリの確保
	cudaMalloc(&pDevSrc, size);
	cudaMalloc(&pDevResult, size);

	//ホストメモリに適当な値を設定
	for (int n = 0; n < length; n++) {
		pHostSrc[n] = n;
	}

	//グローバルメモリをテクスチャメモリ(2次元)にバインド
	cudaBindTexture2D(NULL, tex, pDevSrc, cudaCreateChannelDesc<int>(), sizex, sizey, sizeof(int) * sizex);

	//ホスト->デバイスへ数値を転送
	cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice);

	//カーネル
	dim3 block(16, 16, 1);
	dim3 grid((sizex + 16 - 1) / 16, (sizey + 16 - 1) / 16, 1);
	kernel_mean_texture <<<grid, block >>> (pDevResult, sizex, sizey);

	//デバイス->ホストへ結果を転送
	cudaMemcpy(pHostResult, pDevResult, size, cudaMemcpyDeviceToHost);

	//バインドを解除
	cudaUnbindTexture(tex);

	//デバイスメモリの開放
	cudaFree(pDevSrc);
	cudaFree(pDevResult);
	//ホストメモリの開放
	cudaFreeHost(pHostSrc);
	cudaFreeHost(pHostResult);

}

//処理関数

void Process()

{

int sizex = 8192;

int sizey = 8192;

int length = sizex * sizey;

size_t size = sizeof(int) * length;

//ホストメモリのポインタ

int* pHostSrc; //数値

int* pHostResult; //結果

//デバイスメモリのポインタ

int* pDevSrc; //数値

int* pDevResult; //結果

//ホストメモリの確保

cudaMallocHost(&pHostSrc, size);

cudaMallocHost(&pHostResult, size);

//デバイスメモリの確保

cudaMalloc(&pDevSrc, size);

cudaMalloc(&pDevResult, size);

//ホストメモリに適当な値を設定

for (int n = 0; n < length; n++) {

pHostSrc[n] = n;

}

//グローバルメモリをテクスチャメモリ(2次元)にバインド

cudaBindTexture2D(NULL, tex, pDevSrc, cudaCreateChannelDesc<int>(), sizex, sizey, sizeof(int) * sizex);

//ホスト->デバイスへ数値を転送

cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice);

//カーネル

dim3 block(16, 16, 1);

dim3 grid((sizex + 16 - 1) / 16, (sizey + 16 - 1) / 16, 1);

kernel_mean_texture <<<grid, block >>> (pDevResult, sizex, sizey);

//デバイス->ホストへ結果を転送

cudaMemcpy(pHostResult, pDevResult, size, cudaMemcpyDeviceToHost);

//バインドを解除

cudaUnbindTexture(tex);

//デバイスメモリの開放

cudaFree(pDevSrc);

cudaFree(pDevResult);

//ホストメモリの開放

cudaFreeHost(pHostSrc);

cudaFreeHost(pHostResult);

}

バインド時の引数でcudaCreateChannelDesc<int>を指定することで、tex2D使用時に4Byteを１要素として読み出すことができます。
int を unsigned char や short に置き換えると、1Byte,２Byteを１要素として読み出せるようになります。

速度比較

テクスチャメモリ使用時・未使用時のカーネルの処理時間を比較しました。
なぜか若干遅くなっていますね・・・どのようなケースで高速化されるかもう少し検証が必要そうです。

環境（CPU：Intel Corei7 8700 3.2GHz 　GPU：GeForce GTX1070)

テクスチャメモリ未使用	3.7msec
テクスチャメモリ使用	4.3msec

とりあえず、高速にはなりませんでしたが、メモリ範囲外を読んだときに指定値で自動補完してくれたり、tex2Dのx,y引数に整数以外を指定すると補完値を返すモードがあったりして便利なので使う価値はあるでしょう。

補足

・読み取り専用なので、書き込みには使えません。
・メモリサイズは512byteでアライメントされている必要があります。（512の倍数）
・単一のテクスチャのオブジェクトが（CPU側の）複数スレッドから同時にアクセスされないように注意。
・X、Yどちらの方向のメモリにアクセスしてもキャッシュが効くようです。
・デバッグモードだと結構遅くなります。

作成方法のパターン２はこちら。