NVIDIA HPC SDK v25.7 为采用 GPU 加速的高性能计算（HPC）应用开发者带来了重大突破。该版本凝聚了近两年在统一内存编程方面的持续研发成果，最终构建出一套完整的工具集，能够自动实现 CPU 与 GPU 之间的数据迁移。通过大幅减少传统开发中繁琐的手动数据管理，该版本显著简化了 GPU 开发流程，缩短了 GPU 移植周期，降低了出错风险，同时使开发者能够更灵活地优化科学计算工作负载。
功能详情
在高性能计算（HPC）市场，采用紧密耦合的CPU与GPU架构并实现两者间统一地址空间的NVIDIA一致性计算平台正日益受到青睐。其中，NVIDIA GH200 Grace Hopper超级芯片和NVIDIA GB200 NVL72系统是该架构的代表产品，目前已部署于瑞士国家超级计算中心（CSCS）的ALPS超级计算机以及德国于利希超级计算中心（JSC）的JUPITER系统中。
这些架构不仅凭借 CPU 与 GPU 之间高带宽的 NVLink-C2C 互连实现了卓越的性能，还能显著提升开发者的生产力。得益于共享地址空间，程序员无需再手动管理 CPU 与 GPU 之间的数据传输，数据移动可由 NVIDIA CUDA 驱动程序自动完成。这种简化已在实际项目中展现出显著价值。
“借助统一内存编程，我们能够更快速地将NEMO海洋模型移植到GPU上。与传统方法相比，这种方法还让我们能够灵活地尝试将更多工作负载运行在GPU上。”巴塞罗那超级计算中心（BSC）高级研究工程师Alexey Medvedev表示。
自2023年底NVIDIA推出Grace Hopper超级芯片架构以来，NVIDIA HPC SDK逐步增强了对统一内存编程的支持。其中部分功能已在博客文章《使用NVIDIA Grace Hopper超级芯片简化HPC的GPU编程》中介绍。
NVIDIA HPC SDK 25.7 版本推出了一套完整的工具集，能够简化甚至在多数情况下完全消除 CPU 与 GPU 之间的手动数据传输，显著提升了科学应用开发者的生产力。
数据管理被广泛认为是 GPU 编程中最具挑战性的环节之一。确保 CPU 与 GPU 子程序之间实现正确且高效的数据传输，通常需要经过多轮调试与优化。在许多高性能计算（HPC）应用中，动态分配的数据结构以及复合类型的使用进一步加剧了这一复杂性。以下 Fortran 示例展示了一种常见模式，称为“深度复制”：为了正确并行化访问派生类型中可分配数组成员的循环，必须引入一个额外的循环，专门用于通过 OpenACC 指令管理数据传输。
type mytype
integer, allocatable::arrmem(:)
integer :: scalarmem
end type mytype
! Declare a new array of derived type - mytype.
type(mytype)::base(:)
…
! This loop is exclusively for copying data to the GPU.
!$acc enter data copyin(base(:))
do i=1,N
!$acc enter data copyin(base(i)%arrmem(:))
end do
! Parallel OpenACC loop accessing members of derived type mytype
! in each array element of base.
!$acc parallel loop collapse(2) default(present)
do i=1,N
do j=1,M
base(i)%arrmem(j)=base(i)%scalarmem
end do
end do在现实世界的高性能计算（HPC）应用中，循环通常会访问多个数组，这些数组可能嵌入在派生类型中，也可能在模块或公共块中声明。在大多数实际代码中，管理数据移动往往非常复杂，其难度在许多情况下甚至超过了识别和标注并行循环本身。理解数据依赖关系、确保数据正确传输，并避免内存泄漏或竞争条件，这些都给GPU编程带来了显著的开销。
Grace Hopper 及类似的 GPU 架构采用了统一内存模型，能够消除大部分复杂性。由于 CPU 和 GPU 共享同一个地址空间，通常不再需要显式地进行数据管理。如上所述，此前涉及数组派生类型的示例现在可以直接实现并行化，而无需额外的数据移动指令。
type mytype
integer, allocatable::arrmem(:)
integer :: scalarmem
end type mytype
! Declare a new array of derived type - mytype.
type(mytype)::base(:)
…
! Parallel OpenACC loop accessing members of derived type mytype
! in each array element of base.
!$acc parallel loop collapse(2)
do i=1,N
do j=1,M
base(i)%arrmem(j)=base(i)%scalarmem
end do
end do将某些 C++ 代码卸载到 GPU 后，数据管理的挑战变得更加严峻。广泛使用的面向对象抽象和数据封装通常会限制开发者对内部实现的访问，从而难以将数据复制到 GPU。例如，在没有统一内存的情况下，使用 OpenACC 的代码若涉及 std::vector，便无法正确执行。此外，在循环中使用 std::vector 的下标运算符时，需要访问 std::vector 类中的数据及其所分配的元素数据，而这些数据在内存中与 std::vector 类本身并不连续，进一步增加了数据管理的复杂性。
std::vector<int> v(N);
#pragma acc kernels
for (i = 0; i < v.size(); i++)
v[i] = i;
如果没有统一内存，仅在示例中的内核构造中添加 `copyout(v` 子句是不够的。此时，只有 `type mytype
integer, allocatable::arrmem(:)
integer :: scalarmem
end type mytype
! Declare a new array of derived type - mytype.
type(mytype)::base(:)
…
! This loop is exclusively for copying data to the GPU.
!$acc enter data copyin(base(:))
do i=1,N
!$acc enter data copyin(base(i)%arrmem(:))
end do
! Parallel OpenACC loop accessing members of derived type mytype
! in each array element of base.
!$acc parallel loop collapse(2) default(present)
do i=1,N
do j=1,M
base(i)%arrmem(j)=base(i)%scalarmem
end do
end do0` 对象本身被复制，而其内部包含的元素并不会被自动复制。因此，这类代码通常需要重写为直接操作向量元素的原始指针，从而退化为非面向对象的编程风格。此外，由于缺乏访问标准 C++ 容器元素的统一接口，无法将大多数其他 STL 容器中的数据直接复制到 GPU。
std::vector<int> v(N);
auto ptr = v.data();
#pragma acc kernels copyout(ptr[0:N])
for (i = 0; i < v.size(); i++)
ptr[i] = i;
欧洲海洋建模框架（NEMO）是一个先进的建模系统，广泛应用于海洋与气候科学的研究以及预报服务。
在 BSC 开始将 NEMO 移植到 GPU 之前，我们对公开可用的代码库进行了内部评估，以探索采用统一内存编程可能带来的简化优势。
在 NVIDIA GTC 大会上，我们以这段真实代码作为案例研究，在题为“通过 NVIDIA Grace Hopper 平台加速科学工作流”和“使用 NVIDIA HPC 编译器加速科学计算”的演讲中，展示了在 Grace Hopper 等一致性系统上，开发者的生产力可得到显著提升。
统一内存消除了对显式数据管理代码的需求，使开发者能够更加专注于并行化设计。代码量减少后，开发者在 GPU 移植的早期阶段即可观察到性能提升。本演示基于 ORCA ½ 网格，采用了 NEMO v4.2.0 中的 GYRE_PISCES 基准测试进行验证。
这是一个内存带宽受限的基准测试，最初采用 MPI 在多核 CPU 上实现并行化，其主要性能热点在于主动和被动示踪剂的扩散与输送过程。我们针对这些关键部分，仅通过 OpenACC 对性能敏感区域中的循环进行并行化，而将内存管理交由 CUDA 驱动程序和硬件自动处理。
我们最初将代码移植到 GPU 的策略如下，且未添加 GPU 数据管理相关代码：
- 使用
!$acc parallel loop gang vector collapse()对完全并行的紧密嵌套循环进行并行化。 - 具有交叉交互依赖关系的循环使用
!$acc loop seq进行标注。 - 数组表示法中的操作被封装在
!$acc kernels中。 → 封装在 tg_14 中的是数组表示法的相关操作。 或更自然地表达为： 数组表示法中的各项操作均已封装在 tg_14 中。 （推荐使用后者，语义清晰、结构更流畅。） - 并行循环中调用的外部例程已使用
!$acc routine seq进行标注。
由于 NEMO 的结构特点，许多函数包含多个并行区域，且这些区域通常是连续排列的。因此，每个 OpenACC 并行构造结束时的隐式同步会对性能产生影响。这种同步机制的引入旨在简化编程，并确保并行化代码与顺序代码的执行顺序保持一致，从而与原始的非并行代码行为相符。
为避免这些隐含的障碍，在 parallel 和 kernels 构造中引入了 async clauses，以提升并发性。在异步模式下执行并行区域时，通过 !$acc wait 引入同步机制，确保在后续 MPI 调用之前，GPU 上的计算数据已就绪，或防止局部变量在子例程结束前超出作用域。
以下代码片段来自公共开源 NEMO 资源库中的 trazdf.f90 文件，用于演示前述的 OpenACC 并行化策略。
SUBROUTINE tra_zdf_imp(...)
...
REAL(wp), DIMENSION(ntsi-(nn_hls):ntei+(nn_hls),ntsj-(nn_hls):ntej+(nn_hls),jpk) :: &
& zwi, zwt, zwd, zws
...
DO jn = 1, kjpt
...
!* 3d recurrence: Xk = ( Zk - Sk Xk+1 ) / Tk (result is the after tracer)
! Fully parallel collapsed tightly nested OpenACC loops
!$acc parallel loop gang vector collapse(2) async(1)
DO jj = ntsj, ntej
DO ji = ntsi, ntei
pt(ji,jj,jpkm1,jn,Kaa) = pt(ji,jj,jpkm1,jn,Kaa)/zwt(ji,jj,jpkm1)*tmask(ji,jj,jpkm1)
END DO
END DO
!$acc end parallel
! As above OpenACC parallel loop is with async clause,
! no synchronization with the CPU here
!$acc parallel async(1)
! Sequential OpenACC loop due to vertical dependencies
!$acc loop seq
DO jk = jpk-2, 1, -1
! Fully parallel collapsed tightly nested OpenACC loops
!$acc loop gang vector collapse(2)
DO jj = ntsj, ntej
DO ji = ntsi, ntei
pt(ji,jj,jk,jn,Kaa) = (pt(ji,jj,jk,jn,Kaa)-zws(ji,jj,jk)*pt(ji,jj,jk+1,jn,Kaa)) &
& /zwt(ji,jj,jk)*tmask(ji,jj,jk)
END DO
END DO
END DO
!$acc end parallel
! As above OpenACC parallel region is with async clause,
! no synchronization with the CPU here.
...
END DO
!$acc wait
! As OpenACC wait enforces synchronization with the CPU,
! the CPU waits here for all work to be completed on the GPU.
END SUBROUTINE tra_zdf_imp进一步优化异步执行时，代码的某些部分可能出现数据竞争问题，这一点在 GPU 内核异步访问共享数据而 CPU 已退出本地分配共享数据的函数时尤为明显。幸运的是，2025 年 6 月发布的 OpenACC 3.4 规范在现有数据子句中引入了 capture 修饰符，可有效解决此类竞争条件。更多详情请参阅 ISC 2025 上发布的 OpenACC 3.4 规范。
在 HPC SDK 25.7 中，开发者能够在无需重构应用程序大部分代码的情况下，安全地管理异步操作中的数据。此外，HPC SDK 的新功能不仅致力于消除数据竞争，还着力应对检测数据竞争这一棘手挑战。
在 NEMO 的 GPU 端口中，所有内存管理均由 CUDA 驱动程序自动处理。在 Grace Hopper 等一致性内存平台上，系统分配的内存可被 CPU 和 GPU 同时访问，内存页面的物理位置遵循“先触及”原则。自 CUDA 12.4 起，引入了基于访问计数器的启发式算法，能够自动将被频繁访问的 CPU 内存页面迁移到 GPU 内存中。更多信息请参阅 CUDA Toolkit 12.4，该版本进一步增强了对 NVIDIA Grace Hopper 架构以及机密计算的支持。
这些自动迁移机制有助于提升数据局部性和系统性能。如“使用异构内存管理简化 GPU 应用开发”所述，在基于 x86 架构的集成 GPU 系统上，若 Linux 内核版本足够新，也支持类似功能。
得益于高带宽互连和专用硬件一致性机制，Grace Hopper 系统预计将在性能上优于依赖异构内存管理（HMM）的统一内存系统。然而，基于 HMM 的系统仍在持续演进，通过不断优化内核与驱动程序，旨在降低延迟，并提升页面迁移在更广泛 GPU 加速平台上的效率。
下图展示了 NEMO 海洋模型在单个 Grace Hopper 超级芯片上完成一个时间步的增量移植过程。该过程从 Grace 处理器的多核 CPU 执行开始，逐步将计算任务迁移至 Hopper GPU，仅需如前所述，通过添加 OpenACC 指令对循环进行标注即可实现。
我们对主动和被动示踪剂的水平扩散（步骤1）、示踪剂平流（步骤2）以及垂直扩散与时间过滤（步骤3）进行了代码移植。在每部分移植后的代码于Hopper平台上运行时，相比在Grace核心上的执行，性能提升了约2至5倍；在部分加速的端到端仿真中，整体性能提升约为2倍。
即使在移植的早期阶段，我们已观察到部分 GPU 加速工作负载实现了端到端的性能加速。通过逐步将更多组件卸载到 GPU，可进一步提升仿真性能。统一内存机制简化了内存管理，使频繁访问的 CPU 页面能够自动迁移至 GPU 内存，从而加快 GPU 内核的执行速度。得益于高效的远程访问链路，即使 CPU 组件访问位于 GPU 内存中的数据，其性能仍能保持在较高水平。由于 NEMO 是基于 MPI 的代码，我们在 Grace 平台上采用多进程方式以充分饱和 CPU 内存带宽和 NVLink-C2C 带宽，同时在 Hopper 上启用 MPS（多进程服务）以降低上下文切换开销。
总体而言，以相对较少的投入即实现了显著的性能提升，凸显了统一内存在加速 GPU 科学计算代码方面的巨大潜力。即使在部分 GPU 加速的应用中，开发者也能在移植过程的早期阶段获得可观的加速效果，这极大地改善了开发体验。尽管如此，优化空间依然存在，我们相信通过持续的调优和开发，性能提升还有望进一步增强。
立即下载 NVIDIA HPC SDK，借助 OpenACC 和统一内存加速您的应用程序。有关当前功能、限制以及即将推出的更新详情，请参阅 NVIDIA HPC SDK 文档。
