
在 GPU 编程中,开发者常关注“显存 vs 寄存器”;而在 Ascend NPU 上,真正的性能战场在 Unified Buffer(UB) ——一块仅 256KB(310P)或 512KB(910B) 的片上高速缓存。 若不能高效利用 UB,再精妙的算法也会被内存墙拖垮。CANN 编译器的核心任务之一,就是将数据流精准调度到这有限的片上空间。
Ascend 芯片采用层次化内存设计,以平衡容量与带宽:
存储层级 | 容量 | 带宽 | 特性 |
|---|---|---|---|
Global Memory(HBM/DDR) | 8~64 GB | ~100 GB/s | 主存,CPU/NPU 共享 |
L2 Cache | 4~8 MB | ~800 GB/s | 芯片级缓存,自动管理 |
Unified Buffer(UB) | 256KB / 512KB | >3 TB/s | 手动调度,NPU 计算单元直连 |
⚠️ 关键:UB 是唯一可编程的片上存储,所有 Cube 计算必须从 UB 读取数据。
NPU 的 Cube 单元理论算力高达 256 TFLOPS(FP16),但若数据供给不足,实际利用率可能低于 20%。
UB 的使命:作为 Global Memory 与 Cube 之间的“蓄水池”,通过预取 + 分块,确保计算单元永不“饿死”。
理想数据流:
Global Memory → (DMA) → UB → (Load) → Cube Registers → Compute✅ 目标:让 DMA 与计算重叠(Overlap),隐藏内存延迟。
当 ATC 处理一个算子(如 Conv)时,会执行以下内存规划:
为实现 DMA 与计算重叠,CANN 自动划分 UB 为两个区域:
Cube UB_B UB_A DMA Cube UB_B UB_A DMA #mermaid-svg-wYOGZcKHPjRYXBHU{font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;fill:#333;}@keyframes edge-animation-frame{from{stroke-dashoffset:0;}}@keyframes dash{to{stroke-dashoffset:0;}}#mermaid-svg-wYOGZcKHPjRYXBHU .edge-animation-slow{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 50s linear infinite;stroke-linecap:round;}#mermaid-svg-wYOGZcKHPjRYXBHU .edge-animation-fast{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 20s linear infinite;stroke-linecap:round;}#mermaid-svg-wYOGZcKHPjRYXBHU .error-icon{fill:#552222;}#mermaid-svg-wYOGZcKHPjRYXBHU .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-wYOGZcKHPjRYXBHU .edge-thickness-normal{stroke-width:1px;}#mermaid-svg-wYOGZcKHPjRYXBHU .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-wYOGZcKHPjRYXBHU .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-wYOGZcKHPjRYXBHU .edge-thickness-invisible{stroke-width:0;fill:none;}#mermaid-svg-wYOGZcKHPjRYXBHU .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-wYOGZcKHPjRYXBHU .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-wYOGZcKHPjRYXBHU .marker{fill:#333333;stroke:#333333;}#mermaid-svg-wYOGZcKHPjRYXBHU .marker.cross{stroke:#333333;}#mermaid-svg-wYOGZcKHPjRYXBHU svg{font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-wYOGZcKHPjRYXBHU p{margin:0;}#mermaid-svg-wYOGZcKHPjRYXBHU .actor{stroke:hsl(259.6261682243, 59.7765363128%, 87.9019607843%);fill:#ECECFF;}#mermaid-svg-wYOGZcKHPjRYXBHU text.actor>tspan{fill:black;stroke:none;}#mermaid-svg-wYOGZcKHPjRYXBHU .actor-line{stroke:hsl(259.6261682243, 59.7765363128%, 87.9019607843%);}#mermaid-svg-wYOGZcKHPjRYXBHU .innerArc{stroke-width:1.5;stroke-dasharray:none;}#mermaid-svg-wYOGZcKHPjRYXBHU .messageLine0{stroke-width:1.5;stroke-dasharray:none;stroke:#333;}#mermaid-svg-wYOGZcKHPjRYXBHU .messageLine1{stroke-width:1.5;stroke-dasharray:2,2;stroke:#333;}#mermaid-svg-wYOGZcKHPjRYXBHU #arrowhead path{fill:#333;stroke:#333;}#mermaid-svg-wYOGZcKHPjRYXBHU .sequenceNumber{fill:white;}#mermaid-svg-wYOGZcKHPjRYXBHU #sequencenumber{fill:#333;}#mermaid-svg-wYOGZcKHPjRYXBHU #crosshead path{fill:#333;stroke:#333;}#mermaid-svg-wYOGZcKHPjRYXBHU .messageText{fill:#333;stroke:none;}#mermaid-svg-wYOGZcKHPjRYXBHU .labelBox{stroke:hsl(259.6261682243, 59.7765363128%, 87.9019607843%);fill:#ECECFF;}#mermaid-svg-wYOGZcKHPjRYXBHU .labelText,#mermaid-svg-wYOGZcKHPjRYXBHU .labelText>tspan{fill:black;stroke:none;}#mermaid-svg-wYOGZcKHPjRYXBHU .loopText,#mermaid-svg-wYOGZcKHPjRYXBHU .loopText>tspan{fill:black;stroke:none;}#mermaid-svg-wYOGZcKHPjRYXBHU .loopLine{stroke-width:2px;stroke-dasharray:2,2;stroke:hsl(259.6261682243, 59.7765363128%, 87.9019607843%);fill:hsl(259.6261682243, 59.7765363128%, 87.9019607843%);}#mermaid-svg-wYOGZcKHPjRYXBHU .note{stroke:#aaaa33;fill:#fff5ad;}#mermaid-svg-wYOGZcKHPjRYXBHU .noteText,#mermaid-svg-wYOGZcKHPjRYXBHU .noteText>tspan{fill:black;stroke:none;}#mermaid-svg-wYOGZcKHPjRYXBHU .activation0{fill:#f4f4f4;stroke:#666;}#mermaid-svg-wYOGZcKHPjRYXBHU .activation1{fill:#f4f4f4;stroke:#666;}#mermaid-svg-wYOGZcKHPjRYXBHU .activation2{fill:#f4f4f4;stroke:#666;}#mermaid-svg-wYOGZcKHPjRYXBHU .actorPopupMenu{position:absolute;}#mermaid-svg-wYOGZcKHPjRYXBHU .actorPopupMenuPanel{position:absolute;fill:#ECECFF;box-shadow:0px 8px 16px 0px rgba(0,0,0,0.2);filter:drop-shadow(3px 5px 2px rgb(0 0 0 / 0.4));}#mermaid-svg-wYOGZcKHPjRYXBHU .actor-man line{stroke:hsl(259.6261682243, 59.7765363128%, 87.9019607843%);fill:#ECECFF;}#mermaid-svg-wYOGZcKHPjRYXBHU .actor-man circle,#mermaid-svg-wYOGZcKHPjRYXBHU line{stroke:hsl(259.6261682243, 59.7765363128%, 87.9019607843%);fill:#ECECFF;stroke-width:2px;}#mermaid-svg-wYOGZcKHPjRYXBHU :root{--mermaid-font-family:"trebuchet ms",verdana,arial,sans-serif;} 并行执行 Load next tile Compute on current tile Swap after compute
📌 双缓冲是 CANN 默认启用的高级优化(
--enable_double_buffer=true)。
在关键节点插入 sync,确保数据就绪:
dma_load ub_buf[0], global_addr
sync ; 等待 DMA 完成
mad cube_reg, ub_buf[0] ; 启动计算当你编写 TBE 算子时,需显式控制 UB 使用。
示例:矩阵乘 C = A × B,A: [M, K], B: [K, N]
# 错误:一次性加载整个 A(可能 >256KB)
# 正确:按 K 维度分块
BK = 128 # 根据 UB 容量计算
for ko in range(0, K, BK):
A_tile = A[:, ko:ko+BK] # ~64KB
B_tile = B[ko:ko+BK, :] # ~64KB
C += matmul(A_tile, B_tile)💡 经验公式:
tile_size ≈ sqrt(UB_size / (2 * dtype_size))(FP16 下,256KB UB → tile ≈ 256)
UB 被划分为 32 个 bank,每个 bank 128B。若多个线程同时访问同一 bank,会串行化。
冲突示例:
// 所有线程读取地址 0, 128, 256... → 全部命中 bank 0
float x = ub[128 * tid];规避方法:
s[C].avoid_bank_conflict()。在 GroupNorm 中,均值 mean 和方差 var 可驻留 UB,避免重复从 Global Memory 读取:
# 在 TBE Schedule 中
s[mean].set_scope("local.UB")
s[var].set_scope("local.UB") # 复用 UB 空间📊 实测:UB 复用可减少 30% DMA 传输量。
查看 UB 使用报告
tbe_debug --op=MyAttention --dump_ub_usage输出:
[INFO] Required UB: 312KB > Available: 256KB ❌定位大张量
优化方案:分块计算 Attention
# 按 Query 分块
for qo in range(0, 128, 64):
Q_tile = Q[qo:qo+64]
S = matmul(Q_tile, K) # Score tile
P = softmax(S)
O_tile = matmul(P, V)
output[qo:qo+64] = O_tile验证
tbe_debug --op=MyAttention --dump_ub_usage
# [INFO] Required UB: 198KB ✅✅ 修复后,算子成功编译,性能提升 1.8 倍。
特性 | CANN UB | CUDA Shared Memory |
|---|---|---|
容量 | 256KB~512KB | 48KB~96KB |
管理方式 | 编译器自动 + 手动调度 | 程序员手动分配 |
双缓冲 | 编译器自动插入 | 需手写 __syncthreads() |
Bank 结构 | 32 banks, 128B/bank | 32 banks, 4B/bank |
优化目标 | 最大化 Cube 利用率 | 最大化 SM occupancy |
💡 CANN 更“自动化”,但理解原理仍至关重要。
CANN 正在研发 Auto-Tiling 引擎:
🔮 目标:开发者只需写 Compute,Schedule 全自动。
在算力过剩的时代,带宽决定一切。UB 虽小,却是连接算法与硬件的咽喉要道。
掌握 CANN 的内存调度机制,意味着你不仅能“写出能跑的算子”,更能“写出极致高效的算子”。