关于__syncthreads()同步的问题

《code by example》中说除非线程块中每个线程都执行了__syncthreads(),否则没有线程能执行__syncthreads()之后的指令,即如果一些线程没有执行__syncthreads(),处理器将挂起,程序会崩溃。代码是这样描述的:
(正确的)
while(iHalf!=0)
{
if (cacheIndex<iHalf)
{
cachePerBlock[cacheIndex]+=cachePerBlock[cacheIndex+iHalf];
}

	__syncthreads();
	iHalf/=2;
}

(不正确的)
while(iHalf!=0)
{
if (cacheIndex<iHalf)
{
cachePerBlock[cacheIndex]+=cachePerBlock[cacheIndex+iHalf];
__syncthreads();
}
iHalf/=2;
}

但是我使用两种方式都运行了一遍,其实是可以的,不知是我错了还是书中的描述错了?
附完整代码:
#include <cuda_runtime.h>
#include <stdio.h>

#define N 12346 //N必须为偶数,否则进行归约计算时会计算不到线程计算数组中最后一个数;
#define mathMin(a,b) (a<b?a:b)
#define threadNumPerBlock 256
#define blockNum mathMin(32,(N+threadNumPerBlock-1)/threadNumPerBlock)
#define sum_squares(x) (x*(x+1)(2x+1)/6)

//如果计算的是浮点数,得到的结果会不准确,不知道是为什么;

global void CalcInnerProducts(/float/int* pDevA,/float/int* pDevB,/float/int* pDevC)
{
int tid=threadIdx.x+blockIdx.x*blockDim.x;
int cacheIndex=threadIdx.x;
shared /float/int cachePerBlock[threadNumPerBlock];
/float/int temp=0;

while(tid<N)
{
	temp+=pDevA[tid]*pDevB[tid];
	tid+=blockDim.x*gridDim.x;
}

cachePerBlock[cacheIndex]=temp;
__syncthreads();

int iHalf=blockDim.x/2;
while(iHalf!=0)
{
	if (cacheIndex<iHalf)
	{
		cachePerBlock[cacheIndex]+=cachePerBlock[cacheIndex+iHalf];
	}

	__syncthreads();
	iHalf/=2;
}

if (cacheIndex==0)
{
	pDevC[blockIdx.x]=cachePerBlock[cacheIndex];
}

return;

}

int main()
{
/float/int a[N],b[N],c[blockNum];
/float/int *pDevA,*pDevB,*pDevC;
/float/int fTemp=0,fResult;

for (int i=0;i<N;++i)
{
	a[i]=i;
	b[i]=2*i;
}

cudaMalloc((void**)&pDevA,N*sizeof(float));
cudaMalloc((void**)&pDevB,N*sizeof(float));
cudaMalloc((void**)&pDevC,blockNum*sizeof(float));

cudaMemcpy(pDevA,a,N*sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy(pDevB,b,N*sizeof(float),cudaMemcpyHostToDevice);

CalcInnerProducts<<<blockNum,threadNumPerBlock>>>(pDevA,pDevB,pDevC);

cudaMemcpy(c,pDevC,blockNum*sizeof(float),cudaMemcpyDeviceToHost);


for(int i=0;i<blockNum;++i)
{
	fTemp+=c[i];
}

fResult=2*sum_squares((/*float*/int)(N-1));

if (fTemp==fResult)
{
	printf("Success!\r\n");
}

cudaFree(pDevA);
cudaFree(pDevB);
cudaFree(pDevC);

return 0;

}

楼主您好,您提供的《CUDA By Example》一书存在错误,请以下文为准:

(1)一个warp里只要至少有1个线程执行了__syncthreads(), 则等效整个warp都执行了。
(2)已经死亡的warp将被看成自动执行了所有的__syncthreads()。例如:
global void lingchong(…)
{
if (tid < 512) {…这里有__syncthreads();…}
}
则一个1024的block, 后512个也被认为参与了__syncthreads(). (详见2说明)

此非官方的书的说法有问题,建议直接无视。
特此澄清。
感谢来访。

非常感谢您!

感谢您的来访,

建议时候总是尽量别使用非官方的资料。避免受误导。