[attach]2983[/attach]
重新上图，类似于DIT-FFT运算，只有当每一级运算完成后，把该级运算后的数据作为下一级输入数据，这里只取了8个点，如果点数增加，级数也会增加，同一种颜色代表一个线程要取的两个数据，当级数增加同一种颜色要取的两个数之间的跨度也会增加，要怎么取数据才能满足合并访存啊？
楼主您好，因为此问题已经被总版主，给予过回复，我表示衷心支持。并无法提供其他不同意见和建议。
而回复您的帖子是表示我重视了您的问题，并仔细阅读了。以及您可以继续等他其他会员/版主/NV官方支持/总版主的回答。
如果跨度较大，或许共享内存也无法解决。楼主可以考虑一下texture纹理内存绑定的方式。对于不满足合并访问的情况，纹理内存一般能够带来访存性能的提升！
刚接触cuda不久，最近才开始逛论坛，头一次回复，说得不对的请轻拍。（另外是七天前的帖子，不知道是不是太久远了）
我觉得应该能满足global memory的合并访问。具体分析如下：
1、观察LZ的算法，应该每一级相互操作的两个数据间隔满足 ：
deta=2^(n-1) n为级数
2、由于每一级操作之后，数据需要同步，才能供下一级计算，故每一级应该设置为一个内核运行（当然可以写出一个各级通 用的内核）。
3、第一级的访问可以合并，很显然。
x1 x2 x3 x4 x5 x6 x7 x8
4、第二级的访问，可以这样理解global memory :
x1’ x2’
x3’ x4’
x5’ x6’
x7’ x8’
既将其理解为一个4*2行主存的矩阵（实质还是原来的数据，没有任何i变化）。
内核设置2个tid（当然，这是根据例子中数据少的情况），每个tid先读取 x1’~x2’，满足合并访问；再读取x3’~x4’,依然满足合并访问，依次类推。同一个tid访问的数据刚好是需要做运算的数据，
5. 第三极理解为如下形式：
x1’’ x2’’ x3’’ x4’’
x5’’ x6’’ x7’’ x8’’
用4个tid分两次合并访问读取。
6、当然这种方法在数据量较小和低层次情况下，不能达到合并访问的最大带宽，但是当数据量大、>=6的高层次（2^(6-1)=32 时，可以很好的实现合并访问。故可以根据实际情况，判断数据量。当数据量小、低层次的时候，采用整体读入shared memory的方式访问；数据量大、高层次的时候，采用如上方式访问。
非常感谢楼主给出如此详细的心得，你是这近1年来，本坛的难得的用心的人。真的。
感谢你的回复。
（此信息只表横扫千军对yixi的这种精神的极大赞美和衷心感动。
但不代表横扫千军理解或者不理解，赞同或者反对yixi的观点。）
感谢yixi同学发帖详细讨论~只要是详尽的原创内容，神马时候都不嫌晚的。
简要说下您的几点
1：是否需要每一级都启动kernel，这个要看规模，如果原问题规模以及特定的实现上，能在一个block内搞定，那么直接block内__syncthreads()也可以线程同步。规模大点的话，反复发射kernel即可。
2：您的理解可能有些偏差，比如“内核设置2个tid”是何含义，我不太明白;以及每个“tid先读取…"什么的也不明白。
其实tid一般是用来表示线程编号的，而相邻编号的一个warp的线程是一起访存的。是否合并，需要考虑此时访存的效率如何。
3：本问题的实际情况，其实不妨先用visual profiler跑跑看，看效果如何，再做定夺。
欢迎您常来论坛转转~
感谢版主的点评~
1、我的理解是：对于一大块数据n小块数据的模型，如果仅仅每个小块中需要数据同步，采用block内同步；如果整块数据需要同步，则应该采用同步前后分别发射内核的模式。
2、关于我说的“内核设置两个tid”这个说法；却是说的不好——我也只是根据他举例的数据而言的。像他举例中的数据规模，其实放在一个SM(SMX)中用寄存器去做计算就足够了，所以我觉得他本意应该是大规模数据。而对他数据访问方式的分析(2^(n-1))，对于大数据，刚好可以达到32的倍数——warp的倍数。对于同一个warp的tid，访问内存是顺序的，而且应该尽可能让其满足合并访问。
例如对于同一个swap中的tid:
A[tid]=B[tid}+C[tid];
其实质是每个tid先一起合并访问B[tid]，在一起合并访问C[tid],最后A[tid](当然A、B、C的顺序依赖于编译器汇编结果）。
（这里我也有个小疑问，本来想自己发个帖子问的，现在就在这说一下吧：对于计算能力2.1的机器，貌似都是48SP/SM,而同期的最小执行模型已经是warp。不在是half warp了，那48这个数字不是很尴尬？指令发射神马的，那16个sp怎么办？）
我说的“设置两个tid”本意也是说这个。因为对于高层次，这个tid就会是32的倍数，既是warp。利用这一点就可以合理安排计算，以达到数据能合并访问的目的。
都是linux惹的祸…… warp已经更正过来了……
感谢yixi同学的讨论
1：是的，不过一般而言是从thread的角度考虑的，如果一个block内部的thread需要同步，那么就__syncthreads()；而如果整个kernel里面的线程需要同步，那么只有重启kernel。这里的同步是thread行为的同步，是一群thread在等另外一群。单纯说data的话，data自身是被动的，还需要考虑他们是如何被thread访问的。
假如极端情况，不考虑效率，也可以用一个block甚至一个thread访问和处理大量数据，此时依然只需要__syncthreads()即可。当然这是极端情况，一般不会遇到。
2：我觉得您的叙述里面有两个术语用的不甚合适。第一个是tid，一般习惯上认为tid是thread id的缩写，也就是线程编号的意思。但一般表示线程的时候不说某个tid如何如何。第二个是“swap”，根据您上下文的含义，应当是“warp”（线程束）。
顶楼的问题，简单讲，在级数较高的时候自然是合并的；在级数较低的时候，能否优化为合并访问的情况，取决于编译之后的行为或者cache的行为，有不确定的因素，因此建议跑跑visual profiler看看结果再说。这是个人的建议。
3：关于您提出的，SM 2.1架构下，一个SM里面有48个SP，不是warp线程数的整倍数，会怎么样？简单地说，不会怎么样，会很好地运作，warp scheduler会搞定这些。具体细节是NV不公布的，您也可以继续向NVIDIA原厂支持求教。
此外，如果您还有其他问题需要讨论，欢迎另外开贴。
祝您一切顺利~
修改掉了个别错别字。