企业官网建设流程全解析

CANN Runtime 流内存操作接口详解：aclrtValueWrite 与 aclrtValueWait 的异步内存值写入与等待

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime

本篇文章聚焦 CANN Runtime 提供的流内存操作（Stream Memory Operation）接口aclrtValueWrite与aclrtValueWait，讲解如何在指定 Stream 上异步地向 Device 内存写入数据、以及等待内存数据满足条件后解除阻塞。文章以 11-09 流内存操作 为骨架，结合 仓库源码实现、单元测试 与 多流内存同步样例，帮助你掌握这两个接口的参数语义、比较模式、底层调用链以及基于内存语义实现跨流同步的完整实战方法。

一、流内存操作概述

在 CANN Runtime 中，Stream 上的任务通常按提交顺序执行。当需要跨 Stream 协调算子的执行节奏时，传统方案是使用 Event / Notify 同步机制。而内存语义同步（Memory Semantic Synchronization）提供了另一条更灵活、更底层的路径：它基于通用 Device 内存实现同步，与 Event/Notify 机制最大的不同在于——算子本身可以作为同步参与方，即算子可以在执行过程中与另一条流进行同步（详见 内存语义同步指南）。

流内存操作接口正是实现这种内存语义同步的基础设施，由两个配套接口组成：

两者都以任务（Task）的形式下发到指定的 Stream，由 Stream 调度执行。aclrtValueWait负责"阻塞等待"，aclrtValueWrite（或其他算子对内存的写入）负责"写值唤醒"，一写一等形成完整的同步闭环。

产品支持情况

由于该机制依赖底层硬件指令能力，并非所有产品都支持。两个接口的产品支持情况一致，汇总如下：

说明：在 架构 5162 不支持接口清单 中同样登记了这两个接口，与文档中部分老型号产品不支持的信息相互印证。使用时请先确认目标设备的支持情况。

二、aclrtValueWrite：向指定内存异步写值

2.1 函数原型与功能说明

aclError aclrtValueWrite(void* devAddr, uint64_t value, uint32_t flag, aclrtStream stream)

aclrtValueWrite向指定的 Device 内存地址写入数据，是一个异步接口：接口返回后，写值任务尚未执行完成，实际写入动作发生在 Stream 上任务调度执行时。在内存语义同步场景中，它通常被用来"写值唤醒"另一条 Stream 上阻塞的aclrtValueWait任务（或正在轮询该内存的算子）。

2.2 参数说明

2.3 返回值说明

返回0表示成功（ACL_SUCCESS），返回其他值表示失败，具体错误码请参见 aclError。

2.4 源码实现与调用链

在 memory.cpp 中，aclrtValueWrite的实体实现aclrtValueWriteImpl的调用逻辑非常清晰：

