oneAPI GPU 优化指南 - 分配USM(Unified Shared Memory)

 本章节翻译by weavingtime@formail.com 原文 Unified Shared Memory Allocations (intel.com)

USM(Unified Shared Memory)允许程序使用 C/C++ 指针进行内存访问。 SYCL中有三种分配内存的方式:

malloc_device:

  • 分配的内存只能由指定设备访问,而不能由上下文中的其他设备或主机访问。

  • 数据始终保留在设备上,因此是kernel执行最快的选择。

  • 需要显式复制才能将数据传输到上下文中的主机或其他设备。

malloc_host:

  • 主机和上下文中的任何其他设备都可以访问分配的内存。

  • 数据始终保留在主机上,并通过 PCI 从设备进行访问。

  • 与主机或设备同步数据不需要显式地拷贝。

malloc_shared:

  • 分配的内存只能由主机和指定设备访问.

  • 数据可以在主机和设备之间迁移(由Level Zero驱动程序操作),以实现更快的访问。

  • 主机和设备之间的同步不需要显式复制,但上下文中的其他设备如果要访问则需要显示复制。

三种内存分配方式及其特性总结如下表。

内存分配方式及其特性

内存分配方式

描述

主机是否可以访问

设备是否可以访问

内存所在位置

主机

在主机的内存中分配

是, 可以远程通过PCIe或者fabric link访问

主机

设备

在设备内存中分配

设备

共享

在主机和设备共享的内存中分配

在主机和设备之间动态迁移

在多stack、多 slice GPU 环境中,需要注意的是,设备和共享 USM 分配与根设备相关联。 因此,同一设备上的所有stack和 slice都可以访问它们。 程序应使用根设备进行 malloc_device 和 malloc_shared 分配以避免混淆。

OpenMP中分配USM的API

为了与 SYCL USM 模型保持一致,我们添加了三个新的 OpenMP API 作为Intel的 扩展API,供用户根据应用程序、内存大小和性能要求执行内存分配。 它们的语义和性能特征在以下小节。

主机内存分配

这种主机内存比设备内存更容易使用,因为我们不必在主机和设备之间手动复制数据。 主机内存从主机内存中分配,主机和设备均可访问。 这些内存虽然可以在设备上访问,但无法迁移到设备的附加内存。 相反,部署到设备上的kernel从该内存读取或写入该内存的区域是通过 PCIe 总线或结构链路”远程”完成的。 便利性和性能之间的权衡是我们必须考虑的。尽管主机内存可能会产生更高的访问成本, 但仍然有充分的理由使用它们。 比如很少访问的数据或无法容纳在设备附加内存中的大型数据集。 执行主机内存分配的API是

extern void *omp_target_alloc_host(size_t size, int device_num)

设备内存分配

用户需要设备内存这种方式,以便拥有指向设备附加内存的指针,例如设备上的 (G)DDR 或 HBM。 设备内存分配可以通过部署到设备上运行的kernel来读取或写入,但不能从主机上执行的代码直接访问它们。 在主机上尝试访问设备内存可能会导致数据不正确或程序崩溃。 执行设备内存分配的API是

extern void *omp_target_alloc_device(size_t size, int device_num)

共享内存分配

与主机内存一样,共享内存可以在主机和设备上访问。 它们之间的区别在于,共享内存可以自动在主机内存和设备附加内存之间自由迁移, 无需我们干预。 如果内存已迁移到设备,则在该设备上执行的任何kernel访问该设备内存的性能都将比从主机远程访问它的性能更高。 然而,共享分配有很多好处,它也是有缺点的,例如页面迁移成本和乒乓效应

extern void *omp_target_alloc_shared(size_t size, int device_num)

omp_target_alloc API对USM的支持

用于目标内存分配的OpenMP API是基于 “#pragma omp requires unified shared memory” 实现的,并且使用环境变量 LIBOMPTARGET_USM_HOST_MEM 来支持三种内存分配方式. 其实现逻辑见下面的代码.

