深入浅出谈CUDA技术

所需积分/C币:35 2014-04-11 09:34:47 13.66MB PDF
收藏 收藏
举报

深入浅出谈CUDA技术 最适合cuda初学者的cuda教材
grid 1 block block o/(0)120 ble y1/(1(2 block(1, 0) 每个 thread都有自己的一份 register和 local memory的空间。同一个 block中的每个 thread 则有共享的一份 share memory。此外,所有的 thread(包括不同bock的 thread)都共享一份 global memory、 constant memory、和 texture memory。不同的grid则有各自的 global memory、 constant memory和 texture memory。这些不同的内存的差别,会在之后讨论。 执行模式 由于显示芯片大量并行计算的特性,它处理一些问题的方式,和一般CPU是不同的。主要的特点 包括: 1.内存存取 latency的问题:CPU通常使用 cache来减少存取主内存的次数,以避免内存 latency影响到执行效率。显示芯片则多半没有 cache(或很小),而利用并行化执行的方 式来隐藏内存的 latency(即,当第一个 thread需要等待内存读取结果时,则开始执行第二 个 thread,依此类推) 2.分支指令的问题:CPU通常利用分支预测等方式来减少分支指令造成的 pipeline bubble。显 示芯片则多半使用类似处理内存 latency的方式。不过,通常显示芯片处理分支的效率会比 较差 因此,最适合利用CUDA处理的问题,是可以大量并行化的问题,才能有效隐藏内存的 latency, 并有效利用显示芯片上的大量执行单元。使用CUDA时,同时有上千个 thread在执行是很正常的。 因此,如果不能大量并行化的问题,使用CUDA就没办法达到最好的效率了。 CUDA Toolkit的安装 日前NVD|A提供的 CUDA Toolkit(可从这里下载)支持 Windows(32bits及64bits版本) 及许多不同的 Linux版本 CUDA Toolkit需要配合C/C++ compiler。在 Windows下,目前只支持Ⅴ isual studio7x及 Visua| Studio8(包括免费的 Visual studio c++2005 Express)。 Visual studio6和gcc在 Windows下是不支援的。在 Linux下则只支援gcc 这里简单介绍一下在 Windows下设定并使用CUDA的方式。 「下载及安装 在 Windows下, CUDA Toolkit和 CUDA SDK都是由安装程序的形式安装的。 CUDA Toolkit包 括CUDA的基本工具,而 CUDA SDK则包括许多范例程序以及链接库。基本上要写CUDA的程 序,只需要安装 CUDA Toolkit即可。不过 CUDA SDK仍值得安装,因为里面的许多范例程序和 链接库都相当有用。 CUDA Toolkit安装完后,预设会安装在C:CUDA目录里。其中包括几个目录: bin--工具程序及动态链接库 doc-文件 include-- header檔 lib--链接库档案 open64--基于Open64的 CUDA compiler Src 些原始码 安装程序也会设定一些环境变量,包括 CUDA_BIN_PATH--工具程序的目录,默认为C: CUDAbin CUDA_ INC PATH- header文件的目录,默认为C: CUDAVinc cUDA_LIB_PATH--链接库文件的目录,默认为C: A CUDAlib 在 isual studio中使用cUDA cUDA的主要工具是ηwc,它会执行所需要的程序,将CUDA程序代码编译成执行档(或 object 檔)。在 Visual studio下,我们透过设定 custom build tool的方式,让 sual studio会自动执 行nvcc。 这里以 Visual studio2005为例: 首先,建立一个Wn32 Console模式的 project(在 Application Settings中记得勾选 Empty project),并新增一个档案,例如main.cu。 2.在main.cu上右键单击,并选择 Properties。点选 General,确定Tool的部份是选择 Custom Build tool。 3.选择 Custom Build Step,在 Command line使用以下设定 Release tx: "S(CUDA BIN PATH)\nVcc. exe"-ccbin"S(VCInstallDir)bin "-c DWIN32-D_CONSOLE-D_MBCS-Xcompiler/EHsc, /W3, /nologo, /Wp64, /02, /Zi, /MT "S(CUDA_ INC_ PATH)"-o $(Configuration Name )\(InputName). obj s(InputFile Name) Debug aI: "S(CUDA_BIN_PATH)\nVcc. exe"-ccbin"S(VCInstallDir)bin"-c D DEBUG-DWIN32-D CONSOLE-D MBCS-Xcompiler /EHsc, /W3, /nologo, /p64, /Od, /Zi, /RTC1, /MTd-I"S(CUDA_INC_PATH)"-0 S(ConfigurationName)\(InputName).obj (InputFileName) 4.如果想要使用软件仿真的模式,可以新增两个额外的设定: EmuRelease ar:"S(CUDA BIN PATH)\nVcc. exe"-ccbin"$(VCInstallDir)bin deviceemu-C-DWIN32-D CONSOLE-D MBCS-Xcompiler /EHsc, /3, / nologo, /Wp64, /02, /Zi, /MT-I"S(CUDA INC PATH)"-o S (Configuration Name) s(InputName). obj s(Input FileName) EmuDebug tI:"S(CUDA BIN PATH)\nVcc. exe"-ccbin"$(VCInstallDir)bin' deviceemu-C-D DEBUG-DWIN32-D CONSOLE-D MBCS-Xcompiler /EHsc, /3, /nologo, /Wp64, /Od, /Zi, /RTC1, /MTd-I1"S(CUDA_ INC_ PATH)"-o S(Configuration Name)\(InputName).obj $(Input FileName) 5.对所有的配置文件,在 Custom Build Step的 Outputs中加入 S(Configuration Name) S(InputName). obj 6.选择 project,右键单击选择 Properties,再点选 Linker。对所有的配置文件修改以下设定: General/Enable Incremental linking: no General/Additional Library Directories: S(CUDA_LIBPATH Input/Additional Dependencies: cudart lib 这样应该就可以直接在Ⅴ isual studio的IDE中,编辑CUDA程序后,直接 build以及执行程序 第一个cUDA程序 CUDA目前有两种不同的APl: Runtime AP和 Driver APl,两种AP|各有其适用的范围。由于 runtime ar较容易使用,一开始我们会以 runtime AP|为主 cUDA的初始化 首先,先建立一个档案 first cuda cu。如果是使用 Visual studio的话,则请先按照这里的设定方 式设定 project 要使用 runtime Ap的时候,需要 include cuda runt ime.h。所以,在程序的最前面,加上 inc⊥ude< stdio.h> #include <cuda runt ime. h> 接下来是一个 Init CUDA函式,会呼叫 runtime APl中,有关初始化CUDA的功能: bool InitCUDA( Int counti cudaGetDeviceCount(& count)i if(count 0) fprintf(stderr, "There is no device. \n") return false: int ii for(i =0; i count; i++)t cudaDeviceProp prop; if(cuda GetDeviceProperties(&prop, 1)== cuda Success) I if(prop. major >=1)i break; (主== count) fprintf(stderr, There is no device supporting CUDA 1.x. \n")i return false cuda SetDevice(i)i return true; 这个函式会先呼叫 cudaget Devi cecount函式,取得支持CUDA的装置的数目。如果系统上没 有支持CUDA的装置,则它会传回1,而 device0会是一个仿真的装置,但不支持CUDA1.0以 上的功能。所以,要确定系统上是否有支持CUDA的装置,需要对每个 device呼叫 cudaGetDeviceproperties函式,取得裝置的各项数据,并判断装置支持的CUDA版本 ( prop. major和 prop. mInor分别代表装置支持的版本号码,例如1.0则 prop. major为1而 rop. minor为0)。 透过 cuda GetDeviceProperties函式可以取得许多数据,除了装置支持的CUDA版本之外, 还有装置的名称、内存的大小、最大的 thread数目、执行单元的频率等等。详情可参考NVDA的 CUDA Programming Guide 在找到支持CUDA1.0以上的装置之后,就可以呼叫 cudasetDevice函式,把它设为目前要使 用的装置。 最后是main函式。在main函式中我们直接呼叫刚才的 InitCUDA函式,并显示适当的讯息: int main( if(! InitCUDA())( return 0 printf("CUDA initialized. \n")i return o 这样就可以利用nvcc来 compile这个程序了。使用 Visual studio的话,若按照先前的设定方式, 可以直接 Build Project并执行。 nvcc是CUDA的 compile工具,它会将cu檔拆解出在GPU上执行的部份,及在host上执 行的部份,并呼叫适当的程序进行 compile动作。在GPU执行的部份会透过NVDA提供的 compiler编译成中介码,而host执行的部份则会透过系统上的C++ compiler编译(在 Windows 上使用 Visual c++而在 Linux上使用gcc)。 编译后的程序,执行时如果系统上有支持CUDA的装置,应该会显示 CUDA initialized.的讯息, 否则会显示相关的错误讯息。 利用cUDA进行运算 到日前为止,我们的程序并没有做什么有用的工作。所以,现在我们加入一个简单的动作,就是把 大堆数字,计算出它的平方和。 首先,把程序最前面的 include部份改成: inc⊥ude< stdio.h> #include <stdlib. h> #include <cuda runtime. h> *define DATA SIZE 1048576 int data [DAtA siZe]i 并加入一个新函式 GenerateNumbers void GenerateNumbers(int *number, int size for(int i=0: i size; i++)( number [i]= rand( 10: 这个函式会产生一大堆0~9之间的随机数。 要利用CUDA进行计算之前,要先把数据复制到显卡内存中,才能让显示芯片使用。因此,需要取 得一块适当大小的显卡内存,再把产生好的数据复制进去。在main函式中加入 GenerateNumbers(data, DATA SIZe)i int* gpudatar *resulti cudaMalloc((void**)&gpudata, sizeof (int)* DATA_SIZe)i cudaMalloc((void**) &result, sizeof (int))i cudaMemcpy(gpudata, data, sizeof(int) DATA-SIZE, cudaMemcpy Host ToDevice)i 上面这段程序会先呼叫 GenerateNumbers产生随机数,并呼叫 cudaMa11oc取得一块显卡内 存( result则是用来存取计算结果,在稍后会用到),并透过 cudaMemcpy将产生的随机数复 制到显卡内存中。 cudaMalloc和 cudaMemcpy的用法和一般的 malloc及 memcpy类似,不过 cudaMemcpy则多出一个参数,指示复制内存的方向。在这里因为是从主内存复制到显卡内存,所 以使用 cudaMemcpy HostToDevice。如果是从显卡内存到主内存,则使用 cudaMemcpyDeviceToHost。这在之后会用到 接下来是要写在显示芯片上执行的程序。在CUDA中,在函式前面加上_g1oba1_表示这个函 式是要在显示芯片上执行的。因此,加入以下的函式: global_ static void sumofSquares(int *num, int* result) int sum= 0 Int i for(i= 0; 1 DATA SIZE; i++) sum + num[i]* num[ili *result sum 在显示芯片上执行的程序有一些限制,例如它不能有传回值。其它的限制会在之后提到。 接下来是要让CUDA执行这个函式。在CUDA中,要执行一个函式,使用以下的语法: 函式名称<<b1ock数目, thread数目, shared memory大小>>>(参数,); 呼叫完后,还要把结果从显示芯片复制回主内存上。在main函式中加入以下的程序: sumofSquares<<<l, 1,0>>>(gpudata, result) Int sumi cudaMemcpy(& sum, result, sizeof (int), cudaMemcpyDeviceToHost)i cudaFree(gpudata)i cudaFree(result)i printf("sum: d\n", sum)i 因为这个程序只使用一个 thread,所以 block数目、 thread数日都是1。我们也没有使用到任何 shared memory,所以设为0。编译后执行,应该可以看到执行的结果。 为了确定执行的结果正确,我们可以加上一段以CPU执行的程序代码,来验证结果: sum =0; for (int i =0;1< DATA SIZE; 1++) sum + data[i] data[] printf("sum (CPU): d\n", sum)i 编译后执行,确认两个结果相同。 计算运行时间 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(); for (i =0; 1< DATA SIZE; i++)I sum + num[i]* num [ili *result sum; *time clock()- starti 把main函式中间部份改成: int* gpudata, *resulti clock t* time: cudaMalloc((void**)&gpudata, sizeof(int)* DATA_ SIZE cudaMalloc((void**) result, sizeof (int))i cudaMalloc((void**)&time, sizeof(clock t) cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpy Host ToDevice)i sumOfSquares<<<l, l, 0>>>(gpudata, result int sum; clock t time used; cudaMemcpy(&sum, result, sizeof (int), cudaMemcpyDeviceToHost)i cudaMemcpy(&time_used, time sizeof(clock_ t) cudaMemcpyDeviceToHost)i cudaFree(gpudata)i cudaFree(result) printf("sum: d time: d\n",sum, time_used)i

...展开详情
试读 42P 深入浅出谈CUDA技术
立即下载 低至0.43元/次 身份认证VIP会员低至7折
一个资源只可评论一次,评论内容不能少于5个字
yangchong609 新手,入门教材。
2015-07-24
回复
zhangjjfan 感觉版本很低啊,但是这是唯一的资料了
2014-07-01
回复
924869924 仅有的CUDA的学习资料
2014-06-06
回复
HacyLu 非常好的入门教程
2014-05-30
回复
imperfect00 还可以,可以学习下
2014-05-22
回复
iceroom17 还是感觉有点老了
2014-05-04
回复
上传资源赚积分or赚钱
最新推荐
深入浅出谈CUDA技术 35积分/C币 立即下载
1/42
深入浅出谈CUDA技术第1页
深入浅出谈CUDA技术第2页
深入浅出谈CUDA技术第3页
深入浅出谈CUDA技术第4页
深入浅出谈CUDA技术第5页
深入浅出谈CUDA技术第6页
深入浅出谈CUDA技术第7页
深入浅出谈CUDA技术第8页
深入浅出谈CUDA技术第9页

试读结束, 可继续读4页

35积分/C币 立即下载 >