使用OpenAcc进行GPU并行加速

对于多线程的并行计算,常用的基于编译器的框架主要有OpenMP,其通过编译器指令指导编译器对代码进行并行化处理,具有平台无关且对代码改动小的优点。而随着GPU性能的提升,运算密集的应用场景中大量的运算工作已经开始交由GPU来完成,目前主要的编程工具有OpenCL、CUDA、C++ AMP等,可以通过编写特定的GPU代码将一部分运行转交GPU执行。

OpenAcc则采用类似OpenMP的思路,通过编译器指令指导编译器将代码转化为GPU代码,从而实现使用GPU运行特定的代码段。目前支持OpenAcc的编译器主要有PGI,GCC9,而真正能够比较容易的实现转化为GPU代码的编译器目前只有PGI(GCC需要结合CUDA进行编译才具备OpenACC特性)。

1. PGI的安装

PGI的安装较为简单,在PGI下载页面下载最新版本的社区版本即可,其可以支持Windows、Linux和macOS三种操作系统,GPU只支持CUDA(由于PGI已被nVIDIA收购,在新版本中已经移除了对AMD GPU的支持)。

2. OpenACC的基本语法

OpenAcc的基本语法与OpenMP非常相似,对于简单的循环进行并行化的代码如下

int main() {
    const int N = 1000;
    double A[N];
    //完全交由编译器进行并行优化
    #pragma acc kernels
    for (int i = 0; i < N; i++) {
        A[N] = (double) i * i;
    }
    return 0;
}

kernels表示对下面的循环进行并行化处理,处理的方式完全由编译器决定,编译的命令为

pgcc --acc -O2 main.c

编译后的程序实际是生成了部分CUDA代码,运行过程和CUDA编写的代码没有本质区别,过程为首先将数据从内存复制到GPU的专用内存,GPU完成运算后再将数据复制到内存。如此大部分程序都可以简单的通过增加编译指令完成并行化。

3. 性能对比

使用一个矩阵乘法的例子进行未并行化和并行化的运行速度对比,其代码如下

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <omp.h>
int main() {
    const int N = 5000;
    double A[N][N];
    double B[N][N];
    double C[N][N];
    //初始化矩阵
    for (int i = 0; i < N; i++) {
        for (int j = 0; j < N; j++) {
            A[i][j] = (double) rand() / (double) RAND_MAX;
            B[i][j] = (double) rand() / (double) RAND_MAX;
            C[i][j] = 0.0;
        }
    }
    //矩阵乘法
    double t1 = omp_get_wtime();
    #pragma acc kernels
    for (int i = 0; i < N; i++) {
        for (int j = 0; j < N; j++) {
            for (int k = 0; k < N; k++) {
                C[i][j] += A[i][k] * B[k][j];
            }
        }
    }
    double t2 = omp_get_wtime();
    printf("%f\n", t2 - t1);
    //输出一个值防止因为未使用而被编译器优化到计算的部分
    printf("%f\n", C[N -1][N - 1]);
}

在开启和未开启OpenAcc的情况下运行此代码,其中矩阵乘法部分的运行时间分别为10s和44s,CPU为Ryzen 3600,GPU为GTX1050Ti,其提升还是非常明显的。