关闭

关闭

关闭

封号提示

内容

首页 深入浅出谈CUDA.pdf

深入浅出谈CUDA.pdf

深入浅出谈CUDA.pdf

上传者: 暗黑天使鲁西法 2011-09-19 评分 0 0 0 0 0 0 暂无简介 简介 举报

简介:本文档为《深入浅出谈CUDApdf》,可适用于IT/计算机领域,主题内容包含深入浅出谈CUDAHotballCUDA是NVIDIA的GPGPU模型它使用C语言为基础可以直接以大多数人熟悉的C语言写出在显示芯片上执行的程序而不符等。

深入浅出谈CUDAHotballCUDA是NVIDIA的GPGPU模型它使用C语言为基础可以直接以大多数人熟悉的C语言写出在显示芯片上执行的程序而不需要去学习特定的显示芯片的指令或是特殊的结构。”CUDA是什么?能吃吗?编者注:NVIDIA的GeFoceGTX发布后它的通用计算架构CUDA经过一年多的推广后现在已经在有相当多的论文发表在商业应用软件等方面也初步出现了视频编解码、金融、地质勘探、科学计算等领域的产品是时候让我们对其作更深一步的了解。为了让大家更容易了解CUDA我们征得Hotball的本人同意发表他最近亲自撰写的本文。这篇文章的特点是深入浅出也包含了hotball本人编写一些简单CUDA程序的亲身体验对于希望了解CUDA的读者来说是非常不错的入门文章PCINLIFE对本文的发表没有作任何的删减主要是把一些台湾的词汇转换成大陆的词汇以及作了若干"编者注"的注释。现代的显示芯片已经具有高度的可程序化能力由于显示芯片通常具有相当高的内存带宽以及大量的执行单元因此开始有利用显示芯片来帮助进行一些计算工作的想法即GPGPU。CUDA即是NVIDIA的GPGPU模型。NVIDIA的新一代显示芯片包括GeForce系列及更新的显示芯片都支持CUDA。NVIDIA免费提供CUDA的开发工具(包括Windows版本和Linux版本)、程序范例、文件等等可以在CUDAZone下载。GPGPU的优缺点使用显示芯片来进行运算工作和使用CPU相比主要有几个好处:显示芯片通常具有更大的内存带宽。例如NVIDIA的GeForceGTX具有超过GBs的内存带宽而目前高阶CPU的内存带宽则在GBs左右。显示芯片具有更大量的执行单元。例如GeForceGTX具有个"streamprocessors"频率为GHz。CPU频率通常较高但是执行单元的数目则要少得多。和高阶CPU相比显卡的价格较为低廉。例如目前一张GeForceGT包括MB内存的价格和一颗GHz四核心CPU的价格相若。当然使用显示芯片也有它的一些缺点:显示芯片的运算单元数量很多因此对于不能高度并行化的工作所能带来的帮助就不大。显示芯片目前通常只支持bits浮点数且多半不能完全支持IEEE规格有些运算的精确度可能较低。目前许多显示芯片并没有分开的整数运算单元因此整数运算的效率较差。显示芯片通常不具有分支预测等复杂的流程控制单元因此对于具有高度分支的程序效率会比较差。目前GPGPU的程序模型仍不成熟也还没有公认的标准。例如NVIDIA和AMDATI就有各自不同的程序模型。整体来说显示芯片的性质类似streamprocessor适合一次进行大量相同的工作。CPU则比较有弹性能同时进行变化较多的工作。CUDA架构CUDA是NVIDIA的GPGPU模型它使用C语言为基础可以直接以大多数人熟悉的C语言写出在显示芯片上执行的程序而不需要去学习特定的显示芯片的指令或是特殊的结构。在CUDA的架构下一个程序分为两个部份:host端和device端。Host端是指在CPU上执行的部份而device端则是在显示芯片上执行的部份。Device端的程序又称为"kernel"。通常host端程序会将数据准备好后复制到显卡的内存中再由显示芯片执行device端程序完成后再由host端程序将结果从显卡的内存中取回。由于CPU存取显卡内存时只能透过PCIExpress接口因此速度较慢(PCIExpressx的理论带宽是双向各GBs)因此不能太常进行这类动作以免降低效率。在CUDA架构下显示芯片执行时的最小单位是thread。数个thread可以组成一个block。一个block中的thread能存取同一块共享的内存而且可以快速进行同步的动作。每一个block所能包含的thread数目是有限的。不过执行相同程序的block可以组成grid。不同block中的thread无法存取同一个共享的内存因此无法直接互通或进行同步。因此不同block中的thread能合作的程度是比较低的。不过利用这个模式可以让程序不用担心显示芯片实际上能同时执行的thread数目限制。例如一个具有很少量执行单元的显示芯片可能会把各个block中的thread顺序执行而非同时执行。不同的grid则可以执行不同的程序(即kernel)。Grid、block和thread的关系如下图所示:anna矩形每个thread都有自己的一份register和localmemory的空间。同一个block中的每个thread则有共享的一份sharememory。此外所有的thread(包括不同block的thread)都共享一份globalmemory、constantmemory、和texturememory。不同的grid则有各自的globalmemory、constantmemory和texturememory。这些不同的内存的差别会在之后讨论。执行模式由于显示芯片大量并行计算的特性它处理一些问题的方式和一般CPU是不同的。主要的特点包括:内存存取latency的问题:CPU通常使用cache来减少存取主内存的次数以避免内存latency影响到执行效率。显示芯片则多半没有cache(或很小)而利用并行化执行的方式来隐藏内存的latency(即当第一个thread需要等待内存读取结果时则开始执行第二个thread依此类推)。分支指令的问题:CPU通常利用分支预测等方式来减少分支指令造成的pipelinebubble。显示芯片则多半使用类似处理内存latency的方式。不过通常显示芯片处理分支的效率会比较差。因此最适合利用CUDA处理的问题是可以大量并行化的问题才能有效隐藏内存的latency并有效利用显示芯片上的大量执行单元。使用CUDA时同时有上千个thread在执行是很正常的。因此如果不能大量并行化的问题使用CUDA就没办法达到最好的效率了。CUDAToolkit的安装目前NVIDIA提供的CUDAToolkit(可从这里下载)支持Windows(bits及bits版本)及许多不同的Linux版本。CUDAToolkit需要配合CCcompiler。在Windows下目前只支持VisualStudiox及VisualStudio(包括免费的VisualStudioCExpress)。VisualStudio和gcc在Windows下是不支援的。在Linux下则只支援gcc。这里简单介绍一下在Windows下设定并使用CUDA的方式。下载及安装在Windows下CUDAToolkit和CUDASDK都是由安装程序的形式安装的。CUDAToolkit包括CUDA的基本工具而CUDASDK则包括许多范例程序以及链接库。基本上要写CUDA的程序只需要安装CUDAToolkit即可。不过CUDASDK仍值得安装因为里面的许多范例程序和链接库都相当有用。CUDAToolkit安装完后预设会安装在C:CUDA目录里。其中包括几个目录:•bin工具程序及动态链接库•doc文件•includeheader檔•lib链接库档案•open基于Open的CUDAcompiler•src一些原始码安装程序也会设定一些环境变量包括:•CUDABINPATH工具程序的目录默认为C:CUDAbin•CUDAINCPATHheader文件的目录默认为C:CUDAinc•CUDALIBPATH链接库文件的目录默认为C:CUDAlib在VisualStudio中使用CUDACUDA的主要工具是nvcc它会执行所需要的程序将CUDA程序代码编译成执行档(或object檔)。在VisualStudio下我们透过设定custombuildtool的方式让VisualStudio会自动执行nvcc。这里以VisualStudio为例:首先建立一个WinConsole模式的project(在ApplicationSettings中记得勾选Emptyproject)并新增一个档案例如maincu。在maincu上右键单击并选择Properties。点选General确定Tool的部份是选择CustomBuildTool。选择CustomBuildStep在CommandLine使用以下设定:oRelease模式:"$(CUDABINPATH)nvccexe"ccbin"$(VCInstallDir)bin"cDWINDCONSOLEDMBCSXcompilerEHsc,W,nologo,Wp,O,Zi,MTI"$(CUDAINCPATH)"o$(ConfigurationName)$(InputName)obj$(InpileName)oDebug模式:"$(CUDABINPATH)nvccexe"ccbin"$(VCInstallDir)bin"cDDEBUGDWINDCONSOLEDMBCSXcompilerEHsc,W,nologo,Wp,Od,Zi,RTC,MTdI"$(CUDAINCPATH)"o$(ConfigurationName)$(InputName)obj$(InpileName)如果想要使用软件仿真的模式可以新增两个额外的设定:oEmuRelease模式:"$(CUDABINPATH)nvccexe"ccbin"$(VCInstallDir)bin"deviceemucDWINDCONSOLEDMBCSXcompilerEHsc,W,nologo,Wp,O,Zi,MTI"$(CUDAINCPATH)"o$(ConfigurationName)$(InputName)obj$(InpileName)oEmuDebug模式:"$(CUDABINPATH)nvccexe"ccbin"$(VCInstallDir)bin"deviceemucDDEBUGDWINDCONSOLEDMBCSXcompilerEHsc,W,nologo,Wp,Od,Zi,RTC,MTdI"$(CUDAINCPATH)"o$(ConfigurationName)$(InputName)obj$(InpileName)对所有的配置文件在CustomBuildStep的Outputs中加入$(ConfigurationName)$(InputName)obj。选择project右键单击选择Properties再点选Linker。对所有的配置文件修改以下设定:oGeneralEnableIncrementalLinking:NooGeneralAdditionalLibraryDirectories:$(CUDALIBPATH)oInputAdditionalDependencies:cudartlib这样应该就可以直接在VisualStudio的IDE中编辑CUDA程序后直接build以及执行程序了。第一个CUDA程序CUDA目前有两种不同的API:RuntimeAPI和DriverAPI两种API各有其适用的范围。由于runtimeAPI较容易使用一开始我们会以runetimeAPI为主。CUDA的初始化首先先建立一个档案firstcudacu。如果是使用VisualStudio的话则请先按照这里的设定方式设定project。要使用runtimeAPI的时候需要includecudaruntimeh。所以在程序的最前面加上#include<stdioh>#include<cudaruntimeh>接下来是一个InitCUDA函式会呼叫runtimeAPI中有关初始化CUDA的功能:boolInitCUDA(){intcountcudaGetDeviceCount(count)if(count==){fprintf(stderr,"Thereisnodevicen")returnfalse}intifor(i=i<counti){cudaDeviceProppropif(cudaGetDeviceProperties(prop,i)==cudaSuccess){if(propmajor>=){break}}}if(i==count){fprintf(stderr,"ThereisnodevicesupportingCUDAxn")returnfalse}cudaSetDevice(i)returntrue}这个函式会先呼叫cudaGetDeviceCount函式取得支持CUDA的装置的数目。如果系统上没有支持CUDA的装置则它会传回而device会是一个仿真的装置但不支持CUDA以上的功能。所以要确定系统上是否有支持CUDA的装置需要对每个device呼叫cudaGetDeviceProperties函式取得装置的各项数据并判断装置支持的CUDA版本(propmajor和propminor分别代表装置支持的版本号码例如则propmajor为而propminor为)。透过cudaGetDeviceProperties函式可以取得许多数据除了装置支持的CUDA版本之外还有装置的名称、内存的大小、最大的thread数目、执行单元的频率等等。详情可参考NVIDIA的CUDAProgrammingGuide。在找到支持CUDA以上的装置之后就可以呼叫cudaSetDevice函式把它设为目前要使用的装置。最后是main函式。在main函式中我们直接呼叫刚才的InitCUDA函式并显示适当的讯息:anna线条intmain(){if(!InitCUDA()){return}printf("CUDAinitializedn")return}这样就可以利用nvcc来compile这个程序了。使用VisualStudio的话若按照先前的设定方式可以直接BuildProject并执行。nvcc是CUDA的compile工具它会将cu檔拆解出在GPU上执行的部份及在host上执行的部份并呼叫适当的程序进行compile动作。在GPU执行的部份会透过NVIDIA提供的compiler编译成中介码而host执行的部份则会透过系统上的Ccompiler编译(在Windows上使用VisualC而在Linux上使用gcc)。编译后的程序执行时如果系统上有支持CUDA的装置应该会显示CUDAinitialized的讯息否则会显示相关的错误讯息。利用CUDA进行运算到目前为止我们的程序并没有做什么有用的工作。所以现在我们加入一个简单的动作就是把一大堆数字计算出它的平方和。首先把程序最前面的include部份改成:#include<stdioh>#include<stdlibh>#include<cudaruntimeh>#defineDATASIZEintdataDATASIZE并加入一个新函式GenerateNumbers:voidGenerateNumbers(int*number,intsize){for(inti=i<sizei){numberi=rand()}}这个函式会产生一大堆~之间的随机数。要利用CUDA进行计算之前要先把数据复制到显卡内存中才能让显示芯片使用。因此需要取得一块适当大小的显卡内存再把产生好的数据复制进去。在main函式中加入:GenerateNumbers(data,DATASIZE)int*gpudata,*resultcudaMalloc((void**)gpudata,sizeof(int)*DATASIZE)cudaMalloc((void**)result,sizeof(int))cudaMemcpy(gpudata,data,sizeof(int)*DATASIZE,cudaMemcpyHostToDevice)上面这段程序会先呼叫GenerateNumbers产生随机数并呼叫cudaMalloc取得一块显卡内存(result则是用来存取计算结果在稍后会用到)并透过cudaMemcpy将产生的随机数复制到显卡内存中。cudaMalloc和cudaMemcpy的用法和一般的malloc及memcpy类似不过cudaMemcpy则多出一个参数指示复制内存的方向。在这里因为是从主内存复制到显卡内存所以使用cudaMemcpyHostToDevice。如果是从显卡内存到主内存则使用cudaMemcpyDeviceToHost。这在之后会用到。接下来是要写在显示芯片上执行的程序。在CUDA中在函式前面加上global表示这个函式是要在显示芯片上执行的。因此加入以下的函式:globalstaticvoidsumOfSquares(int*num,int*result){intsum=intifor(i=i<DATASIZEi){sum=numi*numi}*result=sum}在显示芯片上执行的程序有一些限制例如它不能有传回值。其它的限制会在之后提到。接下来是要让CUDA执行这个函式。在CUDA中要执行一个函式使用以下的语法:函式名称<<<block数目,thread数目,sharedmemory大小>>>(参数)呼叫完后还要把结果从显示芯片复制回主内存上。在main函式中加入以下的程序:sumOfSquares<<<,,>>>(gpudata,result)intsumcudaMemcpy(sum,result,sizeof(int),cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)printf("sum:dn",sum)因为这个程序只使用一个thread所以block数目、thread数目都是。我们也没有使用到任何sharedmemory所以设为。编译后执行应该可以看到执行的结果。为了确定执行的结果正确我们可以加上一段以CPU执行的程序代码来验证结果:sum=for(inti=i<DATASIZEi){sum=datai*datai}printf("sum(CPU):dn",sum)编译后执行确认两个结果相同。计算运行时间CUDA提供了一个clock函式可以取得目前的timestamp很适合用来判断一段程序执行所花费的时间(单位为GPU执行单元的频率)。这对程序的优化也相当有用。要在我们的程序中记录时间把sumOfSquares函式改成:globalstaticvoidsumOfSquares(int*num,int*result,clockt*time){intsum=inticlocktstart=clock()for(i=i<DATASIZEi){sum=numi*numi}*result=sum*time=clock()start}把main函式中间部份改成:int*gpudata,*resultclockt*timecudaMalloc((void**)gpudata,sizeof(int)*DATASIZE)cudaMalloc((void**)result,sizeof(int))cudaMalloc((void**)time,sizeof(clockt))cudaMemcpy(gpudata,data,sizeof(int)*DATASIZE,cudaMemcpyHostToDevice)sumOfSquares<<<,,>>>(gpudata,result,time)intsumclockttimeusedcudaMemcpy(sum,result,sizeof(int),cudaMemcpyDeviceToHost)cudaMemcpy(timeused,time,sizeof(clockt),cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)printf("sum:dtime:dn",sum,timeused)编译后执行就可以看到执行所花费的时间了。如果计算实际运行时间的话可能会注意到它的执行效率并不好。这是因为我们的程序并没有利用到CUDA的主要的优势即并行化执行。在下一段文章中会讨论如何进行优化的动作。改良第一个CUDA程序在上一篇文章中我们做了一个计算一大堆数字的平方和的程序。不过我们也提到这个程序的执行效率并不理想。当然实际上来说如果只是要做计算平方和的动作用CPU做会比用GPU快得多。这是因为平方和的计算并不需要太多运算能力所以几乎都是被内存带宽所限制。因此光是把数据复制到显卡内存上的这个动作所需要的时间可能已经和直接在CPU上进行计算差不多了。不过如果进行平方和的计算只是一个更复杂的计算过程的一部份的话那么当然在GPU上计算还是有它的好处的。而且如果数据已经在显卡内存上(例如在GPU上透过某种算法产生)那么使用GPU进行这样的运算还是会比较快的。刚才也提到了由于这个计算的主要瓶颈是内存带宽所以理论上显卡的内存带宽是相当大的。这里我们就来看看倒底我们的第一个程序能利用到多少内存带宽。程序的并行化我们的第一个程序并没有利用到任何并行化的功能。整个程序只有一个thread。在GeForceGT上面在GPU上执行的部份(称为"kernel")大约花费M个频率。GeForceGT的执行单元的频率是GHz因此这表示它花费了约秒的时间。M个bits数字的数据量是MB因此这个程序实际上使用的内存带宽只有MBs左右!这是非常糟糕的表现。为什么会有这样差的表现呢?这是因为GPU的架构特性所造成的。在CUDA中一般的数据复制到的显卡内存的部份称为globalmemory。这些内存是没有cache的而且存取globalmemory所需要的时间(即latency)是非常长的通常是数百个cycles。由于我们的程序只有一个thread所以每次它读取globalmemory的内容就要等到实际读取到数据、累加到sum之后才能进行下一步。这就是为什么它的表现会这么的差。由于globalmemory并没有cache所以要避开巨大的latency的方法就是要利用大量的threads。假设现在有大量的threads在同时执行那么当一个thread读取内存开始等待结果的时候GPU就可以立刻切换到下一个thread并读取下一个内存位置。因此理想上当thread的数目够多的时候就可以完全把globalmemory的巨大latency隐藏起来了。要怎么把计算平方和的程序并行化呢?最简单的方法似乎就是把数字分成若干组把各组数字分别计算平方和后最后再把每组的和加总起来就可以了。一开始我们可以把最后加总的动作由CPU来进行。首先在firstcudacu中在#defineDATASIZE的后面增加一个#define设定thread的数目:#defineDATASIZE#defineTHREADNUM接着把kernel程序改成:globalstaticvoidsumOfSquares(int*num,int*result,clockt*time){constinttid=threadIdxxconstintsize=DATASIZETHREADNUMintsum=inticlocktstartif(tid==)start=clock()for(i=tid*sizei<(tid)*sizei){sum=numi*numi}resulttid=sumif(tid==)*time=clock()start}程序里的threadIdx是CUDA的一个内建的变量表示目前的thread是第几个thread(由开始计算)。以我们的例子来说会有个threads所以同时会有个sumOfSquares函式在执行但每一个的threadIdxx则分别会是~。利用这个变量我们就可以让每一份函式执行时对整个数据不同的部份计算平方和。另外我们也让计算时间的动作只在thread(即threadIdxx=的时候)进行。同样的由于会有个计算结果所以原来存放result的内存位置也要扩大。把main函式中的中间部份改成:int*gpudata,*resultclockt*timecudaMalloc((void**)gpudata,sizeof(int)*DATASIZE)cudaMalloc((void**)result,sizeof(int)*THREADNUM)cudaMalloc((void**)time,sizeof(clockt))cudaMemcpy(gpudata,data,sizeof(int)*DATASIZE,cudaMemcpyHostToDevice)sumOfSquares<<<,THREADNUM,>>>(gpudata,result,time)intsumTHREADNUMclockttimeusedcudaMemcpy(sum,result,sizeof(int)*THREADNUM,cudaMemcpyDeviceToHost)cudaMemcpy(timeused,time,sizeof(clockt),cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)cudaFree(time)可以注意到我们在呼叫sumOfSquares函式时指定THREADNUM为thread的数目。最后在CPU端把计算好的各组数据的平方和进行加总:intfinalsum=for(inti=i<THREADNUMi){finalsum=sumi}printf("sum:dtime:dn",finalsum,timeused)finalsum=for(inti=i<DATASIZEi){sum=datai*datai}printf("sum(CPU):dn",finalsum)编译后执行确认结果和原来相同。这个版本的程序在GeForceGT上执行只需要约Mcycles比前一版程序快了倍!这就是透过大量thread来隐藏latency所带来的效果。不过如果计算一下它使用的内存带宽就会发现其实仍不是很理想大约只有MBs而已。这和GeForceGT所具有的内存带宽是很大的差距。为什么会这样呢?内存的存取模式显卡上的内存是DRAM因此最有效率的存取方式是以连续的方式存取。前面的程序虽然看起来是连续存取内存位置(每个thread对一块连续的数字计算平方和)但是我们要考虑到实际上thread的执行方式。前面提过当一个thread在等待内存的数据时GPU会切换到下一个thread。也就是说实际上执行的顺序是类似thread>thread>thread>因此在同一个thread中连续存取内存在实际执行时反而不是连续了。要让实际执行结果是连续的存取我们应该要让thread读取第一个数字thread读取第二个数字…依此类推。所以我们可以把kernel程序改成如下:globalstaticvoidsumOfSquares(int*num,int*result,clockt*time){constinttid=threadIdxxintsum=inticlocktstartif(tid==)start=clock()for(i=tidi<DATASIZEi=THREADNUM){sum=numi*numi}resulttid=sumif(tid==)*time=clock()start}编译后执行确认结果相同。仅仅是这样简单的修改实际执行的效率就有很大的差别。在GeForceGT上上面的程序执行需要的频率是Mcycles又比前一版程序快了三倍。不过这样仍只有GBs的带宽而已。这是因为我们使用的thread数目还是不够多的原因。理论上个threads最多只能隐藏cycles的latency。但是GPU存取globalmemory时的latency可能高达cycles以上。如果增加thread数目就可以看到更好的效率。例如可以把THREADNUM改成。在GeForceGT上这可以让执行花费的时间减少到Mcycles。有些改进但是仍不够大。不幸的是目前GeForceGT一个block最多只能有个threads所以不能再增加了而且如果thread数目增加太多那么在CPU端要做的最后加总工作也会变多。更多的并行化前面提到了block。在之前介绍呼叫CUDA函式时也有提到"block数目"这个参数。到目前为止我们都只使用一个block。究竟block是什么呢?在CUDA中thread是可以分组的也就是block。一个block中的thread具有一个共享的sharedmemory也可以进行同步工作。不同block之间的thread则不行。在我们的程序中其实不太需要进行thread的同步动作因此我们可以使用多个block来进一步增加thread的数目。首先在#defineDATASIZE的地方改成如下:#defineDATASIZE#defineBLOCKNUM#defineTHREADNUM这表示我们会建立个blocks每个blocks有个threads总共有*=个threads。接着我们把kernel部份改成:globalstaticvoidsumOfSquares(int*num,int*result,clockt*time){constinttid=threadIdxxconstintbid=blockIdxxintsum=intiif(tid==)timebid=clock()for(i=bid*THREADNUMtidi<DATASIZEi=BLOCKNUM*THREADNUM){sum=numi*numi}resultbid*THREADNUMtid=sumif(tid==)timebidBLOCKNUM=clock()}blockIdxx和threadIdxx一样是CUDA内建的变量它表示的是目前的block编号。另外注意到我们把计算时间的方式改成每个block都会记录开始时间及结束时间。main函式部份修改成:int*gpudata,*resultclockt*timecudaMalloc((void**)gpudata,sizeof(int)*DATASIZE)cudaMalloc((void**)result,sizeof(int)*THREADNUM*BLOCKNUM)cudaMalloc((void**)time,sizeof(clockt)*BLOCKNUM*)cudaMemcpy(gpudata,data,sizeof(int)*DATASIZE,cudaMemcpyHostToDevice)sumOfSquares<<<BLOCKNUM,THREADNUM,>>>(gpudata,result,time)intsumTHREADNUM*BLOCKNUMclockttimeusedBLOCKNUM*cudaMemcpy(sum,result,sizeof(int)*THREADNUM*BLOCKNUM,cudaMemcpyDeviceToHost)cudaMemcpy(timeused,time,sizeof(clockt)*BLOCKNUM*,cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)cudaFree(time)intfinalsum=for(inti=i<THREADNUM*BLOCKNUMi){finalsum=sumi}clocktminstart,maxendminstart=timeusedmaxend=timeusedBLOCKNUMfor(inti=i<BLOCKNUMi){if(minstart>timeusedi)minstart=timeusediif(maxend<timeusediBLOCKNUM)maxend=timeusediBLOCKNUM}printf("sum:dtime:dn",finalsum,maxendminstart)基本上我们只是把result的大小变大并修改计算时间的方式把每个block最早的开始时间和最晚的结束时间相减取得总运行时间。这个版本的程序执行的时间减少很多在GeForceGT上只需要约Kcycles相当于GBs左右的带宽。不过它在CPU上执行的部份需要的时间加长了(因为CPU现在需要加总个数字)。为了避免这个问题我们可以让每个block把自己的每个thread的计算结果进行加总。Thread的同步前面提过一个block内的thread可以有共享的内存也可以进行同步。我们可以利用这一点让每个block内的所有thread把自己计算的结果加总起来。把kernel改成如下:globalstaticvoidsumOfSquares(int*num,int*result,clockt*time){externsharedintsharedconstinttid=threadIdxxconstintbid=blockIdxxintiif(tid==)timebid=clock()sharedtid=for(i=bid*THREADNUMtidi<DATASIZEi=BLOCKNUM*THREADNUM){sharedtid=numi*numi}syncthreads()if(tid==){for(i=i<THREADNUMi){shared=sharedi}resultbid=shared}if(tid==)timebidBLOCKNUM=clock()}利用shared声明的变量表示这是sharedmemory是一个block中每个thread都共享的内存。它会使用在GPU上的内存所以存取的速度相当快不需要担心latency的问题。syncthreads()是一个CUDA的内部函数表示block中所有的thread都要同步到这个点才能继续执行。在我们的例子中由于之后要把所有thread计算的结果进行加总所以我们需要确定每个thread都已经把结果写到sharedtid里面了。接下来把main函式的一部份改成:int*gpudata,*resultclockt*timecudaMalloc((void**)gpudata,sizeof(int)*DATASIZE)cudaMalloc((void**)result,sizeof(int)*BLOCKNUM)cudaMalloc((void**)time,sizeof(clockt)*BLOCKNUM*)cudaMemcpy(gpudata,data,sizeof(int)*DATASIZE,cudaMemcpyHostToDevice)sumOfSquares<<<BLOCKNUM,THREADNUM,THREADNUM*sizeof(int)>>>(gpudata,result,time)intsumBLOCKNUMclockttimeusedBLOCKNUM*cudaMemcpy(sum,result,sizeof(int)*BLOCKNUM,cudaMemcpyDeviceToHost)cudaMemcpy(timeused,time,sizeof(clockt)*BLOCKNUM*,cudaMemcpyDeviceToHost)cudaFree(gpudata)cudaFree(result)cudaFree(time)intfinalsum=for(in

用户评论(0)

0/200

精彩专题

上传我的资料

每篇奖励 +2积分

资料评价:

/31
1下载券 下载 加入VIP, 送下载券

意见
反馈

立即扫码关注

爱问共享资料微信公众号

返回
顶部