§cuda driver api的使用

不同于平时使用的runtime api，driver api更底层，使用方式与opencl差不多，编程的主要过程为创建上下文、传输数据、加载内核映像、读取内核函数并执行、拷贝结果等。

§头文件和库

需要包含头文件：#include <cuda.h>

在linux上编译：g++ main.cpp -lcuda

在windows上编译：cl main.cpp -link "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\lib\x64\cuda.lib"

§创建上下文

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15


cuInit(0);
int device_count = 0;
cuDeviceGetCount(&device_count);
if (0 == device_count) {
    std::cerr << "无可用cuda设备!" << std::endl;
    return -1;
}
CUdevice device;
cuDeviceGet(&device, 0);
CUcontext context;
cuCtxCreate(&context, 0, device);

// 其它依赖上下文的操作

cuCtxDestroy(context);

这样的创建方式可以为一个GPU指定多个上下文，也可以由cuDevicePrimaryCtxRetain来获得GPU上唯一的主要上下文，这样可以更好的与runtime api或者其它cuda库合作：

1
2
3
4
5
6
7
8


CUcontext context;
cuDevicePrimaryCtxRetain(&context, device);
cuCtxPushCurrent(context);

// 其它依赖上下文的操作

cuCtxPopCurrent(&context);
cuDevicePrimaryRelease(device);

§数据传输

driver api的数据传输接口与runtime api十分类似，但driver api中指向显存的指针有自己的类型CUdeviceptr，从cuda.h中的定义来看，它应该是unsigned long long类型。分配内存和传输数据的过程如下：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23


// 显存指针
CUdeviceptr d_A = 0;
CUdeviceptr d_B = 0;
CUdeviceptr d_C = 0;

// 分配内存
cuMemAlloc(&d_A, size);
cuMemAlloc(&d_B, size);
cuMemAlloc(&d_C, size);

// 复制数据
cuMemcpyHtoD(d_A, h_A, size);
cuMemcpyHtoD(d_B, h_B, size);

/* 执行内核 */

// 取回结果
cuMemcpyDtoH(h_C, d_C, size);

// 释放内存
cuMemFree(d_A);
cuMemFree(d_B);
cuMemFree(d_C);

其中，h_A和h_B是指向内存的指针，size为需要复制的字节数，可由unsigned long size = N * sizeof(double or float)给出，N是数据个数。

§执行内核代码

内核代码应该事先编译成ptx代码。ptx(Parallel Thread eXecution)代码是编译后的GPU代码的一种中间形式，它可以再次编译为原生的GPU微码(cubin)。

从cuda程序的编译流程可以知道，runtime api写出来的程序需要分离成设备代码与主机代码，使用不同的编译器进行编译。但是使用driver api时，主机代码与设备代码是完全分开的，只需要分别使用nvcc和c++编译器编译即可。

生成ptx代码的方式为命令行：nvcc -ptx kernel.cu -arch compute_50 //-Xcompiler -wd4819 cmake：

1
2
3


add_library(kernel OBJECT kernel.cu)
set_property(TARGET kernel PROPERTY CUDA_PTX_COMPILATION ON)
install(TARGETS kernel OBJECTS DESTINATION ptx)

在执行cmake install之前，生成的ptx代码位于build/CMakeFiles/kernel.dir/kernel.ptx里。

在主机代码中调用核函数显然不可能再交给nvcc处理，因此只能只能通过相应的接口将ptx代码加载进来并发射到GPU上：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13


CUmodule module;
cuModuleLoad(&module, "kernel.ptx");
CUfunction vector_add;
cuModuleGetFunction(&vector_add, module, "AddKernel");

// 启动kernel
int threads_per_block = 1024;
unsigned long blocks_per_grid = (N + threads_per_block - 1) /  threads_per_block;
void *args[] = {&d_A, &d_B, &d_C, &N};
cuLaunchKernel(vector_add,
        blocks_per_grid, 1, 1,
        threads_per_block, 1, 1,
        0, nullptr, args, nullptr);

