Accelerated Ray Tracing （四）

https://developer.nvidia.com/blog/how-query-device-properties-and-handle-errors-cuda-cc/

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

Compute Capability

cudaDeviceProp类型中包含的许多设备属性，特别是major和minor。这些描述了设备的计算能力，这通常在major.minor中给出。并指示第几代架构。

将thread分组为线程block的方式模仿了在GPU上分组线程处理器的方式。这组线程处理器称为流多处理器streaming multiprocessor，在上表中用SM表示。CUDA执行模型在多处理器上执行线程block，一旦执行它们就不会迁移到其他SMs上。受可用资源(on-chip registers和shared memory)和表最后一行所示的限制的限制，多个线程block可以并发驻留在多处理器上。该表中线程和线程块的限制与计算能力相关，而不仅仅是特定的设备:具有相同计算能力的所有设备都有相同的限制。不过，还有其他一些特征，比如每个设备的多处理器数量，这些特征依赖于特定的设备，而不是计算能力。所有这些特征，无论是由特定的设备或其计算能力定义，都可以使用cudaDeviceProp类型获得。

您可以使用nvcc编译器选项-arch=sm_xx为特定的计算能力生成代码，其中xx表示计算能力(没有小数点)。要查看特定版本的nvcc可以为其生成代码的计算能力列表，以及其他与cuda相关的编译器选项，请发出命令nvcc—help并参考-arch条目。

当您为内核指定执行配置时，请记住(并在运行时查询)上表中的限制。这对于第二个执行配置参数(每个线程块的线程数)尤其重要。如果您为每个块指定的线程太少，那么对每个多处理器的线程块的限制将限制可以实现的并行量。如果每个线程块指定了太多线程，那么我们就进入下一节。

Handling CUDA Errors

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

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

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

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);

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

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

Device 同步是昂贵的，因为它会导致整个设备等待，从而破坏程序中此时的任何潜在并发性。所以要小心使用。通常，我只在代码的调试构建中使用预处理器宏来插入异步错误检查，而不是在发布构建中。