理顺什么是关于全局内存在我简单的OpenCL内核会

StackOverflow https://stackoverflow.com/questions/3857981

  •  27-09-2019
  •  | 
  •  

题 

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）这似乎是那么我并不浪费任何全局存储器带宽正确？

BTW，计算探查节目1.00003一GLD和GST效率，这似乎不可思议，我认为这将只是1.0如果我所有的存储和加载被合并。它是如何在1.0以上？

有帮助吗？

解决方案

是你的内存访问模式是相当多的优化。每个halfwarp正在访问16个连续的32位字。此外，该访问被64字节对齐，由于缓冲器自身对齐并为每个halfwarp从startIndex是16的倍数。因此，每个halfwarp将生成一个交易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）硬件

有关的写入到这将产生一个稍微unoptimal图案（以下可以通过ID范围标识的halfwarps和相应的访问模式）：

   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

所以基本上我们正在浪费我们大约有一半的带宽（对于奇halfwarps了不到一倍访问宽度，因为它会产生更多的访问并没有多大帮助，这是不是更快那么浪费更多的字节可以这么说）。

有关的从b中的线程访问所述阵列的只有偶数元素读取，所以对于每个halfwarp所有访问位于一个128字节对齐块（第一个元素是在128B边界，因为该元素的GID是的倍数16 =>索引为32的倍数，4层字节的元件，这意味着地址偏移是128B的倍数）。在整个128B块accesspattern伸展，所以这会为每一个halfwarp 128B传送，再次缩腰一半的带宽。

在为坏读取由C生成的最坏的情况，其中每个线程索引在它自己的128B块，所以每一个线程需要它自己的传输，其中之一一方面是有点系列化场景中的一个（虽然不是很如normaly，由于硬件应该能够重叠传送）。最新更糟糕的是这样的事实，这将传送块32B为每个线程中，带宽的浪费7/8（我们访问4B /线程，32B / 4B = 8，因此，只有1/8的带宽被利用）。由于这是幼稚matrixtransposes的accesspattern，这是非常可取做那些使用本地存储器（从经验来讲）。

<强>计算1.0（G80）

下面将创建一个良好的访问的唯一模式是原始的，在实施例的所有模式将创建完全未聚结的访问，浪费带宽的7/8（32B传递/线程，见上文）。对于G80硬件，其中在一个halfwarp第n个线程不访问的第n个元素创建这样未聚结的访问每个接入

<强>计算2.0（费米）

下面对内存的每次访问创建128B交易（多达必要收集的所有数据，所以16x128B在最坏的情况下），但是这些被缓存起来，使其不那么明显，其中的数据将被转移。就目前而言，假设该高速缓存足够大，以容纳所有数据并没有冲突，所以每一个128B的高速缓存行会最多一次转移。让furthermoe假设halfwarps的序列化执行，所以我们有一个确定性高速缓存占用。

访问B就仍然总是转移128B块（没有其他线程索引在coresponding memoryarea）。访问到c将产生每个线程128B传输（最差accesspattern可能的）。

有关访问到它是下面的（处理它们等读取的时刻）：

   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

因此，对于大型阵列的访问到会浪费几乎没有带宽理论上。 对于这个例子，现实当然是不太好，因为访问到c将毁掉缓存很漂亮

有关探查我假设超过1.0的效率是简单地产生浮点inaccurencies的。

希望帮助

许可以下： CC-BY-SA 和 归因
不隶属于 StackOverflow
scroll top