入门级CUDA程序调试的通用方法——可用于Matlab的MexCuda

Matlab是学术界和业界公认的高级编程语言，其优点在于编程复杂度低，缺点在于for循环的效率极低。为了利用Matlab已有的算法，降低编程复杂度，同时解决for循环效率低的问题，Matlab引入了Mex机制，使得Matlab可以调用C++编写的函数。CUDA是Nvidia推出的并行计算框架，基本遵循了C++语法规则。于是乎出现了Matlab调用CUDA的方法，称之为“MexCuda”。本文就Mex CUDA程序的调试方法简单介绍下经验。

1. CUDA入门——vectorAdd

最简单的CUDA程序莫过于“vectorAdd.cu”，即向量的加法。下面给出了一个向量加法的简单例子：

#include <cstdio>#include <cuda_runtime.h>#define N 5// Kernel 函数__global__ void vectorAdd(const float* a, const float* b, float* c) {    int idx = threadIdx.x;    if (idx < N) {        c[idx] = a[idx] + b[idx];    }}int main() {    // 初始化数据    float a[N], b[N], c[N];    for(int i = 0; i < N; ++i) {        a[i] = i;        b[i] = i + 1;    }    // 将数据复制到 GPU    float* d_a, *d_b, *d_c;    cudaMalloc(&d_a, sizeof(float) * N);    cudaMalloc(&d_b, sizeof(float) * N);    cudaMalloc(&d_c, sizeof(float) * N);    cudaMemcpy(d_a, a, sizeof(float) * N, cudaMemcpyHostToDevice);    cudaMemcpy(d_b, b, sizeof(float) * N, cudaMemcpyHostToDevice);    // 执行Kernel    int blocksPerGrid = 1;    int threadsPerBlock = 64;    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c);    // 将数据从GPU取回    cudaMemcpy(c, d_c, sizeof(float) * N, cudaMemcpyDeviceToHost);    for(int i = 0; i < N; ++i) {       printf("c[%d] = %f\n", i, c[i]);    }    cudaFree(d_a);    cudaFree(d_b);    cudaFree(d_c);    return 0;}

将上面文件保存为vectorAdd.cu，在终端中执行

nvcc -o vectorAdd vectorAdd.cu./vectorAdd

可以得到

TomHeavendeMBP:CUDA tomheaven$ ./vectorAdd c[0] = 1.000000c[1] = 3.000000c[2] = 5.000000c[3] = 7.000000c[4] = 9.000000

2. 检查运行时错误

在上面的例子中，我们只用了1个block，64个thread，而且是最简单的线性排列。那么我们不禁要问，如果程序复杂起来，变量多起来，应该如何查错呢？

首先，引进一段查错代码：

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=false){    if (code != cudaSuccess)    {        printf("GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);        if (abort) exit(code);    } else {        printf("cuda returned code == cudaSuccess\n");    }}

在调用内核语句的后面加上两行查错指令：

vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c);gpuErrchk( cudaPeekAtLastError() );gpuErrchk( cudaDeviceSynchronize() );

编译之后执行可得

TomHeavendeMBP:CUDA tomheaven$ ./vectorAdd cuda returned code == cudaSuccesscuda returned code == cudaSuccessc[0] = 1.000000c[1] = 3.000000c[2] = 5.000000c[3] = 7.000000c[4] = 9.000000

说明内核执行过程没有错误。注意PeekLastError检查的是执行过程有没有直接错误返回，而deviceSynchronize检查的是设备同步状态，包括“共享内存溢”出这样的问题。这个两个检查缺一不可。良好的编程习惯是其他每一个cuda打头的函数都用gpuErrchk包裹起来，这样可以及时发现运行时的错误。

3. 必须注意的问题

下面说说新手最容易犯的错误。用CUDA编程的时候，必须时刻牢记一些硬件限制：

• 单卡的最大线程数：每个维度不超过 65535。当block一维排布时，就是65535。

• 每个Block的最大线程数：1024 （老N卡512）

• 每个Block的最大共享内存：大约48KB。这是非常小的存储空间。

当以上任意一个限制被突破，内核执行都会出现错误。最常见的错误是“CUDA_ILLEGAL_ADDRESS”，即访问了非法显存。但是原因并不是代码逻辑出错——访问了非法地址，而是设备资源不足。

针对以上问题，有如下常用解决方法：

• 如果是单卡最大线程数不够，可以考虑将block设计成二维甚至三维，这样又获得了很多线程资源。

• 如果是单个Block线程过多，缩小“threadsPerBlock”这个参数。

• 第三个问题是比较隐蔽的，新手往往容易忽略。可以自己估算一下每个Block使用的共享内存大小，48KB是很容易超过的。缩小“threadsPerBlock”这个参数来使得每个Block中的线程数减少，从而满足共享内存限制。再不济，不使用“shared”关键字，直接使用内存，当然，这样会牺牲一些效率。功能实现后，优化代码时，想尽办法少用共享内存，比如能用float就尽量不用double。

4. 逻辑调试

前面的方法可以确保程序能编译通过并执行起来。但是程序逻辑到底对不对呢？我们需要一个好的调试方法。我爬了一些帖子和教程，有说用VS的，有说用Nsight的。这些方法是有用，但是有很大局限性，比如不跨平台，不能用于调试Matlab的mexCuda，最大的局限性就是一个字“慢”。

本人在写C++程序的时候就有一个观点：输出调试是最好的调试方法。为什么呢？一是调试很快，不用断点单步。二是对于环境没有要求，只要程序执行了就有输出。这也就意味着无论使用Matlab调用的，还是Python调用的，输出调试都起作用。

那么在CUDA内核中如何用输出调试呢？在内核函数中加“printf”吗？读者可以试试。当然，结果肯定是不奏效。正确的方法就是加一个数组作为调试输出。比如这样：

#define DEBUG// Kernel 函数__global__ void vectorAdd(const float* a, const float* b, float* c, float* info) {    int idx = threadIdx.x;    if (idx < N) {        c[idx] = a[idx] + b[idx];    }#ifdef DEBUG    if (idx == 0) {       info[0] = a[idx];       info[1] = b[idx];       info[2] = c[idx];    }#endif}

在main中添加

float info[N];float * d_info;cudaMalloc(&d_info, sizeof(float) * N);// 执行Kernel...vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, d_info);// ...#ifdef DEBUG    cudaMemcpy(info, d_info, sizeof(float) * N, cudaMemcpyDeviceToHost);    for(int i = 0; i < N; ++i) {       printf("info[%d] = %f\n", i, info[i]);    }#endif

这样我们就获取到第0个线程执行过程中的变量了，通过定义DEBUG宏可以方便的开启或关闭调试信息。这就是Kernel函数输出调试的基本思路。