【NVIDIA HPC SDK】使用NVIDIA GPU并行化C++代码

c++算法并行化加速探索者

已于 2022-11-17 16:50:17 修改

阅读量1.1k

点赞数 1

分类专栏：标准C++算法GPU并行化文章标签： c++ 开发语言 linux 大数据

于 2022-11-10 15:26:10 首次发布

原文链接：https://developer.nvidia.com/blog/accelerating-standard-c-with-gpus-using-stdpar/

版权

标准C++算法GPU并行化专栏收录该内容

1 篇文章 0 订阅

订阅专栏

NVIDIA HPC SDK引入了nvc++编译器，支持使用C++标准并行特性在GPU上加速代码，无需语言扩展。通过-stdpar选项，编译器可以自动在NVIDIA GPU上并行执行标准C++算法。本文介绍了C++ 17的并行执行策略，以及NVIDIA HPC SDK在GPU加速C++代码方面的实现和指导原则，包括对内存、函数指针和异常处理的限制。此外，还展示了LULESH应用程序的GPU并行化例子，展示其性能提升。

摘要由CSDN通过智能技术生成

【NVIDIA HPC SDK】不重构代码的情况下使用NVIDIA GPU并行化C++代码

1.前言
2.c++标准下的并行
3.NVIDIA HPC SDK
4.使用-stdpar选项启用c++并行算法
5.简单例子
6.gpu加速c++并行算法编码guidelines
7.一个大型的例子：LULESH
8.开始学习C++的GPU并行算法
- 8.1支持的NVIDIA GPU
- 8.2支持的CUDA版本
9总结

原文地址：https://developer.nvidia.com/blog/accelerating-standard-c-with-gpus-using-stdpar/

在这里插入图片描述

1.前言

以往来说，如果不使用语言扩展或附加库，是不可能用gpu加速标准c++代码的:

CUDA c++需要在GPU函数上使用 __host__ 和__device__关键字标明，并且在GPU内核启动时使用<<<gridNum,blockNum>>>语法。
OpenACC使用#pragmas来控制GPU加速
Thrust库可以较为简便的写出并行代码，但在内部使用语言扩展，只支持某些的CPU和GPU的backend。应用程序的可移植性受到库可移植性的限制。

在某些场景下，使用上述方法重写代码的性价比还可以。但如果能在不付出成本的情况下获得同样的效果呢?如果你能用你的标准c++代码在GPU上加速呢?

现在可以了！NVIDIA最近发布了nvc++， NVIDIA HPC SDK c++编译器。这是第一个支持gpu加速标准c++的编译器，不用语言扩展、语法、指令或非标准库。您可以编写标准c++，并且可移植到其他编译器和系统，并使用nvc++自动加速高性能NVIDIA gpu。我们构建它是为了让您可以花更少的时间进行移植，而把更多的时间花在真正重要的事情上——用计算科学解决世界上的问题。

2.c++标准下的并行

c++ 11引入了内存模型、并发执行模型和并发库，提供了一种利用多核处理器的标准方法。然而，直到最近，标准c++还缺乏用于并行编程的高级工具。

c++ 17标准引入了更高级别的并行特性，允许用户使用标准库算法的并行化。

这种高级并行性是通过将执行策略作为第一个参数添加到任何支持执行策略的算法中来表示的。大多数现有的标准c++算法都进行了增强，以支持执行策略。c++ 17定义了几个新的并行算法，包括有用的std::reduce和std::transform_reduce。

C++ 定义了4种（串行/并行）execution policies

std::execution::seq：顺序执行。不允许并行。
std::execution::unseq:在被调用线程上向量化执行(此执行策略是在c++ 20中添加的)。(译者认为，这句话是，串行执行程序，但乱序执行的意思)。
std::execution::par:在一个或多个线程上并行执行。
std::execution::par_unseq:在一个或多个线程上并行执行，每个线程都可能向量化(译者认为，与par的区别在于，在并行后乱序执行线程)。

当使用std::execution::seq以外的执行策略时，代码向编译器传递了两件重要的事情:

