oneAPI GPU 优化指南 - USM和buffer对性能影响

翻译by weavingtime@formail.com 原文：Performance Impact of USM and Buffers (intel.com)

SYCL为设备上的内存管理提供了几种选择。本节将简要介绍相关概念以及性能折衷。有关深入解释，请参见 Data Parallel C++。

与其他语言特性一样，规范定义了行为但不定义实现，因此性能特征可能在软件版本和设备之间发生变化。本指南提供最佳的实践指导。

Buffers。buffer是一个可以从设备和主机访问的数据容器。SYCL runtime通过提供分配、读取和写入内存的API来管理内存。runtime负责在主机和设备之间移动数据，并同步对数据的访问。

Unified Shared Memory (USM)。USM与buffer不同， USM允许使用常规指针读写数据，buffer中对数据的访问仅能通过API进行。USM有两种常用变体。设备内存 只能从设备访问，因此需要在主机和设备之间显式移动数据。“共享内存” 可以从设备或主机引用，runtime自动移动内存。

我们通过展示使用三种模型编写的相同示例程序来说明选择之间的折衷。为了突出问题，我们使用一个GPU和主机协同计算的程序示例，这个示例中，程序需要在GPU和主机之间来回传输数据。

我们首先展示下面的串行计算。假设我们希望在GPU上执行第9行的循环，在CPU上执行第14行的循环。两个循环都读写数据数组，因此数据必须在第8行循环的每次迭代中在主机和GPU之间移动。

 1void serial(int stride) {
 2  // Allocate and initialize data
 3  float *data = new float[data_size];
 4  init(data);
 5
 6  timer it;
 7
 8  for (int i = 0; i < time_steps; i++) {
 9    for (int j = 0; j < data_size; j++) {
10      for (int k = 0; k < device_steps; k++)
11        data[j] += 1.0;
12    }
13
14    for (int j = 0; j < data_size; j += stride)
15      data[j] += 1.0;
16  }
17  put_elapsed_time(it);
18
19  check(data);
20
21  delete[] data;
22} // serial

Buffers

下面，我们展示使用buffer管理数据的相同计算。在第3行创建一个buffer并由 init 函数初始化。 init 函数未显示。它接受一个 accessor 或指针。 parallel_for 执行第13行定义的kernel。kernel使用 device_data accessor 读写 buffer_data 中的数据。

注意，代码中并未指定数据的位置。 accessor 表示何时何地需要数据，SYCL runtime将数据移动到 设备（如果需要），然后启动kernel。第21行的 host_accessor 表示将在主机上读/写数据。由于kernel也读/写 buffer_data ，所以 host_accessor 构造函数等待kernel执行完成，并将数据移动到主机后以执行第23行的读/写操作。在下一次循环迭代中， 第11行的 accessor 构造函数等待直到数据被移回设备，这实际上延迟了kernel的启动。

 1void buffer_data(int stride) {
 2  // Allocate buffer, initialize on host
 3  sycl::buffer<float> buffer_data{data_size};
 4  init(sycl::host_accessor(buffer_data, sycl::write_only, sycl::no_init));
 5
 6  timer it;
 7  for (int i = 0; i < time_steps; i++) {
 8
 9    // Compute on device
10    q.submit([&](auto &h) {
11      sycl::accessor device_data(buffer_data, h);
12
13      auto compute = [=](auto id) {
14        for (int k = 0; k < device_steps; k++)
15          device_data[id] += 1.0;
16      };
17      h.parallel_for(data_size, compute);
18    });
19
20    // Compute on host
21    sycl::host_accessor host_data(buffer_data);
22    for (int i = 0; i < data_size; i += stride)
23      host_data[i] += 1.0;
24  }
25  put_elapsed_time(it);
26
27  const sycl::host_accessor h(buffer_data);
28  check(h);
29} // buffer_data

性能注意事项

第15行和第23行上的数据访问看起来像简单的数组引用，但它们是由SYCLruntime使用C++运算符重载实现的。accessor数组引用的效率取决于实现。实际上，设备代码对于重载与直接内存引用相比没有额外开销。 runtime无法预先知道访问buffer的哪一部分，因此必须在kernel开始之前确保所有数据都在设备上。目前是这样，但随着时间的推移可能会改变。

