CUDA并行计算架构及编程.doc
- 文档编号:139715
- 上传时间:2022-10-04
- 格式:DOC
- 页数:7
- 大小:22.50KB
CUDA并行计算架构及编程.doc
《CUDA并行计算架构及编程.doc》由会员分享,可在线阅读,更多相关《CUDA并行计算架构及编程.doc(7页珍藏版)》请在冰豆网上搜索。
CUDA并行架构及编程
摘 要 CUDA是一种由NVIDIA推出的并行计算架构,非常适合大规模数据密集型计算。
CUDA使GPU的超高计算性能在数据处理和并行计算等通用计算领域发挥优势,本文讨论了CUDA的计算架构和基于GPU的CUDAC编程语言,CUDA使GPU流处理器阵列的性能得到充分发挥。
极大地提高了并行计算程序的效率。
关键词 并行计算,GPU通用计算,CUDA
AbstractCUDAisaparallelcomputingarchitectureintroducedbyNVIDIA,itmainlyusedforlargescaledata-intensivecomputing.CUDAmakesGPUahighperformanceinparallelcomputing,dataprocessingandothergeneralcomputing.thispaperdiscussestheCUDAcomputingarchitectureandCUDACprogramminglanguagebasedonGPU,CUDAmakesGPUstreamprocessorarraysfullyusedandGreatlyimprovedtheefficiencyofparallelcomputingprogram.
Keywords parallelcomputing,GPUgeneralpurposecomputation,CUDA
1引言
并行计算是指同时使用多种计算资源解决计算问题的过程。
并行计算科学中主要研究的是空间上的并行问题。
从程序和算法设计的角度来看,并行计算又可分为数据并行和任务并行。
一般来说,GPU更注重于数据并行计算,主要是将一个大任务化解成相同的各个子任务。
早期的GPU研究是通过可编程计算单元为屏幕上的每个像素计算出一个颜色值即渲染问题。
自CUDAC出现后,基于GPU的通用计算已经成为一个新的研究领域。
通常,像素着色器对各种颜色值进行合成并计算出最终的颜色值。
实际上,输入值可以为任意数据,这样不一定非要使用GPU来处理图形,还可以实现某些通用计算。
由于GPU有着很高的计算吞吐量,从而给大规模的数据计算应用提供了一种比CPU更加强大的计算能力。
CUDA是一种由NVIDIA推出的并行计算架构,该架构使GPU能够解决复杂的计算问题。
它包含了CUDA指令集架构以及GPU内部的并行计算引擎[1]。
随着显卡的发展,GPU越来越强大,在计算上已经超越了通用的CPU。
如此强大的芯片如果只是作为显卡会造成计算能力的浪费,因此NVIDIA推出CUDA,让显卡可以用于图像渲染以外的目的。
CUDA的GPU编程语言基于标准的C语言,通过在标准C语言的基础上增加一小部分关键字,任何有C语言基础的用户都很容易地开发CUDA的应用程序。
数以千计的软件开发人员正在使用免费的CUDA软件开发工具来解决各种专业中的问题[2]。
这些解决方案涵盖了石油天然气勘探、产品设计、医学成像以及科学研究等领域。
2CUDA架构
CUDA程序架构分为两部分:
主机和设备。
一般而言,主机指的是CPU及其内存,设备指的是GPU[1]。
在CUDA程序架构中,主程序由CPU来执行,而当遇到数据并行处理的部分,CUDA就会将程序编译成GPU能执行的程序,并传送到GPU。
这种函数在CUDA里叫做核函数。
在GPU中要执行的线程,根据最有效的数据共享来创建线程块,其类型不止一维。
在同一个块里的线程,使用同一个共享内存。
每个线程由线程ID标识。
这是线程块中的线程号。
为了帮助基于线程ID的复杂寻址,还可以将线程块指定为任意大小的一维、二维或高维线程阵列,并使用多个索引分量来标识每个线程[2]。
线程块的大小是也有限制的,所以不能把所有的线程都放到同一个块里。
可以用同样维度和大小的块来组成一个网格做批处理。
执行内核的线程被组织成线程块。
而线程块又组成了线程格,如图1。
线程块内一起协作的线程通过一些快速的共享内存有效地共享数据并同步执行,以协调内存访问。
用户可以在核函数中指定同步点。
线程块中的线程在到达此同步点时挂起。
图1 线程被组织为线程格与线程块 图2 SM单元的组成
在GPU内部,SM代表流处理器,即计算核心。
每个SM中又包含8个标量流处理器(SP)以及少量的其他计算单元[3],如图2。
实际上SP只是执行单元,并不是完整的处理核心。
处理核心必须包含取指、解码、分发逻辑和执行单元。
隶属同一SM的8个SP共用同一套取指和发射单元,也共用一块共享存储器。
一个线程块必须被分配到一个SM中,但是一个SM中同一时刻可以有多个活动线程块等待执行。
这可以更好地利用执行单元的资源,当一个线程块进行同步或者访问显存等高延迟操作时,另一个线程块就可以占用GPU执行资源。
每个线程有一个私有的本地存储器,即每个标量处理器私有的缓存。
每个线程块有一个共享存储器,块内所有的线程都可以访问。
每个线程网格都可以访问设备内存,即GPU片上的全局存储器。
每个线程均可以读写全局存储器。
本地存储器用于储存设备代码中无法放进寄存器的变量[3]。
CUDA架构的所有这些功能都是为了使GPU不仅能执行传统的图形计算,还能高效的执行通用计算。
3CUDAC编程
3.1CUDA并行编程
运行在GPU上的CUDA并行计算函数称为核函数(kernel)。
一个kernel函数并不是一个完整的程序,而是整个程序中的一个可以被并行执行的步骤[2]。
内核函数必须通过_global_函数类型限定符定义,如_global_voidkernel(void)。
并且只能在主机端代码中调用。
在调用时,必须声明内核函数的执行参数。
现考虑给一个N位矢量的每一位加上一个常数。
为完成此操作设计一个add函数,分别用C代码和CUDAC表示为:
#defineN10 #defineN10
voidadd(int*a,intb){ _global_voidadd(int*a,intb)
intindex=0; {
while(index a[index]=a[index]+b; if(index index+=1; } } 在CPU运行的程序中,通过while循环串行的给N位矢量每一位加上常数。 值得注意的是右边的CUDA核函数,在实际运行时,CUDA会产生许多在GPU上执行的线程,每一个线程都会去执行内核这个程序,虽然程序是同一份,但是因为设置了变量blockIdx,这是一个内置变量,在CUDA运行中已经预先定义了这个变量,变量的值是当前执行设备代码线程块的索引,从而取得不同的数据来进行并行计算。 这里考虑的实例是10位矢量的增量运算,当矢量的位数N远远大于10且进行的并不是简单的增量运算时,使用CUDAC就会有明显的差别。 intmain(void){ inthost_a[N]; int*device_a; cudaMalloc((void**)&device_a,N*sizeof(int)); //在设备上分配内存 for(inti=0;i cudaMemcpy(device_a,host_a,N*sizeof(int),cudaMemcpyHostToDevice); add<< cudaMemcpy(host_a,device_a,N*sizeof(int),cudaMemcpyDeviceToHost); cudaFree(device_a); //释放内存 return0; } 在定义了核函数后,可以在主函数中调用它。 add<< CPU串行代码完成的工作包括在核函数启动前进行数据准备和设备初始化的工作。 其中cudaMalloc()函数用来在GPU设备上分配内存,第一个参数是一个指针,指向用于保存新分配内存地址的变量,第二个参数指明分配内存的大小,需要注意的是由这个内存分配函数得到的指针是不能在主机内存上使用的。 cudaMemcpy()通过设置参数在CPU和GPU之间传递数据。 主函数先完成数组初始化的工作,再将数据传入GPU,并行计算完成后再将结果传回CPU。 3.2线程的同步与通信 上一节矢量增值运算的例子,通过创建N个含有一个线程的进程快给矢量的各位各加上一个常数。 上述运算并不需要考虑线程块内线程的同步与通信,但实际中的大多数应用都需要线程之间传递数据。 这也是CUDA所提供的最重要的创新,它使得工作在GPU上的线程可以协作解决问题,允许应用程序更加高效的执行[2]。 在同一个block中的线程通过同步函数_syncthreads()来保证块内的线程同步即线程块中的所有线程都执行到同一位置。 为了在硬件上用很小的代价就能实现_syncthreads()函数,一个block中所有线程的数据都必须交由同一处理核心进行处理[3]。 所以,这导致每个线程块中的线程数量受到处理核心硬件资源的限制。 目前规定每个块里最多只能有512个线程[2]。 在同一个块中的线程通过共享存储器交换数据,为了保证线程块中的各个线程能够有效协作,访问共享存储器的延迟必须很小。 所以在GPU中,共享存储器与执行单元的物理距离必须很小。 在CUDAC中同一个块中的线程通过共享变量来实现通信,使用__shared__来定义一个共享变量。 现利用一个大规模矩阵相乘的例子来说明线程块内线程同步与通信及线程之间的互斥。 矩阵A(MxN)与矩阵B(NxK)的乘法公式为[5]:
- 配套讲稿:
如PPT文件的首页显示word图标,表示该PPT已包含配套word讲稿。双击word图标可打开word文档。
- 特殊限制:
部分文档作品中含有的国旗、国徽等图片,仅作为作品整体效果示例展示,禁止商用。设计者仅对作品中独创性部分享有著作权。
- 关 键 词:
- CUDA 并行 计算 架构 编程