oneAPI GPU 优化指南 - 线程映射和 GPU 占用

TechVideoGPU

已于 2024-01-04 16:25:02 修改

阅读量159

点赞数

分类专栏： oneAPI GPU 优化指南文章标签： oneapi

于 2023-11-09 16:44:48 首次发布

原文链接：https://www.intel.com/content/www/us/en/docs/oneapi/optimization-guide-gpu/2023-1/thread-mapping-and-gpu-occupancy.html

版权

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

46 篇文章 1 订阅

订阅专栏

本章节翻译by chenchensmail@163.com 原文：Thread Mapping and GPU Occupancy (intel.com)

nd_range

线程同步

将 work-group 映射到 Xe-core 以实现最大占用率

work-group 内 work-item 同步的影响

work-group 内本地内存的影响

一个详细的例子

Intel® GPU 占用计算器

SYCL 执行模型公开了 GPU 执行的抽象视图。 SYCL 线程层次结构由 1、2 或 3 维 work-item 网格组成。这些工作项被分组为大小相等的线程组，称为 work-group 。 work-group 中的线程进一步分为大小相等的向量组称为 sub-group （参见下图）。

Work-item

work-item 代表 kernel 函数的并行执行的集合之一。

Sub-group

子组代表一小部分连续的 work-item ，这些 work-item 是作为长度是 8、16、32 或是带有 intel UHD 显卡的 CPU 原生向量长度倍数的 SIMD 向量一起处理的。

Work-group

work-group 是线程层次结构中 1、2 或 3 维线程集。在 SYCL 中，跨 work-item 的同步只能是在同一个 work-group 内带有 barrier 的 work-item 之间。

nd_range

nd_range 将线程层次结构分为 1、2、或 3 维 work-group 网格。它代表了全局范围，每个 work-group 的局部范围。

线程层次结构

上图说明了 ND-Range、work-group、 sub-group 和 work-item 之间的关系。

线程同步

SYCL提供了两种可以调用 kernel 函数的同步机制。两者都仅针对同一个 work-group 内的 work-item。 SYCL 在一个 kernel 里面不提供整个 nd_range 内跨 work-item 的全局同步机制。

``mem_fence`` 在同一个 work-group 中跨 work-item 把内存 fence 插入全局和本地内存访问。
``barrier`` 插入内存 fence 并阻止所有 work-group 中 work-item 的执行，直到所有 work-item 都达到位置。

将 work-group 映射到 Xe-core 以实现最大占用率

本章的其余部分解释如何选择合适的 work-group 大小以最大化 GPU 资源的占用。示例系统是以 Xe-LP GPU 作为执行目标的 Tiger Lake 处理器。这个示例还使用新术语 Xe-core (XC) 表示 Dual Subslice, 和 Xe 矢量引擎 (XVE) 表示执行单元。

我们将使用 Xe-LP Graphics (TGL) GPU 的架构参数总结如下：

Xe-LP (TGL) GPU
	VEs	线程	操作	最大 Work Group 大小
Each Xe-core	16	7×16=112	112×8=896	512
Total	16×6=96	112×6=672	896×6=5376	512

最大的 work-group 大小是由硬件和 GPU 驱动程序约束。你可以使用 device::get_info<cl::sycl::info::device::max_work_group_size>() function 查询最大 work-group 大小。

我们以一个简单的 kernel 开始:

auto command_group =
    [&](auto &cgh) {
      cgh.parallel_for(sycl::range<3>(64, 64, 64), // global range
                       [=](item<3> it) {
                         // (kernel code)
                       })
    }

这个 kernel 包含 262,144 个 work-item, 结构为 3D 范围 64×64×64。它让编译器来选择 work-group 和 sub-group的大小。为了充分利用 GPU slice 中可用的5376并行运算, 编译器必须选择合适的 work-group 大小。

两个最重要的 GPU 资源是:

线程上下文:: kernel 应该有足够数量的线程来利用 GPU 的线程上下文。
SIMD 单元和 SIMD 寄存器:: kernel 应被组织来矢量化 work-item 并采用 SIMD 寄存器。

