计算统一设备架构.docx

《计算统一设备架构.docx》由会员分享，可在线阅读，更多相关《计算统一设备架构.docx（78页珍藏版）》请在三一办公上搜索。

1、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、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 #pr

3、agma 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

4、 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 Ope

5、nGL 互操作性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 执行控制3

6、04.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

7、 每个块的线程数量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 双精度浮点函数6

8、0B.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 atomi

9、cXor()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

10、.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 编程指南，版

11、本 2.0 71 第 1 章简介1.1 CUDA：可伸缩并行编程模型多核 CPU 和多核 GPU 的出现意味着并行系统已成为主流处理器芯片。此外，根据摩尔定律，其并行性将不断扩展。这带来了严峻的挑战，我们需要开发出可透明地扩展并行性的应用软件，以便利用日益增加的处理器内核数量，这种情况正如 3D 图形应用程序透明地扩展其并行性以支持配备各种数量的内核的多核 GPU。CUDA 是一种并行编程模型和软件环境，用于应对这种挑战，同时保证熟悉 C 语言等标准编程语言的程序员能够迅速掌握 CUDA。CUDA 的核心有三个重要抽象概念：线程组层次结构、共享存储器、屏蔽同步（barrier synchron

12、ization），可轻松将其作为 C 语言的最小扩展级公开给程序员。这些抽象提供了细粒度的数据并行化和线程并行化，嵌套于粗粒度的数据并行化和任务并行化之中。它们将指导程序员将问题分解为更小的片段，以便通过协作的方法并行解决。这样的分解保留了语言表达，允许线程在解决各子问题时协作，同时支持透明的可伸缩性，使您可以安排在任何可用处理器内核上处理各子问题：因而，编译后的 CUDA 程序可以在任何数量的处理器内核上执行，只有运行时系统需要了解物理处理器数量。1.2 GPU：高度并行化、多线程、多核处理器市场迫切需要实时、高清晰度的 3D 图形，可编程的 GPU 已发展成为一种高度并行化、多线程、多核的

13、处理器，具有杰出的计算功率和极高的存储器带宽，如图 1-1 所示。图1-1. CPU 和 GPU 的每秒浮点运算次数和存储器带宽CPU 和 GPU 之间浮点功能之所以存在这样的差异，原因就在于 GPU 专为计算密集型、高度并行化的计算而设计，上图显示的正是这种情况，因而，GPU 的设计能使更多晶体管用于数据处理，而非数据缓存和流控制，如图 1-2 所示。图 1-2. GPU 中的更多晶体管用于数据处理更具体地说，GPU 专用于解决可表示为数据并行计算的问题在许多数据元素上并行执行的程序，具有极高的计算密度（数学运算与存储器运算的比率）。由于所有数据元素都执行相同的程序，因此对精密流控制的要求不

14、高；由于在许多数据元素上运行，且具有较高的计算密度，因而可通过计算隐藏存储器访问延迟，而不必使用较大的数据缓存。数据并行处理会将数据元素映射到并行处理线程。许多处理大型数据集的应用程序都可使用数据并行编程模型来加速计算。在 3D 渲染中，大量的像素和顶点集将映射到并行线程。类似地，图像和媒体护理应用程序（如渲染图像的后期处理、视频编码和解码、图像缩放、立体视觉和模式识别等）可将图像块和像素映射到并行处理线程。实际上，在图像渲染和处理领域之外的许多算法也都是通过数据并行处理加速的从普通信号处理或物理仿真一直到数理金融或数理生物学。CUDA 编程模型非常适合公开 GPU 的并行功能。最新一代的 N

15、VIDIA 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 中支持的

16、原子函数。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)

17、;执行内核的每个线程都会被分配一个独特的线程 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 设置为一个包含

18、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,

19、 C);线程的索引及其线程 ID 有着直接的关系：对于一维块来说，两者是相同的；对于大小为 (Dx，Dy) 的二维块来说，索引为 (x,y) 的线程的ID 是 （x + yDx）；对于大小为 （Dx，Dy， Dz） 的三维块来说，索引为 （x, y, z） 的线程的ID 是 （x + yDx + ZDxDy）。一个块内的线程可彼此协作，通过一些共享存储器来共享数据，并同步其执行来协调存储器访问。更具体地说，可以通过调用 _syncthreads()_ 内函数在内核中指定同步点；_syncthreads()_ 起到屏障的作用，块中的所有线程都必须在这里等待处理。为实现有效的协作，共享存储器应该是

