簡單範例
接下來，還是用實例來看吧～這邊是用 CUDA Array 的 texture 來做 transpose 的動作。[程式原始碼下載]
首先，main() 的內容如下：
{ int w = 1920, h = 1200;
// Setup test data
unsigned char *aSrc = new unsigned char[ w * h ],
*aRS1 = new unsigned char[ w * h ],
*aRS2 = new unsigned char[ w * h ];
for( int i = 0; i < w * h ; ++ i ) aSrc[i] = i % 256;
// CPU code
Transpose_CPU( aSrc, aRS1, w, h );
// GPU Code
Transpose_GPU( aSrc, aRS2, w, h );
// check
for( int i = 0; i < w * h; ++ i )
if( aRS1[i] != aRS2[i] )
break;
一樣很簡單，就先宣告出原始資料的 aSrc，還有轉置過的資料 aRS1 和 aRS2;然後在原始資料 aSrc 中，填入一些值。(以此例，aSrc 應該是 1920*1200，aRS1 和 aRS2 應該是 1200 * 1920;不過由於在宣告成一維陣列時沒差別，所以沒特別去修改。)
而接下來，就是分別跑 CPU 版和 GPU 版的程式，並比較兩者的結果了～而 CPU 版的函式 Transpose_CPU() 內容如下：
int w, int h )
{
int x, y, idx1, idx2;
for( y = 0; y < h; ++ y )
for( x = 0; x < w; ++ x )
idx2 = x * h + y;
tImg[idx2] = sImg[idx1];
}
內容應該不用多加解釋了～總之，就是根據方向的不同，採取不同的方法計算出 idx1 和 idx2 兩個記憶體空間的索引值，以此來把資料由 sImg 複製到 tImg，藉此做到轉置的動作。
而 Transpose_GPU() 所在的 .cu 檔，內容則如下：
texture<unsigned char, 2, cudaReadModeElementType> rT;
extern "C" void Transpose_GPU( unsigned char* sImg, unsigned char *tImg, int w, int h );
__global__ void Transpose_Texture( unsigned char* aRS, int w, int h )
{
int idxX = blockIdx.x * blockDim.x + threadIdx.x,
idxY = blockIdx.y * blockDim.y + threadIdx.y;
if( idxX < w && idxY < h )
aRS[ idxX * h + idxY ] = tex2D( rT, idxX, idxY );
}
void Transpose_GPU( unsigned char* sImg, unsigned char *tImg, int w, int h )
{
// compute the size of data
int data_size = sizeof(unsigned char) * w * h;
// part1a. prepare the result data
unsigned char *dImg;
cudaMalloc( (void**)&dImg, data_size );
// part1b. prepare the source data
cudaChannelFormatDesc chDesc = cudaCreateChannelDesc<unsigned char>();
cudaArray* cuArray;
cudaMallocArray(&cuArray, &chDesc, w, h);
cudaMemcpyToArray( cuArray, 0, 0, sImg, data_size, cudaMemcpyHostToDevice );
cudaBindTextureToArray( rT, cuArray );
// part2. run kernel
dim3 block( BLOCK_DIM, BLOCK_DIM ),
grid( ceil( (float)w / BLOCK_DIM), ceil( (float)h / BLOCK_DIM) );
Transpose_Texture<<< grid, block>>>( dImg, w, h );
// part3. copy the data from device
cudaMemcpy( tImg, dImg, data_size, cudaMemcpyDeviceToHost );
// par4. release data
cudaUnbindTexture( rT );
cudaFreeArray( cuArray );
cudaFree( dImg );
}
首先，之前也有提過了，目前的 CUDA 似乎只允許把 texture 宣告在 file-scope，所以一開始就要宣告一個 2D texture 來當輸入資料;說實話，對於這點 Heresy 覺得實在不是很方便。
接下來，直接看 main() 所呼叫的 Transpose_GPU() 吧～他做的內容如下：
- 先把所需要的記憶體大小計算出來
- [part1a] 宣告 dImg，並指派記憶體位址給 dImg 來儲存計算後的結果。
- [part1b] 建立 CUDA array cuArray、派記憶體位址，將資料由 host memory(sImg) 複製到 device memory(cuArray);並透過 cudaBindTextureToArray() 將 rT 和 cuArray 做聯繫。
- [part2] 呼叫 kernel function：Transpose_Texture() 來進行計算。在這邊，thread block 的大小是定義為 BLOCK_DIM*BLOCK_DIM(16*16)，grid 的大小則是根據寬和高來除以 block 的大小。
- [part3] 將結果由 device memory(dImg)複製回 host memory(tImg)。
- [part4] 透過 cudaUnbindTexture() 將 rT 和 sImg 間的聯繫解除，並使用 cudaFreeArray()、cudaFree() 將 device memory 釋放掉。
而本程式的 kernel function Transpose_Texture() 內，則是直接透過 blockIdx、blockDim、threadIdx 這三個變數，計算出二維中的位置，並在 x、y 都沒有超過範圍時，進行資料轉置的複製，把 (idxX, idxY) 的資料，透過 tex2D() 取出，儲存到 aRS[ idxX * h + idxY ]。
到此為止，應該是使用 CUDA 2D texture 最基本的方法了～實際上正如在 part.1 時所提及的，使用 CUDA Array 的 texture 其實還有一些額外的功能可以使用!而除了 high-level 的使用外，也還有 low-level、更細節的功能可以使用～不過這邊就暫時不提了～之後有空再說吧。
