cuda中多重循环的实现.

[makedoge3000]

我有一个双重循环.类似:
for(i=0; i<1024; i++){
for(j=0; j<32; j++)
at1[j] = at0[j] + abs(rand(0));
for(j=0; j<32; j++)
at0[j] = at1[j];
}
我要在cuda中实现它.不过因为外重循环要每次都要等里面的循环完成才能继续，所以不知如何做，希望各位朋友能够多指教．

[makedoge3000]

应是说:每一层的at0都是相关的.它要等上一层完后数据传到下一层才会正确.不然就不正确.因为要等上一层结束,所以老是等不了.所以不知道如何做.

[xuyuanbetter]

路过看看！

[jylirui]

为了加分

[jurenwangzi8]

谢谢分享!

[Cyrosly]

我的意思是是整个循环是只有at1 32个输出,还是全部1024层都有32个单独的输出,
还有:
"应是说:每一层的at0都是相关的.它要等上一层完后数据传到下一层才会正确.不然就不正确.因为要等上一层结束,所以老是等不了.所以不知道如何做."

你这里所说的每一层,是指最外层还是,内层的2个循环.如果1024层每层的32个at*都是单独的切你确定你所示代码逻辑(5L)是正确的,则可以如下处理:

//block layout:(THREADS,1,1)
//grid  layout:(1024*32/THREAD,1)

__global__
void kernel(float* out,const float* in)
{
    const uint tidx=__umul24(blockDim.x,blockIdx.x)+threadIdx.x;
    
    extern __shared__ float smem[];
    
    //load at0 form global mem to shared mem
    smem[threadIdx.x+blockDim.x]=in[tidx];
    //compute each group at1 through one warp
    smem[threadIdx.x]=smem[threadIdx.x+blockDim.x]+__frand(0);
    
    //note:there`s no need to used __syncthreads() function
    //because there`s no relationship between every two warps
    //the computing in every warp depend only itself
    
    const uint laneid=threadIdx.x&31u;
    const uint warpid=threadIdx.x>>5;
    const uint stride=warpid<<5;
    
    if(laneid<16u){
        smem[(laneid<<1)+stride+blockDim.x  ]=smem[laneid+stride    ];
        smem[(laneid<<1)+stride+blockDim.x+1]=smem[laneid+stride+16u];
    }
    
    //output the at1
    out[tidx]=smem[threadIdx.x];
}

[st_wsz]

有依赖性有点麻烦...你内层的两个函数可以分开做么？
要是可以写成两个
for{
for{
...
}
}不影响的话，可以写两个kernel function

[st_wsz]

再仔细看了下5楼的...明显循环都必须按顺序来的，并行效率会很低很低.....

[nino]

问题挺好的

[maggiequ]

[Quote=引用 13 楼 cyrosly 的回复:]
我的意思是是整个循环是只有at1 32个输出,还是全部1024层都有32个单独的输出,
还有:
"应是说:每一层的at0都是相关的.它要等上一层完后数据传到下一层才会正确.不然就不正确.因为要等上一层结束,所以老是等不了.所以不知道如何做."

你这里所说的每一层,是指最外层还是,内层的2个循环.如果1024层每层的32个at*都是单独的切你确定你所示代码逻辑(5L)是正确的,则可以如下处理:
C/C++ code//block layout:(THREADS,1,1)//grid layout:(1024*32/THREAD,1)
__global__void kernel(float*out,constfloat*in)
{constuint tidx=__umul24(blockDim.x,blockIdx.x)+threadIdx.x;extern __shared__float smem[];//load at0 form global mem to shared mem smem[threadIdx.x+blockDim.x]=in[tidx];//compute each group at1 through one warp smem[threadIdx.x]=smem[threadIdx.x+blockDim.x]+__frand(0);//note:there`s no need to used __syncthreads() function//because there`s no relationship between every two warps//the computing in every warp depend only itselfconstuint laneid=threadIdx.x&31u;constuint warpid=threadIdx.x>>5;constuint stride=warpid<<5;if(laneid<16u){
smem[(laneid<<1)+stride+blockDim.x ]=smem[laneid+stride ];
smem[(laneid<<1)+stride+blockDim.x+1]=smem[laneid+stride+16u];
}//output the at1out[tidx]=smem[threadIdx.x];
}
[/Quote]

编程风格值得我学习 ：-）

[wolaizisanchuan]

[Quote=引用 13 楼 cyrosly 的回复:]

laneid=threadIdx.x&31u;constuint warpid=threadIdx.x>>5;constuint stride=warpid<<5;if(laneid<16u){
        smem[(laneid<<1)+stride+blockDim.x  ]=smem[laneid+stride    ];
        smem[(laneid<<1)+stride+blockDim.x+1]=smem[laneid+stride+16u];
    }//output the at1out[tidx]=smem[threadIdx.x];

[/Quote]
这几句看不懂，麻烦解释一下好吗？谢谢

[xingang_ahu]

路过 学习一下

[zwang4000]

不错不错

[yajiedesign]

kernel1 中不写循环 直接
at1[j] = at0[j] + abs(rand(0));
开32个thread
kernel2 中不写循环 直接
at0[2*j] = at1[j];
at0[2*j+1] = at1[j+16];
开16个thread

[newprogrammer09]

学习一个！

[yaoyaoflyfreely]

牛逼啊~~这个问题很好~~

[eblis88]

对于内外层循环依赖的问题，其实最好最快最方便的办法就是改算法。
依赖方式是cpu实现高效的，但gpu确实低下的。

改写算法，增加空间复杂度，更改数据存储结构。

[tinger27]

我也想有这个问题，不过是循环嵌套的，貌似不能

[zqj08]

恩，值得探讨一下，这一块我也有些不太明白

[王维_ICT]

我也有个这样的循环需要在GPU上并行处理

[cust_hf]

看楼主的程序，我也有点晕。

感觉像是两个长度为32的一维数组at0和at1。每次根据at0生成at1，然后根据at1再更新at0。循环处理1024次。如果是这样的话，我发现一个技巧。就是元素对应技巧。

每次处理过后，at0中位于位置indexAfter1的元素总是对应于处理之前的位置indexBefore1。循环1024次之后，位于位置indexAfter1024的元素也有一个对应的处理之前的位置indexBefore1024。所以，可以先1024次处理。最后再根据对应关系调整顺序。


要是万一，我理解得有错。是两个长度为32*1024的一维数组at0和at1，根据at0生成at1,在根据at1更新at0的话。那这就麻烦多了。感觉可以书写两个内核函数，一个负责计算。另外一个负责调整顺序。