1if (exists(“#pragma omp requires unified_shared_memory”)) {
2    if (LIBOMPTARGET_USM_HOST_MEM == 1)
3        return "host memory";
4    else
5        return "shared memory";
6} else {
7     return "device memory";
8}

上一章                                   上级目录 主目录                                                                  下一章   

<think>好的,我需要帮助用户检测当前OpenMP环境是否支持Unified Shared Memory (USM)特性。用户之前已经问过关于requires指令的问题,现在想具体知道如何检测USM的支持情况。 首先,我记得用户的上一次回答中提到了requires指令的使用场景,包括USM。根据引用[2],当使用`#pragma omp requires unified_shared_memory`时,内存分配方式会根据环境变量变化。这可能意味着检测USM支持的方法可能涉及检查编译器是否支持该指令,或者运行时是否存在相关环境变量。 用户想检测当前环境是否支持USM,可能需要几种方法: 1. **编译时检查**:使用预处理指令查看是否定义了相关的宏。例如,某些编译器在支持USM时会定义特定的宏,用户可以通过预处理器命令检查这些宏是否存在。比如用`gfortran`的`-dM -E`选项输出定义的宏,然后过滤查找是否有与USM相关的定义。 2. **运行时检测**:根据引用[2]中的代码逻辑,如果代码中使用了`requires unified_shared_memory`,并且设置了环境变量`LIBOMPTARGET_USM_HOST_MEM`,则内存分配方式会不同。用户可以在程序中通过条件编译或运行时检查环境变量是否存在,或者尝试分配内存并检查其属性,例如使用`omp_target_alloc`并查看返回的指针是否在主机和设备间共享。 3. **编译器文档或版本检查**:查阅编译器文档,确认支持的OpenMP版本和特性。例如,OpenMP 5.0及以上版本支持USM,所以检查编译器是否支持该版本。 另外,用户可能需要考虑不同编译器的差异,比如GCC、Clang或Intel编译器,它们在处理OpenMP扩展时可能有不同的实现。比如,Intel扩展提到了`omp_target_get_context()`,但用户的问题更关注通用的检测方法。 可能的步骤包括: - 预编译时检查编译器是否支持`requires`指令和USM子句。 - 在程序中加入条件代码,根据内存分配的结果判断USM是否生效。 - 使用环境变量来影响内存分配路径,通过输出或日志确定当前使用的内存类型。 还需要注意错误处理,比如如果环境不支持USM,程序可能会在运行时抛出错误,用户可以通过捕获这些错误或检查返回指针的有效性来判断。 总结下来,检测方法可能分为编译时和运行时两种,结合预处理宏、环境变量设置以及动态内存分配测试。</think>#### 一、编译时检测方法 1. **预处理检查** 使用编译器预处理器验证是否支持`unified_shared_memory`特性: ```bash # 检查GCC/GFortran是否支持requires子句 gfortran -fopenmp -dM -E - < /dev/null | grep -i "OPENMP_REQUIRES_UNIFIED_SHARED_MEMORY" ``` 若输出相关宏定义(如`_OPENMP_USM`),则表明编译器支持该特性[^1]。 2. **版本兼容性验证** OpenMP 5.0+ 规范正式支持USM特性。通过检查编译器版本确认支持性: ```bash # 查看GCC支持的OpenMP版本 gfortran -fopenmp -dumpversion # 输出结果需 >= 5.0 ``` #### 二、运行时检测方法 1. **环境变量分析** 根据引用[^2]的实现逻辑,通过设置`LIBOMPTARGET_USM_HOST_MEM`环境变量观察行为差异: ```bash # 设置环境变量(可选值为0/1) export LIBOMPTARGET_USM_HOST_MEM=1 ``` 结合以下代码测试内存分配类型: ```c #include <stdio.h> #include <omp.h> int main() { #pragma omp requires unified_shared_memory void *ptr = omp_target_alloc(1024, omp_get_default_device()); if (ptr == NULL) { printf("USM not supported: 设备内存分配失败\n"); } else { // 检查指针是否可被主机直接访问 #pragma omp target is_device_ptr(ptr) { /* 空操作,仅验证指针可传递性 */ } printf("USM supported: 内存类型 = %s\n", (ptr == omp_get_mapped_ptr(ptr, 0)) ? "统一内存" : "设备内存"); omp_target_free(ptr, omp_get_default_device()); } return 0; } ``` 若输出"统一内存"且无运行时错误,则表明环境支持USM[^2]。 2. **API功能探测** 通过Intel扩展`omp_target_get_context()`查询上下文属性: ```c #include <omp.h> int main() { omp_interop_t ctx; #pragma omp requires unified_shared_memory #pragma omp interop init(targetsync: ctx) int support = omp_get_interop_int(ctx, omp_ipr_memory_support, 0); if (support & omp_ipr_memory_support_unified_shared) { printf("USM特性已激活\n"); } #pragma omp interop destroy(ctx) return 0; } ``` #### 三、综合验证流程 | 步骤 | 方法 | 验证目标 | |------|------|----------| | 1 | 检查编译器版本 | 确认OpenMP ≥5.0支持 | | 2 | 预处理宏检查 | 确认语法支持性 | | 3 | 运行时内存分配测试 | 验证实际功能可用性 | #### 四、常见错误排查 - **编译错误**:若提示`unified_shared_memory子句未定义`,需升级编译器版本或启用实验性OpenMP支持(如Clang需添加`-fopenmp-version=51`)。 - **运行时错误**:若返回`TARGET DEVICE NOT SUPPORTED`,需检查设备驱动和OpenMP运行时库版本。 --- ### 相关问题 1. 如何在混合编程中结合OpenMP USM与CUDA统一内存? 2. 使用`unified_shared_memory`时内存一致性如何保证? 3. OpenMP的USM与SYCL的USM有何异同? [^1]: OpenMP上下文隐式管理机制与Intel扩展接口 : USM实现依赖环境变量与编译指令的交互逻辑
评论
成就一亿技术人!
拼手气红包6.0元
还能输入1000个字符
 
红包 添加红包
表情包 插入表情
 条评论被折叠 查看
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值