Parallelisierung verschachtelte für Schleife mit Cuda haben große Grenze

Question

Ich bin neu in CUDA. Ich versuche, einen CUDA-Kernel zu schreiben, um den folgenden Code auszuführen.

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

      } 

     } 
    } 

der Wert WJH = 1080, WJV = 1920 scale = 255; oz> = 4.This ist, was ich habe zur Zeit, aber mein Code kann nur durchführen, wenn count1 < = 4, wenn oz> 4, es funktioniert nicht, weiß jemand, was soll ich tun? Prost

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

    } 

der Kernel-Parameter ist: dim3 dimBlock (count1,256); dim3 dimGrid (Wjv, Wjh);

Wenn count1> 4, was shuold ich tun, um die verschachtelte für Code zu parallelisieren?

Answer 1

überprüfte ich den Code kurz und es scheint, dass die Berechnung von Ic_img und Ic_real Elemente ist einfach zu parallelisieren (count1, Skala + 1, WJH, WJV untereinander keine Abhängigkeit überhaupt haben). Daher müssen keine Variablen und While-Schleifen im Kernel gemeinsam verwendet werden. Es ist einfach zu implementieren wie unten, wo ein zusätzlicher Parameter int numElements = count1 * (scale + 1) * Wjh * Wjv.

Der Code wird erheblich einfacher zu pflegen und Bugs zu beseitigen, die anfällig für lange Codes wie Ihr Beispiel sind. Wenn src2-Werte sich in der innersten Schleife überhaupt nicht wiederholen, ist die Leistung ebenfalls nahezu optimal. Wenn 'src2' wiederholt werden kann, verwenden Sie einen Ausdruck mit 'atomicAdd', damit die Ergebnisse wie erwartet korrekt sind. Mit atomAdd ist die Leistung möglicherweise nicht optimal, aber mindestens ein korrekt implementierter fehlerfreier Kernel wurde erfolgreich implementiert. Wenn es zu einem Leistungsengpass kommt, modifizieren Sie es, indem Sie verschiedene Implementierungen ausprobieren und ausprobieren.