oneAPI GPU 优化指南 - 归约运算

已于 2023-12-15 11:53:12 修改
阅读量115 收藏
点赞数
分类专栏： oneAPI GPU 优化指南 文章标签： oneapi c++ 1024程序员节
于 2023-12-15 11:41:19 首次发布
原文链接：https://www.intel.com/content/www/us/en/docs/oneapi/optimization-guide-gpu/2023-1/reduction.html
版权

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

归约是并行编程中的一种常见操作，其中将运算符应用于数组的所有元素并生成单个结果。 归约运算符是关联的，在某些情况下也是可交换的。一些归约的例子包括求和、最大值和最小值。 下面展示了一个串行求和归约：

  for (int it = 0; it < iter; it++) {
    sum = 0;
    for (size_t i = 0; i < data_size; ++i) {
      sum += data[i];
    }
  }

归约的时间复杂度与元素数量成线性关系。有几种方法可以使这个过程并行化， 必须注意确保不同被处理元素之间的通信/同步量最小化。 一种简单的并行化归约的方法是使用全局变量，并让线程使用原子操作更新此变量：

    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 为每个块计算本地总和，然后对本地总和进行顺序/树归约。 块的数量取决于平台上存在的计算单元数量。可以使用设备对象上的 get_info<info::device::max_compute_units>() 函数查询这个值：

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      sycl::accessor accum_acc(accum_buf, h, sycl::write_only, sycl::no_init);
      h.parallel_for(num_processing_elements, [=](auto index) {
        size_t glob_id = index[0];
        size_t start = glob_id * BATCH;
        size_t end = (glob_id + 1) * BATCH;
        if (end > N)
          end = N;
        int sum = 0;
        for (size_t i = start; i < end; ++i)
          sum += buf_acc[i];
        accum_acc[glob_id] = sum;
      });
    });

这个 kernel 的性能将优于原子更新共享内存位置的 kernel。但是，它仍然效率低下， 因为编译器无法对循环进行向量化。让编译器生成向量代码的一种方法是修改循环，如下所示：

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

编译器可以将此代码向量化以使性能更好。

在使用 GPU 的情况下，每个物理处理器都有许多线程上下文，称为向量引擎 (VE) 或执行单元 (EU)。 因此，上述代码中线程数等于 VE 数量的情况并未利用所有线程上下文。 即使在每个核心有两个超线程的 CPU 情况下，代码也不会使用所有线程上下文。 通常，最好将工作分成足够多的 work-group 以获得所有线程上下文的完全占用。 这样可以使代码更好地容忍长延迟指令。下表显示了不同设备中每个处理元素可用的线程上下文数量：

设备可用的线程上下文数量

VEs

单个 VE 线程数

总线程数

KBL24724 x 7 = 168
TGL96796 x 7 = 672

下面的代码显示了一个具有足够线程以充分利用可用资源的 kernel。请注意，无法从设备查询可用线程上下文的数量。 因此，根据设备的不同，您可以扩展配置创建的 work-item 数量，以在它们之间分配工作负载。

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

在 GPU 上进行归约操作的一种流行方法是创建多个 work-group 并在每个 work-group 中进行树形归约。 在下面显示的 kernel 中，work-group 中的每个 work-item 都参与归约网络， 最终将该 work-group 中的所有元素求和。然后通过串行归约将所有 work-group 的中间结果求和 （如果这组中间结果足够大，则可以再进行几轮树形归约）。 这种树形归约算法利用了 work-group 中 work-item 之间非常快速的同步操作。 这个 kernel 的性能高度依赖于 kernel 启动的效率，因为启动了大量的 kernel。此外， 下面所写的 kernel 并不是很高效，因为每次循环中实际工作的线程数量呈指数级减少。

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      sycl::accessor accum_acc(accum_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>(num_work_items, work_group_size),
                     [=](sycl::nd_item<1> item) {
                       size_t global_id = item.get_global_id(0);
                       int local_id = item.get_local_id(0);
                       int group_id = item.get_group(0);

                       if (global_id < data_size)
                         scratch[local_id] = buf_acc[global_id];
                       else
                         scratch[local_id] = 0;

                       // Do a tree reduction on items in work-group
                       for (int i = work_group_size / 2; i > 0; i >>= 1) {
                         item.barrier(sycl::access::fence_space::local_space);
                         if (local_id < i)
                           scratch[local_id] += scratch[local_id + i];
                       }

                       if (local_id == 0)
                         accum_acc[group_id] = scratch[0];
                     });
    });

