爲什麼CUDA中存在未定義的外部共享內存行爲？

-3

我想了解在CUDA中共享內存的動態分配。我寫了一個簡單的程序，它使用cudaGetLastError返回錯誤。我在啓動內核時分配了大小爲100個整數元素的共享內存。我試圖訪問101-127索引內存內存元素，並沒有顯示任何錯誤，但訪問第130個元素時，它給出錯誤「未指定的啓動失敗」。我認爲它應該給這個錯誤，同時訪問101元素也因爲我已經分配了大小隻有100元素。爲什麼CUDA中存在未定義的外部共享內存行爲？

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

來源

2014-11-05 krishna

因此，您的問題可以概括爲「爲什麼未定義的行爲未定義在特定的方式」？它可能與分配粒度和硬件設計有關，但爲什麼你要關心？ – talonmies 2014-11-05 05:17:35

是的，我認爲cudagetlasterror在這種情況下應該給出錯誤 – krishna 2014-11-05 05:19:01

當你在每種情況下用'cuda-memcheck'運行測試代碼時會發生什麼？ – 2014-11-05 11:07:48

一般來說，我並不認爲超出訪問在主機上的陣列的端部（在普通的C/C++代碼）將立即觸發故障（例如SEG故障等）

在GPU上，沒有硬件機制可以跟蹤所有分配到字節級別。有一個通用的硬件機制，可以跟蹤分配的內存頁面，並發現訪問是否在有效頁面之外，但粒度並不低於字節或元素級別（並且我不認爲這是如此主機CPU）。

從結構上看，較新的GPU具有更好的硬件訪問跟蹤機制。此外，cuda-memcheck可以對性能進行更嚴格的跟蹤，但性能會大大降低，因爲它正在進行部分基於SW的跟蹤和部分基於HW的跟蹤，可能有點類似於主機上的valgrind等工具。因此，儘管您似乎有一個期望，即任何與分配空間的偏差都會立即觸發故障，但GPU HW本身不支持該功能（並且AFAIK CPU HW也不支持，至少在現代需求分頁的虛擬內存OS的）。通過軟件干預（即cuda-memcheck），情況通常得到改善，但仍會因硬件世代而有所不同。

來源

2014-11-05 13:26:52

爲什麼CUDA中存在未定義的外部共享內存行爲？

回答

相關問題