2

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

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

  1. 为每个坐标创建一个单独的缓冲区，例如sycl::buffer<float> x, y, z;每个大小为 N 的缓冲区。这样访问它们时，我可以使用sycl::id<1>传递给我的内核 lambda 作为索引而无需数学。（我怀疑编译器可能会对此进行优化。）
  2. 为所有这些创建一个打包缓冲区，例如sycl::buffer<float> coords;大小为 3N。当使用sycl::id<1>被调用访问它们时i，我将 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)对齐的访问块。

4

1 回答 1

4

对于 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 个向量。在这种情况下，这将取决于您是否有 size(N, 3)或(3, N). 因为在 SYCL 中，数据布局使得最后一个索引总是最快的，(N, 3)实际上对应于 3) 没有填充问题。(3, N)将类似于 2）但没有跨步内存访问（见下文）

对于 2)，主要的性能问题是，如果 x 位于[3*i]y at[3*i+1]等处，则您正在执行跨步内存访问。对于合并，您希望 x 位于[i]y at[N+i]和 z at 处[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 运行时的角度来看，通常使用单个大缓冲区与使用多个较小缓冲区既有优点也有缺点（例如，许多缓冲区的开销，但任务图优化策略的机会更多）。我希望这些影响是次要的。

于 2020-11-12T01:11:09.980 回答