并行编程实战——CUDA编程的驱动接口与运行时接口的分析

一、说明

在刚刚学习CUDA编程时，向大家介绍过CUDA编程的框架中可以使用两类接口，即Driver API和Runtime API。由于本系列的文章安装应用了Window平台的Runtime API，所以基本的代码应用都是以此为主的。但在需要更紧密贴近硬件的开发中，Driver API则更有优势。
本文将对这两种接口进行整体的对比和分析，让大家有一个清晰明了的认知。

二、Driver API和Runtime API

对于CUDA Driver API与Runtime API，大家可以这样认为：

1. CUDA Driver API
   它是一种低级的，与硬件更接近的，提供了更精细操作控制并允许开发者进行更多细节上的手动开发处理的一种API，如同C/C++之与其它高级语言一样，虽然更复杂但也更灵活、功能也更强大且与上层语言隔离
2. Runtime API
   一种高级的基于Driver API的抽象，所以它更容易学习和应用，提供了大量高级的API用于快速的硬件处理及错误和异常的处理等。一般来说，大多数开发者都是以此为基础进行开发

三、二者的不同

既然对二者的基本内容进行了阐述，那么下面就对二者的具体的不同进行分析：

1. 抽象层次的不同
   Driver API面向底层是一种低级的抽象而Runtime API面向开发者，是一种更高级的抽象。
2. 操作和管理
   Runtime API的操作和管理大多被接口封装不需要开发者手动处理；但Driver API需要显式的手动操作，相对较复杂。
3. 代码的运行效率
   虽然二者的关联性很强，但最终的代码运行效率却因为操作的具体的控制不同，导致可能产生的运行效率不同。在便于编译器优化的部分，Runtime API有优势；在底层控制更细节的地方，Driver API有优势。
4. 内存应用方式
   Driver API对内存管理的精细度更高（比如指定大小和对齐等）但也更复杂；Runtime API则使用封装的接口操作更简单，但管理也比较粗疏
5. 并行化
   在并行的策略上，Runtime API可以通过接口实现高效的并行编程但受限于框架本身的封装；而Driver API可以进行更底层的优化，针对特定的硬件和特定算法，进一步挖掘优化的可能性，特别适合于定制化和高级的编程人员。
   而在并行的同步机制上，Runtime API应用简单，开发方便快捷；Driver API则可以进行细粒度的上下文手动操作，更好的支持了异步编程
6. JIT编译
   在Runtime API上虽然nvcc会自动调用JIT的编译，但其并没有直接调用PTX；而Driver API则支持直接通过PTX动态加载和编译模块
7. 整体的对比
   其实大家反过来从设计者的角度来理解为了什么会有Driver API和Runtime API就更容易理解二者的不同了。就是为了满足不同的目标人群。有需求有能力的来应用Driver API，大多数的开发者推荐使用Runtime API。所以二者的不同就清晰了：Driver API更适合于底层需求或框架开需求的开发；Runtime API适合上层的应用开发。

这时候可能就人要问了，那到底如何进行二者的选择呢？这个非常简单，能问出这种问题的，绝大多数都适合于使用Runtime API开发，他们往往是初学者或上层应用开发者，一般不会和底层进行交互；而使用Driver API的开发者，一般是开发底层应用，如CUDA相关的库或框架、需要进行JIT编译，或者需要进行复杂的异步编程和并行编程的上下文管理等等。

四、例程

通过上面的分析，基本明白了Driver API和Runtime API的不同，那么看二者实现代码的区别：

1. Runtime API：

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
#include <cassert>

__global__ void matrixAdd(float* A, float* B, float* C, int rows, int cols) {
	int row = blockIdx.y * blockDim.y + threadIdx.y;
	int col = blockIdx.x * blockDim.x + threadIdx.x;

	int id = row * cols + col;

	if (row < rows && col < cols) {
		C[id] = A[id] + B[id];
	}
}

int main() {
	int rows = 1 << 10; // 1024
	int cols = 1 << 10; // 1024
	int numElem = rows * cols;
	size_t size = numElem * sizeof(float);

	std::cout << "malloc : " << size / (1024 * 1024) << " MB" << std::endl;

	float *mA = new float[numElem];
	float *mB = new float[numElem];
	float *mC = new float[numElem];

	for (int i = 0; i < numElem; ++i) {
		mA[i] = rand() / (float)RAND_MAX;
		mB[i] = rand() / (float)RAND_MAX;
	}

	float *dmA, *dmB, *dmC;
	cudaMalloc((void**)&dmA, size);
	cudaMalloc((void**)&dmB, size);
	cudaMalloc((void**)&dmC, size);

	cudaMemcpy(dmA, mA, size, cudaMemcpyHostToDevice);
	cudaMemcpy(dmB, mB, size, cudaMemcpyHostToDevice);

	dim3 bSize(16, 16);
	dim3 gSize((cols + bSize.x - 1) / bSize.x,
		(rows + bSize.y - 1) / bSize.y);

	std::cout << "thread block num: [" << gSize.x << " * " << gSize.y
		<< "] , each block have[" << bSize.x << " * " << bSize.y
		<< "] threads " << std::endl;

	matrixAdd << <gSize, bSize >> >(dmA, dmB, dmC, rows, cols);

	cudaError_t err = cudaGetLastError();
	if (err != cudaSuccess) {
		std::cerr << "CUDA exec err: " << cudaGetErrorString(err) << std::endl;
		exit(EXIT_FAILURE);
	}

	cudaMemcpy(mC, dmC, size, cudaMemcpyDeviceToHost);

	cudaDeviceSynchronize();
	cudaDeviceProp prop;
	cudaGetDeviceProperties(&prop, 0);
	std::cout  << prop.name << "exec completed!" << std::endl;

	cudaFree(dmA);
	cudaFree(dmB);
	cudaFree(dmC);
	delete[] mA;
	delete[] mB;
	delete[] mC;

	cudaDeviceReset();

	std::cout << "The program has been executed successfully ！" << std::endl;
	return 0;
}

