1

我将提高 OCL 内核性能，并想阐明内存事务如何工作以及哪种内存访问模式真正更好（以及为什么）。内核被输入了 8 个整数的向量，这些向量被定义为 array: int v[8]，这意味着，在进行任何计算之前，必须将整个向量加载到 GPR 中。所以，我相信这段代码的瓶颈是初始数据加载。

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

目标硬件是 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];

（实际上表现为）

    *(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 和通过操作系统时间测量），但连续存储数据的内核执行速度出奇地慢。两个内核的总时间大致相同。在我看来，两者必须至少以相同的方式表现——要么更慢，要么更快，但这些相反的结果似乎无法解释。

我的第三个问题是：谁能解释这样的结果？还是我做错了什么？（或者完全错误？）

4

2 回答 2

0

好吧，并没有真正回答我所有的问题，但是在浩瀚的互联网上发现的一些信息更清楚地把事情放在一起，至少对我来说是这样（不像上面提到的 AMD 优化指南，它似乎不清楚，有时令人困惑）：

«硬件执行了一些合并，但它很复杂......
经线中的内存访问不一定必须是连续的，但它们落入多少个 32 字节全局内存段（和 128 字节 l1 缓存段）确实很重要。内存控制器可以在单个事务中加载这 32 字节段中的 1、2 或 4 个，但这是通过 128 字节高速缓存行中的高速缓存读取的。
因此，如果 warp 中的每个通道都加载 128 字节范围内的随机字，则不会受到惩罚；这是 1 笔交易，阅读效率很高。但是，如果 warp 中的每个通道以 128 字节的步幅加载 4 字节，那么这非常糟糕：加载了 4096 字节但只使用了 128 字节，导致效率约为 3%。»

因此，就我而言，数据始终是连续的时如何读取/存储数据并不重要，但是加载向量部分的顺序可能会影响编译器随后的命令流（重新）调度。
我还可以想象，较新的 GCN 架构可以进行缓存/合并写入，这就是为什么我的结果与优化指南提示的结果不同的原因。

于 2017-10-12T14:53:40.023 回答
0

查看AMD OpenCL 优化指南中的第 2.1 章。它主要关注老一代卡，但 GCN 架构并没有完全改变，因此仍应适用于您的设备（北极星）。

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

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

此致，

迈克尔

于 2017-10-06T16:50:36.460 回答