深入浅出谈CUDA.docx
- 文档编号:8695880
- 上传时间:2023-02-01
- 格式:DOCX
- 页数:33
- 大小:96.25KB
深入浅出谈CUDA.docx
《深入浅出谈CUDA.docx》由会员分享,可在线阅读,更多相关《深入浅出谈CUDA.docx(33页珍藏版)》请在冰豆网上搜索。
深入浅出谈CUDA
深入浅出谈CUDA
(一):
书签
“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通常利用分支预测等方式来减少分支指令造成的pipeline(流水线)bubble。
显示芯片则多半使用类似处理内存latency的方式。
不过,通常显示芯片处理分支的效率会比较差。
因此,最适合利用CUDA处理的问题,是可以大量并行化的问题,才能有效隐藏内存的latency,并有效利用显示芯片上的大量执行单元。
使用CUDA时,同时有上千个thread在执行是很正常的。
因此,如果不能大量并行化的问题,使用CUDA就没办法达到最好的效率了。
深入浅出谈CUDA
(二):
CUDA是NVIDIA的GPGPU模型,它使用C语言为基础,可以直接以大多数人熟悉的C语言,写出在显示芯片上执行的程序,而不需要去学习特定的显示芯片的指令或是特殊的结构。
”
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使用以下设定:
Release模式:
"$(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)
Debug模式:
"$(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.如果想要使用软件仿真的模式,可以新增两个额外的设定:
EmuRelease模式:
"$(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)
EmuDebug模式:
"$(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(三):
“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()/*bool是什么意思啊?
*/
{ intcount;
cudaGetDeviceCount(&count);/*取得支持cuda的装置数目*/
if(count==0){
fprintf(stderr,"Thereisnodevice.\n");
returnfalse; }
inti;/*可以一次性定义在一起嘛?
?
*/
for(i=0;i cudaDevicePropprop;/*取得支持cuda装置的数据,如支持cuda版本等,见下? */ 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) /*产生size个介于0-9之间的随机数*/ { 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。 这在之后会用到。 (关于malloc及memcpy请点击链接,有说明。 。 后面出现我觉得有必要介绍的函数时都会链接说明文件。 ) 接下来是要写在显示芯片上执行的程序。 在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,很适合用来判断一段程序执行所花费的时间(单位为GPU执行单元的频率)。 这对程序的优化也相当有用。 要在我们的程序中记录时间,把sumOfSquares函式改成: __global__stat
- 配套讲稿:
如PPT文件的首页显示word图标,表示该PPT已包含配套word讲稿。双击word图标可打开word文档。
- 特殊限制:
部分文档作品中含有的国旗、国徽等图片,仅作为作品整体效果示例展示,禁止商用。设计者仅对作品中独创性部分享有著作权。
- 关 键 词:
- 深入浅出 CUDA