Staskaer

Hi, this is Staskaer

0%

CUDA高性能并行计算(4)--CUDA错误检测与程序计时

Posted on Edited on In Valine:

CUDA错误检测与程序计时

错误检测

程序中的错误

前面写的程序其实是存在重大问题的，只是我们刻意选择了一些数据来规避了错误的触发。

其引起错误的关键是这里

1
2
3
4
5
6
7
8

#define N 64
#define TPB 32

...

distanceKernel << <N / TPB, TPB >> > (d_out, d_in, ref);

核函数会创建N/TPB个线程块，每个线程块中有TPB个线程。此处每个线程块内的线程数目是确定的(因为一般都是2的整数次幂)。

而如果N不是2的整数次幂，比如N=63，此时计算出来的N/TPB = 63/32 = 1，也就意味着此时只分配一个有32个线程组成的线程块，是不足以完成整个计算的。

所以我们需要N/TPB向上取整才对，将这个公式改成就完成了向上取整的任务。但是此时又会引入新的问题。

比如此时N取65，那么就会分配64/32+1=3个线程块，而每个线程块内又有32个线程，这样就分配了96个线程用于计算，而每个线程会访问其索引对应的数组位置，但是我们只分配了N=65个浮点数的内存，当后续的线程访问对应的数组下标时，其实引发了数组越界的问题，会导致段错误或者是完成计算但是结果匪夷所思。

此处我们当然可以在核函数内使用如下形式的if来结束掉不必要的线程

1
2
3
4

if (i >= N)
	return;

核函数不可以有返回值，但是能用return，return后不能带参数

检测CUDA运行时的宏函数

上面的错误是比较明显的，所以我们可以直接纠正。但是还有一部分错误我们可能没有办法快速定位(因为CUDA不会主动抛出错误，造成DEBUG上困难)。

此处只讨论运行时错误，因为编译时错误肯定会被检测出来。

但是，虽然我们没法直接定位错误，但是所有以cuda开头的函数其实都是有返回值的，其返回值的类型为cudaError_t，代表了一种错误信息，只有当返回值为cudaSuccess时，才成功调用了函数。

可以创建一个用于在运行时检测CUDA函数错误的宏函数

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

#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)

在调用一个cudafunc时，可以CHECK(cudafunc)这样来调用，当函数出现错误时，就能被宏函数捕捉到。

检查核函数中的错误

上述方法只能用于检测以cuda函数开头的函数是否被正确执行。但是由于核函数不存在返回值，所以核函数需要用别的方法进行检查。

下面给出两个自带的函数：

  • cudaGetLastError()：捕捉上一个错误。
  • cudaDeviceSynchronize()：同步主机与设备。

同步主机与设备

由于核函数是异步调用的，即主机调用完核函数后会直接执行下一条指令，而不会等核函数执行完。

不过由于同步操作相当耗时，所以一般不在循环内层调用。只会在必要的时候调用。

在核函数调用后加上下面这两句就能实现检测核函数错误。

1
2
3
4

CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());

检查内存错误

CUDA提供了CUDA-MEMCHECK工具集，可以在命令行使用cuda-memcheck来使用。

当使用nvcc编译程序后，执行cuda-memcheck ./a.out就可以看到内存错误检查后的结果。

CUDA事件计时

对一个函数或程序进行计时是很常见的事情，毕竟，我们用GPU改写了一个程序后，肯定想定量的计算性能的提升。

在CUDA中，提供了基于CUDA事件的计时方法。在CUDA程序中，CPU端将命令写到命令缓冲区中，GPU端会依次读取命令缓冲区并执行任务，一般的程序中，GPU需要给CPU汇报任务的工作进度。命令缓冲区和“同步信息位置”都位于锁页主机内存上，所以CPU和GPU都能同时读写这些数据。在这个“同步信息位置”的内存上，会设置一个单调递增的整数值（“进度值”），GPU完成一条命令操作之后，就会更新递增这个“同步值”，CPU读取这个“同步值”，就可以知道GPU的工作进度。

CUDA事件可以反映这种硬件能力，cudaEventRecord()函数的作用是将一个命令加入队列，使一个新的同步值写入共享同步位置中，cudaEventQuery()和cudaEventSynchronize()则分别用于检查和等待这个事件的同步值。

其实原理很简单：先用cudaEventRecord()记录下一个起始事件的时刻，然后再记录下结束事件的时刻，然后两者做差就能得到时间了。但是需要注意，事件记录这件事本身也是异步的，因此需要在结束事件的位置进行同步，让CPU等待这个事件被记录。

下面是抽象出来的记录时间的代码片段

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

cudaEvent_t start, stop; // 创建事件
cudaEventCreate(&start); // 初始化事件
cudaEventCreate(&stop); // 初始化事件
cudaEventRecord(start); // 事件开始，用于计时
cudaEventQuery(start); // 刷新队列

/*下面是需要被计时的代码片段*/

/*上面是需要被计时的代码片段*/

cudaEventRecord(stop); // 结束事件记录
cudaEventSynchronize(stop); // 让主机等待事件stop被记录完毕
float times = 0;
cudaEventElapsedTime(&times, start, stop); // 获取起始事件和结束事件的差值，单位ms
printf("time = %f\n", times); 
cudaEventDestroy(start); // 销毁事件
cudaEventDestroy(stop);

如果不在结束事件出进行同步，那么是不会得到正确的时间的。

可以把同步处代码注释掉，然后将获取时间的部分改成CHECK(cudaEventElapsedTime(&times, start, stop));

这样就能看到device not ready的错误了，且输出的时间为0，也是不正确的。

分析代码性能工具

CUDA工具箱中存在一个叫nvprof的程序，可以执行性能分析。但是注意要将cuda下的extras\CUPTI\lib64添加到环境变量，同时此程序必须运行在管理员模式下。

输入nvprof {程序名}即可执行代码性能分析。

影响GPU加速的关键因素

  • 由于调用设备GPU计算需要将数据先传递到GPU中，这会导致额外的耗时，所以小规模计算是不划算的。另外，GPU与CPU的连接桥PCIE总线速度是远远低于GPU显存带宽的，因此不能将数据来回传递。
  • GPU由于没有很多空间去布置高速缓存，因此在GPU中，对内存进行读写是相当耗时的。这导致开销比较低的计算(比如简单加法)时受到内存读写瓶颈的限制，导致性能下降。
  • GPU一个SM能驻留1024个线程，而一个GPU有几个到几十个SM，因此能开到几万至几十万个线程，当我们的并行规模远小于这个数目时，性能就不能达到最优。

因此我们想要获得很好的加速性能就需要满足以下几点：

  • 数据传输比例小，减少主机与设备的数据传输。
  • 提高核函数的运算强度。
  • 增大核函数的并行规模。

总结

  • 每个cuda函数都使用宏函数来检测错误
  • 使用CHECK(cudaGetLastError())来检测核函数错误
  • cuda-memcheck工具来检测内存错误
  • 使用cuda事件用于计时
  • nvprof来进行代码分析