7.OpenACC

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[0:Mw*Nw]) copyout(P[0:Mh*Nw])
for (int i=0; i<Mh; i++) {
    ...
}

is equivalent to:
#pragma acc parallel copyin(M[0:Mh*Mw])
copyin(N[0:Mw*Nw]) copyout(P[0:Mh*Nw])
{
    #pragma acc loop
    for (int i=0; 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=0; i<2048; i++) {
        #pragma acc loop worker
        for (int j=0; j<512; 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 = 23;
}    

这段代码会分配1023*32个thread,每个gang=1024, 对于每个gang来说执行a =23 是冗余的,只需要执行一次即可.再看下面的例子:

#pragma acc parallel num_gangs(32)
{
    Statement 1;
    #pragma acc loop gang
    for (int i=0; i<n; i++) {
        Statement 2;
    }
    Statement 3;
    #pragma acc loop gang
    for (int i=0; i<m; i++) {
        Statement 4;
    }
    Statement 5;
    if (condition) Statement 6;
}

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(32)
{
    Statement 1;
    #pragma acc loop gang
    for (int i=0; i<n; i++) {
        Statement 2;
    }
    Statement 3;
    #pragma acc loop gang
    for (int i=0; i<m; i++) {
        Statement 4;
    }
    Statement 5;
    if (condition) Statement 6;
}

kernel regions 

#pragma acc kernels
{
    #pragma acc loop num_gangs(1024)
    for (int i=0; i<2048; i++) {
        a[i] = b[i];
    }
    #pragma acc loop num_gangs(512)
    for (int j=0; j<2048; j++) {
        c[j] = a[j]*2;
    }
    for (int k=0; k<2048; k++) {
        d[k] = c[k];
    }
}

这段代码和前面的代码比较起来,区别是 acc kernel, 而前面的代码用的是acc parallel.

Kernel 结构主要是描述程序员的意图: 当前程序适合并行,编译器根据这个描述会有非常灵活的表现, 

而parallel则是规定,规定编译器必须把下面的代码段并行操作.