From 75b8af6a8ad71fbd06bca8553c96e5efddade1bf Mon Sep 17 00:00:00 2001 From: learning-chip Date: Fri, 24 Jul 2026 21:59:20 +0800 Subject: [PATCH 1/9] docs: add VMI/PTODSL feature-gap reproducers for scalar GM, vbrc, 1PT, UE8M0, Persistent Co-authored-by: Cursor --- .../vmi/01_scalar_gm_dcache_bypass/README.md | 21 +++++++ .../01_scalar_gm_dcache_bypass/buggy_vmi.pto | 33 ++++++++++ .../desired_vmi.pto | 22 +++++++ .../01_scalar_gm_dcache_bypass/target_mi.pto | 22 +++++++ .../vmi/02_compact_inv_vbrc/README.md | 20 ++++++ .../vmi/02_compact_inv_vbrc/buggy_vmi.pto | 27 ++++++++ .../vmi/02_compact_inv_vbrc/desired_vmi.pto | 18 ++++++ .../vmi/02_compact_inv_vbrc/target_mi.pto | 20 ++++++ .../03_direct_scalar_reduce_store/README.md | 20 ++++++ .../buggy_vmi.pto | 29 +++++++++ .../desired_vmi.pto | 25 ++++++++ .../target_mi.pto | 24 +++++++ .../vmi/04_packed_ue8m0_sf/README.md | 19 ++++++ .../vmi/04_packed_ue8m0_sf/buggy_vmi.pto | 21 +++++++ .../vmi/04_packed_ue8m0_sf/desired_vmi.pto | 23 +++++++ .../vmi/04_packed_ue8m0_sf/target_mi.pto | 22 +++++++ .../05_persistent_stages_block_k/README.md | 20 ++++++ .../buggy_vmi.pto | 48 ++++++++++++++ .../desired_vmi.pto | 62 +++++++++++++++++++ .../target_mi.pto | 34 ++++++++++ docs/feature-gaps/vmi/README.md | 28 +++++++++ 21 files changed, 558 insertions(+) create mode 100644 docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md create mode 100644 docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto create mode 100644 docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.pto create mode 100644 docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/target_mi.pto create mode 100644 docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md create mode 100644 docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto create mode 100644 docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto create mode 100644 docs/feature-gaps/vmi/02_compact_inv_vbrc/target_mi.pto create mode 100644 docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md create mode 100644 docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto create mode 100644 docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.pto create mode 100644 docs/feature-gaps/vmi/03_direct_scalar_reduce_store/target_mi.pto create mode 100644 docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md create mode 100644 docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto create mode 100644 docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.pto create mode 100644 docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.pto create mode 100644 docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md create mode 100644 docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto create mode 100644 docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.pto create mode 100644 docs/feature-gaps/vmi/05_persistent_stages_block_k/target_mi.pto create mode 100644 docs/feature-gaps/vmi/README.md 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..6a9dae3480 --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md @@ -0,0 +1,21 @@ +# 01 — Scalar GM dcache-bypass store/load + +## Algorithm need + +Per-block quantization over groups of shape **32×32** produces one f32 scale +factor (SF) per group. After the scale is computed in UB, that single scalar +must be written to GM (and later re-read) so consumers see a dense SF tensor. + +## Current failure + +PTOAS does not expose a scalar GM path with dcache bypass for these tiny +writes/reads. Authors stage the SF in UB and issue a bulk MTE +(`pto.copy_ubuf_to_gm` / `pto.copy_gm_to_ubuf`) even when `len_burst` is only a +few bytes. That adds MTE setup, alignment padding, and pipeline sync out of +proportion to the payload. + +## Desired PTOAS behavior + +Allow VMI/PTODSL to store or load a scalar (or compact slots=1 vector) directly +against `!pto.ptr` with an explicit **dcache-bypass** attribute, lowering +to a single GM scalar path without a bulk MTE burst. diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto new file mode 100644 index 0000000000..84a624d3a8 --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_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. + +// Compiling-today style workaround: stage one f32 SF in UB, then bulk MTE to GM. + +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, + %scale: !pto.vmi.vreg<1xf32>) { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c0_i64 = arith.constant 0 : i64 + %c1_i64 = arith.constant 1 : i64 + %c4_i64 = arith.constant 4 : i64 + %c32_i64 = arith.constant 32 : i64 + + // Store the scalar SF into a padded UB slot (MTE wants aligned bursts). + pto.vmi.vstore %scale, %sf_ub[%c0], %c1 {group = 1} + : !pto.vmi.vreg<1xf32>, !pto.ptr + + // Bulk MTE for a tiny SF payload (4B useful / 32B burst). + 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/desired_vmi.pto b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.pto new file mode 100644 index 0000000000..867705a8d4 --- /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 (may 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>, + %off: index) { + %c1 = arith.constant 1 : index + // Desired: scalar / slots=1 store to GM, bypassing dcache. + pto.vmi.vstore %scale, %sf_gm[%off], %c1 {group = 1, dcache_bypass} + : !pto.vmi.vreg<1xf32>, !pto.ptr + 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..45273061c6 --- /dev/null +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/target_mi.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. + +// Desired VPTO / pto.mi shape: one GM scalar store, dcache-bypass, no MTE burst. + +module { + func.func @scale_factor_gm_bypass_mi( + %sf_gm: !pto.ptr, + %scale: !pto.vreg<64xf32>, + %off: index) { + %m = pto.pge_b32 "PAT_VL1" : !pto.mask + // Single-point GM store; dcache_bypass on the memory op. + pto.vsts %scale, %sf_gm[%off], %m {dist = "1PT_B32", dcache_bypass} + : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + 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..936b49e907 --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md @@ -0,0 +1,20 @@ +# 02 — Compact inverse `vbrc` (no pad + `BRC`) + +## Algorithm need + +Quantization kernels compute per-group reciprocal scale factors, then broadcast +each inverse across the lanes of its group before multiplying the activation +tile. Authors want a **compact** inverse vector (one slot per group) fanned +out with `pto.vmi.vbrc` / a `BRC` load. + +## Current failure + +VPTO lowering expects reciprocal scales to be materialized as a **padded +full-chunk** layout and reloaded with `dist_mode = "brc"`. Compact inverse +`vbrc` (e.g. 8 valid slots) does not lower cleanly, so kernels pad UB and pay +extra stores plus a broadcast reload. + +## Desired PTOAS behavior + +Lower compact inverse `vbrc` (or a slots-aligned `vload {dist_mode = "brc"}`) +directly to `BRC_B*` without requiring pad-to-full-chunk waste in UB. diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto new file mode 100644 index 0000000000..54fad55bc7 --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto @@ -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. + +// Compiling-today style: pad recip scales to a full chunk, 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_compact: !pto.vmi.vreg<8xf32>) -> !pto.vmi.vreg<64xf32> { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + + // Pad compact 8 inverses into a full 64-lane UB line before BRC. + pto.vmi.vstore %inv_compact, %inv_ub[%c0], %c1 {group = 8} + : !pto.vmi.vreg<8xf32>, !pto.ptr + + // Reload with dist_mode=brc (requires the padded layout today). + %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/desired_vmi.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto new file mode 100644 index 0000000000..73de45971b --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto @@ -0,0 +1,18 @@ +// 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, no pad round-trip (may not lower today). + +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.vreg<64xf32> { + %inv = pto.vmi.vbrc %inv_compact {group = 8} + : !pto.vmi.vreg<8xf32> -> !pto.vmi.vreg<64xf32> + return %inv : !pto.vmi.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..89a8ed37be --- /dev/null +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/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 VPTO / pto.mi: BRC from compact slots, no pad store + reload waste. + +module { + func.func @compact_inv_vbrc_mi( + %inv_ub: !pto.ptr, + %off: index) -> !pto.vreg<64xf32> { + // Direct BRC_B32 from the compact inverse line (8 valid scalars). + %inv = pto.vlds %inv_ub[%off] {dist = "BRC_B32"} + : !pto.ptr -> !pto.vreg<64xf32> + return %inv : !pto.vreg<64xf32> + } +} 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..2661ed9229 --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md @@ -0,0 +1,20 @@ +# 03 — Direct scalar reduce store (`1PT_B32`) + +## Algorithm need + +Residual-mix kernels (MHC-style) accumulate a tile and reduce to **one f32 +scalar per group** (mix gradient / residual coefficient). That compact result +must be stored densely to UB/GM for the next stage. + +## Current failure + +Practical VMI authoring today pads the reduce result to **8 slots** (or +broadcasts with `vbrc` sized for 8) before store, because a true slots=1 / +single-scalar store path is unreliable or rejected. The pad tax shows up as +extra UB traffic and prevents a clean `1PT_B32` store. + +## Desired PTOAS behavior + +`vcadd` / `group_reduce_addf` producing `V<1×f32>` or slots=1 compact vectors +must lower to a direct `pto.vsts {dist = "1PT_B32"}` (or equivalent) with no +mandatory pad-to-8 store. diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto new file mode 100644 index 0000000000..7450c4aa0b --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.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. + +// Compiling-today style: reduce, then pad via vbrc size=8 before 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 + %sum1 = pto.vmi.vcadd %tile, %mask {reassoc} + : !pto.vmi.vreg<64xf32>, !pto.vmi.mask<64xpred> + -> !pto.vmi.vreg<1xf32> + + // Workaround: fan the scalar into 8 slots, then group-store the pad. + %sum8 = pto.vmi.vbrc %sum1 {group = 1} + : !pto.vmi.vreg<1xf32> -> !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/desired_vmi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.pto new file mode 100644 index 0000000000..19ad32753a --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_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. + +// Idiomatic VMI: reduce to one scalar and store directly (may not lower today). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @direct_scalar_reduce_store( + %tile: !pto.vmi.vreg<64xf32>, + %mask: !pto.vmi.mask<64xpred>, + %dst: !pto.ptr, + %off: index) { + %c1 = arith.constant 1 : index + %sum = pto.vmi.vcadd %tile, %mask {reassoc} + : !pto.vmi.vreg<64xf32>, !pto.vmi.mask<64xpred> + -> !pto.vmi.vreg<1xf32> + pto.vmi.vstore %sum, %dst[%off], %c1 {group = 1} + : !pto.vmi.vreg<1xf32>, !pto.ptr + return + } +} 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..c99e9cadcc --- /dev/null +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/target_mi.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. + +// Desired VPTO / pto.mi: vcadd + 1PT_B32 store, no pad-to-8. + +module { + func.func @direct_scalar_reduce_store_mi( + %tile: !pto.vreg<64xf32>, + %mask: !pto.mask, + %dst: !pto.ptr, + %off: index) { + %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..68adfea0ab --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md @@ -0,0 +1,19 @@ +# 04 — Packed UE8M0 scale factors (`pack_factor=2`) + +## Algorithm need + +Per-token MXFP8 quantization stores UE8M0 scale factors with **`pack_factor=2`**: +two consecutive `ui8` exponents share one `ui16` element in GM/UB. This halves +SF bandwidth versus an unpacked `ui8` tensor. + +## Current failure + +VMI/PTODSL paths that try to pack adjacent UE8M0 bytes (reinterpret / element +pack into `ui16`) do not lower reliably. Authors keep **unpacked `ui8` SF +only**, doubling scale-factor traffic on the memory pipe. + +## Desired PTOAS behavior + +Provide a first-class pack path so a logical `ui8` UE8M0 vector lowers through +`PK` (or equivalent) into packed `ui16` storage with `pack_factor=2`, and the +inverse unpack on load. diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto new file mode 100644 index 0000000000..d8b8cc7c8c --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto @@ -0,0 +1,21 @@ +// 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. + +// Compiling-today style: 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) { + // No pack_factor=2 path — write raw ui8 SF tensor. + 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/desired_vmi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.pto new file mode 100644 index 0000000000..987c410365 --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.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. + +// Idiomatic VMI: pack_factor=2 UE8M0 store into ui16 (may not lower today). + +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) { + // Desired: element pack of adjacent UE8M0 bytes into ui16 words (pack_factor=2). + %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/target_mi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.pto new file mode 100644 index 0000000000..0a2e2d1038 --- /dev/null +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.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. + +// Desired VPTO / pto.mi: PK pack of UE8M0 bytes into ui16 storage. + +module { + func.func @packed_ue8m0_sf_store_mi( + %sf_u8: !pto.vreg<256xui8>, + %dst: !pto.ptr, + %off: index) { + %m = pto.pset_b8 "PAT_ALL" : !pto.mask + // Pack adjacent ui8 lanes (pack_factor=2) then store as ui16 / PK path. + pto.vsts %sf_u8, %dst[%off], %m {dist = "PK_B16"} + : !pto.vreg<256xui8>, !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..309dc54f09 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md @@ -0,0 +1,20 @@ +# 05 — Persistent multi-buffer stages and `block_k > 512` + +## Algorithm need + +Long-K quantization / cast-back vector kernels run as **Persistent** loops over +the K dimension. Authors want multi-buffer software pipeline stages (double / +triple buffer) and tile sizes with **`block_k > 512`** so MTE and vector pipes +stay overlapped on large hidden dimensions. + +## Current failure + +Practical PTOAS acceptance for these Persistent VMI modules is capped at +**`stages = 1`** and **`block_k ≤ 512`**. Raising either is rejected or +miscompiles, so kernels cannot express a legal multi-buffer Persistent schedule. + +## Desired PTOAS behavior + +Accept Persistent vector modules with `stages >= 2` and `block_k > 512`, and +lower them to a correct multi-buffer VPTO schedule (distinct UB stage buffers + +pipeline sync) without forcing the single-stage / K≤512 workaround. diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto new file mode 100644 index 0000000000..46f3957274 --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto @@ -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. + +// Compiling-today style Persistent vector module. +// Constraint comment: stages=1 / block_k<=512 required for legal lowering today. + +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/desired_vmi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.pto new file mode 100644 index 0000000000..35c047ae12 --- /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, block_k=1024 (may 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/target_mi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/target_mi.pto new file mode 100644 index 0000000000..dabf41a47a --- /dev/null +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/target_mi.pto @@ -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. + +// Desired VPTO / pto.mi sketch: legal multi-buffer Persistent schedule. + +module attributes { + pto.target_arch = "a5", + 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>, + %off: index) { + %m16 = pto.pset_b16 "PAT_ALL" : !pto.mask + %m32 = pto.pset_b32 "PAT_ALL" : !pto.mask + // Stage 0 buffer pair active this iteration (stage 1 overlaps MTE). + %x0 = pto.vlds %x_ub0[%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_ub0[%off], %m32 {dist = "PK4_B32"} + : !pto.vreg<64xf32>, !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..65bd0ccb00 --- /dev/null +++ b/docs/feature-gaps/vmi/README.md @@ -0,0 +1,28 @@ +# VMI / PTODSL feature gaps + +Minimal reproducers for PTOAS capabilities that VMI/PTODSL authors need when +writing high-throughput **quantization** and **residual-mix** style vector +kernels. Each subdirectory documents one gap with: + +| File | Role | +|---|---| +| `README.md` | Algorithm need, current failure, desired PTOAS behavior | +| `buggy_vmi.pto` | Compiling-today style workaround IR (illustrative) | +| `desired_vmi.pto` | Idiomatic VMI using the missing feature (may not lower today) | +| `target_mi.pto` | Desired VPTO / `pto.mi` shape after lowering | + +These samples are documentation fixtures, not lit tests (no `// RUN:` lines). + +## Gaps + +| # | Gap | Why it matters | +|---|---|---| +| 01 | [Scalar GM dcache bypass](01_scalar_gm_dcache_bypass/) | Per-block quantization writes one f32 scale per 32×32 group to GM. Forcing bulk MTE for a few bytes burns bandwidth and sync. | +| 02 | [Compact inverse `vbrc`](02_compact_inv_vbrc/) | Quant kernels fan reciprocal scales across lanes. Padding to a full chunk before `dist_mode=brc` wastes UB and ops. | +| 03 | [Direct scalar reduce store](03_direct_scalar_reduce_store/) | Residual-mix kernels reduce a tile to one f32 scalar per group. Padding stores to 8 slots pollutes bandwidth and masks `1PT_B32`. | +| 04 | [Packed UE8M0 scale factors](04_packed_ue8m0_sf/) | Per-token MXFP8 uses UE8M0 scales with `pack_factor=2`. Unpacked `ui8` doubles SF traffic vs packed `ui16`. | +| 05 | [Persistent stages / `block_k`](05_persistent_stages_block_k/) | Persistent vector loops need multi-buffer stages and `block_k > 512` for long-K quant/cast tiles. Caps leave the pipe underfed. | + +Closing these gaps improves authoring productivity (fewer workarounds in +PTODSL) and end-to-end performance of quantization + residual-mix kernels +without changing their algorithm. From 3ecb2bb8978df4cd45061cf2df630b7ba92f4304 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Sat, 25 Jul 2026 00:00:47 +0800 Subject: [PATCH 2/9] docs: revise VMI feature-gap repros with lowering dumps and bisheng checks Align gaps with remaining VMI vs ASC BW blockers; add small-L ui8 widen; check in lowered VPTO for working workarounds. --- .../vmi/01_scalar_gm_dcache_bypass/README.md | 46 +++++--- .../01_scalar_gm_dcache_bypass/buggy_vmi.pto | 15 +-- .../desired_vmi.pto | 8 +- .../lowered_vpto.pto | 10 ++ .../01_scalar_gm_dcache_bypass/target_mi.pto | 14 ++- .../vmi/02_compact_inv_vbrc/README.md | 46 +++++--- .../vmi/02_compact_inv_vbrc/buggy_vmi.pto | 14 ++- .../vmi/02_compact_inv_vbrc/desired_vmi.pto | 8 +- .../vmi/02_compact_inv_vbrc/lowered_vpto.pto | 12 +++ .../vmi/02_compact_inv_vbrc/target_mi.pto | 17 +-- .../03_direct_scalar_reduce_store/README.md | 45 +++++--- .../buggy_vmi.pto | 10 +- .../desired_vmi.pto | 13 ++- .../lowered_vpto.pto | 10 ++ .../target_mi.pto | 16 +-- .../vmi/04_packed_ue8m0_sf/README.md | 42 +++++--- .../vmi/04_packed_ue8m0_sf/buggy_vmi.pto | 3 +- .../vmi/04_packed_ue8m0_sf/desired_vmi.pto | 3 +- .../vmi/04_packed_ue8m0_sf/lowered_vpto.pto | 13 +++ .../vmi/04_packed_ue8m0_sf/target_mi.pto | 13 +-- .../05_persistent_stages_block_k/README.md | 47 +++++--- .../buggy_vmi.pto | 3 +- .../desired_vmi.pto | 2 +- .../lowered_vpto.pto | 100 ++++++++++++++++++ .../target_mi.pto | 24 ++--- .../vmi/06_small_l_ui8_widen/README.md | 43 ++++++++ .../vmi/06_small_l_ui8_widen/buggy_vmi.pto | 20 ++++ .../vmi/06_small_l_ui8_widen/desired_vmi.pto | 20 ++++ .../vmi/06_small_l_ui8_widen/lowered_vpto.pto | 9 ++ .../vmi/06_small_l_ui8_widen/target_mi.pto | 27 +++++ docs/feature-gaps/vmi/README.md | 71 +++++++++---- 31 files changed, 555 insertions(+), 169 deletions(-) create mode 100644 docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/lowered_vpto.pto create mode 100644 docs/feature-gaps/vmi/02_compact_inv_vbrc/lowered_vpto.pto create mode 100644 docs/feature-gaps/vmi/03_direct_scalar_reduce_store/lowered_vpto.pto create mode 100644 docs/feature-gaps/vmi/04_packed_ue8m0_sf/lowered_vpto.pto create mode 100644 docs/feature-gaps/vmi/05_persistent_stages_block_k/lowered_vpto.pto create mode 100644 docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md create mode 100644 docs/feature-gaps/vmi/06_small_l_ui8_widen/buggy_vmi.pto create mode 100644 docs/feature-gaps/vmi/06_small_l_ui8_widen/desired_vmi.pto create mode 100644 docs/feature-gaps/vmi/06_small_l_ui8_widen/lowered_vpto.pto create mode 100644 docs/feature-gaps/vmi/06_small_l_ui8_widen/target_mi.pto 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 index 6a9dae3480..c25b68546a 100644 --- a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md @@ -2,20 +2,40 @@ ## Algorithm need -Per-block quantization over groups of shape **32×32** produces one f32 scale -factor (SF) per group. After the scale is computed in UB, that single scalar -must be written to GM (and later re-read) so consumers see a dense SF tensor. +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 failure +## Working-today workaround -PTOAS does not expose a scalar GM path with dcache bypass for these tiny -writes/reads. Authors stage the SF in UB and issue a bulk MTE -(`pto.copy_ubuf_to_gm` / `pto.copy_gm_to_ubuf`) even when `len_burst` is only a -few bytes. That adds MTE setup, alignment padding, and pipeline sync out of -proportion to the payload. +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). -## Desired PTOAS behavior +Reproduce: -Allow VMI/PTODSL to store or load a scalar (or compact slots=1 vector) directly -against `!pto.ptr` with an explicit **dcache-bypass** attribute, lowering -to a single GM scalar path without a bulk MTE burst. +```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 buggy_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`): + +``` +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/buggy_vmi.pto b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto index 84a624d3a8..8ea7f4d7cd 100644 --- a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto @@ -6,25 +6,16 @@ // 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. -// Compiling-today style workaround: stage one f32 SF in UB, then bulk MTE to GM. +// Working-today workaround: 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, - %scale: !pto.vmi.vreg<1xf32>) { - %c0 = arith.constant 0 : index - %c1 = arith.constant 1 : index + %sf_gm: !pto.ptr) { %c0_i64 = arith.constant 0 : i64 %c1_i64 = arith.constant 1 : i64 - %c4_i64 = arith.constant 4 : i64 %c32_i64 = arith.constant 32 : i64 - - // Store the scalar SF into a padded UB slot (MTE wants aligned bursts). - pto.vmi.vstore %scale, %sf_ub[%c0], %c1 {group = 1} - : !pto.vmi.vreg<1xf32>, !pto.ptr - - // Bulk MTE for a tiny SF payload (4B useful / 32B burst). + // 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 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 index 867705a8d4..4de83b0147 100644 --- 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 @@ -6,17 +6,17 @@ // 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 (may not lower today). +// 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>, + %scale: !pto.vmi.vreg<1xf32, #pto.vmi.layout>, %off: index) { %c1 = arith.constant 1 : index - // Desired: scalar / slots=1 store to GM, bypassing dcache. pto.vmi.vstore %scale, %sf_gm[%off], %c1 {group = 1, dcache_bypass} - : !pto.vmi.vreg<1xf32>, !pto.ptr + : !pto.vmi.vreg<1xf32, #pto.vmi.layout>, + !pto.ptr return } } 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 index 45273061c6..002194b6ea 100644 --- 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 @@ -6,17 +6,15 @@ // 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 VPTO / pto.mi shape: one GM scalar store, dcache-bypass, no MTE burst. +// Desired MI: scalar GM store with dcache bypass (no bulk MTE). -module { +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { func.func @scale_factor_gm_bypass_mi( %sf_gm: !pto.ptr, - %scale: !pto.vreg<64xf32>, - %off: index) { - %m = pto.pge_b32 "PAT_VL1" : !pto.mask - // Single-point GM store; dcache_bypass on the memory op. - pto.vsts %scale, %sf_gm[%off], %m {dist = "1PT_B32", dcache_bypass} - : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + %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 index 936b49e907..1163918cb6 100644 --- a/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md @@ -1,20 +1,40 @@ -# 02 — Compact inverse `vbrc` (no pad + `BRC`) +# 02 — Compact inverse `vbrc` (no pad + BRC) ## Algorithm need -Quantization kernels compute per-group reciprocal scale factors, then broadcast -each inverse across the lanes of its group before multiplying the activation -tile. Authors want a **compact** inverse vector (one slot per group) fanned -out with `pto.vmi.vbrc` / a `BRC` load. +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 failure +## Working-today workaround -VPTO lowering expects reciprocal scales to be materialized as a **padded -full-chunk** layout and reloaded with `dist_mode = "brc"`. Compact inverse -`vbrc` (e.g. 8 valid slots) does not lower cleanly, so kernels pad UB and pay -extra stores plus a broadcast reload. +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"}`. -## Desired PTOAS behavior +Reproduce: -Lower compact inverse `vbrc` (or a slots-aligned `vload {dist_mode = "brc"}`) -directly to `BRC_B*` without requiring pad-to-full-chunk waste in UB. +```bash +$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +Compact `vbrc` from `num_groups=8, slots=1` (see `desired_vmi.pto`): + +``` +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/buggy_vmi.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto index 54fad55bc7..e6badf4206 100644 --- a/docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto @@ -6,20 +6,18 @@ // 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. -// Compiling-today style: pad recip scales to a full chunk, then BRC reload. +// Working-today workaround: 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_compact: !pto.vmi.vreg<8xf32>) -> !pto.vmi.vreg<64xf32> { + %inv_padded: !pto.vmi.vreg<8xf32, #pto.vmi.layout>) + -> !pto.vmi.vreg<64xf32> { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index - - // Pad compact 8 inverses into a full 64-lane UB line before BRC. - pto.vmi.vstore %inv_compact, %inv_ub[%c0], %c1 {group = 8} - : !pto.vmi.vreg<8xf32>, !pto.ptr - - // Reload with dist_mode=brc (requires the padded layout today). + 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/desired_vmi.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto index 73de45971b..c2621055ce 100644 --- a/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.pto @@ -6,13 +6,15 @@ // 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, no pad round-trip (may not lower today). +// 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.vreg<64xf32> { + %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.vreg<64xf32> + : !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/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 index 89a8ed37be..f669f53dea 100644 --- a/docs/feature-gaps/vmi/02_compact_inv_vbrc/target_mi.pto +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/target_mi.pto @@ -6,15 +6,20 @@ // 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 VPTO / pto.mi: BRC from compact slots, no pad store + reload waste. +// 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) -> !pto.vreg<64xf32> { - // Direct BRC_B32 from the compact inverse line (8 valid scalars). - %inv = pto.vlds %inv_ub[%off] {dist = "BRC_B32"} - : !pto.ptr -> !pto.vreg<64xf32> - return %inv : !pto.vreg<64xf32> + %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 index 2661ed9229..b87ded3b39 100644 --- a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md @@ -1,20 +1,39 @@ -# 03 — Direct scalar reduce store (`1PT_B32`) +# 03 — Direct scalar / 1PT reduce store ## Algorithm need -Residual-mix kernels (MHC-style) accumulate a tile and reduce to **one f32 -scalar per group** (mix gradient / residual coefficient). That compact result -must be stored densely to UB/GM for the next stage. +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 failure +## Working-today workaround -Practical VMI authoring today pads the reduce result to **8 slots** (or -broadcasts with `vbrc` sized for 8) before store, because a true slots=1 / -single-scalar store path is unreliable or rejected. The pad tax shows up as -extra UB traffic and prevents a clean `1PT_B32` store. +Group-reduce with `{group = 8}` and store the 8-slot result (`buggy_vmi.pto`). +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`. -## Desired PTOAS behavior +Reproduce: -`vcadd` / `group_reduce_addf` producing `V<1×f32>` or slots=1 compact vectors -must lower to a direct `pto.vsts {dist = "1PT_B32"}` (or equivalent) with no -mandatory pad-to-8 store. +```bash +$PTO_TEST_OPT buggy_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`). +**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 workaround emits. + +## Target MI + bisheng status + +`target_mi.pto` asks for `pto.vsts {dist = "1PT_B32"}` after `pto.vcadd`. + +| 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`) | diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto index 7450c4aa0b..43754189d2 100644 --- a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto @@ -6,7 +6,7 @@ // 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. -// Compiling-today style: reduce, then pad via vbrc size=8 before store. +// Working-today workaround: 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( @@ -15,13 +15,9 @@ module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind, %off: index) { %c1 = arith.constant 1 : index - %sum1 = pto.vmi.vcadd %tile, %mask {reassoc} + %sum8 = pto.vmi.vcadd %tile, %mask {group = 8, reassoc} : !pto.vmi.vreg<64xf32>, !pto.vmi.mask<64xpred> - -> !pto.vmi.vreg<1xf32> - - // Workaround: fan the scalar into 8 slots, then group-store the pad. - %sum8 = pto.vmi.vbrc %sum1 {group = 1} - : !pto.vmi.vreg<1xf32> -> !pto.vmi.vreg<8xf32> + -> !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/desired_vmi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.pto index 19ad32753a..c341694d91 100644 --- 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 @@ -6,20 +6,23 @@ // 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 directly (may not lower today). +// 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>, - %mask: !pto.vmi.mask<64xpred>, %dst: !pto.ptr, %off: index) { + %c64 = arith.constant 64 : index %c1 = arith.constant 1 : index - %sum = pto.vmi.vcadd %tile, %mask {reassoc} + %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], %c1 {group = 1} - : !pto.vmi.vreg<1xf32>, !pto.ptr + 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/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/target_mi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/target_mi.pto index c99e9cadcc..6281dc02e1 100644 --- 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 @@ -6,19 +6,21 @@ // 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 VPTO / pto.mi: vcadd + 1PT_B32 store, no pad-to-8. +// Desired MI: vcadd + 1PT_B32 store (ASC ONEPT-style), no pad-to-8. -module { +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) { - %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 + 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 index 68adfea0ab..34f02a7dff 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md @@ -1,19 +1,37 @@ -# 04 — Packed UE8M0 scale factors (`pack_factor=2`) +# 04 — Packed UE8M0 scale factors ## Algorithm need -Per-token MXFP8 quantization stores UE8M0 scale factors with **`pack_factor=2`**: -two consecutive `ui8` exponents share one `ui16` element in GM/UB. This halves -SF bandwidth versus an unpacked `ui8` tensor. +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 failure +## Working-today workaround -VMI/PTODSL paths that try to pack adjacent UE8M0 bytes (reinterpret / element -pack into `ui16`) do not lower reliably. Authors keep **unpacked `ui8` SF -only**, doubling scale-factor traffic on the memory pipe. +Store SF as unpacked `ui8` only (`buggy_vmi.pto`). 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). -## Desired PTOAS behavior +Reproduce: -Provide a first-class pack path so a logical `ui8` UE8M0 vector lowers through -`PK` (or equivalent) into packed `ui16` storage with `pack_factor=2`, and the -inverse unpack on load. +```bash +$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +First-class pack into `ui16` then store (see `desired_vmi.pto`): + +``` +error: custom op 'pto.vmi.vpack' is unknown +``` + +## Target MI + bisheng status + +`target_mi.pto` uses `pto.vsts {dist = "PK_B16"}` into `!pto.ptr`. + +| 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 | diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto index d8b8cc7c8c..1eb22fbaab 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto @@ -6,14 +6,13 @@ // 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. -// Compiling-today style: store UE8M0 scales as unpacked ui8 only. +// Working-today workaround: 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) { - // No pack_factor=2 path — write raw ui8 SF tensor. 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/desired_vmi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.pto index 987c410365..77aed4e86f 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.pto +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.pto @@ -6,14 +6,13 @@ // 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_factor=2 UE8M0 store into ui16 (may not lower today). +// 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) { - // Desired: element pack of adjacent UE8M0 bytes into ui16 words (pack_factor=2). %sf_u16 = pto.vmi.vpack %sf_u8 : !pto.vmi.vreg<64xui8> -> !pto.vmi.vreg<32xui16> pto.vmi.vstore %sf_u16, %dst[%off] 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/target_mi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.pto index 0a2e2d1038..34198771e7 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.pto +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/target_mi.pto @@ -6,17 +6,18 @@ // 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 VPTO / pto.mi: PK pack of UE8M0 bytes into ui16 storage. +// Desired MI: PK_B16 store of packed UE8M0 words. module { func.func @packed_ue8m0_sf_store_mi( - %sf_u8: !pto.vreg<256xui8>, + %sf_u16: !pto.vreg<128xui16>, %dst: !pto.ptr, %off: index) { - %m = pto.pset_b8 "PAT_ALL" : !pto.mask - // Pack adjacent ui8 lanes (pack_factor=2) then store as ui16 / PK path. - pto.vsts %sf_u8, %dst[%off], %m {dist = "PK_B16"} - : !pto.vreg<256xui8>, !pto.ptr, !pto.mask + 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 index 309dc54f09..a65839bff2 100644 --- a/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md @@ -1,20 +1,41 @@ -# 05 — Persistent multi-buffer stages and `block_k > 512` +# 05 — Persistent stages > 1 / `block_k` > 512 ## Algorithm need -Long-K quantization / cast-back vector kernels run as **Persistent** loops over -the K dimension. Authors want multi-buffer software pipeline stages (double / -triple buffer) and tile sizes with **`block_k > 512`** so MTE and vector pipes -stay overlapped on large hidden dimensions. +Long-K quant / cast-back 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. -## Current failure +## Working-today workaround -Practical PTOAS acceptance for these Persistent VMI modules is capped at -**`stages = 1`** and **`block_k ≤ 512`**. Raising either is rejected or -miscompiles, so kernels cannot express a legal multi-buffer Persistent schedule. +`block_k = 512`, single-stage Persistent tile (`buggy_vmi.pto`). Checked-in +dump: `lowered_vpto.pto` — chunked `pto.vlds` / `pto.vcvt` / `pto.vmul` / +`pto.vsts` over the 512-wide tile. -## Desired PTOAS behavior +Reproduce: -Accept Persistent vector modules with `stages >= 2` and `block_k > 512`, and -lower them to a correct multi-buffer VPTO schedule (distinct UB stage buffers + -pipeline sync) without forcing the single-stage / K≤512 workaround. +```bash +$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +`stages = 2` and `block_k = 1024` with dual UB buffers (`desired_vmi.pto`). +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` sketches a stages=2 / block_k=1024 module with UNPK load, +mul, and PK4 store inside `vecscope`. + +| 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 | diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto index 46f3957274..17a2988502 100644 --- a/docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto @@ -6,8 +6,7 @@ // 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. -// Compiling-today style Persistent vector module. -// Constraint comment: stages=1 / block_k<=512 required for legal lowering today. +// Working-today Persistent vector tile: stages=1, block_k=512. module attributes { pto.target_arch = "a5", 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 index 35c047ae12..da24d2cdb3 100644 --- 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 @@ -6,7 +6,7 @@ // 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, block_k=1024 (may not lower today). +// Idiomatic Persistent VMI: stages=2 and block_k=1024 (does not lower today). module attributes { pto.target_arch = "a5", 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/target_mi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/target_mi.pto index dabf41a47a..dc59b4c256 100644 --- 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 @@ -6,29 +6,29 @@ // 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 VPTO / pto.mi sketch: legal multi-buffer Persistent schedule. +// Desired MI sketch: multi-buffer Persistent schedule with larger K tile. 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>, %off: index) { - %m16 = pto.pset_b16 "PAT_ALL" : !pto.mask - %m32 = pto.pset_b32 "PAT_ALL" : !pto.mask - // Stage 0 buffer pair active this iteration (stage 1 overlaps MTE). - %x0 = pto.vlds %x_ub0[%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_ub0[%off], %m32 {dist = "PK4_B32"} - : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + pto.vecscope { + %m32 = pto.pset_b32 "PAT_ALL" : !pto.mask + %x0 = pto.vlds %x_ub0[%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_ub0[%off], %m32 {dist = "PK4_B32"} + : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + } 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..7c9e02d687 --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md @@ -0,0 +1,43 @@ +# 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`. + +## Working-today workaround + +L=256 `pto.vmi.vcvt` ui8→ui16 (`buggy_vmi.pto`). Checked-in dump: +`lowered_vpto.pto` — key ops `pto.vcvt ... {part = "EVEN"}` and `{part = "ODD"}`. + +Reproduce: + +```bash +$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto +``` + +## Desired VMI + current failure + +L=8 group-slot widen (`desired_vmi.pto`): + +``` +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 MI shape the L=256 path already reaches (`vcvt` EVEN + +ui16 store). Small-L VMI should lower here instead of leaving `extui`. + +| 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 | diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/buggy_vmi.pto b/docs/feature-gaps/vmi/06_small_l_ui8_widen/buggy_vmi.pto new file mode 100644 index 0000000000..bcb27cf792 --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/buggy_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. + +// Working-today: 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/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/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/target_mi.pto b/docs/feature-gaps/vmi/06_small_l_ui8_widen/target_mi.pto new file mode 100644 index 0000000000..bd6b9e91f3 --- /dev/null +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/target_mi.pto @@ -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. + +// Desired MI: ui8→ui16 widen via vcvt (EVEN/ODD), then dense ui16 store. +// Small-L VMI should lower to this shape instead of residual pto.vmi.extui. + +module { + func.func @ui8_to_ui16_widen_mi( + %input: !pto.vreg<256xui8>, + %dst: !pto.ptr, + %off: index) { + pto.vecscope { + %m = pto.pset_b8 "PAT_ALL" : !pto.mask + %even = pto.vcvt %input, %m {part = "EVEN"} + : !pto.vreg<256xui8>, !pto.mask -> !pto.vreg<128xui16> + %m16 = pto.pset_b16 "PAT_ALL" : !pto.mask + pto.vsts %even, %dst[%off], %m16 + : !pto.vreg<128xui16>, !pto.ptr, !pto.mask + } + return + } +} diff --git a/docs/feature-gaps/vmi/README.md b/docs/feature-gaps/vmi/README.md index 65bd0ccb00..5b3a1eebc6 100644 --- a/docs/feature-gaps/vmi/README.md +++ b/docs/feature-gaps/vmi/README.md @@ -1,28 +1,59 @@ -# VMI / PTODSL feature gaps +# VMI feature gaps (docs fixtures) -Minimal reproducers for PTOAS capabilities that VMI/PTODSL authors need when -writing high-throughput **quantization** and **residual-mix** style vector -kernels. Each subdirectory documents one gap with: +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 failure, desired PTOAS behavior | -| `buggy_vmi.pto` | Compiling-today style workaround IR (illustrative) | -| `desired_vmi.pto` | Idiomatic VMI using the missing feature (may not lower today) | -| `target_mi.pto` | Desired VPTO / `pto.mi` shape after lowering | - -These samples are documentation fixtures, not lit tests (no `// RUN:` lines). +| `README.md` | Algorithm need, workaround + reproduce, desired failure, target MI / bisheng status | +| `buggy_vmi.pto` | Working-today workaround that **must** lower with `pto-test-opt` + `$PASS` | +| `lowered_vpto.pto` | Checked-in dump from that lower | +| `desired_vmi.pto` | Idiomatic ask (may fail — README pastes the error) | +| `target_mi.pto` | Desired `pto.mi` / VPTO shape; exercised with `ptoas` + `bisheng` | ## Gaps -| # | Gap | Why it matters | +| # | Gap | Why ASC-level BW cares | |---|---|---| -| 01 | [Scalar GM dcache bypass](01_scalar_gm_dcache_bypass/) | Per-block quantization writes one f32 scale per 32×32 group to GM. Forcing bulk MTE for a few bytes burns bandwidth and sync. | -| 02 | [Compact inverse `vbrc`](02_compact_inv_vbrc/) | Quant kernels fan reciprocal scales across lanes. Padding to a full chunk before `dist_mode=brc` wastes UB and ops. | -| 03 | [Direct scalar reduce store](03_direct_scalar_reduce_store/) | Residual-mix kernels reduce a tile to one f32 scalar per group. Padding stores to 8 slots pollutes bandwidth and masks `1PT_B32`. | -| 04 | [Packed UE8M0 scale factors](04_packed_ue8m0_sf/) | Per-token MXFP8 uses UE8M0 scales with `pack_factor=2`. Unpacked `ui8` doubles SF traffic vs packed `ui16`. | -| 05 | [Persistent stages / `block_k`](05_persistent_stages_block_k/) | Persistent vector loops need multi-buffer stages and `block_k > 512` for long-K quant/cast tiles. Caps leave the pipe underfed. | - -Closing these gaps improves authoring productivity (fewer workarounds in -PTODSL) and end-to-end performance of quantization + residual-mix kernels -without changing their algorithm. +| 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 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) | + +Tip findings (high level, no external repo names): residual-mix still pays a +reduce pad-to-8 store tax; per-block Persistent stays on stages=1 / +`block_k≤512`; fused small-L UE8M0 widen hits the L=8 `extui` residual. + +## How to reproduce + +Tool paths used when these dumps were refreshed: + +```bash +PTO_TEST_OPT=/home/jzhuang/work_dir/PTOAS-vmi/build/tools/pto-test-opt/pto-test-opt +PTOAS=/home/jzhuang/work_dir/PTOAS-vmi/build/tools/ptoas/ptoas +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' + +# Working workaround (must succeed; refresh lowered_vpto.pto) +$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto + +# Desired (expect failure for most gaps; paste error into README) +$PTO_TEST_OPT desired_vmi.pto $PASS -o /tmp/desired.pto + +# Target MI +$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 `--vpto-emit-hivm-llvm`). From 0778d50f223155e2225555ead1e2845d393d0632 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Sat, 25 Jul 2026 00:01:13 +0800 Subject: [PATCH 3/9] docs: use relative build paths in VMI feature-gap reproduce notes --- docs/feature-gaps/vmi/README.md | 28 ++++++++++++++++------------ 1 file changed, 16 insertions(+), 12 deletions(-) diff --git a/docs/feature-gaps/vmi/README.md b/docs/feature-gaps/vmi/README.md index 5b3a1eebc6..343579140a 100644 --- a/docs/feature-gaps/vmi/README.md +++ b/docs/feature-gaps/vmi/README.md @@ -31,29 +31,33 @@ reduce pad-to-8 store tax; per-block Persistent stays on stages=1 / ## How to reproduce -Tool paths used when these dumps were refreshed: +Assumes a local `build/` of this tree (or any install with `pto-test-opt` / +`ptoas` on `PATH`) and Ascend `bisheng` available: ```bash -PTO_TEST_OPT=/home/jzhuang/work_dir/PTOAS-vmi/build/tools/pto-test-opt/pto-test-opt -PTOAS=/home/jzhuang/work_dir/PTOAS-vmi/build/tools/ptoas/ptoas -BISHENG=/usr/local/Ascend/cann-9.0.0/tools/bisheng_compiler/bin/bisheng -export PATH="$(dirname $BISHENG):$PATH" +# 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/ # Working workaround (must succeed; refresh lowered_vpto.pto) -$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto +"$PTO_TEST_OPT" buggy_vmi.pto $PASS -o lowered_vpto.pto -# Desired (expect failure for most gaps; paste error into README) -$PTO_TEST_OPT desired_vmi.pto $PASS -o /tmp/desired.pto +# Desired (often fails; paste error into README) +"$PTO_TEST_OPT" desired_vmi.pto $PASS -o /tmp/desired.pto -# Target MI -$PTOAS --pto-arch=a5 --pto-backend=vpto --pto-level=level3 \ +# 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 \ +"$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 `--vpto-emit-hivm-llvm`). +Use `--emit-vpto-llvm-ir` for LLVM IR (not the older `--vpto-emit-hivm-llvm` name). From e1e1b76e7e934bf3a6bb139521ae0f0243cdc199 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Sat, 25 Jul 2026 00:01:50 +0800 Subject: [PATCH 4/9] docs: rename buggy_vmi to current_slow_vmi in feature gaps Name the compiling-today path for what it is: correct but slow. --- .../feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md | 4 ++-- .../{buggy_vmi.pto => current_slow_vmi.pto} | 2 +- docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md | 4 ++-- .../{buggy_vmi.pto => current_slow_vmi.pto} | 2 +- .../vmi/03_direct_scalar_reduce_store/README.md | 6 +++--- .../{buggy_vmi.pto => current_slow_vmi.pto} | 2 +- docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md | 6 +++--- .../{buggy_vmi.pto => current_slow_vmi.pto} | 2 +- .../vmi/05_persistent_stages_block_k/README.md | 6 +++--- .../{buggy_vmi.pto => current_slow_vmi.pto} | 0 docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md | 6 +++--- .../{buggy_vmi.pto => current_slow_vmi.pto} | 0 docs/feature-gaps/vmi/README.md | 8 ++++---- 13 files changed, 24 insertions(+), 24 deletions(-) rename docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/{buggy_vmi.pto => current_slow_vmi.pto} (93%) rename docs/feature-gaps/vmi/02_compact_inv_vbrc/{buggy_vmi.pto => current_slow_vmi.pto} (94%) rename docs/feature-gaps/vmi/03_direct_scalar_reduce_store/{buggy_vmi.pto => current_slow_vmi.pto} (93%) rename docs/feature-gaps/vmi/04_packed_ue8m0_sf/{buggy_vmi.pto => current_slow_vmi.pto} (93%) rename docs/feature-gaps/vmi/05_persistent_stages_block_k/{buggy_vmi.pto => current_slow_vmi.pto} (100%) rename docs/feature-gaps/vmi/06_small_l_ui8_widen/{buggy_vmi.pto => current_slow_vmi.pto} (100%) 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 index c25b68546a..01162f9e15 100644 --- a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md @@ -7,7 +7,7 @@ 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. -## Working-today workaround +## 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 @@ -18,7 +18,7 @@ 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 buggy_vmi.pto $PASS -o lowered_vpto.pto +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto ``` ## Desired VMI + current failure diff --git a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.pto similarity index 93% rename from docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto rename to docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.pto index 8ea7f4d7cd..c154a887bf 100644 --- a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.pto @@ -6,7 +6,7 @@ // 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. -// Working-today workaround: SF already staged in UB; bulk MTE to GM for a tiny payload. +// 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( diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md b/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md index 1163918cb6..46d90051cf 100644 --- a/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md @@ -7,7 +7,7 @@ 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. -## Working-today workaround +## Current slow path Materialize recip scales as `#pto.vmi.layout`, `group_store` them, then reload with `dist_mode = "brc"`. Checked-in dump: @@ -17,7 +17,7 @@ Materialize recip scales as `#pto.vmi.layout`, Reproduce: ```bash -$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto ``` ## Desired VMI + current failure diff --git a/docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto b/docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.pto similarity index 94% rename from docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto rename to docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.pto index e6badf4206..29516a3799 100644 --- a/docs/feature-gaps/vmi/02_compact_inv_vbrc/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.pto @@ -6,7 +6,7 @@ // 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. -// Working-today workaround: pad recip scales to slots=8, store, then BRC reload. +// 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( 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 index b87ded3b39..61862afc21 100644 --- a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md @@ -6,9 +6,9 @@ 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. -## Working-today workaround +## Current slow path -Group-reduce with `{group = 8}` and store the 8-slot result (`buggy_vmi.pto`). +Group-reduce with `{group = 8}` and store the 8-slot result (`current_slow_vmi.pto`). 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`. @@ -16,7 +16,7 @@ key ops `pto.vcgadd` + `pto.vsts` with `PAT_VL8`. Reproduce: ```bash -$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto ``` ## Desired VMI + current failure diff --git a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.pto similarity index 93% rename from docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto rename to docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.pto index 43754189d2..782e54a785 100644 --- a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.pto @@ -6,7 +6,7 @@ // 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. -// Working-today workaround: group-reduce into an 8-slot pad, then group-store. +// 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( diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md index 34f02a7dff..3500e715e6 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md @@ -6,16 +6,16 @@ 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. -## Working-today workaround +## Current slow path -Store SF as unpacked `ui8` only (`buggy_vmi.pto`). Checked-in dump: +Store SF as unpacked `ui8` only (`current_slow_vmi.pto`). 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 buggy_vmi.pto $PASS -o lowered_vpto.pto +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto ``` ## Desired VMI + current failure diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.pto similarity index 93% rename from docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto rename to docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.pto index 1eb22fbaab..207b79eb34 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/buggy_vmi.pto +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.pto @@ -6,7 +6,7 @@ // 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. -// Working-today workaround: store UE8M0 scales as unpacked ui8 only. +// 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( 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 index a65839bff2..94ed6d98e0 100644 --- a/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md @@ -7,16 +7,16 @@ 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. -## Working-today workaround +## Current slow path -`block_k = 512`, single-stage Persistent tile (`buggy_vmi.pto`). Checked-in +`block_k = 512`, single-stage Persistent tile (`current_slow_vmi.pto`). 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 buggy_vmi.pto $PASS -o lowered_vpto.pto +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto ``` ## Desired VMI + current failure diff --git a/docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto b/docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.pto similarity index 100% rename from docs/feature-gaps/vmi/05_persistent_stages_block_k/buggy_vmi.pto rename to docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.pto 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 index 7c9e02d687..da8015ddb0 100644 --- a/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md @@ -10,15 +10,15 @@ compose / larger tiles. Reference for the working L=256 shape: `test/lit/vmi_new/vmi_to_vpto_integer_casts.pto`. -## Working-today workaround +## Current slow path -L=256 `pto.vmi.vcvt` ui8→ui16 (`buggy_vmi.pto`). Checked-in dump: +L=256 `pto.vmi.vcvt` ui8→ui16 (`current_slow_vmi.pto`). Checked-in dump: `lowered_vpto.pto` — key ops `pto.vcvt ... {part = "EVEN"}` and `{part = "ODD"}`. Reproduce: ```bash -$PTO_TEST_OPT buggy_vmi.pto $PASS -o lowered_vpto.pto +$PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto ``` ## Desired VMI + current failure diff --git a/docs/feature-gaps/vmi/06_small_l_ui8_widen/buggy_vmi.pto b/docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.pto similarity index 100% rename from docs/feature-gaps/vmi/06_small_l_ui8_widen/buggy_vmi.pto rename to docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.pto diff --git a/docs/feature-gaps/vmi/README.md b/docs/feature-gaps/vmi/README.md index 343579140a..1c1ab212b5 100644 --- a/docs/feature-gaps/vmi/README.md +++ b/docs/feature-gaps/vmi/README.md @@ -8,8 +8,8 @@ Each subdirectory has: | File | Role | |---|---| -| `README.md` | Algorithm need, workaround + reproduce, desired failure, target MI / bisheng status | -| `buggy_vmi.pto` | Working-today workaround that **must** lower with `pto-test-opt` + `$PASS` | +| `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` | | `lowered_vpto.pto` | Checked-in dump from that lower | | `desired_vmi.pto` | Idiomatic ask (may fail — README pastes the error) | | `target_mi.pto` | Desired `pto.mi` / VPTO shape; exercised with `ptoas` + `bisheng` | @@ -45,8 +45,8 @@ 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/ -# Working workaround (must succeed; refresh lowered_vpto.pto) -"$PTO_TEST_OPT" buggy_vmi.pto $PASS -o lowered_vpto.pto +# 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 From f50526e46ab558ad5b5a77e2266117097845b0e6 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Sat, 25 Jul 2026 00:01:55 +0800 Subject: [PATCH 5/9] docs: finish current_slow_vmi comment wording --- docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md | 2 +- .../vmi/05_persistent_stages_block_k/current_slow_vmi.pto | 2 +- docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.pto | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) 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 index 61862afc21..9a9bc2c874 100644 --- a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md @@ -26,7 +26,7 @@ Idiomatic mask1 store after `vcadd` (mirrors **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 workaround emits. +padded slow path emits. ## Target MI + bisheng status 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 index 17a2988502..4494a4cce1 100644 --- 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 @@ -6,7 +6,7 @@ // 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. -// Working-today Persistent vector tile: stages=1, block_k=512. +// Current slow path: Persistent vector tile: stages=1, block_k=512. module attributes { pto.target_arch = "a5", 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 index bcb27cf792..e32826e4cd 100644 --- 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 @@ -6,7 +6,7 @@ // 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. -// Working-today: L=256 ui8→ui16 vcvt (see test/lit/vmi_new/vmi_to_vpto_integer_casts.pto). +// 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( From 7d3ea333c36ef9c2bd1112e84a6d53976cd64dcc Mon Sep 17 00:00:00 2001 From: learning-chip Date: Sat, 25 Jul 2026 00:14:41 +0800 Subject: [PATCH 6/9] docs: add bf16 UNPK/PK fuse gap for residual-mix fwd BW Cover MHC-style continuous bf16 strips (EVEN-only UNPK today; dist_mode=unpack illegal). Note cast-back as a 05 consumer. --- .../05_persistent_stages_block_k/README.md | 15 ++++-- .../vmi/07_bf16_unpk_pk_fuse/README.md | 49 +++++++++++++++++++ .../07_bf16_unpk_pk_fuse/current_slow_vmi.pto | 33 +++++++++++++ .../vmi/07_bf16_unpk_pk_fuse/desired_vmi.pto | 33 +++++++++++++ .../vmi/07_bf16_unpk_pk_fuse/lowered_vpto.pto | 17 +++++++ .../vmi/07_bf16_unpk_pk_fuse/target_mi.pto | 46 +++++++++++++++++ docs/feature-gaps/vmi/README.md | 10 ++-- 7 files changed, 195 insertions(+), 8 deletions(-) create mode 100644 docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/README.md create mode 100644 docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/current_slow_vmi.pto create mode 100644 docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/desired_vmi.pto create mode 100644 docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/lowered_vpto.pto create mode 100644 docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/target_mi.pto 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 index 94ed6d98e0..42872c2d37 100644 --- a/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md @@ -2,10 +2,17 @@ ## Algorithm need -Long-K quant / cast-back 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. +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 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..e0a4c0781a --- /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`). 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`) 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/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/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 index 1c1ab212b5..86198ae057 100644 --- a/docs/feature-gaps/vmi/README.md +++ b/docs/feature-gaps/vmi/README.md @@ -22,12 +22,14 @@ Each subdirectory has: | 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 capped at stages=1 / `block_k≤512`; stages=2 faults in practice | +| 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 still pays a -reduce pad-to-8 store tax; per-block Persistent stays on stages=1 / -`block_k≤512`; fused small-L UE8M0 widen hits the L=8 `extui` residual. +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` (**05**); fused small-L UE8M0 widen hits L=8 `extui` (**06**). ## How to reproduce From edf90a7db4419884594f86a3c0a09b857997fc39 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Sat, 25 Jul 2026 00:32:28 +0800 Subject: [PATCH 7/9] docs: add ASC/CCE bisheng baselines for VMI gaps 03-06 Add minimal reference_asc_cce.asc kernels that compile on the Ascend CCE path while target_mi HIVM still crashes on bisheng. --- .../03_direct_scalar_reduce_store/README.md | 19 ++++++ .../reference_asc_cce.asc | 20 +++++++ .../vmi/04_packed_ue8m0_sf/README.md | 19 ++++++ .../04_packed_ue8m0_sf/reference_asc_cce.asc | 26 +++++++++ .../05_persistent_stages_block_k/README.md | 20 +++++++ .../reference_asc_cce.asc | 58 +++++++++++++++++++ .../vmi/06_small_l_ui8_widen/README.md | 19 ++++++ .../reference_asc_cce.asc | 21 +++++++ docs/feature-gaps/vmi/README.md | 20 +++++++ 9 files changed, 222 insertions(+) create mode 100644 docs/feature-gaps/vmi/03_direct_scalar_reduce_store/reference_asc_cce.asc create mode 100644 docs/feature-gaps/vmi/04_packed_ue8m0_sf/reference_asc_cce.asc create mode 100644 docs/feature-gaps/vmi/05_persistent_stages_block_k/reference_asc_cce.asc create mode 100644 docs/feature-gaps/vmi/06_small_l_ui8_widen/reference_asc_cce.asc 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 index 9a9bc2c874..85520bf84e 100644 --- a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md @@ -37,3 +37,22 @@ padded slow path emits. | `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/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/04_packed_ue8m0_sf/README.md b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md index 3500e715e6..8a9babcaac 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md @@ -35,3 +35,22 @@ error: custom op 'pto.vmi.vpack' is unknown | `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/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/05_persistent_stages_block_k/README.md b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md index 42872c2d37..95e9d579a0 100644 --- a/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md @@ -46,3 +46,23 @@ mul, and PK4 store inside `vecscope`. | `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` is a stages=2 ping-pong shell (`SetFlag`/`WaitFlag` on +slots `w & 1`) with `block_k=1024` and a VF body of `UNPK_B16` / `vmul` / +`PK_B32`. It compiles to a non-empty `.o` with `bisheng --cce-aicore-only`, +proving multi-stage Persistent 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_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/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/06_small_l_ui8_widen/README.md b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md index da8015ddb0..74a351d691 100644 --- a/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md @@ -41,3 +41,22 @@ ui16 store). Small-L VMI should lower here instead of leaving `extui`. | `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` widens `ui8→ui16` under `PAT_VL8` with `vcvt` / +`PART_EVEN` then stores. It compiles to a non-empty `.o` with +`bisheng --cce-aicore-only`, proving the small-L widen 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_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/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/README.md b/docs/feature-gaps/vmi/README.md index 86198ae057..648907ec1b 100644 --- a/docs/feature-gaps/vmi/README.md +++ b/docs/feature-gaps/vmi/README.md @@ -13,6 +13,7 @@ Each subdirectory has: | `lowered_vpto.pto` | Checked-in dump from that lower | | `desired_vmi.pto` | Idiomatic ask (may fail — README pastes the error) | | `target_mi.pto` | Desired `pto.mi` / VPTO shape; exercised with `ptoas` + `bisheng` | +| `reference_asc_cce.asc` | Gaps **03–06** only: minimal ASC/CCE kernel that `bisheng --cce-aicore-only` accepts (e2e proof while `target_mi` HIVM still crashes) | ## Gaps @@ -63,3 +64,22 @@ cd docs/feature-gaps/vmi/ ``` Use `--emit-vpto-llvm-ir` for LLVM IR (not the older `--vpto-emit-hivm-llvm` name). + +## ASC/CCE reference compile (gaps 03–06) + +When `target_mi` → HIVM IR → `bisheng -x ir` crashes, the sibling +`reference_asc_cce.asc` proves the same instruction shape is legal on the +Ascend CCE path: + +```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 +``` From f43147ba530da198e1de86fb79fc58ead1bba142 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Sat, 25 Jul 2026 00:43:29 +0800 Subject: [PATCH 8/9] docs: expand gap 05 Persistent MI schedule; fix gap 06 small-L MI Express dual-buffer ping-pong and pipe-event sync in target_mi for gap 05 (still fails HIVM/bisheng). Retarget gap 06 MI to PAT_VL8; that shape now builds with bisheng while VMI L=8 legalization remains the gap. --- .../03_direct_scalar_reduce_store/README.md | 3 +- .../vmi/04_packed_ue8m0_sf/README.md | 3 +- .../05_persistent_stages_block_k/README.md | 23 +++--- .../target_mi.pto | 72 ++++++++++++++++--- .../vmi/06_small_l_ui8_widen/README.md | 13 ++-- .../vmi/06_small_l_ui8_widen/target_mi.pto | 20 +++--- docs/feature-gaps/vmi/README.md | 13 ++-- 7 files changed, 106 insertions(+), 41 deletions(-) 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 index 85520bf84e..9803b35beb 100644 --- a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md @@ -30,7 +30,8 @@ padded slow path emits. ## Target MI + bisheng status -`target_mi.pto` asks for `pto.vsts {dist = "1PT_B32"}` after `pto.vcadd`. +`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 | |---|---| diff --git a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md index 8a9babcaac..dea4b74b25 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md @@ -28,7 +28,8 @@ error: custom op 'pto.vmi.vpack' is unknown ## Target MI + bisheng status -`target_mi.pto` uses `pto.vsts {dist = "PK_B16"}` into `!pto.ptr`. +`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 | |---|---| 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 index 95e9d579a0..1f88c85b06 100644 --- a/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md @@ -38,22 +38,27 @@ pto.vmi.ensure_layout cannot materialize this conversion ## Target MI + bisheng status -`target_mi.pto` sketches a stages=2 / block_k=1024 module with UNPK load, -mul, and PK4 store inside `vecscope`. +`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 | +| `ptoas ... --emit-vpto` | OK (schedule + flags preserved) | | `ptoas ... --emit-vpto-llvm-ir` | OK | -| `bisheng ... -c -x ir` → `.o` | **FAIL** — bisheng crashes (exit 139) on the emitted HIVM IR | +| `bisheng ... -c -x ir` → `.o` | **FAIL** — bisheng crashes on the emitted HIVM IR | ## ASC/CCE working baseline -`reference_asc_cce.asc` is a stages=2 ping-pong shell (`SetFlag`/`WaitFlag` on -slots `w & 1`) with `block_k=1024` and a VF body of `UNPK_B16` / `vmul` / -`PK_B32`. It compiles to a non-empty `.o` with `bisheng --cce-aicore-only`, -proving multi-stage Persistent is legal on the Ascend CCE path while -`target_mi` HIVM still crashes. +`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} 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 index dc59b4c256..39407b79f1 100644 --- 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 @@ -6,7 +6,9 @@ // 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 sketch: multi-buffer Persistent schedule with larger K tile. +// 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", @@ -16,19 +18,67 @@ module attributes { } { 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>, - %off: index) { - pto.vecscope { - %m32 = pto.pset_b32 "PAT_ALL" : !pto.mask - %x0 = pto.vlds %x_ub0[%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_ub0[%off], %m32 {dist = "PK4_B32"} - : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + %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 index 74a351d691..25455bcb47 100644 --- a/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md @@ -33,21 +33,24 @@ error: VMI-RESIDUAL-OP: failed to convert all VMI ops/types to VPTO ## Target MI + bisheng status -`target_mi.pto` is the MI shape the L=256 path already reaches (`vcvt` EVEN + -ui16 store). Small-L VMI should lower here instead of leaving `extui`. +`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` | **FAIL** — bisheng crashes (exit 139) on the emitted HIVM IR | +| `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`, proving the small-L widen shape is legal on the -Ascend CCE path while `target_mi` HIVM still crashes. +`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} 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 index bd6b9e91f3..22b9e16a05 100644 --- 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 @@ -6,19 +6,21 @@ // 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: ui8→ui16 widen via vcvt (EVEN/ODD), then dense ui16 store. -// Small-L VMI should lower to this shape instead of residual pto.vmi.extui. +// 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 { - func.func @ui8_to_ui16_widen_mi( - %input: !pto.vreg<256xui8>, +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) { + %off: index) attributes {pto.kernel} { pto.vecscope { - %m = pto.pset_b8 "PAT_ALL" : !pto.mask - %even = pto.vcvt %input, %m {part = "EVEN"} + %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_ALL" : !pto.mask + %m16 = pto.pset_b16 "PAT_VL8" : !pto.mask pto.vsts %even, %dst[%off], %m16 : !pto.vreg<128xui16>, !pto.ptr, !pto.mask } diff --git a/docs/feature-gaps/vmi/README.md b/docs/feature-gaps/vmi/README.md index 648907ec1b..5b43c1f6c8 100644 --- a/docs/feature-gaps/vmi/README.md +++ b/docs/feature-gaps/vmi/README.md @@ -13,7 +13,7 @@ Each subdirectory has: | `lowered_vpto.pto` | Checked-in dump from that lower | | `desired_vmi.pto` | Idiomatic ask (may fail — README pastes the error) | | `target_mi.pto` | Desired `pto.mi` / VPTO shape; exercised with `ptoas` + `bisheng` | -| `reference_asc_cce.asc` | Gaps **03–06** only: minimal ASC/CCE kernel that `bisheng --cce-aicore-only` accepts (e2e proof while `target_mi` HIVM still crashes) | +| `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**) | ## Gaps @@ -30,7 +30,9 @@ Each subdirectory has: 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` (**05**); fused small-L UE8M0 widen hits L=8 `extui` (**06**). +`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 @@ -67,9 +69,10 @@ Use `--emit-vpto-llvm-ir` for LLVM IR (not the older `--vpto-emit-hivm-llvm` nam ## ASC/CCE reference compile (gaps 03–06) -When `target_mi` → HIVM IR → `bisheng -x ir` crashes, the sibling -`reference_asc_cce.asc` proves the same instruction shape is legal on the -Ascend CCE path: +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} From 5771bd3b14a78f5b1dd131e65e7f5447db86b2a0 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Sat, 25 Jul 2026 00:47:55 +0800 Subject: [PATCH 9/9] docs: add PTODSL Python forms for VMI feature-gap fixtures Pair current_slow_vmi.py / desired_vmi.py with each gap IR so the VMI authoring surface is documented alongside the .pto reproducers. --- .../vmi/01_scalar_gm_dcache_bypass/README.md | 2 +- .../current_slow_vmi.py | 25 ++++++++++ .../01_scalar_gm_dcache_bypass/desired_vmi.py | 34 +++++++++++++ .../vmi/02_compact_inv_vbrc/README.md | 2 +- .../02_compact_inv_vbrc/current_slow_vmi.py | 34 +++++++++++++ .../vmi/02_compact_inv_vbrc/desired_vmi.py | 27 +++++++++++ .../03_direct_scalar_reduce_store/README.md | 4 +- .../current_slow_vmi.py | 35 ++++++++++++++ .../desired_vmi.py | 31 ++++++++++++ .../vmi/04_packed_ue8m0_sf/README.md | 4 +- .../04_packed_ue8m0_sf/current_slow_vmi.py | 26 ++++++++++ .../vmi/04_packed_ue8m0_sf/desired_vmi.py | 32 +++++++++++++ .../05_persistent_stages_block_k/README.md | 4 +- .../current_slow_vmi.py | 36 ++++++++++++++ .../desired_vmi.py | 48 +++++++++++++++++++ .../vmi/06_small_l_ui8_widen/README.md | 4 +- .../06_small_l_ui8_widen/current_slow_vmi.py | 27 +++++++++++ .../vmi/06_small_l_ui8_widen/desired_vmi.py | 31 ++++++++++++ .../vmi/07_bf16_unpk_pk_fuse/README.md | 4 +- .../07_bf16_unpk_pk_fuse/current_slow_vmi.py | 31 ++++++++++++ .../vmi/07_bf16_unpk_pk_fuse/desired_vmi.py | 37 ++++++++++++++ docs/feature-gaps/vmi/README.md | 12 +++++ 22 files changed, 478 insertions(+), 12 deletions(-) create mode 100644 docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/current_slow_vmi.py create mode 100644 docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/desired_vmi.py create mode 100644 docs/feature-gaps/vmi/02_compact_inv_vbrc/current_slow_vmi.py create mode 100644 docs/feature-gaps/vmi/02_compact_inv_vbrc/desired_vmi.py create mode 100644 docs/feature-gaps/vmi/03_direct_scalar_reduce_store/current_slow_vmi.py create mode 100644 docs/feature-gaps/vmi/03_direct_scalar_reduce_store/desired_vmi.py create mode 100644 docs/feature-gaps/vmi/04_packed_ue8m0_sf/current_slow_vmi.py create mode 100644 docs/feature-gaps/vmi/04_packed_ue8m0_sf/desired_vmi.py create mode 100644 docs/feature-gaps/vmi/05_persistent_stages_block_k/current_slow_vmi.py create mode 100644 docs/feature-gaps/vmi/05_persistent_stages_block_k/desired_vmi.py create mode 100644 docs/feature-gaps/vmi/06_small_l_ui8_widen/current_slow_vmi.py create mode 100644 docs/feature-gaps/vmi/06_small_l_ui8_widen/desired_vmi.py create mode 100644 docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/current_slow_vmi.py create mode 100644 docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/desired_vmi.py 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 index 01162f9e15..b330173875 100644 --- a/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md +++ b/docs/feature-gaps/vmi/01_scalar_gm_dcache_bypass/README.md @@ -24,7 +24,7 @@ $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`): +`dcache_bypass` (see `desired_vmi.pto` / `desired_vmi.py`): ``` error: 'pto.vmi.group_store' op requires memory destination to be UB-backed 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.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/02_compact_inv_vbrc/README.md b/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md index 46d90051cf..3c7718c25a 100644 --- a/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md +++ b/docs/feature-gaps/vmi/02_compact_inv_vbrc/README.md @@ -22,7 +22,7 @@ $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`): +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: 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.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/03_direct_scalar_reduce_store/README.md b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md index 9803b35beb..3d646037f6 100644 --- a/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md +++ b/docs/feature-gaps/vmi/03_direct_scalar_reduce_store/README.md @@ -8,7 +8,7 @@ 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`). +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`. @@ -22,7 +22,7 @@ $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`). +`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 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.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/04_packed_ue8m0_sf/README.md b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md index dea4b74b25..1522961703 100644 --- a/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md +++ b/docs/feature-gaps/vmi/04_packed_ue8m0_sf/README.md @@ -8,7 +8,7 @@ packed storage and blocks ASC-level BW on that pipe. ## Current slow path -Store SF as unpacked `ui8` only (`current_slow_vmi.pto`). Checked-in dump: +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). @@ -20,7 +20,7 @@ $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`): +First-class pack into `ui16` then store (see `desired_vmi.pto` / `desired_vmi.py`): ``` error: custom op 'pto.vmi.vpack' is unknown 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.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/05_persistent_stages_block_k/README.md b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md index 1f88c85b06..82703bc767 100644 --- a/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md +++ b/docs/feature-gaps/vmi/05_persistent_stages_block_k/README.md @@ -16,7 +16,7 @@ Consumers today: ## Current slow path -`block_k = 512`, single-stage Persistent tile (`current_slow_vmi.pto`). Checked-in +`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. @@ -28,7 +28,7 @@ $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`). +`stages = 2` and `block_k = 1024` with dual UB buffers (`desired_vmi.pto` / `desired_vmi.py`). Layout assignment rejects the 1024-wide multiply: ``` 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.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/06_small_l_ui8_widen/README.md b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md index 25455bcb47..c54b436f89 100644 --- a/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md +++ b/docs/feature-gaps/vmi/06_small_l_ui8_widen/README.md @@ -12,7 +12,7 @@ Reference for the working L=256 shape: ## Current slow path -L=256 `pto.vmi.vcvt` ui8→ui16 (`current_slow_vmi.pto`). Checked-in dump: +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: @@ -23,7 +23,7 @@ $PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto ## Desired VMI + current failure -L=8 group-slot widen (`desired_vmi.pto`): +L=8 group-slot widen (`desired_vmi.pto` / `desired_vmi.py`): ``` error: failed to legalize operation 'pto.vmi.extui' that was explicitly marked illegal 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.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/07_bf16_unpk_pk_fuse/README.md b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/README.md index e0a4c0781a..05570b3ee7 100644 --- a/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/README.md +++ b/docs/feature-gaps/vmi/07_bf16_unpk_pk_fuse/README.md @@ -10,7 +10,7 @@ must lower to the same full unpack/pack form for ASC-level BW. ## Current slow path -Continuous L=64 strip (`current_slow_vmi.pto`). Checked-in dump: +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 @@ -24,7 +24,7 @@ $PTO_TEST_OPT current_slow_vmi.pto $PASS -o lowered_vpto.pto ## Desired VMI + current failure -Explicit unpack load (`desired_vmi.pto`) so authors can request UNPK without +Explicit unpack load (`desired_vmi.pto` / `desired_vmi.py`) so authors can request UNPK without relying on continuous-load heuristics: ``` 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.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/README.md b/docs/feature-gaps/vmi/README.md index 5b43c1f6c8..aef20bd4a0 100644 --- a/docs/feature-gaps/vmi/README.md +++ b/docs/feature-gaps/vmi/README.md @@ -10,11 +10,23 @@ Each subdirectory has: |---|---| | `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 |