oneAPI GPU 优化指南 - 使用更多的 GPU 资源

本章节翻译by chenchensmail@163.com  原文：OpenMP Offload Best Practices (intel.com)

可以通过使用更多可以并行运行的 work-item 来提高部署代码的性能，从而利用更多的 GPU 资源（填满 GPU ）。

注意:

• 循环迭代的 ND-range 划分由编译器和 runtime 启发式算法决定，还取决于 GPU 驱动程序和硬件配置。 因此它会随着时间而改变。但是，基于 LIBOMPTARGET_DEBUG=1 输出确定划分的方法将保持不变。

Collapse 子句

增加循环嵌套中并行性的一种方法是使用 collapse 子句将循环嵌套中的两个或多个循环折叠起来。 折叠会导致更多可以并行运行的迭代次数，从而在 GPU 上使用更多 work-item 。

在下面的示例中，一个由四个完美嵌套循环组成的循环嵌套被部署到 GPU 上。 parallel for 指令表示最外层循环（第 53 行）是并行的。循环中的迭代次数为 BLOCKS ，等于 8。

 1#include <stdio.h>
 2#include <stdlib.h>
 3#include <time.h>
 4
 5#include <math.h>
 6#include <omp.h>
 7
 8#define P 16
 9#define BLOCKS 8
10#define SIZE (BLOCKS * P * P * P)
11
12#define MAX 100
13#define scaled_rand() ((rand() % MAX) / (1.0 * MAX))
14
15#define IDX2(i, j) (i * P + j)
16#define IDX4(b, i, j, k) (b * P * P * P + i * P * P + j * P + k)
17
18int main(void) {
19  double w[SIZE];            /* output */
20  double u[SIZE], dx[P * P]; /* input */
21  int b, i, j, k, l;         /* loop counters */
22  double start, end;         /* timers */
23
24  omp_set_default_device(0);
25
26  /* dummy target region, so as not to measure startup time. */
27  #pragma omp target
28  { ; }
29
30  /* initialize input with random values */
31  srand(0);
32  for (int i = 0; i < SIZE; i++)
33    u[i] = scaled_rand();
34
35  for (int i = 0; i < P * P; i++)
36    dx[i] = scaled_rand();
37
38  /* map data to device */
39  #pragma omp target enter data map(to: u[0:SIZE], dx[0:P * P])
40
41  start = omp_get_wtime();
42
43  /* offload the kernel with no collapse clause */
44  #pragma omp target teams distribute parallel for \
45    private(b, i, j, k, l)
46  for (b = 0; b < BLOCKS; b++) {
47    for (i = 0; i < P; i++) {
48      for (j = 0; j < P; j++) {
49        for (k = 0; k < P; k++) {
50          double ur = 0.;
51          double us = 0.;
52          double ut = 0.;
53
54          for (l = 0; l < P; l++) {
55            ur += dx[IDX2(i, l)] * u[IDX4(b, l, j, k)];
56            us += dx[IDX2(k, l)] * u[IDX4(b, i, l, k)];
57            ut += dx[IDX2(j, l)] * u[IDX4(b, i, j, l)];
58          }
59
60          w[IDX4(b, i, j, k)] = ur * us * ut;
61        }
62      }
63    }
64  }
65
66  end = omp_get_wtime();
67
68  #pragma omp target exit data map(from: w[0:SIZE])
69
70  /* print result */
71  printf("no-collapse-clause: w[0]=%lf time=%lf\n", w[0], end - start);
72
73  return 0;
74}

编译命令:

icx -fiopenmp -fopenmp-targets=spir64 test_no_collapse.cpp

运行命令:

OMP_TARGET_OFFLOAD=MANDATORY ZE_AFFINITY_MASK=0.0 LIBOMPTARGET_DEBUG=1 ./a.out