在 SYCL kernel 中，程序员可以通过以下方式影响工作分配：使用适当的 work-group 大小和 sub-group 大小来构造 kernel, 组织 work-item 以实现高效的向量执行。写高效的向量 kernel 将在单独的部分中介绍。这本章重点介绍work-group 和 sub-group 大小的选择。

线程上下文比 SIMD 向量更容易使用。因此，从选择 work-group 中的线程数开始每个 Xe-core 有 112 个线程上下文，但如果kernel 也是 8 位矢量化的，通常你不能使用所有线程。 (112×8=896>512)。由此，我们可以推导出 work-group 中的最大线程数为 64 (512 / 8)。

SYCL 没有提供直接在 work-group 中设置线程数的机制。但是，你可以使用 work-group 大小和 sub-group 大小设置线程数：

$Workgroupsize = Threads \times Sub-groupSize$

只要在加宽后有足够的 kernel 寄存器数量，你就可以增加 sub-group 大小。请注意，每个 VE 有 128 个 SIMD8 寄存器，因此在简单的 kernel 上有很大的扩展空间。增加 sub-group 大小的效果类似于循环展开：虽然每个 VE 在每个周期仍执行 8 个 32 位操作，每次 work-group 交互的工作量变为两倍/四倍。在 SYCL, 程序员可以使用 intel::reqd_sub_group_size({8|16|32}) 显式指定 sub-group 大小以覆盖编译器的选择。

下表总结了使 Intel® Iris® Xe-LP GPU 占用所有 GPU 资源，线程和 sub-group 大小的选择标准：

保证最大占用率的配置
最大线程数	最小 Sub-group 大小	最大 Sub-group 大小	最大 Work-group 大小	约束
64	8	32	512	$Threads \times Sub-groupSize < = 512$

一般来说，选择更大的 work-group 大小有助于减少 work-group 调度的轮数。增加 sub-group 大小可以减少 work-group 所需的线程数，但会以每个 sub-group 执行的延迟更长和寄存器压力更高为代价。

work-group 内 work-item 同步的影响

我们来看一下需要 work-item 同步的kernel:

 auto command_group =
     [&](auto &cgh) {
       cgh.parallel_for(nd_range(sycl::range(64, 64, 128), // global range
                                 sycl::range(1, R, 128)    // local range
                                 ),
                        [=](sycl::nd_item<3> item) {
                          // (kernel code)
                          // Internal synchronization
                          item.barrier(access::fence_space::global_space);
                          // (kernel code)
                        })
     }

这个 kernel 与前面的例子类似，但需要 work-group barrier 同步。work-item 同步仅适用于同一 work-group 内的 work-item。你必须使用 nd_range 和 nd_item 来选择 work-group 本地范围。 work-group 的所有 work-item 必须分配到同一个 Xe-core, 这会影响 Xe-core 的占用率和 kernel 性能。

在这个内核中, work-group 的本地范围为 range(1, R, 128)。假设 sub-group 大小为 8, 让我们看看变量 R 的值如何影响 VE 占用率。在 R=1 的情况下，本地组范围为 (1, 1, 128), work-group 大小为 128。分配给 work-group 的 Xe-core 只包含 16 个线程，而可用线程上下文中有 112 个线程（即占用率非常低）。但是，系统可以将 7 个 work-group 分派到同一个 Xe-core 上，以达到完全占用率，代价是更高的调度次数。

在 R>4 的情况下， work-group 大小将超过系统支持的最大 work-group 大小 512, kernel 将无法启动。在 R=4 的情况下，一个 Xe-core 只有 57% 的占用率(4/), 三个未使用的线程上下文不足以容纳另一个 work-group, 这样就浪费了可用 VE 容量的 43%。请注意，驱动程序仍然可以将部分 work-group 分派到未使用的 Xe-core 上。但是，由于 kernel 中的 barrier，部分分派的 work-item 在其余 work-group 分派之前无法通过 barrier。在大多数情况下， kernel 的性能不会从部分分派中获得太多好处。因此，通过正确选择 work-group 大小来避免这个问题非常重要。

下表总结了工作组大小、线程数、 Xe-core 利用率和占用率之间的权衡。

