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

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

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