问题是这样的：
有一个32*8的数组:
32列8行
作为输入数据
任务划分为
grid(1,1)
threads(32,8)
i=threadIdx.x;
j=threadIdx.y;
但是这个算法操作数据之间依赖性太强–每一行的数据和上一行的数据有依赖，每一列对前一列的数据有依赖
稍微解释下，依赖顺序如下
（0.0）---->（0，31）
| |
| |
（31,0）–>（31,31）
每一行的第一列数据影响后面一列的数据
每一行的数据影响下一行和其紧邻的数据，比如d[i+j*32]会影响下一行的d[i-1+（j+1）*32]，d[i+（j+1）*32]和d[i+1+（j+1）*32]
所以就设想：
能不能第一行处理完4个元素之后（避免对第二行元素的影响），
然后第二行开始处理数据，这样以此类推
那么能实现的话，总时间就是最后一行处理完的时间
这样，时间就能缩短为原始处理时间的1/4
CPU实现的代码如下：
for(t=0;t<8;t++){
for(s=0;s<32;s++)
{
t_data[s+32*t]|=(A[t_data[s+32*t]]<<4);
if(s<31)
{
t_data[s+1+32*t]|=(B[t_data[s+32*t]]<<4);
}
//右侧一列4个数据sigpass的元素估计
if((t_data[s+32*t]&(0x10))&&(t_data[s+32*t]&(0x1))&&(t<7))
{
t_data[s+32*(t+1)]|=(0x80);
if(s>0)
{
t_data[s-1+32*(t+1)]|=(0x80);
}
if(s<31)
{
t_data[s+1+32*(t+1)]|=(0x80);
}
}
}
}
//A和B数组是根据算法对数据的更新，就是一张包含256个元素的表，没其他的含义，t_data是输入数据
自己cuda方面理论还不够扎实，
如果出了小白错误，望见谅
愿大家新的一年一切都好。。
拜谢。。
特别感谢一下悠闲的小猫的帮助，
看了论坛您好多热情和专业的回复，深表敬意
衷心祝愿工作顺利，一切都好。。
再次感谢。。
[ 本帖最后由 wscuiqiu 于 2011-2-28 09:57 编辑 ]
首先向您道歉，没有仔细阅读您的代码，不过我看了您的描述，我想，您的问题这样分析可以吗？
[
我尝试对您的描述这么分析：
您有M行N列的数据Array0，然后这个算法是这样的：
伪代码表示下：
void F0(int 列号)
{
for(int i=0;i<M;i++)
{
Arg0 = Array0[列号,i-1];
Arg1 = Array0[列号,i];
Result = G0(Arg0,Arg1); //G0函数是您的某个算法
Array0[i,列号] = Result;
}
}
void F1()
{
for(int i=0;i<N;i++)
{
F0(i);
}
}
F1和F0构成了您原始的CPU实现的算法。
然后您今天说，您尝试使用CUDA来并行化这个算法：
(1)那我们首先看到的是您F1()的这个循环不能被直接展开，因为里面的F0(i)不能满足"打乱顺序执行，而结果不变"。
(2)那我们接着看F0(), F0()里面的循环从通用算法的角度来说，也不能展开。
因为(1)中的列之间的更新依赖，和(2)中的行元素间的更新依赖，所以我们不能做任何展开！
额。。。不过这是对于F1调用的F0中调用的G0()是通用算法的角度看的。
不过还是有希望的，就是盼望G0()是个特殊函数，话可以另当别论，如果G0可以满足如下条件:
G0(Arg0,Arg1) == G0(Arg0,G0(Arg0,Arg1))
且G0(Arg0,Arg1) == G0(G0(Arg0,Arg1),Arg1))
且G0(Arg0,Arg1) == G0(G0(Arg0,Arg1),G0(Arg0,Arg1))
这种情况如果可以满足。那么你可以展开F0(), 同时不需要原子操作。。。
接上文：
在上文尝试了直接展开的讨论后，我们尝试另外一种展开方式：
您的CPU算法的伪代码的列处理部分：
void F1()
{
for(int i=0;i<N;i++)
{
F0(i);
}
}
因为F0(K)依赖于F0(K-1), 进一步我们知道F(K)依赖于原始数据Array0(不需要更新),
果我们可以直接得到一个函数F0’(K)可以直接用原始数据Array0来表示的话，那么就可以直接展开了，
这样，
如果可以得到一个可以经过简化的F0’(K), 还是有意义的。。。。
但如果F0’(K)不能简化，还是需要迭代计算0…K-1列，那基本没意思了。。。估计还没有不展开的快。
现在直接分析您附件中的kernel代码：
//您的kernel代码的一些小问题
[
//你原来首先有了3个warp内分支，还需要volatile, 还导致了多次读取-修改-回写。
//建议用下面的等价形式改写，以避免分支、避免volatile、减少读取-修改-回写次数。
//用2条整数运算代替3个warp内分支
unsigned int v = ~(1 << (unsigned int)i);
v = (v >> 31) * 0xf000 + (v & 1) * 0x0f00 + 0x00f0;
d_data[i+j*32] |= v; //减少了2次global访问，减少了2/3的warp内线程互锁等待.
d_data[i+j*32] += i; //减少了2个__syncthreads()
__syncthreads();
这样多好！比起你上一部分！这是你第一部分代码的问题。
[ 本帖最后由 悠闲的小猫 于 2011-2-27 00:22 编辑 ]
还有你代码的剩下的部分问题太多了…大致看了下，你不能这样简单的写。前面说过的，数据有依赖性啊！！
而且你的代码有一点点乱. 
我建议你别发代码了，用自然语言描述一下，然后大家一起和你看看能否解决！
呵呵。。。
[
谢谢你了。。这个给你添麻烦了。。我回头好好总结下。。
恩。。再次感谢。。祝好。。
[
额。没事。。反正今晚我不困。你第二部分上去就有个小问题：
for(int s=0;s-j4<32;s++) s: [0,j4+32)
{
if (s-j4>-1) [j4,…) //你在让你的for()空转么？这个if完全没必要，合并到上个for()里面多好。
{
//…code
}
}
也就是说：
for(int s=j4; s<j4+32; s++) //合并让for()空转的if了。
{
//…code
}
再往下问题就太多了。。。额。。。。
[ 本帖最后由 悠闲的小猫 于 2011-2-27 00:55 编辑 ]
[
这个空转是我对于这个算法设想处理的一部分，想以此来保证这一行处理完4个元素之后，下一行线程再开始做操作。。不知道这样是否可以呢？
[
这个代码是最原始的版本，直接翻译cpu的程序到上面的。。
这里我也做了下修改。。
我当时是用uppper和lower两个变量来改掉这段warp分支的。。
仁兄这段修改的很好，比我那个好得多，受教了。。
不过貌似这里不是对同一个元素的处理，是对d[i-1+j*32]，d[i+j*32]和d[i+1+j*32]这3个元素的更新
[ 本帖最后由 wscuiqiu 于 2011-2-28 10:20 编辑 ]