NVIDIA CUDA统一计算设备架构编程手册
第1章 CUDA 介绍
1.1 作为数据并行计算设备的图像处理器单元
仅仅几年的时间，可编程图像处理器单元已经发展成为绝对的计算主力，如图1-1所示。由于具有由高内存带宽驱动的多个核心，今天的GPU为图像和非图像处理提供了难以置信的资源。
发展背后的主要原因是GPU是特定于计算密集的、高并行的计算，而这正是图像渲染所需要的，因此GPU设计了更多的晶体管专用于数据处理，而非数据高速缓存和流控制，如图1-2所示。
具体来说，GPU专用于解决数据并行计算（同一程序在许多数据元素上并行执行）、具有高运算密度（算术运算与内存操作的比例）的问题。因为同一程序 为每个数据元素执行，所以对高级流控制具有较低的要求；又因为在许多数据元素上执行并具有高运算密度，所以内存访问延迟可以使用计算而非大的数据高速缓存 来隐藏。
数据并行处理将数据元素映射到并行处理线程。处理大型数据集合（比如数组）的许多应用程序可以使用数据并行编程模型来加速计算。在3D渲染中，大型像素和 顶点集合映射到并行线程。同样的，图像和媒体处理应用程序，比如渲染图像的后期处理、视频编码和解码、图像扩展、立体视觉、模式识别，可以将图像块和像素 映射到并行处理线程。事实上，图像渲染和处理以外的许多算法也是可由数据并行处理来加速，范围涉及一般的信号处理或物理模拟，直至财务计算或计算生物学。
但是，直到现在，获得GPU中的所有计算能力并将其有效用于非图像应用程序中仍然是一个难题：
GPU只能通过图像API进行编程，从而把较高的学习曲线强加给新手，并且为非图像应用程序增加了不必要的API开销。
GPU DRAM可以使用一般方式来读取，即GPU程序可以从DRAM的任何部分收集数据元素，但不能使用一般方式来写入，即GPU程序不能将消息分散到DRAM的任何部分，这就大大降低了CPU上易于获取的编程灵活性。
一些应用程序受到DRAM内存带宽的瓶颈限制，未充分利用GPU的计算功能。
本文档描述了一种新颖的硬件和编程模型，直接解决这些问题，并将GPU暴露为一种真正通用的数据并行计算设备。
1.2 CUDA：一种在GPU上进行计算的新架构
CUDA表示Compute Unified Device Architecture（统一计算设备架构），是一种新型的硬件和软件架构，用于将GPU上作为数据并行计算设备在GPU上进行计算的发放和管理，而无 需将其映射到图像API。它可用于GeForce 8系列、Tesla解决方案和一些Quadro解决方案（详细信息请参阅附录A）。操作系统的多任务机制负责管理多个并发运行的CUDA和图像应用程序对 GPU的访问。
CUDA软件堆栈由几层组成，如图1-3所示：硬件驱动器，应用编程接口(API)及其runtime库，还有两个更高层的通用数学库CUFFT和CUBLAS，这两个库在单独的文档中介绍。硬件已经设计为支持轻量级驱动和runtime层，以达到高性能。
CUDA API包括了对C编程语言的扩展，以达到最低的学习曲线（参见第4章）。
CUDA提供了一般的DRAM内存寻址以实现更多的编程灵活性，如图1-4所示：分散和收集内存操作。从编程角度看，这转换为就像在CPU上一样在DRAM的任何位置读取和写入数据的能力。
CUDA提供了具有非常快速的一般读写访问的并行数据高速缓存或芯片共享内存，线程可以使用它来互相共享数据（参见第3章）。如图1-5所示，应用程序可以利用它来最小化对DRAM的过度提取和巡回，从而降低对DRAM内存带宽的依赖程度。
第2章 编程模型
2.1 高度多线程协处理器
通过CUDA编程时，将GPU看作可以并行执行非常多个线程的计算设备(compute device)。它作为主CPU的协处理器或者主机（host）来运作：换句话说，在主机上运行的应用程序中数据并行的、计算密集的部分卸载到此设备上。
更准确地说，多次但在不同数据上独立执行的应用程序部分可以独立放到在此设备上作为许多不同线程执行的函数中。要达到这种效果，可以将这样一个函数编译到设备的指令集合中，并将得到的程序（叫做内核, kernel）下载到设备上。
主机和设备都保留自己的DRAM，分别称为主机内存(host memory)和设备内存(device memory)。用户可以通过优化的API调用将数据从一个DRAM复制到其他DRAM中，而优化的API调用使用了设备的高性能直接内存访问(DMA)引擎。
2.2 线程分批
执行内核的线程批组织为线程块的网格，如2.2.1和2.2.2所述，并参见图2-1。
2.2.1 线程块
线程块是可以一起协作的线程批次，它们通过一些快速的共享内存有效地共享数据，并同步其执行以协调内存访问。更准确地说，用户可以在内核中指定同步点，块中的线程在到达此同步点时挂起。
每个线程由线程ID(thread ID)标识，这是块中的线程号。为了帮助基于线程ID的复杂寻址，应用程序还可以将块指定为任意大小的二维或三维度组，并使用2个或3个组件索引来标识每 个线程。对于大小(Dx,Dy)为的二维块，索引为(x,y)的线程的线程ID为(x+yDx)，对于大小为(Dx,Dy,Dz)的三维块，索引为 (x,y,z)的线程的线程ID为(x+yDx+zDxDy)。
2.2.2 线程块网格
块可以包含的最大线程数是有限制的。但是，执行相同内核的具有相同维度和大小的块可以分批组合到块网格中，以便可以在单个内核调用中启动的线程总数变得更 大。这是以线程协作的降低为代价的，因为同一网格中不同线程块中的线程不能互相通信和同步。此模型允许内核有效运行，而不必在具有不同并行能力的各种设备 上重新编译：如果设备具有非常少的并行能力，则可以顺序运行网格的所有块，如果具有很多并行能力，则可以并行运行网格的所有块，通常是二者组合使用。
每个块由其块ID标识，这是网格中的块号。为了帮助基于块ID的复杂寻址，应用程序还可以将网格指定为任意大小的二维度组，并使用2个组件索引来标识每个块。对于大小(Dx,Dy)为的二维块，索引为(x,y)的块的块ID为(x+yDx)。
主机执行一连串对设备的内核调用。每个内核作为组织为线程块网格的一批线程来执行。
图2-1. 线程分批
2.3 内存模型
在设备上执行的线程只能通过下列内存空间访问设备的DRAM和芯片内存储单元，如图2-2所示：
1、读写每线程寄存器，
2、读写每线程本地内存，
3、读写每块共享内存，
4、读写每网格全局内存，
5、只读每网格常量内存，
6、只读每网格纹理内存。
全局、常量和纹理内存空间可以通过主机读或写，并永久存在于相同应用程序的内核启动中。
全局、常量和纹理内存空间为不同的内存使用进行了优化（参见5.1.2.1、5.1.2.2和5.1.2.3）。纹理内存还为一些特定的数据格式提供不同的寻址模式以及数据筛选（参见4.3.4）。
线程可以通过不同范围的一组内存空间来访问设备的DRAM和芯片内存。
图2-2. 内存模型
第3章 硬件实现
3.1 具有芯片共享内存的一组SIMD多处理器
设备作为一组多处理器(multiprocessors)来实现，如图3-1所示。每个多处理器具有单指令多数据(SIMD)架构：在任何给定的时钟周期，多处理器的每个处理器执行相同的指令，但操作在不同的数据上。
每个多处理器具有下列四种类型的芯片内存储器：
1、每个处理器有一组本地32位寄存器
2、并行数据高速缓存或称为共享内存(shared memory)，由所有处理器共享并实现共享内存空间
3、只读常量高速缓存(constant cache)，由所有处理器共享并加速从常量内存空间的读取，实现为设备内存的只读区域
4、只读纹理高速缓存(texture cache)，由所有处理器共享并加速从纹理内存空间的读取，实现为设备内存的只读区域
本地和全局内存空间为设备内存的读写区域，且无高速缓存。
每个多处理器通过纹理单位(texture unit)访问纹理高速缓存，其中纹理单位实现2.3一节提到的各种寻址模式和数据筛选。
具有芯片共享内存的一组SIMD多处理器。
图3-1. 硬件模型
3.2 执行模型
线程块网格是通过调度块在多处理器上执行来在设备上执行的。每个多处理器一批接一批地处理块批。一个块仅由一个多处理器处理，所以共享内存空间驻留在芯片共享内存中，从而导致非常快的内存访问。
因为多处理器的寄存器和共享内存划分给块批的所有线程，所以每个多处理器一批可以处理多少块取决于给定内核每秒需要多少寄存器以及每块需要多少共享内存。如果每个多处理器没有足够的可用寄存器或共享内存来处理至少一个块，则内核将无法启动。
在一个批次内并被一个多处理器处理的块称为活动(active)块。每个活动块划分到称为warp的SIMD线程组中：其中每个warp包含相同数 量的线程，称为warp大小，并以SIMD方式由多处理器执行。活动warp——比如所有活动块中的所有warp——是分时的：线程调度器（thread scheduler）定期从一个warp切换到另一个warp，以便最大化多处理器计算资源的使用。半warp（half-warp）是一个warp的第 一半或第二半。
块划分为warp的方式始终相同；每个warp包含线程ID连续递增的线程，其中第一个warp包含线程0。2.2.1一节介绍线程ID与块中的线程索引如何相关联。
块中warp的执行顺序没有定义，但其执行可以同步以协调全局或共享内存访问，如2.2.1所述。
线程块网格中块的执行顺序没有定义，且块之间没有同步机制，所以在网格执行期间，来自同一网格的两个不同块中的线程无法通过全局内存安全地互相通信。
如果对于warp的多个线程，由warp执行的非完整（non-atomic）指令写入全局或共享内存中的相同位置，则此位置发生的序列化写入数目及其发 生顺序没有定义，但会保证其中一个写入成功。如果由warp执行的完整（atomic）指令（参见4.4.6）读取、修改并写入warp多个线程的全局内 存中的相同位置，则对此位置的每个读取、修改和写入都会发生，且全部序列化，但发生的顺序没有定义。
3.3 计算能力
设备的计算能力(compute capability)由主要修订号和次要修订号来定义。
具有相同主要修订号的设备具有相同的核心架构。附录A中列出的设备都具有计算能力1.x（其主要修订号为1）。
次要修订号与核心结构的增量改进相对应，其中可能包括新功能。
各种计算能力的技术规范在附录A中给出。
3.4 多个设备
使用多个GPU作为CUDA设备在多GPU系统上由应用程序运行时，仅当这些GPU具有相同的类型时，才能保证工作。但是，如果系统处于SLI模式，则只 有一个GPU可以用作CUDA设备，因为在驱动器堆栈的最低层，所有GPU都熔合在一起。需要在控制面板中关闭SLI模式，CUDA才能将每个GPU看作 单独的设备。
3.5 模式开关
GPU将一些DRAM内存指定为所谓的主表面（primary surface），主表面用于用户查看的显示设备的显示刷新。当用户通过更改显示的分辨率或位深度（使用NVIDIA控制面板或Windows中的显示控 制面板）来启动显示的模式开关（mode switch）时，主表面所需的内存量将随之变化。例如，如果用户将显示分辨率从1280x1024x32位更改为1280x1024x32位时，系统必 须将7.68MB而非5.24MB指定为主表面。（运行时启用了反混淆的全屏图像应用程序需要更多的显示内存作为主表面。）在Windows中，可以启动 显示模式开关的其他事件包括启动全屏DirectX应用程序、按下Alt+Tab将任务切换出全屏DirectX应用程序或按下Ctrl+Alt+Del 锁定计算机。
如果模式开关增加主要表明所需的内存量，系统可能必须抽调指定给CUDA应用程序的内存分配，从而导致这些应用程序的崩溃。
第4章 应用编程接口
4.1 C编程语言扩展
CUDA编程接口的目标是为熟悉C编程语言的用户提供相对简单的路径，以便容易地编写在设备上执行的程序。
它包括：
C语言的最小扩展集合，如4.2所述，允许程序员定位要在设备上执行的部分源码；
运行时runtime库划分为：
主机组件，如4.5所述，在主机上运行，提供函数以控制并访问主机中的一个或多个计算设备；
设备组件，如4.4所述，在设备上运行，并提供特定于设备的函数；
通用组件，如4.3所述，提供内置的向量类型，以及主机和设备代码中都支持的C标准库子集。
必须强调一下，只有支持在设备上运行的C标准库中的函数才是公共运行时runtime组件提供的函数。
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__函数必须具有void返回类型。
对__global__函数的任何调用必须指定其执行配置，如4.2.3所述。
对__global__函数的调用是异步的，这意味着在设备完成其执行之前返回。
__global__函数参数当前通过共享内存传递给设备并限制为256字节。
4.2.2 变量类型限定符
4.2.2.1 __device__
__device__限定符声明驻留在设备上的变量。
下面三节中定义的其他类型限定符中至多一个可以与__device__一起使用，以进一步指定变量属于哪个内存空间。如果其中任何一个都不出现，则变量：
驻留在全局内存空间中，
具有应用程序的生命期，
可通过runtime库从网格中的所有线程中和从主机中访问。
4.2.2.2 __constant__
__constant__限定符，可以与__device__一起使用，声明变量：
驻留在常量内存空间中，
具有应用程序的生命期，
可通过runtime库从网格中的所有线程中和从主机中访问。
4.2.2.3 __shared__
__shared__限定符，可以与__device__一起使用，声明变量：
驻留在线程块的共享内存空间中，
具有块的生命期，
仅可从块内的所有线程中访问。
线程中共享变量的完全顺序一致性，但是在线程中放松的排序。仅在__syncthreads()（参见4.4.2）执行之后，来自其他线程的写入才能保证可见。除非变量声明为挥发，否则只要满足上一语句，编译器就可以优化对共享内存的读写。
将共享内存中的变量声明为外部数据，比如
数组大小在启动时确定（参见4.2.3）。以此方式声明的所有变量在内存中从同一地址开始，从而数组中变量的布局必须通过偏移量明确管理。例如，如果用户想要
位于动态分配的共享内存中，则用户可以使用下列方式声明和初始化数组：
4.2.2.4 限制
这些限定符不允许用于struct和union成员、形参以及在主机上执行的函数内部的本地变量。
__shared__和__constant__变量已经隐含了静态存储。
__device__、__shared__和__constant__变量不能使用extern关键字定义为外部。
__device__和__constant__变量仅允许用于文件范围。
__constant__变量不能从设备中赋值，只能从主机中通过主机运行时runtime函数赋值（参见4.5.2.3和4.5.3.6）。
__shared__变量不能在声明中进行初始化。
在设备代码中声明的不带有其中任何一个限定符的自动变量一般驻留在寄存器中。但是，在一些情况下，编译器可能选择将其放置在本地内存中。这通常适用于将耗 费太多寄存器空间的大型结构或数组，以及编译器无法确定其是否使用常数进行索引的数组。检查ptx汇编代码（通过使用-ptx或-keep选项编译获得） 将告知变量是否已经在第一个编译阶段放置在本地内存中，此时此变量将使用.local助记符声明并使用ld.local和st.local助记符访问。如 果没有，则后续编译阶段可能仍会确定，尽管它们发现此变量在目标架构中耗费了太多寄存器空间。这可以通过使用报告本地内存使用(lmem)的 --ptxas-options=-v选项来检查。
只要编译器能够解析在设备上执行的代码中的指针是指向共享内存空间还是全局内存空间，就支持这些指针，否则，就限制这些指针只能指向在全局内存空间中分配或声明内存。
析取指向在主机上执行的代码中的全局或共享内存的指针或析取指向在设备上执行的代码中的主机内存的指针将导致未定义的行为，通常是分段错误和应用程序终止。
通过提取__device__、__shared__或__constant__变量的指针获得的地址只能用在设备代码中。通过 cudaGetSymbolAddress()（参见4.5.2.3）获得的__device__或__constant__变量的地址只能用在主机代码 中。
4.2.3 执行配置
对__global__函数的任何调用必须为此调用指定执行配置。
执行配置定义将用于在设备上执行函数的网格和块的维度，以及相关联的流（有关流的介绍，参见4.5.1.5）。通过在函数名称和圆括号括起的参数列表之间 插入<<< Dg, Db, Ns, S >>>形式的表达式，来定义执行配置，其中：
Dg是类型dim3（参见4.3.1.2），用于指定网格的维度和大小，因此Dg.x*Dg.y等于要启动的块数；Dg.z未使用；
Db是类型dim3（参见4.3.1.2），用于指定每块的维度和大小，因此Dg.x*Dg.y*Db.z等于每块的线程数；
Ns是类型size_t，用于指定为此调用按块动态分配的共享内存中的字节数以及静态分配的内存；此动态分配的内存由声明为外部数组的任何一个变量使用，如4.2.2.3所述；Ns是默认值为0的可选参数；
S是类型cudaStream_t，用于指定相关联的流；S是默认值为0的可选参数。
例如，函数声明为
必须使用下列方式调用：
执行配置的参数在实际函数参数之前求值，并且与函数参数一样，当前通过共享内存传递给设备。
如果Dg或Db大于附录A.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 限制
不允许提取任何内置变量的地址。
不允许为任何内置变量赋值。
4.2.5 使用NVCC编译
nvcc是用于简化CUDA代码编译过程的编译器驱动程序：它提供简单熟悉的命令行选项，并通过调用用于实现不同编译阶段的工具集合来执行这些选项。
nvcc的基本工作流包括将设备代码与主机代码分开，并将设备代码编译为二进制形式或cubin对象。生成的主机代码是输出，此输出或者作为要使用另一个工具编译的C代码，或者作为最后一个编译阶段直接调用主机编译器的对象代码。
应用程序可以忽略生成的主机代码，并使用CUDA驱动程序API（参见4.5.3）加载并执行设备上的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指令可用于控制任何给定循环的展开。它必须立即放置在循环前，并只应用于此循环。它后面可以跟一个数字，用于指定循环必须展开多少次。
例如，在下列代码示例中：
循环将展开5次。程序员应该确保展开将不影响程序的正确性（在上例中，如果n小于5，则可能影响）。
#pragma unroll 1将禁止编译器展开循环。
如果在#pragma unroll后面不指定任何数字，如果其循环计数是常数，则循环将完全展开，否则根本不展开。
4.3 公共运行时组件
公共runtime运行时组件可以由主机和设备函数使用。
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
这些向量类型是从基本整数和浮点数类型派生而来的。它们是结构体，第1个、第2个、第3个和第4个组件分别可以通过字段x、y、z和w来访问。它们都具有
make_<type name>形式的构造函数；例如，
这将创建类型为int2、值为(x, y)的向量。
4.3.1.2 dim3类型
此类型是基于uint3的整数向量类型，用于指定维度。定义dim3类型的变量时，保留为未指定的任何组件将初始化为1。
4.3.2 数学函数
表B-1包含当前支持的C/C++标准库数学函数的完整列表，以及在设备上执行时各自的误差界。
在主机代码中执行时，给定函数使用可用的Cruntime执行。
4.3.3 时间函数
返回在每个时钟周期递增的计数器的值。
在内核开始和结束时对此计数器取样，求出两次取样之差，并记录每个线程的结果，从而计量设备完全执行每个线程所用的时钟周期数，但这并非设备实际执行线程指令所用的时钟周期数。前者大于后者，因为线程是分时执行的。
4.3.4 纹理类型
CUDA支持GPU本用于图形的纹理硬件以访问纹理内存。从纹理内存而非全局内存读取数据可以具有5.4一节描述的几个性能优点。
纹理内存使用名为纹理拾取(texture fetches)的设备函数来从内核中读取，如4.4.5所述。纹理拾取的第一个参数指定一个名为纹理参考(texture reference)的对象。
纹理参考定义要拾取哪部分纹理内存。它必须通过主机runtime函数（参见4.5.2.6和4.5.3.9）绑定到一些内存区域（称为纹理(texture)），然后才能供内核使用。几个不同的纹理参考可以绑定到同一纹理或在内存中重叠的纹理。
纹理参考具有多个属性。其中之一是其维度，用于指定纹理是使用一个纹理坐标作为一维度组进行寻址，还是使用两个纹理坐标作为二维度组进行寻址。数组的元素称为纹理元素（texel，是“texture element”的简写）。
其他属性定义纹理拾取的输入和输出数据类型，如何解释输入坐标，以及应执行什么处理。
4.3.4.1 纹理参考声明
纹理参考的一些属性不可变，而且必须在编译时已知；它们在声明纹理参考时指定。纹理参考在文件范围内声明为texture类型的变量：
其中：
Type指定拾取纹理时返回的数据类型；Type限制为基本的整数和浮点类型，以及4.3.1.1一节中定义的1-、2-和4-组件向量类型之一。
Dim指定纹理拾取的维度，等于1或2；Dim是可选参数，默认为1；
ReadMode等于cudaReadModeNormalizedFloat或cudaReadModeElementType；如果 ReadMode为cudaReadModeNormalizedFloat，且Type为16位或8位整数类型，则其值实际返回为浮点类型，且整数类型 的完整范围对于无符号整数类型映射为[0.0, 1.0]，对于有符号整数类型映射为[-1.0, 1.0]；例如，值为0xff的无符号8位纹理元素读取为1；如果ReadMode为cudaReadModeElementType，则不执行任何转 换；ReadMode是可选参数，默认为cudaReadModeElementType。
4.3.4.2 Runtime运行时纹理参考属性
纹理参考的其他属性是可变的，可以在运行时通过主机运行时（运行时API参见4.5.2.6，驱动程序API参见4.5.3.9）进行更改。它们指定纹理坐标是否规格化、寻址模式和纹理筛选，详见下文。
默认情况下，使用[0, N)范围内的浮点坐标参考纹理，其中N是与坐标相对应的维度中的纹理大小。例如，大小为64×32的纹理将分别在x和y维度上使用范围[0, 63]和[0, 31]中的坐标来参考。规格化纹理坐标将导致坐标在范围[0.0, 1.0)中而非范围[0, N)中指定，因此，同一64×32纹理将在x和y维度上都使用范围[0, 1)中的规格化坐标来寻址。如果纹理坐标独立于纹理大小更可取，那么规格化纹理坐标对于一些应用程序的要求是一个自然的选择。
寻址模式定义纹理坐标超出范围时要执行的操作。使用非规格化纹理坐标时，超出范围[0, N)的纹理坐标将固定：0以下的值设置为0，大于0或等于N的值设置为N-1。使用规格化纹理坐标时，固定也是默认的寻址模式：小于0.0或大于1.0的 值固定到范围[0.0, 1.0)中。对于规格化坐标，也可以指定“wrap”寻址模式。当纹理包含定期信号时，通常使用wrap寻址。wrap寻址仅使用纹理坐标的小数部分；例 如，1.25当作0.25处理，-1.25当作0.75处理。
只能为配置为返回浮点数据的纹理执行线性纹理筛选。线性纹理筛选在相邻纹理元素之间执行低精度插值。启用线性纹理筛选时，将读取纹理拾取位置周围的纹理元 素，并基于纹理坐标落入纹理元素之间的位置来插入纹理拾取的返回值。对于一维纹理执行简单线性插值，对于二维纹理执行双线性插值。
附录F给出有关纹理拾取的更多详细信息。
4.3.4.3 纹理来自线性内存对来自CUDA数组
纹理可以是线性内存或CUDA数组的任何区域（参见4.5.1.2）。
在线性内存中分配的的纹理：
维度只能等于等于1；
不支持纹理筛选；
只能使用非规格化整数纹理坐标来寻址；
不支持不同的寻址模式：超出范围的纹理访问返回零。
硬件在纹理基地址上强制执行对齐要求。为了让程序员忽略此对齐要求，将纹理参考绑定到设备内存的函数传回一个必须应用到纹理拾取的字节偏移，以便从所需内 存中读取。由CUDA的分配例程返回的基指针符合此对齐约束，因此通过将已分配的指针传递给 cudaBindTexture()/cuTexRefSetAddress()，应用程序可以完全避免偏移。
4.4 设备Runtime组件
设备runtime组件只能在设备函数中使用。
4.4.1 数学函数
对于表B-1中的一些函数，设备runtime运行时组件中存在一些不太准确但运行较快的版本；这些函数具有相同的名称，但加了前辍__（例如__sin(x)）。表B-2中列出这些固有函数及其各自的误差界。
编译器具有一个选项（-use_fast_math）来强制每个函数编译为其不太精确的对应物（如果存在的话）。
4.4.2 同步函数
同步块中所有的线程。当所有线程都达到此点时，执行正常继续。
__syncthreads()用于协调同一块的线程之间的通信。当块中的一些线程访问共享或全局内存中的同一地址时，对于其中的一些内存访问，存在潜在的读后写、写后读或写后写的危害。这些数据危险可以通过同步这些访问之间的线程来避免。
__syncthreads()允许出现在条件代码中，但仅当条件在整个线程块中求值相同时才允许，否则代码执行可能暂挂或产生非预期的副作用。
4.4.3 类型转换函数
下列函数中的后缀指明IEEE-754取整模式：
rn取整为最近的偶数，
rz向零取整，
ru向上取整（到正无穷大），
rd向下取整（到负无穷大）。
使用指定的取整模式将浮点参数转换为整数。
使用指定的取整模式将浮点参数转换为无符号整数。
使用指定的取整模式将整数参数转换为浮点数。
使用指定的取整模式将无符号整数参数转换为浮点数。
4.4.4 类型转换函数
对整数参数执行浮点类型转换，保留值不变。例如，__int_as_float(0xC0000000)等于-2。
对浮点参数执行整数类型转换，保留值不变。例如__float_as_int(1.0f)等于0x3f800000。
4.4.5 纹理函数
4.4.5.1 从设备内存取纹理
从设备内存取纹理时，使用函数的tex1Dfetch()族访问纹理；例如：
这些函数使用纹理坐标x拾取绑定到纹理参考texRef的线性内存区域。不支持任何纹理筛选和寻址模式。对于整数类型，这些函数可以有选择地将整数提升为32位浮点。
除上述函数之外，还支持2元组和4元组；例如：
使用纹理坐标x拾取绑定到纹理参考texRef的线性内存。
4.4.5.2 从CUDA数组取纹理
从CUDA数组取纹理时，使用tex1D()或tex2D()访问纹理：
这些函数使用纹理坐标x和y拾取绑定到纹理参考texRef的CUDA数组。纹理参考的不变属性（编译时）和可变属性（运行时）的组合确定如何解释坐标、在纹理拾取期间执行何种处理、以及纹理拾取传递的返回值（参见4.3.4.1和4.3.4.2）。
4.4.6 原子函数(Atomic Functions)
原子函数仅可用于计算能力1.1的设备。附录C列出了这些函数。
原子函数在驻留于全局内存中的一个32位字上执行读-改-写原子操作。例如，atomicAdd()在全局内存中的同一地址上读取一个32位字，为其加上 一个整数，然后将结果写回同一地址。在保证执行时不受其他线程干扰这种意义上，此操作是原子的。换句话说，只有此操作完成之后，其他线程才可以访问此地 址。
原子操作仅适用于32位有符号的和无符号的整数。
4.5 主机Runtime组件
主机runtime组件只能由主机函数使用。
它提供函数来处理：
设备管理，
上下文管理，
内存管理，
代码模块管理，
执行控制，
上下文参考管理，
OpenGL和Direct3D的互操作性。
它由两个API组成：
名为CUDA驱动程序API的低层API，
名为CUDAruntime API的高层API，在CUDA驱动程序API之上实现。
这些API互斥：应用程序应该使用其中之一。
CUDAruntime通过提供隐式初始化、上下文管理和模块管理使得设备代码管理变得容易。由ncvv生成的C主机代码基于CUDAruntime（参见4.2.5），因此链接到此代码的应用程序必须使用CUDAruntime API。
相反，CUDA驱动程序API需要更多的代码，更难于编程和调试，但是它提供较高级的控制，而且因为它仅处理cubin对象（参见4.2.5），所以是独 立于语言的。特别地，使用CUDA驱动程序API配置和启动内核比较困难，因为执行配置和内核参数必须使用隐式函数调用来指定，而不是使用4.2.3中所 述的执行配置语法来指定。另外，设备仿真（参见4.5.2.7）不使用CUDA驱动程序API。
CUDA驱动程序API通过cuda动态库传递，并且它所有的入口点都带有前缀cu。
CUDAruntime API通过cudart动态库传递，并且它所有的入口点都带有前缀cuda。
4.5.1 常用概念
4.5.1.1 设备
两种API都提供函数来列举系统上可用的设备、查询其属性并为内核执行选择其中之一（runtime API参见4.5.2.2，驱动程序API参见4.5.3.2）。
多个主机线程可以在同一设备上执行设备代码，但在设计上，一个主机线程只能在一个设备上执行设备代码。因此，在多个设备上执行设备代码需要多个主机线程。另外，通过一个主机线程中的runtime创建的任何CUDA资源不能由其他主机线程中的rutime使用。
4.5.1.2 内存
设备内存可以分配为线性内存或CUDA数组。
线性内存在设备上以32位地址空间存在，因此单独分配的实体可以通过指针互相引用，例如在二叉树中。
CUDA数组是为纹理拾取而优化的不透明内存布局（参见4.3.4）。CUDA数组是一维或二维度组，由元素组成，每个元素具有1、2或4个组件，这些组 件可以是有符号或无符号的8、16或32位整数、16位浮点数（当前仅通过驱动程序API支持）或32位浮点数。CUDA数组只能由内核通过纹理拾取来读 取，且只能绑定到具有相同数目包装组件的纹理参考。
线性内存和CUDA数组都可由主机通过内存复制函数（如4.5.2.3和4.5.3.6所述）读取和写入。
主机runtime还提供了函数以分配和释放页面锁定的主机内存——与由malloc()分配的正常可分页主机内存相反（runtime API参见D.5.6和D.5.7，驱动程序API参见E.8.5和E.8.6）。页面锁定内存的一个优点是如果主机内存分配为页面锁定，则主机内存和设 备内存之间的带宽较高——仅用于由分配主机内存的主机线程执行的数据传送。但是，页面锁定内存是稀有资源，所以早在可分页内存中的分配之前，页面锁定内存 中的分配就将开始失败。此外，通过减少可用于操作系统分页的物理内存量，分配太多的页面锁定内存将降低整体系统性能。
4.5.1.3 OpenGL互操作性
OpenGL缓冲对象可以映射到CUDA的地址空间中，从而允许CUDA读取由OpenGL写入的数据，或允许CUDA写入供OpenGL消耗的数据。 4.5.2.7一节描述如果使用runtime API完成此操作，4.5.3.10一节描述如何使用驱动程序API完成此操作。
4.5.1.4 Direct3D互操作性
Direct3D 9.0顶点缓冲可以映射到CUDA的地址空间，从而允许CUDA读取由Direct3D写入的数据，或允许CUDA写入供Direct3D消耗的数据。 4.5.2.8一节描述如果使用runtime API完成此操作，4.5.2.8一节描述如何使用驱动程序API完成此操作。
CUDA上下文和Direct3D设备必须在同一GPU上创建。这可以通过查询与Direct3D使用的适配器相应的CUDA设备来确保这一点，对于 runtime API使用cudaD3D9GetDevice()（参见D.9.7），对于驱动程序API使用cuD3D9GetDevice()（参见 E.11.7）。
Direct3D设备还必须使用D3DCREATE_HARDWARE_VERTEXPROCESSING标记来创建。
CUDA还不支持：
除Direct3D 9.0之外的版本，
除顶点缓冲之外的Direct3D对象。
顺便提一句，当Direct3D和CUDA之间的负载均衡优先于互操作性时，cuda3D9GetDevice()或cuD3D9GetDevice()还可以用于确保Direct3D和CUDA创建在不同的设备上。
4.5.1.5 异步并发执行
为了方便主机和设备之间的并发执行，一些runtime 函数是异步的：在设备已经完成请求的任务之前，控制返回给应用程序。这些函数包括：
内核通过__global__函数或cuGridLaunch()和cuGridLaunchAsync()启动；
执行内存复制并以Async为后缀的函数；
执行设备↔设备内存复制的函数；
设置内存的函数。
一些设备还可以使用内核执行并发地在页面锁定主机内存和设备内存之间执行复制。应用程序可以通过使用 CU_DEVICE_ATTRIBUTE_GPU_OVERLAP调用cuDeviceGetAttribute()来查询此功能（请分别参见 E.2.6）。当前只有不涉及通过cudaMallocPitch()（参见4.5.2.3）或cuMemAllocPitch()（参见 4.5.3.6）分配的CUDA数组或2D数组的内存复制，才支持此功能。
应用程序通过流(streams)并发管理。流是一个顺序执行的操作序列。另一方面，不同的流可以不按顺序执行其操作，或并发执行其操作。
通过创建流对象并将其指定为一序列的内核启动和主机↔内存复制的流参数，可以定义流。4.5.2.4描述如何使用runtime API完成此操作，4.5.3.7介绍如何使用驱动程序API完成此操作。
仅当所有先前的操作（包括属于流部分的操作）完成之后，已经指定零流参数的任何内核启动、内存设置或内存复制才能开始，而且在它完成之前，任何后续操作都不能开始。
runtime API的cudaStreamQuery()和驱动程序API的cuStreamQuery()（请分别参见D.3.2和E.5.2）提供应用程序来确定 流中所有先前的操作是否已经完成。runtime API的cudaStreamSynchronize()和驱动程序cuStreamSynchronize()（请分别参见E.5.2和E.5.3）提 供了一种方法，来明确强制runtime 在流中所有先前的操作完成之前等待。
同样地，使用runtime API的cudaThreadSynchronize()和驱动程序API的cuCtxSynchronize()（请分别参见D.2.1和 E.3.5），应用程序可以强制runtime 在所有先前的设备任务完成之前等待。为了避免不必要的减速，这些函数最适合用于定时目的，或用于隔离失败的启动或内存复制。
通过允许应用程序记录程序中任何点的事件，并查询实际记录这些事件的时间，runtime 还提供了一种方法来密切监控设备的进度并执行确准的定时。当事件之前的所有任务——或可选地，给定流中的所有操作——都已完成时，记录此事件。 4.5.2.5描述如何使用runtime API完成此操作，4.5.3.8描述如何使用驱动程序API完成此操作。
如果页面锁定主机内存分配、设备内存分配、设备内存设置、设备↔设备内存复制或事件记录在不同流中的两个操作之间发生，则这两个操作不能并发执行。
程序员可以通过将CUDA_LAUNCH_BLOCKING环境变量设置为1，全局禁用系统上运行的所有CUDA应用程序的异步执行。此功能只提供用于调试目的，而且绝不能用作让生产软件可靠运行的方法。
4.5.2 Runtime API
4.5.2.1 初始化
runtime API没有任何显式初始化函数；第一次调用runtime 函数时，runtime API初始化。当定时runtime 函数调用时，以及将第一次调用的错误代码解释到runtime 中时，用户一定要记住这一点。
4.5.2.2 设备管理
章节D.1中的函数用于管理呈现在系统中的设备。
cudaGetDeviceCount()和cudaGetDeviceProperties()提供一种用来列举这些设备并检索其属性的方法。
cudaSetDevice()用来选择与主机线程相关的设备：
在调用任何__global__函数或任何附录D中的函数之前必须选择设备。如果显式调用cudaSetDevice()没有执行，将自动选择设备0，且任何随后的显式调用cudaSetDevice()将不起作用。
4.5.2.3 内存管理
D.5中的函数用于分配和释放设备内存，访问为全局内存空间中声明的任何变量分配的内存，并在主机和设备内存之间传送数据。
使用cudaMalloc()或cudaMallocPitch()分配线性内存，使用cudaFree()释放线性内存。
下列代码示例在线性内存中分配了256个浮点元素的数组：
建议使用cudaMallocPitch()进行2D数组的分配，因为它确保了分配适当填补以满足5.1.2.1中描述的对齐要求，从而确保在访问 行地址时或在2D数组和其他设备内存区域执行复制（使用cudaMemcpy2D()函数）时获得最佳性能。返回的pitch（或跨度）必须用于访问数组 元素。下列代码示例分配浮点数值的width×height的2D数组，并显示如何在设备代码中循环处理数组元素：
CUDA数组使用cudaMallocArray()进行分配，使用cudaFreeArray()进行释放。cudaMallocArray()需要使用cudaCreateChannelDesc()创建的格式描述。
下列代码示例分配一个32位浮点组件的高度×高度CUDA数组：
cudaGetSymbolAddress()用于检索指向为全局内存空间中声明的变量分配的地址。已分配内存的大小通过cudaGetSymbolSize()来获得。
D.5一节列出用于在使用cudaMalloc()分配的线性内存、使用cudaMallocPitch()分配的线性内存、CUDA数组和为全局或常量内存空间中声明的变量分配的内存之间复制内存的所有各种函数。
下列代码示例将2D数组复制到在上一代码示例中分配的CUDA数组：
下列代码示例将一些主机内存数组复制到设备内存中：
下列代码示例将一些主机内存数组复制到设备内存中：
下列代码示例将一些主机内存数组复制到常量内存中：
4.5.2.4 流管理
D.3一节中的函数用于创建和销毁流，并确定流的所有操作是否已经完成。
下列代码示例创建两个流：
其中每个流通过下列代码示例定义为一个从主机到设备的内存复制、一个内核启动和一个从设备到主机的内存复制的序列：
每个流将其输入数组hostPtr部分复制到设备内存中的数组inputDevPtr中，通过调用myKernel()处理设备上的 inputDevPtr，并将结果outputDevPtr重新复制到同一hostPtr部分。使用两个流处理处理hostPtr允许一个流的内存复制与 其他流的内核执行相重叠。hostPtr必须指向要发生的任何重叠的页面锁定主机内存：
最后调用cudaThreadSynchronize()以确保在进一步处理之前所有流都已完成。
4.5.2.5 事件管理
D.4一节中的函数用于创建、记录和销毁事件，并查询两个事件之间用去的时间。
下列代码示例创建两个事件：
这些事件可以使用下列方法用于定时上一节的代码示例：
4.5.2.6 纹理参考管理
D.6一节的函数用于管理纹理参考。
由高层API定义的texture类型是一种从由低层API定义的textureReference类型中公共派生出来的结构，如下所示：
normalized指定纹理坐标是否规格化；如果它为非零，则纹理中的所有元素都使用范围[0,1]而非范围[0,width-1]或[0,height-1]中的纹理坐标来寻址，其中width和height是纹理大小；
filterMode指定筛选模式，即当拾取纹理时，如何基于输入纹理坐标来计算返回的值； filterMode等于cudaFilterModePoint或cudaFilterModeLinear；如果它为 cudaFilterModePoint，则返回的值是纹理坐标最接近输入纹理坐标的纹理元素；如果它为cudaFilterModeLinear，则返 回的值是纹理坐标最接近输入纹理坐标的两个（对于一维纹理）或四个（对于二维纹理）纹理元素；cudaFilterModeLinear仅对浮点类型的返 回值有效；
addressMode指定寻址模式，即如何处理超出范围的纹理坐标；addressMode是大小为2的数组，其第一个和第二个元素分别指定第一个 和第二个纹理坐标的寻址模式；寻址模式等于cudaAddressModeClamp，在这种情况下，超出范围的纹理坐标将固定到有效范围，或等于 cudaAddressModeWrap，在这种情况下，超出范围的纹理坐标将包装到有效范围；cudaAddressModeWrap仅支持规格化的纹 理坐标；
hannelDesc描述拾取纹理时返回的值的格式；channelDesc具有下列类型：
其中，x、y、z和w等于返回值的每个组件的位数，f：
如果这些组件为有符号整数类型，则为cudaChannelFormatKindSigned，
如果这些组件为无符号整数类型，则为cudaChannelFormatKindUnsigned，
如果这些组件为浮点类型，则为cudaChannelFormatKindFloat。
normalized、addressMode和filterMode可以直接在主机代码中修改。它们仅适用于绑定到CUDA数组的纹理参考。
必须使用cudaBindTexture()或cudaBindTextureToArray()将纹理参考绑定到纹理之后，内核才可以使用纹理参考从纹理内存中读取。
下列代码示例将纹理参考绑定到devPtr指向的线性内存：
使用低层API：
使用高层API：
以下代码示例将纹理参考绑定到一个CUDA数组 cuArray：
使用低层API：
使用高层API：
将纹理绑定到纹理参考时指定的格式必须与声明纹理参考时指定的参数相匹配；否则，纹理拾取的结果将无定义。
cudaUnbindTexture()用于解除对纹理参考的绑定。
4.5.2.7 OpenGL互操作性
D.8一节中的函数用于控制与OpenGL的互操作性。缓冲对象必须注册到CUDA之后才能映射。此操作使用cudaGLRegisterBufferObject()来完成：
注册之后，内核可以使用由cudaGLMapBufferObject()返回的设备内存地址读取或写入缓冲对象：
使用cudaGLUnmapBufferObject()解除映射，使用cudaGLUnregisterBufferObject()解除注册。
4.5.2.8 Direct3D互操作性
D.9一节中的函数用于控制与Direct3D的互操作性。
与Direct3D的互操作性必须使用cudaD3D9Begin()初始化，使用cudaD3D9End()终止。
在这些调用之间，顶点对象必须注册到CUDA之后才能映射。此操作使用cudaD3D9RegisterVertexBuffer()来完成：
注册之后，内核可以使用由cudaD3D9MapVertexBuffer()返回的设备内存地址读取或写入顶点缓冲：
使用cudaD3D9UnmapVertexBuffer()解除映射，使用cudaD3D9UnregisterVertexBuffer()解除注册。
4.5.2.9 使用设备仿真模式调试
编程环境不包括对设备上运行的代码的任何原生调试支持，但提供了用于调试的设备仿真模式。在此模式下编译应用程序（使用-deviceemu选项）时，设 备代码在主机上编译和运行，从而允许程序员使用主机的原生调试支持来调试应用程序，就像此应用程序是主机应用程序。预处理器宏 __DEVICE_EMULATION__在此模式下定义。 应用程序的所有代码，其中包括使用的任何库，对于设备仿真或设备执行必须一致编译。将为设备仿真编译的代码与为设备执行编译的代码链接在一起将导致在初始 化时返回下列runtime 错误：cudaErrorMixedDeviceExecution。
在设备仿真模式下运行应用程序时，编程模型由runtime 仿真。对于线程块中的每个线程，runtime 在主机上创建一个线程。程序员必须确保：
主机能够运行的最多线程数是每块的最大线程数加上一个主线程。
有足够的内存可用于运行所有线程，并确定每个线程获得256KB的堆栈。
通过设备仿真模式提供的许多功能使其成为一个非常有效的调试工具：
通过使用主机的原生调试支持，程序员可以好似用调试器支持的所有功能，比如设置断点和检查数据。
因为设备代码编译后在主机上运行，所以代码可以使用不能在设备上运行的代码来增加，比如到文件或屏幕的输入和输出操作（printf()等）。
因为所有的数据驻留在主机上，所以任何特定于设备或主机的数据可以从设备或主机代码上读取；同样地，任何设备或主机函数可以从设备或主机代码中调用。
如果错误使用了内部同步，则runtime将检测到死锁情况。
程序员必须切记，设备仿真模式是在仿真设备，而非模拟设备。因此，设备仿真模式在查找算法错误时十分有用，但某些错误难以查找：
当网格中的多个线程可能同时访问某个内存位置时，则在设备仿真模式下运行的结果可能与在设备上runtime的结果不同，因为在仿真模式下，线程顺序执行。
当解参考指向主机上全局内存的指针或指向设备上主机内存的指针时，设备执行几乎肯定以一些未定义的方式失败，而设备仿真则可以生成正确的结果。
大多数时候，在设备上执行时与在设备仿真模式下的主机上执行时，同一浮点计算将不会生成完全相同的结果。这是预期结果，因为一般来说，要让同一浮点计算获得不同的结果，只需使用略有不同的编译器选项，更不要说不同的编译器、不同的指令集或不同的架构。
特别地，一些主机平台将单精度浮点计算的中间结果存储在扩展的精度寄存器中，这可能造成在设备仿真模式下runtime精度有显著差异。当这种情况发生时，程序员可以尝试下列任何方法，但不能保证可行：
将一些浮点变量声明为挥发，以强制单精度存储；
使用gcc的–ffloat-store编译器选项，
使用Visual C++编译器的/Op或/fp编译器选项，
在Linux上使用_FPU_GETCW()和_FPU_SETCW()，或在Windows上使用_controlfp()，以强制一部分代码进行单精度浮点计算，方法是在开始处添加
或
以存储控制字的当前值，并对其进行更改以强制尾数以24位存储，方法是在结尾处使用
以恢复原始控制字。
与计算设备（参见附录A）不同，主机平台通常还支持非规格化的数字。这可能导致设备仿真和设备执行模式之前的结果显著不同，因为一些计算可能在一种情况下生成有限结果，而在另一种情况下生成无限结果。
4.5.3 驱动程序API
驱动程序API是基于句柄的命令式API：大多数对象通过不透明句柄来引用，这些句柄可以指定给函数以操纵对象。
CUDA中的可用对象汇总在表4-1中。
表4-1. CUDA驱动程序API中的可用对象
4.5.3.1 初始化
在调用附录E中的任何函数（参见E.1）之前，需要使用cuInit()进行初始化。
4.5.3.2 设备管理
E.2中的函数用于管理系统中现有的设备。
cuDeviceGetCount()和cuDeviceGet()提供了一种方法来枚举这些设备和E.2中的其他函数以检索其属性：
4.5.3.3 上下文管理
E.3中的函数用于创建、附加和分离CUDA上下文。
CUDA上下文类似于CPU进程。在计算API中执行的所有资源和操作都封装在CUDA上下文中，并且当上下文销毁时，系统将自动清除这些资源。除了模块 和纹理参考等对象之外，每个上下文还具有自己不同的32位地址空间。因此，不同CUDA上下文中的CUdeviceptr值引用不同的内存位置。
上下文具有与主机线程一对一的对应关系。在同一时间，主机线程只能有一个设备上下文。当使用cuCtxCreate()创建上下文时，对于调用主机线程，此上下文就成为当前上下文。
如果有效上下文不是线程的当前上下文，则在上下文中操作的CUDA函数（不涉及设备仿真或上下文管理的大多数函数）将返回CUDA_ERROR_INVALID_CONTEXT。
要促进在同一上下文中操作的第三方授权代码之间的互操作性，驱动程序API维护了由给定上下文的每个不同客户机递增的使用计数。例如，如果加载了三个库使 用相同的CUDA上下文，则每个库必须调用cuCtxAttach()递增使用计数，并在库完成使用上下文时，调用cuCtxDetach()递减使用计 数。当使用计数等于0时，则销毁上下文。对于大多数库，预计应用程序将在加载或初始化库之前创建CUDA上下文；这样，应用程序可以使用其自己的试探法创 建上下文，而库只需在传递给它的上下文上操作。
4.5.3.4 模块管理
E.4中的函数用于加载和卸载模块，并检索指向变量中定义的变量或函数的句柄或指针。
模块是可动态加载的设备代码和数据的包，类似于Windows中的DLL，是nvcc的输出（参见4.2.5）。所有符号（包括函数、全局变量和纹理参考）的名称在模块范围内维护，以便独立第三方写入的模块可以在同一CUDA上下文中互操作。
下列代码示例加载模块并检索指向某个内核的句柄：
4.5.3.5 执行控制
E.7中介绍的函数管理设备上内核的执行。cuFuncSetBlockShape()设置给定函数每块的线程数，以及如何分配其线程ID。 cuFuncSetSharedSize()设置函数的共享内存大小。函数的cuParam*()族用于指定下一次调用cuLaunchGrid()或 cuLaunch()启动内核时将提供给内核的参数。
4.5.3.6 内存管理
E.8中的函数用于分配和释放设备内存，并在主机和设备内存之间传送数据。
线性内存使用cuMemAlloc()或cuMemAllocPitch()进行分配，使用cuMemFree()进行释放。
下列代码示例将具有256个浮点元素的数组分配在线性内存中：
建议在分配2D数组时使用cuMemAllocPitch()，因为这样可以确保分配适当填补以满足对齐要求，如5.1.2.1所述，从而确保在访 问行地址时或执行2D数组和其他设备内存之间的复制（使用cuMemcpy2D()）时达到最佳性能。返回的pitch（或跨度）必须用于访问数组元素。 下列代码示例分配了浮点数值的width×height的2D数组，并显示如何在设备代码中循环处理数据元素：
CUDA数组使用cuArrayCreate()进行创建，使用cuArrayDestroy()进行销毁。
下列代码示例分配了一个32位浮点组件的width×height的CUDA数组：
E.5列出用于在使用cuMemAlloc()分配的线性内存、使用cuMemAllocPitch()分配的线性内存和CUDA数组之间复制内存的所有各种函数。下列示例代码将2D数组复制到在前面的代码示例中分配的CUDA数组中：
下列代码示例将一些主机内存数组复制到设备内存中：
4.5.3.7 流管理
E.5中的函数用于创建和销毁流，并确定流的所有操作是否已经完成。
下列代码示例创建两个流：
其中每个流由下列代码示例定义为一个从主机到设备的内存复制、一个内核启动和一个从设备到主机的内存复制的序列：
每个流将其输入数组hostPtr部分复制到设备内存中的数组inputDevPtr中，通过调用cuFunction处理设备上的 inputDevPtr，并将结果outputDevPtr重新复制给hostPtr的相同部分。使用两个流处理hostPtr允许一个流的内存复制可能 与另一个流的内核执行相重叠。hostPtr必须指向页面锁定主机内存以便任何重叠发生：
最后调用cuCtxSynchronize()以确保在进一步处理之前所有流都已完成。
4.5.3.8 事件管理
E.6中的函数用于创建、记录和销毁事件，并查询两个事件之间用去的时间。
下列代码示例创建两个事件：
这些事件可用于以下列方式定时上一节的代码示例：
4.5.3.9 纹理参考管理
E.9中的函数用于管理纹理参考。
在内核可以使用纹理参考读取纹理内存之前，必须使用cuTexRefSetAddress()或cuTexRefSetArray()将纹理参考绑定到纹理。
如果模块cuModule包含某个定义如下的纹理参考texRef：
4.5.3.10 OpenGL互操作性
E.10中的函数用于控制与OpenGL的互操作性。
与OpenGL的互操作性必须使用cuGLInit()进行初始化。
缓冲对象必须注册到CUDA之后才能映射。此操作使用cuGLRegisterBufferObject()来完成：
4.5.3.11 Direct3D互操作性
D.9中的函数用于控制与Direct3D的互操作性。
与Direct3D的互操作性必须使用cuD3D9Begin()进行初始化，使用cuD3D9End()终止：
在这些调用之间，顶点对象必须注册到CUDA之后才能映射。此操作使用cuD3D9RegisterVertexBuffer()完成：
使用cuD3D9UnmapVertexBuffer()解除映射，使用cuD3D9UnregisterVertexBuffer()解除注册。
第5章 性能指南
5.1 指令性能
要处理一个warp的线程的指令，多处理器必须：
读取warp的每个线程的指令操作数，
执行指令，
写入warp的每个线程的结果。
因此，有效的指令吞吐量取决于名义指令吞吐量以及内存延迟和带宽。它通过下列方式最大化：
最小化具有低吞吐量的指令的使用（参见5.1.1），
最大化每种内存的可用内存带宽（参见5.1.2），
允许线程调度器尽可能地将内存事务与数学计算重叠，这需要：
由线程执行的程序具有高的算术密度，也就是说，每个内存操作具有高的算术操作数；
每个多处理器具有许多活动线程，详见5.2。
5.1.1 指令吞吐量
5.1.1.1 算术指令
要执行warp的一个指令，多处理需要：
4个时钟周期，用于浮点加、浮点乘、浮点乘-加、整数加、位操作、比较、求最小、求最大、类型转换指令；
16个时钟周期，用于倒数、倒数平方根、__log(x)（参见表B-2）。
32位整数乘法使用16个时钟周期，但__mul24和__umul24（参见附录B）提供了4个时钟周期的有符号和无符号24位整数乘法。但是，在将来 的架构中，__[u]mul24将比32位整数乘法慢，所以我们建议提供两个内核由应用程序相应地调用，其中一个使用__[u]mul24，另一个使用一 般的32位整数乘法。
整数除法和模数操作特别昂贵，如果可能的话应该尽可能地避免，或者尽量替换为位操作：如果n是2的幂，则(i/n)等于(i>>log2(n))，(i%n)等于(i&(n-1))；如果n是文本，则编译器将执行这些转换。
其他函数使用更多时钟周期，因为它们实现为多个指令的组合。
浮点平方根实现为倒数平方根与求倒，而非倒数平方根与乘法，所以它对于0和无穷大获得正确的结果。因此，它对于warp使用32个时钟周期。
浮点除法使用36个时钟周期，但__fdividef(x, y)提供了更快的版本，即20个时钟周期（参见附录B）。
__sin(x)、__cos(x)、__exp(x)使用32个时钟周期。
有时候，编译器必须插入转换指令，从而引入附加的执行周期。这种情况包括：
操作在其操作数通常需要转换为int的char或short上的函数，
用作单精度浮点计算输入的双精度浮点常量（不使用任何类型后缀定义），
用作表B-1中定义的数学函数的双精度版本的输入参数的单精度浮点变量。
最后两种情况可以通过下列方式避免：
单精度浮点常量，使用f后缀定义，比如3.141592653589793f、1.0f、0.5f，
数学函数的单精度版本，也使用f后缀定义，比如sinf()、logf()、expf()。
对于单精度代码，我们强烈建议使用浮点类型和单精度数学函数。当在不支持原生双精度的设备（比如计算能力1.x的设备）上编译时，双精度类型默认降级为浮点数，双精度数学函数映射为其单精度对应值。但是，如果将来这些设备将支持双精度，则这些函数将映射为双精度实现。
5.1.1.2 控制流指令
任何流控制指令(if, switch, do, for, while)通过导致同一warp的线程分散，也就是说，按照不同的执行路径执行，可以显著影响有效的指令吞吐量。如果这种情况发生，则不同的执行路径必 须序列化，增加此warp执行的指令总数。当所有不同的执行路径都已完成时，线程将集中到同一执行路径。
当控制流取决于线程ID时，要获得最佳性能，就应该写入控制条件，以便最小化分散的warp数。这可能是因为warp在块中的分布是确定性的，参见 3.2。一个小示例是当控制条件仅取决于(threadIdx / WSIZE)，其中WSIZE是warp大小。在这种情况下，没有任何warp会分散，因为控制条件与warp已完美对齐。
有时，编译器可以通过使用分支预测展开循环或可以优化if 或switch语句，详细说明如下。在这些情况下， warp绝不会分散。程序员也可以使用#pragma unroll指令控制循环展开（参见4.2.5.2）。
当使用分支预测时，其执行取决于控制条件的指令没有一个将跳过。相反，每个指令都与基于控制条件设置为真或假的每线程条件代码或谓词相关，虽然其中每个指令都会调度执行，但只有具有真谓词的指令将实际执行。具有假谓词的指令不写入结果，而且不求地址或读取操作数。
仅当由分支条件控制的指令数小于或等于特定临界值时，编译器才将分支指令替换为谓词指令；如果编译器确定条件可能生成许多分散的warp，则此临界值是7，否则是4。
5.1.1.3 内存指令
内存指令包括从共享或全局内存中读取或写入的任何指令。多处理器使用4个时钟周期来执行warp的一个内存指令。此外，当访问全局内存时，还有400到600个时钟周期的内存延迟。
例如，下列示例代码中的赋值操作符
使用4个时钟周期执行从全局内存中的读取，使用4个时钟周期执行到共享内存的写入，但最重要的是使用400到600个时钟周期从全局内存中读取浮点数。
如果在等待全局内存访问完成时，可以执行足够的独立算术指令，则此全局内存延迟的大部分可以由线程调度器隐藏。
5.1.1.4 同步指令
如果没有任何线程必须等待其他任何线程，则__syncthreads将使用4个时钟周期执行warp。
5.1.2 内存带宽
每个内存空间的有效带宽主要取决于内存访问模式，详见下列小节。
因为设备内存与芯片上内存相比具有较高的延迟和较低的带宽，所以设备内存访问必须最小化。典型的编程模式是将来自设备内存的数据存储到共享内存中；换句话说，就是让块的每个线程：
将设备内存中的数据加载到共享内存中，
与块的所有其他线程同步，以便每个线程可以安全读取由不同线程写入的共享内存位置，
处理共享内存中的数据，
如果必要的话，重新同步以确保共享内存已经由结果更新，
将结果写回到设备内存中。
5.1.2.1 全局内存
全局内存空间没有高速缓存，所以最重要的是按照正确的访问模式获得最大的内存带宽，尤其是已知对设备内存的访问有多昂贵时。
首先，设备能够在单个指令中将32位、64位或128位字从全局内存读取到寄存器中。要将如下赋值：
编译到单个加载指令中，type必须使得sizeof(type)等于4、8或16，且类型为type的变量必须对齐为sizeof(type)个字节（也就是说，让其地址是sizeof(type)的倍数）。
对于4.3.1.1一节中介绍的内置类型，比如float2或float4，对齐要求将自动完成。
对于结构体，大小和对齐要求可以由编译器使用对齐指定符__align__(8)或__align__(16)来强制执行，比如
更准确地说，在每个半warp中，半warp中的线程号N应访问地址
其中，HalfWarpBaseAddress具有类型type*，type满足上述的大小和对齐要求。此 外，HalfWarpBaseAddress应对齐为16*sizeof(type)个字节（比如，是16*sizeof(type)的倍数）。驻留在全 局内存中或由D.5或E.8中的内存分配例程之一返回的变量的任何地址BaseAddress始终对其为至少256个字节，所以为了满足内存对齐约 束，HalfWarpBaseAddress-BaseAddress应是16*sizeof(type)的倍数。
注意，如果半warp满足上述所有要求，即使半warp的一些线程不实际访问内存，每线程内存访问也将合并。
与仅分别履行其每个半warp的合并要求相反，我们建议履行整个warp的合并要求，因为将来的设备将使其成为适当合并的必需操作。
图5-1显示了已合并内存访问的一些示例，而图5-2和图5-3显示了未合并内存访问的一些示例。
已合并64位访问提供了比已合并32位访问稍低的带宽，已合并128位访问提供了比已合并32位访问低很多的带宽。然而，当访问是32位时，尽管未合并访问的带宽比已合并访问的带宽低大约一个数量级，但当访问是64位时，仅低大约4倍，当访问是128时，仅低大约2倍。
左：已合并的float内存访问。
右：已合并的float内存访问（分散warp）。
图5-1. 已合并全局内存访问模式的示例
左：非顺序的float内存访问。
右：未对齐的开始地址。
图5-2. 未合并全局内存访问模式的示例
左：不相邻的float内存访问。
右：未合并的float3内存访问。
图5-3. 未合并全局内存访问模式的示例
常见的全局内存访问模式是当线程ID为tid的每个线程访问位于类型为type*的地址BaseAddress上的数组的一个元素时，使用下列地址：
要获得内存合并，type必须满足上述大小和对齐要求。特别地，这意味着，如果type是大于16个字节的结构体，则应分割为满足这些要求的多个结构体，而且数据应在内存中排列为这些结构体的多个数组，而非类型为type*的单个数组。
另一个常见的全局内存访问模式是当索引为(tx,ty)的每个线程访问位于类型为type*、宽度为width的地址BaseAddress上的2D数组的一个元素时，使用下列地址
在这种情况下，仅当满足下列条件，用户才能获得线程块的所有半warp的内存合并：
线程块的带宽是半个warp大小的倍数；
width是16的倍数。
特别地，这意味着，如果宽度不是16的倍数的数组实际使用向上取整为最接近的16的倍数进行分配，且其行相应地进行填补，则此数组将获得非常有效的访问。 cudaMallocPitch()和cuMemAllocPitch()函数及其相关的内存复制函数（参见D.5和E.8）允许程序员编写不依赖于硬件 的代码来分配符合这些约束的数组。
5.1.2.2 常量内存
常量内存空间具有高速缓存，所以从常量内存中的读取仅在高速缓存缺失时，耗费从设备内存中的一个内存读取，否则它仅耗费从常量高速缓存中的一个读取。
对于半warp的所有线程，只要所有线程读取同一地址，则从常量内存中读取与从寄存器中读取一样快。成本随所有线程读取的不同地址数线性扩展。与仅让每个 半warp中的所有线程读取同一地址相反，我们建议让整个warp的所有线程读取同一地址，因为将来的设备将需要此操作来实现完全的快速读取。
5.1.2.3 纹理内存
纹理内存空间具有高速缓存，所以纹理拾取仅在高速缓存缺失时，耗费从设备内存中的一个内存读取，否则它仅耗费从纹理高速缓存中的一个读取。纹理高速缓存针 对2D空间局部性进行了优化，所以读取紧密相邻的纹理地址的同一warp的线程将达到最佳性能。此外，它还设计用于流水化具有恒定延迟的拾取。比如，高速 缓存命中降低了DRAM带宽需求，但没有降低拾取延迟。
通过纹理拾取读取设备内容可能是从全局或常量内存中读取的设备内存的有利备选方案，详见5.4。
5.1.2.4 共享内存
因为位于芯片上，所以共享内存要比本地和全局内存空间快得多。实际上，对于warp的所有线程，访问共享内存与访问寄存器一样快，只要在线程之间没有任何库冲突，详见下文。
要获得高内存带宽，请将共享内存划分为同样大小的内存模块，命名为库，可以同时访问这些库。因此，由属于n个显式内存库的n个地址组成的任何内存读取或写入请求都可以同时获得服务，最后可收益n倍的有效带宽与单个模块的带宽一样高。
但是如果两个内存请求地址组成同一个内存库，则会导致库冲突和访问必须序列化。硬件将带有库冲突的内存请求按需分成许多单独的无冲突的请求，通过一个因子 来减少有效带宽使其与单独内存请求的数目相等。如果单独内存请求的数目为n，则最初的内存请求据说会导致n种方式的库冲突。
要获得最高性能，因此了解如何将内存地址映射到内存库变得非常重要，这样做的目的是调度内存请求，以便最小化库冲突。
在共享内存空间这种情况下，将库组织为：将连续的32位字分配到连续的库中，且每两个时钟周期每个库都有一个32位带宽。
对于计算能力1.x的设备，warp大小为32，库为16（参见5.1）；将warp的共享内存请求划分为第一半warp的一个请求和第二半warp的一个请求。因此，属于第一半warp的线程和属于同一warp的第二半warp的线程间不会发生任何库冲突。
一个常见的情况是每个线程从按线程ID tid索引的数据中使用某个跨度s来访问一个32位字：
这种情况下，只要s*n是库数m的倍数，或者同等地，只要n是m/d的倍数，其中d是m和s最大公约数，则线程tid和tid+n将访问同一库。因 此，仅当半warp大小小于或等于m/d时，才不会发生库冲突。对于计算能力1.x的设备，仅当d等于1时，或者换句话说，因为m是2的倍数，所以仅当s 是奇数时，才不会发生任何库冲突。
图5-4和图5-5显示了一些无冲突内存访问的示例，而图5-6显示了一些导致库冲突的内存访问示例。
值得一提的其他情况是当每个线程访问大小小于或大于32位的元素时。例如，如果按下列方式访问char的数组，则会发生库冲突。
例如，因为shared[0]、shared[1]、shared[2]和shared[3]属于同一个库。但是，如果按下列方式访问同一数组，则不会发生任何库冲突：
将导致下列结果：
如果type定义如下，则结果为三个单独的无库冲突的内存读取
因为每个成员是使用三个32位字的跨度来访问。
如果type定义如下，则结果为两个单独的无库冲突的内存读取
因为每个成员是使用三个32位字的跨度来访问。
如果type定义如下，则结果为两个单独的无库冲突的内存读取
因为每个成员是使用五个字节的跨度来访问。
最后，共享内存还具有广播机制，当服务一个内存读取请求时，可以读取一个32位字并同时广播到多个线程。当半warp的多个线程从同一32位字内的地址读 取时，这将减少库冲突的数目。更精确地说，由多个地址组成的内存读取请求随时间由多个步来服务——每两个时钟周期一步——每步服务一个这些地址的无冲突子 集，直到所有地址都已服务完毕；在每一步，子集从尚未服务的剩余地址中构建，过程如下：
选择由剩余地址指向的其中一个字作为广播字，
将下列内容包括在子集中：
位于广播字内的所有地址，
由剩余地址指向的每个库的一个地址。
选择哪个字作为广播字以及在每个周期为每个库选择哪个地址均未指定。
常见的无冲突情况是当半warp的所有线程从同一32位字内的地址中读取时。
图5-7显示了一些涉及广播机制的内存读取访问的示例。
左：跨度为一个32位字的线性寻址
右：随机排列
图5-4. 无库冲突的共享内存访问模式示例
跨度为三个32位字的线性寻址。
图5-5. 无库冲突的共享内存访问模式示例
左：跨度为两个32位字的线性寻址将导致2路库冲突。
右：跨度为八个32位字的线性寻址将导致8路库冲突。
图5-6. 有库冲突的共享内存访问模式示例
左：因为所有线程从同一32位字中的地址读取，所以此访问模式是无冲突的。
右：如果在第一步期间选择库5中的字作为广播字，则此访问模式不会导致任何库冲突，否则会导致2路库冲突。
图5-7. 有广播的共享内存读取访问模式示例
5.1.2.5 寄存器
通常，访问寄存器对于每条指令需要零个额外时钟周期，但是由于寄存器读后写依赖关系和寄存器内存库冲突，可能会发生延迟。
由读后写依赖关系导致的延迟可以忽略，只要每个多处理器至少有192个活动线程使其隐藏。
编译器和线程调度器调度指令尽可能最佳，以避免寄存器内存库冲突。当每块中的线程数是64的倍数时，可以获得最佳结果。除了遵循此规则之外，应用程序对这些库冲突没有任何直接的控制。特别地，无需将数据打包为float4或int4类型。
5.2 每块的线程数
给定每网格的线程总数，选择每块的线程数或等同的块数时应该最大化可用计算资源的利用率。这意味着块的数目应该至少与设备中的多处理器的数目一样多。
此外，当每个多处理器仅运行一块时，如果每块没有足够的线程来覆盖加载延迟，则在线程同步期间，以及设备内存读取期间，将强制每个多处理器进入空闲状态。 因此，最好的方法是每个多处理器上允许存在两个或多个活动块，以允许在等待的块和可以运行的块之间出现重叠。要让这种情况发生，不仅块的数目至少应该是设 备中多处理器数目的两倍，而且每块分配的共享内存量至多应该是每个多处理器可用共享内存总量的一半（参见3.2）。更多线程块以管线方式在设备中分流，并 在更大程度上分摊开销。
有了足够大数目的块，每块线程的数目应选择为warp大小的倍数，以避免使用未充满的warp而浪费计算资源，或者更好地，选择为64的倍数，究其原因， 参见5.1.2.5。为每块分配更多线程有利于有效的时间分片，但是每块的线程越多，每线程可用的寄存器就越少。如果内核编译的线程数大于执行配置所允许 的数目，这可能会阻止内核调用继续。当使用--ptxas-options=-v选项编译时，内核编译的寄存器数目（还有本地、共享和常量内存使用）由编 译器报告。
对于计算能力1.x的设备，每线程可用的寄存器数等于：
其中，R是每个多处理器的寄存器总数(参见附录A），B是每个多处理器的活动块数，T是每块的线程数，ceil(T, 32)是T向上取整为32的最近倍数。
每块64个线程是最小的，并且仅当每个多处理器有多个活动块时才有意义。每块有192或256个线程比较好，而且通常允许有足够的寄存器进行编译。
如果用户想将其扩展到将来的设备，则每个网格的块数应该至少是100；1000个块将扩展到几代。
每个多处理器的活动warp数与活动warp（参见附录A）的最大数目的比率称作多处理器占有率。为了最大化占有率，编译器应尝试最小化寄存器使用，而且 程序员需要小心选择执行配置。CUDA软件开发工具包提供了一个电子表格以帮助程序员基于共享内存和寄存器要求来选择线程块大小。
5.3 主机和设备之间的数据传送
设备和设备内存之间的带宽比设备内存和主机内存之间的带宽高得多。因此，用户应该争取最小化主机和和设备之间的数据传送，例如，将更多代码从主机移动到设 备，即使这意味着要使用低并行计算来运行内核。中间数据结构可以在设备内存中创建，由设备操作，销毁，而且永远不会由主机映射，或复制到主机内存。
另外，由于每次传送都会有开销，所以将许多小的传送分批为一次大的传送要比单独执行每一个传送要好得多。
最后，使用页面锁定内存时，可以在主机和设备之间获得较高带宽，详见4.5.1.2。
5.4 纹理拾取与全局或常量内存读取
与从全局或常量内存中读取相比，通过纹理拾取进行设备内存读取具有下列几个优点：
高速缓存，如果在纹理拾取中有位置，则可以潜在地展示较高带宽；
不受内存访问模式的约束，此约束即全局或常量内存读取必须尊重以获得好的性能（参见5.1.2.1和5.1.2.2）；
寻址计算的延迟隐藏得更好，可能会改善执行随机访问数据的应用程序的性能；
打包的数据可以在单个操作中广播到多个独立变量中。
8位和16位整数输入数据可以有选择地转化为[0.0,1.0]或[-1.0,1.0]范围内的32位浮点值（参见4.3.4.1）。
如果纹理是CUDA数组（参见4.3.4.2），则硬件提供了可能对不同应用程序有用的其他能力，尤其是图像处理：
然而，在同一内核调用中，纹理高速缓存与全局内存写不保持一致，从而对已经在同一内核中通过全局写而写入的某个地址的纹理拾取将返回未定义的数据。 换句话说，仅当此内存位置已经由先前的内核调用或内存复制更新时，而不是已经由同一内核调用中的同一个或另一个线程更新时，线程才可以通过纹理安全地读取 某个内存位置。仅当作为内核从线性内存中的拾取无论如何不能写入到CUDA数组中时，这才相关。
5.5 整体性能优化策略
性能优化围绕三个基本策略：
最大化并行执行；
优化内存使用以获得最大内存带宽；
优化指令使用以获得最大指令吞吐量。
最大化并行执行首先应通过暴露尽可能多的数据并行来结构化算法。在算法中，因为一些线程需要同步以便互相共享数据，而破坏了并行性的情况有两种：这些线程 属于同一块，这种情况下，线程应该使用__syncthreads()，并通过同一内核调用中的共享内存来共享数据，或者这些线程属于不同块，这种情况 下，必须使用两个单独的内核调用通过全局内存来共享数据，一个内核调用写入全局内存，另一个从全局内存读取。
暴露算法的并行之后，则需要尽可能有效地将其映射到硬件。通过仔细选择每个内核调用的执行配置来完成此操作，详见5.2。
应用程序还可以通过显式暴露通过流在设备上的并发执行，如4.5.1.5所述，在更高的水平上最大化并行执行，以及最大化主机和设备之间的并发执行。
最优化内存使用首先应最小化具有低带宽的数据传送。这意味着最小化主机和设备之间的数据传送，详见5.3，因为这要比在设备和全局内存之间的数据传送的带 宽低得多。这也意味着通过最大化设备上共享内存的使用来最小化设备和全局内存之间的数据传送，详见5.1.2。有时候，最好的优化甚至可能是通过简单地重 新计算数据来避免任何数据传送，而不管是否需要这样做。
另一方面，通常仅当共享内存访问具有高度的库冲突时才值得优化。
对于优化指令使用，应该最小化具有低吞吐量的算术指令的使用。这包括在不影响最终结果时用精度换速度，比如使用固有函数，而不使用常规函数（固有函数在表 B-2中列出），或使用单精度而不使用双精度。由于设备的SIMD本质，所以要特别注意控制流指令，详见5.1.1.2。
第6章 矩阵乘法的示例
6.1 概述
计算两个维度分别为(wA, hA)和 (wB, wA) 的矩阵A和B的乘积C的任务以下列方式分为多个线程：
每个线程负责计算C 的一个平方子矩阵Csub ；
块内的每个线程负责计算Csub的一个元素。
选择Csub的维度block_size等于16，以便每块的线程数是warp大小的倍数（参见5.2），并且保持低于每块的最大线程数（参见附录A）。
如图6-1所示，Csub 等于两个矩形矩阵的乘积：维度为(wA, block_size)的子矩阵A，与Csub具有相同的行索引，维度为(block_size, wA)的B的子矩阵，与Csub具有相同的列索引。为了适应设备的资源，这两个矩形矩阵可根据需要划分为许多维度为block_size的平方矩阵，并且 Csub计算为这些平方矩阵的乘积之和。其中每个乘积的执行过程是：首先使用一个加载每个矩阵的一个元素的线程，将两个相应的平方矩阵从全局内存加载到共 享内存，然后让每个线程计算乘积的一个元素。每一线程将其中每个乘积的结果累计到寄存器中，执行完毕后，将结果写入全局内存。
通过以这种方式分块计算，我们可以有效利用快速的共享内存，并节省许多全局内存带宽，因为A和B仅从全局内存读取(wA / block_size)次。
尽管如此，编写此示例是为了清楚地说明各种CUDA编程原则，并非是为了为一般的矩阵乘法提供高性能的内核，所以不应如此构造。
每一线程块计算C的一个子矩阵Csub。块内的每一线程计算Csub的一个元素。
图6-1. 矩阵乘法
6.2 源码清单
6.3 源码攻略
源码包含下列两个函数：
Mul()，作为Muld()的包装器的主机函数。
Muld()，在设备上执行矩阵乘法的内核。
6.3.1 Mul()
Mul()接受下列输入：
指向A和B的元素的主机内存的两个指针，
A的高度和宽度，B的宽度，
指向应该写入C的主机内存的指针。
Mul()执行下列操作：
使用cudaMalloc()将足够的全局内存分配到库A、B和C中；
使用cudaMemcpy()将A和B从主机内存复制到全局内存；
调用Muld()在设备上计算C；
使用cudaMemcpy()将C从全局内存复制到主机内存；
使用cudaFree()释放为A、B和C分配的全局内存。
6.3.2 Muld()
除了指针指向设备内存而非主机内存之外，Muld()与Mul()具有相同的输入。
对于每个块，Muld()迭代处理所有需要计算Csub的A和B的子矩阵。在每次迭代中，此函数：
将A的一个子集和B的一个子集从全局内存加载到共享内存中；
同步以确保两个子矩阵都由块内的所有线程完全加载；
计算两个子集的乘积并将其加到上一次迭代期间获得的乘积中；
再次同步以确保在开始下一次迭代之前两个子集的乘积已经完成。
按照5.1.2.1和5.1.2.4所述，编写Muld()是为了最大化内存性能。
的确，假设wA和wB是16的倍数（如5.1.2.1所建议的），则确保了全局内存合并，因为a、b和c都是BLOCK_SIZE的倍数，BLOCK_SIZE等于6。
对于每个半warp，也没有任何共享内存库冲突，所有线程的ty和k都是相同的，tx在0到15之间变化，因此对于内存访问As[ty][tx]、 Bs[ty][tx]和Bs[k][tx]，每个线程都访问一个不同的库，对于内存访问As[ty][k]，每个线程都访问同一个库。
