CUDA 3.0 编程接口

最新推荐文章于 2024-07-25 14:57:48 发布

lycnjupt

最新推荐文章于 2024-07-25 14:57:48 发布

阅读量1k

点赞数 1

CUDA 3.0 编程接口

2010年03月23日17:55 IT168网站原创作者：IT168 风辰译编辑：覃里我要评论(0)

标签： CUDA , CUDA开发 , CUDA技术

　　【IT168 文档】目前可用两种接口写CUDA程序：CUDA C和CUDA驱动API。一个应用典型的只能使用其中一种，但是遵守3.4节描述的限制时，可以同时使用两种。

　　CUDA C将CUDA编程模型作为C的最小扩展集展示出来。任何包含某些扩展的源文件必须使用nvcc 编译，nvcc的概要在3.1节。这些扩展允许程序员像定义C函数一样定义内核和在每次内核调用时，使用新的语法指定网格和块的尺寸。

　　CUDA驱动API是一个低层次的C接口，它提供了从汇编代码或CUDA二进制模块中装载内核，检查内核参数，和发射内核的函数。二进制和汇编代码通常可以通过编译使用C写的内核得到。

　　CUDA C包含运行时API，运行时API和驱动API都提供了分配和释放设备存储器、在主机和内存间传输数据、管理多设备的系统的函数等等。

　　运行时API是基于驱动API的，初始化、上下文和模块管理都是隐式的，而且代码更简明。CUDA C也支持设备模拟，这有利于调试(参见节3.2.8)。

　　相反，CUDA驱动API要求写更多的代码，难于编程和调试，但是易于控制且是语言无关的，因为它处理的是二进制或汇编代码。

　　3.2节接着第二章介绍CUDA C。也引入了CUDA C和驱动API共有的概念：线性存储器、CUDA数组、共享存储器、纹理存储器、分页锁定主机存储器、设备模拟、异步执行和与图形学API互操作。3.3节会介绍有关这些概念的知识和描述它们在驱动API中是怎样表示的。

　　3.1 用nvcc编译

　　内核可以使用PTX编写，PTX就是CUDA指令集架构，PTX参考手册中描述了PTX。通常PTX效率高于像C一样的高级语言。无论是使用PTX还是高级语言，内核都必须使用nvcc编译成二进制代码才能在设备在执行。

　　nvcc是一个编译器驱动，简化了C或PTX的编译流程：它提供了简单熟悉的命令行选项，同时通过调用一系列实现了不同编译步骤的工具集来执行它们。本节简介了nvcc的编译流程和命令选项。完整的描述可在nvcc用户手册中找到。

　　3.1.1 编译流程

　　nvcc可编译同时包含主机代码(有主机上执行的代码)和设备代码(在设备上执行的代码)的源文件。nvcc的基本流程包括分离主机和设备代码并将设备代码编译成汇编形式(PTX)或/和二进制形式(cubin对象)。生成的主机代码要么被输出为C代码供其它工具编译，要么在编译的最后阶段被nvcc调用主机编译器输出为目标代码。

　　应用能够：

　　1.要么在设备上使用CUDA驱动API装载和执行PTX源码或cubin对象(参见3.3节)同时忽略生成的主机代码(如果有);

　　2.要么链接到生成的主机代码;生成的主机代码将PTX代码和/或cubin对象作为已初始化的全局数据数组导入，还将2.1节引入的<<<…>>>语法转化为必要的函数调用以加载和发射每个已编译的内核。

　　应用在运行时装载的任何PTX代码被设备驱动进一步编译成二进制代码。这称为即时编译。即时编译增加了应用装载时间，但是可以享受编译器的最新改进带来的好处。也是当前应用能够在未来的设备上运行的唯一方式，细节参见3.1.4节。

　　3.1.2 二进制兼容性

　　二进制代码是由架构确定的。生成cubin对象时，使用编译器选项-code指定目标架构：例如，用-code=sm_13编译时，为计算能力1.3的设备生成二进制代码。二进制兼容性保证向后兼容，但不保证向前兼容，也不保证跨越主修订号兼容。换句话说，为计算能力为X.y生成的cubin对象只能保证在计算能力为X.z的设备上执行，这里，z>=y。

　　3.1.3 PTX兼容性

　　一些PTX指令只被高计算能力的设备支持。例如，全局存储器上的原子指令只在计算能力1.1及以上的设备上支持;双精度指令只在1.3及以上的设备上支持。将C编译成PTX代码时，-arch编译器选项指定假定的计算能力。因此包含双精度计算的代码，必须使用“-arch=sm_13”(或更高计算能力)编译，否则双精度计算将被降级为单精度计算。

　　为某些特殊计算能力生成的PTX代码始终能够被编译成相等或更高计算能力设备上的二进制代码。

　　3.1.4 应用兼容性

　　为了在特定计算能力的设备上执行代码，应用加载的二进制或PTX代码必须满足如3.1.2节和3.1.3节说明的计算能力兼容性。特别地，为了能在将来更高计算能力(不能产生二进制代码)的架构上执行，应用必须装载PTX代码并为那些设备即时编译。

　　CUDA C应用中嵌入的PTX和二进制代码，由-arch和-code编译器选项或-gencode编译器选项控制，详见nvcc用户手册。例如

nvcc x.cu
–gencode arch = compute_10,code = sm_10
–gencode arch = compute_11,code = \’compute_11,sm_11\’

　　嵌入与计算能力1.0兼容的二进制代码(第一个-gencode选项)和PTX和与计算能力1.1兼容的二进制代码(第二个-gencode选项)。

　　生成的主机代码在运行时自动选择最合适的代码装载并执行，对于上面例子，将会是：

　　1.0二进制代码为计算能力1.0设备，

　　1.1二进制代码为计算能力1.1,1.2,1.3的设备，

　　通过为计算能力2.0或更高的设备编译1.1PTX代码获得的二进制代码。

　　例如，x.cu可有一个使用原子指令的优化代码途径，只能支持计算能力1.1或更高的设备。__CUDA_ARCH__宏可以基于计算能力用于不同的代码途径。它只为设备代码定义。例如，当使用“arch=compte_11”编译时，__CUDA_ARCH__等于110。

　　使用驱动API的应用必须将代码编译成分立的文件，且在运行时显式装载和执行最合适的文件。

