我是新手，刚开始学习cuda，对于合并访问不太理解，请各位前辈解惑，感激不尽。
所谓合并访问就是尽量要求相邻的线程访问相邻的地址空间；
（每一个线程都能自由访问任何地址？为什么相邻的线程要访问相邻的地址空间？这样为什么能够加速？）
在计算能力1.0和1.1的设备上，合并访问能够在一次存储器访问中访问最多达16个数据；在计算能力1.2和1.3的设备上，合并访问一次最多能够访问32个数据。
（一次存储器访问读取16（32）个数据，一次指的是什么？应该不是指一个线程吧？）
还有合并访问这一块一般会提到32、64、128位，这是在说什么？另外也会提到half-warp，它和合并访问有什么关系？
另外附上一个例子，求解释
未合并访问
global void accessFloat3(float3 d_in, float3 d_out)
{
int index = blockIdx.x * blockDim.x + threadIdx.x;
float3 a = d_in[index];
a.x += 2;
a.y += 2;
a.z += 2;
d_out[index] = a;
}
合并访问
global void accessInt3Shared(float g_in, float g_out)
{
int index = blockIdx.x * blockDim.x + threadIdx.x;
shared float s_data[2563];
s_data[threadIdx.x] = g_in[index];
s_data[threadIdx.x+256] = g_in[index+256];
s_data[threadIdx.x+512] = g_in[index+512];
__syncthreads();
float3 a = ((float3)s_data)[threadIdx.x];
a.x += 2;
a.y += 2;
a.z += 2;
((float3)s_data)[threadIdx.x] = a;
__syncthreads();
g_out[index] = s_data[threadIdx.x];
g_out[index+256] = s_data[threadIdx.x+256];
g_out[index+512] = s_data[threadIdx.x+512];
}
LZ您好：
不同硬件架构下，合并访问的要求是不同的，具体到每个特定的硬件架构下的要求细节，请以手册说明为准。
CUDA SIMT机制执行的最小单位是一个warp，half-warp概念是早期1.x版本硬件描述其访存等硬件行为所提出的，只对1.x硬件优化有用，同时也并不与“最小执行单位是warp”相矛盾。由于1.x版本硬件过时已久，基本无需再讨论。所以请您无视所有half-warp的相关说法，只关注主流硬件情况即可。
下面，我将对当前主流版本硬件的合并访问做简单描述，以增进您的理解。
1：在当前的硬件水平下（SM2.X和SM3.X硬件），一个warp的线程在一次global访存的时候，如果能满足局部性（即访问一段特定长度存储空间内部数据），那么粗略地说，是基本满足合并访问的。这里并不要求顺序性，各个线程可以以任意顺序访问这段存储空间内部的数据。
2：至于为何要求满足局部性，原因如下：因为一个warp的线程在执行的时候总是同时的，在访存的时候，也是同时利用SM内部的硬件一起实现（每个CUDA CORE并无能力自身实现global访存）。而SM内部的硬件在访存的时候，只能一次性从global memory中（实际是从cache中）拿到一段连续的数据，如果warp中所有的访存需求都在这一段中，那么只需要一次访存即可获得warp内各个线程所需要的数据，此时是合并的，也是最快的。如若不然，则需要多次从global memory中读取数据，自然就慢了。
3：上述讨论也解释了“一次存储器访问读取32个数据”的含义，因为这恰好就是一个warp每个线程获得自己的数据。（以及1.x硬件，每次访问只能为半个warp的线程服务，所以是16个。1.x硬件不详细展开了）同时需要注意的是，如果您访问的数据时8B的double或者更宽的原生类型，因为数据量实际是倍增的，所以一个warp的访存也需要多次才能实现，但此时依然是合并的。详情请参阅手册。
4：关于“一般会提到32、64、128位”请您提供具体的讨论内容，我无法凭空告诉您原作者的意图。
根据2#的讨论，继续来说一下这两段代码中的访存部分。
您在代码中使用了float3类型，根据手册，这个类型的读写操作将被拆分为3次float读写实现。
在您的第一段代码中，float3读写拆分得到3次float读取/写入。对于每次float读取/写入相邻线程都错开3个float的间距，这样需要整体读取/写入3次才能满足warp内所有的线程的需求，并同时丢弃了2/3的无效数据。
在您的第二段代码中，每次对应的float读取/写入，相邻线程之间是连续的，这自然是合并的。
大致如此，祝您好运~
首先对版主的帮助表示深深的感谢，这么用心仔细。对于版主所说的
《SM内部的硬件在访存的时候，只能一次性从global memory中（实际是从cache中）拿到一段连续的数据，如果warp中所有的访存需求都在这一段中，那么只需要一次访存即可获得warp内各个线程所需要的数据，此时是合并的》
这样一个概念也是清楚的，但是对于实际的拿到一段连续的数据还是不太清楚，比如说
对一个float[32]的数组，一个warp去访问，每个线程得到一个float数据，那也就是说一次性从global中取出这32个数据，然后分配到各个线程中去，这是合并访问。
同样对一个float[64]，每个线程两个float数据，这也是合并访问吧？每个线程错开2个float的间距，这和您上面讲到的float3每个线程错开3个float的间距不满足合并访问有什么区别呢？
另外您上面对于
float3读写拆分得到3次float读取/写入。对于每次float读取/写入相邻线程都错开3个float的间距，这样需要整体读取/写入3次才能满足warp内所有的线程的需求，并同时丢弃了2/3的无效数据。
的讲解，我不是很明白，整体读写3次才能满足warp内的所有县城需求，为什么是3次，这个3次是怎么来的，第一次读的是哪段数据，第二次三次又是哪一段数据，丢弃了2/3的无效数据指的是什么？
对于合并访问的例子，每个线程读写数据是连续的，但是线程内的数据是不连续的
s_data[threadIdx.x]
s_data[threadIdx.x+256]
s_data[threadIdx.x+512]
上面几个数据在线程内部并不连续，这对合并访问没有影响吗？
最后再次怀着感激之心对版主说声谢谢，请原谅我那混沌的大脑。
LZ您好：
1：“对一个float[32]的数组，一个warp去访问，每个线程得到一个float数据，那也就是说一次性从global中取出这32个数据，然后分配到各个线程中去，这是合并访问。同样对一个float[64]，每个线程两个float数据，这也是合并访问吧？每个线程错开2个float的间距”
对您后面这个float[64]，这个其实取决于写法。我来大致写个例子供您参考：
假定tid是算好的线程编号现在考虑一个warp的情况，tid为0~31，cat,dog是各个线程私有的寄存器变量，float是传入的global memory的指针，对应存储空间为64个float数。
如果您kernel写为：
cat=float[tid];
dog=float[tid+32];
考虑这个warp里面各个线程的工作：tid==0的线程在第一句访存的时候访问float[0],tid==1的线程访问float[1]，…tid==31的线程访问float[31]。
所以一个warp访问了float这个数组的前32个数据，这32个数据是连续存放的，而且总长度也满足一次读入的需求，所以可以一次性拿到数据。
给dog赋值的过程类似。
这是两次合并的访问。
如果写成：
cat=float[tid2];
dog=float[tid2+1];
此时考虑这个warp各个线程的行为：
给cat赋值的时候，tid==0的线程访问float[0]，tid==1的线程访问float[2]，…，tid==31的线程访问float[62]。这里面实际需要的都是偶数位置的元素，但是SM从global memory/cache读取的时候，只能读取连续的内容，这样将读取全部的64个元素（需要读两次），再丢弃奇数位置的数据。
对dog赋值的时候，也是类似，需要读取全部的数据，并丢弃偶数位置的数据。
这就是不合并的访问。
需要说明的3点：
a：fermi和kepler有L2 cache，所以后一次读取的时候如果之前读取的64个元素的数据还在cache中，那么后一次读取的代价要小很多，这比1.x硬件是一个显著的改进。1.x硬件无cache，需要老老实实从global memory中读取。
b：后面这种读取可以利用一些小的trick优化为合并访问的，不过这个先不谈。
c：float3 根据手册说法，会转化为3次float的读取，所以您可以根据上述float读取的例子理解为何需要读取3次，并丢弃2/3的数据。
2：“对于合并访问的例子，每个线程读写数据是连续的，但是线程内的数据是不连续的
s_data[threadIdx.x]
s_data[threadIdx.x+256]
s_data[threadIdx.x+512]
上面几个数据在线程内部并不连续，这对合并访问没有影响吗？ ”
其实这个恰恰是传统的单线程串行思维和SIMT多线程并行行为的差别所在。
您可以这样简单考虑：
对于传统的单线程的代码，您需要维护代码访存中连续访存的连续性/局部性，这样有助于CPU一次性从memory/cahce中拿到连续的数据并充分利用。
但是对于CUDA 的SIMT执行方式，最小的执行单位是一个warp，SM一样是一次性从global memory/cache中拿到一段连续的数据，但是这段数据是给一个warp的线程使用的，而不是给一个线程前后多个连续访存使用的。
所以，这要求这个warp的线程对于同一次访存，访问一段连续的存储空间（以及有一定对齐要求，这个不细说），当前的主流GPU可以不要求线程访存是顺序的，即可以tid==0的线程访问float[5]，tid==1的线程访问float[18]这种，并且无额外开销。
大致如此，祝您编码顺利~
感谢版主的悉心帮助，感觉已经明白了，thank you！