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

上面的内核是一个向量加法,每个循环执行十次。我已经使用编程指南和堆栈溢出来弄清楚全局内存是如何工作的,但是如果我以一种好的方式访问全局内存,我仍然无法通过查看我的代码来弄清楚。我以连续的方式访问它,并且我以一致的方式猜测。卡是否为数组 a、b 和 c 加载了 128kb 的全局内存块?然后它是否为每处理 32 个 gid 索引加载一次每个数组的 128kb 块?(4*32=128) 看来我并没有浪费任何全局内存带宽,对吧?

顺便说一句,计算分析器显示 gld 和 gst 效率为 1.00003,这看起来很奇怪,我认为如果我所有的存储和负载都合并,它只会是 1.0。1.0以上怎么样?

4

1 回答 1

12

是的,您的内存访问模式非常理想。每个 halfwarp 访问 16 个连续的 32 位字。此外,访问是 64 字节对齐的,因为缓冲区本身是对齐的,并且每个半束的 startindex 是 16 的倍数。所以每个半束将生成一个 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

所以基本上我们浪费了大约一半的带宽(奇数半扭曲的访问宽度小于两倍的访问宽度并没有太大帮助,因为它会产生更多的访问,这并不比浪费更多的字节更快)。

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

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

计算 1.0 (G80)

这里唯一可以创建良好访问的模式是原始模式,示例中的所有模式都将创建完全未合并的访问,浪费 7/8 的带宽(32B 传输/线程,见上文)。对于 G80 硬件,半经线中的第 n 个线程不访问第 n 个元素的每次访问都会创建这种未合并的访问

计算 2.0(费米)

在这里,每次访问内存都会创建 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 的效率仅仅是浮点不准确的结果。

希望有帮助

于 2010-10-04T18:46:37.187 回答