深入浅出谈CUDA
发表时间：2008-11-21
GPGPU((非图像计算)模型，它使用C语言为基础，“CUDA是NVIDIA的GPGPU
直接以大多数人熟悉的C语言，写出在显示芯片上执行的程序，而不需要去学定的显示芯片的指令或是特殊的结构。”
GPU就是早期的图形运算处理器，但是随着发展，拥有较高带宽的GPU被人发出来进行除开图形运算的其它运算，提高程序的执行效率即GPGPU；CUDA是什么？
编者注：NVIDIA的GeFoce8800GTX发布后，它的通用计算架构CUDA经过多的推广后，现在已经在有相当多的论文发表，在商业应用软件等方面也初步出视频编解码、金融、地质勘探、科学计算等领域的产品，是时候让我们对其作更步的了解。为了让大家更容易了解CUDA，我们征得Hotball的本人同意，发表近亲自撰写的本文。这篇文章的特点是深入浅出，也包含了hotball本人编写一单CUDA程序的亲身体验，对于希望了解CUDA的读者来说是非常不错的入门PCINLIFE对本文的发表没有作任何的删减，主要是把一些台湾的词汇转换成大词汇以及作了若干"编者注"的注释。
现代的显示芯片已经具有高度的可程序化能力，由于显示芯片通常具有相当高的带宽，以及大量的执行单元，因此开始有利用显示芯片来帮助进行一些计算工作法，即GPGPU。CUDA即是NVIDIA的GPGPU模型。
NVIDIA的新一代显示芯片，包括GeForce8系列及更新的显示芯片都支持C NVIDIA免费提供CUDA的开发工具（包括Windows版本和Linux版本）序范例、文件等等，可以在CUDA Zone下载。
GPGPU的优缺点
使用显示芯片来进行运算工作，和使用CPU相比，主要有几个好处：
1.显示芯片通常具有更大的内存带宽（前端总线）。例如，NVIDIA的GeF
Ge 8800GTX具有超过50GB/s的内存带宽，而目前高阶CPU的内存带宽10GB/s左右。
对于带宽的解释,详细见文档什么是带宽。
2.显示芯片具有更大量的执行单元。例如GeForce8800GTX具有128个
"stream processors"，频率为1.35GHz。CPU频率通常较高，但是执元的数目则要少得多。
3.和高阶CPU相比，显卡的价格较为低廉。例如目前一张GeForce8800G
8800括512MB内存的价格，和一颗2.4GHz四核心CPU的价格相若。
当然，使用显示芯片也有它的一些缺点：
1.显示芯片的运算单元数量很多，因此对于不能高度并行化的工作，所能带来
助就不大。
2.显示芯片目前通常只支持32bits浮点数，且多半不能完全支持IEEE7
格，有些运算的精确度可能较低。目前许多显示芯片并没有分开的整数运元，因此整数运算的效率较差。
3.显示芯片通常不具有分支预测等复杂的流程控制单元，因此对于具有高度分
程序，效率会比较差。
4.目前GPGPU的程序模型仍不成熟，也还没有公认的标准。例如NVIDI
AMD/ATI就有各自不同的程序模型。
整体来说，显示芯片的性质类似stream processor，适合一次进行大量相同的CPU则比较有弹性，能同时进行变化较多的工作。
CUDA架构
CUDA是NVIDIA的GPGPU模型，它使用C语言为基础，可以直接以大多熟悉的C语言，写出在显示芯片上执行的程序，而不需要去学习特定的显示芯指令或是特殊的结构。
在CUDA的架构下，一个程序分为两个部份：host端和device端。Host端在CPU上执行的部份，而device端则是在显示芯片上执行的部份。Device 程序又称为"kernel（核心）"。通常host端程序会将数据准备好后，复制到显内存中，再由显示芯片执行device端程序，完成后再由host端程序将结果从的内存中取回。
由于CPU存取显卡内存时只能透过PCI Express接口，因此速度较慢（PC Express x16的理论带宽是双向各4GB/s），因此不能太常进行这类动作，以低效率。
在CUDA架构下，显示芯片执行时的最小单位是thread
thread((线程)。数个thread block((块)。一个block中的thread能存取同一块共享的内存，而且组成一个block
快速进行同步的动作。
每一个block所能包含的thread数目是有限的。不过，执行相同程序的blo grid((格子)。不同block中的thread无法存取同一个共享的内存，可以组成grid
无法直接互通或进行同步。因此，不同block中的thread能合作的程度是比较不过，利用这个模式，可以让程序不用担心显示芯片实际上能同时执行的thre 目限制。例如，一个具有很少量执行单元的显示芯片，可能会把各个block中thread顺序执行，而非同时执行。不同的grid则可以执行不同的程序（即ke Grid、block和thread的关系，如下图所示：
每个thread都有自己的一份register和local memory的空间。同一个bl b 中的每个thread则有共享的一份share memory。此外，所有的thread（包同block的thread）都共享一份global memory、constant memory、和te memory。不同的grid则有各自的global memory、constant memory和te memory。这些不同的内存的差别，会在之后讨论。
执行模式
由于显示芯片大量并行计算的特性，它处理一些问题的方式，和一般CPU是不主要的特点包括：
1.内存存取latency的问题：CPU通常使用cache来减少存取主内存的
而利用并行化执行的方式来隐藏内存的latency（即，当第一个thread 等待内存读取结果时，则开始执行第二个thread，依此类推）。
2.分支指令的问题：CPU通常利用分支预测等方式来减少分支指令造成的
pipeline bubble。显示芯片则多半使用类似处理内存latency的方式。
通常显示芯片处理分支的效率会比较差。
因此，最适合利用CUDA处理的问题，是可以大量并行化的问题，才能有效隐存的latency，并有效利用显示芯片上的大量执行单元。使用CUDA时，同时千个thread在执行是很正常的。因此，如果不能大量并行化的问题，使用CUD 没办法达到最好的效率了。
CUDA Toolkit的安装
目前NVIDIA提供的CUDA Toolkit（可从这里下载）支持Windows（32b 64bits版本）及许多不同的Linux版本。
CUDA Toolkit需要配合C/C++compiler。在Windows下，目前只支持Vi V Studio7.x及Visual Studio8（包括免费的Visual Studio C++2005Expre Visual Studio6和gcc在Windows下是不支援的。在Linux下则只支援这里简单介绍一下在Windows下设定并使用CUDA的方式。
下载及安装
在Windows下，CUDA Toolkit和CUDA SDK都是由安装程序的形式安装
链接库。基本上要写CUDA的程序，只需要安装CUDA Toolkit即可。不过C SDK仍值得安装，因为里面的许多范例程序和链接库都相当有用。
CUDA Toolkit安装完后，预设会安装在C:\CUDA目录里。其中包括几个目录
?bin--工具程序及动态链接库
?doc--文件
?include--header檔
?lib--链接库档案
?open64--基于Open64的CUDA compiler
?src--一些原始码
安装程序也会设定一些环境变量，包括：
?CUDA_BIN_PATH--工具程序的目录，默认为C:\CUDA\bin
?CUDA_INC_PATH--header文件的目录，默认为C:\CUDA\inc ?CUDA_LIB_PATH--链接库文件的目录，默认为C:\CUDA\lib 在Visual Studio中使用CUDA
CUDA的主要工具是nvcc，它会执行所需要的程序，将CUDA程序代码编译行档(或object檔)。在Visual Studio下，我们透过设定custom build to 方式，让Visual Studio会自动执行nvcc。
这里以Visual Studio2005为例：
1.首先，建立一个Win32Console模式的project（在Application Settin
记得勾选Empty project），并新增一个档案，例如main.cu。
2.在main.cu上右键单击，并选择Properties。点选General，确定To
部份是选择Custom Build Tool。
3.选择Custom Build Step，在Command Line使用以下设定：
o Release模式："$(CUDA_BIN_PATH)\nvcc.exe"-ccbin "$(VCInstallDir)bin"-c-DWIN32-D_CONSOLE-D_MBCS-Xcompile /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT-I"$(CUDA_INC_PATH)"-o
$(ConfigurationName)\$(InputName).obj$(InputFileName)
o 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.如果想要使用软件仿真的模式，可以新增两个额外的设定：
o EmuRelease模式："$(CUDA_BIN_PATH)\nvcc.exe"-ccbin "$(VCInstallDir)bin"-deviceemu-c-DWIN32-D_CONSOLE-D_MBC -Xcompiler/EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT-I"$(CUDA_INC_PA -o$(ConfigurationName)\$(InputName).obj$(InputFileName)
o EmuDebug模式："$(CUDA_BIN_PATH)\nvcc.exe"-ccbin "$(VCInstallDir)bin"-deviceemu-c-D_DEBUG-DWIN32-D_CONSO
-I"$(CUDA_INC_PATH)"-o$(ConfigurationName)\$(InputName).obj
$(InputFileName)
5.对所有的配置文件，在Custom Build Step的Outputs中加入
$(ConfigurationName)\$(InputName).obj。
6.选择project，右键单击选择Properties，再点选Linker。对所有的配置
修改以下设定：
o General/Enable Incremental Linking：No
o General/Additional Library Directories：$(CUDA_LIB_PATH) o Input/Additional Dependencies：cudart.lib
这样应该就可以直接在Visual Studio的IDE中，编辑CUDA程序后，直接以及执行程序了。
第一个CUDA程序
interface))：Runtim CUDA目前有两种不同的API(application programmers interface
和Driver API，两种API各有其适用的范围。由于runtime API较容易使用，始我们会以runetime API为主。
CUDA的初始化
首先，先建立一个档案first_cuda.cu。如果是使用Visual Studio的话，则请照这里的设定方式设定project。
要使用runtime API的时候，需要include cuda_runtime.h。所以，在程序前面，加上
#include
#include
接下来是一个InitCUDA(初始化（=initialization）)函式，会呼叫runtime AP 有关初始化CUDA的功能：
bool InitCUDA()
{
int count;
cudaGetDeviceCount(&count);
if(count==0){
fprintf(stderr,"There is no device.\n");
return false;
}
int i;
for(i=0;i cudaDeviceProp prop; if(cudaGetDeviceProperties(&prop,i)==cudaSucces break; } } } if(i==count){ fprintf(stderr,"There is no device supporting CUD 1.x.\n"); return false; } cudaSetDevice(i); return true; } 这个函式会先呼叫cudaGetDeviceCount函式，取得支持CUDA的装置的如果系统上没有支持CUDA的装置，则它会传回1，而device0会是一个仿装置，但不支持CUDA1.0以上的功能。所以，要确定系统上是否有支持CUD 装置，需要对每个device呼叫cudaGetDeviceProperties函式，取得装各项数据，并判断装置支持的CUDA版本（prop.major和prop.minor分别装置支持的版本号码，例如1.0则prop.major为1而prop.minor为0） 透过cudaGetDeviceProperties函式可以取得许多数据，除了装置支持的CUDA版本之外，还有装置的名称、内存的大小、最大的thread数目、执行单频率等等。详情可参考NVIDIA的CUDA Programming Guide。 在找到支持CUDA1.0以上的装置之后，就可以呼叫cudaSetDevice函式，设为目前要使用的装置。 最后是main函式。在main函式中我们直接呼叫刚才的InitCUDA函式，并适当的讯息： int main() { if(!InitCUDA()){ return0; } printf("CUDA initialized.\n"); return0; } 这样就可以利用nvcc来compile这个程序了。使用Visual Studio的话，若先前的设定方式，可以直接Build Project并执行。 nvcc是CUDA的compile工具，它会将.cu檔拆解出在GPU GPU((图形处理器 Unit))上执行的部份，及在host上执行的部份，并呼叫适Graphic Processing Unit 程序进行compile动作。在GPU执行的部份会透过NVIDIA提供的compi 译成中介码，而host执行的部份则会透过系统上的C++compiler编译（在Windows上使用Visual C++而在Linux上使用gcc）。 编译后的程序，执行时如果系统上有支持CUDA的装置，应该会显示CUDA initialized.的讯息，否则会显示相关的错误讯息。 利用CUDA进行运算 到目前为止，我们的程序并没有做什么有用的工作。所以，现在我们加入一个简动作，就是把一大堆数字，计算出它的平方和。 首先，把程序最前面的include部份改成： #include #include stdlib.h>> #include #define DATA_SIZE1048576 并加入一个新函式GenerateNumbers： void GenerateNumbers(int*number,int size) { for(int i=0;i number[i]=rand()%10; } } 这个函式会产生一大堆0~9之间的随机数。 要利用CUDA进行计算之前，要先把数据复制到显卡内存中，才能让显示芯片因此，需要取得一块适当大小的显卡内存，再把产生好的数据复制进去。在ma 式中加入： 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产生随机数，并呼叫cudaMallo cudaMemcpy将产生的随机数复制到显卡内存中。cudaMalloc和cudaMem 的用法和一般的malloc及memcpy类似，不过cudaMemcpy则多出一个指示复制内存的方向。在这里因为是从主内存复制到显卡内存，所以使用cudaMemcpyHostToDevice。如果是从显卡内存到主内存，则使用cudaMemcpyDeviceToHost。这在之后会用到。 接下来是要写在显示芯片上执行的程序。在CUDA中，在函式前面加上 __global__表示这个函式是要在显示芯片上执行的。因此，加入以下的函式__global__static void sumOfSquares(int*num,int*result) { int sum=0; int i; for(i=0;i sum+=num[i]*num[i]; } *result=sum; } 在显示芯片上执行的程序有一些限制，例如它不能有传回值。其它的限制会在之到。 接下来是要让CUDA执行这个函式。在CUDA中，要执行一个函式，使用以语法： 函式名称<< 呼叫完后，还要把结果从显示芯片复制回主内存上。在main函式中加入以下的sumOfSquares<<<1,1,0>>>(gpudata,result); int sum; cudaMemcpy(&sum,result,sizeof(int), cudaMemcpyDeviceToHost); cudaFree(gpudata); cudaFree(result); printf("sum:%d\n",sum); 因为这个程序只使用一个thread，所以block数目、thread数目都是1。我没有使用到任何shared memory，所以设为0。编译后执行，应该可以看到执结果。 为了确定执行的结果正确，我们可以加上一段以CPU执行的程序代码，来验证 sum=0; for(int i=0;i sum+=data[i]*data[i]; } printf("sum(CPU):%d\n",sum); 编译后执行，确认两个结果相同。 计算运行时间 CUDA提供了一个clock函式，可以取得目前的timestamp，很适合用来判段程序执行所花费的时间（单位为GPU执行单元的频率）。这对程序的优化也有用。要在我们的程序中记录时间，把sumOfSquares函式改成： __global__static void sumOfSquares(int*num,int*result, clock_t*time) { int sum=0; int i; clock_t start=clock(); sum+=num[i]*num[i]; } *result=sum; *time=clock()-start; } 把main函式中间部份改成： int*gpudata,*result; clock_t*time; cudaMalloc((void**)&gpudata,sizeof(int)*DATA_SIZE) cudaMalloc((void**)&result,sizeof(int)); cudaMalloc((void**)&time,sizeof(clock_t)); cudaMemcpy(gpudata,data,sizeof(int)*DATA_SIZE, cudaMemcpyHostToDevice); sumOfSquares<<<1,1,0>>>(gpudata,result,time); int sum; clock_t time_used; cudaMemcpy(&sum,result,sizeof(int), cudaMemcpyDeviceToHost); cudaMemcpyDeviceToHost); cudaFree(gpudata); cudaFree(result); printf("sum:%d time:%d\n",sum,time_used); 编译后执行，就可以看到执行所花费的时间了。 如果计算实际运行时间的话，可能会注意到它的执行效率并不好。这是因为我们序并没有利用到CUDA的主要的优势，即并行化执行。在下一段文章中，会讨何进行优化的动作。 改良第一个CUDA程序 在上一篇文章中，我们做了一个计算一大堆数字的平方和的程序。不过，我们也这个程序的执行效率并不理想。当然，实际上来说，如果只是要做计算平方和的用CPU做会比用GPU快得多。这是因为平方和的计算并不需要太多运算能力以几乎都是被内存带宽所限制。因此，光是把数据复制到显卡内存上的这个动作需要的时间，可能已经和直接在CPU上进行计算差不多了。 不过，如果进行平方和的计算，只是一个更复杂的计算过程的一部份的话，那么在GPU上计算还是有它的好处的。而且，如果数据已经在显卡内存上（例如在上透过某种算法产生），那么，使用GPU进行这样的运算，还是会比较快的。 刚才也提到了，由于这个计算的主要瓶颈是内存带宽，所以，理论上显卡的内存是相当大的。这里我们就来看看，倒底我们的第一个程序，能利用到多少内存带程序的并行化 我们的第一个程序，并没有利用到任何并行化的功能。整个程序只有一个thre GeForce 8800GT 上面，在GPU 上执行的部份（称为"kernel kernel""）大约花费64064 频率。GeForce 8800GT 的执行单元的频率是1.5GHz ，因此这表示它花费了约秒的间。1M 个32bits 数字的数据量是4MB ，因此，这个程序实际上使用的内宽，只有9.3MB/s 左右！这是非常糟糕的表现。 为什么会有这样差的表现呢？这是因为GPU 的架构特性所造成的。在CUDA 一般的数据复制到的显卡内存的部份，称为global memory 。这些内存是没有的，而且，存取global memory 所需要的时间（即latency ）是非常长的，通数百个cycles 。由于我们的程序只有一个thread ，所以每次它读取global me 的内容，就要等到实际读取到数据、累加到sum 之后，才能进行下一步。这就什么它的表现会这么的差。 由于global memory 并没有cache ，所以要避开巨大的latency 的方法，就利用大量的threads 。假设现在有大量的threads 在同时执行，那么当一个t 读取内存，开始等待结果的时候，GPU 就可以立刻切换到下一个thread ，并读
相关主题
文本预览