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

テクスチャメモリ(TextureMemory)の使用方法のパターンその２を解説します。
こちらで解説する方法だと、動的にテクスチャのオブジェクトを生成したり、その他の設定ができます。

その１はこちら

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

テクスチャメモリ(TextureMemory)の解説を行います。スレッド内で連続したメモリ領域にアクセスした場合、キャッシュが効くのでグローバルメモリよりも高速化できる可能性があります。

コード
設定

コード

２Dで呼び出す場合のコード例を示します。

以下はカーネル呼び出し側のコードです。

//デバイスメモリのポインタ
BYTE* pDevSrc;		//数値
BYTE* pDevResult;	//結果

・・・メモリ確保や転送などは省略

cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeClamp; //配列外の読み出しモードx
texDesc.addressMode[1] = cudaAddressModeClamp; //配列外の読み出しモードy
texDesc.filterMode = cudaFilterModePoint;
texDesc.normalizedCoords = 0;
texDesc.readMode = cudaReadModeElementType;
texDesc.sRGB = 0;

cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<BYTE>(); //1要素のサイズ
resDesc.res.pitch2D.pitchInBytes = sizex * sizeof(BYTE); //配列X方向のバイト数
resDesc.res.pitch2D.width = sizex; //配列Xサイズ
resDesc.res.pitch2D.height = sizey; //配列yサイズ
resDesc.res.pitch2D.devPtr = pDevSrc; //デバイスメモリのポインタ
resDesc.resType = cudaResourceTypePitch2D;

//テクスチャオブジェクト　生成
cudaTextureObject_t texObj;
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

・・・

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

・・・

//テクスチャオブジェクト　破棄
cudaDestroyTextureObject(texObj);

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

BYTE* pDevSrc; //数値

BYTE* pDevResult; //結果

・・・メモリ確保や転送などは省略

cudaTextureDesc texDesc;

memset(&texDesc, 0, sizeof(texDesc));

texDesc.addressMode[0] = cudaAddressModeClamp; //配列外の読み出しモードx

texDesc.addressMode[1] = cudaAddressModeClamp; //配列外の読み出しモードy

texDesc.filterMode = cudaFilterModePoint;

texDesc.normalizedCoords = 0;

texDesc.readMode = cudaReadModeElementType;

texDesc.sRGB = 0;

cudaResourceDesc resDesc;

memset(&resDesc, 0, sizeof(resDesc));

resDesc.res.pitch2D.desc = cudaCreateChannelDesc<BYTE>(); //1要素のサイズ

resDesc.res.pitch2D.pitchInBytes = sizex * sizeof(BYTE); //配列X方向のバイト数

resDesc.res.pitch2D.width = sizex; //配列Xサイズ

resDesc.res.pitch2D.height = sizey; //配列yサイズ

resDesc.res.pitch2D.devPtr = pDevSrc; //デバイスメモリのポインタ

resDesc.resType = cudaResourceTypePitch2D;

//テクスチャオブジェクト　生成

cudaTextureObject_t texObj;

cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

・・・

//カーネル

dim3 block(16, 16, 1);

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

kernel_mean_texture2 << <grid, block >> > (texObj, pDevResult, sizex, sizey);

・・・

//テクスチャオブジェクト　破棄

cudaDestroyTextureObject(texObj);

カーネルはこのようになります。


__global__ void kernel_mean_texture2(cudaTextureObject_t texObj, BYTE *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 < 0 || idy < 0 || 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<BYTE>(texObj, idx + x - 1, idy + y - 1); //２次元のテクスチャメモリにアクセス
		}
	}

	pResult[idy * sizex + idx] = sum / 9; //結果を書き込み

	return;
}

__global__ void kernel_mean_texture2(cudaTextureObject_t texObj, BYTE *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 < 0 || idy < 0 || 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<BYTE>(texObj, idx + x - 1, idy + y - 1); //２次元のテクスチャメモリにアクセス

}

pResult[idy * sizex + idx] = sum / 9; //結果を書き込み

return;

}

このコードでは、カーネル内で配列範囲外が読み込まれますが、ソフトが落ちるということはありません。

設定

cudaTextureDesc 　

addressMode　：　テクスチャの範囲外を指定したときの読み出しモード
　・cudaAddressModeClamp　：　直近の有効な要素の値にする座標が正規化されていないこと
　・cudaAddressModeBorder　：　範囲外なら０を返す？座標が正規化されていないこと

filterMode　：　テクスチャ読み出し時のフィルタモード
　・ cudaFilterModePoint ：最近傍の座標の値を読み出す
　・ cudaFilterModeLinear：座標で小数点以下を指定したとき線形補完した値を読み出す
　　これを選択するとreadMode で正規化する必要がありそう

readMode　：　テクスチャ読み出し時のモード
　・ cudaReadModeElementType：通常通り読み出す
　・ cudaReadModeNormalizedFloat：０～１に正規化して読み出す。
　　1要素が1Byteなら255を掛ける必要がある

読み出しモードなどを変えたとき、カーネルでtex2D<BYTE> を tex2D<FLOAT>などに変えてやる必要があるので注意してください。