CUDA 是否有可能进行不会导致错误的越界访问？

Question

所以我在 CUDA 中遇到了一个非常奇怪的行为。

我在 GPU 上有多个阵列，假设一个是 float a[100];

我有一个如下所示的函数：

float access(int i) {
    if (i >= 100) printf("i is out of bound\n");
    if (i+blockDim.x >= 100) printf("i+blockDim.x is out of bound\n");
    return a[i] + a[i + blockDim.x];
}

在执行过程中，第二条语句被多次打印。然而，即使我到处都使用 cudaGetLastError()，函数仍然没有捕获到错误！

程序永远不会崩溃。 nvprof 也没有显示任何错误。

我最好的猜测如下：

float a[100]; // byte 0, 4, 8, ..., 396
...
float z[100]; // byte 400, 404, 408, ...

这意味着当我访问 a[100] 时，我实际上是在访问 z[0]，所以它不会崩溃。

还有其他可能的原因吗？

编辑：在我的实际程序中越界访问很远，[1000]等cuda-memcheck可以检测到问题

Answer 1

稍微越界访问数组通常不会导致明显的运行时错误（尽管您的代码仍然可能计算不正确）。 GPU 运行时机制不会测试字节级别的访问有效性。

如果访问越界，最终会遇到运行时错误。

根据我对 CPU 主机代码的经验，同样的说法也是正确的（试试吧）。

正因为如此，在CPU端你可以使用像valgrind这样的工具来捕捉这样的错误，而在GPU端你可以使用像cuda-memcheck这样的工具来捕获这样的错误捕获此类错误。

CUDA 是否有可能进行不会导致错误的越界访问？

CUDA is it possible to have out-of-bound access that results in no error?

c++

cuda

undefined-behavior

indexoutofboundsexception