OpenACC:
openacc 可以用于fortran, c 和 c++程序,可以运行在CPU或者GPU设备.
openacc的代码就是在原有的C语言基础上进行修改,通过添加:
compiler directives 编译器指令(pragmas): #pragma 来标示.
cuda 中有 __syncthreads()来进行线程同步,目前的OpenAcc还没有线程同步机制.
OpenAcc device model
OpenAcc excute model
parallel loops
下面地一段代码和第二段代码是等效的,在OpenAcc中一个parallel区域有一个单个loop组成.
#pragma acc parallel loop copyin(M[0:Mh*Mw])
copyin(N[:Mw*Nw]) copyout(P[:Mh*Nw])
for (int i=; i<Mh; i++) {
...
} is equivalent to:
#pragma acc parallel copyin(M[0:Mh*Mw])
copyin(N[:Mw*Nw]) copyout(P[:Mh*Nw])
{
#pragma acc loop
for (int i=; i<Mh; i++) {
...
}
}
}
copyin对应拷贝内存从host到device,
copyout对应拷贝内存从device到host
gangs and workers
gangs可以类比成cuda的block,
workers可以类比成thread
#pragma acc parallel num_gangs(1024) num_workers(32)
{
#pragma acc loop gang
for (int i=; i<; i++) {
#pragma acc loop worker
for (int j=; j<; j++) {
foo(i,j);
}
}
}
这段代码会分配: 1024*32 = 32K 个thread, 这两个循环题一共是执行2048*512 = 1M, 所以每个thread执行foo()函数 1M/32K = 32 次.
再看另外一个代码:
#pragma acc parallel copyout(a) num_gangs(1024) num_workers(32)
{
a = ;
}
这段代码会分配1023*32个thread,每个gang=1024, 对于每个gang来说执行a =23 是冗余的,只需要执行一次即可.再看下面的例子:
#pragma acc parallel num_gangs(32)
{
Statement ;
#pragma acc loop gang
for (int i=; i<n; i++) {
Statement ;
}
Statement ;
#pragma acc loop gang
for (int i=; i<m; i++) {
Statement ;
}
Statement ;
if (condition) Statement ;
}
gang有32个,statement2的循环次数是n,statement4循环次数是m, 最终到底分配多少个thread取决于编译器,有可能m>n,则分配m个,当然实际情况可能更加复杂,
statement1, 3, 5,6 对于32gang来说是冗余的,情况和上面的相同,可以看出OpenAcc中的冗余是对于gang来说的,下面的这种写法可以消除这种冗余:
#pragma acc parallel num_gangs(1)
num_workers()
{
Statement ;
#pragma acc loop gang
for (int i=; i<n; i++) {
Statement ;
}
Statement ;
#pragma acc loop gang
for (int i=; i<m; i++) {
Statement ;
}
Statement ;
if (condition) Statement ;
}
kernel regions
#pragma acc kernels
{
#pragma acc loop num_gangs(1024)
for (int i=; i<; i++) {
a[i] = b[i];
}
#pragma acc loop num_gangs(512)
for (int j=; j<; j++) {
c[j] = a[j]*;
}
for (int k=; k<; k++) {
d[k] = c[k];
}
}
这段代码和前面的代码比较起来,区别是 acc kernel, 而前面的代码用的是acc parallel.
Kernel 结构主要是描述程序员的意图: 当前程序适合并行,编译器根据这个描述会有非常灵活的表现,
而parallel则是规定,规定编译器必须把下面的代码段并行操作.