1、Mat lab使用GPU并行加速方法GPU具有十分強(qiáng)大的數(shù)值計(jì)算能力，它使用大規(guī)模并行方式進(jìn)行加速。Matlab是十分重要的數(shù)學(xué)語言，矩陣計(jì)算十分方便。但是Matlab是解釋型語言， 執(zhí)行相對(duì)較慢。我們可以使用GPU對(duì)Matlab進(jìn)行加速。Matlab調(diào)用GPU加速方 法很多，主要有：1在GPU上執(zhí)行重載的MATLAB函數(shù)1.1最簡(jiǎn)單的編程模式對(duì)GPU上已加載數(shù)據(jù)的Matlab函數(shù)直接調(diào)用。Matlab已經(jīng)重載了很多GPU 標(biāo)準(zhǔn)函數(shù)。優(yōu)點(diǎn)用戶可以決定何時(shí)在Matlab工作區(qū)和GPU之間移動(dòng)數(shù)據(jù)或創(chuàng)建存儲(chǔ)在 GPU內(nèi)存中的數(shù)據(jù)，以盡可能減少主機(jī)與設(shè)備間數(shù)據(jù)傳輸?shù)拈_銷。用戶可在同一函數(shù)調(diào)用中將在G

2、PU上加載的數(shù)據(jù)和Matlab工作區(qū)中的數(shù) 據(jù)混合，以實(shí)現(xiàn)最優(yōu)的靈活性與易用性。這種方法提供了一個(gè)簡(jiǎn)單的接口，讓用戶可以在GPU上直接執(zhí)行標(biāo)準(zhǔn)函 數(shù)，從而獲得性能提升，而無需花費(fèi)任何時(shí)間開發(fā)專門的代碼。缺點(diǎn)在這種情況下，用戶不得對(duì)函數(shù)進(jìn)行任何更改，只能指定何時(shí)從GPU內(nèi)存 移動(dòng)和檢索數(shù)據(jù)，這兩種操作分別通過gpuArray和gather命令來完成。11試開始CFU上矩陣乘法執(zhí)行的時(shí)間為。.5姬刑L秒！ 在CFU上總執(zhí)行時(shí)間為。.777216秒！CFU上矩障乘法執(zhí)行的時(shí)間為Q. 口。1434秒！在GFU上總執(zhí)行時(shí)間為L(zhǎng) 9S4727秒！GFU加速比。,598541/0. 001434=417.

3、4858341.2在Matlab中定義GPU內(nèi)核用戶可以定義Matlab函數(shù)，執(zhí)行對(duì)GPU上的數(shù)據(jù)的標(biāo)量算術(shù)運(yùn)算。使用這 種方法，用戶可以擴(kuò)展和自定義在GPU上執(zhí)行的函數(shù)集，以構(gòu)建復(fù)雜應(yīng)用程序并 實(shí)現(xiàn)性能加速。這種方式需要進(jìn)行的內(nèi)核調(diào)用和數(shù)據(jù)傳輸比上述方法少。優(yōu)點(diǎn)這種編程模式允許用算術(shù)方法定義要在GPU上執(zhí)行的復(fù)雜內(nèi)核，只需使用 Matlab語言即可。使用這種方法，可在GPU上執(zhí)行復(fù)雜的算術(shù)運(yùn)算，充分利用數(shù)據(jù)并行化并 最小化與內(nèi)核調(diào)用和數(shù)據(jù)傳輸有關(guān)的開銷。缺點(diǎn)在這種情況下，用戶不得對(duì)函數(shù)進(jìn)行任何更改，只能指定何時(shí)從GPU內(nèi)存 移動(dòng)和檢索數(shù)據(jù)以及使用arrayfun命令調(diào)用函數(shù)。函數(shù)會(huì)在GPU矢

