oneAPI GPU 优化指南 - 避免在主机和设备之间来回移动数据

TechVideoGPU

已于 2024-02-19 23:27:47 修改

阅读量32

点赞数

分类专栏： oneAPI GPU 优化指南文章标签： oneapi

于 2024-02-19 23:12:12 首次发布

oneAPI GPU 优化指南专栏收录该内容

46 篇文章 1 订阅

订阅专栏

本章节翻译by weavingtime@formail.com 原文 Avoiding Moving Data Back and Forth between Host and Device (intel.com)

在主机和设备之间移动数据的成本相当高，特别是在独立加速器（比如独显）的情况下。因此，尽可能避免主机和设备之间的数据传输非常重要。在某些情况下，可能需要将加速器上的kernel计算的数据带到主机并对其进行一些操作，然后将其发送回设备以进行进一步处理。在这种情况下，我们最终将产生设备到主机传输的开销，然后再次产生主机到设备传输的开销。

考虑以下示例，其中一个kernel通过某种操作（在本例中为vector相加）将数据生成到新vector中。然后，通过对每个值应用函数，将这个新vector转换为第三个vector，最后将这个第三个vector作为输入传到另一个kernel中以进行一些额外的计算。在许多领域中这种形式的计算非常常见，在这些领域中，算法是迭代的，并且一个计算的输出需要作为另一个计算的输入。例如，在机器学习中，模型被构造为计算层，一层的输出被输入到下一层。

 1double myFunc1(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
 2               AlignedVector<int> &c, AlignedVector<int> &d,
 3               AlignedVector<int> &res, int iter) {
 4  sycl::range num_items{a.size()};
 5  VectorAllocator<int> alloc;
 6  AlignedVector<int> sum(a.size(), alloc);
 7
 8  const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
 9  sycl::buffer a_buf(a, props);
10  sycl::buffer b_buf(b, props);
11  sycl::buffer c_buf(b, props);
12  sycl::buffer d_buf(b, props);
13  sycl::buffer res_buf(res, props);
14  sycl::buffer sum_buf(sum.data(), num_items, props);
15
16  Timer timer;
17  for (int i = 0; i < iter; i++) {
18    // kernel1
19    q.submit([&](auto &h) {
20      // Input accessors
21      sycl::accessor a_acc(a_buf, h, sycl::read_only);
22      sycl::accessor b_acc(b_buf, h, sycl::read_only);
23      // Output accessor
24      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
25
26      h.parallel_for(num_items,
27                     [=](auto id) { sum_acc[id] = a_acc[id] + b_acc[id]; });
28    });
29
30    {
31      sycl::host_accessor h_acc(sum_buf);
32      for (int j = 0; j < a.size(); j++)
33        if (h_acc[j] > 10)
34          h_acc[j] = 1;
35        else
36          h_acc[j] = 0;
37    }
38
39    // kernel2
40    q.submit([&](auto &h) {
41      // Input accessors
42      sycl::accessor sum_acc(sum_buf, h, sycl::read_only);
43      sycl::accessor c_acc(c_buf, h, sycl::read_only);
44      sycl::accessor d_acc(d_buf, h, sycl::read_only);
45      // Output accessor
46      sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init);
47
48      h.parallel_for(num_items, [=](auto id) {
49        res_acc[id] = sum_acc[id] * c_acc[id] + d_acc[id];
50      });
51    });
52    q.wait();
53  }
54  double elapsed = timer.Elapsed() / iter;
55  return (elapsed);
56} // end myFunc1

你可以创建一个 kernel3 来在设备上执行此函数，而不是将数据传送到主机并将函数应用于数据然后将其发送回设备，如以下示例所示。内核 kernel3 对 kernel1 和 kernel2 之间的 accum_buf 中的中间数据进行操作，避免了设备和主机之间数据传输的往返。

 1double myFunc2(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
 2               AlignedVector<int> &c, AlignedVector<int> &d,
 3               AlignedVector<int> &res, int iter) {
 4  sycl::range num_items{a.size()};
 5  VectorAllocator<int> alloc;
 6  AlignedVector<int> sum(a.size(), alloc);
 7
 8  const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
 9  sycl::buffer a_buf(a, props);
10  sycl::buffer b_buf(b, props);
11  sycl::buffer c_buf(b, props);
12  sycl::buffer d_buf(b, props);
13  sycl::buffer res_buf(res, props);
14  sycl::buffer sum_buf(sum.data(), num_items, props);
15
16  Timer timer;
17  for (int i = 0; i < iter; i++) {
18    // kernel1
19    q.submit([&](auto &h) {
20      // Input accessors
21      sycl::accessor a_acc(a_buf, h, sycl::read_only);
22      sycl::accessor b_acc(b_buf, h, sycl::read_only);
23      // Output accessor
24      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
25
26      h.parallel_for(num_items,
27                     [=](auto i) { sum_acc[i] = a_acc[i] + b_acc[i]; });
28    });
29
30    // kernel3
31    q.submit([&](auto &h) {
32      sycl::accessor sum_acc(sum_buf, h, sycl::read_write);
33      h.parallel_for(num_items, [=](auto id) {
34        if (sum_acc[id] > 10)
35          sum_acc[id] = 1;
36        else
37          sum_acc[id] = 0;
38      });
39    });
40
41    // kernel2
42    q.submit([&](auto &h) {
43      // Input accessors
44      sycl::accessor sum_acc(sum_buf, h, sycl::read_only);
45      sycl::accessor c_acc(c_buf, h, sycl::read_only);
46      sycl::accessor d_acc(d_buf, h, sycl::read_only);
47      // Output accessor
48      sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init);
49
50      h.parallel_for(num_items, [=](auto i) {
51        res_acc[i] = sum_acc[i] * c_acc[i] + d_acc[i];
52      });
53    });
54    q.wait();
55  }
56  double elapsed = timer.Elapsed() / iter;
57  return (elapsed);
58} // end myFunc2