aclError aclrtValueWriteImpl(void* devAddr, uint64_t value, uint32_t flag, aclrtStream stream) { ACL_PROFILING_REG(acl::AclProfType::AclrtValueWrite); ACL_LOG_INFO("start to execute aclrtValueWrite"); ACL_REQUIRES_NOT_NULL_WITH_INPUT_REPORT(devAddr); // 校验 devAddr 非空 ACL_REQUIRES_RTS_OK(rtsValueWrite(devAddr, value, flag, static_cast<rtStream_t>(stream))); return ACL_SUCCESS; }

从源码结构可以看出完整的调用链：

aclrtValueWrite (对外 API，头文件声明于 include/external/acl/acl_rt.h) └─> aclrtValueWriteImpl (src/acl/aclrt_impl/memory.cpp) ├─ 参数校验：devAddr 为空时上报输入错误并返回 ACL_ERROR_INVALID_PARAM ├─ 性能打点：ACL_PROFILING_REG(AclrtValueWrite) 接入 Profiling 采集 └─ rtsValueWrite(devAddr, value, flag, stream) 下发任务到 RTS 层

其中rtsValueWrite是 ACL 层对 RTS（Runtime System）层的封装调用，任务最终由驱动在 Stream 上异步调度执行。接口声明位于 acl_rt.h。

三、aclrtValueWait：等待内存值满足条件后解除阻塞

3.1 函数原型与功能说明

aclError aclrtValueWait(void* devAddr, uint64_t value, uint32_t flag, aclrtStream stream)

aclrtValueWait等待指定内存中的数据满足一定条件后解除阻塞，同样是一个异步接口。调用方将"等待任务"下发到 Stream，任务在执行期间会持续比较内存值与给定条件，条件满足前后续任务保持阻塞，满足后解除阻塞继续执行。它是内存语义同步中"等待侧"的核心接口。

3.2 参数说明

3.3 flag 比较模式详解

flag参数决定了内存值与value的比较逻辑，共支持 4 种模式，宏定义位于 acl_rt.h：

其中ACL_STREAM_WAIT_VALUE_EQ是最常用的模式，语义最直观——"等到内存值精确等于目标值"。GEQ模式常用于计数/水位类场景（等到计数达到阈值），AND/NOR模式适合利用位标志实现多个信号量的组合等待。需要特别注意的是，以上比较均基于无符号 64 位内存值进行（GEQ模式内部按int64_t做差值判断），使用时需结合业务语义合理取值。

3.4 返回值说明

返回0表示成功，返回其他值表示失败，具体错误码请参见 aclError。

3.5 源码实现与调用链

与写值接口对称，aclrtValueWaitImpl实现在 memory.cpp 中：

aclError aclrtValueWaitImpl(void* devAddr, uint64_t value, uint32_t flag, aclrtStream stream) { ACL_PROFILING_REG(acl::AclProfType::AclrtValueWait); ACL_LOG_INFO("start to execute aclrtValueWait"); ACL_REQUIRES_NOT_NULL_WITH_INPUT_REPORT(devAddr); // 校验 devAddr 非空 ACL_REQUIRES_RTS_OK(rtsValueWait(devAddr, value, flag, static_cast<rtStream_t>(stream))); return ACL_SUCCESS; }

调用链与aclrtValueWrite完全对称：

aclrtValueWait (对外 API，头文件声明于 include/external/acl/acl_rt.h) └─> aclrtValueWaitImpl (src/acl/aclrt_impl/memory.cpp) ├─ 参数校验：devAddr 为空时返回 ACL_ERROR_INVALID_PARAM ├─ 性能打点：ACL_PROFILING_REG(AclrtValueWait) └─ rtsValueWait(devAddr, value, flag, stream) 下发等待任务到 RTS 层

四、完整实战：基于流内存操作实现跨流同步

4.1 场景一：双线程多流内存同步样例

仓库中的 9_multistream_sync_memory 样例 完整演示了这两个接口的配合使用。其核心流程是：

1. 线程 A 在 Stream A 上下发aclrtValueWait，等待devPtrA处的内存值等于 100；
2. 线程 B 在 Stream B 上下发aclrtValueWrite，向devPtrA写入 100；
3. 写入完成前，Stream A 上的等待任务持续阻塞；写入完成后，等待条件满足，Stream A 解除阻塞继续执行。

样例核心代码（main.cpp）中的等待线程与写入线程如下：

// 等待线程：在 streamA 上等待 devPtr 处的内存值等于 valueCompare int ThreadWait(aclrtStream stream, int32_t deviceId, void* devPtr, uint64_t valueCompare, const char* filePath) { aclrtSetDevice(deviceId); CHECK_ERROR(aclrtValueWait(devPtr, valueCompare, ACL_STREAM_WAIT_VALUE_EQ, stream)); // 等待任务解除阻塞后，同步等待该 Stream 上所有任务完成 CHECK_ERROR(aclrtSynchronizeStream(stream)); return 0; } // 写入线程：在 streamB 上向 devPtr 写入 valueWrite int ThreadWrite(aclrtStream stream, int32_t deviceId, void* devPtr, uint64_t valueWrite, const char* filePath) { aclrtSetDevice(deviceId); CHECK_ERROR(aclrtValueWrite(devPtr, valueWrite, 0, stream)); CHECK_ERROR(aclrtSynchronizeStream(stream)); return 0; }

主流程则完整演示了"申请内存 → 创建双流 → 启动等待线程 → 启动写入线程"的标准步骤：

int32_t main() { aclInit(nullptr); int32_t deviceId = 0; aclrtSetDevice(deviceId); // 1. 在 Device 上申请同步用内存 uint64_t size = 1 * 1024 * 1024; void* devPtrA; CHECK_ERROR(aclrtMalloc(&devPtrA, size, ACL_MEM_MALLOC_HUGE_FIRST)); // 2. 创建两条独立 Stream aclrtStream streamA = nullptr; CHECK_ERROR(aclrtCreateStream(&streamA)); aclrtStream streamB = nullptr; CHECK_ERROR(aclrtCreateStream(&streamB)); // 3. 先启动等待线程，再启动写入线程 uint64_t valueCompare = 100; std::thread threadA(ThreadWait, streamA, deviceId, devPtrA, valueCompare, filePath); (void)usleep(waitTime); // 确保等待任务已下发并进入阻塞 uint64_t valueWrite = 100; std::thread threadB(ThreadWrite, streamB, deviceId, devPtrA, valueWrite, filePath); threadB.join(); threadA.join(); aclrtDestroyStreamForce(streamA); aclrtDestroyStreamForce(streamB); aclrtFree(devPtrA); aclrtResetDeviceForce(deviceId); aclFinalize(); return 0; }

编译运行方式（详见样例 README）：

# 1. 切换到样例目录 cd ${git_clone_path}/example/1_basic_features/memory/9_multistream_sync_memory # 2. 设置环境变量（${install_root} 替换为 CANN 安装根目录，默认 /usr/local/Ascend） source ${install_root}/cann/set_env.sh source ${git_clone_path}/example/set_sample_env.sh # 3. 运行样例 bash run.sh

运行输出的关键日志验证了同步时序的正确性：Stream A: wait ... to meet the condition表明等待任务已阻塞，随后Stream B: write data写入值后，Stream A: the data ... has met the condition表明等待被成功唤醒，且等待线程读取到的标志位123证明写入线程执行前等待线程确实处于阻塞状态。

4.2 场景二：算子参与的内存语义同步

流内存操作接口更强大的能力体现在算子级同步：等待侧既可以是aclrtValueWait任务，也可以是算子核函数内部的轮询逻辑。在 内存语义同步指南 中给出了完整的 Host 侧参考示例：

// 创建 Stream aclrtStream stream1; aclrtStream stream2; aclrtCreateStream(&stream1); aclrtCreateStream(&stream2); // 申请 Device 内存 void* syncMem; aclrtMalloc(&syncMem, sizeof(uint64_t), ACL_MEM_MALLOC_NORMAL_ONLY); // 在 stream1 上下发 wait 任务，阻塞等待直到 syncMem 指向内存中的值为 1 aclrtValueWait(syncMem, 1, ACL_STREAM_WAIT_VALUE_EQ, stream1); // 在 stream2 上下发 myKernel1，该 kernel 内部向 syncMem 指向内存写 1，解除 stream1 上 wait 任务的阻塞 myKernel1<<<numBlocks, nullptr, stream2>>>(syncMem); // 在 stream1 上下发 myKernel2，该 kernel 内部轮询等待直到 syncMem 指向内存的值为 2 myKernel2<<<numBlocks, nullptr, stream1>>>(syncMem); // 在 stream2 上下发 write 任务，向 syncMem 写入 2，解除 stream1 上 myKernel2 的阻塞 aclrtValueWrite(syncMem, 2, 0, stream2);

对应的 Device 侧核函数实现要点（源码参考）：

• 写侧核函数：直接向同步内存写入值后，调用dcci指令做缓存一致性刷新（__gm__ uint64_t* flag = reinterpret_cast<__gm__ uint64_t*>(syncMem); *flag = 1; dcci(flag, 0, 2);），确保写值对其他执行单元可见；
• 读侧核函数：将内存声明为volatile后轮询等待目标值，轮询间隙调用dcci保持缓存一致性（while (*flag != 2) { dcci(flag, 0, 2); }）。

这个示例揭示了两个关键实践细节：其一，同步内存必须通过aclrtMalloc等接口申请，且内存值按 64 bit 读写；其二，跨执行单元的可见性依赖缓存一致性指令（dcci），Host 侧aclrtValueWrite与 Device 侧核函数写入在硬件层面共用同一套内存语义。此外，由于内存语义同步基于通用 Device 内存实现，可以使用aclrtMemset/aclrtMemsetAsync对同步内存做初始化和清除。

五、单元测试验证与错误处理

仓库在 tests/ut/acl/testcase/acl_runtime_unittest.cpp 中为两个接口提供了完整的 UT 用例，覆盖了成功路径与非法参数路径：

// 非法参数：devAddr 传 nullptr，应返回 ACL_ERROR_INVALID_PARAM TEST_F(UTEST_ACL_Runtime, aclrtValueWrite_failed_with_invalid_args) { auto ret = aclrtValueWrite(nullptr, 100, 0, nullptr); EXPECT_EQ(ret, ACL_ERROR_INVALID_PARAM); // 底层 rtsValueWrite 返回错误时，错误码应被透传 EXPECT_CALL(MockFunctionTest::aclStubInstance(), rtsValueWrite(_, _, _, _)) .WillOnce(Return(ACL_ERROR_RT_PARAM_INVALID)); auto devAddr = reinterpret_cast<void*>(0x1000U); ret = aclrtValueWrite(devAddr, 100, 0, nullptr); EXPECT_EQ(ret, ACL_ERROR_RT_PARAM_INVALID); } // 成功路径 TEST_F(UTEST_ACL_Runtime, aclrtValueWrite_success) { auto devAddr = reinterpret_cast<void*>(0x1000U); const auto ret = aclrtValueWrite(devAddr, 100, 0, nullptr); EXPECT_EQ(ret, ACL_SUCCESS); }

aclrtValueWait的测试用例结构完全对称（aclrtValueWait_failed_with_invalid_args与aclrtValueWait_success）。从测试可以总结出错误处理要点：

1. devAddr为空时，接口在 ACL 层即返回ACL_ERROR_INVALID_PARAM，不会下发到 RTS；
2. 底层rtsValueWrite/rtsValueWait返回错误时，错误码原样透传给调用方（如ACL_ERROR_RT_PARAM_INVALID）；
3. stream传NULL表示使用默认 Stream，是合法入参；
4. 成功时统一返回ACL_SUCCESS（0）。

六、使用注意事项与最佳实践

1. 内存对齐与位宽：devAddr必须为 8 字节对齐，有效位宽为 64 bit，且需通过aclrtMalloc等接口在 Device 侧提前申请，不能使用未经对齐或越界的内存地址。
2. 产品兼容性：当前仅 Ascend 950PR/950DT、Atlas A3 系列、Atlas A2 系列产品支持该接口，老型号 Atlas 训练/推理系列产品及 IPV350 不支持，跨版本开发时需按目标设备做能力判断。
3. 同步内存的生命周期：同步所用内存可用aclrtMemset/aclrtMemsetAsync初始化与清除；多线程场景下，等待线程与写入线程需各自调用aclrtSetDevice绑定设备（参见样例 main.cpp）。
4. 阻塞语义理解：aclrtValueWait阻塞的是其所在 Stream 上后续任务的执行，而非 Host 线程；若需要 Host 侧同步等待，需配合aclrtSynchronizeStream使用。
5. 与 Event/Notify 的选型：如果同步双方都是 Host 侧任务，Event/Notify 即可满足；当需要算子在执行过程中参与同步（如核函数内轮询等待其他流写值）时，应使用内存语义同步方案。
6. flag 参数：aclrtValueWrite的flag为预留参数，必须固定传 0；aclrtValueWait的flag则需从 4 种比较模式宏中取值，非法值会导致任务行为不确定。

七、小结

aclrtValueWrite与aclrtValueWait是 CANN Runtime 实现内存语义同步的一对基础接口：前者在 Stream 上异步写入 64 位内存值，后者在 Stream 上按 4 种比较模式等待内存值满足条件后解除阻塞。二者通过 memory.cpp 中的实体实现下发至 RTS 层执行，支持与核函数内部的dcci+ 轮询逻辑配合，实现"算子参与"的跨流同步。结合 9_multistream_sync_memory 样例 与 内存语义同步指南 中的参考代码，你可以在实际项目中按"申请对齐内存 → 下发等待任务 → 下发写值任务 → 同步等待收尾"的标准流程，构建可靠的跨流协作逻辑。

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime