AMD CDNA™2 GPU 中的寄存器压力

于 2024-10-01 10:30:12 发布
阅读量368 收藏 8
点赞数 8
分类专栏： # ROCm 文章标签： 人工智能 学习
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/eidolon_foot/article/details/142668178
版权

Register pressure in AMD CDNA™2 GPUs — ROCm Blogs

注意： 此博客以前是  AMD实验室笔记 博客系列的一部分。

GPU kernel 中的寄存器压力对高性能计算 (HPC) 应用程序的整体性能有着巨大的影响。理解和控制寄存器的使用可以让开发者精心设计代码，以最大化硬件资源。下面的博客文章专注于一个实际演示，展示如何应用2022年8月23日的 OLCF培训讲座 中解释的建议。这里是 培训存档，你也可以在那里找到幻灯片。我们只专注于使用 ROCm 5.4 的 AMD CDNA™2 架构 (MI200系列GPU)。

寄存器和占用率

通用寄存器是传统处理器中最快三的内存类型。在大多数情况下，传统处理器和加速器（如GPU）中的算术逻辑单元（ALU）是唯一可以直接访问寄存器的组件。不幸的是，寄存器是一种稀缺且昂贵的资源，编译器会尽力“优化”局部变量分配到硬件寄存器以供ALU操作时使用。

当我们使用“优化”一词时，我们总是要澄清优化过程的目标。实际上，常规CPU和加速器（如GPU）由于其本质不同，它们执行程序和实现高性能的方法也不同。一方面，传统CPU是面向延迟设计的，旨在尽可能多地执行单个串行线程的指令。另一方面，GPU是面向吞吐量设计的，旨在尽可能利用独立线程之间的并行性。

在AMD GPU中，在同一个计算单元（CU）上运行的高数量并发波前（wavefront）使得GPU能够隐藏全局内存访问时间，比执行计算操作所需的时间更长，而此时其他波前正在执行操作。

“占用率”这个术语表示在同一个CU上能同时运行的最大波前数量。一般来说，较高的占用率有助于通过其他操作隐藏昂贵的内存访问操作，从而实现更好的性能，但这并不是总能如此。

在图1中，我们展示了CDNA™2架构中CU的示意图。矢量通用寄存器（VGPRs）用于存储在波前中各个工作项不同的数据，即每个工作项的数据是不同的。它们是CU中最通用的寄存器，可以直接由矢量ALU（VALU）操作。VALU负责执行CU中的大部分工作，包括浮点运算（FLOPs）、内存加载、整数和逻辑操作等。

标量通用寄存器（SGPRs）代表一组用于存储在编译时已知的波前中各工作项相同数据的寄存器。SGPRs由标量ALU（SALU）操作，而SALU与VALU不同，只能用于有限的操作集，如整数和逻辑操作。

本地数据共享（LDS）是一种快速的CU上管理的内存，可以在块内所有工作项之间高效共享数据。

../../_images/gcn_compute_unit.png

图1：CDNA™2架构中CU的示意图

理想情况下，我们希望始终有尽可能高的占用率。实际上，占用率受硬件设计选择和由运行在卡上的内核（如HIP、OpenCL等）所决定的资源限制。例如，每个基于AMD CDNA™2的GPU的CU有四组波前缓冲区，每个执行单元（EU，也称为图1中的SIMD单元）一个波前缓冲区，每个CU有四个EU。每个EU最多可以管理*八个*波前。这意味着CDNA™2中的物理占用率上限是每个CU 32个波前。

内核所需的寄存器数量是最常见的占用率限制因素之一。另一个常见的限制因素是LDS。以下表格总结了在AMD CDNA™2基于的GPU上，作为内核使用的VGPRs数量的函数，可以达到的最大占用率。

../../_images/occupancy_vgpr.JPG

表1：MI200中与VGPRs使用相关的占用率

寄存器溢出

寄存器分配是将GPU内核的局部变量和表达式结果分配到硬件可用寄存器的过程。它由编译器在编译时进行，并受到指令调度等其他阶段的影响。找到该问题的最优解决方案是NP难题，因此必须采用启发式技术在合理时间内找到接近最优的解决方案。

编译器试图通过减少寄存器的需求来应用启发式技术以最大化占用率，遵循表1。当请求的寄存器数量变得过高时，性能会因“寄存器压力”而受到惩罚，这导致较低的占用率和临时存储器的使用。

有时，编译器可能认为，即使请求的寄存器数量超过表1中报告的限制，达到更高的占用率也是有益的。例如，应用程序需要134个寄存器，但编译器仅分配128个，其余的放在临时存储器中。通过在临时存储器中保存一些变量，可以实现更高的占用率：这是一个线程私有的局部存储器，由全局存储器支持，速度比寄存器存储器慢得多。这种技术被称为“寄存器溢出”。

虽然观察到变量被分配到临时存储器可能是高寄存器压力的前兆，但应该从更广泛的性能背景中考虑这一点。事实上，通过保存少量寄存器来实现更高的占用率，与没有任何临时存储器使用但占用率较低相比，可以提供显著的性能提升。

