오류 : CUDA 동기화의 BFS

Question

다음 코드에 오류가 발생하여 그래프 가중치 중 일부가 덮어 쓰여지고 있지만 Xa 배열 (어느 것이 이미 방문했는지를 유지)과 함께 발생해서는 안됩니다. __syncthreads() 함수 ... 누군가 도움을 줄 수 있습니까?

struct Node 
{ 
    int begin;  // begining of the substring 
    int num; // size of the sub-string 
}; 

__global__ void BFS (Node *Va, int *Ea, bool *Fa, bool *Xa, int *Ca, bool *parada) 
{ 
    int tid = threadIdx.x; 

    if (Fa[tid] == true && Xa[tid] == false) 
    { 
     Fa[tid] = false; 
     __syncthreads(); 

     // Va begin is where it's edges' subarray begins, Va is it's 
     // number of elements 
     for (int i = Va[tid].begin; i < (Va[tid].begin + Va[tid].num); i++) 
     {   
      int nid = Ea[i]; 

      if (Xa[nid] == false) 
      { 
       Ca[nid] = Ca[tid] + 1; 
       Fa[nid] = true; 
       *parada = true; 
      } 
     }  
     Xa[tid] = true;    
    } 
} 

// The BFS frontier corresponds to all the nodes being processed 
// at the current level. 
int main() 
{ 

    //descrição do grafo 
    struct Node node[4]; 
    node[0].begin=0; 
    node[0].num=2; 
    node[1].begin=1; 
    node[1].num=0; 
    node[2].begin=2; 
    node[2].num=2; 
    node[3].begin=1; 
    node[3].num=0; 
    int edges[]={1,2,3,1}; 

    bool frontier[4]={false}; 
    bool visited[4]={false}; 
    int custo[4]={0}; 

    int source=0; 
    frontier[source]=true; 

    Node* Va; 
    cudaMalloc((void**)&Va,sizeof(Node)*4); 
    cudaMemcpy(Va,node,sizeof(Node)*4,cudaMemcpyHostToDevice); 

    int* Ea; 
    cudaMalloc((void**)&Ea,sizeof(Node)*4); 
    cudaMemcpy(Ea,edges,sizeof(Node)*4,cudaMemcpyHostToDevice); 

    bool* Fa; 
    cudaMalloc((void**)&Fa,sizeof(bool)*4); 
    cudaMemcpy(Fa,frontier,sizeof(bool)*4,cudaMemcpyHostToDevice); 

    bool* Xa; 
    cudaMalloc((void**)&Xa,sizeof(bool)*4); 
    cudaMemcpy(Xa,visited,sizeof(bool)*4,cudaMemcpyHostToDevice); 

    int* Ca; 
    cudaMalloc((void**)&Ca,sizeof(int)*4); 
    cudaMemcpy(Ca,custo,sizeof(int)*4,cudaMemcpyHostToDevice); 

    dim3 threads(4,1,1); 

    bool para; 
    bool* parada; 
    cudaMalloc((void**)&parada,sizeof(bool)); 
    printf("\n"); 
    int n=1; 
    do{ 
     para=false; 
     cudaMemcpy(parada,&para,sizeof(bool),cudaMemcpyHostToDevice);  
     BFS <<<1,threads>>>(Va,Ea,Fa,Xa,Ca,parada);  
     CUT_CHECK_ERROR("kernel1 execution failed"); 
     cudaMemcpy(&para,parada,sizeof(bool),cudaMemcpyDeviceToHost); 



     printf("Run number: %d >> ",n); 
     cudaMemcpy(custo,Ca,sizeof(int)*4,cudaMemcpyDeviceToHost); 
     for(int i=0;i<4;i++) 
      printf("%d ",custo[i]); 
     printf("\n"); 
     n++; 

    }while(para); 


    printf("\nFinal:\n"); 
    cudaMemcpy(custo,Ca,sizeof(int)*4,cudaMemcpyDeviceToHost); 

    for(int i=0;i<4;i++) 
     printf("%d ",custo[i]); 
    printf("\n"); 

} 

Answer 1

해당 장치 코드에는 여러 가지 중대한 결함이 있습니다. 먼저, 메모리 레이스가 Xa과 Ca입니다. 두 번째로, 조건부로 수행 된 __syncthreads() 호출이 있습니다. 이는 불법이며 호출 주변의 분기가 발생할 수있는 스레드 워프에 의해 실행되는 경우 커널이 멈출 수 있습니다.

원자력 메모리 액세스 기능을 사용하여 게시 된 코드에서 최악의 pf 읽기 후 쓰기 레이스를 제거하는 경우에도 사용중인 알고리즘의 구조가 CUDA에서 정확하지 않을 수 있습니다. 원자 적 메모리 액세스를 사용하면 코드를 효과적으로 직렬화하고 많은 비용을 들일 수 있습니다.

CUDA의 광범위한 첫 번째 검색은 아직 해결되지 않은 문제가 아닙니다. 구현을 검토 할 때 좋은 사례가 많이 있습니다. 아직 보지 않았다면 High Performance and Scalable GPU Graph Traversal을 권하고 싶습니다. 그 저자의 구현 코드는 here에서 다운로드 할 수 있습니다.