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內核中的無限循環

回答

相關問題