Vector与Cube不能并发执行，需靠流水同步；必须显式插入同步点、严格区分buffer用途、确保layout匹配，并通过dump中间寄存器调试。
Vector 和 Cube 单元不能“并发执行”而是靠流水同步
昇腾 AICore 的 Vector 与 Cube 并非传统意义上的多线程并行，它们共享同一核内调度器，不支持任意时刻各自独立取指执行。所谓“协同”，本质是数据通路级的流水衔接：Cube 输出结果暂存 L0C（或 UB），Vector 紧接着从该位置读取并做逐元素计算。若跳过显式同步，Vector 可能读到未就绪的脏数据，导致结果随机错误，且极难复现。
实操建议：
- 必须调用
__bang_sync_thread()或__sync_threads()（取决于 API 层级）在 Cube 计算完成、Vector 启动前插入同步点； - 在
matmul_leakyrelu_basic_api示例中，Mmad指令后紧跟Fixpipe+__bang_sync_thread()，就是为确保 L0C 数据落稳； - 高阶 API（如
matmul::Matmul）会自动插入同步，但仅限其封装范围内；一旦你在 Vector 侧手动加了额外 Load/Store，仍需自己补同步； - 避免把同步写在条件分支里——编译器可能优化掉，造成部分线程不同步。
内存搬运路径决定协同效率，不是越“近”越好
常见误区是认为把 Cube 输出直接喂给 Vector 就一定快，实际上 L0C → Vector 输入寄存器之间仍有隐式格式转换开销。若 Cube 输出是 FRACTAL_NZ 格式，而 Vector 需要 ND 布局，中间必须经过 Copy 或 Cast，这步若没显式安排，运行时会 fallback 到低效路径甚至报错 Invalid memory layout for vector op。
实操建议：
- 查清两个单元对 Tensor layout 的硬性要求：
matmul::Matmul默认输出FRACTAL_NZ，__vector_op类函数通常要求ND或DEFAULT； - 在
basic_api中，必须用DataCopy+Cast显式转 Layout，例如：DataCopy(dst_ub, src_l0c); Cast(dst_ub_nd, dst_ub);； - advanced_api 虽隐藏了 layout 转换，但若你传入的 input Tensor layout 不符合约定（比如
weight未预转FRACTAL_NZ），编译期就会报Layout mismatch in matmul binding； - double buffer 机制只缓解搬运延迟，不消除 layout 不匹配问题——错的布局搬得再快也是错的。
Vector 计算不能“覆盖”Cube 的寄存器资源
Cube 单元使用 L0A/L0B/L0C 寄存器块做矩阵分块计算，Vector 单元也依赖同一片片上存储（UB）做中间结果暂存。若 Vector 在 Cube 还未释放 L0C 时就往同一地址写入，会破坏 Cube 的累加过程，引发数值溢出或 NaN 扩散。这不是竞态，而是物理寄存器冲突。
实操建议：
- 严格区分 buffer 用途：L0C 专供 Cube 写出，Vector 输入应绑定独立 UB buffer（如
ub_vec_in），而非复用l0c_out； - 在
basic_api的 kernel 中，检查所有__gm__, __l1__, __l0a__, __l0b__, __l0c__, __ub__地址是否互斥——尤其注意__ub__buffer 是否被 Cube 的LoadData和 Vector 的__vector_add同时引用； - 用
ascendc_check_memory_usage工具离线分析 buffer 分配图，比 runtime 报错Out of L0C memory更早发现问题； - 不要依赖“Cube 先写完就自动释放”的假设——L0C 生命周期由指令流显式控制，不是作用域自动管理。
调试时别只看最终输出，要抓中间寄存器快照
Vector+Cube 协同出错，90% 的 case 是中间值异常，但 LOG_DEBUG 打印全局内存（GM）结果往往掩盖了问题。比如 LeakyRelu 输入本该是 2.1，却因 Cube 未同步就提前被 Vector 读成 0，最终输出全零，但你只在 GM 打 log，看不出哪一环断了。
实操建议：
- 用
Ascend Debugger的dump_tensor功能，在 L0C 写完后、Vector 读取前，dump 出 L0C buffer 内容，确认数值和 layout 正确； - CPU 孪生模式下，可加断点在
Fixpipe后，用print_tensor查看模拟的 L0C 内存； - 避免在 Vector 计算内部插 log——
LOG_DEBUG本身会打乱流水节奏，可能掩盖或诱发同步问题； - 最有效的验证方式：单独跑 Cube kernel 输出到 GM，再单独跑 Vector kernel 读该 GM 地址，两者结果一致，才能说明协同链路里是协同逻辑而非单体 bug。
真正卡住人的从来不是“怎么写同步”，而是“同步该插在哪、buffer 该分几块、layout 该转几次”。这些细节不在文档首页，而在 matmul_leakyrelu_basic_api 的每一行注释和 tilelang 编译生成的 .asm 文件里。
C++免费学习笔记（深入）：立即使用
在学习笔记中，你将探索 C++ 的入门与实战技巧！
