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 \n",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\n",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\n",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\n",ired);
return 0;
}
在reduction(;)第一个参数为数学符号,第二个参数为最后赋予值的变量