Бесконечный цикл в ядре CUDA - PullRequest
Новая
       5

Бесконечный цикл в ядре CUDA

0 голосов
/ 21 ноября 2011

У меня есть ядро ​​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 или есть какая-то скрытая ошибка, которую я не вижу?Потому что я не понимаю, как код может войти в бесконечный цикл.

1 Ответ

4 голосов
/ 22 ноября 2011

Циклы в 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__, вы можете отладить его, вызвав его из ЦП и эмулировав, что будет делать запуск ядра: 

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

Он должен висеть на хосте точно так же, как на устройстве. Вставьте printf внутрь, чтобы узнать, где. Конечно, вы должны быть уверены, что сделали копии своих массивов на стороне хоста, такие как centroids, потому что хост не может разыменовывать указатели на память устройства.

Несмотря на то, что printf доступно для использования из __device__ функций с более новым оборудованием, вы можете предпочесть этот подход, потому что вызовы printf из ядра не фиксируются до после ядро удаляется. Если ядро ​​никогда не выходит из системы (как, по-видимому, в вашем случае), ваши выходные данные отладки никогда не появятся на экране.

...