更希望程序并行执行，但不强制要求其并行执行。某些不太适合并行的代码实现可能会忽略你的希望，按顺序运行算法。但是，某些符合并行实现的代码会遵循你的希望，并在谨慎的情况下并行运行算法。
该算法可以安全地并行运行。对于std::execution::par和std::execution::par_unseq策略，任何你编写的代码——比如传递给算法的迭代器、lambda或函数对象——在不同线程上跑的时候，不能有数据竞争（data races）现象。如果使用std::execution::unseq和std::execution::par_unseq策略，就要保证在一个线程中乱序调用某个迭代器/lambda或函数对象时，你写的代码都不能引入数据竞争或死锁。有关更多关于有可能产生死锁的情况，请查阅：forward progress guarantees provided by the parallel policies或查阅： CppCon 2018: Bryce Adelstein Lelbach “The C++ Execution Model”.（Youtube，需要科学上网）

c++标准赋予编译器很大的自由，可以选择是否、何时以及如何并行执行算法。转发进程保证尊重程序员的意愿。例如，std::execution::unseq可以通过向量化实现，而std::execution::par可以通过CPU线程池实现。也可以在GPU上执行并行算法，这对于具有足够并行性的调用来说是一个很好的选择，可以利用现代GPU处理器的处理能力和内存带宽加速程序。

3.NVIDIA HPC SDK

NVIDIA HPC SDK是一套综合的编译器、库和工具，用于GPU加速HPC建模和仿真应用。随着对运行Linux的NVIDIA gpu和x86-64、OpenPOWER或Arm cpu的支持，NVIDIA HPC SDK为构建跨平台、性能可移植和可扩展的HPC应用程序提供了经过验证的工具和技术。

NVIDIA HPC SDK包括新的NVIDIA HPC c++编译器：nvc++。nvc++支持c++ 17, c++标准并行(stdpar)用于CPU和GPU。

NVC++可以用并行执行策略std::execution::par或std::execution::par_unseq编译标准c++代码，以便在NVIDIA gpu上执行。一个nvc++命令行选项-stdpar用于启用gpu加速的c++并行算法。Lambdas(包括泛型Lambdas)在并行算法调用中得到完全支持。启用GPU加速不需要任何语言扩展或非标准库。所有数据在主机内存和GPU设备内存之间的移动都是在CUDA统一内存的控制下隐式和自动执行的。

用nvc++自动使用GPU加速c++并行算法是很容易的。然而，其中也存在一些限制和不足，我们将在后面的文章中解释。

4.使用-stdpar选项启用c++并行算法

c++并行算法的GPU加速是通过nvc++的-stdpar命令行选项启用的。如果指定了-stdpar，几乎所有使用并行执行策略的算法都会被编译，以便在NVIDIA GPU上并行运行:

nvc++ -stdpar program.cpp -o program

5.简单例子

下面是几个简单的示例，让您了解c++并行算法是如何工作的。

在c++早期，对存储在适当容器中的对象进行排序是比较容易的，只需调用一个函数，例如:

std::sort(employees.begin(), employees.end(),
            CompareByLastName());

假设CompareByLastName是线程安全的(大多数比较函数都是)，那么用c++并行算法并行处理这种排序是很简单的。使用<execution>并在函数调用中添加执行策略:

std::sort(std::execution::par,
           employees.begin(), employees.end(),
           CompareByLastName());

使用std::accumulate算法计算容器中所有元素的和也很简单。在c++ 17之前，在计算和的同时以某种方式转换数据有点麻烦。例如，要计算员工的平均年龄，可以编写以下代码:

int ave_age =
    std::transform_reduce(std::execution::par_unseq,
                          employees.begin(), employees.end(),
                          0, std::plus<int>(),
                          [](const Employee& emp){
                              return emp.age();
                          })
    / employees.size();

（译者注：原文给的例子知识比较简单的示例语句，跑不起来，我将在这里放一个github链接，分享一个能跑起来的nvc++示例）

6.gpu加速c++并行算法编码guidelines

gpu不仅仅是拥有更多线程的cpu。为了利用GPU上的大量并行性，GPU编程模型通常会对在GPU上执行的代码进行一些限制。

c++并行算法的nvc++实现也有限制。在这篇文章中，我们列出了在文章发表时（2020.08.04）存在的主要限制。NVIDIA正在跨硬件和软件团队合作，为未来的新版本消除尽可能多的这些限制。

