GPU高效能程式設計CUDA實戰
書中摘錄+稍微的補充
1、CPU及繫系統的記憶體稱為主機，將GPU及其記憶體稱為裝置，在GPU裝置上執行的函式通常稱為核函式(kernel)；
2、CPU並行執行緒結構
《1》、執行緒 –> 執行緒塊 –>（執行緒格）grid ；
《2》、關鍵的內建變數：
threadIdx：3維向量，標識執行緒。每個執行緒塊內的threadIdx是唯一的；
blockIdx：3維向量，標識執行緒塊，每個grid內的blockIdx是唯一的；
blockDim：3維向量，儲存執行緒塊每個維度執行緒的數量資訊；
gridDim: 3維向量，儲存執行緒格每個維度執行緒塊的數量資訊
實際執行核函式的執行緒的tid計算(一維情況)： tid = threadIdx.x + blockIdx.x * blockDim.x
3、核函式帶有修飾符: _ global _
__global__ void kernel(int a)
{
a;
}核函式的呼叫帶有修飾符<<<>>>,在主機端呼叫kernel如下：
kennel_name<<<Dg, Db, Ns, S>>>([kernel arguments])Dg: dim3型別，表示使用的grid的緯度和大小資訊，既是裝置在執行該函式時使用的並行執行緒塊的數量；
Db: dim3型別，儲存為每個執行緒塊的緯度及大小資訊，既是裝置在執行該函式時每個執行緒塊內的執行緒數量；
Ns: size_t型別，可選項，預設為0；
S: cudaStream_t型別，cuda使用的流。可選項，預設是0；
例子：
__global__ void add(int a, int b, int *c)
{
*c = a + b;
}
int main(void)
{
int *dev_c;
.....
add<<<1,1>>>(2,7,dev_c);
.....
return 0;
}4、CUDA架構上GPU標準記憶體（又稱linear memory）分配（還有一些特殊記憶體如常量記憶體、頁鎖定記憶體等不涉及）
《1》、標準記憶體分配與釋放：cudaMalloc()/cudaFree()
[還有其它的分配方式如： cudaMallocPitch() 、 cudaMalloc3D() ]
《2》、主機與裝置之間記憶體資料拷貝函式：cudaMemcpy()
cudaMemcpy函式需要傳進標誌位來識別資料的流動方向
cudaMemcpyHostToDevice — 資料從主機到裝置
cudaMemcpyDeviceToHost — 資料從裝置到主機
// Allocate vectors in device memory
float* d_A;
cudaMalloc(&d_A, size);
float* d_B;
cudaMalloc(&d_B, size);
float* d_C;
cudaMalloc(&d_C, size);
// Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);5、裝置指標的使用限制
裝置指標：指向GPU記憶體地址的指標；
限制：只能在GPU上操作其所指記憶體，不能再主機上操作裝置指標所指內容
，具體如下：
《1》、可以將cudaMalloc()分配的指標傳達給在裝置上執行的函式；
《2》、可以在裝置程式碼中使用cudeMalloc()分配的指標進行記憶體讀/寫操作；
《3》、可以將cudaMalloc()分配的指標傳遞給在主機上執行的函式；
《4》、不能再主機程式碼使用cudaMalloc()分配的指標進行記憶體讀/寫操作；
6、共享記憶體與同步
《1》、共享記憶體宣告關鍵字：_share_
共享記憶體能被執行緒塊內的所有線性訪問,編譯器將為每個執行緒塊生成共享記憶體的一個副本。
《2》、共享記憶體快取區位於物理GPU上而不是駐留在GPU之外，故在訪問共享記憶體時的延遲要遠遠低於訪問普通緩衝區的延遲；
《3》、共享記憶體使用時需要進行執行緒之間的同步: __syncthreads()函式
該函式將確保執行緒塊中的每個執行緒都執行完__syncthreads()前面的語句後，才會執行下一條語句。
const int threadsPerBlock = 256;
__global__ void test(float *a , float* b, float *c)
{
__share__ float cache[threadsPerBlock];
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int cacheIndex = threadIdx.x;
//共享記憶體快取中的偏移就等於線性索引，因為每個執行緒塊都擁有該共享記憶體的私有副本
float temp = 0;
while（tid < N）
{
temp += a[i] + b[i];
tid += blockDim.x * gridDim.x;
}
cache[cacheIndex] = temp;
//對執行緒塊中的執行緒進行同步
__syncthreads();
}執行緒發散(Thread Divergence): 當某些執行緒需要執行一條指令而其他執行緒不需要執行是，這種情況稱為執行緒發散。
* 如果__syncthreads()位於執行緒發散的分支，那GPU程式會一直保持等待，不會結束
*
7、兩類只讀記憶體
a、常量記憶體（constant memory spaces）– 只讀
《1》、 常量記憶體一般用於儲存在核函式執行期間不會發生變化的資料，NVIDIA硬體提供了64KB的常量記憶體，並且對常量記憶體採取了不同於標準全域性記憶體的處理方式。在某些情況中，用常量記憶體來替換全域性記憶體能有效的減少記憶體頻寬；
減少頻寬原因：
對常量記憶體的單次讀操作可以廣播到其它的“鄰近（Nearby）”執行緒，這將節約15次讀取操作；
常量記憶體的資料將快取起來，因此對相同的地址的連續讀操作將不會產生額外的記憶體通訊量。
“鄰近（Nearby）”執行緒：執行緒束概念【一個包含32個執行緒的集合–cuda中定義】
《2》、常量記憶體宣告修飾符關鍵字: _constant_
#define SPHERES 20
struct Sphere {
int a;
float b;
};
__constant__ Sphere s[SPHERES];
__global__ void test()
{
//do something
}
int main( void )
{
// allocate temp memory, initialize it, copy to constant
// memory on the GPU, then free our temp memory
Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES );
for (int i=0; i<SPHERES; i++) {
temp_s[i].r = rnd( 1.0f );
temp_s[i].g = rnd( 1.0f );
temp_s[i].b = rnd( 1.0f );
temp_s[i].x = rnd( 1000.0f ) - 500;
temp_s[i].y = rnd( 1000.0f ) - 500;
temp_s[i].z = rnd( 1000.0f ) - 500;
temp_s[i].radius = rnd( 100.0f ) + 20;
}
cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere) * SPHERES);
free( temp_s );
return 0;
}
《3》、常量記憶體的拷貝使用cudaMemcpyToSymbol()函式；
b、紋理記憶體( texture and surface memory spaces)【書上簡單介紹，詳細可看《cuda c programing guide》的3.11部分】
《1》、只讀記憶體，在特定的訪問模式中，紋理記憶體同樣能夠提升效能並減少記憶體流量。紋理記憶體時專門為那些在記憶體訪問模式中存在大量空間區域性性(Spatial Locality)的圖形應用程式而設計的。在某個計算應用程式中，這意味著這一個執行緒讀取的位置可能與鄰近執行緒讀取的位置“非常接近”。
紋理記憶體分3種：一維/二維/三維紋理記憶體；
《2》、宣告texture型別的引用API：texture
texture<float> texConstSrc;
texture<int> textIn;texture 引用只能被宣告為全域性靜態變數，不能將其作為引數傳遞給函式。
例子：
// these exist on the GPU side
texture<float> texConstSrc;
__global__ void copy_const_kernel( float *iptr ) {
// map from threadIdx/BlockIdx to pixel position
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;
float c = tex1Dfetch(texConstSrc,offset);//讀取一維紋理記憶體的方式
if (c != 0)
iptr[offset] = c;
}
int main( void )
{
int imageSize = 1024;
float* dev_inSrc = NULL;
cudaMalloc( (void**)&dev_inSrc, imageSize );
cudaBindTexture( NULL, texConstSrc,dev_inSrc, imageSize );//必須繫結
return 0;
｝至此CUDA 記憶體包括：全域性記憶體，共享記憶體，常量記憶體，紋理記憶體 共四類。
8、事件機制用於測量效能（耗時）
cuda中的事件本質上是一個GPU時間戳，這個時間戳實在使用者指定的事件點上的記錄。GPU本身支援記錄時間戳，避免了很多與CPU定時器一起統計的麻煩。
事件的建立及使用：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0); //第二引數指定使用的流，預設0
//kernel function
cudaEventRecord(stop,0);
cudaEventSynchronize(stop); //同步
float elapsedTime = 0.0f;
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Cost Time: %3.1f ms \n", elapsedTime); //毫秒級9、計算功能集(Compute Capability)：NVIDIA GPU支援的各種功能的統稱。
10、原子操作(像Linux 下的原子操作)
11、頁鎖定記憶體(Page-locked memory)
《1》、頁鎖定記憶體的一個重要屬性：作業系統不會把這塊記憶體分頁，也不會將其交換到磁碟上，從而確保了該記憶體始終駐留在實體記憶體中。該記憶體的實體地址是能訪問的，因為這塊記憶體將不會被破壞或重新定位。
《2》、GPU知道記憶體的實體地址，因此可以通過DMA技術來在GPU和主機頁鎖定記憶體之間複製資料。在GPU和主機之間進行DMA資料傳輸時，使用頁鎖定記憶體會比使用標準記憶體的效能快很多（大約2倍）。
使用頁鎖定記憶體的一個不足是：使用頁鎖定記憶體會更快的耗盡系統記憶體。所以使用頁鎖定記憶體應該有針對性不能隨意使用，一般來說僅對cudaMemcpy()函式調中的源記憶體或目標記憶體才使用頁鎖定記憶體
。
《3》、通過cudaHostAlloc函式分配， cudaFreeHost()釋放；
float cuda_host_alloc_test( int size, bool up ) {
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime;
HANDLE_ERROR( cudaEventCreate( &start ) );
HANDLE_ERROR( cudaEventCreate( &stop ) );
HANDLE_ERROR( cudaHostAlloc( (void**)&a,size * sizeof( *a ),cudaHostAllocDefault ) );
HANDLE_ERROR( cudaMalloc( (void**)&dev_a,size * sizeof( *dev_a ) ) );
HANDLE_ERROR( cudaEventRecord( start, 0 ) );
for (int i=0; i<100; i++) {
if (up)
HANDLE_ERROR( cudaMemcpy( dev_a, a,size * sizeof( *a ),cudaMemcpyHostToDevice ) );
else
HANDLE_ERROR( cudaMemcpy( a, dev_a,size * sizeof( *a ),cudaMemcpyDeviceToHost ) );
}
HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
HANDLE_ERROR( cudaEventSynchronize( stop ) );
HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,start, stop ) );
HANDLE_ERROR( cudaFreeHost( a ) );
HANDLE_ERROR( cudaFree( dev_a ) );
HANDLE_ERROR( cudaEventDestroy( start ) );
HANDLE_ERROR( cudaEventDestroy( stop ) );
return elapsedTime;
}
12、cuda流
《1》、cuda流表示一個GPU操作佇列，並且該佇列中的操作將以指定的順序執行。cuda流在加速應用程式方面起著重要作用，可以將每個流視為GPU上的一個任務，並且這些任務是可以並行執行的。將操作新增到流的順序也就是它們的執行順序。
《2》、裝置重疊功能（Device Overlap）:支援裝置重疊功能的GPU能夠在執行一個CUDA C核函式的同事，還能在裝置與主機之間執行復制操作。
《3》、從邏輯上看，不同流之間是相互獨立的，但事實上這種理解並不完全符合GPU的佇列機制。在硬體中並沒有流的概念而是包含一個或多個引擎來執行記憶體複製操作，以及一個引擎來執行核函式。
程式設計的時候需要按照硬體處理邏輯來進行，否則無法實現並行。