並行嵌套for循環與cuda有很大的限制

-2

我是CUDA的新手。我正在嘗試編寫一個CUDA內核來執行下面的一段代碼。並行嵌套for循環與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，Wjv = 1920，比例= 255;盎司> = 4，本是我現在，但我的代碼只能執行時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，我做什麼並行化嵌套的代碼？

來源

2017-07-25 Jany

CUDA與C無關。 – Olaf

我簡要地檢查了代碼，它似乎Ic_img和Ic_real元素的計算是簡單的並行（COUNT1，規模+ 1，五交化，Wjv沒有依賴於所有彼此之間）。因此，內核中不需要共享變量和while循環;這很容易實現，如下所示，其中一個額外的參數int numElements = count1 *（scale + 1）* Wjh * Wjv。

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

該代碼將明顯更容易維護並消除像您的示例那樣容易出現長代碼的錯誤。如果src2值在最內層循環中完全不重複，則性能也接近最佳。如果'src2'可能重複，請使用帶有'atomicAdd'的表達式，以便結果如預期的那樣正確;使用atomicAdd，性能可能不是最佳的，但至少有一個正確實現的無bug內核已成功實現。如果它導致性能瓶頸，那麼通過嘗試和嘗試一些不同的實現來修改它。

來源

2017-07-26 07:35:14 Tom

並行嵌套for循環與cuda有很大的限制

回答

相關問題