本节涵盖的主题包括:

c++并行算法和device函数注释
c++并行算法与CUDA Unified Memory
c++并行算法和函数指针
随机访问迭代器 /迭代器（Random-access iterators）
与c++标准库的互用性
GPU code中的没有异常

6.1 c++并行算法和device函数注释

与CUDA c++不同，函数不需要任何__device__注释或其他特殊标记来编译用于GPU执行。nvc++编译器遍历每个源文件的调用图，并自动推断出GPU执行时必须编译哪些函数。

但是，只有当编译器可以在调用函数的同一源文件中看到函数定义时，这种方法才有效。这对于大多数内联函数和模板函数都是正确的，但是当函数定义在不同的源文件中或从外部库链接时可能会失效。

6.2 c++并行算法与CUDA Unified Memory

大多数GPU编程模型要求数据对象在CPU内存和GPU内存之间的移动由用户手动管理。例如，CUDA提供cudaMemcpy，而OpenACC有#pragma acc数据指令来控制数据在CPU内存和GPU内存之间的移动。

nvc++依赖CUDA统一内存（CUDA Unified Memory）在CPU和GPU内存之间进行所有数据移动。通过CUDA设备驱动程序和NVIDIA GPU硬件的支持，CUDA统一内存管理器根据使用情况自动移动某些类型的数据。

目前，只有nvc++编译的CPU代码中动态分配在堆上的数据才能被自动管理。GPU代码中动态分配的内存只能从GPU代码中可见，CPU无法访问。

同样，CPU和GPU堆栈内存以及大多数系统中用于全局对象的内存不能被自动管理。没有被nvc++编译的带有-stdpar选项的CPU代码的动态分配不会被CUDA统一内存自动管理，即使它在CPU堆上。

因此，在c++并行算法调用中，任何被解引用的指针和任何被引用的对象都必须引用CPU堆。撤销对CPU堆栈或全局对象的指针的引用会导致GPU代码中的内存冲突。

例如，std::vector使用动态分配的内存，当使用stdpar时，它可以从GPU访问。在c++并行算法中迭代std::vector的内容按预期执行:

std::vector<int> v = ...;
std::sort(std::execution::par,
          v.begin(), v.end()); // Okay, accesses heap memory.

与之相反，std::array不执行动态分配。它的内容存储在std::array对象本身中，该对象通常位于CPU栈上。除非std::array对象本身被分配到堆上，否则迭代std::array的内容将不起作用:

std::array<int, 1024> a = ...;
std::sort(std::execution::par,
          a.begin(), a.end()); // Fails, array is on a CPU stack.

在上述示例中，函数参数a是通过引用传参的。运行在GPU上的lambda代码体中的代码尝试访问位于CPU栈内存中的a。这将导致内存冲突和未定义的行为。在这种情况下，通过更改lambda值来按值传参可以解决这个问题:

void saxpy(float* x, float* y, int N, float a) {
  std::transform(std::execution::par_unseq, x, x + N, y, y,
                 [=](float xi, float yi){ return a * xi + yi; });
}

通过这个一个字符（[=]）的更改，lambda生成a的副本，然后将其复制到GPU，并且不会尝试从GPU代码引用CPU栈内存。

6.3 c++并行算法和函数指针

为了在CPU或GPU上运行，编译的函数必须编译成两个不同的版本，一个带有CPU机器指令，另一个带有GPU机器指令。

在当前HPC SDK版本中，函数指针要么指向函数的CPU版本，要么指向函数的GPU版本。如果你试图在CPU和GPU代码之间传递函数指针，这会导致问题。你可能会无意中将一个指向函数的CPU版本的指针传递给GPU代码。在未来，可能会支持自动地、无缝地跨CPU和GPU代码边界使用函数指针，但目前的实现并不支持。

函数指针不能传递给在GPU上运行的c++并行算法，也不能通过GPU代码中的函数指针调用函数。例如，以下代码示例将无法正常工作:

void square(int& x) { x = x * x; }
void square_all(std::vector<int>& v) {
  std::for_each(std::execution::par_unseq,
                v.begin(), v.end(), &square);
}

