CUDA内核中的无限循环

我有一个CUDA内核，每个线程遍历一棵树。正因为如此，我有一个循环，直到线程到达叶子为止。在树下的每一步都会检查应该选择哪个孩子。CUDA内核中的无限循环

的代码如下：

__global__ void search(float* centroids, float* features, int featureCount, int *votes) 
{ 
    int tid = threadIdx.x + blockIdx.x * blockDim.x; 

    if(tid < featureCount) 
    { 
     int index = 0; 
     while (index < N) 
     { 
      votes[tid] = index; 
      int childIndex = index * CHILDREN + 1; 
      float minValue = FLT_MAX; 

      if(childIndex >= (N-CHILDREN)) break; 

      for(int i = 0; i < CHILDREN; i++) 
      { 
       int centroidIndex = childIndex + i; 
       float value = distance(centroids, features, centroidIndex, tid); 
       if(value < minValue) 
       { 
        minValue = value; 
        index = childIndex + i; 
       } 
      } 
     } 
     tid += blockDim.x * gridDim.x; 
    } 
} 

__device__ float distance(float* a, float* b, int aIndex, int bIndex) 
{ 
    float sum = 0.0f; 
    for(int i = 0; i < FEATURESIZE; i++) 
    { 
     float val = a[aIndex + i] - b[bIndex + i]; 
     sum += val * val; 
    } 

    return sum; 
}

此代码进入无限循环。这是我觉得奇怪的。如果我改变distance方法返回一个常量，它就起作用（即遍历树中的左边）。

我错过了CUDA中的循环，还是只有一些隐藏的错误，我看不到？因为我没有看到代码如何进入无限循环。

来源

2011-11-21 Mads Andersen

有一个隐藏的bug :)你可以尝试调试这个代码，在主机上执行它并检查哪个''tid'''导致无限循环。 –

通过在主机上执行并检查哪个tid导致无限循环，你是什么意思？我只能从设备代码中获得tid :)我尝试使用nvidea的“cuPrintf”，但我不确定我能否信任它。 –

CUDA C++中的循环与C++中的循环具有相同的语义，因此在代码中必定存在一处错误。调试它的一个策略是在主机上这样做。首先，由于您的代码是标量（例如，它不包含对__syncthreads的调用），因此可以将它重构为__host__ __device__函数。

distance不包含特定CUDA的标识符或功能，这样你就可以简单地追加__host__：（依赖于CUDA专用标识threadIndex等）

__host__ __device__ float distance(float* a, float* b, int aIndex, int bIndex);

重构你的search功能，葫芦tid它为参数之外，并使它成为一个__host__ __device__功能：

__host__ __device__ void search(int tid, float* centroids, float* features, int featureCount, int *votes) 
{ 
    if(tid < featureCount) 
    { 
    int index = 0; 
    while (index < N) 
    { 
     votes[tid] = index; 
     int childIndex = index * CHILDREN + 1; 
     float minValue = FLT_MAX; 

     if(childIndex >= (N-CHILDREN)) break; 

     for(int i = 0; i < CHILDREN; i++) 
     { 
     int centroidIndex = childIndex + i; 
     float value = distance(centroids, features, centroidIndex, tid); 
     if(value < minValue) 
     { 
      minValue = value; 
      index = childIndex + i; 
     } 
     } 
    } 
    } 
}

现在写一个__global__功能，什么也不做，除了计算tid并调用search：

for(int tid = 0; tid < featureCount; ++tid) 
{ 
    search(tid, centroids, features, featureCount, votes); 
}

它应该：

__global__ void search_kernel(float *centroids, float features, int featureCount, int *votes) 
{ 
    int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    search(tid, centroids, features, featureCount, votes); 
}

因为search现在__host__ __device__，您可以通过从CPU调用它，模仿什么内核启动会做调试它完全像在设备上一样挂在主机上。在里面贴一个printf找出在哪里。当然，您需要确保您的阵列的主机端副本（如centroids），因为主机无法取消引用指向设备内存的指针。

即使printf可从与新硬件__device__功能使用，原因你可能更喜欢这种方法，从内核到printf电话不承诺直到后内核退休。如果内核从不退出（因为它明显不在你的情况下），那么你的调试输出将永远不会出现在屏幕上。

来源

2011-11-21 20:12:27

谢谢！我没有意识到这一点。这将有助于很多调试我认为:) –

CUDA内核中的无限循环

回答

相关问题