【新手】一个简单的cuda程序,循环次数不对,请帮忙

调用方法:matchByArr1<<<100, 100, 0>>>(feature1,feature2,countx);

方法代码:
if (bid == 0 && tid ==0)
{
num = 0;//统计每个BLOCK的执行次数
}
for ( unsigned long l1 = (bidfeaturelen + tidpointlen); l1 < (bid*featurelen + (tid+5)*pointlen); )
{
for (long l3 = 0; l3 < pointNum; l3++)
{
for (long l4 = 0; l4 < pointlen; l4++)
{
num++;
}
}
l1 = l1 + pointlen;
}
__syncthreads();
count[bid] = num;//返回每个BLOCK的执行次数

执行结果部分如下:
block0 loop count:712318
block1 loop count:712697
block2 loop count:713206
block3 loop count:712934
block4 loop count:712223
block5 loop count:713206
block6 loop count:713206
block7 loop count:713206
block8 loop count:665875
block9 loop count:712999
block10 loop count:713150
block11 loop count:713206
block12 loop count:713206
block13 loop count:713407
block14 loop count:713206
block15 loop count:713206

问题:
理想中每个BLOCK的执行次数应该是:1005pointNum*pointlen次。
为啥每个BLOCK循环的次数不一样呢?

LZ您好,您的代码没有贴全,多个变量看不到定义,不过仍将尝试无责任猜测一下:

您代码中的累加变量num没有给出定义,我按照3种情况讨论:

1:num是每个线程私有的变量。那么,第一个block的第一个线程对自己的num初始化为0,其他线程将直接使用某个未知的初始化值。
下面三重循环累加计数此时是可以的,但是num的增量为pointlenpointNum5,是一个线程的累加次数。同时num的具体值还受到初值的影响。
在同步后,每个block内的所有线程都将自己的num写入同一个global 的地址count[bid],最终结果是随机的。并且不是一个block所有线程的循环次数的累加。

2:如果您的num是一个block共享的,那么,第一个block的第一个thread对本block的num进行了初始化为0 ,其他block直接不初始化就黑上。
而且,第一个block在初始化的时候,并没有使用__syncthreads()来保证0号线程对num的赋值真正完成以后才读取,所以一些线程读到的num的初始值可能是错误的。
然后您在循环累加的过程中,使用了普通累加而不是原子加法,那么会导致同一block内一群线程向同一地址不断地读取——加一——回写——再读取。此时因为没有使用原子加法,会导致线程b读到的不保证是其他线程已经完好回写的值。这样的累加是不正确的。
最后在同步之后,num不管如何,算是有了一个值,然后这一个block里面的线程每人把这个值回写到count[bid]里面一次,做了很多无用功。

3:如果您的num是全局唯一的,那么
您使用block0的thread0初始化,这完全不保证其他block能正确获得这一初始化值(因为block间执行顺序无法保证),block0内部,因为没有__syncthreads(),也不保证其他线程能读取到正确的初始值。

然后循环累加的时候,和2:里面一样,因为没有使用原子操作,这个累加也是不正确的。

最后每个block的线程把一个不确定数值的num写会本block的count[bid],请注意,在回写的时候,num的值还可能被其他block的线程改写,这既不是block的循环数量,也不是一个thread的,也不是整个grid的,而且即便初始化正确,即便使用了原子操作累加,这个num的值也会变化的,也没有意义。

综上,稍微符合LZ实现意图的是2:的做法,但是依然BUG多多。加上LZ只给出了不完整的代码,所以仅作上述猜测,聊以娱乐。

祝LZ好运!

经好友提醒,__threadfence()无法实现需求,特此编辑,只留下__syncthreads()

global static void matchByArr1(int *feature1, int *feature2, unsigned long *count)
{
//使用100个block,每个block中含有100个线程,每个线程处理5组128的数据分别与500个待比较的点计算
const int tid = threadIdx.x;
const int bid = blockIdx.x;

const long featurelen = 1;
const int pointlen = 128;
const int pointNum = 500;

__shared__ unsigned  long num ;
if (bid == 0 && tid ==0)
{
	num = 0;
}

//取每个图像的特征,准备对比
for ( unsigned long l1 = (bid*featurelen + tid*pointlen); l1 < ((bid)*featurelen  + (tid+5)*pointlen); )
{
	for (long l3 = 0; l3 < pointNum; l3++)
	{
		for (long l4 = 0; l4 < pointlen; l4++)
		{
			{
				num++;
			}					
		}
	}				
	l1 = l1 + pointlen;
}
__syncthreads();
count[bid] = num;

}

这个是全部的代码,我的本意是sum只有在bid为0且tid也为0的情况下才会进行初始化,然后该block下的其它线程可以都进行直接的累加。 而且代码最后也加上了__syncthreads(); 为啥还是出现每个block的num值不一样呢?

非常感谢。

根据您的代码,您实际上是上述推断中的2:的情况。

其实上面都已经说过了,不过我再结合您的代码说一下吧。
1:根据您的代码,您有100个block,每个block有自己的shared memory中存放的num变量。
那么,您只给block0的num初始化( if (bid == 0 && tid ==0)),其他所有block都无法初始化自己的num。

2:即便是block0,您没有使用__syncthreads()来保证num初始化赋0值的操作已经确实完成,所以其他线程读取到的num值依然可能是初始化之前的。

3:您在循环累加的时候没有使用原子操作。

上述3点问题足以造成您观察到的结果。

另外__syncthreads()的作用范围是一个block内部。它只能同步一个block内部的线程,block之间是无作用的。

大致如上,请您理清思路,修改您的代码。

祝您编码顺利~


经好友提醒,__threadfence()无法实现需求,特此编辑,只留下__syncthreads().

非常感谢。:lol

:lol:lol:lol