优化 OpenCL 缓冲区写入？

Question

我在写入缓冲区时遇到了主要瓶颈。

我想做的很简单。 首先，我使用了两个全局 id（我使用的是 image2d）。每个线程读取 9 个像素值，位置 (x,y) 处的像素及其 8 个相邻像素，基本上是一个 3x3 方形块。 这项工作由每个线程完成。现在，我计算了一些值，我想将每个线程的结果写入输出缓冲区。

每个线程产生 64 个值，我将它们写入输出缓冲区，这意味着输出缓冲区的大小为 (rows*cols*64)。 我还想支持最多支持 640 个值的计算，但显然每个线程都无法将 640 个值写入缓冲区，因为需要 VRAM。

我必须说线程写入不同的位置，没有覆盖，也就是说会有 64*number_of_threads = 64*global_id(0)*global_id(1) = 64*rows*cols 值。

这是我代码中的一个主要瓶颈，我的意思是写入 64 个值，我认为这与内存带宽有关，但我不太确定。

我该怎么做才能让每个线程高效地计算 64 个值并将其写入输出缓冲区？这不可能吗？

我的GPU是rx 480 4gb，我知道(rows*cols*64)大小有时候可能太大了，放不下VRAM，但是就算放好，写起来也慢，我觉得带宽很显卡高？

还有其他两个输出缓冲区，但它们的大小要小得多，所以我们可以忽略它们。

总而言之，这段代码的作用是 1) 读取一个9像素的方块，中间的一个是当前值。

2) 将 8 个邻居乘以当前值，我们得到每个像素的 8 个值。

3) 将 8 个邻居写入邻居缓冲区。

4) 将 8*8 值写入 Rx 缓冲区。此缓冲区“模拟” x_* x_^T 结果，即相邻值的 (8x1)x(1x8) 矩阵相乘。

请注意，我正在以“转置形式”写入输出缓冲区，即位置 (x,y) 处的每个线程在 (y,x)、(y+1,x) 处连续写入 64 个值。 ..(y+63,x) 这是因为：

1) 这是最快的方法！我写成 (x,y) -> (x+1,y),...(x+63,y) 的版本肯定更慢。

2) 我需要这种形式，因为我使用的是 ArrayFire 库，该库需要加载缓冲区，但它会以行优先顺序消耗缓冲区并将内容以列优先顺序放入其数组中，即这样就不需要转置数组（这将使用大量的 vram 副本）

Answer 1

首先，由于您没有明确提及，我会指出，如果可以，请使用您的 GPU 制造商的分析工具验证您的瓶颈。即使某件事看起来是一个瓶颈，它也可能是一个红鲱鱼。

但是，听起来全局内存写入可能是您的内核中的一个问题。由于您没有提供任何细节，我只能指出一些需要注意的一般事项：

1。内存布局

您设置中的每个工作项似乎在内存中连续写入了 64 个值。这意味着每个工作项都将写入不同的缓存行，这几乎肯定不是最佳的。如果您可以更改输出的内存布局，请尝试安排它，使工作项同时写入 相邻 内存位置。

例如，您目前可能有：

uint output_index = 64 * (get_global_size(0) * get_global_id(1) + get_global_id(0));
for (unsigned i = 0; i < 64; ++i)
{
    output[output_index + i] = calculation(inputs, i);
}

这里，work-item (0, 0) 会先写入 item 0，然后是 item 1，然后是 item 2，而 work-item (1, 0) 会先写入 item 64，然后是 65，依此类推。

如果工作项 (0, 0) 在工作项 (1, 0) 写入索引 1 的同时写入索引 0，通常会更快，依此类推。因此，如果可以，请尝试对输出数组进行布局，以使值维度具有更高阶的步幅，这样您就可以编写：

uint stride = get_global_size(0) * get_global_size(1);
uint output_index = (get_global_size(0) * get_global_id(1) + get_global_id(0));
for (unsigned i = 0; i < 64; ++i)
{
    output[output_index] = calculation(inputs, i);
    output_index += stride;
}

2。使用本地内存作为中间体

如果更改全局内存的布局不是一个选项，您可以改为将结果写入本地内存，该顺序对全局内存来说效率低下，然后有效地将其从本地内存复制到全局内存。高效，我的意思是工作组中的相邻工作项应该再次写入相邻的全局内存位置。您可以显式执行此操作，也可以在内核中使用 async_work_group_copy 函数。

3。压缩

如果有某种方法可以更节省空间来表示您的 64 个值，那将有很大帮助，尤其是当您随后将结果发送回主机 CPU 时。例如，如果精度不是那么重要并且范围受到限制并且您当前正在使用floats，您可以尝试使用half（16 位）浮点值，或short/ushort 16-位整数值，精度略高但范围更小。或者，如果您的值以某种方式相关，您可以使用其他一些表示形式，例如共享指数。

4。 GPU 上的前向计算

如果您当前在主机 CPU 上使用计算结果，您可能会受到 PCIe 带宽的限制，该带宽远低于 GPU 到 VRAM 的带宽。在这种情况下，请考虑将您正在执行的任何进一步计算转移到 GPU 上，即使 CPU 实现本身目前还不是瓶颈。避免从 VRAM 复制到系统 RAM 可能会给您带来更大的提升。

更好的是，如果您可以完全避免将此结果写入全局内存，例如通过在同一个内核中执行向前计算，可能在将中间结果存储在本地内存中以与工作组共享它之后，那么您可以完全避免内存瓶颈。

5。阅读 OpenCL 优化指南

您可能还可以针对您的工作负载执行其他优化。由于您没有详细说明您正在做什么，我们无法轻易猜测这些优化。 GPU 制造商发布了 OpenCL 优化指南，请确保您阅读并理解它们，看看您是否可以将任何建议应用到您的任务中。

Answer 2

根据我的经验，在写入内核中的缓冲区时，我没有发现任何类型的瓶颈，我相信这就是您所指的。

在每个线程中写入这 64 个值应该不是问题，您的瓶颈可能在其他地方。这可能是在将内核入队之前，同时准备缓冲区参数时所做的事情。