基于国产加速器海光DCU&GPGPU的OpenMP Offload多卡编程实践

目录

一、一个最简单的OpenMP Offload程序示例

二、 改进：通过omp_target_memcpy自主的控制数据传输

三、优化Offload计算中线程分组以及与loop计算任务的对应关系

四、编译与执行示例

五、讨论

---

随着OpenMP Offload技术的发展，特别是相关的OpenACC技术证明了该方向的显著编程优势以及性能方面劣势的补足，该技术方向已经逐渐脱离了简单的导语编程模式，逐渐具备了跟原生编程技术同样的编程控制力与性能，本文简单的介绍下基于DCU的OpenMP Offload编程经验，特别是在针对多卡计算等场景，给出一个编程示范，并讨论两个技术建议。

一、一个最简单的OpenMP Offload程序示例

本示例中，简单的计算一个数组 x 乘以一个标量 a 的结果，并记为 y 。具体的代码见本文最后面的附件中。
核心的代码如下所示：

#pragma omp target device(0) map(tofrom: x[0:n],y[0:n])
{
#pragma omp parallel for 
for (int i=0;i<n;i++){
    y[i] = a * x[i];
}

其中，pragma omp target device(0) map(tofrom: x[0:n],y[0:n],a) 依次的表示，启动一个OpenMP Offload target任务，使用0号device，将主机内存中的 x、y 变量在计算前复制到设备中并在计算完成后立刻复制回主机内存。其中，变量a没有被指定，但会被自动识别进行双向数据传输，其实x、y也是不必被指定的，这体现了OpenMP语法确实可以简化与隐藏很多的细节操作。

第二行的 pragma omp parallel for 表示对以下for循环在设备汇总执行并行的计算。

这是一个最简单的示例，代码也很少，因为OpenMP技术隐藏了背后的很多操作细节，提供了简洁的编程，但同时也限制了功能，特别是限制了优化性能所必须的细节控制能力，最显著的是，我们可能需要一些数据按需的复制到内存，同时再多次Offload后，再回传到主机内存，因此我们需要自主的控制数据的传输。

二、 改进：通过omp_target_memcpy自主的控制数据传输

控制数据传输并不总是需要omp_target_memcpy函数，实际上，可以通过如下target enter date map 语句来在设备上建立同样变量名的数据，并通过target update来随时的双向传输数据：

#pragma omp target enter data map(alloc: x[10*idev:10*idev+10], y[10*idev:10*idev+10], a) device(idev)
#pragma omp target update from(y[c:c+CHUNKSZ])  device(idev)

该方法应该是一个好的技术趋势，但是其隐藏了设备变量，不利于清晰的完成多设备关联的数据控制，因此本文使用更加底层的函数omp_target_memcpy来示例。

为了做主句与设备间的数据传输，首先需要建立设备上的变量，并分配空间：

double *y_d0;
double *x_d0;
double *a_d0;
x_d0 = (double *) omp_target_alloc( sizeof(double) * n/ndev, 0);
y_d0 = (double *) omp_target_alloc( sizeof(double) * n/ndev, 0);
a_d0 = (double *) omp_target_alloc( sizeof(double) , 0);
double *y_d1;
double *x_d1;
double *a_d1;
x_d1 = (double *) omp_target_alloc( sizeof(double) * n/ndev, 1);
y_d1 = (double *) omp_target_alloc( sizeof(double) * n/ndev, 1);
a_d1 = (double *) omp_target_alloc( sizeof(double) , 1);

如上所示，omp_target_alloc 可以方便的帮助分配数据空间，之后可以使用omp_target_memcpy来显示的在任意需要时拷贝数据。

omp_target_memcpy( x_d0, &x[ipart*CHUNKSZ], sizeof(double)*CHUNKSZ, sizeof(double)*deviC*CHUNKSZ, 0,  idev,   hdev); 
omp_target_memcpy( &x[ipart*CHUNKSZ], x_d0, sizeof(double)*CHUNKSZ, 0, sizeof(double)*deviC*CHUNKSZ,  hdev,   idev); 
omp_target_memcpy( a_d0, &a, sizeof(double), 0, 0,  idev,   hdev); 

以上分别是，将x从主机内存hdev拷贝到设备idev，将x从设备idev拷贝到主机内存hdev，将a变量从主机内存hdev拷贝到设备idev，其中hdev是通过 int hdev = omp_get_initial_device() 函数来获取，idev分别是从0开始的GPU卡编号。

此外，程序在标记OpenMP计算时，还需要通过is_device_ptr指定x_d0, y_d0, a_d0是设备指针，完整的代码段为：

>    #pragma omp target device(idev) is_device_ptr( x_d0, y_d0, a_d0 ) 
>    { 
>    #pragma omp teams num_teams(2) thread_limit(5)
>    { 
>      #pragma omp distribute parallel for dist_schedule(static,10) schedule(static,1) 
>      for (i = deviC*ChunkSize; i < deviC*ChunkSize+ChunkSize; i++){
>        y_d0[i] = a_d0[0] * x_d0[i];
>      }
>    }
>    }

其中，变量后缀d0表达了该变量为设备0中的指针，最后一行将于设备中新计算得到的y_d0数据传输回主机内存。

三、优化Offload计算中线程分组以及与loop计算任务的对应关系

线程会被分组，单组的线程数量由thread_limit限制，组数量由num_teams限制，更具体的规则不被OpenMP规则所明确，由各个实现决定，一般情况下此两个数据指定了程序使用的组数量与线程数量。每个组具有一个主线程，可以独立执行任务，例如使用#pragma omp master可以开启一段计算，team内的除主线程之外的其他线程均不会参与计算。
如下是一个较完整的示例：

>  #pragma omp target device(idev) is_device_ptr( x_d0, y_d0, a_d0 ) 
>  { 
>  #pragma omp teams num_teams(2) thread_limit(5)
>    { 
>    #pragma omp distribute parallel for dist_schedule(static,10) schedule(static,1)  
>    for (i = deviC*ChunkSize; i < deviC*ChunkSize+ChunkSize; i++){
>        y_d0[i] = a_d0[0] * x_d0[i];
>    }
>  }

其中，distribute指示将loop任务分配给各个组，dist_schedule指示了loop任务分配给各个组的分配策略，schedule指示了将team组中的loop任务分配给各个组内threads的分配方式。

这些配置参数显著的提升了开发者对计算的控制力，特别是调整这部分划分与分配规则，可以提升对数据的访问效率，对计算速度会起到很大的影响，应该根据具体问题进行具体的配置。

四、编译与执行示例

下载附件后，首先在昆山中心申请一个交互式节点，载入DTK环境compiler/rocm/dtk-22.10.1，然后采用hipcc来编译，编译命令为 hipcc -target x86_64-pc-linux-gnu -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 main.c -o multi-dev.x，其已经被封装到Makefile中，然后./multi-dev.x来执行。
简单的操作如下：

$ salloc -p kshdtest   -J test --time=0-0:30:00 -N 1 -n 1 --cpus-per-task=2  --gres=dcu:2 --qos=partition_hgdcutest srun --pty bash

$ module purge 
$ module load compiler/rocm/dtk-22.10.1 2>&1 
$ make clean
$ make 
$ export OMP_TARGET_OFFLOAD=MANDATORY
$ export LIBOMPTARGET_KERNEL_TRACE=1 
$ make test

其中，在make test执行测试之前，通过配置OMP_TARGET_OFFLOAD环境变量可以避免代码在CPU中执行，通过设置LIBOMPTARGET_KERNEL_TRACE环境变量，可以使输出更详细的Kernel执行信息，包括启用的team数，以及每个team中的thread数。

五、讨论

本文演示了一个简单的OpenMP Offload程序，不过其中采用了较底层的函数API，可以为开发者提供接近原生编程语言的控制力。
在扩展到多卡的过程中，需要对不同卡上的数据分配独立的变量，以后缀“_d0”、“_d1”来区分，以便用is_device_ptr来标记设备指针，这导致如下的多卡编程方式：

>  if (idev==0){
>    #pragma omp target device(idev) is_device_ptr( x_d0, y_d0, a_d0 ) 
>    ...
>  }
>  else if (idev==1) {
>    #pragma omp target device(idev) is_device_ptr( x_d1, y_d1, a_d1 ) 
>    ...
>  }

在成熟的开发中，我们希望对不同卡进行统一的编程，使用一个指针数组变量 x_d[idev] 来统一描述多个设备上的数据，例如：

for (int idev=0;idev<ndev;idev++) {
      x_d[idev] = (double *) omp_target_alloc( sizeof(double) * n/ndev, idev);
      y_d[idev] = (double *) omp_target_alloc( sizeof(double) * n/ndev, idev);
      a_d[idev] = (double *) omp_target_alloc( sizeof(double) , idev);
...
      #pragma omp target device(idev) is_device_ptr( x_d[idev], y_d[idev], a_d[idev] )
...
}

如上模式需要底层工具的支持。

此外，当前的DCU开发支持的OpenMP4.5规范中，没有异步的内存拷贝API，导致代码编写过程中需要将更多数据存在GPU中，这对一部分逻辑处理过程不太友好，也是OpenMP与OpenACC主流的编译器实现之间的关键区别，特别是导致一些隐式数据传输的offload代码之间具有较大的速度差异。OpenMP5.1中规范了新的API：omp_target_memcpy_async 来提供异步拷贝的支持，弥补了OpenMP之于OpenACC的最大短板，也同样会大大缩小OpenMP跟原生的GPU开发之间的差异。

未来，在改进如上两个技术支持后，OpenMP Offload技术应该可以扩展其应用范围，获得用户的支持，因为当前最先进的GPU单节点已经具备非常强大的性能，足够大多数用户的使用。特别是，CPU算例相比GPU算例的巨大差距，导致CPU算力可以被弃用，OpenMP也将可以基于GPU-GPU直接传输的通信库做到多节点的互联计算。

#OpenMP# #Offload#

多设备源码示例下载见我的资源