合理化我的简单 OpenCL 内核中有关全局内存的情况

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以上怎么样？

是的，您的内存访问模式几乎是最佳的。每个 halfwarp 访问 16 个连续的 32 位字。此外，访问是 64 字节对齐的，因为缓冲区本身是对齐的，并且每个 halfwarp 的起始索引是 16 的倍数。因此每个 halfwarp 将生成一个 64 字节事务。因此，您不应该通过未合并的访问来浪费内存带宽。

既然您在上一个问题中要求提供示例，那么让我们修改此代码以用于其他（不太理想的访问模式（因为循环实际上没有做任何事情，我将忽略它）：

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) 硬件上的工作原理

对于对 a 的写入，这将生成稍微不太理想的模式（遵循由其 id 范围和相应的访问模式标识的 halfwarp）：

   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

所以基本上我们浪费了大约一半的带宽（奇数 halfwarp 的访问宽度加倍并没有多大帮助，因为它会生成更多的访问，可以说这并不比浪费更多字节更快）。

对于从 b 读取，线程仅访问数组的偶数元素，因此对于每个 halfwarp，所有访问都位于 128 字节对齐的块中（第一个元素位于 128B 边界，因为对于该元素，gid 是 16 的倍数=>索引是32的倍数，对于4字节元素，这意味着地址偏移量是128B的倍数）。访问模式延伸到整个 128B 块，因此这将为每个 halfwarp 执行 128B 传输，再次减少一半的带宽。

从 c 中读取会产生最坏的情况之一，其中每个线程都在自己的 128B 块中索引，因此每个线程都需要自己的传输，一方面有点序列化场景（尽管不像正常情况那么糟糕，因为硬件应该能够重叠传输）。更糟糕的是，这将为每个线程传输 32B 块，浪费 7/8 的带宽（我们访问 4B/线程，32B/4B=8，因此只利用了 1/8 的带宽）。由于这是朴素矩阵转置的访问模式，因此强烈建议使用本地内存进行这些访问模式（根据经验）。

计算 1.0 (G80)

这里唯一能够创建良好访问的模式是原始模式，示例中的所有模式都将创建完全未合并的访问，浪费 7/8 的带宽（32B 传输/线程，见上文）。对于 G80 硬件，halfwarp 中第 n 个线程不访问第 n 个元素的每次访问都会创建此类未合并的访问

计算 2.0（费米）

在这里，每次对内存的访问都会创建 128B 事务（收集所有数据所需的数量，因此在最坏的情况下为 16x128B），但是这些事务被缓存，使得数据将传输到何处不太明显。目前我们假设缓存足够大，可以容纳所有数据并且不存在冲突，因此每个 128B 缓存行最多会传输一次。让我们进一步假设 halfwarp 是串行执行的，因此我们有确定性的缓存占用。

对 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 的效率只是浮点不准确的结果。

希望有帮助