CUDA编程

http://www.cnblogs.com/stewart/archive/2013/01/05/2846860.html

目录：

1.什么是CUDA

2.为什么要用到CUDA

3.CUDA环境搭建

4.第一个CUDA程序

5. CUDA编程

5.1. 基本概念

5.2. 线程层次结构

5.3. 存储器层次结构

5.4. 运行时API

5.4.1. 初始化

5.4.2. 设备管理

5.4.3. 存储器管理

5.4.3.1. 共享存储器

5.4.3.2. 常量存储器

5.4.3.3. 线性存储器

5.4.3.4. CUDA数组

5.4.4. 流管理

5.4.5. 事件管理

5.4.6. 纹理参考管理

5.4.6.1. 纹理声明

5.4.6.2. 纹理绑定

5.4.6.3. 纹理获取

5.4.7. OpenGL互操作

5.4.8. Direct3D互操作

5.5. 驱动API

5.5.1. 初始化

5.5.2. 设备管理

5.5.3. 上下文管理

5.5.4. 模块管理

5.5.5. 执行控制

5.5.6. 存储器管理

5.5.7. 流管理

5.5.8. 事件管理

5.5.9. 纹理参考管理

5.5.10. OpenGL互操作

5.5.11. Direct3D互操作

5.6. 性能优化

5.7. NVCC编译器

5.8. 设备模拟

5.9. 其他

参考文献

 

1.什么是CUDA

       CUDA(Compute Unified Device Architecture)，统一计算架构，是NVidia推出的并行计算平台。NVidia官方对其的解释是：一个并行计算平台和简单（简洁）地使用图像处理单元（GPU）进行通用计算的编程模型。利用GPU的能力在计算性能上有惊人的提升。

       简单地说CUDA是便于程序员利用NVidia GPU进行通用计算的开发环境及工具，目前支持C/C++语言，将来还会支持Fortran语言。

 

2.为什么要用到CUDA

       CPU主频要比GPU高2-3倍左右，但是通常情况下GPU核心的数量要比CPU多2-3个数量级以上。因此GPU的计算能力要远大于CPU，充分发挥GPU的计算能力，可以有成倍的性能提升。   

       早期利用GPU的计算能力是使用着色器和着色语言（GLSL等）。目前广泛使用的是CUDA和OpenCL。CUDA是针对NVidia GPU硬件设备设计的，而 OpenCL是针对跨平台设计的。因此CUDA可充分发挥NVidia GPU的计算性能。

       CUDA可以直接使用C/C++语言来开发GPU程序，省去了程序员重新学一种新语言的麻烦。

 

3.CUDA环境搭建

       CUDA环境主要分为四点：硬件（GPU设备）、操作系统、C/C++编译器和CUDA工具包。

       硬件（GPU设备），必须是支持CUDA的GPU。可到NVidia官网查询支持CUDA的GPU设备，具体地址为：http://www.nvidia.com/object/cuda_home_new.html 。

       操作系统，支持Microsoft Windows、Mac OS X和Linux。

       C/C++编译器，对不同的操作系统有不同的要求。

       CUDA工具包，NVidia提供了不同操作系统对应的CUDA Toolkit，可从https://developer.nvidia.com/cuda-downloads 下载对应的版本。

       本文只以Microsoft Windows为例介绍如何搭建CUDA环境。

       准备材料：

       ·一台装有支持CUDA GPU的电脑。

       ·Microsoft Windows操作系统（Microsoft Windows XP,Vista,7,or 8 or Windows Server 2003 or 2008）。

       ·CUDA工具包（相应操作系统）。下载地址：https://developer.nvidia.com/cuda-downloads

       ·C/C++编译器：Microsoft Visual Studio 2008 或 2010，或者对应版本的Microsoft Visual C++ Express产品。

       安装步骤：

       ·在装有支持CUDA GPU的电脑上安装Microsoft Windows操作系统（一般情况下都已经完成这步骤）。

       ·安装C/C++编译器，可只安装其中的C++编译器部分。

       ·安装CUDA工具包。（CUDA工具包中有NVidia GPU的驱动程序，尚未安装的请选择安装。）

       安装验证：

       Windows XP系统：进入 C:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\CUDA Samples\v5.0\bin\win32\Release 目录运行deviceQuery.exe文件。

       Windows Vista, Windows 7, Windows 8, Windows Server 2003, and Windows Server 2008系统：进入 C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.0\bin\win32\Release 目录运行deviceQuery.exe文件。

       如果安装正确，执行deviceQuery.exe文件会得到GPU设备的相应信息。如果没有安装支持CUDA的GPU也会得出GPU的信息，其中CUDA Capability Major/Minor version number信息为9999.9999。

       Microsoft Windows上更详细的安装信息请查看：

