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

Is it more efficient in SYCL to use one buffer or multiple buffers?

本文关键字：缓冲区 SYCL 有效 一个      更新时间：2024-09-23

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

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

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

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

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

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

請參閱這裡對(舊的)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 运行时的角度来看，一般来说，使用单个大缓冲区与多个较小的缓冲区既有优点也有缺点(例如，许多缓冲区的开销，但任务图优化策略的机会更多)。我希望这些影响是次要的。

相关文章：