
テクスチャメモリ(TextureMemory) とは、グローバルメモリと同様に扱えるメモリです。
使用するメリットですが、スレッド内で連続したメモリ領域にアクセスした場合、 キャッシュが効くようなので グローバルメモリよりも高速化できる可能性があります。
ちょっとしたコードの追加で使えるようになるので、試しに使ってみましょう。
API
テクスチャメモリをバインドする際に使用するAPIです。以下は2次元で読み出すときに使います。
1次元や3次元で使用できるものもありますが、今回は省略します。
API名 | 解説 |
cudaBindTexture2D | デバイスメモリにテクスチャメモリをバインドします。 |
cudaUnbindTexture | テクスチャメモリのバインドを解除します。 |
tex2D | カーネル内で使用します。テクスチャメモリからデータを取得します。 |
コード例
コード例を示します。
こちらで示したコードを、テクスチャメモリ化しています。

カーネルは以下の通り。
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 |
//テクスチャのオブジェクト <int型,2次元> グローバルで宣言 texture<int, cudaTextureType2D> tex; //テクスチャメモリで平均を取得するカーネル __global__ void kernel_mean_texture(int *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 < 1 || idy < 1 || 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(tex, idx + x - 1, idy + y - 1); //2次元のテクスチャメモリにアクセス } } pResult[idy * sizex + idx] = sum / 9; //結果を書き込み 書き込みにはテクスチャメモリを使えない return; } |
処理関数は以下の通り。
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 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 |
//処理関数 void Process() { int sizex = 8192; int sizey = 8192; int length = sizex * sizey; size_t size = sizeof(int) * length; //ホストメモリのポインタ int* pHostSrc; //数値 int* pHostResult; //結果 //デバイスメモリのポインタ int* pDevSrc; //数値 int* pDevResult; //結果 //ホストメモリの確保 cudaMallocHost(&pHostSrc, size); cudaMallocHost(&pHostResult, size); //デバイスメモリの確保 cudaMalloc(&pDevSrc, size); cudaMalloc(&pDevResult, size); //ホストメモリに適当な値を設定 for (int n = 0; n < length; n++) { pHostSrc[n] = n; } //グローバルメモリをテクスチャメモリ(2次元)にバインド cudaBindTexture2D(NULL, tex, pDevSrc, cudaCreateChannelDesc<int>(), sizex, sizey, sizeof(int) * sizex); //ホスト->デバイスへ数値を転送 cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice); //カーネル dim3 block(16, 16, 1); dim3 grid((sizex + 16 - 1) / 16, (sizey + 16 - 1) / 16, 1); kernel_mean_texture <<<grid, block >>> (pDevResult, sizex, sizey); //デバイス->ホストへ結果を転送 cudaMemcpy(pHostResult, pDevResult, size, cudaMemcpyDeviceToHost); //バインドを解除 cudaUnbindTexture(tex); //デバイスメモリの開放 cudaFree(pDevSrc); cudaFree(pDevResult); //ホストメモリの開放 cudaFreeHost(pHostSrc); cudaFreeHost(pHostResult); } |
バインド時の引数でcudaCreateChannelDesc<int>を指定することで、tex2D使用時に4Byteを1要素として読み出すことができます。
int を unsigned char や short に置き換えると、1Byte,2Byteを1要素として読み出せるようになります。
速度比較
テクスチャメモリ使用時・未使用時のカーネルの処理時間を比較しました。
なぜか若干遅くなっていますね・・・どのようなケースで高速化されるかもう少し検証が必要そうです。
環境(CPU:Intel Corei7 8700 3.2GHz GPU:GeForce GTX1070)
テクスチャメモリ未使用 | 3.7msec |
テクスチャメモリ使用 | 4.3msec |
とりあえず、高速にはなりませんでしたが、メモリ範囲外を読んだときに指定値で自動補完してくれたり、tex2Dのx,y引数に整数以外を指定すると補完値を返すモードがあったりして便利なので使う価値はあるでしょう。
補足
・読み取り専用なので、書き込みには使えません。
・メモリサイズは512byteでアライメントされている必要があります。(512の倍数)
・単一のテクスチャのオブジェクトが(CPU側の)複数スレッドから同時にアクセスされないように注意。
・X、Yどちらの方向のメモリにアクセスしてもキャッシュが効くようです。
・デバッグモードだと結構遅くなります。
作成方法のパターン2はこちら。