http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-microsoft-windows/index.html 。

       Mac OS X的安装：

http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-mac-os-x/index.html 。

       Linux的安装：

http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-linux/index.html 。

 

4.第一个CUDA程序

       在Microsoft Windows系统上，如果成功搭建了CUDA环境，则在Microsoft Visual Studio中已经集成了CUDA的开发组件。

       以下以Windows 7 + Microsoft Visual Studio 2008为例，创建第一个CUDA程序。

       打开Microsoft Visual Studio 2008，依次：File->New->Project->NVIDIA->CUDA->CUDA 5.0 Runtime，输入相应的项目名称确定即可。

       默认会生成一个kernel.cu文件，内容如下：

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

#include

void addWithCuda(int *c, const int *a, const int *b, size_t size);

__global__ void addKernel(int *c, const int *a, const int *b)

{

    int i = threadIdx.x;

    c[i] = a[i] + b[i];

}

int main()

{

    const int arraySize = 5;

    const int a[arraySize] = { 1, 2, 3, 4, 5 };

    const int b[arraySize] = { 10, 20, 30, 40, 50 };

    int c[arraySize] = { 0 };

    // Add vectors in parallel.

    addWithCuda(c, a, b, arraySize);

    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",

        c[0], c[1], c[2], c[3], c[4]);

    // cudaThreadExit must be called before exiting in order for profiling and

    // tracing tools such as Nsight and Visual Profiler to show complete traces.

    cudaThreadExit();

    return 0;

}

// Helper function for using CUDA to add vectors in parallel.

void addWithCuda(int *c, const int *a, const int *b, size_t size)

{

    int *dev_a = 0;

    int *dev_b = 0;

    int *dev_c = 0;

    // Choose which GPU to run on, change this on a multi-GPU system.

    cudaSetDevice(0);

    // Allocate GPU buffers for three vectors (two input, one output)    .

    cudaMalloc((void**)&dev_c, size * sizeof(int));

    cudaMalloc((void**)&dev_a, size * sizeof(int));

    cudaMalloc((void**)&dev_b, size * sizeof(int));

    // Copy input vectors from host memory to GPU buffers.

    cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);

    cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

    // Launch a kernel on the GPU with one thread for each element.

    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

    // cudaThreadSynchronize waits for the kernel to finish, and returns

    // any errors encountered during the launch.

    cudaThreadSynchronize();

    // Copy output vector from GPU buffer to host memory.

    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_c);

    cudaFree(dev_a);

    cudaFree(dev_b);

}

代码1

       这是一个将两个一维数组相加的例子。

       其中addKernel是内核函数，它的计算过程是在GPU上实现的，用函数类型限定符__global__限制，且函数类型为void型。

       cuda_runtime.h头文件包括了运行时API和其参数的定义。（如果使用驱动API则使用cuda.h头文件）。

       device_launch_parameters.h头文件包含了内核函数的5个变量threadIdx、blockDim、blockIdx、gridDim和wrapSize。

       对其中CUDA运行时API函数的解释：

       ·cudaSetDevice()：选择设备（GPU）。（可以不使用，不使用的情况下，默认选择设备0）

       ·cudaMalloc()：动态分配显存。

       ·cudaMemcpy()：设备与主机之内的数据拷贝。

       ·cudaThreadSynchronize()：同步所有设备上的线程，等待所有线程结束。

       ·cudaFree():释放由cudaMalloc分配的显存。

       ·cudaThreadExit():结束CUDA上下文环境，释放其中的资源。

       这些函数的具体介绍在 http://docs.nvidia.com/cuda/cuda-runtime-api/index.html 中。

 

