PR功能描述 / 为什么需要这个合入:

新增 vocab-parallel embedding_lookup 算子，按 README 与 docs/adding_an_operator.md 的组织方式落入 tensor 业务域：

• src/cann_ops_tilelang/tensor/embedding_lookup.py：公共接口与契约校验。
• src/cann_ops_tilelang/tensor/embedding_lookup_kernel.py：get_embedding_lookup_kernel 与调度配置。
• 从 cann_ops_tilelang 包入口导出 embedding_lookup。

语义与接口：

• 对本地权重分片 [vocab_start_index, vocab_end_index) 做 gather；区间外 id 输出零行，因此各 TP rank 的分片结果相加可重建全表。
• 支持 int32/int64 ids（任意 rank，含 rank-0）、bfloat16/float16/float32 权重、连续输入；支持可选 out=；空输入不启动
  kernel。
• 非连续输入、非法 dtype、区间与权重行数不一致、out 形状/dtype/别名不合法等均给出明确异常，不做静默拷贝或回退。

Kernel 设计：

• AIV persistent 网格；满网格时每 tile 一次块 MTE3 存储，非整除的尾部行由前 tail_rows 个 core 各存一行，输出严格为
  input_ids.shape + (H,)，无 host 侧 reshape/pad/slice。
• miss 行不读权重（clamp 到 L2 行的读是纯浪费），直接在 UB 中清零。
• 行宽为 256B 向量寄存器整数倍时用 SimdVF 矢量清零，否则回退 SimtVF；UB 预算仍按 216 KiB（248 KiB 物理 − 32 KiB SimtVF
  context）规划。

为什么需要：生产推理使用 TP 切分的词表 embedding lookup；现有 tensor/gather 不提供分片区间零行语义和精确输出形状/尾行处理，本 PR
补齐该算子。

该PR关联的issue

fixes #（如无对应 issue，可填写 N/A）

希望检视人员了解:

• 正确性用 torch.equal 位精确比较；reference 使用 torch.where 产生 +0.0，避免 miss 行 ±0.0 的字节差异。
• per-row 路径对最后一个不满 tile 的 ids copy 用 valid_rows 截断，避免越界读；块存储只在网格填满且尾行可分配时启用。
• 性能证据见下方「性能证据」章节：测试环境、计时/带宽口径、benchmark 数据、流水利用率与复现命令。
• benchmark 基线与 NPU 显存画像不随 PR 提交，可用 --refresh-baseline / --refresh-mem 在本地生成。

性能证据

测试环境

计时与带宽口径

• 时延：仓库统一口径 tilelang.profiler.bench.do_bench(backend='msprof')（warmup=0、rep=30 ms），只计设备端 kernel 时间；msprof 对多次 kernel 时长取平均，轮间 flush L2。
• 带宽：(ids + 输出 + 命中行读) / 时延；命中行读 = shard_rows / vocab_size × num_tokens × hidden × element_size（本 rank 外的 id 在 UB 置零，不读表）。
• 口径核对：与 msprof op 的 Memory.csv 实测 HBM 流量对比，h7168 / h4096 / tp8_h4096 / fp32 / small 五个场景偏差 0.2–0.4%。
• 带宽基准：同结构（persistent + 多版本流水、无 gather）纯 copy kernel 实测 ≈3015 GB/s（标称 HBM 4 TB/s 的 75%），作为 100% 利用率。

实测性能（仓库 benchmark）

• 全表场景达到同构 copy 上限的 95–99%（同配置重复测量 315–327 µs，物理地址导致 ±4%）；
• 同 shape 下比 torch 原生 embedding 快约 6–10%；
• TP=8 场景带宽低于 copy 上限，是因为 87.5% 的 token 只写零、读流量本身很小，属写主导负载。

流水利用率（msprof op，56 block）

• 全命中 / fp32 场景读路径饱和（MTE2 98–99%），trace 中 MTE2 事件最大间隔 21.8 ns，流水无空泡；写路径被读完全掩盖；
• TP=8 场景转为写瓶颈（MTE3 84–90%），VEC 14–19% 为 miss 行置零（87.5% 的行），标量约 13–19%；
• 剩余差距来自随机 14 KiB 行的读效率（MTE2 active BW ≈26.5 GB/s/core，顺序读 ≈53.6 GB/s/core），不是调度或空泡问题。

采集口径说明

• PipeTimeline 不与 PipeUtilization 写在同一条 msprof op 命令里：56 block 超过 108 个 sub-block 的插桩上限，计数会失真（实测 MTE2 记为 0%、时长翻倍）。
• trace.json 只采样部分 lane（h7168 为 12/56），核利用率以 PipeUtilization.csv 的 Block Dim（56）为准。

复现命令

# 时延/带宽（仓库 benchmark）
OPS_TILELANG_TEST_LEVEL=2 python -m pytest tests/tensor/test_embedding_lookup.py -m benchmark --run-benchmark -q

# 流水占比（driver 侧）
msprof op --kernel-name=embedding_lookup_kernel_kernel --aic-metrics=PipeUtilization,Memory,L2Cache --output=/tmp/oprof <launch_script>

建议把性能基线固定到同一台机器：同配置下全表场景仍有 ±4% 的物理地址抖动，跨机器对比会引入额外噪声。

改动类型 / Change Type

• Bug 修复 / Bug Fix
• 新功能 / New Feature
• 性能优化 / Performance
• 代码重构 / Refactoring
• 文档更新 / Documentation
• 测试相关 / Test
• 其它 / Other

测试信息 / Testing

• 单元测试通过 / UT passed
  OPS_TILELANG_TEST_LEVEL=0/1/2 pytest tests/tensor/test_embedding_lookup.py -m 'not benchmark'
  → level 2 全量 176 passed（4 benchmark deselected）
• 集成测试通过 / ST passed
• 人工验证通过 / Manual verified
  Ascend 950DT 实机：benchmark 4/4 通过（基线文件不随 PR 提交）；CPU 用例 25 passed；
  ruff 0.14.14 check + format --check 全仓通过。

复现命令：

OPS_TILELANG_TEST_LEVEL=2 pytest tests/tensor/test_embedding_lookup.py -m 'not benchmark'
pytest tests/tensor/test_embedding_lookup.py -m benchmark --run-benchmark
./scripts/check.sh

检查清单 / Checklist

- [x] 代码符合规范 / Code follows style guide
- [x] 测试添加并通过 / Tests added and passed
- [x] 文档已更新 / Docs updated if needed（接口 docstring 写明支持范围、限制与空输入语义）
- [x] 无硬编码敏感信息 / No secrets hardcoded
- [x] 提交信息符合规范 / Commit message follows convention