OpenMP引导语卸载到GPU上运算

OpenMP是由OpenMP Architecture Review Board牵头提出的，并已被广泛接受，用于共享内存并行系统的多处理器程序设计的一套指导性编译处理方案（Compiler Directive），最开始OpenMP并不支持GPU卸载。在OpenMP 4.0（2013年发布）中，规范提供了一组指令来指示编译器和运行时将代码块卸载到设备，包括GPU设备，这也标志着OpenMP正式支持GPU。

void kunpeng_test()
{
    int len = 16;
    int x[len], y[len], z[len];
    for(int i = 0; i < len; i++)
    {
        x[i] = i;
        y[i] = i * 2;
        z[i] = 0;
     }
    #pragma omp target map(to: x[0:len], y[0:len], len) map(tofrom: z[0:len])
    for(int index = 0, index < len; index++)
    {
         z[index] = x[index] * y[index];
    }
}

其中to表示数据从Host设备传递到Device设备，from表示数据从Device设备传递到Host设备。

假设Host设备为鲲鹏服务器，Device设备为GPU设备，那么上述的数据传输如图1所示。

图1 数据传输示意图

OpenMP将卸载到硬件设备和在硬件设备上进行并行操作是不同的引导语，当然为了方便起见，OpenMP定义了复合引导语供用户使用，如下述代码就是结合了卸载和并行操作。

void kunpeng_test()
{
    int len = 16;
    int x[len], y[len], z[len];
    for(int i = 0; i < len; i++)
    {
        x[i] = i;
        y[i] = i * 2;
        z[i] = 0;
     }
    #pragma omp target teams distribute parallel for simd num_teams(num_blocks) map(to: x[0:len], y[0:len], len) map(tofrom: z[0:len])
    for(index = 0, index < len; index++)
    {
         z[index] = x[index] * y[index];
    }
}

Lammps GPU版本采用OpenMP进行卸载任务到GPU。