并行嵌套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有很大的限制

回答

相关问题