oneAPI GPU 优化指南 - 本地 barrier 与全局原子操作

本章节翻译by chenshusmail@163.com 原文：Local Barriers vs Global Atomics (intel.com)

原子操作可以让 kernel 中的多个 work-item 操作共享资源。 barrier 可以让 work-group 中的 work-item 同步。通过谨慎使用 kernel 启动和本地 barrier， 可以实现全局原子操作的功能。根据架构和涉及的数据量，这两者之一可能有更好的性能表现。

在下面的示例中，我们尝试对向量中相对较少的元素求和。这项任务可以通过不同的方式实现。 下面显示的第一个 kernel 仅使用一个 work-item 遍历向量的所有元素并将它们相加。

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
      h.parallel_for(data_size, [=](auto index) {
        int glob_id = index[0];
        if (glob_id == 0) {
          int sum = 0;
          for (int i = 0; i < N; i++)
            sum += buf_acc[i];
          sum_acc[0] = sum;
        }
      });
    });

在下面显示的 kernel 中，使用全局原子操作解决了相同的问题， 其中每个 work-item 使用它需要累积的值更新全局变量。尽管这里有很多并行操作， 但全局变量的争用非常频繁，在大多数情况下，其性能不会很好。

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);

      h.parallel_for(data_size, [=](auto index) {
        size_t glob_id = index[0];
        auto v = sycl::atomic_ref<int, sycl::memory_order::relaxed,
                                  sycl::memory_scope::device,
                                  sycl::access::address_space::global_space>(
            sum_acc[0]);
        v.fetch_add(buf_acc[glob_id]);
      });
    });

在下面的 kernel 中，每个 work-item 负责累积向量中的多个元素。 这种累积是并行完成的，然后更新到 work-group 中所有 work-item 共享的数组中。 此时，work-group 中的所有 work-item 使用 barrier 在彼此之间同步， 以在共享内存中将中间结果减少到最终结果。 这个 kernel 显式地创建了一个 work-group， 并将向量中所有元素的责任分配给 work-group 中的 work-item。 尽管它没有使用机器在线程数量方面的全部能力，但有时这种并行性对于小问题规模来说足够了。

    Timer timer;
    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      sycl::accessor sum_acc(sum_buf, h, sycl::write_only, sycl::no_init);
      sycl::local_accessor<int, 1> scratch(work_group_size, h);
      h.parallel_for(sycl::nd_range<1>{work_group_size, work_group_size},
                     [=](sycl::nd_item<1> item) {
                       size_t loc_id = item.get_local_id(0);
                       int sum = 0;
                       for (int i = loc_id; i < data_size; i += num_work_items)
                         sum += buf_acc[i];
                       scratch[loc_id] = sum;
                       for (int i = work_group_size / 2; i > 0; i >>= 1) {
                         item.barrier(sycl::access::fence_space::local_space);
                         if (loc_id < i)
                           scratch[loc_id] += scratch[loc_id + i];
                       }
                       if (loc_id == 0)
                         sum_acc[0] = scratch[0];
                     });
    });

这三个 kernel 的性能在不同平台之间差异很大，开发人员需要选择适合其应用程序和硬件的技术。

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