轉自：http://luofl1992.is-programmer.com/posts/38830.html

CUDA編程中，習慣稱CPU為Host，GPU為Device。編程中最開始接觸的東西恐怕是并行架構，諸如Grid、Block的區別會讓人一頭霧水，我所看的書上所講述的內容比較抽象，對這些概念的內容沒有細講，于是在這里作一個整理。

Grid、Block和Thread的關系

Thread? ：并行運算的基本單位（輕量級的線程）
Block?? ：由相互合作的一組線程組成。一個block中的thread可以彼此同步，快速交換數據，最多可以同時512個線程。
Grid???? ：一組Block，有共享全局內存
Kernel ：在GPU上執行的程序，一個Kernel對應一個Grid。

其結構如下圖所示：

1 /* 2 另外：Block和Thread都有各自的ID，記作blockIdx（1D，2D），threadIdx（1D，2D，3D） 3 Block和Thread還有Dim，即blockDim與threadDim. 他們都有三個分量x，y，z 4 線程同步：void __syncthreads(); 可以同步一個Block內的所有線程 5 總結來說，每個 thread 都有自己的一份 register 和 local memory 的空間。 6 一組thread構成一個 block，這些 thread 則共享有一份shared memory。 7 此外，所有的 thread(包括不同 block 的 thread)都共享一份 8 global memory、constant memory、和 texture memory。 9 不同的 grid 則有各自的 global memory、constant memory 和 texture memory。 10 */

?

存儲層次
1 per-thread register 1 cycle 2 per-thread local memory slow 3 per-block shared memory 1 cycle 4 per-grid global memory 500 cycle,not cached!! 5 constant and texture memories 500 cycle, but cached and read-only 6 分配內存：cudaMalloc，cudaFree，它們分配的是global memory 7 Hose-Device數據交換：cudaMemcpy 變量類型
1 __device__ // GPU的global memory空間，grid中所有線程可訪問 2 __constant__ // GPU的constant memory空間，grid中所有線程可訪問 3 __shared__ // GPU上的thread block空間，block中所有線程可訪問 4 local // 位于SM內，僅本thread可訪問 5 // 在編程中，可以在變量名前面加上這些前綴以區分。 數據類型
1 // 內建矢量類型： 2 int1，int2，int3，int4，float1，float2， float3，float4 ... 3 // 紋理類型： 4 texture<Type, Dim, ReadMode>texRef; 5 // 內建dim3類型：定義grid和block的組織方法。例如： 6 dim3 dimGrid(2, 2); 7 dim3 dimBlock(4, 2, 2); 8 // CUDA函數CPU端調用方法 9 kernelFoo<<<dimGrid, dimBlock>>>(argument); 函數定義
1 __device__ // 執行于Device，僅能從Device調用。限制，不能用&取地址；不支持遞歸；不支持static variable；不支持可變長度參數 2 __global__ // void： 執行于Device，僅能從Host調用。此類函數必須返回void 3 __host__ // 執行于Host，僅能從Host調用，是函數的默認類型 4 // 在執行kernel函數時，必須提供execution configuration，即<<<....>>>的部分。 5 // 例如： 6 __global__ void KernelFunc(...); 7 dim3 DimGrid(100, 50); // 5000 thread blocks 8 dim3 DimBlock(4, 8, 8); // 256 threads per block 9 size_t SharedMemBytes = 64; // 64 bytes of shared memory 10 KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>(...);

?