　　nvcc用户手册为-arch,-code和-gencode编译器选项列出了多种简写。如“arch=sm_13”是“arch=compute_13 ?code=compute_13,sm_13”的简写(等价于“-gencode arch=compute_13,code=\’compute_13,sm_13\’”)。

　　3.2 CUDA C

　　CUDA C为熟悉C语言的用户提供了一个简单途径，让他们能够轻易的写出能够在设备上执行的程序。

　　CUDA C包含了一个C语言的最小扩展集和一个运行时库。语言核心扩展在第二章已经介绍了。本节继续介绍运行时。所有扩展的完整的描述可在附录B找到，CUDA运行时的完整描述可在CUDA参考手册中找到。

　　cudart动态库是运行时的实现，它所有的入口点前缀都是cuda。

　　运行时没有显式的初始化函数;在初次调用运行时函数(更精确地，不在参考手册中设备和版本管理节中的任何函数)时初始化。在计算运行时函数调用时间和解析初次调用运行时产生的错误码时必须牢记这点。

　　一旦运行时在主机线程中初始化，在主机线程中通过一些运行时函数调用分配的任何资源(存储器，流，事件等)只在当前主机线程的上下文中有效。因此只有在这个主机线程中调用的运行时函数(存储器拷贝，内核发射等)才能操作这些资源。这是因为CUDA上下文(参见3.3.1节)作为初始化的一部分建立且成为主机线程的当前上下文，且不能成为其它主机线程的当前上下文。

　　在多设备的系统中，内核默认在0号设备上执行，详见3.2.3节。

　　3.2.1 设备存储器

　　正如2.4节所提到的，CUDA编程模型假定系统包含主机和设备，它们各有自己独立的存储器。内核不能操作设备存储器，所以运行时提供了分配，释放，拷贝设备存储器和在设备和主机间传输数据的函数。

　　设备存储器可被分配为线性存储器或CUDA数组。

　　CUDA数组是不透明的存储器层次，为纹理获取做了优化。它们的细节在3.2.4节。

　　计算能力1.x的设备，其线性存储器存在于32位地址空间内，计算能力2.0的设备，其线性存储器存在于40位地址空间内，所以独立分配的实体能够通过指针引用，如，二叉树。

　　典型地，线性存储器使用cudaMalloc()分配，通过cudaFree()释放，使用cudaMemcpy()在设备和主机间传输。在2.1节的向量加法代码中，向量要从主机存储器复制到设备存储器：