多种配置的利用率
Work-items	组大小	线程数	Xe-core 利用率	Xe-core 占用率
64×64×128=524288	(R=1) 128	16	16/112=14%	100% with 7 work-groups
64×64×128=524288	(R=2) 128×2	2×16=32	32/112=28.6%	86% with 3 work-groups
64×64×128=524288	(R=3) 128×4	3×16=48	48/112=42.9%	86% with 2 work-groups
64×64×128=524288	(R=4) 128×4	4×16=64	64/112=57%	57% maximum
64×64×128=524288	(R>4) 640+			Fail to launch

work-group 内本地内存的影响

让我们看一个 kernel 为工作组分配本地内存的例子：

 auto command_group =
     [&](auto &cgh) {
       // local memory variables shared among work items
       sycl::accessor<int, 1, sycl::access::mode::read_write,
                      sycl::access::target::local>
           myLocal(sycl::range(R), cgh);
       cgh.parallel_for(nd_range(sycl::range<3>(64, 64, 128), // global range
                                 sycl::range<3>(1, R, 128)    // local range
                                 ),
                        [=](ngroup<3> myGroup) {
                          // (work group code)
                          myLocal[myGroup.get_local_id()[1]] = ...
                        })
     }

由于 work-group 本地变量在其 work-item 之间共享，所以它们在 Xe-core 的 SLM 中分配。因此，这个 work-group 必须分配给一个 Xe-core, 与组内同步相同。此外，你还必须权衡不同组大小选项下的本地变量大小，以使本地变量适合 Xe-core 的 128KB SLM 容量限制。

一个详细的例子

在结束本节之前，让我们来看一下简单向量加法示例的变体的硬件占用率。使用 TGL 平台的 Intel® Iris® Xe 图形作为底层硬件，并使用 Xe-LP (TGL) GPU 中指定的资源参数。

 int VectorAdd1(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();
   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) {
       for (int j = 0; j < iter; j++)
         sum_acc[i] = a_acc[i] + b_acc[i];
     });
   });
   q.wait();
   auto end = std::chrono::steady_clock::now();
   std::cout << "VectorAdd1 completed on device - took " << (end - start).count()
             << " u-secs\n";
   return ((end - start).count());
 } // end VectorAdd1

上面的 VectorAdd1 让编译器选择 work-group 大小和 SIMD 宽度。在这种情况下，编译器选择了一个 work-group 大小为 512 和一个宽度为 32 的 SIMD, 因为 kernel 的寄存器压力很低。

 int 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 = groups;
   size_t wg_size = 512;
   // get the max wg_sie instead of 512 size_t wg_size = 512;
   auto start = std::chrono::steady_clock::now();
   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(32)]] {
           size_t grp_id = index.get_group()[0];
           size_t loc_id = index.get_local_id();
           size_t start = grp_id * mysize;
           size_t end = start + mysize;
           for (int j = 0; j < iter; j++)
             for (size_t i = start + loc_id; i < end; i += wg_size) {
               sum_acc[i] = a_acc[i] + b_acc[i];
             }
         });
   });
   q.wait();
   auto end = std::chrono::steady_clock::now();
   std::cout << "VectorAdd2<" << groups << "> completed on device - took "
             << (end - start).count() << " u-secs\n";
   return ((end - start).count());
 } // end VectorAdd2

上面的 VectorAdd2 示例明确指定了 work-group 大小为 51, SIMD 宽度为 32, 以及变量数量的 work-group 作为函数参数组。

将线程数除以 GPU 中可用线程上下文的数量，可以给出 GPU 硬件占用率的估计。下表使用 TGL Intel® Iris® Xe 架构参数计算了上述两个 kernel 在各种参数下的 GPU 硬件占用率。