单阶段归约效率不高，因为它会给主机端留下很多工作负载。增加一个阶段将减少主机端的工作量，并显著提高性能。 可以看到，在下面的 kernel 中，第一阶段计算的中间结果被用作第二阶段的输入。这可以推广为形成多级归约， 直到结果足够小，以便在主机端执行。

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      sycl::accessor accum_acc(accum1_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>(num_work_items1, work_group_size),
                     [=](sycl::nd_item<1> item) {
                       size_t global_id = item.get_global_id(0);
                       int local_id = item.get_local_id(0);
                       int group_id = item.get_group(0);

                       if (global_id < data_size)
                         scratch[local_id] = buf_acc[global_id];
                       else
                         scratch[local_id] = 0;

                       // Do a tree reduction on items in work-group
                       for (int i = work_group_size / 2; i > 0; i >>= 1) {
                         item.barrier(sycl::access::fence_space::local_space);
                         if (local_id < i)
                           scratch[local_id] += scratch[local_id + i];
                       }

                       if (local_id == 0)
                         accum_acc[group_id] = scratch[0];
                     });
    });
    q.submit([&](auto &h) {
      sycl::accessor buf_acc(accum1_buf, h, sycl::read_only);
      sycl::accessor accum_acc(accum2_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>(num_work_items2, work_group_size),
                     [=](sycl::nd_item<1> item) {
                       size_t global_id = item.get_global_id(0);
                       int local_id = item.get_local_id(0);
                       int group_id = item.get_group(0);

                       if (global_id < static_cast<size_t>(num_work_items2))
                         scratch[local_id] = buf_acc[global_id];
                       else
                         scratch[local_id] = 0;

                       // Do a tree reduction on items in work-group
                       for (int i = work_group_size / 2; i > 0; i >>= 1) {
                         item.barrier(sycl::access::fence_space::local_space);
                         if (local_id < i)
                           scratch[local_id] += scratch[local_id + i];
                       }

                       if (local_id == 0)
                         accum_acc[group_id] = scratch[0];
                     });
    });

SYCL 还支持内置的归约操作，应在适当的情况下使用，因为其实现经过了对底层架构的优化。 下面的 kernel 显示了如何在编译器中使用内置的归约运算符。

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      auto sumr = sycl::reduction(sum_buf, h, sycl::plus<>());
      h.parallel_for(sycl::nd_range<1>{data_size, 256}, sumr,
                     [=](sycl::nd_item<1> item, auto &sumr_arg) {
                       int glob_id = item.get_global_id(0);
                       sumr_arg += buf_acc[glob_id];
                     });
    });

进一步的优化是阻止对输入向量的访问并使用 SLM 来存储中间结果。下面展示了这个 kernel。 在这个 kernel 中，每个 work-item 操作一定数量的向量元素， 然后 work-group 中的一个线程通过线性遍历包含中间结果的 SLM 将所有这些元素归约为一个结果。

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      sycl::accessor accum_acc(accum_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>{num_work_items, work_group_size},
                     [=](sycl::nd_item<1> item) {
                       size_t glob_id = item.get_global_id(0);
                       size_t group_id = item.get_group(0);
                       size_t loc_id = item.get_local_id(0);
                       int offset = ((glob_id >> log2workitems_per_block)
                                     << log2elements_per_block) +
                                    (glob_id & mask);
                       int sum = 0;
                       for (int i = 0; i < elements_per_work_item; ++i)
                         sum +=
                             buf_acc[(i << log2workitems_per_block) + offset];
                       scratch[loc_id] = sum;
                       // 串行归约操作
                       item.barrier(sycl::access::fence_space::local_space);
                       if (loc_id == 0) {
                         int sum = 0;
                         for (int i = 0; i < work_group_size; ++i)
                           sum += scratch[i];
                         accum_acc[group_id] = sum;
                       }
                     });
    });

下面的 kernel 与上面的类似，只是使用树形归约来减少 work-group 中所有 work-item 的中间结果。 在大多数情况下，这似乎并没有对性能产生很大的影响。

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      sycl::accessor accum_acc(accum_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>{num_work_items, work_group_size},
                     [=](sycl::nd_item<1> item) {
                       size_t glob_id = item.get_global_id(0);
                       size_t group_id = item.get_group(0);
                       size_t loc_id = item.get_local_id(0);
                       int offset = ((glob_id >> log2workitems_per_block)
                                     << log2elements_per_block) +
                                    (glob_id & mask);
                       int sum = 0;
                       for (int i = 0; i < elements_per_work_item; ++i)
                         sum +=
                             buf_acc[(i << log2workitems_per_block) + offset];
                       scratch[loc_id] = sum;
                       // tree reduction
                       item.barrier(sycl::access::fence_space::local_space);
                       for (int i = work_group_size / 2; i > 0; i >>= 1) {
                         item.barrier(sycl::access::fence_space::local_space);
                         if (loc_id < static_cast<size_t>(i))
                           scratch[loc_id] += scratch[loc_id + i];
                       }
                       if (loc_id == 0)
                         accum_acc[group_id] = scratch[0];
                     });
    });

