特此聲明：這個(gè)內(nèi)容我是轉(zhuǎn)別人的

我只摘錄一些我需要的東西，若是想看原文的，請(qǐng)點(diǎn)擊下面的鏈接

原文：http://hi.baidu.com/coolrainbow/item/de05efc83151671a50505878

?

?

?

1 硬件架構(gòu)
CUDA編程中，習(xí)慣稱CPU為Host，GPU為Device。

?

?

2 并行模型
Thread：并行基本單位
Block：相互合作的一組線程?？梢员舜送?#xff0c;快速交換數(shù)據(jù)，最多可以512個(gè)線程
Grid：一組Block，有共享全局內(nèi)存
Kernel：在GPU上執(zhí)行的程序，一個(gè)Kernel對(duì)應(yīng)一個(gè)Grid

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

?

?

3 存儲(chǔ)層次
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
分配內(nèi)存：cudaMalloc，cudaFree，它們分配的是global memory
Hose-Device數(shù)據(jù)交換：cudaMemcpy

?

4 變量類型
__device__：GPU的global memory空間，grid中所有線程可訪問
__constant__：GPU的constant memory空間，grid中所有線程可訪問
__shared__：GPU上的thread block空間，block中所有線程可訪問
local：位于SM內(nèi)，僅本thread可訪問

在編程中，可以在變量名前面加上這些前綴以區(qū)分。

?

?

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

?

?

6 CUDA函數(shù)定義

__device__：執(zhí)行于Device，僅能從Device調(diào)用。
限制：不能用&取地址；不支持遞歸；不支持static variable；不支持可變長(zhǎng)度參數(shù)
__global__ void： 執(zhí)行于Device，僅能從Host調(diào)用。此類函數(shù)必須返回void
__host__：執(zhí)行于Host，僅能從Host調(diào)用

在執(zhí)行kernel函數(shù)時(shí)，必須提供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包含一些數(shù)學(xué)函數(shù)，如sin，pow等。每一個(gè)函數(shù)包含有兩個(gè)版本，例如正弦函數(shù)sin，一個(gè)普通版本sin，另一個(gè)不精確但速度極快的__sin版本。

?

?

8 內(nèi)置變量
gridDim, blockIdx, blockDim, threadIdx, wrapsize. 這些內(nèi)置變量不允許賦值的

?

?

9 編寫程序

目前CUDA僅能良好的支持C，在編寫含有CUDA代碼的程序時(shí)，首先要導(dǎo)入頭文件cuda_runtime_api.h。文件名后綴為.cu，使用nvcc編譯器編譯。本來(lái)想在這里給出些源碼的，但是源碼教程，以后單獨(dú)開一個(gè)文章在說吧。

?

?

?

?

這部分是一些枯燥的硬件知識(shí)的總結(jié)，但是對(duì)優(yōu)化CUDA程序有著至關(guān)重要的作用，在后面的文章里，我將盡量結(jié)合實(shí)例來(lái)講解這些東西

?

1 GPU硬件

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

?

ii 多個(gè)SP組成Streaming Multiprocessor(SM)。每一個(gè)SM執(zhí)行一個(gè)block。每個(gè)SM包含
8個(gè)SP；
2個(gè)special function unit(SFU)：這里面有4個(gè)FPU可以進(jìn)行超越函數(shù)和插值計(jì)算
MultiThreading Issue Unit：分發(fā)線程指令
具有指令和常量緩存。
包含shared memory

?

?

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

?

?

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

?

i CPU以順序結(jié)構(gòu)執(zhí)行代碼，GPU以threads blocks組織并發(fā)執(zhí)行的代碼，即無(wú)數(shù)個(gè)threads同時(shí)執(zhí)行

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

iii Threads block聲明：
可以包含有1到512個(gè)并發(fā)線程，具有唯一的blockID，可以是1,2,3D
同一個(gè)block中的線程執(zhí)行同一個(gè)程序，不同的操作數(shù)，可以同步，每個(gè)線程具有唯一的ID

?

3 線程硬件原理

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

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

iii Block是獨(dú)立執(zhí)行的，每個(gè)Block內(nèi)的threads是可協(xié)同的。

iv 每個(gè)線程由SM中的一個(gè)SP執(zhí)行。當(dāng)然，由于SM中僅有8個(gè)SP，768個(gè)線程是以warp為單位執(zhí)行的，每個(gè)warp包含32個(gè)線程，這是基于線程指令的流水線特性完成的。Warp是SM基本調(diào)度單位，實(shí)際上，一個(gè)Warp是一個(gè)32路SIMD指令。基本單位是half-warp。
如，SM滿負(fù)載工作有768個(gè)線程，則共有768/32=24個(gè)warp，每一瞬時(shí)，只有一組warp在SM中執(zhí)行。
Warp全部線程是執(zhí)行同一個(gè)指令，每個(gè)指令需要4個(gè)clock cycle，通過復(fù)雜的機(jī)制執(zhí)行。

v 一個(gè)thread的一生：
Grid在GPU上啟動(dòng)；block被分配到SM上；SM把線程組織為warp；SM調(diào)度執(zhí)行warp；執(zhí)行結(jié)束后釋放資源；block繼續(xù)被分配....

4 線程存儲(chǔ)模型

i Register and local memory：線程私有，對(duì)程序員透明。
每個(gè)SM中有8192個(gè)register，分配給某些block，block內(nèi)部的thread只能使用分配的寄存器。線程數(shù)多，每個(gè)線程使用的寄存器就少了。

ii shared memory：block內(nèi)共享，動(dòng)態(tài)分配。如__shared__ float region[N]。
shared memory 存儲(chǔ)器是被劃分為16個(gè)小單元，與half-warp長(zhǎng)度相同，稱為bank，每個(gè)bank可以提供自己的地址服務(wù)。連續(xù)的32位word映射到連續(xù)的bank。
對(duì)同一bank的同時(shí)訪問稱為bank conflict。盡量減少這種情形。

?

iii Global memory：沒有緩存！容易稱為性能瓶頸，是優(yōu)化的關(guān)鍵！
一個(gè)half-warp里面的16個(gè)線程對(duì)global memory的訪問可以被coalesce成整塊內(nèi)存的訪問，如果：
數(shù)據(jù)長(zhǎng)度為4,8或16bytes；地址連續(xù)；起始地址對(duì)齊；第N個(gè)線程訪問第N個(gè)數(shù)據(jù)。
Coalesce可以大大提升性能

?

uncoalesced

Coalesced方法：如果所有線程讀取同一地址，不妨使用constant memory；如果為不規(guī)則讀取可以使用texture內(nèi)存
如果使用了某種結(jié)構(gòu)體，其大小不是4 8 16的倍數(shù)，可以通過__align(X)強(qiáng)制對(duì)齊，X=4 8 16

與50位技術(shù)專家面對(duì)面20年技術(shù)見證，附贈(zèng)技術(shù)全景圖

總結(jié)

以上是生活随笔為你收集整理的cuda初步认识的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。