CUDA C编程（三十）OpenACC的使用

  OpenACC是CUDA的一个补充编程模型，使用基于编译器指令的API，具有高性能、可编程性和跨平台可移植性。接下来将会介绍OpenACC的概念和方法，重点介绍CUDA和OpenACC之间的关系。
  OpenACC的线程模型与CUDA的线程模型类似，但添加了一个并行的维度。OpenACC可以分为gang、worker和vector3个并行层次。在上层，gang类类似于CUDA线程块。一个gang可包含一个或多个执行的线程，在每个gang内部每个gang都包含一个或多个worker。在CUDA中，一个worker类似于线程中的一个线程束。在CUDA中，一个worker类似于线程中的一个线程束。每个worker都有一个向量宽度，由一个或多个同时执行指令的向量元素组成。每个向量元素都类似一个CUDA线程，因为它是一个单一的执行流。OpenACC和CUDA线程模型之间的主要区别在于，OpenACC在编程模型中直接指出了worker的概念(即线程束)，而在CUDA中并没有明确建立线程束。
  OpenACC平台模型与CUDA类似，但它使用不同的术语和略有不同的抽象概念。OpenACC的目标是建立一个具有单线程的主机程序平台，在该主机程序中，将内核交付给多处理单元(PU),在此平台上，每个PU一次只运行一个gang。每个PU可以同时执行多个独立的并发执行线程(worker)。每个执行线程可以执行具有一定向量宽度的向量运行。每个worker里的并行以及一个跨向量操作的并行被称为向量并行。当在GPU上使用OpenACC时，一个PU就类似于一个SM。
  根据任务是否通过gang、worker、vector并行执行，OpenACC执行被分成几种模式。现在，假设在一个OpenACC程序的并行计算区域中，创建了G个gang，其中每个gang包含W个worker，每个worker的向量宽度为V。那么，总共有G×W×V个执行线程处理这个并行区域。
  当开始执行并行区域时，gang以gang冗余模式执行，这有利于在并行执行开始前对gang的状态进行初始化。在gang冗余模式中，每个gang的worker中只有一个活跃vector元素，其他worker和vector元素是闲置的，因此只有G个活跃的执行线程。此外，每个gang都执行相同的运算，所以在这个阶段没有通过gang并行任务。在CUDA中，gang冗余并行将作为执行相同计算的线程块里的一个线程来实现:

__global__ void kernel(...){
   
   
   if(threadIdx.x == 0){
   
   
      foo();
   }
}

  在OpenACC并行区域的某些地方，应用程序可能通过gang转换为并行执行。在这种情况下，程序以gang分裂模式执行。在分裂模式下，每个gang中仍然只有一个活跃地vector元素和一个活动的worker，但每个活跃的vector元素执行不同的并行区域。因此，该计算区域的任务被分散到各个gang中。在CUDA中，gang分裂模式将作为一个线程来实现，在这个线程里每个线程块处理分离的数据点。对于向量加法，在gang分裂模式下执行的CUDA内核如下所示：

__global__ void kernel(int *in1, int *in2, int *out, int N){
   
   
   if(threadIdx.x == 0){
   
   
      int i;
      for(i = blockIdx.x; i < N; i += gridDim.x){
   
   
         out[i] = in1[i] + in2[i];
   }
}

  那么对于worker并行和vector并行呢？当在一个gang中只有一个活跃的worker时，程序处于单一worker模式。当worker中只有一个活跃的vector元素时，程序处于单一vector模式。因此，gang冗余模式和gang分裂模式也可以被称为单一worker模式和单一vector模式。
  在worker分裂模式下，并行区域的工作被划分到多个gang和多个worker中。使用所有gang里的所有worker可以提供G×W路并行。在CUDA中，worker分裂模式通过每个线程束中的第一个线程来实现：

__global__ void kernel(int *in1, int *in2, int *out, int N){
   
   
   if(threadIdx.x % warpSize == 0){
   
   
      int warpId = threadIdx.x / warpSize;
      int warpsPerBlock = blockDim.x / warpSize;
      int i;
      for(i = blockIdx.x * warpsPerBlock + warpId; i < N; i += gridDim.x * warpsPerBlock){
   
   
         out[i] = in1[i] + in2[i];
   }
}

  在vector分裂模式下，工作任务在gang、worker和vector通道上进行划分，同时提供G×W×V路并行。这个模式与编写的CUDA内核模式最为相似。这些不同的OpenACC模式，使得一个应用程序的并行性可在代码的并行区域内进行动态调整。
  当使用OpenACC时，由程序员用编译器指令指定并行区域，或是并行运行。编译器指令还可以指定使用何种类型的并行处理。编译器指令是一行源代码。用C/C++编写，开头为#pragma。OpenACC指令使用acc关键字作为唯一标识，这意味着所有OpenACC指令都是以#pragma acc开头的。
  尽管编译器指令是程序源代码的一部分，但它们对编译器生成的可执行文件的影响是不一定的。如果编译器无法识别或不支持#pragma的这种类型，那么在编译时就会忽略#pragma的存在。目前，PGI、Cray和CAPS编译器都支持OpenACC指令。
  使用OpenACC编译器能对程序的OpenACC指令做出解释，也能对源代码进行自动分析以自动生成加速器源代码。因此，只需要添加几行指令代码，就可以在GPU上自动执行应用程序了。如下是OpenACC向量加法的实现：

#pragma acc kernels
for(int i = 0; i < N; i++){
   
   
   C[i] = A[i] + B[i];