OpenACC 入门指南

1. 概述

OpenACC 应用程序编程接口 (API) 是一组编译器指令和运行时例程，允许软件开发人员在标准 Fortran、C++ 和 C 程序中指定循环和代码区域，这些循环和代码区域应通过卸载到加速器（如 GPU）或通过在主机 CPU 的所有核心上执行来并行执行。OpenACC API 由行业联盟设计和维护。有关 OpenACC API 的更多信息，请参阅 OpenACC 网站。

本入门指南提供了关于如何使用 NVIDIA HPC 编译器中的 OpenACC 指令支持来编写、构建和运行程序的示例。

1.1. 系统先决条件

在 NVIDIA HPC 编译器中使用 OpenACC 实现需要以下条件：

运行 Linux 的基于 x86-64 或 Arm 服务器 CPU 的系统。

对于面向 GPU，需要支持 CUDA 的 NVIDIA GPU 和已安装的 CUDA 设备驱动程序，CUDA 版本 8.0 或更高版本。有关获取和安装 CUDA 设备驱动程序的更多信息，请参阅 NVIDIA CUDA 网页。

1.2. 准备您的系统

要使用 NVIDIA OpenACC 编译器，请按照以下步骤操作：

从 NVIDIA 下载网页下载最新的 NVIDIA HPC SDK 25.1 软件包。
将下载的软件包安装到 /opt/nvidia/hpc_sdk 或您选择的其他目录中。
将 /opt/nvidia/hpc_sdk/target/25.1/compilers/bin 目录添加到您的路径中，其中 target 是 Linux_x86_64 或 Linux_aarch64 之一。

调用 nvaccelinfo 命令以查看您的 GPU 和驱动程序是否已正确安装且可用。您应该看到类似以下内容的输出：

CUDA Driver Version:           10020
NVRM version:                  NVIDIA UNIX x86_64 Kernel Module  440.33.01  Wed Nov 13 00:00:22 UTC 2019

Device Number:                 0
Device Name:                   Tesla V100-PCIE-16GB
Device Revision Number:        7.0
Global Memory Size:            16945512448
Number of Multiprocessors:     80
Concurrent Copy and Execution: Yes
Total Constant Memory:         65536
Total Shared Memory per Block: 49152
Registers per Block:           65536
Warp Size:                     32
Maximum Threads per Block:     1024
Maximum Block Dimensions:      1024, 1024, 64
Maximum Grid Dimensions:       2147483647 x 65535 x 65535
Maximum Memory Pitch:          2147483647B
Texture Alignment:             512B
Clock Rate:                    1380 MHz
Execution Timeout:             No
Integrated Device:             No
Can Map Host Memory:           Yes
Compute Mode:                  default
Concurrent Kernels:            Yes
ECC Enabled:                   Yes
Memory Clock Rate:             877 MHz
Memory Bus Width:              4096 bits
L2 Cache Size:                 6291456 bytes
Max Threads Per SMP:           2048
Async Engines:                 7
Unified Addressing:            Yes
Managed Memory:                Yes
Concurrent Managed Memory:     Yes
Preemption Supported:          Yes
Cooperative Launch:            Yes
  Multi-Device:                Yes
NVIDIA Default Target:         cc70

这会告诉您驱动程序版本、GPU 类型（如果您有多个 GPU，则为多个 GPU）、可用内存、您应该使用的命令行标志以面向此 GPU（在本例中为 -gpu=cc70）等等。

1.3. 支持文档和示例

您可能需要查阅 OpenACC 2.7 规范，该规范可在 OpenACC 网站上找到。简单示例出现在将 OpenACC 与 HPC 编译器结合使用中。

2. 将 OpenACC 与 NVIDIA HPC 编译器结合使用

通过将 -⁠acc 标志添加到编译器命令行来启用 OpenACC 指令。默认情况下，NVIDIA HPC 编译器会将 OpenACC 区域并行化并卸载到 NVIDIA GPU。您可以指定 -⁠acc=multicore 以针对多核 CPU 进行并行化，或指定 -⁠acc=host 以生成将在主机 CPU 上串行运行的可执行文件。

GPU 定向和代码生成的许多方面可以通过将 -⁠gpu 标志添加到编译器命令行来控制。默认情况下，NVIDIA HPC 编译器将面向编译主机上安装的 NVIDIA GPU。您可以指定 -⁠gpu=cc70 以专门面向 Volta GPU，或指定多个计算能力 (ccXY) 以生成针对多代 NVIDIA GPU 优化的 GPU 可执行文件。

有关 -⁠acc 和 -⁠gpu 编译器选项的完整子选项列表，请参阅编译器手册页。此版本包括对几乎所有 OpenACC 2.7 规范的支持。有关此版本中支持的功能的详细信息，请参阅已实现的功能。

2.1. CUDA 版本

NVIDIA HPC 编译器使用 NVIDIA CUDA 工具包中的组件来构建在 NVIDIA GPU 上执行的程序。NVIDIA HPC SDK 将 CUDA 工具包组件放入 HPC SDK 安装子目录中；HPC SDK 当前捆绑了两个最近发布的工具包版本。

您可以在 HPC 编译器支持的任何系统上编译用于 NVIDIA GPU 的程序。您只能在具有 NVIDIA GPU 和已安装 NVIDIA CUDA 驱动程序的系统上运行该程序。NVIDIA HPC SDK 产品不包含 CUDA 设备驱动程序。您必须从 NVIDIA 下载并安装适当的 CUDA 驱动程序。

NVIDIA HPC SDK 实用程序 nvaccelinfo 将驱动程序版本作为其输出的第一行打印。您可以使用它来查找系统上安装的 CUDA 驱动程序的版本。

NVIDIA HPC SDK 25.1 包括来自以下 CUDA 工具包版本的组件：

CUDA 11.8

CUDA 12.4

如果您正在没有安装 CUDA 驱动程序的系统上编译用于 GPU 执行的程序，则编译器会根据 localrc 文件中包含的 DEFCUDAVERSION 变量的值来选择要使用的 CUDA 工具包版本，该文件是在 HPC SDK 安装期间创建的。

如果您正在有安装 CUDA 驱动程序的系统上编译用于 GPU 执行的程序，则编译器会检测 CUDA 驱动程序的版本，并从 HPC SDK 捆绑的工具包中选择适当的 CUDA 工具包版本。

编译器在 /opt/nvidia/hpc_sdk/target/25.1/cuda 目录中查找与系统上安装的 CUDA 驱动程序版本匹配的 CUDA 工具包版本。如果找不到完全匹配的版本，编译器将搜索最接近的匹配项。对于 CUDA 驱动程序版本 11.2 到 11.8，编译器将使用 CUDA 11.8 工具包。对于 CUDA 驱动程序版本 12.0 及更高版本，编译器将使用最新的 CUDA 12.x 工具包。

您可以使用编译器选项更改编译器 CUDA 工具包版本的默认选择。将 cudaX.Y 子选项添加到 -⁠gpu，其中 X.Y 表示 CUDA 版本。使用编译器选项会更改编译器一次调用的 CUDA 工具包版本。例如，要使用 CUDA 11.8 工具包编译 OpenACC C 文件，您将使用：

nvc -acc -gpu=cuda11.8

2.2. 计算能力

编译器可以为 NVIDIA GPU 计算能力 3.5 到 8.6 生成代码。编译器构建与编译中使用的系统上找到的 GPU 支持的计算能力匹配的默认计算能力列表。如果未检测到 GPU，则编译器会为每个受支持的计算能力生成代码。

您可以使用命令行选项或 rcfile 覆盖默认值。

要使用命令行选项更改默认值，请为 -⁠gpu 选项提供以逗号分隔的计算能力列表。

要使用 rcfile 更改默认值，请将 DEFCOMPUTECAP 值设置为安装目录的 bin 目录中的 siterc 文件中以空格分隔的计算能力列表：

set DEFCOMPUTECAP=60 70;

或者，如果您没有更改 siterc 文件的权限，您可以将 DEFCOMPUTECAP 定义添加到主目录中的单独的 .mynvrc 文件中。

设备代码的生成可能很耗时，因此您可能会注意到，随着计算能力数量的增加，编译时间也会增加。

2.3. PTX JIT 编译

从 HPC SDK 22.9 开始，所有编译器在可重定位设备代码模式下都启用了对 PTX JIT 编译的支持。这意味着使用 -gpu=rdc 构建的应用程序（即，启用了可重定位设备代码，这是默认模式）由于嵌入的 PTX 代码而向前兼容更新的 GPU。当应用程序在比编译时指定的架构更新的 GPU 架构上运行时，嵌入的 PTX 代码会被动态编译。

对 PTX JIT 编译的支持是自动启用的，这意味着您无需更改现有项目的编译器调用命令行。

使用场景

例如，您可以编译面向 Ampere GPU 的应用程序，而无需担心 Hopper GPU 架构。一旦应用程序在 Hopper GPU 上运行，它将无缝地使用嵌入的 PTX 代码。
在 CUDA Fortran 中，或者在启用 CUDA 互操作模式的情况下，您可以混合使用使用 CUDA NVCC 编译器编译的对象文件，其中包含 PTX 代码。来自 NVCC 的 PTX 代码将与 HPC SDK 编译器生成的对象文件中包含的 PTX 代码一起由 JIT 编译器处理。使用 CUDA NVCC 编译器时，必须使用 NVCC --relocatable-device-code true 开关显式启用可重定位设备代码生成，如 CUDA 编译器驱动程序指南中所述。有关 CUDA 互操作性的信息，请参阅 https://docs.nvda.net.cn/hpc-sdk/compilers/hpc-compilers-user-guide/index.html#openmp-interop-cuda。CUDA Fortran 编程指南可在此处找到：https://docs.nvda.net.cn/hpc-sdk/compilers/cuda-fortran-prog-guide。

默认情况下，编译器将选择与代码正在编译的系统上的 GPU 匹配的计算能力。对于将在编译代码的系统上运行的代码，我们建议让编译器设置计算能力。

