キューを持つforループをネストした並列化には大きな制限があります

Question

CUDAを初めて使用しています。私は次のコードを実行するためにCUDAカーネルを作成しようとしています。

for(int oz=0;oz<count1;oz++) 
    { 
     for(int ox=0;ox<scale+1;ox++) 
     { 

      for(int xhn=0;xhn<Wjh;xhn++) 
      { 
       for(int yhn=0;yhn<Wjv;yhn++) 
       { 
        //int numx=xhn+ox*Wjh; 
        int numx=oz*(scale+1)*Wjh+ox*Wjh+xhn; 
        int src2=yhn+xhn*Wjv; 
        Ic_real[src2]=Ic_real[src2]+Sr[oz*(scale+1)*Wjv+ox*Wjv+yhn]*Hr_table[numx]-Si[oz*(scale+1)*Wjv+ox*Wjv+yhn]*Hi_table[numx]; 
        Ic_img[src2]=Ic_img[src2]+Sr[oz*(scale+1)*Wjv+ox*Wjv+yhn]*Hi_table[numx]+Si[oz*(scale+1)*Wjv+ox*Wjv+yhn]*Hr_table[numx]; 
       } 

      } 

     } 
    } 

値Wjh = 1080、1920 Wjv =、規模= 255;オンス> = 4.Thisは私が現在持っているものですが、私のコードのみを実行することができたときにCOUNT1 < = 4、もしオンス> 4、それは動作しません、誰が私は何をすべきか知っていますか？歓声

__global__ void lut_kernel(float *Sr,float *Si,dim3 size,int Wjh,int Wjv,float *vr,float *vi, 
          float *hr,float *hi,float *Ic_re,float *Ic_im) 
{  
    __shared__ float cachere[threadPerblock]; 
    __shared__ float cacheim[threadPerblock]; 
    int blockId=blockIdx.x + blockIdx.y * gridDim.x; 
    int cacheIndex=threadIdx.y*blockDim.x+threadIdx.x; 
    int z=threadIdx.x; 
    int x=threadIdx.y; 
    int tid1=threadIdx.y*blockDim.x+threadIdx.x; 

    //int tid= blockId * (blockDim.x * blockDim.y) 
        // + (threadIdx.y * blockDim.x) + threadIdx.x; 
    int countnum=0; 
    float re=0.0f; 
    float im=0.0f; 
    float re_value=0.0f; 
    float im_value=0.0f; 
    if (z<4 && x<256) 
    {  

     int src2=z*(scale+1)*Wjh+x*Wjh+blockIdx.y; 
     re=Sr[z*(scale+1)*Wjv+x*Wjv+blockIdx.x]*hr[src2]-Si[z*(scale+1)*Wjv+x*Wjv+blockIdx.x]*hi[src2]; 
     im=Sr[z*(scale+1)*Wjv+x*Wjv+blockIdx.x]*hi[src2]+Si[z*(scale+1)*Wjv+x*Wjv+blockIdx.x]*hr[src2]; 


     } 
     cachere[cacheIndex]=re; 
     cacheim[cacheIndex]=im; 

     __syncthreads(); 

     int index=threadPerblock/2; 
     while(index!=0) 
     { 
      if(cacheIndex<index) 
      { 
       cachere[cacheIndex]+=cachere[cacheIndex+index]; 
       cacheim[cacheIndex]+=cacheim[cacheIndex+index]; 
      } 
      index/=2; 

     } 
     if(cacheIndex==0) 
     { 
     Ic_re[blockId]=cachere[0]; 
     Ic_im[blockId]=cacheim[0]; 
     //printf("Ic= %d,blockId= %d\n",Ic_re[blockId],blockId); 
     } 

    } 

カーネルパラメータがある： DIM3 dimBlock（count1,256）。 dim3 dimGrid（Wjv、Wjh）;

lut_kernel<<<dimGrid,dimBlock>>>(d_Sr,d_Si,size,Wjh,Wjv,dvr_table,dvi_table,dhr_table,dhi_table,dIc_re,dIc_im); 

count1> 4の場合、ネストしたコードを並列化するにはどうすればよいですか？

Answer 1

は、私は簡単にコードをチェックし、（COUNT1、規模Ic_imgとIc_real要素の計算が並列化しやすいようです+が互いの間にまったく依存関係を持たないWjh、Wjv 1、 ）。したがって、カーネル内の変数やwhileループを共有する必要はありません。以下のように実装するのは簡単です。余分なパラメータは、numElements = count1 *（scale + 1）* Wjh * Wjvです。

int i = blockDim.x * blockIdx.x + threadIdx.x; 
if (i < numElements) { 
    //.... 
} 

このコードは、あなたの例のような長いコードの傾向があるバグを管理し、除去することが非常に簡単になります。最も内側のループでsrc2の値がまったく繰り返されない場合、パフォーマンスは最適に近いものになります。 'src2'が繰り返される場合は、 'atomicAdd'の式を使用して、結果が期待どおり正しいようにします。 atomicAddでは、パフォーマンスは最適ではないかもしれませんが、正しく実装されたバグのないカーネルが少なくとも1つはうまく実装されています。パフォーマンスのボトルネックが発生する場合は、いくつかの異なる実装を試して試してみてください。