share memory的赋值何时生效?

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都不行。

大致如此,您可以继续补充您的看法和信息。

祝您好运~