| 异步错误报告：当 printf 无法执行时 | 您所在的位置：网站首页 › flash自动停用 › 异步错误报告：当 printf 无法执行时 |
异步错误报告：当 printf 无法执行时
|
有些编程情况要求异步报告“软”错误。虽然printf可以是一个有用的工具，它可以增加寄存器的使用并影响性能。在这篇文章中，我们提出了一个替代方案，包括一个头库，用于在 GPU 上生成自定义错误和警告消息，而无需对内核进行硬停止。 错误报告往往会影响性能。虽然有些错误必须立即处理，但其他错误可以以警告和软错误的形式出现，稍后可以报告和解决。 对于 GPU 来说，这通常是一个很好的策略，因为不同的内核可以在不同的流上启动。如果出现任何错误，您可以异步查询并解决。 例如，在一些物理模拟代码中，可能存在物理上不可行的数值解决方案，例如负质量。您可能需要改变运行参数以获得可行的解决方案，如设置较小的时间步长。 虽然有时可以创建误差估计器，但在极少数情况下，估计器仍可能失败。 在 GPU 的上下文中， CUDA 用户可能倾向于检查偶尔出现的不可行解决方案，然后使用printf以在屏幕上提醒最终用户。此解决方案有几个潜在的缺点： 如果有几个流异步运行，那么输出可能会变得复杂。发生错误时，必须重新启动某些操作。额外的调试反馈没有那么有用。 在寄存器受限内核的情况下，您可能希望增加占用率，使用printf不分青红皂白地可能会迫使编译器将许多寄存器专用于代码的一个分支，而该分支只能在偶尔触发。 您对何时查询错误和何时报告错误的控制较少。我们在一些情况下遇到了这个错误报告问题，并使用atomicCAS以帮助高性能地检测软错误。然后，我们使用固定系统内存来协调主机端查询和软错误报告。 我们在一个仅限标头的小型库中提供了此解决方案，该库提供了基础设施，以便您可以将此异步错误报告解决方案放入代码中。模板的使用使您能够自定义错误报告有效负载，而我们的库处理创建和映射系统固定和设备端错误信息。 此外，我们的库使用 lambda 函数为您在 GPU 内核中触发错误提供了足够的灵活性。它为查询和报告错误提供了灵活的帮助功能。 
图 1 。不可预测的表面
工作量示例
作为演示，我们使用下面的内核来模拟一个工作负载，该工作负载可以产生平滑变化的结果，但在罕见的边缘情况下除外。内核生成一些介于 0 和 7210 之间的伪随机整数。然后，它将该整数传递到一个函数中，该函数在 100 左右达到峰值。在极少数情况下，此内核会生成 1e6 。在剩下的时间里，这些值都小于 1 . 0 。 #include #include #include #include __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和DeviceStatus，为错误有效负载以及设备端和固定状态分配分配和销毁系统固定内存。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 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(blockIdx.x), .thread = static_cast(threadIdx.x), .idx = idx, .val = val }; }); } out[idx] = val; } }类型的错误负载RandomSpikeError在用户提供的 lambda 函数中直接在设备上设置 in pined 内存 函数 report _ first _ error 的定义如下： template inline __device__ void report_first_error( MappedErrorType & error_dat, FunctionType func){ if(atomicCAS(reinterpret_cast(error_dat.deviceData.device_status), static_cast(ATOMIC_NO_ERROR), static_cast(ATOMIC_ERROR_REPORTED)) == static_cast(ATOMIC_NO_ERROR) ) { func(*error_dat.deviceData.host_data); __threadfence_system(); error_dat.synchronizeStatus(); } }正如您所看到的，使用atomicCAS其中首先执行设备侧状态。如果成功，则执行用户提供的 lambda 函数并将其写入固定内存。之后，使用系统范围的线程围栏来保证在将主机固定状态与设备侧状态同步之前已执行该功能。 然后，主机可以使用查询并报告错误MappedErrorType直接地 int reportError( MappedErrorType & error_dat) { int retval = NO_ERROR; if (error_dat.checkErrorReported()) { auto & error = error_dat.get(); retval = error.code; std::cerr |
【本文地址】