CANN 昇腾 AICPU 设备端 Tiling 下沉实战基于 aicpu_device_tiling 样例的双 Stream 内核协同编程指南【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples本文基于 CANN 昇腾开源样例仓库cann-samples中的Samples/0_Introduction/04_aicpu/00_device_tiling样例系统讲解AI CPU 算子负责 Tiling 计算、AI Core 算子负责实际计算的设备端 Tiling 下沉编程范式。文中将完整继承样例 README 的目录结构、编译选项与运行结果说明并结合仓库中的main.asc、aicpu_tiling.aicpu、aicore_kernel.asc、kernel_args.h等源码逐层剖析双 Stream 协同、Event 同步、模板分支选择等底层实现机制。读完本文你将掌握在 ASCAscend C编程框架下通过内核调用符...下发 AI CPU 与 AI Core 算子、并利用aclrtRecordEvent/aclrtStreamWaitEvent实现跨流同步的完整实战方法。一、技术背景为什么要用 AI CPU 做 Tiling 下沉在昇腾 NPU 的算子开发中Tiling分片/切分指根据输入张量的形状、数据类型、可用片上内存如 UB/L1等运行时信息将一个大任务切分为多个可在硬件上执行的子任务并确定每个子任务的循环次数、搬运粒度、核数分配等参数的过程。传统方案中 Tiling 通常在Host 端完成Host 先读取输入 shape 等元信息计算 Tiling 参数再连同数据一起下发到设备端。该方式存在两点开销Host 与 Device 之间需要额外同步Tiling 计算期间内核无法启动需要 Host 侧感知完整的算子上下文在某些图模式或下沉场景下并不方便。本样例展示的设备端 Tiling 下沉Tiling Offload则是把 Tiling 计算搬到设备侧的AI CPU上执行Host 只负责把数据指针等必要参数传给 AI CPU 算子AI CPU 算子运行在昇腾设备上、直接算出 Tiling 参数并写入设备内存随后 AI Core 算子读取这些参数执行真正的计算。这样既省去了 Host↔Device 的往返同步也让 Tiling 计算与 AI Core 计算共享同一份设备端数据视图是算子下沉场景下的关键能力。从样例分类看本样例位于 Samples/0_Introduction/04_aicpu 目录下是 AICPU 算子入门路径的第一个样例对应描述为使用 AI CPU 算子进行 Tiling 下沉计算的实现方法。二、样例概览与目录结构样例根目录为 Samples/0_Introduction/04_aicpu/00_device_tiling其下唯一的子目录aicpu_device_tiling即为完整工程。目录结构如下├── aicpu_device_tiling │ ├── CMakeLists.txt // 编译工程文件 │ ├── aicore_kernel.asc // AI Core算子实现 │ ├── kernel_args.h // tiling结构体头文件 │ ├── main.asc // AI CPU算子与AI Core算子调用 │ ├── aicpu_tiling.aicpu // AI CPU算子实现 │ └── README.md // 样例说明文档各文件职责一目了然文件类型职责kernel_args.h头文件定义 AI CPU 与 AI Core 共享的 Tiling 结构体与内核参数结构体aicpu_tiling.aicpuAI CPU 算子源码在设备端计算并写入 Tiling 参数aicore_kernel.ascAI Core 算子源码读取 Tiling 参数并据此选择模板实例打印 Hello Worldmain.ascHost 端入口创建双 Stream、分配设备内存、按序下发两类算子并同步CMakeLists.txt构建脚本同时编译 ASC、AICPU、CXX 三种语言并链接为可执行文件从构建脚本可以看出该工程的一个特点project(kernel_samples LANGUAGES ASC AICPU CXX)声明了三种语言add_executable(demo aicore_kernel.asc main.asc aicpu_tiling.aicpu)将 AI Core 算子、Host 代码与 AI CPU 算子编译进同一个可执行文件并通过set_target_properties(demo PROPERTIES LINKER_LANGUAGE ASC)指定链接语言为 ASC——这意味着 AI CPU 算子与 AI Core 算子可以在同一个编译工程、同一个可执行程序中协同工作这是设备端 Tiling 下沉落地的基础设施。三、支持的产品与 CANN 版本样例 README 明确了如下产品兼容性矩阵运行时请据此选择正确的 NPU 架构与 CANN 版本产品CANN软件版本Ascend 950PR/Ascend 950DT CANN 9.1.0Atlas A3 训练系列产品/Atlas A3 推理系列产品 CANN 9.0.0Atlas A2 训练系列产品/Atlas A2 推理系列产品 CANN 9.0.0编译时需通过CMAKE_ASC_ARCHITECTURES指定 NPU 架构详见下文编译选项说明其取值与产品的对应关系为dav-2201Atlas A2 训练/推理系列产品、Atlas A3 训练/推理系列产品dav-3510Ascend 950PR / Ascend 950DT。四、核心实现原理剖析结合源码本小节将四个源码文件串成一条完整的数据流共享结构体定义 → AI CPU 写入 Tiling → AI Core 消费 Tiling → Host 编排时序。4.1 共享 Tiling 结构体kernel_args.hkernel_args.h 定义了两类结构体均置于KernelInfo命名空间中struct TilingInfo { int8_t type; int8_t mode; int8_t len; }; struct KernelArgs { uint32_t* xDevice; uint32_t* yDevice; uint32_t* zDevice; TilingInfo* ti; // Parameters shared with aicore are used for synchronizing tiling selection };TilingInfo是设备端 Tiling 参数本体由 AI CPU 算子写入、AI Core 算子读取。本样例仅用type/mode/len三个字段模拟真实场景下的分片决策如数据类型选择、计算模式、数据长度KernelArgs是内核参数打包结构xDevice/yDevice/zDevice为设备端数据指针ti指向 Tiling 参数所在设备内存。关键点在于该头文件被aicpu_tiling.aicpu与aicore_kernel.asc同时 include保证了两个算子看到的是同一份内存布局这是设备端跨算子传递 Tiling 参数的前提。4.2 AI CPU 算子设备端计算并写入 Tilingaicpu_tiling.aicpuaicpu_tiling.aicpu 是整个流程的生产者__global__ __aicpu__ uint32_t MyAicpuKernel(KernelInfo::KernelArgs args) { AscendC::printf(MyAicpuKernel inited\n); args.ti-type 1; args.ti-mode 2; args.ti-len 4; AscendC::DataStoreBarrier(); AscendC::printf(MyAicpuKernel inited type %u mode %u len %u end!\n, args.ti-type, args.ti-mode, args.ti-len); return 0; }要点解读__global__ __aicpu__修饰符声明这是一个在 AI CPU 上运行的内核函数与 AI Core 内核的__global__ __aicore__修饰符相区别Tiling 计算发生的位置样例中以直接赋值type1, mode2, len4代替真实的 Tiling 计算逻辑真实场景中此处会根据 shape、内存容量等条件计算分片写入目标是指向设备内存的args.tiAscendC::DataStoreBarrier()数据存储屏障。AI CPU 与 AI Core 是设备上不同的执行单元AI CPU 写入的 Tiling 参数必须经过屏障保证其先于后续操作可见AI Core 算子才能可靠地读到正确值——这是设备端 Tiling 下沉正确性的关键一步内核打印的两行日志MyAicpuKernel inited与MyAicpuKernel inited type 1 mode 2 len 4 end!与 README 中的执行结果完全对应。4.3 AI Core 算子按 Tiling 参数选择模板aicore_kernel.ascaicore_kernel.asc 是流程的消费者实现了基于 Tiling 参数的模板分支选择template typename T, int8_t mode, int8_t len __aicore__ void hello_world_impl(__gm__ uint8_t* m) { if constexpr (std::is_same_vT, float) { AscendC::printf(Hello World: float mode %u len %u m %d.\n, mode, len, *m); } else if constexpr (std::is_same_vT, int) { AscendC::printf(Hello World: int mode %u len %u m %d.\n, mode, len, *m); } } template typename T, int8_t mode, int8_t len __mix__(1, 2) __global__ __aicore__ void hello_world(__gm__ uint8_t* m, __gm__ uint8_t* TilingPtr) { __gm__ struct KernelInfo::TilingInfo* ti (__gm__ struct KernelInfo::TilingInfo*)TilingPtr; // Select different templates based on tiling value 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); } } extern C void hello_world_do(uint32_t numBlocks, void* stream, uint8_t* m, uint8_t* ti) { hello_worldint, 10, 201, 0, stream(m, ti); }要点解读从设备内存读取 Tiling内核入口将TilingPtr强转为__gm__ KernelInfo::TilingInfo*直接解引用设备端__gm__Global Memory指针读取 AI CPU 写入的type/mode/len模板分支选择根据ti-type/mode/len的实际值选择不同的hello_world_implT, mode, len模板实例。由于 Tiling 参数在运行时才确定而模板在编译期实例化这里实际上是用运行时 if 判断 编译期if constexpr的组合把运行期决策映射到编译期实例——这正是 Tiling 参数驱动不同计算路径的典型写法__mix__(1, 2)混合调度声明该内核同时占用1 个 Cube 执行单元和 2 个 Vector 执行单元因此 README 中特别说明 Hello World 日志会打印 3 次详见执行结果解读extern C包装hello_world_do以 C 符号导出供main.asc中的 C Host 代码调用其内部真正以1, 0, stream下发 AI Core 内核。4.4 Host 端编排双 Stream Event 同步main.ascmain.asc 是主机端入口完整演示了AI CPU 与 AI Core 在不同 Stream 上启动、以 Event 建立依赖的编排方式整体流程如下int32_t main(int argc, char const* argv[]) { aclInit(nullptr); int32_t deviceId 0; aclrtSetDevice(deviceId); aclrtStream aicpu_stream nullptr; aclrtStream aicore_stream nullptr; aclrtCreateStream(aicpu_stream); aclrtCreateStream(aicore_stream); aclrtEvent event; aclrtCreateEventExWithFlag(event, ACL_EVENT_SYNC); void* zDevice; void* ti; void* zHost; aclrtMalloc((void**)zDevice, 4096, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc((void**)ti, 4096, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMallocHost((void**)zHost, 4096); aclrtMemset((void*)ti, 4096, 0, 4096); aclrtMemset((void*)zHost, 4096, 10, 4096); aclrtMemcpy(zDevice, 4096, zHost, 4096, ACL_MEMCPY_HOST_TO_DEVICE); struct KernelInfo::KernelArgs args {0}; args.xDevice (uint32_t*)zDevice; args.yDevice args.xDevice 1; args.zDevice args.yDevice 1; args.ti (KernelInfo::TilingInfo*)ti; MyAicpuKernel1, 0, aicpu_stream(args); aclrtRecordEvent(event, aicpu_stream); aclrtStreamWaitEvent(aicore_stream, event); hello_world_do(1, aicore_stream, (uint8_t*)zDevice, (uint8_t*)ti); aclrtSynchronizeStream(aicpu_stream); aclrtSynchronizeStream(aicore_stream); ... }按执行顺序拆解步骤代码作用1aclInit/aclrtSetDevice(0)初始化 ACL 运行时并指定设备2aclrtCreateStream× 2创建aicpu_stream与aicore_stream两条独立流3aclrtCreateEventExWithFlag(event, ACL_EVENT_SYNC)创建同步事件用于跨流同步4aclrtMalloc/aclrtMemset/aclrtMemcpy分配设备内存数据区zDevice、Tiling 区ti并初始化Tiling 区清零、数据区全部置为 10这也是最终日志中m 10的来源args.xDevice/yDevice/zDevice依次指向同一段内存中的相邻 uint32 单元5MyAicpuKernel1, 0, aicpu_stream(args)在aicpu_stream上下发 AI CPU 内核执行 Tiling 计算并写入ti6aclrtRecordEvent(event, aicpu_stream)在 aicpu_stream 中记录 event标记AI CPU 内核已完成7aclrtStreamWaitEvent(aicore_stream, event)阻塞aicore_stream直到该 event 完成——保证 AI Core 内核读取到的 Tiling 参数是 AI CPU 写好的最终值8hello_world_do(1, aicore_stream, ...)在 aicore_stream 上下发 AI Core 内核9aclrtSynchronizeStream× 2 / 资源释放等待两条流全部执行完成释放内存与流、复位设备、aclFinalize这正是 README 所述AI CPU 算子与 AI Core 算子在不同 stream 上进行 launch……使用aclrtRecordEvent在指定 stream 中记录 event使用aclrtStreamWaitEvent阻塞指定的 stream直到指定的 event 完成的完整源码级呈现。五、编译运行与结果验证5.1 配置环境变量请根据当前环境上 CANN 开发套件包的安装方式配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。同时请确保当前环境已安装与目标产品匹配的 CANN 版本参考支持的产品与 CANN 版本一节。5.2 编译与执行在样例目录Samples/0_Introduction/04_aicpu/00_device_tiling/aicpu_device_tiling下执行mkdir -p build cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # 编译工程 ./demo # 执行编译生成的可执行程序执行样例5.3 编译选项说明选项可选值说明CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品 与 Atlas A3 训练系列产品/Atlas A3 推理系列产品dav-3510 对应 Ascend 950PR/Ascend 950DT该选项在 CMakeLists.txt 中以CACHE STRING形式声明默认值为dav-2201并作为--npu-arch编译参数传入 ASC 编译器target_compile_options中--npu-arch${CMAKE_ASC_ARCHITECTURES}。请在编译前按实际产品选择对应架构。另外构建脚本中还包含一个实验性编译选项# Experimental option. Future compatibility and support are not guaranteed. --cce-aicpu-launch-with-interface该选项启用后AI CPU 算子得以通过内核调用符...的接口方式被直接 launch这也是main.asc中MyAicpuKernel1, 0, aicpu_stream(args)能工作的原因。注释明确标注其为实验性特性未来兼容性不受保证实战项目中请关注 CANN 版本的演进说明。5.4 执行结果解读执行成功时输出如下MyAicpuKernel inited MyAicpuKernel inited type 1 mode 2 len 4 end! Hello World: int mode 2 len 4 m 10. Hello World: int mode 2 len 4 m 10. Hello World: Hello World: int mode 2 len 4 m 10.逐行对应关系前两行来自AI CPU 算子MyAicpuKernel先是进入时的MyAicpuKernel inited写入type1, mode2, len4并执行DataStoreBarrier后打印MyAicpuKernel inited type 1 mode 2 len 4 end!后三行来自AI Core 算子AI Core 读取 Tiling 参数命中type1 mode2 len4分支实例化hello_world_implint, 2, 4打印Hello World: int mode 2 len 4 m 10.m 10对应数据区被初始化为 10。其中打印 3 次的原因正如 README 所解释__mix__(1, 2)会启动1 个 Cube 执行单元和 2 个 Vector 执行单元每个执行单元都会执行一次内核体因此日志共出现 3 次。六、关键机制深度解析6.1 内核调用符...的统一抽象本样例最值得注意的一点是AI CPU 算子与 AI Core 算子均使用内核调用符...进行调用README 原文。传统上 AI CPU 算子通常通过独立的 AICPU 接口下发而本工程通过实验性编译选项--cce-aicpu-launch-with-interface统一了二者的调用语法MyAicpuKernel1, 0, aicpu_stream(args); // AI CPU 内核aicpu_stream hello_worldint, 10, 201, 0, stream(m, ti); // AI Core 内核hello_world_do 内部aicore_stream这种统一抽象让 Host 代码可以用同一套心智模型管理两类算子块数本样例均为 1、共享内存0、所属 Stream以及参数列表。6.2 Event 跨流同步谁先谁后的保证AI CPU 与 AI Core 运行在不同执行单元上两条 Stream 之间没有天然的执行顺序。样例用aclrtRecordEventaclrtStreamWaitEvent建立显式依赖aclrtRecordEvent(event, aicpu_stream)在 aicpu_stream 上打点当该点之前的所有任务即MyAicpuKernel完成后 event 即被触发aclrtStreamWaitEvent(aicore_stream, event)aicore_stream 会阻塞等待该 event从而保证 AI Core 内核启动时AI CPU 写入的 Tiling 参数已经就绪。这一机制避免了aclrtSynchronizeStream式的全量同步只让依赖方aicore_stream等待必要的事件两条流仍可并行推进其他无关任务是下沉 流式并行场景的标准做法。6.3 DataStoreBarrier跨执行单元的数据可见性aicpu_tiling.aicpu在写完 Tiling 参数后调用了AscendC::DataStoreBarrier()。结合源码结构可以推断AI CPU 与 AI Core 是不同的执行单元AI CPU 对设备内存的写入需要经过屏障才能保证对后续跨执行单元的访问可见、有序。若不设置该屏障AI Core 算子可能读到 Tiling 参数的中间态导致模板分支选择错误——这是设备端 Tiling 下沉中极易踩坑的正确性问题样例以极简形式给出了标准示范。6.4 模板分支选择运行期决策 → 编译期实例AI Core 算子中Tiling 参数在运行期由 AI CPU 算出但内核要执行的代码路径float还是int、mode/len取何值在编译期就需要确定。样例的解法是两层配合外层运行期if (ti-type ... ...)对 Tiling 值做分支判断内层hello_world_implT, mode, len用if constexpr (std::is_same_vT, float)在编译期消除不匹配的分支。于是同一份内核源码可以针对多种 Tiling 组合分别实例化出专用版本既保证了正确性又避免了运行期分支开销——这一Tiling 驱动模板实例化的模式在真实算子如各类 elementwise、矩阵类算子中非常常见。6.5 构建基础设施ASC AICPU 混合工程从 CMakeLists.txt 可以看到支撑本样例的完整构建链路find_package(ASC REQUIRED)与find_package(AICPU REQUIRED)分别引入 ASCAscend C 内核与 AICPUAI CPU 算子的构建支持project(... LANGUAGES ASC AICPU CXX)一个工程同时编译三种语言源码LINKER_LANGUAGE ASC可执行文件的链接以 ASC 工具链为主导保证内核与 Host 代码正确链接链接pthread与dl满足 ACL 运行时与动态加载需求。对于希望把本样例改造成真实算子的开发者这一 CMake 骨架可以直接复用。七、从样例到实战注意事项与扩展方向结合上文源码分析将本样例迁移到真实算子开发时建议关注以下几点Tiling 计算逻辑替换样例中用固定赋值模拟 Tiling真实场景应在MyAicpuKernel中根据输入 shape、数据类型、UB/L1 容量等运行时信息计算分片参数并扩展TilingInfo结构体承载更多字段注意 AI CPU 与 AI Core 共享的头文件要保持同步修改同步点不可省略AI CPU 写 Tiling 后的DataStoreBarrier()与 Host 端的aclrtStreamWaitEvent分别保证了写入可见与跨流有序二者缺一不可内存布局与指针语义样例中xDevice/yDevice/zDevice指向同一段设备内存的相邻单元真实场景建议按需独立分配并明确各指针的__gm__语义避免别名混淆多核与混合调度__mix__(1, 2)展示了 Cube/Vector 混合占用语法真实算子应结合计算特征选择执行单元组合并用块数参数样例中为 1控制核数扩展版本与架构匹配编译选项CMAKE_ASC_ARCHITECTURES必须与目标产品匹配dav-2201/dav-3510CANN 版本需满足第一节表格中的最低要求实验性编译选项--cce-aicpu-launch-with-interface的长期兼容性需以对应 CANN 版本发布说明为准。八、小结本样例以最小的代码量约 150 行完整呈现了AI CPU 计算 Tiling AI Core 消费 Tiling的设备端 Tiling 下沉全链路共享结构体统一两端的内存视图AI CPU 内核写入参数并做数据屏障AI Core 内核按参数做模板分支Host 端以双 Stream Event 建立正确的执行时序。其核心工程要点——__global__ __aicpu__/__global__ __aicore__内核声明、...统一调用符、aclrtRecordEvent/aclrtStreamWaitEvent跨流同步、AscendC::DataStoreBarrier()数据可见性保证、if constexpr模板实例化——构成了昇腾算子 Tiling 下沉开发的通用方法论可直接迁移到更复杂的真实算子工程中。【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 
长河君 · 2026/9/18 4:41:54 · 阅读 1.2k CANN 昇腾 AICPU 设备端 Tiling 下沉实战：基于 aicpu_device_tiling 样例的双 Stream 内核协同编程指南
CANN 昇腾 AICPU 设备端 Tiling 下沉实战基于 aicpu_device_tiling 样例的双 Stream 内核协同编程指南【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples本文基于 CANN 昇腾开源样例仓库cann-samples中的Samples/0_Introduction/04_aicpu/00_device_tiling样例系统讲解AI CPU 算子负责 Tiling 计算、AI Core 算子负责实际计算的设备端 Tiling 下沉编程范式。文中将完整继承样例 README 的目录结构、编译选项与运行结果说明并结合仓库中的main.asc、aicpu_tiling.aicpu、aicore_kernel.asc、kernel_args.h等源码逐层剖析双 Stream 协同、Event 同步、模板分支选择等底层实现机制。读完本文你将掌握在 ASCAscend C编程框架下通过内核调用符...下发 AI CPU 与 AI Core 算子、并利用aclrtRecordEvent/aclrtStreamWaitEvent实现跨流同步的完整实战方法。一、技术背景为什么要用 AI CPU 做 Tiling 下沉在昇腾 NPU 的算子开发中Tiling分片/切分指根据输入张量的形状、数据类型、可用片上内存如 UB/L1等运行时信息将一个大任务切分为多个可在硬件上执行的子任务并确定每个子任务的循环次数、搬运粒度、核数分配等参数的过程。传统方案中 Tiling 通常在Host 端完成Host 先读取输入 shape 等元信息计算 Tiling 参数再连同数据一起下发到设备端。该方式存在两点开销Host 与 Device 之间需要额外同步Tiling 计算期间内核无法启动需要 Host 侧感知完整的算子上下文在某些图模式或下沉场景下并不方便。本样例展示的设备端 Tiling 下沉Tiling Offload则是把 Tiling 计算搬到设备侧的AI CPU上执行Host 只负责把数据指针等必要参数传给 AI CPU 算子AI CPU 算子运行在昇腾设备上、直接算出 Tiling 参数并写入设备内存随后 AI Core 算子读取这些参数执行真正的计算。这样既省去了 Host↔Device 的往返同步也让 Tiling 计算与 AI Core 计算共享同一份设备端数据视图是算子下沉场景下的关键能力。从样例分类看本样例位于 Samples/0_Introduction/04_aicpu 目录下是 AICPU 算子入门路径的第一个样例对应描述为使用 AI CPU 算子进行 Tiling 下沉计算的实现方法。二、样例概览与目录结构样例根目录为 Samples/0_Introduction/04_aicpu/00_device_tiling其下唯一的子目录aicpu_device_tiling即为完整工程。目录结构如下├── aicpu_device_tiling │ ├── CMakeLists.txt // 编译工程文件 │ ├── aicore_kernel.asc // AI Core算子实现 │ ├── kernel_args.h // tiling结构体头文件 │ ├── main.asc // AI CPU算子与AI Core算子调用 │ ├── aicpu_tiling.aicpu // AI CPU算子实现 │ └── README.md // 样例说明文档各文件职责一目了然文件类型职责kernel_args.h头文件定义 AI CPU 与 AI Core 共享的 Tiling 结构体与内核参数结构体aicpu_tiling.aicpuAI CPU 算子源码在设备端计算并写入 Tiling 参数aicore_kernel.ascAI Core 算子源码读取 Tiling 参数并据此选择模板实例打印 Hello Worldmain.ascHost 端入口创建双 Stream、分配设备内存、按序下发两类算子并同步CMakeLists.txt构建脚本同时编译 ASC、AICPU、CXX 三种语言并链接为可执行文件从构建脚本可以看出该工程的一个特点project(kernel_samples LANGUAGES ASC AICPU CXX)声明了三种语言add_executable(demo aicore_kernel.asc main.asc aicpu_tiling.aicpu)将 AI Core 算子、Host 代码与 AI CPU 算子编译进同一个可执行文件并通过set_target_properties(demo PROPERTIES LINKER_LANGUAGE ASC)指定链接语言为 ASC——这意味着 AI CPU 算子与 AI Core 算子可以在同一个编译工程、同一个可执行程序中协同工作这是设备端 Tiling 下沉落地的基础设施。三、支持的产品与 CANN 版本样例 README 明确了如下产品兼容性矩阵运行时请据此选择正确的 NPU 架构与 CANN 版本产品CANN软件版本Ascend 950PR/Ascend 950DT CANN 9.1.0Atlas A3 训练系列产品/Atlas A3 推理系列产品 CANN 9.0.0Atlas A2 训练系列产品/Atlas A2 推理系列产品 CANN 9.0.0编译时需通过CMAKE_ASC_ARCHITECTURES指定 NPU 架构详见下文编译选项说明其取值与产品的对应关系为dav-2201Atlas A2 训练/推理系列产品、Atlas A3 训练/推理系列产品dav-3510Ascend 950PR / Ascend 950DT。四、核心实现原理剖析结合源码本小节将四个源码文件串成一条完整的数据流共享结构体定义 → AI CPU 写入 Tiling → AI Core 消费 Tiling → Host 编排时序。4.1 共享 Tiling 结构体kernel_args.hkernel_args.h 定义了两类结构体均置于KernelInfo命名空间中struct TilingInfo { int8_t type; int8_t mode; int8_t len; }; struct KernelArgs { uint32_t* xDevice; uint32_t* yDevice; uint32_t* zDevice; TilingInfo* ti; // Parameters shared with aicore are used for synchronizing tiling selection };TilingInfo是设备端 Tiling 参数本体由 AI CPU 算子写入、AI Core 算子读取。本样例仅用type/mode/len三个字段模拟真实场景下的分片决策如数据类型选择、计算模式、数据长度KernelArgs是内核参数打包结构xDevice/yDevice/zDevice为设备端数据指针ti指向 Tiling 参数所在设备内存。关键点在于该头文件被aicpu_tiling.aicpu与aicore_kernel.asc同时 include保证了两个算子看到的是同一份内存布局这是设备端跨算子传递 Tiling 参数的前提。4.2 AI CPU 算子设备端计算并写入 Tilingaicpu_tiling.aicpuaicpu_tiling.aicpu 是整个流程的生产者__global__ __aicpu__ uint32_t MyAicpuKernel(KernelInfo::KernelArgs args) { AscendC::printf(MyAicpuKernel inited\n); args.ti-type 1; args.ti-mode 2; args.ti-len 4; AscendC::DataStoreBarrier(); AscendC::printf(MyAicpuKernel inited type %u mode %u len %u end!\n, args.ti-type, args.ti-mode, args.ti-len); return 0; }要点解读__global__ __aicpu__修饰符声明这是一个在 AI CPU 上运行的内核函数与 AI Core 内核的__global__ __aicore__修饰符相区别Tiling 计算发生的位置样例中以直接赋值type1, mode2, len4代替真实的 Tiling 计算逻辑真实场景中此处会根据 shape、内存容量等条件计算分片写入目标是指向设备内存的args.tiAscendC::DataStoreBarrier()数据存储屏障。AI CPU 与 AI Core 是设备上不同的执行单元AI CPU 写入的 Tiling 参数必须经过屏障保证其先于后续操作可见AI Core 算子才能可靠地读到正确值——这是设备端 Tiling 下沉正确性的关键一步内核打印的两行日志MyAicpuKernel inited与MyAicpuKernel inited type 1 mode 2 len 4 end!与 README 中的执行结果完全对应。4.3 AI Core 算子按 Tiling 参数选择模板aicore_kernel.ascaicore_kernel.asc 是流程的消费者实现了基于 Tiling 参数的模板分支选择template typename T, int8_t mode, int8_t len __aicore__ void hello_world_impl(__gm__ uint8_t* m) { if constexpr (std::is_same_vT, float) { AscendC::printf(Hello World: float mode %u len %u m %d.\n, mode, len, *m); } else if constexpr (std::is_same_vT, int) { AscendC::printf(Hello World: int mode %u len %u m %d.\n, mode, len, *m); } } template typename T, int8_t mode, int8_t len __mix__(1, 2) __global__ __aicore__ void hello_world(__gm__ uint8_t* m, __gm__ uint8_t* TilingPtr) { __gm__ struct KernelInfo::TilingInfo* ti (__gm__ struct KernelInfo::TilingInfo*)TilingPtr; // Select different templates based on tiling value 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); } } extern C void hello_world_do(uint32_t numBlocks, void* stream, uint8_t* m, uint8_t* ti) { hello_worldint, 10, 201, 0, stream(m, ti); }要点解读从设备内存读取 Tiling内核入口将TilingPtr强转为__gm__ KernelInfo::TilingInfo*直接解引用设备端__gm__Global Memory指针读取 AI CPU 写入的type/mode/len模板分支选择根据ti-type/mode/len的实际值选择不同的hello_world_implT, mode, len模板实例。由于 Tiling 参数在运行时才确定而模板在编译期实例化这里实际上是用运行时 if 判断 编译期if constexpr的组合把运行期决策映射到编译期实例——这正是 Tiling 参数驱动不同计算路径的典型写法__mix__(1, 2)混合调度声明该内核同时占用1 个 Cube 执行单元和 2 个 Vector 执行单元因此 README 中特别说明 Hello World 日志会打印 3 次详见执行结果解读extern C包装hello_world_do以 C 符号导出供main.asc中的 C Host 代码调用其内部真正以1, 0, stream下发 AI Core 内核。4.4 Host 端编排双 Stream Event 同步main.ascmain.asc 是主机端入口完整演示了AI CPU 与 AI Core 在不同 Stream 上启动、以 Event 建立依赖的编排方式整体流程如下int32_t main(int argc, char const* argv[]) { aclInit(nullptr); int32_t deviceId 0; aclrtSetDevice(deviceId); aclrtStream aicpu_stream nullptr; aclrtStream aicore_stream nullptr; aclrtCreateStream(aicpu_stream); aclrtCreateStream(aicore_stream); aclrtEvent event; aclrtCreateEventExWithFlag(event, ACL_EVENT_SYNC); void* zDevice; void* ti; void* zHost; aclrtMalloc((void**)zDevice, 4096, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc((void**)ti, 4096, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMallocHost((void**)zHost, 4096); aclrtMemset((void*)ti, 4096, 0, 4096); aclrtMemset((void*)zHost, 4096, 10, 4096); aclrtMemcpy(zDevice, 4096, zHost, 4096, ACL_MEMCPY_HOST_TO_DEVICE); struct KernelInfo::KernelArgs args {0}; args.xDevice (uint32_t*)zDevice; args.yDevice args.xDevice 1; args.zDevice args.yDevice 1; args.ti (KernelInfo::TilingInfo*)ti; MyAicpuKernel1, 0, aicpu_stream(args); aclrtRecordEvent(event, aicpu_stream); aclrtStreamWaitEvent(aicore_stream, event); hello_world_do(1, aicore_stream, (uint8_t*)zDevice, (uint8_t*)ti); aclrtSynchronizeStream(aicpu_stream); aclrtSynchronizeStream(aicore_stream); ... }按执行顺序拆解步骤代码作用1aclInit/aclrtSetDevice(0)初始化 ACL 运行时并指定设备2aclrtCreateStream× 2创建aicpu_stream与aicore_stream两条独立流3aclrtCreateEventExWithFlag(event, ACL_EVENT_SYNC)创建同步事件用于跨流同步4aclrtMalloc/aclrtMemset/aclrtMemcpy分配设备内存数据区zDevice、Tiling 区ti并初始化Tiling 区清零、数据区全部置为 10这也是最终日志中m 10的来源args.xDevice/yDevice/zDevice依次指向同一段内存中的相邻 uint32 单元5MyAicpuKernel1, 0, aicpu_stream(args)在aicpu_stream上下发 AI CPU 内核执行 Tiling 计算并写入ti6aclrtRecordEvent(event, aicpu_stream)在 aicpu_stream 中记录 event标记AI CPU 内核已完成7aclrtStreamWaitEvent(aicore_stream, event)阻塞aicore_stream直到该 event 完成——保证 AI Core 内核读取到的 Tiling 参数是 AI CPU 写好的最终值8hello_world_do(1, aicore_stream, ...)在 aicore_stream 上下发 AI Core 内核9aclrtSynchronizeStream× 2 / 资源释放等待两条流全部执行完成释放内存与流、复位设备、aclFinalize这正是 README 所述AI CPU 算子与 AI Core 算子在不同 stream 上进行 launch……使用aclrtRecordEvent在指定 stream 中记录 event使用aclrtStreamWaitEvent阻塞指定的 stream直到指定的 event 完成的完整源码级呈现。五、编译运行与结果验证5.1 配置环境变量请根据当前环境上 CANN 开发套件包的安装方式配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。同时请确保当前环境已安装与目标产品匹配的 CANN 版本参考支持的产品与 CANN 版本一节。5.2 编译与执行在样例目录Samples/0_Introduction/04_aicpu/00_device_tiling/aicpu_device_tiling下执行mkdir -p build cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # 编译工程 ./demo # 执行编译生成的可执行程序执行样例5.3 编译选项说明选项可选值说明CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品 与 Atlas A3 训练系列产品/Atlas A3 推理系列产品dav-3510 对应 Ascend 950PR/Ascend 950DT该选项在 CMakeLists.txt 中以CACHE STRING形式声明默认值为dav-2201并作为--npu-arch编译参数传入 ASC 编译器target_compile_options中--npu-arch${CMAKE_ASC_ARCHITECTURES}。请在编译前按实际产品选择对应架构。另外构建脚本中还包含一个实验性编译选项# Experimental option. Future compatibility and support are not guaranteed. --cce-aicpu-launch-with-interface该选项启用后AI CPU 算子得以通过内核调用符...的接口方式被直接 launch这也是main.asc中MyAicpuKernel1, 0, aicpu_stream(args)能工作的原因。注释明确标注其为实验性特性未来兼容性不受保证实战项目中请关注 CANN 版本的演进说明。5.4 执行结果解读执行成功时输出如下MyAicpuKernel inited MyAicpuKernel inited type 1 mode 2 len 4 end! Hello World: int mode 2 len 4 m 10. Hello World: int mode 2 len 4 m 10. Hello World: Hello World: int mode 2 len 4 m 10.逐行对应关系前两行来自AI CPU 算子MyAicpuKernel先是进入时的MyAicpuKernel inited写入type1, mode2, len4并执行DataStoreBarrier后打印MyAicpuKernel inited type 1 mode 2 len 4 end!后三行来自AI Core 算子AI Core 读取 Tiling 参数命中type1 mode2 len4分支实例化hello_world_implint, 2, 4打印Hello World: int mode 2 len 4 m 10.m 10对应数据区被初始化为 10。其中打印 3 次的原因正如 README 所解释__mix__(1, 2)会启动1 个 Cube 执行单元和 2 个 Vector 执行单元每个执行单元都会执行一次内核体因此日志共出现 3 次。六、关键机制深度解析6.1 内核调用符...的统一抽象本样例最值得注意的一点是AI CPU 算子与 AI Core 算子均使用内核调用符...进行调用README 原文。传统上 AI CPU 算子通常通过独立的 AICPU 接口下发而本工程通过实验性编译选项--cce-aicpu-launch-with-interface统一了二者的调用语法MyAicpuKernel1, 0, aicpu_stream(args); // AI CPU 内核aicpu_stream hello_worldint, 10, 201, 0, stream(m, ti); // AI Core 内核hello_world_do 内部aicore_stream这种统一抽象让 Host 代码可以用同一套心智模型管理两类算子块数本样例均为 1、共享内存0、所属 Stream以及参数列表。6.2 Event 跨流同步谁先谁后的保证AI CPU 与 AI Core 运行在不同执行单元上两条 Stream 之间没有天然的执行顺序。样例用aclrtRecordEventaclrtStreamWaitEvent建立显式依赖aclrtRecordEvent(event, aicpu_stream)在 aicpu_stream 上打点当该点之前的所有任务即MyAicpuKernel完成后 event 即被触发aclrtStreamWaitEvent(aicore_stream, event)aicore_stream 会阻塞等待该 event从而保证 AI Core 内核启动时AI CPU 写入的 Tiling 参数已经就绪。这一机制避免了aclrtSynchronizeStream式的全量同步只让依赖方aicore_stream等待必要的事件两条流仍可并行推进其他无关任务是下沉 流式并行场景的标准做法。6.3 DataStoreBarrier跨执行单元的数据可见性aicpu_tiling.aicpu在写完 Tiling 参数后调用了AscendC::DataStoreBarrier()。结合源码结构可以推断AI CPU 与 AI Core 是不同的执行单元AI CPU 对设备内存的写入需要经过屏障才能保证对后续跨执行单元的访问可见、有序。若不设置该屏障AI Core 算子可能读到 Tiling 参数的中间态导致模板分支选择错误——这是设备端 Tiling 下沉中极易踩坑的正确性问题样例以极简形式给出了标准示范。6.4 模板分支选择运行期决策 → 编译期实例AI Core 算子中Tiling 参数在运行期由 AI CPU 算出但内核要执行的代码路径float还是int、mode/len取何值在编译期就需要确定。样例的解法是两层配合外层运行期if (ti-type ... ...)对 Tiling 值做分支判断内层hello_world_implT, mode, len用if constexpr (std::is_same_vT, float)在编译期消除不匹配的分支。于是同一份内核源码可以针对多种 Tiling 组合分别实例化出专用版本既保证了正确性又避免了运行期分支开销——这一Tiling 驱动模板实例化的模式在真实算子如各类 elementwise、矩阵类算子中非常常见。6.5 构建基础设施ASC AICPU 混合工程从 CMakeLists.txt 可以看到支撑本样例的完整构建链路find_package(ASC REQUIRED)与find_package(AICPU REQUIRED)分别引入 ASCAscend C 内核与 AICPUAI CPU 算子的构建支持project(... LANGUAGES ASC AICPU CXX)一个工程同时编译三种语言源码LINKER_LANGUAGE ASC可执行文件的链接以 ASC 工具链为主导保证内核与 Host 代码正确链接链接pthread与dl满足 ACL 运行时与动态加载需求。对于希望把本样例改造成真实算子的开发者这一 CMake 骨架可以直接复用。七、从样例到实战注意事项与扩展方向结合上文源码分析将本样例迁移到真实算子开发时建议关注以下几点Tiling 计算逻辑替换样例中用固定赋值模拟 Tiling真实场景应在MyAicpuKernel中根据输入 shape、数据类型、UB/L1 容量等运行时信息计算分片参数并扩展TilingInfo结构体承载更多字段注意 AI CPU 与 AI Core 共享的头文件要保持同步修改同步点不可省略AI CPU 写 Tiling 后的DataStoreBarrier()与 Host 端的aclrtStreamWaitEvent分别保证了写入可见与跨流有序二者缺一不可内存布局与指针语义样例中xDevice/yDevice/zDevice指向同一段设备内存的相邻单元真实场景建议按需独立分配并明确各指针的__gm__语义避免别名混淆多核与混合调度__mix__(1, 2)展示了 Cube/Vector 混合占用语法真实算子应结合计算特征选择执行单元组合并用块数参数样例中为 1控制核数扩展版本与架构匹配编译选项CMAKE_ASC_ARCHITECTURES必须与目标产品匹配dav-2201/dav-3510CANN 版本需满足第一节表格中的最低要求实验性编译选项--cce-aicpu-launch-with-interface的长期兼容性需以对应 CANN 版本发布说明为准。八、小结本样例以最小的代码量约 150 行完整呈现了AI CPU 计算 Tiling AI Core 消费 Tiling的设备端 Tiling 下沉全链路共享结构体统一两端的内存视图AI CPU 内核写入参数并做数据屏障AI Core 内核按参数做模板分支Host 端以双 Stream Event 建立正确的执行时序。其核心工程要点——__global__ __aicpu__/__global__ __aicore__内核声明、...统一调用符、aclrtRecordEvent/aclrtStreamWaitEvent跨流同步、AscendC::DataStoreBarrier()数据可见性保证、if constexpr模板实例化——构成了昇腾算子 Tiling 下沉开发的通用方法论可直接迁移到更复杂的真实算子工程中。【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考