内存语义同步【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime内存语义同步机制允许用户基于通用Device内存实现同步。与Event/Notify同步机制不同基于内存语义的同步机制还支持算子作为同步参与方即算子可以在执行过程中与另一条流进行同步。下图展示了在算子执行过程中与另一条流之间的同步过程内存语义同步相关接口的调用代码示例如下不可以直接拷贝编译运行仅供参考Device示例代码算子核函数实现代码extern C __global__ __aicore__ void myKernel1(GM_ADDR syncMem) { // 算子逻辑 ...... // 向syncMem所指向内存写1 __gm__ uint64_t* flag reinterpret_cast__gm__ uint64_t*(syncMem); *flag 1; dcci(flag, 0, 2); // 算子逻辑 ...... } extern C __global__ __aicore__ void myKernel2(GM_ADDR syncMem) { // 算子逻辑 ... // 轮询阻塞直到syncMem所指向内存值为2 __gm__ volatile uint64_t* flag reinterpret_cast__gm__ uint64_t*(syncMem); dcci(flag, 0, 2); while (*flag ! 2) { dcci(flag, 0, 2); } // 算子逻辑 ...... }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任务的阻塞状态 myKernel1numBlocks, nullptr, stream2(syncMem); // 在stream1上下发myKernel2该kernel内部轮询等待直到syncMem所指向内存的值为2 myKernel2numBlocks, nullptr, stream1(syncMem); // 在stream2上下发write任务该任务往syncMem所指向内存写为2从而解除stream1上myKernel2的阻塞状态 aclrtValueWrite(syncMem, 2, 0, stream2);说明因为内存语义同步机制是基于通用Device内存实现所以可以通过aclrtMemset/aclrtMemsetAsync初始化和清除同步所用的内存。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考