Architecture

分层

框架或用户代码
        │
        ▼
cann_ops_tilelang 公共 API
  - 参数语义
  - device/dtype/shape/stride 校验
  - 输出内存分配
        │
        ▼
Ascend kernel generator
  - TileLang JIT 特化参数
  - block/core/pipeline 配置
  - GM/UB 数据搬运
  - SimdVF/SimtVF/Cube 计算
        │
        ▼
TileLang Ascend lowering / CANN toolchain
        │
        ▼
Ascend NPU

模块边界

公共 API

公共 API 必须只接受和返回 PyTorch Tensor、标量或简单配置对象。API 层负责:

  • 确认运行设备为 Ascend NPU;
  • 校验 shape、dtype、stride、layout 和可选输出;
  • 计算不会依赖设备 IR 的 launch 参数;
  • 调用 kernel generator 并执行 kernel;
  • 对空 Tensor 给出定义良好的行为。

Kernel generator

Kernel generator 使用 @tilelang.jit,参数分成两类:

  • 编译期参数:dtype、block size、是否启用某个融合分支;
  • 运行期动态参数:token 数、序列长度等可通过 T.dynamic 表达的维度。

不要把高基数、频繁变化的运行期值错误地放进编译期参数,否则会导致编译缓存膨胀。

配置层

config.py 只处理 Ascend 环境探测、核心数和调试开关。配置查询使用惰性调用,包导入阶段不得访问设备。

测试层

测试辅助工具和 PyTorch reference 全部位于 tests/,不属于运行时包,也不会进入 wheel。

测试框架由测试目录根部的 tests/conftest.py 统一装载插件,每个算子测试分为:

  1. CPU 可执行的导入和参数校验;
  2. NPU 正确性测试;
  3. 显式启用的 benchmark;
  4. NPU 峰值显存画像和 large_npu_mem 自动分组;
  5. 按测试文件分片、按芯片集群选择的性能基线。

xdist worker 在导入 TileLang 前通过 ASCEND_RT_VISIBLE_DEVICES 绑定 NPU。普通测试和大/未知显存测试分两阶段执行;benchmark 独占每张 NPU。随机种子由 --seed 与 node ID 共同生成,保证并行调度变化时每个用例仍可复现。

PyTorch reference 层

每个算子域的设备无关 PyTorch reference 与对应测试放在一起,例如 tests/engram/engram_gate_ref.py。reference 只作为数值语义、梯度和测试 oracle,不参与生产 kernel dispatch,也不构成 cann_ops_tilelang 的公开接口。Engram gate 的前向和反向均以该实现为基准。

示例调用链

engram_fused_weight(weight_hidden, weight_embed) 的调用链:

cann_ops_tilelang.engram_fused_weight
  -> validate BF16/NPU/2-D/contiguous inputs
  -> choose block size and persistent scheduling from total elements
  -> get_engram_fused_weight_kernel(hidden_size, hc_mult)
  -> allocate an FP32 output
  -> launch flattened tensors
  -> compare with the FP32 PyTorch reference

Engram gate 训练链路:

engram_fused_weight(weight_hidden, weight_embed)
  -> engram_gate_fwd(..., save_for_backward=True)
  -> output + (dot, gate_score, rstd_x, rstd_k)
  -> engram_gate_bwd(...)
  -> grad_hidden_states + grad_kv + grad_w_partial
  -> engram_grad_w_reduce(...)
  -> update grad_weight_hidden and grad_weight_embed in place

MHC 训练链路由显式前向和反向原语组成:

expand_to_mhc_forward(...)
  -> mhc_fn_normw_merge_forward(...)
  -> mhc_gemm_with_sqrsum_forward(...)
  -> mhc_reduce_partials_and_rmsnorm_forward(...)
  -> mhc_pre_split_mixes_forward(...)
  -> mhc_sinkhorn_forward(...)
  -> mhc_pre_apply_mix_forward(...)
  -> upper-layer computation
  -> mhc_post_forward(...)

mhc_post_backward(...)
  -> upper-layer backward
  -> mhc_pre_apply_mix_backward(...)
  -> mhc_sinkhorn_backward(...)
  -> mhc_pre_split_mixes_backward(...)
  -> mhc_reduce_partials_and_rmsnorm_backward(...)
  -> mhc_gemm_with_sqrsum_backward(...)
  -> mhc_fn_normw_merge_backward(...)
  -> expand_to_mhc_backward(...)

mhc_head_compute_mix_forward 和 mhc_head_compute_mix_backward 处理 head mixing 系数,mhc_multilayer_recompute 融合多层 pre/post mix 重计算。mhc_pre_big_fuse 融合 reduce/RMSNorm、split mixes、Sinkhorn 和 pre apply mix 的前向路径。

Engram fused weight 的 torch.compile 接入

engram_fused_weight 支持固定形状推理的 npugraph_ex 编译路径:

import torch
import torch_npu

from cann_ops_tilelang import engram_fused_weight


def model(weight_hidden, weight_embed):
    return engram_fused_weight(-weight_hidden, weight_embed) + 1


compiled_model = torch.compile(
    model, backend='npugraph_ex', fullgraph=True, dynamic=False,
)
with torch.inference_mode():
    a = torch.randn(4, 4096, device='npu', dtype=torch.bfloat16)
    b = torch.randn_like(a)
    result = compiled_model(a, b)

编译时,公共接口调用 cann_ops_tilelang::engram_fused_weight 自定义算子;其 fake 实现只校验 shape、dtype、device 和布局,并推导同形状的 FP32 输出。实际执行仍使用原有 TileLang 内核。 公共接口只接收两个输入张量,直接返回 FP32 结果,不修改输入。普通 eager 调用保持原有路径。

本仓库要求 PyTorch >= 2.10;此接入使用 torch.library.custom_op,并需要提供 npugraph_ex 的配套 torch_npu。实机验证环境为 PyTorch / torch_npu 2.10。 输入仍须为连续的二维 BF16 NPU 张量,非空元素数须能被 256 整除;空输入也受支持。 这里按形状编译,形状变化可能触发重新编译,未注册 autograd 公式。 首次执行包含 TileLang JIT 和图准备开销。

以下测试包含原有 eager 用例和一个固定形状的 compile 基本功能用例,验证 npugraph_ex 编译执行的结果与 FP32 参考结果一致:

OMP_NUM_THREADS=1 OPENBLAS_NUM_THREADS=1 OPS_TILELANG_TEST_LEVEL=0 \
    python -m pytest tests/engram/test_engram_fused_weight.py -m 'not benchmark' -v

扩展方向

  • src/cann_ops_tilelang/quant/:公共 scaling-factor layout 和 FP8/FP4 类型系统;
  • src/cann_ops_tilelang/moe/:在现有 TopK、routing 和 expert mapping 基础上扩展 expert dispatch/combine;
  • src/cann_ops_tilelang/norm/:RMSNorm、LayerNorm 和融合残差;
  • src/cann_ops_tilelang/attention/:RoPE、KV Cache、prefill/decode;
  • tests/<operator>/benchmark/:在算子测试目录内按测试文件存储 JSONL 性能基线。
  • tests/<operator>/npu_mem_profiles/:在算子测试目录内按测试文件存储 JSON 峰值显存画像。