本章节翻译by chenshusmail@163.com 原文：Registerization and Avoiding Register Spills (intel.com)
目录
寄存器和性能
寄存器是内存层次结构中最快的存储器。尽可能长时间地将数据保存在寄存器中对性能至关重要。 然而，寄存器空间有限，比内存空间小得多。例如， 当前一代 Intel® GPU 每个 XVE 线程有 128 个通用寄存器，每个默认宽度为 32 字节。 尽管编译器旨在将尽可能多的变量分配给寄存器， 但有限数量的寄存器只能在执行过程中的某一时刻分配给一小组变量。 一个特定的寄存器可以在不同时间保存不同的变量，因为不同时间需要不同的变量集合。 如果没有足够的寄存器来保存所有变量，则寄存器可能会溢出， 或者当前在寄存器中的某些变量可以移动到内存中以腾出空间给其他变量。
在 SYCL 中，编译器将寄存器分配给 work-item 中的私有变量。 一个 sub-group 中的多个 work-item 被打包到一个 XVE 线程中。 默认情况下，编译器使用寄存器压力作为选择 SIMD 宽度或 sub-group 大小的参考因素之一。 如果没有显式指定 sub-group 大小，高寄存器压力可能导致较小的 sub-group 大小（例如 8 而不是 16）， 这也可能导致寄存器溢出或导致某些变量无法提升到寄存器。
如果 sub-group 大小或 SIMD 宽度不是硬件支持的最大值，则硬件可能无法充分利用。 寄存器溢出可能导致性能显著下降，特别是当溢出发生在热点循环内部时。 当变量未提升到寄存器时，对这些变量的访问会导致内存访问显著增加。
尽管编译器使用智能算法在寄存器中分配变量并最小化寄存器溢出， 但开发人员的优化可以帮助编译器做得更好，并且通常会产生很大的性能差异。
优化技术
以下技术可以减小寄存器压力：
-
尽可能缩短私有变量的生命周期。
尽管编译器调度指令并优化变量的距离，但在某些情况下， 在源码中将加载和使用相同变量移动得更近或删除某些依赖关系可以帮助编译器做得更好。
-
避免过度的循环展开。
循环展开可以让编译器看到更多指令调度优化的机会，从而提高性能。 然而，展开引入的临时变量可能会增加寄存器分配的压力并导致寄存器溢出。 实践中的好方法是比较带有和不带有循环展开以及不同次数展开的性能，以决定是否应该展开循环或展开多少次。
-
优先选择使用 USM 指针。
Buffer accessor 方式的访问占用的空间比 USM 指针多。 如果可以在 USM 指针和 buffer accessor 方式之间选择，请选择使用 USM 指针。
-
对于计算开销小的值，按需重新计算，而不是将它一直保留为变量，否则它会被长时间保存在寄存器中。
-
避免使用大数组或大结构，或者将大结构的数组拆分为多个小结构的数组。
例如，一个
sycl::float4的数组：sycl::float4 v[8];可以拆分为 4 个
float数组:float x[8]; float y[8]; float z[8]; float w[8];所有或部分 4 个
float数组比sycl::float4数组更有可能被分配到寄存器中。 -
将一个大循环拆分为多个小循环，以减少同时活跃的变量数量。
-
如果可能，选择较小尺寸的数据类型。
-
不要将私有变量声明为 volatile。
-
在 sub-group 中共享寄存器。
-
如果可能，使用 sub-group 块来加载/存储。
-
尽可能使用共享本地内存（SLM）。
这里列举了一些技术，并不详尽。
本章的其余部分将展示如何在实际示例中应用这些技术，重点是最后五项。
选择较小的数据类型
constexpr int BLOCK_SIZE = 256;
constexpr int NUM_BINS = 32;
std::vector<unsigned long> hist(NUM_BINS, 0);
sycl::buffer<unsigned long, 1> mbuf(input.data(), N);
sycl::buffer<unsigned long, 1> hbuf(hist.data(), NUM_BINS);
auto e = q.submit([&](auto &h) {
sycl::accessor macc(mbuf, h, sycl::read_only);
auto hacc = hbuf.get_access<sycl::access::mode::atomic>(h);
h.parallel_for(
sycl::nd_range(sycl::range{N / BLOCK_SIZE}, sycl::range{64}),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
int group = it.get_group()[0];
int gSize = it.get_local_range()[0];
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
int sgGroup = sg.get_group_id()[0];
unsigned long
histogram[NUM_BINS]; // 直方图 bin 占用太多存储空间，无法提升到寄存器中
for (int k = 0; k < NUM_BINS; k++) {
histogram[k] = 0;
}
for (int k = 0; k < BLOCK_SIZE; k++) {
unsigned long x =
sg.load(macc.get_pointer() + group * gSize * BLOCK_SIZE +
sgGroup * sgSize * BLOCK_SIZE + sgSize * k);
#pragma unroll
for (int i = 0; i < 8; i++) {
unsigned int c = x & 0x1FU;
histogram[c] += 1;
x = x >> 8;
}
}
for (int k = 0; k < NUM_BINS; k++) {
hacc[k].fetch_add(histogram[k]);
}
});
});这个例子计算了单区间（bin）大小为 32 的直方图。每个 work-item 都有 32 个私有的 unsigned long 数据类型的 bin。 由于所需存储空间较大，私有 bin 无法放入寄存器中，导致整体性能较差。
当 BLOCK_SIZE 为 256 时，每个私有 bin 的最大值不会超过 unsigned integer 的最大值。 我们可以使用 unsigned integer 类型来代替 unsigned long 类型，以减小寄存器的压力， 使私有 bin 能够放入寄存器中。这个简单的改变会带来显著的性能提升。
constexpr int BLOCK_SIZE = 256;
constexpr int NUM_BINS = 32;
std::vector<unsigned long> hist(NUM_BINS, 0);
sycl::buffer<unsigned long, 1> mbuf(input.data(), N);
sycl::buffer<unsigned long, 1> hbuf(hist.data(), NUM_BINS);
auto e = q.submit([&](auto &h) {
sycl::accessor macc(mbuf, h, sycl::read_only);
auto hacc = hbuf.get_access<sycl::access::mode::atomic>(h);
h.parallel_for(
sycl::nd_range(sycl::range{N / BLOCK_SIZE}, sycl::range{64}),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
int group = it.get_group()[0];
int gSize = it.get_local_range()[0];
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
int sgGroup = sg.get_group_id()[0];
unsigned int histogram[NUM_BINS]; // 使用较小的数据类型，直方图 bin 占用的存储空间会更少
for (int k = 0; k < NUM_BINS; k++) {
histogram[k] = 0;
}
for (int k = 0; k < BLOCK_SIZE; k++) {
unsigned long x =
sg.load(macc.get_pointer() + group * gSize * BLOCK_SIZE +
sgGroup * sgSize * BLOCK_SIZE + sgSize * k);
#pragma unroll
for (int i = 0; i < 8; i++) {
unsigned int c = x & 0x1FU;
histogram[c] += 1;
x = x >> 8;
}
}
for (int k = 0; k < NUM_BINS; k++) {
hacc[k].fetch_add(histogram[k]);
}
});
});不要将私有变量声明为 volatile
现在我们对代码示例进行了一些小改动：
constexpr int BLOCK_SIZE = 256;
constexpr int NUM_BINS = 32;
std::vector<unsigned long> hist(NUM_BINS, 0);
sycl::buffer<unsigned long, 1> mbuf(input.data(), N);
sycl::buffer<unsigned long, 1> hbuf(hist.data(), NUM_BINS);
auto e = q.submit([&](auto &h) {
sycl::accessor macc(mbuf, h, sycl::read_only);
auto hacc = hbuf.get_access<sycl::access::mode::atomic>(h);
h.parallel_for(sycl::nd_range(sycl::range{N / BLOCK_SIZE}, sycl::range{64}),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
int group = it.get_group()[0];
int gSize = it.get_local_range()[0];
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
int sgGroup = sg.get_group_id()[0];
volatile unsigned int
histogram[NUM_BINS]; // 被 volatile 标记的变量不会被分配到任何寄存器中
for (int k = 0; k < NUM_BINS; k++) {
histogram[k] = 0;
}
for (int k = 0; k < BLOCK_SIZE; k++) {
unsigned long x = sg.load(
macc.get_pointer() + group * gSize * BLOCK_SIZE +
sgGroup * sgSize * BLOCK_SIZE + sgSize * k);
#pragma unroll
for (int i = 0; i < 8; i++) {
unsigned int c = x & 0x1FU;
histogram[c] += 1;
x = x >> 8;
}
}
for (int k = 0; k < NUM_BINS; k++) {
hacc[k].fetch_add(histogram[k]);
}
});
});私有直方图数组被定义为 volatile 数组。volatile 变量不会被提升到寄存器中， 因为它们的值可能在两次不同的加载操作之间发生变化。
实际上，私有直方图数组没有必要是 volatile 的，因为它只能被本地执行线程访问。 事实上，如果一个私有变量真的需要是 volatile 的，那么它就不再是私有的了。
在 sub-group 中共享寄存器
现在我们将 bin 的大小增加到 256：
constexpr int BLOCK_SIZE = 256;
constexpr int NUM_BINS = 256;
std::vector<unsigned long> hist(NUM_BINS, 0);
sycl::buffer<unsigned long, 1> mbuf(input.data(), N);
sycl::buffer<unsigned long, 1> hbuf(hist.data(), NUM_BINS);
auto e = q.submit([&](auto &h) {
sycl::accessor macc(mbuf, h, sycl::read_only);
auto hacc = hbuf.get_access<sycl::access::mode::atomic>(h);
h.parallel_for(
sycl::nd_range(sycl::range{N / BLOCK_SIZE}, sycl::range{64}),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
int group = it.get_group()[0];
int gSize = it.get_local_range()[0];
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
int sgGroup = sg.get_group_id()[0];
unsigned int
histogram[NUM_BINS]; // 直方图 bin 占用太多存储空间，无法提升到寄存器中
for (int k = 0; k < NUM_BINS; k++) {
histogram[k] = 0;
}
for (int k = 0; k < BLOCK_SIZE; k++) {
unsigned long x =
sg.load(macc.get_pointer() + group * gSize * BLOCK_SIZE +
sgGroup * sgSize * BLOCK_SIZE + sgSize * k);
#pragma unroll
for (int i = 0; i < 8; i++) {
unsigned int c = x & 0x1FU;
histogram[c] += 1;
x = x >> 8;
}
}
for (int k = 0; k < NUM_BINS; k++) {
hacc[k].fetch_add(histogram[k]);
}
});
});当 bin 的大小增加到 256 时，即使使用较小的数据类型 unsigned integer，性能也会下降。 每个 work-item 中私有 bin 的所需的存储空间太大，无法放入寄存器中。
如果 sub-group 的大小为 16，我们知道 16 个 work-item 被打包到一个 EU 线程中。 我们也知道同一 sub-group 中的 work-item 可以非常高效地相互通信和共享数据。 如果同一 sub-group 中的 work-item 共享私有 bin， 那么整个 sub-group 只需要 256 个私有 bin，或者每个 work-item 只需要 16 个私有 bin。
为了在 sub-group 中共享 bin，每个 work-item 将其输入数据广播给同一 sub-group 中的每个 work-item。 由拥有相应 bin 的 work-item 进行数据更新。
constexpr int BLOCK_SIZE = 256;
constexpr int NUM_BINS = 256;
std::vector<unsigned long> hist(NUM_BINS, 0);
sycl::buffer<unsigned long, 1> mbuf(input.data(), N);
sycl::buffer<unsigned long, 1> hbuf(hist.data(), NUM_BINS);
auto e = q.submit([&](auto &h) {
sycl::accessor macc(mbuf, h, sycl::read_only);
auto hacc = hbuf.get_access<sycl::access::mode::atomic>(h);
h.parallel_for(
sycl::nd_range(sycl::range{N / BLOCK_SIZE}, sycl::range{64}),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
int group = it.get_group()[0];
int gSize = it.get_local_range()[0];
auto sg = it.get_sub_group();
int sgSize = sg.get_local_range()[0];
int sgGroup = sg.get_group_id()[0];
unsigned int
histogram[NUM_BINS / 16]; // 直方图 bin 占用太多存储空间，无法提升到寄存器中
for (int k = 0; k < NUM_BINS / 16; k++) {
histogram[k] = 0;
}
for (int k = 0; k < BLOCK_SIZE; k++) {
unsigned long x =
sg.load(macc.get_pointer() + group * gSize * BLOCK_SIZE +
sgGroup * sgSize * BLOCK_SIZE + sgSize * k);
// sub-group 的大小是 16
#pragma unroll
for (int j = 0; j < 16; j++) {
unsigned long y = sycl::group_broadcast(sg, x, j);
#pragma unroll
for (int i = 0; i < 8; i++) {
unsigned int c = y & 0xFF;
// (c & 0xF) 是 bin 所在的 work-item
// (c >> 4) 是 bin 的索引
if (sg.get_local_id()[0] == (c & 0xF)) {
histogram[c >> 4] += 1;
}
y = y >> 8;
}
}
}
for (int k = 0; k < NUM_BINS / 16; k++) {
hacc[16 * k + sg.get_local_id()[0]].fetch_add(histogram[k]);
}
});
});使用 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}, sycl::range{32}),
[=](sycl::nd_item<1> it) {
int i = it.get_global_linear_id();
data[i] = data2[i];
});
});语句
data[i] = data2[i];中的内存加载和存储被向量化，每个向量通道都有自己的地址。假设 SIMD 宽度或 sub-group 大小为 16， 则 16 个通道的地址的总寄存器空间为 128 字节。如果每个 GRF 寄存器宽度为 32 字节，则需要 4 个 GRF 寄存器用于存储地址。
注意到地址是连续的，我们可以使用 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}, sycl::range{32}),
[=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
auto sg = it.get_sub_group();
int 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]);
x = sg.load(global_ptr(&(data2[base + 0])));
sg.store(global_ptr(&(data[base + 0])), x);
});
});语句
x = sg.load(global_ptr(&(data2[base + 0]))); sg.store(global_ptr(&(data[base + 0])), x);中每个都加载/存储一个连续的内存块，编译器将这 2 条语句编译为特殊的内存块加载/存储指令。 由于它是一个连续的内存块，我们只需要该块的起始地址。因此， 每个块加载/存储的地址只使用了 8 字节而不是 128 字节的实际寄存器空间，或着最多也就用到 1 个寄存器。
使用共享本地内存（SLM）
如果 bin 的数量很大，例如大于 1024，即使私有 bin 在同一 sub-group 中共享， 也不会有足够的寄存器空间用于存储私有 bin。为了减少内存访问，本地 bin 可以分配在 SLM 中， 并由同一 work-group 中的 work-item 共享。请参阅“共享本地内存（SLM）”章节，并查看直方图示例中是如何完成的。
906
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
