内存语义同步

内存语义同步机制允许用户基于通用Device内存实现同步。与Event/Notify同步机制不同，基于内存语义的同步机制还支持算子作为同步参与方，即算子可以在执行过程中与另一条流进行同步。

下图展示了在算子执行过程中与另一条流之间的同步过程：

内存语义同步相关接口的调用代码示例如下，不可以直接拷贝编译运行，仅供参考：

• Device示例代码（算子核函数实现代码）

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24

  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示例代码

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18

  // 创建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内存实现，所以可以通过aclrtMemset/aclrtMemsetAsync初始化和清除同步所用的内存。