在 SYCL 中使用一个缓冲区还是多个缓冲区更有效？
Is it more efficient in SYCL to use one buffer or multiple buffers?
假设我有一个数据数组，例如一个大小为 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)对齐的访问块。
对于 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 运行时的角度来看，通常使用单个大缓冲区与多个较小缓冲区既有优点也有缺点（例如，许多缓冲区的开销，但任务图优化策略的机会更多）。我希望这些影响是次要的。
假设我有一个数据数组，例如一个大小为 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)对齐的访问块。
对于 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 运行时的角度来看，通常使用单个大缓冲区与多个较小缓冲区既有优点也有缺点（例如，许多缓冲区的开销，但任务图优化策略的机会更多）。我希望这些影响是次要的。