TEGRA 3 登場!

NVIDIA 在近期發佈了全新的手機運算核心: Tegra 3 ! (Project Kal-El)

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

[轉錄] CUDA速成篇(下)

CUDA速成的最後一篇

[轉錄] CUDA速成篇(中)

延續上一篇...

[轉錄] CUDA速成篇(上)

最近在複習CUDA的一些相關知識，就把這篇在PTT上，由a5000ml大大發的教學文拿出來看。寫得非常淺顯易懂，非常適合新手或是有經驗人士閱讀學習 :)

測量CUDA kernel的執行時間

以下是我找到一些測試kernel function的方法:

1. cutCreateTimer : 單位 ms ->據說在GPU上比較準(我自己實際測好像也是如此)

#include <cutil_inline.h>
....

// create and start timer
     unsigned int timer = 0;
     cutilCheckError(cutCreateTimer(&timer));
     cutilCheckError(cutStartTimer(timer));

....

// stop and destroy timer

    cutilCheckError(cutStopTimer(timer));

    printf("GPU Processing time: %f (ms) \n", cutGetTimerValue(timer));

    cutilCheckError(cutDeleteTimer(timer));

2. cuda event: 單位 ms ->還在測試準度...感覺沒上面來的精準

.....

cudaEvent_t gpu0,gpu1;
cudaEventCreate(&gpu0);

cudaEventCreate(&gpu1);
...
// Start
cudaEventRecord(gpu0, 0);

.... // your program
// stop
cudaEventRecord(gpu1, 0);
cudaEventSynchronize(gpu1); // not sure if we have to add this one....

float gpu_time;
cudaEventElapsedTime(&gputime, gpu0, gpu1);
printf(" %f " ,gpu_time);

CUDA 3.2設定

NVIDIA CUDA 3.2終於從RC變成正式版了。有興趣可以前往此網址下載:
http://developer.nvidia.com/object/cuda_3_2_downloads.html
關於詳細的更新內容，就上網查看囉~
不過這次SDK資料夾有點改變...害我之前好多程式都不能跑@@這邊附上今天改寫好的路徑，這樣下次就不會忘了~

3.2版:

$(CUDA_INC_PATH);C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 3.2\C\common\inc;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 3.2\shared\inc
$(CUDA_LIB_PATH);C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 3.2\C\common\lib;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 3.2\shared\lib
cudart.lib cutil32D.lib shrUtils32D.lib

有衝突時使用(cpp or c):
C/C++ Code Generation: Runtime Library (/MTD)


3.2以前的版本就把在路徑裡面的" 3.2"移除就可以了~

PS: VS2010好像還是不太好設定一些CUDA的參數路徑....還是先用VS2008吧~