核函数如下：define N 4096*4096
global void Add(float *a, float *b, float c)
{
int tid=threadIdx.x+blockIdx.xblockdim.x；
while(tid<N)
{
c[tid]=a[tid]+b[tid];
tid+=blockDim.x*gridDim.x;
}
}
上面的程序，在核函数执行的时候是不是合并访问全局存储器啊？是不是每个线程都要访问一次全局存储器？16个线程一起访问？我想如果将全局内存上面的数据导入到共享内存上，然后在计算的时候访问共享内存，那速度岂不更快，但是共享内存不够大啊。
合并访问最简单的说，就是连续的32个线程，访问连续的内容。
在2.x上，默认您需要至少每个线程都访问连续的4B才能取得较好的效果。
在3.x上，默认您需要至少每个线程都访问连续的2B才能取得较好的效果。
对于您的 int tid=threadIdx.x+blockIdx.x*blockdim.x；来说，
然后使用[tid], 这个访问是合并的。
我的设备是2.0的，针对我上面的这个例子，你说的每个线程都访问连续的4B是什么意思呢？目前，我貌似只知道去编程，我发现很多内在的东西不懂啊！
int tid=threadIdx.x+blockIdx.x*blockdim.x；
while(tid<N)
{
c[tid]=a[tid/2]+b[tid-2];
tid+=blockDim.x*gridDim.x;
}
这算不算是合并访问呢？问的有点基础啊，还请大师不耐烦的解释啊
首先说，在上文的基础上，
[tid - 2]也是合并的。
其次，[tid / 2]这个也算合并的，但对cache的利用率可能会下降（也可能不会）。
不懂就多看看CUDA Programming Guide, 这个在最开头的一章就有详细讲述，
您阅读基本手册，比直接询问要好。
（因为论坛上我的回答不能像手册那样的详细，那样的娓娓道来。为何？手册可以用几百页说明，但是回帖却只能尽量简短。您觉得呢？）
谢谢指点，我先自己看，不懂的再请教您。
LZ你好，首先纠正一点，在2.x之后，GPU的最小执行单元是warp(32 threads)不再是half-warp(16 threads)了。(友情提示：LZ要是刚入门CUDA，建议阅读CUDA官方手册）
其次，关于合并访问，手册上面有详细说明：1.x必须同一个warp内的threads顺序访问对齐的显存块才能合并访问，比较严格。2.x可以同一个warp内的threads乱序访问对齐的显存块，也能合并访问。在有L1 cache的情况下，内存块大小为128B，即32*4B。
具体请参看手册。（如果我的回答给LZ带来了更大的疑惑，那么请暂时无视，并不影响LZ编程）
你好，你讲的这些我多少还是懂点的，这里的4B是不是就是线程访问的宽度，譬如线程0需要float a[0],线程1需要float a[1],其实在线程0访问内存的时候已经把a[1]也读进来了。我理解的对吗？求大神解释，我新手，大神勿笑话！！
LZ您好：
1：4B指的是4 BYTE，指线程访存的时候访问的数据类型的宽度，比如float类型就是4B的。
2：一个warp总是一起行动的，如果这一个warp的访存在满足某些条件（如顺序访问一段连续的空间），那么在硬件的支持下，可以一次完成，谓之合并访问。如果不满足一些条件，那么会分为多次完成，此时是非合并访问。
以及运行一个线程上的指令的一个SP自身是没有能力读取global memory的，需要整个warp为单位利用SM上的硬件单元去访问。
大致如此，希望此解释能为您带来些许清晰的认识。
祝您好运~~
soga，这样啊，了解了解。那么假如我例子中的float全部改成double，那整体的访问时间是不是耗的更长啊，因为每个double数据都得分两次去读啊？我是这样理解的，不知道能差多少。
LZ您好：
1：float换成double的话，如果都满足合并访问，那么在访存阶段所用的时间后者是前者的两倍。
2：以及，1：的说法并不是因为每个double数据分两次读取，而是因为数据量增加了一倍。这个很好理解，因为double数据每个数据是8B的，相当于两个4B的float，而且海斯连续存放的。而访存的时候硬件一次只能读取一段连续的空间，因此是每次访问半数的double，分两次读取完。而不是每次访问全数double的前一半或者后一半。
3：以及，double情况下，计算速度也会变慢，具体变慢程度和您的卡有关，可能是几分之一，也可能是十分之一左右。以及double数据也更耗费寄存器资源。
大致如此，供您参考。
祝您好运~~