20、接近各处理器核心的低延迟存储器，如 L1 缓存，_syncthreads()_ 应是轻量级的，一个块中的所有线程都必须位于同一个处理器核心中。因而，一个处理器核心的有限存储器资源制约了每个块的线程数量。在 NVIDIA Tesla 架构中，一个线程块最多可以包含 512 个线程。但一个内核可能由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数量。这些块将组织为一个一维或二维线程块网格，如图 2-1 所示。该网格的维度由 语法的第一个参数指定。网格内的每个块多可由一个一维或二维索引标识，可通过内置的 blockIdx 变量在内核中访问此索引。可以通过内置的 blockDim

21、变量在内核中访问线程块的维度。此时，之前的示例代码应修改为：_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) / dimBl

22、ock.x, (N + dimBlock.y 1) / dimBlock.y); matAdd(A, B, C);我们随机选择了大小为 16x16 的线程块（即包含 256 个线程），此外创建了一个网格，它具有足够的块，可将每个线程作为一个矩阵元素，这与之前完全相同。线程块需要独立执行：必须能够以任意顺序执行、能够并行或顺序执行。这种独立性需求允许跨任意数量的核心安排线程块，从而使程序员能够编写出可伸缩的代码。一个网格内的线程块数量通常是由所处理的数据大小限定的，而不是由系统中的处理器数量决定的，前者可能远远超过后者的数量。图 2-1. 线程块网格2.2 存储器层次结构CUDA 线程可在执行过

23、程中访问多个存储器空间的数据，如图 2-2 所示。每个线程都有一个私有的本地存储器。每个线程块都有一个共享存储器，该存储器对于块内的所有线程都是可见的，并且与块具有相同的生命周期。最终，所有线程都可访问相同的全局存储器。此外还有两个只读的存储器空间，可由所有线程访问，这两个空间是固定存储器空间和纹理存储器空间。全局、固定和纹理存储器空间经过优化，适于不同的存储器用途（参见第 5.1.2.1、5.1.2.3 和 5.1.2.4）。纹理存储器也为某些特殊的数据格式提供了不同的寻址模式以及数据过滤（参见第 4.3.4）。对于同一个应用程序启动的内核而言，全局、固定和纹理存储器空间都是持久的。图 2-

24、2. 存储器层次结构2.3 主机和设备如图 2-3 所示，CUDA 假设 CUDA 线程可在物理上独立的设备上执行，此类设备作为运行 C 语言程序的主机的协同处理器操作。例如，当内核在 GPU 上执行，而 C 语言程序的其他部分在 CPU 上执行时，就是这样一种情况。此外，CUDA 还假设主机和设备均维护自己的 DRAM，分别称为主机存储器和设备存储器。因而，一个程序通过调用 CUDA 运行时来管理对内核可见的全局、固定和纹理存储器空间（详见第 4 章）。这包括设备存储器分配和取消分配，还包括主机和设备存储器之间的数据传输。串行代码在主机上执行，而并行代码在设备上执行。图2-3. 异构编程2.

25、4 软件栈CUDA 软件栈包含多个层，如图 2-4 所示：设备驱动程序、应用程序编程接口（API）及其运行时、两个较高级别的通用数学库，即 CUFFT 和 CUBLAS，这些内容均在其他文档中介绍。图 2-4. 计算统一设备架构软件栈2.5 计算能力一个设备的计算能力（compute capability）由主要修订号和次要修订号定义。具有相同主要修订号的设备属于相同的核心架构。附录 A 中列举的设备均为计算能力是 1.x 的设备（其主要修订号为 1）。次要修订号对应于核心架构的增量式改进，可能包含新特性。附录 A 提供了各种计算能力的技术规范。第 3 章GPU 实现NVIDIA 于 2006

26、 年 11 月引入的 Tesla 统一图形和计算架构扩展了 GPU，超越了图形领域，其强大的多线程处理器阵列已经成为高效的统一平台，同时适用于图形和通用并行计算应用程序。通过扩展处理器和存储器分区的数量，Tesla 架构就延伸了市场覆盖率，从高性能发烧级 GeForce GTX 280 GPU 和专业 Quadr 与 Tesla 计算产品，一直到多种主流经济型 GeForce GPU（在附录 A 中可查看所有支持 CUDA 的 GPU 的列表）。其计算特性支持利用 CUDA 在 C 语言中直观地编写 GPU 核心程序。Tesla 架构具有在笔记本电脑、台式机、工作站和服务器上的广泛可用性，配以

