CUDA用于大量数据的超级计算.docx
- 文档编号:10418149
- 上传时间:2023-02-11
- 格式:DOCX
- 页数:24
- 大小:577.26KB
CUDA用于大量数据的超级计算.docx
《CUDA用于大量数据的超级计算.docx》由会员分享,可在线阅读,更多相关《CUDA用于大量数据的超级计算.docx(24页珍藏版)》请在冰豆网上搜索。
CUDA用于大量数据的超级计算
CUDA——用于大量数据的超级计算:
第一部分
发表时间:
2008-06-20
您是否有兴趣在使用高级语言(比如C)编程时,通过标准多核处理器将性能提升几个数量级?
您是否也希望拥有跨多个设备的伸缩能力?
RobFarber
CUDA使您能够在开发GPU上运行的软件时使用熟悉的编程概念。
RobFarber是西北太平洋国家实验室(PacificNorthwestNationalLaboratory)的高级科研人员。
他在多个国家实验室进行大量数据平行计算的研究,并且是好几个新创企业的创始人之一。
可以发邮件到rob.farber@pnl.gov联系他。
您是否有兴趣在使用高级语言(比如C)编程时,通过标准多核处理器将性能提升几个数量级?
您是否也希望拥有跨多个设备的伸缩能力?
很多人(包括我自己)都通过使用NVIDIA的CUDA(ComputeUnifiedDeviceArchitecture,即计算统一设备架构的简称)获得了这种高性能和可伸缩性,以编写廉价的多线程GPU程序。
我特别强调“编程”是因为CUDA是为您的工作服务的架构,它不会迫使您的工作适应有限的一组性能库。
使用CUDA,您可以发挥才能设计软件以便在多线程硬件上获得最佳性能——并从中获得乐趣,因为计算正确的映射是很有意思的,而且软件开发环境十分合理和直观。
本文是这一系列文章的第一篇,介绍了CUDA的功能(通过使用代码)和思维过程,帮助您将应用程序映射到多线程硬件(比如GPU)以获得较大的性能提升。
当然,并不是所有问题都可以有效映射到多线程硬件,因此我会介绍哪些可以进行有效映射,哪些不能,而且让您对哪些映射可以运行良好有个常识性的了解。
“CUDA编程”和“GPGPU编程”并不相同(尽管CUDA运行在GPU上)。
以前,为GPU编写软件意味着使用GPU语言编程。
我的一个朋友曾将这一过程描述为将数据从您的肘部拉到眼前。
CUDA允许您使用熟悉的编程概念开发能在GPU上运行的软件。
它可以通过将软件直接编译到硬件(例如,GPU汇编语言)避免图形层API的性能开销,这样可以提供更好的性能。
您可以任选一种CUDA设备。
图1和图2分别显示了运行在笔记本和台式机的离散GPU上的CUDA多体模拟(N-bodysimulation)程序。
CUDA真的可以将应用程序性能提高一到两个数量级——或者这只是一种与事实不符的夸张?
CUDA是一种相当新的技术,但是在一些书面作品和Internet中已经有很多实例突出介绍了这种技术在使用当前商用GPU硬件时对性能的极大提升。
表1和表2显示了NVIDIA和BeckmanInstitute网站上相关内容的总结。
CUDA的核心是让程序员能够使数千线程保持工作状态。
目前这一代NVIDIAGPU能够有效地支持大量线程,因此它们可以将应用程序性能提高一到两个数量级。
这些图形处理器的价位有高有低,几乎可供任何人使用。
较新的主板将通过提供更大的内存带宽、异步数据传输、原子操作和双精度浮点计算等多项硬件技术改进扩展CUDA的功能。
随着技术的不断进步,CUDA软件环境将不断扩展,最终GPU和“多核”处理器之间的区别也会逐渐消失。
作为开发人员,我们可以预计,具有成千上万活动线程的应用程序将变得很常见而且CUDA将会运行在多个平台上,包括一般用途的处理器。
在20世纪80年代,作为LosAlamosNationalLaboratory的科学家,我有幸使用了拥有多达65,536个平行处理器的ThinkingMachines超级计算机。
CUDA被证明是天生用于现代大量平行(即高线程)环境的框架。
它的性能优势非常明显。
我的一段生产代码,现在用CUDA编写并且运行在NVIDIAGPU上,与2.6-Ghz四核Opteron系统相比,具有明显的线形伸缩和几乎两个数量级的速度提升。
启用CUDA的图形处理器作为宿主计算机内的联合处理器运行。
这意味着每个GPU都被认为有其自己的内存和处理元素,它们是与宿主计算机分开的。
要进行有效的工作,数据必须在宿主计算机的内存空间和CUDA设备之间传输。
因此,性能结果必须包括IO时间才更有意义。
同事们喜欢将其称为“诚实的数据”,因为它们会更准确地反映将要交付生产的性能应用程序。
我认为与现有技术相比,一到两个数量级的性能提升是一个巨变,可以在很大程度上改变计算的某些方面。
例如,以前需要一年的计算任务现在只要几天就可以完成,几个小时的计算突然变得可交互了,因为使用新技术它们可以在几秒钟内完成,过去不易处理的实时处理任务现在变得极易处理。
最后,它为具有正确技能集和能力的顾问和工程师们提供了良好的机会,使他们可以编写高线程(大量平行)软件。
那么,对于您来说,这种计算能力又能给您的职业、应用程序或实时处理需求带来哪些好处呢?
开始不需要任何成本,可以轻松地从CUDAZone主页(找到GetCUDA)上下载CUDA。
之后,按照用于您特定操作系统的安装说明进行操作。
甚至不需要图形处理器,因为可以使用软件模拟器在当前笔记本或工作站上运行,可以立即开始工作。
当然,通过启用CUDA的GPU运行可以获得更好的性能。
或许您的计算机上已经有了一个启用CUDA的GPU。
在CUDAZone主页上查看支持CUDA的GPU链接(支持CUDA的GPU包括共享的片上内存和线程管理)。
如果购买新的图形处理卡,我建议您阅读本系列文章,因为我将介绍各种硬件特征(比如内存带宽、注册数量、原子操作等等)如何影响应用程序性能,这将有助您选择适合应用程序的硬件。
另外,CUDAZone论坛提供了关于CUDA各个方面的丰富信息,包括购买哪些硬件。
一旦安装,CUDAToolkit将提供一个合理的C语言程序开发工具集,它包括:
• nvccC编译器
• GPU的CUDAFFT和BLAS库
• 性能分析器
• alpha版本(截至2008年3月)的GPU的gdb调试器
• CUDA运行时驱动程序(现在还可以在标准的NVIDIAGPU驱动程序中得到)
• CUDA编程手册
nvccC编译器完成了大部分将C代码转换成将运行在GPU或模拟器上的可执行程序的工作。
幸好,汇编语言编程不要求达到很高的性能。
今后的文章将介绍从其它高级语言,包括C++、FORTRAN和Python使用CUDA的内容。
我假设您熟悉C/C++。
不需要有平行编程或CUDA经验。
这与现有CUDA文档是一致的。
创建和运行CUDAC语言程序与创建和运行其它C编程环境的工作流是一样的。
面向Windows和Linux环境的明确构建和运行说明在CUDA文档中。
简言之,这一工作流就是:
1. 使用最喜欢的编辑器创建或编辑CUDA程序。
注意:
CUDAC语言程序的后缀为.cu。
2. 使用nvcc编译程序创建可执行程序(NVIDIA提供了带有示例的完整makefiles。
通常用于CUDA设备时您只需键入make,用于模拟器时只需键入makeemu=1)。
3. 运行可执行程序。
列表1是一个带您入门的简单CUDA程序。
它只是一个简单的程序,调用CUDAAPI将数据移入和移出CUDA设备。
并没有添加新内容,以免在学习如何使用工具构建和运行CUDA程序时发生混淆。
在下一篇文章中,我将介绍如何开始使用CUDA设备执行一些工作。
请尝试使用这些开发工具。
对初学者的一些建议:
可以使用printf语句看看在模拟器下运行时(使用makeemu=1构建可执行程序)GPU上会发生什么。
还可以随意试验调试器的alpha版本。
相关文章
MicrosoftTalks'EmbeddedSystems'
TheROIofDistributedDevelopment
Cat:
AFunctionalStack-BasedLittleLanguage
IEEEIntroducesEntry-LevelSoftwareDeveloperCertification
GettingWiredforTerahertzComputing
排名前5位的文章
ScalingScrum
InterruptPolitely
IdentifyingTopDevelopers
AnInterviewwithBjarneStroustrup
OptimizingMath-IntensiveApplicationswithFixed-PointArithmetic
CUDA——用于大量数据的超级计算:
第二部分
发表时间:
2008-06-23
本文在第一个示例的基础上添加了几行代码,以便在CUDA设备上进行简单的计算——特别是在浮点数组中以1为增量增加每个元素。
令人惊喜的是,该示例已经提供了使用CUDA解决很多问题的基本框架(“将数据移动到支持CUDA的设备、进行计算并获取结果”)!
在本系列文章的第一部分,我展示了第一个简单的CUDA(ComputeUnifiedDeviceArchitecture,计算统一设备架构)程序——moveArrays.cu,使您熟悉用于构建和执行程序的CUDA工具。
对于C程序员而言,这个程序只是调用了CUDAAPI将数据移入和移出CUDA设备。
并没有添加新内容,以免在学习如何使用工具构建和运行CUDA程序时发生混淆。
本文在第一个示例的基础上添加了几行代码,以便在CUDA设备上进行简单的计算——特别是在浮点数组中以1为增量增加每个元素。
令人惊喜的是,该示例已经提供了使用CUDA解决很多问题的基本框架(“将数据移动到支持CUDA的设备、进行计算并获取结果”)!
在开始更高级的话题之前,您首先需要了解:
• 什么是内核?
内核是一个可以从主机调用又可以在CUDA设备上执行的函数——同时由多个线程平行执行。
• 主机如何调用内核?
这涉及指定内核名称及执行配置。
就本专栏文章而言,执行配置仅指定义在运行CUDA设备内核时组中的平行线程数及要使用的组数。
这实际上是个很重要的话题,我们将在以后的栏目中详细介绍。
• 如何同步内核和主机代码。
在列表1的最上面(incrementArrays.cu),我们可以看到主机例程的例子,incrementArrayOnHost和我们的第一个内核,incrementArraysOnDevice。
主机函数incrementArrayOnHost只是对数组元素数的简单循环,以1为增量增加每个数组元素。
此函数用于在此代码末尾进行比较,以验证内核在CUDA设备上进行了正确的计算。
在列表1稍下面的位置是我们的第一个CUDA内核,incrementArrayOnDevice。
CUDA提供了几个对C语言的扩展。
该函数类型限定符__global__将函数声明为CUDA设备上的可执行内核,只能从主机调用。
所有内核必须声明返回类型为void。
内核incrementArrayOnDevice与incrementArrayOnHost进行相同的计算。
仔细查看incrementArrayOnDevice会发现里面没有循环!
这是因为该函数是由CUDA设备上的一组线程同时执行的。
但是,每个线程都具有一个唯一的ID,可以用于计算不同的数组索引或制定控制决策(比如,如果数组索引超过数组大小则不进行任何操作)。
这使得incrementArrayOnDevice的计算变得非常简单,如同计算寄存器变量idx中的唯一ID一样,然后使用该变量唯一地引用数组中的每个元素并以1为增量递增。
因为线程数可能超过数组大小,因此先将idx与N相比较(N是向内核中传递的一个参数,用于指定数组中的元素数),看一下是否需要进行一些操作。
那么内核是如何调用,执行配置又是如何指定的呢?
控制按顺序流过源代码,从main开始到包含列表1中Part2of2语句的注释下面。
//incrementArray.cu
#include
#include
#include
voidincrementArrayOnHost(float*a,intN)
{
inti;
for(i=0;i } __global__voidincrementArrayOnDevice(float*a,intN) { intidx=blockIdx.x*blockDim.x+threadIdx.x; if(idx } intmain(void) { float*a_h,*b_h; //pointerstohostmemory float*a_d; //pointertodevicememory inti,N=10; size_tsize=N*sizeof(float); //allocatearraysonhost a_h=(float*)malloc(size); b_h=(float*)malloc(size); //allocatearrayondevice cudaMalloc((void**)&a_d,size); //initializationofhostdata for(i=0;i //copydatafromhosttodevice cudaMemcpy(a_d,a_h,sizeof(float)*N,cudaMemcpyHostToDevice); //docalculationonhost incrementArrayOnHost(a_h,N); //docalculationondevice: //Part1of2.Computeexecutionconfiguration intblockSize=4; intnBlocks=N/blockSize+(N%blockSize==0? 0: 1); //Part2of2.CallincrementArrayOnDevicekernel incrementArrayOnDevice<< //Retrieveresultfromdeviceandstoreinb_h cudaMemcpy(b_h,a_d,sizeof(float)*N,cudaMemcpyDeviceToHost); //checkresults for(i=0;i //cleanup free(a_h);free(b_h);cudaFree(a_d); 列表1: incrementArrays.cu. 这将排队启动支持CUDA的设备上的incrementArrayOnDevice,并说明添加到C语言的另一个CUDA,即对CUDA内核的异步调用。 该调用指定了内核的名称和封闭在三角括号"<<<"and">>>"之间的执行配置。 注意指定执行配置的两个参数: nBlocks和blockSize,将在下面对它们进行介绍。 任何对内核调用的参数都通过标准C语言参数列表提供,该列表包含以标准C语言样式"("and")"分界的函数。 在本示例中,指向设备全局内存的指针a_d(它包含数组元素)和N(数组元素数)都被传递到内核。 因为CUDA设备是空闲的,内核立即开始根据执行配置和函数参数运行。 同时,内核启动后,主机继续执行代码的下一行。 此时,CUDA设备和主机同时运行他们各自的程序。 在incrementArrays.cu中,主机立即调用cudaMemcpy,它会等待设备上的所有线程完成(例如,从incrementArrayOnDevice返回),之后它将修改后的数组拉回主机。 该程序在主机系统进行完串行比较后完成,以验证我们在平行CUDA设备上通过incrementArrayOnDevice得到的结果与在主机上通过串行版incrementArrayOnHost得到的结果相同。 在内核启动时要通过执行配置(在本例中通过包含在三角括号"<<<"and">>>"之间的变量nBlocks和blockSize)确定几个变量,这几个变量对任何内核都可用。 nBlocks和blockSize背后的思想是非常精妙的,它使开发人员能够解决硬件限制而无需重新编译应用程序——这是通过CUDA开发商业软件的本质特征。 块中的线程有能力彼此通信和同步,这将在今后的专栏文章中介绍。 这是一个绝佳的软件特性,只是从硬件的角度来看花费较大。 比起较便宜的(旧)设备,这更需要较昂贵的(未来的)设备来支持更大的每块线程数。 创建网格抽象可以使开发人员只需考虑(无需重新编译)区分硬件的能力而不用管价格和年代。 网格实际上将具有同样维度和大小的块对同一内核的调用汇成一批处理,然后乘以一个因子nBlocks(该因子是可以在单个内核调用中启动的线程数)。 性能较差的设备可能只能同时运行一个或几个线程块,而性能强大(较贵和未来的)的设备可能一次运行很多个线程块。 使用网格抽象设计软件需要在同时运行的多个独立线程之间进行平衡,并需要块内具有大量能够彼此合作的线程。 请认识到与两种线程相关的成本。 当然,不同算法会提出不同的要求,但在可能时,要尽量使用较大数量的线程块。 在支持CUDA的设备上的内核中,有几个可用的内置变量,它们是通过内核调用的执行配置设置的。 它们是: • blockIdx包含网格内的块索引。 • threadIdx包含块内的线程索引。 • blockDim包含块内的线程数。 这些变量是包含整数变量组件的结构。 例如,块有x-、y-和z-整数组件,因为它们是三维的。 而网格只有x-和y-组件,因为它们是二维的。 本示例只使用了这些变量的x-组件,因为我们移动到CUDA设备的数组是一维的(今后的专栏文章将介绍二维和三维配置能力的功效,以及如何利用这种功效)。 我们的示例内核使用这些内置的变量,通过下面的语句来计算线程索引idx: intidx=blockIdx.x*blockDim.x+threadIdx.x; 变量nBlocks和blockSize分别是网格中块的数量和每个块中的线程数。 在本示例中,它们就在主机代码的内核调用前初始化: intblockSize=4; intnBlocks=N/blockSize+(N%blockSize==0? 0: 1); 当N不能被blockSize整除时,nBlocks计算中的最后一项会加上一个额外的块,这意味着在有些情况下,块中的某些线程将不会进行任何有用的工作。 很明显,本例人为地使其简化了,因为它假设数组大小小于能够被包含在4个线程块中的线程数。 这很显然过于简单了,但它让我们能够通过简单的代码了解对incrementArrayOnDevice的内核调用。 还有很重要的一点需要强调,每个线程都能够访问设备上的整个数组a_d。 在内核启动时没有固有的数据分区。 这由程序员在编写内核时根据要识别和利用的计算的数据平行情况来决定。 图1说明了如何计算idx和如何引用数组a_d。 (如果前面的文本有任何不清楚的地方,我建议向incrementArrayOnDevice添加一个printf语句,以便将idx和用于计算它的相关变量打印出来。 为仿真器编译程序,"makeemu=1",运行它看看会发生什么。 一定要指定到仿真器可执行程序的正确路径来查看printf输出。 ) 图1 同样,内核调用是不同步的——在内核启动后,控制立即返回到主机CPU。 之前所有CUDA调用结束后,内核将在CUDA设备上运行。 不同步的内核调用是重叠主机和设备上的计算的极好方式。 在本例中,对incrementArrayOnHost的调用可以放在对incrementArrayOnDevice的调用之后,以重叠主机和设备上的计算来获得更好的性能。 主机和设备可以同时计算,这取决于内核完成计算需要的时间量。 在继续阅读下一篇专栏文章之前,我建议: • 尝试改变N和nBlocks的值。 看一下当它们超过设备能力时会发生什么。 • 想想如何引入循环来处理任意大小的数组。 • 区分不同类型的支持CUDA的设备的内存(比如,全局内存、寄存器、共享内存和持久内存)。 看一下CUDA占用率计算器,以及nvcc选项-cubin或--ptxas-options=-v,来决定内核中使用的寄存器的数量。 更多信息 有关本系列文章的后续部分,请参见CUDA,SupercomputingfortheMasses: Part3。 有关本系列文章的前一部分,请参见CUDA,SupercomputingfortheMasses: Part1。 单击这里CUDA,获得有关CUDA的更多信息,单击这里NVIDIA,获得有关NVIDIA的更多信息。 相关文章 MaximizeLocality,MinimizeContention QtandWindowsCE NanosoccerTakestheField FightingBackAgainstSpam-ZombieHordes Dr.Dobb'sAgileNewsletter05/08 排名前5位的文章 CUDA,SupercomputingfortheMasses: Part3 CUDA,SupercomputingfortheMasses: Part1 CUDA,SupercomputingfortheMasses: Part2 JForth: ImplementingForthinJava Tunny,ColossusandAda: KeepinganOpenMind CUDA——用于大量数据的超级计算: 第三部分 发表时间: 2008-06-23 本专栏和以下几个专栏文章将利用一个简单的数组反向应用程序来扩充您的知识,将重点介绍共享内存的性能影响。 我将和CUDA剖析工具一起
- 配套讲稿:
如PPT文件的首页显示word图标,表示该PPT已包含配套word讲稿。双击word图标可打开word文档。
- 特殊限制:
部分文档作品中含有的国旗、国徽等图片,仅作为作品整体效果示例展示,禁止商用。设计者仅对作品中独创性部分享有著作权。
- 关 键 词:
- CUDA 用于 大量 数据 超级 计算