本章节翻译by chenshusmail@163.com 原文:Sub-Groups and SIMD Vectorization (intel.com)
目录
Sub-Group 大小 vs. 最大 Sub-Group 大小
ND-Range kernel 的索引空间被划分为 work-group,sub-group 和 work-item 。 work-item 是最基本的单位。一组 work-item 组成一个 sub-group, 一组 sub-group 组成一个 work-group。 work-item 和 work-group 到硬件向量引擎 (VE) 的映射是依赖于实现的。 所有的 work-group 都同时运行,但根据资源的可用性,它们可能被安排在不同的时间运行。 work-group 的执行可能会或者不会被抢占,这取决于底层硬件的能力。 同一 work-group 中的 work-item 保证同时运行。 同一 sub-group 中的 work-item 可能有额外的调度保证,并且可以访问额外的功能。
sub-group 是全局索引空间中连续的 work-item 的集合,它们在同一个 VE 线程中执行。 当设备编译器编译 kernel 时,多个 work-item 通过向量化被打包到一个 sub-group 中, 以便生成的 SIMD 指令流可以同时执行多个 work-item 的任务。 正确地将 work-item 划分为 sub-group 可以产生很大的性能差异。
让我们从一个简单的例子开始,来说明 sub-group:
q.submit([&](auto &h) {
sycl::stream out(65536, 256, h);
h.parallel_for(sycl::nd_range(sycl::range{32}, sycl::range{32}),
[=](sycl::nd_item<1> it) {
int groupId = it.get_group(0);
int globalId = it.get_global_linear_id();
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
int sgGroupId = sg.get_group_id()[0];
int sgId = sg.get_local_id()[0];
out << "globalId = " << sycl::setw(2) << globalId
<< " groupId = " << groupId
<< " sgGroupId = " << sgGroupId << " sgId = " << sgId
<< " sgSize = " << sycl::setw(2) << sgSize
<< sycl::endl;
});
});
这个例子的输出可能是这样的:
Device: Intel(R) Gen12HP
globalId = 0 groupId = 0 sgGroupId = 0 sgId = 0 sgSize = 16
globalId = 1 groupId = 0 sgGroupId = 0 sgId = 1 sgSize = 16
globalId = 2 groupId = 0 sgGroupId = 0 sgId = 2 sgSize = 16
globalId = 3 groupId = 0 sgGroupId = 0 sgId = 3 sgSize = 16
globalId = 4 groupId = 0 sgGroupId = 0 sgId = 4 sgSize = 16
globalId = 5 groupId = 0 sgGroupId = 0 sgId = 5 sgSize = 16
globalId = 6 groupId = 0 sgGroupId = 0 sgId = 6 sgSize = 16
globalId = 7 groupId = 0 sgGroupId = 0 sgId = 7 sgSize = 16
globalId = 16 groupId = 0 sgGroupId = 1 sgId = 0 sgSize = 16
globalId = 17 groupId = 0 sgGroupId = 1 sgId = 1 sgSize = 16
globalId = 18 groupId = 0 sgGroupId = 1 sgId = 2 sgSize = 16
globalId = 19 groupId = 0 sgGroupId = 1 sgId = 3 sgSize = 16
globalId = 20 groupId = 0 sgGroupId = 1 sgId = 4 sgSize = 16
globalId = 21 groupId = 0 sgGroupId = 1 sgId = 5 sgSize = 16
globalId = 22 groupId = 0 sgGroupId = 1 sgId = 6 sgSize = 16
globalId = 23 groupId = 0 sgGroupId = 1 sgId = 7 sgSize = 16
globalId = 8 groupId = 0 sgGroupId = 0 sgId = 8 sgSize = 16
globalId = 9 groupId = 0 sgGroupId = 0 sgId = 9 sgSize = 16
globalId = 10 groupId = 0 sgGroupId = 0 sgId = 10 sgSize = 16
globalId = 11 groupId = 0 sgGroupId = 0 sgId = 11 sgSize = 16
globalId = 12 groupId = 0 sgGroupId = 0 sgId = 12 sgSize = 16
globalId = 13 groupId = 0 sgGroupId = 0 sgId = 13 sgSize = 16
globalId = 14 groupId = 0 sgGroupId = 0 sgId = 14 sgSize = 16
globalId = 15 groupId = 0 sgGroupId = 0 sgId = 15 sgSize = 16
globalId = 24 groupId = 0 sgGroupId = 1 sgId = 8 sgSize = 16
globalId = 25 groupId = 0 sgGroupId = 1 sgId = 9 sgSize = 16
globalId = 26 groupId = 0 sgGroupId = 1 sgId = 10 sgSize = 16
globalId = 27 groupId = 0 sgGroupId = 1 sgId = 11 sgSize = 16
globalId = 28 groupId = 0 sgGroupId = 1 sgId = 12 sgSize = 16
globalId = 29 groupId = 0 sgGroupId = 1 sgId = 13 sgSize = 16
globalId = 30 groupId = 0 sgGroupId = 1 sgId = 14 sgSize = 16
globalId = 31 groupId = 0 sgGroupId = 1 sgId = 15 sgSize = 16
在这个例子中,每个 sub-group 有16个 work-item,也就是 sub-group 的大小为 16。 这意味着每个线程同时执行 16 个 work-item,32 个 work-item 由两个 VE 线程执行。
默认情况下,编译器使用设备特定信息和一些启发式方法来选择 sub-group 大小。 用户可以使用 kernel 属性 intel::reqd_sub_group_size
来指定最大的 sub-group 大小,从而覆盖编译器的选择。 有时候,显式地指定 sub-group 大小可能有助于提高性能,但并不总是如此。
q.submit([&](auto &h) {
sycl::stream out(65536, 256, h);
h.parallel_for(sycl::nd_range(sycl::range{32}, sycl::range{32}),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(32)]] {
int groupId = it.get_group(0);
int globalId = it.get_global_linear_id();
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
int sgGroupId = sg.get_group_id()[0];
int sgId = sg.get_local_id()[0];
out << "globalId = " << sycl::setw(2) << globalId
<< " groupId = " << groupId
<< " sgGroupId = " << sgGroupId << " sgId = " << sgId
<< " sgSize = " << sycl::setw(2) << sgSize
<< sycl::endl;
});
});
输出将是:
Device: Intel(R) Gen12HP
globalId = 0 groupId = 0 sgGroupId = 0 sgId = 0 sgSize = 32
globalId = 1 groupId = 0 sgGroupId = 0 sgId = 1 sgSize = 32
globalId = 2 groupId = 0 sgGroupId = 0 sgId = 2 sgSize = 32
globalId = 3 groupId = 0 sgGroupId = 0 sgId = 3 sgSize = 32
globalId = 4 groupId = 0 sgGroupId = 0 sgId = 4 sgSize = 32
globalId = 5 groupId = 0 sgGroupId = 0 sgId = 5 sgSize = 32
globalId = 6 groupId = 0 sgGroupId = 0 sgId = 6 sgSize = 32
globalId = 7 groupId = 0 sgGroupId = 0 sgId = 7 sgSize = 32
globalId = 8 groupId = 0 sgGroupId = 0 sgId = 8 sgSize = 32
globalId = 9 groupId = 0 sgGroupId = 0 sgId = 9 sgSize = 32
globalId = 10 groupId = 0 sgGroupId = 0 sgId = 10 sgSize = 32
globalId = 11 groupId = 0 sgGroupId = 0 sgId = 11 sgSize = 32
globalId = 12 groupId = 0 sgGroupId = 0 sgId = 12 sgSize = 32
globalId = 13 groupId = 0 sgGroupId = 0 sgId = 13 sgSize = 32
globalId = 14 groupId = 0 sgGroupId = 0 sgId = 14 sgSize = 32
globalId = 15 groupId = 0 sgGroupId = 0 sgId = 15 sgSize = 32
globalId = 16 groupId = 0 sgGroupId = 0 sgId = 16 sgSize = 32
globalId = 17 groupId = 0 sgGroupId = 0 sgId = 17 sgSize = 32
globalId = 18 groupId = 0 sgGroupId = 0 sgId = 18 sgSize = 32
globalId = 19 groupId = 0 sgGroupId = 0 sgId = 19 sgSize = 32
globalId = 20 groupId = 0 sgGroupId = 0 sgId = 20 sgSize = 32
globalId = 21 groupId = 0 sgGroupId = 0 sgId = 21 sgSize = 32
globalId = 22 groupId = 0 sgGroupId = 0 sgId = 22 sgSize = 32
globalId = 23 groupId = 0 sgGroupId = 0 sgId = 23 sgSize = 32
globalId = 24 groupId = 0 sgGroupId = 0 sgId = 24 sgSize = 32
globalId = 25 groupId = 0 sgGroupId = 0 sgId = 25 sgSize = 32
globalId = 26 groupId = 0 sgGroupId = 0 sgId = 26 sgSize = 32
globalId = 27 groupId = 0 sgGroupId = 0 sgId = 27 sgSize = 32
globalId = 28 groupId = 0 sgGroupId = 0 sgId = 28 sgSize = 32
globalId = 29 groupId = 0 sgGroupId = 0 sgId = 29 sgSize = 32
globalId = 30 groupId = 0 sgGroupId = 0 sgId = 30 sgSize = 32
globalId = 31 groupId = 0 sgGroupId = 0 sgId = 31 sgSize = 32
有效的 sub-group 大小与设备相关。您可以查询设备获取相关信息:
std::cout << "Sub-group Sizes: ";
for (const auto &s :
q.get_device().get_info<sycl::info::device::sub_group_sizes>()) {
std::cout << s << " ";
}
std::cout << std::endl;
支持的有效 sub-group 大小可能包括:
Device: Intel(R) Gen12HP
Subgroup Sizes: 8 16 32
接下来,我们将展示如何使用 sub-group 来提高性能。
向量化和内存访问
Intel® 的图形设备具有多个 VE。每个 VE 都是一个多线程的 SIMD 处理器。 编译器生成 SIMD 指令,将一个 sub-group 中的多个 work-item 打包,以便在 VE 线程中同时执行。 编译器选择的 SIMD 宽度(即 sub-group 大小)基于设备特征和启发式方法, 或由 kernel 显式指定,可以是 8、16 或 32。
在给定 SIMD 宽度的情况下,最大化 SIMD 通道利用率可以获得最佳的指令性能。 如果一个或多个通道(或 kernel 实例或 work-item)发生分歧,线程在路径合并之前执行两个分支路径, 增加动态指令计数。 SIMD 分歧对性能产生负面影响。 编译器努力最小化分歧,但如果可能的话,最好在源代码中避免分歧。
work-item 中内存的访问方式会影响 sub-group 中内存的访问方式或 SIMD 通道的利用方式。 在 work-item 中访问连续内存通常不是最佳选择。例如:
constexpr int N = 1024 * 1024;
int *data = sycl::malloc_shared<int>(N, q);
auto e = q.submit([&](auto &h) {
h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
[=](sycl::nd_item<1> it) {
int i = it.get_global_linear_id();
i = i * 16;
for (int j = i; j < (i + 16); j++) {
data[j] = -1;
}
});
});
q.wait();
这个简单的 kernel 初始化了一个 1024 x 1024 的整数数组。每个 work-item 初始化 16 个连续的整数。 假设编译器选择的 sub-group 大小为 16,则每个 sub-group 或线程中初始化 256 个整数。 然而,16 个 SIMD 通道中的存储是分散的。
与其在一个 work-item 中初始化 16 个连续的整数,不如在一个 SIMD 指令中初始化 16 个连续的整数更有效率。
constexpr int N = 1024 * 1024;
int *data = sycl::malloc_shared<int>(N, q);
auto e = q.submit([&](auto &h) {
h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
[=](sycl::nd_item<1> it) {
int i = it.get_global_linear_id();
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
i = (i / sgSize) * sgSize * 16 + (i % sgSize);
for (int j = 0; j < sgSize * 16; j += sgSize) {
data[i + j] = -1;
}
});
});
我们在示例中使用了内存写入,但同样的技术也适用于内存读取。
constexpr int N = 1024 * 1024;
int *data = sycl::malloc_shared<int>(N, q);
int *data2 = sycl::malloc_shared<int>(N, q);
memset(data2, 0xFF, sizeof(int) * N);
auto e = q.submit([&](auto &h) {
h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
[=](sycl::nd_item<1> it) {
int i = it.get_global_linear_id();
i = i * 16;
for (int j = i; j < (i + 16); j++) {
data[j] = data2[j];
}
});
});
这个 kernel 将一个 1024 x 1024 的整数数组复制到另一个相同大小的整数数组中。 每个 work-item 复制 16 个连续的整数。然而,从 data2 中读取的数据是聚集的, 而存储到 data 中的数据是分散的。更改代码, 使每个 sub-group 中而不是每个 work-item 读取和存储连续的整数会更有效率。
constexpr int N = 1024 * 1024;
int *data = sycl::malloc_shared<int>(N, q);
int *data2 = sycl::malloc_shared<int>(N, q);
memset(data2, 0xFF, sizeof(int) * N);
auto e = q.submit([&](auto &h) {
h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
[=](sycl::nd_item<1> it) {
int i = it.get_global_linear_id();
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
i = (i / sgSize) * sgSize * 16 + (i % sgSize);
for (int j = 0; j < sgSize * 16; j += sgSize) {
data[i + j] = data2[i + j];
}
});
});
最大化内存带宽利用率
在上面的示例中,每个 work-item 在每次循环迭代中加载和存储 1 个整数或 4 个字节 (data[i + j] = data2[i + j];
), 或者每个向量化内存操作加载/存储 64 个字节(假设 sub-group 大小为 16),使内存带宽不饱和。
增加每个 work-item 在一次内存操作中加载或存储的有效负载或数据大小将导致更好的带宽利用率。
constexpr int N = 1024 * 1024;
int *data = sycl::malloc_shared<int>(N, q);
int *data2 = sycl::malloc_shared<int>(N, q);
memset(data2, 0xFF, sizeof(int) * N);
auto e = q.submit([&](auto &h) {
h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
[=](sycl::nd_item<1> it) {
int i = it.get_global_linear_id();
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
i = (i / sgSize) * sgSize * 16 + (i % sgSize) * 4;
for (int j = 0; j < 4; j++) {
sycl::vec<int, 4> x;
sycl::vec<int, 4> *q =
(sycl::vec<int, 4> *)(&(data2[i + j * sgSize * 4]));
x = *q;
sycl::vec<int, 4> *r =
(sycl::vec<int, 4> *)(&(data[i + j * sgSize * 4]));
*r = x;
}
});
});
每个 work-item 在每次循环迭代中加载/存储一个 sycl::vec<int,4>
而不是一个整数。 读取/写入 256 个连续字节的内存(假设 sub-group 大小为 16), 每个向量化内存操作的有效负载增加了四倍。
最大带宽因硬件而异。可以使用 Intel® VTune Profiler 来测量带宽并找到最佳大小。
然而,使用向量类型可能会增加寄存器压力。建议仅在寄存器不溢出时使用长向量。 请参阅“寄存器化和避免寄存器溢出”一章,了解避免寄存器溢出的技巧。
内存块加载和存储
Intel® 图形具有针对内存块加载/存储优化的指令。 因此,如果 sub-group 中的 work-item 访问连续的内存块, 您可以使用 sub-group 块访问函数来利用块加载/存储指令。
constexpr int N = 1024 * 1024;
int *data = sycl::malloc_shared<int>(N, q);
int *data2 = sycl::malloc_shared<int>(N, q);
memset(data2, 0xFF, sizeof(int) * N);
auto e = q.submit([&](auto &h) {
h.parallel_for(
sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
auto sg = it.get_sub_group();
sycl::vec<int, 8> x;
using global_ptr =
sycl::multi_ptr<int, sycl::access::address_space::global_space>;
int base = (it.get_group(0) * 32 +
sg.get_group_id()[0] * sg.get_local_range()[0]) *
16;
x = sg.load<8>(global_ptr(&(data2[base + 0])));
sg.store<8>(global_ptr(&(data[base + 0])), x);
x = sg.load<8>(global_ptr(&(data2[base + 128])));
sg.store<8>(global_ptr(&(data[base + 128])), x);
});
});
这个示例也使用 sycl::vec
来提高性能,但是 x
中的整数在内存中并不连续, 而是以 sub-group 大小为步长!
您可能已经注意到,显式指定了 sub-group 大小 16。 当您使用 sub-group 函数时,最好始终覆盖编译器的选择, 以确保 sub-group 大小始终符合您的预期。
数据共享
由于 sub-group 中的 work-item 在同一线程中执行,因此在 work-item 之间共享数据更有效率, 即使数据对每个 work-item 是私有的。 在 sub-group 中共享数据比使用共享本地内存或 SLM 在 work-group 中共享数据更有效率。 在 sub-group 中的 work-item 之间共享数据的一种方法是使用 shuffle 函数。
constexpr size_t BLOCK_SIZE = 16;
sycl::buffer<uint, 2> m(matrix.data(), sycl::range<2>(N, N));
auto e = q.submit([&](auto &h) {
sycl::accessor marr(m, h);
sycl::local_accessor<uint, 2> barr1(
sycl::range<2>(BLOCK_SIZE, BLOCK_SIZE), h);
sycl::local_accessor<uint, 2> barr2(
sycl::range<2>(BLOCK_SIZE, BLOCK_SIZE), h);
h.parallel_for(
sycl::nd_range<2>(sycl::range<2>(N / BLOCK_SIZE, N),
sycl::range<2>(1, BLOCK_SIZE)),
[=](sycl::nd_item<2> it) [[intel::reqd_sub_group_size(16)]] {
int gi = it.get_group(0);
int gj = it.get_group(1);
auto sg = it.get_sub_group();
uint sgId = sg.get_local_id()[0];
uint bcol[BLOCK_SIZE];
int ai = BLOCK_SIZE * gi;
int aj = BLOCK_SIZE * gj;
for (uint k = 0; k < BLOCK_SIZE; k++) {
bcol[k] = sg.load(marr.get_pointer() + (ai + k) * N + aj);
}
uint tcol[BLOCK_SIZE];
for (uint n = 0; n < BLOCK_SIZE; n++) {
if (sgId == n) {
for (uint k = 0; k < BLOCK_SIZE; k++) {
tcol[k] = sg.shuffle(bcol[n], k);
}
}
}
for (uint k = 0; k < BLOCK_SIZE; k++) {
sg.store(marr.get_pointer() + (ai + k) * N + aj, tcol[k]);
}
});
});
这个 kernel 转置一个 16 x 16 的矩阵。它看起来比前面的例子更复杂,但思想很简单: 一个 sub-group 加载一个 16 x 16 的子矩阵,然后使用 sub-group shuffle 函数对子矩阵进行转置。 只有一个子矩阵,子矩阵就是矩阵,所以只需要一个 sub-group。更大的矩阵,比如 4096 x 4096, 可以使用相同的技术进行转置:每个 sub-group 加载一个子矩阵,然后使用 sub-group shuffle 函数对子矩阵进行转置。 这留给读者作为练习。
SYCL 提供了多种 sub-group shuffle 函数的变体。每种变体都针对特定设备上的特定用途进行了优化。 使用这些优化过的函数(如果它们符合您的需求)会是更好的选择,而不是创建您自己的函数。
Sub-Group 大小 vs. 最大 Sub-Group 大小
到目前为止,在我们的示例中,work-group 大小可以被 sub-group 大小整除, work-group 大小和 sub-group 大小(由用户指定或由编译器自动选择)都是 2 的幂。 如果 work-group 大小可以被最大 sub-group 大小整除且两个大小都是 2 的幂, 则 sub-group 大小和最大 sub-group 大小相同。但是,如果 work-group 大小不能被 sub-group 大小整除会发生什么呢? 请看以下示例:
auto e = q.submit([&](auto &h) {
sycl::stream out(65536, 128, h);
h.parallel_for(sycl::nd_range<1>(7, 7),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
int i = it.get_global_linear_id();
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
int sgMaxSize = sg.get_max_local_range()[0];
int sId = sg.get_local_id()[0];
int j = data[i];
int k = data[i + sgSize];
out << "globalId = " << i << " sgMaxSize = " << sgMaxSize
<< " sgSize = " << sgSize << " sId = " << sId
<< " j = " << j << " k = " << k << sycl::endl;
});
});
q.wait();
这个示例的输出是这样的:
globalId = 0 sgMaxSize = 16 sgSize = 7 sId = 0 j = 0 k = 7
globalId = 1 sgMaxSize = 16 sgSize = 7 sId = 1 j = 1 k = 8
globalId = 2 sgMaxSize = 16 sgSize = 7 sId = 2 j = 2 k = 9
globalId = 3 sgMaxSize = 16 sgSize = 7 sId = 3 j = 3 k = 10
globalId = 4 sgMaxSize = 16 sgSize = 7 sId = 4 j = 4 k = 11
globalId = 5 sgMaxSize = 16 sgSize = 7 sId = 5 j = 5 k = 12
globalId = 6 sgMaxSize = 16 sgSize = 7 sId = 6 j = 6 k = 13
sub-group 大小为七,尽管最大 sub-group 大小仍然为 16!最大 sub-group 大小实际上是 SIMD 宽度, 所以它不会改变,但是 sub-group 中的 work-item 少于八个,所以 sub-group 大小为七。 因此,当您的 work-group 大小不能被最大 sub-group 大小整除时要小心。 最后一个 work-item 较少的 sub-group 可能需要特殊处理。