CUDA负数组索引测试答案

【问题标题】：CUDA negative array index testCUDA负数组索引测试
【发布时间】：2020-07-26 07:03:14
【问题描述】：

因为在 C 中可以索引一个负数组位置并超出数组边界，所以这段代码可以编译并“工作”。

__global__ void do_something_bad(int * in_a){
   in_a[-1] = 666; // assign a value to an out of bounds memory location
}

我的假设是上面的代码正在执行以下操作（如果这个假设是错误的，请告诉我）：

GPU memory before:
[0x00 = usually unused memory][0x01= Start of in_a][0x02 = in_a] ....
GPU memory after:
[0x00 = 666][0x01= Start of in_a][0x02 = in_a] ....

总而言之，in_a 数组被设置值之前的内存。 in_a 之前的内存可能包含其他重要数据，但是当我测试它时它不包含任何重要数据，因此不会给我错误或测试失败。

仅供参考：我正在使用 pycuda，并且正在对我的代码进行单元测试。

我正试图避免由于上述原因而产生无声的不可预测的错误。当然，在现实世界中，示例 -1 会被计算出来，并且我已将代码简化为我想要解决的问题。

如何识别此错误并强制我的单元测试可以检测到可检测的问题？

【问题讨论】：

您的示例是未定义的行为。为 UB 设计单元测试是徒劳的
我看不出这与cuda 有什么关系。你会如何用 C/C++ 回答这个问题？ CUDA 主要尝试坚持 C++。（顺便说一句，valgrind - C++ 或 cuda-memcheck - CUDA C++ 应该能够捕捉到这种类型的错误。我不知道这是否符合您对单元测试的定义或您想要的任何内容完成。）
有完全标准的设计模式可以避免您想象的问题，而不会出现重大的性能问题。就像我说的，真的没有办法设计单元测试来检测未定义的行为，因为它是未定义的行为。可能静态分析可能能够检测出常见的错误习语，但即使这样也非常有限，这让我印象深刻，就像在风车上的标题
cuda-memcheck 不提供运行时检查，除非您在运行时使用它。它不能替代一般的运行时边界检查。它在单元测试场景中可能很有用，其中单元测试特别激发了越界条件。 cuda-memcheck 可以告诉您您的代码是否实际索引超出范围。它的实用性将取决于您如何构建单元测试和运行时边界检查（如果有）。在一般情况下，它并不是对运行时边界检查的建议。你也不会在那里使用 valgrind。
如果您有合适的方法（即您知道和喜欢的方法）用于 C++ 检查您想要的类型，我会考虑也只在 CUDA C++ 中实现（或在此处描述）或在代码审查中获得一些反馈）。我已经多次看到像你这样的论点（“当只有少数边缘情况会出现问题时，强制每个线程检查其入站感觉是错误的”），我只是不买它们，除非伴随着科学分析（例如，分析表明此类活动是一个重要的性能问题）。负索引检查的成本应该很低

标签： cuda pycuda

【解决方案1】：

为了避免内核中出现静默错误，如果您不在 MacOS 上，我会使用 assertion。像这样的：

#include <assert.h>

__global__ void do_something_bad(int* in_a){
  int indx;
  indx = 0; // A valid index
  assert(indx >= 0); // Lets the kernel continue
  in_a[indx] = 666;
  indx = -1; // An invalid index
  assert(indx >= 0); // Sends an error to stderr 
  in_a[indx] = 666; // This never gets executed  
}

int main(){
  int *a;
  cudaMalloc((void **)&a, 10*sizeof(int));
  do_something_bad<<<1,1>>>(a);
  cudaDeviceSynchronize();
}

但是，这可能会影响程序的性能。来自编程指南：

断言用于调试目的。它们会影响性能和因此建议在生产代码中禁用它们。他们可以通过定义 NDEBUG 预处理器在编译时禁用包含 assert.h 之前的宏

【讨论】：

作为附加说明，内核代码中的失败断言会破坏 CUDA 上下文，因此在此处给出的示例代码中，失败的断言将是运行时 API 在下一次同步时返回的可检测的 CUDA 错误观点。此错误是一个粘性错误（损坏的 CUDA 上下文），因此： 1. 在此过程中使用 CUDA 将无法进行任何其他操作。 2. 在内核启动后，有问题的数据（例如a[indx]）将永远不会对主机代码可见。