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。因此，假設以下的資料：
那麼，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 的時候，使用以下的方式：
那就不會有任何 bank conflict，可以達到最高的效率。但是，如果是以下的方式：
那麼，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。例如：
這樣不會造成 bank conflict，因為所有的 thread 都讀取同一個位址的資料。
很多時候 shared memory 的 bank conflict 可以透過修改資料存放的方式來解決。例如，以下的程式：
...
int number = data[16 * tid];
會造成嚴重的 bank conflict，為了避免這個問題，可以把資料的排列方式稍加修改，把存取方式改成：
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 的存取。例如：
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 的條件。例如，以下的程式：
...
__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)__ 的指示，例如：
這會讓 compiler 在 vec3d 後面加上一個空的 4 bytes，以補齊 16 bytes。另一個方法，是把資料結構轉換成三個連續的陣列，例如：
{
output[tid] = x[tid] * x[tid] + y[tid] * y[tid] +
z[tid] * z[tid];
}
如果因為其它原因使資料結構無法這樣調整，也可以考慮利用 shared memory 在 GPU 上做結構的調整。例如：
{
__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 讀取的問題。
