技术开发 频道

CUDA Texture Part.3 CUDA Array

  【IT168 文档】在 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);
0
相关文章