這裡我們會簡單介紹，NVIDIA 目前支援 CUDA 的 GPU，其在執行 CUDA 程式的部份（基本上就是其 shader 單元）的架構。這裡的資料是綜合 NVIDIA 所公布的資訊，以及 NVIDIA 在各個研討會、學校課程等所提供的資料，因此有可能會有不正確的地方。主要的資料來源包括 NVIDIA 的 CUDA Programming Guide 1.1、NVIDIA 在 Supercomputing '07 介紹 CUDA 的 session，以及 UIUC 的 CUDA 課程。 GPU 的基本介紹目前 NVIDIA 推出的顯示晶片，支援 CUDA 的是 G80 系列的顯示晶片。其中 G80 顯示晶片支援 CUDA 1.0 版，而 G84、G86、G92、G94、G96 則支援 CUDA 1.1 版。基本上，除了最早的 GeForce 8800 Ultra/GTX 及 320MB/640MB 版本的 GeForce 8800GTS、Tesla 等顯示卡是 CUDA 1.0 版之外，其它 GeForce 8 系列及 9 系列顯示卡都支援 CUDA 1.1。詳細情形可以參考 CUDA Programming Guide 1.1 的 Appendix A。 所有目前支援 CUDA 的 NVIDIA 顯示晶片，其 shader 部份都是由多個 multiprocessors 組成。每個 multiprocessor 裡包含了八個 stream processors，其組成是四個四個一組，也就是說實際上可以看成是有兩組 4D 的 SIMD 處理器。此外，每個 multiprocessor 還具有 8192 個暫存器，16KB 的 share memory，以及 texture cache 和 constant cache。大致上如下圖所示： 
詳細的 multiprocessor 資訊，都可以透過 CUDA 的 cudaGetDeviceProperties() 函式或 cuDeviceGetProperties() 函式取得。不過，目前還沒有辦法直接取得一個顯示晶片中有多少 multiprocessor 的資訊。 在 CUDA 中，大部份基本的運算動作，都可以由 stream processor 進行。每個 stream processor 都包含一個 FMA（fused-multiply-add）單元，可以進行一個乘法和一個加法。比較複雜的運算則會需要比較長的時間。 執行過程在執行 CUDA 程式的時候，每個 stream processor 就是對應一個 thread。每個 multiprocessor 則對應一個 block。從之前的文章中，可以注意到一個 block 經常有很多個 thread（例如 256 個），遠超過一個 multiprocessor 所有的 stream processor 數目。這又是怎麼回事呢？ 實際上，雖然一個 multiprocessor 只有八個 stream processor，但是由於 stream processor 進行各種運算都有 latency，更不用提記憶體存取的 latency，因此 CUDA 在執行程式的時候，是以 warp 為單位。目前的 CUDA 裝置，一個 warp 裡面有 32 個 threads，分成兩組 16 threads 的 half-warp。由於 stream processor 的運算至少有 4 cycles 的 latency，因此對一個 4D 的 stream processors 來說，一次至少執行 16 個 threads（即 half-warp）才能有效隱藏各種運算的 latency。 由於 multiprocessor 中並沒有太多別的記憶體，因此每個 thread 的狀態都是直接保存在 multiprocessor 的暫存器中。所以，如果一個 multiprocessor 同時有愈多的 thread 要執行，就會需要愈多的暫存器空間。例如，假設一個 block 裡面有 256 個 threads，每個 thread 用到 20 個暫存器，那麼總共就需要 256x20 = 5,120 個暫存器才能保存每個 thread 的狀態。目前 CUDA 裝置中每個 multiprocessor 有 8,192 個暫存器，因此，如果每個 thread 使用到 16 個暫存器，那就表示一個 multiprocessor 同時最多只能維持 512 個 thread 的執行。如果同時進行的 thread 數目超過這個數字，那麼就會需要把一部份的資料儲存在顯示記憶體中，就會降低執行的效率了。 Shared memory目前 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，為了避免這個問題，可以把資料的排列方式稍加修改，把存取方式改成： int row = tid / 16;
int column = tid % 16;
data[row * 17 + column] = global_data[tid];
...
int number = data[17 * tid]; 這樣就不會造成 bank conflict 了。 Global memory由於 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 讀取的問題。 TextureCUDA 支援 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 要來得快。 運算單元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 等等。這些函式的速度較快，但精確度不如標準的函式。詳細的資料可以參考 CUDA Programming Guide 1.1 的 Appendix B。 和主記憶體間的資料傳輸在 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 的記憶體，將會影響到作業系統對記憶體的管理，可能會減低系統的效率。 | 
