1、比如，0号线程给一个share memory赋值后，1 - 31号线程（位于同一wrap的其他线程）何时才嫩看到更改的值？
2、通过线程同步函数sycnthread()可以强制让赋值生效。有没有其他的途径，还是只有这一种手段？wrap vote函数是否具有类似的效果？
楼主您好：
（1）除了__syncthreads()外，最简单的方式是将用volatile修饰shared memory中的数组，这样只要写入语句结束，warp里的其他线程都可以看到。
（2）如果不仅仅是在warp内，而是在block内，推荐使用__syncthreads()替代volatile.
（3）warp vote显然无此效果（它本来就不是用来干shared memory的，无关的，它只能用来交换1bit的数据，在warp内的)。
如果是warp内需要交换4B数据（例如一个float), 您还可以用warp shuffle, 这样可以规避shared memory操作，但您需要购买3.0+(包含)的卡。
感谢来访。
谢谢了。
用threadfence_block()可以嗎?我看說明書的感覺似乎可以
不可以。
你还是得用__syncthreads()或者volatile(分别针对block内和warp内通过shared的交换）。
前者(block内交换）是因为你总是需要同步的（光threadfence无法同步）才能读取，你再用个threadfence是在那里画蛇添足。
而warp内的则可以直接读取，无需用threadfence进行延迟一定周期，shared memory写完(我指真的执行了写入了，例如STS指令后立刻跟随LDS这种指令）后就能读的。
所以通过shared memory交换只需要分别__syncthreads(), 或者volatile即可(分别针对2种情况）。
喔喔我似乎了解了！
如果該線程做__threadfence_block()是他等到他存到shared的結果可以被看見後他就往下走了
但是我是要等待大家都把資料存完後才往下走，而等待大家都做到__syncthreads()時保證大家都做完存入shared的指令也可以讓其他人看見，所以是要用__syncthreads()才對
也難怪__threadfence類的指令只出現在reduction相關的程式碼配合著atomic_func使用。
汇报一下：
在我的程序里。使用wrap shuffle和volatile速度一样，考虑到wrap shuffle可以节省share memory，我选用的是 wrap shuffle。
__syncthreads()会导致结果错误，我的程序无法保证block内的每个线程执行的__syncthreads()次数一样，可能是这个原因导致错误。
感谢您的心得反馈：
但果断warp shuffle的效率和直接带有volatile关键字的shared memory上的交换效率不同的。
我这里说明一下，避免您的结论误导其他会员（但这不代表对您的批判）：
（1）warp shuffle和shared memory均需要通过LSU执行相应指令的（请注意warp shuffle指令不是被SP执行的）
（2）warp shuffle的理论吞吐率是32操作/cycle/smx(或者按warp计算是1操作/cycle/smx)。warp shuffle在它的理论吞吐率下，可以在一个周期内完成32个4B数据的交换。而shared memory需要经过至少2步的写入-读取操作，甚至还需要进行地址计算的操作。
因为上文(1)(2)给出的的执行单元信息和吞吐率数据，warp shuffle一般比shared memory的交换效率要高2倍～3倍（分别对应无需使用SP进行地址计算和需要使用SP进行地址计算）。
您的测试没有得到速度上的提升，可能是因为您的kernel的数据交换操作不是瓶颈，从而无法凸显出来。但对于一些原来需要shared memory进行交换的算法，在升级到3.x的硬件后，如果可以使用warp shuffle，往往会得到较高的性能提升。
非常感谢您的心得反馈，此文完全不是对您的批评，只是简单的阐述一些客观资料，避免其他会员误解您，从而造成对您不满。
（此文信息来自非官方公开资料，仅供参考）
谢谢了，应该是数据交换在我的程序里不是瓶颈，所以影响不大。
不错，又学到了
lianuosun您好：
本版禁止无意义顶不相关的帖子的，请您注意。
不好意思，学校bbs泡惯了，忘了这茬了，以后注意。要不，这俩灌水删了吧。
我想問一下
loop1
shared ← global
loop2 i from 0~blockDim-1
reg ← shared[idx]
compute …
end of loop2
end of loop1
一般我的程式流程大概是這樣，在loop1先取好loop2要用的資料到shared中
然後在loop2計算時從shared拿資料到reg中再計算
照版主的說法warp shuffle可以提供32*4B的資料，那這樣不就跟shared最大提供128B的資料是一樣的帶寬
若不幸產生bank conflict 的話那帶寬則會在縮小些，那warp shuffle發生一些事情讓帶寬縮小
假若沒有的話那資料寫到reg去中在loop2計算時再用warp shuffle交換資料則會更有效率
又如果blockDim>warpSize時則這樣的方法就不能用，必須再回到用shared的方式存資料
iHakka您好：
不是很明白您的意思，不过先回答一下：
1：warp shuffle本身用于一个warp内部交换数据，和shared memory相比，这个机制可以一次性交换数据完毕，而shared memory需要先写入shared memory再读取出来，因此warp shuffle要快一些。这个前面横扫斑竹有详细解释。
2：bank conflict是使用shared memory的时候的概念，而warp shuffle没有这个概念。
3：由于您原文只有很简单的使用shared memory的示例代码以及我没看明白您这里的表述，所以无法继续回答，您可以根据warp shuffle的特点，根据您的需要判断一下是否可以使用。
4：warp shuffle本身就是一个warp内部交换数据的，超过一个warp都不行。
大致如此，您可以继续补充您的看法和信息。
祝您好运~