轉：cuda 初學大全

1 硬體架構 CUDA編程中，習慣稱CPU為Host，GPU為Device。

 

 

2 並行模型 Thread：並行基本單位 Block：相互合作的一組線程。可以彼此同步，快速交換資料，最多可以512個線程 Grid：一組Block，有共用全域記憶體 Kernel：在GPU上執行的程式，一個Kernel對應一個Grid

Block和Thread都有各自的ID，記作blockIdx（1D，2D），threadIdx（1D，2D，3D） Block和Thread還有Dim，即blockDim與threadDim. 他們都有三個分量x，y，z 線程同步：void __syncthreads(); 可以同步一個Block內的所有線程

 

 

3 儲存層次 per-thread register 1 cycle per-thread local memory slow per-block shared memory 1 cycle per-grid global memory 500 cycle,not cached!! constant and texture memories 500 cycle, but cached and read-only 分配記憶體：cudaMalloc，cudaFree，它們分配的是global memory Hose-Device資料交換：cudaMemcpy

 

4 變數類型 __device__：GPU的global memory空間，grid中所有線程可訪問 __constant__：GPU的constant memory空間，grid中所有線程可訪問 __shared__：GPU上的thread block空間，block中所有線程可訪問 local：位於SM內，僅本thread可訪問

在編程中，可以在變數名前面加上這些首碼以區分。

 

 

5 資料類型 內建向量：int1，int2，int3，int4，float1，float2， float3，float4 ... 紋理類型：texture<Type, Dim, ReadMode>texRef; 內建dim3類型：定義grid和block的組織方法。 例如： dim3 dimGrid(2, 2); dim3 dimBlock(4, 2, 2); kernelFoo<<<dimGrid, dimBlock>>>(argument);

 

 

6 CUDA函數定義

__device__：執行於Device，僅能從Device調用。 限制：不能用&取地址；不支援遞迴；不支援static variable；不支援可變長度參數 __global__ void： 執行於Device，僅能從Host調用。此類函數必須返回void __host__：執行於Host，僅能從Host調用
在執行kernel函數時，必須提供execution configuration，即<<<....>>>的部分。

例如： __global__ void KernelFunc(...); dim3 DimGrid(100, 50); // 5000 thread blocks dim3 DimBlock(4, 8, 8); // 256 threads per block size_t SharedMemBytes = 64; // 64 bytes of shared memory KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>(...);

 

 

7 CUDA包含一些數學函數，如sin，pow等。每一個函數包含有兩個版本，例如正弦函數sin，一個普通版本sin，另一個不精確但速度極快的__sin版本。

 

 

8 內建變數 gridDim, blockIdx, blockDim, threadIdx, wrapsize. 這些內建變數不允許賦值的

 

 

9 編寫程式

目前CUDA僅能良好的支援C，在編寫含有CUDA代碼的程式時，首先要匯入標頭檔cuda_runtime_api.h。檔案名稱尾碼為.cu，使用nvcc編譯器編譯。本來想在這裡給出些源碼的，但是源碼教程，以後單獨開一個文章在說吧。

 

 

 

 

這部分是一些枯燥的硬體知識的總結，但是對最佳化CUDA程式有著至關重要的作用，在後面的文章裡，我將盡量結合執行個體來講解這些東西

 

1 GPU硬體

i GPU一個最小單元稱為Streaming Processor(SP)，全流水線單事件無序微處理器，包含兩個ALU和一個FPU，多組寄存器檔案（register file，很多寄存器的組合），這個SP沒有cache。事實上，現代GPU就是一組SP的array，即SPA。每一個SP執行一個thread

 

ii 多個SP組成Streaming Multiprocessor(SM)。每一個SM執行一個block。每個SM包含 8個SP； 2個special function unit(SFU)：這裡面有4個FPU可以進行超越函數和插值計算 MultiThreading Issue Unit：分發線程指令 具有指令和常量緩衝。 包含shared memory

 

 

iii Texture Processor Cluster(TPC) ：包含某些其他單元的一組SM

 

 

2 Single-Program Multiple-Data （SPMD）模型

 

i CPU以順序結構執行代碼，GPU以threads blocks組織並發執行的代碼，即無數個threads同時執行

ii 回顧一下CUDA的概念： 一個kernel程式執行在一個grid of threads blocks之中 一個threads block是一批相互合作的threads：可以用過__syncthreads同步；通過shared memory共用變數，不同block的不能同步。

iii Threads block聲明： 可以包含有1到512個並發線程，具有唯一的blockID，可以是1,2,3D 同一個block中的線程執行同一個程式，不同的運算元，可以同步，每個線程具有唯一的ID

 

3 線程硬體原理

i GPU通過Global block scheduler來調度block，根據硬體架構分配block到某一個SM。每個SM最多分配8個block，每個SM最多可接受768個thread（可以是一個block包含512個thread，也可以是3個block每個包含256個thread（3*256=768！））。同一個SM上面的block的尺寸必須相同。每個線程的調度與ID由該SM管理。

ii SM滿負載工作效率最高！考慮某個Block，其尺寸可以為8*8,16*16,32*32 8*8：每個block有64個線程，由於每個SM最多處理768個線程，因此需要768/64=12個block。但是由於SM最多8個block，因此一個SM實際執行的線程為8*64=512個線程。 16*16：每個block有256個線程，SM可以同時接受三個block，3*256=768，滿負載：） 32*32：每個block有1024個線程，SM無法處理！

iii Block是獨立執行的，每個Block內的threads是可協同的。

iv 每個線程由SM中的一個SP執行。當然，由於SM中僅有8個SP，768個線程是以warp為單位執行的，每個warp包含32個線程，這是基於線程指令的流水線特性完成的。Warp是SM基本調度單位，實際上，一個Warp是一個32路SIMD指令。基本單位是half-warp。 如，SM滿負載工作有768個線程，則共有768/32=24個warp，每一瞬時，只有一組warp在SM中執行。 Warp全部線程是執行同一個指令，每個指令需要4個clock cycle，通過複雜的機制執行。

v 一個thread的一生： Grid在GPU上啟動；block被分配到SM上；SM把線程組織為warp；SM調度執行warp；執行結束後釋放資源；block繼續被分配....

4 線程儲存模型

i Register and local memory：線程私人，對程式員透明。 每個SM中有8192個register，分配給某些block，block內部的thread只能使用分配的寄存器。線程數多，每個線程使用的寄存器就少了。

ii shared memory：block內共用，動態分配。如__shared__ float region[N]。 shared memory 儲存空間是被劃分為16個小單元，與half-warp長度相同，稱為bank，每個bank可以提供自己的地址服務。連續的32位word映射到連續的bank。 對同一bank的同時訪問稱為bank conflict。盡量減少這種情形。

 

iii Global memory：沒有緩衝！容易稱為效能瓶頸，是最佳化的關鍵！ 一個half-warp裡面的16個線程對global memory的訪問可以被coalesce成整塊記憶體的訪問，如果： 資料長度為4,8或16bytes；地址連續；起始地址對齊；第N個線程訪問第N個資料。 Coalesce可以大大提升效能

 

uncoalesced

Coalesced方法：如果所有線程讀取同一地址，不妨使用constant memory；如果為不規則讀取可以使用texture記憶體 如果使用了某種結構體，其大小不是4 8 16的倍數，可以通過__align(X)強制對齊，X=4 8 16

原文：http://www.cnblogs.com/yangs/archive/2012/07/28/2613269.html