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