CUDA 커널의 이상한 동작

Question

프리픽스 합 계산을위한 코드를 작성 중입니다. 여기서 상기 커널의 커널 내

__global__ void prescan(int *indata,int *outdata,int n,long int *sums) 
{ 
    extern __shared__ int temp[]; 

    int tid=threadIdx.x; 
    int offset=1,start_id,end_id; 
    int *global_sum=&temp[n+2]; 

    if(tid==0) 
    { 
     temp[n]=blockDim.x*blockIdx.x; 
     temp[n+1]=blockDim.x*(blockIdx.x+1)-1; 
     start_id=temp[n]; 
     end_id=temp[n+1]; 
     //cuPrintf("Value of start %d and end %d\n",start_id,end_id); 

    } 
    __syncthreads(); 
    start_id=temp[n]; 
    end_id=temp[n+1]; 
    temp[tid]=indata[start_id+tid]; 
    temp[tid+1]=indata[start_id+tid+1]; 


    for(int d=n>>1;d>0;d>>=1) 
    { 
     __syncthreads(); 
     if(tid<d) 
     { 
      int ai=offset*(2*tid+1)-1; 
      int bi=offset*(2*tid+2)-1; 

      temp[bi]+=temp[ai]; 
     } 
     offset*=2; 
    } 

    if(tid==0) 
    { 
     sums[blockIdx.x]=temp[n-1]; 
     temp[n-1]=0; 
     cuPrintf("sums %d\n",sums[blockIdx.x]); 
    } 
    for(int d=1;d<n;d*=2) 
    { 
     offset>>=1; 
     __syncthreads(); 
     if(tid<d) 
     { 
      int ai=offset*(2*tid+1)-1; 
      int bi=offset*(2*tid+2)-1; 
      int t=temp[ai]; 
      temp[ai]=temp[bi]; 
      temp[bi]+=t; 
     } 
    } 

    __syncthreads(); 

    if(tid==0) 
    { 
     outdata[start_id]=0; 
    } 

    __threadfence_block(); 
    __syncthreads(); 
    outdata[start_id+tid]=temp[tid]; 
    outdata[start_id+tid+1]=temp[tid+1]; 
    __syncthreads(); 

    if(tid==0) 
    { 
     temp[0]=0; 
     outdata[start_id]=0; 

    } 

    __threadfence_block(); 
    __syncthreads(); 

    if(blockIdx.x==0 && threadIdx.x==0) 
    { 
     for(int i=1;i<gridDim.x;i++) 
     { 
      sums[i]=sums[i]+sums[i-1]; 
     } 
    } 

    __syncthreads(); 
    __threadfence(); 

    if(blockIdx.x==0 && threadIdx.x==0) 
    { 
     for(int i=0;i<gridDim.x;i++) 
     { 
      cuPrintf("****sums[%d]=%d ",i,sums[i]); 
     } 
    } 


    __syncthreads(); 
    __threadfence(); 


    if(blockIdx.x!=gridDim.x-1) 
    { 
     int tid=(blockIdx.x+1)*blockDim.x+threadIdx.x; 
     if(threadIdx.x==0) 
      cuPrintf("Adding %d \n",sums[blockIdx.x]); 
     outdata[tid]+=sums[blockIdx.x]; 

    } 
    __syncthreads(); 

} 

이고, 합계 어레이 블록 당 프리픽스 합을 축적되고 그리고 제 실이 합 어레이의 프리픽스 합을 계산한다. 이제 장치 쪽에서이 합계 배열을 인쇄하면 올바른 결과가 표시됩니다.

cuPrintf ("% d \ n"합계 [blockIdx.x]);

이 줄은 오래된 값을 취하고 있음을 나타냅니다. 그 이유는 무엇일까요?

Answer 1

코드가 다중 블록 접두사 합계의 유효한 구현이 아닙니다. 이러한 부분 블록 합계가 메모리에 기록되기 전에 부분 블록 합계의 접두사 합계를 계산하기 위해 블록 0의 단일 스레드를 사용하려고합니다. __syncthreads()은 단일 블록의 스레드 만 동기화하며 블록을 통과하는 스레드는 동기화하지 않습니다. 따라서이 코드에서 :

__threadfence_block(); 
__syncthreads(); 

if(blockIdx.x==0 && threadIdx.x==0) 
{ 
    for(int i=1;i<gridDim.x;i++) 
    { 
     sums[i]=sums[i]+sums[i-1]; 
    } 
} 

블록 0에서이 코드를 실행하기 전에 모든 블록이 합계 [blockIdx.x]를 계산하지 않을 수 있습니다. 사실, 장치에서 동시에 실행할 수있는 것보다 많은 블록을 실행하면이 코드에 도달 할 때 모든 블록이 시작되지 않을 수도 있습니다.

이 코드를 올바르게 작성하려면이 코드 앞에 커널을 종료하고 블록 접두사 합계 결과를 계산하기 위해 새 커널을 시작하고 그 결과를 각 스레드 블록에 추가해야합니다.