oneAPI GPU 优化指南 - 优化主机和加速器之间的内存移动-CSDN博客

本文介绍了如何通过使用SYCL的`use_host_ptr`属性来优化主机与GPU加速器之间的内存移动，减少复制开销，特别是在共享内存和独立设备环境下。文章通过示例展示了如何控制内存分配以避免不必要的复制，并强调了内存对齐的重要性。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

本章节翻译by weavingtime@formail.com 原文 Optimizing Memory Movement Between Host and Accelerator (intel.com)

可以使用属性创建buffer来控制它们的分配方式。其中一个属性是use_host_ptr，它告诉runtime，如果可能，应该直接使用主机内存而不是复制。这避免了在主机内存和buffer内存之间来回复制buffer内容，从而在buffer创建和销毁期间节省一些时间。

再举一个例子，当GPU和CPU具有共享内存时，可以通过共享内存page来避免复制内存。但是为了实现内存page共享，分配的内存需要具有一些属性，例如在page边界上对齐。对于独立设备，可能无法实现这种好处，因为加速器的任何访问主机内存操作都必须通过PCIe或其他比加速器的内存更慢的接口进行。

下面的代码展示了如何在主机上、buffer内以及kernel中加速器设备上打印内存地址。

 1int VectorAdd0(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
 2               AlignedVector<int> &sum, int iter) {
 3  sycl::range num_items{a.size()};
 4
 5  const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
 6
 7  for (int i = 0; i < iter; i++) {
 8    sycl::buffer a_buf(a, props);
 9    sycl::buffer b_buf(b, props);
10    sycl::buffer sum_buf(sum.data(), num_items, props);
11    {
12      sycl::host_accessor a_host_acc(a_buf);
13      std::cout << "add0: buff memory address =" << a_host_acc.get_pointer()
14                << "\n";
15      std::cout << "add0: address of vector a = " << a.data() << "\n";
16    }
17    q.submit([&](auto &h) {
18      // Input accessors
19      sycl::accessor a_acc(a_buf, h, sycl::read_only);
20      sycl::accessor b_acc(b_buf, h, sycl::read_only);
21      // Output accessor
22      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
23      sycl::stream out(1024 * 1024, 1 * 128, h);
24
25      h.parallel_for(num_items, [=](auto i) {
26        if (i[0] == 0)
27          out << "add0:  dev addr = " << a_acc.get_pointer() << "\n";
28        sum_acc[i] = a_acc[i] + b_acc[i];
29      });
30    });
31  }
32  q.wait();
33  return (0);
34}

当运行此程序时，可以看到当为集成GPU设备设置 use_host_ptr 属性时，所有三个地址（主机上、buffer中和加速器上）都相同。但对于独立的GPU设备（独显）， buffer和设备地址将不同。另请注意，在第1行中，所有参数都没有被声明为 const 。如果这些参数被声明为 const ，则在创建buffer时它们将被复制并分配新内存，而不是重用主机vector中的内存。下面的代码片段演示了这一点。当执行此代码时，我们看到与传入vector相关联的地址与buffer中存在的内存以及加速器设备中存在的内存不同。

 1int VectorAdd1(sycl::queue &q, const AlignedVector<int> &a,
 2               const AlignedVector<int> &b, AlignedVector<int> &sum, int iter) {
 3  sycl::range num_items{a.size()};
 4
 5  const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
 6
 7  for (int i = 0; i < iter; i++) {
 8    sycl::buffer a_buf(a, props);
 9    sycl::buffer b_buf(b, props);
10    sycl::buffer sum_buf(sum.data(), num_items, props);
11    {
12      sycl::host_accessor a_host_acc(a_buf);
13      std::cout << "add1: buff memory address =" << a_host_acc.get_pointer()
14                << "\n";
15      std::cout << "add1: address of vector aa = " << a.data() << "\n";
16    }
17    q.submit([&](auto &h) {
18      // Input accessors
19      sycl::accessor a_acc(a_buf, h, sycl::read_only);
20      sycl::accessor b_acc(b_buf, h, sycl::read_only);
21      // Output accessor
22      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
23      sycl::stream out(16 * 1024, 16 * 1024, h);
24
25      h.parallel_for(num_items, [=](auto i) {
26        if (i[0] == 0)
27          out << "add1: dev addr = " << a_acc.get_pointer() << "\n";
28        sum_acc[i] = a_acc[i] + b_acc[i];
29      });
30    });
31  }
32  q.wait();
33  return (0);
34}

由于在创建buffer时设置了 use_host_ptr 属性，并且buffer在page边界上对齐，因此kernel vectorAdd3 不会产生将buffer内容复制到加速器设备的开销（针对集成GPU设备）。如果buffer指向的内存未在page边界上对齐，则将分配新内存以便在page边界上对齐，并将buffer内容复制到该内存中。然后通过将buffer中的新内存与加速器共享（对于不共享任何内存的加速器，从主机上的buffer复制内容到设备上）或使用页表避免来复制可在设备上使用的内存（对于共享内存的加速器）。

 1int VectorAdd2(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
 2               AlignedVector<int> &sum, int iter) {
 3  sycl::range num_items{a.size()};
 4
 5  const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
 6
 7  auto start = std::chrono::steady_clock::now();
 8  for (int i = 0; i < iter; i++) {
 9    sycl::buffer a_buf(a, props);
10    sycl::buffer b_buf(b, props);
11    sycl::buffer sum_buf(sum.data(), num_items, props);
12    q.submit([&](auto &h) {
13      // Input accessors
14      sycl::accessor a_acc(a_buf, h, sycl::read_only);
15      sycl::accessor b_acc(b_buf, h, sycl::read_only);
16      // Output accessor
17      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
18
19      h.parallel_for(num_items,
20                     [=](auto i) { sum_acc[i] = a_acc[i] + b_acc[i]; });
21    });
22  }
23  q.wait();
24  auto end = std::chrono::steady_clock::now();
25  std::cout << "Vector add2 completed on device - took "
26            << (end - start).count() << " u-secs\n";
27  return ((end - start).count());
28}

下面的kernel将产生在主机和buffer之间以及从buffer到加速器之间复制内存内容的开销。

 1int VectorAdd3(sycl::queue &q, const AlignedVector<int> &a,
 2               const AlignedVector<int> &b, AlignedVector<int> &sum, int iter) {
 3  sycl::range num_items{a.size()};
 4
 5  auto start = std::chrono::steady_clock::now();
 6  for (int i = 0; i < iter; i++) {
 7    sycl::buffer a_buf(a);
 8    sycl::buffer b_buf(b);
 9    sycl::buffer sum_buf(sum.data(), num_items);
10    auto e = q.submit([&](auto &h) {
11      // Input accessors
12      sycl::accessor a_acc(a_buf, h, sycl::read_only);
13      sycl::accessor b_acc(b_buf, h, sycl::read_only);
14      // Output accessor
15      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
16
17      h.parallel_for(num_items,
18                     [=](auto i) { sum_acc[i] = a_acc[i] + b_acc[i]; });
19    });
20  }
21  q.wait();
22  auto end = std::chrono::steady_clock::now();
23  std::cout << "Vector add3 completed on device - took "
24            << (end - start).count() << " u-secs\n";
25  return ((end - start).count());
26}

必须注意在创建buffer和将buffer中的内存传递给kernel时避免不必要的复制。即使加速器与主机共享内存，也必须满足一些额外条件才能避免这些额外复制。