CUDA内核中的无限循环

Question

我有一个CUDA内核，其中每个线程遍历树。因此，我有一个while循环，直到线程到达一个叶子。在树的每一步中，它都会检查它应该选择跟随哪个子节点。

代码如下:

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

这段代码进入一个无限循环。这就是我觉得奇怪的地方。如果我改变距离方法返回一个常数，它的工作(即。

我错过了一些循环在CUDA或只是有一些隐藏的bug我看不见?因为我不明白代码怎么能进入无限循环

Answer 1

CUDA c++中的循环与c++中的循环具有相同的语义，因此您的代码中一定存在某个错误。调试它的一种策略是在主机上进行调试。

首先，因为你的代码是标量的(例如，它不包含对__syncthreads的调用)，你可以将它重构成__host__ __device__函数。

distance不包含cuda特定的标识符或函数，因此您可以简单地在__host__前面添加:

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

要重构search函数，将tid(取决于cuda特定的标识符threadIndex等)提升到它外面的一个参数中，并使其成为__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之外什么都不做:

__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调用它来调试它，模拟内核启动的操作:

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

它应该像挂在设备上一样挂在主机上。在里面放一个printf看看是哪里。当然，您需要确保在主机端复制您的数组，例如centroids，因为主机不能解引用指向设备内存的指针。

尽管printf可以在较新的硬件上从__device__函数中使用，但您可能更喜欢这种方法的原因是，从内核调用printf直到内核退役之后的才会提交。如果内核永远不会退出(就像你的情况一样)，那么你的调试输出将永远不会出现在屏幕上。