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

最新推荐文章于 2025-11-24 15:27:16 发布

翻译最新推荐文章于 2025-11-24 15:27:16 发布 · 139 阅读

文章标签：

#oneapi #算法

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

46 篇文章

订阅专栏

本章节翻译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}