4、量的各個(gè)元素上執(zhí)行，充分利用數(shù)據(jù)并行化。GPUMQdBl2KBrnel.m 笠 | GPUMacfBl2.m + 1定火即U kernel:zurction v=FUMo de1ZKerne1(k j-H1+k, (H-.dk. * f (l+x, f (1-Hs, k (1+,.S - L ei.L%在晌tL布中同憐-clear;-tic;eM牛成教精-N-GOOO;9 -占=rand(N*N“1):-fprintf。迥試松；toLi crvth 行L2 一tie：L3 - Ycpu=&PUIodel2Kernel (A);L-4 cpuTiai9=toc ;L5 -fprlntf C11

5、XnCFUlh.行時(shí)間為：Vf 利;！. cpuTime):LG17GPU 撕行L日-Agpu-gpuluay W);LS 一tic;-Ygpuarrayfuii(FUModel2EexneL Agpu);gpuTine=toc ;一Y2Cpu=2&TheT(lEDLl):-fprint： nGPUih行時(shí)間為：秒！ n ? gpuTimeJ :242o - fpxintf nGFU加 i吏比%/%-Enn|? ? cpuTiiis gpulims, cpuTine/gpurine)刪試開始CFUttl行時(shí)間0.072501!GPU執(zhí)行時(shí)間為:0. 000727!GPlrtnl 比Q. 072

6、501/0. 000727=99. 7252941.3直接從Mat lab調(diào)用CUDA代碼為了進(jìn)一步擴(kuò)展在GPU上執(zhí)行的集合函數(shù)，可以從CUDA代碼中創(chuàng)建一個(gè) Matlab可調(diào)用的GPU內(nèi)核。第三種編程模式可以讓用戶輕松地從Matlab直接調(diào) 用已有CUDA代碼，使非CUDA專家同樣能夠進(jìn)行代碼重用。優(yōu)點(diǎn)這種編程模式提供了直接從Matlab進(jìn)行CUDA代碼測(cè)試的整體解決方案，無需使用GPU在環(huán)配置進(jìn)行基于文件的數(shù)據(jù)交換。用戶還可以直接從Matlab控制有關(guān)線程塊大小和共享內(nèi)存的參數(shù)。 缺點(diǎn)用戶需要會(huì)CUDA編碼。頊 GPUMcjWimm 閥+ 1%直接從Mat ME調(diào)用CUHMt甜蝙- Cud

7、a2Mex G-PUModel2Kerne 1. cu );蘋-aSI涵中定義CUDA keroel%定義iernelttL行屬性-k=parallel. gpu. CUDAKernelf GPUModel2Keinel. ptK, 3 GPUMode 12Kerne 1. cu )-k.GxidSize= 12, ?12:10-k. ThreadBlack5ise= 132,321:12-r=30oa；-A=rand(N, 1):15M IS希CUD-A_gpu=gpuArray (A):-Y_gpu=gpuArr ay (zeros (Nn 1);-Y_cpu=gather (f eval

8、 (k, Y_gpu, A_gpu, N):202、Mat lab與CUDA C混合編程用Matlab與C/C+混合編程，采用動(dòng)態(tài)鏈接庫(kù)的方式產(chǎn)生可以供Matlab調(diào) 用的.dll文件。該方法使用CUDA C/C+語言編寫在GPU上執(zhí)行的代碼，將之編 譯成.dll文件，然后使用C/C+語言編寫mexFunction函數(shù)，在函數(shù)中加載使用 CUDA的.dll文件，使用Matlab或者VC+編譯mexFunction為另一個(gè).dll文件。 最后在Matlab中調(diào)用含有mexFunction的.dll文件，執(zhí)行GPU加速。47矩陣乘法的kbH以函數(shù)4S E lobal void kernelfflo

9、at *Hj -float *Nj float *Pj int width)49r泌計(jì)算巳h中元素的行素弓51int Row = b-l-ackldsjy*BLOCK WIDTH 4- threadldxyj52計(jì)具巳W中元素的列素弓53int lockIdK.k*BLOCIC WIDTH 4- threadldK.k;54“Pv嘰如存摘線程計(jì)算得到的P律的值試shared float Rvalue = 0.;56每個(gè)線程計(jì)尊于矩釁的卜個(gè)元素57for (int k =k width; k+)5859Rvalue 4-= r.Row*width 4- k * Nk*vidth 4- Col;6

10、0卜 -SL把矩陣寫入邵u內(nèi)存，與線程寫一62P ll*width 十 CqI = Pvaluej6365矩眸乘法全局函數(shù)66 日vcid raatnul(float *虬 float *吼 float 叩/ int widthES6970717273747576H79BSI8283848586878889int sise = width*jidth*izeof(float) ifloat *Md j, *Mdj *Pdj/ ,在石PLJ內(nèi)存中分配存儲(chǔ)空間oudaHallDC(voidsize)jcudaMalloc( (viid *)S_Ndj size) jcudaMalloc( (vid

11、size) j將虬喂制到GPU內(nèi)存中cudaMemcpy(Md j si z-Cj cudaMemcpyHos tToDevLce) j cudaMemcpy( Nd j N. siz-Cj cudaMemcpyHostToDevLce) jW|d imB loc k (B LOC K_W IDTH t aLOtKJ-JL&THj 1);dhn3 ddntGrid(width 7 BLOCK_WIDTHj width / BLOCK-WIDTrtj啟動(dòng)計(jì)算線程_kernel (Ndj Pdj width 把P從研旗制到主機(jī)內(nèi)宿中cudaMencpy(Pj Pdj 5iz-Cj cudaMemc

12、pyDeviceToHast)j /rafcspu 內(nèi)存cudaFree(Mcl);cudaFreetMd);cudaFreetFcl);121314151617 IB92&21222324252627259 魂 313233343536373S9void mexFunction(int nlh5j msArray *plhsj int nrhSj const mxArray*prh5) - -if (nrhs != 2 iiieKErrMsgTKt。泌成有兩饞入?yún)?shù),力if (nihs != 1me kE rrMs gTx (必須有一個(gè)輸出卷敷.)；size_t md = niKGetNun

13、iberOf&iinensionstprhsf);size_t nd = mxGetMuinberOfDiinensions(prhsL);if (md = 2 I I nd != 2mekErrMs呂Txt (如輻入矩陣必須是二矩陣.)Jsize_t m = niKGetHfprhs Q) t size_t p = imxGetNprhs 0 j size_t q = mKGetM(prhs1); size_t n = nixGetNtprhs 1); if (p != q me kE rrMs gT Kt 矩陣鍬jj不匹配,11 )；hplhs0 = micCreateDoub-leMat

14、rixnij n msREAL);size_t ni5 = m*p;double *px = nucGetPr(p-rhs&)fliat 4)c=new float ms i44849 |5&51525354555657585560616253636465| 661 e7lES69701717273747576777B 73 E| oilE2|S3B4|翡86B79 9float *K=new float as;f?r fint i = i m; i-n-)for (int j = 0; j p; *十) -vi*P + j =4- i5Jsize_t ns = q*

