轉自：http://blog.sina.com.cn/s/blog_48b9e1f90100fm5f.html
結合lec07_intro_cuda.pptx學習
內存類型
CGMA: Compute to Global Memory Access ratio
Constant memory只允許device只讀，比global memory 能夠提供更快更多的並行數據訪問路徑給kernel。
Register和local memory是線程私有的。Shared memory是同一個block中的線程共享的。
Table 1顯示了cuda聲明變量的語法。Scope表示變量能夠被訪問的線程范圍。包括thread：線程單獨訪問，每個thread都有一個變量，如果kernel聲明一個scope為thread的變量y，在啟動x個線程后，就會有x個版本的變量y。block：被block中的所有thread訪問，grid：被grid中的所有線程訪問。
Lifetime是變量的生存期。注意：如果生存期為kernel，那么在kernel不同的啟動之間，變量的值是不會被保存下來的。每次啟動一次kernel都要對變量進行初始化。生存期為application的變量，必須在所有函數體外進行聲明，變量可以在程序執行中保存下來並可以被所有kernel訪問。
非數組自動變量：除了在kernel和device函數中聲明的數組外，其他所有自動變量都在寄存器中。這些變量稱為scalar變量，scope是單獨的線程。當一個kernel聲明了一個自動變量，系統會為執行這個kernel函數的所有線程copy這個變量。線程終止后，所有變量也就不存在了。
自動數組變量：存在global memory中，對它們的訪問需要長延遲。他們的scope也是單獨的線程。因此，對這種變量盡量避免使用。
（__device__）__shared__修飾的變量，表示CUDA中的共享變量。共享變量的scope是block，block中的所有線程都可以看到共享變量的同一個版本。Lifetime是kernel，kernel結束，共享變量內存也就不存在了。對共享內存的訪問非常快而且是高度並行的。CUDA編程者通常用共享內存來保留一部分在kernel中用的多的全局內存數據。
(__device__) __constant__修飾的變量表示常數變量constant variable。Constant variable必須在函數體外進行聲明。Scope是grids, lifetime是整個應用程序的執行。Constant variable常用於為kernel function提供輸入值，存儲在global memory中但被cached。一個程序constant variable最大可以使65536個字節。
__device__修飾的變量是global variable，存儲在global memory中。對global memory的訪問非常慢。由於global variable對所有kernel中的所有線程都是可見的。因此，global variable可以作為跨block的線程之間的協同方法。但是，如果不終止目前的kernel，無法保證線程之間數據的一致性。因此，global variable通常作為kernel function之間的信息傳遞。
指針只能用於指向global memory的數據對象，不用於device memory。指針有兩種典型用法：第一，如果一個對象由host function分配，指向此對象的指針被cudaMalloc()初始化並能夠作為參數傳遞給kernel function。第二，在global memory中聲明的變量的地址可以分配給一個指針變量。例如，
float * ptr=&GlobalVar。
減少全局內存通信的策略
由於全局內存大而慢，共享內存小而快。常用的策略是把數據划分成片tile，每一片適合共享內存的使用。對這些tile的kernel計算可以獨立的進行。
__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width)
{
1. __shared__float Mds[TILE_WIDTH][TILE_WIDTH];
2. __shared__float Nds[TILE_WIDTH][TILE_WIDTH];
3. int bx = blockIdx.x; int by = blockIdx.y;
4. int tx = threadIdx.x; int ty = threadIdx.y;
// Identify the row and column of the Pd element to work on
5. int Row = by * TILE_WIDTH + ty;
6. int Col = bx * TILE_WIDTH + tx;
7. float Pvalue = 0;
// Loop over the Md and Nd tiles required to compute the Pd element
8. for (int m = 0; m < Width/TILE_WIDTH; ++m) {
// Coolaborative loading of Md and Nd tiles into shared memory
9. Mds[ty][tx] = Md[Row][m*TILE_WIDTH + tx];
10. Nds[ty][tx] = Nd[m*TILE_WIDTH + ty][Col];
11. __Syncthreads();
12. for (int k = 0; k < TILE_WIDTH; ++k)
13. Pvalue += Mds[ty][k] * Nds[k][tx];
14. }
15. Pd[Row][Col] = Pvalue;
}
硬件限制：
GeForce 8800GTX每個SM有8K個寄存器，整個處理器有128K個寄存器。一個SM最多有768個線程。如果要達到這個線程最大數，每個線程只能用8K/768=10個寄存器。如果每個線程要用11個寄存器，那么線程數就會減少。例如，如果一個block有256個線程，那么每個SM中只有1/3的線程同時存在。
共享內存也會限制線程數目。在GeForce 8800 GTX中，每個SM有16K bytes大小的共享內存。而共享內存是block使用的。每個SM最多有8個block，所以，如果一個SM中有8個block，那么每個block最多能夠使用2K字節的共享內存。以矩陣乘為例，若tile大小為16*16，那么，每個block需要16*16*4=1K字節存儲Mds,需要1K字節存儲Nds。因此，一個block需要2K字節的共享內存。根據共享內存16K 字節的限制，最多有8個block可以同時存在於一個SM中，這也是硬件限制上的最大block數目了。若tile size是32*32，那么每個block需要8K 字節共享內存，那么一個SM只能有2個block。
注意：不時關注CUDA主頁，關注CUDA相關的最新研究