企业官网建设流程全解析

昇腾 AI CPU 算子开发：3 步打通设备端 Tiling 下沉

【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples

先抛个问题

Host 端算完 Tiling、回传结果、再随内核下发——一来一回，轻算子被拖成了重流水。更反直觉：Tiling 逻辑往往比 compute 还复杂，Host CPU 反倒成了瓶颈。

读完这篇，你能把 Tiling 下沉到 AI CPU 上跑，用跨流 event 把 AI Core 的启动卡在 Tiling 之后，一次跑通。

它到底在干什么

用餐厅类比：备菜厨师（AI CPU）写好便签（Tiling 值），服务员（stream + event）递到灶台，主厨（AI Core）看便签决定用哪套锅（模板分支）。便签就夹在一块共享内存里，两方都能看见。

先看一眼传统流程——下沉之前 Host 是怎么干的：

图里"调用 Tiling 函数"那一步就在 Host 侧完成，tiling 数据跟着KernelLaunch一起下发到 NPU。本样例把这一整块挪进 NPU 内部：AI CPU 先把 Tiling 算出来，写进设备内存，AI Core 从同一块内存读。

时序：

1. Host 初始化：aclInit+aclrtSetDevice
2. 建两条 stream：aicpu_stream跑 Tiling、aicore_stream跑 compute
3. 建一个ACL_EVENT_SYNCevent，跨流依赖的锚点
4. 设备侧 alloc 一块ti（Tiling 结果容器）+ 一块zDevice（输入）
5. Host 侧memset填 10、H2D 拷到zDevice
6. 组装KernelArgs（x/y/z 指针 +ti指针）
7. AI CPU 内核MyAicpuKernel<<<1, 0, aicpu_stream>>>(args)下发
8. aclrtRecordEvent(event, aicpu_stream)在 AICPU 流上打标
9. aclrtStreamWaitEvent(aicore_stream, event)让 AI Core 流等待
10. AI Core 内核通过hello_world_do封装的<<<>>>下发
11. 两条流各自aclrtSynchronizeStream，释放资源，收尾

产品与版本：

编译架构：

代码骨架：从入口到内核

第 1 步：初始化与资源准备

先起来、两条流分开跑，这是跨核编排的语法基础。

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);

ACL_EVENT_SYNC表明这个 event 用来做 stream 间同步——不是普通的完成通知，而是"跨流依赖"的锚点。两条流分开，就是为了下一步让 AI Core 卡在 Tiling 之后。

⚠️ 事件必须是ACL_EVENT_SYNC才支持跨流 wait，用默认 event 会踩坑。

第 2 步：设备内存 + 数据准备

Tiling 结果的容器（ti）要清零，输入（zDevice）要从 Host 拷进去。

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);

ti是 Tiling 结果的落地位置，AI CPU 写、AI Core 读；zHost是 Host 侧数据源，填 10 后 H2D 到zDevice。

⚠️ti必须清零，否则 AI Core 可能读到脏值——这块内存不是首次分配就不保证干净。

第 3 步：组装共享参数包

把设备指针打包成KernelArgs，两个内核都从它拿。

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;

x/y/z在本样例里没实际用，只是占位；真正跨核传值的是args.ti这一条指针。

第 4 步：AI CPU 内核下发 + 跨流 event 同步

Tiling 计算放在aicpu_stream，AI Core 计算放在aicore_stream，中间用 event 串起来。

MyAicpuKernel<<<1, 0, aicpu_stream>>>(args); aclrtRecordEvent(event, aicpu_stream); aclrtStreamWaitEvent(aicore_stream, event);

<<<1, 0, stream>>>是内核调用符，三个参数分别是 blocks、共享内存、stream。第二行在aicpu_stream上打 event，第三行让aicore_stream阻塞等这个 event 完成——AI CPU 的 Tiling 写完之后，AI Core 才允许启动。

⚠️ 顺序不能反：先 Record、再 Wait。反过来aicore_stream会先跑到 AI Core，然后 event 才打出来，同步就失效了。

第 5 步：AI Core 内核下发

AI Core 侧用extern "C"封装了一个 host 可调用的函数，内部再走<<<>>>。

hello_world_do(1, aicore_stream, (uint8_t*)zDevice, (uint8_t*)ti);

hello_world_do是 aicore_kernel.asc 里的extern "C"函数，内部会调用hello_world<int, 10, 20><<<1, 0, stream>>>(m, ti)。

第 6 步：收尾同步 + 释放

