= partialSum;
__threadfence();
unsigned int value = atomicInc(&count ,gridDim.x);
isLastBlockDone = (value ==griddim.x-1)
}
__syncthreads();
if(isLastBlockDone)
{
floa totalSum = calclateTotalSum(rusult);
if(threadIdx.x ==0)
{
result[0] = totalSum;
count = 0;
}
}
}
书上说：如果在存储部分和与计数器递增之间不进行fence操作，则计数器的值可能在存储子序列之和之前就已经递增了（这个情况怎么会发生，计数器加1不是在result赋值之后么？）。此时最后一个block读到的子序列和可能还没更新，造成计算结果错误。
[ 本帖最后由 hn0505 于 2011-4-13 21:57 编辑 ]
最近正在看，也有些疑惑，不是很理解
楼主您好，这个问题很好理解:
在CUDA里面，不同线程间的数据读写会彼此影响，这种影响的作用效果根据不同的线程组织单位和不同的读写对象是不同。
在不考虑2.x的优化的情况下，
(1)在同一个warp内的线程读写shared/global,
读写global和shared是立刻对本warp内的其他线程立刻可见的。
(2)在同一个block内的不同warp内线程读写shared/global,
这种读写必须使用__syncthreads(), 或者__threadfence()来实现不同的读写可见效果。
(3)在同一个grid内的不同block内的线程读写shared/gloabl,
这种读写必须使用__threadfence*()来实现一定的读写可见效果。
(4)任何线程组织单位内的原子操作总是可见的。
所以:
[
这种说法的本意是对的，但是说的不对。这本书的意思我想应该是，如果在存储部分和与计数器递增之间不使用__threadfence()，则计数器的值可能， 在存储子序列之和对所有的grid内线程都可见之前就已经递增了。您所说的"这个情况怎么会发生，计数器加1不是在result赋值之后么？", 前面的写入中间结果的操作的确是在atomicInc()之前，但是他们的对别的线程单位的读写效果可见性，生效时间不是同的。
也就是说，因为您的代码是这个grid内的最后一个block在进行中间结果加和。
而同时，因为atomicInc的读写效果是立即的，如果您这里没有使用__threadfence(), 那么当进入了最后一个block，开始求和的时候，由于前面的存储中间结果的其他block的写入效果，对于本最后一个block中的线程来说，没有立刻生效。
也就是说，最后block中的线程有一定可能看不到这种写入的值。所以您必须在让最后一个block之前，用__threadfence()来确保所有的block中的第0个线程的写入效果全grid都可见。而您的代码，这种地方只能写在atomicInc之前了。
总之，读写是你看到的，但是生效是隐藏在后面的。
[ 本帖最后由 悠闲的小猫 于 2011-4-13 21:10 编辑 ]
i can not understand
but it is a goood ideal
您好!感谢您的详细回复。看了几遍有点理解了。
这里还有个问题，程序中有一个__syncthreads();
这样的话，似乎那个__threadfence（）显得多余？
[ 本帖最后由 hn0505 于 2011-4-13 22:05 编辑 ]
共同学习，共同进步，呵呵
[
楼主您好。根据手册的说法。__syncthreads()可以让同一个block的线程们在这个点同步，同时的附加效果是，本block内的所有线程的读写对彼此立刻有效。
看到您后面是不同的block得到的值，由最后一个block来求总和。所以我觉得只有一个__syncthreads()是不够的。
[
谢谢您的耐心答复，受益匪浅！应该是您说的那样，__syncthreads()只保证了对块内的可见，并保证没有对其他块的可见，threadfence保证了这点。
[
呵呵。这不是我说的，那本《Programming Guide》说的是只能使用thread fence类的方法来保证可见性，但是《Fermi Compatiblity Guide》又指出，__syncthreads()有这种隐含的效果(作为一种副作用)。