5. CUDA编程

5.1. 基本概念

       CUDA编程中需要注意一些基本概念，分别为：主机、设备、运行时API、驱动API、warp、bank、函数类型限定符、变量类型限定符、thread、block、grid、计算能力、SIMT、内置变量、纹理、CUDA数组等。

       主机：可理解为CPU与内存的组合。

       设备：可理解为GPU与显存的组合。

       运行时API：是指CUDA运行时API是在驱动API的基础上封装而成的，简化了CUDA的开发。

       驱动API：是指CUDA驱动API，相比运行时API更接近于设备，可灵活运用设备的特性开发CUDA，可实现运行时API无法实现的功能。

       warp：多处理器激活、管理、调度和执行并行任务的单位。计算能力2.x的设备warp为32个线程。未来的设备可能不同，可以通过内置变量warpSize查询。

       bank：为了获得较高的存储器带宽，共享存储器被划分为多个大小相等的存储器模块，称为存储体，这些存储体就叫bank，可同步访问。

       函数类型限定符：是CUDA C中特有的，用来修饰是主机函数，设备调用的设备函数，还是主机调用的设备函数。有__device__、__global__、__host__。

       变量类型限定符：是用来修饰设备变量的。有__device__、__constant__、__shared__。

       thread：设备中的线程，与主机中的线程是同一个概念。

       block：线程块，由一组线程组成。一个线程块中的所以线程会在同一个多处理器上执行，一个多处理器上可同时执行多个线程块。

       grid：有所有线程块组成的网格。

       计算能力：是NVidia GPU不同架构的计算能力。

       SIMT：单指令多线程，与单指令多数据（SIMD）类似。一条指令多个线程一同执行，实现程序的并行化。

       内置变量：有threadIdx、blockDim、blockIdx、gridDim、warpSize。其中threadIdx指此线程在线程块中的位置；blockDim指线程块维度；blockIdx指该线程块在网格中的位置；gridDim指线程块网格维度；warpSize指一个warp多少个线程。

       纹理：本文主要涉及到的是纹理参考、纹理绑定、纹理获取。

       CUDA数组：区别于线性存储器，对数据进行了对齐等的处理，包括一维、二维和三维。其中的数据为：一元、二元或四元组。

 

5.2. 线程层次结构

       CUDA线程的层次结构，由小到大依次为线程、线程块、线程块网格。一维、二维或三维的线程组组成一个线程块，一维、二维或三维的线程块组组成一个线程块网格。

       下图是由二维的线程块组组成的线程块网络，其中线程块是由二维的线程组组成。

       图1

       NVidia GPU的硬件结构是，一组流处理器组成一个多处理器，一个或多个多处理器组成一个GPU。其中流处理器，可以理解为处理计算的核心单元。多处理器类似于多核CPU。NVidia GPU从DX10（DirectX10）开始出现了Tesla、Fermi、Kepler架构，不同的架构多处理器中流处理器数量都有差别。

 

5.3. 存储器层次结构

       CUDA存储器有：寄存器、共享存储器、常量存储器、本地存储器、全局存储器、纹理存储器等。其中寄存器和本地存储器是线程私有的，共享存储器是对线程块中的所有线程可见，常量存储器、全局存储器和纹理存储器是对网格中所有线程可见。

       下图解释了存储器的层次结构：

       图2

 

5.4. 运行时API

       运用运行时API开发CUDA程序需要了解：初始化、设备管理、存储器管理、流管理、事件管理、纹理参考管理、OpenGL互操作和Direct3D互操作。

       运行时API文档地址为：http://docs.nvidia.com/cuda/cuda-runtime-api/index.html 。

 

5.4.1. 初始化

       运行时API不存在显示初始化函数，初始化会在首次调用运行时函数时完成。虽然不需要调用初始化函数进行初始化，但是退出时需要调用退出函数cudaThreadExit()释放资源。

 

