OpenACC 中parallel 和kernels的区别

Kernels构件

Kernels构件源于PGI Accelerator模型的region构件。嵌套kernels构件里的循环可能会被编译器转换成能在GPU上高效并行的部分。在这个过程中有三步。 

1：判断并行中遇到的循环。

2：把抽象的并行转换成硬件上的并行。对于NVIDIA CUDA GPU， 它会把并行的循环映射到grid层次(blockIdx) 或 thread层次(threadIdx)。OpenACC申明， gang 对应grid， vector 对应thread。编译器可能会通过strip-mining(一种拆分循环利用缓存的技术)把一层的循环映射到多层。

3：编译器生成并优化代码。

 

在kernels构件中，编译器用自动并行技术识别并行的循环。这个识别能被指令(directives)和条款(clauses)增强，比如说loop independent。使用-Minfo标志可以让PGI 编译器显示编译信息。你能看到类似Loop is parallelizable 如果编译器认为这个循环可以被并行化。。

 

在第二步中，PGI编译器使用目标硬件的模型去选择循环被映射成vector并行还是gang。比如说，循环中有多个步进为1的数组时，更多的会被映射成vector(thread)并行。在NVIDIA GPU上，这种映射更倾向于生成能同时运行的代码。同样能用 loop gang 明确生成grid上的并行，或者用loop vector(64)生成64个thread的block。但这会造成移植上的困难。

第三步底层代码生成和优化。

Parallel 构件

Parallel构件源于OpenMp的parallel构件。OpenMP会立即产生很多多余的线程，当运行到一个循环是，一个线程运行一部分。

同样的，OpenACC的parallel构件也会产生多余的gangs。不同的是OpenACC的parallel在结束时没有同步。就像带了nowait参数的OpenMP一样。

Kernels和parallel主要的不同是， 整个parallel构件会被变成一个kernel。比如说在CUDA中，会变成一个kernel<<<grid, block>>>();

 

举例分析：

单个循环

         在单个循环中kernels和parallel几乎一样。

#pragma acc kernels loop

        for( i = 0; i < n; ++i )

            a[i] = b[i] + c[i];

等价于

 #pragma acc kernels
    {
        for( i = 0; i < n; ++i )
            a[i] = b[i] + c[i];
    }

如果a,b,c是指针的话，编译器没法消除指针的歧义，可能不会生成并行代码。可以用restrict属性/加-Msafeptr(不推荐)/jia independent子语

#pragma acc kernels loop independent
        for( i = 0; i < n; ++i )
            a[i] = b[i] + c[i];

如果用parallel构件，就相当于告诉编译器两件事

1. 把接下来的循环映射成kernel
2. 把每一个步长分到gangs(grid)上

#pragma acc parallel loop
        for( i = 0; i < n; ++i )
            a[i] = b[i] + c[i];

这会造成每一个block里都只有一个thread的情况。

下面一种和上面完全一致

 #pragma acc parallel
    {
        #pragma acc loop
        for( i = 0; i < n; ++i )
            a[i] = b[i] + c[i];
    }

但是如果把里面的loop子语去掉

#pragma acc parallel
    {
        for( i = 0; i < n; ++i )
            a[i] = b[i] + c[i];
    }

循环就会在所有gang中运行。这是另一种浪费。

         嵌套循环：

  !$acc kernels loop
    do j = 1, m
        do i = 1, n
            a(j,i) = b(j,i) * alpha + c(i,j) * beta
        enddo
enddo
 

编译器发现j被更多的数组用作stride-1索引，他会选择把j循环作为vector并行，把i循环作为gangs并行。

但是，编译器在这里有点自由。他可能会生成二维的grid，把i，j作为在gangs中多个block的索引。或者生成二维的block，把i，j作为vector中多个threads的索引。但是他总是寻找好性能的实现。

 !$acc parallel loop
    do j = 1, m
        do i = 1, n
            a(j,i) = b(j,i) * alpha + c(i,j) * beta
        enddo
    enddo