15、n;double *py = mKGetPr(prh51);float *y = new floatns;for fint i = #_; i 霜 i-n-)for (int j = 0; j double *pz = nxGetPr(plhs;0) j5ize_t Z5 = m*njfLoat *z = new floatzs;TCHAR *pMame = _TfCU&Adlltest .dll) jCHAR *pFuncMane = mstmul;HINSTANCE hDll = LoadLibrary(pNanie) j /iffiiiDCLlA dll test j dllHINSTA

16、NCE hDll = LoadLitrary(pMame)j/UDAdlltest.dll if (hDll) 荻得函教指針DLLFUNC dllFun = (DLL FUNC)Get ProcA ddress(hDll, pFuncMame); if (dllFun)dllFun (KjyjZj (int)m) ;，/執(zhí)行函數(shù) for (int i = -9; i n; i+4-) , , , fc r (int j = % j nj j-+) “一 pz j *m + i = z in + ji - else/何能由于函數(shù)名措誤riexErrHsgTxt(Can ncrt find the

17、 function in dll! *); _FreeLibrary(hDlL) iZ/TitPSClIDAdlltest. dllelse(meKErrHsgTK(J Load dll faM.!”);魚 CUDAdlltest.dll2015/4/24 15:56Application erten&.14 KB肴 matVScuda2-015/4/24 16:28MATLAB Code1 KB闖 TestDLLdll2015/4/24 1&刀Applicaiion ectens.15 KB2-01W24 1 女召M(fèi)ATLAB MEX15 KBmatVScuda.m 明 +1。比較CUIU匚與旅it 1戒鼓奉一clearclc-size2660;-X=rand(size) ;|一Y=r and (size);一tic;-M=K*Y：-u.atlabTLine=tcc:一fprint(? iiMatlab time - %fn? ! mat lab rime);1314 一tic;IE 一I=t&2tp (X, V):一cudalime=toc;一pr int(? nCUDA time =,cudaTine);18-r-M-N-r=r.*T；一r=sumiswn(T);2223 -fpr intf C nSUM( (mat 1 abRes-cudaRes) 2) = %fn