计算统一设备架构.docx
《计算统一设备架构.docx》由会员分享，可在线阅读，更多相关《计算统一设备架构.docx（92页珍藏版）》请在冰豆网上搜索。
计算统一设备架构
NVIDIACUDA
计算统一设备架构
编程指南
版本2.0
6/7/2008
图表目录
图1-1.CPU和GPU的每秒浮点运算次数和存储器带宽
图1-2.GPU中的更多晶体管用于数据处理...............................................................................2
图2-1.线程块网格.........................................................................................................................5
图2-2.存储器层次结构..................................................................................................................6
图2-3.异构编程.............................................................................................................................7
图2-4.计算统一设备架构软件栈...............................................................................................8
图3-1.硬件模型............................................................................................................................10
图4-1.库上下文管理..................................................................................................................30
图5-1.接合后的存储器访问模式示例.....................................................................................39
图5-2.未为计算能力是1.0或1.1的设备接合的全局存储器访问模式示例........................40
图5-3.未为计算能力是1.0或1.1的设备接合的全局存储器访问模式示例........................41
图5-4.计算能力为1.2或更高的设备的全局存储器访问示例............................................42
图5-5.无存储体冲突的共享存储器访问模式示例.................................................................45
图5-6.无存储体冲突的共享存储器访问模式示例.................................................................46
图5-7.有存储体冲突的共享存储器访问模式示例..................................................................47
图5-8.使用广播机制的共享存储器读取访问模式示例............................................................48
图6-1.矩阵乘法...........................................................................................................................52
第1章简介
1.1CUDA：
可伸缩并行编程模型
多核CPU和多核GPU的出现意味着并行系统已成为主流处理器芯片。
此外，根据摩尔定律，其并行性将不断扩展。
这带来了严峻的挑战，我们需要开发出可透明地扩展并行性的应用软件，以便利用日益增加的处理器内核数量，这种情况正如3D图形应用程序透明地扩展其并行性以支持配备各种数量的内核的多核GPU。
CUDA是一种并行编程模型和软件环境，用于应对这种挑战，同时保证熟悉C语言等标准编程语言的程序员能够迅速掌握CUDA。
CUDA的核心有三个重要抽象概念：
线程组层次结构、共享存储器、屏蔽同步（barriersynchronization），可轻松将其作为C语言的最小扩展级公开给程序员。
这些抽象提供了细粒度的数据并行化和线程并行化，嵌套于粗粒度的数据并行化和任务并行化之中。
它们将指导程序员将问题分解为更小的片段，以便通过协作的方法并行解决。
这样的分解保留了语言表达，允许线程在解决各子问题时协作，同时支持透明的可伸缩性，使您可以安排在任何可用处理器内核上处理各子问题：
因而，编译后的CUDA程序可以在任何数量的处理器内核上执行，只有运行时系统需要了解物理处理器数量。
1.2GPU：
高度并行化、多线程、多核处理器
市场迫切需要实时、高清晰度的3D图形，可编程的GPU已发展成为一种高度并行化、多线程、多核的处理器，具有杰出的计算功率和极高的存储器带宽，如图1-1所示。
图1-1.CPU和GPU的每秒浮点运算次数和存储器带宽
CPU和GPU之间浮点功能之所以存在这样的差异，原因就在于GPU专为计算密集型、高度并行化的计算而设计，上图显示的正是这种情况，因而，GPU的设计能使更多晶体管用于数据处理，而非数据缓存和流控制，如图1-2所示。
图1-2.GPU中的更多晶体管用于数据处理
更具体地说，GPU专用于解决可表示为数据并行计算的问题——在许多数据元素上并行执行的程序，具有极高的计算密度（数学运算与存储器运算的比率）。
由于所有数据元素都执行相同的程序，因此对精密流控制的要求不高；由于在许多数据元素上运行，且具有较高的计算密度，因而可通过计算隐藏存储器访问延迟，而不必使用较大的数据缓存。
数据并行处理会将数据元素映射到并行处理线程。
许多处理大型数据集的应用程序都可使用数据并行编程模型来加速计算。
在3D渲染中，大量的像素和顶点集将映射到并行线程。
类似地，图像和媒体护理应用程序（如渲染图像的后期处理、视频编码和解码、图像缩放、立体视觉和模式识别等）可将图像块和像素映射到并行处理线程。
实际上，在图像渲染和处理领域之外的许多算法也都是通过数据并行处理加速的——从普通信号处理或物理仿真一直到数理金融或数理生物学。
CUDA编程模型非常适合公开GPU的并行功能。
最新一代的NVIDIAGPU基于Tesla架构（在附录A中可以查看所有支持CUDA的GPU列表），支持CUDA编程模型，可显著加速CUDA应用程序。
1.3文档结构
本文档分为以下几个章节：
⏹第1章是CUDA和GPU的简介。
⏹第2章概述CUDA编程模型。
⏹第3章介绍GPU实现。
⏹第4章介绍CUDAAPI和运行时。
⏹第5章提供如何实现最高性能的一些指南。
⏹第6章通过一些简单的示例代码概况之前各章的内容。
⏹附录A提供各种设备的技术规X。
⏹附录B列举CUDA中支持的数学函数。
⏹附录C列举CUDA中支持的原子函数。
⏹附录D详细说明纹理获取。
第2章编程模型
CUDA允许程序员定义称为内核（kernel）的C语言函数，从而扩展了C语言，在调用此类函数时，它将由N个不同的CUDA线程并行执行N次，这与普通的C语言函数只执行一次的方式不同。
在定义内核时，需要使用_global_声明说明符，使用一种全新的<<<…>>>语法指定每次调用的CUDA线程数：
//Kerneldefinition
__global__voidvecAdd（float*A,float*B,float*C）
{
}
intmain（）
{
//Kernelinvocation
vecAdd<<<1,N>>>（A,B,C）;
}
执行内核的每个线程都会被分配一个独特的线程ID，可通过内置的threadIdx变量在内核中访问此ID。
以下示例代码将大小为N的向量A和向量B相加，并将结果存储在向量C中：
__global__voidvecAdd（float*A,float*B,float*C）
{
inti=threadIdx.x;
C[i]=A[i]+B[i];
}
intmain（）
{
//Kernelinvocation
vecAdd<<<1,N>>>（A,B,C）;
}
执行vecAdd（）的每个线程都会执行一次成对的加法运算。
2.1线程层次结构
为方便起见，我们将threadIdx设置为一个包含3个组件的向量，因而可使用一维、二维或三维缩影标识线程，构成一维、二维或三维线程块。
这提供了一种自然的方法，可为一个域中的各元素调用计算，如向量、矩阵或字段。
下面的示例代码将大小为NxN的矩阵A和矩阵B相加，并将结果存储在矩阵C中：
__global__voidmatAdd（floatA[N][N],floatB[N][N],
floatC[N][N]）
{
inti=threadIdx.x;
intj=threadIdx.y;
C[i][j]=A[i][j]+B[i][j];
}
intmain（）
{
//Kernelinvocation
dim3dimBlock（N,N）;
matAdd<<<1,dimBlock>>>（A,B,C）;
}
线程的索引及其线程ID有着直接的关系：
对于一维块来说，两者是相同的；对于大小为（Dx，Dy）的二维块来说，索引为（x,y）的线程的ID是（x+yDx）；对于大小为（Dx，Dy，Dz）的三维块来说，索引为（x,y,z）的线程的ID是（x+yDx+ZDxDy）。
一个块内的线程可彼此协作，通过一些共享存储器来共享数据，并同步其执行来协调存储器访问。
更具体地说，可以通过调用_syncthreads（）_内函数在内核中指定同步点；_syncthreads（）_起到屏障的作用，块中的所有线程都必须在这里等待处理。
为实现有效的协作，共享存储器应该是接近各处理器核心的低延迟存储器，如L1缓存，_syncthreads（）_应是轻量级的，一个块中的所有线程都必须位于同一个处理器核心中。
因而，一个处理器核心的有限存储器资源制约了每个块的线程数量。
在NVIDIATesla架构中，一个线程块最多可以包含512个线程。
但一个内核可能由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数量。
这些块将组织为一个一维或二维线程块网格，如图2-1所示。
该网格的维度由<<<…>>>语法的第一个参数指定。
网格内的每个块多可由一个一维或二维索引标识，可通过内置的blockIdx变量在内核中访问此索引。
可以通过内置的blockDim变量在内核中访问线程块的维度。
此时，之前的示例代码应修改为：
__global__voidmatAdd（floatA[N][N],floatB[N][N],
floatC[N][N]）
{
inti=blockIdx.x*blockDim.x+threadIdx.x;
intj=blockIdx.y*blockDim.y+threadIdx.y;
if（iC[i][j]=A[i][j]+B[i][j];
}
intmain（）
{
//Kernelinvocation
dim3dimBlock（16,16）;
dim3dimGrid（（N+dimBlock.x–1）/dimBlock.x,
（N+dimBlock.y–1）/dimBlock.y）;
matAdd<<>>（A,B,C）;
}
我们随机选择了大小为16x16的线程块（即包含256个线程），此外创建了一个网格，它具有足够的块，可将每个线程作为一个矩阵元素，这与之前完全相同。
线程块需要独立执行：
必须能够以任意顺序执行、能够并行或顺序执行。
这种独立性需求允许跨任意数量的核心安排线程块，从而使程序员能够编写出可伸缩的代码。
一个网格内的线程块数量通常是由所处理的数据大小限定的，而不是由系统中的处理器数量决定的，前者可能远远超过后者的数量。
图2-1.线程块网格
2.2存储器层次结构
CUDA线程可在执行过程中访问多个存储器空间的数据，如图2-2所示。
每个线程都有一个私有的本地存储器。
每个线程块都有一个共享存储器，该存储器对于块内的所有线程都是可见的，并且与块具有相同的生命周期。
最终，所有线程都可访问相同的全局存储器。
此外还有两个只读的存储器空间，可由所有线程访问，这两个空间是固定存储器空间和纹理存储器空间。
全局、固定和纹理存储器空间经过优化，适于不同的存储器用途（参见第5.1.2.1、5.1.2.3和5.1.2.4）。
纹理存储器也为某些特殊的数据格式提供了不同的寻址模式以及数据过滤（参见第4.3.4）。
对于同一个应用程序启动的内核而言，全局、固定和纹理存储器空间都是持久的。
图2-2.存储器层次结构
2.3主机和设备
如图2-3所示，CUDA假设CUDA线程可在物理上独立的设备上执行，此类设备作为运行C语言程序的主机的协同处理器操作。
例如，当内核在GPU上执行，而C语言程序的其他部分在CPU上执行时，就是这样一种情况。
此外，CUDA还假设主机和设备均维护自己的DRAM，分别称为主机存储器和设备存储器。
因而，一个程序通过调用CUDA运行时来管理对内核可见的全局、固定和纹理存储器空间（详见第4章）。
这包括设备存储器分配和取消分配，还包括主机和设备存储器之间的数据传输。
串行代码在主机上执行，而并行代码在设备上执行。
图2-3.异构编程
2.4软件栈
CUDA软件栈包含多个层，如图2-4所示：
设备驱动程序、应用程序编程接口（API）及其运行时、两个较高级别的通用数学库，即CUFFT和CUBLAS，这些内容均在其他文档中介绍。
图2-4.计算统一设备架构软件栈
2.5计算能力
一个设备的计算能力（putecapability）由主要修订号和次要修订号定义。
具有相同主要修订号的设备属于相同的核心架构。
附录A中列举的设备均为计算能力是1.x的设备（其主要修订号为1）。
次要修订号对应于核心架构的增量式改进，可能包含新特性。
附录A提供了各种计算能力的技术规X。
第3章GPU实现
NVIDIA于2006年11月引入的Tesla统一图形和计算架构扩展了GPU，超越了图形领域，其强大的多线程处理器阵列已经成为高效的统一平台，同时适用于图形和通用并行计算应用程序。
通过扩展处理器和存储器分区的数量，Tesla架构就延伸了市场覆盖率，从高性能发烧级GeForceGTX280GPU和专业Quadr与Tesla计算产品，一直到多种主流经济型GeForceGPU（在附录A中可查看所有支持CUDA的GPU的列表）。
其计算特性支持利用CUDA在C语言中直观地编写GPU核心程序。
Tesla架构具有在笔记本电脑、台式机、工作站和服务器上的广泛可用性，配以C语言编程能力和CUDA软件，使这种架构成为最优秀的超级计算平台。
这一章介绍了CUDA编程模型与Tesla架构的映射。
3.1具有芯片共享存储器的一组SIMT多处理器
Tesla架构的构建以一个可伸缩的多线程流处理器（SM）阵列为中心。
当主机CPU上的CUDA程序调用内核网格时，网格的块将被枚举并分发到具有可用执行容量的多处理器上。
一个线程块的线程在一个多处理器上并发执行。
在线程块终止时，将在空闲多处理器上启动新块。
多处理器包含8个标量处理器（SP）核心、两个用于先验（transcendental）的特殊函数单元、一个多线程指令单元以及芯片共享存储器。
多处理器会在硬件中创建、管理和执行并发线程，而调度开销保持为0。
它可通过一条内部指令实现_syncthreads（）_屏障同步。
快速的屏障同步与轻量级线程创建和零开销的线程调度相结合，有效地为细粒度并行化提供了支持，举例来说，您可以为各数据元素（如图像中的一个像素、语音中的一个语音元素、基于网格的计算中的一个单元）分配一个线程，从而对问题进行细粒度分解。
为了管理运行各种不同程序的数百个线程，多处理器利用了一种称为SIMT（单指令、多线程）的新架构。
多处理器会将各线程映射到一个标量处理器核心，各标量线程使用自己的指令地址和寄存器状态独立执行。
多处理器SIMT单元以32个并行线程为一组来创建、管理、调度和执行线程，这样的线程组称为warp块。
（此术语源于第一种并行线程技术weaving。
半warp块可以是一个warp块的第一半或第二半。
）构成SIMTwarp块的各个线程在同一个程序地址一起启动，但也可随意分支、独立执行。
为一个多处理器指定了一个或多个要执行的线程块时，它会将其分成warp块，并由SIMT单元进行调度。
将块分割为warp块的方法总是相同的，每个warp块都包含连续的线程，递增线程ID，第一个warp块中包含线程0。
第2.1节介绍了线程ID与块中的线程索引之间的关系。
每发出一条指令时，SIMT单元都会选择一个已准备好执行的warp块，并将下一条指令发送到该warp块的活动线程。
Warp块每次执行一条通用指令，因此在warp块的全部32个线程均认可其执行路径时，可达到最高效率。
如果一个warp块的线程通过独立于数据的条件分支而分散，warp块将连续执行所使用的各分支路径，而禁用未在此路径上的线程，完成所有路径时，线程重新汇聚到同一执行路径下。
分支仅在warp块内出现，不同的warp块总是独立执行的——无论它们执行的是通用的代码路径还是彼此无关的代码路径。
SIMT架构类似于SIMD（单指令、多数据）向量组织方法，共同之处是使用单指令来控制多个处理元素。
一项主要差别在于SIMD向量组织方法会向软件公开SIMD宽度，而SIMT指令指定单一线程的执行和分支行为。
与SIMD向量机不同，SIMT允许程序员为独立、标量线程编写线程级的并行代码，还允许为协同线程编写数据并行代码。
为了确保正确性，程序员可忽略SIMT行为，但通过维护很少需要使一个warp块内的线程分支的代码，即可实现显著的性能提升。
在实践中，这与传统代码中的超高速缓冲存储器线作用相似：
在以正确性为目标进行设计时，可忽略超高速缓冲存储器线的大小，但如果以峰值性能为目标进行设计，在代码结构中就必须考虑其大小。
另一方面，向量架构要求软件将负载并入向量，并手动管理分支。
如图3-1所示，每个多处理器都有一个属于以下四种类型之一的芯片存储器：
⏹每个处理器上有一组本地32位寄存器；
⏹并行数据缓存或共享存储器，由所有标量处理器核心共享，共享存储器空间就位于此处；
⏹只读固定缓存，由所有标量处理器核心共享，可加速从固定存储器空间进行的读取操作（这是设备存储器的一个只读区域）；
⏹一个只读纹理缓存，由所有标量处理器核心共享，加速从纹理存储器空间进行的读取操作（这是设备存储器的一个只读区域），每个多处理器都会通过实现不同寻址模型和数据过滤的纹理单元访问纹理缓存，相关内容请参见第4.3.4节。
本地和全局存储器空间是设备存储器的读/写区域，不应缓存。
一个多处理器一次可处理的块数量取决于每个线程有多少个寄存器、每个块需要多少共享存储器来支持给定的内核，这是因为多处理器的寄存器和共享存储器对于一批块的所有线程来说都是分离的。
如果没有足够的寄存器或共享存储器可供多处理器用于处理至少一个块，内核将启动失败。
一个多处理器可并发执行多达8个线程块。
如果warp块执行的非原子指令为warp块的多个线程写入全局或共享存储器中的同一位置，针对此位置的串行化写入操作的数量和这些写入操作所发生的顺序将无法确定，但其中一项操作必将成功。
如果warp块执行原子指令来为warp块的多个线程读取、修改和写入全局存储器中的同一位置，则针对该位置的每一项读取、修改或写入操作都将发生，且均为串行化操作，但这些操作所发生的顺序无法确定。
具有芯片共享存储器的一组SIMT多处理器
图3-1.硬件模型
3.2多个设备
若要多GPU系统上运行的应用程序将多个GPU作为CUDA设备使用，则这些GPU必须具有相同的类型。
但如果系统采用的是SLI模式，则仅有一个GPU可用作CUDA设备，因为所有GPU都将在驱动程序栈的最低级别融合。
要使CUDA能够将各GPU视为独立设备，需要在CUDA的控制面板内关闭SLI模式。
3.3模式切换
GPU将部分DRAM存储器专门用于处理所谓的主表面（primarysurface），它用于刷新显示设备，用户将查看该设备的输出。
当用户通过更改显示器的分辨率或位深度（使用NVIDIA的控制面板或Windows的显示控制面板）发起模式切换时，主表面所需的存储器数量而言会随之改变。
例如，如果用户将显示器的分辨率从1280x1024x32位更改为1600x1200x32位，系统必须为主表面分配7.68MB的存储器，而不是5.24MB。
（使用防锯齿设置运行的全屏图形应用程序可能需要为主表面分配更多显示存储器。
）在Windows上，其他事件也可能会启动显示模式切换，包括启动全屏DirectX应用程序、按Alt+Tab键从全屏DirectX应用程序中切换出来或者按Ctrl+Alt+Del键锁定计算机。
如果模式切换增加了主表面所需的存储器数量，系统可能就必须挪用分配给CUDA应用程序的存储器，从而导致此类应用程序崩溃。
第4章应用程序编程接口
4.1C编程语言的扩展
CUDA编程接口的目标是为熟悉C编程语言的用户提供相对简单的途径，使之可轻松编写由设备执行的程序。
它包含：
⏹C语言的最小扩展集，如4.2节所述，这允许程序员使源代码的某些部分可在设备上执行；
⏹一个运行时库，可分割为：
●一个主机组件，如4.5节所述，运行在主机上，提供函数来通过主机控制和访问一个或多个计算设备；
●一个设备组件，如4.4
升级会员 