CUDA编程指南2.0

NVIDIA CUDA 计算统一设备架构 编程指南 版本 2.0 6 / 7 / 2008 目    录 第 1 章    简介    1 1.1  CUDA：可伸缩并行编程模型    1 1.2  GPU：高度并行化、多线程、多核处理器    1 1.3  文档结构    3 第2章    编程模型    4 2.1  线程层次结构    4 2.2  存储器层次结构    6 2.3  主机和设备    6 2.4  软件栈    7 2.5  计算能力    8 第 3 章    GPU 实现    9 3.1  具有芯片共享存储器的一组 SIMT 多处理器    9 3.2  多个设备    11 3.3  模式切换    11 第 4 章    应用程序编程接口    12 4.1  C 编程语言的扩展    12 4.2  语言扩展    12 4.2.1  函数类型限定符    12 4.2.1.1  _device_    12 4.2.1.2  _global_    13 4.2.1.3  _host_    13 4.2.1.4  限制    13 4.2.2  变量类型限定符    13 4.2.2.1  _device_    13 4.2.2.2  _constant_    13 4.2.2.3  _shared_    14 4.2.2.4  限制    14 4.2.3  执行配置    15 4.2.4  内置变量    15 4.2.4.1  gridDim    15 4.2.4.2  blockIdx    15 4.2.4.3  blockDim    15 4.2.4.4  threadIdx    15 4.2.4.5  warpSize    16 4.2.4.6  限制    16 4.2.5  使用 NVCC 进行编译    16 4.2.5.1  _noinline_    16 4.2.5.2  #pragma unroll    16 4.3  通用运行时组件    17 4.3.1  内置向量类型    17 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    17 4.3.1.2  dim3 类型    17 4.3.2  数学函数    17 4.3.3  计时函数    17 4.3.4  纹理类型    18 4.3.4.1  纹理参考声明    18 4.3.4.2  运行时纹理参考属性    18 4.3.4.3  来自线性存储器的纹理与来自 CUDA 数组的纹理    19 4.4  设备运行时组件    19 4.4.1  数学函数    19 4.4.2  同步函数    19 4.4.3  纹理函数    19 4.4.3.1  来自线性存储器的纹理    19 4.4.3.2  来自 CUDA 数组的纹理    20 4.4.4  原子函数    20 4.4.5  warp vote 函数    20 4.5  主机运行时组件    21 4.5.1  一般概念    21 4.5.1.1  设备    21 4.5.1.2  存储器    22 4.5.1.3  OpenGL 互操作性    22 4.5.1.4  Direct3D 互操作性    22 4.5.1.5  异步并发执行    22 4.5.2  运行时 API    23 4.5.2.1  初始化    23 4.5.2.2  设备管理    23 4.5.2.3  存储器管理    24 4.5.2.4  流管理    25 4.5.2.5  事件管理    25 4.5.2.6  纹理参考管理    25 4.5.2.7  OpenGL 互操作性    27 4.5.2.8  Direct3D 互操作性    27 4.5.2.9  使用设备模拟模式进行调试    28 4.5.3  驱动程序 API    29 4.5.3.1  初始化    29 4.5.3.2  设备管理    29 4.5.3.3  上下文管理    29 4.5.3.4  模块管理    30 4.5.3.5  执行控制    30 4.5.3.6  存储器管理    31 4.5.3.7  流管理    32 4.5.3.8  事件管理    32 4.5.3.9  纹理参考管理    33 4.5.3.10  OpenGL 互操作性    33 4.5.3.11  Direct3D 互操作性    33 第 5 章    性能指南    35 5.1  指令性能    35 5.1.1    指令吞吐量    35 5.1.1.1  数学指令    35 5.1.1.2  控制流指令    36 5.1.1.3  存储器指令    36 5.1.1.4  同步指令    37 5.1.2  存储器带宽    37 5.1.2.1  全局存储器    37 5.1.2.2  本地存储器    43 5.1.2.3  固定存储器    43 5.1.2.4  纹理存储器    43 5.1.2.5  共享存储器    43 5.1.2.6  寄存器    48 5.2  每个块的线程数量    49 5.3  主机和设备间的数据传输    49 5.4  纹理获取与全局或固定存储器读取的对比    50 5.5  整体性能优化战略    50 第 6 章    矩阵乘法示例    52 6.1  概述    52 6.2  源代码清单    53 6.3  源代码说明    54 6.3.1  Mul()    54 6.3.2  Muld()    54 附录 A    技术规范    56 A.1  一般规范    56 A.1.1  计算能力 1.0 的规范    56 A.1.2  计算能力 1.1 的规范    57 A.1.3  计算能力 1.2 的规范    57 A.1.4  计算能力 1.3 的规范    57 A.2  浮点标准    57 附录 B    标准数学函数    59 B.1  一般运行时组件    59 B.1.1  单精度浮点函数    59 B.1.2  双精度浮点函数    60 B.1.3  整型函数    62 B.2  设备运行时组件    62 B.2.1  单精度浮点函数    62 B.2.2  双精度浮点函数    63 B.2.3  整型函数    64 附录 C    原子函数    65 C.1  数学函数    65 C.1.1  atomicAdd()    65 C.1.2  atomicSub()    65 C.1.3  atomicExch()    65 C.1.4  atomicMin()    65 C.1.5  atomicMax()    66 C.1.6  atomicInc()    66 C.1.7  atomicDec()    66 C.1.8  atomicCAS()    66 C.2  位逻辑函数    66 C.2.1  atomicAnd()    66 C.2.2  atomicOr()    67 C.2.3  atomicXor()    67 附录 D    纹理获取    68 D.1  最近点取样    68 D.2  线性过滤    69 D.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. 矩阵乘法........................................................... .............. .............. ....................................52 第 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  文档结构 本文档分为以下几个章节： ? 第 1 章是 CUDA 和 GPU 的简介。 ? 第 2 章概述 CUDA 编程模型。 ? 第 3 章介绍 GPU 实现。 ? 第 4 章介绍 CUDA API 和运行时。 ? 第 5 章提供如何实现最高性能的一些指南。 ? 第 6 章通过一些简单的示例代码概况之前各章的内容。 ? 附录 A 提供各种设备的技术规范。 ? 附录 B 列举 CUDA 中支持的数学函数。 ? 附录 C 列举 CUDA 中支持的原子函数。 ? 附录 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<<<1, N>>>(A, B, C); } 执行内核的每个线程都会被分配一个独特的线程 ID，可通过内置的 threadIdx 变量在内核中访问此 ID。以下示例代码将大小为 N 的向量 A 和向量 B 相加，并将结果存储在向量 C 中： __global__ void vecAdd(float* A, float* B, float* C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main() { // Kernel invocation vecAdd<<<1, N>>>(A, B, C); } 执行 vecAdd( ) 的每个线程都会执行一次成对的加法运算。 2.1  线程层次结构 为方便起见，我们将 threadIdx 设置为一个包含 3 个组件的向量，因而可使用一维、二维或三维缩影标识线程，构成一维、二维或三维线程块。这提供了一种自然的方法，可为一个域中的各元素调用计算，如向量、矩阵或字段。下面的示例代码将大小为 NxN 的矩阵 A 和矩阵 B 相加，并将结果存储在矩阵 C 中： __global__ void matAdd(float A[N][N], float B[N][N], float C[N][N]) { int i = threadIdx.x; int j = threadIdx.y; C[i][j] = A[i][j] + B[i][j]; } int main() { // Kernel invocation dim3 dimBlock(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()_ 应是轻量级的，一个块中的所有线程都必须位于同一个处理器核心中。因而，一个处理器核心的有限存储器资源制约了每个块的线程数量。在 NVIDIA Tesla 架构中，一个线程块最多可以包含 512 个线程。 但一个内核可能由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数量。这些块将组织为一个一维或二维线程块网格，如图 2-1 所示。该网格的维度由 <<<…>>> 语法的第一个参数指定。网格内的每个块多可由一个一维或二维索引标识，可通过内置的 blockIdx 变量在内核中访问此索引。可以通过内置的 blockDim 变量在内核中访问线程块的维度。此时，之前的示例代码应修改为： __global__ void matAdd(float A[N][N], float B[N][N], float C[N][N]) { int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; if (i < N && j < N) C[i][j] = A[i][j] + B[i][j]; } 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 所示，每个多处理器都有一个属于以下四种类型之一的芯片存储器： ? 每个处理器上有一组本地 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 存储器专门用于处理所谓的主表面（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 编程语言的用户提供相对简单的途径，使之可轻松编写由设备执行的程序。 它包含： ? C 语言的最小扩展集，如 4.2 节所述，这允许程序员使源代码的某些部分可在设备上执行； ? 一个运行时库，可分割为： ● 一个主机组件，如 4.5 节所述，运行在主机上，提供函数来通过主机控制和访问一个或多个计算设备； ● 一个设备组件，如 4.4 节所述，运行在设备上，提供特定于设备的函数； ● 一个通用组件，如 4.3 节所述，提供内置向量类型和 C 标准库的一个子集，主机和设备代码中都将支持此子集。 有必要强调，C 标准库中支持在设备上运行的函数只有通用运行时组件所提供的函数。    4.2  语言扩展 对 C 编程语言的扩展共有四重： ? 函数类型限定符，指定函数是在主机上还是设备上执行，以及函数是可通过主机还是可通过设备调用（参见第 4.2.1 节）； ? 变量类型限定符，指定一个变量在设备上的存储器位置（参见第 4.2.2 节）； ? 一条新指令，指定如何通过主机在设备上执行内核（参见第 4.2.3 节）； ? 四个内置变量，指定网格和块维度以及块和线程索引（参见第 4.2.4 节）。 包含这些扩展的所有源文件都必须使用 CUDA 编译器 nvcc 进行编译，4.2.5 节简单介绍了相关内容。关于 nvcc 的具体介绍将在其他文档中提供。 这些扩展均具有一些限制，下面几个小节将分别加以介绍。如果违背了这些限制，nvcc 将发出错误或警报信息，但有些违规情况无法检测到。 4.2.1  函数类型限定符 4.2.1.1  _device_ 使用 _device_ 限定符声明的函数具有以下特征： ? 在设备上执行； ? 仅可通过设备调用。 4.2.1.2  _global_ 使用 _global_ 限定符可将函数声明为内核。此类函数： ? 在设备上执行； ? 仅可通过主机调用。 4.2.1.3  _host_ 使用 _host_ 限定符声明的函数具有以下特征： ? 在主机上执行； ? 仅可通过主机调用。 仅使用 _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_ 限定符一起使用，以更具体地指定变量属于哪个存储器空间。如果未出现其他任何限定符，则变量具有以下特征： ? 位于全局存储器空间中； ? 与应用程序具有相同的生命周期； ? 可通过网格内的所有线程访问，也可通过运行时库从主机访问。 4.2.2.2  _constant_ _constant_ 限定符可选择与 _device_ 限定符一起使用，所声明的变量具有以下特征： ? 位于固定存储器空间中； ? 与应用程序具有相同的生命周期； ? 可通过网格内的所有线程访问，也可通过运行时库从主机访问。 4.2.2.3  _shared_ _shared_ 限定符可选择与 _device_ 限定符一起使用，所声明的变量具有以下特征： ? 位于线程块的共享存储器空间中； ? 与块具有相同的生命周期； ? 尽可通过块内的所有线程访问。 只有在 _syncthreads()_（参见第 4.4.2 节）的执行写入之后，才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量，否则只要之前的语句完成，编译器即可随意优化共享存储器的读写操作。 将共享存储器中的变量声明为外部数组时，例如： extern __shared__ float shared[]; 数组的大小将在启动时确定（参见第 4.2.3 节）。所有变量均以这种形式声明，在存储器中的同一地址开始，因此数组中的变量布局必须通过偏移显式管理。例如，如果一名用户希望在动态分配的共享存储器内获得与以下代码对应的内容： short array0[128]; float array1[64]; int array2[256]; 则应通过以下方法声明和初始化数组： extern __shared__ char array[]; __device__ void func() // __device__ or __global__ function { short* array0 = (short*)array; float* array1 = (float*)&array0[128]; int* array2 = (int*)&array1[64]; } 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 节介绍）。可通过在函数名称和括号参数列表之间插入 <<>> 形式的表达式来指定，其中： ? Dg 的类型为 dim3（参见第 4.3.1.2 节），指定网格的维度和大小，Dg.x * Dg.y 等于所启动的块数量，Dg.z 无用； ? Db 的类型为 dim3（参见第 4.3.1.2 节），指定各块的维度和大小，Db.x * Db.y * Db.z 等于各块的线程数量； ? Ns 的类型为 size_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为外部数组的其他任何变量使用（参见第 4.2.2.3 节），Ns 是一个可选参数，默认值为 0； ? S 的类型为 cudaStream_t，指定相关流；S 是一个可选参数，默认值为 0。 举例来说，一个函数的声明如下： __global__ void Func(float* parameter); 必须通过如下方法来调用此函数： Func<<< Dg, Db, Ns >>>(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  限制 ? 不允许接受任何内置变量的地址。 ? 不允许为任何内置变量赋值。 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 5 for (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）将纹理作为一维数组寻址、使用两个纹理坐标作为二维数组寻址，还是使用三个纹理坐标作为三维数组寻址。数组的元素称为 texel，即“texture elements（纹理元素）”的简写。 其他属性定义纹理获取的输入和输出数据类型，并指定如何介绍输入坐标、应进行怎样的处理。 4.3.4.1  纹理参考声明 纹理参考的部分属性是不变的，在编译时必须为已知，这些属性是在声明纹理参考时指定的。纹理参考在文件作用域内声明，形式为 texture 类型的变量： texture texRef; 其中： ? Type 指定获取纹理时所返回的数据类型；Type 仅限于基本整型、单精度浮点类型和第 4.3.1.1 节定义的 1 组件、2 组件 和 4 组件向量类型； ? Dim 指定纹理参考的维度，其值为 1、2 或 3；Dim 是一个可选的参数，默认值为 1； ? ReadMode 等于 cudaReadModeNormalizedFloat 或 cudaReadModeElementType；如果是cudaReadModeNormalizedFloat，且 Type 为 16 位或 8 位整型类型，则值将作为浮点类型返回，对于所有整型数据而言，无符号整型将映射为 [0.0, 1.0]，有符号整型将映射为 [-1.0, 1.0]，例如，一个值为 0xff 的无符号 8 位纹理元素将被读取为 1；如果是 cudaReadModeElementType，则不执行任何转换操作；ReadMode 是一个可选的参数，默认值为 cudaReadModeElementType。 4.3.4.2  运行时纹理参考属性 纹理参考的其他属性是可变的，可通过主机运行时在运行时更改（第 4.5.2.6 节介绍了运行时 API，第 4.5.3.9 介绍了驱动程序 API）。它们指定纹理坐标是否为规范化的，以及寻址模式和纹理过滤，下面将介绍相关内容。 默认情况下，使用 [0, N） 范围内的浮点坐标引用纹理，其中的 N 是纹理在对应于坐标的维度中的大小。例如，有一个大小为 64x32 的纹理，在 x 和 y 维度引用此纹理时坐标分别处于 [0, 63] 和 [0, 31] 范围内。规范化的纹理坐标将在 [0.0, 1.0) 的范围内指定，而非 [0, N)，因此在规范化的坐标内，同一 64x32 纹理的寻址范围在 x 和 y 维度均为 [0, 1)。一般情况下，纹理坐标与纹理大小无关，规范化的纹理坐标通常足以满足一些应用程序的需求。 寻址模式定义在纹理坐标超出范围时将出现怎样的情况。在使用非规范化纹理坐标时，超出 [0, N) 范围的纹理坐标将被调整：小于 0 的值被设置为 0，大于或等于 N 的值被设置为 N-1。在使用规范化纹理坐标时，默认寻址模式也是调整坐标：小于 0.0 或大于 1.0 的值将被调整到范围 [0.0, 1.0) 内。对于规范化坐标，“warp 块”的寻址模式也可指定。Warp 块寻址往往在纹理包含周期信号时使用。它仅使用纹理坐标的一部分，例如，1.25 被视为 0.25，-1.25 被视为 0.75. 线性纹理过滤只能对配置为返回浮点数据的纹理进行。这将在相邻 texel 间执行低精度插值。在启用时，位于纹理获取位置周围的 texel 将被读取，纹理获取的返回值将根据纹理坐标在 texel 间的位置进行插值。对于一维纹理执行简单的线性插值，而对于二维纹理则执行双线性插值。 附录 D 提供了关于纹理获取的更多细节。 4.3.4.3  来自线性存储器的纹理与来自 CUDA 数组的纹理 纹理可以是线性存储器或 CUDA 数组的任意区域（参见第 4.5.1.2 节）。 在线性存储器内分配的纹理： ? 维度仅能为 1； ? 不支持纹理过滤； ? 仅可使用非规范化整型纹理坐标寻址； ? 不支持多种寻址模式：超出范围的纹理访问将返回零。 硬件会对纹理基址实施对齐要求。为了抽象这种来自程序员的对齐要求，绑定设备存储器上的纹理参考的函数将传回一个字节偏移，必须将其应用到纹理获取，之后才能读取所需的存储器。CUDA 分配例程返回的基址指针符合这种对齐限制，因此应用程序可通过向 cudaBindTexture()/cuTexRefSetAddress() 传递所分配的指针来完全避免偏移。 4.4  设备运行时组件 设备运行时组件仅可用于设备函数。 4.4.1  数学函数 对于 B.1 节介绍的部分函数而言，设备运行时组件中存在准确性略低而速度更快的版本；其名称相同，但带有一个_前缀（如_sinf(x)）。B.2 节列举了这些内部函数，还列举了它们的对应误差范围。 编译器有一个 (-use_fast_math) 选项，用于强制要求所有函数编译其准确性略低的版本（如果存在）。 4.4.2  同步函数 void __syncthreads(); 同步块中的所有线程。一旦所有线程均达到此同步点，执行将正常恢复。 _syncthreads() 用于调整同一个块的线程之间的通信。在一个块内的某些线程访问共享或全局存储器中的相同地址时，部分访问操作可能存在写入后读取、读取后写入或写入后写入之类的风险。可通过在这些访问操作间同步线程来避免这些数据风险。 _syncthreads() 允许在条件代码中使用，但仅当条件估值在整个线程块中都相同时才允许使用，否则代码执行将有可能挂起，或者出现意料之外的副作用。 4.4.3  纹理函数 4.4.3.1  来自线性存储器的纹理 对于来自线性存储器的纹理，通过 tex1Dfetch() 系列函数访问纹理，示例如下： template Type tex1Dfetch( texture texRef, int x); float tex1Dfetch( texture texRef, int x); float tex1Dfetch( texture texRef, int x); float tex1Dfetch( texture texRef, int x); float tex1Dfetch( texture texRef, int x); 这些函数会使用纹理坐标 x 获取绑定到纹理参考 texRef 的线性存储器区域。不支持纹理过滤和寻址模式。对于整型来说，这些函数可选择将整型转变为单精度浮点类型。 除了上述函数以外，还支持 2 元组和 4 元组，示例如下： float4 tex1Dfetch( texture texRef, int x); 以上示例将使用纹理坐标 x 获取绑定到纹理参考 texRef 的线性存储器。 4.4.3.2  来自 CUDA 数组的纹理 对于来自 CUDA 数组的纹理，可通过 tex1D()、tex2D()、tex3D() 访问纹理： template Type tex1D(texture texRef, float x); template Type tex2D(texture texRef, float x, float y); template Type tex3D(texture texRef, float x, float y, float z); 这些函数将使用纹理坐标 x、y 和 z 获取绑定到纹理参考 texRef 的 CUDA 数组。纹理参考的不变（编译时）和可变（运行时）属性相互结合，共同确定坐标的解释方式、在纹理获取过程中发生的处理以及纹理获取所提供的返回值（参见第 4.3.4.1 和第 4.3.4.2 节）。 4.4.4  原子函数 原子函数对位于全局或共享存储器内的一个 32 位或 64 位字执行读取-修改-写入原子操作。例如，atomicAdd() 将在全局或共享存储器内的某个地址读取 32 位字，将其与一个整型相加，并将结果写回同一地址。之所以说这样的操作是原子的，是因为它可在不干扰其他线程的前提下执行。换句话说，在操作完成中，其他任何线程都无法访问此地址。 附录 C 列举了受支持的所有原子函数。如附录所述，并非所有设备都支持这些函数。具体来说，计算能力为 1.0 的设备不支持任何原子函数。 原子操作仅适用于有符号和无符号整型（但 atomicExch() 是一个例外情况，它支持单精度浮点数字）。 4.4.5  warp vote 函数 只有计算能力为 1.2 或更高的设备支持 Warp vote 函数。 int __all(int predicate); 为 warp 块内的所有线程计算 predicate，当且仅当所有线程的 predicate 均非零时返回非零值。 int __any(int predicate); 为 warp 块内的所有线程计算 predicate，当且仅当任意线程的 predicate 非零时返回非零值。 4.5  主机运行时组件 只有主机函数才能使用主机运行时组件。 它提供了具有以下功能的函数： ? 设备管理； ? 上下文管理； ? 存储器管理； ? 代码模块管理； ? 执行控制； ? 纹理参考管理； ? 与 OpenGL 和 Direct3D 的互操作性。 它包含两个 API： ? 一个称为 CUDA 驱动程序 API 的低级 API； ? 一个称为 CUDA 运行时 API 的高级 API，它是在 CUDA 驱动程序 API 的基础之上实现的。 这些 API 是互斥的：一个应用程序仅能使用其中之一。 CUDA 运行时提供了隐式初始化、上下文管理和模块管理，从而简化了设备代码管理。Nvcc 生成的 C 主机代码基于 CUDA 运行时（请参见第 4.2.5 节），因此链接到此代码的应用程序必须使用 CUDA 运行时 API。 相反，CUDA 驱动程序 API 需要的代码数量更多，编程和调试更加困难，但提供了更出色的控制级别，此外还具有独立于语言的特点，因为它仅处理 cubin 对象（请参见第 4.2.5 节）。具体来说，使用 CUDA 驱动程序 API 配置和启动内核的难度更大，因为执行配置和内核参数必须通过显式函数调用来指定，而不能利用第 4.2.3 节介绍的执行配置语法。此外，设备模拟（请参见第 4.5.2.9 节）不适用于 CUDA 驱动程序 API。 CUDA 驱动程序 API 是通过 nvcuda 动态库提供的，其所有入口点都带有 cu 前缀。 CUDA 运行时 API 是通过 cudart 动态库提供的，其所有入口点都带有 cuda 前缀。 4.5.1  一般概念 4.5.1.1  设备 两种 API 都提供了可枚举系统上可用设备、查询其属性、为内核执行选择一个设备的函数（运行时 API 的相关内容请参见第 4.5.2.2 节，驱动程序 API 的相关内容请参见第 4.5.3.2 节）。 多个主机线程可在同一个设备上执行设备代码，但根据设计，主机线程只能在一个设备上执行设备代码。因而，需要多个主机线程在多个设备上执行设备代码。此外，通过一个主机线程的运行时创建的 CUDA 资源无法由来自其他主机线程的运行时使用。 4.5.1.2  存储器 设备存储器可指派为线性存储器或 CUDA 数组。 线性存储器位于 32 位地址空间内的设备上，因此，举例来说，独立分配的实体可通过二进制树内的指针引用另外一个实体。 CUDA 数组是不透明的存储器布局，专为纹理获取而优化（参见第 4.3.4 节）。它们可以是一维、二维或三维的，由元素组成，均包含 1、2、4 个组件，这些组件可以是有符号或无符号的 8 位、16 位 或 32 位整型，也可以是 16 位浮点（当前仅有驱动程序 API 支持）或 32 位浮点。CUDA 数组仅可由内核通过纹理获取读取，仅可绑定到具有相同数量的打包组件的纹理参考。 主机可通过第 4.5.2.3 和第 4.5.3.6 节介绍的存储器复制函数读取和写入线性存储器和 CUDA 数组。 主机运行时还提供了函数来分配和释放分页锁定的主机存储器—与 malloc() 分配的普通可分页存储器恰好相反。分页锁定存储器的优势之一在于，如果将主机存储器指派为分页锁定存储器，主机存储器和设备存储器之间的带宽较高——但仅针对分配主机存储器的主机线程所执行的数据传输。分页锁定的存储器是一种稀缺资源，因此分页锁定存储器中的分配将先于可分页存储器的分配而出错。此外，由于减少了操作系统可用于分页的物理存储器数量，分配过多的分页锁定存储器将降低整体系统性能。 4.5.1.3  OpenGL 互操作性 OpenGL 缓冲对象可映射到 CUDA 的地址空间，从而使 CUDA 能够读取 OpenGL 写入的数据或使 CUDA 能够写入入数据供 OpenGL 使用。第 4.5.2.7 节描述了如何通过运行时 API 实现此目标，第 4.5.3.10 节描述了如何通过驱动程序 API 实现此目标。 4.5.1.4  Direct3D 互操作性 Direct3D 资源可映射到 CUDA 的地址空间，从而使 CUDA 能够读取 Direct3D 写入的数据，或者使 CUDA 能够写入数据供 Direct3D 使用。第 4.5.2.8 节描述了如何通过运行时 API 实现此目标，同一节还介绍了如何通过驱动程序 API 实现此目标。 对于可映射哪些资源的限制条件，请参见 cudaD3D9RegisterResource() 和 cuD3D9RegisterResource() 参考手册。 CUDA 上下文一次仅可与一个 Direct3D 设备互操作，CUDA 上下文和 Direct3D 设备必须是在同一个 GPU 上创建的。此外，Direct3D 设备在创建时必须使用 D3DCREATE_HARDWARE_VERTEXPROCESSING 标记。 Direct3D 互操作性目前仅支持 Direct3D 9.0。 4.5.1.5  异步并发执行 为了促进主机和设备之间的并发执行，某些运行时函数是异步的：控制将在设备完成所请求的任务之前返回应用程序。此类函数包括： ? 通过 _global_ 函数或 cuLaunchGrid() 和 cuLaunchGridAsync() 启动的内核； ? 执行存储器复制和带有 Async 后缀的函数； ? 执行设备与设备间双向存储器复制的函数； ? 设置存储器的函数。 某些设备还可在分页锁定的主机存储器和设备存储器之间执行复制，且与内核执行并发地执行此类复制操作。如果使用的是运行时 API，应用程序可通过调用 cudaGetDeviceProperties() 并检查 deviceOverlap 来查询此功能；如果使用的是驱动程序 API，则可通过使用 CU_DEVICE_ATTRIBUTE_GPU_OVERLAP 调用 cuDeviceGetAttribute() 来查询。此功能当前仅支持不涉及 CUDA 数组或通过 cudaMallocPitch()（请参见第 4.5.2.3 节）或 cuMemAllocPitch()（请参见第 4.5.3.6 节）分配的二维数组的存储器复制。 应用程序通过流管理并发。流就是按顺序执行的一系列操作。另一方面，不同的流可采用针对其他流来说不正确的顺序来执行其操作，也可并发执行。 流的定义方法是创建流对象，并将其指定为内核启动和主机与设备间内存双向复制序列的流参数。第 4.5.2.4 节描述了如何通过运行时 API 实现此目标，第 4.5.3.7 节描述了如何通过驱动程序 API 实现此目标。 不带流参数或使用零作为流参数的任何内核启动、存储器设置或存储器复制函数都仅能在此前的操作完成后开始，包括作为流的一部分的操作，而在此类函数完成之前，无法启动任何后续操作。提供了无流参数的内核启动，无 Async 后缀的存储器复制将被指派给默认零流。 用于运行时 API 的 cudaStreamQuery() 和用于驱动程序 API 的 cuStreamQuery() 为应用程序提供了一种方法，使其能够了解一个流中的之前全部操作是否已完成。用于运行时 API 的 cudaStreamSnchronize() 和用于驱动程序 API 的 cuStreamSynchronize() 提供了一种方法，可显式强制要求运行时等待流中之前的所有操作完成。 类似地，用于运行时 API 的 cudaThreadSynchronize() 和用于驱动程序 API 的 cuCtxSynchronize() 使应用程序能够强制要求运行时等待所有流中的所有之前设备任务完成。为了避免不必要的减速，最好将这些函数用于计时或者隔离失败的启动或存储器复制操作。用于运行时 API 的 cudaStreamDestroy() 和用于驱动程序 API 的 cuStreamDestroy() 将等待给定流内之前的所有任务完成，之后再销毁流，并将控制返回给主机线程。 运行时还提供了一种方法，可密切监控设备的进度并执行准确的计时，它允许应用程序异步记录程序内任意点的事件，并查询这些事件的实际记录事件。一个事件将在先于该事件的所有任务（也可以是给定流中的所有操作）均已完成时记录。零流中的事件将在设备完成来自所有流的之前任务/操作后记录。第 4.5.2.5 节将描述如何通过运行时 API 实现此目标，第 4.5.3.8 节将描述如何通过驱动程序 API 实现此目标。 如果主机线程在来自不同流的两项操作之间调用了分页锁定的主机存储器分配、设备存储器分配、设备存储器设置、设备与设备之间的双向存储器复制或零流的任何 CUDA 操作，则这两项操作无法并行执行。 程序员可将 CUDA_LAUNCH_BLOCKING 环境变量设置为 1，从而为系统上运行的所有 CUDA 应用程序全局禁用异步执行。此特性仅为调试提供，不应用于使生产软件更可靠运行之目的。 4.5.2  运行时 API 4.5.2.1  初始化 运行时 API 不存在显式初始化函数；它将在运行时函数被首次调用时初始化。需要牢记，在计时运行时调用发生时，或解释来自第一次运行时调用的错误码时，即将进行初始化。 4.5.2.2  设备管理 cudaGetDeviceCount() 和 cudaGetDeviceProperties() 提供了一种方法，用于枚举这些设备并检索其属性： int deviceCount; cudaGetDeviceCount(&deviceCount); int device; for (device = 0; device < deviceCount; ++device) { cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp, device); } cudaSetDevice() 用于选择与主机线程相关的设备： cudaSetDevice(device); 必须首先选择设备，之后才能调用 _global_ 函数或任何来自运行时 API 的函数。如果未通过显式调用 cudaSetDevice() 完成此任务，将自动选中设备 0，随后对 cudaSetDevice() 的任何显式调用都将无效。 4.5.2.3  存储器管理 线性存储器是使用 cudaMalloc() 或 cudaMallocPitch() 分配的，使用 cudaFree() 释放。 以下示例代码将在线性存储器中分配一个包含 256 个浮点元素的数组： float* devPtr; cudaMalloc((void**)&devPtr, 256 * sizeof(float)); 建议在分配二维数组时使用 cudaMallocPitch()，因为它能确保合理填充已分配的存储器，满足第 5.1.2.1 节介绍的对齐要求，从而确保访问行地址或执行二维数组与设备存储器的其他区域之间的复制（使用 cudaMemcpy2D()）时获得最优性能。所返回的间距（或步幅）必须用于访问数组元素。以下代码示例将分配一个 widthxheight 的二维浮点值数组，并显示如何在设备代码中循环遍历数组元素： // host code float* devPtr; int pitch; cudaMallocPitch((void**)&devPtr, &pitch, width * sizeof(float), height); myKernel<<<100, 512>>>(devPtr, pitch); // device code __global__ void myKernel(float* devPtr, int pitch) { for (int r = 0; r < height; ++r) { float* row = (float*)((char*)devPtr + r * pitch); for (int c = 0; c < width; ++c) { float element = row[c]; } } } CUDA 数组是使用 cudaMallocArray() 分配的，使用 cudaFreeArray() 释放。cudaMallocArray() 需要使用 cudaCreateChannelDesc() 创建的格式描述。 以下代码示例分配了一个 widthxheight 的 CUDA 数组，包含一个 32 位的浮点组件： cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); cudaArray* cuArray; cudaMallocArray(&cuArray, &channelDesc, width, height); cudaGetSymbolAddress() 用于检索指向为全局存储器空间中声明的变量分配的存储器的地址。所分配存储器的大小是通过 cudaGetSymbolSize() 获取的。 参考手册列举了用于在 cudaMalloc() 分配的线性存储器、cudaMallocPitch() 分配的线性存储器、CUDA 数组和为全局或固定存储器空间中声明的变量分配的存储器之间复制存储器的所有函数。 下面的代码示例将二维数组复制到之前代码示例中分配的 CUDA 数组中： cudaMemcpy2DToArray(cuArray, 0, 0, devPtr, pitch, width * sizeof(float), height, cudaMemcpyDeviceToDevice); 下面的代码示例将一些主机存储器数组复制到设备存储器中： float data[256]; int size = sizeof(data); float* devPtr; cudaMalloc((void**)&devPtr, size); cudaMemcpy(devPtr, data, size, cudaMemcpyHostToDevice); 下面的代码示例将一些主机存储器数组复制到固定存储器中： __constant__ float constData[256]; float data[256]; cudaMemcpyToSymbol(constData, data, sizeof(data)); 4.5.2.4  流管理 以下代码示例创建两个流： cudaStream_t stream[2]; for (int i = 0; i < 2; ++i) cudaStreamCreate(&stream[i]); 这些流均通过以下代码示例定义为一个序列，包括一次从主机到设备的存储器复制、一次内核启动、一次从设备到主机的存储器复制： for (int i = 0; i < 2; ++i) cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]); for (int i = 0; i < 2; ++i) myKernel<<<100, 512, 0, stream[i]>>> (outputDevPtr + i * size, inputDevPtr + i * size, size); for (int i = 0; i < 2; ++i) cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]); cudaThreadSynchronize(); 两个流均会将其输入数组 hostPtr 的一部分复制到设备存储器的 inputDevPtr 数组中，通过调用 myKernel() 处理设备上的 inputDevPtr，并将结果 outputDevPtr 复制回 hostPtr 的相同部分。使用两个流处理 hostPtr 允许一个流的存储器复制与另外一个流的内核执行相互重叠。hostPtr 必须指向分页锁定的主机存储器，这样才能出现重叠： float* hostPtr; cudaMallocHost((void**)&hostPtr, 2 * size); 最后调用了 cudaThreadSynchronize()，目的是在进一步处理之前确定所有流均已完成。cudaStreamSynchronize() 可用于同步主机与特定流，允许其他流继续在该设备上执行。通过调用 cudaStreamDestroy() 可释放流。 4.5.2.5  事件管理 下面的代码示例创建了两个事件： cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); 这些事件可用于为上一节的代码示例计时，方法如下： cudaEventRecord(start, 0); for (int i = 0; i < 2; ++i) cudaMemcpyAsync(inputDev + i * size, inputHost + i * size, size, cudaMemcpyHostToDevice, stream[i]); for (int i = 0; i < 2; ++i) myKernel<<<100, 512, 0, stream[i]>>> (outputDev + i * size, inputDev + i * size, size); for (int i = 0; i < 2; ++i) cudaMemcpyAsync(outputHost + i * size, outputDev + i * size, size, cudaMemcpyDeviceToHost, stream[i]); cudaEventRecord(stop, 0); cudaEventSynchronize(stop); float elapsedTime; cudaEventElapsedTime(&elapsedTime, start, stop); cudaEventDestroy(start); cudaEventDestroy(stop); 4.5.2.6  纹理参考管理 高级 API 定义的 texture 类型是一种公开继承自低级 API 定义的 textureReference 类型的纹理，如下： struct textureReference { int normalized; enum cudaTextureFilterMode filterMode; enum cudaTextureAddressMode addressMode[3]; struct cudaChannelFormatDesc channelDesc; } ? Normalized 指定纹理坐标是否为规范化形式；如果非零，纹理中的所有元素都将使用 [0, 1] 范围内的纹理坐标寻址，而非 [0,width-1]、[0,height-1] 或 [0,depth-1]，其中的 width、height 和 depth 是纹理大小； ? filterMode 指定过滤模式，即获取根据输入的纹理坐标计算的纹理时如何返回值；filterMode 等于 cudaFilterModePoint 或 cudaFilterModeLinear；如果等于 cudaFilterModePoint，则所返回的值为纹理坐标最接近输入纹理坐标的 texel；如果等于 cudaFilterModeLinear，则所返回的值为纹理坐标最接近输入纹理坐标的两个（针对一维纹理）、四个（针对二维纹理）或八个（针对三维纹理）texel 的线性插值； 对于浮点型的返回值，cudaFilterModeLinear 是惟一的有效值。 ? addressMode 指定寻址模式，表明如何处理超出范围的纹理坐标；addressMode 是一个大小为 3 的数组，其第一个、第二个和第三个元素分别指定第一个、第二个和第三个纹理坐标的寻址模式；寻址模式可等于 cudaAddressModeClamp，此时超出范围的纹理坐标将被调整到有效范围之内，也可等于 cudaAddressModeWrap，此时超出范围的纹理坐标将被限定到有效范围之内； 对于规范化的纹理坐标，仅支持 cudaAddressModeWrap。 ? channelDesc 描述获取纹理时所返回的值的格式；channelDesc 的类型如下： struct cudaChannelFormatDesc { int x, y, z, w; enum cudaChannelFormatKind f; }; 其中 x、y、z 和 w 等于返回值各组件的位数，而 f 为： ● cudaChannelFormatKindSigned，在这些组件是有符号整型时； ● cudaChannelFormatKindUnsigned，在这些组件是无符号整型时； ● cudaChannelFormatKindFloat，在这些组件是浮点类型时。 normalized、addressMode 和 filterMode 可直接在主机代码中修改。它们仅适用于绑定到 CUDA 数组的纹理参考。 在内核使用纹理参考从纹理存储器中读取之前，必须使用 cudaBindTexture() 或 cudaBindTextureToArray() 将纹理参考绑定到纹理。 以下代码示例将一个纹理参考绑定到 devPtr 指向的线性存储器： ? 使用低级 API： texture texRef; textureReference* texRefPtr; cudaGetTextureReference(&texRefPtr, “texRef”); cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); cudaBindTexture(0, texRefPtr, devPtr, &channelDesc, size); ? 使用高级 API： texture texRef; cudaBindTexture(0, texRef, devPtr, size); 以下代码示例将一个纹理参考绑定到 CUDA 数组 cuArray： ? 使用低级 API： texture texRef; textureReference* texRefPtr; cudaGetTextureReference(&texRefPtr, “texRef”); cudaChannelFormatDesc channelDesc; cudaGetChannelDesc(&channelDesc, cuArray); cudaBindTextureToArray(texRef, cuArray, &channelDesc); ? 使用高级 API： texture texRef; cudaBindTextureToArray(texRef, cuArray); 将纹理绑定到纹理参考时指定的格式必须与声明纹理参考时指定的参数相匹配；否则纹理获取的结果将无法确定。 cudaUnbindTexture() 用于解除纹理参考的绑定。 4.5.2.7  OpenGL 互操作性 首先必须将一个缓冲对象注册到 CUDA，之后才能进行映射。可通过 cudaGLRegisterBufferObject() 完成： GLuint bufferObj; cudaGLRegisterBufferObject(bufferObj); 注册完成后，内核即可使用 cudaGLMapBufferObject() 返回的设备存储器地址读取或写入缓冲对象： GLuint bufferObj; float* devPtr; cudaGLMapBufferObject((void**)&devPtr, bufferObj); 解除映射是通过 cudaGLUnmapBufferObject() 完成的，可使用 cudaGLUnregisterBufferObject() 取消注册。 4.5.2.8  Direct3D 互操作性 Direct3D 互操作性要求在执行其他任何运行时调用之前通过 cudaD3D9SetDirect3DDevice() 指定 Direct3D 设备。 随后即可使用 cudaD3D9RegisterResource() 将 Direct3D 资源注册到 CUDA： LPDIRECT3DVERTEXBUFFER9 buffer; cudaD3D9RegisterResource(buffer, cudaD3D9RegisterFlagsNone); LPDIRECT3DSURFACE9 surface; cudaD3D9RegisterResource(surface, cudaD3D9RegisterFlagsNone); cudaD3D9RegisterResource() 可能具有较高的开销，通常仅为每个资源调用一次。使用 cudaD3D9UnregisterVertexBuffer() 可取消注册。将资源注册到 CUDA 之后，即可在需要时分别使用  cudaD3D9MapResources() 和 cudaD3D9UnmapResources() 任意多次地映射和解除映射。内核可使用 cudaD3D9ResourceGetMappedPointer() 返回的设备存储器地址和 cudaD3D9ResourceGetMappedSize()、cudaD3D9ResourceGetMappedPitch() 及 cudaD3D9ResourceGetMappedPitchSlice() 返回的大小和间距信息来读取和写入已映射的资源。通过 Direct3D 访问已映射的资源将导致不确定的结果。 下面的代码示例使用 0  填充了一个缓冲区： void* devPtr; cudaD3D9ResourceGetMappedPointer(&devPtr, buffer); size_t size; cudaD3D9ResourceGetMappedSize(&size, buffer); cudaMemset(devPtr, 0, size); 在下面的代码示例中，每个线程都访问大小为 (width, height) 的二维表面的一个像素，像素格式为 float4： // host code void* devPtr; cudaD3D9ResourceGetMappedPointer(&devPtr, surface); size_t pitch; cudaD3D9ResourceGetMappedPitch(&pitch, surface); dim3 Db = dim3(16, 16); dim3 Dg = dim3((width+Db.x–1)/Db.x, (height+Db.y–1)/Db.y); myKernel<<>>((unsigned char*)devPtr, width, height, pitch); // device code __global__ void myKernel(unsigned char* surface, int width, int height, size_t pitch) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x >= width || y >= height) return;