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

研究生筆記: Survey Paper的重要性

Hello 各位或是我XD 今天要跟分享一個大家都懂，但又時常犯的錯誤。就是對於Paper Survey 的忽視...

有些時候我們做研究，以為這項東西沒人做就貿然投入下去弄，到最後時常發生這東西已經有人做過或是發過paper的情況...我就是一個活生生的例子 (老馬了...)
所以大家在做研究前，一定要！務必要！做相關研究的paper survey！看過，做過筆記後，確定沒有再下去弄！

另外別人跟你說的未必都是對的，必須你問過，或是驗證過才能相信，對於別人嘴砲要有抵抗力！千萬別亂相信！這就跟詐騙集團的例子很像，你以為真的沒人弄，真的這麼easy嗎？請求正過後再說，不然免談！

這邊分享給各位研究生及自己做一個警惕，希望日後研究不會在發生這種慘事(不要再老馬囉 >.^)

媲美Matlab的Freeware: Octave (附加MAC安裝教學)

相信很多學工程或是數學計算的人對Matlab都不會感到陌生。只是對於Matlab的取得方式我想不是正版的話就是去找一些神奇的版本...

這邊推薦一套類似Matlab的軟體：Octave。Octave 跟 Matlab很非常像，只是不用錢而已 XD
有興趣的朋友可以去這個網站看看：http://www.gnu.org/software/octave/

Windows 的安裝很簡單(點到底就對了)，這邊就不提了。Linux的話其實也很好裝，如果是Ubuntu的話就apt-cache search看看吧～這邊我就用MAC OSX 做一個簡單的安裝教學。

Octave on MAC tutorial:

1. 首先到 http://octave.sourceforge.net/ 去下載For OSX的Octave

2. 打開下載的檔案，然後把Octave drag to Application 就ok了～(沒錯就這麼簡單XD)

3. 為了讓Octave也可以plot graphic，必須安裝gnuplot。但不用重新下載，他就已經在octave原本的dmg裡面了！打開Extras，然後一樣把gunplot drag to Application就ok了！

4. 但這邊的gnuplot有一個bug，會導致你在畫圖的時候出現這個訊息：


dyld: Library not loaded: /usr/local/lib/libaquaterm.1.0.0.dylib
  Referenced from: /usr/local/bin/gnuplot
  Reason: image not found
dyld: Library not loaded: /usr/local/lib/libaquaterm.1.0.0.dylib
  Referenced from: /usr/local/bin/gnuplot
  Reason: image not found
error: you must have gnuplot installed to display graphics; if you have gnuplot installed in a non-standard location, see the 'gnuplot_binary' function

看起來很嚇人對吧？其實是gnuplot link library不對所導致。但沒關係，把它改對就好：
     a. 打開你的terminal（如果有人沒用過的話就用MAC search terminal）
     b. cd /Applications/Gnuplot.app/Contents/Resources/bin (這行指令是換到Gnuplot裡面的bin資料夾)
     c. vim gnuplot (用vim文字處理器去編輯gnuplot)
     d. 到第34行把 DYLD_LIBRARY_PATH="${ROOT}/lib:${DYLD_LIBRARY_PATH}" 改成 DYLD_LIBRARY_PATH="${ROOT}/lib"。存檔後跳出
     e.重開octave再試一次，應該就ok了！


這邊我只有寫在MAC上安裝的流程跟遇到問題的解法，如果有其他人對於MAC安裝上還有問題的在一起互相交流吧！

PS:另外最新版的Gnuplot好像都有用Aquaterm了，所以就不用另外再安裝囉～