// Device code
__global__ void VecAdd( float * A, float * B, float * C, int N) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N)
C[i] = A[i] + B[i];
}
// Host code
int main(){
int N = ...;
size_t size = N * sizeof ( float );
// Allocate input vectors h_A and h_B in host memory
float * h_A = ( float * )malloc(size);
float * h_B = ( float * )malloc(size);
// Initialize input vectors
...
// Allocate vectors in device memory
float * d_A;
cudaMalloc(( void ** ) & d_A, size);
float * d_B;
cudaMalloc(( void ** ) & d_B, size);
float * d_C;
cudaMalloc(( void ** ) & d_C, size);
// Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// Invoke kernel
int threadsPerBlock = 256 ;

　　线性存储器也可以通过cudaMallocPitch()和cudaMalloc3D()分配。在分配2D或3D数组的时候，推荐使用，因为这些分配增加了合适的填充以满足5.3.2.1节的对齐要求，在按行访问时或者在二维数组和设备存储器的其它区域间复制(用cudaMemcpy2D()和cudaMemcpy3D()函数)时，保证了最佳性能。返回的步长(pitch,stride)必须用于访问数组元素。下面的代码分配了一个尺寸为width*height的二维浮点数组，同时演示了怎样在设备代码中遍历数组元素。

// 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];

演示

// Host code
cudaPitchedPtr devPitchedPtr;
cudaExtent extent = make_cudaExtent( 64 , 64 , 64 );
cudaMalloc3D( & devPitchedPtr, extent);
MyKernel <<< 100 , 512 >>> (devPitchedPtr, extent);
// Device code
__global__ void MyKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent) {
char * devPtr = devPitchedPtr.ptr;
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
for ( int z = 0 ; z < extent.depth; ++ z) {
char * slice = devPtr + z * slicePitch;
for ( int y = 0 ; y < extent.height; ++ y) {
float * row = ( float * )(slice + y * pitch);
for ( int x = 0 ; x < extent.width; ++ x) {

　　参考手册列出了在cudaMalloc()分配的线性存储器，cudaMallocPitch()或cudaMalloc3D()分配的线性存储器，CUDA数组和为声明在全局存储器和常量存储器空间分配的存储器之间拷贝的所有各种函数。

　　下面的例子代码复制了一些主机存储器数组到常量存储器中：

__constant__ float constData[ 256 ];
float data[ 256 ];
cudaMemcpyToSymbol(constData, data, sizeof (data));

　　为声明在全局存储器空间的变量分配的存储器的地址，可以使用cudaGetSymbolAddress()函数检索到。分配的存储器的尺寸可以通过cudaGetSymbolSize()函数获得。

　　3.2.2 共享存储器

　　共享存储器使用__shared__限定词分配，详见B.2节。

　　正如在2.2节提到的，共享存储器应当比全局存储器更快，详见5.3.2.3节。任何用访问共享存储器取代访问全局存储器的机会应当被发掘，如下面的矩阵相乘例子展示的那样。

　　下面的代码是矩阵相乘的一个直接的实现，没有利用到共享存储器。每个线程读入A的一行和B的一列，然后计算C中对应的元素，如图3-1所示。这样，A读了B.width次，B读了A.height次。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.width + col)
typedef struct {
int width;
int height;
float * elements;
} Matrix;
// Thread block size
#define BLOCK_SIZE 16
// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel( const Matrix, const Matrix, Matrix);
// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul( const Matrix A, const Matrix B, Matrix C) {
// Load A and B to device memory
Matrix d_A;
d_A.width = A.width; d_A.height = A.height;
size_t size = A.width * A.height * sizeof ( float );
cudaMalloc(( void ** ) & d_A.elements, size);
cudaMemcpy(d_A.elements, A.elements, size,

　　下面的例子代码利用了共享存储器实现矩阵相乘。本实现中，每个线程块负责计算一个小方阵Csub，Csub是C的一部分，而块内的每个线程计算Csub的一个元素。如图3-2所示。Csub等于两个长方形矩阵的乘积：A的子矩阵尺寸是(A.width,block_size)，行索引与Csub相同，B的子矩阵的尺寸是(block_size,A.width)，列索引与Csub相同。为了满足设备的资源，两个长方形的子矩阵分割为尺寸为block_size的方阵，Csub是这些方阵积的和。每次乘法的计算是这样的，首先从全局存储器中将二个对应的方阵载入共享存储器中，载入的方式是一个线程载入一个矩阵元素，然后一个线程计算乘积的一个元素。每个线程积累每次乘法的结果并写入寄存器中，结束后，再写入全局存储器。

　　采用这种将计算分块的方式，利用了快速的共享存储器，节约了许多全局存储器带宽，因为在全局存储器中，A只被读了(B.width/block_size)次同时B读了(A.height/block_size)次。

　　前面代码中的Matrix 类型增加了一个stride域，这样子矩阵能够用同样的类型有效表示。__device__函数(见B.1.1节)用于读写元素和从矩阵中建立子矩阵。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct {
int width;
int height;
int stride;
float * elements;
} Matrix;
// Get a matrix element
__device__ float GetElement( const Matrix A, int row, int col) {
return

// Read C from device memory
cudaMemcpy(C.elements, d_C.elements, size,
cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A.elements);
cudaFree(d_B.elements);
cudaFree(d_C.elements);
}

　　3.2.3 多设备

　　主机系统上可以有多个设备。可以枚举这些设备，也可以查询他们的属性，可以选择它们中的一个执行内核。

　　多个主机线程可以在同一个设备上执行设备代码，但是设计成在某个既定时间，一个主机线程只能在一个设备上执行设备代码。这样，多个主机线程在多个设备上执行设备代码。在某个主机线程内，使用CUDA运行时创建的任何CUDA资源不能被其它线程使用。

　　下面的例子代码枚举了系统中的所有设备同时检索了它们的属性。也确定了支持CUDA的设备的数目。

int deviceCount;
cudaGetDeviceCount( & deviceCount);
int device;
for (device = 0 ; device < deviceCount; ++ device) {
cudaDeviceProp deviceProp;
cudaGetDeviceProperties( & deviceProp, device);
if (dev == 0 ) {
if (deviceProp.major == 9999 && deviceProp.minor == 9999 )
printf( " There is no device supporting CUDA.\n " );
else if (deviceCount == 1 )
printf( " There is 1 device supporting CUDA\n " );
else

　　默认情况下，只要一个非运行时设备管理函数调用(例外参见3.6节)，主机线程隐式的使用0号设备。可以通过调用cudaSetDevice()函数来启用其它的设备。一旦设备启用，无论是显式的还是隐式的，其后对cudaSetDevice()的调用都会失败，除非调用了cudaThreadExit()。cudaThreadExit()清理所有与主机调用线程相关的运行相关的资源。随后的运行时API调用将重新初始化运行时。

　　3.2.4 纹理存储器

　　CUDA支持纹理硬件的一个子集，GPU为图形使用这个子集访问纹理存储器。如5.3.2.5节所示，从纹理存储器而不是全局存储器中读数据有许多性能好处。

　　如B.8节所示，在内核中，调用纹理获取设备函数读纹理存储器。纹理获取的第一个参数指定的对象称为纹理参考。

　　纹理参考定义了被获取的纹理存储器部分。如3.2.4.3节所述，纹理参考在被内核使用之前，必须使用运行时函数绑定到存储器的某个区域，这个区域称为纹理。多种不同的纹理参考可能绑定到同一纹理或者绑定到存储器重叠的纹理。

　　纹理参考有许多属性。其中之一就是维数，维数指定纹理是作为一维的数组使用一个纹理坐标、二维数组使用两个纹理坐标、还是三维数组使用三维坐标来寻址。数组的元素称为texels，是纹理参考元素的简称。

　　其它属性定义纹理获取的输入输出数据类型，也包括怎样解释输入坐标和要做那些处理。

　　纹理可以是线性存储器的任何一个区域或者一个CUDA数组。

　　CUDA数组是为纹理获取优化的不透明的存储器层次。它们可以是一维的，二维的或三维的，也可由多个元素组成，每个元素可有1，2或4个组件，这些组件可能是有符号或无符号8，16或32位整形，16位浮点(目前只在驱动API中支持)，或32位浮点。CUDA数组只能在内核中通过纹理获取读取，且只能绑定到和已打包的组件数目相同的纹理参考。

　　2.2.4.1 纹理参考声明

　　纹理参考的一些属性不可变并且在编译时必须知道;它们在声明纹理参考时指定。纹理参考必须在文件域内声明，变量类型为texture;

texture < Type, Dim, ReadMode > texRef;

　　其中：

　　Type指定纹理获取时的返回的数据类型，Type限制为基本的整形和单精度浮点型和B.3.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的纹理元素映射为1;如果ReadMode是cudaReadModeElementType，不会进行转换;ReadMode是个可选参数，默认为cudaReadModeElementType。

　　3.2.4.2 运行时纹理参考属性

　　纹理参考的其它属性是可变的，并且能够在运行时通过主机运行时改变。这些属性指定纹理坐标是否归一化、寻址模式和纹理滤波，细节如下。

　　默认情况下，纹理使用[0，N)范围内的浮点坐标引用，其中N是坐标对应维度的尺寸。例如，尺寸为64*32的纹理可引用的坐标范围是x维[0，63]和y维[0，31]。归一化的纹理坐标范围指定为[0.0，1.0)而不是[0，N)，所以同样的64*32纹理的归一化坐标x维和y维可寻址范围都是[0,1)。归一化的纹理坐标天然的符合某些应用的要求，如果为了让纹理坐标独立于纹理尺寸，就更可取了。

　　寻址模式定义了当纹理坐标越界时发生了什么了。当使用非归一化纹理坐标时，纹理坐标在[0，N)范围之外的被钳位(clamp)：小于0的设置为0而大于等于N的设置为N-1。钳位也是使用归一化纹理坐标时默认的寻址模式：小于0.0或大于1.0钳位到[0.0，1.0)范围。对于归一化坐标，也可以指定为循环寻址模式。一般在纹理有周期性信号时使用循环模式。循环模式只使用纹理坐标的小数部分;如1.25和0.25等同，-1.25和0.75等同。

　　线性纹理滤波只能对返回值为浮点型的纹理配置起作用。它在周围的纹理元素点上执行低精度插值。如果启用滤波，纹理获取点周围的点被读取，纹理获取点的返回值基于那些纹理坐标落入那些元素中间的元素进行插值。对于一维的纹理进行简单的线性插值，而二维纹理使用双线性插值。

　　附录F给出了纹理获取的细节。

　　3.2.4.3 纹理绑定

　　如参考手册中所解释的，运行时API有一个低级的C风格的接口和一个高级的C++风格的接口。texture类型是在高级API中定义的一个结构体，公有继承自在低级API中定义的textrueReference类型。textureReference定义如下：

　　1、normalized指定纹理坐标是否归一化;如果非零，纹理中所有元素可寻址的纹理坐标范围是[0，1]，而不是[0，width-1]，[0，height-1]，或[0，depth-1]，其中width, height和depth是纹理尺寸。

　　2、filterMode指定滤波模式，即纹理获取时，如何根据输入的纹理坐标计算返回值;filterMode 等于cudaFilterModePoint或cudaFilterModeLinear;如果是cudaFilterModePoint，则所返回的值为纹理坐标最接近输入纹理坐标的纹理元素;如果等于 cudaFilterModeLinear，则所返回的值为纹理坐标最接近输入纹理坐标的两个(针对一维纹理)、四个(针对二维纹理)或八个(针对三维纹理)纹理元素的线性插值;对于浮点型的返回值，cudaFilterModeLinear 是惟一的有效值。

　　3、addressMode 指定寻址模式，即如何处理越界的纹理坐标;addressMode 是一个尺寸为 3 的数组，其第一个、第二个和第三个元素各自指定第一个、第二个和第三个纹理坐标的寻址模式;寻址模式可等于 cudaAddressModeClamp，此时越界的纹理坐标将被钳位到有效范围之内，也可等于 cudaAddressModeWrap，此时越界的纹理坐标将被环绕到有效范围之内;cudaAddressModeWrap仅支持归一化的纹理坐标。

　　4、channelDesc 描述获取纹理时返回值的格式;channelDesc类型定义如下：

struct cudaChannelFormatDesc {
int x, y, z, w;
enum cudaChannelFormatKind f;
};

　　其中 x、y、z 和 w 是返回值各组件的位数，而 f 为：

　　a. cudaChannelFormatKindSigned，如果这些组件是有符号整型;

　　b. cudaChannelFormatKindUnsigned，如果这些组件是无符号整型;

　　c. cudaChannelFormatKindFloat，如果这些组件是浮点类型。

　　normalized、addressMode 和 filterMode 可直接在主机代码中修改。

　　在内核中使用纹理参考从纹理存储器中读取数据之前，必须使用 cudaBindTexture() 或 cudaBindTextureToArray() 将纹理参考绑定到纹理。cudaUnbindTexture()用于解绑定纹理参考。

　　下面的代码将纹理参考绑定到devPtr指针指向的线性存储器：

　　使用低级API：

texture < float , 2 , cudaReadModeElementType > texRef;
textureReference * texRefPtr;
cudaGetTextureReference( & texRefPtr, “texRef”);
cudaChannelFormatDesc channelDesc =
cudaCreateChannelDesc < float > ();

　　使用高级API

texture < float , 2 , cudaReadModeElementType > texRef;
cudaChannelFormatDesc channelDesc =
cudaCreateChannelDesc < float > ();

　　下面的代码将纹理绑定到CUDA数组cuArray：

　　使用低级API

　　使用高级API

texture < float , 2 , cudaReadModeElementType > texRef;
cudaBindTextureToArray(texRef, cuArray);

　　声明纹理参考时指定的参数必须与将纹理绑定到纹理参考时指定的格式匹配;否则纹理获取的结果没有定义。

　　下面的代码在内核中应用了一些简单的转换。

// 2D float texture
texture < float , 2 , cudaReadModeElementType > texRef;
// Simple transformation kernel
__global__ void transformKernel( float * output,
int width, int height, float theta) {
// Calculate normalized texture coordinates
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
float u = x / ( float )width;
float v = y / ( float )height;
// Transform coordinates
u -= 0.5f ; v -= 0.5f ;
float tu = u * cosf(theta) – v * sinf(theta) + 0.5f ;
float tv = v * cosf(theta) + u * sinf(theta) + 0.5f ;
// Read from texture and write to global memory
output[y * width + x] = tex2D(tex, tu, tv);
}
// Host code
int main()
{
// Allocate CUDA array in device memory
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(
32 , 0 , 0 , 0 , cudaChannelFormatKindFloat);
cudaArray * cuArray;
cudaMallocArray( & cuArray, & channelDesc, width, height);
// Copy to device memory some data located at address h_data
// in host memory
cudaMemcpyToArray(cuArray, 0 , 0 ,h_data,size,
cudaMemcpyHostToDevice);
// Set texture parameters
texRef.addressMode[ 0 ] = cudaAddressModeWrap;
texRef.addressMode[ 1 ] = cudaAddressModeWrap;
texRef.filterMode = cudaFilterModeLinear;
texRef.normalized = true ;
// Bind the array to the texture
cudaBindTextureToArray(texRef, cuArray, channelDesc);
// Allocate result of transformation in device memory

　　3.2.5 分页锁定主机存储器

　　运行时提供了分配和释放分页锁定主机存储器(也称为pinned)的函数cudaHostAlloc()和cudaFreeHost()，分页锁定主机存储器与常规的使用malloc()分配的可分页的主机存储器不同。

　　使用分页锁定主机存储器有许多优点：

　　1、如3.2.6节提到的，在某些设备上，设备存储器和分页锁定主机存储器间数据拷贝可与内核执行并发进行;

　　2、在一些设备上，分页锁定主机内存可映射到设备地址空间，减少了和设备间的数据拷贝，详见3.2.5.3节;

　　3、在有前端总线的系统上，如果主机存储器是分页锁定的，主机存储器和设备存储器间的带宽会高些，如果再加上3.2.5.2节所描述的写结合(write-combining)的话，带宽会更高。

　　然而分页锁定主机存储器是稀缺资源，所以可分页内存分配得多的话，分配会失败。另外由于减少了系统可分页的物理存储器数量，分配太多的分页锁定内存会降低系统的整体性能。

　　SDK中的simple zero-copy例子中有分页锁定API的详细文档。

　　3.2.5.1可分享存储器(portable memory)

　　一块分页锁定存储器可被任何主机线程使用，但是默认的情况下，只有分配它的线程可以使用它。为了让所有线程可以使用它，可以在使用cudaHostAlloc()分配时传入cudaHostAllocPortable标签。

　　3.2.5.2 写结合存储器

　　默认情况下，分页锁定主机存储器是可缓存的。可以在使用cudaHostAlloc()分配时传入cudaHostAllocWriteCombined标签使其被分配为写结合的。写结合存储器没有一级和二级缓存资源，所以应用的其它部分就有更多的缓存可用。另外写结合存储器在通过PCI-e总线传输时不会被监视(snoop)，这能够获得高达40%的传输加速。

　　从主机读取写结合存储器极其慢，所以写结合存储器应当只用于那些主机只写的存储器。

　　3.2.5.3 被映射存储器

　　在一些设备上，在使用cudaHostAlloc()分配时传入cudaHostAllocMapped标签可分配一块被映射到设备地址空间的分页锁定主机存储器。这块存储器有两个地址：一个在主机存储器上，一个在设备存储器上。主机指针是从cudaHostAlloc()返回的，设备指针可通过cudaHostGetDevicePointer()函数检索到，可以使用这个设备指针在内核中访问这块存储器。

　　从内核中直接访问主机存储器有许多优点：

　　a. 无须在设备上分配存储器，也不用在这块存储器和主机存储器间显式传输数据;数据传输是在内核需要的时候隐式进行的。

　　b. 无须使用流(参见3.2.6.4节)重叠数据传输和内核执行;数据传输和内核执行自动重叠。

　　由于被映射分页锁定存储器在主机和设备间共享，应用必须使用流或事件(参见3.2.6节)来同步存储器访问以避免任何潜在的读后写，写后读，或写后写危害。

　　一块分页锁定存储器可同时分配为被映射的和可分享的(见3.2.5.1节)，这种情况下，每个要映射这块存储器的主机线程必须调用cudaHostGetDevicePointer()检索设备指针，因为每个主机线程持有的设备指针一般不同。

　　为了在给定的主机线程中能够检索到被映射分页锁定存储器的设备指针，必须在调用任何CUDA运行时函数前调用cudaSetDeviceFlags()，并传入cudaDeviceMapHost标签。否则，cudaHostGetDevicePointer()将会返回错误。

　　如果设备不支持被映射分页锁定存储器，cudaHostGetDevicePointer()将会返回错误。

　　应用可能会查询设备是否支持映射分页锁定主机存储器，可以使用函数cudaGetDeviceProperties()检查canMapHostMemory属性。

　　注意：从主机和其它设备的角度看，操作被映射分页锁定存储器的原子函数(5.4.3和B.10节)不是原子的。

　　3.2.6 异步并发执行

　　3.2.6.1 主机和设备间异步执行

　　为了易于使用主机和设备间的异步执行，一些函数是异步的：在设备完全完成任务前，控制已经返回给主机线程了。它们是：

　　1、内核发射;

　　2、存储器拷贝函数中带有Async后缀的;

　　3、设备间数据拷贝函数;

　　4、设置设备存储器的函数;

　　程序员可通过将CUDA_LAUNCH_BLOCKING环境变量设置为1来全局禁用所有运行在系统上的应用的异步内核发射。提供这个特性只是为了调试，永远不能作为使软件产品运行得可靠的方式。

　　当应用通过CUDA调试器或CUDA profiler运行时，所有的内核发射都是同步的。

　　3.2.6.2 数据传输和内核执行重叠

　　一些计算能力1.1或更高的设备可在内核执行时，在分页锁定存储器和设备存储器之间拷贝数据。应用可以通过调用cudaGetDeviceProperties()函数检查deviceOverlap属性查询这种能力。这种能力目前只支持不涉及CUDA数组和使用cudaMallocPitch()分配的二维数组的存储器拷贝(参见3.2.1节)。

　　3.2.6.3 并发内核执行

　　一些计算能力2.0的设备可并发执行多个内核。应用可以通过调用cudaGetDeviceProperties()函数检查concurrentKernels属性以查询这种能力。

　　设备最大可并发执行的内核数目是4。

　　来自不同CUDA上下文的内核不能并发执行。

　　使用了许多纹理或大量本地存储器的内核和其它内核并发执行的可能性比较小。

　　3.2.6.4 并发数据传输

　　在计算能力2.0的设备上，从主机分页锁定存储器复制数据到设备存储器和从设备存储器复制数据到主机分页锁定存储器，这两个操作可并发执行。

　　3.2.6.5 流

　　应用通过流管理并发。流是一系列顺序执行的命令。另外，流之间相对无序的或并发的执行它们的命令;这种行为是没有保证的，而且不能作为正确性的的保证(如内核间的通信没有定义)。

　　可以通过创建流对象来定义流，且可指定它作为一系列内核发射和设备主机间存储器拷贝的流参数。下面的代码创建了两个流且在分页锁定存储器中分配了一个名为hostPtr的浮点数组。

cudaStream_t stream[ 2 ];
for ( int i = 0 ; i < 2 ; ++ i)
cudaStreamCreate( & stream[i]);
float * hostPtr;

　　下面的代码定义的每个流是一个由一次主机到设备的传输，一次内核发射，一次设备到主机的传输组成的系列。

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,

　　每个流将它的hostPtr输入数组的部分拷贝到设备存储器数组inputdevPtr，调用MyKernel()内核处理inputDevPtr，然后将结果outputDevPtr传输回hostPtr同样的部分。使用两个流处理hostPtr允许一个流的传输和另一个流的执行重叠。为了使用重叠hostPtr必须指向分页锁定主机存储器。

　　最后调用的cudaThreadSynchronize()保证所有的流在进一步执行前已经完成。这个函数强制运行时等待所有流中的任务都完成。cudaStreamSynchronize()强制运行时等待某个流中的任务都完成。可用于同步主机和特定流，同时允许其它流继续执行。cudaStreamQuery()用于查询流中的所有之前的命令是否已经完成。为了避免不必要的性能损失，这些函数最好用于计时或隔离失败的发射或存储器拷贝。

　　调用cudaStreamDestroy()来释放流。

for ( int i = 0 ; i < 2 ; ++ i)
cudaStreamDestroy(stream[i]);

　　cudaStreamDestroy()等待指定流中所有在前的任务完成，然后释放流并将控制权返回给主机线程。

　　如果是下面情况，来自不同流的两个命令也不能并发，或分页锁定主机存储器分配，设备存储器分配，设备存储器设置，设备之间拷贝，或主机线程在0号流中调用的任何它们之间的CUDA命令(包含没有指定任何流参数的内核发射和设备和主机间存储器拷贝)。

　　任何需要依赖检测以确定内核发射是否完成的操作会阻塞CUDA上下文中后面任何流中所有的内核发射直至被检测的内核发射完成。需要依赖检测的操作包括同一个流中的一些其它类似被检查的发射的命令和流中的任何cudaStreamQuery()调用。因此，应用应当遵守这些指导以提升潜在的内核并发执行：

　　1、所有独立操作应当在依赖操作之前发出。

　　2、任何类型同步尽量延后。

　　G.4.1节描述的一级缓存和共享存储器的配置切换为所有未完成的内核发射插入了一个设备端同步栅栏。

　　3.2.6.6 事件

　　通过在应用的任意点上异步地记载事件和查询事件是否真正被记载，运行时提供了精密地监测设备运行进度的方式，事件同时也提供精确计时功能。当事件记载点前面，事件指定的流中的所有任务或者命令全部完成时，事件被记载。只有记载点之前所有的流中的任务/命令都已完成，0号流的事件才会记载。

　　下面的代码创建了两个事件：

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 );

　　可用这种方式销毁事件

cudaEventDestroy(start);
cudaEventDestroy(stop);

　　3.2.6.7 同步调用

　　直到设备真正完成任务，同步函数调用的控制权才会返回给主机线程。在主机线程执行任何其它CUDA调用前，通过调用cudaSetDeviceFlags()并传入指定标签(参见参考手册)可以指定主机线程的让步，阻塞，或自旋状态。

　　3.2.7 图形学互操作性

　　一些OpenGL和Direct3D的资源可被映射到CUDA地址空间，要么使CUDA可以读OpenGL或Direct3D写的数据，要么使CUDA写数据供OpenGL或Direct3D消费。

　　资源必须先在CUDA中注册，才能被3.2.7.1和3.2.7.2提到的函数映射。这些函数返回一个指向cudaGraphicsResource类型结构体的CUDA图形资源。资源注册是潜在高消耗的，因此通常每个资源只注册一次。可以使用cudaGraphicsUnregisterResource()解注册CUDA图形资源。

　　一旦资源被注册到CUDA，就可以按需要被任意次的映射和解映射，映射和解映射使用cudaGraphicsMapResources()和cudaGraphicsUnmapResources()。可以使用cudaGraphicsResourceSetMapFlags()来指定资源用处(只读，只写)，CUDA驱动可以据此优化资源管理。

　　可以获得cudaGraphicsResourceGetMappedPointer()为缓冲区返回的设备地址空间和cudaGraphicsSubResourceGetMappedArray()为CUDA数组返回的设备地址空间，内核通过读写这些空间读写被映射资源。

　　通过OpenGL或Direct3D访问被映射到CUDA的OpenGL或Direct3D的资源，其结果未定义。

　　3.2.7.1和3.2.7.2节给出了每种图形API的特性和一些代码例子。

　　3.2.7.1 OpenGL互操作性

　　和OpenGL互操作要求在其它任何运行时函数调用前，使用cudaGLSetGLDevice()指定CUDA设备。注意cudaSetDevice()和cudaGLSetDevice()是相互排斥的。

　　可以被映射到CUDA地址空间的OpenGL资源有OpenGL缓冲区、纹理和渲染缓存对象。

　　使用cudaGraphicsGLRegisterBuffer()注册缓冲对象。在CUDA中，缓冲对象表现为设备指针，因此可以在内核中读写或通过cudaMemcpy()调用。

　　纹理或渲染缓存对象使用cudaGraphicsGLRegisterImage()注册。在CUDA中，它们表现为CUDA数组，可绑定到纹理参考，可被内核读写或通过cudaMemcpy2D()调用。cudaGraphicsGLRegisterImage()使用内置的float类型(例如，GL_RGBA_FLOAT32)和非归一化整数(例如GL_RGBA8UI)支持所有纹理格式。请注意，由于GL_RGBA8UI是OpenGL3.0纹理格式，只能被着色器写，不能被固定功能的管线写。

　　下面的代码使用内核动态的修改一个存储在顶点缓冲对象中的二维width*height顶点网格。

GLuint positionsVBO;
struct cudaGraphicsResource * positionsVBO_CUDA;

int main()
{
// Explicitly set device
cudaGLSetGLDevice( 0 );
// Initialize OpenGL and GLUT
...
glutDisplayFunc(display);
// Create buffer object and register it with CUDA
glGenBuffers( 1 , positionsVBO);
glBindBuffer(GL_ARRAY_BUFFER, & vbo);
unsigned int size = width * height * 4 * sizeof( float );
glBufferData(GL_ARRAY_BUFFER, size, 0 , GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, 0 );
cudaGraphicsGLRegisterBuffer( & positionsVBO_CUDA, positionsVBO,

void display()
{
// Map buffer object for writing from CUDA
float4 * positions;
cudaGraphicsMapResources( 1 , & positionsVBO_CUDA, 0 );
size_t num_bytes;
cudaGraphicsResourceGetMappedPointer(( void ** ) & positions, & num_bytes,

　　在Windows系统上和对于Quadro显卡，可以用cudaWGLGetDevice()检索关联到wglEnumGpusNV()返回的句柄的CUDA设备。Quadro显卡与OpenGL的互操作性能比GeForce和Tesla要好。在一个多GPU的系统中，在Quadro GPU上运行OpenGL渲染，在其它的GPU进行CUDA计算。

　　3.2.7.2 Direct3D互操作性

　　Direct3D互操作性支持Direct3D 9，Direct3D 10，和Direct3D 11。

　　一个CUDA上下文一次只能和一个Direct3D设备互操作，且CUDA上下文和Direct3D设备必须在同一个GPU上创建，而且Direct3D设备必须使用D3DCREATE_HARDWARE_VERTEXPROCESSING标签创建。

　　和Direct3D的互操作性要求：在任何其它的运行时函数调用前，使用cudaD3D9SetDirect3DDevice()，cudaD3D10SetDirect3DDevice() 和cudaD3D11SetDirect3DDevice()指定Direct3D设备。可用cudaD3D9GetDevice(),cudaD3D10GetDevice() 和 cudaD3D11GetDevice()检索关联到一些适配器的CUDA设备。

　　可以被映射到CUDA地址空间的Direct3D资源有Direct3D缓冲区，纹理和表面。可以使用cudaGraphicsD3D9RegisterResource()， cudaGraphicsD3D10RegisterResource()和cudaGraphicsD3D11RegisterResource()注册这些资源。

　　下面的代码使用内核动态的修改一个存储在顶点缓冲对象中的二维width*height网格顶点。

　　Direct3D 9版本

IDirect3D9 * D3D;
IDirect3DDevice9 * device;
struct CUSTOMVERTEX {
FLOAT x, y, z;
DWORD color;
};
IDirect3DVertexBuffer9 * positionsVB;
struct cudaGraphicsResource * positionsVB_CUDA;

int main() {
// Initialize Direct3D
D3D = Direct3DCreate9(D3D_SDK_VERSION);
// Get a CUDA-enabled adapter
unsigned int adapter = 0 ;
for (; adapter < g_pD3D -> GetAdapterCount(); adapter ++ ) {
D3DADAPTER_IDENTIFIER9 adapterId;
g_pD3D -> GetAdapterIdentifier(adapter, 0 , & adapterId);
int dev;
if (cudaD3D9GetDevice( & dev, adapterId.DeviceName) == cudaSuccess)
break ;
}

// Create device
...
D3D -> CreateDevice(adapter, D3DDEVTYPE_HAL, hWnd, D3DCREATE_HARDWARE_VERTEXPROCESSING, & params , & device);
// Register device with CUDA
cudaD3D9SetDirect3DDevice(device);
// Create vertex buffer and register it with CUDA
unsigned int size = width * height * sizeof (CUSTOMVERTEX);
device -> CreateVertexBuffer(size, 0 , D3DFVF_CUSTOMVERTEX, D3DPOOL_DEFAULT, & positionsVB, 0 );
cudaGraphicsD3D9RegisterResource( & positionsVB_CUDA, positionsVB, c

void Render() {
// Map vertex buffer for writing from CUDA
float4 * positions;
cudaGraphicsMapResources( 1 , & positionsVB_CUDA, 0 );
size_t num_bytes;
cudaGraphicsResourceGetMappedPointer(( void ** ) & positions, & num_bytes, positionsVB_CUDA));
// Execute kernel
dim3 dimBlock( 16 , 16 , 1 );
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1 );
createVertices <<< dimGrid, dimBlock >>> (positions, time, width, height);
// Unmap vertex buffer
cudaGraphicsUnmapResources( 1 , & positionsVB_CUDA, 0 );
// Draw and present
...
}

void releaseVB() {
cudaGraphicsUnregisterResource(positionsVB_CUDA);
positionsVB -> Release();
}

__global__ void createVertices(float4 * positions, float time, unsigned int width, unsigned int height) {
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
// Calculate uv coordinates
float u = x / ( float )width;
float v = y / ( float )height;

　　Direct3D 10版本：

ID3D10Device * device;
struct CUSTOMVERTEX {
FLOAT x, y, z;
DWORD color;
};
ID3D10Buffer * positionsVB;
struct cudaGraphicsResource * positionsVB_CUDA;
int main() {
// Get a CUDA-enabled adapter
IDXGIFactory * factory;
CreateDXGIFactory(__uuidof(IDXGIFactory), ( void ** ) & factory);
IDXGIAdapter * adapter = 0 ;
for (unsigned int i = 0 ; ! adapter; ++ i) {
if (FAILED(factory -> EnumAdapters(i, & adapter))
break ;

factory -> Release();
// Create swap chain and device
D3D10CreateDeviceAndSwapChain(adapter, D3D10_DRIVER_TYPE_HARDWARE, 0 , D3D10_CREATE_DEVICE_DEBUG, D3D10_SDK_VERSION, & swapChainDesc, & swapChain, & device);
adapter -> Release();
// Register device with CUDA
cudaD3D10SetDirect3DDevice(device);
// Create vertex buffer and register it with CUDA
unsigned int size = width * height * sizeof (CUSTOMVERTEX);
D3D10_BUFFER_DESC bufferDesc;
bufferDesc.Usage = D3D10_USAGE_DEFAULT;
bufferDesc.ByteWidth = size;
bufferDesc.BindFlags = D3D10_BIND_VERTEX_BUFFER;
bufferDesc.CPUAccessFlags = 0 ;
bufferDesc.MiscFlags = 0 ;
device -> CreateBuffer( & bufferDesc, 0 , & positionsVB);
cudaGraphicsD3D10RegisterResource( & positionsVB_CUDA, positionsVB, cudaGraphicsRegisterFlagsNone);
cudaGraphicsResourceSetMapFlags(positionsVB_CUDA, cudaGraphicsMapFlagsWriteDiscard);
/ Launch rendering loop
while (...) {
Render();
}
}
void Render() {
// Map vertex buffer for writing from CUDA
float4 * positions;
cudaGraphicsMapResources( 1 , & positionsVB_CUDA, 0 );
size_t num_bytes;
cudaGraphicsResourceGetMappedPointer(( void ** ) & positions, & num_bytes, positionsVB_CUDA));
// Execute kernel
dim3 dimBlock( 16 , 16 , 1 );
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1 );
createVertices <<< dimGrid, dimBlock >>> (positions, time, width, height);
// Unmap vertex buffer
cudaGraphicsUnmapResources( 1 , & positionsVB_CUDA, 0 );
// Draw and present
}
void releaseVB() {
cudaGraphicsUnregisterResource(positionsVB_CUDA);
positionsVB -> Release();
}
__global__ void createVertices(float4 * positions, float time, unsigned int width, unsigned int height) {
unsigned int x =

　　Direct3D 11 版本：

ID3D11Device * device;
struct CUSTOMVERTEX {
FLOAT x, y, z;
DWORD color;
};
ID3D11Buffer * positionsVB;
struct cudaGraphicsResource * positionsVB_CUDA;

int main() {
// Get a CUDA-enabled adapter
IDXGIFactory * factory;
CreateDXGIFactory(__uuidof(IDXGIFactory), ( void ** ) & factory);
IDXGIAdapter * adapter = 0 ;
for (unsigned int i = 0 ; ! adapter; ++ i) {
if (FAILED(factory -> EnumAdapters(i, & adapter))
break ;
int dev;
if (cudaD3D11GetDevice( & dev, adapter) == cudaSuccess)
break ;
adapter -> Release();
}
factory -> Release();
// Create swap chain and device
…
sFnPtr_D3D11CreateDeviceAndSwapChain(adapter, D3D11_DRIVER_TYPE_HARDWARE, 0 , D3D11_CREATE_DEVICE_DEBUG, featureLevels, 3 , D3D11_SDK_VERSION, & swapChainDesc, & swapChain, & device, & featureLevel, & deviceContext);
adapter -> Release();
// Register device with CUDA
cudaD3D11SetDirect3DDevice(device);
// Create vertex buffer and register it with CUDA
unsigned int size = width * height * sizeof (CUSTOMVERTEX);
D3D11_BUFFER_DESC bufferDesc;
bufferDesc.Usage = D3D11_USAGE_DEFAULT;
bufferDesc.ByteWidth = size;
bufferDesc.BindFlags = D3D11_BIND_VERTEX_BUFFER;
bufferDesc.CPUAccessFlags = 0 ;
bufferDesc.MiscFlags = 0 ;
device -> CreateBuffer( & bufferDesc, 0 , & positionsVB);
cudaGraphicsD3D11RegisterResource( & positionsVB_CUDA,
positionsVB,
cudaGraphicsRegisterFlagsNone);
cudaGraphicsResourceSetMapFlags(positionsVB_CUDA,

// Execute kernel
dim3 dimBlock( 16 , 16 , 1 );
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1 );
createVertices <<< dimGrid, dimBlock >>> (positions, time, width, height);
// Unmap vertex buffer
cudaGraphicsUnmapResources( 1 , & positionsVB_CUDA, 0 );
// Draw and present
...
}

void releaseVB() {
cudaGraphicsUnregisterResource(positionsVB_CUDA);
positionsVB -> Release();
}
__global__ void createVertices(float4 * positions, float time, unsigned int width, unsigned int height){
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
// Calculate uv coordinates
float u = x / ( float )width;
float v = y / ( float )height;
u = u * 2.0f - 1.0f ;
v = v * 2.0f - 1.0f ;

　　3.2.8 错误处理

　　所有的运行时函数都返回错误码，但对于异步函数(参见3.2.6节)，由于会在任务结束前返回，因此错误码不能报告异步调用的错误;错误码只报告在任务执行之前的错误，典型的错误有关参数有效性;如果异步调用出错，错误将会在后面某个无关的函数调用中出现。

　　唯一能够检查异步调用出错的方式是通过在异步调用函数后面使用cudaThreadSynchronize()同步(或使用3.2.6节介绍的其它同步机制)，然后检查cudaThreadSynchronize()的返回值。

　　运行时为每个主机线程维护着一个初始化为cudaSuccess的错误变量，每次错误发生(可以是参数不正确或异步错误)时，该变量会被错误码重写。cudaGetLastError()会返回这个变量，并将它重新设置为cudaSuccess。

　　内核发射不返回任何错误码，所以应当在内核发射后立刻调用cudaGetLastError()检测发射前错误。为保证cudaGetLastError()返回的错误值不是由于内核发射前错误导致的，必须保证运行时错误变量在内核发射前被设置为cudaSuccess，可以通过在内核发射前调用cudaGetLastError()实现。内核发射是异步的，因此为了检测异步错误，应用必须在内核发射和cudaGetLastError()之间同步。

　　注意cudaStreamQuery()可能返回cudaErrorNotReady，而由于cudaEventQuery()没有考虑错误，因此不会被cudaGetLastError()报告。

　　3.2.9 使用设备模拟模式调试

　　CUDA-GDB可以用于调试计算能力大于1.0的设备(参看 CUDA-GDB用户手册获得支持平台)。编译器和运行时也支持模拟模式调试，模拟模式可以在没有CUDA支持的设备上运行。当使用模拟模式编译应用时(使用-deviceemu选项)，设备代码为主机编译，并在主机上运行，允许程序员使用主机上的本地调试器，就好像是主机程序一样。预处理宏__DEVICE_EMULATION__是为这种模式定义的。使用了任何库的所有应用代码必须一致的用设备模拟模式或设备执行模式编译。链接设备模拟模式编译的代码和设备执行模式编译的代码会返回下面的运行时初始化错误：cudaErrorMixedDeviceExecution。

　　在设备模拟模式下运行应用，运行时模拟编程模型。对于块内的每个线程，运行时在主机上建立一个线程。程序员要确定：

　　1、主机能够运行每个块的最大线程数，外加一主线程。

　　2、足够的可用存储器用于运行所有线程，已知每个线程都要占据256 KB栈空间。

　　设备模拟模式提供的许多特性让他成为一种有效的调试工具：

　　1、通过使用主机的本地调试支持，程序员可利用调试器支持的所有特性，例如断点设置和数据检查。

　　2、由于设备代码编译后在主机上运行，故可调用无法在设备上运行的代码扩展功能，如文件输入和输出操作或输出到屏幕(printf() 等)。

　　3、由于所有数据都位于主机上，因而可从设备或主机代码可读取任何特定于设备或主机的数据;类似地，设备和主机可调用任何设备或主机函数。

　　4、万一误用了内置同步函数，运行时将检测到死锁情况。

　　程序员必须牢记，设备模拟模式的目的在于模拟设备，而非仿真。因此，设备模拟模式在查找算法错误时非常有用，但是很难发现某些错误：

　　1、竞争条件在设备模拟模式中体现的可能很小，因为同时执行的线程数量要比实际运行在设备上的少得多。

　　2、在主机上解引用指向全局存储器的指针或在设备上解引用指向主机存储器的指针时，设备执行几乎必然以某种没有定义的方式失败，而设备模拟能得到正确的结果。

　　3、大多数时候，相同的浮点计算在设备上执行时所得到的结果，与在设备模拟模式中得到的结果并不完全相同。这是可预见的，因为一般只要使用略有差异的编译器选项，相同的浮点计算就会得到不同的结果，更不用说不同的编译器、不同的指令集或不同的架构了。