27、 C 语言编程能力和 CUDA 软件，使这种架构成为最优秀的超级计算平台。这一章介绍了 CUDA 编程模型与 Tesla 架构的映射。3.1 具有芯片共享存储器的一组 SIMT 多处理器Tesla 架构的构建以一个可伸缩的多线程流处理器（SM）阵列为中心。当主机 CPU 上的 CUDA 程序调用内核网格时，网格的块将被枚举并分发到具有可用执行容量的多处理器上。一个线程块的线程在一个多处理器上并发执行。在线程块终止时，将在空闲多处理器上启动新块。多处理器包含 8 个标量处理器（SP）核心、两个用于先验（transcendental）的特殊函数单元、一个多线程指令单元以及芯片共享存储器。多处理器会

28、在硬件中创建、管理和执行并发线程，而调度开销保持为0。它可通过一条内部指令实现 _syncthreads()_ 屏障同步。快速的屏障同步与轻量级线程创建和零开销的线程调度相结合，有效地为细粒度并行化提供了支持，举例来说，您可以为各数据元素（如图像中的一个像素、语音中的一个语音元素、基于网格的计算中的一个单元）分配一个线程，从而对问题进行细粒度分解。为了管理运行各种不同程序的数百个线程，多处理器利用了一种称为 SIMT（单指令、多线程）的新架构。多处理器会将各线程映射到一个标量处理器核心，各标量线程使用自己的指令地址和寄存器状态独立执行。多处理器 SIMT 单元以 32 个并行线程为一组来创建、

29、管理、调度和执行线程，这样的线程组称为 warp 块。（此术语源于第一种并行线程技术 weaving。半 warp 块可以是一个 warp 块的第一半或第二半。）构成 SIMT warp 块的各个线程在同一个程序地址一起启动，但也可随意分支、独立执行。为一个多处理器指定了一个或多个要执行的线程块时，它会将其分成 warp 块，并由 SIMT 单元进行调度。将块分割为 warp 块的方法总是相同的，每个 warp 块都包含连续的线程，递增线程 ID，第一个 warp 块中包含线程 0。第 2.1 节介绍了线程 ID 与块中的线程索引之间的关系。每发出一条指令时，SIMT 单元都会选择一个已准备好

30、执行的 warp 块，并将下一条指令发送到该 warp 块的活动线程。Warp 块每次执行一条通用指令，因此在 warp 块的全部 32 个线程均认可其执行路径时，可达到最高效率。如果一个 warp 块的线程通过独立于数据的条件分支而分散，warp 块将连续执行所使用的各分支路径，而禁用未在此路径上的线程，完成所有路径时，线程重新汇聚到同一执行路径下。分支仅在 warp 块内出现，不同的 warp 块总是独立执行的无论它们执行的是通用的代码路径还是彼此无关的代码路径。SIMT 架构类似于 SIMD（单指令、多数据）向量组织方法，共同之处是使用单指令来控制多个处理元素。一项主要差别在于 SIMD

31、 向量组织方法会向软件公开 SIMD 宽度，而 SIMT 指令指定单一线程的执行和分支行为。与 SIMD 向量机不同，SIMT 允许程序员为独立、标量线程编写线程级的并行代码，还允许为协同线程编写数据并行代码。为了确保正确性，程序员可忽略 SIMT 行为，但通过维护很少需要使一个 warp 块内的线程分支的代码，即可实现显著的性能提升。在实践中，这与传统代码中的超高速缓冲存储器线作用相似：在以正确性为目标进行设计时，可忽略超高速缓冲存储器线的大小，但如果以峰值性能为目标进行设计，在代码结构中就必须考虑其大小。另一方面，向量架构要求软件将负载并入向量，并手动管理分支。如图 3-1 所示，每个多处

32、理器都有一个属于以下四种类型之一的芯片存储器：n 每个处理器上有一组本地 32 位寄存器；n 并行数据缓存或共享存储器，由所有标量处理器核心共享，共享存储器空间就位于此处；n 只读固定缓存，由所有标量处理器核心共享，可加速从固定存储器空间进行的读取操作（这是设备存储器的一个只读区域）；n 一个只读纹理缓存，由所有标量处理器核心共享，加速从纹理存储器空间进行的读取操作（这是设备存储器的一个只读区域），每个多处理器都会通过实现不同寻址模型和数据过滤的纹理单元访问纹理缓存，相关内容请参见第 4.3.4 节。本地和全局存储器空间是设备存储器的读/写区域，不应缓存。一个多处理器一次可处理的块数量取决于每

