Почему существует неопределенное поведение внешней памяти в 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++) немедленно вызвать неисправность (например, сегментный вина и т.д.)

На графическом процессоре нет аппаратного механизма, который отслеживает все распределения до уровня байта. Существует общий аппаратный механизм, который позволяет отслеживать выделенные страницы памяти и обнаруживать, находится ли доступ за пределами допустимой страницы, но степень детализации не ниже уровня байта или элемента (и я не думаю, что это так центральный процессор хоста).

Архитектурные, более новые графические процессоры имеют лучшие механизмы отслеживания доступа hw. Кроме того, cuda-memcheck может выполнять более тесное отслеживание доступа, за счет значительного снижения производительности, поскольку он выполняет частично отслеживание на основе SW и частично отслеживание на основе HW, возможно, несколько похожее на такой инструмент, как valgrind на хосте.

Таким образом, хотя у вас есть ожидание, что любое отклонение от выделенного пространства сразу вызовет ошибку, сам GPU HW не поддерживает это (и CPU HW от AFAIK тоже не работает, по крайней мере, в современной виртуальной памяти с запросом по запросу ОС). С программным вмешательством (т. Е. Cuda-memcheck) ситуация улучшается в целом, но все равно будет меняться в зависимости от генерации HW.