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》