cuda linux 算力_CUDA 3.0 编程接口

【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__voidVecAdd(float*A,float*B,float*C,intN) {inti=blockDim.x*blockIdx.x+threadIdx.x;if(i

C[i]=A[i]+B[i];

}//Host codeintmain(){intN=...;

size_t size=N*sizeof(float);//Allocate input vectors h_A and h_B in host memoryfloat*h_A=(float*)malloc(size);float*h_B=(float*)malloc(size);//Initialize input vectors...//Allocate vectors in device memoryfloat*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 memorycudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);//Invoke kernelintthreadsPerBlock=256;

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

//Host codefloat*devPtr;intpitch;

cudaMallocPitch((void**)&devPtr,&pitch,

width*sizeof(float), height);

MyKernel<<<100,512>>>(devPtr, pitch);//Device code__global__voidMyKernel(float*devPtr,intpitch) {for(intr=0; r

演示

//Host codecudaPitchedPtr devPitchedPtr;

cudaExtent extent=make_cudaExtent(64,64,64);

cudaMalloc3D(&devPitchedPtr, extent);

MyKernel<<<100,512>>>(devPitchedPtr, extent);//Device code__global__voidMyKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent) {char*devPtr=devPitchedPtr.ptr;

size_t pitch=devPitchedPtr.pitch;

size_t slicePitch=pitch*extent.height;for(intz=0; z

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

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

__constant__floatconstData[256];floatdata[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)typedefstruct{intwidth;intheight