功能

新增 cann_ops_tilelang.hisa_indexer，为 prefill 阶段的每个 query 选取可见 KV token。支持单序列和打包多序列，输出相对于各行可见窗口起点的 INT32 索引。

Q/K 为 FP8 E4M3FN，反量化系数、头权重和打分累加为 FP32。粗排将块分数舍入到 BF16，优先保留可见首尾块。精排使用 FP32 分数，同分时取较小的候选槽位，无效槽位填 -1。8192 个候选选 2048 个 token 时使用 radix-4 专用内核。

公开入口支持 eager 和单流、固定形状的 torch.compile 推理。图模式注册为 cann_ops_tilelang::hisa_indexer，沿用相同的签名和选择规则。最终索引由 PyTorch 的 NPU 分配器创建，TileLang 内核直接写入。接口与调用示例

验证

• 在昇腾 NPU 上完成 HISA 定向 L2。默认 AscendC 后端与隔离的 PTO 环境各通过 34 项功能、精度与图接入测试。独立 PyTorch 参考从原始 Q、K、scale、weights 和窗口计算最终索引，覆盖边界块、同分、空窗口和尾部形状。
• 三条 token top-k 路径均验证了 npugraph_ex 捕获、原地址换值和新地址换输入重放，逐次检查后端事件与参考索引。另在默认后端用完整 8192 单序列形状验证了三次图输出与 eager 的索引集合一致。用例还覆盖 schema/fake、空张量和与上下游 Torch 操作的组合。
• 默认后端六组 eager benchmark 与已有 pytest/msprof 基线逐组比较，全部命中，零回退。保留原 JSONL 与显存画像，新增图接入用例的显存峰值。Ruff、格式检查与打包通过。
• 已核对公开 CANN 相近索引器的粗选语义。8192 单序列的四阶段设备流水图、各流水线活跃率及 HBM 实测口径见性能剖析。

环境与复现

依赖 CANN、PyTorch、torch_npu 和 Ascend TileLang。测试覆盖默认 AscendC 后端及配套的 PTO 工具链。

OPS_TILELANG_TEST_LEVEL=2 TILELANG_DEFAULT_TARGET=auto python -m pytest -s tests/attention/test_dsa_hisa.py -k hisa_compile
OPS_TILELANG_TEST_LEVEL=2 TILELANG_DEFAULT_TARGET=auto python -m pytest tests/attention/test_dsa_hisa.py -m 'not benchmark'
# 在配套的 PTO 环境中运行
OPS_TILELANG_TEST_LEVEL=2 TILELANG_DEFAULT_TARGET=pto python -m pytest tests/attention/test_dsa_hisa.py -m 'not benchmark'
TILELANG_DEFAULT_TARGET=auto python -m pytest tests/attention/test_dsa_hisa.py -m benchmark --run-benchmark \
  --benchmark-regression-threshold=0.10 --benchmark-min-delta-us=1.0