在请求的寄存器数远高于硬件可用寄存器数的情况下，性能将在低占用率（最坏情况下每个计算单元1个波前）和高成本访问被“溢出”到临时存储器的寄存器变量之间受到影响。 

如何减少寄存器压力

如前所述，编译器通过采用启发式技术来最大化占用率，从而最小化某些GPU内核所需的寄存器数量。这些启发式技术有时未能接近最优解决方案，程序员需要重新构建代码以减少寄存器压力并提高性能。

在本节中，我们将介绍如何识别寄存器压力问题以及如何缓解它。

首先，可以通过两种方式检测GPU内核使用的寄存器数量：1) 使用 -Rpass-analyze=kernel-resource-usage 标志编译包含内核的文件，该标志将在编译时打印每个内核的资源使用情况；其中包括SGPRs（标量通用寄存器）、VGPRs（矢量通用寄存器）、ScratchSize（临时内存大小）、VGPR/SGPR溢出、占用率和LDS（本地数据共享）使用情况。2) 使用 --save-temps 进行编译，并检查 hip-amdgcn-amd-amdhsa-gfx90a.s 文件中的 .vgpr_spill_count。`-Rpass-analyze=kernel-resource-usage` 标志报告的所有信息也可以在这个文件中找到。

一旦评估/确认了寄存器压力情况，可以将以下几种技术应用于代码以减少寄存器压力。

  1. 为每个内核设置 __launch_bounds__ 限定符。默认情况下，编译器假设每个内核的块大小为1024个工作项。定义 __launch_bounds__ 后，编译器可以适当地分配寄存器，从而潜在地降低寄存器压力。

  2. 将变量定义/赋值移动到接近它们使用的位置。在GPU内核顶部定义一个或多个变量，并在底部使用它们，这会迫使编译器将这些变量存储在寄存器或临时内存中，直到它们被使用，从而影响使用这些寄存器处理更多性能关键变量的可能性。将定义/赋值移动到它们第一次使用的位置将有助于启发式技术对其余代码做出更有效的选择。

  3. 避免在堆栈上分配数据。在堆栈上分配的内存，例如 double array[10]，默认情况下会存在于临时内存中，编译器可能会将其优化存储到寄存器中。如果你的应用程序使用在堆栈上分配的内存，看到临时内存的使用不应感到惊讶。

  4. 避免将大对象作为内核参数传递。函数参数在堆栈上分配，可能作为优化存储到寄存器中。有时，将这些参数存储为 constant 可能会有所帮助。

  5. 避免编写包含许多函数调用（包括数学函数和断言）的大型内核。当前，编译器总是内联设备函数，包括数学函数和断言。拥有许多这些函数调用会引入额外的代码和潜在的更高寄存器压力。例如，将 pow(var,2.0) 替换为简单的 var*var 可以显著减少寄存器压力。

  6. 控制循环展开。可以通过在编译时已知迭代次数的循环上添加 #pragma unroll 命令来实现循环展开。通过这样做，所有迭代都将完全展开，从而减少检查循环退出条件的开销。然而，循环展开会增加寄存器压力，因为需要同时存储更多变量。在寄存器压力成为问题的情况下，应限制使用循环展开。请注意，Clang编译器在展开循环方面往往比其他编译器更加字面化。

  7. 手动溢出到LDS。作为最后的手段，可以使用一些LDS内存手动存储变量，可能是那些生命周期最长的变量，从而为每个线程节省几个寄存器。

示例

接下来，我们将重点讨论以下代码：