当默认值不起作用时，我们建议为应用程序预期运行的计算能力范围编译应用程序，例如，使用 -gpu=ccall 编译器选项。当在支持这些计算能力之一的系统上运行应用程序时，CUDA 驱动程序次要版本允许低于编译时使用的 CUDA 工具包版本，如 CUDA 版本部分所述。

性能考虑因素

PTX JIT 编译在发生时，可能会对应用程序产生启动开销。JIT 编译器保留生成的设备代码的缓存副本，这减少了后续运行的开销。有关 JIT 编译器如何工作的详细信息，请参阅 CUDA 编程指南。

已知限制

一般来说，为了使 PTX JIT 编译工作，部署系统上安装的 CUDA 驱动程序必须至少是与用于编译应用程序的 CUDA 工具包匹配的版本。此要求比 CUDA 版本部分中解释的要求更严格。

例如，如该部分中所述，当系统中安装的 CUDA 驱动程序至少为 11.2 时，编译器将使用作为 HPC SDK 工具包一部分提供的 CUDA 11.8 工具包。但是，虽然 CUDA 11.2 驱动程序通常足以运行应用程序，但它将无法编译 CUDA 11.8 工具包生成的 PTX 代码。这意味着任何预期使用 PTX JIT 编译的部署系统都必须至少安装 CUDA 11.8 驱动程序。有关 CUDA 驱动程序与 CUDA 工具包的兼容性的更多信息，请参阅 CUDA 兼容性指南。

当应用程序预计在比编译时指定的更新的 GPU 架构上运行时，我们建议在部署系统上安装与用于构建应用程序的 CUDA 工具包匹配的 CUDA 驱动程序。实现此目的的一种方法是在编译时使用 NVHPC_CUDA_HOME 环境变量来提供特定的 CUDA 工具包。

以下是一些关于如何诊断和修复 PTX 版本不兼容性的示例。作为一般规则，如果 CUDA 驱动程序由于 PTX 不兼容而无法运行应用程序，则应用程序将终止并显示指示原因的错误消息。在大多数情况下，OpenACC 和 OpenMP 应用程序将建议编译器标志以面向当前的 CUDA 安装。

OpenACC

考虑这个程序，我们将为 Volta GPU 编译它，并尝试在 Ampere GPU 上运行，系统上安装了 CUDA 11.5

#include <stdio.h>
#define N 1000
int array[N];
int main() {
#pragma acc parallel loop copy(array[0:N])
   for(int i = 0; i < N; i++) {
      array[i] = 3.0;
   }
   printf("Success!\n");
}

当我们构建程序时，HPC SDK 将选择包含在默认设置中的 CUDA 11.8 工具包。当我们尝试运行它时，它会失败，因为使用 11.8 生成的代码不适用于 11.5 驱动程序

$ nvc -acc -gpu=cc70 app.c
$ ./a.out
Accelerator Fatal Error: This file was compiled: -acc=gpu -gpu=cc70
Rebuild this file with -gpu=cc80 to use NVIDIA Tesla GPU 0
 File: /tmp/app.c
 Function: main:3
 Line: 3

从错误消息中可以看出，系统无法在当前系统上执行 Volta GPU 指令。嵌入的 Volta PTX 无法编译，这意味着 CUDA 驱动程序不兼容。修复此问题的一种方法是在编译时使用已安装的 CUDA 11.5 工具包

$ export NVHPC_CUDA_HOME=/usr/local/cuda-11.5
$ nvc -acc -gpu=cc70 app.c
$ ./a.out
Success!

OpenMP

同样，OpenMP 程序将编译但无法运行

#include <stdio.h>
#define N 1000
int array[N];
int main() {
#pragma omp target loop
   for(int i = 0; i < N; i++) {
      array[i] = 0;

   }
   printf("Success!\n");
}

$ nvc -mp=gpu -gpu=cc70 app.c
$ ./a.out
Accelerator Fatal Error: Failed to find device function 'nvkernel_main_F1L3_2'! File was compiled with: -gpu=cc70
Rebuild this file with -gpu=cc80 to use NVIDIA Tesla GPU 0
 File: /tmp/app.c
 Function: main:3
 Line: 3

我们还可以通过使 NVHPC_CUDA_HOME 指向匹配的 CUDA 工具包位置来修复它

$ export NVHPC_CUDA_HOME=/usr/local/cuda-11.5
$ nvc -acc -gpu=cc70 app.c
$ ./a.out
Success!

C++

与 OpenACC 和 OpenMP 应用程序在 PTX JIT 遇到不充分的 CUDA 驱动程序版本时简单地终止相反，当存在 PTX 不兼容性时，C++ 应用程序会抛出系统异常

#include <vector>
#include <algorithm>
#include <execution>
#include <iostream>
#include <assert.h>
int main() {
  std::vector<int> x(1000, 0);
  x[1] = -20;
  auto result = std::count(std::execution::par, x.begin(), x.end(), -20);
  assert(result == 1);
  std::cout << "Success!" << std::endl;
}

$ nvc++ -stdpar -gpu=cc70 app.cpp
$ ./a.out
terminate called after throwing an instance of 'thrust::system::system_error'
  what():  after reduction step 1: cudaErrorUnsupportedPtxVersion: the provided PTX was compiled with an unsupported toolchain.
Aborted (core dumped)

异常消息包含对不兼容 PTX 的直接引用，这反过来意味着 CUDA 工具包和 CUDA 驱动程序版本之间不匹配。

我们可以通过设置 NVHPC_CUDA_HOME 类似地修复它

$ export NVHPC_CUDA_HOME=/usr/local/cuda-11.5
$ nvc++ -stdpar -gpu=cc70 app.cpp
$ ./a.out
Success!

2.4. OpenACC 中的 C 结构体

NVIDIA C++ 和 C 编译器支持在 OpenACC 区域中使用结构体的静态数组和指向结构体动态数组的指针。

typedef struct{
    float x, y, z;
 }point;

extern point base[1000];

void vecaddgpu( point *restrict r, int n ){
    #pragma acc parallel loop present(base) copyout(r[0:n])
    for( int i = 0; i < n; ++i ){
        r[i].x = base[i].x;
        r[i].y = sqrtf( base[i].y*base[i].y + base[i].z*base[i].z );
        r[i].z = 0;
    }
}

指向标量结构体的指针被视为单元素数组，应塑形为 r[0:1]。

typedef struct{
    base[1000];
    int n;
    float *x, *y, *z;
 }point;

extern point A;

void vecaddgpu(){
    #pragma acc parallel loop copyin(A) \
        copyout(A.x[0:A.n], A.y[0:A.n], A.z[0:A.n])
    for( int i = 0; i < A.n; ++i ){
        A.x[i] = A.base[i];
        A.y[i] = sqrtf( A.base[i] );
        A.z[i] = 0;
    }
}

在此示例中，结构体 A 被复制到设备，这会复制静态数组成员 A.base 和标量 A.n。然后，动态成员 A.x、A.y 和 A.z 被复制到设备。结构体 A 应在其动态成员之前复制，方法是将结构体放在较早的数据子句中，或者在封闭的数据区域或动态数据生命周期中在设备上复制或创建它。如果动态成员被复制时结构体在设备上不存在，则对设备上动态成员（例如 A.x[i]）的访问将无效，因为指针 A.x 将不会更新。

指向结构体的指针被视为单元素数组。如果结构体还包含指针成员，则应将结构体复制到设备，然后创建或复制指针成员

typedef struct{
    int n;
    float *x, *y, *z;
 }point;

void vecaddgpu( point *A, float* base ){
    #pragma acc parallel loop copyin(A[0:1]) \
        copyout(A->x[0:A->n], A->y[0:A->n], A->z[0:A->n]) \
        present(base[0:A->n])
    for( int i = 0; i < A->n; ++i ){
        A->x[i] = base[i];
        A->y[i] = sqrtf( base[i] );
        A->z[i] = 0;
    }
}

将包含指针的结构体复制回主机时要小心。在设备上，指针成员将使用设备指针进行更新。如果这些指针被复制回主机结构体，则指针在主机上将无效。

在设备上创建或复制结构体时，将分配整个结构体。不支持分配结构体的子集，或者仅为单个成员分配空间。

结构体和指针成员也可以使用动态数据指令进行管理

typedef struct{
    int n;
    float *x, *y, *z;
 }point;

void move_to_device( point *A ){
    #pragma acc enter data copyin(A[0:1])
    #pragma acc enter data create(A->x[0:A->n], A->y[0:A->n], A->z[0:A->n])
}

void move_from_device( point* A ){
    #pragma acc enter data copyout(A->x[0:A->n], A->y[0:A->n], A->z[0:A->n])
    #pragma acc enter data delete(A[0:1])
}

void vecaddgpu( point *A, float* base ){
    #pragma acc parallel loop present(A[0:1]) \
        present(A->x[0:A->n], A->y[0:A->n], A->z[0:A->n]) \
        present(base[0:A->n])
    for( int i = 0; i < A->n; ++i ){
        A->x[i] = base[i];
        A->y[i] = sqrtf( base[i] );
        A->z[i] = 0;
    }
}

2.5. OpenACC 中的 C++ 类

NVIDIA C++ 编译器支持在 OpenACC 区域中使用 C++ 类，包括静态数组类成员、指向动态数组的成员指针以及成员函数和运算符。通常，类本身也必须复制到设备内存，方法是将类变量放在类外部的数据子句中，或者在类内部的数据子句中使用适当塑形的 this[0:1] 引用。整个类将在设备内存中分配。

// my managed vector datatype
template<typename elemtype> class myvector{
    elemtype* data;
    size_t size;
public:
    myvector( size_t size_ ){ // constructor
        size = size_;
        data = new elemtype[size];
    }
    todev(){ // move to device
        #pragma acc enter data copyin(this[0:1], data[0:size])
    }
    fromdev(){ // remove from device
        #pragma acc exit data delete( data[0:size], this[0:1])
    }
    void updatehost(){ // update host copy of data
        #pragma acc update self( data[0:size] )
    }
    void updatedev(){ // update device copy of data
        #pragma acc update device( data[0:size] )
    }
    ~myvector(){ // destructor from host
        delete[] data;
    }
    inline elemtype & operator[] (int i) const { return data[i]; }
    // other member functions
};

