原文地址
https://developer.nvidia.com/blog/how-query-device-properties-and-handle-errors-cuda-cc/
在cuda c/c++系列的第三篇博文中，我们讨论了各种支持cuda的gpu的各种特性，如何从cuda c/c++程序中查询设备属性，以及如何处理错误
查询设备属性
在我们上一篇关于性能指标的文章中，我们讨论了如何计算gpu的理论峰值贷款。这种计算方法是用了我们从产品文献中获得的gpu的内存时钟频率和总线接口宽度。一下cuda c++代码演示了一种更通用的方法，通过查询链接的设备（单个或多个）以获取所需信息，来计算理论峰值带宽。
#include <stdio.h> int main(void) { 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); } return 0; }
此代码调用函数cudaGetDeviceCount()，该函数在参数nDevices中返回连接在该系统上的支持cuda的设备的数量。然后在一个循环中，我们计算每个设备的理论峰值带宽。循环体使用cudaGetDeviceProperties()来填充比阿娘prop的字段，它是结构体cudaDeviceProp的一个实例。该程序只使用了cudaDeviceProp众多成员变量中的三个：name，memoryClockRate和memoryBusWidth
当我编译，并在具有单个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工具包提供的deviceQuery代码示例已经做到了这一点
计算性能（compute capability）
我们将在本系列的后续文章中讨论cudaDeviceProp类型中包含的许多设备属性，但我想在这里提到两个重要的领域，major和minor（主要和次要？）。这些描述了设备的计算能力，通常以major，minor格式给出，并指示架构生成。Tesla产品线中第一款支持cuda的设备是Tesla C870，他的计算能力是1.0.第一个具有双精度能力的gpuTesla C1060，具有1.3的计算性能。Fermi架构的gpu，比如上面使用的C2050，计算能力为2.X，kepler架构的gpu，计算性能为3.X。许多与执行配置相关的限制会因计算能力而异，如下表所示。
在本系列的第一篇文章中，我们提到将线程分组到block中模仿了线程处理器在gpu上的分组方式。这组线程处理器成为流多处理器，在上表中用SM表示。cuda执行模型在多处理器上发布线程块，一旦发布，他们就不会迁移到其他SM上。
根据可用资源（片上寄存器和共享内存）和表最后一行中显示的限制，多线程块可以同时驻留在（一个）SM上。此表中对线程和线程块的限制与计算能力相关，而不仅仅是特定设备：具有相同计算能力的所有设备都具有相同的限制。还有一些其他的特性（例如每个设备的多处理器数量）取决于特定设备而不是计算能力。所有的这些特性，无论是由特定设备还是其他计算能力定义，都可以使用cudaDeviceProp类型获取。
你可以使用nvcc编译器选项-arch=sm_xx为特定计算能力的设备生成代码，其中xx表示计算能力（不带小数点）。要查看特定版本的nvcc可以为其生成代码的compute capability列表，以及其他与cuda相关的编译器选项，请使用命令nvcc --help并参考-arch条目。
当你为内核指定执行配置时，请记住（并在运行时查询）上表中的限制。这对于第二个执行配置参数尤为重要：block中的thread个数。如果你为每个block指定的线程太少，则每个多处理器的线程块上限将限制可实现的并行量。如果你为每个block指定了太多的thread，那么我们将在下面讲解。
处理cuda错误
所有的cuda c runtime api函数都有一个返回值，可用于检查执行期间发生的错误。在上面查看gpu属性的例子中，我们可以检查cudaGetDeviceCount()是否成功执行，像这样：
cudaError_t err = cudaGetDeviceCount(&nDevices); if (err != cudaSuccess) printf("%s\n", cudaGetErrorString(err));
我们检查cudaGetDeviceCount()返回值是不是cudaSuccess。如果有错误，那我们调用函数cudaGetErrorString()来获取描述错误的字符串。
处理内核错误有点复杂，因为内核相对于主机异步执行。为了帮助检查内核执行以及其他异步操作的错误，cuda运行时会维护一个错误变量，每次发生错误时都会覆盖该变量。函数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));
这个代码把同步错误和异步错误都检查了。无效的执行配置参数，例如每个block的线程过多，反映在cudaGetLatsError()返回的errSync值中（同步错误）。异步错误是控制权返回host端后，device端发生的错误，例如越界内存访问，需要一个同步机制，比如cudaDeviceSynchronize()，他会阻塞host线程，知道所有先前发出的命令都完成。任何异步的错误都由cudaDeviceSynchronize()返回。我们还可以通过修改最后一条语句调用cudaGetLastError()。我们还可以通过修改最后一条语句调用cudaGetLastError()来检查异步错误并重置runtime错误状态
if (errAsync != cudaSuccess) printf("Async kernel error: %s\n", cudaGetErrorString(cudaGetLastError());
设备同步的代价很高，因为它会导致整个设备等待，从而破坏程序中此时任何的潜在并发性。所以请小心使用它。通常，我使用预处理器宏仅在我的代码的debug版本中和插入异步错误检查，而不是在release版本中。
总结
现在你知道如何查询cuda设备属性并处理cuda c程序中的错误。这些是编写健壮的cuda应用程序非常重要的概念。
在本系列的前三篇文章中，我们涵盖了编写cuda c程序的一些基础知识，重点介绍了基本的编程模型和编写简单例子的语法。我们在第二篇文章中讨论了计时代码和性能指标，但我们还没有使用这些工具来优化我们的代码。在下篇文章中，我们将研究优化host端和device端之间的数据传输。