Part.1 簡(jiǎn)介

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


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


如果在 CUDA 中把 device memory 的資料，當(dāng)作使用 texture 的話，那資料會(huì)變成是唯讀的，要透過特殊的函式來(lái)讀取，沒有辦法進(jìn)行修改；不過相對(duì)的，和 global memory 或 constant memory 比起來(lái)，也有不少優(yōu)點(diǎn)～（詳細(xì)資料請(qǐng)參考《CUDA Programming Guide 1.1》的 5.4）

• 有快取，在某些狀況下頻寬會(huì)比較大??

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

• 位置計(jì)算的延遲被隱藏得更好，對(duì)於隨機(jī)存取資料的效能可能更好??

• 被封包過的資料可以在一個(gè)運(yùn)算中被個(gè)別的變數(shù)使用??

• 8bit 和 16bit 整數(shù)可以簡(jiǎn)單的轉(zhuǎn)換成 [0.0, 1.0] 或 [-1.0, 1.0]

除了上述的優(yōu)點(diǎn)外，如果使用 CUDA Array 來(lái)當(dāng) texture 的話，更可以套用 filter、使用內(nèi)建的功能來(lái)做內(nèi)插取值、並且設(shè)定要用 clamping 或 repeat 模式處理邊界。


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

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

不過，在 CUDA 中，對(duì)於 texture 的控制，有 low-level 和 high level 兩種，Heresy 在這就先針對(duì)比較簡(jiǎn)單的 high-level 方法來(lái)做一些簡(jiǎn)單的說明。


CUDA 的 texture 型別
在 CUDA 中，有提供名為 texture 的 template 型別，他的形式是：
texture<Type, Dim, ReadMode> texRef;復(fù)制代碼
其中：

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

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

• ReadMode
  讀取 texture 的模式，有 cudaReadModeNormalizedFloat 和 cudaReadModeElementType 兩種模式。當(dāng)模式是 cudaReadModeElementType 時(shí)，資料會(huì)以原來(lái)的方式讀取出來(lái)；當(dāng)模式是 cudaReadModeNormalizedFloat、且資料型別是整數(shù)型別時(shí)，他則會(huì)對(duì)資料進(jìn)行 normalize，傳回 [0,1] 或 [-1,1] 之間的數(shù)（視原始型別是否為 unsigned 決定）。

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

texture<int, 1, cudaReadModeElementType> texRef;復(fù)制代碼
而 texRef 的資料，則還要再透過 BindTexture 的函式，來(lái)和 device 上的變數(shù)做連結(jié)，這樣才算完成 texture 的使用前準(zhǔn)備。
不過另外一個(gè)要注意的是，目前的 CUDA 似乎只允許將 texture reference 宣告成 globalvariable，而無(wú)法將它宣告在函式內(nèi)，用參數(shù)的方法傳遞給 kernel function（Heresy 這樣寫會(huì)使得 nvcc產(chǎn)生內(nèi)部錯(cuò)誤，要寫在 file-scope 是參考 ISI 的課程後才知道的）。



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

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

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

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

轉(zhuǎn)載于:https://blog.51cto.com/general/239036

總結(jié)

以上是生活随笔為你收集整理的CUDA Texture Memory的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。