DataStoreBarrier产品支持情况产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√Atlas 200I/500 A2 推理产品xAtlas 推理系列产品AI CorexAtlas 推理系列产品Vector CorexAtlas 训练系列产品x功能说明数据同步屏障指令该指令会阻塞当前线程执行确保所有先前的写内存操作完成并对其它硬件单元可见后才继续执行后续指令。用于AI CPU与AI Core多核之间的数据一致性。需要包含的头文件#include aicpu_api.h函数原型DataStoreBarrier(void)参数说明无返回值说明无约束说明该接口仅支持通过...调用并在异构编译场景使用。调用示例在AI CPU算子Kernel侧实现代码中调用AscendC::DataStoreBarrier()确保AI CPU算子对Tiling数据的修改写入内存使得AI Core算子能够正确读取Tiling数据struct TilingInfo { uint64_t lock; // AI CPU/AI Core之间同步的锁 int8_t type; int8_t mode; int8_t len; }; struct KernelArgs { uint32_t *xDevice; uint32_t *yDevice; uint32_t *zDevice; TilingInfo *ti; // 与AI Core共享的参数用于同步tiling选择 }; templatetypename T, int8_t mode, int8_t len __aicore__ void hello_world_impl(GM_ADDR m) { if constexpr (std::is_same_vT, float) { AscendC::printf(Hello World: float mode %u len %u.\n, mode, len); } else if constexpr (std::is_same_vT, int) { AscendC::printf(Hello World: int mode %u len %u.\n, mode, len); } } // AI Core算子总入口 // tilingInfo: 和AI CPU算子共同传递的参数用于数据共享 templatetypename T, int8_t mode, int8_t len __mix__(1,2) __global__ __aicore__ void hello_world(GM_ADDR m, GM_ADDR TilingPtr) { __gm__ struct KernelInfo::TilingInfo *ti (__gm__ struct KernelInfo::TilingInfo *)TilingPtr; AscendC::GlobalTensoruint64_t lock; lock.SetGlobalBuffer(reinterpret_cast__gm__ uint64_t *(ti-lock)); if ASCEND_IS_AIV { if (AscendC::GetBlockIdx() 0) { while (*reinterpret_castvolatile __gm__ uint64_t*(lock.GetPhyAddr(0)) 0) { // 下沉模式AI Core等待AICPU tiling计算完成 AscendC::DataCacheCleanAndInvaliduint64_t, AscendC::CacheLine::SINGLE_CACHE_LINE, AscendC::DcciDst::CACHELINE_OUT(lock); //直接访问Global Memory获取最新数据 } } } // 上面是1个核等待AI CPU tiling计算完成这里进行核间同步 AscendC::SyncAllfalse(); // 根据tiling参数值选择不同模板 if (ti-type 0 ti-mode 1 ti-len 2) { hello_world_implfloat, 1, 2(m); } else if (ti-type 1 ti-mode 2 ti-len 4) { hello_world_implint, 2, 4(m); } // 执行完留一个核释放lock if ASCEND_IS_AIV { if (AscendC::GetBlockIdx() 0) { lock.SetValue(0, 0); // 刷新 lock AscendC::DataCacheCleanAndInvaliduint64_t, AscendC::CacheLine::SINGLE_CACHE_LINE, AscendC::DcciDst::CACHELINE_OUT(lock); //刷新Dcache同步与GM之间的数据 } } } extern C __global__ __aicpu__ uint32_t MyAicpuKernel(void *arg) { KernelArgs* cfg (KernelArgs*)arg; AscendC::printf(MyAicpuKernel inited!\n); cfg-ti-lock 1; cfg-ti-type 1; cfg-ti-mode 2; cfg-ti-len 4; AscendC::DataStoreBarrier(); // 对tilingInfo进行写同步 AscendC::printf(MyAicpuKernel inited type %u mode %u len %u end!\n, cfg-ti-type, cfg-ti-mode, cfg-ti-len); return 0; }创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
