1、Matlab使用GPU并行加速方法GPU具有十分强大的数值计算能力,它使用大规模并行方式进行加速。Matlab是十分重要的数学语言,矩阵计算十分方便。但是Matlab是解释型语言,执行相对较慢。我们可以使用GPU对Matlab进行加速。Matlab调用GPU加速方法很多,主要有:1在GPU上执行重载的MATLAB函数1.1最简单的编程模式对GPU上已加载数据的Matlab函数直接调用。Matlab已经重载了很多GPU标准函数。优点用户可以决定何时在Matlab工作区和GPU之间移动数据或创建存储在GPU内存中的数据,以尽可能减少主机与设备间数据传输的开销。用户可在同一函数调用中将在G
3、4858341.2在Matlab中定义GPU内核用户可以定义Matlab函数,执行对GPU上的数据的标量算术运算。使用这种方法,用户可以扩展和自定义在GPU上执行的函数集,以构建复杂应用程序并实现性能加速。这种方式需要进行的内核调用和数据传输比上述方法少。优点这种编程模式允许用算术方法定义要在GPU上执行的复杂内核,只需使用Matlab语言即可。使用这种方法,可在GPU上执行复杂的算术运算,充分利用数据并行化并最小化与内核调用和数据传输有关的开销。缺点在这种情况下,用户不得对函数进行任何更改,只能指定何时从GPU内存移动和检索数据以及使用arrayfun命令调用函数。函数会在GPU矢
4、量的各个元素上执行,充分利用数据并行化。GPUMQdBl2KBrnel.m笠|GPUMacfBl2.m+1定火即Ukernel:zurctionv=FUMode1ZKerne1(kj-H1+k,(H-.dk.*f(l+x,f(1-Hs,k(1+,.S-Lei.L%在晌tL布中同怜-clear;-tic;eM牛成教精-N-GOOO;9-占=rand(N*N“1):-fprintf。迥试松;toLicrvth行L2一tie:L3-Ycpu=&PUIodel2Kernel(A);L-4cpuTiai9=toc;L5-fprlntfC11
6、501/0.000727=99.7252941.3直接从Matlab调用CUDA代码为了进一步扩展在GPU上执行的集合函数,可以从CUDA代码中创建一个Matlab可调用的GPU内核。第三种编程模式可以让用户轻松地从Matlab直接调用已有CUDA代码,使非CUDA专家同样能够进行代码重用。优点这种编程模式提供了直接从Matlab进行CUDA代码测试的整体解决方案,无需使用GPU在环配置进行基于文件的数据交换。用户还可以直接从Matlab控制有关线程块大小和共享内存的参数。缺点用户需要会CUDA编码。顼GPUMcjWimm阀+1%直接从MatME调用CUHMt甜蝙-Cud
7、a2MexG-PUModel2Kerne1.cu);苹-aSI涵中定义CUDAkeroel%定义iernelttL行属性-k=parallel.gpu.CUDAKernelfGPUModel2Keinel.ptK,3GPUMode12Kerne1.cu)-k.GxidSize=12,12:10-k.ThreadBlack5ise=132,321:12-r=30oa;-A=rand(N,1):15MIS希CUD-A_gpu=gpuArray(A):-Y_gpu=gpuArray(zeros(Nn1);-Y_cpu=gather(feval
8、(k,Y_gpu,A_gpu,N):202、Matlab与CUDAC混合编程用Matlab与C/C+混合编程,采用动态链接库的方式产生可以供Matlab调用的.dll文件。该方法使用CUDAC/C+语言编写在GPU上执行的代码,将之编译成.dll文件,然后使用C/C+语言编写mexFunction函数,在函数中加载使用CUDA的.dll文件,使用Matlab或者VC+编译mexFunction为另一个.dll文件。最后在Matlab中调用含有mexFunction的.dll文件,执行GPU加速。47矩阵乘法的kbH以函数4SElobalvoidkernelfflo
9、at*Hj-float*Njfloat*Pjintwidth)49r泌计算巳h中元素的行素弓51intRow=b-l-ackldsjy*BLOCKWIDTH4-threadldxyj52计具巳W中元素的列素弓53intlockIdK.k*BLOCICWIDTH4-threadldK.k;54“Pv叽如存摘线程计算得到的P律的值试sharedfloatRvalue=0.;56每个线程计尊于矩衅的卜个元素57for(intk=kwidth;k+)5859Rvalue4-=r.Row*width4-k*Nk*vidth4-Col;6
10、0卜-SL把矩阵写入邵u内存,与线程写一62Pll*width十CqI=Pvaluej6365矩眸乘法全局函数66日vcidraatnul(float*虬float*吼float叩/intwidthES6970717273747576H79BSI8283848586878889intsise=width*jidth*izeof(float)ifloat*Mdj,*Mdj*Pdj/,在石PLJ内存中分配存储空间oudaHallDC(voidsize)jcudaMalloc((viid*)S_Ndjsize)jcudaMalloc((vid
11、size)j将虬喂制到GPU内存中cudaMemcpy(Mdjsiz-CjcudaMemcpyHostToDevLce)jcudaMemcpy(NdjN.siz-CjcudaMemcpyHostToDevLce)jW|dimBlock(BLOCK_WIDTHtaLOtKJ-JL&THj1);dhn3ddntGrid(width7BLOCK_WIDTHjwidth/BLOCK-WIDTrtj启动计算线程_kernel(NdjPdjwidth把P从研旗制到主机内宿中cudaMencpy(PjPdj5iz-CjcudaMemc
12、pyDeviceToHast)j/rafcspu内存cudaFree(Mcl);cudaFreetMd);cudaFreetFcl);121314151617IB92&21222324252627259魂313233343536373S9voidmexFunction(intnlh5jmsArray*plhsjintnrhSjconstmxArray*prh5)--if(nrhs!=2iiieKErrMsgTKt。泌成有两馋入参数,力if(nihs!=1mekErrMsgTx(必须有一个输出卷敷.);size_tmd=niKGetNun
13、iberOf&iinensionstprhsf);size_tnd=mxGetMuinberOfDiinensions(prhsL);if(md=2IInd!=2mekErrMs吕Txt(如辐入矩阵必须是二矩阵.)Jsize_tm=niKGetHfprhsQ)tsize_tp=imxGetNprhs0jsize_tq=mKGetM(prhs1);size_tn=nixGetNtprhs1);if(p!=qmekErrMsgTKt矩阵锹jj不匹配,11);hplhs0=micCreateDoub-leMat
14、rixnijnmsREAL);size_tni5=m*p;double*px=nucGetPr(p-rhs&)fliat4)c=newfloatmsi44849|5&51525354555657585560616253636465|661e7lES69701717273747576777B73E|oilE2|S3B4|翡86B799float*K=newfloatas;frfinti=im;i-n-)for(intj=0;jp;*十)-vi*P+j=4-i5Jsize_tns=q*
15、n;double*py=mKGetPr(prh51);float*y=newfloatns;forfinti=#_;i霜i-n-)for(intj=0;jdouble*pz=nxGetPr(plhs;0)j5ize_tZ5=m*njfLoat*z=newfloatzs;TCHAR*pMame=_TfCU&Adlltest.dll)jCHAR*pFuncMane=mstmul;HINSTANCEhDll=LoadLibrary(pNanie)j/iffiiiDCLlAdlltestjdllHINSTA
16、NCEhDll=LoadLitrary(pMame)j/UDAdlltest.dllif(hDll)荻得函教指针DLLFUNCdllFun=(DLLFUNC)GetProcAddress(hDll,pFuncMame);if(dllFun)dllFun(KjyjZj(int)m);,/执行函数for(inti=-9;in;i+4-),,,fcr(intj=%jnjj-+)“一pzj*m+i=zin+ji-else/何能由于函数名措误riexErrHsgTxt(Canncrtfindthe
17、functionindll!*);_FreeLibrary(hDlL)iZ/TitPSClIDAdlltest.dllelse(meKErrHsgTK(JLoaddllfaM.!”);鱼CUDAdlltest.dll2015/4/2415:56Applicationerten&.14KB肴matVScuda2-015/4/2416:28MATLABCode1KB闯TestDLLdll2015/4/241&刀Applicaiionectens.15KB2-01W241女召MATLABMEX15KBmatVScuda.m明+1。比较CUIU匚与旅it1戒鼓奉一clearclc-size2660;-X=rand(size);|一Y=rand(size);一tic;-M=K*Y:-u.atlabTLine=tcc:一fprint(iiMatlabtime-%fn!matlabrime);1314一tic;IE一I=t&2tp(X,V):一cudalime=toc;一print(nCUDAtime=,cudaTine);18-r-M-N-r=r.*T;一r=sumiswn(T);2223-fprintfCnSUM((mat1abRes-cudaRes)2)=%fn