Driver API ：

const char* ptxCode = R"(
.version 8.0
.target sm_50
.address_size 64

.visible .entry matrixAdd(
    .param .u64 matrixAdd_param_0,
    .param .u64 matrixAdd_param_1,
    .param .u64 matrixAdd_param_2,
    .param .u32 matrixAdd_param_3,
    .param .u32 matrixAdd_param_4
)
{
    .reg .pred      %p<3>;
    .reg .b32       %r<16>;
    .reg .b64       %rd<16>;

    ld.param.u64    %rd1, [matrixAdd_param_0];
    ld.param.u64    %rd2, [matrixAdd_param_1];
    ld.param.u64    %rd3, [matrixAdd_param_2];
    ld.param.u32    %r1, [matrixAdd_param_3];
    ld.param.u32    %r2, [matrixAdd_param_4];
    mov.u32         %r3, %ctaid.y;
    mov.u32         %r4, %ntid.y;
    mov.u32         %r5, %tid.y;
    mad.lo.s32      %r6, %r3, %r4, %r5;
    mov.u32         %r7, %ctaid.x;
    mov.u32         %r8, %ntid.x;
    mov.u32         %r9, %tid.x;
    mad.lo.s32      %r10, %r7, %r8, %r9;
    setp.ge.s32     %p1, %r6, %r1;
    setp.ge.s32     %p2, %r10, %r2;
    or.pred         %p3, %p1, %p2;
    @%p3 bra        $L__BB0_2;
    mul.lo.s32      %r11, %r6, %r2;
    add.s32         %r12, %r11, %r10;
    mul.wide.s32    %rd4, %r12, 4;
    add.s64         %rd5, %rd1, %rd4;
    ld.global.f32   %f1, [%rd5];
    add.s64         %rd6, %rd2, %rd4;
    ld.global.f32   %f2, [%rd6];
    add.f32         %f3, %f1, %f2;
    add.s64         %rd7, %rd3, %rd4;
    st.global.f32   [%rd7], %f3;

$L__BB0_2:
    ret;
}
)";
#include <cuda.h>
#include <iostream>
#include <cstdlib>
#include <cstring>

int main() {
    CUdevice device;
    CUcontext context;
    CUmodule module;
    CUfunction function;
    CUdeviceptr dA, dB, dC;

    int rows = 1 << 10; // 1024
    int cols = 1 << 10; // 1024
    int numElem = rows * cols;
    size_t size = numElem * sizeof(float);

    std::cout << "malloc : " << size / (1024 * 1024) << " MB" << std::endl;

    float *mA = new float[numElem];
    float *mB = new float[numElem];
    float *mC = new float[numElem];

    for (int i = 0; i < numElem; ++i) {
        mA[i] = rand() / (float)RAND_MAX;
        mB[i] = rand() / (float)RAND_MAX;
    }

    // 1. 初始化Driver API
    cuInit(0);

    // 2. 获取设备
    cuDeviceGet(&device, 0);

    // 3. 创建上下文
    cuCtxCreate(&context, 0, device);

    // 4. 从PTX代码字符串加载模块
    cuModuleLoadDataEx(&module, ptxCode, 0, 0, 0);

    // 5. 获取函数句柄
    cuModuleGetFunction(&function, module, "matrixAdd");

    // 6. 分配设备内存
    cuMemAlloc(&dA, size);
    cuMemAlloc(&dB, size);
    cuMemAlloc(&dC, size);

    // 7. 将数据从主机复制到设备
    cuMemcpyHtoD(dA, mA, size);
    cuMemcpyHtoD(dB, mB, size);

    // 8. 设置网格和块维度
    dim3 bSize(16, 16);
    dim3 gSize((cols + bSize.x - 1) / bSize.x,
               (rows + bSize.y - 1) / bSize.y);

    std::cout << "thread block num: [" << gSize.x << " * " << gSize.y
              << "] , each block have[" << bSize.x << " * " << bSize.y
              << "] threads " << std::endl;

    // 9. 准备内核参数
    void* kernelParams[] = { &dA, &dB, &dC, &rows, &cols };

    // 10. 启动内核
    cuLaunchKernel(function,
                                   gSize.x, gSize.y, 1,    // 网格维度
                                   bSize.x, bSize.y, 1,    // 块维度
                                   0, 0,                   // 共享内存和流
                                   kernelParams, 0);      // 参数和额外参数

    // 11. 同步上下文，等待内核完成
    cuCtxSynchronize();

    // 12. 将结果从设备复制回主机
    cuMemcpyDtoH(mC, dC, size);

    // 13. 获取设备属性并打印完成信息
    char deviceName[256];
    cuDeviceGetName(deviceName, sizeof(deviceName), device);
    std::cout << deviceName << " exec completed!" << std::endl;

    // 14. 清理资源
    cuMemFree(dA));
    cuMemFree(dB));
    cuMemFree(dC));
    cuModuleUnload(module);
    cuCtxDestroy(context);

    delete[] mA;
    delete[] mB;
    delete[] mC;

    std::cout << "The program has been executed successfully!" << std::endl;
    return 0;
}

仍然使用的原来的代码，大家可以对比着看一下。

五、总结

正如前面反复提到的，一般开情况下，开发者使用的是运行时接口，但这不代码驱动接口应用就少，毕竟其功能更强大、灵活。另外，为了“二者得兼”，CUDA框架还提供了二者的混合编程。所以在某些情况下，就可以在兼顾双方的情况来实现具体的功能。