feat(attention): add FP8 HISA prefill indexer for Ascend 950-ops-tilelang-AtomGit

Star257

Star257

Open

feat(attention): add FP8 HISA prefill indexer for Ascend 950 #67

jins666Add on 1 day agoMerged fromjins666/ops-tilelang： feat/dsa-hisa-indexer into cann/ops-tilelang： master

feat(attention): add FP8 HISA prefill indexer for Ascend 950 #67

Open

jins666Add on 1 day ago

Merged fromjins666/ops-tilelang feat/dsa-hisa-indexerintocann/ops-tilelang master

Total: 18 files changed+3108-4

Adocs/dsa_hisa.md+87-0

Msrc/cann_ops_tilelang/__init__.py+6-0

Msrc/cann_ops_tilelang/attention/__init__.py+6-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/__init__.py+7-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/block_sparse_mqa_fp8.py+578-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/block_topk.py+75-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/fp8_block_mean_pooling.py+95-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/hisa.py+242-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/pool_mqa_fp8.py+537-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/tilelang_utils.py+106-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/token_topk.py+376-0

Asrc/cann_ops_tilelang/attention/dsa_hisa/token_topk_radix4.py+189-0

Atests/attention/benchmark/test_dsa_hisa.jsonl+6-0

Atests/attention/dsa_hisa_frozen_scorer.py+137-0

Atests/attention/dsa_hisa_ref.py+59-0

Atests/attention/npu_mem_profiles/test_dsa_hisa.json+104-0

Atests/attention/test_dsa_hisa.py+494-0

File content is under review, please try again later