两条流各自同步，再按"先同步 → 再释放 → 后销毁"的顺序收尾。

aclrtSynchronizeStream(aicpu_stream); aclrtSynchronizeStream(aicore_stream); aclrtFree(zDevice); aclrtFree(ti); aclrtFreeHost(zHost); aclrtDestroyStream(aicpu_stream); aclrtDestroyStream(aicore_stream); aclrtResetDevice(deviceId); aclFinalize();

先同步是为了保证两条流上的任务都跑完，再释放资源才不会踩"释放正在使用的内存"这种坑。

跨核共享结构体：TilingInfo与KernelArgs

两个内核对"设备内存里 Tiling 值长什么样"必须有一致的理解，这就是 kernel_args.h 存在的意义。

namespace 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; }; }

• TilingInfo是 Tiling 的计算产物，本样例里type=1, mode=2, len=4作为 AI Core 端模板分派的判别条件；
• KernelArgs是双算子共享的参数包，ti指针指向设备内存里的 TilingInfo，AI CPU 写、AI Core 读；
• 头文件被.aicpu与.asc两个内核同时 include，任何一方改字段都必须双端同步——这是 Tiling 下沉的正确性底线。

内核内部：谁写谁读

写入 → 屏障 → 读取 → 分派，四步串起来。

写入（aicpu_tiling.aicpu）：AI CPU 直接把三个字段写到ti指向的设备内存。

args.ti->type = 1; args.ti->mode = 2; args.ti->len = 4;

屏障：DataStoreBarrier保证写入对后续读取方可见，没有它 AI Core 可能读到未刷新的旧值。

AscendC::DataStoreBarrier();

读取（同一份 aicore_kernel.asc）：AI Core 入口把TilingPtr强转成__gm__ TilingInfo*，从全局内存读三个字段。

__gm__ struct KernelInfo::TilingInfo* ti = (__gm__ struct KernelInfo::TilingInfo*)TilingPtr;

分派：根据 Tiling 值命中不同模板实例。

if (ti->type == 0 && ti->mode == 1 && ti->len == 2) { hello_world_impl<float, 1, 2>(m); } else if (ti->type == 1 && ti->mode == 2 && ti->len == 4) { hello_world_impl<int, 2, 4>(m); }

如果把 Tiling 值从 (1, 2, 4) 改成 (0, 1, 2)，命中的是float, 1, 2这条分支——但编译产物里另一条 (1, 2, 4) 的分支仍然存在，只是运行时不会走到。这就是"tiling 决定 kernel 变体"的最小示范：编译期实例化所有变体，运行时按 Tiling 值挑一个。

__mix__(1, 2)让内核同时占用 1 个 Cube 执行单元和 2 个 Vector 执行单元，所以Hello World日志会打 3 次。

亲手跑一遍

1. 环境准备

source ${install_path}/cann/set_env.sh

${install_path}是 CANN 包安装目录，默认/usr/local/Ascend。

2. 编译

mkdir -p build && cd build cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 .. make -j

3. 运行

./demo

预期输出：

MyAicpuKernel inited # AI CPU 内核入口 MyAicpuKernel inited type 1 mode 2 len 4 end! # Tiling 写入完成 + DataStoreBarrier Hello World: int mode 2 len 4 m 10. # AI Core Cube 单元打印 Hello World: int mode 2 len 4 m 10. # AI Core Vector 单元打印 Hello World: int mode 2 len 4 m 10. # AI Core Vector 单元打印

如果输出和这里不一样，大概率是CMAKE_ASC_ARCHITECTURES和实际硬件不匹配——编译时选的架构和aclrtSetDevice指定的设备不是一家人。

带走这三句话

契约先行：跨核共享的 Tiling 结构体必须由.aicpu和.asc两个内核共同 include 的同一头文件定义，字段变更双端同步，这是 Tiling 下沉的正确性底线。

顺序靠 event，可见性靠 barrier：不同 stream 上的两个内核用aclrtRecordEvent+aclrtStreamWaitEvent建立依赖；AI CPU 写完 Tiling 后必须DataStoreBarrier，AI Core 才能读到新值。

Tiling 决定 kernel 变体：AI Core 侧用 Tiling 值驱动模板分支，编译期实例化所有变体、运行时挑一个——这就是"一次下发、按 Tiling 自适应"的最小可运行范式。

延伸阅读：

• Samples/0_Introduction/04_aicpu/README.md
• Samples/0_Introduction/04_aicpu/00_device_tiling/aicpu_device_tiling/README.md

【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples