cuda - CUDA负数组索引测试

问题描述

由于在 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 会被计算出来，并且我已经将代码简化为我想要解决的问题。

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

标签： cudapycuda

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

#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();
}

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

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

cuda - CUDA负数组索引测试

问题描述

解决方案

推荐阅读