DataStoreBarrier【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况Ascend 950PR/Ascend 950DT支持Atlas A3 训练系列产品/Atlas A3 推理系列产品支持Atlas A2 训练系列产品/Atlas A2 推理系列产品支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持功能说明数据同步屏障指令该指令会阻塞当前线程执行确保所有先前的写内存操作完成并对其它硬件单元可见后才继续执行后续指令。用于AI CPU与AI Core多核之间的数据一致性。需要包含的头文件#include aicpu_api.h函数原型DataStoreBarrier(void)参数说明无返回值说明无约束说明该接口仅支持通过...调用并在异构编译场景使用。调用示例在AI CPU算子Kernel侧实现代码中调用AscendC::DataStoreBarrier()确保AI CPU算子对Tiling数据的修改写入内存使得AI Core算子能够正确读取Tiling数据namespace KernelInfo { 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 %d len %d.\n, mode, len); } else if constexpr (std::is_same_vT, int) { AscendC::printf(Hello World: int mode %d len %d.\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) { KernelInfo::KernelArgs* cfg (KernelInfo::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 %d mode %d len %d end!\n, cfg-ti-type, cfg-ti-mode, cfg-ti-len); return 0; }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考