數學函數
1 CUDA包含一些數學函數，如sin，pow等。每一個函數包含有兩個版本， 2 例如正弦函數sin，一個普通版本sin，另一個不精確但速度極快的__sin版本。 內置變量
1 /* 2 gridDim, blockIdx, blockDim, 3 threadIdx, wrapsize. 4 這些內置變量不允許賦值的 5 */ 編寫程序
1 /* 2 目前CUDA僅能良好的支持C，在編寫含有CUDA代碼的程序時， 3 首先要導入頭文件cuda_runtime_api.h。文件名后綴為.cu，使用nvcc編譯器編譯。 4 目前最新的CUDA版本為5.0，可以在官方網站下載最新的工具包，網址為： 5 https://developer.nvidia.com/cuda-downloads 6 該工具包內包含了ToolKit、樣例等，安裝起來比原先的版本也方便了很多。 7 */ 相關擴展
1 1 GPU硬件 2 // i GPU一個最小單元稱為Streaming Processor(SP)，全流水線單事件無序微處理器， 3 包含兩個ALU和一個FPU，多組寄存器文件（register file，很多寄存器的組合）， 4 這個SP沒有cache。事實上，現代GPU就是一組SP的array，即SPA。 5 每一個SP執行一個thread 6 7 // ii 多個SP組成Streaming Multiprocessor(SM)。 8 每一個SM執行一個block。每個SM包含8個SP； 9 2個special function unit(SFU)： 10 這里面有4個FPU可以進行超越函數和插值計算 11 MultiThreading Issue Unit：分發線程指令 12 具有指令和常量緩存。 13 包含shared memory 14 15 // iii Texture Processor Cluster(TPC) ：包含某些其他單元的一組SM 16 17 2 Single-Program Multiple-Data （SPMD）模型 18 19 // i CPU以順序結構執行代碼， 20 GPU以threads blocks組織并發執行的代碼，即無數個threads同時執行 21 22 // ii 回顧一下CUDA的概念： 23 一個kernel程序執行在一個grid of threads blocks之中 24 一個threads block是一批相互合作的threads： 25 可以用過__syncthreads同步； 26 通過shared memory共享變量，不同block的不能同步。 27 28 // iii Threads block聲明： 29 可以包含有1到512個并發線程，具有唯一的blockID，可以是1,2,3D 30 同一個block中的線程執行同一個程序，不同的操作數，可以同步，每個線程具有唯一的ID 31 32 3 線程硬件原理 33 34 // i GPU通過Global block scheduler來調度block， 35 根據硬件架構分配block到某一個SM。 36 每個SM最多分配8個block，每個SM最多可接受768個thread 37 （可以是一個block包含512個thread， 38 也可以是3個block每個包含256個thread（3*256=768！））。 39 同一個SM上面的block的尺寸必須相同。每個線程的調度與ID由該SM管理。 40 41 // ii SM滿負載工作效率最高！考慮某個Block，其尺寸可以為8*8,16*16,32*32 42 8*8：每個block有64個線程， 43 由于每個SM最多處理768個線程，因此需要768/64=12個block。 44 但是由于SM最多8個block，因此一個SM實際執行的線程為8*64=512個線程。 45 16*16：每個block有256個線程，SM可以同時接受三個block，3*256=768，滿負載 46 32*32：每個block有1024個線程，SM無法處理！ 47 48 // iii Block是獨立執行的，每個Block內的threads是可協同的。 49 50 // iv 每個線程由SM中的一個SP執行。 51 當然，由于SM中僅有8個SP，768個線程是以warp為單位執行的， 52 每個warp包含32個線程，這是基于線程指令的流水線特性完成的。 53 Warp是SM基本調度單位，實際上，一個Warp是一個32路SIMD指令 54 。基本單位是half-warp。 55 如，SM滿負載工作有768個線程，則共有768/32=24個warp 56 ，每一瞬時，只有一組warp在SM中執行。 57 Warp全部線程是執行同一個指令， 58 每個指令需要4個clock cycle，通過復雜的機制執行。 59 60 // v 一個thread的一生： 61 Grid在GPU上啟動； 62 block被分配到SM上； 63 SM把線程組織為warp； 64 SM調度執行warp； 65 執行結束后釋放資源； 66 block繼續被分配.... 67 68 4 線程存儲模型 69 70 // i Register and local memory：線程私有，對程序員透明。 71 每個SM中有8192個register，分配給某些block， 72 block內部的thread只能使用分配的寄存器。 73 線程數多，每個線程使用的寄存器就少了。 74 75 // ii shared memory：block內共享，動態分配。 76 如__shared__ float region[N]。 77 shared memory 存儲器是被劃分為16個小單元， 78 與half-warp長度相同，稱為bank，每個bank可以提供自己的地址服務。 79 連續的32位word映射到連續的bank。 80 對同一bank的同時訪問稱為bank conflict。 81 盡量減少這種情形。 82 83 // iii Global memory：沒有緩存！容易稱為性能瓶頸，是優化的關鍵！ 84 一個half-warp里面的16個線程對global memory的訪問可以被coalesce成整塊內存的訪問，如果： 85 數據長度為4,8或16bytes；地址連續；起始地址對齊；第N個線程訪問第N個數據。 86 Coalesce可以大大提升性能。 87 88 // uncoalesced 89 Coalesced方法：如果所有線程讀取同一地址， 90 不妨使用constant memory； 91 如果為不規則讀取可以使用texture內存 92 如果使用了某種結構體，其大小不是4 8 16的倍數， 93 可以通過__align(X)強制對齊，X=4 8 16

?

?example：//想象為二維數組

int row = blockIdx.y * blockDim.y + threadIdx.y; //thread 所在的行

int col = blockIdx.x * blockDim.x + threadIdx.x;?//thread 所在的列

blockDim是指每個block的size，blockDim.y 相當于block的height，blockDim.x相當于block的width。

blockIdx.y是指block在grid中豎排的位置，同理，blockIdx.x是指block在grid中橫排的位置。

?

?

?

轉載于:https://www.cnblogs.com/qingsunny/p/3384779.html

總結

以上是生活随笔為你收集整理的GPU CUDA 经典入门指南的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。