CUDA负数组索引测试

Question

因为在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之前的内存可能包含其他重要数据，但是当我对其进行测试时，其中不包含任何重要数据，因此没有错误或测试失败。

FYI：我正在使用pycuda并正在对我的代码进行单元测试。

由于上述原因，我试图避免产生无提示的不可预测的错误。当然，在现实世界中，示例-1是经过计算的，我已将代码简化为仅要解决的问题。

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

Answer 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之前的宏。