oneAPI GPU 优化指南 - 应避免在循环内声明buffer

本章节翻译by weavingtime@formail.com 原文 访问buffer的不同模式 — oneAPI GPU Optimization Guide

当在 for 循环内重复启动kernel时，可以通过在循环外部声明buffer来防止重复分配和释放buffer。 在循环内声明buffer会引入重复的主机到设备和设备到主机内存复制。

在以下示例中，kernel在 for 循环内重复启动。 buffer C用作临时数组，用于保存迭代中的值， 并且在一次迭代中分配的值不会在任何其他迭代中使用。 由于buffer C 是在 for 循环内声明的， 因此它在每次循环迭代中都会被分配和释放。 除了buffer的分配和释放之外，与buffer关联的内存在每次迭代中从主机重复 地传输到设备以及从设备传输到主机。

 1#include <CL/sycl.hpp>
 2#include <stdio.h>
 3
 4constexpr int N = 25;
 5constexpr int STEPS = 100000;
 6
 7int main() {
 8
 9  int AData[N];
10  int BData[N];
11  int CData[N];
12
13  sycl::queue Q;
14
15  // Create 2 buffers, each holding N integers
16  sycl::buffer<int> ABuf(&AData[0], N);
17  sycl::buffer<int> BBuf(&BData[0], N);
18
19  Q.submit([&](auto &h) {
20    // Create device accessors.
21    // The property no_init lets the runtime know that the
22    // previous contents of the buffer can be discarded.
23    sycl::accessor aA(ABuf, h, sycl::write_only, sycl::no_init);
24    sycl::accessor aB(BBuf, h, sycl::write_only, sycl::no_init);
25    h.parallel_for(N, [=](auto i) {
26      aA[i] = 10;
27      aB[i] = 20;
28    });
29  });
30
31  for (int j = 0; j < STEPS; j++) {
32    sycl::buffer<int> CBuf(&CData[0], N);
33    Q.submit([&](auto &h) {
34      // Create device accessors.
35      sycl::accessor aA(ABuf, h);
36      sycl::accessor aB(BBuf, h);
37      sycl::accessor aC(CBuf, h);
38      h.parallel_for(N, [=](auto i) {
39        aC[i] = (aA[i] < aB[i]) ? -1 : 1;
40        aA[i] += aC[i];
41        aB[i] -= aC[i];
42      });
43    });
44  } // end for
45
46  // Create host accessors.
47  const sycl::host_accessor haA(ABuf);
48  const sycl::host_accessor haB(BBuf);
49  printf("%d %d\n", haA[N / 2], haB[N / 2]);
50
51  return 0;
52}

更好的方法是在 for 循环之前声明buffer C，以便仅分配和释放一次，从而通过避免主机和设备之间的多余数据传输来提高性能。 以下kernel显示了此更改。

 1#include <CL/sycl.hpp>
 2#include <stdio.h>
 3
 4constexpr int N = 25;
 5constexpr int STEPS = 100000;
 6
 7int main() {
 8
 9  int AData[N];
10  int BData[N];
11  int CData[N];
12
13  sycl::queue Q;
14
15  // Create 3 buffers, each holding N integers
16  sycl::buffer<int> ABuf(&AData[0], N);
17  sycl::buffer<int> BBuf(&BData[0], N);
18  sycl::buffer<int> CBuf(&CData[0], N);
19
20  Q.submit([&](auto &h) {
21    // Create device accessors.
22    // The property no_init lets the runtime know that the
23    // previous contents of the buffer can be discarded.
24    sycl::accessor aA(ABuf, h, sycl::write_only, sycl::no_init);
25    sycl::accessor aB(BBuf, h, sycl::write_only, sycl::no_init);
26    h.parallel_for(N, [=](auto i) {
27      aA[i] = 10;
28      aB[i] = 20;
29    });
30  });
31
32  for (int j = 0; j < STEPS; j++) {
33    Q.submit([&](auto &h) {
34      // Create device accessors.
35      sycl::accessor aA(ABuf, h);
36      sycl::accessor aB(BBuf, h);
37      sycl::accessor aC(CBuf, h);
38      h.parallel_for(N, [=](auto i) {
39        aC[i] = (aA[i] < aB[i]) ? -1 : 1;
40        aA[i] += aC[i];
41        aB[i] -= aC[i];
42      });
43    });
44  } // end for
45
46  // Create host accessors.
47  const sycl::host_accessor haA(ABuf);
48  const sycl::host_accessor haB(BBuf);
49  printf("%d %d\n", haA[N / 2], haB[N / 2]);
50
51  return 0;
52}

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