CUDA TEXTURE MEMORY

Part.1 簡介

在執行 CUDA 程式前，都要把資料先從 Host 的記憶體，複製一份到 device 的記憶體中；一般來說，這樣的部分，都是使用 device 的 global memory 來直接進行存取。不過實際上，有的時候還有別的選擇的～在《nVidia CUDA 簡介》中一文就有提到，除了 global memory 外，還可以透過 constant memory 或 texture memory 的形式，來對 device memory 資料的存取。


texture 是一般 graphics 裡的名詞，2D texture 大致上可以理解為一張圖片，一般應該是翻譯成材質（維基百科是稱為「紋理」）；而由於在傳統的 render pipeline 中，texture 佔了很重要的地位，所以在顯示卡的部分，也會對這部分做特殊的最佳化。而在 nVidia 的 CUDA 中，也把 texture 這項元素保留下來了！


如果在 CUDA 中把 device memory 的資料，當作使用 texture 的話，那資料會變成是唯讀的，要透過特殊的函式來讀取，沒有辦法進行修改；不過相對的，和 global memory 或 constant memory 比起來，也有不少優點～（詳細資料請參考《CUDA Programming Guide 1.1》的 5.4）

• 有快取，在某些狀況下頻寬會比較大  

• 不像 global/constant memory 要依照某些存取模式下才有比較好的效能  

• 位置計算的延遲被隱藏得更好，對於隨機存取資料的效能可能更好  

• 被封包過的資料可以在一個運算中被個別的變數使用  

• 8bit 和 16bit 整數可以簡單的轉換成 [0.0, 1.0] 或 [-1.0, 1.0]

除了上述的優點外，如果使用 CUDA Array 來當 texture 的話，更可以套用 filter、使用內建的功能來做內插取值、並且設定要用 clamping 或 repeat 模式處理邊界。


而在 CUDA 中，要使用 texture 的話，是要使用所謂的「texture reference」；下面就大概來介紹 CUDA 中 texture 的使用方法。首先，他的大致流程會是：

• Host 的 texture 建立部分  
  • 宣告出 texture reference  
  • 透過 Bind Texture 的函式，將 texture reference 和現有的 device memory 上的變數（linear memory 或 CUDA array）做連結
• Device 使用  
  • 透過 CUDA 提供的 texture 讀取函式（tex1Dfetch, tex1D, tex2D）來讀取 texture 的內容
• Host 刪除 texture  
  • 呼叫 unbind texture 的函式，將 texture reference 的資源釋放

不過，在 CUDA 中，對於 texture 的控制，有 low-level 和 high level 兩種，Heresy 在這就先針對比較簡單的 high-level 方法來做一些簡單的說明。


CUDA 的 texture 型別

在 CUDA 中，有提供名為 texture 的 template 型別，他的形式是：

2. texture<Type, Dim, ReadMode> texRef;

复制代码

其中：

• Type
  texture 中元素的資料型別；可以是一般的基本的 int, float 型別，也可以是 CUDA 中提供的 vector 型別。

• Dim
  代表這個 texture 的維度，在 CUDA 中只有 1 或 2 兩種值，並不支援 3D texture。

• ReadMode
  讀取 texture 的模式，有 cudaReadModeNormalizedFloat 和 cudaReadModeElementType 兩種模式。當模式是 cudaReadModeElementType 時，資料會以原來的方式讀取出來；當模式是 cudaReadModeNormalizedFloat、且資料型別是整數型別時，他則會對資料進行 normalize，傳回 [0,1] 或 [-1,1] 之間的數（視原始型別是否為 unsigned 決定）。

舉個例子，如果我們要宣告一個資料型別是 int 的 1D texture，就可以寫成：

2. 
   
3. texture<int, 1, cudaReadModeElementType> texRef;

复制代码

而 texRef 的資料，則還要再透過 BindTexture 的函式，來和 device 上的變數做連結，這樣才算完成 texture 的使用前準備。
不過另外一個要注意的是，目前的 CUDA 似乎只允許將 texture reference 宣告成 globalvariable，而無法將它宣告在函式內，用參數的方法傳遞給 kernel function（Heresy 這樣寫會使得 nvcc產生內部錯誤，要寫在 file-scope 是參考 ISI 的課程後才知道的）。

兩種不同的 texture 資料

在 CUDA 中，可以接受兩種資料：linear memory 和 CUDA array。其中，linear memory 就是之前提過，用 cudaMalloc()cudaMallocArray()，來宣告出一塊連續的 1D/2D 陣列。 宣告出的連續記憶體空間（一維陣列）；而 CUDA array 則是透過 由兩種不同的資料所建立出來的 texture reference，在使用上有一些不同的性質：

	Texture with linear memory	Texture with CUDA Array
維度	只有一維	一維或二維
讀取的索引	整數	整數或浮點數
filter	不支援	支援 (cudaFilterModeLinear / cudaFilterModePoint)
邊界值	n/a	cudaAddressModeClamp / cudaAddressModeWrap
讀取函式	tex1Dfetch()	tex1D() / tex2D()

透過 CUDA array 的 texture，可以讓 CUDA 直接幫忙做好內插（雖然目前看來只有線性內插）、邊值處理的問題，其實在很多時候都是相當方便的～

而由於使用 linear memory 和 CUDA array 的 texture 在使用上有不少的差異，所以接下來就分開寫吧