33、个线程有多少个寄存器、每个块需要多少共享存储器来支持给定的内核，这是因为多处理器的寄存器和共享存储器对于一批块的所有线程来说都是分离的。如果没有足够的寄存器或共享存储器可供多处理器用于处理至少一个块，内核将启动失败。一个多处理器可并发执行多达 8 个线程块。如果 warp 块执行的非原子指令为 warp 块的多个线程写入全局或共享存储器中的同一位置，针对此位置的串行化写入操作的数量和这些写入操作所发生的顺序将无法确定，但其中一项操作必将成功。如果 warp 块执行原子指令来为 warp 块的多个线程读取、修改和写入全局存储器中的同一位置，则针对该位置的每一项读取、修改或写入操作都将发生，且均为

34、串行化操作，但这些操作所发生的顺序无法确定。具有芯片共享存储器的一组 SIMT 多处理器图 3-1. 硬件模型3.2 多个设备若要多 GPU 系统上运行的应用程序将多个 GPU 作为 CUDA 设备使用，则这些 GPU 必须具有相同的类型。但如果系统采用的是 SLI 模式，则仅有一个 GPU 可用作 CUDA 设备，因为所有 GPU 都将在驱动程序栈的最低级别融合。要使 CUDA 能够将各 GPU 视为独立设备，需要在 CUDA 的控制面板内关闭 SLI 模式。3.3 模式切换GPU 将部分 DRAM 存储器专门用于处理所谓的主表面（primary surface），它用于刷新显示设备，用户将

35、查看该设备的输出。当用户通过更改显示器的分辨率或位深度（使用 NVIDIA 的控制面板或 Windows 的显示控制面板）发起模式切换时，主表面所需的存储器数量而言会随之改变。例如，如果用户将显示器的分辨率从 1280x1024x32 位更改为 1600x1200x32 位，系统必须为主表面分配 7.68 MB 的存储器，而不是 5.24 MB。（使用防锯齿设置运行的全屏图形应用程序可能需要为主表面分配更多显示存储器。）在 Windows 上，其他事件也可能会启动显示模式切换，包括启动全屏 DirectX 应用程序、按 Alt+Tab 键从全屏 DirectX 应用程序中切换出来或者按 Ctr

36、l+Alt+Del 键锁定计算机。如果模式切换增加了主表面所需的存储器数量，系统可能就必须挪用分配给 CUDA 应用程序的存储器，从而导致此类应用程序崩溃。第 4 章应用程序编程接口4.1 C 编程语言的扩展CUDA 编程接口的目标是为熟悉 C 编程语言的用户提供相对简单的途径，使之可轻松编写由设备执行的程序。它包含：n C 语言的最小扩展集，如 4.2 节所述，这允许程序员使源代码的某些部分可在设备上执行；n 一个运行时库，可分割为：l 一个主机组件，如 4.5 节所述，运行在主机上，提供函数来通过主机控制和访问一个或多个计算设备；l 一个设备组件，如 4.4 节所述，运行在设备上，提供特定

37、于设备的函数；l 一个通用组件，如 4.3 节所述，提供内置向量类型和 C 标准库的一个子集，主机和设备代码中都将支持此子集。有必要强调，C 标准库中支持在设备上运行的函数只有通用运行时组件所提供的函数。4.2 语言扩展对 C 编程语言的扩展共有四重：n 函数类型限定符，指定函数是在主机上还是设备上执行，以及函数是可通过主机还是可通过设备调用（参见第 4.2.1 节）；n 变量类型限定符，指定一个变量在设备上的存储器位置（参见第 4.2.2 节）；n 一条新指令，指定如何通过主机在设备上执行内核（参见第 4.2.3 节）；n 四个内置变量，指定网格和块维度以及块和线程索引（参见第 4.2.4

38、节）。包含这些扩展的所有源文件都必须使用 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

39、_使用 _host_ 限定符声明的函数具有以下特征：n 在主机上执行；n 仅可通过主机调用。仅使用 _host_ 限定符声明函数等同于不使用 _host_、_device_ 或 _global_ 限定符声明函数，这两种情况下，函数都将仅为主机进行编译。但 _host_ 限定符也可与 _device_ 限定符一起使用，此时函数将为主机和设备进行编译。4.2.1.4 限制_device_ 和 _global_ 函数不支持递归。_device_ 和 _global_ 函数的函数体内无法声明静态变量。_device_ 和 _global_ 函数不得有数量可变的参数。_device_ 函数的地址无法获取

40、，但支持 _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_ 限定符一起使用