CUDA0 编程接口.docx

CUDA0 编程接口CUDA 3.0 编程接口 目前可用两种接口写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) Ci = Ai + Bi; / 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 = rowc; 演示 / 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 constData256; float data256; 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 CUDAn"); 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 texture<float, 2, cudaReadModeElementType> texRef; textureReference* texRefPtr; cudaGetTextureReference(&texRefPtr, “texRef”); cudaChannelFormatDesc channelDesc; cudaGetChannelDesc(&channelDesc, cuArray); 使用高级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 outputy * 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.addressMode0 = cudaAddressModeWrap; texRef.addressMode1 = 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分配时传入