在下面的示例中，this 指针在 data 之前复制到设备，因此设备上指向 data 的指针将得到更新。这称为“附加”操作；class myvector 指针 data 附加到 data 向量的设备副本。

另一个类始终创建设备数据以及主机数据

// my managed host+device vector datatype
template<typename elemtype> class hdvector{
    elemtype* data;
    size_t size;
public:
    hdvector( size_t size_ ){ // constructor
        size = size_;
        data = new elemtype[size];
        #pragma acc enter data copyin(this[0:1]) create(data[0:size])
    }
    void updatehost(){ // update host copy of data
        #pragma acc update self( data[0:size] )
    }
    void updatedev(){ // update device copy of data
        #pragma acc update device( data[0:size] )
    }
    ~hdvector(){ // destructor from host
        #pragma acc exit data delete( data[0:size], this[0:1] )
        delete[] data;
    }
    inline elemtype & operator[] (int i) const { return data[i]; }
    // other member functions
};

构造函数复制类，因此将复制 size 值，并创建（分配）data 向量。

一个稍微复杂的类包括一个复制构造函数，该构造函数复制数据指针而不是数据副本

#include <openacc.h>
// my managed vector datatype
template<typename elemtype> class dupvector{
    elemtype* data;
    size_t size;
    bool iscopy;
public:
    dupvector( size_t size_ ){ // constructor
        size = size_;
        data = new elemtype[size];
        iscopy = false;
        #pragma acc enter data copyin(this[0:1]) create(data[0:size])
    }
    dupvector( const dupvector &copyof ){ // copy constructor
        size = copyof.size;
        data = copyof.data;
        iscopy = true;
        #pragma acc enter data copyin(this[0:1])
        acc_attach( (void**)&data );
    }
    void updatehost(){ // update host copy of data
        #pragma acc update self( data[0:size] )
    }
    void updatedev(){ // update device copy of data
        #pragma acc update device( data[0:size] )
    }
    ~dupvector(){ // destructor from host
        if( !iscopy ){
            #pragma acc exit data delete( data[0:size] )
            delete[] data;
        }
        #pragma acc exit data delete( this[0:1] )
    }
    inline elemtype & operator[] (int i) const { return data[i]; }
    // other member functions
};

请注意在复制构造函数中调用 OpenACC 运行时例程 acc_attach。此例程采用指针的地址，转换该指针的地址以及指针的内容，并将转换后的内容存储到设备上的转换后的地址中。在本例中，它将从设备上的原始类复制的数据指针附加到设备上此类副本。

在类外部的代码中，可以按预期在计算子句中引用数据

dupvector<float> v = new dupvector<float>(n);
dupvector<float> x = new dupvector<float>(n);
...
#pragma acc parallel loop present(v,x)
 for( int i = 0; i < n; ++i ) v[i] += x[i];

上面的示例显示了并行循环构造中对 v 和 x 类的引用。operator[] 通常会被内联。如果它未内联或禁用内联，则编译器将注意到该运算符是从 OpenACC 计算区域内调用的，并编译该运算符的设备版本。这实际上与在运算符上方暗示 #pragma acc routine seq 相同。对于 C++ 中的任何函数（无论是类成员函数还是独立函数）也是如此：如果该函数是从计算区域内调用的，或者从计算区域内调用的函数调用的，并且没有 #pragma acc routine，则编译器会将其视为以 #pragma acc routine seq 为前缀。当您编译文件并启用 -Minfo=accel 时，您将看到以下消息：

T1 &dupvector<T1>::operator [](int) const [with T1=float]:
          35, Generating implicit acc routine seq

在上面的示例中，循环上限是简单的标量变量 n，而不是更自然的类成员 v.size。在 NVIDIA C++ 编译器的当前实现中，并行循环或内核循环的循环上限必须是简单变量，而不是类成员。

类变量出现在并行构造的 present 子句中。计算构造的正常默认值是编译器将对类的引用视为 present_or_copy。但是，如果类实例不存在，则仅复制类本身不会复制动态数据成员，因此不会提供必要的行为。因此，当在计算构造中引用类对象时，应将类放在 present 子句中。

可以在并行循环中显式调用类成员函数

template<typename elemtype> class dupvector{
    ...
    void inc1( int i, elemtype y ){
        data[i] += y;
    }
}
...
#pragma acc parallel loop present(v,x)
    for( int i = 0; i < n; ++i ) v.inc1( i, x[i] );

如上所述，当启用优化时，编译器通常会内联 inc1，但由于它是从计算区域内调用的，因此也会编译该函数的设备版本。

计算构造本身可以包含计算构造

template<typename elemtype> class dupvector{
    ...
    void inc2( dupvector<elemtype> &y ){
        int n = size;
        #pragma acc parallel loop gang vector present(this,y)
        for( int i = 0; i < n; ++i ) data[i] += y[i];
    }
}
...
    v.inc2( x );

再次注意循环上限 n，以及 present 子句中的 this 和 y 类。第三个示例将并行构造放在例程周围，但循环本身在例程内。正确执行此操作需要您在例程定义之前放置适当的 acc routine，以在正确的并行级别调用例程。

template<typename elemtype> class dupvector{
    ...
    #pragma acc routine gang
    void inc3( dupvector<elemtype> &y ){
        int n = size;
        #pragma acc loop gang vector
        for( int i = 0; i < n; ++i ) data[i] += y[i];
    }
}
...
    #pragma acc parallel
        v.inc3( x );

当从主机代码调用 inc3 时，它将在主机上运行，从而递增主机值。当从 OpenACC 并行构造中调用时，它将递增设备值。

2.6. OpenACC 中的 Fortran 派生类型

NVIDIA Fortran 编译器支持在 OpenACC 区域中使用派生类型的静态和可分配数组。

module mpoint
type point
    real :: x, y, z
end type
type(point) :: base(1000)
end module

subroutine vecaddgpu( r, n )
 use mpoint
 type(point) :: r(:)
 integer :: n
 !$acc parallel loop present(base) copyout(r(:))
 do i = 1, n
  r(i)%x = base(i)%x
  r(i)%y = sqrt( base(i)%y*base(i)%y + base(i)%z*base(i)%z )
  r(i)%z = 0
 enddo
end subroutine

您可以显式引用派生类型的数组成员，包括派生类型内的静态数组和可分配数组。在任何一种情况下，都必须将整个派生类型放置在设备内存中，方法是将派生类型本身放在适当的数据子句中。在当前的实现中，派生类型变量本身必须出现在数据子句中，至少是 present 子句，对于任何直接使用派生类型变量的计算构造。

module mpoint
type point
    real :: base(1000)
    integer :: n
    real, allocatable, dimension(:) :: x, y, z
end type

type(point) :: A
end module

subroutine vecaddgpu()
 integer :: i
 !$acc parallel loop copyin(A) copyout(A%x,A%y,A%z)
 do i = 1, n
  A%x(i) = A%base(i)
  A%y(i) = sqrt( A%base(i) )
  A%z(i) = 0
 enddo
end subroutine

在此示例中，派生类型 A 被复制到设备，这会复制静态数组成员 A%base 和标量 A%n。然后，可分配数组成员 A%x、A%y 和 A%z 被复制到设备。派生类型变量 A 应在其可分配数组成员之前复制，方法是将派生类型放在较早的数据子句中，或者在封闭的数据区域或动态数据生命周期中在设备上复制或创建它。如果可分配数组成员被复制时派生类型在设备上不存在，则对设备上可分配成员（例如 A%x(i)）的访问将无效，因为派生类型变量中的隐藏指针和描述符值将不会更新。

将包含可分配成员的派生类型复制回主机时要小心。在设备上，可分配成员将更新为指向设备内存。如果整个派生类型被复制回主机，则可分配成员在主机上将无效。

在设备上创建或复制派生类型时，将分配整个派生类型。不支持分配派生类型的子集，或者仅为单个成员分配空间。

派生类型和可分配成员也可以使用动态数据指令进行管理

module mpoint
 type point
    integer :: n
    real, dimension(:), allocatable :: x, y, z
 end type
contains
 subroutine move_to_device( A )
  type(point) :: A
  !$acc enter data copyin(A)
  !$acc enter data create(A%x, A%y, A%z)
 end subroutine

 subroutine move_off_device( A )
  type(point) :: A
  !$acc exit data copyout(A%x, A%y, A%z)
  !$acc exit data delete(A)
 end subroutine
end module

subroutine vecaddgpu( A, base )
 use mpoint
 type(point) :: A
 real :: base(:)
 integer :: i
 !$acc parallel loop present(A,base)
 do i = 1, n
  A%x(i) = base(i)
  A%y(i) = sqrt( base(i) )
  A%z(i) = 0
 enddo
end subroutine

2.7. Fortran I/O

NVIDIA Fortran 编译器包括对 GPU 设备代码中 PRINT 语句的有限支持。Fortran GPU 运行时库在 CUDA Fortran 和 OpenACC 之间共享，用于 NVIDIA GPU 目标，它会缓冲输出并在一次操作中打印整行。支持整数、字符、逻辑、实数和复数数据类型。

底层 CUDA printf 实现限制了内核启动中打印语句的数量为 4096 个。用户在使用此功能时应考虑此限制。

2.7.1. OpenACC PRINT 示例

这是一个在 OpenACC 计算区域内打印字符串、整数、逻辑和实型数据的简短示例

program t
integer(4) a(10000)
a = [ (1+i,i=1,10000) ]
!$acc kernels
do i = 1, 10000
  if (a(i)/3000*3000.eq.a(i)) print *," located ",i,a(i),i.gt.5000,a(i)/5.0
end do
!$acc end kernels
end

2.8. OpenACC 原子操作支持

NVIDIA OpenACC 编译器完全实现了符合 OpenACC 规范的原子操作支持。例如