占用率
项目占用率	Work-groups	Work-items	Work-group 大小	SIMD	线程 Work-group	线程数	占用率
VectorAdd1	53760	27.5M	512	32	16	860K	100%
VectorAdd2<1>	1	512	512	32	16	16	16/672 = 2.4%
VectorAdd2<2>	2	1024	512	32	16	32	32/672 = 4.8%
VectorAdd2<3>	3	1536	512	32	16	48	48/672 = 7.1%
VectorAdd2<4>	4	2048	512	32	16	64	64/672 = 9.5%
VectorAdd2<5>	5	2560	512	32	16	80	80/672 = 11.9%
VectorAdd2<6>	6	3072	512	32	16	96	96/672 = 14.3%
VectorAdd2<7>	7	3584	512	32	16	112	112/672 = 16.7%
VectorAdd2<8>	8	4096	512	32	16	128	128/672 = 19%
VectorAdd2<12>	12	6144	512	32	16	192	192/672 = 28.6%
VectorAdd2<16>	16	8192	512	32	16	256	256/672 = 38.1%
VectorAdd2<20>	20	10240	512	32	16	320	320/672 = 47.7%
VectorAdd2<24>	24	12288	512	32	16	384	384/672 = 57.1%
VectorAdd2<28>	28	14336	512	32	16	448	448/672 = 66.7%
VectorAdd2<32>	32	16384	512	32	16	512	512/672 = 76.2%
VectorAdd2<36>	36	18432	512	32	16	576	576/672 = 85.7%
VectorAdd2<40>	40	20480	512	32	16	640	640/672 = 95.2%
VectorAdd2<42>	42	21504	512	32	16	672	672/672 = 100%
VectorAdd2<44>	44	22528	512	32	16	704	100% then 4.7%
VectorAdd2<48>	48	24576	512	32	16	768	100% then 14.3%

下面是 VectorAdd2 的 VTune 分析器图表，其中包含各种 work-group 大小，证实了我们估计的准确性。网格视图中的数字与估计略有不同，因为网格视图在整个执行过程中给出了平均值。

VTune 中显示的 VectorAdd2 占用率

下面的时间轴视图给出了一段时间内的占用率。请注意，由于每个线程完成执行的时间不同，因此占用率指标在内核执行的大部分时间内是准确的，并且朝向结束时逐渐减少。

VectorAdd2 时间线视图

下面显示的内核 VectorAdd3 与上面的 kernel 类似，但有两个重要的区别。

1. 它可以使用 work-group 数量、work-group 大小和 sub-group 大小作为模板参数进行实例化。这使我们能够进行实验，以研究 sub-group 和 work-group 数量对线程占用率的影响。

2. kernel 内部完成的工作量大大增加，以确保这些 kernel 驻留在执行单元中，以长时间执行工作。

 template <int groups, int wg_size, int sg_size>
 int 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 = groups;
   auto start = std::chrono::steady_clock::now();
   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(sg_size)]] {
           size_t grp_id = index.get_group()[0];
           size_t loc_id = index.get_local_id();
           size_t start = grp_id * mysize;
           size_t end = start + mysize;
           for (int j = 0; j < iter; j++)
             for (size_t i = start + loc_id; i < end; i += wg_size) {
               sum_acc[i] = a_acc[i] + b_acc[i];
             }
         });
   });
   q.wait();
   auto end = std::chrono::steady_clock::now();
   std::cout << "VectorAdd3<" << groups << "> completed on device - took "
             << (end - start).count() << " u-secs\n";
   return ((end - start).count());
 } // end VectorAdd3

kernel VectorAdd4 与上面的 kernel VectorAdd3 类似，只是在 kernel 执行的开始和结束处有一个 barrier 同步。这个 barrier 在功能上是不需要的，但会显著影响线程在硬件上的调度方式。

 template <int groups, int wg_size, int sg_size>
 int 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 = groups;
   auto start = std::chrono::steady_clock::now();
   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(sg_size)]] {
           index.barrier(sycl::access::fence_space::local_space);
           size_t grp_id = index.get_group()[0];
           size_t loc_id = index.get_local_id();
           size_t start = grp_id * mysize;
           size_t end = start + mysize;
           for (int j = 0; j < iter; j++) {
             for (size_t i = start + loc_id; i < end; i += wg_size) {
               sum_acc[i] = a_acc[i] + b_acc[i];
             }
           }
         });
   });
   q.wait();
   auto end = std::chrono::steady_clock::now();
   std::cout << "VectorAdd4<" << groups << "> completed on device - took "
             << (end - start).count() << " u-secs\n";
   return ((end - start).count());
 } // end VectorAdd4

