openmp并行编程_[编程实践] OpenMP 与 NVIDIA

引言

OpenMP 是基于指令的并行编程的主要标准。NVIDIA 于 2011 年加入 OpenMP，致力于围绕并行加速器的讨论。NVIDIA在2012年 OpenMP 4.0 中提出了针对加速器的TEAMS构造，并于2013 年发布了对加速器的支持。

OpenMP 轻量易用，可轻松构建多线程程序。本文介绍如何利用 OpenMP 执行 GPU 并行计算。

NVIDIA在 OpenMP 上支持的编译器

• Clang
• XL
• GCC
• Cray Compiler Environment

注意，Apple 在 macOS 上提供的 Clang 并不自带 OpenMP 支持，GNU GCC for macOS 支持 OpenMP，一般需要在 g++ 编译参数中添加 -fopenmp.

OpenMP PARALLEL FOR

在一组线程中并行执行 for 循环的迭代。这是 OpenMP 在 CPU 上执行多线程计算的基本用法之一。

#pragma omp parallel for
for (int i = 0; i < N; i++) {
    p[i] = v1[i] * v2[i];
}

OpenMP TARGET 指令

• target 指令提供一种机制，将线程执行的指令从 CPU 移动到 target，并重新分配所需的数据。它针对一个线程组的线程执行 for 循环，线程组在下节介绍。
• 几乎所有的 OpenMP 可以被用在 target 区域内，但只有其中一部分在 GPU 上执行有意义。
• target 区域中所有的标量都被设置为 firstprivate，即多线程执行前首先通过 target 区域外的代码对标量进行初始化。
• terget 区域中所有的数组都会在 CPU 和 device 之间拷贝交换。

TEAM 指令解决的问题

一个关键问题是，OpenMP 原本是为共享内存并行计算机的多线程设计的，所以 parallel 指令仅创建一个并行层。

线程必须要同步执行，意味着在 GPU 上它们需要一个线程块(thread block)。

team 指令被用于表达第二个层次的可扩展并行化。

OpenMP TEAMS 指令

为了更好地利用 GPU 资源，通过 TEAMS 指令来使用大量线程组。

• 代码对于一个线程组或者多个线程组是可移植的。
• 程序员不需要思考如何将循环分解！
• 产生1个或更多具有相同线程数的线程组
• 在每一组的主线程上继续执行（冗余）
• 线程组之间没有同步

OpenMP TEAMS - DISTRIBUTE 指令

将循环的迭代分发到线程组的主线程。

• 迭代以静态方式分发。
• 对线程组执行的顺序没有保证。
• 对线程组是否同时执行没有保证。
• 在线程组内不产生并行化。

注意，使用 OpenMP 操作 GPU 时，总是使用 teams 和 distribute 来使整个 GPU 并行化。

扩展并行化到内循环

但是，上述并行化都是在外循环，我们能否将并行化扩展到内循环呢？两种方案：

• 将线程组分发进一步分解
• collapse 语句

在应用线程组和线程并行化到两层循环前，分解两层循环：

以上方法的性能对比：

TARGET DATA 指令

• 每次循环都在 CPU 和 GPU 之间交换数据是不高效的。
• target data 指令和 map 语句允许我们控制数据移动。

map()…

• to – 在 GPU 上创建空间并复制输入数据
• from – 在 GPU 上创建空间并复制输出数据
• tofrom – 在 GPU 上创建空间并复制输入输出数据
• alloc – 在 GPU 上创建空间，不复制数据

在循环之外移动数据，以便在两个 target 区域共享数据。

主机回退 Host Fallback

if子句将在何处运行循环的决定推迟到运行时，并强制同时构建主机和设备版本。

CUDA 互操作性

• OpenMP 是高级语言，有时为了达到最佳性能，低层优化是必须的。
• CUDA内核或加速库是一个典型例子
• use_device_ptr映射类型允许将OpenMP设备数组传递给CUDA或加速库。
• is_device_ptr 映射子句允许在OpenMP的 target 区域内使用CUDA数组。

以下是使用示例。

OpenMP 任务

• OpenMP 任务允许程序员表示独立的工作块，并允许运行时调度它们。
• 所有的 OpenMP target 区域都是任务
  • 默认情况，与主机同步
  • nowait 子句实现异步
  • 允许 depend 子句来与其他任务交互
• 使用OpenMP的nowait和depend子句，可以进行异步数据传输和内核启动来提高系统利用率

异步流水线的例子：

OpenMP 任务图与 CUDA 流

• CUDA Streams
  • 简单映射到硬件
  • 开发者将依赖显式映射到 CUDA 流
• OpenMP 任务图
  • 可能开销更大
  • 任务图必须在运行时映射到流
  • 开发者在不同任务间表达依赖

OpenMP 在 GPU 计算的最佳实践

• 总是使用 teams 和 distribute 指令来充分利用并行化
• collapse 循环以提升并行效率
• 使用 target data 指令和 map 子句来减少 CPU 和 GPU 之间的数据流动
• 可能情况下，使用加速库
• 使用 OpenMP 进行异步计算，更好利用整个系统
• 使用主机回退来生成主机和设备代码

其他