const char programSource[] =
"__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
"{"
" int gid = get_global_id(0);"
"for(int i=0; i<10; i++){"
" a[gid] = b[gid] + c[gid];}"
"}";
上面的内核是每个循环执行十次的向量加法。我已经使用了编程指南和堆栈溢出来确定全局内存的工作方式，但是如果我以一种很好的方式访问全局内存，我仍然无法通过查看代码来确定。我正在以一种连续的方式访问它，并且以一种一致的方式猜测。该卡是否为阵列a，b和c加载128kb的全局内存块？然后，是否每处理32个gid索引就为每个数组加载128kb的块一次？ （4 * 32 = 128）看来我没有浪费任何全局内存带宽了吗？
顺便说一句，计算机性能分析器显示的gld和gst效率为1.00003，这似乎很奇怪，我认为如果我所有的存储和负载都合并在一起，那么效率仅为1.0。高于1.0的情况如何？
最佳答案
是的，您的内存访问模式几乎是最佳的。每个半扭曲都访问16个连续的32位字。此外，访问是64字节对齐的，因为缓冲区本身是对齐的，并且每个半扭曲的startindex是16的倍数。因此，每个半扭曲将生成一个64Byte事务。因此，您不应该通过无节制的访问来浪费内存带宽。
由于您在最后一个问题中询问了示例，因此让我们将此代码修改为其他代码（最佳访问模式较差（由于循环实际上并没有做任何事情，因此我将忽略该代码））：
kernel void vecAdd(global int* a, global int* b, global int* c)
{
int gid = get_global_id(0);
a[gid+1] = b[gid * 2] + c[gid * 32];
}
首先，让我们看看这在计算1.3（GT200）硬件上如何工作
对于写操作，这将生成一个稍微不理想的模式（遵循由其id范围和相应的访问模式标识的半扭曲）：
gid | addr. offset | accesses | reasoning
0- 15 | 4- 67 | 1x128B | in aligned 128byte block
16- 31 | 68-131 | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
32- 47 | 132-195 | 1x128B | in aligned 128byte block
48- 63 | 196-256 | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
因此，基本上，我们浪费了大约一半的带宽（奇数半扭曲的访问宽度再少一倍，则无济于事，因为它会产生更多访问，而这并不浪费更多字节那么快）。
对于从b的读取，线程仅访问数组的偶数元素，因此对于每个半扭曲，所有访问都位于128byte对齐的块中（第一个元素位于128B边界，因为对于该元素，gid是16的倍数=>索引是32的倍数（对于4个字节元素），这意味着地址偏移量是128B的倍数）。 accesspattern扩展了整个128B块，因此这将对每个半扭曲执行一次128B传输，这又使一半带宽受限。
对c的读取会产生最坏的情况之一，其中每个线程在其自己的128B块中进行索引，因此每个线程都需要自己的传输，这一方面有点像是序列化的情况（尽管不如正常情况那么糟糕，因为硬件应该能够使传输重叠）。更糟糕的是，这将为每个线程传输32B块，浪费了7/8的带宽（我们访问4B /线程，32B / 4B = 8，因此仅使用1/8的带宽）。由于这是朴素矩阵转置的访问模式，因此强烈建议使用本地内存（根据经验来讲）。
计算1.0（G80）
在这里，唯一可以创建良好访问权限的模式是原始模式，示例中的所有模式都将创建完全不分时段的访问，浪费7/8的带宽（32B传输/线程，请参见上文）。对于G80硬件，每一次半warp中的第n个线程都不访问nth元素的访问都会创建这样的未分批访问
计算2.0（Fermi）
在这里，每次对内存的访问都会创建128B事务（收集所有数据所需的事务很多，因此在最坏的情况下为16x128B），但是会缓存这些事务，从而使数据传输位置变得不那么明显。目前，假设高速缓存足够大以容纳所有数据，并且没有冲突，因此每条128B高速缓存行将最多传输一次。让我们进一步假设半扭曲的序列化执行，因此我们具有确定性的缓存占用率。
对b的访问仍将始终传输128B块（核心响应存储区中没有其他线程索引）。对c的访问将为每个线程生成128B传输（可能是最差的访问模式）。
对于访问a的操作如下（对它们的读取类似于现在进行处理）：
gid | offset | accesses | reasoning
0- 15 | 4- 67 | 1x128B | bringing 128B block to cache
16- 31 | 68-131 | 1x128B | offsets 68-127 already in cache, bring 128B for 128-131 to cache
32- 47 | 132-195 | - | block already in cache from last halfwarp
48- 63 | 196-259 | 1x128B | offsets 196-255 already in cache, bringing in 256-383
因此，对于大型阵列，理论上对a的访问几乎不会浪费带宽。
对于此示例，实际情况当然不尽如人意，因为对c的访问将很好地破坏缓存
对于探查器，我认为效率超过1.0只是浮点精度的结果。
希望能有所帮助