テクスチャメモリ

テクスチャメモリ†

CUDAではグローバルメモリ，ローカルメモリ，シェアードメモリの他に特殊なメモリとして，コンスタントメモリ，テクスチャメモリがある． これらのメモリはオンチップのキャッシュが用意されており，このキャッシュが聞く限りはかなり高速に実行できる． ここでは，テクスチャメモリの使用方法について述べる．

CUDA Array(cudaArray)を使用†

CUDA Arrayを使ったテクスチャメモリの使用手順は，

1. テクスチャリファレンスの作成(カーネル関数から読めるようにグローバル変数として)
2. CUDA Arrayの確保とホストからのデータ転送
3. テクスチャパラメータの設定
4. CUDA Arrayをテクスチャにバインド
5. カーネル内でtex2Dなどの参照関数を用いて値を取得

テクスチャリファレンスの作成†

テクスチャの属性のいくつかはコンパイル時に既知でなければならないため， テンプレートを使ってテクスチャリファレンス変数宣言時に指定する．

texture<DataType, Type, ReadMode> texRef;

ここで，

• DataTypeはデータ型で基本整数型，単精度浮動小数点型(float)，および，これらの1,2,4要素のベクトル型のみ指定できる．
• Typeはテクスチャの形状を表し，1D,2D,3Dテクスチャの場合，cudaTextureType1D, cudaTextureType2D, cudaTextureType3Dをそれぞれ指定する． そのほかに，1D,2Dレイヤーテクスチャ(もしくはtexture array,OpenGLではGL_TEXTURE_2D_ARRAYなど)用にcudaTextureType1DLayered, cudaTextureType2DLayeredが用意されている．デフォルトは，cudaTextureType1D．(以前(CUDA2?)はこの項は単に次元(1,2,3)を指定するだけだったが，レイヤテクスチャの追加で変わったようである． ただ，texture_type.hを見るとcudaTextureType1D,2D,3Dにはそれぞれ1,2,3が割り当てられているので互換性は保たれている．)
• ReadModeはオプション引数であり，cudaReadModeElementType(デフォルト)か cudaReadModeNormalizedFloat(8ビットまたは16ビットのintで有効であり， 符号付で[-1,1]、符号なしで[0,1]を返す)のどちらかを指定する．

例えば，2次元の浮動小数点型のテクスチャでは，

texture<float, cudaTextureType2D, cudaReadModeElementType> g_Tex;

となる．テクスチャリファレンス変数をカーネル関数内で参照するために，グローバル変数としている．

CUDA Arrayの確保とホストからのデータ転送†

カーネルを呼び出す前に，テクスチャにデータを転送する(ホストコード)． ここではCUDA Arrayを使った方法でテータ転送する． まず，テクスチャ用のデバイスメモリ割り当てのために，テクスチャのデータ構造を示す型(cudaChannelFormatDesc)の変数を宣言する．

cudaChannelFormatDesc cdesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

cudaCreateChannelDesc関数の最初の4引数はテクスチャの各チャネルのビット数(0,8,16,32)を指定する． ここでは，上記のテクスチャリファレンス変数の宣言でTypeにfloat型を設定した場合を示している． float2やfloat4ならば，32, 32, 0, 0 や32, 32, 32, 32 とする． 最後の引数は型の種類を示し，cudaChannelFormatKindSigned (符号付整数型の場合)，cudaChannelFormatKindUnsigned (符号なし整数型の場合)， cudaChannelFormatKindFloat (浮動小数点型の場合)のいずれかを指定する．

次に，デバイスメモリにCUDA Array(cudaArray)を確保し，ホストメモリからデータを転送する．

cudaArray *cu_array;
cutilSafeCall(cudaMallocArray(&cu_array, &cdesc, width, height));
cutilSafeCall(cudaMemcpyToArray(cu_array, 0, 0, hData, size, cudaMemcpyHostToDevice));

ここで，width, heightが2次元配列のサイズ，size=width*height*sizeof(float)はメモリ上のサイズ， hDataはホストメモリのデータポインタである． cudaMallocArray関数でデバイスメモリ確保時にcudaChannelFormatDesc型の変数を渡している．

CUDA Arrayを解放するときは以下のようにする．

cutilSafeCall(cudaFreeArray(cu_array));

テクスチャパラメータの設定†

使用するテクスチャのパラメータを設定する(ホストコード)． 上記で宣言した，texture型は，低レベルAPIで以下のように定義されるtextureReference型のpublic派生構造体として， 高レベルAPI内で定義されている．

struct textureReference
{
  int                          normalized;
  enum cudaTextureFilterMode   filterMode;
  enum cudaTextureAddressMode  addressMode[3];
  struct cudaChannelFormatDesc channelDesc;
};

ここで，それぞれの変数は，

• normalized : テクスチャ座標を正規化するかどうか．0ならばテクスチャ座標は[0,width), [0,height), [0,depth)となり， 0以外ならば[0,1]となる．
• filterMode : フィルタリングモード(テクスチャをフェッチしたときの返値を入力テクスチャ座標に基づきどのように計算するか)を指定する． cudaFilterModePoint か cudaFilterModeLinear(返値が浮動小数点型の場合のみ)を指定できる． cudaFilterModePointは入力テクスチャ座標の最近傍テクセルの値を返し， cudaFilterModeLinearは入力テクスチャ座標の近傍テクセル(1Dで2，2Dで4，3Dで8テクセル)を線形補間した値を返す．
• addressMode : アドレッシングモード(テクスチャ座標の範囲外アクセス時の挙動)を指定する． addressMode配列のそれぞれの要素はテクスチャ座標の各次元のアドレッシングモードに対応する． addressModeには，cudaAddressModeClamp(範囲外テクスチャ座標を境界にクランプ)，cudaAddressModeWrap(範囲外テクスチャ座標を繰り返す，正規化テクスチャ座標のみ) を指定できる．
• channelDesc : テクスチャのデータ構造．上記参照．

