CUDA 学习笔记（六）_openacc-优快云博客

前言

本小节介绍的是OpenACC，原书中写到的介绍是：OpenACC is a complementary programming model to CUDA that uses a compiler directives-based API designed for high performance, programmability, and portability across many platforms. 通篇读下来后，博主认为该工具主要的作用在于使用简单的语法实现CUDA底层调用gpu硬件的加速，可以理解为对CUDA编程的封装（？

想要具体学习OpenACC可以去官网查看他的文档及API说明：https://www.openacc.org/resources，博主仅针对原书的介绍进行简单的总结

所有涉及到的代码可从博主的代码仓库中查看：https://github.com/palLiHua/cuda

第八章 CUDA相关库及OpenACC

8.4 OpenACC介绍

概述

OpenACC的线程模型包括gang,worker,vector，类比于CUDA编程中的blocks,warp,thread.与CUDA最大的区别是OpenACC显示暴露了warp，而在CUDA编程中代码逻辑层面是看不到warp的。博主这里说明得可能不够清晰，但学过前面前三章的童鞋应该能明白我在说什么，原书中的介绍是：OpenACC exposes the concept of workers (that is, warps) directly in the programming model, whereas in CUDA you do not explicitly create warps of threads.

OpenACC执行的模式主要有三种，假设有G个gangs,W个workers,每个vector有V个elements，三种模式及其运行情况如下：

1. gang-redundant mode

这种模式下每一个gang中只有一个worker工作，每一个worker又只有一个element工作，即共有G个线程在同时运行。类比于CUDA编程的话，就是创建了G个block，但每个block只有一个thread工作，等价于如下代码：

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

2. gang-partitioned mode

这种模式下线程工作模式同gang-redundant相同，即同时有G个线程并行运行，那就有疑问了，和gang-redundant有什么区别呢？原文解释是：In gang-partitioned mode, there is still only one active vector element and one active worker in each gang, but each active vector element is executing a different piece of the parallel region.

给出的代码示例是：

__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]; 
        }
    } 
}

查了一下资料以及问询GPT，个人认为和gang-redundant的区别体现在gang-redundant是所有线程做着同样的工作及处理同样的数据，换句话说就是在重复性工作，没有任何意义（可能？），而gang-partitioned则是每个gang在处理着不同的数据。

3. worker-single mode

在该模式下又可以分为worker-partitioned mode和vector-partitioned mode。有了上面的了解，不难理解worker-partitioned mode就是每个gang下所有worker都开始并行运行了，因此总共会有G*W个线程同时工作，等价于：

__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-partitioned mode下就会有G*W*V个线程同时执行。

使用OpenACC计算指令

主要可以使用#pragma acc kernels和#pragma acc parallel两种方式以启用gpu加速。

#pragma acc kernels

使用该指令通常会使coder更轻松，包括数据传输、并行程度均有编译器自行决定，但代价自然就是性能的下降。具体的使用方法可以参考acc.c，相较于CUDA需要自己手写核函数，OpenACC仅需要一行代码就可以实现gpu加速，非常便捷。使用的编译器是PGI，博主下载的是英伟达提供的HPC，可从https://developer.nvidia.com/hpc-sdk处下载，使用指令pgcc -acc -Minfo=accel acc.c -o test进行编译，其中-Minfo=accel可以用于查看编译的一些细节，顺利的话你会得到类似于如下的输出：

从红框圈出部分不难看出，两处循环使用了gpu进行加速，前两行代码也不难猜出，要在gpu上运行的话，它自动帮我们完成了数据host to device以及device to host传输的工作。

使用这种方式还可以加上if条件语句以及async(id)用于异步运行，具体用法参考官方文档及例子。需要值得注意的一点是，if条件不满足情况下，使用pgcc -acc -Minfo=accel编译仍会有上图中的输出，但具体有没有使用gpu加速其实是在运行过程中才决定的。

#pragma acc parallel

使用parallel的好处是你可以决定程序运行的并发程度，例如使用num_gangs(int)指定gang的数量，同样也可以加入if、async、wait等指令，与kernels最大的区别就是你能控制的部分更多了。在执行循环时需要加上#pragma acc loop，默认情况下是gang-redundant mode，如果要增加并发度，那么就要再加上gang用以表明使用gang-partitioned mode，代码如下：

#pragma acc parallel {
    int b = a + c; 
#pragma acc loop gang
    for (i = 0; i < N; i++) 
    { 
        ...
    } 
}

那么同理加上worker/vector就能切换至worker-partitioned mode或vector-partitioned mode

使用OpenACC数据指令

上面介绍的过程中，我们都默认数据传输是自动进行的，该小节介绍的就是我们如何手工控制数据的传输。有两种方式，第一种是使用data copyin/ data copyout，具体如何使用可参考data.c，这里仅贴出核心代码

#pragma acc data copyin(A[0:N], B[0:N]) copyout(C[0:N], D[0:N])
    {
#pragma acc parallel
        {
#pragma acc loop
            for (i = 0; i < N; i++)
            {
                C[i] = A[i] + B[i];
            }

#pragma acc loop
            for (i = 0; i < N; i++)
            {
                D[i] = C[i] * A[i];
            }
        }
    }

这里要提到的一点是，为写代码方便，上述代码等价于：

#pragma acc parallel copyin(A[0:N], B[0:N]) copyout(C[0:N], D[0:N])
    {
#pragma acc loop
        for (i = 0; i < N; i++)
        {
            C[i] = A[i] + B[i];
        }

#pragma acc loop
        for (i = 0; i < N; i++)
        {
            D[i] = C[i] * A[i];
        }
    }

这种方式存在的一个问题是，数据传输会block程序的进行，即不能异步进行，那么这时候就需要第二种方式闪亮登场。

在第二种方式下，仅需要把指令改为#pragma acc enter data和#pragma acc exit data，如果需要异步进行，则仅需要加上async(int)，可以参考原书中的这段代码：

#pragma acc enter data copyin(B[0:N]) async(0)

do_some_heavy_work(C);

#pragma acc kernels async(1) wait(0)
{
    for (i = 0; i < N; i++) 
    {
        A[i] = do_work(B[i]);
    }
}

#pragma acc exit data copyout(A[0:N]) async(2) wait(1)

#pragma wait(2)

OpenACC Runtime API

除了上述介绍的compiler directives，OpenACC还提供了一系列函数指令，包含在openacc.h头文件中，这里不作解释，可自行查看文档

OpenACC和CUDA一同使用

CUDA编程的优势在于能自主控制的部分更多，而OpenACC自然在于其编程的便利性，这里就以一个例子来介绍下如何同时使用这两个工具，代码可参考cuda_openacc.cu，这个例子中使用curand产生随机的两个矩阵，然后使用3个for循环进行矩阵相乘操作（使用OpenACC），最后将生成的矩阵进行每个元素加和的操作（使用cublas库的cublasSasum先进行每行求和，最后再将vector加和）

在编译的时候，原书给出的指令是pgcpp -acc cuda-openacc.cu -o cuda-openacc -Minfo=accel -L${CUDA_HOME}/lib64 -lcurand -lcublas –lcudart，但博主在执行过程中发现根本没有pgcpp，查看了一下英伟达提供的指令：