5.4.2. 设备管理

       有些电脑上可能有多块设备，因此对于不同的要求选择合适的设备。设备管理主要是获取设备信息和选择执行设备。

       主要有三个函数：

       ·cudaGetDeviceCount()：得到电脑上设备的个数。

       ·cudaGetDeviceProperties()：获得对应设备的信息。

       ·cudaSetDevice()：设置CUDA上下文对应的设备。

       运行__global__函数前需要提前选择设备，如果不调用cudaSetDevice()函数，则默认使用0号设备。

       上面三个函数的具体用法请查看CUDA运行时API文档。

 

5.4.3. 存储器管理

       共享存储器、常量存储器、线性存储器和CUDA数组的使用是存储器管理的主要部分。

 

5.4.3.1. 共享存储器

       共享存储器，使用__shared__变量限定符修饰，可静态或动态分配共享存储器。

       静态分配共享存储器，是在设备代码中直接分配共享存储器的大小，如下代码：

#define SHARED_MEM 16

__global__ void kernel(…)

{
       __shared__ int shared[SHARED_MEM];
}

void main()
{
       kernel<<>>(…);
}

代码2

       动态分配共享存储器，是在主机代码中使用内核函数的第三个特定参数传入分配共享存储器的大小，如下代码：

#define SHARED_MEM 16

__global__ void kernel(…)
{
       extern __shared__ int shared[];
}

void main()
{
       int nSharedMem = (int)SHARED_MEM;
       kernel<<sizeof(int)>>>(…);
}

代码3

 

5.4.3.2. 常量存储器

       常量存储器，使用__constant__变量限定符修饰。使用常量存储器，是由于其在设备上有片上缓存，比全局存储器读取效率高很多。

       使用常量存储器时会涉及的运行时API函数主要有：

       ·cudaMemcpyToSymbol()

       ·cudaMemcpyFromSymbol()

       ·cudaGetSymbolAddress()

       ·cudaGetSymbolSize()

       主机代码中使用cudaGetSymbolAddress()获取__constant__或__device__定义的变量地址。设备代码中可通过提取__device__、__shared__或__constant__变量的指针获取变量地址。

 

5.4.3.3. 线性存储器

       线性存储器是使用cudaMalloc()、cudaMallocPitch()或cudaMalloc3D()分配的，使用cudaFree()释放。二维的时候建议使用cudaMallocPitch()分配，cudaMallocPitch()函数对对齐进行了调整。这三个分配函数对应cudaMemset()、cudaMemset2D()、cudaMemset3D()三个memset函数和cudaMemcpy()、cudaMemcpy2D()、cudaMemcpy3D()三个memcpy函数。

 

5.4.3.4. CUDA数组

       CUDA数组是使用cudaMallocArray()、cudaMalloc3DArray()分配的，使用cudaFreeArray()释放。

       相关memcpy函数请查阅CUDA运行时API文档。

       具体使用可查阅CUDA编程指南：

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html 。

 

5.4.4. 流管理

       主机设备之间的内存拷贝与内核在设备上执行是异步的。在不使用流的情况下，是这样执行的：设备先从主机上拷贝内存，拷贝完成之后，再在设备上执行内核代码计算，最后当内核执行完毕，再把设备上的内存拷贝到主机上。当使用两个流的情况下，0号流执行内核代码的同时1号流拷贝主机内存到设备，1号流执行的同时0号流拷贝设备内存到主机（具体的实现并不一定如此，这里是为了说明流的作用简单做了假设）。两个流的情况下，部分内存拷贝和内置执行是同时进行的（异步的），比同步的内存拷贝和内核执行节省了时间。

       与流有关的函数有：

       ·cudaStreamCreate()：流的创建；

       ·cudaStreamDestroy()：流的销毁；

       ·cudaStreamSynchronize()：流同步；

       ·*Async：与流相关的其他函数。

       内核<<<…>>>的第四个参数为哪个流。

       CUDA编程指南中有对流具体实现的讲解。

 

