CUDA编程接口的目标是为熟悉C编程语言的用户提供相对简单的路径，以便容易地编写在设备上执行的程序。
它包括：
主机组件，如4.5所述，在主机上运行，提供函数以控制并访问主机中的一个或多个计算设备；
设备组件，如4.4所述，在设备上运行，并提供特定于设备的函数；
通用组件，如4.3所述，提供内置的向量类型，以及主机和设备代码中都支持的C标准库子集。
必须强调一下，只有支持在设备上运行的C标准库中的函数才是公共运行时runtime组件提供的函数。
C编程语言的扩展有四个部分：
函数类型限定符，用于指定函数是在主机上还是在设备上执行，以及可以从主机中还是设备中调用（参见4.2.1）；
变量类型限定符，用于指定变量在设备上的内存位置（参见4.2.2）；
新指令，用于指定如何从主机中的设备上执行内核（参见4.2.3）；
四个内置变量，用于指定网格和块维度，以及块和线程索引（参见4.2.4）。
包含这些扩展的每个源文件必须使用CUDA编译器nvcc编译，4.2.5中有简单介绍。nvcc的详细介绍可以参见单独的文档。
其中每个扩展附带下文各节中描述的一些限制。违反这些限制时，nvcc将给出错误或警告，但其中一些违规无法发现。
__device__限定符声明函数：
在设备上执行，
只能从设备中调用。
__global__限定符将函数声明为内核。这种函数：
在设备上执行，
只能从主机中调用。
__host__限定符声明函数：
在主机上执行，
只能从主机中调用。
它等同于仅使用__host__声明函数，或不使用__host__、__device__或__global__限定符任意一个声明函数；不管是哪一种情况，函数仅为主机编译。
但是，__host__限定符还可以与__device__限定符结合使用，此时，函数同时为主机和设备编译。
__device__和__global__函数不支持迭代。
__device__和__global__函数不能在函数体内声明静态变量。
__device__和__global__函数不能具有可变个参数。
__device__函数不能取其地址；相反，__global__函数的函数指针则受支持。
__global__和__host__限定符不能一起使用。
__global__函数必须具有void返回类型。
对__global__函数的任何调用必须指定其执行配置，如4.2.3所述。
对__global__函数的调用是异步的，这意味着在设备完成其执行之前返回。
__global__函数参数当前通过共享内存传递给设备并限制为256字节。
__device__限定符声明驻留在设备上的变量。
下面三节中定义的其他类型限定符中至多一个可以与__device__一起使用，以进一步指定变量属于哪个内存空间。如果其中任何一个都不出现，则变量：
驻留在全局内存空间中，
具有应用程序的生命期，
可通过runtime库从网格中的所有线程中和从主机中访问。
__constant__限定符，可以与__device__一起使用，声明变量：
驻留在常量内存空间中，
具有应用程序的生命期，
可通过runtime库从网格中的所有线程中和从主机中访问。
__shared__限定符，可以与__device__一起使用，声明变量：
驻留在线程块的共享内存空间中，
具有块的生命期，
仅可从块内的所有线程中访问。
线程中共享变量的完全顺序一致性，但是在线程中放松的排序。仅在__syncthreads()（参见4.4.2）执行之后，来自其他线程的写入才能保证可见。除非变量声明为挥发，否则只要满足上一语句，编译器就可以优化对共享内存的读写。
将共享内存中的变量声明为外部数据，比如
数组大小在启动时确定（参见4.2.3）。以此方式声明的所有变量在内存中从同一地址开始，从而数组中变量的布局必须通过偏移量明确管理。例如，如果用户想要
位于动态分配的共享内存中，则用户可以使用下列方式声明和初始化数组：
这些限定符不允许用于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__变量的地址只能用在主机代码中。
对__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大于设备上可用的最大共享内存量减去静态分配、函数参数和执行配置所需的共享内存量，则函数调用将失败。
此变量的类型为dim3（参见4.3.1.2），包含网格的维度。
此变量的类型为uint3（参见4.3.1.1），包含网格中的块索引。
此变量的类型为dim3（参见4.3.1.2），包含块的维度。
此变量的类型为uint3（参见4.3.1.1），包含块中的线程索引。
不允许提取任何内置变量的地址。
不允许为任何内置变量赋值。
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引入了两个编译器指令，见下文所述。
默认情况下，__device__函数始终为内联。但是，__noinline__函数限定符可用作编译器的提示以尽量不内联函数。函数体必须仍位于调用此函数的同一文件中。
对于带有指针参数的函数和具有大型参数列表的函数，编译器将不考虑__noinline__限定符。
默认情况下，编译器展开带有已知循环计数的小循环。但是，#pragma unroll指令可用于控制任何给定循环的展开。它必须立即放置在循环前，并只应用于此循环。它后面可以跟一个数字，用于指定循环必须展开多少次。
例如，在下列代码示例中：
循环将展开5次。程序员应该确保展开将不影响程序的正确性（在上例中，如果n小于5，则可能影响）。
#pragma unroll 1将禁止编译器展开循环。
如果在#pragma unroll后面不指定任何数字，如果其循环计数是常数，则循环将完全展开，否则根本不展开。
公共runtime运行时组件可以由主机和设备函数使用。
这些向量类型是从基本整数和浮点数类型派生而来的。它们是结构体，第1个、第2个、第3个和第4个组件分别可以通过字段x、y、z和w来访问。它们都具有
make_<type name>形式的构造函数；例如，
这将创建类型为int2、值为(x, y)的向量。
此类型是基于uint3的整数向量类型，用于指定维度。定义dim3类型的变量时，保留为未指定的任何组件将初始化为1。
表B-1包含当前支持的C/C++标准库数学函数的完整列表，以及在设备上执行时各自的误差界。
在主机代码中执行时，给定函数使用可用的Cruntime执行。
返回在每个时钟周期递增的计数器的值。
在内核开始和结束时对此计数器取样，求出两次取样之差，并记录每个线程的结果，从而计量设备完全执行每个线程所用的时钟周期数，但这并非设备实际执行线程指令所用的时钟周期数。前者大于后者，因为线程是分时执行的。
CUDA支持GPU本用于图形的纹理硬件以访问纹理内存。从纹理内存而非全局内存读取数据可以具有5.4一节描述的几个性能优点。
纹理内存使用名为纹理拾取(texture fetches)的设备函数来从内核中读取，如4.4.5所述。纹理拾取的第一个参数指定一个名为纹理参考(texture reference)的对象。
纹理参考定义要拾取哪部分纹理内存。它必须通过主机runtime函数（参见4.5.2.6和4.5.3.9）绑定到一些内存区域（称为纹理(texture)），然后才能供内核使用。几个不同的纹理参考可以绑定到同一纹理或在内存中重叠的纹理。
纹理参考具有多个属性。其中之一是其维度，用于指定纹理是使用一个纹理坐标作为一维度组进行寻址，还是使用两个纹理坐标作为二维度组进行寻址。数组的元素称为纹理元素（texel，是“texture element”的简写）。
其他属性定义纹理拾取的输入和输出数据类型，如何解释输入坐标，以及应执行什么处理。
纹理参考的一些属性不可变，而且必须在编译时已知；它们在声明纹理参考时指定。纹理参考在文件范围内声明为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。
纹理参考的其他属性是可变的，可以在运行时通过主机运行时（运行时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给出有关纹理拾取的更多详细信息。
纹理可以是线性内存或CUDA数组的任何区域（参见4.5.1.2）。
在线性内存中分配的的纹理：
维度只能等于等于1；
不支持纹理筛选；
只能使用非规格化整数纹理坐标来寻址；
不支持不同的寻址模式：超出范围的纹理访问返回零。
硬件在纹理基地址上强制执行对齐要求。为了让程序员忽略此对齐要求，将纹理参考绑定到设备内存的函数传回一个必须应用到纹理拾取的字节偏移，以便从所需内存中读取。由CUDA的分配例程返回的基指针符合此对齐约束，因此通过将已分配的指针传递给cudaBindTexture()/cuTexRefSetAddress()，应用程序可以完全避免偏移。
设备runtime组件只能在设备函数中使用。
对于表B-1中的一些函数，设备runtime运行时组件中存在一些不太准确但运行较快的版本；这些函数具有相同的名称，但加了前辍__（例如__sin(x)）。表B-2中列出这些固有函数及其各自的误差界。
编译器具有一个选项（-use_fast_math）来强制每个函数编译为其不太精确的对应物（如果存在的话）。
同步块中所有的线程。当所有线程都达到此点时，执行正常继续。
__syncthreads()用于协调同一块的线程之间的通信。当块中的一些线程访问共享或全局内存中的同一地址时，对于其中的一些内存访问，存在潜在的读后写、写后读或写后写的危害。这些数据危险可以通过同步这些访问之间的线程来避免。
__syncthreads()允许出现在条件代码中，但仅当条件在整个线程块中求值相同时才允许，否则代码执行可能暂挂或产生非预期的副作用。
下列函数中的后缀指明IEEE-754取整模式：
rn取整为最近的偶数，
rz向零取整，
ru向上取整（到正无穷大），
rd向下取整（到负无穷大）。
使用指定的取整模式将浮点参数转换为整数。
使用指定的取整模式将浮点参数转换为无符号整数。
使用指定的取整模式将整数参数转换为浮点数。
使用指定的取整模式将无符号整数参数转换为浮点数。
对整数参数执行浮点类型转换，保留值不变。例如，__int_as_float(0xC0000000)等于-2。
对浮点参数执行整数类型转换，保留值不变。例如__float_as_int(1.0f)等于0x3f800000。
从设备内存取纹理时，使用函数的tex1Dfetch()族访问纹理；例如：
这些函数使用纹理坐标x拾取绑定到纹理参考texRef的线性内存区域。不支持任何纹理筛选和寻址模式。对于整数类型，这些函数可以有选择地将整数提升为32位浮点。
除上述函数之外，还支持2元组和4元组；例如：
使用纹理坐标x拾取绑定到纹理参考texRef的线性内存。
从CUDA数组取纹理时，使用tex1D()或tex2D()访问纹理：
这些函数使用纹理坐标x和y拾取绑定到纹理参考texRef的CUDA数组。纹理参考的不变属性（编译时）和可变属性（运行时）的组合确定如何解释坐标、在纹理拾取期间执行何种处理、以及纹理拾取传递的返回值（参见4.3.4.1和4.3.4.2）。
原子函数仅可用于计算能力1.1的设备。附录C列出了这些函数。
原子函数在驻留于全局内存中的一个32位字上执行读-改-写原子操作。例如，atomicAdd()在全局内存中的同一地址上读取一个32位字，为其加上一个整数，然后将结果写回同一地址。在保证执行时不受其他线程干扰这种意义上，此操作是原子的。换句话说，只有此操作完成之后，其他线程才可以访问此地址。
原子操作仅适用于32位有符号的和无符号的整数。
主机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。
两种API都提供函数来列举系统上可用的设备、查询其属性并为内核执行选择其中之一（runtime API参见4.5.2.2，驱动程序API参见4.5.3.2）。
多个主机线程可以在同一设备上执行设备代码，但在设计上，一个主机线程只能在一个设备上执行设备代码。因此，在多个设备上执行设备代码需要多个主机线程。另外，通过一个主机线程中的runtime创建的任何CUDA资源不能由其他主机线程中的rutime使用。
设备内存可以分配为线性内存或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）。页面锁定内存的一个优点是如果主机内存分配为页面锁定，则主机内存和设备内存之间的带宽较高——仅用于由分配主机内存的主机线程执行的数据传送。但是，页面锁定内存是稀有资源，所以早在可分页内存中的分配之前，页面锁定内存中的分配就将开始失败。此外，通过减少可用于操作系统分页的物理内存量，分配太多的页面锁定内存将降低整体系统性能。
OpenGL缓冲对象可以映射到CUDA的地址空间中，从而允许CUDA读取由OpenGL写入的数据，或允许CUDA写入供OpenGL消耗的数据。4.5.2.7一节描述如果使用runtime API完成此操作，4.5.3.10一节描述如何使用驱动程序API完成此操作。
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创建在不同的设备上。
为了方便主机和设备之间的并发执行，一些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应用程序的异步执行。此功能只提供用于调试目的，而且绝不能用作让生产软件可靠运行的方法。
runtime API没有任何显式初始化函数；第一次调用runtime 函数时，runtime API初始化。当定时runtime 函数调用时，以及将第一次调用的错误代码解释到runtime 中时，用户一定要记住这一点。
章节D.1中的函数用于管理呈现在系统中的设备。
cudaGetDeviceCount()和cudaGetDeviceProperties()提供一种用来列举这些设备并检索其属性的方法。
cudaSetDevice()用来选择与主机线程相关的设备：
在调用任何__global__函数或任何附录D中的函数之前必须选择设备。如果显式调用cudaSetDevice()没有执行，将自动选择设备0，且任何随后的显式调用cudaSetDevice()将不起作用。
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数组：
下列代码示例将一些主机内存数组复制到设备内存中：
下列代码示例将一些主机内存数组复制到设备内存中：
下列代码示例将一些主机内存数组复制到常量内存中：
D.3一节中的函数用于创建和销毁流，并确定流的所有操作是否已经完成。
下列代码示例创建两个流：
其中每个流通过下列代码示例定义为一个从主机到设备的内存复制、一个内核启动和一个从设备到主机的内存复制的序列：
每个流将其输入数组hostPtr部分复制到设备内存中的数组inputDevPtr中，通过调用myKernel()处理设备上的inputDevPtr，并将结果outputDevPtr重新复制到同一hostPtr部分。使用两个流处理处理hostPtr允许一个流的内存复制与其他流的内核执行相重叠。hostPtr必须指向要发生的任何重叠的页面锁定主机内存：
最后调用cudaThreadSynchronize()以确保在进一步处理之前所有流都已完成。
D.4一节中的函数用于创建、记录和销毁事件，并查询两个事件之间用去的时间。
下列代码示例创建两个事件：
这些事件可以使用下列方法用于定时上一节的代码示例：
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()用于解除对纹理参考的绑定。
D.8一节中的函数用于控制与OpenGL的互操作性。缓冲对象必须注册到CUDA之后才能映射。此操作使用cudaGLRegisterBufferObject()来完成：
注册之后，内核可以使用由cudaGLMapBufferObject()返回的设备内存地址读取或写入缓冲对象：
使用cudaGLUnmapBufferObject()解除映射，使用cudaGLUnregisterBufferObject()解除注册。
D.9一节中的函数用于控制与Direct3D的互操作性。
与Direct3D的互操作性必须使用cudaD3D9Begin()初始化，使用cudaD3D9End()终止。
在这些调用之间，顶点对象必须注册到CUDA之后才能映射。此操作使用cudaD3D9RegisterVertexBuffer()来完成：
注册之后，内核可以使用由cudaD3D9MapVertexBuffer()返回的设备内存地址读取或写入顶点缓冲：
使用cudaD3D9UnmapVertexBuffer()解除映射，使用cudaD3D9UnregisterVertexBuffer()解除注册。
编程环境不包括对设备上运行的代码的任何原生调试支持，但提供了用于调试的设备仿真模式。在此模式下编译应用程序（使用-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）不同，主机平台通常还支持非规格化的数字。这可能导致设备仿真和设备执行模式之前的结果显著不同，因为一些计算可能在一种情况下生成有限结果，而在另一种情况下生成无限结果。
驱动程序API是基于句柄的命令式API：大多数对象通过不透明句柄来引用，这些句柄可以指定给函数以操纵对象。
CUDA中的可用对象汇总在表4-1中。
.BMP)
在调用附录E中的任何函数（参见E.1）之前，需要使用cuInit()进行初始化。
E.2中的函数用于管理系统中现有的设备。
cuDeviceGetCount()和cuDeviceGet()提供了一种方法来枚举这些设备和E.2中的其他函数以检索其属性：
E.3中的函数用于创建、附加和分离CUDA上下文。
CUDA上下文类似于CPU进程。在计算API中执行的所有资源和操作都封装在CUDA上下文中，并且当上下文销毁时，系统将自动清除这些资源。除了模块和纹理参考等对象之外，每个上下文还具有自己不同的32位地址空间。因此，不同CUDA上下文中的CUdeviceptr值引用不同的内存位置。
上下文具有与主机线程一对一的对应关系。在同一时间，主机线程只能有一个设备上下文。当使用cuCtxCreate()创建上下文时，对于调用主机线程，此上下文就成为当前上下文。
如果有效上下文不是线程的当前上下文，则在上下文中操作的CUDA函数（不涉及设备仿真或上下文管理的大多数函数）将返回CUDA_ERROR_INVALID_CONTEXT。
要促进在同一上下文中操作的第三方授权代码之间的互操作性，驱动程序API维护了由给定上下文的每个不同客户机递增的使用计数。例如，如果加载了三个库使用相同的CUDA上下文，则每个库必须调用cuCtxAttach()递增使用计数，并在库完成使用上下文时，调用cuCtxDetach()递减使用计数。当使用计数等于0时，则销毁上下文。对于大多数库，预计应用程序将在加载或初始化库之前创建CUDA上下文；这样，应用程序可以使用其自己的试探法创建上下文，而库只需在传递给它的上下文上操作。
E.4中的函数用于加载和卸载模块，并检索指向变量中定义的变量或函数的句柄或指针。
模块是可动态加载的设备代码和数据的包，类似于Windows中的DLL，是nvcc的输出（参见4.2.5）。所有符号（包括函数、全局变量和纹理参考）的名称在模块范围内维护，以便独立第三方写入的模块可以在同一CUDA上下文中互操作。
下列代码示例加载模块并检索指向某个内核的句柄：
E.7中介绍的函数管理设备上内核的执行。cuFuncSetBlockShape()设置给定函数每块的线程数，以及如何分配其线程ID。cuFuncSetSharedSize()设置函数的共享内存大小。函数的cuParam*()族用于指定下一次调用cuLaunchGrid()或cuLaunch()启动内核时将提供给内核的参数。
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数组中：
下列代码示例将一些主机内存数组复制到设备内存中：
E.5中的函数用于创建和销毁流，并确定流的所有操作是否已经完成。
下列代码示例创建两个流：
其中每个流由下列代码示例定义为一个从主机到设备的内存复制、一个内核启动和一个从设备到主机的内存复制的序列：
每个流将其输入数组hostPtr部分复制到设备内存中的数组inputDevPtr中，通过调用cuFunction处理设备上的inputDevPtr，并将结果outputDevPtr重新复制给hostPtr的相同部分。使用两个流处理hostPtr允许一个流的内存复制可能与另一个流的内核执行相重叠。hostPtr必须指向页面锁定主机内存以便任何重叠发生：
最后调用cuCtxSynchronize()以确保在进一步处理之前所有流都已完成。
E.6中的函数用于创建、记录和销毁事件，并查询两个事件之间用去的时间。
下列代码示例创建两个事件：
这些事件可用于以下列方式定时上一节的代码示例：
E.9中的函数用于管理纹理参考。
在内核可以使用纹理参考读取纹理内存之前，必须使用cuTexRefSetAddress()或cuTexRefSetArray()将纹理参考绑定到纹理。
如果模块cuModule包含某个定义如下的纹理参考texRef：
E.10中的函数用于控制与OpenGL的互操作性。
与OpenGL的互操作性必须使用cuGLInit()进行初始化。
缓冲对象必须注册到CUDA之后才能映射。此操作使用cuGLRegisterBufferObject()来完成：
D.9中的函数用于控制与Direct3D的互操作性。
与Direct3D的互操作性必须使用cuD3D9Begin()进行初始化，使用cuD3D9End()终止：
在这些调用之间，顶点对象必须注册到CUDA之后才能映射。此操作使用cuD3D9RegisterVertexBuffer()完成：
使用cuD3D9UnmapVertexBuffer()解除映射，使用cuD3D9UnregisterVertexBuffer()解除注册。
|
|
|
|
|
|
|
|
|