libomptarget.so 调试信息（在环境变量 LIBOMPTARGET_DEBUG=1 时在运行时发出） 显示了循环迭代的 ND-range 划分以及如何通过使用 collapse 子句来增加并行性。 在输出中， Lb 和 Ub 分别指划分中每个维度的并行循环下限和上限。

没有 collapse 子句， LIBOMPTARGET_DEBUG=1 输出显示了关于第 50 行的 target 区域的以下信息。

Libomptarget --> Launching target execution __omp_offloading_3d_9b5f515d__Z4main_l45 with pointer 0x000000000143d5d8 (index=1).
Target LEVEL0 RTL --> Executing a kernel 0x000000000143d5d8...
Target LEVEL0 RTL --> Assumed kernel SIMD width is 32
Target LEVEL0 RTL --> Preferred group size is multiple of 64
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 7, Stride = 1
Target LEVEL0 RTL --> Group sizes = {1, 1, 1}
Target LEVEL0 RTL --> Group counts = {8, 1, 1}

请注意，没有 collapse 子句，并行循环迭代次数 = 8 ，因为最外层循环(BLOCKS)的上限 = 8。 在这种情况下，我们最终得到 8 个 work-group ，每个 work-group 有一个 work-item (work-group 数量 = 8 x 1 x 1 = 8 ，每个 work-group 大小= 1 x 1 x 1 = 1个 work-item)。 kernel 使用 SIMD 32 进行矢量化，这意味着每个 work-group 中的 32 个 work-item 被合并成一个 sub-group。 由于我们每个 work-group 只有一个 work-item ，因此每个 work-group 只有一个 sub-group ，其中只有一个 SIMD 通道处于 active 状态。

我们可以通过在 parallel for 指令上添加一个 collapse 子句来增加并行性，从而增加 GPU 上使用的 work-item 数量。 我们首先添加 collapse(2) 子句，如下面修改后的示例所示。

49  /* offload the kernel with collapse clause */
50  #pragma omp target teams distribute parallel for collapse(2) \
51    private(b, i, j, k, l)
52  for (b = 0; b < BLOCKS; b++) {
53    for (i = 0; i < P; i++) {
54      for (j = 0; j < P; j++) {
55        for (k = 0; k < P; k++) {
56          double ur = 0.;
57          double us = 0.;
58          double ut = 0.;
59
60          for (l = 0; l < P; l++) {
61            ur += dx[IDX2(i, l)] * u[IDX4(b, l, j, k)];
62            us += dx[IDX2(k, l)] * u[IDX4(b, i, l, k)];
63            ut += dx[IDX2(j, l)] * u[IDX4(b, i, j, l)];
64          }
65
66          w[IDX4(b, i, j, k)] = ur * us * ut;
67        }
68      }
69    }
70  }

当使用 collapse(2) 时， LIBOMPTARGET_DEBUG=1 输出显示了以下划分。

Libomptarget --> Launching target execution __omp_offloading_3d_9b5f515f__Z4main_l45 with pointer 0x00000000017f45d8 (index=1).
Target LEVEL0 RTL --> Executing a kernel 0x00000000017f45d8...
Target LEVEL0 RTL --> Assumed kernel SIMD width is 32
Target LEVEL0 RTL --> Preferred group size is multiple of 64
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 15, Stride = 1
Target LEVEL0 RTL --> Level 1: Lb = 0, Ub = 7, Stride = 1
Target LEVEL0 RTL --> Group sizes = {1, 1, 1}
Target LEVEL0 RTL --> Group counts = {16, 8, 1}

请注意，使用 collapse(2) 时，并行循环迭代次数 = BLOCKS x P = 8 x 16 = 128。 在这种情况下，我们最终得到 128 个 work-group ，每个 work-group 有 1 个 work-item (总 work-group 数量 = 16 x 8 x 1 = 128 ，每个 work-group 大小 = 1 x 1 x 1 = 1个 work-item)。 kernel 使用 SIMD 32 进行矢量化，这意味着每个 work-group 中的 32 个 work-item 被合并成一个 sub-group。 由于我们每个 work-group 只有一个 work-item ，因此每个 work-group 只有一个 sub-group ，其中只有一个 SIMD 通道处于 active 状态。

