向量化数据的OpenCL（AMD GCN）全局内存访问模式：跨步与连续

我将改进OCL内核性能，并且想要阐明内存事务如何工作以及哪种内存访问模式真的更好（以及为什么）。内核提供了8个整数的向量，这些整数被定义为数组：int v [8]，这意味着，在进行任何计算之前，必须将整个向量加载到GPR中。所以，我相信这个代码的瓶颈是初始数据加载。向量化数据的OpenCL（AMD GCN）全局内存访问模式：跨步与连续

首先，我考虑一些理论基础知识。

目标硬件是Radeon RX 480/580，具有256位GDDR5存储器总线，突发读/写事务有8个字粒度，因此，一个存储器事务读取2048位或256字节。也就是说，我相信，什么CL_DEVICE_MEM_BASE_ADDR_ALIGN是指：

Alignment (bits) of base address: 2048.

因此，我的第一个问题：什么是128字节的缓存行的物理意义？它是否保留通过单个突发读取获取的数据部分，但没有真正请求？如果我们要求32或64个字节，剩余部分会发生什么 - 因此，剩余的部分超过了缓存行的大小？（我想，它将被丢弃 - 然后，哪个部分：头部，尾部...？）

现在回到我的内核，我认为缓存不起作用在我的情况下，因为一个突发读取64个整数 - >理论上一个内存事务可以一次输入8个工作项，没有额外的数据要读取，并且内存总是合并的。

但是，我可以把我的数据有两个不同的访问模式：

1）连续

a[i] = v[get_global_id(0) * get_global_size(0) + i];

（至极实际上perfomed为）

*(int8*)a = *(int8*)v;

2）交错

a[i] = v[get_global_id(0) + i * get_global_size(0)];

我期望在我的情况下连续会更快，因为如上所述，一个内存事务可以完全填充8个工作项与数据。但是，我不知道，计算单元中的调度程序如何在物理上工作：是否需要为所有SIMD通道准备好所有数据，或者只需要4个并行SIMD元素的第一部分就足够了？尽管如此，我认为只要CU可以独立执行命令流，它就足够聪明，可以首先完全提供至少一个CU的数据。在第二种情况下，我们需要执行8 * global_size/64个事务来获取完整的向量。

所以，我的第二个问题：我的假设是正确的吗？

现在，练习。

实际上，我将整个任务分成两个内核，因为一个部分的注册压力低于另一个，因此可以使用更多的工作项目。所以我首先用模式演示了如何在内核之间转换时存储的数据（使用vload8/vstore8或转换为int8给出相同的结果），结果有点奇怪：以相邻方式读取数据的内核工作速度快10％ CodeXL和操作系统时间测量），但连续存储数据的内核执行速度惊人地更慢。两个内核的总体时间大致相同。在我的想法中，两者必须至少以相同的方式行事 - 要么更慢，要么更快，但是这些相反的结果似乎无法解释。

而我的第三个问题是：任何人都可以解释这样的结果吗？或者可能是我做错了什么？（或完全错误？）

来源

2017-10-06 qpdb

看看AMD OpenCL Optimization Guide的2.1章。它主要关注老一代卡，但GCN架构没有完全改变，因此仍然适用于您的设备（polaris）。

一般而言，AMD卡具有多个内存控制器，在每个时钟周期内分配内存请求。例如，如果您以列专业而非行专业逻辑访问您的值，则性能会更差，因为请求会发送到同一个内存控制器。（按列专业我的意思是矩阵的一列由当前时钟周期中执行的所有工作项一起访问，这就是你所说的合并vs交错）。如果您在一个时钟周期内访问一行元素（意味着合并）（意味着所有工作项访问同一行内的值），那么这些请求应分配给不同的内存控制器，而不是相同。

关于对齐和缓存行大小，我想知道这是否真的有助于提高性能。如果我处于你的情况，我会尝试看看我是否可以优化算法本身，或者如果我经常访问这些值，将它们复制到本地内存是有意义的。但是，如果没有任何有关你的内核执行的知识，很难再说。

最好的问候，

迈克尔

来源

2017-10-06 16:50:36 MichaelE1000

谢谢你的回答。但是我不是在谈论合并vs交错访问。可能是我的着作不那么清楚，但访问总是合并 - 区别仅在于读取数据矢量明智与元素明智。为了澄清一点，我纠正了这个问题。 – qpdb

@qpdb从内核的角度来看，你称之为连续的东西是连续的，并且在给定的周期内从内存的角度进行交错，因此读取每个工作项的第一个元素可以缓存剩余的数据。但是在写作时，没有这种行为，所以它变得更慢。由于读/写调度器（或者读/写组合的任何部分）可以提供n个工作项，这些工作项可以在大量相邻元素上统一读写，所以称为“交错”的内容在内存的给定周期内实际上是连续的。 –

顺便说一句，再次感谢指出文档。我从那里了解到：“南岛设备不支持合并写入;但是，工作组内的连续地址提供了最佳性能。”这个信息看起来很奇怪，因为我的实验给出了完全不同的结果。还是我完全相反地理解“合并”的整个概念？ – qpdb

嗯，不是真的回答所有我的问题，但在互联网的浩瀚发现了一些信息，把东西放在一起更清晰的方式，至少对我来说（不同于上述AMD优化指南，这似乎不清楚，有时混淆）：

«硬件执行一些合并，但它很复杂...
变形内存访问不一定是连续的，但它涉及多少32字节的全局内存段（和128字节的l1高速缓存段）。内存控制器可以在单个事务中加载这些32字节段中的1,2或4个，但是这是通过128字节缓存行中的缓存读取的。
因此，如果warp中的每个通道加载128字节范围内的随机字，则不存在惩罚;这是1笔交易，阅读完全有效。但是，如果warp中的每条通道都加载了4个字节，并且步长为128个字节，那么这非常糟糕：加载了4096个字节，但仅使用了128个字节，因此效率为〜3％。情况并非如此，数据读取/存储的方式始终是连续的，但矢量部分的加载顺序可能会影响编译器后续的命令流（重新）调度。
我也可以想象，较新的GCN架构可以执行缓存/合并写入，这就是为什么我的结果与该“优化指南”提示的结果不同。

来源

2017-10-12 14:53:40 qpdb

向量化数据的OpenCL（AMD GCN）全局内存访问模式：跨步与连续

回答

相关问题