oneAPI GPU 优化指南 - 指针别名和 Restrict 指令

TechVideoGPU

已于 2023-12-08 13:16:28 修改

阅读量34

点赞数

分类专栏： oneAPI GPU 优化指南文章标签： oneapi c++ 1024程序员节

于 2023-12-08 12:51:25 首次发布

原文链接：https://www.intel.com/content/www/us/en/docs/oneapi/optimization-guide-gpu/2023-1/pointer-aliasing-and-the-restrict-directive.html

版权

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

46 篇文章 1 订阅

订阅专栏

本章节翻译by chenshusmail@163.com 原文：Pointer Aliasing and the Restrict Directive (intel.com)

Kernel 通常对使用指针作为入参的元素数组进行操作。当编译器无法确定这些指针是否存在别名时，它会保守地认为它们存在别名，在这种情况下，它不会对这些指针进行重排操作。请看以下向量加法示例，其中每次循环迭代都有两次 load 和一次 store。

size_t VectorAdd(sycl::queue &q, const IntArray &a, const IntArray &b,
                 IntArray &sum, int iter) {
  sycl::range num_items{a.size()};

  sycl::buffer a_buf(a);
  sycl::buffer b_buf(b);
  sycl::buffer sum_buf(sum.data(), num_items);

  auto start = std::chrono::steady_clock::now();
  for (int i = 0; i < iter; i++) {
    auto e = q.submit([&](auto &h) {
      // Input accessors
      sycl::accessor a_acc(a_buf, h, sycl::read_only);
      sycl::accessor b_acc(b_buf, h, sycl::read_only);
      // Output accessor
      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);

      h.parallel_for(num_items,
                     [=](auto i) { sum_acc[i] = a_acc[i] + b_acc[i]; });
    });
  }
  q.wait();
  auto end = std::chrono::steady_clock::now();
  std::cout << "Vector add completed on device - took " << (end - start).count()
            << " u-secs\n";
  return ((end - start).count());
} // end VectorAdd

在这个示例中，程序员将有关向量长度和 work-group 数量的所有选择权留给编译器。在大多数情况下，编译器能很好地选择这些参数以获得良好性能。在某些情况下，获得更高性能的代码的更好方法是显式地指定 work-group 的数量和大小，向编译器提供提示以获得良好性能。

下面的 kernel 被编写为每个 work-item 处理数组的多个元素，并显式选择 work-group 数量和大小。第25行上的 intel::kernel_args_restrict 告诉编译器该 kernel 中的缓冲区访问器彼此不别名。这将允许编译器提升 load 和 store 次序，从而为指令完成提供更多时间并获得更好的指令调度。第27行上的 pragma 指示编译器循环展开系数为2。

size_t VectorAdd2(sycl::queue &q, const IntArray &a, const IntArray &b,
                  IntArray &sum, int iter) {
  sycl::range num_items{a.size()};

  sycl::buffer a_buf(a);
  sycl::buffer b_buf(b);
  sycl::buffer sum_buf(sum.data(), num_items);
  // size_t num_groups =
  // q.get_device().get_info<sycl::info::device::max_compute_units>(); size_t
  // wg_size =
  // q.get_device().get_info<sycl::info::device::max_work_group_size>();
  size_t num_groups = 1;
  size_t wg_size = 16;
  auto start = std::chrono::steady_clock::now();
  for (int i = 0; i < iter; i++) {
    q.submit([&](auto &h) {
      // Input accessors
      sycl::accessor a_acc(a_buf, h, sycl::read_only);
      sycl::accessor b_acc(b_buf, h, sycl::read_only);
      // Output accessor
      sycl::accessor sum_acc(sum_buf, 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) [[intel::reqd_sub_group_size(
                         16)]] [[intel::kernel_args_restrict]] {
                       size_t loc_id = index.get_local_id();
        // unroll with a directive
#pragma unroll(2)
                       for (size_t i = loc_id; i < mysize; i += wg_size) {
                         sum_acc[i] = a_acc[i] + b_acc[i];
                       }
                     });
    });
  }
  q.wait();
  auto end = std::chrono::steady_clock::now();
  std::cout << "Vector add2 completed on device - took "
            << (end - start).count() << " u-secs\n";
  return ((end - start).count());
} // end VectorAdd2

下面的 kernel 手动展开了循环而不是使用编译器指令（编译器可能会也可能不会根据其内部启发式成本模型遵守该指令）。展开的优点是执行的指令更少，因为循环不必迭代那么多次，从而节省了比较和分支的指令。