__global__ void kernel (double *phi, double *laplacian_phi,
      double *grad_phi_x, double *grad_phi_y, double *grad_phi_z,
      double *f0, double *f1, double *f2, double *f3, double *f4,
      double *f5, double *f6,
      double *g0, double *g1, double *g2, double *g3, double *g4,
      double *g5, double *g6, double* g7, double *g8, double *g9,
      double *g10, double *g11, double *g12, double *g13, double *g14,
      double *g15, double *g16, double *g17, double *g18,
      int nx, int ny, int nz, int ldx, int ldy, int current, int next,
      double k, double alpha, double phi2, double gamma,
      double itauphi, double itauphi1, double ieta,
      double itaurho, double grav,
      double eg1, double eg2, double eg0, double egc0, double egc1, double egc2)
{
  int i = (threadIdx.x + blockIdx.x * blockDim.x);
  int j = (threadIdx.y + blockIdx.y * blockDim.y);
  int z = (threadIdx.z + blockIdx.z * blockDim.z);

  int m, current_pos;

  double mu_phi, current_phi, current_phi_2;
  double rho;
  double fx, fy, fz;
  double uf, ux, uy, uz, v;
  double af, ag, cf;
  double eg1ag, eg2ag, eg1rho, eg2rho;
  double tmp1, tmp2;

  if(i <= nx && j <= ny && z <= nz)
    {
      m = i + ldx * (j + ldy * z);
      current_pos = m + current;

      current_phi = phi[m];
      current_phi_2 = pow(current_phi,2.0);

      rho = g0[m] + g1[current_pos] + g2[current_pos] + g3[current_pos] + g4[current_pos] +
  g5[current_pos] + g6[current_pos] + g7[current_pos] + g8[current_pos]  + g9[current_pos] +
  g10[current_pos] + g11[current_pos] + g12[current_pos] + g13[current_pos] + g14[current_pos] +
  g15[current_pos] + g16[current_pos] + g17[current_pos] + g18[current_pos];

      mu_phi = alpha * current_phi * ( current_phi_2 - phi2 ) - k * laplacian_phi[m];

      fx = mu_phi * grad_phi_x[m];
      fy = mu_phi * grad_phi_y[m];
      fz = mu_phi * grad_phi_z[m];

      ux = ( g1[current_pos] - g2[current_pos] + g7[current_pos] - g8[current_pos] + g9[current_pos] -
       g10[current_pos] + g11[current_pos] - g12[current_pos] + g13[current_pos] - g14[current_pos] +
       0.50 * fx ) * 1.0/rho;
      uy = ( g3[current_pos] - g4[current_pos] + g7[current_pos] - g8[current_pos] - g9[current_pos] +
       g10[current_pos] + g15[current_pos] - g16[current_pos] + g17[current_pos] - g18[current_pos] +
       0.50 * fy ) * 1.0/rho;
      uz = ( g5[current_pos] - g6[current_pos] + g11[current_pos] - g12[current_pos] - g13[current_pos] +
       g14[current_pos] + g15[current_pos] - g16[current_pos] - g17[current_pos] + g18[current_pos] +
       0.50 * fz ) * 1.0/rho;

      af = 0.50 * gamma * mu_phi * itauphi;
      cf = itauphi * ieta * current_phi;

      f0[m] = itauphi1 * f0[m] + -3.0 * gamma * mu_phi * itauphi + itauphi * current_phi;

      f1[current_pos] = itauphi1 * f1[current_pos] + af + cf * ux;
      f2[current_pos] = itauphi1 * f2[current_pos] + af - cf * ux;
      f3[current_pos] = itauphi1 * f3[current_pos] + af + cf * uy;
      f4[current_pos] = itauphi1 * f4[current_pos] + af - cf * uy;
      f5[current_pos] = itauphi1 * f5[current_pos] + af + cf * uz;
      f6[current_pos] = itauphi1 * f6[current_pos] + af - cf * uz;

      ag  = 3.0 * current_phi * mu_phi + rho;
      eg1ag = eg1 * ag;
      eg2ag = eg2 * ag;
      eg1rho = eg1 * rho;
      eg2rho = eg2 * rho;
      v  = 1.50 * ( ux*ux + uy*uy + uz*uz );
      uf = ux * fx + uy * fy + uz * fz;

      g0[m] = itaurho * g0[m] + eg0 * ( (rho - 6.0 * current_phi * mu_phi) - rho*v ) - egc0*uf;

      tmp1 = eg1ag + eg1rho*( 0.50*ux*ux - v ) + egc1*( ux*fx - uf );
      tmp2 = eg1rho*ux + egc1*fx;

      g1[m+next + 1] = itaurho * g1[current_pos] + tmp1 + tmp2;
      g2[m+next - 1] = itaurho * g2[current_pos] + tmp1 - tmp2;

      tmp1 = eg1ag + eg1rho*( 0.50 * uy * uy - v ) + egc1 * ( uy * fy - uf );
      tmp2 = eg1rho * uy + egc1 * fy;

      g3[m+next + ldx] = itaurho * g3[current_pos] + tmp1 + tmp2;
      g4[m+next - ldx] = itaurho * g4[current_pos] + tmp1 - tmp2;

      tmp1 = eg1ag + eg1rho*( 0.50 * uz * uz - v ) + egc1 * ( uz * fz - uf );
      tmp2 = eg1rho * uz + egc1 * fz;

      g5[m+next + ldx*ldy] = itaurho * g5[current_pos] + tmp1 + tmp2;
      g6[m+next - ldx*ldy] = itaurho * g6[current_pos] + tmp1 - tmp2;

      tmp1 = eg2ag + eg2rho * ( 0.50 * ( ux + uy ) * ( ux + uy ) - v ) +
  egc2 * ( ( ux + uy ) * ( fx + fy ) - uf );

      tmp2 = eg2rho * ( ux + uy ) + egc2 * ( fx + fy );

      g7[m+next + 1 + ldx] = itaurho * g7[current_pos] + tmp1 + tmp2;
      g8[m+next - 1 - ldx] = itaurho * g8[current_pos] + tmp1 - tmp2;

      tmp1 = eg2ag + eg2rho * ( 0.50 * ( ux - uy ) * ( ux - uy ) - v ) +
  egc2 * ( ( ux - uy )*( fx - fy ) - uf );
      tmp2 = eg2rho * ( ux - uy ) + egc2 * ( fx - fy );

      g9[m+next + 1 - ldx]  = itaurho * g9[current_pos]  + tmp1 + tmp2;
      g10[m+next - 1 + ldx] = itaurho * g10[current_pos] + tmp1 - tmp2;

      tmp1 = eg2ag + eg2rho * ( 0.50 * ( ux + uz ) * ( ux + uz ) - v ) +
  egc2 * ( ( ux + uz ) * ( fx + fz ) - uf );
      tmp2 = eg2rho * ( ux + uz ) + egc2 * ( fx + fz );

      g11[m+next + 1 + ldx*ldy] = itaurho * g11[current_pos] + tmp1 + tmp2;
      g12[m+next - 1 - ldx*ldy] = itaurho * g12[current_pos] + tmp1 - tmp2;

      tmp1 = eg2ag + eg2rho * ( 0.50 * ( ux - uz ) * ( ux - uz ) - v ) +
  egc2 * ( ( ux - uz ) * ( fx - fz ) - uf );
      tmp2 = eg2rho * ( ux - uz ) + egc2 * ( fx - fz );

      g13[m+next + 1 - ldx*ldy] = itaurho * g13[current_pos] + tmp1 + tmp2;
      g14[m+next - 1 + ldx*ldy] = itaurho * g14[current_pos] + tmp1 - tmp2;

      tmp1 = eg2ag + eg2rho * ( 0.50 * ( uy + uz ) * ( uy + uz ) - v ) +
  egc2 * ( ( uy + uz ) * ( fy + fz ) - uf );
      tmp2 = eg2rho * ( uy + uz ) + egc2 * ( fy + fz );

      g15[m+next + ldx + ldx*ldy] = itaurho * g15[current_pos] + tmp1 + tmp2;
      g16[m+next - ldx - ldx*ldy] = itaurho * g16[current_pos] + tmp1 - tmp2;

      tmp1 = eg2ag + eg2rho * ( 0.50 * ( uy - uz ) * ( uy - uz ) - v ) +
  egc2 * ( ( uy - uz ) * ( fy - fz ) - uf );
      tmp2 = eg2rho * ( uy - uz ) + egc2 * ( fy - fz );

      g17[m+next + ldx - ldx*ldy] = itaurho * g17[current_pos] + tmp1 + tmp2;
      g18[m+next - ldx + ldx*ldy] = itaurho * g18[current_pos] + tmp1 - tmp2;

    }
}