5.4.5. 事件管理

       由于部分CUDA运行时函数的执行与主机代码是异步的。在一块代码中，CUDA运行时函数执行没有结束就直接执行其后的主机代码了，主机并不知道已经执行到哪个CUDA运行时函数了。事件的引入就是为了解决这一问题，在CUDA运行时函数已经执行完毕后记入下事件，查询此事件是否记录就能知道那个CUDA运行时函数已经执行完毕。在CUDA运行时函数前后记入事件就能获得此函数执行的时间。

       与事件有关的函数有：

       ·cudaEventCreate()：事件的创建；

       ·cudaEventDestroy()：事件的销毁；

       ·cudaEventRecord()；记录事件；

       ·cudaEventSynchronize()：事件同步；

       ·cudaEventElapsedTime()：计算两事件的时间差。

       具体的实现请查询CUDA编程指南。

 

5.4.6. 纹理参考管理

       纹理参考的实现是由纹理声明、纹理绑定、纹理获取完成的。

 

5.4.6.1. 纹理声明

       纹理声明是在文件域中声明纹理变量，供主机使用CUDA函数绑定纹理和设备获取纹理。纹理声明为：

       texture texRef;

       其中：

       ·DateType：纹理元的格式，有float、unsigned char、signed char、unsigned short、signed short及它们的2元和4元组。

       ·Type：纹理参考格式，有cudaTextureType1D、cudaTextureType2D、cudaTextureType3D、cudaTextureType1DLayered、cudaTextureType2DLayered。是可选参数，默认为cudaTextureType1D。

       ·ReadMode：读取模式，有cudaReadModeElementType、cudaReadModeNormalizedFloat。是可选参数，默认为cudaReadModeElementType。为cudaReadModeNormalizedFloat时，纹理元数据使用了单位化的映射，映射到了[0.0, 1.0]或[-1.0, 1.0]。为cudaReadModeElementType时，不进行任何映射变换。

 

5.4.6.2. 纹理绑定

       纹理绑定是将分配的线性存储器或CUDA数组绑定到纹理存储器。CUDA运行时纹理绑定API分为高级和低级两种绑定类型。

       下面是CUDA编程文档上绑定线性存储器的例子：

       低级API：

texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;

textureReference* texRefPtr;

cudaGetTextureReference(&texRefPtr, texRef);

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();

size_t offset;

cudaBindTexture2D(&offset, texRefPtr, devPtr, &channelDesc, width, height, pitch);

代码4

       高级API：

texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();

size_t offset;

cudaBindTexture2D(&offset, texRef, devPtr, channelDesc, width, height, pitch);

代码5

       其中devPtr是由cudaMallocPitch()分配的线性存储器指针；width、height、pitch是cudaMallocPitch()使用或获得的变量。

       绑定线程存储器还可以使用cudaBindTexture()。

       绑定CUDA数组使用cudaBindTextureToArray()，CUDA编程文档上有如何使用的具体介绍。

 

5.4.6.3. 纹理获取

       纹理获取是在内核中获得纹理中某一个坐标对应的纹理元值。

       纹理获取的函数有：tex1Dfetch()、tex1D()、tex2D()、tex3D()、tex1Dlayered()、tex2Dlayered()、texCubemap()、texCubemapLayered()、tex2Dgather()。具体使用请查询CUDA编程文档。

 

5.4.7. OpenGL互操作

       OpenGL（Open Graphics Library），一个图形硬件API。

       OpenGL与CUDA互操作，主要是缓冲对象的注册与取消注册、映射与取消映射。对应的函数有：

       ·cudaGLRegisterBufferObject()：缓冲对象注册；

       ·cudaGLUnregisterBufferObject()：取消缓冲对象注册；

       ·cudaGLMapBufferObject()：映射缓冲对象；

       ·cudaGLUnmapBufferObject()：取消映射。

       cudaGLMapBufferObject()映射缓冲对象后，CUDA可以使用其返回的设备存储器地址读取和写入缓冲对象。

       CUDA关于OpenGL互操作的具体介绍请查询CUDA编程文档及运行时API。

       OpenGL部分的知识，请查看：

 

