CUDA - 如何在CUDA C/C++中查询设备属性和处理错误

原文链接：How to Query Device Properties and Handle Errors in CUDA C/C++

---

在CUDA C/C++系列的第三篇文章中，我们讨论各种支持CUDA的GPU的各种特性、如何在CUDA C/C++程序中查询设备属性，以及如何处理错误。

查询设备属性

在我们上一篇关于性能指标的文章中，我们讨论了如何计算GPU的理论峰值带宽。该计算使用了产品手册中的GPU的内存时钟速率和总线接口宽度。下面的CUDA C++代码演示了一种更通用的方法，通过向连接的（多个）设备查询所需信息来计算理论峰值带宽。

#include <stdio.h> 

int main() {
  int nDevices;

  cudaGetDeviceCount(&nDevices);
  for (int i = 0; i < nDevices; i++) {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, i);
    printf("Device Number: %d\n", i);
    printf("  Device name: %s\n", prop.name);
    printf("  Memory Clock Rate (KHz): %d\n",
           prop.memoryClockRate);
    printf("  Memory Bus Width (bits): %d\n",
           prop.memoryBusWidth);
    printf("  Peak Memory Bandwidth (GB/s): %f\n\n",
           2.0*prop.memoryClockRate*(prop.memoryBusWidth/8)/1.0e6);
  }
}

此代码使用函数cudaGetDeviceCount()，该函数在参数nDevices中返回连接到此系统的具有CUDA功能的设备的数量。然后在循环中计算每个设备的理论峰值带宽。循环的主体使用cudaGetDeviceProperties()来填充变量prop的字段，变量prop是结构体cudaDeviceProp的一个实例。该程序只使用了cudaDeviceProp众多成员中的三个：name、memoryClockRate和memoryBusWidth。

当我（使用CUDA nvcc编译器的任何最新版本，例如4.2或5.0rc）编译并在装有一个NVIDIA Tesla C2050的机器上运行此代码时，我会得到以下结果。

Device Number: 0
  Device name: Tesla C2050
  Memory Clock Rate (KHz): 1500000
  Memory Bus Width (bits): 384
  Peak Memory Bandwidth (GB/s): 144.00

这与我们在上一篇文章中计算的理论峰值带宽值相同。当在我的笔记本电脑上编译并运行相同的代码时，我会得到以下输出。

Device Number: 0
  Device name: NVS 4200M
  Memory Clock Rate (KHz): 800000
  Memory Bus Width (bits): 64
  Peak Memory Bandwidth (GB/s): 12.800000

cudaDeviceProp结构体中还有许多其他字段，用于描述各种类型内存的数量、线程块大小的限制以及GPU的许多其他特性。我们可以扩展上面的代码来打印出所有这样的数据，NVIDIA CUDA Toolkit提供的deviceQuery示例代码已经做了这件事。

计算能力

我们将在本系列后面的文章中讨论 cudaDeviceProp类型中包含的许多设备属性，但我想在这里提一下两个重要的字段，major和minor。这两个字段描述了设备的计算能力，通常以major.minor格式给出，也表示了哪一代架构。Tesla产品线中第一款支持CUDA的设备是Tesla C870，其计算能力为1.0。第一批双精度GPU，如Tesla C1060，计算能力为1.3。Fermi架构的GPU，例如上面使用的Tesla C2050，具有2.x的计算能力，Kepler架构的GPU计算能力为3.x。根据计算能力的不同，许多与执行配置相关的限制也不同，如下表所示。

	Tesla C870	Tesla C1060	Tesla C2050	Tesla K10	Tesla K20
Compute Capability	1.0	1.3	2.0	3.0	3.5
Max Threads per Thread Block	512	512	1024	1024	1024
Max Threads per SM	768	1024	1536	2048	2048
Max Thread Blocks per SM	8	8	8	16	16

在本系列的第一篇文章中，我们提到将线程分组为线程块模仿了GPU上线程处理器的分组方式。一组线程处理器被称为一个streaming multiprocessor，上表中表示为SM。CUDA执行模型在SM上发布线程块，一旦发布，就不会移动到其他SM中。

