CUDA 的 texture 型別
在 CUDA 中，有提供名為 texture 的 template 型別，他的形式是：
其中：
- 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，就可以寫成：
而 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 在使用上有不少的差異，所以接下來就分開寫吧～