5.4.8. Direct3D互操作

       Direct3D是Microsoft自己的3D图形API。

       Direct3D与CUDA互操作，主要是Direct3D设备的设置、资源的注册、资源映射、映射后信息获取、取消映射、取消注册。对应的函数有（以Direct3D9为例）：

       ·cudaD3D9SetDirect3DDevice()：Direct3D设备的设置；

       ·cudaD3D9RegisterResource ()：注册资源；

       ·cudaD3D9MapResources()：资源映射；

       ·cudaD3D9ResourceGetMappedPointer()：获取资源映射后的CUDA设备存储器地址；

       ·cudaD3D9ResourceGetMappedSize()：获取大小；

       ·cudaD3D9ResourceGetMappedPitch()：获取间隔；

       ·cudaD3D9UnmapResources()：取消映射；

       ·cudaD3D9UnregisterResource()：取消注册。

       具体如何使用请查询CUDA编程文档。

       Direct3D部分的知识，请查询MSDN。

 

5.5. 驱动API

       驱动API是比运行时API更底层的一套接口，运行时API是在驱动API的基础上封装而成。驱动API是一种基于句柄、命令式的API：大多数对象都通过不透明的句柄引用。

       以下列了主要的句柄：

对象	句柄	描述
设备	CUdevice	支持CUDA的设备
上下文	CUcontext	大致等同于CPU进程
模块	CUmodule	大致等同于动态库
函数	CUfunction	内核
堆存储器	CUdeviceptr	设备存储器的指针
CUDA数组	CUarray	设备上一维或二维数据的不透明容器，可通过纹理参考读取
纹理参考	CUtexref	描述如何解释纹理存储器数据的对象

       运用驱动API开发CUDA程序需要了解：初始化、设备管理、上下文管理、模块管理、执行控制、存储器管理、流管理、事件管理、纹理参考管理、OpenGL互操作、Direct3D互操作。

       驱动API文档地址为：http://docs.nvidia.com/cuda/cuda-driver-api/index.html 。

 

5.5.1. 初始化

       驱动API与运行时API不同，需要在调用任何驱动API函数（不包括初始化函数）之前初始化。初始化函数为cuInit()。

 

5.5.2. 设备管理

       驱动API与运行时API不同，不需要设置设备，而是直接使用得到的设备句柄操作设备。

       设备管理的主要函数有：

       ·cuDeviceGetCount()：获得主机上设备总数；

       ·cuDeviceGet()：获得对应设备句柄；

       ·cuDeviceGetProperties()：获得设备信息。

       具体解释及其中参数信息请查阅驱动API文档。

 

5.5.3. 上下文管理

       CUDA上下文类似于CPU进程。在驱动程序API中执行的所有资源和操作都封装在CUDA上下文内在该上下文被销毁时，系统将自动清除这些资源。除了模块和纹理参考之类的对象之外，每个上下文都有自己独特的32位地址空间。

       一个主机线程只能有一个当前设备上下文。每个主机线程都有一个当前上下文堆栈，并为每个上下文维护一个使用计数。

       上下文管理的主要函数有：

       ·cuCtxCreate()：创建上下文；

       ·cuCtxDestroy()：销毁上下文；

       ·cuCtxPopCurrent()：使当前上下文离开当前主机线程上下文堆栈；

       ·cuCtxPushCurrent()：压入上下文到当前主机线程上下文堆栈；

       ·cuCtxAttach()：增加一个上下文计数；

       ·cuCtxDetach()：消耗一个上下文计数（当上下文使用计数为0时，自动销毁上下文）。

       具体使用请查询CUDA编程文档及驱动API文档。

 

5.5.4. 模块管理

       模块是可独立加载的设备代码和数据包，类似于windows中的DLL。所有符号的名称（包括函数、全局变量和纹理参考）均在模块范围内维护，从而使独立的第三方编写的模块可在相同的CUDA上下文中进行互操作。

       模块管理的主要函数有：

       ·cuModuleLoad()：模块加载；

       ·cuModuleGetFunction()：得到模块中相应函数。

       具体使用请查询CUDA编程文档及驱动API文档。

 

