CUDA 30 编程接口.docx
- 文档编号:10806355
- 上传时间:2023-02-23
- 格式:DOCX
- 页数:32
- 大小:68.17KB
CUDA 30 编程接口.docx
《CUDA 30 编程接口.docx》由会员分享,可在线阅读,更多相关《CUDA 30 编程接口.docx(32页珍藏版)》请在冰豆网上搜索。
CUDA30编程接口
CUDA3.0编程接口
目前可用两种接口写CUDA程序:
CUDAC和CUDA驱动API。
一个应用典型的只能使用其中一种,但是遵守3.4节描述的限制时,可以同时使用两种。
CUDAC将CUDA编程模型作为C的最小扩展集展示出来。
任何包含某些扩展的源文件必须使用nvcc编译,nvcc的概要在3.1节。
这些扩展允许程序员像定义C函数一样定义内核和在每次内核调用时,使用新的语法指定网格和块的尺寸。
CUDA驱动API是一个低层次的C接口,它提供了从汇编代码或CUDA二进制模块中装载内核,检查内核参数,和发射内核的函数。
二进制和汇编代码通常可以通过编译使用C写的内核得到。
CUDAC包含运行时API,运行时API和驱动API都提供了分配和释放设备存储器、在主机和内存间传输数据、管理多设备的系统的函数等等。
运行时API是基于驱动API的,初始化、上下文和模块管理都是隐式的,而且代码更简明。
CUDAC也支持设备模拟,这有利于调试(参见节3.2.8)。
相反,CUDA驱动API要求写更多的代码,难于编程和调试,但是易于控制且是语言无关的,因为它处理的是二进制或汇编代码。
3.2节接着第二章介绍CUDAC。
也引入了CUDAC和驱动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.3PTX兼容性
一些PTX指令只被高计算能力的设备支持。
例如,全局存储器上的原子指令只在计算能力1.1及以上的设备上支持;双精度指令只在1.3及以上的设备上支持。
将C编译成PTX代码时,-arch编译器选项指定假定的计算能力。
因此包含双精度计算的代码,必须使用“-arch=sm_13”(或更高计算能力)编译,否则双精度计算将被降级为单精度计算。
为某些特殊计算能力生成的PTX代码始终能够被编译成相等或更高计算能力设备上的二进制代码。
3.1.4应用兼容性
为了在特定计算能力的设备上执行代码,应用加载的二进制或PTX代码必须满足如3.1.2节和3.1.3节说明的计算能力兼容性。
特别地,为了能在将来更高计算能力(不能产生二进制代码)的架构上执行,应用必须装载PTX代码并为那些设备即时编译。
CUDAC应用中嵌入的PTX和二进制代码,由-arch和-code编译器选项或-gencode编译器选项控制,详见nvcc用户手册。
例如
nvccx.cu
–gencodearch=compute_10,code=sm_10
–gencodearch=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”的简写(等价于“-gencodearch=compute_13,code=\’compute_13,sm_13\’”)。
3.2CUDAC
CUDAC为熟悉C语言的用户提供了一个简单途径,让他们能够轻易的写出能够在设备上执行的程序。
CUDAC包含了一个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节的向量加法代码中,向量要从主机存储器复制到设备存储器:
// Devicecode
__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];
}
// Hostcode
int main(){
int N = ...;
size_tsize = N * sizeof(float);
// Allocateinputvectorsh_Aandh_Binhostmemory
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
// Initializeinputvectors
...
// Allocatevectorsindevicememory
float* d_A;
cudaMalloc((void**)&d_A,size);
float* d_B;
cudaMalloc((void**)&d_B,size);
float* d_C;
cudaMalloc((void**)&d_C,size);
// Copyvectorsfromhostmemorytodevicememory
cudaMemcpy(d_A,h_A,size,cudaMemcpyHostToDevice);
cudaMemcpy(d_B,h_B,size,cudaMemcpyHostToDevice);
// Invokekernel
int threadsPerBlock = 256;
线性存储器也可以通过cudaMallocPitch()和cudaMalloc3D()分配。
在分配2D或3D数组的时候,推荐使用,因为这些分配增加了合适的填充以满足5.3.2.1节的对齐要求,在按行访问时或者在二维数组和设备存储器的其它区域间复制(用cudaMemcpy2D()和cudaMemcpy3D()函数)时,保证了最佳性能。
返回的步长(pitch,stride)必须用于访问数组元素。
下面的代码分配了一个尺寸为width*height的二维浮点数组,同时演示了怎样在设备代码中遍历数组元素。
// Hostcode
float* devPtr; int pitch;
cudaMallocPitch((void**)&devPtr, &pitch,
width * sizeof(float),height);
MyKernel<<<100, 512>>>(devPtr,pitch);
// Devicecode
__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];
演示
// Hostcode
cudaPitchedPtrdevPitchedPtr;
cudaExtentextent = make_cudaExtent(64, 64, 64);
cudaMalloc3D(&devPitchedPtr,extent);
MyKernel<<<100, 512>>>(devPitchedPtr,extent);
// Devicecode
__global__ void MyKernel(cudaPitchedPtrdevPitchedPtr,cudaExtentextent){
char* devPtr = devPitchedPtr.ptr;
size_tpitch = devPitchedPtr.pitch;
size_tslicePitch = 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次。
// Matricesarestoredinrow-majororder:
// M(row,col)=*(M.elements+row*M.width+col)
typedef struct {
int width;
int height;
float* elements;
}Matrix;
// Threadblocksize
#define BLOCK_SIZE16
// Forwarddeclarationofthematrixmultiplicationkernel
__global__ void MatMulKernel(const Matrix, const Matrix,Matrix);
// Matrixmultiplication-Hostcode
// MatrixdimensionsareassumedtobemultiplesofBLOCK_SIZE
void MatMul(const MatrixA, const MatrixB,MatrixC){
// LoadAandBtodevicememory
Matrixd_A;
d_A.width = A.width;d_A.height = A.height;
size_tsize = 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节)用于读写元素和从矩阵中建立子矩阵。
// Matricesarestoredinrow-majororder:
// M(row,col)=*(M.elements+row*M.stride+col)
typedef struct {
int width;
int height;
int stride;
float* elements;
}Matrix;
// Getamatrixelement
__device__ float GetElement(const MatrixA, int row, int col){
return
// ReadCfromdevicememory
cudaMemcpy(C.elements,d_C.elements,size,
cudaMemcpyDeviceToHost);
// Freedevicememory
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){
cudaDevicePropdeviceProp;
cudaGetDeviceProperties(&deviceProp,device);
if (dev == 0){
if (deviceProp.major == 9999 && deviceProp.minor == 9999)
printf("ThereisnodevicesupportingCUDA.\n");
else if (deviceCount == 1)
printf("Thereis1devicesupportingCUDA\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指定纹理获取时的返回的数据类型,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运行时纹理参考属性
纹理参考的其它属性是可变的,并且能够在运行时通过主机运行时改变。
这些属
- 配套讲稿:
如PPT文件的首页显示word图标,表示该PPT已包含配套word讲稿。双击word图标可打开word文档。
- 特殊限制:
部分文档作品中含有的国旗、国徽等图片,仅作为作品整体效果示例展示,禁止商用。设计者仅对作品中独创性部分享有著作权。
- 关 键 词:
- CUDA 30 编程接口 编程 接口
![提示](https://static.bdocx.com/images/bang_tan.gif)