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

【CUDA】テクスチャメモリの使用方法 その1
テクスチャメモリ(TextureMemory)の解説を行います。
スレッド内で連続したメモリ領域にアクセスした場合、 キャッシュが効くので グローバルメモリよりも高速化できる可能性があります。
コード
2Dで呼び出す場合のコード例を示します。
以下はカーネル呼び出し側のコードです。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 |
//デバイスメモリのポインタ 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); |
カーネルはこのようになります。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 |
__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); //2次元のテクスチャメモリにアクセス } } pResult[idy * sizex + idx] = sum / 9; //結果を書き込み return; } |
このコードでは、カーネル内で配列範囲外が読み込まれますが、ソフトが落ちるということはありません。
設定
cudaTextureDesc
addressMode : テクスチャの範囲外を指定したときの読み出しモード
・cudaAddressModeClamp : 直近の有効な要素の値にする 座標が正規化されていないこと
・cudaAddressModeBorder : 範囲外なら0を返す? 座標が正規化されていないこと
filterMode : テクスチャ読み出し時のフィルタモード
・ cudaFilterModePoint : 最近傍の座標の値を読み出す
・ cudaFilterModeLinear:座標で小数点以下を指定したとき線形補完した値を読み出す
これを選択するとreadMode で正規化する必要がありそう
readMode : テクスチャ読み出し時のモード
・ cudaReadModeElementType: 通常通り読み出す
・ cudaReadModeNormalizedFloat:0~1に正規化して読み出す。
1要素が1Byteなら255を掛ける必要がある
読み出しモードなどを変えたとき、カーネルでtex2D<BYTE> を tex2D<FLOAT>などに変えてやる必要があるので注意してください。
