简化我的OpenCL核函数中全局内存的运行过程。

Question

简化我的OpenCL核函数中全局内存的运行过程。

4

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];}"
        "}";

上面的内核是每个循环进行十次向量加法。我已经使用编程指南和堆栈溢出来了解全局内存的工作原理，但是仍然无法通过查看我的代码确定是否以良好的方式访问全局内存。我正在以连续的方式访问它，并猜测以对齐的方式访问它。卡片是否会加载128kb的全局内存块用于数组a、b和c？然后它是否为每32个gid索引处理一次每个数组的128kb块？（4*32=128）这似乎意味着我没有浪费任何全局内存带宽，对吗？

另外，计算分析器显示gld和gst效率为1.00003，这似乎很奇怪，如果我所有的存储和加载都是协同的，我认为它只会是1.0。为什么它超过1.0呢？

- smuggledPancakes

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Grizzly · Accepted Answer

是的，你的内存访问模式几乎是最优的。每个半warp正在访问16个连续的32位字。此外，访问是64字节对齐的，因为缓冲区本身是对齐的，每个半warp的起始索引是16的倍数。因此，每个半warp将生成一个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范围识别半线程束和相应的访问模式）：

   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

基本上，我们浪费了大约一半的带宽（奇数半warp的访问宽度不到两倍并没有帮助太多，因为它会生成更多访问，这不如浪费更多字节更快）。

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

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

Compute 1.0（G80）

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

Compute 2.0（Fermi）

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

对于访问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的效率仅仅是浮点不准确性的结果。

希望这可以帮到你。