或者

!$acc parallel
    do j = 1, m
        !$acc loop
        do i = 1, n
            a(j,i) = b(j,i) * alpha + c(i,j) * beta
        enddo
    enddo

这样编译器能把i或j循环转成vector并行。

 !$acc parallel loop
    do i = 1, n
        do j = 1, m
            a(j,i) = b(j,i) * alpha + c(i,j) * beta
        enddo
    enddo

这样编译器更可能会把i作为gangs并行，j作为vector并行。

不紧凑的嵌套循环(Non-tight Nested Loop)

#pragma acc kernels loop
    for( i = 0; i < nrows; ++i ){
        double val = 0.0;
        int nstart = rowindex[i];
        int nend = rowindex[i+1];
        #pragma acc loop vector reduction(+:val)
        for( n = nstart; n < nend; ++n )
            val += m[n] * v[colndx[n]];
        r[i] = val;
    }

编译器会为外循环生成gang并行，内循环vector并行

#pragma acc kernels loop
    for( i = 0; i < nrows; ++i ){
        double val = 0.0;
        int nstart = rowindex[i];
        int nend = rowindex[i+1];
        for( n = nstart; n < nend; ++n )
            val += m[n] * v[colndx[n]];
        r[i] = val;
    }

如果没有loop导语， 编译器可能会把外循环在gangs和vector上并行， 或者自动把内循环vector化。

 #pragma acc parallel loop
    for( i = 0; i < nrows; ++i ){
        double val = 0.0;
        int nstart = rowindex[i];
        int nend = rowindex[i+1];
        #pragma acc loop vector reduction(+:val)
        for( n = nstart; n < nend; ++n )
            val += m[n] * v[colndx[n]];
        r[i] = val;
    }

这样，编译器就不需要分析了。或者可以去掉loop vector reduction而让编译器自动寻找并行的机会。

相邻的循环

 !$acc kernels
    do i = 1, n
        a(i) = mob(i)*charge(i)
    enddo
    do i = 2, n-1
        c(i) = 0.5*(a(i-1) + a(i+1))
    enddo
!$acc end kernels

这两个循环会被独自的分析，调度和编译。A会在第二个循环开始前结束。这和下面一样。

 !$acc kernels loop
    do i = 1, n
        a(i) = mob(i)*charge(i)
    enddo
    !$acc kernels loop
    do i = 2, n-1
        c(i) = 0.5*(a(i-1) + a(i+1))
    enddo
 

但是，写成parallel构件就会很不同。

!$acc parallel
    !$acc loop
    do i = 1, n
        a(i) = mob(i)*charge(i)
    enddo
    !$acc loop
    do i = 2, n-1
        c(i) = 0.5*(a(i-1) + a(i+1))
    enddo
    !$acc end parallel

两个循环会同时被并行化，而且没有同步操作。A可能还没完成，第二个循环就开始了。

同样的问题可能会在reduction中出现

 sum = 0.0
    !$acc kernels
    !$acc loop reduction(+:sum)
    do i = 1, n
        sum = sum + v(i)
    enddo
    do i = 1,n
        v(i) = v(i) / sum
    enddo
    !$acc end kernels

因为用的是kernel，第一个规约会在第二个循环开始前结束。

如果写成parallel，这个依赖关系就又破了

   sum = 0.0
    !$acc parallel
    !$acc loop reduction(+:sum)
    do i = 1, n
        sum = sum + v(i)
    enddo
    !$acc loop
    do i = 1,n
        v(i) = v(i) / sum
    enddo
    !$acc end parallel

事实上，多数教程都建议不要把reduction 和acc loop parallel 一起用。

总结：

         Kernels 和parallel构件都用于解决一样的问题。区别是 kernels构件更隐含一些，给编译器更多的自由性选择。Parallel更加严格，需要程序员更多的分析。