【ARM-Linux开发】OpenACC并行编程实战笔记

今年运气比较好，学了cuda之后，了解到了gpu的另两种使用语言opencl和openacc，
opencl（Open Computing Language ，开放计算语言）是面向异构系统的并行编程语言的免费标准，支持多种设备，包含CPU（多核多线程CPU），GPU（NVIDIA，AMD），数字信号处理器（居然还支持DSP），但缺点是对源代码进行并行改进的代码量较大；
OpenACC与cudac和opencl不同，不需要学习相对更底层的东西，不需要对代码进行很大的改进，在代码中间加上相应的指令，再用相应的编译器进行编译就能对源程序进行加速，因为是编译器自动转换为并行处理的语言所以效率比不上用cuda或着用OpenCL对源代码进行改进的效率，而且现在OpenACC只支持C/C++,Fortran（比较幸运的是之前支持OpenACC的编译器PGI只能免费试用1个月，购买要正版1w多，今年刚刚开放社区版即免费版本，这么好的东西不试一下太可惜了）
OpenACC指令包含导语和子语两部分如：

#pragma acc loop independent

中#pragma acc loop 是导语，independent是子语，导语的作用是告诉编译器接下来代码中大致要转换为怎样的并行代码（实现什么功能），子语的作用是帮助编译器更精确地改代码，具体的作用可以在用的过程中理解；
要使用OpenACC的指令要使用相应的编译器，比如gcc不支持OpenACC
用以下代码来验证：

#include<stdio.h>
#ifdef _OPENACC
#include<openacc.h>
#endif
int main()
{
    #ifdef OPENACC
        printf("Number of device :%d
",acc_get_num_devices(acc_device_not_host));
    #else
        printf("OpenACC is not support.
");
    #endif
    return;
}

如果用gcc进行编译
gcc test.c -o test.c
./test.exe
会出现 OpenACC is not support
用支持OpenACC的PGI编译器进行编译：
pgcc -acc test.c -o test.exe
./test.exe
会出现Number of device ：1
支持OpenACC的设备为一个
像cuda一样先学习循环数组进行

#include<stdio.h>
#define N 256
int main()
{
    int i,a[N],b[N],c[N];
    for(i=0;i<N;i++)
    {
        a[i]=0;
        b[i]=c[i]=i;
    }
    #pragma acc kernels
    for(i=0;i<N;i++)
    {
        a[i]=b[i]+c[i];
    }
    printf("a[N-1]=%d 
",a[N-1]);
    return 0;

}

这里通过在循环前面加上#pragma acc kernels指令来将下面的循环改为并行处理。
通过pgcc进行编译后执行可以得到结果；
pgcc -acc -Minfo klc.c -o klc.exe
通过在选项-Minfo可以返回一些编译信息：
设置PGI编译器环境的变量：export PGI_ACC_NOTITY=1
将环境变量告诉编译器可以得到运行程序时输出的一些CUDA内核配置
./klc.exe
launch CUDA kernel file=…
function =main line=12 device=0,threadid=1 num_gangs=2 num_workers=1 vector_length=128 grid=2 block=128
OpenACC中gangs，workers，vectors类似于CUDA中的grids，blocks，threads来表示线程数，线程块数，不同的是在CUDA中这些量可以表示为三维的结构，而在OpenACC中表示为一维，其中gangs对应blocks，workers、vectors对应threads；
读《OpenACC并行编程实战》后记

OpenACC中常用的一些导语与子语：
#pragma acc kernels
如同上一篇代码所显示，在代码前之间加上，编译器发现这一指令时会自动将接下来代码中可以改动的改成并行

#include<stdio.h>
#define N 256
int main()
{
    int i,a[N],b[N],c[N];
    for(i=0;i<N;i++)
    {
        a[i]=0;
        b[i]=c[i]=i;
    }
    #pragma acc kernels
    for(i=0;i<N;i++)
    {
        a[i]=b[i]+c[i];
    }
    printf("a[N-1]=%d 
",a[N-1]);
    return 0;

}

一重循环嵌套启用一个或多个gangs和相应的vectors来实现多线程，
二重循环嵌套和三重循环嵌套时增加gangs和works来实现多线程；
如果想看经过编译后的代码，可以在编译的时候选用选项nollvm和keepgpu
pgcc -acc -Minfo -ta=tesla：nollvm，keepgpu test.c
可能得到中间代码文件 test.n001.gpu ,其中tesla为显卡的架构
#pragma acc loop
用loop相对于前面的kernel，可以更加准确地指导编译器的并行化工作
loop导语直接跟着循环语句
loop在使用时会自动检测数据的依赖性，当数据相互依赖时会将数据串行运行如下面例子：

#include<stdio.h>
#define N 1024
int main()
{
    int i,a[N],b[N],c[N];
    for(i=0;i<N;i++)
    {
        a[i]=0;
        b[i]=c[i]=i;
    }
    #pragma acc kernels
    {
        #pragma acc loop
        for(i=0;i<N;i++)
            a[i]=b[i]+c[i];
        #pragma acc loop
        for(i=0;i<N;i++)
            b[i]=b[i-1];
    }
    printf("b[2]=%d
",b[2]);
    return 0;
}

显然第一个loop下面的循环中的数据不是相互依赖的可以转化为并行，
第二个loop下面的循环中数据是相互依赖的，所以只能以串行的方式进行
最后返回值为：

b[2]=0;

independent子语告诉编译器该循环的迭代步是相互独立的，强制允许生成并行代码

#include<stdio.h>
#define N 1024
int main()
{
    int i,a[N],b[N],c[N];
    for(i=0;i<N;i++)
    {
        a[i]=0;
        b[i]=c[i]=i;
    }
    #pragma acc kernels
    {
        #pragma acc loop
        for(i=0;i<N;i++)
            a[i]=b[i]+c[i];
        #pragma acc loop independent
        for(i=0;i<N;i++)
            b[i]=b[i-1];
    }
    printf("b[2]=%d
",b[2]);
    return 0;

编译器将不检测循环内数据的依赖性而选择并行处理，最终结果为：

b[2]=1

用independent子语时编译器可能会误解原程序想表达的意思，所以要注意；
reduction子语：
reduction子语常用语一些计算的求和，乘积等,以求和为例：
s=∑ni=1ai=s1+s2=∑n1i=1ai+∑ni=n1+1ai
将数据分成两部分相加，最后再赋给s，相乘也一样
在c/c++中reduction子语试用于int，float，doubl，complex，char，wchar_t，适用于：+，*，max，min，&，|，%，&&，||
用法如下面例子：

#include<stdio.h>
#define N 101
int main()
{
    int a[N],i,ired;
    for(i=0;i<N;i++)
        a[i]=i;
    ired=0;
    #pragma acc parallel
    {
    #pragma acc loop reduction(+;ired)
    for(i=0;i<N;i++)
        ired+=a[i];
    }
    printf("ired=%d
",ired);
    return 0;
}

在reduction（；）第一个参数为数学符号，第二个参数为最后赋予值的变量