5.5.5. 执行控制

       执行控制是指，执行和控制设备代码（内核）。驱动API内核的执行，不同运行时API一样方便，需要设置额外设置grid、block和参数等，还是使用特定的launch函数。

       执行控制的主要函数有：

       ·cuFuncSetCacheConfig()：设置函数对应的cache偏好（是设置cache多还是共享内存多）；

       ·cuFuncSetSharedMemConfig()：设置共享内存bank的大小；

       ·cuLaunchKernel()：launch函数；

       ·cuFuncSetBlockShape()：设置block的函数；

       ·cuFuncSetSharedSize()：设置共享内存大小；

       ·cuLaunch()：launch函数；

       ·cuLaunchGrid()：launch函数；

       ·cuParamSetSize()：设置内核函数参数的长度；

       ·cuParamSet*()：设置内核函数参数。

       cuParam*()系列函数用于指定在下一次调用launch函数来启动内核时为内核提供的参数。其第二个参数指定参数在参数堆栈中的偏移。这个偏移量必须与参数类型的对齐要求相匹配。

       具体使用请查询CUDA编程文档及驱动API文档。

 

5.5.6. 存储器管理

       驱动API存储器管理与运行时API类似，只是API接口不同。

       存储器管理的主要函数有：

       ·cuMemAlloc()：分配线性存储器；

       ·cuMemAllocPitch()：分配线性存储器；

       ·cuMemFree()：释放线性存储器；

       ·cuArrayCreate()：创建数组；

       ·cuArrayDestroy()：销毁数组；

       ·cuMemcpy()：数据拷贝；

       ·cuMemcpy2D()：数据拷贝；

       ·cuMemcpy3D()：数据拷贝；

       ·cuMemcpyHtoD()：数据拷贝，从主机拷贝到设备。

       具体使用请查询驱动API文档。

 

5.5.7. 流管理

       驱动API流管理与运行时API类似，只是API接口不同。

       流管理的主要函数有：

       ·cuStreamCreate()：流创建；

       ·cuStreamDestroy()：流销毁；

       ·cuStreamQuery()：流查询；

       ·cuStreamSynchronize()：同步流；

       ·cuCtxSynchronize()：同步上下文。

       具体使用请查询驱动API文档。

 

5.5.8. 事件管理

       驱动API事件管理与运行时API类似，只是API接口不同。

       事件管理的主要函数有：

       ·cuEventCreate()：事件创建；

       ·cuEventDestroy()：事件销毁；

       ·cuEventElapsedTime()：计算两事件的时间差；

       ·cuEventQuery()：查询事件；

       ·cuEventRecord()：记录事件；

       ·cuEventSynchronize()：事件同步。

       具体使用请查询驱动API文档。

 

5.5.9. 纹理参考管理

       驱动API纹理参考管理与运行时API类似，只是API接口不同。

       纹理参考管理的主要函数有：

       ·cuTexRefCreate()：创建参考纹理；

       ·cuTexRefDestroy()：销毁纹理参考；

       ·cuTexRefSetAddress()：绑定纹理参考；

       ·cuTexRefSetArray()：绑定纹理参考；

       ·…：其他一系列与纹理参考有关的函数。

       具体使用请查询驱动API文档。

 

5.5.10. OpenGL互操作

       驱动API必须使用cuGLInit()初始化与OpenGL的互操作性，其他与运行时API类似。

       OpenGL互操作的主要参数有：

       ·cuGLInit()：初始化OpenGL互操作性；

       ·cuGLRegisterBufferObject()：注册缓冲对象；

       ·cuGLUnregisterBufferObject()：取消注册缓冲对象；

       ·cuGLMapBufferObject()：绑定缓冲对象；

       ·cuGLUnmapBufferObject()：取消绑定缓冲对象。

       具体使用请查询驱动API文档。

 

