oneAPI GPU 优化指南 - 在设备端同时执行多个 kernel

本章节翻译by chenshusmail@163.com 原文：Executing Multiple Kernels on the Device at the Same Time (intel.com)

SYCL 有两种 queue，程序员可以创建并使用它们来提交 kernel 进行执行。

• 顺序的 queue

  在这里，kernel 按照它们被提交到 queue 的顺序进行执行。

• 无序的 queue

  在这里，kernel 可以按照任意顺序执行（受到它们之间的依赖性约束）。

创建按顺序或无序的 queue 的选择是在 queue 构建时通过属性 sycl::property::queue::in_order()。默认情况下，当没有指定属性时，queue 是无序的。

在以下示例中，每次迭代提交了三个 kernel。这些 kernel 中的每一个只使用一个 work-group， 包含 256 个 work-item。这些 kernel 特意创建为一个 group，以确保它们不会使用整个机器。 这样做是为了展示并行执行 kernel 的好处。

int multi_queue(sycl::queue &q, const IntArray &a, const IntArray &b) {
  IntArray s1, s2, s3;

  sycl::buffer a_buf(a);
  sycl::buffer b_buf(b);
  sycl::buffer sum_buf1(s1);
  sycl::buffer sum_buf2(s2);
  sycl::buffer sum_buf3(s3);

  size_t num_groups = 1;
  size_t wg_size = 256;
  auto start = std::chrono::steady_clock::now();
  for (int i = 0; i < iter; i++) {
    q.submit([&](sycl::handler &h) {
      sycl::accessor a_acc(a_buf, h, sycl::read_only);
      sycl::accessor b_acc(b_buf, h, sycl::read_only);
      sycl::accessor sum_acc(sum_buf1, h, sycl::write_only, sycl::no_init);

      h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
                     [=](sycl::nd_item<1> index) {
                       size_t loc_id = index.get_local_id();
                       sum_acc[loc_id] = 0;
                       for (int j = 0; j < 1000; j++)
                         for (size_t i = loc_id; i < array_size; i += wg_size) {
                           sum_acc[loc_id] += a_acc[i] + b_acc[i];
                         }
                     });
    });
    q.submit([&](sycl::handler &h) {
      sycl::accessor a_acc(a_buf, h, sycl::read_only);
      sycl::accessor b_acc(b_buf, h, sycl::read_only);
      sycl::accessor sum_acc(sum_buf2, h, sycl::write_only, sycl::no_init);

      h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
                     [=](sycl::nd_item<1> index) {
                       size_t loc_id = index.get_local_id();
                       sum_acc[loc_id] = 0;
                       for (int j = 0; j < 1000; j++)
                         for (size_t i = loc_id; i < array_size; i += wg_size) {
                           sum_acc[loc_id] += a_acc[i] + b_acc[i];
                         }
                     });
    });
    q.submit([&](sycl::handler &h) {
      sycl::accessor a_acc(a_buf, h, sycl::read_only);
      sycl::accessor b_acc(b_buf, h, sycl::read_only);
      sycl::accessor sum_acc(sum_buf3, h, sycl::write_only, sycl::no_init);

      h.parallel_for(sycl::nd_range<1>(num_groups * wg_size, wg_size),
                     [=](sycl::nd_item<1> index) {
                       size_t loc_id = index.get_local_id();
                       sum_acc[loc_id] = 0;
                       for (int j = 0; j < 1000; j++)
                         for (size_t i = loc_id; i < array_size; i += wg_size) {
                           sum_acc[loc_id] += a_acc[i] + b_acc[i];
                         }
                     });
    });
  }
  q.wait();
  auto end = std::chrono::steady_clock::now();
  std::cout << "multi_queue completed on device - took "
            << (end - start).count() << " u-secs\n";
  // check results
  return ((end - start).count());
} // end multi_queue

在底层 queue 是按顺序的情况下，这些 kernel 不能并行执行，即使机器中有足够的资源， 且 kernel 之间没有依赖关系，它们也必须顺序执行。这可以从所有 kernel 的总执行时间较长中看出。 下面显示了 queue 的创建和 kernel 的提交。

  sycl::property_list q_prop{sycl::property::queue::in_order()};
  std::cout << "In order queue: Jitting+Execution time\n";
  sycl::queue q1(sycl::default_selector_v, q_prop);
  multi_queue(q1, a, b);
  usleep(500 * 1000);
  std::cout << "In order queue: Execution time\n";
  multi_queue(q1, a, b);

当 queue 是无序的时候，总的执行时间会大大降低，这表明机器能够同时执行来自 queue 的不同的 kernel。 下面显示了 queue 的创建和 kernel 的调用。

  sycl::queue q2(sycl::default_selector_v);
  std::cout << "Out of order queue: Jitting+Execution time\n";
  multi_queue(q2, a, b);
  usleep(500 * 1000);
  std::cout << "Out of order queue: Execution time\n";
  multi_queue(q2, a, b);

在 kernel 的扩展性不强，因此无法有效利用全机器计算资源的情况下， 最好只通过适当选择 work-group/work-item 值来分配所需的计算单元，并尝试同时执行多个 kernel。

以下的时间线视图显示了按顺序和无序 queue 执行的 kernel（这是使用 onetrace 工具收集的， 该工具可在 https://github.com/intel/pti-gpu/tree/master/tools/onetrace 获取）。 在这里，可以清楚地看到提交给无序 queue 的 kernel 正在并行执行。另一件需要注意的事情是， 并非所有三个 kernel 都一直并行执行。并行执行多少个 kernel 受到多个因素的影响， 例如硬件资源的可用性、kernel 提交之间的时间间隔等。

执行按顺序和无序 queue 的 kernel 的时间线

也可以通过使用 device class 的 create_sub_devices 函数将单个设备静态分区为多个 sub-device。 这为程序员提交 kernel 到适当的 sub-device 提供了更多的控制。然而，设备到 sub-device 的分区是静态的， 所以 runtime 将无法适应应用程序的动态负载，因为它没法灵活地将 kernel 从一个 sub-device 移动到另一个。

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