テクスチャメモリ†CUDAではグローバルメモリ,ローカルメモリ,シェアードメモリの他に特殊なメモリとして,コンスタントメモリ,テクスチャメモリがある. これらのメモリはオンチップのキャッシュが用意されており,このキャッシュが聞く限りはかなり高速に実行できる. ここでは,テクスチャメモリの使用方法について述べる. CUDA Array(cudaArray)を使用†CUDA Arrayを使ったテクスチャメモリの使用手順は,
テクスチャリファレンスの作成†テクスチャの属性のいくつかはコンパイル時に既知でなければならないため, テンプレートを使ってテクスチャリファレンス変数宣言時に指定する. texture<DataType, Type, ReadMode> texRef; ここで,
例えば,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; }; ここで,それぞれの変数は,
テクスチャパラメータの指定例は以下. 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; } } |