还有其他方法可以优化此示例。例如， kernel3 中的裁剪操作可以合并到 kernel1 的计算中，如下所示。这就是kernel融合，并且还有一个好处就是不用启动第三个内核。 SYCL编译器无法进行这种优化。在机器学习等某些特定领域，图编译器可以对机器学习模型进行操作并融合操作，这和kernel融合有一样的作用。

 1double myFunc3(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
 2               AlignedVector<int> &c, AlignedVector<int> &d,
 3               AlignedVector<int> &res, int iter) {
 4  sycl::range num_items{a.size()};
 5  VectorAllocator<int> alloc;
 6  AlignedVector<int> sum(a.size(), alloc);
 7
 8  const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
 9  sycl::buffer a_buf(a, props);
10  sycl::buffer b_buf(b, props);
11  sycl::buffer c_buf(b, props);
12  sycl::buffer d_buf(b, props);
13  sycl::buffer res_buf(res, props);
14  sycl::buffer sum_buf(sum.data(), num_items, props);
15
16  Timer timer;
17  for (int i = 0; i < iter; i++) {
18    // kernel1
19    q.submit([&](auto &h) {
20      // Input accessors
21      sycl::accessor a_acc(a_buf, h, sycl::read_only);
22      sycl::accessor b_acc(b_buf, h, sycl::read_only);
23      // Output accessor
24      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
25
26      h.parallel_for(num_items, [=](auto i) {
27        int t = a_acc[i] + b_acc[i];
28        if (t > 10)
29          sum_acc[i] = 1;
30        else
31          sum_acc[i] = 0;
32      });
33    });
34
35    // kernel2
36    q.submit([&](auto &h) {
37      // Input accessors
38      sycl::accessor sum_acc(sum_buf, h, sycl::read_only);
39      sycl::accessor c_acc(c_buf, h, sycl::read_only);
40      sycl::accessor d_acc(d_buf, h, sycl::read_only);
41      // Output accessor
42      sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init);
43
44      h.parallel_for(num_items, [=](auto i) {
45        res_acc[i] = sum_acc[i] * c_acc[i] + d_acc[i];
46      });
47    });
48    q.wait();
49  }
50  double elapsed = timer.Elapsed() / iter;
51  return (elapsed);
52} // end myFunc3

We can take this kernel fusion one level further and fuse both kernel1 and kernel2 as shown in the code below. This gives very good performance since it avoids the intermediate accum_buf completely, saving memory in addition to launching an additional kernel. Most of the performance benefit in this case is due to improvement in locality of memory references.

我们可以将此kernel融合进一步提升一层，融合 kernel1 和 kernel2 ，如下面的代码所示。这个改动提供了非常好的性能，因为它完全避免了中间的 accum_buf，除了启动额外的kernel之外还节省了内存。在这种情况下，大部分性能优势是来自内存引用局部性的改进。

 1double myFunc4(sycl::queue &q, AlignedVector<int> &a, AlignedVector<int> &b,
 2               AlignedVector<int> &c, AlignedVector<int> &d,
 3               AlignedVector<int> &res, int iter) {
 4  sycl::range num_items{a.size()};
 5  VectorAllocator<int> alloc;
 6
 7  const sycl::property_list props = {sycl::property::buffer::use_host_ptr()};
 8  sycl::buffer a_buf(a, props);
 9  sycl::buffer b_buf(b, props);
10  sycl::buffer c_buf(b, props);
11  sycl::buffer d_buf(b, props);
12  sycl::buffer res_buf(res, props);
13
14  Timer timer;
15  for (int i = 0; i < iter; i++) {
16    // kernel1
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      sycl::accessor c_acc(c_buf, h, sycl::read_only);
22      sycl::accessor d_acc(d_buf, h, sycl::read_only);
23      // Output accessor
24      sycl::accessor res_acc(res_buf, h, sycl::write_only, sycl::no_init);
25
26      h.parallel_for(num_items, [=](auto i) {
27        int t = a_acc[i] + b_acc[i];
28        if (t > 10)
29          res_acc[i] = c_acc[i] + d_acc[i];
30        else
31          res_acc[i] = d_acc[i];
32      });
33    });
34    q.wait();
35  }
36  double elapsed = timer.Elapsed() / iter;
37  return (elapsed);
38} // end myFunc4

TechVideoGPU

关注

0
点赞
踩
0

收藏

觉得还不错? 一键收藏
0
评论
oneAPI GPU 优化指南 - 避免在主机和设备之间来回移动数据

然后，通过对每个值应用函数，将这个新vector转换为第三个vector，最后将这个第三个vector作为输入传到另一个kernel中以进行一些额外的计算。在许多领域中这种形式的计算非常常见，在这些领域中，算法是迭代的，并且一个计算的输出需要作为另一个计算的输入。在某些情况下，可能需要将加速器上的kernel计算的数据带到主机并对其进行一些操作，然后将其发送回设备以进行进一步处理。在这种情况下，我们最终将产生设备到主机传输的开销，然后再次产生主机到设备传输的开销。的计算中，如下所示。
复制链接

扫一扫