size_t VectorAdd3(sycl::queue &q, const IntArray &a, const IntArray &b,
                  IntArray &sum, int iter) {
  sycl::range num_items{a.size()};

  sycl::buffer a_buf(a);
  sycl::buffer b_buf(b);
  sycl::buffer sum_buf(sum.data(), num_items);
  size_t num_groups = 1;
  size_t wg_size = 16;
  auto start = std::chrono::steady_clock::now();
  for (int i = 0; i < iter; i++) {
    q.submit([&](auto &h) {
      // Input accessors
      sycl::accessor a_acc(a_buf, h, sycl::read_only);
      sycl::accessor b_acc(b_buf, h, sycl::read_only);
      // Output accessor
      sycl::accessor sum_acc(sum_buf, 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)
                         [[intel::reqd_sub_group_size(16)]] {
                           // Manual unrolling
                           size_t loc_id = index.get_local_id();
                           for (size_t i = loc_id; i < mysize; i += 32) {
                             sum_acc[i] = a_acc[i] + b_acc[i];
                             sum_acc[i + 16] = a_acc[i + 16] + b_acc[i + 16];
                           }
                         });
    });
  }
  q.wait();
  auto end = std::chrono::steady_clock::now();
  std::cout << "Vector add3 completed on device - took "
            << (end - start).count() << " u-secs\n";
  return ((end - start).count());
} // end VectorAdd3

下面的 kernel 显示了如何对 load 和 store 进行重新排序，以便在对它们进行任何操作之前所有 load 都被发出。通常，GPU 中的每个线程都可以有许多未完成的 load 。好的方式是在对它们进行任何操作之前发出 load 。这样才可能在实际需要数据进行计算之前完成 load 操作。

size_t VectorAdd4(sycl::queue &q, const IntArray &a, const IntArray &b,
                  IntArray &sum, int iter) {
  sycl::range num_items{a.size()};

  sycl::buffer a_buf(a);
  sycl::buffer b_buf(b);
  sycl::buffer sum_buf(sum.data(), num_items);
  size_t num_groups = 1;
  size_t wg_size = 16;
  auto start = std::chrono::steady_clock::now();
  for (int i = 0; i < iter; i++) {
    q.submit([&](auto &h) {
      // Input accessors
      sycl::accessor a_acc(a_buf, h, sycl::read_only);
      sycl::accessor b_acc(b_buf, h, sycl::read_only);
      // Output accessor
      sycl::accessor sum_acc(sum_buf, 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)
                         [[intel::reqd_sub_group_size(16)]] {
                           // Manual unrolling
                           size_t loc_id = index.get_local_id();
                           for (size_t i = loc_id; i < mysize; i += 32) {
                             int t1 = a_acc[i];
                             int t2 = b_acc[i];
                             int t3 = a_acc[i + 16];
                             int t4 = b_acc[i + 16];
                             sum_acc[i] = t1 + t2;
                             sum_acc[i + 16] = t3 + t4;
                           }
                         });
    });
  }
  q.wait();
  auto end = std::chrono::steady_clock::now();
  std::cout << "Vector add4 completed on device - took "
            << (end - start).count() << " u-secs\n";
  return ((end - start).count());
} // end VectorAdd4

下面的 kernel 使用了 restrict 指令，它向编译器提供了一个提示，即循环内访问的向量之间没有别名，编译器可以像在前面的示例中手动完成的那样，将 load 提升到 store 之前。

size_t VectorAdd5(sycl::queue &q, const IntArray &a, const IntArray &b,
                  IntArray &sum, int iter) {
  sycl::range num_items{a.size()};

  sycl::buffer a_buf(a);
  sycl::buffer b_buf(b);
  sycl::buffer sum_buf(sum.data(), num_items);
  size_t num_groups = 1;
  size_t wg_size = 16;
  auto start = std::chrono::steady_clock::now();
  for (int i = 0; i < iter; i++) {
    q.submit([&](auto &h) {
      // Input accessors
      sycl::accessor a_acc(a_buf, h, sycl::read_only);
      sycl::accessor b_acc(b_buf, h, sycl::read_only);
      // Output accessor
      sycl::accessor sum_acc(sum_buf, 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) [[intel::reqd_sub_group_size(
                         16)]] [[intel::kernel_args_restrict]] {
                       // compiler needs to hoist the loads
                       size_t loc_id = index.get_local_id();
                       for (size_t i = loc_id; i < mysize; i += 32) {
                         sum_acc[i] = a_acc[i] + b_acc[i];
                         sum_acc[i + 16] = a_acc[i + 16] + b_acc[i + 16];
                       }
                     });
    });
  }
  q.wait();
  auto end = std::chrono::steady_clock::now();
  std::cout << "Vector add5 completed on device - took "
            << (end - start).count() << " u-secs\n";
  return ((end - start).count());
} // end VectorAdd5

TechVideoGPU

关注

0
点赞
踩
0

收藏

觉得还不错? 一键收藏
0
评论
oneAPI GPU 优化指南 - 指针别名和 Restrict 指令

在某些情况下，获得更高性能的代码的更好方法是显式地指定 work-group 的数量和大小，向编译器提供提示以获得良好性能。这将允许编译器提升 load 和 store 次序，从而为指令完成提供更多时间并获得更好的指令调度。展开的优点是执行的指令更少，因为循环不必迭代那么多次，从而节省了比较和分支的指令。下面的 kernel 使用了 restrict 指令，它向编译器提供了一个提示，即循环内访问的向量之间没有别名，编译器可以像在前面的示例中手动完成的那样，将 load 提升到 store 之前。
复制链接

扫一扫