oneAPI GPU 优化指南 - 移除条件检查

翻译已于 2023-12-08 11:48:55 修改 · 193 阅读

0 ·

本内容遵循CC 4.0 BY-SA版权协议

原文链接：https://www.intel.com/content/www/us/en/docs/oneapi/optimization-guide-gpu/2023-1/removing-conditional-checks.html

GEO检测

标签

#oneapi

于 2023-11-17 11:52:08 首次发布

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

46 篇文章

订阅专栏

本文讨论了如何通过移除条件检查和使用关系函数（如min()和max()）来优化SIMD向量计算性能，以提升共享本地内存（SLM）中的卷积操作效率。作者通过示例展示了在SYCL中如何实现这一优化过程。

本章节翻译by chenshusmail@163.com 原文：Removing Conditional Checks (intel.com)

填充缓冲区以移除条件检查

用关系函数替换条件检查

在子组（sub-group）和 SIMD 向量化中, 我们了解到 SIMD 分歧会对性能产生负面影响。如果 sub-group 中的所有 work-item 执行相同的指令，则 SIMD 通道得到最大程度的利用。如果一个或多个 work-item 走了不同的路径，则两条路径都必须执行才能合并。

分歧是由条件检查引起的，尽管并非所有条件检查都会导致分歧。一些条件检查，即使它们不会导致 SIMD 分歧，仍然可能是性能隐患。总的来说，移除条件检查有助于提高性能。

填充缓冲区以移除条件检查

来看看共享本地内存（SLM）中的卷积示例:

    sycl::buffer<int> ibuf(input.data(), N);
    sycl::buffer<int> obuf(output.data(), N);
    sycl::buffer<int> kbuf(kernel.data(), M);

    auto e = q.submit([&](auto &h) {
      sycl::accessor iacc(ibuf, h, sycl::read_only);
      sycl::accessor oacc(obuf, h);
      sycl::accessor kacc(kbuf, h, sycl::read_only);

      h.parallel_for(sycl::nd_range<1>(sycl::range{N}, sycl::range{256}),
                     [=](sycl::nd_item<1> it) {
                       int i = it.get_global_linear_id();
                       int group = it.get_group()[0];
                       int gSize = it.get_local_range()[0];

                       int t = 0;
                       int _M = static_cast<int>(M);
                       int _N = static_cast<int>(N);

                       if ((group == 0) || (group == _N / gSize - 1)) {
                         if (i < _M / 2) {
                           for (int j = _M / 2 - i, k = 0; j < _M; ++j, ++k) {
                             t += iacc[k] * kacc[j];
                           }
                         } else {
                           if (i + _M / 2 >= _N) {
                             for (int j = 0, k = i - _M / 2;
                                  j < _M / 2 + _N - i; ++j, ++k) {
                               t += iacc[k] * kacc[j];
                             }
                           } else {
                             for (int j = 0, k = i - _M / 2; j < _M; ++j, ++k) {
                               t += iacc[k] * kacc[j];
                             }
                           }
                         }
                       } else {
                         for (int j = 0, k = i - _M / 2; j < _M; ++j, ++k) {
                           t += iacc[k] * kacc[j];
                         }
                       }

                       oacc[i] = t;
                     });
    });

嵌套的 if-then-else 条件检查是必要的，以检查输入中的第一个和最后一个 128 个元素，这样索引就不会越界。如果我们在输入数组前后填充足够多的 0，这些条件检查就可以被安全地移除：

  std::vector<int> input(N + M / 2 + M / 2);
  std::vector<int> output(N);
  std::vector<int> kernel(M);

  srand(2009);
  for (size_t i = M / 2; i < N + M / 2; ++i) {
    input[i] = rand();
  }

  for (size_t i = 0; i < M / 2; ++i) {
    input[i] = 0;
    input[i + N + M / 2] = 0;
  }

  for (size_t i = 0; i < M; ++i) {
    kernel[i] = rand();
  }

  {
    sycl::buffer<int> ibuf(input.data(), N + M / 2 + M / 2);
    sycl::buffer<int> obuf(output.data(), N);
    sycl::buffer<int> kbuf(kernel.data(), M);

    auto e = q.submit([&](auto &h) {
      sycl::accessor iacc(ibuf, h, sycl::read_only);
      sycl::accessor oacc(obuf, h);
      sycl::accessor kacc(kbuf, h, sycl::read_only);

      h.parallel_for(sycl::nd_range(sycl::range{N}, sycl::range{256}),
                     [=](sycl::nd_item<1> it) {
                       int i = it.get_global_linear_id();
                       int t = 0;

                       for (size_t j = 0; j < M; ++j) {
                         t += iacc[i + j] * kacc[j];
                       }

                       oacc[i] = t;
                     });
    });
    q.wait();

    size_t kernel_ns = (e.template get_profiling_info<
                            sycl::info::event_profiling::command_end>() -
                        e.template get_profiling_info<
                            sycl::info::event_profiling::command_start>());
    std::cout << "Kernel Execution Time Average: total = " << kernel_ns * 1e-6
              << " msec" << std::endl;
  }

用关系函数替换条件检查

另一种移除条件检查的方法是用关系函数替换它们，特别是内置关系函数。强烈建议使用内置函数（如果有）。 SYCL 提供了丰富的内置关系函数，如 select(), min(), max()。在许多情况下，您可以使用这些函数替换条件检查并获得更好的性能。

再看看上边的卷积示例。if-then-else 条件检查可以用内置函数 min() and max() 来替换。

    sycl::buffer<int> ibuf(input.data(), N);
    sycl::buffer<int> obuf(output.data(), N);
    sycl::buffer<int> kbuf(kernel.data(), M);

    auto e = q.submit([&](auto &h) {
      sycl::accessor iacc(ibuf, h, sycl::read_only);
      sycl::accessor oacc(obuf, h);
      sycl::accessor kacc(kbuf, h, sycl::read_only);

      h.parallel_for(sycl::nd_range(sycl::range{N}, sycl::range{256}),
                     [=](sycl::nd_item<1> it) {
                       int i = it.get_global_linear_id();
                       int t = 0;
                       int startj = sycl::max<int>(M / 2 - i, 0);
                       int endj = sycl::min<int>(M / 2 + N - i, M);
                       int startk = sycl::max<int>(i - M / 2, 0);
                       for (int j = startj, k = startk; j < endj; j++, k++) {
                         t += iacc[k] * kacc[j];
                       }
                       oacc[i] = t;
                     });
    });