double *a, *b, *c;
. . .
#pragma acc loop vector
    for (int k = 0; k < n; ++k)
    {
        #pragma acc atomic
        c[i*n+j] += a[i*n+k]*b[k*n+j];
    }

NVIDIA 编译器还包括对 CUDA 风格原子操作的支持。CUDA 原子操作名称可以在 Fortran、C 和 C++ 的加速区域中使用。例如

. . .
#pragma acc loop gang
    for (j = 0; j < n1 * n2; j += n2) {
        k = 0;
        #pragma acc loop vector reduction(+:k)
            for (i = 0; i < n2; i++)
                k = k + a[j + i];
            atomicAdd(x, k);
    }

2.9. OpenACC Declare Data 指令用于全局变量和 Fortran 模块变量

编译器支持带有 copyin、create 和 device_resident 子句的 OpenACC declare 指令，用于 C 全局变量和 Fortran 模块变量。这主要用于 OpenACC routine 指令和单独编译。declare 子句中的数据将在程序连接到设备时在设备上静态分配。copyin 子句中的数据将在此时从主机数据初始化。当程序到达其第一个数据或计算结构，或者当它调用 OpenACC acc_init 例程时，程序将连接到设备。

在 C 语言中，下面的示例使用了一个全局结构体和一个全局数组指针

struct{
    float a, b;
}coef;
float* x;
#pragma acc declare create(coef,x)
. . .
#pragma acc routine seq
void modxi( int i ){
    x[i] *= coef.a;
}
. . .
void initcoef( float a, float b ){
    coef.a = a;
    coef.b = b;
    #pragma acc update device(coef)
}
. . .
void allocx( int n ){
    x = (float*)malloc( sizeof(float)*n );
    #pragma acc enter data create(x[0:n])
}
. . .
void modx( int s, int e ){
    #pragma acc parallel loop
    for( int i = s; i < e; ++i ) modxi(i);
}

declare create(coef,x) 将在设备上静态分配结构体 coef 和指针 x 的副本。在 initcoef 例程中，系数在主机上赋值，update 指令将这些值复制到设备。allocx 例程在主机上为 x 向量分配空间，然后使用非结构化数据指令在设备上也分配该空间；因为 x 指针已经静态存在于设备上，所以 x 的设备副本也将使用指向设备数据的指针进行更新。最后，并行循环调用例程 modxi，该例程引用全局 x 指针和 coef 结构体。当在主机上调用时，此例程将访问主机上的全局 x 和 coef，而在设备上调用时（例如在此并行循环中），此例程将访问设备上的全局 x 指针和 coef 结构体。

如果 modxi 例程在单独的文件中，则 coef 和 x 的声明将具有 extern 属性，但其他代码将相同，如下所示。请注意，即使变量声明为 extern，此文件中仍然需要 acc declare create 指令，以告知编译器这些变量在设备上作为外部变量可用。

extern struct{
    float a, b;
}coef;
extern float* x;
#pragma acc declare create(coef,x)
. . .
#pragma acc routine seq
void modxi( int i ){
    x[i] *= coef.a;
}

由于全局变量存在于设备内存中，因此它也存在于 OpenACC 运行时 *present* 表中，该表跟踪主机对象和设备对象之间的对应关系。这意味着指向全局变量的指针可以作为参数传递给另一个文件中的例程，该例程在 present 子句中使用该指针。在以下示例中，调用例程使用了一个小的、静态大小的全局系数数组

float xcoef[11] = { 1.0, 2.0, 1.5, 3.5, ... 9.0 };
#pragma acc declare copyin(xcoef)
. . .
extern void test( float*, float*, float*, n );
. . .
void caller( float* x, float* y, int n ){
    #pragma acc data copy( x[0:n], y[0:n] )
    {
        . . .
        test( x, y, xcoef, n );
        . . .
    }
}

declare copyin 指令告诉编译器生成代码，以便在程序连接到设备时从主机数组初始化设备数组。在另一个文件中，定义了过程 test，并且它的所有数组参数都将已经存在于设备上；x 和 y 是因为数据结构，而 xcoef 是因为它静态地存在于设备上。

void test( float* xx, float* yy, float* cc, int n ){
    #pragma acc data present( xx[0:n], y[00:n], cc[0:11] )
    {
        . . .
        #pragma acc parallel loop
        for( int i = 5; i < n-5; ++i ){
            float t = 0.0;
            for( int j = -5; j <= 5; ++j ){
                t += cc[j+5]*yy[i+j];
            }
            xx[i] /= t;
        }
        . . .
    }
}

在 Fortran 中，模块作用域中 declare 指令中出现的模块固定大小变量和数组，以及模块可分配数组，将在 CPU 上以及设备代码中全局可用。出现在 declare create、declare copyin 或 declare device_resident 中的模块可分配数组，当它们出现在 allocate 语句中时，将在主机内存和设备内存中分配。编译器管理指向数据的实际指针以及包含每个维度数组下限和上限的描述符，并且指针的设备副本将设置为指向设备内存中的数组。

以下示例模块包含一个固定大小数组和一个可分配数组，它们都出现在 declare create 子句中。静态数组 xstat 将在加速计算区域或例程内的任何时间可用。

module staticmod
 integer, parameter :: maxl = 100000
 real, dimension(maxl) :: xstat
 real, dimension(:), allocatable :: yalloc
 !$acc declare create(xstat,yalloc)
end module

此模块可以在另一个文件中使用，该文件分配了 yalloc 数组。当可分配数组 yalloc 被分配时，它将在主机和设备内存中都进行分配，然后将在加速计算区域或例程中的任何时间可用。

subroutine allocit(n)
 use staticmod
 integer :: n
 allocate( yalloc(n) )
end subroutine

在另一个模块中，这些数组可以在计算区域或加速例程中使用

module useit
 use staticmod
contains
 subroutine computer( n )
  integer :: n
  integer :: i
  !$acc parallel loop
   do i = 1, n
    yalloc(i) = iprocess( i )
   enddo
 end subroutine
 real function iprocess( i )
  !$acc routine seq
  integer :: i
  iprocess = yalloc(i) + 2*xstat(i)
 end function
end module

2.10. OpenACC 错误处理

OpenACC 规范提供了一种机制，允许您拦截在 GPU 上执行期间触发的错误，并在程序退出前执行特定的例程作为响应。例如，如果 MPI 进程在 GPU 上分配内存时失败，应用程序可能希望调用 MPI_Abort 以在程序退出前关闭所有其他进程。本节解释了如何利用此功能。

为了拦截错误，应用程序必须向 OpenACC 运行时提供回调例程。为了提供回调，应用程序使用指向回调例程的指针调用 acc_set_error_routine。

接口如下，其中 err_msg 包含错误的描述

typedef void (*exitroutinetype)(char *err_msg);
extern void acc_set_error_routine(exitroutinetype callback_routine);

当 OpenACC 运行时检测到运行时错误时，它将调用 callback_routine。

注意

此功能与错误恢复不同。如果回调例程返回到应用程序，则行为绝对是未定义的。

让我们使用一个示例更深入地了解此功能。

以下面的 MPI 程序为例，并使用两个进程运行它。进程 0 尝试在 GPU 上分配一个大型数组，然后向第二个进程发送消息以确认操作成功。进程 1 等待确认并在收到确认后终止。

#include <stdio.h>
#include <stdlib.h>
#include "mpi.h"

#define N 2147483648

