Matlab使用GPU并行加速方法

1、Mat lab使用GPU并行加速方法GPU具有十分强大的数值计算能力，它使用大规模并行方式进行加速。Matlab是十分重要的数学语言，矩阵计算十分方便。但是Matlab是解释型语言， 执行相对较慢。我们可以使用GPU对Matlab进行加速。Matlab调用GPU加速方 法很多，主要有：1在GPU上执行重载的MATLAB函数1.1最简单的编程模式对GPU上已加载数据的Matlab函数直接调用。Matlab已经重载了很多GPU 标准函数。优点用户可以决定何时在Matlab工作区和GPU之间移动数据或创建存储在 GPU内存中的数据，以尽可能减少主机与设备间数据传输的开销。用户可在同一函数调用中将在G

2、PU上加载的数据和Matlab工作区中的数 据混合，以实现最优的灵活性与易用性。这种方法提供了一个简单的接口，让用户可以在GPU上直接执行标准函 数，从而获得性能提升，而无需花费任何时间开发专门的代码。缺点在这种情况下，用户不得对函数进行任何更改，只能指定何时从GPU内存 移动和检索数据，这两种操作分别通过gpuArray和gather命令来完成。11试开始CFU上矩阵乘法执行的时间为。.5姬刑L秒！ 在CFU上总执行时间为。.777216秒！CFU上矩障乘法执行的时间为Q. 口。1434秒！在GFU上总执行时间为L 9S4727秒！GFU加速比。,598541/0. 001434=417.

3、4858341.2在Matlab中定义GPU内核用户可以定义Matlab函数，执行对GPU上的数据的标量算术运算。使用这 种方法，用户可以扩展和自定义在GPU上执行的函数集，以构建复杂应用程序并 实现性能加速。这种方式需要进行的内核调用和数据传输比上述方法少。优点这种编程模式允许用算术方法定义要在GPU上执行的复杂内核，只需使用 Matlab语言即可。使用这种方法，可在GPU上执行复杂的算术运算，充分利用数据并行化并 最小化与内核调用和数据传输有关的开销。缺点在这种情况下，用户不得对函数进行任何更改，只能指定何时从GPU内存 移动和检索数据以及使用arrayfun命令调用函数。函数会在GPU矢

4、量的各个元素上执行，充分利用数据并行化。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执行时间为:0. 000727!GPlrtnl 比Q. 072

6、501/0. 000727=99. 7252941.3直接从Mat lab调用CUDA代码为了进一步扩展在GPU上执行的集合函数，可以从CUDA代码中创建一个 Matlab可调用的GPU内核。第三种编程模式可以让用户轻松地从Matlab直接调 用已有CUDA代码，使非CUDA专家同样能够进行代码重用。优点这种编程模式提供了直接从Matlab进行CUDA代码测试的整体解决方案，无需使用GPU在环配置进行基于文件的数据交换。用户还可以直接从Matlab控制有关线程块大小和共享内存的参数。 缺点用户需要会CUDA编码。顼 GPUMcjWimm 阀+ 1%直接从Mat ME调用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+混合编程，采用动态链接库的方式产生可以供Matlab调 用的.dll文件。该方法使用CUDA C/C+语言编写在GPU上执行的代码，将之编 译成.dll文件，然后使用C/C+语言编写mexFunction函数，在函数中加载使用 CUDA的.dll文件，使用Matlab或者VC+编译mexFunction为另一个.dll文件。 最后在Matlab中调用含有mexFunction的.dll文件，执行GPU加速。47矩阵乘法的kbH以函数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内存，与线程写一62P ll*width 十 CqI = Pvaluej6365矩眸乘法全局函数66 日vcid raatnul(float *虬 float *吼 float 叩/ int widthES6970717273747576H79BSI8283848586878889int sise = width*jidth*izeof(float) ifloat *Md j, *Mdj *Pdj/ ,在石PLJ内存中分配存储空间oudaHallDC(voidsize)jcudaMalloc( (viid *)S_Ndj size) jcudaMalloc( (vid

11、size) j将虬喂制到GPU内存中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从研旗制到主机内宿中cudaMencpy(Pj Pdj 5iz-Cj cudaMemc

12、pyDeviceToHast)j /rafcspu 内存cudaFree(Mcl);cudaFreetMd);cudaFreetFcl);121314151617 IB92&21222324252627259 魂 313233343536373S9void mexFunction(int nlh5j msArray *plhsj int nrhSj const mxArray*prh5) - -if (nrhs != 2 iiieKErrMsgTKt。泌成有两馋入参数,力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) ;，/执行函数 for (int i = -9; i n; i+4-) , , , fc r (int j = % j nj j-+) “一 pz j *m + i = z in + ji - else/何能由于函数名措误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