上述示例将一个指向CPU版本的函数square指针传递给并行的for_each算法调用。当算法被并行化并转移到GPU上时，程序无法解析GPU版本的函数square指针。

通常可以通过使用函数对象（function object）来解决这个问题，函数对象是一个带有函数调用操作符的对象。函数对象的调用操作符在编译时解析为函数的GPU版本，而不是像上一个例子中那样在运行时解析为函数的错误CPU版本。例如，以下代码示例可以工作:

struct squared {
  void operator()(int& x) const { x = x * x; }
};
void square_all(std::vector<int>& v) {
  std::for_each(std::execution::par_unseq,
                v.begin(), v.end(), squared{});
}

另一种有效的的解决方案是将函数更改为lambda函数，因为lambda是作为一个无名称的函数对象实现的:

void square_all(std::vector<int>& v) {
  std::for_each(std::execution::par_unseq, v.begin(), v.end(),
                [](int& x) { x = x * x; });
}

如果函数太大，不能转换为函数对象或lambda，那么应该可以将对该函数的调用封装在lambda中:

void compute(int& x) {
// Assume lots and lots of code here.
}
void compute_all(std::vector<int>& v) {
  std::for_each(std::execution::par_unseq, v.begin(), v.end(),
                [](int& x) { compute(x); });
}

本例中没有使用函数指针。

不幸的是，通过函数指针调用函数的限制意味着目前不支持将多态对象从CPU代码传递到gpu加速的c++并行算法，因为虚函数是使用函数指针实现的。

6.4 随机访问迭代器

c++ Standard 要求传递给大多数c++并行算法的迭代器必须是前向迭代器。然而，gpu上的stdpar只能与随机访问迭代器一起工作。将前向迭代器或双向迭代器传递给gpu加速的c++并行算法会导致编译错误。将指向堆或标准库随机访问迭代器的原始指针传递给算法具有最佳性能，但大多数其他随机访问迭代器都可以正常工作。

6.5 与c++标准库的互用性

c++标准库的大部分内容可以与gpu上的stdpar一起使用。

std::atomic<T> 只要T是一个4字节或8字节的整数类型，GPU代码中的对象就可以运行。std::atomic<T>对象可以从CPU和GPU代码访问，只要该对象在堆上。
操作浮点类型的数学函数——例如sin、cos、log和<cmath>中声明的大多数其他函数——可以在GPU代码中使用，并解析为CUDA c++程序中使用的相同实现。
std::complex, std::tuple, std::pair, std::optional, std::variant和<type_traits>，都被支持，并在GPU代码中按预期工作。

c++标准库中GPU代码不支持的部分包括：I/O函数，以及通常访问CPU操作系统的任何函数。作为一种特殊情况，基本的printf调用可以在GPU代码中使用，并利用CUDA c++中使用的相同实现。

6.6 GPU code中没有exception

与大多数其他GPU编程模型一样，在装载到GPU的c++并行算法调用中不支持抛出和捕获c++异常。

与其他一些GPU编程模型(其中try/catch块和throw表达式是编译错误)不同的是，异常代码确实会编译，但具有非标准的行为。Catch子句会被忽略，而throw表达式在执行时会中止GPU内核。CPU代码中的异常工作不受限制。

7.一个大型的例子：LULESH

劳伦斯利弗莫尔国家实验室开发了LULESH流体动力学小应用程序，用于对编译器进行压力测试，并对其全尺寸生产流体动力学应用程序的预期性能进行建模。它大约有9000行c++代码，其中2800行是应该并行化的核心计算。该应用程序已经移植到许多不同的并行编程模型，包括MPI、OpenMP、OpenACC、CUDA c++、RAJA和Kokkos。

我们将LULESH移植到c++并行算法中，并在LULESH’s GitHub repository.中提供了该端口。要编译它，需要安装NVIDIA HPC SDK，check out LULESH的2.0.2-dev branch，转到正确的目录，并运行make。

git clone --branch 2.0.2-dev https://github.com/LLNL/LULESH.git
cd LULESH/stdpar/build
make run

（译者注：上面例子我在跑的时候出现如下异常