下面的 kernel 使用阻塞技术，然后使用编译器归约运算符进行最终归约。这在测试过的大多数平台上都能获得良好的性能。

    q.submit([&](auto &h) {
      sycl::accessor buf_acc(buf, h, sycl::read_only);
      auto sumr = sycl::reduction(sum_buf, h, sycl::plus<>());
      h.parallel_for(sycl::nd_range<1>{num_work_items, work_group_size}, sumr,
                     [=](sycl::nd_item<1> item, auto &sumr_arg) {
                       size_t glob_id = item.get_global_id(0);
                       int offset = ((glob_id >> log2workitems_per_block)
                                     << log2elements_per_block) +
                                    (glob_id & mask);
                       int sum = 0;
                       for (int i = 0; i < elements_per_work_item; ++i)
                         sum +=
                             buf_acc[(i << log2workitems_per_block) + offset];
                       sumr_arg += sum;
                     });
    });

这个 kernel 使用了一种完全不同的内存访问技术。它使用 sub-group 加载来以向量形式生成中间结果。 然后将这个中间结果带回主机并在那里执行最终归约。在某些情况下， 最好创建另一个 kernel 来在单个 work-group 中减少这个结果，这样可以通过高效的 barrier 来执行树形归约。

    q.submit([&](auto &h) {
      const sycl::accessor buf_acc(buf, h);
      sycl::accessor accum_acc(accum_buf, h, sycl::write_only, sycl::no_init);
      sycl::local_accessor<sycl::vec<int, 8>, 1l> scratch(work_group_size, h);
      h.parallel_for(
          sycl::nd_range<1>{num_work_items, work_group_size},
          [=](sycl::nd_item<1> item) [[intel::reqd_sub_group_size(16)]] {
            size_t group_id = item.get_group(0);
            size_t loc_id = item.get_local_id(0);
            sycl::sub_group sg = item.get_sub_group();
            sycl::vec<int, 8> sum{0, 0, 0, 0, 0, 0, 0, 0};
            using global_ptr =
                sycl::multi_ptr<int, sycl::access::address_space::global_space>;
            int base = (group_id * work_group_size +
                        sg.get_group_id()[0] * sg.get_local_range()[0]) *
                       elements_per_work_item;
            for (int i = 0; i < elements_per_work_item / 8; ++i)
              sum += sg.load<8>(global_ptr(&buf_acc[base + i * 128]));
            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 < static_cast<size_t>(i))
                scratch[loc_id] += scratch[loc_id + i];
            }
            if (loc_id == 0)
              accum_acc[group_id] = scratch[0];
          });
    });

这里提供并讨论了归约操作的不同实现，这些实现可能具有不同的性能特征，具体取决于加速器的架构。 另一个需要注意的重要事项是， 将归约结果通过 PCIe 接口（对于离散 GPU ）带回主机所需的时间几乎与在设备上实际执行整个归约所需的时间相同。 这表明应尽可能避免主机和设备之间的数据传输或将 kernel 执行与数据传输重叠。

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

