阅读数 6,208

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.000000
c[1] = 3.000000
c[2] = 5.000000
c[3] = 7.000000
c[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 == cudaSuccess
cuda returned code == cudaSuccess
c[0] = 1.000000
c[1] = 3.000000
c[2] = 5.000000
c[3] = 7.000000
c[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函数输出调试的基本思路。

TomHeaven的博客

个人IT技术文章

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

1. CUDA入门——vectorAdd

2. 检查运行时错误

3. 必须注意的问题

4. 逻辑调试