int main(int argc, char **argv)
{
  int rank, size;

  MPI_Init(&argc, &argv);

  MPI_Comm_rank(MPI_COMM_WORLD, &rank);
  MPI_Comm_size(MPI_COMM_WORLD, &size);

  int ack;
  if(rank == 0) {
    float *a = (float*) malloc(sizeof(float) * N);

#pragma acc enter data create(a[0:N])
#pragma acc parallel loop independent
    for(int i = 0; i < N; i++) {
      a[i] = i *0.5;
    }
#pragma acc exit data copyout(a[0:N])
    printf("I am process %d, I have initialized a vector of size %ld bytes on the GPU. Sending acknowledgment to process 1.", rank, N);
    ack = 1;
    MPI_Send(&ack, 1, MPI_INT, 1, 0, MPI_COMM_WORLD);
  } else if(rank == 1) {
    MPI_Recv(&ack, 1, MPI_INT, 0, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
    printf("I am process %d, I have received the acknowledgment from process 0 that data in the GPU has been initialized.\n", rank, N);
    fflush(stdout);
  }

  // do some more work

  MPI_Finalize();

  return 0;
}

我们使用以下命令编译程序

$ mpicc -acc -o error_handling_mpi error_handling_mpi.c

如果我们使用两个 MPI 进程运行此程序，则输出将如下所示

$ mpirun -n 2 ./error_handling_mpi
Out of memory allocating -8589934592 bytes of device memory
total/free CUDA memory: 11995578368/11919294464
Present table dump for device[1]:
NVIDIA Tesla GPU 0, compute capability 3.7, threadid=1
...empty...
call to cuMemAlloc returned error 2: Out of memory

-------------------------------------------------------
Primary job terminated normally, but 1 process returned
a non-zero exit code.. Per user-direction, the job has been aborted.
-------------------------------------------------------
--------------------------------------------------------------------------
mpirun detected that one or more processes exited with non-zero status,
thus causing the job to be terminated.

进程 0 在 GPU 上分配内存时失败，并因错误而意外终止。在这种情况下，mpirun 能够识别其中一个进程失败，因此它关闭了剩余的进程并终止了应用程序。像这样简单的双进程程序很容易调试。但在实际应用中，对于数百或数千个进程，进程过早退出可能会导致应用程序无限期挂起。因此，理想的情况是捕获进程的失败，控制其他进程的终止，并提供有用的错误消息。

我们可以使用 OpenACC 错误处理功能来改进之前的程序，并在 MPI 进程失败的情况下正确终止应用程序。

在以下示例代码中，我们添加了一个错误处理回调例程，如果进程在 GPU 上执行时遇到错误，该例程将关闭其他进程。进程 0 尝试在 GPU 中分配一个大型数组，如果操作成功，进程 0 将向进程 1 发送确认。进程 0 调用 OpenACC 函数 acc_set_error_routine 将函数 handle_gpu_errors 设置为错误处理回调例程。此例程打印一条消息并调用 MPI_Abort 以关闭所有 MPI 进程。如果进程 0 成功在 GPU 上分配了数组，则进程 1 将收到确认。否则，如果进程 0 失败，它将自行终止并触发对 handle_gpu_errors 的调用。然后，进程 1 由回调例程中执行的代码终止。

#include <stdio.h>
#include <stdlib.h>
#include "mpi.h"

#define N 2147483648


typedef void (*exitroutinetype)(char *err_msg);
extern void acc_set_error_routine(exitroutinetype callback_routine);

void handle_gpu_errors(char *err_msg) {
  printf("GPU Error: %s", err_msg);
  printf("Exiting...\n\n");
  MPI_Abort(MPI_COMM_WORLD, 1);
  exit(-1);
}


int main(int argc, char **argv)
{
  int rank, size;

  MPI_Init(&argc, &argv);

  MPI_Comm_rank(MPI_COMM_WORLD, &rank);
  MPI_Comm_size(MPI_COMM_WORLD, &size);

  int ack;
  if(rank == 0) {
    float *a = (float*) malloc(sizeof(float) * N);


    acc_set_error_routine(&handle_gpu_errors);


#pragma acc enter data create(a[0:N])
#pragma acc parallel loop independent
    for(int i = 0; i < N; i++) {
      a[i] = i *0.5;
    }
#pragma acc exit data copyout(a[0:N])
    printf("I am process %d, I have initialized a vector of size %ld bytes on the GPU. Sending acknowledgment to process 1.", rank, N);
    fflush(stdout);
    ack = 1;
    MPI_Send(&ack, 1, MPI_INT, 1, 0, MPI_COMM_WORLD);
  } else if(rank == 1) {
    MPI_Recv(&ack, 1, MPI_INT, 0, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
    printf("I am process %d, I have received the acknowledgment from process 0 that data in the GPU has been initialized.\n", rank, N);
    fflush(stdout);
  }

  // more work

  MPI_Finalize();

  return 0;
}

同样，我们使用以下命令编译程序

$ mpicc -acc -o error_handling_mpi error_handling_mpi.c

我们使用两个 MPI 进程运行程序，并获得以下输出

$ mpirun -n 2 ./error_handling_mpi
Out of memory allocating -8589934592 bytes of device memory
total/free CUDA memory: 11995578368/11919294464
Present table dump for device[1]:
NVIDIA Tesla GPU 0, compute capability 3.7, threadid=1
...empty...
GPU Error: call to cuMemAlloc returned error 2: Out of memory
Exiting...

--------------------------------------------------------------------------
MPI_ABORT was invoked on rank 0 in communicator MPI_COMM_WORLD
with errorcode 1.

这次 GPU 上的错误被应用程序拦截，应用程序通过错误处理回调例程对其进行了管理。在这种情况下，该例程打印了一些关于问题的信息，并调用了 MPI_Abort 以终止剩余进程并避免应用程序出现任何意外行为。

2.11. C 示例

OpenACC 最简单的 C 示例是在 GPU 上进行向量加法

#include <stdio.h>
#include <stdlib.h>
void vecaddgpu( float *restrict r, float *a, float *b, int n ){
    #pragma acc kernels loop copyin(a[0:n],b[0:n]) copyout(r[0:n])
    for( int i = 0; i < n; ++i ) r[i] = a[i] + b[i];
}

int main( int argc, char* argv[] ){
    int n; /* vector length */
    float * a; /* input vector 1 */
    float * b; /* input vector 2 */
    float * r; /* output vector */
    float * e; /* expected output values */
    int i, errs;
    if( argc > 1 ) n = atoi( argv[1] );
    else n = 100000; /* default vector length */
    if( n <= 0 ) n = 100000;
    a = (float*)malloc( n*sizeof(float) );
    b = (float*)malloc( n*sizeof(float) );
    r = (float*)malloc( n*sizeof(float) );
    e = (float*)malloc( n*sizeof(float) );
    for( i = 0; i < n; ++i ){
         a[i] = (float)(i+1);
         b[i] = (float)(1000*i);
    }
    /* compute on the GPU */
    vecaddgpu( r, a, b, n );
    /* compute on the host to compare */
    for( i = 0; i < n; ++i ) e[i] = a[i] + b[i];
    /* compare results */
    errs = 0;
    for( i = 0; i < n; ++i ){
        if( r[i] != e[i] ){
           ++errs;
       }
   }
   printf( "%d errors found\n", errs );
   return errs;
}

此示例的重要部分是例程 vecaddgpu，它包含一个用于循环的 OpenACC 指令。此 (#pragma acc) 指令告诉编译器为以下循环 (kernels loop) 生成一个内核，在 GPU 上执行之前，为向量 a 和 b 分配 n 个元素并从主机内存复制到 GPU 内存，从 a[0] 和 b[0] 开始 (copyin(a[0:n],b[0:n]))；并在 GPU 上执行之前，为向量 r 分配 n 个元素，并将这些 n 个元素从 GPU 内存复制到主机内存，从 r[0] 开始 (copyout(r[0:n]))。

如果您将此示例键入到文件 a1.c 中，则可以使用命令 nvc -acc a1.c 构建它。-acc 标志启用对 OpenACC 编译指示的识别，并包含 OpenACC 运行时库。此命令生成通常的 a.out 可执行文件，您可以通过正常运行 a.out 来运行程序。您应该看到以下输出

0 errors found

如果相反，您得到以下输出，则您的硬件安装或 GPU 驱动程序存在问题。

libcuda.so not found, exiting
Please check that the CUDA driver is installed and the shared object
is in the install directory or on your ``LD_LIBRARY_PATH``.

您可以通过设置环境变量来启用其他输出。如果您将环境变量 NVCOMPILER_ACC_NOTIFY 设置为 1，则每次在 GPU 上运行内核时，运行时都会打印一行输出。对于此程序，您可能会得到如下所示的输出

launch CUDA kernel file=/user/guest/a1.c function=vecaddgpu
line=5 device=0 threadid=1 num_gangs=782 num_workers=1
vector_length=128 grid=782 block=128
0 errors found

额外的输出告诉您，程序为第 5 行的循环启动了一个内核，CUDA 网格大小为 782，线程块大小为 128。

如果您将环境变量 NVCOMPILER_ACC_NOTIFY 设置为 3，则输出还将包括有关数据传输的信息

upload CUDA data file=/user/guest/a1.c function=vecaddgpu
line=4 device=0 threadid=1 variable=a bytes=400000
upload CUDA data file=/user/guest/a1.c function=vecaddgpu
line=4 device=0 threadid=1 variable=b bytes=400000
launch CUDA kernel file=/user/guest/a1.c function=vecaddgpu
line=5 device=0 threadid=1 num_gangs=782 num_workers=1 vector_length=128 grid=782 block=128
download CUDA data file=/user/guest/a1.c function=vecaddgpu
line=6 device=0 threadid=1 variable=r bytes=400000
0 errors found

如果您将环境变量 NVCOMPILER_ACC_TIME 设置为 1，则运行时会汇总主机和 GPU 之间数据移动以及 GPU 上计算所花费的时间。在 Linux 上，您可能需要设置 LD_LIBRARY_PATH 环境变量以包含 /opt/nvidia/hpc_sdk/Linux_x86_64/25.1/compilers/lib 目录，或 Arm Server 目标的相应目录。OpenACC 可执行文件动态加载共享对象以实现此分析功能，并且库的路径必须可用。

对于此程序，您可能会得到类似于以下的输出

0 errors found

Accelerator Kernel Timing data
/user/guest/a1.c
  vecaddgpu  NVIDIA  devicenum=0
    time(us): 167
    4: compute region reached 1 time
        5: kernel launched 1 time
            grid: [782]  block: [128]
             device time(us): total=5 max=5 min=5 avg=5
            elapsed time(us): total=700 max=700 min=700 avg=700
    4: data region reached 2 times
        4: data copyin transfers: 2
             device time(us): total=110 max=67 min=43 avg=55
        6: data copyout transfers: 1
             device time(us): total=52 max=52 min=52 avg=52

这告诉您程序进入了一个加速区域，并在该区域总共花费了大约 167 微秒。它将两个数组复制到设备，启动了一个内核，并将一个数组带回主机。

当您编写自己的 OpenACC 程序时，您可能还会发现启用编译器反馈很有用。这可以通过 -⁠Minfo 标志启用。如果您使用命令 nvc -acc -fast -Minfo a1.c 编译此程序，您将获得以下输出

 vecaddgpu:
      4, Generating copyin(a[:n])
         Generating copyout(r[:n])
         Generating copyin(b[:n])
      5, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
          5, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
      5, Loop not fused: no successor loop
         Generated 2 alternate versions of the loop
         Generated vector simd code for the loop
         Generated 2 prefetch instructions for the loop
         Generated vector simd code for the loop
         Generated 2 prefetch instructions for the loop
         Generated vector simd code for the loop
         Generated 2 prefetch instructions for the loop
main:
     21, Loop not fused: function call before adjacent loop
         Loop not vectorized: data dependency
         Loop unrolled 16 times
         Generated 1 prefetches in scalar loop
     28, Loop not fused: dependence chain to sibling loop
         Generated 2 alternate versions of the loop
         Generated vector and scalar versions of the loop; pointer conflict tests determine which is executed
         Generated 2 prefetch instructions for the loop
         Generated vector and scalar versions of the loop; pointer conflict tests determine which is executed
         Generated 2 prefetch instructions for the loop
         Generated vector and scalar versions of the loop; pointer conflict tests determine which is executed
         Generated 2 prefetch instructions for the loop
         Loop unrolled 16 times
         Generated 1 prefetches in scalar loop
     31, Loop not fused: function call before adjacent loop

此输出给出了循环使用的 *schedule*；在本例中，schedule 是 gang,vector(128)。这意味着循环的迭代被分解为 128 个向量，并且这些向量由 GPU 的 SM 或计算单元并行执行。

此输出很重要，因为它告诉您何时将获得并行执行或顺序执行。如果您从例程 vecaddgpu 的虚拟参数 *r* 的声明中删除 restrict 关键字，则 -⁠Minfo 输出会告诉您，通过指针 r 的存储和通过指针 a 和 b 的获取之间可能存在依赖关系

        5, Complex loop carried dependence of b->,a-> prevents parallelization
   Loop carried dependence of r-> prevents parallelization
   Loop carried backward dependence of r-> prevents vectorization
   Accelerator serial kernel generated
   Accelerator kernel generated
   Generating Tesla code
    5, #pragma acc loop seq
5, Complex loop carried dependence of b->,a-> prevents parallelization
   Loop carried dependence of r-> prevents parallelization
   Loop carried backward dependence of r-> prevents vectorization
   Loop not fused: no successor loop
   Generated 2 alternate versions of the loop
   Generated vector and scalar versions of the loop; pointer conflict tests determine which is executed
   ...

标量内核在一个线程块的一个线程上运行，这比相同的并行内核慢大约 1000 倍。对于这个简单的程序，总时间主要由 GPU 初始化决定，因此您可能不会注意到时间上的差异，但在生产模式下，您需要并行内核执行才能获得可接受的性能。

对于我们的第二个示例，我们稍微修改了程序，方法是将 kernels 编译指示上的数据子句替换为 present 子句，并在对 vecaddgpu 例程的调用周围添加一个数据结构。数据结构将数据移动到主程序中的 GPU。vecaddgpu 例程中的 present 子句告诉编译器使用已在 GPU 上分配的数据的 GPU 副本。如果您在设置了 NVCOMPILER_ACC_TIME 的 GPU 上运行此程序，您会看到内核区域现在没有与之关联的数据移动。相反，数据移动都与主程序中的数据结构相关联。

#include <stdio.h>
#include <stdlib.h>

void vecaddgpu( float *restrict r, float *a, float *b, int n ){
       #pragma acc kernels loop present(r,a,b)
       for( int i = 0; i < n; ++i ) r[i] = a[i] + b[i];
}

int main( int argc, char* argv[] ){
       int n; /* vector length */
       float * a; /* input vector 1 */
       float * b; /* input vector 2 */
       float * r; /* output vector */
       float * e; /* expected output values */
       int i, errs;

       if( argc > 1 ) n = atoi( argv[1] );
       else n = 100000; /* default vector length */
       if( n <= 0 ) n = 100000;
       a = (float*)malloc( n*sizeof(float) );
       b = (float*)malloc( n*sizeof(float) );
       r = (float*)malloc( n*sizeof(float) );
       e = (float*)malloc( n*sizeof(float) );
       for( i = 0; i < n; ++i ){
           a[i] = (float)(i+1);
           b[i] = (float)(1000*i);
       }
       /* compute on the GPU */
       #pragma acc data copyin(a[0:n],b[0:n]) copyout(r[0:n])
       {
           vecaddgpu( r, a, b, n );
       }
       /* compute on the host to compare */
       for( i = 0; i < n; ++i ) e[i] = a[i] + b[i];
       /* compare results */
       errs = 0;
       for( i = 0; i < n; ++i ){
           if( r[i] != e[i] ){
           ++errs;
           }
       }
       printf( "%d errors found\n", errs );
       return errs;
}

2.12. Fortran 示例

OpenACC 最简单的 Fortran 示例是在 GPU 上进行向量加法。

2.12.1. 在 GPU 上进行向量加法

本节包含两个在 GPU 上进行向量加法的 Fortran 示例

module vecaddmod
  implicit none
 contains
  subroutine vecaddgpu( r, a, b, n )
   real, dimension(:) :: r, a, b
   integer :: n
   integer :: i
!$acc kernels loop copyin(a(1:n),b(1:n)) copyout(r(1:n))
   do i = 1, n
    r(i) = a(i) + b(i)
   enddo
  end subroutine
end module

program main
  use vecaddmod
  implicit none
  integer :: n, i, errs, argcount
  real, dimension(:), allocatable :: a, b, r, e
  character*10 :: arg1
  argcount = command_argument_count()
  n = 1000000  ! default value
  if( argcount >= 1 )then
   call get_command_argument( 1, arg1 )
   read( arg1, '(i)' ) n
   if( n <= 0 ) n = 100000
  endif
  allocate( a(n), b(n), r(n), e(n) )
  do i = 1, n
   a(i) = i
   b(i) = 1000*i
  enddo
  ! compute on the GPU
  call vecaddgpu( r, a, b, n )
  ! compute on the host to compare
  do i = 1, n
   e(i) = a(i) + b(i)
  enddo
  ! compare results
  errs = 0
  do i = 1, n
   if( r(i) /= e(i) )then
     errs = errs + 1
   endif
  enddo
  print *, errs, ' errors found'
  if( errs ) call exit(errs)
end program

此示例的重要部分是子例程 vecaddgpu，它包含一个用于循环的 OpenACC 指令。此 (!$acc) 指令告诉编译器为以下循环 (kernels loop) 生成一个内核，在 GPU 上执行之前，为向量 a 和 b 分配 n 个元素并从主机内存复制到 GPU 内存，从 a(1) 和 b(1) 开始 (copyin(a(1:n),b(1:n))；并在 GPU 上执行之前，为向量 r 分配 n 个元素，并将这些 n 个元素从 GPU 内存复制到主机内存，从 r(1) 开始 (copyout(r(1:n))。

如果您将此示例键入到文件 f1.f90 中，则可以使用命令 nvfortran -acc f1.f90 构建它。-acc 标志启用对 OpenACC 编译指示的识别，并包含 OpenACC 运行时库。此命令生成通常的 a.out 可执行文件，您可以通过正常运行 a.out 来运行程序。您应该看到以下输出

0  errors found

如果相反，您得到以下输出，则您的硬件安装或 CUDA 驱动程序存在问题。

libcuda.so not found, exiting
Please check that the CUDA driver is installed and the shared object
is in the install directory or on your LD_LIBRARY_PATH.

您可以通过设置环境变量来启用其他输出。如果您将环境变量 NVCOMPILER_ACC_NOTIFY 设置为 1，则每次在 GPU 上运行内核时，运行时都会打印一行输出。对于此程序，您可能会得到如下所示的输出

launch CUDA kernel file=/user/guest/f1.f90 function=vecaddgpu
line=9 device=0 threadid=1 num_gangs=7813 num_workers=1
vector_length=128 grid=7813 block=128
            0  errors found

额外的输出告诉您，程序为第 9 行的循环启动了一个内核，CUDA 网格大小为 7813，线程块大小为 128。如果您将环境变量 NVCOMPILER_ACC_NOTIFY 设置为 3，则输出还将包括有关数据传输的信息

upload CUDA data  file=/user/guest/f1.f90 function=vecaddgpu
line=8 device=0 threadid=1 variable=a bytes=4000000
upload CUDA data  file=/user/guest/f1.f90 function=vecaddgpu
line=8 device=0 threadid=1 variable=b bytes=4000000
launch CUDA kernel  file=/user/guest/f1.f90 function=vecaddgpu
line=9 device=0 threadid=1 num_gangs=7813 num_workers=1 vector_length=128 grid=7813 block=128
download CUDA data  file=/user/guest/f1.f90 function=vecaddgpu
line=12 device=0 threadid=1 variable=r bytes=4000000
            0  errors found

如果您将环境变量 NVCOMPILER_ACC_TIME 设置为 1，则运行时会汇总主机和 GPU 之间数据移动以及 GPU 上计算所花费的时间。对于此程序，您可能会得到类似于以下的输出

            0  errors found

Accelerator Kernel Timing data
/home/ams/tat/example-f/f1.f90
  vecaddgpu  NVIDIA  devicenum=0
    time(us): 1,040
    8: compute region reached 1 time
        9: kernel launched 1 time
            grid: [7813]  block: [128]
             device time(us): total=19 max=19 min=19 avg=19
            elapsed time(us): total=738 max=738 min=738 avg=738
    8: data region reached 2 times
        8: data copyin transfers: 2
             device time(us): total=689 max=353 min=336 avg=344
        12: data copyout transfers: 1
             device time(us): total=332 max=332 min=332 avg=332

这告诉您程序进入了一个加速区域，并在该区域总共花费了大约 1 毫秒。它将两个数组复制到设备，启动了一个内核，并将一个数组带回主机。

当您编写自己的 OpenACC 程序时，您可能还会发现启用编译器反馈很有用。这可以通过 -Minfo 标志启用。

如果您使用命令 nvfortran -acc -fast -Minfo f1.f90 编译此程序，您将获得以下输出

vecaddgpu:
      8, Generating copyin(a(:n))
         Generating copyout(r(:n))
         Generating copyin(b(:n))
      9, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
          9, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
      9, Loop not fused: no successor loop
         Generated 2 alternate versions of the loop
         Generated vector simd code for the loop
         Generated 2 prefetch instructions for the loop
         Generated vector simd code for the loop
         Generated 2 prefetch instructions for the loop
         Generated vector simd code for the loop
         Generated 2 prefetch instructions for the loop
main:
     29, Loop not fused: function call before adjacent loop
         Loop not vectorized: may not be beneficial
         Unrolled inner loop 8 times
     36, Loop not fused: function call before adjacent loop
         2 loops fused

此输出给出了循环使用的 schedule；在本例中，schedule 是 gang, vector(128)。这意味着循环的迭代被分解为 128 个向量，并且这些向量由 GPU 的 SM 并行执行。此输出很重要，因为它告诉您何时将获得并行执行或顺序执行。

对于我们的第二个示例，我们稍微修改了程序，方法是将 kernels 编译指示上的数据子句替换为 present 子句，并在对 vecaddgpu 子例程的调用周围添加一个数据结构。数据结构将数据移动到主程序中的 GPU。vecaddgpu 子例程中的 present 子句告诉编译器使用已在 GPU 上分配的数据的 GPU 副本。如果您在设置了 NVCOMPILER_ACC_TIME 的 GPU 上运行此程序，您会看到内核区域现在没有与之关联的数据移动。相反，数据移动都与主程序中的数据结构相关联。

在 Fortran 程序中，如果编译器可以从声明中推断出数组边界，或者如果数组是假定形状虚拟参数或可分配数组，则不必在数据子句中指定数组边界。

module vecaddmod
  implicit none
 contains
  subroutine vecaddgpu( r, a, b, n )
   real, dimension(:) :: r, a, b
   integer :: n
   integer :: i
!$acc kernels loop present(r,a,b)
   do i = 1, n
    r(i) = a(i) + b(i)
   enddo
  end subroutine
end module

program main
  use vecaddmod
  implicit none
  integer :: n, i, errs, argcount
  real, dimension(:), allocatable :: a, b, r, e
  character*10 :: arg1
  argcount = command_argument_count()
  n = 1000000  ! default value
  if( argcount >= 1 )then
  call get_command_argument( 1, arg1 )
   read( arg1, '(i)' ) n
   if( n <= 0 ) n = 100000
  endif
  allocate( a(n), b(n), r(n), e(n) )
  do i = 1, n
   a(i) = i
   b(i) = 1000*i
  enddo
  ! compute on the GPU
!$acc data copyin(a,b) copyout(r)
  call vecaddgpu( r, a, b, n )
!$acc end data
  ! compute on the host to compare
  do i = 1, n
   e(i) = a(i) + b(i)
  enddo
  ! compare results
  errs = 0
  do i = 1, n
   if( r(i) /= e(i) )then
     errs = errs + 1
   endif
  enddo
  print *, errs, ' errors found'
  if( errs ) call exit(errs)
end program

2.12.2. 利用多个设备的多线程程序

这个简单的示例演示了如何运行利用多个设备的多线程主机程序。

program tdot
! Compile with "nvfortran -mp -acc tman.f90 -lblas
! Set OMP_NUM_THREADS environment variable to run with
! up to 2 threads, currently.
!
use openacc
use omp_lib
!
integer, parameter :: N = 10000
real*8 x(N), y(N), z
integer, allocatable :: offs(:)
real*8, allocatable :: zs(:)
real*8 ddot

! Max at 2 threads for now
nthr = omp_get_max_threads()
if (nthr .gt. 2) nthr = 2
call omp_set_num_threads(nthr)

! Run on host
call random_number(x)
call random_number(y)
z = ddot(N,x,1,y,1)
print *,"Host Serial",z

! Attach each thread to a device
!$omp PARALLEL private(i)
     i = omp_get_thread_num()
     call acc_set_device_num(i, acc_device_nvidia)
!$omp end parallel

! Break up the array into sections
nsec = N / nthr
allocate(offs(nthr),zs(nthr))
offs = (/ (i*nsec,i=0,nthr-1) /)
zs = 0.0d0

! Decompose the problem across devices
!$omp PARALLEL private(i,j,z)
     i = omp_get_thread_num() + 1
     z = 0.0d0
     !$acc kernels loop &
        copyin(x(offs(i)+1:offs(i)+nsec),y(offs(i)+1:offs(i)+nsec))
     do j = offs(i)+1, offs(i)+nsec
         z = z + x(j) * y(j)
     end do
     zs(i) = z
!$omp end parallel
z = sum(zs)
print *,"Multi-Device Parallel",z
end

程序开始时，每个线程都调用 acc_set_device_num，以便每个线程都使用不同的 GPU。在计算 OpenMP 并行区域内，每个线程将其所需的数据复制到其 GPU 并继续。

2.13. 故障排除提示和已知限制

此版本的 NVIDIA HPC SDK 编译器实现了 OpenACC 2.7 规范的大部分功能。有关尚未实现的功能的说明，请参阅第 3 章“已实现的功能”。

Linux CUDA 驱动程序将关闭空闲 GPU 的电源。这意味着，如果您使用没有连接显示器的 GPU，或 NVIDIA 仅计算 GPU，并且没有打开的 CUDA 上下文，则 GPU 将断电，直到需要时才会重新启动。由于重新启动 GPU 可能需要长达一秒钟的时间，因此在启动程序时可能会遇到明显的延迟。当您在环境变量 NVCOMPILER_ACC_TIME 设置为 1 的情况下运行程序时，此时间将显示为初始化时间。如果您正在运行许多测试，或者想要将实际时间与初始化时间隔离，则可以在后台运行 NVIDIA 实用程序 nvcudainit。此实用程序打开一个 CUDA 上下文并保持打开状态，直到您终止它或让它完成。

NVIDIA OpenACC 编译器支持 async 子句和 wait 指令。当您使用异步计算或数据移动时，您有责任确保程序具有足够的同步来解决主机和 GPU 之间的任何数据竞争。如果您的程序使用 async 子句并生成了错误的答案，则可以通过在运行程序之前将环境变量 NVCOMPILER_ACC_SYNCHRONOUS 设置为 1 来测试 async 子句是否导致问题。此操作会导致 OpenACC 运行时忽略 async 子句并在同步模式下运行程序。

3. 已实现的功能

本节概述了 NVIDIA HPC SDK 编译器中当前实现的 OpenACC 功能，并列出了已知限制。

3.1. OpenACC 规范合规性

NVIDIA HPC SDK 编译器包括对 OpenACC 2.7 规范的大部分功能的支持。以下 OpenACC 2.7 功能不受支持

Declare link

嵌套并行

将 cache 子句变量引用限制为缓存区域内的变量。

reduction 子句中的子数组和复合变量

self 子句

数据结构上的 default 子句

3.2. 默认值

默认的 ACC_DEVICE_TYPE 是 acc_device_nvidia，就像 -acc 编译器选项默认情况下以 NVIDIA GPU 为目标一样。设备类型 acc_device_default 和 acc_device_not_host 的行为与 acc_device_nvidia 相同。可以使用环境变量或调用 acc_set_device_type() 来更改设备类型。

对于 acc_device_nvidia 类型，默认的 ACC_DEVICE_NUM 是 0，这与 CUDA 设备编号系统一致。有关更多信息，请参阅准备系统中的 nvaccelinfo 输出。可以使用环境变量或调用 acc_set_device_num 来更改设备编号。

3.3. 环境变量

本节总结了 NVIDIA OpenACC 支持的环境变量。这些环境变量是用户可设置的环境变量，用于控制启用加速器的程序在执行时的行为。这些环境变量必须遵守以下规则

环境变量的名称必须为大写。

环境变量的值不区分大小写，并且可能具有前导和尾随空格。

如果在程序启动后环境变量的值发生更改，即使程序本身修改了这些值，其行为也是实现定义的。

下表包含当前支持的环境变量，并提供了每个变量的简要说明。

表 1. 支持的环境变量
使用此环境变量…	执行此操作…
NVCOMPILER_ACC_CUDA_PROFSTOP	设置为 1（或任何正值）以告知运行时环境在退出时插入 ‘atexit(cuProfilerStop)’ 调用。在配置文件不完整或发出消息以调用 cudaProfilerStop() 的情况下，可能需要此行为。
NVCOMPILER_ACC_DEVICE_NUM	设置要使用的默认设备编号。`NVCOMPILER_ACC_DEVICE_NUM`。指定在执行加速区域时要使用的默认设备编号。此环境变量的值必须是介于零和连接到主机的设备数量之间的非负整数。
ACC_DEVICE_NUM	遗留名称。已被 `NVCOMPILER_ACC_DEVICE_NUM` 取代。
NVCOMPILER_ACC_DEVICE_TYPE	设置要用于 OpenACC 区域的默认设备类型。`NVCOMPILER_ACC_DEVICE_TYPE`。指定当程序已编译为使用多种不同类型的设备时，在执行加速区域时要使用的加速器设备。此环境变量的值是实现定义的，在 NVIDIA OpenACC 实现中，可以是字符串 NVIDIA、MULTICORE 或 HOST
ACC_DEVICE_TYPE	遗留名称。已被 `NVCOMPILER_ACC_DEVICE_TYPE` 取代。
NVCOMPILER_ACC_GANGLIMIT	对于 NVIDIA CUDA 设备，这定义了内核将启动的最大 gangs（CUDA 线程块）数量。
NVCOMPILER_ACC_NOTIFY	在没有参数的情况下，将为每个内核启动和/或数据传输向 stderr 写入调试消息。当设置为整数值时，该值用作位掩码以打印有关以下内容的信息 1：内核启动 2：数据传输 4：区域进入/退出 8：等待操作或与设备的同步 16：设备内存分配和释放
NVCOMPILER_ACC_PROFLIB	使用新的分析器动态库接口启用第三方工具接口。
NVCOMPILER_ACC_SYNCHRONOUS	禁用异步启动和数据移动。
NVCOMPILER_ACC_TIME	启用轻量级分析器以测量数据移动和加速器内核执行时间，并在程序执行结束时打印摘要。

3.4. OpenACC Fortran API 扩展

本节总结了 NVIDIA Fortran 编译器中实现的 OpenACC API 扩展。

3.4.1. acc_malloc

acc_malloc 函数返回一个设备指针，类型为 type(c_devptr) 的变量，指向设备上新分配的内存。如果无法分配数据，则此函数返回 C_NULL_DEVPTR。

NVIDIA Fortran 中有一种受支持的调用格式

type(c_devptr) function acc_malloc (bytes)

bytes 是一个整数，用于指定请求的字节数

3.4.2. acc_free

acc_free 子例程释放先前由 acc_malloc 分配的内存。它接受一个参数，可以是派生类型 type(c_devptr) 实例中包含的设备指针，或者为了方便起见，可以是 CUDA Fortran 设备数组。在 NVIDIA Fortran 中，使用 F90 allocate 语句分配的 CUDA Fortran 设备数组调用 acc_free（或 cudaFree）会导致未定义的行为。

NVIDIA Fortran 中有两种受支持的调用格式

subroutine acc_free ( devptr )

devptr 是派生类型 type(c_devptr) 的实例

subroutine acc_free ( dev )

dev 是 CUDA Fortran 设备数组

3.4.3. acc_map_data

acc_map_data 例程将主机数据关联（映射）到设备数据。第一个参数是主机数组、连续主机数组节，或 type(c_ptr) 中包含的地址。第二个参数必须是 type(c_devptr) 中包含的设备地址，例如从 acc_malloc 或 acc_deviceptr 返回的地址，或者 CUDA Fortran 设备数组。

NVIDIA Fortran 中有四种受支持的调用格式

subroutine acc_map_data ( host, dev, bytes )

host 是主机变量、数组或起始数组元素
dev 是 CUDA Fortran 设备变量、数组或起始数组元素
bytes 是一个整数，用于指定映射长度（以字节为单位）

subroutine acc_map_data ( host, dev )

host 是主机数组或连续主机数组段
dev 是一个 CUDA Fortran 设备数组或数组段，它与 host 一致

subroutine acc_map_data ( host, devptr, bytes )

host 是主机变量、数组或起始数组元素
devptr 是派生类型 type(c_devptr) 的实例
bytes 是一个整数，用于指定映射长度（以字节为单位）

subroutine acc_map_data ( ptr, devptr, bytes )

ptr 是派生类型 (c_ptr) 的实例
devptr 是派生类型 type(c_devptr) 的实例
bytes 是一个整数，用于指定映射长度（以字节为单位）

3.4.4. acc_unmap_data

acc_unmap_data 例程从指定的主机数据中取消映射（或解除关联）设备数据。

NVIDIA Fortran 中有一种受支持的调用格式

subroutine acc_unmap_data ( host )

host 是一个主机变量，它在先前调用 acc_map_data 时被映射到设备数据

3.4.5. acc_deviceptr

acc_deviceptr 函数返回设备指针，类型为 type(c_devptr)，它映射到主机地址。输入参数是一个主机变量或数组元素，它在当前设备上具有活动生命周期。如果数据不存在，此函数返回 C_NULL_DEVPTR。

NVIDIA Fortran 中有一种受支持的调用格式

type(c_devptr) function acc_deviceptr ( host )

host 是任何类型、种类和秩的主机变量或数组元素

3.4.6. acc_hostptr

acc_hostptr 函数返回主机指针，类型为 type(c_ptr)，它映射到设备地址。输入参数是一个设备地址，例如从 acc_malloc 或 acc_deviceptr 返回的地址，或者是一个 CUDA Fortran 设备数组。

NVIDIA Fortran 中有两种受支持的调用格式

type(c_ptr) function acc_hostptr ( dev )

dev 是 CUDA Fortran 设备数组

type(c_ptr) function acc_hostptr ( devptr )

devptr 是派生类型 type(c_devptr) 的实例

3.4.7. acc_is_present

acc_is_present 函数返回 .true. 或 .false.，具体取决于主机变量或数组区域是否在设备上存在。

NVIDIA Fortran 中有两种受支持的调用格式

logical function acc_is_present ( host )

host 是固有类型的连续数组段

logical function acc_is_present ( host, bytes )

host 是任何类型、种类和秩的主机变量
bytes 是一个整数，指定要检查的数据长度

3.4.8. acc_memcpy_to_device

acc_memcpy_to_device 例程将数据从本地内存复制到设备内存。源地址是主机数组、连续数组段或类型 (c_ptr) 中包含的地址。目标地址必须是设备地址，例如从 acc_malloc 或 acc_deviceptr 返回的地址，或者是一个 CUDA Fortran 设备数组。

NVIDIA Fortran 中有四种受支持的调用格式

subroutine acc_memcpy_to_device ( dev, src, bytes )

dev 是一个 CUDA Fortran 设备变量、数组或起始数组元素
src 是一个主机变量、数组或起始数组元素
bytes 是一个整数，指定要复制的长度（以字节为单位）

subroutine acc_memcpy_to_device ( dev, src )

dev 是一个 CUDA Fortran 设备数组或连续数组段
src 是一个主机数组或数组段，它与 dev 一致

subroutine acc_memcpy_to_device ( devptr, src, bytes )

devptr 是派生类型 type(c_devptr) 的实例
src 是一个主机变量、数组或起始数组元素
bytes 是一个整数，指定要复制的长度（以字节为单位）

subroutine acc_memcpy_to_device ( devptr, ptr, bytes )

其中 devptr 是派生类型 (c_devptr) 的实例
ptr 是派生类型 (c_ptr) 的实例
bytes 是一个整数，指定要复制的长度（以字节为单位）

3.4.9. acc_memcpy_from_device

acc_memcpy_from_devic``e 例程将数据从设备内存复制到本地内存。源地址必须是设备地址，例如从 ``acc_malloc、 acc_deviceptr 或 CUDA Fortran 设备数组返回的地址。源地址是主机数组、连续数组段或类型 (c_ptr) 中包含的地址。

NVIDIA Fortran 中有四种受支持的调用格式

subroutine acc_memcpy_from_device ( dest, dev, bytes )

dest 是一个主机变量、数组或起始数组元素
dev 是一个 CUDA Fortran 设备变量、数组或起始数组元素
bytes 是一个整数，指定要复制的长度（以字节为单位）

subroutine acc_memcpy_from_device ( dest, dev )

dest 是一个主机数组或连续数组段
dev 是一个 CUDA Fortran 设备数组或数组段，它与 dest 子例程一致

subroutine acc_memcpy_from_device ( dest, devptr, bytes )

其中 dest 是一个主机变量、数组或起始数组元素。
devptr 是派生类型 (c_devptr) 的实例。
bytes 是一个整数，指定要复制的长度（以字节为单位）

subroutine acc_memcpy_from_device ( ptr, devptr, bytes )

ptr 是派生类型 (c_ptr) 的实例
devptr 是派生类型 type(c_devptr) 的实例
bytes 是一个整数，指定要复制的长度（以字节为单位）

3.4.10. acc_get_cuda_stream

acc_get_cuda_stream 函数返回 CUDA 流值，该值对应于 OpenACC 异步队列。输入参数是一个异步编号或预定义值，例如 acc_async_sync。此调用仅在 NVIDIA 平台上受支持。

NVIDIA Fortran 中有一种受支持的调用格式

integer(acc_handle_kind) function acc_get_cuda_stream ( async )

async 是用户定义或预定义的异步值

3.4.11. acc_set_cuda_stream

acc_set_cuda_stream 子例程在当前设备上为 OpenACC 异步队列设置 CUDA 流值。输入参数是一个异步编号和一个流。此调用仅在 NVIDIA 平台上受支持。

NVIDIA Fortran 中有一种受支持的调用格式

subroutine acc_set_cuda_stream ( async, stream )

async 和 stream 是 acc_handle_kind 类型的整数

3.5. 已知限制

本节包括 NVIDIA HPC SDK 编译器实现的 OpenACC API 中的已知限制。

3.5.1. ACC 例程指令限制

外部变量可能不与 acc routine 过程一起使用。
仅当 NVIDIA GPU 支持计算能力 3.0 或更高版本时，才支持在带有 acc routine 的过程中的归约。
尚不支持 Fortran 假定形状参数。

3.5.2. C++ 和 OpenACC 限制

可以出现在 OpenACC 数据构造和计算区域中的数据存在限制

OpenACC 数据子句中不支持可变长度数组；VLA 不是 C⁠+⁠+ 标准的一部分。

当需要构造函数和析构函数的类类型的变量出现在数据子句中时，它们的行为不正确。

异常在计算区域中不被处理。

成员变量在 host_data 构造的 use_device 子句中未完全支持；此放置可能会在运行时导致错误。

3.5.3. 其他限制

在调用 acc_shutdown 后，不支持以另一个加速器设备为目标。

3.6. 与优化的交互

本节讨论程序员应注意的与编译器优化的交互。

3.6.1. 与内联的交互

程序内联可以通过多种方式启用。用户控制的内联使用 -Minline 标志或使用 -Mextract=lib: 和 -Minline=lib: 标志启用。对于 C 和 C++，编译器控制的内联使用 -Mautoinline 或 -fast 标志启用。过程间分析也可以使用 -Mipa=inline 选项来控制内联。内联是一种性能优化，通过消除过程调用的开销，并通过专门化和优化调用站点的内联过程代码。

当包含计算构造（acc parallel 或 acc kernels）的过程内联到 acc data 构造中时，编译器将使用数据构造子句来优化主机和设备之间的数据移动。在某些情况下，当某些变量的主机和设备副本不同时，这可能会产生不同的答案。例如，数据构造可以为标量变量或包含标量变量的 Fortran 公共块指定数据子句。内联过程中的计算构造现在将看到标量变量存在于设备上，并将使用该变量的设备副本。在内联之前，计算构造可能已对该标量变量使用默认的 firstprivate 行为，这将使用该变量的主机值。

声明

注意

所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“材料”），均“按原样”提供。NVIDIA 对这些材料不作任何明示、暗示、法定或其他方面的保证，并且明确声明不承担所有关于不侵权、适销性和针对特定用途适用性的暗示保证。

所提供的信息据信是准确可靠的。但是，NVIDIA 公司对使用此类信息造成的后果或因使用此类信息而可能导致的侵犯第三方专利或其他权利的行为不承担任何责任。未通过暗示或其他方式授予 NVIDIA 公司任何专利权下的许可。本出版物中提及的规范如有更改，恕不另行通知。本出版物取代并替换之前提供的所有其他信息。未经 NVIDIA 公司的明确书面批准，NVIDIA 公司产品不得用作生命维持设备或系统中的关键组件。

商标

NVIDIA、NVIDIA 徽标、CUDA、CUDA-X、GPUDirect、HPC SDK、NGC、NVIDIA Volta、NVIDIA DGX、NVIDIA Nsight、NVLink、NVSwitch 和 Tesla 是 NVIDIA Corporation 在美国和其他国家/地区的商标和/或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。