Comportement étrange du noyau CUDA

Question

J'écris un code pour calculer la somme du préfixe. Voici mon noyau

__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(); 

} 

Dans le noyau ci-dessus, tableau sommes accumulera somme de préfixe par bloc et puis premier filet calcule la somme de préfixe de ce tableau de somme. Maintenant, si j'imprime ce tableau sum à partir du périphérique, il affichera des résultats corrects dans

cuPrintf ("Ajouter% d \ n", sommes [blocIdx.x]);

cette ligne imprime qu'elle prend une ancienne valeur. Quelle pourrait être la raison?

Answer 1

Votre code n'est pas une implémentation valide de la somme des préfixes multi-blocs. Vous essayez d'utiliser un seul thread du bloc 0 pour calculer la somme du préfixe des sommes partielles du bloc, avant que ces sommes partielles ne soient écrites en mémoire. __syncthreads() synchronise uniquement les threads dans un seul bloc, pas entre blocs. Donc, dans ce code:

__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]; 
    } 
} 

Tous les blocs ne sont pas garantis avoir calculé leurs sommes [blockIdx.x] avant le bloc 0 exécute ce code. En fait, si vous lancez plus de blocs que ceux qui peuvent s'exécuter simultanément sur le périphérique, tous les blocs ne sont pas garantis avoir même démarré lorsque vous atteignez ce code. Pour corriger ce code, vous devez terminer le noyau avant ce code et lancer un nouveau noyau pour calculer le résultat de la somme du préfixe de bloc, et un autre pour ajouter le résultat à chaque bloc de threads.