本章节翻译by chenshusmail@163.com 原文:Small Register Mode vs. Large Register Mode (intel.com)
目录
SYCL 寄存器分配模式 - 配置单个 kernel 的属性
在 Intel® 数据中心 GPU Max 系列中,有两种通用寄存器(GRF)模式可用 - 小 GRF 模式和大 GRF 模式。 有两种方法可以控制 Intel® 图形编译器(IGC)如何在这两种模式之间进行选择: (1)命令行控制,(2)配置单个 kernel 的属性。在本章中,我们提供了一个分步指南, 说明用户如何在 SYCL 和 OpenMP 后端中进行这种控制。
GRF - 概述
Intel® 数据中心 GPU Max 系列产品支持两种 GRF 模式:小 GRF 模式和大 GRF 模式。 每个执行单元(EU)总共有 64 KB 的寄存器存储空间。在小 GRF 模式下, 一个 EU 中的单个硬件线程可以访问 128 个 GRF 寄存器,每个寄存器宽度为 64B。 在这种模式下,每个 EU 可用的硬件线程数为 8。在大 GRF 模式下, 一个 EU 中的单个硬件线程可以访问 256 个 GRF 寄存器,每个寄存器宽度为 64B。 在这种模式下,每个 EU 可用的硬件线程数为 4。
命令行下的 GRF 模式规范
以下是用户可以提供的后端编译选项列表,以指导 IGC 图形编译器对 GRF 模式进行选择。
-
-ze-opt-large-register-file:
强制 IGC 为所有 kernel 选择大寄存器文件模式
-
-ze-opt-intel-enable-auto-large-GRF-mode:
允许 IGC 根据启发式规则在每个 kernel 的基础上选择小/大 GRF 模式
-
Default:
IGC 为所有 kernel 选择小 GRF 模式
OpenMP - GRF 模式选择(AOT)
以下是可以用来在 OpenMP 后端的 AOT 编译期间指定所需后端选项的各种命令。这里, test.cpp
可以是任何有效的程序:
icpx -fiopenmp -fopenmp-targets=spir64_gen -Xopenmp-target-backend
"-device pvc -options -ze-opt-large-register-file" test.cpp
// IGC 将强制所有 kernel 使用大 GRF 模式
icpx -fiopenmp -fopenmp-targets=spir64_gen -Xopenmp-target-backend
"-device pvc -options -ze-intel-enable-auto-large-GRF-mode" test.cpp
// IGC 将使用编译器的启发式规则为每个 kernel 选择小/大 GRF 模式
icpx -fiopenmp -fopenmp-targets=spir64_gen -Xopenmp-target-backend "-device pvc" test.cpp
// IGC 将自动为所有 kernel 使用小 GRF 模式
OpenMP – GRF 模式选择(JIT)
对于 OpenMP 后端,我们使用以下环境变量来指定后端选项 - LIBOMPTARGET_LEVEL0_COMPILATION_OPTIONS。 以下是可用的选项:
LIBOMPTARGET_LEVEL0_COMPILATION_OPTIONS="-ze-opt-large-register-file"
// IGC 将强制所有 kernel 使用大 GRF 模式
LIBOMPTARGET_LEVEL0_COMPILATION_OPTIONS
="-ze-intel-enable-auto-large-GRF-mode"
// IGC 将使用编译器的启发式规则为每个 kernel 选择小/大 GRF 模式
env-var 未设置
// IGC 将自动为所有 kernel 使用小 GRF 模式
SYCL – GRF 模式选择(AOT)
以下是可以用来在 SYCL 后端的 AOT 编译期间指定所需后端选项的命令。这里, test.cpp
可以是任何有效的 SYCL 程序:
icpx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend
"-device pvc -options -ze-opt-large-register-file" test.cpp
// IGC 将强制所有 kernel 使用大 GRF 模式
icpx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend
"-device pvc -options -ze-intel-enable-auto-large-GRF-mode" test.cpp
// IGC 将使用编译器的启发式规则为每个 kernel 选择小/大 GRF 模式
icpx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device pvc" test.cpp
// IGC 将自动为所有 kernel 使用小 GRF 模式
SYCL – GRF 模式选择(JIT)
以下是可以用来在 SYCL 后端的 JIT 编译期间指定所需后端选项的命令。这里, test.cpp
可以是任何有效的 SYCL 程序:
icpx -fsycl -fsycl-targets=spir64 -Xsycl-target-backend
"-options -ze-opt-large-register-file" test.cpp
// IGC 将强制所有 kernel 使用大 GRF 模式
icpx -fsycl -fsycl-targets=spir64 -Xsycl-target-backend "-options -ze-intel-enable-auto-large-GRF-mode" test.cpp
// IGC 将使用编译器的启发式规则为每个 kernel 选择小/大 GRF 模式
icpx -fsycl -fsycl-targets=spir64 -Xsycl-target-backend "" test.cpp
// IGC 将自动为所有 kernel 使用小 GRF 模式
SYCL 寄存器分配模式 - 配置单个 kernel 的属性
让我们从 SYCL 入门指南中的 kernel 示例代码开始。
cgh.parallel_for<class FillBuffer>(
NumOfWorkItems, [=](sycl::id<1> WIid) {
// 用索引填充 buffer
Accessor[WIid] = (sycl::cl_int)WIid.get(0);
});
目前,没有指定 GRF 分配模式。要指定大 GRF 模式,需要包含以下头文件:
#include <sycl/ext/intel/experimental/kernel_properties.hpp>
其次,在 kernel 的函数调用树中添加以下函数调用:
set_kernel_properties(kernel_properties::use_large_grf);
这是修改后的 kernel:
cgh.parallel_for<class FillBuffer>(
NumOfWorkItems, [=](sycl::id<1> WIid) {
set_kernel_properties(kernel_properties::use_large_grf);
// 用索引填充 buffer
Accessor[WIid] = (sycl::cl_int)WIid.get(0);
});
针对 GRF 模式选择进行性能调优
本节讨论了 GRF 模式选择对 device 代码性能的影响。本节中展示的示例使用 OpenMP 卸载模型和 JIT 编译流程。 控制 GRF 模式选择的两个主要特征是:(1)内核代码的寄存器压力(2)并行执行线程的数量。 下面是一个包含 OpenMP 卸载区域的代码片段,接下来将用它来分析。
#pragma omp target teams distribute thread_limit(ZDIM *NX1 *NX1)
for (int e = 0; e < nelt; e++) {
double s_u[NX1 * NX1 * NX1];
double s_D[NX1 * NX1];
// 这里的三个数组使用 SLM
double s_ur[NX1 * NX1 * NX1];
double s_us[NX1 * NX1 * NX1];
double s_ut[NX1 * NX1 * NX1];
#pragma omp parallel for
for (int inner = 0; inner < innerub; inner++) {
int k = inner / (NX1 * NX1);
int j = (inner - k * NX1 * NX1) / NX1;
int i = inner - k * NX1 * NX1 - j * NX1;
if (k == 0)
s_D[I2(i, j)] = D[I2(i, j)];
for (; k < NX1; k += ZDIM) {
s_u[I3(i, j, k)] = u[I4(i, j, k, e)];
}
}
#pragma omp parallel for
for (int inner = 0; inner < innerub; inner++) {
int k = inner / (NX1 * NX1);
int j = (inner - k * NX1 * NX1) / NX1;
int i = inner - k * NX1 * NX1 - j * NX1;
double r_G00, r_G01, r_G02, r_G11, r_G12, r_G22;
for (; k < NX1; k += ZDIM) {
double r_ur, r_us, r_ut;
r_ur = r_us = r_ut = 0;
#ifdef FORCE_UNROLL
#pragma unroll NX1
#endif
for (int m = 0; m < NX1; m++) {
r_ur += s_D[I2(i, m)] * s_u[I3(m, j, k)];
r_us += s_D[I2(j, m)] * s_u[I3(i, m, k)];
r_ut += s_D[I2(k, m)] * s_u[I3(i, j, m)];
}
const unsigned gbase = 6 * I4(i, j, k, e);
r_G00 = g[gbase + 0];
r_G01 = g[gbase + 1];
r_G02 = g[gbase + 2];
s_ur[I3(i, j, k)] = r_G00 * r_ur + r_G01 * r_us + r_G02 * r_ut;
r_G11 = g[gbase + 3];
r_G12 = g[gbase + 4];
s_us[I3(i, j, k)] = r_G01 * r_ur + r_G11 * r_us + r_G12 * r_ut;
r_G22 = g[gbase + 5];
s_ut[I3(i, j, k)] = r_G02 * r_ur + r_G12 * r_us + r_G22 * r_ut;
}
}
#pragma omp parallel for
for (int inner = 0; inner < innerub; inner++) {
int k = inner / (NX1 * NX1);
int j = (inner - k * NX1 * NX1) / NX1;
int i = inner - k * NX1 * NX1 - j * NX1;
for (; k < NX1; k += ZDIM) {
double wr = 0.0;
for (int m = 0; m < NX1; m++) {
double s_D_i = s_D[I2(m, i)];
double s_D_j = s_D[I2(m, j)];
double s_D_k = s_D[I2(m, k)];
wr += s_D_i * s_ur[I3(m, j, k)] + s_D_j * s_us[I3(i, m, k)] +
s_D_k * s_ut[I3(i, j, m)];
}
w[I4(i, j, k, e)] = wr;
}
}
}
这里有两个参数可以修改:(1)第 36 行内循环的展开因子;(2)第 1 行指定的 OpenMP team 数量。 循环展开因子可用于控制寄存器压力。展开因子越大,寄存器压力越高。OpenMP team 数量可用于控制并行线程数量。 在本讨论中,device 端的 kernel 执行时间被用作性能指标。这里没有提供测试结果的数值, 因为它们可能会因用户环境和 device 设置而有所不同。以下是一些观察到的现象:
- 当关闭循环展开时,发现使用小 GRF 模式性能更好。这意味着在寄存器压力不够高时,无法从使用大 GRF 模式中获得任何好处。
- 当打开循环展开时,发现使用大 GRF 模式性能更好。这意味着当寄存器压力很高,需要大 GRF 模式来承受这种压力。
- 在使用小 GRF 模式时,增加 team 数量更容易为较大(寄存器压力较高)的工作负载带来更好的性能表现。