【更新】应该是关于原子操作的问题

问题是这样的:
有一个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 编辑 ]

还有你代码的剩下的部分问题太多了…大致看了下,你不能这样简单的写。前面说过的,数据有依赖性啊!!
而且你的代码有一点点乱. :frowning: 我建议你别发代码了,用自然语言描述一下,然后大家一起和你看看能否解决!
呵呵。。。

[

谢谢你了。。这个给你添麻烦了。。我回头好好总结下。。
恩。。再次感谢。。祝好。。

[

额。没事。。反正今晚我不困。你第二部分上去就有个小问题:
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 编辑 ]

[

你仔细看看。