在執行 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 型別，他的形式是：
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，就可以寫成：
texture<int, 1, cudaReadModeElementType> texRef;而 texRef 的資料，則還要再透過 BindTexture 的函式，來和 device 上的變數做連結，這樣才算完成 texture 的使用前準備。
不過另外一個要注意的是，目前的 CUDA 似乎只允許將 texture reference 宣告成 global variable，而無法將它宣告在函式內，用參數的方法傳遞給 kernel function（Heresy 這樣寫會使得 nvcc 產生內部錯誤，要寫在 file-scope 是參考 ISI 的課程後才知道的）。
兩種不同的 texture 資料
在 CUDA 中，可以接受兩種資料：linear memory 和 CUDA array。其中，linear memory 就是之前提過，用 cudaMalloc() 宣告出的連續記憶體空間（一維陣列）；而 CUDA array 則是透過 cudaMallocArray()，來宣告出一塊連續的 1D/2D 陣列。
由兩種不同的資料所建立出來的 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 在使用上有不少的差異，所以接下來就分開寫吧∼
- Texture bound to Linear Memory
- Texture bound to CUDA Array