5.5.11. Direct3D互操作

       驱动API Direct3D互操作性要求在创建CUDA上下文时指定Direct3D设备。通过使用cuD3D9CtxCreate()而非cuCtxCreate()创建CUDA上下文即可实现此目标。其他与运行时API类型。

       Direct3D互操作的主要函数有：

       ·cuD3D9CtxCreate()：创建与Direct3D互操作的CUDA上下文；

       ·cuD3D9RegisterResource()：注册资源；

       ·cuD3D9UnregisterResource()：取消注册资源；

       ·cuD3D9MapResources()：绑定资源；

       ·cuD3D9UnmapResources()：取消绑定资源；

       ·cuD3D9ResourceGetMappedPointer()：获取资源映射后的CUDA设备存储器地址；

       ·cuD3D9ResourceGetMappedSize()：获取大小；

       ·cuD3D9ResourceGetMappedPitch()：获取间隔。

       具体使用请查询驱动API文档。

 

5.6. 性能优化

       性能优化主要有：warp中减少控制指令、合理使用共享内存、防止共享内存bank冲突、单个线程中寄存器使用的量、block中线程数、常量存储器的合理利用、线程对全局存储器的合理访问等。

       多处理器是以warp为单位处理线程的，有控制指令时，会执行完所有的控制指令对应的指令后才会继续执行下面的指令。举个例子，if/else语句两个方向的线程在同一个warp中，线程1执行if方向，线程2执行else方向，它们可能的执行顺序是这样的：

       ·线程1执行if方向，线程2等待

       ·线程1执行if方向完毕等待，线程2执行else方向

       ·线程2执行else方向完毕，线程1、2共同执行后面的指令

       共享内存属于片上缓存比全局存储器读写速度更快。把一部分全局存储器上的数据放入共享内存中处理可有效提高性能。共享存储器的访问速度和寄存器差不多，大约读写4B的数据需要两个时钟周期。共享存储器的读取是以半warp为单位的，当半warp中的多个线程访问数组元素处于同一个bank时会发生bank冲突。但假如线程1、2访问bank1中的同一块4字节数据，其他的线程访问互不冲突的bank时不会有bank冲突。当半warp中所有的线程都访问同一个bank中同一块4字节的数据时也不会发生bank冲突，称为广播访问，此时只访问一次bank。每个多处理器中的共享存储器大小是有限的，应按照block的大小分配合适的共享存储器。Block的大小会影响多处理器每次激活的block数。

       每个多处理器寄存器数量是有些的，而且在每个线程中寄存器是线程私有的。按照每个多处理器激活的线程数，合理分配寄存器。如果每个线程分配太多线程，则每个多处理器同时激活的线程数就会减少，从而影响并行效果。

       Block中的线程数（NThread）也会影响每个多处理器同时激活的线程数。每个多处理器有最大同时激活线程数（NMThread），且每个多处理器有最大同时激活block数（NMBlock）。Block中的线程数满足：NThread >= NMThread / NMBlock会激活在一个多处理器中可激活的所有block。在其他资源可充分利用的情况下，多处理器上同时激活的线程数越多，效率越高。

       常量存储器也是带片上缓存的存储器。充分利用常量存储器可有效提升性能。

       最新的设备，全局存储器都带有片上缓存。可以利用多处理器处理线程的特性合理访问全局存储器的数据，可使更多数据命中。

 

5.7. NVCC编译器

       NVCC编译器会分离源码中设备代码和主机代码，主机代码交由一般的C/C++编译器（gcc等）编译，设备代码由NVCC编译。

       具体编译命令请查阅NVCC文档：

http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html

 

5.8. 设备模拟

       对于未装支持CUDA的设备或者调试时，可使用设备模拟。设备模拟是将设备执行的代码由主机模拟执行，设备代码并不是在设备上执行，而是主机上模拟出多个线程执行。设备模拟的结果和实际设备实际的结果可能不同。

 

5.9. 其他

       CUDA并不支持windows的默认远程登入客户端（mstsc）登入远程主机执行设备。需要远程登入主机执行CUDA设备，可使用VNC工具。

 

参考文献：

《GPGPU编程技术——从GLSL、CUDA到OpenCL》——仇德元

http://blogs.nvidia.com/2012/09/what-is-cuda-2/

http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-microsoft-windows/index.html

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

http://docs.nvidia.com/cuda/cuda-runtime-api/index.html