为了展示线程的调度方式，上述两个 kernel 将以 8 个 work-group、 sub-group 大小为 8 和 work-group 大小为320 的方式调用, 如下所示。根据 work-group 大小和 sub-group 大小的选择，每个 work-group 需要由硬件调度 40 个线程。

  Initialize(sum);
  VectorAdd3<8, 320, 8>(q, a, b, sum, 10000);
  Initialize(sum);
  VectorAdd4<8, 320, 8>(q, a, b, sum, 10000);

下面的 VTune 图表显示了 VectorAdd3 和 VectorAdd4 kernel 的测量 GPU 占用率。

VectorAdd3, VectorAdd4 kernel 的GPU 占用率

对于 VectorAdd3 kernel , 占用率有两个阶段：在 TGL 机器上，占用率为 33.3% (224 个线程占用)和 14.3%(96 个线程占用)。由于总共有八个 work-group, 每个 work-group 有 40 个线程，因此有两个 Xe-core （每个核心有 112 个线程）可以调度六个 work-group 的线程。这意味着在第一阶段中，每个来自四个 work-group 的 40 个线程被调度，并且每个来自另外两个 work-group 的 32 个线程都被调度。然后在第二阶段中，剩余两个 work-group 的 40 个线程都被调度执行。

如 VectorAdd4 kernel 所示，占用率有三个阶段： 45.3% (304 个线程)、39.3% (264 个线程)和 11.9% (80 个线程)。在第一阶段中，所有八个 work-group 都在 3 个 Xe-core 上一起调度，其中两个 Xe-core 分别获得 112 个线程 (80 个来自两个 work-group 和 32 个来自一个工作组), 而另一个 Xe-core 获得 80 个线程(来自两个 work-group)。在第二阶段中，一个 work-group 完成了执行，这给我们带来了 (304-40=264) 的占用率。在最后一个阶段中，剩余的两个 work-group 的八个线程被调度执行。

当使用与 Xe-core 中的线程数成倍数的 work-group 大小并且有更多的 work-group 运行相同的 kernel 时，硬件得到了很好的利用，实现了接近 100% 的占用率，如下所示。

  Initialize(sum);
  VectorAdd3<24, 224, 8>(q, a, b, sum, 10000);
  Initialize(sum);
  VectorAdd4<24, 224, 8>(q, a, b, sum, 10000);

由于我们有更多的线程，而且 work-group 大小是 Xe-core 中的线程数的倍数，因此此 kernel 执行具有不同的线程占用率。这在 VTune 时间轴上的线程占用率指标中显示如下。

VTune 中的线程占用指标

请注意，上面的时间表是根据不同的占用率数字猜测出来的，因为我们还没有一种方法来检查每个片段的占用率数字。

你可以使用上面的 kernel 运行不同的实验，以更好地了解 GPU 硬件如何在执行单元上调度软件线程。除了大量的 work-group 之外，还要注意 work-group 和 sub-group 大小，以确保有效利用 GPU 硬件。

Intel® GPU 占用计算器

总结一下, SYCL work-group 通常分派到一个 Xe-core。 work-group 中的所有 work-item 都共享一个 Xe-core 的 SLM, 用于 work-group 内部线程 barrier 和内存 fence 同步。如果有足够的 VE ALU、SLM 和线程上下文来容纳它们，则可以将多个 work-group 分派到同一个 Xe-core 上。

你可以通过充分利用所有可用的 Xe-core 来实现更高的性能。影响 kernel GPU 占用率的参数是 work-group 大小和 SIMD sub-group 大小，这也确定了 work-group 中的线程数。

可以使用 Intel® GPU Occupancy Calculator 计算给定 kernel 及其 work-group 参数在 Intel® GPU 上的占用率。

TechVideoGPU

关注

0
点赞
踩
1

收藏

觉得还不错? 一键收藏
0
评论
oneAPI GPU 优化指南 - 线程映射和 GPU 占用

SYCL 执行模型公开了 GPU 执行的抽象视图。SYCL 线程层次结构由 1、2 或 3 维 work-item 网格组成。这些工作项被分组为大小相等的线程组，称为 work-group。work-group 中的线程进一步分为大小相等的向量组称为 sub-group （参见下图）。Work-itemwork-item 代表 kernel 函数的并行执行的集合之一。Sub-group。
复制链接

扫一扫