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에서 루프를 놓친 적이 있습니까? 아니면 볼 수없는 숨겨진 버그가 있습니까? 왜냐하면 코드가 무한 루프로 어떻게 움직일 수 있는지 알지 못하기 때문입니다.

Answer 1

CUDA C++의 루프는 C++에서와 동일한 의미를 가지므로 코드 어딘가에 버그가 있어야합니다. 디버깅을위한 하나의 전략은 호스트에서 그렇게하는 것입니다.

먼저 코드가 스칼라 (예 : __syncthreads에 대한 호출을 포함하지 않음)이므로 __host__ __device__ 함수로 리팩토링 할 수 있습니다. 당신은 단순히 __host__을 붙일 수

distance

, 아니 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__ 함수를 작성하는 제외하고 아무것도 search을 tid을 계산하지 않고 전화 않습니다

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

그것은해야을 지금 __host__ __device__입니다

__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 때문에, 당신은 커널 출시 어떻게 할 것인지 모방의 CPU에서 호출하여 디버깅 할 수 있습니다 장치에서와 똑같이 호스트에서 멈추십시오. 내부에 printf을 붙이면 어디 있는지 알 수 있습니다. 물론 호스트가 centroids과 같은 배열의 호스트 쪽 복사본을 만들도록해야합니다. 호스트가 장치 메모리에 대한 포인터를 역 참조 할 수 없기 때문입니다.

printf는 새로운 하드웨어와 __device__ 기능에서 사용하는 데 사용할 수 있지만, 이유는이 방법을 선호 할 수는 커널에서 printf에 대한 호출은 커널이 은퇴 후 때까지 범하지 않는다는 것입니다. 커널이 결코 은퇴하지 않는다면 (분명히 여러분의 경우는 아님) 디버깅 결과가 결코 화면에 나타나지 않을 것입니다.