深入浅出的谈CUDA.docx

上传人:b****3 文档编号:10335796 上传时间:2023-05-25 格式:DOCX 页数:28 大小:68.54KB
下载 相关 举报
深入浅出的谈CUDA.docx_第1页
第1页 / 共28页
深入浅出的谈CUDA.docx_第2页
第2页 / 共28页
深入浅出的谈CUDA.docx_第3页
第3页 / 共28页
深入浅出的谈CUDA.docx_第4页
第4页 / 共28页
深入浅出的谈CUDA.docx_第5页
第5页 / 共28页
深入浅出的谈CUDA.docx_第6页
第6页 / 共28页
深入浅出的谈CUDA.docx_第7页
第7页 / 共28页
深入浅出的谈CUDA.docx_第8页
第8页 / 共28页
深入浅出的谈CUDA.docx_第9页
第9页 / 共28页
深入浅出的谈CUDA.docx_第10页
第10页 / 共28页
深入浅出的谈CUDA.docx_第11页
第11页 / 共28页
深入浅出的谈CUDA.docx_第12页
第12页 / 共28页
深入浅出的谈CUDA.docx_第13页
第13页 / 共28页
深入浅出的谈CUDA.docx_第14页
第14页 / 共28页
深入浅出的谈CUDA.docx_第15页
第15页 / 共28页
深入浅出的谈CUDA.docx_第16页
第16页 / 共28页
深入浅出的谈CUDA.docx_第17页
第17页 / 共28页
深入浅出的谈CUDA.docx_第18页
第18页 / 共28页
深入浅出的谈CUDA.docx_第19页
第19页 / 共28页
深入浅出的谈CUDA.docx_第20页
第20页 / 共28页
亲,该文档总共28页,到这儿已超出免费预览范围,如果喜欢就下载吧!
下载资源
资源描述

深入浅出的谈CUDA.docx

《深入浅出的谈CUDA.docx》由会员分享,可在线阅读,更多相关《深入浅出的谈CUDA.docx(28页珍藏版)》请在冰点文库上搜索。

深入浅出的谈CUDA.docx

深入浅出的谈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,很适合用来判断一段程

展开阅读全文
相关资源
猜你喜欢
相关搜索
资源标签

当前位置:首页 > 解决方案 > 学习计划

copyright@ 2008-2023 冰点文库 网站版权所有

经营许可证编号:鄂ICP备19020893号-2