小编sap*_*ous的帖子

假设我有一组数据，例如大小为 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)对齐的访问块。