在老板的要求下，本博主從2012年上高性能計算課程開始接觸CUDA編程笤闯，隨后將該技術(shù)應(yīng)用到了實際項目中，使處理程序加速超過1K溉贿，可見基于圖形顯示器的并行計算對于追求速度的應(yīng)用來說無疑是一個理想的選擇畦浓。還有不到一年畢業(yè)，怕是畢業(yè)后這些技術(shù)也就隨畢業(yè)而去张症，準(zhǔn)備這個暑假開辟一個CUDA專欄，從入門到精通鸵贬，步步為營俗他，順便分享設(shè)計的一些經(jīng)驗教訓(xùn)，希望能給學(xué)習(xí)CUDA的童鞋提供一定指導(dǎo)阔逼。個人能力所及兆衅，錯誤難免，歡迎討論嗜浮。
PS：申請專欄好像需要先發(fā)原創(chuàng)帖超過15篇羡亩。。危融。算了畏铆，先寫夠再申請吧，到時候一并轉(zhuǎn)過去吉殃。
NVIDIA于2006年推出CUDA（Compute Unified Devices Architecture）辞居，可以利用其推出的GPU進行通用計算楷怒，將并行計算從大型集群擴展到了普通顯卡，使得用戶只需要一臺帶有Geforce顯卡的筆記本就能跑較大規(guī)模的并行處理程序瓦灶。
使用顯卡的好處是鸠删，和大型集群相比功耗非常低，成本也不高贼陶，但性能很突出刃泡。以我的筆記本為例，Geforce 610M碉怔，用DeviceQuery程序測試烘贴，可得到如下硬件參數(shù)：
計算能力達48X0.95 = 45.6 GFLOPS。而筆記本的CPU參數(shù)如下：
CPU計算能力為（4核）：2.5G*4 = 10GFLOPS眨层，可見庙楚，顯卡計算性能是4核i5 CPU的4~5倍，因此我們可以充分利用這一資源來對一些耗時的應(yīng)用進行加速趴樱。
好了馒闷，工欲善其事必先利其器，為了使用CUDA對GPU進行編程叁征，我們需要準(zhǔn)備以下必備工具：
1. 硬件平臺纳账，就是顯卡，如果你用的不是NVIDIA的顯卡捺疼，那么只能說抱歉疏虫，其他都不支持CUDA。
2. 操作系統(tǒng)啤呼，我用過windows XP卧秘，Windows 7都沒問題，本博客用Windows7官扣。
3. C編譯器翅敌，建議VS2008，和本博客一致惕蹄。
4. CUDA編譯器NVCC蚯涮，可以免費免注冊免license從官網(wǎng)下載CUDA ToolkitCUDA下載，最新版本為5.0卖陵，本博客用的就是該版本遭顶。
5. 其他工具（如Visual Assist，輔助代碼高亮）
準(zhǔn)備完畢泪蔫，開始安裝軟件棒旗。VS2008安裝比較費時間，建議安裝完整版（NVIDIA官網(wǎng)說Express版也可以）鸥滨，過程不必詳述嗦哆。CUDA Toolkit 5.0里面包含了NVCC編譯器谤祖、設(shè)計文檔、設(shè)計例程老速、CUDA運行時庫粥喜、CUDA頭文件等必備的原材料。
安裝完畢橘券，我們在桌面上發(fā)現(xiàn)這個圖標(biāo)：
不錯额湘，就是它，雙擊運行旁舰，可以看到一大堆例程锋华。我們找到Simple OpenGL這個運行看看效果：
點右邊黃線標(biāo)記處的Run即可看到美妙的三維正弦曲面，鼠標(biāo)左鍵拖動可以轉(zhuǎn)換角度箭窜，右鍵拖動可以縮放毯焕。如果這個運行成功，說明你的環(huán)境基本搭建成功磺樱。
出現(xiàn)問題的可能：
1. 你使用遠程桌面連接登錄到另一臺服務(wù)器纳猫，該服務(wù)器上有顯卡支持CUDA，但你遠程終端不能運行CUDA程序竹捉。這是因為遠程登錄使用的是你本地顯卡資源芜辕，在遠程登錄時看不到服務(wù)器端的顯卡，所以會報錯：沒有支持CUDA的顯卡块差！解決方法：1. 遠程服務(wù)器裝兩塊顯卡益老，一塊只用于顯示弓乙，另一塊用于計算郊酒；2.不要用圖形界面登錄确封，而是用命令行界面如telnet登錄。
2.有兩個以上顯卡都支持CUDA的情況鹉动，如何區(qū)分是在哪個顯卡上運行诗舰？這個需要你在程序里控制，選擇符合一定條件的顯卡训裆，如較高的時鐘頻率、較大的顯存蜀铲、較高的計算版本等边琉。詳細操作見后面的博客。
好了记劝，先說這么多变姨，下一節(jié)我們介紹如何在VS2008中給GPU編程。
書接上回厌丑，我們既然直接運行例程成功了定欧，接下來就是了解如何實現(xiàn)例程中的每個環(huán)節(jié)渔呵。當(dāng)然，我們先從簡單的做起砍鸠，一般編程語言都會找個helloworld例子扩氢，而我們的顯卡是不會說話的，只能做一些簡單的加減乘除運算爷辱。所以录豺，CUDA程序的helloworld，我想應(yīng)該最合適不過的就是向量加了饭弓。
打開VS2008双饥，選擇File->New->Project，彈出下面對話框弟断，設(shè)置如下：
之后點OK咏花，直接進入工程界面。
工程中阀趴，我們看到只有一個.cu文件昏翰，內(nèi)容如下：
[cpp]view plaincopy
#include?"cuda_runtime.h"
#include?"device_launch_parameters.h"
#include?
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize);
__global__voidaddKernel(int*c,constint*a,constint*b)
{
inti?=?threadIdx.x;
c[i]?=?a[i]?+?b[i];
}
intmain()
{
constintarraySize?=?5;
constinta[arraySize]?=?{?1,?2,?3,?4,?5?};
constintb[arraySize]?=?{?10,?20,?30,?40,?50?};
intc[arraySize]?=?{?0?};
//?Add?vectors?in?parallel.
cudaError_t?cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"addWithCuda?failed!");
return1;
}
printf("{1,2,3,4,5}?+?{10,20,30,40,50}?=?{%d,%d,%d,%d,%d}\n",
c[0],?c[1],?c[2],?c[3],?c[4]);
//?cudaThreadExit?must?be?called?before?exiting?in?order?for?profiling?and
//?tracing?tools?such?as?Nsight?and?Visual?Profiler?to?show?complete?traces.
cudaStatus?=?cudaThreadExit();
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaThreadExit?failed!");
return1;
}
return0;
}
//?Helper?function?for?using?CUDA?to?add?vectors?in?parallel.
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize)
{
int*dev_a?=?0;
int*dev_b?=?0;
int*dev_c?=?0;
cudaError_t?cudaStatus;
//?Choose?which?GPU?to?run?on,?change?this?on?a?multi-GPU?system.
cudaStatus?=?cudaSetDevice(0);
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?");
gotoError;
}
//?Allocate?GPU?buffers?for?three?vectors?(two?input,?one?output)????.
cudaStatus?=?cudaMalloc((void**)&dev_c,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_a,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_b,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
//?Copy?input?vectors?from?host?memory?to?GPU?buffers.
cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
//?Launch?a?kernel?on?the?GPU?with?one?thread?for?each?element.
addKernel<<<1,?size>>>(dev_c,?dev_a,?dev_b);
//?cudaThreadSynchronize?waits?for?the?kernel?to?finish,?and?returns
//?any?errors?encountered?during?the?launch.
cudaStatus?=?cudaThreadSynchronize();
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n",?cudaStatus);
gotoError;
}
//?Copy?output?vector?from?GPU?buffer?to?host?memory.
cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*sizeof(int),?cudaMemcpyDeviceToHost);
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
returncudaStatus;
}
可以看出，CUDA程序和C程序并無區(qū)別舍咖，只是多了一些以"cuda"開頭的一些庫函數(shù)和一個特殊聲明的函數(shù)：
[cpp]view plaincopy
__global__voidaddKernel(int*c,constint*a,constint*b)
{
inti?=?threadIdx.x;
c[i]?=?a[i]?+?b[i];
}
這個函數(shù)就是在GPU上運行的函數(shù)矩父，稱之為核函數(shù)，英文名Kernel Function排霉，注意要和操作系統(tǒng)內(nèi)核函數(shù)區(qū)分開來窍株。
我們直接按F7編譯，可以得到如下輸出：
[html]view plaincopy
1>------?Build?started:?Project:?cuda_helloworld,?Configuration:?Debug?Win32?------
1>Compiling?with?CUDA?Build?Rule...
1>"C:\Program?Files\NVIDIA?GPU?Computing?Toolkit\CUDA\v5.0\\bin\nvcc.exe"??-G-gencode=arch=compute_10,code=\"sm_10,compute_10\"-gencode=arch=compute_20,code=\"sm_20,compute_20\"??--machine?32?-ccbin?"C:\Program?Files?(x86)\Microsoft?Visual?Studio?9.0\VC\bin"????-Xcompiler?"/EHsc?/W3?/nologo?/O2?/Zi???/MT??"??-I"C:\Program?Files\NVIDIA?GPU?Computing?Toolkit\CUDA\v5.0\\include"-maxrregcount=0--compile?-o?"Debug/kernel.cu.obj"?kernel.cu
1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.gpu
1>tmpxft_000000ec_00000000-14_kernel.compute_10.cudafe2.gpu
1>tmpxft_000000ec_00000000-5_kernel.compute_20.cudafe1.gpu
1>tmpxft_000000ec_00000000-17_kernel.compute_20.cudafe2.gpu
1>kernel.cu
1>kernel.cu
1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.cpp
1>tmpxft_000000ec_00000000-24_kernel.compute_10.ii
1>Linking...
1>Embedding?manifest...
1>Performing?Post-Build?Event...
1>copy?"C:\Program?Files\NVIDIA?GPU?Computing?Toolkit\CUDA\v5.0\\bin\cudart*.dll"?"C:\Users\DongXiaoman\Documents\Visual?Studio?2008\Projects\cuda_helloworld\Debug"
1>C:\Program?Files\NVIDIA?GPU?Computing?Toolkit\CUDA\v5.0\\bin\cudart32_50_35.dll
1>C:\Program?Files\NVIDIA?GPU?Computing?Toolkit\CUDA\v5.0\\bin\cudart64_50_35.dll
1>已復(fù)制?????????2?個文件攻柠。
1>Build?log?was?saved?at?"file://c:\Users\DongXiaoman\Documents\Visual?Studio?2008\Projects\cuda_helloworld\cuda_helloworld\Debug\BuildLog.htm"
1>cuda_helloworld?-?0?error(s),?105?warning(s)
==========?Build:?1?succeeded,?0?failed,?0?up-to-date,?0skipped==========
可見球订，編譯.cu文件需要利用nvcc工具。該工具的詳細使用見后面博客瑰钮。
直接運行冒滩，可以得到結(jié)果圖如下：
如果顯示正確，那么我們的第一個程序宣告成功浪谴！
剛?cè)腴TCUDA开睡，跑過幾個官方提供的例程，看了看人家的代碼苟耻，覺得并不難篇恒，但自己動手寫代碼時，總是不知道要先干什么凶杖，后干什么胁艰，也不知道從哪個知識點學(xué)起。這時就需要有一本能提供指導(dǎo)的書籍或者教程，一步步跟著做下去腾么，直到真正掌握奈梳。
一般講述CUDA的書，我認為不錯的有下面這幾本：
初學(xué)者可以先看美國人寫的這本《GPU高性能編程CUDA實戰(zhàn)》解虱，可操作性很強攘须，但不要期望能全看懂（Ps：里面有些概念其實我現(xiàn)在還是不怎么懂），但不影響你進一步學(xué)習(xí)饭寺。如果想更全面地學(xué)習(xí)CUDA阻课，《GPGPU編程技術(shù)》比較客觀詳細地介紹了通用GPU編程的策略，看過這本書艰匙，可以對顯卡有更深入的了解限煞，揭開GPU的神秘面紗。后面《OpenGL編程指南》完全是為了體驗圖形交互帶來的樂趣员凝，可以有選擇地看署驻；《GPU高性能運算之CUDA》這本是師兄給的，適合快速查詢（感覺是將官方編程手冊翻譯了一遍）一些關(guān)鍵技術(shù)和概念健霹。
有了這些指導(dǎo)材料還不夠旺上，我們在做項目的時候，遇到的問題在這些書上肯定找不到糖埋，所以還需要有下面這些利器：
這里面有很多工具的使用手冊宣吱，如CUDA_GDB，Nsight瞳别，CUDA_Profiler等征候，方便調(diào)試程序；還有一些有用的庫祟敛，如CUFFT是專門用來做快速傅里葉變換的疤坝，CUBLAS是專用于線性代數(shù)（矩陣、向量計算）的馆铁，CUSPASE是專用于稀疏矩陣表示和計算的庫跑揉。這些庫的使用可以降低我們設(shè)計算法的難度，提高開發(fā)效率埠巨。另外還有些入門教程也是值得一讀的历谍，你會對NVCC編譯器有更近距離的接觸。
好了辣垒，前言就這么多扮饶，本博主計劃按如下順序來講述CUDA：
1.了解設(shè)備
2.線程并行
3.塊并行
4.流并行
5.線程通信
6.線程通信實例：規(guī)約
7.存儲模型
8.常數(shù)內(nèi)存
9.紋理內(nèi)存
10.主機頁鎖定內(nèi)存
11.圖形互操作
12.優(yōu)化準(zhǔn)則
13.CUDA與MATLAB接口
14.CUDA與MFC接口
前面三節(jié)已經(jīng)對CUDA做了一個簡單的介紹，這一節(jié)開始真正進入編程環(huán)節(jié)乍构。
首先，初學(xué)者應(yīng)該對自己使用的設(shè)備有較為扎實的理解和掌握，這樣對后面學(xué)習(xí)并行程序優(yōu)化很有幫助哥遮，了解硬件詳細參數(shù)可以通過上節(jié)介紹的幾本書和官方資料獲得岂丘，但如果仍然覺得不夠直觀，那么我們可以自己動手獲得這些內(nèi)容眠饮。
以第二節(jié)例程為模板奥帘，我們稍加改動的部分代碼如下：
[cpp]view plaincopy
//?Add?vectors?in?parallel.
cudaError_t?cudaStatus;
intnum?=?0;
cudaDeviceProp?prop;
cudaStatus?=?cudaGetDeviceCount(&num);
for(inti?=?0;i
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);
這個改動的目的是讓我們的程序自動通過調(diào)用cuda API函數(shù)獲得設(shè)備數(shù)目和屬性，所謂“知己知彼仪召，百戰(zhàn)不殆”寨蹋。
cudaError_t 是cuda錯誤類型，取值為整數(shù)扔茅。
cudaDeviceProp為設(shè)備屬性結(jié)構(gòu)體已旧，其定義可以從cuda Toolkit安裝目錄中找到，我的路徑為：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\driver_types.h召娜，找到定義為：
[cpp]view plaincopy
/**
*?CUDA?device?properties
*/
struct__device_builtin__?cudaDeviceProp
{
charname[256];/**<?ASCII?string?identifying?device?*/
size_ttotalGlobalMem;/**<?Global?memory?available?on?device?in?bytes?*/
size_tsharedMemPerBlock;/**<?Shared?memory?available?per?block?in?bytes?*/
intregsPerBlock;/**<?32-bit?registers?available?per?block?*/
intwarpSize;/**<?Warp?size?in?threads?*/
size_tmemPitch;/**<?Maximum?pitch?in?bytes?allowed?by?memory?copies?*/
intmaxThreadsPerBlock;/**<?Maximum?number?of?threads?per?block?*/
intmaxThreadsDim[3];/**<?Maximum?size?of?each?dimension?of?a?block?*/
intmaxGridSize[3];/**<?Maximum?size?of?each?dimension?of?a?grid?*/
intclockRate;/**<?Clock?frequency?in?kilohertz?*/
size_ttotalConstMem;/**<?Constant?memory?available?on?device?in?bytes?*/
intmajor;/**<?Major?compute?capability?*/
intminor;/**<?Minor?compute?capability?*/
size_ttextureAlignment;/**<?Alignment?requirement?for?textures?*/
size_ttexturePitchAlignment;/**<?Pitch?alignment?requirement?for?texture?references?bound?to?pitched?memory?*/
intdeviceOverlap;/**<?Device?can?concurrently?copy?memory?and?execute?a?kernel.?Deprecated.?Use?instead?asyncEngineCount.?*/
intmultiProcessorCount;/**<?Number?of?multiprocessors?on?device?*/
intkernelExecTimeoutEnabled;/**<?Specified?whether?there?is?a?run?time?limit?on?kernels?*/
intintegrated;/**<?Device?is?integrated?as?opposed?to?discrete?*/
intcanMapHostMemory;/**<?Device?can?map?host?memory?with?cudaHostAlloc/cudaHostGetDevicePointer?*/
intcomputeMode;/**<?Compute?mode?(See?::cudaComputeMode)?*/
intmaxTexture1D;/**<?Maximum?1D?texture?size?*/
intmaxTexture1DMipmap;/**<?Maximum?1D?mipmapped?texture?size?*/
intmaxTexture1DLinear;/**<?Maximum?size?for?1D?textures?bound?to?linear?memory?*/
intmaxTexture2D[2];/**<?Maximum?2D?texture?dimensions?*/
intmaxTexture2DMipmap[2];/**<?Maximum?2D?mipmapped?texture?dimensions?*/
intmaxTexture2DLinear[3];/**<?Maximum?dimensions?(width,?height,?pitch)?for?2D?textures?bound?to?pitched?memory?*/
intmaxTexture2DGather[2];/**<?Maximum?2D?texture?dimensions?if?texture?gather?operations?have?to?be?performed?*/
intmaxTexture3D[3];/**<?Maximum?3D?texture?dimensions?*/
intmaxTextureCubemap;/**<?Maximum?Cubemap?texture?dimensions?*/
intmaxTexture1DLayered[2];/**<?Maximum?1D?layered?texture?dimensions?*/
intmaxTexture2DLayered[3];/**<?Maximum?2D?layered?texture?dimensions?*/
intmaxTextureCubemapLayered[2];/**<?Maximum?Cubemap?layered?texture?dimensions?*/
intmaxSurface1D;/**<?Maximum?1D?surface?size?*/
intmaxSurface2D[2];/**<?Maximum?2D?surface?dimensions?*/
intmaxSurface3D[3];/**<?Maximum?3D?surface?dimensions?*/
intmaxSurface1DLayered[2];/**<?Maximum?1D?layered?surface?dimensions?*/
intmaxSurface2DLayered[3];/**<?Maximum?2D?layered?surface?dimensions?*/
intmaxSurfaceCubemap;/**<?Maximum?Cubemap?surface?dimensions?*/
intmaxSurfaceCubemapLayered[2];/**<?Maximum?Cubemap?layered?surface?dimensions?*/
size_tsurfaceAlignment;/**<?Alignment?requirements?for?surfaces?*/
intconcurrentKernels;/**<?Device?can?possibly?execute?multiple?kernels?concurrently?*/
intECCEnabled;/**<?Device?has?ECC?support?enabled?*/
intpciBusID;/**<?PCI?bus?ID?of?the?device?*/
intpciDeviceID;/**<?PCI?device?ID?of?the?device?*/
intpciDomainID;/**<?PCI?domain?ID?of?the?device?*/
inttccDriver;/**<?1?if?device?is?a?Tesla?device?using?TCC?driver,?0?otherwise?*/
intasyncEngineCount;/**<?Number?of?asynchronous?engines?*/
intunifiedAddressing;/**<?Device?shares?a?unified?address?space?with?the?host?*/
intmemoryClockRate;/**<?Peak?memory?clock?frequency?in?kilohertz?*/
intmemoryBusWidth;/**<?Global?memory?bus?width?in?bits?*/
intl2CacheSize;/**<?Size?of?L2?cache?in?bytes?*/
intmaxThreadsPerMultiProcessor;/**<?Maximum?resident?threads?per?multiprocessor?*/
};
后面的注釋已經(jīng)說明了其字段代表意義运褪，可能有些術(shù)語對于初學(xué)者理解起來還是有一定困難，沒關(guān)系玖瘸，我們現(xiàn)在只需要關(guān)注以下幾個指標(biāo)：
name：就是設(shè)備名稱秸讹；
totalGlobalMem：就是顯存大小雅倒；
major,minor：CUDA設(shè)備版本號璃诀，有1.1, 1.2, 1.3, 2.0, 2.1等多個版本；
clockRate：GPU時鐘頻率蔑匣；
multiProcessorCount：GPU大核數(shù)劣欢，一個大核（專業(yè)點稱為流多處理器，SM殖演，Stream-Multiprocessor）包含多個小核（流處理器氧秘，SP，Stream-Processor）
編譯趴久，運行丸相，我們在VS2008工程的cudaGetDeviceProperties()函數(shù)處放一個斷點，單步執(zhí)行這一函數(shù)彼棍，然后用Watch窗口灭忠，切換到Auto頁，展開+座硕，在我的筆記本上得到如下結(jié)果：
可以看到弛作，設(shè)備名為GeForce 610M，顯存1GB华匾，設(shè)備版本2.1（比較高端了映琳，哈哈），時鐘頻率為950MHz（注意950000單位為kHz），大核數(shù)為1萨西。在一些高性能GPU上（如Tesla有鹿，Kepler系列），大核數(shù)可能達到幾十甚至上百谎脯，可以做更大規(guī)模的并行處理葱跋。
PS：今天看SDK代碼時發(fā)現(xiàn)在help_cuda.h中有個函數(shù)實現(xiàn)從CUDA設(shè)備版本查詢相應(yīng)大核中小核的數(shù)目，覺得很有用源梭，以后編程序可以借鑒娱俺，摘抄如下：
[cpp]view plaincopy
//?Beginning?of?GPU?Architecture?definitions
inlineint_ConvertSMVer2Cores(intmajor,intminor)
{
//?Defines?for?GPU?Architecture?types?(using?the?SM?version?to?determine?the?#?of?cores?per?SM
typedefstruct
{
intSM;//?0xMm?(hexidecimal?notation),?M?=?SM?Major?version,?and?m?=?SM?minor?version
intCores;
}?sSMtoCores;
sSMtoCores?nGpuArchCoresPerSM[]?=
{
{?0x10,??8?},//?Tesla?Generation?(SM?1.0)?G80?class
{?0x11,??8?},//?Tesla?Generation?(SM?1.1)?G8x?class
{?0x12,??8?},//?Tesla?Generation?(SM?1.2)?G9x?class
{?0x13,??8?},//?Tesla?Generation?(SM?1.3)?GT200?class
{?0x20,?32?},//?Fermi?Generation?(SM?2.0)?GF100?class
{?0x21,?48?},//?Fermi?Generation?(SM?2.1)?GF10x?class
{?0x30,?192},//?Kepler?Generation?(SM?3.0)?GK10x?class
{?0x35,?192},//?Kepler?Generation?(SM?3.5)?GK11x?class
{???-1,?-1?}
};
intindex?=?0;
while(nGpuArchCoresPerSM[index].SM?!=?-1)
{
if(nGpuArchCoresPerSM[index].SM?==?((major?<<?4)?+?minor))
{
returnnGpuArchCoresPerSM[index].Cores;
}
index++;
}
//?If?we?don't?find?the?values,?we?default?use?the?previous?one?to?run?properly
printf("MapSMtoCores?for?SM?%d.%d?is?undefined.??Default?to?use?%d?Cores/SM\n",?major,?minor,?nGpuArchCoresPerSM[7].Cores);
returnnGpuArchCoresPerSM[7].Cores;
}
//?end?of?GPU?Architecture?definitions
可見，設(shè)備版本2.1的一個大核有48個小核废麻，而版本3.0以上的一個大核有192個小核荠卷！
前文說到過，當(dāng)我們用的電腦上有多個顯卡支持CUDA時脑溢，怎么來區(qū)分在哪個上運行呢僵朗？這里我們看一下addWithCuda這個函數(shù)是怎么做的。
[cpp]view plaincopy
cudaError_t?cudaStatus;
//?Choose?which?GPU?to?run?on,?change?this?on?a?multi-GPU?system.
cudaStatus?=?cudaSetDevice(0);
if(cudaStatus?!=?cudaSuccess)?{
fprintf(stderr,"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?");
gotoError;
}
使用了cudaSetDevice(0)這個操作屑彻，0表示能搜索到的第一個設(shè)備號验庙，如果有多個設(shè)備，則編號為0,1,2...社牲。
再看我們本節(jié)添加的代碼粪薛，有個函數(shù)cudaGetDeviceCount(&num)，這個函數(shù)用來獲取設(shè)備總數(shù)搏恤，這樣我們選擇運行CUDA程序的設(shè)備號取值就是0,1,...num-1违寿，于是可以一個個枚舉設(shè)備，利用cudaGetDeviceProperties(&prop)獲得其屬性,然后利用一定排序熟空、篩選算法藤巢，找到最符合我們應(yīng)用的那個設(shè)備號opt，然后調(diào)用cudaSetDevice(opt)即可選擇該設(shè)備息罗。選擇標(biāo)準(zhǔn)可以從處理能力掂咒、版本控制、名稱等各個角度出發(fā)迈喉。后面講述流并發(fā)過程時绍刮，還要用到這些API。
如果希望了解更多硬件內(nèi)容可以結(jié)合http://www.geforce.cn/hardware獲取挨摸。
多線程我們應(yīng)該都不陌生孩革，在操作系統(tǒng)中，進程是資源分配的基本單元得运，而線程是CPU時間調(diào)度的基本單元（這里假設(shè)只有1個CPU）膝蜈。
將線程的概念引申到CUDA程序設(shè)計中锅移，我們可以認為線程就是執(zhí)行CUDA程序的最小單元，前面我們建立的工程代碼中彬檀，有個核函數(shù)概念不知各位童鞋還記得沒有帆啃，在GPU上每個線程都會運行一次該核函數(shù)。
但GPU上的線程調(diào)度方式與CPU有很大不同窍帝。CPU上會有優(yōu)先級分配，從高到低诽偷，同樣優(yōu)先級的可以采用時間片輪轉(zhuǎn)法實現(xiàn)線程調(diào)度坤学。GPU上線程沒有優(yōu)先級概念，所有線程機會均等报慕，線程狀態(tài)只有等待資源和執(zhí)行兩種狀態(tài)深浮，如果資源未就緒，那么就等待眠冈；一旦就緒飞苇，立即執(zhí)行。當(dāng)GPU資源很充裕時蜗顽，所有線程都是并發(fā)執(zhí)行的布卡，這樣加速效果很接近理論加速比；而GPU資源少于總線程個數(shù)時雇盖，有一部分線程就會等待前面執(zhí)行的線程釋放資源忿等，從而變?yōu)榇谢瘓?zhí)行。
代碼還是用上一節(jié)的吧崔挖，改動很少贸街，再貼一遍：
[cpp]view plaincopy
#include?"cuda_runtime.h"???????????//CUDA運行時API
#include?"device_launch_parameters.h"
#include?
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize);
__global__voidaddKernel(int*c,constint*a,constint*b)
{
inti?=?threadIdx.x;
c[i]?=?a[i]?+?b[i];
}
intmain()
{
constintarraySize?=?5;
constinta[arraySize]?=?{?1,?2,?3,?4,?5?};
constintb[arraySize]?=?{?10,?20,?30,?40,?50?};
intc[arraySize]?=?{?0?};
//?Add?vectors?in?parallel.
cudaError_t?cudaStatus;
intnum?=?0;
cudaDeviceProp?prop;
cudaStatus?=?cudaGetDeviceCount(&num);
for(inti?=?0;i
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"addWithCuda?failed!");
return1;
}
printf("{1,2,3,4,5}?+?{10,20,30,40,50}?=?{%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
//?cudaThreadExit?must?be?called?before?exiting?in?order?for?profiling?and
//?tracing?tools?such?as?Nsight?and?Visual?Profiler?to?show?complete?traces.
cudaStatus?=?cudaThreadExit();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadExit?failed!");
return1;
}
return0;
}
//?重點理解這個函數(shù)
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize)
{
int*dev_a?=?0;//GPU設(shè)備端數(shù)據(jù)指針
int*dev_b?=?0;
int*dev_c?=?0;
cudaError_t?cudaStatus;//狀態(tài)指示
//?Choose?which?GPU?to?run?on,?change?this?on?a?multi-GPU?system.
cudaStatus?=?cudaSetDevice(0);//選擇運行平臺
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?");
gotoError;
}
//?分配GPU設(shè)備端內(nèi)存
cudaStatus?=?cudaMalloc((void**)&dev_c,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_a,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_b,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
//?拷貝數(shù)據(jù)到GPU
cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
//?運行核函數(shù)
????addKernel<<<1,?size>>>(dev_c,?dev_a,?dev_b);
//?cudaThreadSynchronize?waits?for?the?kernel?to?finish,?and?returns
//?any?errors?encountered?during?the?launch.
cudaStatus?=?cudaThreadSynchronize();//同步線程
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n",?cudaStatus);
gotoError;
}
//?Copy?output?vector?from?GPU?buffer?to?host?memory.
cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*sizeof(int),?cudaMemcpyDeviceToHost);//拷貝結(jié)果回主機
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
Error:
cudaFree(dev_c);//釋放GPU設(shè)備端內(nèi)存
cudaFree(dev_a);
cudaFree(dev_b);
returncudaStatus;
}
紅色部分即啟動核函數(shù)的調(diào)用過程，這里看到調(diào)用方式和C不太一樣狸相。<<<>>>表示運行時配置符號薛匪，里面1表示只分配一個線程組（又稱線程塊、Block）脓鹃，size表示每個線程組有size個線程（Thread）逸尖。本程序中size根據(jù)前面?zhèn)鬟f參數(shù)個數(shù)應(yīng)該為5，所以運行的時候将谊，核函數(shù)在5個GPU線程單元上分別運行了一次冷溶，總共運行了5次。這5個線程是如何知道自己“身份”的尊浓？是靠threadIdx這個內(nèi)置變量逞频，它是個dim3類型變量，接受<<<>>>中第二個參數(shù)栋齿，它包含x,y,z 3維坐標(biāo)苗胀，而我們傳入的參數(shù)只有一維襟诸，所以只有x值是有效的。通過核函數(shù)中int i = threadIdx.x;這一句基协，每個線程可以獲得自身的id號歌亲，從而找到自己的任務(wù)去執(zhí)行。
同一版本的代碼用了這么多次澜驮，有點過意不去陷揪，于是這次我要做較大的改動
，大家要擦亮眼睛杂穷，拭目以待悍缠。
塊并行相當(dāng)于操作系統(tǒng)中多進程的情況，上節(jié)說到耐量，CUDA有線程組（線程塊）的概念飞蚓，將一組線程組織到一起，共同分配一部分資源廊蜒，然后內(nèi)部調(diào)度執(zhí)行趴拧。線程塊與線程塊之間，毫無瓜葛山叮。這有利于做更粗粒度的并行著榴。我們將上一節(jié)的代碼改為塊并行版本如下：
下節(jié)我們介紹塊并行。
[cpp]view plaincopy
#include?"cuda_runtime.h"
#include?"device_launch_parameters.h"
#include?
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize);
__global__voidaddKernel(int*c,constint*a,constint*b)
{
inti?=?blockIdx.x;
????c[i]?=?a[i]?+?b[i];
}
intmain()
{
constintarraySize?=?5;
constinta[arraySize]?=?{?1,?2,?3,?4,?5?};
constintb[arraySize]?=?{?10,?20,?30,?40,?50?};
intc[arraySize]?=?{?0?};
//?Add?vectors?in?parallel.
cudaError_t?cudaStatus;
intnum?=?0;
cudaDeviceProp?prop;
cudaStatus?=?cudaGetDeviceCount(&num);
for(inti?=?0;i
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"addWithCuda?failed!");
return1;
}
printf("{1,2,3,4,5}?+?{10,20,30,40,50}?=?{%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
//?cudaThreadExit?must?be?called?before?exiting?in?order?for?profiling?and
//?tracing?tools?such?as?Nsight?and?Visual?Profiler?to?show?complete?traces.
cudaStatus?=?cudaThreadExit();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadExit?failed!");
return1;
}
return0;
}
//?Helper?function?for?using?CUDA?to?add?vectors?in?parallel.
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize)
{
int*dev_a?=?0;
int*dev_b?=?0;
int*dev_c?=?0;
cudaError_t?cudaStatus;
//?Choose?which?GPU?to?run?on,?change?this?on?a?multi-GPU?system.
cudaStatus?=?cudaSetDevice(0);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?");
gotoError;
}
//?Allocate?GPU?buffers?for?three?vectors?(two?input,?one?output)????.
cudaStatus?=?cudaMalloc((void**)&dev_c,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_a,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_b,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
//?Copy?input?vectors?from?host?memory?to?GPU?buffers.
cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
//?Launch?a?kernel?on?the?GPU?with?one?thread?for?each?element.
???addKernel<<>>(dev_c,?dev_a,?dev_b);
//?cudaThreadSynchronize?waits?for?the?kernel?to?finish,?and?returns
//?any?errors?encountered?during?the?launch.
cudaStatus?=?cudaThreadSynchronize();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n",?cudaStatus);
gotoError;
}
//?Copy?output?vector?from?GPU?buffer?to?host?memory.
cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*sizeof(int),?cudaMemcpyDeviceToHost);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
returncudaStatus;
}
和上一節(jié)相比聘芜，只有這兩行有改變兄渺，<<<>>>里第一個參數(shù)改成了size，第二個改成了1汰现，表示我們分配size個線程塊挂谍，每個線程塊僅包含1個線程，總共還是有5個線程瞎饲。這5個線程相互獨立口叙，執(zhí)行核函數(shù)得到相應(yīng)的結(jié)果，與上一節(jié)不同的是嗅战，每個線程獲取id的方式變?yōu)閕nt i = blockIdx.x妄田；這是線程塊ID。
于是有童鞋提問了驮捍，線程并行和塊并行的區(qū)別在哪里疟呐？
線程并行是細粒度并行，調(diào)度效率高东且；塊并行是粗粒度并行启具，每次調(diào)度都要重新分配資源，有時資源只有一份珊泳，那么所有線程塊都只能排成一隊鲁冯，串行執(zhí)行。
那是不是我們所有時候都應(yīng)該用線程并行薯演，盡可能不用塊并行撞芍？
當(dāng)然不是，我們的任務(wù)有時可以采用分治法跨扮，將一個大問題分解為幾個小規(guī)模問題序无，將這些小規(guī)模問題分別用一個線程塊實現(xiàn)，線程塊內(nèi)可以采用細粒度的線程并行衡创，而塊之間為粗粒度并行愉镰，這樣可以充分利用硬件資源，降低線程并行的計算復(fù)雜度钧汹。適當(dāng)分解，降低規(guī)模录择，在一些矩陣乘法审丘、向量內(nèi)積計算應(yīng)用中可以得到充分的展示员帮。
實際應(yīng)用中，常常是二者的結(jié)合。線程塊贡珊、線程組織圖如下所示。
多個線程塊組織成了一個Grid恃鞋，稱為線程格（經(jīng)歷了從一位線程洞慎，二維線程塊到三維線程格的過程，立體感很強傲饨浴）须误。
好了，下一節(jié)我們介紹流并行仇轻，是更高層次的并行京痢。
前面我們沒有講程序的結(jié)構(gòu)，我想有些童鞋可能迫不及待想知道CUDA程序到底是怎么一個執(zhí)行過程篷店。好的祭椰，這一節(jié)在介紹流之前，先把CUDA程序結(jié)構(gòu)簡要說一下疲陕。
CUDA程序文件后綴為.cu方淤，有些編譯器可能不認識這個后綴的文件，我們可以在VS2008的Tools->Options->Text Editor->File Extension里添加cu后綴到VC++中蹄殃，如下圖：
一個.cu文件內(nèi)既包含CPU程序（稱為主機程序）携茂，也包含GPU程序（稱為設(shè)備程序）。如何區(qū)分主機程序和設(shè)備程序窃爷？根據(jù)聲明邑蒋，凡是掛有“__global__”或者“__device__”前綴的函數(shù)姓蜂，都是在GPU上運行的設(shè)備程序，不同的是__global__設(shè)備程序可被主機程序調(diào)用医吊，而__device__設(shè)備程序則只能被設(shè)備程序調(diào)用钱慢。
沒有掛任何前綴的函數(shù)，都是主機程序卿堂。主機程序顯示聲明可以用__host__前綴束莫。設(shè)備程序需要由NVCC進行編譯，而主機程序只需要由主機編譯器（如VS2008中的cl.exe草描，Linux上的GCC）览绿。主機程序主要完成設(shè)備環(huán)境初始化，數(shù)據(jù)傳輸?shù)缺貍溥^程穗慕，設(shè)備程序只負責(zé)計算饿敲。
主機程序中，有一些“cuda”打頭的函數(shù)逛绵，這些都是CUDA Runtime API怀各，即運行時函數(shù)，主要負責(zé)完成設(shè)備的初始化术浪、內(nèi)存分配瓢对、內(nèi)存拷貝等任務(wù)。我們前面第三節(jié)用到的函數(shù)cudaGetDeviceCount()胰苏，cudaGetDeviceProperties()硕蛹，cudaSetDevice()都是運行時API。這些函數(shù)的具體參數(shù)聲明我們不必一一記下來硕并，拿出第三節(jié)的官方利器就可以輕松查詢法焰，讓我們打開這個文件：
打開后，在pdf搜索欄中輸入一個運行時函數(shù)鲤孵，例如cudaMemcpy壶栋，查到的結(jié)果如下：
可以看到，該API函數(shù)的參數(shù)形式為普监，第一個表示目的地贵试，第二個表示來源地，第三個參數(shù)表示字節(jié)數(shù)凯正，第四個表示類型毙玻。如果對類型不了解，直接點擊超鏈接廊散，得到詳細解釋如下：
可見桑滩，該API可以實現(xiàn)從主機到主機、主機到設(shè)備允睹、設(shè)備到主機运准、設(shè)備到設(shè)備的內(nèi)存拷貝過程幌氮。同時可以發(fā)現(xiàn)，利用該API手冊可以很方便地查詢我們需要用的這些API函數(shù)胁澳，所以以后編CUDA程序一定要把它打開该互，隨時準(zhǔn)備查詢，這樣可以大大提高編程效率韭畸。
好了宇智，進入今天的主題：流并行。
前面已經(jīng)介紹了線程并行和塊并行胰丁，知道了線程并行為細粒度的并行随橘，而塊并行為粗粒度的并行，同時也知道了CUDA的線程組織情況锦庸，即Grid-Block-Thread結(jié)構(gòu)机蔗。一組線程并行處理可以組織為一個block，而一組block并行處理可以組織為一個Grid甘萧，很自然地想到蜒车，Grid只是一個網(wǎng)格，我們是否可以利用多個網(wǎng)格來完成并行處理呢幔嗦？答案就是利用流。
流可以實現(xiàn)在一個設(shè)備上運行多個核函數(shù)沥潭。前面的塊并行也好邀泉，線程并行也好，運行的核函數(shù)都是相同的（代碼一樣钝鸽，傳遞參數(shù)也一樣）汇恤。而流并行，可以執(zhí)行不同的核函數(shù)拔恰，也可以實現(xiàn)對同一個核函數(shù)傳遞不同的參數(shù)因谎，實現(xiàn)任務(wù)級別的并行。
CUDA中的流用cudaStream_t類型實現(xiàn)颜懊，用到的API有以下幾個：cudaStreamCreate(cudaStream_t * s)用于創(chuàng)建流财岔，cudaStreamDestroy(cudaStream_t s)用于銷毀流，cudaStreamSynchronize()用于單個流同步河爹，cudaDeviceSynchronize()用于整個設(shè)備上的所有流同步匠璧，cudaStreamQuery()用于查詢一個流的任務(wù)是否已經(jīng)完成。具體的含義可以查詢API手冊咸这。
下面我們將前面的兩個例子中的任務(wù)改用流實現(xiàn)夷恍，仍然是{1,2,3,4,5}+{10,20,30,40,50} = {11,22,33,44,55}這個例子。代碼如下：
[cpp]view plaincopy
#include?"cuda_runtime.h"
#include?"device_launch_parameters.h"
#include?
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize);
__global__voidaddKernel(int*c,constint*a,constint*b)
{
inti?=?blockIdx.x;
c[i]?=?a[i]?+?b[i];
}
intmain()
{
constintarraySize?=?5;
constinta[arraySize]?=?{?1,?2,?3,?4,?5?};
constintb[arraySize]?=?{?10,?20,?30,?40,?50?};
intc[arraySize]?=?{?0?};
//?Add?vectors?in?parallel.
cudaError_t?cudaStatus;
intnum?=?0;
cudaDeviceProp?prop;
cudaStatus?=?cudaGetDeviceCount(&num);
for(inti?=?0;i
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"addWithCuda?failed!");
return1;
}
printf("{1,2,3,4,5}?+?{10,20,30,40,50}?=?{%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
//?cudaThreadExit?must?be?called?before?exiting?in?order?for?profiling?and
//?tracing?tools?such?as?Nsight?and?Visual?Profiler?to?show?complete?traces.
cudaStatus?=?cudaThreadExit();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadExit?failed!");
return1;
}
return0;
}
//?Helper?function?for?using?CUDA?to?add?vectors?in?parallel.
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize)
{
int*dev_a?=?0;
int*dev_b?=?0;
int*dev_c?=?0;
cudaError_t?cudaStatus;
//?Choose?which?GPU?to?run?on,?change?this?on?a?multi-GPU?system.
cudaStatus?=?cudaSetDevice(0);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?");
gotoError;
}
//?Allocate?GPU?buffers?for?three?vectors?(two?input,?one?output)????.
cudaStatus?=?cudaMalloc((void**)&dev_c,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_a,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_b,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
//?Copy?input?vectors?from?host?memory?to?GPU?buffers.
cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
??cudaStream_t?stream[5];
for(inti?=?0;i<5;i++)
{
cudaStreamCreate(&stream[i]);//創(chuàng)建流
}
//?Launch?a?kernel?on?the?GPU?with?one?thread?for?each?element.
for(inti?=?0;i<5;i++)
{
addKernel<<<1,1,0,stream[i]>>>(dev_c+i,?dev_a+i,?dev_b+i);//執(zhí)行流
}
cudaDeviceSynchronize();
//?cudaThreadSynchronize?waits?for?the?kernel?to?finish,?and?returns
//?any?errors?encountered?during?the?launch.
cudaStatus?=?cudaThreadSynchronize();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n",?cudaStatus);
gotoError;
}
//?Copy?output?vector?from?GPU?buffer?to?host?memory.
cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*sizeof(int),?cudaMemcpyDeviceToHost);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
Error:
for(inti?=?0;i<5;i++)
{
cudaStreamDestroy(stream[i]);//銷毀流
}
????cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
returncudaStatus;
}
注意到媳维，我們的核函數(shù)代碼仍然和塊并行的版本一樣酿雪，只是在調(diào)用時做了改變遏暴，<<<>>>中的參數(shù)多了兩個，其中前兩個和塊并行指黎、線程并行中的意義相同朋凉，仍然是線程塊數(shù)（這里為1）、每個線程塊中線程數(shù)（這里也是1）袋励。第三個為0表示每個block用到的共享內(nèi)存大小侥啤，這個我們后面再講；第四個為流對象茬故，表示當(dāng)前核函數(shù)在哪個流上運行盖灸。我們創(chuàng)建了5個流，每個流上都裝載了一個核函數(shù)磺芭，同時傳遞參數(shù)有些不同赁炎，也就是每個核函數(shù)作用的對象也不同。這樣就實現(xiàn)了任務(wù)級別的并行钾腺，當(dāng)我們有幾個互不相關(guān)的任務(wù)時徙垫，可以寫多個核函數(shù)，資源允許的情況下放棒，我們將這些核函數(shù)裝載到不同流上姻报，然后執(zhí)行，這樣可以實現(xiàn)更粗粒度的并行间螟。
好了吴旋，流并行就這么簡單，我們處理任務(wù)時厢破，可以根據(jù)需要荣瑟，選擇最適合的并行方式。
我們前面幾節(jié)主要介紹了三種利用GPU實現(xiàn)并行處理的方式：線程并行摩泪，塊并行和流并行笆焰。在這些方法中，我們一再強調(diào)见坑，各個線程所進行的處理是互不相關(guān)的嚷掠，即兩個線程不回產(chǎn)生交集，每個線程都只關(guān)注自己的一畝三分地荞驴，對其他線程毫無興趣叠国，就當(dāng)不存在。戴尸。粟焊。。
當(dāng)然，實際應(yīng)用中项棠，這樣的例子太少了悲雳，也就是遇到向量相加、向量對應(yīng)點乘這類才會有如此高的并行度香追，而其他一些應(yīng)用合瓢，如一組數(shù)求和，求最大（型傅洹）值晴楔，各個線程不再是相互獨立的，而是產(chǎn)生一定關(guān)聯(lián)峭咒，線程2可能會用到線程1的結(jié)果税弃，這時就需要利用本節(jié)的線程通信技術(shù)了。
線程通信在CUDA中有三種實現(xiàn)方式：
1. 共享存儲器凑队；
2. 線程?同步则果；
3. 原子操作；
最常用的是前兩種方式漩氨，共享存儲器西壮，術(shù)語Shared Memory，是位于SM中的特殊存儲器叫惊。還記得SM嗎款青，就是流多處理器，大核是也霍狰。一個SM中不僅包含若干個SP（流處理器可都，小核），還包括一部分高速Cache蚓耽，寄存器組，共享內(nèi)存等旋炒，結(jié)構(gòu)如圖所示：
從圖中可看出步悠，一個SM內(nèi)有M個SP，Shared Memory由這M個SP共同占有瘫镇。另外指令單元也被這M個SP共享鼎兽，即SIMT架構(gòu)（單指令多線程架構(gòu)），一個SM中所有SP在同一時間執(zhí)行同一代碼铣除。
為了實現(xiàn)線程通信谚咬，僅僅靠共享內(nèi)存還不夠，需要有同步機制才能使線程之間實現(xiàn)有序處理尚粘。通常情況是這樣：當(dāng)線程A需要線程B計算的結(jié)果作為輸入時择卦，需要確保線程B已經(jīng)將結(jié)果寫入共享內(nèi)存中，然后線程A再從共享內(nèi)存中讀出。同步必不可少秉继，否則祈噪，線程A可能讀到的是無效的結(jié)果，造成計算錯誤尚辑。同步機制可以用CUDA內(nèi)置函數(shù)：__syncthreads()辑鲤；當(dāng)某個線程執(zhí)行到該函數(shù)時，進入等待狀態(tài)杠茬，直到同一線程塊（Block）中所有線程都執(zhí)行到這個函數(shù)為止月褥，即一個__syncthreads()相當(dāng)于一個線程同步點，確保一個Block中所有線程都達到同步瓢喉，然后線程進入運行狀態(tài)宁赤。
綜上兩點，我們可以寫一段線程通信的偽代碼如下：
[cpp]view plaincopy
//Begin
ifthisisthreadB
write?something?to?Shared?Memory;
endif
__syncthreads();
ifthisisthreadA
read?something?from?Shared?Memory;
endif
//End
上面代碼在CUDA中實現(xiàn)時灯荧，由于SIMT特性礁击，所有線程都執(zhí)行同樣的代碼，所以在線程中需要判斷自己的身份逗载，以免誤操作哆窿。
注意的是，位于同一個Block中的線程才能實現(xiàn)通信厉斟，不同Block中的線程不能通過共享內(nèi)存挚躯、同步進行通信，而應(yīng)采用原子操作或主機介入擦秽。
對于原子操作码荔，如果感興趣可以翻閱《GPU高性能編程CUDA實戰(zhàn)》第九章“原子性”。
本節(jié)完感挥。下節(jié)我們給出一個實例來看線程通信的代碼怎么設(shè)計缩搅。
接著上一節(jié)，我們利用剛學(xué)到的共享內(nèi)存和線程同步技術(shù)触幼，來做一個簡單的例子硼瓣。先看下效果吧：
很簡單，就是分別求出1~5這5個數(shù)字的和置谦，平方和堂鲤，連乘積。相信學(xué)過C語言的童鞋都能用for循環(huán)做出同上面一樣的效果媒峡，但為了學(xué)習(xí)CUDA共享內(nèi)存和同步技術(shù)瘟栖，我們還是要把簡單的東西復(fù)雜化(^_^)。
簡要分析一下谅阿，上面例子的輸入都是一樣的半哟，1,2,3,4,5這5個數(shù)酬滤，但計算過程有些變化，而且每個輸出和所有輸入都相關(guān)镜沽，不是前幾節(jié)例子中那樣敏晤，一個輸出只和一個輸入有關(guān)。所以我們在利用CUDA編程時缅茉，需要針對特殊問題做些讓步嘴脾，把一些步驟串行化實現(xiàn)。
輸入數(shù)據(jù)原本位于主機內(nèi)存蔬墩，通過cudaMemcpy API已經(jīng)拷貝到GPU顯存（術(shù)語為全局存儲器译打，Global Memory），每個線程運行時需要從Global Memory讀取輸入數(shù)據(jù)拇颅，然后完成計算奏司，最后將結(jié)果寫回Global Memory。當(dāng)我們計算需要多次相同輸入數(shù)據(jù)時樟插，大家可能想到韵洋，每次都分別去Global Memory讀數(shù)據(jù)好像有點浪費，如果數(shù)據(jù)很大黄锤，那么反復(fù)多次讀數(shù)據(jù)會相當(dāng)耗時間搪缨。索性我們把它從Global Memory一次性讀到SM內(nèi)部，然后在內(nèi)部進行處理鸵熟，這樣可以節(jié)省反復(fù)讀取的時間副编。
有了這個思路，結(jié)合上節(jié)看到的SM結(jié)構(gòu)圖流强，看到有一片存儲器叫做Shared Memory痹届，它位于SM內(nèi)部，處理時訪問速度相當(dāng)快（差不多每個時鐘周期讀一次）打月，而全局存儲器讀一次需要耗費幾十甚至上百個時鐘周期队腐。于是，我們就制定A計劃如下：
線程塊數(shù)：1奏篙，塊號為0柴淘；（只有一個線程塊內(nèi)的線程才能進行通信，所以我們只分配一個線程塊报破，具體工作交給每個線程完成）
線程數(shù)：5，線程號分別為0~4千绪；（線程并行充易，前面講過）
共享存儲器大小：5個int型變量大休┬汀（5 * sizeof(int））盹靴。
步驟一：讀取輸入數(shù)據(jù)。將Global Memory中的5個整數(shù)讀入共享存儲器，位置一一對應(yīng)稿静，和線程號也一一對應(yīng)梭冠，所以可以同時完成。
步驟二：線程同步改备，確保所有線程都完成了工作控漠。
步驟三：指定線程，對共享存儲器中的輸入數(shù)據(jù)完成相應(yīng)處理悬钳。
代碼如下：
[cpp]view plaincopy
#include?"cuda_runtime.h"
#include?"device_launch_parameters.h"
#include?
cudaError_t?addWithCuda(int*c,constint*a,size_tsize);
__global__voidaddKernel(int*c,constint*a)
{
inti?=?threadIdx.x;
extern__shared__intsmem[];
???smem[i]?=?a[i];
__syncthreads();
if(i?==?0)//?0號線程做平方和
{
c[0]?=?0;
for(intd?=?0;?d?<?5;?d++)
{
c[0]?+=?smem[d]?*?smem[d];
}
}
if(i?==?1)//1號線程做累加
{
c[1]?=?0;
for(intd?=?0;?d?<?5;?d++)
{
c[1]?+=?smem[d];
}
}
if(i?==?2)//2號線程做累乘
{
c[2]?=?1;
for(intd?=?0;?d?<?5;?d++)
{
c[2]?*=?smem[d];
}
}
}
intmain()
{
constintarraySize?=?5;
constinta[arraySize]?=?{?1,?2,?3,?4,?5?};
intc[arraySize]?=?{?0?};
//?Add?vectors?in?parallel.
cudaError_t?cudaStatus?=?addWithCuda(c,?a,?arraySize);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"addWithCuda?failed!");
return1;
}
printf("\t1+2+3+4+5?=?%d\n\t1^2+2^2+3^2+4^2+5^2?=?%d\n\t1*2*3*4*5?=?%d\n\n\n\n\n\n",?c[1],?c[0],?c[2]);
//?cudaThreadExit?must?be?called?before?exiting?in?order?for?profiling?and
//?tracing?tools?such?as?Nsight?and?Visual?Profiler?to?show?complete?traces.
cudaStatus?=?cudaThreadExit();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadExit?failed!");
return1;
}
return0;
}
//?Helper?function?for?using?CUDA?to?add?vectors?in?parallel.
cudaError_t?addWithCuda(int*c,constint*a,size_tsize)
{
int*dev_a?=?0;
int*dev_c?=?0;
cudaError_t?cudaStatus;
//?Choose?which?GPU?to?run?on,?change?this?on?a?multi-GPU?system.
cudaStatus?=?cudaSetDevice(0);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?");
gotoError;
}
//?Allocate?GPU?buffers?for?three?vectors?(two?input,?one?output)????.
cudaStatus?=?cudaMalloc((void**)&dev_c,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_a,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
//?Copy?input?vectors?from?host?memory?to?GPU?buffers.
cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
//?Launch?a?kernel?on?the?GPU?with?one?thread?for?each?element.
????addKernel<<<1,?size,?size?*sizeof(int),?0>>>(dev_c,?dev_a);
//?cudaThreadSynchronize?waits?for?the?kernel?to?finish,?and?returns
//?any?errors?encountered?during?the?launch.
cudaStatus?=?cudaThreadSynchronize();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n",?cudaStatus);
gotoError;
}
//?Copy?output?vector?from?GPU?buffer?to?host?memory.
cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*sizeof(int),?cudaMemcpyDeviceToHost);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
returncudaStatus;
}
從代碼中看到執(zhí)行配置<<<>>>中第三個參數(shù)為共享內(nèi)存大醒谓荨（字節(jié)數(shù)），這樣我們就知道了全部4個執(zhí)行配置參數(shù)的意義默勾。恭喜碉渡，你的CUDA終于入門了！
CUDA從入門到精通（十）：性能剖析和Visual Profiler
入門后的進一步學(xué)習(xí)的內(nèi)容母剥，就是如何優(yōu)化自己的代碼滞诺。我們前面的例子沒有考慮任何性能方面優(yōu)化，是為了更好地學(xué)習(xí)基本知識點环疼，而不是其他細節(jié)問題习霹。從本節(jié)開始，我們要從性能出發(fā)考慮問題秦爆，不斷優(yōu)化代碼序愚，使執(zhí)行速度提高是并行處理的唯一目的。
測試代碼運行速度有很多方法等限，C語言里提供了類似于SystemTime()這樣的API獲得系統(tǒng)時間爸吮，然后計算兩個事件之間的時長從而完成計時功能。在CUDA中望门，我們有專門測量設(shè)備運行時間的API形娇，下面一一介紹。
翻開編程手冊《CUDA_Toolkit_Reference_Manual》筹误，隨時準(zhǔn)備查詢不懂得API桐早。我們在運行核函數(shù)前后，做如下操作：
[cpp]view plaincopy
cudaEvent_t?start,?stop;??//事件對象
cudaEventCreate(&start);??//創(chuàng)建事件
cudaEventCreate(&stop);???????//創(chuàng)建事件
cudaEventRecord(start,?stream);???//記錄開始
myKernel<<>>(parameter?list);//執(zhí)行核函數(shù)
cudaEventRecord(stop,stream);?//記錄結(jié)束事件
cudaEventSynchronize(stop);???//事件同步厨剪，等待結(jié)束事件之前的設(shè)備操作均已完成
floatelapsedTime;
cudaEventElapsedTime(&elapsedTime,start,stop);//計算兩個事件之間時長（單位為ms）
核函數(shù)執(zhí)行時間將被保存在變量elapsedTime中哄酝。通過這個值我們可以評估算法的性能。下面給一個例子祷膳，來看怎么使用計時功能陶衅。
前面的例子規(guī)模很小，只有5個元素直晨，處理量太小不足以計時搀军，下面將規(guī)模擴大為1024膨俐，此外將反復(fù)運行1000次計算總時間，這樣估計不容易受隨機擾動影響罩句。我們通過這個例子對比線程并行和塊并行的性能如何焚刺。代碼如下：
[cpp]view plaincopy
#include?"cuda_runtime.h"
#include?"device_launch_parameters.h"
#include?
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize);
__global__voidaddKernel_blk(int*c,constint*a,constint*b)
{
inti?=?blockIdx.x;
c[i]?=?a[i]+?b[i];
}
__global__voidaddKernel_thd(int*c,constint*a,constint*b)
{
inti?=?threadIdx.x;
c[i]?=?a[i]+?b[i];
}
intmain()
{
constintarraySize?=?1024;
inta[arraySize]?=?{0};
intb[arraySize]?=?{0};
for(inti?=?0;i
{
a[i]?=?i;
b[i]?=?arraySize-i;
}
intc[arraySize]?=?{0};
//?Add?vectors?in?parallel.
cudaError_t?cudaStatus;
intnum?=?0;
cudaDeviceProp?prop;
cudaStatus?=?cudaGetDeviceCount(&num);
for(inti?=?0;i
{
cudaGetDeviceProperties(&prop,i);
}
cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"addWithCuda?failed!");
return1;
}
//?cudaThreadExit?must?be?called?before?exiting?in?order?for?profiling?and
//?tracing?tools?such?as?Nsight?and?Visual?Profiler?to?show?complete?traces.
cudaStatus?=?cudaThreadExit();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadExit?failed!");
return1;
}
for(inti?=?0;i
{
if(c[i]?!=?(a[i]+b[i]))
{
printf("Error?in?%d\n",i);
}
}
return0;
}
//?Helper?function?for?using?CUDA?to?add?vectors?in?parallel.
cudaError_t?addWithCuda(int*c,constint*a,constint*b,size_tsize)
{
int*dev_a?=?0;
int*dev_b?=?0;
int*dev_c?=?0;
cudaError_t?cudaStatus;
//?Choose?which?GPU?to?run?on,?change?this?on?a?multi-GPU?system.
cudaStatus?=?cudaSetDevice(0);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?");
gotoError;
}
//?Allocate?GPU?buffers?for?three?vectors?(two?input,?one?output)????.
cudaStatus?=?cudaMalloc((void**)&dev_c,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_a,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
cudaStatus?=?cudaMalloc((void**)&dev_b,?size?*sizeof(int));
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMalloc?failed!");
gotoError;
}
//?Copy?input?vectors?from?host?memory?to?GPU?buffers.
cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*sizeof(int),?cudaMemcpyHostToDevice);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
cudaEvent_t?start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
for(inti?=?0;i<1000;i++)
{
//??????addKernel_blk<<>>(dev_c,?dev_a,?dev_b);
addKernel_thd<<<1,size>>>(dev_c,?dev_a,?dev_b);
}
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
floattm;
cudaEventElapsedTime(&tm,start,stop);
printf("GPU?Elapsed?time:%.6f?ms.\n",tm);
//?cudaThreadSynchronize?waits?for?the?kernel?to?finish,?and?returns
//?any?errors?encountered?during?the?launch.
cudaStatus?=?cudaThreadSynchronize();
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n",?cudaStatus);
gotoError;
}
//?Copy?output?vector?from?GPU?buffer?to?host?memory.
cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*sizeof(int),?cudaMemcpyDeviceToHost);
if(cudaStatus?!=?cudaSuccess)
{
fprintf(stderr,"cudaMemcpy?failed!");
gotoError;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
returncudaStatus;
}
addKernel_blk是采用塊并行實現(xiàn)的向量相加操作，而addKernel_thd是采用線程并行實現(xiàn)的向量相加操作门烂。分別運行线欲，得到的結(jié)果如下圖所示：
線程并行：
塊并行：
可見性能竟然相差近16倍甚疟！因此選擇并行處理方法時点骑，如果問題規(guī)模不是很大需忿，那么采用線程并行是比較合適的，而大問題分多個線程塊處理時氓润，每個塊內(nèi)線程數(shù)不要太少赂乐，像本文中的只有1個線程，這是對硬件資源的極大浪費咖气。一個理想的方案是挨措，分N個線程塊，每個線程塊包含512個線程崩溪，將問題分解處理浅役，效率往往比單一的線程并行處理或單一塊并行處理高很多。這也是CUDA編程的精髓伶唯。
上面這種分析程序性能的方式比較粗糙觉既，只知道大概運行時間長度，對于設(shè)備程序各部分代碼執(zhí)行時間沒有一個深入的認識乳幸，這樣我們就有個問題瞪讼，如果對代碼進行優(yōu)化，那么優(yōu)化哪一部分呢粹断？是將線程數(shù)調(diào)節(jié)呢符欠，還是改用共享內(nèi)存？這個問題最好的解決方案就是利用Visual Profiler瓶埋。下面內(nèi)容摘自《CUDA_Profiler_Users_Guide》
“Visual Profiler是一個圖形化的剖析工具希柿，可以顯示你的應(yīng)用程序中CPU和GPU的活動情況，利用分析引擎幫助你尋找優(yōu)化的機會养筒≡罚”
其實除了可視化的界面，NVIDIA提供了命令行方式的剖析命令：nvprof晕粪。對于初學(xué)者挤悉，使用圖形化的方式比較容易上手，所以本節(jié)使用Visual Profiler兵多。
打開Visual Profiler尖啡，可以從CUDA Toolkit安裝菜單處找到。主界面如下：
我們點擊File->New Session剩膘，彈出新建會話對話框衅斩，如下圖所示：
其中File一欄填入我們需要進行剖析的應(yīng)用程序exe文件，后面可以都不填（如果需要命令行參數(shù)怠褐，可以在第三行填入）畏梆，直接Next，見下圖：
第一行為應(yīng)用程序執(zhí)行超時時間設(shè)定奈懒，可不填奠涌；后面三個單選框都勾上，這樣我們分別使能了剖析磷杏，使能了并發(fā)核函數(shù)剖析溜畅，然后運行分析器。
點Finish极祸，開始運行我們的應(yīng)用程序并進行剖析慈格、分析性能。
上圖中遥金，CPU和GPU部分顯示了硬件和執(zhí)行內(nèi)容信息浴捆，點某一項則將時間條對應(yīng)的部分高亮，便于觀察稿械，同時右邊詳細信息會顯示運行時間信息选泻。從時間條上看出，cudaMalloc占用了很大一部分時間美莫。下面分析器給出了一些性能提升的關(guān)鍵點页眯，包括：低計算利用率（計算時間只占總時間的1.8%，也難怪茂嗓，加法計算復(fù)雜度本來就很低呀２鸵稹）；低內(nèi)存拷貝/計算交疊率（一點都沒有交疊述吸，完全是拷貝——計算——拷貝）忿族；低存儲拷貝尺寸（輸入數(shù)據(jù)量太小了，相當(dāng)于你淘寶買了個日記本蝌矛，運費比實物價格還高５琅）；低存儲拷貝吞吐率（只有1.55GB/s）入撒。這些對我們進一步優(yōu)化程序是非常有幫助的隆豹。
我們點一下Details，就在Analysis窗口旁邊茅逮。得到結(jié)果如下所示：
通過這個窗口可以看到每個核函數(shù)執(zhí)行時間璃赡，以及線程格判哥、線程塊尺寸，占用寄存器個數(shù)碉考，靜態(tài)共享內(nèi)存塌计、動態(tài)共享內(nèi)存大小等參數(shù)，以及內(nèi)存拷貝函數(shù)的執(zhí)行情況侯谁。這個提供了比前面cudaEvent函數(shù)測時間更精確的方式锌仅，直接看到每一步的執(zhí)行時間，精確到ns墙贱。
在Details后面還有一個Console热芹，點一下看看。
這個其實就是命令行窗口惨撇，顯示運行輸出伊脓。看到加入了Profiler信息后魁衙，總執(zhí)行時間變長了（原來線程并行版本的程序運行時間只需4ms左右）丽旅。這也是“測不準(zhǔn)定理”決定的，如果我們希望測量更細微的時間纺棺，那么總時間肯定是不準(zhǔn)的榄笙；如果我們希望測量總時間，那么細微的時間就被忽略掉了祷蝌。
后面Settings就是我們建立會話時的參數(shù)配置茅撞，不再詳述。
通過本節(jié)巨朦，我們應(yīng)該能對CUDA性能提升有了一些想法米丘，好，下一節(jié)我們將討論如何優(yōu)化CUDA程序糊啡。
http://blog.csdn.net/kkk584520/article/details/9413973
http://blog.csdn.net/kkk584520/article/details/9414191
http://blog.csdn.net/kkk584520/article/details/9415199
http://blog.csdn.net/kkk584520/article/details/9417251
http://blog.csdn.net/kkk584520/article/details/9420793
http://blog.csdn.net/kkk584520/article/details/9428389
http://blog.csdn.net/kkk584520/article/details/9428859
http://blog.csdn.net/kkk584520/article/details/9449635
http://blog.csdn.net/kkk584520/article/details/9472695
