在 SYCL 中使用一个缓冲区还是多个缓冲区更有效？

Posted 2023-03-10

【中文标题】在 SYCL 中使用一个缓冲区还是多个缓冲区更有效？

【英文标题】：Is it more efficient in SYCL to use one buffer or multiple buffers?

【发布时间】：2020-11-11 04:57:33

【问题描述】：

假设我有一个数据数组，例如一个大小为 N 的 3D 向量数组。假设我的 SYCL 内核的每次迭代都专门或主要只关注一个向量。作为一般规则，以下哪一种将其​​分解为连续缓冲区的方法更有效——或者这很重要？

我意识到目标设备对此影响很大，所以我们假设它是一个独立的 GPU（即数据确实必须复制到不同的内存芯片，并且该设备没有像 FPGA 这样的疯狂架构——我主要通过 CUDA 瞄准 GTX 1080，但我希望当代码编译到 OpenCL 或我们使用另一个现代 GPU 时，答案可能相似。

为每个坐标创建一个单独的缓冲区，例如sycl::buffer<float> x, y, z;，每个大小为 N。这样访问它们时，我可以使用传递给我的内核 lambda 的 sycl::id<1> 作为索引而无需数学运算。 （我怀疑编译器可能会对此进行优化。） 为所有这些创建一个打包缓冲区，例如sycl::buffer<float> coords; 尺寸为 3N。当使用名为i 的sycl::id<1> 访问它们时，我将x 坐标设为buffer_accessor[3*i]，将y 坐标设为buffer_accessor[3*i+1]，将z 坐标设为buffer_accessor[3*i+2]。 （我不知道编译器是否可以对此进行优化，我不确定对齐问题是否会起作用。） 使用结构体创建一个解压缓冲区，例如struct Coord float x,y,z; ; sycl::buffer<Coord> coords;。由于对齐填充，内存使用量增加了相当惊人的成本，在这个例子中增加了 33%——这也会增加将缓冲区复制到设备所需的时间。但权衡是您可以在不操作 sycl::id<1> 的情况下访问数据，运行时只需处理一个缓冲区，并且设备上不应该出现任何缓存行对齐效率低下的问题。 使用大小为 (N,3) 的二维缓冲区并仅在第一个维度的范围内进行迭代。这是一个不太灵活的解决方案，我不明白为什么在不迭代所有维度时要使用多维缓冲区，除非为此用例内置了大量优化。

我找不到任何关于数据架构的指南来获得对这类事情的直觉。现在（4）看起来很傻，（3）涉及不可接受的内存浪费，我正在使用（2）但想知道我是否不应该使用（1）来避免 id 操作和 3*sizeof(float)对齐的访问块。

【参考方案1】：

对于 GPU 上的内存访问模式，首先要了解合并的概念。基本上这意味着在某些情况下，设备将合并相邻工作项的内存访问，而不是发出一个大内存访问。这对性能非常重要。 发生合并时的详细要求因 GPU 供应商而异（甚至在一个供应商的 GPU 代之间）。但通常，要求往往是这样的

一定数量的相邻工作项访问相邻的数据元素。例如。 SYCL 子组/CUDA warp 中的所有工作项都可以访问后续数据元素。 第一个工作项访问的数据元素可能必须对齐，例如到缓存行。

在此处查看（旧）NVIDIA GPU 的说明：https://developer.nvidia.com/blog/how-access-global-memory-efficiently-cuda-c-kernels/

考虑到这一点，3) 不仅浪费内存容量，还浪费内存带宽，如果你有类似 my_accessor[id].x 的东西，你就有一个跨步的内存访问，可以防止合并。

对于 4)，我不确定我是否理解正确。我假设您的意思是具有 3 个元素的维度控制您是否访问 x/y/z，而具有 N 的维度描述了第 n 个向量。在这种情况下，这取决于您的大小是(N, 3) 还是(3, N)。因为在 SYCL 中，数据布局使得最后一个索引总是最快的，(N, 3) 实际上对应于 3) 没有填充问题。 (3, N) 类似于 2) 但没有跨步内存访问（见下文）

对于 2)，主要的性能问题是，如果 x 位于 [3*i]，y 位于 [3*i+1] 等，则您正在执行跨步内存访问。对于合并，您希望 x 位于 [i]，y 位于[N+i] 和 z 在[2N+i]。 如果你有类似的东西

float my_x = data[i]; // all N work items perform coalesced access for x
float my_y = data[i+N];
float my_z = data[i+2N];

你有一个很好的内存访问模式。根据您对 N 的选择以及设备合并内存访问的对齐要求，您可能会因为对齐而遇​​到 y 和 z 的性能问题。

我认为您需要向索引添加偏移量这一事实不会显着影响性能。

对于 1)，您主要可以保证所有数据都很好地对齐并且访问将合并。正因为如此，我希望它能够在所提出的方法中表现最好。

从 SYCL 运行时的角度来看，一般来说，使用单个大缓冲区与使用多个较小缓冲区既有优点也有缺点（例如，许多缓冲区的开销，但任务图优化策略的机会更多）。我希望这些影响是次要的。