diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md new file mode 100644 index 0000000000..b330173875 --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md @@ -0,0 +1,41 @@ +# 01 — Scalar GM dcache-bypass store/load + +## Algorithm need + +Per-block quantization over **32×32** groups produces one f32 scale factor (SF) +per group. That scalar must land in GM (and later be re-read). Today the tiny +write still goes through bulk MTE, so setup/alignment/sync dominate and keep +VMI short of ASC bandwidth on the SF path. + +## Current slow path + +Stage the SF in UB, then `pto.copy_ubuf_to_gm` with a padded burst (here 32B +for a 4B payload). Checked-in dump: `lowered_vpto.pto` — key op is +`pto.copy_ubuf_to_gm` (no scalar GM path). + +Reproduce: + +```bash +PTO_TEST_OPT=.../pto-test-opt +PASS='-vmi-lower-unified-to-legacy -vmi-mask-granularity-assignment -vmi-layout-assignment -vmi-to-vpto' +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +Direct `pto.vmi.vstore` of a slots=1 vector to `!pto.ptr` with +`dcache_bypass` (see `desired_vmi.pto` / `desired_vmi.py`): + +``` +error: 'pto.vmi.group_store' op requires memory destination to be UB-backed +``` + +## Target MI + bisheng status + +`target_mi.pto` uses `pto.store_scalar ... {dcache_bypass}` on GM. + +| Step | Result | +|---|---| +| `ptoas ... --emit-vpto` | OK | +| `ptoas ... --emit-vpto-llvm-ir` | OK | +| `bisheng ... -c -x ir` → `.o` | OK | diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.pto b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.pto new file mode 100644 index 0000000000..c154a887bf --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.pto @@ -0,0 +1,24 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: SF already staged in UB; bulk MTE to GM for a tiny payload. + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @scale_factor_bulk_mte( + %sf_ub: !pto.ptr, + %sf_gm: !pto.ptr) { + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c32_i64 = arith.constant 32 : i64 + // 4B useful SF / 32B burst — MTE setup dominates. + pto.copy_ubuf_to_gm %sf_ub, %sf_gm, + %c0_i64, %c1_i64, %c32_i64, %c0_i64, %c32_i64, %c32_i64 + : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + return + } +} diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.py b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.py new file mode 100644 index 0000000000..d8b4823c88 --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.py @@ -0,0 +1,25 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``current_slow_vmi.pto`` — bulk UB→GM MTE for a tiny SF. + +Canonical IR: ``current_slow_vmi.pto``. Emit: ``scale_factor_bulk_mte.compile().mlir_text()``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def scale_factor_bulk_mte(sf_gm: pto.ptr(pto.f32, "gm")): + # SF staged in UB; 4B useful / 32B burst — MTE setup dominates. + sf_ub = pto.alloc_tile(shape=[1, 8], dtype=pto.f32) + pto.mte_ub_gm(sf_ub.as_ptr(), sf_gm, 32, nburst=(1, 32, 32)) + + +if __name__ == "__main__": + print(scale_factor_bulk_mte.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.pto b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.pto new file mode 100644 index 0000000000..4de83b0147 --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.pto @@ -0,0 +1,22 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Idiomatic VMI: direct scalar GM store with dcache bypass (does not lower today). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @scale_factor_gm_bypass( + %sf_gm: !pto.ptr, + %scale: !pto.vmi.vreg<1xf32, #pto.vmi.layout>, + %off: index) { + %c1 = arith.constant 1 : index + pto.vmi.vstore %scale, %sf_gm[%off], %c1 {group = 1, dcache_bypass} + : !pto.vmi.vreg<1xf32, #pto.vmi.layout>, + !pto.ptr + return + } +} diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.py b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.py new file mode 100644 index 0000000000..95d71b98bb --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.py @@ -0,0 +1,34 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``desired_vmi.pto`` — scalar GM store with dcache bypass. + +Canonical IR: ``desired_vmi.pto`` (``{group = 1, dcache_bypass}``). +PTODSL ``vstore`` does not yet expose ``dcache_bypass``; IR remains authoritative. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def scale_factor_gm_bypass(sf_gm: pto.ptr(pto.f32, "gm")): + ub = pto.alloc_tile(shape=[1, 8], dtype=pto.f32) + off = pto.const(0, dtype=pto.index) + scale = pto.vmi.vload(ub.as_ptr(), off, size=1) + # Desired IR also sets dcache_bypass on this store. + pto.vmi.vstore( + scale, + sf_gm, + off, + group=1, + stride=pto.const(1, dtype=pto.index), + ) + + +if __name__ == "__main__": + print(scale_factor_gm_bypass.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/lowered_vpto.pto b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/lowered_vpto.pto new file mode 100644 index 0000000000..369643383d --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/lowered_vpto.pto @@ -0,0 +1,10 @@ +module attributes {pto.kernel_kind = #pto.kernel_kind, pto.target_arch = "a5"} { + func.func @scale_factor_bulk_mte(%arg0: !pto.ptr, %arg1: !pto.ptr) { + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c32_i64 = arith.constant 32 : i64 + pto.copy_ubuf_to_gm %arg0, %arg1, %c0_i64, %c1_i64, %c32_i64, %c0_i64, %c32_i64, %c32_i64 : !pto.ptr, !pto.ptr, i64, i64, i64, i64, i64, i64 + return + } +} + diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/target_mi.pto b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/target_mi.pto new file mode 100644 index 0000000000..002194b6ea --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/target_mi.pto @@ -0,0 +1,20 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: scalar GM store with dcache bypass (no bulk MTE). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @scale_factor_gm_bypass_mi( + %sf_gm: !pto.ptr, + %scale: f32) { + %c0 = arith.constant 0 : index + pto.store_scalar %scale, %sf_gm[%c0] {dcache_bypass} + : !pto.ptr, f32 + return + } +} diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md b/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md new file mode 100644 index 0000000000..3c7718c25a --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md @@ -0,0 +1,40 @@ +# 02 — Compact inverse `vbrc` (no pad + BRC) + +## Algorithm need + +Quant kernels compute per-group reciprocal scales, then broadcast each inverse +across the group's lanes before the multiply. A **compact** inverse (one slot +per group, no pad) fan-out is what ASC-level BW wants; padding + BRC reload +adds UB traffic and ops on the hot path. + +## Current slow path + +Materialize recip scales as `#pto.vmi.layout`, +`group_store` them, then reload with `dist_mode = "brc"`. Checked-in dump: +`lowered_vpto.pto` — key ops `pto.vsts` with `PAT_VL8`, then +`pto.vlds {dist = "BRC_B32"}`. + +Reproduce: + +```bash +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +Compact `vbrc` from `num_groups=8, slots=1` (see `desired_vmi.pto` / `desired_vmi.py`): + +``` +error: VMI-LAYOUT-CONTRACT: pto.vmi.group_broadcast has no registered layout support: +operand#0=...num_groups = 8, slots = 1... result#0=...contiguous... +``` + +## Target MI + bisheng status + +`target_mi.pto` is a `BRC_B32` load (plus a consumer store inside `vecscope`). + +| Step | Result | +|---|---| +| `ptoas ... --emit-vpto` | OK | +| `ptoas ... --emit-vpto-llvm-ir` | OK | +| `bisheng ... -c -x ir` → `.o` | OK | diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.pto new file mode 100644 index 0000000000..29516a3799 --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.pto @@ -0,0 +1,25 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: pad recip scales to slots=8, store, then BRC reload. + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @padded_inv_brc( + %inv_ub: !pto.ptr, + %inv_padded: !pto.vmi.vreg<8xf32, #pto.vmi.layout>) + -> !pto.vmi.vreg<64xf32> { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + pto.vmi.vstore %inv_padded, %inv_ub[%c0], %c1 {group = 8} + : !pto.vmi.vreg<8xf32, #pto.vmi.layout>, + !pto.ptr + %inv_brc = pto.vmi.vload %inv_ub[%c0] {dist_mode = "brc"} + : !pto.ptr -> !pto.vmi.vreg<64xf32> + return %inv_brc : !pto.vmi.vreg<64xf32> + } +} diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.py b/docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.py new file mode 100644 index 0000000000..c7b0dd8046 --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.py @@ -0,0 +1,34 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``current_slow_vmi.pto`` — pad recip to slots=8, store, BRC reload. + +Canonical IR: ``current_slow_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def padded_inv_brc(): + inv_ub = pto.alloc_tile(shape=[1, 8], dtype=pto.f32) + off = pto.const(0, dtype=pto.index) + inv_padded = pto.vmi.vload(inv_ub.as_ptr(), off, size=8) + pto.vmi.vstore( + inv_padded, + inv_ub.as_ptr(), + off, + group=8, + stride=pto.const(1, dtype=pto.index), + ) + inv_brc = pto.vmi.vload(inv_ub.as_ptr(), off, size=64, dist_mode="brc") + _ = inv_brc + + +if __name__ == "__main__": + print(padded_inv_brc.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto new file mode 100644 index 0000000000..c2621055ce --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto @@ -0,0 +1,20 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Idiomatic VMI: compact inverse vbrc from slots=1 (no pad round-trip). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @compact_inv_vbrc( + %inv_compact: !pto.vmi.vreg<8xf32, #pto.vmi.layout>) + -> !pto.vmi.vreg<64xf32> { + %inv = pto.vmi.vbrc %inv_compact {group = 8} + : !pto.vmi.vreg<8xf32, #pto.vmi.layout> + -> !pto.vmi.vreg<64xf32> + return %inv : !pto.vmi.vreg<64xf32> + } +} diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.py b/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.py new file mode 100644 index 0000000000..4b2657359e --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.py @@ -0,0 +1,27 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``desired_vmi.pto`` — compact inverse ``vbrc`` (no pad round-trip). + +Canonical IR: ``desired_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def compact_inv_vbrc(): + inv_ub = pto.alloc_tile(shape=[1, 8], dtype=pto.f32) + off = pto.const(0, dtype=pto.index) + inv_compact = pto.vmi.vload(inv_ub.as_ptr(), off, size=8) + inv = pto.vmi.vbrc(inv_compact, size=64, group=8) + _ = inv + + +if __name__ == "__main__": + print(compact_inv_vbrc.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/lowered_vpto.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/lowered_vpto.pto new file mode 100644 index 0000000000..38b9283247 --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/lowered_vpto.pto @@ -0,0 +1,12 @@ +module attributes {pto.kernel_kind = #pto.kernel_kind, pto.target_arch = "a5"} { + func.func @padded_inv_brc(%arg0: !pto.ptr, %arg1: !pto.vreg<64xf32>) -> !pto.vreg<64xf32> { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %0 = pto.pge_b32 "PAT_VL8" : !pto.mask + pto.vsts %arg1, %arg0[%c0], %0 : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + %c0_0 = arith.constant 0 : index + %result = pto.vlds %arg0[%c0] {dist = "BRC_B32"} : !pto.ptr -> !pto.vreg<64xf32> + return %result : !pto.vreg<64xf32> + } +} + diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/target_mi.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/target_mi.pto new file mode 100644 index 0000000000..f669f53dea --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/target_mi.pto @@ -0,0 +1,25 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: BRC_B32 from the compact inverse line (no pad store + reload). + +module { + func.func @compact_inv_vbrc_mi( + %inv_ub: !pto.ptr, + %off: index, + %dst: !pto.ptr) { + pto.vecscope { + %inv = pto.vlds %inv_ub[%off] {dist = "BRC_B32"} + : !pto.ptr -> !pto.vreg<64xf32> + %m = pto.pset_b32 "PAT_ALL" : !pto.mask + pto.vsts %inv, %dst[%off], %m + : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + } + return + } +} diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md new file mode 100644 index 0000000000..3d646037f6 --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md @@ -0,0 +1,59 @@ +# 03 — Direct scalar / 1PT reduce store + +## Algorithm need + +Residual-mix post-bwd reduces a tile to **one f32 scalar per group**. ASC uses +a ONEPT/1PT-style store. If VMI still pads that result to 8 slots before store, +extra UB traffic shows up on the residual-mix path and keeps VMI below ASC BW. + +## Current slow path + +Group-reduce with `{group = 8}` and store the 8-slot result (`current_slow_vmi.pto` / `current_slow_vmi.py`). +High-level kernels also keep a pad-to-8 habit because dense adjacent size-1 +stores have been unreliable in practice. Checked-in dump: `lowered_vpto.pto` — +key ops `pto.vcgadd` + `pto.vsts` with `PAT_VL8`. + +Reproduce: + +```bash +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +Idiomatic mask1 store after `vcadd` (mirrors +`test/lit/vmi_new/vmi_to_vpto_reduce_addf_store.pto` in `desired_vmi.pto` / `desired_vmi.py`). +**This lowers today** to `pto.vcadd` + `pto.vsts` under a `PAT_VL1` mask +(no explicit `dist = "1PT_B32"` in the dump). The remaining gap is practical: +residual-mix authoring still pads to 8, and ASC's ONEPT path is not what the +padded slow path emits. + +## Target MI + bisheng status + +`target_mi.pto` asks for `pto.vcadd` then `pto.vsts {dist = "1PT_B32"}` — +the MI form of ASC `ONEPT_B32` (matches `reference_asc_cce.asc`). + +| Step | Result | +|---|---| +| `ptoas ... --emit-vpto` | OK | +| `ptoas ... --emit-vpto-llvm-ir` | OK | +| `bisheng ... -c -x ir` → `.o` | **FAIL** — bisheng crashes (exit 139) on the emitted HIVM IR (`vcadd` / `vstsx1`) | + +## ASC/CCE working baseline + +`reference_asc_cce.asc` is a minimal hand-written Ascend CCE kernel: +`vcadd` then `vsts(..., ONEPT_B32)`. It compiles to a non-empty `.o` with +`bisheng --cce-aicore-only` (same CCE path ASC uses), proving the instruction +pattern is legal while `target_mi` HIVM still crashes. + +```bash +BISHENG=${BISHENG:-/usr/local/Ascend/cann-9.0.0/tools/bisheng_compiler/bin/bisheng} +ASCEND=${ASCEND:-/usr/local/Ascend/cann-9.0.0} +"$BISHENG" -O2 -fPIC -std=c++17 --npu-arch=dav-3510 --cce-aicore-only -c \ + reference_asc_cce.asc -o /tmp/ref_gap03.o \ + -I"$ASCEND/include" \ + -I"$ASCEND/compiler/tikcpp/tikcfw" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/impl" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/interface" +test -s /tmp/ref_gap03.o +``` diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.pto new file mode 100644 index 0000000000..782e54a785 --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.pto @@ -0,0 +1,25 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: group-reduce into an 8-slot pad, then group-store. + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @padded_reduce_store( + %tile: !pto.vmi.vreg<64xf32>, + %mask: !pto.vmi.mask<64xpred>, + %dst: !pto.ptr, + %off: index) { + %c1 = arith.constant 1 : index + %sum8 = pto.vmi.vcadd %tile, %mask {group = 8, reassoc} + : !pto.vmi.vreg<64xf32>, !pto.vmi.mask<64xpred> + -> !pto.vmi.vreg<8xf32> + pto.vmi.vstore %sum8, %dst[%off], %c1 {group = 8} + : !pto.vmi.vreg<8xf32>, !pto.ptr + return + } +} diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.py b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.py new file mode 100644 index 0000000000..9eea0b1c67 --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.py @@ -0,0 +1,35 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``current_slow_vmi.pto`` — group-reduce into 8-slot pad + store. + +Canonical IR: ``current_slow_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def padded_reduce_store(): + tile_ub = pto.alloc_tile(shape=[1, 64], dtype=pto.f32) + dst = pto.alloc_tile(shape=[1, 8], dtype=pto.f32) + off = pto.const(0, dtype=pto.index) + mask = pto.vmi.create_mask(pto.const(64, dtype=pto.index), size=64) + tile = pto.vmi.vload(tile_ub.as_ptr(), off, size=64) + sum8 = pto.vmi.vcadd(tile, mask, group=8, reassoc=True) + pto.vmi.vstore( + sum8, + dst.as_ptr(), + off, + group=8, + stride=pto.const(1, dtype=pto.index), + ) + + +if __name__ == "__main__": + print(padded_reduce_store.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.pto new file mode 100644 index 0000000000..c341694d91 --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.pto @@ -0,0 +1,28 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Idiomatic VMI: reduce to one scalar and store with a 1-lane mask +// (mirrors test/lit/vmi_new/vmi_to_vpto_reduce_addf_store.pto). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @direct_scalar_reduce_store( + %tile: !pto.vmi.vreg<64xf32>, + %dst: !pto.ptr, + %off: index) { + %c64 = arith.constant 64 : index + %c1 = arith.constant 1 : index + %mask64 = pto.vmi.create_mask %c64 : index -> !pto.vmi.mask<64xpred> + %mask1 = pto.vmi.create_mask %c1 : index -> !pto.vmi.mask<1xpred> + %sum = pto.vmi.vcadd %tile, %mask64 {reassoc} + : !pto.vmi.vreg<64xf32>, !pto.vmi.mask<64xpred> + -> !pto.vmi.vreg<1xf32> + pto.vmi.vstore %sum, %dst[%off], %mask1 + : !pto.vmi.vreg<1xf32>, !pto.ptr, !pto.vmi.mask<1xpred> + return + } +} diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.py b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.py new file mode 100644 index 0000000000..df48713a05 --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.py @@ -0,0 +1,31 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``desired_vmi.pto`` — scalar reduce + 1-lane store. + +Canonical IR: ``desired_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def direct_scalar_reduce_store(): + tile_ub = pto.alloc_tile(shape=[1, 64], dtype=pto.f32) + # 32B-aligned tile; store uses a 1-lane mask (ONEPT / 1PT shape). + dst = pto.alloc_tile(shape=[1, 8], dtype=pto.f32) + off = pto.const(0, dtype=pto.index) + mask64 = pto.vmi.create_mask(pto.const(64, dtype=pto.index), size=64) + mask1 = pto.vmi.create_mask(pto.const(1, dtype=pto.index), size=1) + tile = pto.vmi.vload(tile_ub.as_ptr(), off, size=64) + total = pto.vmi.vcadd(tile, mask64, reassoc=True) + pto.vmi.vstore(total, dst.as_ptr(), off, mask1) + + +if __name__ == "__main__": + print(direct_scalar_reduce_store.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/lowered_vpto.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/lowered_vpto.pto new file mode 100644 index 0000000000..8cf259476b --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/lowered_vpto.pto @@ -0,0 +1,10 @@ +module attributes {pto.kernel_kind = #pto.kernel_kind, pto.target_arch = "a5"} { + func.func @padded_reduce_store(%arg0: !pto.vreg<64xf32>, %arg1: !pto.mask, %arg2: !pto.ptr, %arg3: index) { + %c1 = arith.constant 1 : index + %0 = pto.vcgadd %arg0, %arg1 : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %1 = pto.pge_b32 "PAT_VL8" : !pto.mask + pto.vsts %0, %arg2[%arg3], %1 : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + return + } +} + diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/reference_asc_cce.asc b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/reference_asc_cce.asc new file mode 100644 index 0000000000..8733d79f9f --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/reference_asc_cce.asc @@ -0,0 +1,20 @@ +// ASC/CCE working baseline for gap 03: vcadd + ONEPT_B32 store. +// Compiles with bisheng --cce-aicore-only (see README). +#include "kernel_operator.h" + +__simd_vf__ inline void vf_onept_reduce(__ubuf__ float* src, __ubuf__ float* dst) { + vector_bool v = pset_b32(PAT_ALL); + vector_f32 x; + vlds(x, src, 0, NORM); + vector_f32 y; + vcadd(y, x, v, MODE_ZEROING); + // ONEPT_B32: write the reduced scalar (1×f32), not a padded VL8 tile. + vsts(y, dst, 0, ONEPT_B32, v); +} + +extern "C" __global__ __vector__ void ref_gap03_onept_reduce(__gm__ float* /*gm_in*/, + __gm__ float* /*gm_out*/) { + AscendC::InitSocState(); + __ubuf__ float* ub = (__ubuf__ float*)0; + vf_onept_reduce(ub, ub + 64); +} diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/target_mi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/target_mi.pto new file mode 100644 index 0000000000..6281dc02e1 --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/target_mi.pto @@ -0,0 +1,26 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: vcadd + 1PT_B32 store (ASC ONEPT-style), no pad-to-8. + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @direct_scalar_reduce_store_mi( + %tile: !pto.vreg<64xf32>, + %mask: !pto.mask, + %dst: !pto.ptr, + %off: index) { + pto.vecscope { + %sum = pto.vcadd %tile, %mask {reassoc} + : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %store_m = pto.pge_b32 "PAT_VL1" : !pto.mask + pto.vsts %sum, %dst[%off], %store_m {dist = "1PT_B32"} + : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + } + return + } +} diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md new file mode 100644 index 0000000000..1522961703 --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md @@ -0,0 +1,57 @@ +# 04 — Packed UE8M0 scale factors + +## Algorithm need + +Per-token MXFP8 stores UE8M0 scales with **pack_factor=2**: two `ui8` exponents +share one `ui16` element. Unpacked `ui8` SF doubles scale-factor traffic versus +packed storage and blocks ASC-level BW on that pipe. + +## Current slow path + +Store SF as unpacked `ui8` only (`current_slow_vmi.pto` / `current_slow_vmi.py`). Checked-in dump: +`lowered_vpto.pto` — key op `pto.vsts ... {dist = "PK4_B32"}` into `!pto.ptr` +(byte store of the unpacked tensor, not a packed `ui16` SF layout). + +Reproduce: + +```bash +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +First-class pack into `ui16` then store (see `desired_vmi.pto` / `desired_vmi.py`): + +``` +error: custom op 'pto.vmi.vpack' is unknown +``` + +## Target MI + bisheng status + +`target_mi.pto` uses `pto.vsts {dist = "PK_B16"}` into `!pto.ptr` +(packed SF word store; ASC baseline also exercises `ONEPT_B16` / `PK_B16`). + +| Step | Result | +|---|---| +| `ptoas ... --emit-vpto` | OK | +| `ptoas ... --emit-vpto-llvm-ir` | OK | +| `bisheng ... -c -x ir` → `.o` | **FAIL** — bisheng crashes (exit 139) on the emitted HIVM IR | + +## ASC/CCE working baseline + +`reference_asc_cce.asc` stores packed `ui16` SF with `ONEPT_B16` and `PK_B16`. +It compiles to a non-empty `.o` with `bisheng --cce-aicore-only`, proving the +packed-SF store shape is legal on the Ascend CCE path while `target_mi` HIVM +still crashes. + +```bash +BISHENG=${BISHENG:-/usr/local/Ascend/cann-9.0.0/tools/bisheng_compiler/bin/bisheng} +ASCEND=${ASCEND:-/usr/local/Ascend/cann-9.0.0} +"$BISHENG" -O2 -fPIC -std=c++17 --npu-arch=dav-3510 --cce-aicore-only -c \ + reference_asc_cce.asc -o /tmp/ref_gap04.o \ + -I"$ASCEND/include" \ + -I"$ASCEND/compiler/tikcpp/tikcfw" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/impl" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/interface" +test -s /tmp/ref_gap04.o +``` diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.pto new file mode 100644 index 0000000000..207b79eb34 --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.pto @@ -0,0 +1,20 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: store UE8M0 scales as unpacked ui8 only. + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @unpacked_ue8m0_sf_store( + %sf_u8: !pto.vmi.vreg<64xui8>, + %dst: !pto.ptr, + %off: index) { + pto.vmi.vstore %sf_u8, %dst[%off] + : !pto.vmi.vreg<64xui8>, !pto.ptr + return + } +} diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.py b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.py new file mode 100644 index 0000000000..962d8c4f43 --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.py @@ -0,0 +1,26 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``current_slow_vmi.pto`` — unpacked ui8 SF store. + +Canonical IR: ``current_slow_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def unpacked_ue8m0_sf_store(): + dst = pto.alloc_tile(shape=[1, 64], dtype=pto.ui8) + off = pto.const(0, dtype=pto.index) + sf_u8 = pto.vmi.vload(dst.as_ptr(), off, size=64) + pto.vmi.vstore(sf_u8, dst.as_ptr(), off) + + +if __name__ == "__main__": + print(unpacked_ue8m0_sf_store.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.pto new file mode 100644 index 0000000000..77aed4e86f --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.pto @@ -0,0 +1,22 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Idiomatic VMI: pack adjacent UE8M0 bytes into ui16 (pack_factor=2). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @packed_ue8m0_sf_store( + %sf_u8: !pto.vmi.vreg<64xui8>, + %dst: !pto.ptr, + %off: index) { + %sf_u16 = pto.vmi.vpack %sf_u8 + : !pto.vmi.vreg<64xui8> -> !pto.vmi.vreg<32xui16> + pto.vmi.vstore %sf_u16, %dst[%off] + : !pto.vmi.vreg<32xui16>, !pto.ptr + return + } +} diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.py b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.py new file mode 100644 index 0000000000..4f7f22fc7d --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.py @@ -0,0 +1,32 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``desired_vmi.pto`` — pack UE8M0 into ui16 then store. + +Canonical IR: ``desired_vmi.pto`` (``pto.vmi.vpack`` — unknown / not bound today). +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def packed_ue8m0_sf_store(): + src = pto.alloc_tile(shape=[1, 64], dtype=pto.ui8) + dst = pto.alloc_tile(shape=[1, 32], dtype=pto.ui16) + off = pto.const(0, dtype=pto.index) + sf_u8 = pto.vmi.vload(src.as_ptr(), off, size=64) + # Desired surface: pack_factor=2 ui8 → ui16 (not on PTODSL yet). + sf_u16 = pto.vmi.vpack(sf_u8) + pto.vmi.vstore(sf_u16, dst.as_ptr(), off) + + +if __name__ == "__main__": + try: + print(packed_ue8m0_sf_store.compile().mlir_text()) + except Exception as exc: # noqa: BLE001 — document the gap + print(f"desired emit failed (expected until pto.vmi.vpack lands): {exc}") diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/lowered_vpto.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/lowered_vpto.pto new file mode 100644 index 0000000000..3593145105 --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/lowered_vpto.pto @@ -0,0 +1,13 @@ +module attributes {pto.kernel_kind = #pto.kernel_kind, pto.target_arch = "a5"} { + func.func @unpacked_ue8m0_sf_store(%arg0: !pto.vreg<256xui8>, %arg1: !pto.ptr, %arg2: index) { + %c0 = arith.constant 0 : index + %0 = pto.vzunpack %arg0, %c0 : !pto.vreg<256xui8> -> !pto.vreg<128xui16> + %c0_0 = arith.constant 0 : index + %1 = pto.vzunpack %0, %c0_0 : !pto.vreg<128xui16> -> !pto.vreg<64xui32> + %2 = pto.vbitcast %1 : !pto.vreg<64xui32> -> !pto.vreg<256xui8> + %3 = pto.pge_b32 "PAT_VL64" : !pto.mask + pto.vsts %2, %arg1[%arg2], %3 {dist = "PK4_B32"} : !pto.vreg<256xui8>, !pto.ptr, !pto.mask + return + } +} + diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/reference_asc_cce.asc b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/reference_asc_cce.asc new file mode 100644 index 0000000000..700292f7b2 --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/reference_asc_cce.asc @@ -0,0 +1,26 @@ +// ASC/CCE working baseline for gap 04: packed ui16 SF store (ONEPT_B16 / PK_B16). +// Compiles with bisheng --cce-aicore-only (see README). +#include "kernel_operator.h" + +// Packed UE8M0 path stores one ui16 per pair of ui8 exponents. +__simd_vf__ inline void vf_onept_u16(__ubuf__ uint16_t* src, __ubuf__ uint16_t* dst) { + vector_bool v = pset_b16(PAT_ALL); + vector_u16 x; + vlds(x, src, 0, NORM); + vsts(x, dst, 0, ONEPT_B16, v); +} + +__simd_vf__ inline void vf_pk_b16(__ubuf__ uint16_t* src, __ubuf__ uint16_t* dst) { + vector_bool v = pset_b16(PAT_ALL); + vector_u16 x; + vlds(x, src, 0, NORM); + vsts(x, dst, 0, PK_B16, v); +} + +extern "C" __global__ __vector__ void ref_gap04_packed_sf(__gm__ uint16_t* /*gm_in*/, + __gm__ uint16_t* /*gm_out*/) { + AscendC::InitSocState(); + __ubuf__ uint16_t* ub = (__ubuf__ uint16_t*)0; + vf_onept_u16(ub, ub + 64); + vf_pk_b16(ub, ub + 128); +} diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.pto new file mode 100644 index 0000000000..34198771e7 --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.pto @@ -0,0 +1,23 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: PK_B16 store of packed UE8M0 words. + +module { + func.func @packed_ue8m0_sf_store_mi( + %sf_u16: !pto.vreg<128xui16>, + %dst: !pto.ptr, + %off: index) { + pto.vecscope { + %m = pto.pset_b16 "PAT_ALL" : !pto.mask + pto.vsts %sf_u16, %dst[%off], %m {dist = "PK_B16"} + : !pto.vreg<128xui16>, !pto.ptr, !pto.mask + } + return + } +} diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md new file mode 100644 index 0000000000..82703bc767 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md @@ -0,0 +1,73 @@ +# 05 — Persistent stages > 1 / `block_k` > 512 + +## Algorithm need + +Long-K quant / **cast-back (dequant)** Persistent loops want multi-buffer +stages (e.g. stages=2) and `block_k > 512` so MTE and vector stay overlapped. +Caps at stages=1 / `block_k ≤ 512` leave the pipe underfed versus ASC on large +hidden sizes. Raising stages in the host schedule has also faulted in practice. + +Consumers today: + +- **Per-block quant** — stuck on stages=1 and `block_k ≤ 512` on large shapes. +- **Cast-back / dequant** — sweet spot is stages=1 + tile_k=512; tile_k=1024 + regresses until larger K / deeper Persistent is legal and fast. Closing this + gap is the main lever to push cast-back toward ASC BW. + +## Current slow path + +`block_k = 512`, single-stage Persistent tile (`current_slow_vmi.pto` / `current_slow_vmi.py`). Checked-in +dump: `lowered_vpto.pto` — chunked `pto.vlds` / `pto.vcvt` / `pto.vmul` / +`pto.vsts` over the 512-wide tile. + +Reproduce: + +```bash +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +`stages = 2` and `block_k = 1024` with dual UB buffers (`desired_vmi.pto` / `desired_vmi.py`). +Layout assignment rejects the 1024-wide multiply: + +``` +error: VMI-UNSUPPORTED: pto.vmi.mulf ... deinterleaved = 2 ... requires ... deinterleaved = 4; +pto.vmi.ensure_layout cannot materialize this conversion +``` + +## Target MI + bisheng status + +`target_mi.pto` expresses the full stages=2 / `block_k=1024` Persistent +schedule (aligned with `desired_vmi.pto` + ASC HardEvent shell): + +- dual UB args (`x_ub0/1`, `y_ub0/1`) selected by `slot = i % 2` +- prime / drain `pto.set_flag` / `pto.wait_flag` on `EVENT_ID0/1` +- in-loop `pto.wait_flag_dyn` / `pto.set_flag_dyn` on `%slot` +- VF body: `UNPK_B16` load → `vmul` → `PK4_B32` store + +| Step | Result | +|---|---| +| `ptoas ... --emit-vpto` | OK (schedule + flags preserved) | +| `ptoas ... --emit-vpto-llvm-ir` | OK | +| `bisheng ... -c -x ir` → `.o` | **FAIL** — bisheng crashes on the emitted HIVM IR | + +## ASC/CCE working baseline + +`reference_asc_cce.asc` is the same schedule on the Ascend CCE path: +stages=2 ping-pong (`SetFlag`/`WaitFlag` on slots `w & 1`), `block_k=1024`, +VF body `UNPK_B16` / `vmul` / `PK_B32`. It compiles to a non-empty `.o` with +`bisheng --cce-aicore-only`, proving the pattern is legal while `target_mi` +HIVM still crashes. + +```bash +BISHENG=${BISHENG:-/usr/local/Ascend/cann-9.0.0/tools/bisheng_compiler/bin/bisheng} +ASCEND=${ASCEND:-/usr/local/Ascend/cann-9.0.0} +"$BISHENG" -O2 -fPIC -std=c++17 --npu-arch=dav-3510 --cce-aicore-only -c \ + reference_asc_cce.asc -o /tmp/ref_gap05.o \ + -I"$ASCEND/include" \ + -I"$ASCEND/compiler/tikcpp/tikcfw" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/impl" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/interface" +test -s /tmp/ref_gap05.o +``` diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.pto new file mode 100644 index 0000000000..4494a4cce1 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.pto @@ -0,0 +1,47 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: Persistent vector tile: stages=1, block_k=512. + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind + // stages = 1 + // block_k = 512 +} { + func.func @persistent_quant_tile_k512( + %x_ub: !pto.ptr, + %y_ub: !pto.ptr, + %scale_slot: !pto.vmi.vreg<1xf32>, + %block_count: index) + attributes {pto.kernel} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c512 = arith.constant 512 : index + + pto.vecscope { + %scale = pto.vmi.vbrc %scale_slot + : !pto.vmi.vreg<1xf32> -> !pto.vmi.vreg<512xf32> + scf.for %i = %c0 to %block_count step %c1 { + %off = arith.muli %i, %c512 : index + %x_bf16 = pto.vmi.vload %x_ub[%off] + : !pto.ptr -> !pto.vmi.vreg<512xbf16> + %x_f32 = pto.vmi.vcvt %x_bf16 + : !pto.vmi.vreg<512xbf16> -> !pto.vmi.vreg<512xf32> + %scaled = pto.vmi.vmul %x_f32, %scale + : !pto.vmi.vreg<512xf32>, !pto.vmi.vreg<512xf32> + -> !pto.vmi.vreg<512xf32> + %y = pto.vmi.vcvt %scaled {saturate = "SAT"} + : !pto.vmi.vreg<512xf32> -> !pto.vmi.vreg<512xf8E4M3FN> + pto.vmi.vstore %y, %y_ub[%off] + : !pto.vmi.vreg<512xf8E4M3FN>, !pto.ptr + } + } + return + } +} diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.py b/docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.py new file mode 100644 index 0000000000..2c0f4e96b2 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.py @@ -0,0 +1,36 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``current_slow_vmi.pto`` — stages=1, block_k=512 Persistent tile. + +Canonical IR: ``current_slow_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def persistent_quant_tile_k512(*, BLOCK_COUNT: pto.const_expr = 4): + # stages = 1, block_k = 512 + x_ub = pto.alloc_tile(shape=[1, 512], dtype=pto.bf16) + y_ub = pto.alloc_tile(shape=[1, 512], dtype=pto.f8e4m3) + scale_ub = pto.alloc_tile(shape=[1, 8], dtype=pto.f32) + off0 = pto.const(0, dtype=pto.index) + scale_slot = pto.vmi.vload(scale_ub.as_ptr(), off0, size=1) + scale = pto.vmi.vbrc(scale_slot, size=512) + for i in range(BLOCK_COUNT): + off = i * 512 + x_bf16 = pto.vmi.vload(x_ub.as_ptr(), off, size=512) + x_f32 = pto.vmi.vcvt(x_bf16, pto.f32) + scaled = pto.vmi.vmul(x_f32, scale) + y = pto.vmi.vcvt(scaled, pto.f8e4m3, saturate="SAT") + pto.vmi.vstore(y, y_ub.as_ptr(), off) + + +if __name__ == "__main__": + print(persistent_quant_tile_k512.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.pto new file mode 100644 index 0000000000..da24d2cdb3 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.pto @@ -0,0 +1,62 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Idiomatic Persistent VMI: stages=2 and block_k=1024 (does not lower today). + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind, + pto.pipeline.stages = 2 : i32, + pto.block_k = 1024 : i32 +} { + func.func @persistent_quant_tile_k1024( + %x_ub0: !pto.ptr, + %x_ub1: !pto.ptr, + %y_ub0: !pto.ptr, + %y_ub1: !pto.ptr, + %scale_slot: !pto.vmi.vreg<1xf32>, + %block_count: index) + attributes {pto.kernel} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c1024 = arith.constant 1024 : index + + pto.vecscope { + %scale = pto.vmi.vbrc %scale_slot + : !pto.vmi.vreg<1xf32> -> !pto.vmi.vreg<1024xf32> + scf.for %i = %c0 to %block_count step %c1 { + %off = arith.muli %i, %c1024 : index + %stage = arith.remui %i, %c2 : index + %use1 = arith.cmpi eq, %stage, %c1 : index + %x_ub = scf.if %use1 -> (!pto.ptr) { + scf.yield %x_ub1 : !pto.ptr + } else { + scf.yield %x_ub0 : !pto.ptr + } + %y_ub = scf.if %use1 -> (!pto.ptr) { + scf.yield %y_ub1 : !pto.ptr + } else { + scf.yield %y_ub0 : !pto.ptr + } + %x_bf16 = pto.vmi.vload %x_ub[%off] + : !pto.ptr -> !pto.vmi.vreg<1024xbf16> + %x_f32 = pto.vmi.vcvt %x_bf16 + : !pto.vmi.vreg<1024xbf16> -> !pto.vmi.vreg<1024xf32> + %scaled = pto.vmi.vmul %x_f32, %scale + : !pto.vmi.vreg<1024xf32>, !pto.vmi.vreg<1024xf32> + -> !pto.vmi.vreg<1024xf32> + %y = pto.vmi.vcvt %scaled {saturate = "SAT"} + : !pto.vmi.vreg<1024xf32> -> !pto.vmi.vreg<1024xf8E4M3FN> + pto.vmi.vstore %y, %y_ub[%off] + : !pto.vmi.vreg<1024xf8E4M3FN>, !pto.ptr + } + } + return + } +} diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.py b/docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.py new file mode 100644 index 0000000000..cdee3d6176 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.py @@ -0,0 +1,48 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``desired_vmi.pto`` — stages=2, block_k=1024 dual-buffer tile. + +Canonical IR: ``desired_vmi.pto`` (module attrs ``pto.pipeline.stages`` / +``pto.block_k``). Layout assignment rejects the 1024-wide multiply today. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def persistent_quant_tile_k1024(*, BLOCK_COUNT: pto.const_expr = 4): + # stages = 2, block_k = 1024 — dual UB ping-pong + x_ub0 = pto.alloc_tile(shape=[1, 1024], dtype=pto.bf16) + x_ub1 = pto.alloc_tile(shape=[1, 1024], dtype=pto.bf16) + y_ub0 = pto.alloc_tile(shape=[1, 1024], dtype=pto.f8e4m3) + y_ub1 = pto.alloc_tile(shape=[1, 1024], dtype=pto.f8e4m3) + scale_ub = pto.alloc_tile(shape=[1, 8], dtype=pto.f32) + off0 = pto.const(0, dtype=pto.index) + scale_slot = pto.vmi.vload(scale_ub.as_ptr(), off0, size=1) + scale = pto.vmi.vbrc(scale_slot, size=1024) + for i in range(BLOCK_COUNT): + off = i * 1024 + if (i % 2) == 1: + x_ptr = x_ub1.as_ptr() + y_ptr = y_ub1.as_ptr() + else: + x_ptr = x_ub0.as_ptr() + y_ptr = y_ub0.as_ptr() + x_bf16 = pto.vmi.vload(x_ptr, off, size=1024) + x_f32 = pto.vmi.vcvt(x_bf16, pto.f32) + scaled = pto.vmi.vmul(x_f32, scale) + y = pto.vmi.vcvt(scaled, pto.f8e4m3, saturate="SAT") + pto.vmi.vstore(y, y_ptr, off) + + +if __name__ == "__main__": + try: + print(persistent_quant_tile_k1024.compile().mlir_text()) + except Exception as exc: # noqa: BLE001 — document the gap + print(f"desired emit/lower may fail (layout / stages): {exc}") diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/lowered_vpto.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/lowered_vpto.pto new file mode 100644 index 0000000000..7bcff4e15e --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/lowered_vpto.pto @@ -0,0 +1,100 @@ +module attributes {pto.kernel_kind = #pto.kernel_kind, pto.target_arch = "a5"} { + func.func @persistent_quant_tile_k512(%arg0: !pto.ptr, %arg1: !pto.ptr, %arg2: !pto.vreg<64xf32>, %arg3: index) attributes {pto.kernel} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c512 = arith.constant 512 : index + pto.vecscope { + %0 = pto.pset_b32 "PAT_ALL" : !pto.mask + %1 = pto.vdup %arg2, %0 {position = "LOWEST"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %2 = pto.pset_b32 "PAT_ALL" : !pto.mask + %3 = pto.vdup %arg2, %2 {position = "LOWEST"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %4 = pto.pset_b32 "PAT_ALL" : !pto.mask + %5 = pto.vdup %arg2, %4 {position = "LOWEST"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %6 = pto.pset_b32 "PAT_ALL" : !pto.mask + %7 = pto.vdup %arg2, %6 {position = "LOWEST"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %8 = pto.pset_b32 "PAT_ALL" : !pto.mask + %9 = pto.vdup %arg2, %8 {position = "LOWEST"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %10 = pto.pset_b32 "PAT_ALL" : !pto.mask + %11 = pto.vdup %arg2, %10 {position = "LOWEST"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %12 = pto.pset_b32 "PAT_ALL" : !pto.mask + %13 = pto.vdup %arg2, %12 {position = "LOWEST"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %14 = pto.pset_b32 "PAT_ALL" : !pto.mask + %15 = pto.vdup %arg2, %14 {position = "LOWEST"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + scf.for %arg4 = %c0 to %arg3 step %c1 { + %16 = arith.muli %arg4, %c512 : index + %result = pto.vlds %arg0[%16] : !pto.ptr -> !pto.vreg<128xbf16> + %c128 = arith.constant 128 : index + %17 = arith.addi %16, %c128 : index + %result_0 = pto.vlds %arg0[%17] : !pto.ptr -> !pto.vreg<128xbf16> + %c256 = arith.constant 256 : index + %18 = arith.addi %16, %c256 : index + %result_1 = pto.vlds %arg0[%18] : !pto.ptr -> !pto.vreg<128xbf16> + %c384 = arith.constant 384 : index + %19 = arith.addi %16, %c384 : index + %result_2 = pto.vlds %arg0[%19] : !pto.ptr -> !pto.vreg<128xbf16> + %20 = pto.pset_b16 "PAT_ALL" : !pto.mask + %21 = pto.vcvt %result, %20 {part = "EVEN"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %22 = pto.vcvt %result_0, %20 {part = "EVEN"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %23 = pto.vcvt %result_1, %20 {part = "EVEN"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %24 = pto.vcvt %result_2, %20 {part = "EVEN"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %25 = pto.vcvt %result, %20 {part = "ODD"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %26 = pto.vcvt %result_0, %20 {part = "ODD"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %27 = pto.vcvt %result_1, %20 {part = "ODD"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %28 = pto.vcvt %result_2, %20 {part = "ODD"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %low, %high = pto.vintlv %21, %25 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_3, %high_4 = pto.vintlv %22, %26 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_5, %high_6 = pto.vintlv %23, %27 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_7, %high_8 = pto.vintlv %24, %28 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_9, %high_10 = pto.vdintlv %low, %high : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_11, %high_12 = pto.vdintlv %low_3, %high_4 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_13, %high_14 = pto.vdintlv %low_9, %low_11 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_15, %high_16 = pto.vdintlv %high_10, %high_12 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_17, %high_18 = pto.vdintlv %low_5, %high_6 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_19, %high_20 = pto.vdintlv %low_7, %high_8 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_21, %high_22 = pto.vdintlv %low_17, %low_19 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %low_23, %high_24 = pto.vdintlv %high_18, %high_20 : !pto.vreg<64xf32>, !pto.vreg<64xf32> -> !pto.vreg<64xf32>, !pto.vreg<64xf32> + %29 = pto.pset_b32 "PAT_ALL" : !pto.mask + %30 = pto.vmul %low_13, %1, %29 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %31 = pto.pset_b32 "PAT_ALL" : !pto.mask + %32 = pto.vmul %low_21, %3, %31 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %33 = pto.pset_b32 "PAT_ALL" : !pto.mask + %34 = pto.vmul %low_15, %5, %33 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %35 = pto.pset_b32 "PAT_ALL" : !pto.mask + %36 = pto.vmul %low_23, %7, %35 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %37 = pto.pset_b32 "PAT_ALL" : !pto.mask + %38 = pto.vmul %high_14, %9, %37 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %39 = pto.pset_b32 "PAT_ALL" : !pto.mask + %40 = pto.vmul %high_22, %11, %39 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %41 = pto.pset_b32 "PAT_ALL" : !pto.mask + %42 = pto.vmul %high_16, %13, %41 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %43 = pto.pset_b32 "PAT_ALL" : !pto.mask + %44 = pto.vmul %high_24, %15, %43 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %45 = pto.pset_b32 "PAT_ALL" : !pto.mask + %46 = pto.pset_b8 "PAT_ALL" : !pto.mask + %47 = pto.vcvt %30, %45 {part = "P0", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %48 = pto.vcvt %34, %45 {part = "P1", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %49 = pto.vcvt %38, %45 {part = "P2", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %50 = pto.vcvt %42, %45 {part = "P3", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %51 = pto.vor %47, %48, %46 : !pto.vreg<256xf8E4M3FN>, !pto.vreg<256xf8E4M3FN>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %52 = pto.vor %51, %49, %46 : !pto.vreg<256xf8E4M3FN>, !pto.vreg<256xf8E4M3FN>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %53 = pto.vor %52, %50, %46 : !pto.vreg<256xf8E4M3FN>, !pto.vreg<256xf8E4M3FN>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %54 = pto.pset_b8 "PAT_ALL" : !pto.mask + %55 = pto.vcvt %32, %45 {part = "P0", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %56 = pto.vcvt %36, %45 {part = "P1", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %57 = pto.vcvt %40, %45 {part = "P2", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %58 = pto.vcvt %44, %45 {part = "P3", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %59 = pto.vor %55, %56, %54 : !pto.vreg<256xf8E4M3FN>, !pto.vreg<256xf8E4M3FN>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %60 = pto.vor %59, %57, %54 : !pto.vreg<256xf8E4M3FN>, !pto.vreg<256xf8E4M3FN>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %61 = pto.vor %60, %58, %54 : !pto.vreg<256xf8E4M3FN>, !pto.vreg<256xf8E4M3FN>, !pto.mask -> !pto.vreg<256xf8E4M3FN> + %62 = pto.pset_b8 "PAT_ALL" : !pto.mask + pto.vsts %53, %arg1[%16], %62 : !pto.vreg<256xf8E4M3FN>, !pto.ptr, !pto.mask + %63 = pto.pset_b8 "PAT_ALL" : !pto.mask + %c256_25 = arith.constant 256 : index + %64 = arith.addi %16, %c256_25 : index + pto.vsts %61, %arg1[%64], %63 : !pto.vreg<256xf8E4M3FN>, !pto.ptr, !pto.mask + } + } + return + } +} + diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/reference_asc_cce.asc b/docs/feature-gaps/vmi/05_persistent_stages_block_k/reference_asc_cce.asc new file mode 100644 index 0000000000..a7989be4b7 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/reference_asc_cce.asc @@ -0,0 +1,58 @@ +// ASC/CCE working baseline for gap 05: stages=2 ping-pong + UNPK/vmul/PK body. +// Compiles with bisheng --cce-aicore-only (see README). +#include "kernel_operator.h" + +__simd_vf__ inline void vf_scale_mul(__ubuf__ bfloat16_t* x_ub, __ubuf__ float* scale_ub, + __ubuf__ bfloat16_t* out_ub) { + vector_bool v16 = pset_b16(PAT_ALL); + vector_bool v32 = pset_b32(PAT_ALL); + vector_bf16 xb; + vlds(xb, x_ub, 0, UNPK_B16); + vector_f32 xf; + vcvt(xf, xb, v16, PART_EVEN, MODE_ZEROING); + vector_f32 sc; + vlds(sc, scale_ub, 0, BRC_B32); + vector_f32 yf; + vmul(yf, xf, sc, v32, MODE_ZEROING); + vector_bf16 yb; + vcvt(yb, yf, v16, ROUND_R, RS_ENABLE, PART_EVEN, MODE_ZEROING); + vsts(yb, out_ub, 0, PK_B32, v16); +} + +extern "C" __global__ __vector__ void ref_gap05_stages2(__gm__ bfloat16_t* /*x*/, + __gm__ float* /*sf*/, + __gm__ bfloat16_t* /*out*/) { + AscendC::InitSocState(); + __ubuf__ uint8_t* ub = (__ubuf__ uint8_t*)0; + + AscendC::SetFlag(0); + AscendC::SetFlag(1); + AscendC::SetFlag(0); + AscendC::SetFlag(1); + + constexpr int kBlockK = 1024; + constexpr int kStages = 2; + for (int w = 0; w < 4; ++w) { + int slot = w & 1; + AscendC::WaitFlag(slot); + // Producer would MTE2 into the ping-pong UB slot for this stage. + AscendC::SetFlag(slot); + AscendC::WaitFlag(slot); + AscendC::WaitFlag(slot); + + __ubuf__ bfloat16_t* x_ub = (__ubuf__ bfloat16_t*)(ub + slot * kBlockK * 2); + __ubuf__ float* sc_ub = (__ubuf__ float*)(ub + kStages * kBlockK * 2); + __ubuf__ bfloat16_t* o_ub = (__ubuf__ bfloat16_t*)(ub + slot * kBlockK * 2 + 4096); + vf_scale_mul(x_ub, sc_ub, o_ub); + + AscendC::SetFlag(slot); + AscendC::SetFlag(slot); + AscendC::WaitFlag(slot); + AscendC::SetFlag(slot); + } + + AscendC::WaitFlag(0); + AscendC::WaitFlag(1); + AscendC::WaitFlag(0); + AscendC::WaitFlag(1); +} diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/target_mi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/target_mi.pto new file mode 100644 index 0000000000..39407b79f1 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/target_mi.pto @@ -0,0 +1,84 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: stages=2 / block_k=1024 Persistent schedule — dual UB ping-pong, +// pipe-event sync on slot = i % 2, and UNPK / mul / PK4 VF body per tile. +// Mirrors the ASC/CCE reference_asc_cce.asc HardEvent shell. + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind, + pto.pipeline.stages = 2 : i32, + pto.block_k = 1024 : i32 +} { + func.func @persistent_quant_tile_k1024_mi( + %x_ub0: !pto.ptr, + %x_ub1: !pto.ptr, + %y_ub0: !pto.ptr, + %y_ub1: !pto.ptr, + %scale: !pto.vreg<64xf32>, + %block_count: index) attributes {pto.kernel} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c1024 = arith.constant 1024 : index + + // Prime dual-slot events (ASC: SetFlag MTE3_V / V_MTE2 on slots 0 and 1). + pto.set_flag["PIPE_MTE3", "PIPE_V", "EVENT_ID0"] + pto.set_flag["PIPE_MTE3", "PIPE_V", "EVENT_ID1"] + pto.set_flag["PIPE_V", "PIPE_MTE2", "EVENT_ID0"] + pto.set_flag["PIPE_V", "PIPE_MTE2", "EVENT_ID1"] + + scf.for %i = %c0 to %block_count step %c1 { + %slot = arith.remui %i, %c2 : index + %use1 = arith.cmpi eq, %slot, %c1 : index + %off = arith.muli %i, %c1024 : index + + // Wait producer-ready / consumer-done for this ping-pong slot. + pto.wait_flag_dyn [#pto.pipe, #pto.pipe, %slot] + // (Host MTE2 would fill the selected UB slot here.) + pto.set_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.wait_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.wait_flag_dyn [#pto.pipe, #pto.pipe, %slot] + + %x_ub = scf.if %use1 -> (!pto.ptr) { + scf.yield %x_ub1 : !pto.ptr + } else { + scf.yield %x_ub0 : !pto.ptr + } + %y_ub = scf.if %use1 -> (!pto.ptr) { + scf.yield %y_ub1 : !pto.ptr + } else { + scf.yield %y_ub0 : !pto.ptr + } + + pto.vecscope { + %m32 = pto.pset_b32 "PAT_ALL" : !pto.mask + %x0 = pto.vlds %x_ub[%off] {dist = "UNPK_B16"} + : !pto.ptr -> !pto.vreg<64xf32> + %y0 = pto.vmul %x0, %scale, %m32 + : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask + -> !pto.vreg<64xf32> + pto.vsts %y0, %y_ub[%off], %m32 {dist = "PK4_B32"} + : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + } + + pto.set_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.set_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.wait_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.set_flag_dyn [#pto.pipe, #pto.pipe, %slot] + } + + // Drain both slots. + pto.wait_flag["PIPE_MTE3", "PIPE_V", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE3", "PIPE_V", "EVENT_ID1"] + pto.wait_flag["PIPE_V", "PIPE_MTE2", "EVENT_ID0"] + pto.wait_flag["PIPE_V", "PIPE_MTE2", "EVENT_ID1"] + return + } +} diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md new file mode 100644 index 0000000000..c54b436f89 --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md @@ -0,0 +1,65 @@ +# 06 — Small-L `ui8→ui16` widen + +## Algorithm need + +Fused quant→dequant needs a **small-L** UE8M0 widen (`ui8→ui16` at L=8) on the +compact SF path. L=256 widen already works; L=8 illegalizes to residual +`pto.vmi.extui`, so fused round-trip cannot match ASC BW and falls back to +compose / larger tiles. + +Reference for the working L=256 shape: +`test/lit/vmi_new/vmi_to_vpto_integer_casts.pto`. + +## Current slow path + +L=256 `pto.vmi.vcvt` ui8→ui16 (`current_slow_vmi.pto` / `current_slow_vmi.py`). Checked-in dump: +`lowered_vpto.pto` — key ops `pto.vcvt ... {part = "EVEN"}` and `{part = "ODD"}`. + +Reproduce: + +```bash +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +L=8 group-slot widen (`desired_vmi.pto` / `desired_vmi.py`): + +``` +error: failed to legalize operation 'pto.vmi.extui' that was explicitly marked illegal +... +error: VMI-RESIDUAL-OP: failed to convert all VMI ops/types to VPTO +``` + +## Target MI + bisheng status + +`target_mi.pto` is the small-L MI shape fused quant→dequant needs: load, +`vcvt` `{part = "EVEN"}` under `PAT_VL8`, then ui16 store. Small-L VMI +(`desired_vmi.pto`) should lower here instead of leaving `extui`. + +| Step | Result | +|---|---| +| `ptoas ... --emit-vpto` | OK | +| `ptoas ... --emit-vpto-llvm-ir` | OK | +| `bisheng ... -c -x ir` → `.o` | **OK** | + +Remaining gap is **VMI legalization** of L=8 (desired → residual `extui`), not +HIVM/bisheng on this MI shape. + +## ASC/CCE working baseline + +`reference_asc_cce.asc` widens `ui8→ui16` under `PAT_VL8` with `vcvt` / +`PART_EVEN` then stores. It compiles to a non-empty `.o` with +`bisheng --cce-aicore-only` (CCE twin of the MI path above). + +```bash +BISHENG=${BISHENG:-/usr/local/Ascend/cann-9.0.0/tools/bisheng_compiler/bin/bisheng} +ASCEND=${ASCEND:-/usr/local/Ascend/cann-9.0.0} +"$BISHENG" -O2 -fPIC -std=c++17 --npu-arch=dav-3510 --cce-aicore-only -c \ + reference_asc_cce.asc -o /tmp/ref_gap06.o \ + -I"$ASCEND/include" \ + -I"$ASCEND/compiler/tikcpp/tikcfw" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/impl" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/interface" +test -s /tmp/ref_gap06.o +``` diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.pto b/docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.pto new file mode 100644 index 0000000000..e32826e4cd --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.pto @@ -0,0 +1,20 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: L=256 ui8→ui16 vcvt (see test/lit/vmi_new/vmi_to_vpto_integer_casts.pto). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @ui8_to_ui16_l256( + %input: !pto.vmi.vreg<256xui8, #pto.vmi.layout>) + -> !pto.vmi.vreg<256xui16, #pto.vmi.layout> { + %wide = pto.vmi.vcvt %input + : !pto.vmi.vreg<256xui8, #pto.vmi.layout> + -> !pto.vmi.vreg<256xui16, #pto.vmi.layout> + return %wide : !pto.vmi.vreg<256xui16, #pto.vmi.layout> + } +} diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.py b/docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.py new file mode 100644 index 0000000000..5f15167bc0 --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.py @@ -0,0 +1,27 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``current_slow_vmi.pto`` — L=256 ui8→ui16 ``vcvt``. + +Canonical IR: ``current_slow_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def ui8_to_ui16_l256(): + ub = pto.alloc_tile(shape=[1, 256], dtype=pto.ui8) + off = pto.const(0, dtype=pto.index) + inp = pto.vmi.vload(ub.as_ptr(), off, size=256) + wide = pto.vmi.vcvt(inp, pto.ui16) + _ = wide + + +if __name__ == "__main__": + print(ui8_to_ui16_l256.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/desired_vmi.pto b/docs/feature-gaps/vmi/06_small_l_ui8_widen/desired_vmi.pto new file mode 100644 index 0000000000..453480165a --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/desired_vmi.pto @@ -0,0 +1,20 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Idiomatic small-L widen: L=8 ui8→ui16 (needed for fused quant→dequant UE8M0). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @ui8_to_ui16_l8( + %input: !pto.vmi.vreg<8xui8, #pto.vmi.layout>) + -> !pto.vmi.vreg<8xui16, #pto.vmi.layout> { + %wide = pto.vmi.vcvt %input + : !pto.vmi.vreg<8xui8, #pto.vmi.layout> + -> !pto.vmi.vreg<8xui16, #pto.vmi.layout> + return %wide : !pto.vmi.vreg<8xui16, #pto.vmi.layout> + } +} diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/desired_vmi.py b/docs/feature-gaps/vmi/06_small_l_ui8_widen/desired_vmi.py new file mode 100644 index 0000000000..c0ce08b2f4 --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/desired_vmi.py @@ -0,0 +1,31 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``desired_vmi.pto`` — L=8 ui8→ui16 widen. + +Canonical IR: ``desired_vmi.pto``. VMI legalization leaves residual ``extui``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def ui8_to_ui16_l8(): + # 32B-aligned tile; load/convert only L=8 lanes (compact SF widen). + ub = pto.alloc_tile(shape=[1, 32], dtype=pto.ui8) + off = pto.const(0, dtype=pto.index) + inp = pto.vmi.vload(ub.as_ptr(), off, size=8) + wide = pto.vmi.vcvt(inp, pto.ui16) + _ = wide + + +if __name__ == "__main__": + try: + print(ui8_to_ui16_l8.compile().mlir_text()) + except Exception as exc: # noqa: BLE001 — document the gap + print(f"desired emit/lower may fail (L=8 widen): {exc}") diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/lowered_vpto.pto b/docs/feature-gaps/vmi/06_small_l_ui8_widen/lowered_vpto.pto new file mode 100644 index 0000000000..e60f96b4a8 --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/lowered_vpto.pto @@ -0,0 +1,9 @@ +module attributes {pto.kernel_kind = #pto.kernel_kind, pto.target_arch = "a5"} { + func.func @ui8_to_ui16_l256(%arg0: !pto.vreg<256xui8>) -> (!pto.vreg<128xui16>, !pto.vreg<128xui16>) { + %0 = pto.pset_b8 "PAT_ALL" : !pto.mask + %1 = pto.vcvt %arg0, %0 {part = "EVEN"} : !pto.vreg<256xui8>, !pto.mask -> !pto.vreg<128xui16> + %2 = pto.vcvt %arg0, %0 {part = "ODD"} : !pto.vreg<256xui8>, !pto.mask -> !pto.vreg<128xui16> + return %1, %2 : !pto.vreg<128xui16>, !pto.vreg<128xui16> + } +} + diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/reference_asc_cce.asc b/docs/feature-gaps/vmi/06_small_l_ui8_widen/reference_asc_cce.asc new file mode 100644 index 0000000000..37249243ec --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/reference_asc_cce.asc @@ -0,0 +1,21 @@ +// ASC/CCE working baseline for gap 06: small-L ui8→ui16 widen (PART_EVEN). +// Compiles with bisheng --cce-aicore-only (see README). +#include "kernel_operator.h" + +__simd_vf__ inline void vf_ui8_widen_l8(__ubuf__ uint8_t* src, __ubuf__ uint16_t* dst) { + // L=8-shaped mask: the compact SF widen VMI wants on the fused path. + vector_bool v = pset_b16(PAT_VL8); + vector_u8 x; + vlds(x, src, 0, NORM); + vector_u16 y; + vcvt(y, x, v, PART_EVEN, MODE_ZEROING); + vsts(y, dst, 0, NORM_B16, v); +} + +extern "C" __global__ __vector__ void ref_gap06_ui8_widen(__gm__ uint8_t* /*gm_in*/, + __gm__ uint16_t* /*gm_out*/) { + AscendC::InitSocState(); + __ubuf__ uint8_t* ub8 = (__ubuf__ uint8_t*)0; + __ubuf__ uint16_t* ub16 = (__ubuf__ uint16_t*)(ub8 + 256); + vf_ui8_widen_l8(ub8, ub16); +} diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/target_mi.pto b/docs/feature-gaps/vmi/06_small_l_ui8_widen/target_mi.pto new file mode 100644 index 0000000000..22b9e16a05 --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/target_mi.pto @@ -0,0 +1,29 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: small-L (L=8) ui8→ui16 widen via vcvt PART_EVEN, then ui16 store. +// Fused quant→dequant needs this shape; VMI currently illegalizes to extui. + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @ui8_to_ui16_widen_l8_mi( + %src: !pto.ptr, + %dst: !pto.ptr, + %off: index) attributes {pto.kernel} { + pto.vecscope { + %m8 = pto.pset_b8 "PAT_VL8" : !pto.mask + %in = pto.vlds %src[%off] {dist = "NORM"} + : !pto.ptr -> !pto.vreg<256xui8> + %even = pto.vcvt %in, %m8 {part = "EVEN"} + : !pto.vreg<256xui8>, !pto.mask -> !pto.vreg<128xui16> + %m16 = pto.pset_b16 "PAT_VL8" : !pto.mask + pto.vsts %even, %dst[%off], %m16 + : !pto.vreg<128xui16>, !pto.ptr, !pto.mask + } + return + } +} diff --git a/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/README.md b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/README.md new file mode 100644 index 0000000000..05570b3ee7 --- /dev/null +++ b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/README.md @@ -0,0 +1,49 @@ +# 07 — bf16 UNPK / PK fuse (full strip) + +## Algorithm need + +Residual-mix **forward** strips (mix-apply / post combine) are dense bf16↔f32 +lanes: load bf16, widen, scale/accumulate, narrow, store. ASC hits +`UNPK_B16` + part-aware widen and `PK_B32` narrow store on both halves of the +unpacked register. VMI authors write continuous `vload` + `vcvt`; that path +must lower to the same full unpack/pack form for ASC-level BW. + +## Current slow path + +Continuous L=64 strip (`current_slow_vmi.pto` / `current_slow_vmi.py`). Checked-in dump: +`lowered_vpto.pto` — already emits `pto.vlds {dist = "UNPK_B16"}` and +`pto.vsts {dist = "PK_B32"}`, but only feeds **`part = "EVEN"`** after UNPK +(half of the unpacked load is unused). That half-utilized form is what residual-mix +fwd pays today versus ASC’s EVEN+ODD strip. + +Reproduce: + +```bash +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +Explicit unpack load (`desired_vmi.pto` / `desired_vmi.py`) so authors can request UNPK without +relying on continuous-load heuristics: + +``` +error: failed to legalize operation 'pto.vmi.vload' that was explicitly marked illegal + %x = pto.vmi.vload %src[%off] {dist_mode = "unpack"} +... +error: VMI-RESIDUAL-OP: failed to convert all VMI ops/types to VPTO +``` + +Also want continuous L=64 (or L=128) strips to lower with **both** EVEN and ODD +parts filled, matching `target_mi.pto`. + +## Target MI + bisheng status + +`target_mi.pto` — `UNPK_B16` load, EVEN+ODD `vcvt`, dual `vmul`, `vor` + +`PK_B32` store. + +| Step | Result | +|---|---| +| `ptoas ... --emit-vpto` | OK | +| `ptoas ... --emit-vpto-llvm-ir` | OK | +| `bisheng ... -c -x ir` → `.o` | **OK** | diff --git a/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/current_slow_vmi.pto b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/current_slow_vmi.pto new file mode 100644 index 0000000000..788d3cdde6 --- /dev/null +++ b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/current_slow_vmi.pto @@ -0,0 +1,33 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: continuous L=64 bf16 load + vcvt + mul + narrow store. +// Lowers today, but only consumes EVEN of UNPK_B16 (half of the unpacked load). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @bf16_strip_continuous( + %src: !pto.ptr, + %dst: !pto.ptr, + %scale: f32, + %off: index) { + %x = pto.vmi.vload %src[%off] + : !pto.ptr -> !pto.vmi.vreg<64xbf16> + %xf = pto.vmi.vcvt %x + : !pto.vmi.vreg<64xbf16> -> !pto.vmi.vreg<64xf32> + %s = pto.vmi.vbrc %scale + : f32 -> !pto.vmi.vreg<64xf32> + %y = pto.vmi.vmul %xf, %s + : !pto.vmi.vreg<64xf32>, !pto.vmi.vreg<64xf32> + -> !pto.vmi.vreg<64xf32> + %yo = pto.vmi.vcvt %y {saturate = "SAT"} + : !pto.vmi.vreg<64xf32> -> !pto.vmi.vreg<64xbf16> + pto.vmi.vstore %yo, %dst[%off] + : !pto.vmi.vreg<64xbf16>, !pto.ptr + return + } +} diff --git a/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/current_slow_vmi.py b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/current_slow_vmi.py new file mode 100644 index 0000000000..4df6f8c033 --- /dev/null +++ b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/current_slow_vmi.py @@ -0,0 +1,31 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``current_slow_vmi.pto`` — continuous bf16 strip. + +Canonical IR: ``current_slow_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def bf16_strip_continuous(): + src = pto.alloc_tile(shape=[1, 64], dtype=pto.bf16) + dst = pto.alloc_tile(shape=[1, 64], dtype=pto.bf16) + off = pto.const(0, dtype=pto.index) + x = pto.vmi.vload(src.as_ptr(), off, size=64) + xf = pto.vmi.vcvt(x, pto.f32) + scale = pto.vmi.vbrc(pto.f32(1.0), size=64) + y = pto.vmi.vmul(xf, scale) + yo = pto.vmi.vcvt(y, pto.bf16, saturate="SAT") + pto.vmi.vstore(yo, dst.as_ptr(), off) + + +if __name__ == "__main__": + print(bf16_strip_continuous.compile().mlir_text()) diff --git a/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/desired_vmi.pto b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/desired_vmi.pto new file mode 100644 index 0000000000..a5b3b19c73 --- /dev/null +++ b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/desired_vmi.pto @@ -0,0 +1,33 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired: explicit unpack load so authors can request UNPK_B16 without +// relying on continuous-load heuristics (illegal / residual today). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @bf16_strip_unpack( + %src: !pto.ptr, + %dst: !pto.ptr, + %scale: f32, + %off: index) { + %x = pto.vmi.vload %src[%off] {dist_mode = "unpack"} + : !pto.ptr -> !pto.vmi.vreg<64xbf16> + %xf = pto.vmi.vcvt %x + : !pto.vmi.vreg<64xbf16> -> !pto.vmi.vreg<64xf32> + %s = pto.vmi.vbrc %scale + : f32 -> !pto.vmi.vreg<64xf32> + %y = pto.vmi.vmul %xf, %s + : !pto.vmi.vreg<64xf32>, !pto.vmi.vreg<64xf32> + -> !pto.vmi.vreg<64xf32> + %yo = pto.vmi.vcvt %y {saturate = "SAT"} + : !pto.vmi.vreg<64xf32> -> !pto.vmi.vreg<64xbf16> + pto.vmi.vstore %yo, %dst[%off] + : !pto.vmi.vreg<64xbf16>, !pto.ptr + return + } +} diff --git a/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/desired_vmi.py b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/desired_vmi.py new file mode 100644 index 0000000000..4319b3092b --- /dev/null +++ b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/desired_vmi.py @@ -0,0 +1,37 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``desired_vmi.pto`` — explicit ``dist_mode=\"unpack\"`` load. + +Canonical IR: ``desired_vmi.pto``. Illegal / residual today. +PTODSL ``vload(..., dist_mode=\"unpack\")`` currently requires a one-step +``to_dtype`` widen; the IR keeps bf16 — see the ``.pto`` sibling. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def bf16_strip_unpack(): + src = pto.alloc_tile(shape=[1, 64], dtype=pto.bf16) + dst = pto.alloc_tile(shape=[1, 64], dtype=pto.bf16) + off = pto.const(0, dtype=pto.index) + # Mirrors desired_vmi.pto: dist_mode="unpack" (illegal / residual today). + x = pto.vmi.vload(src.as_ptr(), off, size=64, dist_mode="unpack") + xf = pto.vmi.vcvt(x, pto.f32) + scale = pto.vmi.vbrc(pto.f32(1.0), size=64) + y = pto.vmi.vmul(xf, scale) + yo = pto.vmi.vcvt(y, pto.bf16, saturate="SAT") + pto.vmi.vstore(yo, dst.as_ptr(), off) + + +if __name__ == "__main__": + try: + print(bf16_strip_unpack.compile().mlir_text()) + except Exception as exc: # noqa: BLE001 — document the gap + print(f"desired emit failed (expected until unpack bf16 lands): {exc}") diff --git a/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/lowered_vpto.pto b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/lowered_vpto.pto new file mode 100644 index 0000000000..c5d9179ab0 --- /dev/null +++ b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/lowered_vpto.pto @@ -0,0 +1,17 @@ +module attributes {pto.kernel_kind = #pto.kernel_kind, pto.target_arch = "a5"} { + func.func @bf16_strip_continuous(%arg0: !pto.ptr, %arg1: !pto.ptr, %arg2: f32, %arg3: index) { + %result = pto.vlds %arg0[%arg3] {dist = "UNPK_B16"} : !pto.ptr -> !pto.vreg<128xbf16> + %0 = pto.pset_b16 "PAT_ALL" : !pto.mask + %1 = pto.vcvt %result, %0 {part = "EVEN"} : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %2 = pto.pset_b32 "PAT_ALL" : !pto.mask + %3 = pto.vdup %arg2, %2 : f32, !pto.mask -> !pto.vreg<64xf32> + %4 = pto.pset_b32 "PAT_ALL" : !pto.mask + %5 = pto.vmul %1, %3, %4 : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %6 = pto.pset_b32 "PAT_ALL" : !pto.mask + %7 = pto.vcvt %5, %6 {part = "EVEN", rnd = "R", sat = "SAT"} : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<128xbf16> + %8 = pto.pge_b32 "PAT_VL64" : !pto.mask + pto.vsts %7, %arg1[%arg3], %8 {dist = "PK_B32"} : !pto.vreg<128xbf16>, !pto.ptr, !pto.mask + return + } +} + diff --git a/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/target_mi.pto b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/target_mi.pto new file mode 100644 index 0000000000..ba808d3e68 --- /dev/null +++ b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/target_mi.pto @@ -0,0 +1,46 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: UNPK_B16 load, BOTH EVEN+ODD widen, compute, PK_B32 store. +// ASC residual-mix strips use this full unpack/pack form; continuous VMI +// today only feeds EVEN after UNPK_B16. + +module { + func.func @bf16_strip_unpk_pk_mi( + %src: !pto.ptr, + %dst: !pto.ptr, + %scale: f32, + %off: index) { + pto.vecscope { + %raw = pto.vlds %src[%off] {dist = "UNPK_B16"} + : !pto.ptr -> !pto.vreg<128xbf16> + %m16 = pto.pset_b16 "PAT_ALL" : !pto.mask + %even = pto.vcvt %raw, %m16 {part = "EVEN"} + : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %odd = pto.vcvt %raw, %m16 {part = "ODD"} + : !pto.vreg<128xbf16>, !pto.mask -> !pto.vreg<64xf32> + %m32 = pto.pset_b32 "PAT_ALL" : !pto.mask + %s = pto.vdup %scale, %m32 : f32, !pto.mask -> !pto.vreg<64xf32> + %ye = pto.vmul %even, %s, %m32 + : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %yo = pto.vmul %odd, %s, %m32 + : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<64xf32> + %ne = pto.vcvt %ye, %m32 {part = "EVEN", rnd = "R", sat = "SAT"} + : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<128xbf16> + %no = pto.vcvt %yo, %m32 {part = "ODD", rnd = "R", sat = "SAT"} + : !pto.vreg<64xf32>, !pto.mask -> !pto.vreg<128xbf16> + %packed = pto.vor %ne, %no, %m16 + : !pto.vreg<128xbf16>, !pto.vreg<128xbf16>, !pto.mask + -> !pto.vreg<128xbf16> + %store_m = pto.pge_b32 "PAT_VL64" : !pto.mask + pto.vsts %packed, %dst[%off], %store_m {dist = "PK_B32"} + : !pto.vreg<128xbf16>, !pto.ptr, !pto.mask + } + return + } +} diff --git a/docs/feature-gaps/vmi/README.md b/docs/feature-gaps/vmi/README.md new file mode 100644 index 0000000000..aef20bd4a0 --- /dev/null +++ b/docs/feature-gaps/vmi/README.md @@ -0,0 +1,100 @@ +# VMI feature gaps (docs fixtures) + +Minimal reproducers for PTOAS capabilities that still block VMI from reaching +ASC bandwidth on quantization and residual-mix style vector kernels. These are +**documentation fixtures**, not lit tests (no `// RUN:` lines). + +Each subdirectory has: + +| File | Role | +|---|---| +| `README.md` | Algorithm need, current slow path + reproduce, desired failure, target MI / bisheng status | +| `current_slow_vmi.pto` | Compiles and runs today, but slow / non-optimal; **must** lower with `pto-test-opt` + `$PASS` | +| `current_slow_vmi.py` | PTODSL Python form of the slow path (`@pto.jit` + `pto.vmi.*`) | +| `lowered_vpto.pto` | Checked-in dump from that lower | +| `desired_vmi.pto` | Idiomatic ask (may fail — README pastes the error) | +| `desired_vmi.py` | PTODSL Python form of the idiomatic ask (may fail to emit/lower) | +| `target_mi.pto` | Desired `pto.mi` / VPTO shape; exercised with `ptoas` + `bisheng` | +| `reference_asc_cce.asc` | Gaps **03–06**: minimal ASC/CCE kernel that `bisheng --cce-aicore-only` accepts (e2e CCE proof; required where `target_mi` HIVM still crashes — **03–05**) | + +Python fixtures use **PTODSL** (`from ptodsl import pto`). Canonical semantics stay in +the `.pto` siblings when a Python surface lag exists (e.g. missing `dcache_bypass` +or `pto.vmi.vpack`). Emit MLIR with: + +```bash +# From repo root, with PTODSL + MLIR Python bindings on PYTHONPATH: +python docs/feature-gaps/vmi//current_slow_vmi.py +python docs/feature-gaps/vmi//desired_vmi.py +``` + +## Gaps + +| # | Gap | Why ASC-level BW cares | +|---|---|---| +| 01 | [Scalar GM dcache bypass](01_scalar_gm_dcache_bypass/) | Tiny per-block SF writes still force bulk MTE | +| 02 | [Compact inverse `vbrc`](02_compact_inv_vbrc/) | Padded recip + BRC reload tax on quant | +| 03 | [Direct scalar / 1PT reduce store](03_direct_scalar_reduce_store/) | Residual-mix post-bwd still pads to 8; ASC uses ONEPT/1PT | +| 04 | [Packed UE8M0 SF](04_packed_ue8m0_sf/) | Unpacked SF doubles traffic | +| 05 | [Persistent stages / `block_k`](05_persistent_stages_block_k/) | Per-block + cast-back capped at stages=1 / `block_k≤512`; stages=2 faults in practice | +| 06 | [Small-L ui8→ui16 widen](06_small_l_ui8_widen/) | L=8 `vcvt` illegalizes to residual `extui` (blocks fused quant→dequant) | +| 07 | [bf16 UNPK / PK fuse](07_bf16_unpk_pk_fuse/) | Residual-mix **fwd** strips: continuous load only feeds EVEN after UNPK; `dist_mode=unpack` illegal | + +Tip findings (high level, no external repo names): residual-mix **bwd** still +pays a reduce pad-to-8 store tax (**03**); residual-mix **fwd** needs full +UNPK/PK strip fuse (**07**); per-block / cast-back Persistent stay on stages=1 / +`block_k≤512` — MI schedule+flags emit but HIVM/bisheng still crash (**05**); +fused small-L UE8M0 widen hits L=8 VMI `extui` while the MI `PAT_VL8` shape +already compiles (**06**). + +## How to reproduce + +Assumes a local `build/` of this tree (or any install with `pto-test-opt` / +`ptoas` on `PATH`) and Ascend `bisheng` available: + +```bash +# From repo root after a normal build: +PTO_TEST_OPT=${PTO_TEST_OPT:-$PWD/build/tools/pto-test-opt/pto-test-opt} +PTOAS=${PTOAS:-$PWD/build/tools/ptoas/ptoas} +# Example Ascend install; override if yours differs: +BISHENG=${BISHENG:-/usr/local/Ascend/cann-9.0.0/tools/bisheng_compiler/bin/bisheng} +export PATH="$(dirname "$BISHENG"):$PATH" + +PASS='-vmi-lower-unified-to-legacy -vmi-mask-granularity-assignment -vmi-layout-assignment -vmi-to-vpto' +cd docs/feature-gaps/vmi/ + +# Current slow path (must succeed; refresh lowered_vpto.pto) +"$PTO_TEST_OPT" current_slow_vmi.pto $PASS -o lowered_vpto.pto + +# Desired (often fails; paste error into README) +"$PTO_TEST_OPT" desired_vmi.pto $PASS -o /tmp/desired.pto + +# Target MI → VPTO IR → HIVM LLVM → device object +"$PTOAS" --pto-arch=a5 --pto-backend=vpto --pto-level=level3 \ + --emit-vpto -o /tmp/t.emit.pto target_mi.pto +"$PTOAS" --pto-arch=a5 --pto-backend=vpto --pto-level=level3 \ + --emit-vpto-llvm-ir -o /tmp/t.ll target_mi.pto +"$BISHENG" --target=hiipu64-hisilicon-cce -march=dav-c310-vec \ + --cce-aicore-arch=dav-c310-vec --cce-aicore-only -c -x ir /tmp/t.ll -o /tmp/t.o +``` + +Use `--emit-vpto-llvm-ir` for LLVM IR (not the older `--vpto-emit-hivm-llvm` name). + +## ASC/CCE reference compile (gaps 03–06) + +For gaps where `target_mi` → HIVM IR → `bisheng -x ir` still crashes (**03–05**), +the sibling `reference_asc_cce.asc` proves the same instruction / schedule +shape is legal on the Ascend CCE path. Gap **06** keeps a CCE twin even though +its small-L `target_mi` now builds with `-x ir`: + +```bash +BISHENG=${BISHENG:-/usr/local/Ascend/cann-9.0.0/tools/bisheng_compiler/bin/bisheng} +ASCEND=${ASCEND:-/usr/local/Ascend/cann-9.0.0} +cd docs/feature-gaps/vmi/ # 03..06 +"$BISHENG" -O2 -fPIC -std=c++17 --npu-arch=dav-3510 --cce-aicore-only -c \ + reference_asc_cce.asc -o /tmp/ref.o \ + -I"$ASCEND/include" \ + -I"$ASCEND/compiler/tikcpp/tikcfw" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/impl" \ + -I"$ASCEND/compiler/tikcpp/tikcfw/interface" +test -s /tmp/ref.o +```