テクスチャメモリ
をテンプレートにして作成
[
トップ
|
新規
|
一覧
|
検索
|
最終更新
|
ヘルプ
]
開始行:
----
#contents
----
*テクスチャメモリ [#y2a18b12]
CUDAではグローバルメモリ，ローカルメモリ，シェアードメモ...
これらのメモリはオンチップのキャッシュが用意されており，...
ここでは，テクスチャメモリの使用方法について述べる．
**CUDA Array(cudaArray)を使用 [#uf97a47b]
CUDA Arrayを使ったテクスチャメモリの使用手順は，
+テクスチャリファレンスの作成(カーネル関数から読めるよう...
+CUDA Arrayの確保とホストからのデータ転送
+テクスチャパラメータの設定
+CUDA Arrayをテクスチャにバインド
+カーネル内でtex2Dなどの参照関数を用いて値を取得
**テクスチャリファレンスの作成 [#eccc10a8]
テクスチャの属性のいくつかはコンパイル時に既知でなければ...
テンプレートを使ってテクスチャリファレンス変数宣言時に指...
#code(C){{
texture<DataType, Type, ReadMode> texRef;
}}
ここで，
-DataTypeはデータ型で基本整数型，単精度浮動小数点型(float...
-Typeはテクスチャの形状を表し，1D,2D,3Dテクスチャの場合，...
そのほかに，1D,2Dレイヤーテクスチャ(もしくはtexture array...
ただ，texture_type.hを見るとcudaTextureType1D,2D,3Dにはそ...
-ReadModeはオプション引数であり，cudaReadModeElementType(...
cudaReadModeNormalizedFloat(8ビットまたは16ビットのintで...
符号付で[-1,1]、符号なしで[0,1]を返す)のどちらかを指定す...
例えば，2次元の浮動小数点型のテクスチャでは，
#code(C){{
texture<float, cudaTextureType2D, cudaReadModeElementType...
}}
となる．テクスチャリファレンス変数をカーネル関数内で参照...
***CUDA Arrayの確保とホストからのデータ転送 [#h6f09fd1]
カーネルを呼び出す前に，テクスチャにデータを転送する(ホス...
ここではCUDA Arrayを使った方法でテータ転送する．
まず，テクスチャ用のデバイスメモリ割り当てのために，テク...
#code(C){{
cudaChannelFormatDesc cdesc = cudaCreateChannelDesc(32, 0...
}}
cudaCreateChannelDesc関数の最初の4引数はテクスチャの各チ...
ここでは，上記のテクスチャリファレンス変数の宣言でTypeにf...
float2やfloat4ならば，32, 32, 0, 0 や32, 32, 32, 32 とす...
最後の引数は型の種類を示し，cudaChannelFormatKindSigned (...
cudaChannelFormatKindFloat (浮動小数点型の場合)のいずれか...
次に，デバイスメモリにCUDA Array(cudaArray)を確保し，ホス...
#code(C){{
cudaArray *cu_array;
cutilSafeCall(cudaMallocArray(&cu_array, &cdesc, width, h...
cutilSafeCall(cudaMemcpyToArray(cu_array, 0, 0, hData, si...
}}
ここで，width, heightが2次元配列のサイズ，size=width*heig...
hDataはホストメモリのデータポインタである．
cudaMallocArray関数でデバイスメモリ確保時にcudaChannelFor...
CUDA Arrayを解放するときは以下のようにする．
#code(C){{
cutilSafeCall(cudaFreeArray(cu_array));
}}
***テクスチャパラメータの設定 [#z303d20a]
使用するテクスチャのパラメータを設定する(ホストコード)．
上記で宣言した，texture型は，低レベルAPIで以下のように定...
高レベルAPI内で定義されている．
#code(C){{
struct textureReference
{
int normalized;
enum cudaTextureFilterMode filterMode;
enum cudaTextureAddressMode addressMode[3];
struct cudaChannelFormatDesc channelDesc;
};
}}
ここで，それぞれの変数は，
-normalized : テクスチャ座標を正規化するかどうか．0ならば...
0以外ならば[0,1]となる．
-filterMode : フィルタリングモード(テクスチャをフェッチし...
cudaFilterModePoint か cudaFilterModeLinear(返値が浮動小...
cudaFilterModePointは入力テクスチャ座標の最近傍テクセルの...
cudaFilterModeLinearは入力テクスチャ座標の近傍テクセル(1D...
-addressMode : アドレッシングモード(テクスチャ座標の範囲...
addressMode配列のそれぞれの要素はテクスチャ座標の各次元の...
addressModeには，cudaAddressModeClamp(範囲外テクスチャ座...
を指定できる．
-channelDesc : テクスチャのデータ構造．上記参照．
テクスチャパラメータの指定例は以下．
#code(C){{
g_Tex.addressMode[0] = cudaAddressModeWrap;
g_Tex.addressMode[1] = cudaAddressModeWrap;
g_Tex.filterMode = cudaFilterModeLinear;
g_Tex.normalized = true;
}}
***CUDA Arrayをテクスチャにバインド [#k9895964]
CUDA Arrayとテクスチャリファレンスをバインドする(ホストコ...
#code(C){{
cutilSafeCall(cudaBindTextureToArray(g_Tex, cu_array, cde...
}}
第一引数でテクスチャリファレンス，第二引数でCUDA Array，...
バインド時のテクスチャフォーマット(cdesc)はテクスチャリフ...
***カーネル内でtex2Dなどの参照関数を用いて値を取得 [#y5e0...
カーネルを呼び出し，カーネル内でテクスチャフェッチ関数で...
テクスチャを参照して出力するだけのカーネルを以下に示す．
#code(C){{
__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);
}
}}
*テクスチャメモリ使用例 [#r7c8e1a6]
**行列乗算 [#mc143b4b]
[[CUDAで行列演算：乗算]]をテクスチャメモリを使用した例．
ホストコード
#code(C){{
void CuMulMatrixTexture(const Matrix hA, const Matrix hB,...
{
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,...
size = hA.width*hA.height*sizeof(float);
cutilSafeCall(cudaMallocArray(&caA, &cdesc0, hA.width, h...
cutilSafeCall(cudaMemcpyToArray(caA, 0, 0, hA.elements, ...
cudaChannelFormatDesc cdesc1 = cudaCreateChannelDesc(32,...
size = hB.width*hB.height*sizeof(float);
cutilSafeCall(cudaMallocArray(&caB, &cdesc1, hB.width, h...
cutilSafeCall(cudaMemcpyToArray(caB, 0, 0, hB.elements, ...
// テクスチャパラメータ
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...
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(t...
cutilCheckError(cutDeleteTimer(timer));
// デバイスからホストへ結果を転送
size = dC.width*dC.height*sizeof(float);
cutilSafeCall(cudaMemcpy(hC.elements, dC.elements, size,...
// デバイスメモリ解放
cutilSafeCall(cudaFreeArray(caA));
cutilSafeCall(cudaFreeArray(caB));
cutilSafeCall(cudaFree(dC.elements));
}
}}
テクスチャ定義とカーネル関数
#code(C){{
texture<float, cudaTextureType2D, cudaReadModeElementType...
texture<float, cudaTextureType2D, cudaReadModeElementType...
__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;
}
}
}}
終了行:
----
#contents
----
*テクスチャメモリ [#y2a18b12]
CUDAではグローバルメモリ，ローカルメモリ，シェアードメモ...
これらのメモリはオンチップのキャッシュが用意されており，...
ここでは，テクスチャメモリの使用方法について述べる．
**CUDA Array(cudaArray)を使用 [#uf97a47b]
CUDA Arrayを使ったテクスチャメモリの使用手順は，
+テクスチャリファレンスの作成(カーネル関数から読めるよう...
+CUDA Arrayの確保とホストからのデータ転送
+テクスチャパラメータの設定
+CUDA Arrayをテクスチャにバインド
+カーネル内でtex2Dなどの参照関数を用いて値を取得
**テクスチャリファレンスの作成 [#eccc10a8]
テクスチャの属性のいくつかはコンパイル時に既知でなければ...
テンプレートを使ってテクスチャリファレンス変数宣言時に指...
#code(C){{
texture<DataType, Type, ReadMode> texRef;
}}
ここで，
-DataTypeはデータ型で基本整数型，単精度浮動小数点型(float...
-Typeはテクスチャの形状を表し，1D,2D,3Dテクスチャの場合，...
そのほかに，1D,2Dレイヤーテクスチャ(もしくはtexture array...
ただ，texture_type.hを見るとcudaTextureType1D,2D,3Dにはそ...
-ReadModeはオプション引数であり，cudaReadModeElementType(...
cudaReadModeNormalizedFloat(8ビットまたは16ビットのintで...
符号付で[-1,1]、符号なしで[0,1]を返す)のどちらかを指定す...
例えば，2次元の浮動小数点型のテクスチャでは，
#code(C){{
texture<float, cudaTextureType2D, cudaReadModeElementType...
}}
となる．テクスチャリファレンス変数をカーネル関数内で参照...
***CUDA Arrayの確保とホストからのデータ転送 [#h6f09fd1]
カーネルを呼び出す前に，テクスチャにデータを転送する(ホス...
ここではCUDA Arrayを使った方法でテータ転送する．
まず，テクスチャ用のデバイスメモリ割り当てのために，テク...
#code(C){{
cudaChannelFormatDesc cdesc = cudaCreateChannelDesc(32, 0...
}}
cudaCreateChannelDesc関数の最初の4引数はテクスチャの各チ...
ここでは，上記のテクスチャリファレンス変数の宣言でTypeにf...
float2やfloat4ならば，32, 32, 0, 0 や32, 32, 32, 32 とす...
最後の引数は型の種類を示し，cudaChannelFormatKindSigned (...
cudaChannelFormatKindFloat (浮動小数点型の場合)のいずれか...
次に，デバイスメモリにCUDA Array(cudaArray)を確保し，ホス...
#code(C){{
cudaArray *cu_array;
cutilSafeCall(cudaMallocArray(&cu_array, &cdesc, width, h...
cutilSafeCall(cudaMemcpyToArray(cu_array, 0, 0, hData, si...
}}
ここで，width, heightが2次元配列のサイズ，size=width*heig...
hDataはホストメモリのデータポインタである．
cudaMallocArray関数でデバイスメモリ確保時にcudaChannelFor...
CUDA Arrayを解放するときは以下のようにする．
#code(C){{
cutilSafeCall(cudaFreeArray(cu_array));
}}
***テクスチャパラメータの設定 [#z303d20a]
使用するテクスチャのパラメータを設定する(ホストコード)．
上記で宣言した，texture型は，低レベルAPIで以下のように定...
高レベルAPI内で定義されている．
#code(C){{
struct textureReference
{
int normalized;
enum cudaTextureFilterMode filterMode;
enum cudaTextureAddressMode addressMode[3];
struct cudaChannelFormatDesc channelDesc;
};
}}
ここで，それぞれの変数は，
-normalized : テクスチャ座標を正規化するかどうか．0ならば...
0以外ならば[0,1]となる．
-filterMode : フィルタリングモード(テクスチャをフェッチし...
cudaFilterModePoint か cudaFilterModeLinear(返値が浮動小...
cudaFilterModePointは入力テクスチャ座標の最近傍テクセルの...
cudaFilterModeLinearは入力テクスチャ座標の近傍テクセル(1D...
-addressMode : アドレッシングモード(テクスチャ座標の範囲...
addressMode配列のそれぞれの要素はテクスチャ座標の各次元の...
addressModeには，cudaAddressModeClamp(範囲外テクスチャ座...
を指定できる．
-channelDesc : テクスチャのデータ構造．上記参照．
テクスチャパラメータの指定例は以下．
#code(C){{
g_Tex.addressMode[0] = cudaAddressModeWrap;
g_Tex.addressMode[1] = cudaAddressModeWrap;
g_Tex.filterMode = cudaFilterModeLinear;
g_Tex.normalized = true;
}}
***CUDA Arrayをテクスチャにバインド [#k9895964]
CUDA Arrayとテクスチャリファレンスをバインドする(ホストコ...
#code(C){{
cutilSafeCall(cudaBindTextureToArray(g_Tex, cu_array, cde...
}}
第一引数でテクスチャリファレンス，第二引数でCUDA Array，...
バインド時のテクスチャフォーマット(cdesc)はテクスチャリフ...
***カーネル内でtex2Dなどの参照関数を用いて値を取得 [#y5e0...
カーネルを呼び出し，カーネル内でテクスチャフェッチ関数で...
テクスチャを参照して出力するだけのカーネルを以下に示す．
#code(C){{
__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);
}
}}
*テクスチャメモリ使用例 [#r7c8e1a6]
**行列乗算 [#mc143b4b]
[[CUDAで行列演算：乗算]]をテクスチャメモリを使用した例．
ホストコード
#code(C){{
void CuMulMatrixTexture(const Matrix hA, const Matrix hB,...
{
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,...
size = hA.width*hA.height*sizeof(float);
cutilSafeCall(cudaMallocArray(&caA, &cdesc0, hA.width, h...
cutilSafeCall(cudaMemcpyToArray(caA, 0, 0, hA.elements, ...
cudaChannelFormatDesc cdesc1 = cudaCreateChannelDesc(32,...
size = hB.width*hB.height*sizeof(float);
cutilSafeCall(cudaMallocArray(&caB, &cdesc1, hB.width, h...
cutilSafeCall(cudaMemcpyToArray(caB, 0, 0, hB.elements, ...
// テクスチャパラメータ
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...
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(t...
cutilCheckError(cutDeleteTimer(timer));
// デバイスからホストへ結果を転送
size = dC.width*dC.height*sizeof(float);
cutilSafeCall(cudaMemcpy(hC.elements, dC.elements, size,...
// デバイスメモリ解放
cutilSafeCall(cudaFreeArray(caA));
cutilSafeCall(cudaFreeArray(caB));
cutilSafeCall(cudaFree(dC.elements));
}
}}
テクスチャ定義とカーネル関数
#code(C){{
texture<float, cudaTextureType2D, cudaReadModeElementType...
texture<float, cudaTextureType2D, cudaReadModeElementType...
__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;
}
}
}}
ページ名:
![[PukiWiki]](image/pukiwiki.png "[PukiWiki]")