テクスチャパラメータの指定例は以下．

g_Tex.addressMode[0] = cudaAddressModeWrap;
g_Tex.addressMode[1] = cudaAddressModeWrap;
g_Tex.filterMode     = cudaFilterModeLinear;
g_Tex.normalized     = true;

CUDA Arrayをテクスチャにバインド†

CUDA Arrayとテクスチャリファレンスをバインドする(ホストコード)．

cutilSafeCall(cudaBindTextureToArray(g_Tex, cu_array, cdesc));

第一引数でテクスチャリファレンス，第二引数でCUDA Array，最後の引数でテクスチャフォーマット(cudaChannelFormatDesc)を指定する． バインド時のテクスチャフォーマット(cdesc)はテクスチャリファレンスを宣言したときに指定したパラメータと一致しなければならない．

カーネル内でtex2Dなどの参照関数を用いて値を取得†

カーネルを呼び出し，カーネル内でテクスチャフェッチ関数で値を取得する(デバイスコード)． テクスチャを参照して出力するだけのカーネルを以下に示す．

__global__ 
void KernelFunc(float* odata, int w, int h) 
{
    // calculate normalized texture coordinates
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    float u = x/(float)w;
    float v = y/(float)h;

    odata[y*w+x] = tex2D(g_Tex, u, v);
}

テクスチャメモリ使用例†

行列乗算†

CUDAで行列演算：乗算をテクスチャメモリを使用した例．

ホストコード

void CuMulMatrixTexture(const Matrix hA, const Matrix hB, Matrix hC)
{
	cudaArray *caA, *caB;
	Matrix dC;
	dC.width = dC.stride = hC.width;	dC.height = hC.height;

	int size;

	// 結果格納用デバイスメモリの確保
	size = dC.width*dC.height*sizeof(float);
	cutilSafeCall(cudaMalloc((void**)&dC.elements, size));

	// CUDA Arrayの確保とホストからのデータ転送
	cudaChannelFormatDesc cdesc0 = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
	size = hA.width*hA.height*sizeof(float);
	cutilSafeCall(cudaMallocArray(&caA, &cdesc0, hA.width, hA.height));
	cutilSafeCall(cudaMemcpyToArray(caA, 0, 0, hA.elements, size, cudaMemcpyHostToDevice));

	cudaChannelFormatDesc cdesc1 = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
	size = hB.width*hB.height*sizeof(float);
	cutilSafeCall(cudaMallocArray(&caB, &cdesc1, hB.width, hB.height));
	cutilSafeCall(cudaMemcpyToArray(caB, 0, 0, hB.elements, size, cudaMemcpyHostToDevice));

	// テクスチャパラメータ
	g_TexA.addressMode[0] = cudaAddressModeWrap;
	g_TexA.addressMode[1] = cudaAddressModeWrap;
	g_TexA.filterMode = cudaFilterModePoint;
	g_TexA.normalized = true;	// 正規化されたテクスチャ座標でアクセス
	g_TexB.addressMode[0] = cudaAddressModeWrap;
	g_TexB.addressMode[1] = cudaAddressModeWrap;
	g_TexB.filterMode = cudaFilterModePoint;
	g_TexB.normalized = true;	// 正規化されたテクスチャ座標でアクセス

	// CUDA Arrayをテクスチャにバインド
	cutilSafeCall(cudaBindTextureToArray(g_TexA, caA, cdesc0));
	cutilSafeCall(cudaBindTextureToArray(g_TexB, caB, cdesc1));

	dim3 block(BLOCK_SIZE, BLOCK_SIZE);
	dim3 grid((dC.width+block.x-1)/block.x, (dC.height+block.y-1)/block.y);

	unsigned int timer = 0;
	cutilCheckError(cutCreateTimer(&timer));
	cutilCheckError(cutStartTimer(timer));

	matrixMulTexture<<< grid, block >>>(hA.width, dC);

	// カーネル実行エラーのチェック
	cutilCheckMsg("Kernel execution failed");

    cutilSafeCall(cudaThreadSynchronize());
	cutilCheckError(cutStopTimer(timer));
	printf("Processing time: %f (ms) \n", cutGetTimerValue(timer));
	cutilCheckError(cutDeleteTimer(timer));

	// デバイスからホストへ結果を転送
	size = dC.width*dC.height*sizeof(float);
	cutilSafeCall(cudaMemcpy(hC.elements, dC.elements, size, cudaMemcpyDeviceToHost));

	// デバイスメモリ解放
	cutilSafeCall(cudaFreeArray(caA));
	cutilSafeCall(cudaFreeArray(caB));
	cutilSafeCall(cudaFree(dC.elements));
}

テクスチャ定義とカーネル関数

texture<float, cudaTextureType2D, cudaReadModeElementType> g_TexA;
texture<float, cudaTextureType2D, cudaReadModeElementType> g_TexB;

__global__ 
void matrixMulTexture(int Awidth, Matrix C)
{
	int row = blockIdx.y*blockDim.y+threadIdx.y;
	int col = blockIdx.x*blockDim.x+threadIdx.x;

	if(row < C.height && col < C.width){
		// 正規化テクスチャ座標
		float u = (float)col/(float)C.width;
		float v = (float)row/(float)C.height;

		float x = 0.0f;
		int aw = Awidth;
		for(int k = 0; k < aw; ++k){
			float w = (float)k/(float)aw;
			x += tex2D(g_TexA, w, v)*tex2D(g_TexB, u, w);
		}

		C.elements[row*C.width+col] = x;
	}
}