引言
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 进行异步计算,更好利用整个系统
- 使用主机回退来生成主机和设备代码