c++ - 为什么 CUDA 中的外部共享内存存在未定义的行为？

Question

我正在尝试了解 CUDA 中共享内存的动态分配。我编写了一个简单的程序，它使用 cudaGetLastError 返回错误。我在启动内核时分配了大小为 100 个整数元素的共享内存。我尝试访问 101-127 个索引内存元素，它没有显示任何错误，但在访问第 130 个元素时，它给出错误“未指定的启动失败”。我认为在访问 101 个元素时应该会出现此错误，因为我只为 100 个元素分配了大小。

#include <cuda.h>
#include <stdio.h>
__global__ void xyz(int offset)
{
 extern __shared__ int array[];
 array[101]=offset; 
 printf("%d\n", array[101]);
}

int main()
{
 dim3 grid(1,1,1);
 dim3 block(100,1,1);
 int offset=50;
 xyz<<<grid,block,sizeof(int)*100>>>(offset);
 cudaDeviceSynchronize();
 cudaError_t err=cudaGetLastError();
 if(err!=cudaSuccess)
 {
    printf("Error is =%s\n",cudaGetErrorString(err));
 }
 return 0;
}

Answer 1

一般来说，我认为访问主机上数组末尾之外的内容(在普通 C/C++ 代码中)不会立即触发错误(例如 seg 错误等)

在 GPU 上，没有硬件机制可以跟踪字节级别的所有分配。有一种通用的硬件机制可以跟踪分配的内存页面，并发现访问是否在有效页面之外，但粒度并没有下降到字节或元素级别(我不认为这是这样的)主机 CPU 之一)。

从架构上来说，较新的 GPU 具有更好的硬件访问跟踪机制。此外，cuda-memcheck 可以对访问进行更严格的跟踪，但代价是显着降低性能，因为它执行部分基于软件的跟踪和部分基于硬件的跟踪，可能有点类似于主机上的 valgrind 等工具。

因此，尽管您似乎期望任何与分配空间的偏差都会立即触发故障，但 GPU 硬件本身并不支持这一点(并且 AFAIK CPU 硬件也不支持，至少在现代按需分页虚拟内存操作系统中) )。通过软件干预(即 cuda-memcheck)，情况总体上有所改善，但仍会因硬件生成而异。