malloc: cuMemMallocManaged returns error code 2 for new pool allocation
malloc: cuMemMallocManaged returns error code 2 for new pool allocation
malloc: cuMemMallocManaged returns error code 2 for new pool allocation
malloc: cuMemMallocManaged returns error code 2 for new pool allocation
malloc: cuMemMallocManaged returns error code 2 for new pool allocation
malloc: cuMemMallocManaged returns error code 2 for new pool allocation
malloc: cuMemMallocManaged returns error code 2 for new pool allocation
malloc: cuMemMallocManaged returns error code 2 for new pool allocation
malloc: cuMemMallocManaged returns error code 2 for new pool allocation
(null): call to cuMemAllocManaged returned error 2: Out of memory
make: *** [Makefile:33: run] Aborted

意思大概是说爆内存了，但是我是64G内存，GPU显存也有48G，查了nvidia-smi没有显示显存不足，free命令也查了内存是够的，暂待研究吧，解决方法是将LULESH\stdpar\build\Makefile 的第33行改为

./lulesh2.0 -p -s 80 #原数是150，改成80

即可正常运行）

虽然LULESH太大了，无法在本文中显示完整的源代码，但这里有一些关键部分演示了stdpar的使用。
LULESH代码中有许多具有大量的循环，并且没有循环前后的依赖，因此具备并行化的条件。其中大部分循环都可以用std::execution::par策略轻松地转换为对std::for_each_n的调用，其中传递给std::for_each_n的lambda主体与原始循环主体是相同的。

函数CalcMonotonicQRegionForElems就是这方面的一个例子。如果用OpenMP编写，其循环头看起来像下面的代码例子。

