OpenACC——#pragma acc kernels

#pragma acc kernels：

1、kernels构件将需要进行并行化的区域，编译成一系列可以在设备上运行的CUDA kernel，这也是为什么起名叫“kernels”的原因吧。

2、kernels区域结束时，会有一个隐式的同步障碍，线程到此等待，直到所有的线程到达此处。

3、将kernels区域映射成CUDA内核时，线程的网格维度和块维度等参数由编译器自行进行选择，编译器可以根据设备环境自动选择合适的参数。

4、kernels并不是完全按照制导语句进行并行化，如果发现问题，比如存在数据依赖，会拒绝并行化。

1、单重循环：

代码示例：

#include <iostream>
using namespace std;
#define N 256
int main()
{
	int i, a[N], b[N], c[N];
	for(i = 0; i < N; i++)
	{
		a[i] = 0;
		b[i] = c[i] = i;
	}
	#pragma acc kernels
	for(i=0; i < N; i++)
	{
		a[i] = b[i] + c[i];
	}
	cout << "a[N-1]=" << a[N-1] << endl;
   return 0;
}

编译结果：

在编译命令是加入-Minfo可以显示编译信息

从编译信息可以看出，kernels之后的for循环确实被并行化了，#pragma acc loop gang，vector(128)是编译器选择的并行模式及参数。

2、区域内双循环：

代码示例：

#include <iostream>
using namespace std;
#define N 1024
int main()
{
	int i, a[N], b[N], c[N];
	for(i = 0; i < N; i++)
	{
		a[i] = 0;
		b[i] = c[i] = i;
	}
	#pragma acc kernels
    {
        for(i=0; i < N; i++)
        {
            a[i] = b[i] + c[i];
        }
        for(i=0; i <= N/2; i++)
        {
            b[i] = 2 * a[i];
        }
    }
	cout << "b[N/2]=" << b[N/2] << endl;
   return 0;
}

反馈信息：

在环境变量中加入下面代码，可以查看相关的CUDA内核配置

export PGI_ACC_NOTIFY=1

上面的代码的反馈信息如下：

 从上面的信息可以看出，两个循环都被并行化，且不同循环可以根据需要灵活的自行设置参数。

注意：对于区域内的双循环，要将两个循环用大括号括起来，划定这个并行区域，否则只会将第一个循环并行化。

3、双重循环

示例代码：

#include <iostream>
using namespace std;
#define N 256
#define M 128
int main()
{
	int i, j, a[M][N], b[M][N], c[M][N];
	for(i = 0; i < M; i++)
	{
        for(j = 0; j < N; j++)
        {
            a[i][j] = -1;
            b[i][j] = c[i][j] = i+j;
        }
	}
	#pragma acc kernels
    for(i=0; i < M; i++)
    {
        for(j = 0; j < N; j++)
        {
            a[i][j] = b[i][j] + c[i][j];
        }
    }
        
	cout << "a[M-1][N-1]=" << a[M-1][N-1] << endl;
   return 0;
}

编译信息：

从中可以看出两个循环都被并行化。执行的配置信息自行查看吧，这里就不展示了。