本章节翻译by chenchensmail@163.com 原文：Thread Mapping and GPU Occupancy (intel.com)
目录
将 work-group 映射到 Xe-core 以实现最大占用率
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 的架构参数 总结如下：
| 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 大小设置线程数：
只要在加宽后有足够的 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 | |
一般来说，选择更大的 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 上的占用率。
3068
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
