sap*_*ous 3 c++ memory-management sycl dpc++
假设我有一组数据，例如大小为 N 的 3D 向量数组。假设我的 SYCL 内核的每次迭代专门或主要只涉及一个向量。一般来说，以下哪种将其分解为连续缓冲区的方法更有效——或者重要吗？
我意识到目标设备对此影响很大，所以我们假设它是一个独立的 GPU（即数据确实必须复制到不同的内存芯片，并且该设备没有像 FPGA 这样的疯狂架构 - 我是主要针对通过 CUDA 的 GTX 1080，但我预计当代码编译为 OpenCL 或我们使用另一个现代 GPU 时答案可能类似。
sycl::buffer<float> x, y, z;，每个大小为 N。这样，在访问它们时，我可以使用sycl::id<1>传递给内核的 lambda 作为索引，无需进行数学运算。（我怀疑编译器可能能够对此进行优化。）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]。（我不知道编译器是否可以对此进行优化，并且我不确定对齐问题是否会发挥作用。）struct Coord { float x,y,z; }; sycl::buffer<Coord> coords;. 由于对齐填充，这会增加内存使用量，在本例中增加 33%，这会带来相当惊人的成本，这也会增加将缓冲区复制到设备所需的时间。但代价是您可以在不操作 的情况下访问数据sycl::id<1>，运行时只需处理一个缓冲区，并且设备上不应该出现任何缓存行对齐效率低下的情况。我找不到任何关于数据架构的指南来获得对此类事情的直觉。现在（4）看起来很愚蠢，（3）涉及不可接受的内存浪费，我正在使用（2），但想知道我是否不应该使用（1）来避免 id 操作和 3*sizeof(float)对齐的访问块。
对于 GPU 上的内存访问模式，首先要了解合并的概念。基本上，这意味着在某些情况下，设备将合并相邻工作项的内存访问，而不是发出一个大的内存访问。这对于性能来说非常重要。发生合并时的详细要求因 GPU 供应商（甚至同一供应商的 GPU 代之间）而异。但通常情况下，要求往往是这样的
请参阅此处（旧版）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 为 at [3*i]、 y at[3*i+1]等，则您正在执行跨步内存访问。对于合并，您希望 x 为 at [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 运行时的角度来看，一般来说，使用单个大缓冲区与使用多个较小缓冲区相比，既有优点也有缺点（例如，许多缓冲区的开销，但任务图优化策略的机会更多）。我预计这些影响是次要的。