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 loopfor( 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 independentfor( i = 0; i < n; ++i )a[i] = b[i] + c[i];
如果用parallel构件，就相当于告诉编译器两件事
- 把接下来的循环映射成kernel
- 把每一个步长分到gangs(grid)上
#pragma acc parallel loopfor( i = 0; i < n; ++i )a[i] = b[i] + c[i];
这会造成每一个block里都只有一个thread的情况。
下面一种和上面完全一致
#pragma acc parallel{#pragma acc loopfor( 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 loopdo j = 1, mdo i = 1, na(j,i) = b(j,i) * alpha + c(i,j) * betaenddo enddo
编译器发现j被更多的数组用作stride-1索引，他会选择把j循环作为vector并行，把i循环作为gangs并行。
但是，编译器在这里有点自由。他可能会生成二维的grid，把i，j作为在gangs中多个block的索引。或者生成二维的block，把i，j作为vector中多个threads的索引。但是他总是寻找好性能的实现。
!$acc parallel loopdo j = 1, mdo i = 1, na(j,i) = b(j,i) * alpha + c(i,j) * betaenddoenddo
或者
!$acc paralleldo j = 1, m!$acc loopdo i = 1, na(j,i) = b(j,i) * alpha + c(i,j) * betaenddoenddo
这样编译器能把i或j循环转成vector并行。
!$acc parallel loopdo i = 1, ndo j = 1, ma(j,i) = b(j,i) * alpha + c(i,j) * betaenddoenddo
这样编译器更可能会把i作为gangs并行，j作为vector并行。
不紧凑的嵌套循环(Non-tight Nested Loop)
#pragma acc kernels loopfor( 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 loopfor( 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 loopfor( 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 kernelsdo i = 1, na(i) = mob(i)*charge(i)enddodo i = 2, n-1c(i) = 0.5*(a(i-1) + a(i+1))enddo !$acc end kernels
这两个循环会被独自的分析，调度和编译。A会在第二个循环开始前结束。这和下面一样。
!$acc kernels loopdo i = 1, na(i) = mob(i)*charge(i)enddo!$acc kernels loopdo i = 2, n-1c(i) = 0.5*(a(i-1) + a(i+1))enddo
但是，写成parallel构件就会很不同。
!$acc parallel!$acc loopdo i = 1, na(i) = mob(i)*charge(i)enddo!$acc loopdo i = 2, n-1c(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, nsum = sum + v(i)enddodo i = 1,nv(i) = v(i) / sumenddo!$acc end kernels
因为用的是kernel，第一个规约会在第二个循环开始前结束。
如果写成parallel，这个依赖关系就又破了
sum = 0.0!$acc parallel!$acc loop reduction(+:sum)do i = 1, nsum = sum + v(i)enddo!$acc loopdo i = 1,nv(i) = v(i) / sumenddo!$acc end parallel
事实上，多数教程都建议不要把reduction 和acc loop parallel 一起用。
总结：
Kernels 和parallel构件都用于解决一样的问题。区别是 kernels构件更隐含一些，给编译器更多的自由性选择。Parallel更加严格，需要程序员更多的分析。
后退问题)
解决办法)
