GPU儲存模型

1.global memory

在 CUDA 中，一般的資料複製到的顯卡記憶體的部份，稱為 global memory 。這些記憶體是沒有 cache 的，而且，存取 global memory 所需要的時間（即 latency）是非常長的，通常是數百個 cycles。由於 global memory 並沒有 cache ，所以要避開巨大的 latency 的方法，就是要利用大量的 threads。假設現在有大量的threads 在同時執行，那麼當一個 thread 讀取記憶體，開始等待結果的時候，GPU 就可以立刻切換到下一個 thread，並讀取下一個記憶體位置。因此，理想上當 thread 的數目夠多的時候，就可以完全把 global memory 的巨大 latency 隱藏起來了。

2.記憶體存模數式

顯卡上的記憶體是 DRAM，因此最有效率的存取方式，是以連續的方式存取。當一個 thread 在等待記憶體的資料時，GPU 會切換到下一個 thread。也就是說，實際上執行的順序是類似 thread 0 -> thread 1 -> thread 2 -> ...

理論上 256 個 threads 最多隻能隱藏 256 cycles 的 latency。但是 GPU 存取 global memory 時的 latency 可能高達 500 cycles 以上。

3.塊

一個 block 中的 thread ，具有一個共用的 shared memory ，也可以進行同步工作。不同 block  之間的 thread 則不行。

一個 block  內的 thread 可以有共用的記憶體，也可以進行同步。

利用 __shared__ 聲明的變數表示這是 shared memory ，是一個 block 中每個 thread 都共用的記憶體。它會使用在 GPU 上的記憶體，所以存取的速度相當快，不需要擔心 latency 的問題。

__syncthreads() 是一個 CUDA 的內建函式，表示 block 中所有的 thread 都要同步到這個點，才能繼續執行。

4.體繫結構(以下皆以G80為準)

G80多核流處理器(Streaming Multiprocessor  SM)只有一個。每個SM中含8個流處理器(Streaming Processor SP)。

在SM中，warp是線程調度的單位，每個warp包含32個線程。

G80中，每個SM中可駐留8個塊，每個SM中最多可駐留768個線程，於是每個SM最多可駐留768/32=24個warp。

每個SM中寄存器的大小為8KB，每個SM的共用儲存空間大小為16KB，共用儲存空間是分配給塊使用的。