在SYCL中使用一个或多个缓冲区更有效吗？

如何解决在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。当使用名为sycl::id<1>的{{1}}访问它们时，我然后将x坐标捕获为i，将y坐标捕获为buffer_accessor[3*i]，将z坐标捕获为buffer_accessor[3*i+1] 。（我不知道编译器是否可以优化它，也不确定对齐问题是否会起作用。）
使用结构创建一个未打包的缓冲区，例如buffer_accessor[3*i+2]。由于对齐填充，这会增加内存使用率，这在本例中增加了33％，这是相当惊人的成本，这也会增加将缓冲区复制到设备所需的时间。但是要权衡的是，您无需操纵struct Coord { float x,z; }; sycl::buffer<Coord> coords;就可以访问数据，运行时只需要处理一个缓冲区，并且设备上不应存在任何缓存行对齐效率低下的情况。
使用大小为（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在{{ 1}}，y在[i]，z在[N+i]。如果你有类似的东西

[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];的选择以及设备对合并的内存访问的对齐要求，由于对齐，您可能会遇到y和z的性能问题。

我不希望您需要在索引中添加偏移量这一事实会严重影响性能。

对于1），您将主要获得所有数据完全对齐且访问将合并的保证。因此，我希望它能在所介绍的方法中表现最好。

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