优化主机与加速器间内存移动：use_host_ptr属性与内存管理
本文介绍了如何通过使用SYCL的`use_host_ptr`属性来优化主机与GPU加速器之间的内存移动，减少复制开销，特别是在共享内存和独立设备环境下。文章通过示例展示了如何控制内存分配以避免不必要的复制，并强调了内存对齐的重要性。
本章节翻译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时避免不必要的复制。即使加速器与主机共享内存，也必须满足一些额外条件才能避免这些额外复制。
扫一扫
24
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