我们注意到，使用大量的双精度变量来存储暂时的数学运算结果以及有意义的物理数量，可能会导致该核函数的性能受寄存器压力影响。为了验证这个假设，我们可以如下编译核函数，以获取核函数的资源使用情况：

hipcc --offload-arch=gfx90a lbm.cpp -Rpass-analysis=kernel-resource-usage -c
lbm.cpp:16:1: remark: Function Name: _Z6kernelPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiiiiiddddddddddddddd [-Rpass-analysis=kernel-resource-usage]
{
^
lbm.cpp:16:1: remark:     SGPRs: 98 [-Rpass-analysis=kernel-resource-usage]
lbm.cpp:16:1: remark:     VGPRs: 102 [-Rpass-analysis=kernel-resource-usage]
lbm.cpp:16:1: remark:     AGPRs: 0 [-Rpass-analysis=kernel-resource-usage]
lbm.cpp:16:1: remark:     ScratchSize [bytes/lane]: 0 [-Rpass-analysis=kernel-resource-usage]
lbm.cpp:16:1: remark:     Occupancy [waves/SIMD]: 4 [-Rpass-analysis=kernel-resource-usage]
lbm.cpp:16:1: remark:     SGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm.cpp:16:1: remark:     VGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm.cpp:16:1: remark:     LDS Size [bytes/block]: 0 [-Rpass-analysis=kernel-resource-usage]

虽然没有寄存器溢出，但我们注意到占用率仅为每SIMD单元4个波浪（wave）；约为最佳可实现情况的一半。查看之前显示的占用率表格，我们发现需要将使用的VGPRs数量从102减少到96或以下，以达到每个SIMD单元5个波浪的占用率。

优化1：移除不必要的数学函数调用

查看以下代码，我们注意到变量 current_phi 进行平方时使用了 pow 函数。

 if(i <= nx && j <= ny && z <= nz)
    {
      m = i + ldx * (j + ldy * z);
      current_pos = m + current;

      current_phi = phi[m];
      current_phi_2 = pow(current_phi,2.0);

如前所述，编译器会内联所有对设备函数的调用，包括数学函数。一个可能的优化是，将通用函数 pow 替换为用于对变量进行平方的具体代码，如下所示：

 if(i <= nx && j <= ny && z <= nz)
    {
      m = i + ldx * (j + ldy * z);
      current_pos = m + current;

      current_phi = phi[m];
      current_phi_2 = current_phi * current_phi;

重新编译新代码，我们观察到这些改变将VGPRs使用量从102减少到100：

hipcc --offload-arch=gfx90a lbm_nopow_1.cpp -Rpass-analysis=kernel-resource-usage -c
lbm_nopow_1.cpp:16:1: remark: Function Name: _Z6kernelPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiiiiiddddddddddddddd [-Rpass-analysis=kernel-resource-usage]
{
^
lbm_nopow_1.cpp:16:1: remark:     SGPRs: 98 [-Rpass-analysis=kernel-resource-usage]
lbm_nopow_1.cpp:16:1: remark:     VGPRs: 100 [-Rpass-analysis=kernel-resource-usage]
lbm_nopow_1.cpp:16:1: remark:     AGPRs: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_nopow_1.cpp:16:1: remark:     ScratchSize [bytes/lane]: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_nopow_1.cpp:16:1: remark:     Occupancy [waves/SIMD]: 4 [-Rpass-analysis=kernel-resource-usage]
lbm_nopow_1.cpp:16:1: remark:     SGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_nopow_1.cpp:16:1: remark:     VGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_nopow_1.cpp:16:1: remark:     LDS Size [bytes/block]: 0 [-Rpass-analysis=kernel-resource-usage]

虽然减少的幅度看起来并不显著，但这为后续优化提供了更大的空间。

优化 n.2：将变量定义尽量靠近其首次使用的位置

一旦一个变量被定义，它的值会存储在寄存器中以备将来使用。如果在内核开头定义变量并在结尾使用，会显著增加寄存器的使用量。一个可能带来显著收益的优化是查找这些变量定义与其首次使用之间“距离较远”的情况，并手动重新排列代码。

经过快速的目视检查，我们可以看到数组位置 f[m] 的定义并不依赖于 ux、`uy` 或 uz，与其他数组 f1 到 f6 不同。

      mu_phi = alpha * current_phi * ( current_phi_2 - phi2 ) - k * laplacian_phi[m];

      fx = mu_phi * grad_phi_x[m];
      fy = mu_phi * grad_phi_y[m];
      fz = mu_phi * grad_phi_z[m];

      ux = ( g1[current_pos] - g2[current_pos] + g7[current_pos] - g8[current_pos] + g9[current_pos] -
       g10[current_pos] + g11[current_pos] - g12[current_pos] + g13[current_pos] - g14[current_pos] +
       0.50 * fx ) * 1.0/rho;
      uy = ( g3[current_pos] - g4[current_pos] + g7[current_pos] - g8[current_pos] - g9[current_pos] +
       g10[current_pos] + g15[current_pos] - g16[current_pos] + g17[current_pos] - g18[current_pos] +
       0.50 * fy ) * 1.0/rho;
      uz = ( g5[current_pos] - g6[current_pos] + g11[current_pos] - g12[current_pos] - g13[current_pos] +
       g14[current_pos] + g15[current_pos] - g16[current_pos] - g17[current_pos] + g18[current_pos] +
       0.50 * fz ) * 1.0/rho;

      af = 0.50 * gamma * mu_phi * itauphi;
      cf = itauphi * ieta * current_phi;

      f0[m] = itauphi1 * f0[m] + -3.0 * gamma * mu_phi * itauphi + itauphi * current_phi;

      f1[current_pos] = itauphi1 * f1[current_pos] + af + cf * ux;
      f2[current_pos] = itauphi1 * f2[current_pos] + af - cf * ux;
      f3[current_pos] = itauphi1 * f3[current_pos] + af + cf * uy;
      f4[current_pos] = itauphi1 * f4[current_pos] + af - cf * uy;
      f5[current_pos] = itauphi1 * f5[current_pos] + af + cf * uz;
      f6[current_pos] = itauphi1 * f6[current_pos] + af - cf * uz;

将 f[m] 的定义移动到 ux 定义之前：

      mu_phi = alpha * current_phi * ( current_phi_2 - phi2 ) - k * laplacian_phi[m];

      f0[m] = itauphi1 * f0[m] + -3.0 * gamma * mu_phi * itauphi + itauphi * current_phi;

      fx = mu_phi * grad_phi_x[m];
      fy = mu_phi * grad_phi_y[m];
      fz = mu_phi * grad_phi_z[m];

      ux = ( g1[current_pos] - g2[current_pos] + g7[current_pos] - g8[current_pos] + g9[current_pos] -
       g10[current_pos] + g11[current_pos] - g12[current_pos] + g13[current_pos] - g14[current_pos] +
       0.50 * fx ) * 1.0/rho;
      uy = ( g3[current_pos] - g4[current_pos] + g7[current_pos] - g8[current_pos] - g9[current_pos] +
       g10[current_pos] + g15[current_pos] - g16[current_pos] + g17[current_pos] - g18[current_pos] +
       0.50 * fy ) * 1.0/rho;
      uz = ( g5[current_pos] - g6[current_pos] + g11[current_pos] - g12[current_pos] - g13[current_pos] +
       g14[current_pos] + g15[current_pos] - g16[current_pos] - g17[current_pos] + g18[current_pos] +
       0.50 * fz ) * 1.0/rho;

      af = 0.50 * gamma * mu_phi * itauphi;
      cf = itauphi * ieta * current_phi;

      f1[current_pos] = itauphi1 * f1[current_pos] + af + cf * ux;
      f2[current_pos] = itauphi1 * f2[current_pos] + af - cf * ux;
      f3[current_pos] = itauphi1 * f3[current_pos] + af + cf * uy;
      f4[current_pos] = itauphi1 * f4[current_pos] + af - cf * uy;
      f5[current_pos] = itauphi1 * f5[current_pos] + af + cf * uz;
      f6[current_pos] = itauphi1 * f6[current_pos] + af - cf * uz;

我们注意到新的 VGPRs 使用量为 96，这使我们在 SIMD 上的占用从四个波提高到五个波：

hipcc --offload-arch=gfx90a lbm_rearrage_2.cpp -Rpass-analysis=kernel-resource-usage -c
lbm_rearrage_2.cpp:16:1: remark: Function Name: _Z6kernelPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiiiiiddddddddddddddd [-Rpass-analysis=kernel-resource-usage]
{
^
lbm_rearrage_2.cpp:16:1: remark:     SGPRs: 94 [-Rpass-analysis=kernel-resource-usage]
lbm_rearrage_2.cpp:16:1: remark:     VGPRs: 96 [-Rpass-analysis=kernel-resource-usage]
lbm_rearrage_2.cpp:16:1: remark:     AGPRs: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_rearrage_2.cpp:16:1: remark:     ScratchSize [bytes/lane]: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_rearrage_2.cpp:16:1: remark:     Occupancy [waves/SIMD]: 5 [-Rpass-analysis=kernel-resource-usage]
lbm_rearrage_2.cpp:16:1: remark:     SGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_rearrage_2.cpp:16:1: remark:     VGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_rearrage_2.cpp:16:1: remark:     LDS Size [bytes/block]: 0 [-Rpass-analysis=kernel-resource-usage]

使用 restrict 关键字的说明

在类似 C++ 的 C 语言类型中，别名是实现高性能的主要限制之一。为了避免这个问题，C99 标准引入了“受限指针”：一种用户告知编译器不同的对象指针类型和函数参数数组不指向重叠内存区域的方式。这允许编译器执行更激进的优化，否则由于别名问题，这些优化可能无法进行。使用受限指针可能会增加寄存器压力，因为编译器会尝试通过将更多数据存储在寄存器中来重用数据。在 AMD 硬件上并非总是如此，有时使用 restrict 有助于减少 SGPR 和 VGPR 的压力。作为经验法则，在函数参数上使用 restrict 会倾向于减少 SGPR 的使用，同时可能会增加 VGPR 的使用。

例如，让我们在 g14 数组上添加 restrict 关键字，因为它在其余代码中多次重复使用，重新使用可能会带来更高的性能。

__global__ void kernel (double *  phi, double *  laplacian_phi,
              double *  grad_phi_x, double * grad_phi_y, double *  grad_phi_z,
              double *  f0, double *  f1, double *  f2, double *  f3, double *  f4,
              double *  f5, double *  f6,
              double *  g0, double *  g1, double *  g2, double *  g3, double *  g4,
              double *  g5, double *  g6, double*  g7, double *  g8, double *  g9,
              double *  g10, double *  g11, double *  g12, double *  g13, double * __restrict__ g14,
              double *  g15, double *  g16, double *  g17, double *  g18,
              int nx, int ny, int nz, int ldx, int ldy, int current, int next,
              double k, double alpha, double phi2, double gamma,
              double itauphi, double itauphi1, double ieta,
              double itaurho, double grav,
              double eg1, double eg2, double eg0, double egc0, double egc1, double egc2)

结果是 SGPR 和 VGPR 的寄存器压力都减少了：

lbm_2_restrict.cpp:16:1: remark: Function Name: _Z6kernelPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiiiiiddddddddddddddd [-Rpass-analysis=kernel-resource-usage]
{
^
lbm_2_restrict.cpp:16:1: remark:     SGPRs: 86 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     VGPRs: 94 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     AGPRs: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     ScratchSize [bytes/lane]: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     Occupancy [waves/SIMD]: 5 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     SGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     VGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     LDS Size [bytes/block]: 0 [-Rpass-analysis=kernel-resource-usage]

通过将 restrict 添加到变量 g7，我们观察到 SGPR 的使用量进一步减少，但 VGPR 略有增加，但仍保持在每 SIMD 5 波的占用率：

lbm_2_restrict.cpp:16:1: remark: Function Name: _Z6kernelPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiiiiiddddddddddddddd [-Rpass-analysis=kernel-resource-usage]
{
^
lbm_2_restrict.cpp:16:1: remark:     SGPRs: 78 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     VGPRs: 96 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     AGPRs: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     ScratchSize [bytes/lane]: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     Occupancy [waves/SIMD]: 5 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     SGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     VGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage]
lbm_2_restrict.cpp:16:1: remark:     LDS Size [bytes/block]: 0 [-Rpass-analysis=kernel-resource-usage]

结论

在这篇文章中，我们高层次地描述了在AMD的CDNA™2架构上运行的HPC应用和算法中，寄存器压力的性质和影响。我们还提供了一组被证明可以有效减少寄存器压力并提高占用率的规则。需要强调的是，这篇博客文章中所展示的结果仅能在使用ROCm 5.4的基于CDNA™2的GPU上完全复制。由于编译器及其启发式方法的不断变化，当使用不同于5.4版本的ROCm时，代码示例的结果可能会有所不同。我们鼓励读者尝试这些代码示例，并在不同的ROCm版本上评估每次更改后的性能。

配套代码示例

作者们感谢Justin Chang、Maria Ruiz Varela和Gina Sitaraman的有益评论和建议。如果您有任何问题或意见，请在GitHub的讨论版联系我们。

109702008
关注
专栏目录
AMD GPU ISA 阅读
数字人生
06-09 877
在这篇博客中，我们讨论了如何阅读 AMDGCN 架构的 ISA。我们讨论了一些基本的指令类型，它们与处理器子单元和内存层次结构的关系。为了让读者熟悉 ISA，我们使用了一些例子。虽然这作为对 AMDGCN ISA 阅读的良好介绍，但鼓励读者参考特定的 AMDGCN ISA 文档。如果您有任何问题或评论，请在 GitHub 讨论区与我们联系。
大型语言模型在AMD GPU上的推理优化
数字人生
06-18 1060
在这篇博客中，我们介绍了几种软件优化技术，用于在AMD CDNA2 GPUs上部署最先进的大型语言模型（LLMs）。这些包括PyTorch 2编译、Flash Attention v2、`paged_attention`、PyTorch TunableOp以及多GPU推理。这些优化技术已经被AI社区广泛采纳。使用这些优化，根据批量大小和输入序列长度，你可以享受高达三倍的即开即用加速。
cpu和gpu已过时，npu和apu的时代开始
热门推荐
振华OPPO的博客世界
07-25 2万+
🌎CPU是中央处理器。其实就是机器的“大脑”，也是布局谋略、发号施令、控制行动的“总司令官”。CPU的结构主要包括运算器（ALU，ArithmeticandLogicUnit）、控制单元（CU,ControlUnit）、寄存器（Register）、高速缓存器（Cache）和它们之间通讯的数据、控制及状态的总线（BUS）。🌏GPU一般指图形处理器，又称为显示芯片、视觉处理器，是一种专门在个人电脑和一些移动设备上做图像和图形相关运算工作的微处理器。...
linux越狱amd卡代码,Linux 5.12代码达到2880万行 AMDGPU驱动近300万行
weixin_33019803的博客
05-14 121
受美国严重冰雪风暴天气的影响，Linus Torvalds 刚刚艰难地在 2 月底放出了 Linux 5.12 的首个“冰冻荒原”候选(RC1)更新。Phoronix 指出，合并窗口之后的 Linux 内核源代码树，包含了 6 万多个文件和 2881 万行的源码。而树上最大的内核驱动，依然来自于 AMDGPU，预计可在下一内核发布前后达成 300 万行。当然，Linux 5.12 合并窗口并不是近...
AMD 锐龙9 7950X芯片调研
qq_62617798的博客
07-18 1470
对于AMD R9 7950X芯片的调研
linux内核5万行代码,Linux 5.12代码达到2880万行 AMDGPU驱动近300万行
weixin_42118701的博客
04-29 281
原标题：Linux 5.12代码达到2880万行 AMDGPU驱动近300万行受美国严重冰雪风暴天气的影响，Linus Torvalds 刚刚艰难地在 2 月底放出了 Linux 5.12 的首个“冰冻荒原”候选(RC1)更新。 Phoronix 指出，合并窗口之后的 Linux 内核源代码树，包含了 6 万多个文件和 2881 万行的源码。 而树上最大的内核驱动，依然来自于 AMDGPU，预计可...
GPU深度报告，三大巨头，十四个国内玩家一文看懂【物联网智商精选】
人工智能学家
03-18 6526
来源: 智东西编辑：智东西内参关于GPU，你想知道的都在这里了。GPU是Graphics Processing Unit（图形处理器）的简称，它是一种专门在个人电脑、工作站、游戏机和一些移...
团头鲂Kiss2/Kiss2r cDNA 序列全长的克隆及其在早期发育中的表达分析
01-02
团头鲂Kiss2/Kiss2r cDNA 序列全长的克隆及其在早期发育中的表达分析，温久福，赵鸿昊，为了检测Kiss/Kissr系统是否在鱼类胚胎和个体早期发育过程中具有一定的作用，本文克隆获得了团头鲂的Kiss2及其受体(Kiss2r)...
论文研究 - 金墨鱼肝脏中几丁质酶的分布及几丁质酶同工酶的cDNA克隆
06-01
从esculenta的肝脏中获得了两个编码几丁质酶的全长cDNA（SeChi-1：1484 bp； SeChi-2：1748 bp）。 SeChi-1包含1377 bp的开放阅读框（ORF），编码459个氨基酸，SeChi-2包含1656 bp的ORF，编码552个氨基酸。 根据...
人肝cDNA文库中中期因子相互作用蛋白的筛选及鉴定
01-04
另外，研究中提到的pEGFP-MK、pCMV-TnT等质粒，以及它们所携带的MK完整编码序列、NKD2的全长基因等，都说明了研究者在前期做了大量的基因克隆和表达工作。 在研究结果部分，作者指出通过体外蛋白沉降实验，证实了3...
多头绒泡菌G2期cDNA表达文库的构建 (2005年)
05-22
2. **反转录**：使用逆转录酶将mRNA反转录成cDNA。 3. **文库构建**：将合成的cDNA克隆到适当的载体中，并转化至宿主细胞内。 4. **文库筛选与鉴定**：通过PCR扩增或其他方法筛选含有特定序列的克隆，并进行测序验证...
[论文精读]AI-Guardian: Defeating Adversarial Attacks using Backdoors
m0_52911108的博客
09-27 1143
本文提出了AI-Guardian框架，本质是让模型预测结果受到后门控制，而很难被攻击者操纵，从而抵御对抗攻击。在模型训练阶段，AI-Guardian会针对原始训练样本生成一批后门样本（经过后门处理的样本），按标签双射关系为后门样本分配对应标签，实现后门注入；在模型预测阶段，AI-Guardian先对输入数据进行后门处理，然后将其输入到模型，根据模型输出结果、标签双射关系返回对输入数据真实标签的预测结果。
Linux嵌入式有发展吗，以及对uboot，kernel，rootfs的领悟
rjszcb的博客
09-30 419
上一份工作是2022,11—2023，11年底，汽车公司，底层修修改改，编译镜像，修改后，客户给东西，重新编译给他，客制化配置启动参数，环境，支持应用部门，不懂怎么使用某些驱动。uboot到底是干嘛的，想想，emmc的镜像文件是怎么被读到ddr的，emmc为啥可以被分区，烧录镜像，tftp网络命令为啥可以运行，启动设备时，为啥lcd可以看到图像，kernel还没启动，怎么出图像的。不管做应用开发，驱动开发也好，知道，系统原理，驱动框架，进程管理调度，内核竞争，信号，中断，锁，这些就够了。
Python精选200Tips：171-175
AnFany
09-26 1221
Python深度学习的项目
遥感影像-实例分割数据集：iSAID 从切图到YOLO格式数据集制作详细介绍
最新发布
GIS潮流
09-30 601
开源数据集isaid标注包含实例分割，但是原始影像太大，很吃显存，一般显卡无法用原始影像直接训练，所以需要对影像进行裁剪，并生成对应的标签，因为想用yolo系列跑模型，所以将标签需要转为txt格式。
Her来了，OpenAI正式发布高级语音模式
shadowcz007的博客
09-25 467
OpenAI宣布将Advanced Voice Mode（AVM）扩展到更多付费用户，包括ChatGPT的Plus和Teams用户。该功能旨在使ChatGPT的语音交互更加自然，并新增五种新声音和改进的口音识别能力。虽然Sky声音因法律争议被移除，但AVM的整体设计和性能得到了提升。新版本不仅能在用户打断时立即停止回答，还能根据用户语音中的情感调整响应。此外，视频和屏幕共享功能暂未推出。AVM目前...
持续学习与创新能力的双重提升
thinkers
09-25 499
  ---
numpy is not available
xfsong2012的博客
09-27 325
在测试第一个程序的时候，出现Numpy is not available的错误，根据。安装的时候，安装的是最新版的numpy，检查numpy的版本为2.1.1版。我的pytorch版本为1.12.1，选择不高于1.21的版本即可。的办法知道是因为numpy与pytorch版本不一致造成的。2、安装需要的numpy版本，比如1.21版。
python 实现linear discriminant analysis线性判别分析算法
luthane的博客
09-26 506
线性判别分析（Linear Discriminant Analysis，简称LDA）是一种用于降维和分类的监督学习算法。它通过最大化类间散度和最小化类内散度来找到最佳投影方向，使得不同类别的数据在新的空间中尽可能分开，同一类别的数据尽可能靠近。LDA的基本思想LDA的基本思想是将多维数据映射到低维空间，同时保留数据之间的类别差异。具体来说，LDA通过计算类内散度矩阵（描述同一类别内部数据的分布情况）和类间散度矩阵（描述不同类别之间的差异），然后求解这两个矩阵的广义特征值问题，找到最佳的投影方向。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

109702008

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值