從CUDA遞歸函數中的結構訪問數據失敗，導致非法內存訪問

Question

我的任務是使用CUDA在高維數據結構中實現最近鄰居搜索。我正在使用遞歸搜索，以便在樹中進行有效的搜索。但首先我試圖在遞歸函數中訪問數據結構中的3d點並計算它們之間的距離。但也有遇到過，當我計算點之間的距離在數據結構的一個非法的內存訪問：

float dist = (a[0]-b[0])*(a[0]-b[0]) + (a[1]-b[1])*(a[1]-b[1]) + (a[2]-b[2])*(a[2]-b[2]); 

下面是完整的小例子：

struct Node { 
    int divfeat; 
    float divval; 

    int child1; 
    int child2; 

    Node() { 
     child1 = NULL; 
     child2 = NULL; 
    } 
}; 

__global__ 
void gpuKernel(int veclen_ /*dimension*/, int size_ /*number of points*/, Node* devpool_ /*tree*/, float* devdataset_ /*points*/) 
{ 
    gpuRecursive(0, veclen_, size_, devpool_, devdataset_); 
} 

/* problems with recusive function */ 
__device__ 
void gpuRecursive(int index_,int veclen_, int size_, Node* devpool_, float* devdataset_) 
{ 
    /* if current Node has a valid children, call gpuRecursive for this child */ 
    if (devpool_[index_].child1) { gpuRecursive(devpool_[index_].child1, veclen_, size_, devpool_, devdataset_); } 
    if (devpool_[index_].child2) { gpuRecursive(devpool_[index_].child2, veclen_, size_, devpool_, devdataset_); 

    /* if current node is a leaf do anything */ 
    if (!devpool_[index_].child1 && !devpool_[index_].child2) { 
     if (devpool_[index_].divfeat != size_){ 
      float* a = &devdataset_[devpool_[index_].divfeat*veclen_]; 
      float* b = &devdataset_[devpool_[index_].divfeat*veclen_]; 

      /* when computing dist an error occcurs */ 
      float dist = (a[0]-b[0])*(a[0]-b[0]) + (a[1]-b[1])*(a[1]-b[1]) + (a[2]-b[2])*(a[2]-b[2]); 
     } 
    } 
} 

但是當我使用一個正常的功能沒有任何遞歸訪問是有效的I可以計算任意兩個點之間的距離：

__device__ 
void gpuRegular(int index_,int veclen_, int size_ Node* devpool_, float* devdataset_) 
{ 
    for (int i = 0; i< size_; i++) { 
     ElementType* a = &devdataset_[i*veclen_]; 
     ElementType* b = &devdataset_[(i+1)*veclen_]; 

     /* when computing dist all works fine */ 
     float dist = (a[0]-b[0])*(a[0]-b[0]) + (a[1]-b[1])*(a[1]-b[1]) + (a[2]-b[2])*(a[2]-b[2]); 
    } 
} 

另外，也可以在遞歸函數和c，以限定兩個陣列注意它們之間的距離。這意味着在使用遞歸函數時，問題出在與全局內存通信的任何地方？

任何人都可以解釋我在CUDA中使用遞歸函數時的這種行爲。有沒有可能繞過這個問題，並通過遞歸實現來解決搜索問題？

Answer 1

[我想發表評論，但我的聲望不允許。雖然我希望我能想出這個問題。]

我的建議是有一個節點divfeat > size_。由於您可能有類似float* devdataset_ = new float[veclen_ * size_];的行，因此會導致遇到非法內存訪問。

你應該改變

if (devpool_[index_].divfeat != size_) 

到

if (devpool_[index_].divfeat < size_) 

這將是類似於在gpuRegular內核，以避免非法的內存訪問的for條件。但是，當然，這隻會隱藏以前在程序中定義節點的錯誤，並且您將不會遇到任何錯誤，但會得到錯誤的結果。

因此，您應該更好地檢查定義Node的代碼行，包括divfeat變量，這不僅可以避免非法內存訪問，還可以獲得正確的結果。

此外，float* a和float* b在gpuRecursive和gpuRegular中的不同定義讓我很好奇你在你的程序中真正做了什麼。 （正如你所提到的，你不是簡單地複製和粘貼你的代碼，而是爲你的問題寫了新的代碼。）在遞歸內核a和b中都存儲相同的地址，但是在常規內核中b將地址存儲到以下數據節點。

Answer 2

我發現當你有一定數量的點時，或者當你有一定深度的底層樹時，距離的遞歸計算會失敗。當您在點雲上使用簡單循環來計算距離時，計算不會失敗。當你不使用遞歸時，程序不會失敗。

float* a = (float*)malloc(sizeof(float) * 3); 
float* b = (float*)malloc(sizeof(float) * 3); 
a[0] = 0.2; a[1] = 0.5; a[2] = 0.4; 
b[0] = 0.4; b[1] = 0.7; b[2] = 0.1; 

但是當你使用的是某種類型的遞歸距離計算的程序不會失敗，以及：

float* a 
float result = float(); 
float diff0; 
float* last = a + veclen_; 
while (a < last) { 
    diff0 = (float)(*a++ - *b++); 
    result += diff0 * diff0; 
} 

以下距離的計算也犯規失敗：

float result = float(); 
result = (a[0] - b[0])*(a[0] - b[0]) + (a[1] - b[1])*(a[1] - b[1]) + (a[2] - b[2])*(a[2] - b[2]); 

但是使用這種類型的距離計算在遞歸中失敗：

float result = float(); 
float diff0,diff1,diff2; 
float* last = a + 3; 
while (a < last - 2) { 
    diff0 = (DistanceType)(a[0] - b[0]); 
    diff1 = (DistanceType)(a[1] - b[1]); 
    diff2 = (DistanceType)(a[2] - b[2]); 

    result += diff0 * diff0 + diff1 * diff1 + diff2 * diff2; 

    a += 3; 
    b += 3; 
}