Matlab使用GPU并行加速方法

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

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

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

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

6、501/0. 000727=99. 7252941.3直接從Mat lab調(diào)用CUDA代碼為了進一步擴展在GPU上執(zhí)行的集合函數(shù)，可以從CUDA代碼中創(chuàng)建一個 Matlab可調(diào)用的GPU內(nèi)核。第三種編程模式可以讓用戶輕松地從Matlab直接調(diào) 用已有CUDA代碼，使非CUDA專家同樣能夠進行代碼重用。優(yōu)點這種編程模式提供了直接從Matlab進行CUDA代碼測試的整體解決方案，無需使用GPU在環(huán)配置進行基于文件的數(shù)據(jù)交換。用戶還可以直接從Matlab控制有關(guān)線程塊大小和共享內(nèi)存的參數(shù)。 缺點用戶需要會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+混合編程，采用動態(tài)鏈接庫的方式產(chǎn)生可以供Matlab調(diào) 用的.dll文件。該方法使用CUDA C/C+語言編寫在GPU上執(zhí)行的代碼，將之編 譯成.dll文件，然后使用C/C+語言編寫mexFunction函數(shù)，在函數(shù)中加載使用 CUDA的.dll文件，使用Matlab或者VC+編譯mexFunction為另一個.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泌計算巳h中元素的行素弓51int Row = b-l-ackldsjy*BLOCK WIDTH 4- threadldxyj52計具巳W中元素的列素弓53int lockIdK.k*BLOCIC WIDTH 4- threadldK.k;54“Pv嘰如存摘線程計算得到的P律的值試shared float Rvalue = 0.;56每個線程計尊于矩釁的卜個元素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)存中分配存儲空間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啟動計算線程_kernel (Ndj Pdj width 把P從研旗制到主機內(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 (必須有一個輸出卷敷.)；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 女召MATLAB 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