使用 OpenMP 高效异构并行编程

在某些情况下，将计算任务卸载到 GPU 等加速器意味着主机 CPU 在卸载的计算完成之前会处于空闲状态。然而，同时使用 CPU 和 GPU 资源可以提高应用程序的性能。在利用异构并行性的 OpenMP® 程序中，可以使用 master 子句来利用 CPU 和 GPU 的同时执行。在本文中，我们将向您展示如何使用 OpenMP 进行 CPU+GPU 异步计算。

SPEC ACCEL 514.pomriq MRI 重建基准测试是用 C 语言编写的，并使用 OpenMP 进行并行化。它可以将一些计算卸载到加速器上以进行异构并行执行。在本文中，我们将计算任务分配给主机 CPU 和独立的 Intel® GPU，以确保两个处理器都保持忙碌状态。我们还将使用 Intel VTune™ Profiler 来测量 CPU 和 GPU 的利用率并分析性能。

我们将研究异构并行开发的五个阶段和性能调优

1. 寻找合适的代码区域进行并行化
2. 并行化这些区域，使 CPU 和 GPU 都保持忙碌状态
3. 找到最佳的工作分配系数
4. 使用该分配系数启动异构并行应用程序
5. 测量性能改进。

最初，并行区域仅在 GPU 上运行，而 CPU 处于空闲状态（图 1）。正如您所看到的，只有“OMP 主线程”在 CPU 上执行，而 GPU 完全被 ComputeQ 卸载内核占用（GPU 执行单元→EU 阵列→活动）。

在检查代码后，我们决定复制每个数组和每个执行区域，以便在 GPU 上执行第一个副本，在 CPU 上执行第二个副本。主线程使用 OpenMP target 指令将工作卸载到 GPU。这在图 2 中示意性地显示。nowait 指令避免了在 CPU 和 GPU 上运行的线程之间不必要的同步。它们还改善了线程之间的负载平衡。

CPU 和 GPU 之间的工作分配平衡由从 STDIN 读取的 part 变量调节（图 3）。此变量是将要卸载到 GPU 的工作负载的百分比乘以 numX。其余工作将在 CPU 上完成。 OpenMP 异构并行实现的一个示例显示在图 4 中。

    float part;
    char *end;
    part   = strtof(argv[2], &end);

int TILE_SIZE = numX*part;

Qrc = (float*) memalign(alignment, numX, * sizeof (float));
Qrg = (float*) memalign(alignment, TILE_SIZE, * sizeof (float));

#pragma omp parallel
  {
	    #pragma omp master
	    #pragma omp target teams distribute parallel for nowait private (expАrg, cosArg, sinArg)
	    for (indexX = 0; indexX < TILE_SIZE; indexX++) {
	    	float QrSum = 0.0;
	    	float QiSum = 0.0;
	    	#pragma omp simd private (expАrg, cosArg, sinArg) reduction(+: Orsum, QiSum)
	    	for (indexK = 0; indexk_< numK; indexK++) {
	    		expАrg = PIX2 * (GkVals[indexk].kx * xg[indexX] + GkVals[index]. Ky * yg[indexX] + GkVals[index].kz * zg[indexX]);
	    		cosArg = cosf(expArg);
	    		sinArg = sinf(expArg);
	    		float phi = GkVals[indexk]. PhiMag;
	    		Q-Sum += phi * cosArg; Qisum += phi * sinArg;
	    	}
	    	Qrg[indexX] += QrSum;
	    	Qig[indexX] += QiSum;
	    }
	    #pragma omp for private(expАrg, cosArg, sinArg)
	    for (indexX = TILE_SIZE; indexX < numk; indexX++) {
	    	float Qrsum = 0.0;
	    	float QiSum = 0.0;
	    	#pragma omp simd private (expАrg, cosArg, sinArg) reduction(+: Orsum, QiSum)
	    	for (indexK = 0; indexk < numk; indexK++) {
	    		expАrg = PIX2 * (ckVals[indexk].kx * xc[indexX] + CkVals[index]. Ky * yc[indexX] + CkVals[indexK] .kz * zc[indexX]);
	    		cosArg = cosf(expArg);
	    		sinArg = sinf(expArg);
	    		float phi = CkVals[indexk].PhiMag;
	    		Qrsum += phi * cosArg;
	    		Qisum += phi * sinArg;
	    	}
	    	Qrc[indexX] += QrSum;
	    	Qic[indexX] += QiSum;
	    }
  }

Intel® oneAPI DPC++/C++ 编译器与以下命令行选项一起使用

‑O3 ‑Ofast ‑xCORE‑AVX512 ‑mprefer‑vector‑width=512 ‑ffast‑math
‑qopt‑multiple‑gather‑scatter‑by‑shuffles ‑fimf‑precision=low
‑fiopenmp ‑fopenmp‑targets=spir64="‑fp‑model=precise"

表 1 显示了不同 CPU 到 GPU 工作比率（即，上面描述的 part 变量）的性能。对于我们的系统和工作负载，0.65 的卸载比率在 CPU 和 GPU 之间提供了最佳的负载平衡，从而实现了处理器资源的最佳利用率。来自 Intel VTune Profiler 的配置文件显示，工作在 CPU 和 GPU 之间分布更均匀，并且两个处理器都得到有效利用（图 5）。当“OMP 主线程”提交卸载的内核（main: 237）以在 GPU 上执行时，其他“OMP 工作线程”在 CPU 上处于活动状态。

卸载部分	总时间，秒	GPU 时间，秒
0.00	61.2	0.0
0.20	51.6	8.6
0.40	41.0	16.8
0.60	31.5	24.7
0.65	28.9	26.7
0.80	34.8	32.6
1.00	43.4	40.7

表 1. 对应于不同卸载工作量（即，part 变量）的热点时间

图 6 显示了不同 part 值的运行时间。请记住，part 为零意味着没有工作被卸载到 GPU。 part 为一意味着所有工作都被卸载。很明显，跨 CPU 和 GPU 的均衡工作分配比任何极端情况都更能带来更好的性能。

OpenMP 在 CPU+GPU 系统上提供真正的异步异构执行。从我们的计时结果和 VTune 配置文件中可以清楚地看出，在 OpenMP 并行区域中让 CPU 和 GPU 保持忙碌状态可以带来最佳性能。我们鼓励您尝试这种方法。