然而目前对于 host_accessor 情况并非如此。runtime不会将所有数据移动到主机。数组引用是由更复杂的代码来实现，并且比原生C++数组引用慢得多。虽然引用少量数据是可以接受的，但应该在计算密集型算法中避免使用 host_accessor ， 否则会付出巨大的性能代价。

另一个问题是并发性。 host_accessor 可能会阻止引用同一buffer的kernel启动，即使 accessor未被用于读/写数据。 请将包含 host_accessor 的范围限制在最小可能范围内。在本例中，第4行上的主机accessor在init函数返回后销毁，第21行上的主机accessor在每次循环迭代结束时销毁。

共享内存

接下来我们展示使用共享内存实现相同算法。数据在程序示例的第2行分配。这里不需要accessor，因为USM分配的数据可以使用常规指针进行引用。 因此，第10行和第15行上的数组引用可以使用简单索引实现。第12行上的 parallel_for 以 parallel_for 结束，以确保内核在主机访问第15行上的 data 之前完成。 与buffer类似，在启动内核之前，SYCL runtime确保所有数据都驻留在设备上。共享内存与buffer一样，除非被引用，否则不会被复制到主机。主机首次引用数据时，会发生操作系统page fault，然后从设备复制一页数据到主机，并继续执行。对同一页面上的数据进行的后续引用得以全速执行。当启动kernel时，所有驻留在主机上的页面都会刷新回设备。

 1void shared_usm_data(int stride) {
 2  float *data = sycl::malloc_shared<float>(data_size, q);
 3  init(data);
 4
 5  timer it;
 6
 7  for (int i = 0; i < time_steps; i++) {
 8    auto compute = [=](auto id) {
 9      for (int k = 0; k < device_steps; k++)
10        data[id] += 1.0;
11    };
12    q.parallel_for(data_size, compute).wait();
13
14    for (int k = 0; k < data_size; k += stride)
15      data[k] += 1.0;
16  }
17  q.wait();
18  put_elapsed_time(it);
19
20  check(data);
21
22  sycl::free(data, q);
23} // shared_usm_data

性能注意事项

与buffer相比，数据引用是简单的指针并且性能良好。然而，在将数据带到主机时处理page fault除了传输数据的成本外还会产生额外开销。对应用程序的影响取决于引用模式。稀疏随机访问具有最高开销，而通过数据进行线性扫描则受到page fault的影响较小。

由于所有同步都是显式且由开发者控制，因此并发性对于设计良好的程序不是问题。

设备内存

下面可以找到具有设备分配的相同程序。使用设备分配，数据只能在设备上直接访问，并且必须显式地复制到主机，如第21行所示。设备和主机之间所有同步都是显式的。 第21行以 wait 结束，因此主机代码直到异步复制完成才会执行。队列定义在这里未显示，但实际使用了顺序队列，因此第21行上的 memcpy 等待第18行上的 parallel_for 完成。

 1void device_usm_data(int stride) {
 2  // Allocate and initialize host data
 3  float *host_data = new float[data_size];
 4  init(host_data);
 5
 6  // Allocate device data
 7  float *device_data = sycl::malloc_device<float>(data_size, q);
 8
 9  timer it;
10
11  for (int i = 0; i < time_steps; i++) {
12    // Copy data to device and compute
13    q.memcpy(device_data, host_data, sizeof(float) * data_size);
14    auto compute = [=](auto id) {
15      for (int k = 0; k < device_steps; k++)
16        device_data[id] += 1.0;
17    };
18    q.parallel_for(data_size, compute);
19
20    // Copy data to host and compute
21    q.memcpy(host_data, device_data, sizeof(float) * data_size).wait();
22    for (int k = 0; k < data_size; k += stride)
23      host_data[k] += 1.0;
24  }
25  q.wait();
26  put_elapsed_time(it);
27
28  check(host_data);
29
30  sycl::free(device_data, q);
31  delete[] host_data;
32} // device_usm_data

性能注意事项

数据移动和同步都是显示的并且在开发者的完全控制之下。 数组引用是主机上的数组引用，因此它既没有共享内存的page fault开销， 也没有与buffer相关的重载开销。 共享内存仅以内存页为单位传输主机实际引用的数据。 理论上，设备内存允许任意颗粒度的按需移动。 实际上，细粒度、异步的数据移动可能很复杂， 大多数开发者选择移动整个数据结构一次。 显式数据移动和同步的要求使代码更加复杂，但设备内存可以提供最佳的性能。

上一章                                   上级目录 主目录                                                                  下一章