深入浅出的谈CUDA.docx
《深入浅出的谈CUDA.docx》由会员分享,可在线阅读,更多相关《深入浅出的谈CUDA.docx(28页珍藏版)》请在冰豆网上搜索。
深入浅出的谈CUDA
深入浅出谈cuda(转载)
猫猫观点:
cuda即显卡加速程序端口,他的优点就是用c语言作为基础,这点很符合中国的现状,中国多数电脑编程高手都懂得c语言的意思,这也是cuda这个接口伟大的地方,但是由于目前显卡gpgpu的先天的不足,目前还没有功能模块化线程增强指令,所以,提高利用gpgpu的利用率变成了各个编程高手绞尽脑汁的问题,此外windows目前有很多程序并不能利用到gpgpu的处理能力而实现加速,因此猫猫的另一个观点就是如何无缝连接windows程序所能调用的资源,让整个windows的系统环境都能得益于gpgpu的加速能力而获得整体的效率提升,这时候gpgpu就已经普及了!
希望这个补丁能尽早出来,让windows环境用户受益!
“CUDA是NVIDIA的GPGPU模型,它使用C语言为基础,可以直接以大多数人熟悉的C语言,写出在显示芯片上执行的程序,而不需要去学习特定的显示芯片的指令或是特殊的结构。
”
CUDA是什么?
能吃吗?
编者注:
NVIDIA的GeFoce8800GTX发布后,它的通用计算架构CUDA经过一年多的推广后,现在已经在有相当多的论文发表,在商业应用软件等方面也初步出现了视频编解码、金融、地质勘探、科学计算等领域的产品,是时候让我们对其作更深一步的了解。
为了让大家更容易了解CUDA,我们征得Hotball的本人同意,发表他最近亲自撰写的本文。
这篇文章的特点是深入浅出,也包含了hotball本人编写一些简单CUDA程序的亲身体验,对于希望了解CUDA的读者来说是非常不错的入门文章,PCINLIFE对本文的发表没有作任何的删减,主要是把一些台湾的词汇转换成大陆的词汇以及作了若干"编者注"的注释。
现代的显示芯片已经具有高度的可程序化能力,由于显示芯片通常具有相当高的内存带宽,以及大量的执行单元,因此开始有利用显示芯片来帮助进行一些计算工作的想法,即GPGPU。
CUDA即是NVIDIA的GPGPU模型。
NVIDIA的新一代显示芯片,包括GeForce8系列及更新的显示芯片都支持CUDA。
NVIDIA免费提供CUDA的开发工具(包括Windows版本和Linux版本)、程序范例、文件等等,可以在CUDAZone下载。
GPGPU的优缺点
使用显示芯片来进行运算工作,和使用CPU相比,主要有几个好处:
1.显示芯片通常具有更大的内存带宽。
例如,NVIDIA的GeForce8800GTX具有超过50GB/s的内存带宽,而目前高阶CPU的内存带宽则在10GB/s左右。
2.显示芯片具有更大量的执行单元。
例如GeForce8800GTX具有128个"streamprocessors",频率为1.35GHz。
CPU频率通常较高,但是执行单元的数目则要少得多。
3.和高阶CPU相比,显卡的价格较为低廉。
例如目前一张GeForce8800GT包括512MB内存的价格,和一颗2.4GHz四核心CPU的价格相若。
当然,使用显示芯片也有它的一些缺点:
1.显示芯片的运算单元数量很多,因此对于不能高度并行化的工作,所能带来的帮助就不大。
2.显示芯片目前通常只支持32bits浮点数,且多半不能完全支持IEEE754规格,有些运算的精确度可能较低。
目前许多显示芯片并没有分开的整数运算单元,因此整数运算的效率较差。
3.显示芯片通常不具有分支预测等复杂的流程控制单元,因此对于具有高度分支的程序,效率会比较差。
4.目前GPGPU的程序模型仍不成熟,也还没有公认的标准。
例如NVIDIA和AMD/ATI就有各自不同的程序模型。
整体来说,显示芯片的性质类似streamprocessor,适合一次进行大量相同的工作。
CPU则比较有弹性,能同时进行变化较多的工作。
CUDA架构
CUDA是NVIDIA的GPGPU模型,它使用C语言为基础,可以直接以大多数人熟悉的C语言,写出在显示芯片上执行的程序,而不需要去学习特定的显示芯片的指令或是特殊的结构。
在CUDA的架构下,一个程序分为两个部份:
host端和device端。
Host端是指在CPU上执行的部份,而device端则是在显示芯片上执行的部份。
Device端的程序又称为"kernel"。
通常host端程序会将数据准备好后,复制到显卡的内存中,再由显示芯片执行device端程序,完成后再由host端程序将结果从显卡的内存中取回。
由于CPU存取显卡内存时只能透过PCIExpress接口,因此速度较慢(PCIExpressx16的理论带宽是双向各4GB/s),因此不能太常进行这类动作,以免降低效率。
在CUDA架构下,显示芯片执行时的最小单位是thread。
数个thread可以组成一个block。
一个block中的thread能存取同一块共享的内存,而且可以快速进行同步的动作。
每一个block所能包含的thread数目是有限的。
不过,执行相同程序的block,可以组成grid。
不同block中的thread无法存取同一个共享的内存,因此无法直接互通或进行同步。
因此,不同block中的thread能合作的程度是比较低的。
不过,利用这个模式,可以让程序不用担心显示芯片实际上能同时执行的thread数目限制。
例如,一个具有很少量执行单元的显示芯片,可能会把各个block中的thread顺序执行,而非同时执行。
不同的grid则可以执行不同的程序(即kernel)。
Grid、block和thread的关系,如下图所示:
每个thread都有自己的一份register和localmemory的空间。
同一个block中的每个thread则有共享的一份sharememory。
此外,所有的thread(包括不同block的thread)都共享一份globalmemory、constantmemory、和texturememory。
不同的grid则有各自的globalmemory、constantmemory和texturememory。
这些不同的内存的差别,会在之后讨论。
执行模式
由于显示芯片大量并行计算的特性,它处理一些问题的方式,和一般CPU是不同的。
主要的特点包括:
1.内存存取latency的问题:
CPU通常使用cache来减少存取主内存的次数,以避免内存latency影响到执行效率。
显示芯片则多半没有cache(或很小),而利用并行化执行的方式来隐藏内存的latency(即,当第一个thread需要等待内存读取结果时,则开始执行第二个thread,依此类推)。
2.分支指令的问题:
CPU通常利用分支预测等方式来减少分支指令造成的pipelinebubble。
显示芯片则多半使用类似处理内存latency的方式。
不过,通常显示芯片处理分支的效率会比较差。
因此,最适合利用CUDA处理的问题,是可以大量并行化的问题,才能有效隐藏内存的latency,并有效利用显示芯片上的大量执行单元。
使用CUDA时,同时有上千个thread在执行是很正常的。
因此,如果不能大量并行化的问题,使用CUDA就没办法达到最好的效率了。
CUDAToolkit的安装
目前NVIDIA提供的CUDAToolkit(可从这里下载)支持Windows(32bits及64bits版本)及许多不同的Linux版本。
CUDAToolkit需要配合C/C++compiler。
在Windows下,目前只支持VisualStudio7.x及VisualStudio8(包括免费的VisualStudioC++2005Express)。
VisualStudio6和gcc在Windows下是不支援的。
在Linux下则只支援gcc。
这里简单介绍一下在Windows下设定并使用CUDA的方式。
下载及安装
在Windows下,CUDAToolkit和CUDASDK都是由安装程序的形式安装的。
CUDAToolkit包括CUDA的基本工具,而CUDASDK则包括许多范例程序以及链接库。
基本上要写CUDA的程序,只需要安装CUDAToolkit即可。
不过CUDASDK仍值得安装,因为里面的许多范例程序和链接库都相当有用。
CUDAToolkit安装完后,预设会安装在C:
\CUDA目录里。
其中包括几个目录:
∙bin--工具程序及动态链接库
∙doc--文件
∙include--header檔
∙lib--链接库档案
∙open64--基于Open64的CUDAcompiler
∙src--一些原始码
安装程序也会设定一些环境变量,包括:
∙CUDA_BIN_PATH--工具程序的目录,默认为C:
\CUDA\bin
∙CUDA_INC_PATH--header文件的目录,默认为C:
\CUDA\inc
∙CUDA_LIB_PATH--链接库文件的目录,默认为C:
\CUDA\lib
在VisualStudio中使用CUDA
CUDA的主要工具是nvcc,它会执行所需要的程序,将CUDA程序代码编译成执行档(或object檔)。
在VisualStudio下,我们透过设定custombuildtool的方式,让VisualStudio会自动执行nvcc。
这里以VisualStudio2005为例:
1.首先,建立一个Win32Console模式的project(在ApplicationSettings中记得勾选Emptyproject),并新增一个档案,例如main.cu。
2.在main.cu上右键单击,并选择Properties。
点选General,确定Tool的部份是选择CustomBuildTool。
3.选择CustomBuildStep,在CommandLine使用以下设定:
oRelease模式:
"$(CUDA_BIN_PATH)\nvcc.exe"-ccbin"$(VCInstallDir)bin"-c-DWIN32-D_CONSOLE-D_MBCS-Xcompiler/EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT-I"$(CUDA_INC_PATH)"-o$(ConfigurationName)\$(InputName).obj$(InputFileName)
oDebug模式:
"$(CUDA_BIN_PATH)\nvcc.exe"-ccbin"$(VCInstallDir)bin"-c-D_DEBUG-DWIN32-D_CONSOLE-D_MBCS-Xcompiler/EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd-I"$(CUDA_INC_PATH)"-o$(ConfigurationName)\$(InputName).obj$(InputFileName)
4.如果想要使用软件仿真的模式,可以新增两个额外的设定:
oEmuRelease模式:
"$(CUDA_BIN_PATH)\nvcc.exe"-ccbin"$(VCInstallDir)bin"-deviceemu-c-DWIN32-D_CONSOLE-D_MBCS-Xcompiler/EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT-I"$(CUDA_INC_PATH)"-o$(ConfigurationName)\$(InputName).obj$(InputFileName)
oEmuDebug模式:
"$(CUDA_BIN_PATH)\nvcc.exe"-ccbin"$(VCInstallDir)bin"-deviceemu-c-D_DEBUG-DWIN32-D_CONSOLE-D_MBCS-Xcompiler/EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd-I"$(CUDA_INC_PATH)"-o$(ConfigurationName)\$(InputName).obj$(InputFileName)
5.对所有的配置文件,在CustomBuildStep的Outputs中加入$(ConfigurationName)\$(InputName).obj。
6.选择project,右键单击选择Properties,再点选Linker。
对所有的配置文件修改以下设定:
oGeneral/EnableIncrementalLinking:
No
oGeneral/AdditionalLibraryDirectories:
$(CUDA_LIB_PATH)
oInput/AdditionalDependencies:
cudart.lib
这样应该就可以直接在VisualStudio的IDE中,编辑CUDA程序后,直接build以及执行程序了。
“CUDA是NVIDIA的GPGPU模型,它使用C语言为基础,可以直接以大多数人熟悉的C语言,写出在显示芯片上执行的程序,而不需要去学习特定的显示芯片的指令或是特殊的结构。
”
第一个CUDA程序
CUDA目前有两种不同的API:
RuntimeAPI和DriverAPI,两种API各有其适用的范围。
由于runtimeAPI较容易使用,一开始我们会以runetimeAPI为主。
CUDA的初始化
首先,先建立一个档案first_cuda.cu。
如果是使用VisualStudio的话,则请先按照这里的设定方式设定project。
要使用runtimeAPI的时候,需要includecuda_runtime.h。
所以,在程序的最前面,加上
#include
#include
接下来是一个InitCUDA函式,会呼叫runtimeAPI中,有关初始化CUDA的功能:
boolInitCUDA()
{
intcount;
cudaGetDeviceCount(&count);
if(count==0){
fprintf(stderr,"Thereisnodevice.\n");
returnfalse;
}
inti;
for(i=0;i cudaDevicePropprop;
if(cudaGetDeviceProperties(&prop,i)==cudaSuccess){
if(prop.major>=1){
break;
}
}
}
if(i==count){
fprintf(stderr,"ThereisnodevicesupportingCUDA1.x.\n");
returnfalse;
}
cudaSetDevice(i);
returntrue;
}
这个函式会先呼叫cudaGetDeviceCount函式,取得支持CUDA的装置的数目。
如果系统上没有支持CUDA的装置,则它会传回1,而device0会是一个仿真的装置,但不支持CUDA1.0以上的功能。
所以,要确定系统上是否有支持CUDA的装置,需要对每个device呼叫cudaGetDeviceProperties函式,取得装置的各项数据,并判断装置支持的CUDA版本(prop.major和prop.minor分别代表装置支持的版本号码,例如1.0则prop.major为1而prop.minor为0)。
透过cudaGetDeviceProperties函式可以取得许多数据,除了装置支持的CUDA版本之外,还有装置的名称、内存的大小、最大的thread数目、执行单元的频率等等。
详情可参考NVIDIA的CUDAProgrammingGuide。
在找到支持CUDA1.0以上的装置之后,就可以呼叫cudaSetDevice函式,把它设为目前要使用的装置。
最后是main函式。
在main函式中我们直接呼叫刚才的InitCUDA函式,并显示适当的讯息:
intmain()
{
if(!
InitCUDA()){
return0;
}
printf("CUDAinitialized.\n");
return0;
}
这样就可以利用nvcc来compile这个程序了。
使用VisualStudio的话,若按照先前的设定方式,可以直接BuildProject并执行。
nvcc是CUDA的compile工具,它会将.cu檔拆解出在GPU上执行的部份,及在host上执行的部份,并呼叫适当的程序进行compile动作。
在GPU执行的部份会透过NVIDIA提供的compiler编译成中介码,而host执行的部份则会透过系统上的C++compiler编译(在Windows上使用VisualC++而在Linux上使用gcc)。
编译后的程序,执行时如果系统上有支持CUDA的装置,应该会显示CUDAinitialized.的讯息,否则会显示相关的错误讯息。
利用CUDA进行运算
到目前为止,我们的程序并没有做什么有用的工作。
所以,现在我们加入一个简单的动作,就是把一大堆数字,计算出它的平方和。
首先,把程序最前面的include部份改成:
#include
#include
#include
#defineDATA_SIZE1048576
intdata[DATA_SIZE];
并加入一个新函式GenerateNumbers:
voidGenerateNumbers(int*number,intsize)
{
for(inti=0;i number[i]=rand()%10;
}
}
这个函式会产生一大堆0~9之间的随机数。
要利用CUDA进行计算之前,要先把数据复制到显卡内存中,才能让显示芯片使用。
因此,需要取得一块适当大小的显卡内存,再把产生好的数据复制进去。
在main函式中加入:
GenerateNumbers(data,DATA_SIZE);
int*gpudata,*result;
cudaMalloc((void**)&gpudata,sizeof(int)*DATA_SIZE);
cudaMalloc((void**)&result,sizeof(int));
cudaMemcpy(gpudata,data,sizeof(int)*DATA_SIZE,
cudaMemcpyHostToDevice);
上面这段程序会先呼叫GenerateNumbers产生随机数,并呼叫cudaMalloc取得一块显卡内存(result则是用来存取计算结果,在稍后会用到),并透过cudaMemcpy将产生的随机数复制到显卡内存中。
cudaMalloc和cudaMemcpy的用法和一般的malloc及memcpy类似,不过cudaMemcpy则多出一个参数,指示复制内存的方向。
在这里因为是从主内存复制到显卡内存,所以使用cudaMemcpyHostToDevice。
如果是从显卡内存到主内存,则使用cudaMemcpyDeviceToHost。
这在之后会用到。
接下来是要写在显示芯片上执行的程序。
在CUDA中,在函式前面加上__global__表示这个函式是要在显示芯片上执行的。
因此,加入以下的函式:
__global__staticvoidsumOfSquares(int*num,int*result)
{
intsum=0;
inti;
for(i=0;i sum+=num[i]*num[i];
}
*result=sum;
}
在显示芯片上执行的程序有一些限制,例如它不能有传回值。
其它的限制会在之后提到。
接下来是要让CUDA执行这个函式。
在CUDA中,要执行一个函式,使用以下的语法:
函式名称<<>>(参数...);
呼叫完后,还要把结果从显示芯片复制回主内存上。
在main函式中加入以下的程序:
sumOfSquares<<<1,1,0>>>(gpudata,result);
intsum;
cudaMemcpy(&sum,result,sizeof(int),cudaMemcpyDeviceToHost);
cudaFree(gpudata);
cudaFree(result);
printf("sum:
%d\n",sum);
因为这个程序只使用一个thread,所以block数目、thread数目都是1。
我们也没有使用到任何sharedmemory,所以设为0。
编译后执行,应该可以看到执行的结果。
为了确定执行的结果正确,我们可以加上一段以CPU执行的程序代码,来验证结果:
sum=0;
for(inti=0;i sum+=data[i]*data[i];
}
printf("sum(CPU):
%d\n",sum);
编译后执行,确认两个结果相同。
计算运行时间
CUDA提供了一个clock函式,可以取得目前的timestamp,很适合用来判断一段程