#pragma omp parallel for firstprivate(qlc_monoq, qqc_monoq, \
                  monoq_limiter_mult, monoq_max_slope, ptiny)
  for ( Index_t i = 0 ; i < domain.regElemSize(r); ++i ) {

而stdpar中循环头如下：

std::for_each_n(
  std::execution::par, counting_iterator(0), domain.regElemSize(r),
    [=, &domain](Index_t i) {

循环体在本例中几乎有200行长，成为lambda的主体，但其他方面没有变化。
在一些地方，明确的for循环被改为使用C++并行算法，以更好地表达代码的意图，如函数CalcPressureForElems。

#pragma omp parallel for firstprivate(length)
for (Index_t i = 0; i < length ; ++i) {
  Real_t c1s = Real_t(2.0)/Real_t(3.0) ;
  bvc[i] = c1s * (compression[i] + Real_t(1.));
  pbvc[i] = c1s;
}

函数重写为如下所示：

constexpr Real_t cls = Real_t(2.0) / Real_t(3.0);
std::transform(std::execution::par,
  compression, compression + length, bvc,
  [=](Real_t compression_i) {
    return cls * (compression_i + Real_t(1.0));
  });
std::fill(std::execution::par, pbvc, pbvc + length, cls);

在一些情况下，代码清晰度的提高，对生产力的提高是非常显著的。函数CalcHydroConstraintForElems从一组值中找出最小值。

该函数的OpenMP版本长52行，使用一个临时数组来保存每个线程找到的最小值。这个函数可以用标准C++编写，并在GPU上并行运行，只需调用std::transform_reduce，让编译器来处理所有做并行还原的复杂问题。这一改动将之前的52行缩减为以下12行。

dthydro = std::transform_reduce(std::execution::par,
  counting_iterator(0), counting_iterator(length),
  dthydro, [](Real_t a, Real_t b) { return a < b ? a : b; },
  [=, &domain](Index_t i) {
     Index_t indx = regElemlist[i];
     if (domain.vdov(indx) == Real_t(0.0)) {
      return std::numeric_limits<Real_t>::max();
     } else {
      return dvovmax /
             (std::abs(domain.vdov(indx)) + Real_t(1.e-20));
     }
  });

当LULESH的C++并行算法版本为单个A100 GPU运行时，应用程序的运行速度几乎是在双插槽Skylake系统的所有40个CPU核心上编译运行时的7倍，如图1所示。A100上的并行算法版本的性能与A100上的OpenACC版本几乎相同。
Figure 1.LULESH的不同实现的相对性能。左栏是标准C++并行算法版本，在40核双插槽Skylake系统的所有CPU核心上运行。其他两列是在单个A100 GPU上运行的。中间一列是标准C++并行算法版本，与左列的代码相同。右栏是LULESH的OpenACC版本。

8.开始学习C++的GPU并行算法

要想开始使用，请下载NVIDIA HPC SDK并将其安装在基于x86-64、OpenPOWER或Arm CPU的系统上，该系统运行支持的Linux版本。

英伟达HPC SDK可以免费下载，其中包括所有英伟达注册开发者的永久使用许可，包括访问未来发布的更新版本。在你的系统上安装了NVIDIA HPC SDK之后，NVC++编译器就可以在/opt/nvidia/hpc_sdk目录结构下使用。

要在Linux/x86-64系统上使用编译器，包括NVC++，请将目录/opt/nvidia/hpc_sdk/Linux_x86_64/20.5/compilers/bin添加到你的路径中。（译者注：其中20.5应替换为你所下载的HPC SDK版本号，这个blog是我见到的唯一一个路径对的，泪目）
在基于OpenPOWER或Arm CPU的系统上，将Linux_x86_64分别替换为Linux_ppc64le或Linux_aarch64。

8.1支持的NVIDIA GPU

NVC++编译器可以自动将C++并行算法卸载到基于Volta、Turing或Ampere架构（及以后）的NVIDIA GPU上。这些架构包括独立的线程调度和针对CUDA统一内存的硬件优化等功能，这些功能是专门为支持像C++并行算法这样的高性能、通用的并行编程模型而设计。

NVC++编译器对Pascal架构上的C++并行算法提供了有限的支持，Pascal架构不具备正确支持std::execution::par策略所需的独立线程调度。当为Pascal架构（-gpu=cc60）编译时，NVC++为CPU编译具有std::execution::par策略的算法。只有具有std::execution::par_unseq策略的算法可以在Pascal GPU上运行。

默认情况下，NVC++会自动检测并为编译器所运行的系统上安装的GPU类型生成GPU代码。要为特定的GPU架构生成代码，这在应用程序被编译并在不同的系统上运行时可能是必要的，添加-gpu=ccXX命令行选项。目前，编译器只能生成针对一种GPU架构的可执行文件。在一次编译中使用多个-gpu=ccXX选项会导致编译器出现错误。
（译者注：这个ccXX代表计算架构，Compile for compute capability X.Y; supported values: 35,50,60,61,62,70,72,75,80,86,87,90）

8.2支持的CUDA版本

NVC++编译器是建立在CUDA库和技术之上的，它使用CUDA将C++并行算法加速NVIDIA GPU。运行NVC++编译的应用程序的GPU加速系统必须安装有CUDA 10.1或更新的设备驱动程序。

英伟达HPC SDK编译器附带有集成的CUDA工具链、头文件和库，可在编译过程中使用，因此不需要在系统上安装CUDA工具包。

当指定-stdpar时，NVC++使用与编译时系统上安装的CUDA驱动相匹配的CUDA工具链版本进行编译。要使用不同版本的CUDA工具链进行编译，请使用-gpu=cudaX.Y选项。例如，使用-gpu=cuda11.0选项来指定你的程序应该使用11.0工具链为CUDA 11.0系统进行编译。

9总结

最近发布的NVIDIA HPC SDK中包含的NVC++编译器使用户首次能够使用完全标准和完全可移植的C++结构对NVIDIA GPU进行编程。以适当的执行策略为工具的C++并行算法调用会自动并行化并卸载到NVIDIA GPU上。由此产生的编程环境提高了GPU程序员的工作效率，为新用户提供了进入GPU计算的便捷通道，并在各种标准C++实现中实现了性能移植。

c++算法并行化加速探索者

关注

1
点赞
踩
7

收藏

觉得还不错? 一键收藏
0
评论
【NVIDIA HPC SDK】使用NVIDIA GPU并行化C++代码

最近发布的NVIDIA HPC SDK中包含的NVC++编译器使用户首次能够使用完全标准和完全可移植的C++结构对NVIDIA GPU进行编程。以适当的执行策略为工具的C++并行算法调用会自动并行化并卸载到NVIDIA GPU上。由此产生的编程环境提高了GPU程序员的工作效率，为新用户提供了进入GPU计算的便捷通道，并在各种标准C++实现中实现了性能移植。
复制链接

扫一扫