根据可用资源（片上寄存器和共享内存）和表格最后一行所示的限制，多个线程块可以同时驻留在一个SM上。该表中对线程和线程块的限制与计算能力相关，而不仅仅是与特定设备相关：具有相同计算能力的设备都有相同的限制。然而，还有其他特征，例如每个设备的SM数量，这取决于特定设备，而不是计算能力。所有这些特性，无论是由特定设备还是其计算能力定义的，都可以通过cudaDeviceProp获得。

您可以使用nvcc编译器选项-arch=sm_xx为特定的计算能力生成代码，其中xx表示计算能力（无小数点）。要查看特定版本的nvcc可以为哪些计算能力生成代码，以及其他与CUDA相关的编译器选项，使用命令nvcc --help并查看-arch条目。

当您为内核指定执行配置时，请记住（并在运行时查询）上表中的限制。这对于第二个执行配置参数特别重要：每个线程块的线程数。如果每个块指定的线程太少，那么GPU的并行能力将受到每个SM的最大线程块数量的限制。如果每个线程块指定的线程太多，这就是本文下一段要讲的内容了。

处理CUDA错误

所有CUDA C Runtime API函数都有一个返回值，可用于检查执行过程中发生的错误。在上面的例子中，我们可以检查cudaGetDeviceCount()是否成功完成，如下所示：

cudaError_t err = cudaGetDeviceCount(&nDevices);
  if (err != cudaSuccess) printf("%s\n", cudaGetErrorString(err));

我们检查以确保cudaGetDeviceCount()返回cudaSuccess。如果出现错误，那么我们调用函数cudaGetErrorString()来获取描述错误的字符串。

处理内核错误有点复杂，因为内核相对于主机异步执行。为了帮助检查内核执行以及其他异步操作的错误，CUDA runtime维护一个错误变量，该变量在每次发生错误时都会被覆盖。函数cudaPeekAtLastError()返回此变量的值，函数cudaGetLastError()返回该变量的值并将其重置为cudaSuccess。

我们可以检查本系列第一篇文章中使用的saxpy内核中的错误，如下所示。

saxpy<<<(N+255)/256, 256>>>(N, 2.0, d_x, d_y);
cudaError_t errSync  = cudaGetLastError();
cudaError_t errAsync = cudaDeviceSynchronize();
if (errSync != cudaSuccess) 
  printf("Sync kernel error: %s\n", cudaGetErrorString(errSync));
if (errAsync != cudaSuccess)
  printf("Async kernel error: %s\n", cudaGetErrorString(errAsync));

此代码检查了同步和异步错误。cudaGetLastError()返回的errSync值反映了无效的执行配置参数，例如每个线程块的线程太多。当控制权返回到主机后在设备上发生的异步错误，如越界内存访问，需要一种同步机制，如cudaDeviceSynchronize()，它会阻塞主机线程，直到之前发出的所有命令都完成为止。cudaDeviceSynchronize()会返回任何异步错误。通过修改最后一条语句调用cudaGetLastError()，我们还可以检查异步错误同时重置运行时错误状态。

if (errAsync != cudaSuccess)
  printf("Async kernel error: %s\n", cudaGetErrorString(cudaGetLastError());

设备同步是昂贵的，因为它会导致整个设备处于等待，从而破坏该时刻程序的任何并发能力。所以要小心使用。通常，我只在代码的debug build阶段使用预处理器宏插入异步错误检查，而不会在release build阶段使用。

总结

现在您知道了如何查询CUDA设备属性以及如何处理CUDA C和C++程序中的错误。对于编写健壮的CUDA应用程序来说，这些都是非常重要的概念。

在本系列的前三篇文章中，我们介绍了编写CUDA C/C++程序的一些基础知识，重点介绍了基本的编程模型和编写简单示例的语法。我们在第二篇文章中讨论了计时代码和性能指标，但我们还没有在优化代码时使用这些工具。在下一篇文章中，我们将在文章中优化主机和设备之间的数据传输。