本节书摘来自华章出版社《OpenACC并行编程实战》一 书中的第3章，第3.3节，作者何沧平，更多章节内容可以访问云栖社区“华章计算机”公众号查看。
3.3 计算构件kernels
OpenACC有两个计算构件（Compute Construct）：parallel和kernels，用来将循环并行化。两个构件的目标是一样的，但行为有较大的区别。初学者应优先使用简单的kernels构件，熟练以后可以使用parallel构件，众多的子语还能够精细控制并行化方案。
想象这样一些场景：刚接触OpenACC，细节了解不多，想尽快上手体验下效果；虽然了解精细的调优技术，但是要在有限的时间内改造一个大程序，因此不再追求最优性能，希望编译器能够自动选择一些参数，尽快完成任务；程序将要在多种设备上运行，为了在所有设备上都能跑出较好的性能，参数不能写死，需要编译器根据设备环境自动选择合适的参数。此时，kernels构件正合适。
编译器将把kernels构件区域编译成一系列在加速器设备上执行的CUDA kernel，这也是kernels构件名字是复数形式的原因。
本节列出kernels构件的所有子语，但只讲解一部分，剩余子语会在后面陆续讲解。阅读后续章节时，或者编写实际代码时，可以回到此处或2.2节查阅完整的语法，而不必翻遍全书查找零碎的语法。
C和C++中，kernels导语的语法是：
#pragma acc kernels [子语列表] 换行
结构块在Fortran中的语法是：
!$acc kernels [子语列表]
结构块
!$acc end kernels这里的子语是下列中的一个：
async[(整数表达式)]
wait[(整数表达式列表)]
num_gangs(整数表达式)
num_workers(整数表达式)
vector_length(整数表达式)
device_type(设备类型列表)
if(条件)
copy(变量列表)
copyin(变量列表)
copyout(变量列表)
create(变量列表)
present(变量列表)
deviceptr(变量列表)
default(none|present)编译器将kernels区域内的代码分割为一系列的加速器内核（kernel）。通常，每个非嵌套循环成为一个单独的内核。当程序遇到一个kernels构件时，它在设备上按顺序启动这一系列内核。对不同的内核，gang的数量、每个gang包含多少个worker、vector的长度以及三者的组织方式都可能不同。
如果没有使用async子语，kernels区域结束时将有一个隐式障碍，本地线程不再向前执行，直至所有的内核都执行完毕。
用一些小例子说明kernels构件的行为特征。这些例子本身没有应用价值，只为演示语法和程序行为。
3.3.1 构件内有1个循环
例3.5和例3.6中的循环执行一维数组相加操作，很简单，不再描述代码本身的含义。
例3.5只在串行代码上添加第11行这一个导语，告诉编译器将第12～13行的循环并行化，即映射成一个CUDA内核，线程网格维数、线程块维数等参数由编译器自行选择。第15行输出检验信息，正确结果应该是2N。用pgcc编译代码：
$ pgcc -acc -Minfo k1c.c -o k1c.exe
这里的选项-Minfo要求编译器反馈一些编译信息：
main:
11, Generating copyout(a[:])
Generating copyin(b[:],c[:])
12, Loop is parallelizable
Accelerator kernel generated
Generating Tesla code
12, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */从反馈信息的最后一行可以看出，例3.5第12行的for循环确实被并行化了，编译器自行选择gang模式并行，向量（verctor）长度是128。
为了观察运行时采用的参数，设置PGI编译器的环境变量：export PGI_ACC_NOTIFY=1
环境变量PGI_ACC_NOTIFY告诉编译器，运行程序时输出CUDA内核的执行配置。运行生成的可执行程序，得到下列运行时反馈信息：
$ ./k1c.exe
launch CUDA kernel file=C:\cygwin64\home\he\accbook\kernels\k1c.c function=main line=12 device=0 threadid=1 num_gangs=2 num_workers=1 vector_length=128 grid=2 block=128这里的num_gangs是gang数量，num_workers是worker数量，vector_length是向量长度，这三个是OpenACC的术语。grid和block是CUDA术语，block对应gang，grid与OpenACC术语无确定对应关系。例3.5中的对应关系如图3.1和图3.2。
用OpenACC术语来说，长度为256的一维循环分成割成gang-0和gang-1，每个gang只包含一个worker，即worker-0，每个worker里有一个vector，vector长度是128。
用CUDA术语来说，长度为256的一维循环映射到一个网格上，这个网格包含2个线程块，每个线程块包含128个线程。
例3.5中的N值为256时，编译器选取的向量长度为128。调整N的数值，根据对第12行循环选取的向量长度，可以推测编译器的选取规律，见表3.1。
从表3.1中可知，向量长度是32的整数倍，最大长度是128。当迭代步数量位于两个向量长度之间时，采用较大的向量长度。表3.1中的32对应CUDA编程模型里的幅（Warp）。此时的设备是英伟达显卡GeForce GT 420M，在其他型号的设备上，向量长度的跳跃间隔和最大长度可能会有所不同。
Fortran版例3.6的不同之处是导语要成对出现，其他跟C版相同。
3.3.2 构件内2个循环
如果kernels构件内包含2个及以上的循环，编译器也会自动并行化。自动选取的向量长度取决于设备本身和循环的迭代次数，请看例3.7和例3.8。
语法上需要注意，例3.7中第13～16行的两循环用花括号包围起来形成一个结构块，从而第11行的kernels导语能作用到两个循环上。如果没有这两个花括号，kernels导语只对第13～14行上的循环起作用。编译器的反馈信息表明两个循环都被并行化了：
13, Loop is parallelizable
Accelerator kernel generated
Generating Tesla code
13, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
15, Loop is parallelizable
Accelerator kernel generated
Generating Tesla code
15, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */运行时反馈信息如下：
launch CUDA kernel file=C:\cygwin64\home\he\accbook\kernels\k2c.c function=main line=13 device=0 threadid=1 num_gangs=8 num_workers=1 vector_length=128 grid=8 block=128
launch CUDA kernel file=C:\cygwin64\home\he\accbook\kernels\k2c.c function=main line=15 device=0 threadid=1 num_gangs=5 num_workers=1 vector_length=128 grid=5 block=128由此运行时反馈可知，第13行的循环使用8个gang，每个只包含1个worker，每个worker包含一个vector，vector长度为128。而第15行的循环使用5个gang，其他与第13行相同。这个例子显示了一个重要特征：kernels构件中的不同循环可以使用不同的并行化参数。原因也很容易理解，每个循环映射成一个独立的CUDA内核，每个内核当然可以灵活选择执行配置参数。
3.3.3 构件内二重嵌套循环
例3.9让两个二维数组b和c相加，结果存入数组c，对应元素的相加操作使用kernels构件并行化。
例3.9中第6～8行对数组初始化，第14行输出一个数组元素验证结果的正确性。编译反馈信息表明，二维数组的并行化方案也是二维的。第11行的i循环对应gang剖分，而第12行的j循环映射为gang、vector两层循环。
编译反馈如下：
11, Loop is parallelizable
12, Loop is parallelizable
Accelerator kernel generated
Generating Tesla code
11, #pragma acc loop gang /* blockIdx.y */
12, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */运行时反馈如下：
launch CUDA kernel file=C:\cygwin64\home\he\accbook\kernels\k3c.c function=main line=12 device=0 threadid=1 num_gangs=256 num_workers=1 vector_length=128 grid=2x128 block=128用OpenACC术语来说，如图3.3，128×256个迭代步被划分成256个gang，每个gang包含1个worker，每个worker包含1个vector，vector长度为128。
用CUDA术语来说，如图3.4，128×256个迭代步映射到一个2×128的网格上，共2×128个一维线程块，每个线程块包含128个线程，即blockDim.x的值为128。
Fortran版二重嵌套循环代码的映射关系大致相同，见例3.10。需要注意，为了说明kernels构件的行为，将原本可以使用数组整体操作的运算，改用展开的循环操作。
编译反馈如下：
14, Loop is parallelizable
15, Loop is parallelizable
Accelerator kernel generated
Generating Tesla code
14, !$acc loop gang ! blockidx%y
15, !$acc loop gang, vector(128) ! blockidx%x threadidx%x运行时反馈如下：
launch CUDA kernel
file=C:\cygwin64\home\he\accbook\kernels\k3f.f90 function=main line=15 device=0 threadid=1 num_gangs=256 num_workers=1 vector_length=128 grid=1x256 block=128 因为C数组是按行存储，而Fortran数组是按列存储，所以例3.10第14～15行将列循环放在了外层。编译反馈显示，例3.10的第15行对应数组行循环，例3.9的第11行也对应数组的行循环，它们均被映射到CUDA的y指标方向，这说明编译器还是很智能的。当然，映射关系也要行列对调。
用OpenACC术语来说，如图3.5所示，数组的每列对应1个gang，每个gang包含1个worker，每个worker包含1个vector，vector的长度对应数组的行数，是128。
用CUDA术语来说，如图3.6，数组的每列对应1个一维线程块，每个线程块包含128个线程，即blockDim.x的值为128。
3.3.4 kernels构件内三重嵌套循环
对三重循环，编译器的翻译方式就不一样了。仍然沿用数组相加的例子，稍加修改就得到例3.11和例3.12。
例3.11的编译反馈显示，第13～15行的3个循环都被并行化了，但只有第14～15行上的2个循环映射到了线程上：
13, Loop is parallelizable
14, Loop is parallelizable
15, Loop is parallelizable
Accelerator kernel generated
Generating Tesla code
14, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */
15, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */例3.11的运行时反馈如下：
launch CUDA kernel file=C:\cygwin64\home\he\accbook\kernels\k4c.c function=main line=15 device=0 threadid=1 num_gangs=8 num_workers=4 vector_length=32 grid=1x8 block=32x4用OpenACC术语来说，如图3.7，例3.11的第14～15行的32×32个迭代步划分为8个gang，每个gang包含4个worker，每worker包含1个长度为32的vector。虽然对应的数组元素按二维形式组织，但是图3.7中的线程动用了三个并行层次。这是因为OpenACC的三个并行层次gang、worker和vector都是一维，不像CUDA C中的网络、线程块可以按一维、二维、三维形式组织。
用CUDA术语来说，如图3.8所示，例3.11的第14～15行的32×32个迭代步划分为8个block，每个block包含32×4个线程。
例3.11第13行的循环怎么并行化的呢？猜想是循环展开，写了32遍二维并行代码。到底怎么并行化的？不知道哇。编译的时候可以用选项nollvm和keepgpu保留下来的中间代码，像这样：
$ pgcc -acc -Minfo -ta=tesla:nollvm,keepgpu k4c.c
得到的中间代码文件k4c.n001.gpu是机器自动生成的，也是给机器看的，不是给人看的，缩排列出如下，想读懂的话要投入很多时间和耐心，读者随意感受下就好。
#include "cuda_runtime.h"
#include "pgi_cuda_runtime.h"
#include "k4c.n001.h"
extern "C" __global__ __launch_bounds__(128) void
main_15_gpu(
signed char* __restrict__ p2/* c */,
signed char* __restrict__ p3/* b */,
signed char* p4/* a */)
{
int _i_1;int i19s;int i21s;int e26;int e28;int e30;int e31;int e33;int j123;
int j121;int j119;int j84;
e30 = ((int)threadIdx.x)*(4);e26 = ((int)blockIdx.x)*(128);
e28 = ((int)threadIdx.y)*(128);i19s = 0;
e31 = (((((int)blockIdx.y)*(512))+(e26))+(e28))+(e30);j123 = 0;
_BB_11: ;
i21s = 0;j121 = 0;
_BB_12: ;
e33 = ((e31)+((i19s)*(128)))+((i21s)*(4));
if( (((-32)+(j121))>=0)) goto _BB_13;
if( (((-32)+(j123))>=0)) goto _BB_13;
if( (((((((int)blockIdx.y)*(4))+((int)threadIdx.y))+(j123))+(-32))>=0)) goto _BB_13;
if( (((((((int)blockIdx.x)*(32))+((int)threadIdx.x))+(j121))+(-32))>=0)) goto _BB_13;
j84 = -32;j119 = 0;
_BB_31: ;
_i_1 = (e33)+(j119);
*( int*)((p4/* a */)+((long long)(_i_1))) = (*( int*)((p2/* c */)+((long long)((((e31)+((i19s)*(128)))+((i21s)*(4)))+(j119)))))+(*( int*)((p3/* b */)+((long long)(_i_1))));
j84 = (j84)+(1);j119 = (j119)+(4096);
if( ((j84)<0)) goto _BB_31;
_BB_13: ;
_i_1 = ((int)gridDim.x)*(32);i21s = (_i_1)+(i21s);j121 = (j121)+(_i_1);
if( (((-32)+(j121))<0)) goto _BB_12;
_i_1 = ((int)gridDim.y)*(4);i19s = (_i_1)+(i19s);j123 = (j123)+(_i_1);
if( (((-32)+(j123))<0)) goto _BB_11;
}