cuda 學習筆記

1. __global__: Alerts the compiler that a function should be compiled to run on a device instead of the host.

2.The trick is actually in calling the device code from the host code which is handled by the CUDA compiler and runtime.

3.The hardware limits the number of blocks in a single launch to 65535. Similarly, the hardware limits the number of threads per block

  with which we can launch a kernel by the maxThreadPerBlock field of the device properties structure.

4.Cuda C compiler treats variables in shared memory differently than typical variables. It creates a copy of the variable for each block that you launch on the GPU. Every thread in that block shares the memory, but threads cannot see or modify the copy of this variable that is seen within other blocks.
5. Without synchronization, we have created a race condition where the correctness of the execution results depends on the nondeterministic details of the hardware.

6.一種關於__syncthreads() 的無窮等待的狀況。關鍵是認識到：The Cuda Architecture guarantees that no thread will advance to an instruction beyond the __syncthreads() until every thread in the block has executed __syncthreads(). For example:

      if(cacheIndex<i) {

       cache[cacheIndex] += cache[cacheIndex  + i];

      __syncthreads();

}

 若是這樣寫的話，由於不符合條件的線程不會執行到__syncthreads();而cuda runtime 會一直等待全部線程（包括這種不符合條件的線程）執行到__syncthreads();就會出現無窮等待的狀況。

 

7.Reading from constant memory can conserve memory bandwidth when compared to reading the same data from global memory.

 8. The trickiest part of using events arises as a consequence of the fact that some of the calls we make in CUDA C are actually asynchronous. For example, when we launch the kernel in our ray tracer, the GPU begins executing our code, but the CPU continues executing the next line of our program before the GPU finishes.v 

 

9.爲何Kernel函數沒有返回值？ 在主機提交Kernel一段時間後，Kernel纔開始在GPU上實際投入運行。這種異步調用機制致使Kernel沒法返回函數值。

 

10.Cuda 編程的三個基本法則：

    1）將數據放入並始終存儲於GPGPU

    2) 交給GPGPU足夠更多的任務

    3）注重GPGPU上的數據重用，以免帶寬限制

 

11.可用於CUDA計算的GPGPU包含片上與片外兩大類存儲器。流多處理器片上存儲器是速度最快，可擴展性最佳的備受青睞的存儲器。這些存儲器容量有限，一般僅在KB級別。板載的全局內存是一個共享的存儲系統，能夠被GPU上的全部流多處理器訪問。該內存容量能夠達到GB級別，是目前最大，使用最廣泛，但也是GPU上最慢的存儲器。

12.一個定義爲volatile的變量是說這個變量可能會被意想不到地改變，這樣，編譯器就不會去假設這個變量的值了。精確的說就是，優化器在用到這個變量的時候必須每次都當心的從新讀取這個變量的值，而不是使用保存在寄存器裏的備份。

13. 在GF100芯片上的流多處理器上，寄存器溢出後會將多餘的局部數據存放於L1緩存中，因爲L1緩存帶寬較高，程序仍可保持高性能，這就凸顯了L1緩存的重要性。但須要知道，因爲寄存器溢出和棧（最多會消耗1KB的L1緩存資源）佔用了L1緩存，這將致使其餘緩存數據更容易被擠出緩存，所以就會下降緩存的命中率，影響程序性能。

14.C編譯器容許main()沒有參數，或者有兩個參數（有些實現容許更多的參數，但這將是對標準的擴展）。有兩個參數時，第一個參數是命令行中的字符串數。按照慣例，這個int參數被稱爲argc（表明argument count），系統使用空格判斷一個字符串結束，另外一個字符串開始。第二個參數是一個指向字符串的指針數組。命令行中的每一個字符串被存儲到內存中，而且分配一個指針指向它。按照慣例，這個指針數組被稱爲argv(表明argument value)。若是能夠（有些操做系統不容許這樣作），把程序自己的名字賦值給argv[0],接着，把隨後的第一個字符串賦給argv[1]，等等。

  int main(int argc, char **argv) 這種對argv的聲明和char *argv[]是等價的。這意味着argv是一個指向「指向字符的指針」的指針。

15.對於結構體，和數組不一樣的是，一個結構的名字不是該結構的地址，必須使用&運算符。

16.關於編譯和連接的一點總結：在編譯時，編譯器只檢測程序語法，和函數，變量是否被聲明，若是函數未被聲明，編譯器會給出一個警告，可是能夠生成object file,而在連接程序時，連接器會在全部的Object文件中尋找函數的實現，若是找不到，就會報連接錯誤碼（Linker Error）