这样的加载方式是即时编译的，拥有较好的兼容性，只要真实架构大代不低于指定的虚拟架构，都能为它生成代码。如果目标GPU在编码时就已经确定，还可以直接编译成二进制GPU代码，以获得更好的性能。

编译成二进制GPU代码的命令为：nvcc -cubin kernel.cu -arch=compute_50 -code=sm_50 //-Xcompiler -wd4819 或者生成多个目标GPU的二进制代码：nvcc -cubin kernel.cu -gencode arch=compute_50,code=sm_50 -gencode arch=compute_61,code=61 //-Xcompiler -wd4819 。

二进制代码的加载方式与ptx代码的加载方式一样：

1
2


CUmodule module;
cuModuleLoad(&module, "kernel.cubin");

但是需要注意，kernel.cubin的真实架构必须和运行程序的GPU一致，同时，否则将引起错误：device kernel image is invalid。

内核的执行需要设置块数、块内线程数、参数以及共享内存，主要接口为cuLauchKernel，完整的接口声明如下：

1
2
3
4
5


CUresult cuLaunchKernel(CUfunction f,
		unsigned int  gridDimX, unsigned int  gridDimY, unsigned int  gridDimZ,
		unsigned int  blockDimX, unsigned int  blockDimY, unsigned int  blockDimZ,
		unsigned int  sharedMemBytes, CUstream hStream,
		void **kernelParams, void **extra)；

在gridDimX $\times$ gridDimY $\times$ gridDimZ网格块上调用内核f，每个块包含blockDimX $\times$ blockDimY $\times$ blockDimZ线程。sharedMemBytes设置每个线程块可用的动态共享内存量。

f的内核参数可以用以下两种方式之一指定：

(1）可以通过kernelParams指定内核参数。如果f有N个参数，那么kernelParams需要是N个指针的数组。 kernelParams[0]到kernelParams[N-1]中的每一个都必须指向将从中复制实际内核参数的内存区域。不需要指定内核参数的数量及其偏移量和大小，因为直接从内核的映像中检索该信息。

(2）内核参数也可以打包到单个缓冲区中，通过extra参数传入。这需要知道每个内核参数的大小和它们在缓冲区内的对齐/填充信息。

注意：必须使用3.2及以上的工具链编译内核代码，使其包含内核参数信息，如果没有内核参数，则cuLaunchKernel将返回CUDA_ERROR_INVALID_IMAGE。

§编写内核

内核函数必须放在独立的文件中，后缀名为.cu，写法与使用runtime api一般无二，向量加法的一个示例为：

1
2
3
4
5
6


#include <device_launch_parameters.h>
extern "C" __global__ void AddKernel(double *A, double *B, double *C) {
    unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N)
        C[i] = A[i] + B[i];
}

注意：

头文件可以不包含，编译出来的ptx文件中没有对头文件的处理，即无论有没有第一行编译出来的ptx文件完全一样。但是包含头文件可以让编辑器更好的工作。
extern "C"不能省略，否则编译出的代码中函数名字不正确，以上面的代码为例子，如果不加extern "C"，编译出来的代码中函数名字为_Z9AddKernelPKdS0_Pdi。

§错误处理

错误处理的两个主要API为：

1
2
3
4


// 获取错误代码枚举名称的字符串表示形式.
CUresult cuGetErrorName( CUresult error, const char** pStr );
// 获取错误代码的字符串描述.
CUresult cuGetErrorString( CUresult error, const char** pStr );

可以把它们包装成一个宏，在出错时予以提示：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15


#define DEBUG
#ifndef DEBUG
#define CheckError(call) do { call; } while(0)
#else
#define CheckError(call) do {\
    const CUresult status = call;\
    if (status != CUDA_SUCCESS) {\
        const char *error_string = nullptr;\
        cuGetErrorString(status, &error_string);\
        std::cerr << "于文件" << __FILE__ << "第" << __LINE__ << "行发生错误！\n";\
        std::cerr << error_string << std::endl;\
        exit(-1);\
    }\
} while (0)
#endif

使用方式为

1

CheckError(cuModuleLoad(&module, "kernel.ptx"));

cuda driver api的使用