有些编程情况要求异步报告“软”错误。虽然printf可以是一个有用的工具，它可以增加寄存器的使用并影响性能。在这篇文章中，我们提出了一个替代方案，包括一个头库，用于在 GPU 上生成自定义错误和警告消息，而无需对内核进行硬停止。
错误报告往往会影响性能。虽然有些错误必须立即处理，但其他错误可以以警告和软错误的形式出现，稍后可以报告和解决。
对于 GPU 来说，这通常是一个很好的策略，因为不同的内核可以在不同的流上启动。如果出现任何错误，您可以异步查询并解决。
例如，在一些物理模拟代码中，可能存在物理上不可行的数值解决方案，例如负质量。您可能需要改变运行参数以获得可行的解决方案，如设置较小的时间步长。
虽然有时可以创建误差估计器，但在极少数情况下，估计器仍可能失败。
在 GPU 的上下文中， CUDA 用户可能倾向于检查偶尔出现的不可行解决方案，然后使用printf以在屏幕上提醒最终用户。此解决方案有几个潜在的缺点：
- 如果有几个流异步运行，那么输出可能会变得复杂。发生错误时，必须重新启动某些操作。额外的调试反馈没有那么有用。
- 在寄存器受限内核的情况下，您可能希望增加占用率，使用
printf不分青红皂白地可能会迫使编译器将许多寄存器专用于代码的一个分支，而该分支只能在偶尔触发。 - 您对何时查询错误和何时报告错误的控制较少。
我们在一些情况下遇到了这个错误报告问题，并使用atomicCAS以帮助高性能地检测软错误。然后，我们使用固定系统内存来协调主机端查询和软错误报告。
我们在一个仅限标头的小型库中提供了此解决方案，该库提供了基础设施，以便您可以将此异步错误报告解决方案放入代码中。模板的使用使您能够自定义错误报告有效负载，而我们的库处理创建和映射系统固定和设备端错误信息。
此外，我们的库使用 lambda 函数为您在 GPU 内核中触发错误提供了足够的灵活性。它为查询和报告错误提供了灵活的帮助功能。
工作量示例
作为演示，我们使用下面的内核来模拟一个工作负载，该工作负载可以产生平滑变化的结果，但在罕见的边缘情况下除外。内核生成一些介于 0 和 7210 之间的伪随机整数。然后，它将该整数传递到一个函数中，该函数在 100 左右达到峰值。在极少数情况下，此内核会生成 1e6 。在剩下的时间里，这些值都小于 1 . 0 。
#include <iostream>
#include <stdio.h>
#include <assert.h>
#include <cuda.h>
__global__ void randomSpikeKernel(float* out, int sz)
// Generate a pseudo-random number
// Pass it into f(x) = 1/(x-100+1e-6)
// Write result to out
{
for (int idx = threadIdx.x + blockIdx.x * blockDim.x;
idx < sz;
idx += blockDim.x * gridDim.x)
{
const int A = 187;
const int M = 7211;
int ival = ((idx + A) * A) % M;
ival = (ival*A) % M;
ival = (ival*A) % M;
float val = 1.f/(ival-100+1e-6);
//assert(val < 10000);
out[idx] = val;
}
}
我们对的呼叫进行了评论assert，可在 GPU 或 CPU 上调用的函数，该函数会立即停止执行并返回错误。这是一种无法恢复的错误解决方案。
在许多情况下，最好让内核运行并稍后报告软错误。如果发生任何错误，您可能有兴趣得到通知，但不想停止工作。你的第一反应可能是添加一个 printf 语句，如下所示：
__global__ void randomSpikeKernelwError(float* out, int sz)
// Generate a pseudo-random number
// Pass it into f(x) = 1/(x-100+1e-6)
// Write result to out
// In the case of a large value (>1e5) print and error, but continue
{
for (int idx = threadIdx.x + blockIdx.x * blockDim.x;
idx < sz;
idx += blockDim.x * gridDim.x)
{
const int A = 187;
const int M = 7211;
int ival = ((idx + A) * A) % M;
ival = (ival*A) % M;
ival = (ival*A) % M;
float val = 1.f/(ival-100+1e-6);
if (val >= 10000) {
printf("val (%f) out of range for idx = %d\n",
val, idx);
}
out[idx] = val;
}
}
这通常是一个可以接受的解决方案。但对于占用率受寄存器使用限制的内核来说，这可能会产生不希望的后果。即使printf语句很少被执行，编译器必须分配寄存器以防万一。
寄存器是仅在线程中使用的快速内存。寄存器中的数据可以低延迟读取和写入，但一个线程中的寄存器对任何其他线程都不可见。您可以通过添加-Xptxas=-v到编译行，或者使用 NVIDIA Nsight Compute 来评测内核。
按照如下方式编译以前的代码：
nvcc -c -arch=sm_80 -Xptxas=-v kernel.cu
在编译过程中，您会看到以下消息：
ptxas info : 36 bytes gmem
ptxas info : Compiling entry function '_Z17randomSpikeKernelPfi' for 'sm_80'
ptxas info : Function properties for _Z17randomSpikeKernelPfi
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 14 registers, 364 bytes cmem[0], 8 bytes cmem[2]
ptxas info : Compiling entry function '_Z23randomSpikeKernelwErrorPfi' for 'sm_80'
ptxas info : Function properties for _Z23randomSpikeKernelwErrorPfi
16 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 24 registers, 364 bytes cmem[0], 8 bytes cmem[2]
第一个内核没有错误报告，使用了 14 个寄存器。第二个内核添加了printf语句，寄存器计数跳到 24 。
由于每个 SM 都有固定的内核寄存器空间，内核对寄存器的更高要求可能会限制每个 SM 上可以活动的线程块的数量。这可能会导致暴露的延迟和较差的性能。此示例内核仅用于演示，不太可能出现寄存器压力问题。
在 HPC 代码中，与高寄存器数作斗争是很常见的。在编译过程中，寄存器计数通常会受到人为限制，但这可能会对性能产生其他负面影响。我们将在下一节对此进行讨论。
此外，通过依赖控制台流来报告信息，您放弃了对何时查询和报告潜在软错误的一些控制。在报告同时运行的几个内核或设备功能的软错误时，使用共享控制台流可能会使输出复杂化。
小的改进
如果寄存器压力降低了内核的性能，因为printf消耗了保留但通常未使用的额外寄存器，一个潜在的解决方案是告诉编译器通过设置-maxregcount编译标志或使用__launch_bounds__在代码中。
这会限制寄存器的数量，从而溢出多余的寄存器。只有在极少数情况下，你才会获得表演上的成功printf发生。这是缓解登记册压力问题的一个重要提示，但它们可能是一个生硬的工具，并干扰其他减少登记册的工作。
拟议的替代方案：比较和交换
报告此类错误的更好方法是使用atomicCAS起到异步屏障的作用，以检测软错误的第一个实例。
CASatomicCAS代表比较和交换，也称为比较交换.
atomicCAS获取一个内存位置、一个比较值和一个新值，并且只有当内存位置与比较值匹配时才将该值写入内存位置。如果从存储器读取的值等于所提供的比较值，atomicCAS将新值写入内存位置。否则，它将保持值不变。在任何一种情况下，它都会返回最初从内存位置读取的值。
最重要的是，如果线程之间存在争用，那么一次只有一个线程进行完全读取、比较和交换。剩下的线程从内存中读取更改后的值，然后跳过写入。 CUDA 支持atomicCAS用于 32 位有符号整数和 16 位、 32 位或 64 位无符号整数。
在这个解决方案中，您使用atomicCAS以确保只有一个线程可以写入错误消息。在清除错误之前，不会报告以后检测到的所有错误。这避免了写入错误消息的不同线程之间的竞争条件，并与本机 CUDA 错误的行为相匹配。
当检测到错误时，应用程序通常必须记录一些额外的数据——行号、错误代码等等。在本例中，您可以写入这些额外的数据，称为错误“有效负载”zero copy，作为从 GPU 内核到系统固定的 CPU 内存的直接写入。因为软错误的有效负载通常很小，所以可以跳过显式内存拷贝直接写入有效负载。
您还可以在系统固定内存中跟踪此错误的状态。这使 CPU 主机知道在 GPU 上生成的错误。使用__threadfence_system以提供系统范围的屏障，以确保在状态标志改变之前有效载荷被完全写入。这使主机能够异步查询状态。当主机看到状态发生变化时，可以确保错误负载包含适当的数据。
库
由于解决方案的设置和初始化可能有点麻烦，我们提供了templated header-only library这简化了这个过程，并使您能够指定自定义的错误有效载荷。
我们引入了两种基本模板类型，PinnedMemory<ErrorType>和DeviceStatus<T>，为错误有效负载以及设备端和固定状态分配分配和销毁系统固定内存。DeviceStatus还有一个仅限主机的状态 getter ，使您能够使用查询固定状态cuda::atomics.
与您交互的主要类是MappedErrorType，使用PinnedMemory和DeviceStatus类，以便轻松地协调状态和有效负载组件。MappedErrorType处理底层类型的初始化、异步查询错误、异步查询有效负载、清除错误以及同步设备端和主机固定状态。
以下代码示例显示类型为的错误RandomSpikeError可以使用结构进行记录RandomSpikeError.
struct RandomSpikeError
{
int code;
int line;
int filenum;
int block;
int thread;
// payload information
int idx;
float val;
};
__global__ void randomSpikeKernelFinal(float* out, int sz, MappedErrorType<RandomSpikeError> device_error_data)
// This kernel generates a pseudo-random number
// then puts it into 1/num-100+1e-6. That curve is
// sharply peaked at num=100 where the value is 1e6.
// In the case of a large value, you want to report an
// error without stopping the kernel.
{
for (int idx = threadIdx.x + blockIdx.x * blockDim.x;
idx < sz;
idx += blockDim.x * gridDim.x)
{
const int A = 187;
const int M = 7211;
int ival = ((idx + A) * A) % M;
ival = (ival*A) % M;
ival = (ival*A) % M;
float val = 1.f/(ival-100+1e-6);
if (val >= 10000) {
report_first_error(device_error_data, [&] (auto &error){
error = RandomSpikeError {
.code = LARGE_VALUE_ERROR,
.line = __LINE__,
.filenum = 0,
.block = static_cast<int>(blockIdx.x),
.thread = static_cast<int>(threadIdx.x),
.idx = idx,
.val = val
};
});
}
out[idx] = val;
}
}
类型的错误负载RandomSpikeError在用户提供的 lambda 函数中直接在设备上设置 in pined 内存
函数 report _ first _ error 的定义如下：
template <typename ErrorType, typename FunctionType>
inline __device__ void report_first_error(
MappedErrorType<ErrorType> & error_dat,
FunctionType func){
if(atomicCAS(reinterpret_cast<int*>(error_dat.deviceData.device_status),
static_cast<int>(ATOMIC_NO_ERROR),
static_cast<int>(ATOMIC_ERROR_REPORTED)) ==
static_cast<int>(ATOMIC_NO_ERROR) ) {
func(*error_dat.deviceData.host_data);
__threadfence_system();
error_dat.synchronizeStatus();
}
}
正如您所看到的，使用atomicCAS其中首先执行设备侧状态。如果成功，则执行用户提供的 lambda 函数并将其写入固定内存。之后，使用系统范围的线程围栏来保证在将主机固定状态与设备侧状态同步之前已执行该功能。
然后，主机可以使用查询并报告错误MappedErrorType<RandomSpikeError>直接地
int reportError( MappedErrorType<RandomSpikeError> & error_dat)
{
int retval = NO_ERROR;
if (error_dat.checkErrorReported()) {
auto & error = error_dat.get();
retval = error.code;
std::cerr << "ERROR " << error.code
<< ", line " << error.line
<< ". block " << error.block
<< ", thread " << error.thread;
if (retval == LARGE_VALUE_ERROR)
std::cerr << ", value = " << error.val;
std::cerr << std::endl;
}
return retval;
}
auto async_err = reportError(mapped_error);
if (async_err != NO_ERROR) std::cout << "ERROR! " << "code: " << async_err << std::endl;
else std::cout << "No error" << std::endl;
由于错误可能异步发生，您的应用程序可能必须正确同步或等待特定事件，以确保内核已经完成。这与本机 CUDA 错误的行为类似。
把它们放在一起
虽然我们的库简化了许多必要的工作，但以下是幕后发生的事情，以便您可以根据需要扩展和调整错误报告。
在执行内核之前，我们初始化MappedErrorType<T>主机端和设备端状态自动初始化的对象 (ATOMIC_NO_ERROR = 0) . 当在内核中检测到错误时，report_first_error使用atomicCAS标记设备端状态 (ATOMIC_ERROR_REPORTED=1) 然后执行用户提供的 lambda 函数以在应用前面提到的线程围栏和主机设备状态同步之前将有效载荷写入系统固定存储器。
线程只能在以下情况下写入错误数据atomicCAS退货ATOMIC_NO_ERROR，这意味着没有其他线程已经记录到错误。除非您将状态重置为ATOMIC_NO_ERROR，没有记录此错误的其他实例。接收的线程ATOMIC_NO_ERROR写入其错误代码和相关数据。
为了清除数据，我们提供了clear(cudaStream_t)将状态设置为的方法ATOMIC_NO_ERROR主机端和设备端状态。
为了检查主机是否出现错误reportError使用checkErrorReported，仅检查主机侧状态是否设置为ATOMIC_ERROR_REPORTED。然后我们打电话get在错误类型的有效负载上 (struct RandomSpikeError) 并读取错误信息。
在内核执行期间检测到错误既不会停止内核，也不会停止主机。与本机 CUDA 错误一样，主机可能会在检测到此内核中的错误之前启动几个内核。
int main(void)
{
…
// Create pinned flags/data and device-side atomic flag for CAS
auto mapped_error =
CASError::MappedErrorType<RandomSpikeError>();
auto mapped_error2 = CASError::MappedErrorType<OtherError>();
…
int async_err; // error query result
// Allocate memory and a stream
float *out, *h_out;
h_out = (float*)malloc(sizeof(float)*MAX_IDX);
cudaMalloc((void**)&out, sizeof(float)*MAX_IDX);
cudaStream_t stream; cudaStreamCreate(&stream);
CASError::checkCuda(
cudaEventCreate(&finishedRandomSpikeKernel) );
// Launch the kernel. This launch causes a
// LARGE_VALUE_ERROR
randomSpikeKernel<<<100,32,0,stream>>>(out, MAX_IDX);
randomSpikeKernelFinal<<<100,32,0,stream>>>(out, MAX_IDX,
mapped_error);
CASError::checkCuda(
cudaEventRecord(finishedRandomSpikeKernel, stream) );
// Check the error message from err_data
async_err = reportError(mapped_error);
if (async_err != NO_ERROR) std::cout << "ERROR! " << "code: "
<< async_err << std::endl;
else std::cout << "No error" << std::endl;
// Launch another kernel
otherKernel<<<100,32,0,stream>>>(out, MAX_IDX,
mapped_error2);
…
async_err = reportError(mapped_error2, stream);
if (async_err != NO_ERROR)
std::cout << "ERROR! " << "code: " << async_err <<
std::endl;
else std::cout << "No error" << std::endl;
std::cout << "Launch memcpy" << std::endl;
cudaMemcpyAsync(h_out, out, sizeof(float)*MAX_IDX,
cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);
async_err = reportError(mapped_error);
if (async_err != NO_ERROR) std::cout << "ERROR! " << "code: "
<< async_err << std::endl;
else std::cout << "No error" << std::endl;
mapped_error.clear(stream);
async_err = reportError(mapped_error2, stream);
if (async_err != NO_ERROR) std::cout << "ERROR! " << "code: "
<< async_err << std::endl;
else std::cout << "No error" << std::endl;
int final_err = reportError(mapped_error);
if (final_err != NO_ERROR) std::cout << "ERROR! " << "code: "
<< final_err << std::endl;
else std::cout << "No error" << std::endl;
// Free memory, stream
cudaFree(out);
free(h_out);
cudaStreamDestroy(stream);
return 0;
}
在测试中，您首先启动生成错误的内核。然后检查主机线程上的错误。在此检查之前，您没有同步主机和设备。当同步影响性能时，您可能需要在同步之前排队等待更多的 GPU 工作。
您正在检查此示例代码中的错误，以演示异步错误报告。如果您不介意性能命中，并且希望按顺序报告错误，请添加cudaStreamSynchronize在呼叫之前呼叫reportError.
检查错误后，启动另一个内核，otherKernel，然后再次检查错误。使用将生成的数据复制回主机cudaMemcpyAsync。同步流以确保主机上的数据是正确的，并再次检查是否有错误。现在，您一定会发现自己的错误。
接下来，清除错误，并检查第二种类型的错误，也保证会被捕获。最后，为了显示错误已被清除，请最后一次检查错误。
当编译并执行此代码时，您可能会看到以下输出：
No error No error Launch memcpy ERROR 2, line 144. block 92, thread 20, value = 1e+06 ERROR! code: 2 ERROR 3, line 171, file /tmp/devblog/main.cu. block 25, thread 8 ERROR! code: 3 No error
错误是在的 GPU 执行期间生成的randomSpikeKernelFinal，但由于您没有在调用之间同步主机和设备，因此主机线程能够对内核和memcpy立即执行，而无需等待第一个 CUDA 内核完成。直到流同步之后， CPU 才检测到并报告错误。
由于您有两种不同类型的错误，因此可以分别捕获和清除每一种错误。否则，您只报告您观察到的每种类型的第一个错误。
收益
使用编译时-Xptxas=-v，您可以看到此输出（添加了突出显示）：
ptxas info : Compiling entry function '_Z17randomSpikeKernelPfi' for 'sm_70' ptxas info : Function properties for _Z17randomSpikeKernelPfi 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 14 个寄存器, 364 bytes cmem[0], 8 bytes cmem[2] ptxas info : Compiling entry function '_Z23randomSpikeKernelwErrorPfi' for 'sm_70' ptxas info : Function properties for _Z23randomSpikeKernelwErrorPfi 16 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 24 个寄存器, 364 bytes cmem[0], 8 bytes cmem[2] ptxas info : Compiling entry function '_Z22randomSpikeKernelFinalPfi' for 'sm_70' ptxas info : Function properties for _Z22randomSpikeKernelFinalPfi 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 18 个寄存器, 364 bytes cmem[0], 8 bytes cmem[2]
第一个内核没有错误报告。第二个报告错误使用printf第三个也是最后一个内核使用了前面描述的新方法。下表显示了早期输出的寄存器计数。
| 内核 | 错误报告方法 | 寄存器 |
randomSpikeKernel |
没有一个 | 14 |
randomSpikeKernelwError |
printf |
24 |
randomSpikeKernelFinal |
atomicCAS |
18 |
检查错误并使用报告atomicCAS与使用时的 10 个新寄存器相比，增加了 4 个寄存器printf在这种小情况下，寄存器计数不太可能影响性能。对于寄存器使用是一个性能问题的内核，这种新的错误报告可以产生显著的影响。
一个真实世界的例子
下面是一个例子，说明了这种新方法可以在实际代码中产生的差异。
我们在 hpMusic 中对这个库进行了野外测试， hpMusic 是一个高阶计算流体动力学模拟代码示例。在基线代码中，一个内核printf用于报告罕见软错误的语句使用了 248 个寄存器。通过评论printf（无错误报告），ncu报告了内核的 148 个寄存器。
最后，通过访问我们的图书馆，ncu报告编译的内核也使用了 150 个寄存器。由于这些内核受寄存器约束，因此通过避免printf在这个性能关键的内核中，它对运行时产生了重大影响。
| 内核变化 | 寄存器 | 占用率 | 内核运行时间（毫秒） |
| 输出函数 | 248 | 11 . 97% | 293 . 5 |
| 无报告 | 148 | 17 . 83% | 243 . 6 |
| 打印备选方案 | 150 | 17 . 80% | 239 . 5 |
| printf 和 launch _ bounds | 168 | 17 . 25% | 299 |
虽然 hpMusic 开发人员是领域专家，他们也编写 GPU 应用程序，但他们对使用printf在寄存器约束内核中。
结论
如果您正在报告软错误或其他不常见的内核信息，下载 headers以及这篇文章中的例子，并自己尝试一下。我们总是对反馈感兴趣，所以请发送消息，让我们知道它是如何工作的！
