本章节翻译by chenshusmail@163.com 原文：Shared Local Memory (intel.com)
目录
通常，work-item 需要共享数据并相互通信。一方面，所有 work-group 中的所有 work-item 都可以访问全局内存， 因此数据共享和通信可以通过全局内存进行。但是，由于其较低的带宽和较高的延迟， 通过全局内存进行共享和通信效率较低。另一方面， 在矢量引擎（VE）线程中同时执行的 sub-group 中的 work-item 可以非常高效地共享数据并相互通信， 但 sub-group 中的 work-item 数量通常较小，数据共享和通信的范围非常有限。 具有更高带宽和更低延迟的内存，可供更大范围的 work-item 访问，对于 work-item 之间的数据共享通信非常合适。 Intel® GPU 中的共享本地内存（SLM）就是为此而设计。
Intel GPU 的每个 Xe-core 都有自己的 SLM。对 SLM 的访问仅限于 Xe-core 中的 VE 或计划在同一 Xe-core 的 VE 上执行的同一 work-group 中的 work-item。它局限于一个 Xe-core（或 work-group）， 并由同一 Xe-core（或同一 work-group 中的 work-item）中的 VE 共享，因此称为 SLM。 由于它位于每个 Xe-core 芯片上，因此 SLM 比全局内存具有更高的带宽和更低的延迟。 由于它可供一个 work-group 中所有 work-item 访问，因此 SLM 可以容纳数百个 work-item 之间的数据共享和通信， 具体取决于 work-group 大小。
将 SLM 视为由 work-group 管理的缓存通常很有帮助。当一个 work-group 启动时， 该 work-group 中的 work-item 可以显式地将数据从全局内存加载到 SLM 中。 在整个 work-group 生命周期内，数据将保留在 SLM 中以便更快地访问。在 work-group 完成之前， SLM 中的数据可以由 work-item 显式写回全局内存。在 work-group 完成执行后，SLM 中的数据也消失并且无效。 SLM 与全局内存之间的数据一致性是程序员负责处理。正确地使用 SLM 可以产生显著的性能差异。
SLM 的大小和 work-group 的大小
由于它在芯片器件上，SLM 的大小有限。一个 work-group 可用的内存量取决于设备，并且可以通过查询设备获得，例如：
std::cout << "Local Memory Size: "
<< q.get_device().get_info<sycl::info::device::local_mem_size>()
<< std::endl;输出可能是这样：
Local Memory Size: 65536输出数值的单位是字节。因此，这个 GPU 设备每个 work-group 有65,536字节或 64KB 的SLM。
了解一个 work-group 可以拥有的最大 SLM 大小非常重要。在许多情况下， 可供一个 work-group 使用的 SLM 总大小和 work-group 中 work-item 数量是非常数的函数关系。 最大 SLM 大小可以限制 work-group 中 work-item 的总数，即 work-group 大小。例如， 如果最大 SLM 大小为 64KB，每个 work-item 需要512字节的 SLM，则最大 work-group 大小不能超过128。
存储体（bank）冲突
SLM 被划分为大小相等的内存 bank，可以同时被访问以获得高带宽。bank 总数取决于设备。在撰写本文时， 64个连续字节以4字节（32位）粒度存储在16个连续 bank 中。对不同 bank 的访问请求可以并行处理， 但对同一 bank 中不同地址的请求会导致 bank 冲突并被序列化。bank 冲突会对性能产生不利影响。 参考这个例子：
constexpr int N = 32;
int *data = sycl::malloc_shared<int>(N, q);
auto e = q.submit([&](auto &h) {
sycl::local_accessor<int, 1> slm(sycl::range(32 * 64), 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();
int j = it.get_local_linear_id();
slm[j * 16] = 0;
it.barrier(sycl::access::fence_space::local_space);
for (int m = 0; m < 1024 * 1024; m++) {
slm[j * 16] += i * m;
it.barrier(sycl::access::fence_space::local_space);
}
data[i] = slm[j * 16];
});
});如果 bank 数量为 16，则上述示例中的所有 work-item 将从同一 bank 的不同地址读取并写入。内存带宽是全带宽的1/16。
下一个示例没有 SLM bank 冲突并实现了全内存带宽，因为每个 work-item 都从不同 bank 的不同地址读取并写入。
constexpr int N = 32;
int *data = sycl::malloc_shared<int>(N, q);
auto e = q.submit([&](auto &h) {
sycl::local_accessor<int, 1> slm(sycl::range(32 * 64), 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();
int j = it.get_local_linear_id();
slm[j] = 0;
it.barrier(sycl::access::fence_space::local_space);
for (int m = 0; m < 1024 * 1024; m++) {
slm[j] += i * m;
it.barrier(sycl::access::fence_space::local_space);
}
data[i] = slm[j];
});
});数据共享和 work-group 的 barrier
让我们回顾一下 寄存器化和避免寄存器溢出 章节中256个 bin 的直方图示例。
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) is the workitem in which the bin resides
// (c >> 4) is the bin index
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]);
}
});
});这个示例已经被优化为使用 int 数据类型而不是 long，并在 sub-group 中共享寄存器， 以便 private 直方图的 bin 可以放入寄存器以获得最佳性能。 如果您需要更大的 bin 大小（例如 1024），则 private 直方图的 bin 将不可避免地溢出到全局内存。
只要每个 bin 都以原子方式更新，直方图的 bin 就可以被同一个 work-group 中的 work-item 共享。
constexpr int NUM_BINS = 1024;
constexpr int BLOCK_SIZE = 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);
sycl::accessor hacc(hbuf, h, sycl::read_write);
sycl::local_accessor<unsigned int, 1> local_histogram(sycl::range(NUM_BINS),
h);
h.parallel_for(
sycl::nd_range(sycl::range{N / BLOCK_SIZE}, sycl::range{64}),
[=](sycl::nd_item<1> it) {
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];
int factor = NUM_BINS / gSize;
int local_id = it.get_local_id()[0];
if ((factor <= 1) && (local_id < NUM_BINS)) {
sycl::atomic_ref<unsigned int, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::local_space>
local_bin(local_histogram[local_id]);
local_bin.store(0);
} else {
for (int k = 0; k < factor; k++) {
sycl::atomic_ref<unsigned int, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::local_space>
local_bin(local_histogram[gSize * k + local_id]);
local_bin.store(0);
}
}
it.barrier(sycl::access::fence_space::local_space);
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 (std::uint8_t shift : {0, 16, 32, 48}) {
constexpr unsigned long mask = 0x3FFU;
sycl::atomic_ref<unsigned int, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::local_space>
local_bin(local_histogram[(x >> shift) & mask]);
local_bin += 1;
}
}
it.barrier(sycl::access::fence_space::local_space);
if ((factor <= 1) && (local_id < NUM_BINS)) {
sycl::atomic_ref<unsigned int, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::local_space>
local_bin(local_histogram[local_id]);
sycl::atomic_ref<unsigned long, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::global_space>
global_bin(hacc[local_id]);
global_bin += local_bin.load();
} else {
for (int k = 0; k < factor; k++) {
sycl::atomic_ref<unsigned int, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::local_space>
local_bin(local_histogram[gSize * k + local_id]);
sycl::atomic_ref<unsigned long, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::global_space>
global_bin(hacc[gSize * k + local_id]);
global_bin += local_bin.load();
}
}
});
});当启动 work-group 时，work-group 中的每个 work-item 都会将 SLM 中的直方图 bin 中的一部分初始化为 0 （上述示例中第 24-38 行的代码）。您可以指定一个 work-item 来初始化所有直方图的 bin， 但通常将这项任务分配给 work-group 中的所有 work-item 会更高效。
第 39 行初始化后的 work-group barrier 保证了所有直方图 bin 在任何 work-item 更新任何 bin 之前都被初始化为 0。
由于 SLM 中的直方图 bin 是被所有 work-item 共享的，因此任何 work-item 对任何 bin 的更新都必须是原子 (atomic) 操作。
一旦完成了 work-group 中的本地直方图，就会更新全局直方图。但是，在读取本地 SLM bin 以更新全局 bin 之前， 再次在第 43 行调用 work-group barrier，以确保所有 work-item 都完成了它们的工作。
当共享 SLM 数据时，通常需要使用 work-group barrier 进行 work-item 同步。barrier 有耗时，并且耗时可能会随着 work-group 大小的增加而增加。 一种好方法是尝试不同的 work-group 大小以找到最适合您应用程序的大小。
您可以在示例文件夹中找到一个带有 256 个 bin 的直方图的 SLM 版本示例。您可以将其性能与使用寄存器版本的性能进行比较。 您可能会得到一些令人惊讶的结果，并思考如何进行进一步优化。
使用 SLM 作为缓存
有时您可能会发现让应用程序自己来管理一些热数据的缓存，比让硬件自动执行的性能更为理想。 通过应用程序直接管理数据缓存，每当需要数据时，您都确切地知道数据在哪里以及访问它的成本。SLM 可以被用于此目的。
参看以下 1-D 卷积示例：
sycl::buffer<int> ibuf(input.data(), N);
sycl::buffer<int> obuf(output.data(), N);
sycl::buffer<int> kbuf(kernel.data(), M);
auto e = q.submit([&](auto &h) {
sycl::accessor iacc(ibuf, h, sycl::read_only);
sycl::accessor oacc(obuf, h);
sycl::accessor kacc(kbuf, h, sycl::read_only);
h.parallel_for(sycl::nd_range<1>(sycl::range{N}, sycl::range{256}),
[=](sycl::nd_item<1> it) {
int i = it.get_global_linear_id();
int group = it.get_group()[0];
int gSize = it.get_local_range()[0];
int t = 0;
int _M = static_cast<int>(M);
int _N = static_cast<int>(N);
if ((group == 0) || (group == _N / gSize - 1)) {
if (i < _M / 2) {
for (int j = _M / 2 - i, k = 0; j < _M; ++j, ++k) {
t += iacc[k] * kacc[j];
}
} else {
if (i + _M / 2 >= _N) {
for (int j = 0, k = i - _M / 2;
j < _M / 2 + _N - i; ++j, ++k) {
t += iacc[k] * kacc[j];
}
} else {
for (int j = 0, k = i - _M / 2; j < _M; ++j, ++k) {
t += iacc[k] * kacc[j];
}
}
}
} else {
for (int j = 0, k = i - _M / 2; j < _M; ++j, ++k) {
t += iacc[k] * kacc[j];
}
}
oacc[i] = t;
});
});该示例使用 257 个元素的 kernel 数组对 8192 x 8192 个元素的整数数组进行卷积，并将结果写入输出数组。 每个 work-item 卷积一个元素。但是，要卷积一个元素，需要多达 256 个相邻元素的值。
注意到，每个输入元素都被多个 work-item 使用，您可以将整个 work-group 所需的所有输入元素预加载到 SLM 中。 之后，当需要某一个元素时，它可以被从 SLM 而不是全局内存中加载。
sycl::buffer<int> ibuf(input.data(), N);
sycl::buffer<int> obuf(output.data(), N);
sycl::buffer<int> kbuf(kernel.data(), M);
auto e = q.submit([&](auto &h) {
sycl::accessor iacc(ibuf, h, sycl::read_only);
sycl::accessor oacc(obuf, h);
sycl::accessor kacc(kbuf, h, sycl::read_only);
sycl::local_accessor<int, 1> ciacc(sycl::range(256 + (M / 2) * 2), h);
h.parallel_for(
sycl::nd_range(sycl::range{N}, sycl::range{256}),
[=](sycl::nd_item<1> it) {
int i = it.get_global_linear_id();
int group = it.get_group()[0];
int gSize = it.get_local_range()[0];
int local_id = it.get_local_id()[0];
int _M = static_cast<int>(M);
ciacc[local_id + M / 2] = iacc[i];
if (local_id == 0) {
if (group == 0) {
for (int j = 0; j < _M / 2; ++j) {
ciacc[j] = 0;
}
} else {
for (int j = 0, k = i - _M / 2; j < _M / 2; ++j, ++k) {
ciacc[j] = iacc[k];
}
}
}
if (local_id == gSize - 1) {
if (group == static_cast<int>(it.get_group_range()[0]) - 1) {
for (int j = gSize + _M / 2; j < gSize + _M / 2 + _M / 2; ++j) {
ciacc[j] = 0;
}
} else {
for (int j = gSize + _M / 2, k = i + 1;
j < gSize + _M / 2 + _M / 2; ++j, ++k) {
ciacc[j] = iacc[k];
}
}
}
it.barrier(sycl::access::fence_space::local_space);
int t = 0;
for (int j = 0, k = local_id; j < _M; ++j, ++k) {
t += ciacc[k] * kacc[j];
}
oacc[i] = t;
});
});当 work-group 启动时，每个 work-item 所需的所有输入元素都被加载到 SLM 中。除第一个和最后一个 work-item 外， 每个 work-item 都将一个元素加载到 SLM 中。第一个 work-item 将第一个元素左侧的邻居加载到 SLM 中， 最后一个 work-item 将最后一个元素右侧的邻居加载到 SLM 中。如果不存在邻居，则在 SLM 中将该元素填充为0。
在每个 work-item 开始卷积之前，调用 local barrier 以确保所有输入元素都已加载到 SLM 中。
每个 work-item 中的卷积都可以畅通无阻。所有相邻元素都从更快的 SLM 而不是全局内存中加载。
排查 SLM 错误
当 kernel 使用的 __private 或者 __local 内存超过模拟器默认支持的数量时，可能会发生 CL_OUT_OF_RESOURCES 错误。发生这种情况时，您会看到类似于以下内容的错误信息：
$ ./myapp
:
Problem size: c(150,600) = a(150,300) * b(300,600)
terminate called after throwing an instance of
'cl::sycl::runtime_error'
what(): Native API failed. Native API returns:
-5 (CL_OUT_OF_RESOURCES) -5 (CL_OUT_OF_RESOURCES)
Aborted (core dumped)
$.或者如果使用的是 onetrace:
$ onetrace -c ./myapp
:
>>>> [6254070891] zeKernelSuggestGroupSize: hKernel = 0x263b7a0 globalSizeX = 163850 globalSizeY = 1 globalSizeZ = 1 groupSizeX = 0x7fff94e239f0 groupSizeY = 0x7fff94e239f4 groupSizeZ = 0x7fff94e239f8
<<<< [6254082074] zeKernelSuggestGroupSize [922 ns] -> ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY(0x1879048195)
terminate called after throwing an instance of 'cl::sycl::runtime_error'
what(): Native API failed. Native API returns: -5 (CL_OUT_OF_RESOURCES) -5 (CL_OUT_OF_RESOURCES)
Aborted (core dumped)
$要查看复制到 SLM 的内存量以及实际硬件限制，请设置 debug 环境变量：
export PrintDebugMessages=1
export NEOReadDebugKeys=1输出将被更改为：
$ ./myapp
:
Size of SLM (656384) larger than available (131072)
terminate called after throwing an instance of 'cl::sycl::runtime_error'
what(): Native API failed. Native API returns: -5 (CL_OUT_OF_RESOURCES) -5 (CL_OUT_OF_RESOURCES)
Aborted (core dumped)
$或者如果使用的是 onetrace:
$ onetrace -c ./myapp
:
>>>> [317651739] zeKernelSuggestGroupSize: hKernel = 0x2175ae0 globalSizeX = 163850 globalSizeY = 1 globalSizeZ = 1 groupSizeX = 0x7ffd9caf0950 groupSizeY = 0x7ffd9caf0954 groupSizeZ = 0x7ffd9caf0958
Size of SLM (656384) larger than available (131072)
<<<< [317672417] zeKernelSuggestGroupSize [10325 ns] -> ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY(0x1879048195)
terminate called after throwing an instance of 'cl::sycl::runtime_error'
what(): Native API failed. Native API returns: -5 (CL_OUT_OF_RESOURCES) -5 (CL_OUT_OF_RESOURCES)
Aborted (core dumped)
$确定所需 SLM 大小后，尝试为 CL_CONFIG_CPU_FORCE_PRIVATE_MEM_SIZE 或者 CL_CONFIG_CPU_FORCE_LOCAL_MEM_SIZE 环境变量设置更大的值, 如 oneAPI 编程指南中的 Emulator Environment Variables 章节所述。
扫一扫
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
