cuda共享記憶體，全域性記憶體，紋理等的解釋
開始閱讀粗大資料感覺文獻1的描述講得明白，摘錄到這裡方便他人瞭解。
增加簡單排版後，摘錄
1.共享記憶體
目前 CUDA 裝置中，每個 multiprocessor 有 16KB 的 shared memory。 Shared memory 分成
16 個 bank。如果同時每個 thread 是存取不同的 bank，就不會產生任何問題，存取 shared
memory 的速度和存取暫存器相同。不過，如果同時有兩個（或更多個） threads 存取同一個
bank 的資料，就會發生 bank conflict，這些 threads 就必須照順序去存取，而無法同時存取
shared memory 了。
Shared memory 是以 4 bytes 為單位分成 banks。因此，假設以下的資料：
__shared__ int data[128];
那麼， data[0] 是 bank 0、 data[1] 是 bank 1、 data[2] 是 bank 2、 …、 data[15] 是 bank 15，
而 data[16] 又回到 bank 0。由於 warp 在執行時是以 half-warp 的方式執行，因此分屬於不
同的 half warp 的 threads，不會造成 bank conflict。
因此，如果程式在存取 shared memory 的時候，使用以下的方式：
int number = data[base + tid];
那就不會有任何 bank conflict，可以達到最高的效率。但是，如果是以下的方式：
int number = data[base + 4 * tid];
那麼， thread 0 和 thread 4 就會存取到同一個 bank， thread 1 和 thread 5 也是同樣，這樣就
會造成 bank conflict。在這個例子中，一個 half warp 的 16 個 threads 會有四個 threads 存
取同一個 bank，因此存取 share memory 的速度會變成原來的 1/4。
一個重要的例外是，當多個 thread 存取到同一個 shared memory 的地址時， shared memory 可
以將這個地址的 32 bits 資料「廣播」到所有讀取的 threads，因此不會造成 bank conflict。例
如：
int number = data[3];
這樣不會造成 bank conflict，因為所有的 thread 都讀取同一個地址的資料。
很多時候 shared memory 的 bank conflict 可以透過修改資料存放的方式來解決。例如，以下
的程式：
data[tid] = global_data[tid];
...
int number = data[16 * tid];
會造成嚴重的 bank conflict，為了避免這個問題，可以把資料的排列方式稍加修改，把存取方
data[row
* 17 + column] = global_data[tid];
...
int number = data[17 * tid];
這樣就不會造成 bank conflict 了。
編者注：share memory在NVIDIA的文件中其實還有不同的叫法，例如PDC（Parallel
Data
Cache）、PBSM（per-block
share memory）。
2.全域性記憶體
由於
multiprocessor 並沒有對 global memory 做 cache（如果每個 multiprocessor 都有自己的
global memory cache，將會需要 cache coherence protocol，會大幅增加 cache 的複雜度），所
以 global memory 存取的 latency 非常的長。除此之外，前面的文章中也提到過 global
memory 的存取，要儘可能的連續。這是因為 DRAM 存取的特性所造成的結果。
更精確的說， global memory 的存取，需要是 "coalesced"。所謂的 coalesced，是表示除了連
續之外，而且它開始的地址，必須是每個 thread 所存取的大小的 16 倍。例如，如果每個
thread 都讀取 32 bits 的資料，那麼第一個 thread 讀取的地址，必須是 16*4 = 64 bytes 的倍
數。
如果有一部份的 thread 沒有讀取記憶體，並不會影響到其它的 thread 速行 coalesced 的存取。
例如：
if(tid != 3) {
int number = data[tid];
}
雖然 thread 3 並沒有讀取資料，但是由於其它的 thread 仍符合 coalesced 的條件（假設 data
的地址是 64 bytes 的倍數），這樣的記憶體讀取仍會符合 coalesced 的條件。
在目前的 CUDA 1.1 裝置中，每個 thread 一次讀取的記憶體資料量，可以是 32 bits、 64 bits、
或 128 bits。不過， 32 bits 的效率是最好的。 64 bits 的效率會稍差，而一次讀取 128 bits 的
效率則比一次讀取 32 bits 要顯著來得低（但仍比 non-coalesced 的存取要好）。
如果每個 thread 一次存取的資料並不是 32 bits、 64 bits、或 128 bits，那就無法符合 coalesced
的條件。例如，以下的程式：
struct vec3d { float x, y, z; };
...
__global__ void func(struct vec3d* data, float* output)
{
output[tid] = data[tid].x * data[tid].x +
data[tid].y * data[tid].y +
data[tid].z * data[tid].z;
}
並不是 coalesced 的讀取，因為 vec3d 的大小是 12 bytes，而非 4 bytes、 8 bytes、或 16 bytes。
要解決這個問題，可以使用 __align(n)__ 的指示，例如：
struct __align__(16) vec3d { float x, y, z; };
這會讓 compiler 在 vec3d 後面加上一個空的 4 bytes，以補齊 16 bytes。另一個方法，是把
資料結構轉換成三個連續的陣列，例如：
__global__ void func(float* x, float* y, float* z, float* output)
{
output[tid] = x[tid] * x[tid] + y[tid] * y[tid] +
z[tid] * z[tid];
}
如果因為其它原因使資料結構無法這樣調整，也可以考慮利用 shared memory 在 GPU 上做
結構的調整。例如：
__global__ void func(struct vec3d* data, float* output)
{
__shared__ float temp[THREAD_NUM * 3];
const float* fdata = (float*) data;
temp[tid] = fdata[tid];
temp[tid + THREAD_NUM] = fdata[tid + THREAD_NUM];
temp[tid + THREAD_NUM*2] = fdata[tid + THREAD_NUM*2];
__syncthreads();
output[tid] = temp[tid*3] * temp[tid*3] +
temp[tid*3+1] * temp[tid*3+1] +
temp[tid*3+2] * temp[tid*3+2];
}
在上面的例子中，我們先用連續的方式，把資料從 global memory 讀到 shared memory。由於
shared memory 不需要擔心存取順序（但要注意 bank conflict 問題，參照前一節），所以可以
避開 non-coalesced 讀取的問題。
3.紋理
CUDA 支援 texture。在 CUDA 的 kernel 程式中，可以利用顯示晶片的 texture 單元，讀取
texture 的資料。使用 texture 和 global memory 最大的差別在於 texture 只能讀取，不能寫入，
而且顯示晶片上有一定大小的 texture cache。因此，讀取 texture 的時候，不需要符合 coalesced
的規則，也可以達到不錯的效率。此外，讀取 texture 時，也可以利用顯示晶片中的 texture
filtering 功能（例如 bilinear filtering），也可以快速轉換資料型態，例如可以直接將 32 bits
RGBA 的資料轉換成四個 32 bits 浮點數。
顯示晶片上的 texture cache 是針對一般繪圖應用所設計，因此它仍最適合有區塊性質的存取
動作，而非隨機的存取。因此，同一個 warp 中的各個 thread 最好是讀取地址相近的資料，
對於已經能符合 coalesced 規則的資料，使用 global memory 通常會比使用 texture 要來得
快。
4. 運算單元
Stream processor 裡的運算單元，基本上是一個浮點數的 fused multiply-add
單元，也就是說
它可以進行一次乘法和一次加法，如下所示：
a = b * c + d;
compiler 會自動把適當的加法和乘法運算，結合成一個 fmad 指令。
除了浮點數的加法及乘法之外，整數的加法、位運算、比較、取最小值、取最大值、及以型
態的轉換（浮點數轉整數或整數轉浮點數）都是可以全速進行的。整數的乘法則無法全速進
行，但 24 bits 的乘法則可以。在 CUDA 中可以利用內建的 __mul24 和 __umul24 函式來
進行 24 bits 的整數乘法。
浮點數的除法是利用先取倒數，再相乘的方式計算，因此精確度並不能達到 IEEE 754 的規範
（最大誤差為 2 ulp）。內建的 __fdividef(x,y) 提供更快速的除法，和一般的除法有相同的精
確度，但是在 2216<
y < 2218時會得到錯誤的結果。
此外 CUDA 還提供了一些精確度較低的內部函式，包括 __expf、__logf、__sinf、__cosf、__powf
等等。
5.和主記憶體之間資料傳輸
在 CUDA 中， GPU 不能直接存取主記憶體，只能存取顯示卡上的顯示記憶體。因此，會需要將數
據從主記憶體先複製到顯示卡記憶體中，進行運算後，再將結果從顯示卡記憶體中複製到主記憶體中。這
些複製的動作會限於 PCI Express 的速度。使用 PCI Express x16 時， PCI Express 1.0 可以提
供雙向各 4GB/s 的頻寬，而 PCI Express 2.0 則可提供 8GB/s 的頻寬。當然這都是理論值。
從一般的記憶體複製資料到顯示卡記憶體的時候，由於一般的記憶體可能隨時會被作業系統搬動，因
此 CUDA 會先將資料複製到一塊內部的記憶體中，才能利用 DMA 將資料複製到顯示卡記憶體中。
如果想要避免這個重複的複製動作，可以使用 cudaMallocHost 函式，在主記憶體中取得一塊
page locked 的記憶體。不過，如果要求太大量的 page locked 的記憶體，將會影響到
-------------------------------------------非直接轉載----------------------------------------
2016-03-10新增
最近看資料， 在 計算能力1.0 ，1.2 的時代，全域性記憶體(8個指令週期)和快取讀取(2個還是多少，但是不超過8個。)讀取時間其實差距不是非常大。之所以有很大差異是，在另一部分延遲開銷。全域性記憶體還有一部分400~600指令週期的開銷。這部分開銷從哪裡來的還不清楚。後續閱讀深入再補充。
參考文獻：
1.《深入淺出談CUDA》