[CUDA] Share memory conflict 筆記 （From google到的資料）

這是很久之前google到有關CUDA shared memory的資料，備份給自己一下XDD


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 了。