TechVideoGPU
关注
  • 0
    点赞
  • 踩
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
面向GPU计算平台的归约算法的性能优化研究.pdf
09-25
面向GPU计算平台的归约算法的性能优化研究.pdf
6.CUDA编程手册中文版---附录A&B
热门推荐
专注于人工智能领域的小何尚
04-16 1万+
附录A 支持GPU设备列表 https://developer.nvidia.com/cuda-gpus 列出了所有支持 CUDA 的设备及其计算能力。 可以使用运行时查询计算能力、多处理器数量、时钟频率、设备内存总量和其他属性（参见参考手册）。 附录B 对C++扩展的详细描述 B.1 函数执行空间说明符 函数执行空间说明符表示函数是在主机上执行还是在设备上执行，以及它是可从主机调用还是从设备调用。 B.1.1 __global__ __global__ 执行空间说明符将函数声明为内核。 它的功能是： 在
尹成学院区块链 Go 学习大纲-取得大纲试看视频联系微信yinchengak48
尹成的技术博客
08-21 4614
所处阶段 主讲内容 技术要点 学习目标 第一阶段Go语言开发入门实战 1.Go语言介绍及开发环境搭建 1.Go语言是什么 2.Go语言优势 3.Go语言适合来做什么 4.Go语言安装和设置 5.标准命令概述 6.第一个Go语言程序 课程设计理念： 夯实基础，Go语言作为...
oneAPI GPU 优化指南 - 内核（kernel）
TechVideoGPU的博客
11-17 74
你可以在一个迭代空间上提交一个 kernel，这样就可以对指定的数据对象进行计算。在 oneAPI 卸载模型中，kernel 是执行计算的单元。本节将介绍与 kernel 编程、提交和执行相关的内容。
oneAPI GPU 优化指南
TechVideoGPU的博客
09-16 206
oneAPI GPU Optimization Guide 的中文版
使用 TensorFlow 2.0 进行分布式训练
一个幽灵
05-12 7707
概览 tf.distribute.Strategy 是一个用于分布式训练的TensorFlow API ，横跨多GPU、多机器或TPU. 通过这个API，经过少量改造，就可以让现存的模型和训练代码支持分布式训练。 tf.distribute.Strategy 的设计目标是: 容易使用，支持多种用户，包括研究员、机器学习工程师等等 提供良好的性能、开箱即用 更容易更换策略 tf.distribute.Strategy 能被用于高级API比如 Keras, 也能用于自定义训练循环 (基于TensorFlow
基于映射-归约分布式架构的城市轨道交通信号系统检测平台.pdf
08-08
#资源达人分享计划#
CUDA专家手册 GPU编程权威指南
02-24
《CUDA专家手册：GPU编程权威指南》详细讨论CUDA的硬件和软件，包括CUDA 5.0和开普勒架构的最新特性。每个CUDA开发人员，不论新手还是高手，都可以在这里找到感兴趣的内容并即时上手。新晋的CUDA开发者将理解硬件...
First_order1:区间类型2模糊系统的Maclaurin级数展开复杂度-归约集中心类型归约+去模糊化-matlab开发
05-31
这是 Maclaurin 级数展开复杂度的代码 - 降低集中心类型 - 归约 + 区间类型 2 模糊系统的去模糊化。 可以在以下位置找到这种类型减少器+除颤器的描述： Khanesar, MA 和 Mendel, JM，2016 年，7 月。 区间类型 2 ...
3-SAT归约到独立集问题
01-06
3-SAT归约到独立集问题 【3-SAT ≤p\leq_p≤p​ 独立集】 要证明3-SAT问题可以归约到独立集，就需要证明，有一个关于独立集的黑盒子，通过解3-SAT实例，能够解3-SAT问题。 图4为从3-SAT到独立集归约的一个实例。 图...
c++父类指针指向子类
最新发布
qq_43700779的博客
05-11 105
父类指针指向子类，调用的函数是谁的
初识C语言——第十四天
硬件菜鸟的成长之路
05-07 690
【代码】初识C语言——第十四天。
UE5C++ FString做为参数取值时报错error：C4840
YanisWu的博客
05-11 89
UE5 C++ 的FString类型取值问题。
webassembly入门详解（C++）
啊渊的专栏
05-05 1026
webassembly C++ windows下开发环境搭建。
C++ STL的锁介绍
05-07 232
C++ STL的锁介绍
C++ list 介绍
2301_77203593的博客
05-08 650
构造空list构造可包含n个T类型数据的list构造list，用first到last区间中的元素构造。注意可以是别的容器中的元素。（例如可以用vectorint<int>int中的一段数据构造一个新的listint<int>int。还可以用string等等）。拷贝构造。用已有的一个list构造新的list。注：前提是原链表和链表x都为有序链表，默认升序。将有序链表x按照默认的升序合并至原链表中，合并后x变为空链表。
effective c++ 和 more effective c++中知识点
zengchenAAA的博客
05-11 254
尽可能使用 const在 C++ 编程中，尽可能使用。
直方图-数值归约技术介绍
06-11
直方图是一种数据可视化的方式，它将数据分成一些离散的区间，并统计每个区间中数据的数量或频率。直方图可以帮助我们更直观地了解数据的分布情况，例如数据的集中程度、分散程度和偏态程度等。直方图也可以用于数值归约技术。 数值归约是一种数据预处理技术，它旨在减少数据集的规模，同时保留数据的重要特征。直方图可以用于数值归约，具体做法是将原始数据分成若干个区间，并对每个区间进行一些统计计算，例如计算每个区间的均值、中位数或众数等。然后将这些统计结果作为新的数据集，从而实现数据的归约。 数值归约技术可以提高数据挖掘算法的效率，同时可以去除数据集中的噪声和冗余信息，从而提高数据挖掘的准确性和可靠性。直方图作为一种常用的数据可视化工具，可以帮助我们更好地理解数据的分布情况，并为数值归约提供有效的工具和方法。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值