在 CUDA Texture 文章的第一篇大概講了一下 texture 在 CUDA 裡的基本概念,而第二篇則是講了 linear memory 的 texture,接下來,自然就是 CUDA Array 的 texture 了~
CUDA Array
CUDA array 在 cuda 中是一個特殊的資料型別,叫做 cudaArray,在 CUDA 中,他應該是專門給 texture 用的一種型別;要對他做記憶體的管裡,則是要透過 cudaMallocArray()、cudaFreeArray()、cudaMemcpyToArray() 等函式。此外,由於 cudaArray 本身並非 template 的型別,所以在透過 cudaMallocArray() 來配置記憶體時,也要透過 cudaChannelFormatDesc 這個特殊的資料型別,來設定他的資料型別。
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>(); cudaArray* cuArray; cudaMallocArray(&cuArray, &channelDesc, width, height);
上面就是一個簡單的例子,所宣告出來的 cuArray,就是一個內部資料是 float,大小是 width * height 的 CUDA Array。其中,cudaChannelFormatDesc 是一個用來描述 fetch 一個 texture 時,回傳值的資料的型別;而要產生對應型別的資料,只要使用他的 template function:
template<class T> struct cudaChannelFormatDesc cudaCreateChannelDesc<T>();而在 cudaMallocArray() 的使用上,也只需要給四個參數:cudaArray**、cudaChannelFormatDesc*、寬、高。
而相較於一般 linear memory 是用 cudaMemcpy() 來將資料由 host memory 複製到 device memory,CUDA Array 要改用 cudaMemcpyToArray() 來做複製的動作;這個函式的形式為:
cudaError_t cudaMemcpyToArray(struct cudaArray* dstArray, size_t dstX, size_t dstY, const void* src, size_t count, enum cudaMemcpyKind kind);這個函式會把來源資料 src 複製到 dstArray 中;而 cudaMemcpyKind 則是用來指定複製的方向,有 cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice 四種值。
第五個參數 count 是代表要複製的資料量;而 dstX 和 dstY 則是代表要由 src 的左上角 (dstX, dstY) 的位置開始複製資料,一般來說應該都是給 0。
Texture with CUDA array
對於使用 CUDA array 的 texture,是要使用 cudaBindTextureToArray() 這個函式來把 CUDA array 和 texture 聯繫起來;使用上,只要給他 texture 和 cudaArray 當作參數就可以了~其函式形式為:
template<class T, int dim, enum cudaTextureReadMode readMode> cudaError_t cudaBindTextureToArray( const struct texture<T, dim, readMode>& texRef, const struct cudaArray* cuArray);而要解除 texture 和 CUDA array 的關係,使用的函式和 linear memory 時是一樣的,都是 cudaUnbindTexture()。
而在存取上,和 linear memory 的 texture 時的 tex1Dfetch() 不同,是要使用 tex1D()、tex2D() 這兩個函式,分別是用在 1D 和 2D 的 texture。其形式分別為:
template<class Type, enum cudaTextureReadMode readMode> Type tex1D(texture<Type, 1, readMode> texRef, float x); template<class Type, enum cudaTextureReadMode readMode> Type tex2D(texture<Type, 2, readMode> texRef, float x, float y);
簡單範例
接下來,還是用實例來看吧~這邊是用 CUDA Array 的 texture 來做 transpose 的動作。[程式原始碼下載]
首先,main() 的內容如下:
void main( int argc, char** argv ) { 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] ) { printf( "Error!!!!n" ); break; } }一樣很簡單,就先宣告出原始資料的 aSrc,還有轉置過的資料 aRS1 和 aRS2;然後在原始資料 aSrc 中,填入一些值。(以此例,aSrc 應該是 1920*1200,aRS1 和 aRS2 應該是 1200 * 1920;不過由於在宣告成一維陣列時沒差別,所以沒特別去修改。)
而接下來,就是分別跑 CPU 版和 GPU 版的程式,並比較兩者的結果了~而 CPU 版的函式 Transpose_CPU() 內容如下:
void Transpose_CPU( unsigned char* sImg, unsigned char *tImg, int w, int h ) { int x, y, idx1, idx2; for( y = 0; y < h; ++ y ) for( x = 0; x < w; ++ x ) { idx1 = y * w + x; idx2 = x * h + y; tImg[idx2] = sImg[idx1]; } }內容應該不用多加解釋了~總之,就是根據方向的不同,採取不同的方法計算出 idx1 和 idx2 兩個記憶體空間的索引值,以此來把資料由 sImg 複製到 tImg,藉此做到轉置的動作。
而 Transpose_GPU() 所在的 .cu 檔,內容則如下:
#define BLOCK_DIM 16 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、更細節的功能可以使用~不過這邊就暫時不提了~之後有空再說吧。 :p
[…] Texture bound to CUDA Array […]
讚讚
[…] h_volume,他是一個 uchar 的陣列。而和 2D CUDA Array 時一樣,先透過一個 cudaChannelFormatDesc 來描述 cudaArray […]
讚讚
[…] 3D Texture 的支援了~在 CUDA Array 的 texture 方面,維度正式支援 3D 了!也因此,也多了 tex3D() […]
讚讚
[…] ← CUDA Texture Part.3 CUDA Array 未來 […]
讚讚
[…] CUDA Texture Part.3 CUDA Array (2008/1/29) 介紹使用 CUDA array 的 2D texture 的基本方法。 […]
讚讚