oneAPI GPU 优化指南 - 小寄存器模式与大寄存器模式

TechVideoGPU

已于 2023-12-08 11:41:26 修改

阅读量80

点赞数

分类专栏： oneAPI GPU 优化指南文章标签： oneapi

于 2023-11-24 11:09:08 首次发布

原文链接：https://www.intel.com/content/www/us/en/docs/oneapi/optimization-guide-gpu/2023-1/small-register-mode-vs-large-register-mode.html

版权

oneAPI GPU 优化指南专栏收录该内容

46 篇文章 1 订阅

订阅专栏

本章节翻译by chenshusmail@163.com 原文：Small Register Mode vs. Large Register Mode (intel.com)

GRF - 概述

命令行下的 GRF 模式规范

OpenMP - GRF 模式选择（AOT）

OpenMP – GRF 模式选择（JIT）

SYCL – GRF 模式选择（AOT）

SYCL – GRF 模式选择（JIT）

SYCL 寄存器分配模式 - 配置单个 kernel 的属性

针对 GRF 模式选择进行性能调优

在 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 数量更容易为较大（寄存器压力较高）的工作负载带来更好的性能表现。

TechVideoGPU

关注

0
点赞
踩
0

收藏

觉得还不错? 一键收藏
1
评论
oneAPI GPU 优化指南 - 小寄存器模式与大寄存器模式

Intel® 数据中心 GPU Max 系列产品支持两种 GRF 模式：小 GRF 模式和大 GRF 模式。每个执行单元（EU）总共有 64 KB 的寄存器存储空间。在小 GRF 模式下，一个 EU 中的单个硬件线程可以访问 128 个 GRF 寄存器，每个寄存器宽度为 64B。在这种模式下，每个 EU 可用的硬件线程数为 8。在大 GRF 模式下，一个 EU 中的单个硬件线程可以访问 256 个 GRF 寄存器，每个寄存器宽度为 64B。在这种模式下，每个 EU 可用的硬件线程数为 4。
复制链接

扫一扫