计算统一设备架构

NVIDIA CUDA计算统一设备架构编程指南版本 2.06 / 7 / 2008目 录第 1 章简介11.1 CUDA：可伸缩并行编程模型11.2 GPU：高度并行化、多线程、多核处理器11.3 文档结构3第2章编程模型42.1 线程层次结构42.2 存储器层次结构62.3 主机和设备62.4 软件栈72.5 计算能力8第 3 章GPU 实现93.1 具有芯片共享存储器的一组 SIMT 多处理器93.2 多个设备113.3 模式切换11第 4 章应用程序编程接口124.1 C 编程语言的扩展124.2 语言扩展124.2.1 函数类型限定符124.2.1.1 _device_124.2.1.2 _global_134.2.1.3 _host_134.2.1.4 限制134.2.2 变量类型限定符134.2.2.1 _device_134.2.2.2 _constant_134.2.2.3 _shared_144.2.2.4 限制144.2.3 执行配置154.2.4 内置变量154.2.4.1 gridDim154.2.4.2 blockIdx154.2.4.3 blockDim154.2.4.4 threadIdx154.2.4.5 warpSize164.2.4.6 限制164.2.5 使用 NVCC 进行编译164.2.5.1 _noinline_164.2.5.2 #pragma unroll164.3 通用运行时组件174.3.1 内置向量类型174.3.1.1 char1、uchar1、char2、uchar2、char3、uchar3、char4、uchar4、short1、ushort1、short2、ushort2、short3、ushort3、short4、ushort4、int1、uint1、int2、uint2、int3、uint3、int4、uint4、long1、ulong1、long2、ulong2、long3、ulong3、long4、ulong4、float1、float2、float3、float4、double2174.3.1.2 dim3 类型174.3.2 数学函数174.3.3 计时函数174.3.4 纹理类型184.3.4.1 纹理参考声明184.3.4.2 运行时纹理参考属性184.3.4.3 来自线性存储器的纹理与来自 CUDA 数组的纹理194.4 设备运行时组件194.4.1 数学函数194.4.2 同步函数194.4.3 纹理函数194.4.3.1 来自线性存储器的纹理194.4.3.2 来自 CUDA 数组的纹理204.4.4 原子函数204.4.5 warp vote 函数204.5 主机运行时组件214.5.1 一般概念214.5.1.1 设备214.5.1.2 存储器224.5.1.3 OpenGL 互操作性224.5.1.4 Direct3D 互操作性224.5.1.5 异步并发执行224.5.2 运行时 API234.5.2.1 初始化234.5.2.2 设备管理234.5.2.3 存储器管理244.5.2.4 流管理254.5.2.5 事件管理254.5.2.6 纹理参考管理254.5.2.7 OpenGL 互操作性274.5.2.8 Direct3D 互操作性274.5.2.9 使用设备模拟模式进行调试284.5.3 驱动程序 API294.5.3.1 初始化294.5.3.2 设备管理294.5.3.3 上下文管理294.5.3.4 模块管理304.5.3.5 执行控制304.5.3.6 存储器管理314.5.3.7 流管理324.5.3.8 事件管理324.5.3.9 纹理参考管理334.5.3.10 OpenGL 互操作性334.5.3.11 Direct3D 互操作性33第 5 章性能指南355.1 指令性能355.1.1指令吞吐量355.1.1.1 数学指令355.1.1.2 控制流指令365.1.1.3 存储器指令365.1.1.4 同步指令375.1.2 存储器带宽375.1.2.1 全局存储器375.1.2.2 本地存储器435.1.2.3 固定存储器435.1.2.4 纹理存储器435.1.2.5 共享存储器435.1.2.6 寄存器485.2 每个块的线程数量495.3 主机和设备间的数据传输495.4 纹理获取与全局或固定存储器读取的对比505.5 整体性能优化战略50第 6 章矩阵乘法示例526.1 概述526.2 源代码清单536.3 源代码说明546.3.1 Mul()546.3.2 Muld()54附录 A技术规范56A.1 一般规范56A.1.1 计算能力 1.0 的规范56A.1.2 计算能力 1.1 的规范57A.1.3 计算能力 1.2 的规范57A.1.4 计算能力 1.3 的规范57A.2 浮点标准57附录 B标准数学函数59B.1 一般运行时组件59B.1.1 单精度浮点函数59B.1.2 双精度浮点函数60B.1.3 整型函数62B.2 设备运行时组件62B.2.1 单精度浮点函数62B.2.2 双精度浮点函数63B.2.3 整型函数64附录 C原子函数65C.1 数学函数65C.1.1 atomicAdd()65C.1.2 atomicSub()65C.1.3 atomicExch()65C.1.4 atomicMin()65C.1.5 atomicMax()66C.1.6 atomicInc()66C.1.7 atomicDec()66C.1.8 atomicCAS()66C.2 位逻辑函数66C.2.1 atomicAnd()66C.2.2 atomicOr()67C.2.3 atomicXor()67附录 D纹理获取68D.1 最近点取样68D.2 线性过滤69D.3 表查找69图表目录图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. 矩阵乘法. . . .52CUDA 编程指南，版本 2.0 vi 编号：时间：2021年x月x日书山有路勤为径，学海无涯苦作舟页码：第75页 共81页第 1 章简介1.1 CUDA：可伸缩并行编程模型多核 CPU 和多核 GPU 的出现意味着并行系统已成为主流处理器芯片。此外，根据摩尔定律，其并行性将不断扩展。这带来了严峻的挑战，我们需要开发出可透明地扩展并行性的应用软件，以便利用日益增加的处理器内核数量，这种情况正如 3D 图形应用程序透明地扩展其并行性以支持配备各种数量的内核的多核 GPU。CUDA 是一种并行编程模型和软件环境，用于应对这种挑战，同时保证熟悉 C 语言等标准编程语言的程序员能够迅速掌握 CUDA。CUDA 的核心有三个重要抽象概念：线程组层次结构、共享存储器、屏蔽同步（barrier synchronization），可轻松将其作为 C 语言的最小扩展级公开给程序员。这些抽象提供了细粒度的数据并行化和线程并行化，嵌套于粗粒度的数据并行化和任务并行化之中。它们将指导程序员将问题分解为更小的片段，以便通过协作的方法并行解决。这样的分解保留了语言表达，允许线程在解决各子问题时协作，同时支持透明的可伸缩性，使您可以安排在任何可用处理器内核上处理各子问题：因而，编译后的 CUDA 程序可以在任何数量的处理器内核上执行，只有运行时系统需要了解物理处理器数量。1.2 GPU：高度并行化、多线程、多核处理器市场迫切需要实时、高清晰度的 3D 图形，可编程的 GPU 已发展成为一种高度并行化、多线程、多核的处理器，具有杰出的计算功率和极高的存储器带宽，如图 1-1 所示。图1-1. CPU 和 GPU 的每秒浮点运算次数和存储器带宽CPU 和 GPU 之间浮点功能之所以存在这样的差异，原因就在于 GPU 专为计算密集型、高度并行化的计算而设计，上图显示的正是这种情况，因而，GPU 的设计能使更多晶体管用于数据处理，而非数据缓存和流控制，如图 1-2 所示。图 1-2. GPU 中的更多晶体管用于数据处理更具体地说，GPU 专用于解决可表示为数据并行计算的问题在许多数据元素上并行执行的程序，具有极高的计算密度（数学运算与存储器运算的比率）。由于所有数据元素都执行相同的程序，因此对精密流控制的要求不高；由于在许多数据元素上运行，且具有较高的计算密度，因而可通过计算隐藏存储器访问延迟，而不必使用较大的数据缓存。数据并行处理会将数据元素映射到并行处理线程。许多处理大型数据集的应用程序都可使用数据并行编程模型来加速计算。在 3D 渲染中，大量的像素和顶点集将映射到并行线程。类似地，图像和媒体护理应用程序（如渲染图像的后期处理、视频编码和解码、图像缩放、立体视觉和模式识别等）可将图像块和像素映射到并行处理线程。实际上，在图像渲染和处理领域之外的许多算法也都是通过数据并行处理加速的从普通信号处理或物理仿真一直到数理金融或数理生物学。CUDA 编程模型非常适合公开 GPU 的并行功能。最新一代的 NVIDIA GPU 基于 Tesla 架构（在附录 A 中可以查看所有支持 CUDA 的 GPU 列表），支持 CUDA 编程模型，可显著加速 CUDA 应用程序。1.3 文档结构本文档分为以下几个章节：n 第 1 章是 CUDA 和 GPU 的简介。n 第 2 章概述 CUDA 编程模型。n 第 3 章介绍 GPU 实现。n 第 4 章介绍 CUDA API 和运行时。n 第 5 章提供如何实现最高性能的一些指南。n 第 6 章通过一些简单的示例代码概况之前各章的内容。n 附录 A 提供各种设备的技术规范。n 附录 B 列举 CUDA 中支持的数学函数。n 附录 C 列举 CUDA 中支持的原子函数。n 附录 D 详细说明纹理获取。第2章编程模型CUDA 允许程序员定义称为内核（kernel）的 C 语言函数，从而扩展了 C 语言，在调用此类函数时，它将由 N 个不同的 CUDA 线程并行执行 N 次，这与普通的 C 语言函数只执行一次的方式不同。在定义内核时，需要使用 _global_ 声明说明符，使用一种全新的 语法指定每次调用的 CUDA 线程数：/ Kernel definition_global_ void vecAdd(float* A, float* B, float* C)int main() / Kernel invocation vecAdd(A, B, C);执行内核的每个线程都会被分配一个独特的线程 ID，可通过内置的 threadIdx 变量在内核中访问此 ID。以下示例代码将大小为 N 的向量 A 和向量 B 相加，并将结果存储在向量 C 中：_global_ void vecAdd(float* A, float* B, float* C) int i = threadIdx.x; Ci = Ai + Bi;int main() / Kernel invocation vecAdd(A, B, C);执行 vecAdd( ) 的每个线程都会执行一次成对的加法运算。2.1 线程层次结构为方便起见，我们将 threadIdx 设置为一个包含 3 个组件的向量，因而可使用一维、二维或三维缩影标识线程，构成一维、二维或三维线程块。这提供了一种自然的方法，可为一个域中的各元素调用计算，如向量、矩阵或字段。下面的示例代码将大小为 NxN 的矩阵 A 和矩阵 B 相加，并将结果存储在矩阵 C 中：_global_ void matAdd(float ANN, float BNN,float CNN) int i = threadIdx.x; int j = threadIdx.y; Cij = Aij + Bij;int main() / Kernel invocation dim3 dimBlock(N, N); matAdd(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()_ 应是轻量级的，一个块中的所有线程都必须位于同一个处理器核心中。因而，一个处理器核心的有限存储器资源制约了每个块的线程数量。在 NVIDIA Tesla 架构中，一个线程块最多可以包含 512 个线程。但一个内核可能由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数量。这些块将组织为一个一维或二维线程块网格，如图 2-1 所示。该网格的维度由 语法的第一个参数指定。网格内的每个块多可由一个一维或二维索引标识，可通过内置的 blockIdx 变量在内核中访问此索引。可以通过内置的 blockDim 变量在内核中访问线程块的维度。此时，之前的示例代码应修改为：_global_ void matAdd(float ANN, float BNN,float CNN) int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; if (i N & j N) Cij = Aij + Bij;int main() / Kernel invocation dim3 dimBlock(16, 16); dim3 dimGrid(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 计算能力一个设备的计算能力（compute capability）由主要修订号和次要修订号定义。具有相同主要修订号的设备属于相同的核心架构。附录 A 中列举的设备均为计算能力是 1.x 的设备（其主要修订号为 1）。次要修订号对应于核心架构的增量式改进，可能包含新特性。附录 A 提供了各种计算能力的技术规范。第 3 章GPU 实现NVIDIA 于 2006 年 11 月引入的 Tesla 统一图形和计算架构扩展了 GPU，超越了图形领域，其强大的多线程处理器阵列已经成为高效的统一平台，同时适用于图形和通用并行计算应用程序。通过扩展处理器和存储器分区的数量，Tesla 架构就延伸了市场覆盖率，从高性能发烧级 GeForce GTX 280 GPU 和专业 Quadr 与 Tesla 计算产品，一直到多种主流经济型 GeForce GPU（在附录 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 块的第一半或第二半。）构成 SIMT warp 块的各个线程在同一个程序地址一起启动，但也可随意分支、独立执行。为一个多处理器指定了一个或多个要执行的线程块时，它会将其分成 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 所示，每个多处理器都有一个属于以下四种类型之一的芯片存储器：n 每个处理器上有一组本地 32 位寄存器；n 并行数据缓存或共享存储器，由所有标量处理器核心共享，共享存储器空间就位于此处；n 只读固定缓存，由所有标量处理器核心共享，可加速从固定存储器空间进行的读取操作（这是设备存储器的一个只读区域）；n 一个只读纹理缓存，由所有标量处理器核心共享，加速从纹理存储器空间进行的读取操作（这是设备存储器的一个只读区域），每个多处理器都会通过实现不同寻址模型和数据过滤的纹理单元访问纹理缓存，相关内容请参见第 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 存储器专门用于处理所谓的主表面（primary surface），它用于刷新显示设备，用户将查看该设备的输出。当用户通过更改显示器的分辨率或位深度（使用 NVIDIA 的控制面板或 Windows 的显示控制面板）发起模式切换时，主表面所需的存储器数量而言会随之改变。例如，如果用户将显示器的分辨率从 1280x1024x32 位更改为 1600x1200x32 位，系统必须为主表面分配 7.68 MB 的存储器，而不是 5.24 MB。（使用防锯齿设置运行的全屏图形应用程序可能需要为主表面分配更多显示存储器。）在 Windows 上，其他事件也可能会启动显示模式切换，包括启动全屏 DirectX 应用程序、按 Alt+Tab 键从全屏 DirectX 应用程序中切换出来或者按 Ctrl+Alt+Del 键锁定计算机。如果模式切换增加了主表面所需的存储器数量，系统可能就必须挪用分配给 CUDA 应用程序的存储器，从而导致此类应用程序崩溃。第 4 章应用程序编程接口4.1 C 编程语言的扩展CUDA 编程接口的目标是为熟悉 C 编程语言的用户提供相对简单的途径，使之可轻松编写由设备执行的程序。它包含：n C 语言的最小扩展集，如 4.2 节所述，这允许程序员使源代码的某些部分可在设备上执行；n 一个运行时库，可分割为：l 一个主机组件，如 4.5 节所述，运行在主机上，提供函数来通过主机控制和访问一个或多个计算设备；l 一个设备组件，如 4.4 节所述，运行在设备上，提供特定于设备的函数；l 一个通用组件，如 4.3 节所述，提供内置向量类型和 C 标准库的一个子集，主机和设备代码中都将支持此子集。有必要强调，C 标准库中支持在设备上运行的函数只有通用运行时组件所提供的函数。4.2 语言扩展对 C 编程语言的扩展共有四重：n 函数类型限定符，指定函数是在主机上还是设备上执行，以及函数是可通过主机还是可通过设备调用（参见第 4.2.1 节）；n 变量类型限定符，指定一个变量在设备上的存储器位置（参见第 4.2.2 节）；n 一条新指令，指定如何通过主机在设备上执行内核（参见第 4.2.3 节）；n 四个内置变量，指定网格和块维度以及块和线程索引（参见第 4.2.4 节）。包含这些扩展的所有源文件都必须使用 CUDA 编译器 nvcc 进行编译，4.2.5 节简单介绍了相关内容。关于 nvcc 的具体介绍将在其他文档中提供。这些扩展均具有一些限制，下面几个小节将分别加以介绍。如果违背了这些限制，nvcc 将发出错误或警报信息，但有些违规情况无法检测到。4.2.1 函数类型限定符4.2.1.1 _device_使用 _device_ 限定符声明的函数具有以下特征：n 在设备上执行；n 仅可通过设备调用。4.2.1.2 _global_使用 _global_ 限定符可将函数声明为内核。此类函数：n 在设备上执行；n 仅可通过主机调用。4.2.1.3 _host_使用 _host_ 限定符声明的函数具有以下特征：n 在主机上执行；n 仅可通过主机调用。仅使用 _host_ 限定符声明函数等同于不使用 _host_、_device_ 或 _global_ 限定符声明函数，这两种情况下，函数都将仅为主机进行编译。但 _host_ 限定符也可与 _device_ 限定符一起使用，此时函数将为主机和设备进行编译。4.2.1.4 限制_device_ 和 _global_ 函数不支持递归。_device_ 和 _global_ 函数的函数体内无法声明静态变量。_device_ 和 _global_ 函数不得有数量可变的参数。_device_ 函数的地址无法获取，但支持 _global_ 函数的函数指针。_global_ 和 _host_ 限定符无法一起使用。_global_ 函数的返回类型必须为空。对 _global_ 函数的任何调用都必须按第 4.2.3 节介绍的方法指定其执行配置。_global_ 函数的调用是异步的，也就是说它会在设备执行完成之前返回。_global_ 函数参数将同时通过共享存储器传递给设备，且限制为 256 字节。4.2.2 变量类型限定符4.2.2.1 _device_device_ 限定符声明位于设备上的变量。在接下来的三节中介绍的其他类型限定符中，最多只能有一种可与 _device_ 限定符一起使用，以更具体地指定变量属于哪个存储器空间。如果未出现其他任何限定符，则变量具有以下特征：n 位于全局存储器空间中；n 与应用程序具有相同的生命周期；n 可通过网格内的所有线程访问，也可通过运行时库从主机访问。4.2.2.2 _constant_constant_ 限定符可选择与 _device_ 限定符一起使用，所声明的变量具有以下特征：n 位于固定存储器空间中；n 与应用程序具有相同的生命周期；n 可通过网格内的所有线程访问，也可通过运行时库从主机访问。4.2.2.3 _shared_shared_ 限定符可选择与 _device_ 限定符一起使用，所声明的变量具有以下特征：n 位于线程块的共享存储器空间中；n 与块具有相同的生命周期；n 尽可通过块内的所有线程访问。只有在 _syncthreads()_（参见第 4.4.2 节）的执行写入之后，才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量，否则只要之前的语句完成，编译器即可随意优化共享存储器的读写操作。将共享存储器中的变量声明为外部数组时，例如：extern _shared_ float shared;数组的大小将在启动时确定（参见第 4.2.3 节）。所有变量均以这种形式声明，在存储器中的同一地址开始，因此数组中的变量布局必须通过偏移显式管理。例如，如果一名用户希望在动态分配的共享存储器内获得与以下代码对应的内容：short array0128;float array164;int array2256;则应通过以下方法声明和初始化数组：extern _shared_ char array;_device_ void func() / _device_ or _global_ function short* array0 = (short*)array; float* array1 = (float*)&array0128; int* array2 = (int*)&array164;4.2.2.4 限制不允许为在主机上执行的函数内的 struct 和 union 成员、形参和局部变量使用这些限定符。_shared_ 和 _constant_ 变量具有隐含的静态存储。_device_、_shared_ 和 _constant_ 变量无法使用 extern 关键字定义为外部变量。_device_ 和 _constant_ 变量仅允许在文件作用域内使用。不可为设备或从设备指派 _constant_ 变量，仅可通过主机运行时函数从主机指派（参见第 4.5.2.3 节和第 4.5.3.6 节）。_shared_ 变量的声明中不可包含初始化。在设备代码中声明、不带任何限定符的自动变量通常位于寄存器中。但在某些情况下，编译器可能选择将其置于本地存储器中。如果使用占用了过多寄存器空间的大型结构或数组，或者编译器无法确定其是否使用固定数量索引的数组，则往往会出现这种情况。检查 ptx 汇编代码（通过使用 ptx 或 keep 选项编译获得）即可在初次编译过程中确定一个变量是否位于本地存储器中，因为它将使用 .local 助记符声明，可使用 ld.local 和 st.local 助记符访问。如果不是这样，在后续编译阶段仍能确定是否占用了目标架构的过多寄存器空间。可通过使用 -ptxas- options =-v 选项编译来进行检查，这将报告本地存储器的使用情况（lmem）。只要编译器能够确定在设备上执行的代码中的指针指向的是共享存储器空间还是全局存储器空间，此类指针即受支持，否则将仅限于指向在全局存储器空间中分配或声明的存储器。如果取消在主机上执行的代码中全局或共享存储器指针，或者在设备上执行的代码中主机存储器指针的引用，将导致不确定的行为，往往会出现分区错误和应用程序终止。通过获取 _device_、_shared_ 或 _constant_ 变量的地址而获得的地址仅可在设备代码中使用。通过 cudaGetSymbolAddress() （参见第 4.5.23 节）获取的 _device_ 或 _constant_ 变量的地址仅可在主机代码中使用。4.2.3 执行配置对 _global_ 函数的任何调用都必须指定该调用的执行配置。执行配置定义将用于在该设备上执行函数的网格和块的维度，以及相关的流（有关流的内容将在第 4.5.1.5 节介绍）。可通过在函数名称和括号参数列表之间插入 形式的表达式来指定，其中：n Dg 的类型为 dim3（参见第 4.3.1.2 节），指定网格的维度和大小，Dg.x * Dg.y 等于所启动的块数量，Dg.z 无用；n Db 的类型为 dim3（参见第 4.3.1.2 节），指定各块的维度和大小，Db.x * Db.y * Db.z 等于各块的线程数量；n Ns 的类型为 size_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为外部数组的其他任何变量使用（参见第 4.2.2.3 节），Ns 是一个可选参数，默认值为 0；n S 的类型为 cudaStream_t，指定相关流；S 是一个可选参数，默认值为 0。举例来说，一个函数的声明如下：_global_ void Func(float* parameter);必须通过如下方法来调用此函数：Func(parameter);执行配置的参数将在实际函数参数之前被评估，与函数参数相同，通过共享存储器同时传递给设备。如果 Dg 或 Db 大于设备允许的最大大小（参见附录 A.1.1），或 Ns 大于设备上可用的共享存储器最大值，或者小于静态分配、函数参数和执行配置所需的共享存储器数量，则函数将失败。4.2.4 内置变量4.2.4.1 gridDim此变量的类型为 dim3（参见第 4.3.1.2 节），包含网格的维度。4.2.4.2 blockIdx此变量的类型为 uint3（参见第 4.3.1.1 节），包含网格内的块索引。4.2.4.3 blockDim此变量的类型为 dim3（参见第 4.3.1.2 节），包含块的维度。4.2.4.4 threadIdx此变量的类型为 uint3（参见第 4.3.1.1 节），包含块内的线程索引。4.2.4.5 warpSize此变量的类型为 int，包含以线程为单位的 warp 块大小。4.2.4.6 限制n 不允许接受任何内置变量的地址。n 不允许为任何内置变量赋值。4.2.5 使用 NVCC 进行编译Nvcc 是一种可简化 CUDA 代码编译过程的编译器驱动程序：它提供了简单、熟悉的命令行选项，通过调用实现不同编译阶段的工具集合来执行它们。Nvcc 的基本工作流在于将设备代码与主机代码分离开来，并将设备代码编译为二进制形式或 cubin 对象。所生成的主机代码将作为需要使用其他工具编译的 C 代码输出，或通过在最后一个编译阶段中调用主机编译器直接作为对象代码输出。应用程序可忽略所生成的主机代码，使用 CUDA 驱动程序 API在设备上加载并执行 cubin 对象，也可链接到所生成的主机代码，其中包含 cubin 对象，其形式为全局初始化数据数组，包含将第 4.2.3 节所述执行配置语法转换为必要的 CUDA 运行启动代码的转换，目的在于加载和启动编译后的各内核（参见第 4.5.2 节）。编译器的前端根据 C+ 语法规则处理 CUDA 源文件。主机代码支持完整的 C+ 语法。但设备代码仅支持 C+ 的 C 子集，类、继承、基本块内的变量声明等 C+ 特殊特性不受支持。由于使用了 C+ 语法规则，因此若未经过强制类型转换，无法将空指针（例如 malloc() 所返回的空指针）指派给非空指针。关于 nvcc 工作流和命令选项的详细说明将在其他文档中提供。Nvcc 引入了两个编译器指令，下面几节将加以介绍。4.2.5.1 _noinline_默认情况下，_device_ 函数总是内嵌的。_noinline_ 函数限定符可用于指示编译器尽可能不要内嵌该函数。函数体必须位于所调用的同一个文件内。如果函数具有指针参数或者具有较大的参数列表，则编译器不会遵从 _noinline_ 限定符。4.2.5.2 #pragma unroll默认情况下，编译器将展开具有已知行程计数的小循环。#pragma unroll 指令可用于控制任何给定循环的展开操作。它必须紧接于循环之前，而且仅应用于该循环。可选择在其后接一个数字，指定必须展开多少次循环。例如，在下面的代码示例中：#pragma unroll 5for (int i = 0; i n; +i)循环将展开 5 次。程序员需要负责确保展开操作不会影响程序的正确性（在上面的示例中，如果 n 小于 5，则程序的正确性将受到影响）。#pragma unroll 1 将阻止编译器展开一个循环。如果在 #pragma unroll 后未指定任何数据，如果其行程计数为常数，则该循环将完全展开，否则将不会展开。4.3 通用运行时组件主机和设备函数均可使用通用运行时组件。4.3.1 内置向量类型4.3.1.1 char1、uchar1、char2、uchar2、char3、uchar3、char4、uchar4、short1、ushort1、short2、ushort2、short3、ushort3、short4、ushort4、int1、uint1、int2、uint2、int3、uint3、int4、uint4、long1、ulong1、long2、ulong2、long3、ulong3、long4、ulong4、float1、float2、float3、float4、double2这些向量类型继承自基本整形和浮点类型。它们均为结构体，第 1、2、3、4 个组件分别可通过字段 x、y、z 和 w 访问。它们均附带形式为 make_ 的构造函数，示例如下：int2 make_int2(int x, int y);这将创建一个类型为 int2 的向量，值为 (x, y)。4.3.1.2 dim3 类型此类型是一种整形向量类型，基于用于指定维度的 uint3。在定义类型为 dim3 的变量时，未指定的任何组件都将初始化为 1。4.3.2 数学函数B.1 节包含了当前支持的 C/C+ 标准库数学函数的完整列表，还分别给出了在设备上执行时的误差范围。在主机代码中执行时，给定函数将在可用的前提下使用 C 运行时实现。4.3.3 计时函数clock_t clock();在设备代码中执行时，返回随每一次时钟周期而递增的每个多处理器计数器的值。在内核启动和结束时对此计数器取样，确定两次取样的差别，然后为每个线程记录下结果，这为各线程提供一种度量方法，可度量设备为了完全执行线程而占用的时钟周期数，但不是设备在执行线程指令时而实际使用的时钟周期数。前一个数字要比后一个数字大得多，因为线程是分时的。4.3.4 纹理类型CUDA 支持 GPU 用于图形的纹理硬件子集，使之可访问纹理存储器。从纹理存储器而非全局存储器读取数据可带来多方面的性能收益，请参见第 5.4 节。内核使用称为纹理获取（texture fetch）的设备函数读取纹理存储器，请参见第 4.4.3 节。纹理获取的第一个参数指定称为纹理参考的对象。纹理参考定义获取哪部分的纹理存储器。必须通过主机运行时函数（参见第 4.5.2.6 和第 4.5.3.9 节）将其绑定到存储器的某些区域（即纹理），之后才能供内核使用。多个不同的纹理参考可绑定到同一个纹理，也可绑定到在存储器中存在重叠的纹理。纹理参考有一些属性。其中之一就是其维度，指定纹理是使用一个纹理坐标（texture coordinate）将纹理作为一维数组寻址、使用两个纹理坐标作为二维数组寻址，还是使用三个纹理坐标作为三维数组寻址。数