另一方面，如果我们使用 collapse(3) 子句，则 LIBOMPTARGET_DEBUG=1 输出显示以下划分。

Libomptarget --> Launching target execution __omp_offloading_3d_9b5f5160__Z4main_l45 with pointer 0x0000000001728d08 (index=1).
Target LEVEL0 RTL --> Executing a kernel 0x0000000001728d08...
Target LEVEL0 RTL --> Assumed kernel SIMD width is 32
Target LEVEL0 RTL --> Preferred group size is multiple of 64
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 15, Stride = 1
Target LEVEL0 RTL --> Level 1: Lb = 0, Ub = 15, Stride = 1
Target LEVEL0 RTL --> Level 2: Lb = 0, Ub = 7, Stride = 1
Target LEVEL0 RTL --> Group sizes = {8, 1, 1}
Target LEVEL0 RTL --> Group counts = {2, 16, 8}

使用 collapse(3) 时，所产生的并行循环迭代次数 = BLOCKS x P x P = 8 x 16 x 16 = 2048。 在这种情况下，我们有 256 个 work-group ，每个 work-group 有 8 个 work-item （总 work-group 数量 = 2 x 16 x 8 = 256 ，每个 work-group 大小 = 8 x 1 x 1 = 8 个 work-item ）。 kernel 使用 SIMD 32 进行矢量化，这意味着每个 work-group 中的 32 个 work-item 被合并成一个 sub-group。 由于我们每个 work-group 只有 8 个 work-item ，因此我们只有一个 sub-group，其中只有 8 个 SIMD 通道处于 active 状态。

如果我们使用 collapse(4) 子句而不是 collapse(3) 子句，则 LIBOMPTARGET_DEBUG=1 输出显示以下划分。

Target LEVEL0 RTL --> Executing a kernel 0x0000000001aab5d8...
Target LEVEL0 RTL --> Assumed kernel SIMD width is 32
Target LEVEL0 RTL --> Preferred group size is multiple of 64
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 32767, Stride = 1
Target LEVEL0 RTL --> Group sizes = {64, 1, 1}
Target LEVEL0 RTL --> Group counts = {512, 1, 1}

使用 collapse(4) 时，所产生的并行循环迭代次数= BLOCKS x P x P x P = 8 x 16 x 16 x 16 = 32768。 在这种情况下，我们有 512 个 work-group ，每个 work-group 有 64 个 work-item (总 work-group 数量 = 512 x 1 x 1 =512 ，每个 work-group 大小 = 64 x 1 x 1 = 64 个 work-item)。 kernel 使用 SIMD32 进行矢量化，这意味着每 32 个 work-item 被合并成一个 sub-group。 因此每个 work-group 都有 2 个 sub-group。

使用 collapse 子句显著减少了循环嵌套的运行时间。在特定 GPU 上 runtime （仅 1 堆栈），各版本的性能如下:

没有 collapse 版本 : 0.002430 秒
collapse(2) 版本 : 0.000839 秒
collapse(3) 版本 : 0.000321 秒
collapse(4) 版本 : 0.000325 秒

上述时间显示添加 collapse(3) 或 collapse(4) 子句可提供约 7.5 倍的性能提升。(0.000321 秒对比 0.002430 秒)。

注意事项:

• 在 GPU 上， collapse 子句可能根本不会导致任何实际的循环折叠， 但该子句向编译器和 runtime 传达了循环嵌套中的并行度，并用于确定 ND-range 划分。

• 要利用矢量加载和存储，请建议不要将循环嵌套中的最内层循环包含在折叠中以便进行矢量化。 当最内层循环具有单位步长且迭代次数至少与 SIMD 宽度一样大时，将获得最佳性能。

上一章                                    主目录​​    上级目录                                                               下一章