From 8b84f9f6b08126b8856aaa990a29855376f6e230 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Mon, 27 Jul 2026 15:09:36 +0000 Subject: [PATCH 1/3] Add repro for VMI Persistent stages vs CCE ping-pong --- docs/repro/pipelining/.gitignore | 10 + docs/repro/pipelining/README.md | 83 ++++++++ .../pipelining/fixtures/current_slow_vmi.pto | 47 +++++ .../pipelining/fixtures/current_slow_vmi.py | 36 ++++ .../repro/pipelining/fixtures/desired_vmi.pto | 62 ++++++ docs/repro/pipelining/fixtures/desired_vmi.py | 48 +++++ .../pipelining/fixtures/lowered_vpto.pto | 100 +++++++++ .../pipelining/fixtures/reference_asc_cce.asc | 58 ++++++ docs/repro/pipelining/fixtures/target_mi.pto | 84 ++++++++ docs/repro/pipelining/scripts/check_stages.sh | 194 ++++++++++++++++++ docs/repro/pipelining/scripts/env.sh | 23 +++ 11 files changed, 745 insertions(+) create mode 100644 docs/repro/pipelining/.gitignore create mode 100644 docs/repro/pipelining/README.md create mode 100644 docs/repro/pipelining/fixtures/current_slow_vmi.pto create mode 100644 docs/repro/pipelining/fixtures/current_slow_vmi.py create mode 100644 docs/repro/pipelining/fixtures/desired_vmi.pto create mode 100644 docs/repro/pipelining/fixtures/desired_vmi.py create mode 100644 docs/repro/pipelining/fixtures/lowered_vpto.pto create mode 100644 docs/repro/pipelining/fixtures/reference_asc_cce.asc create mode 100644 docs/repro/pipelining/fixtures/target_mi.pto create mode 100755 docs/repro/pipelining/scripts/check_stages.sh create mode 100755 docs/repro/pipelining/scripts/env.sh diff --git a/docs/repro/pipelining/.gitignore b/docs/repro/pipelining/.gitignore new file mode 100644 index 0000000000..a95b67ec69 --- /dev/null +++ b/docs/repro/pipelining/.gitignore @@ -0,0 +1,10 @@ +build/ +sim_outputs/ +*.so +*.o +log_ca/ +instr.bin +__pycache__/ +*.pyc +.pytest_cache/ +*.egg-info/ diff --git a/docs/repro/pipelining/README.md b/docs/repro/pipelining/README.md new file mode 100644 index 0000000000..53912f0160 --- /dev/null +++ b/docs/repro/pipelining/README.md @@ -0,0 +1,83 @@ +# pipelining — VMI Persistent stages vs CCE ping-pong repro + +Self-contained compile repro for PTOAS gap 05: 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. Ascend CCE can express stages=2 ping-pong +with `block_k=1024`; VMI on tag **vmi-v0.1.3** cannot reach the same schedule +— layout assignment rejects wide tiles, and the hand-written MI path crashes +bisheng even after ptoas emits LLVM IR. + +Fixtures are copied from `docs/feature-gaps/vmi/05_persistent_stages_block_k/`. + +## What breaks + +| Path | stages | block_k | Fault mode | +|------|-------:|--------:|------------| +| CCE (`fixtures/reference_asc_cce.asc`) | 2 | 1024 | **OK** — bisheng `--cce-aicore-only` produces a non-empty `.o` | +| VMI desired (`fixtures/desired_vmi.*`) | 2 | 1024 | **Layout reject** — `pto.vmi.mulf` on 1024-wide tile: deinterleaved=2 vs contiguous | +| VMI target MI (`fixtures/target_mi.pto`) | 2 | 1024 | **Bisheng crash** — ptoas `--emit-vpto` and LLVM IR OK; object compile exit 70 | +| VMI current slow (`fixtures/current_slow_vmi.*`) | 1 | 512 | **Frontend OK** — ptodsl emits MLIR; idiomatic VMI `.pto` hits same layout issue | +| VMI lowered MI (`fixtures/lowered_vpto.pto`) | 1 | 512 | **OK to LLVM IR** — chunked 128-wide `vlds`/`vmul`/`vsts` path lowers | + +Run `./scripts/check_stages.sh` to refresh the table below. + +## Check results (vmi-v0.1.3, 2026-07-27) + +Measured by `./scripts/check_stages.sh`: + +| Step | Result | +|------|--------| +| `reference_asc_cce.asc` → bisheng `--cce-aicore-only` | **PASS** | +| `desired_vmi.py` → MLIR frontend | **PASS** (frontend only) | +| `desired_vmi.py` MLIR → LLVM IR | **FAIL** — layout reject on 1024-wide vmul | +| `desired_vmi.pto` → `--emit-vpto` | **FAIL** — same layout reject | +| `target_mi.pto` → `--emit-vpto` | **PASS** | +| `target_mi.pto` → LLVM IR → bisheng `.o` | **FAIL** — bisheng crash (exit 70) | +| `current_slow_vmi.py` → MLIR frontend | **PASS** (frontend only) | +| `lowered_vpto.pto` → LLVM IR | **PASS** — stages=1 chunked MI path | + +Full log: `sim_outputs/check_stages/compile_results.txt` + +## Layout + +``` +pipelining/ + fixtures/ IR dumps and reference ASC from gap 05 + scripts/ env.sh, check_stages.sh +``` + +CCE stages=2 baseline lives in `fixtures/reference_asc_cce.asc` (compile-only +is enough for this repro; on-device overlap profiling is TODO below). + +## Prerequisites + +- CANN 9.0 (`bisheng`) +- Built PTOAS at tag **vmi-v0.1.3** (branch `zjw/pipelining_issue`) + +Push remote hint: https://github.com/learning-chip/PTOAS.git + +## Run + +From this directory (`docs/repro/pipelining/`): + +```bash +source scripts/env.sh +./scripts/check_stages.sh +``` + +## Pin + +- PTOAS tag: **vmi-v0.1.3** +- Branch: `zjw/pipelining_issue` +- `PTOAS_ROOT` defaults to the repo root (`../../..` from this repro directory). + +## TODO + +- On-device run and **msopprof** bandwidth study for stages=2 CCE vs stages=1 + VMI on large hidden sizes. +- Guide: https://gitcode.com/Ascend/msopprof/blob/master/docs/en/user_guide/msopprof_user_guide.md + +## Done when + +VMI can compile and run stages=2 / `block_k=1024` Persistent tiles with +dual-buffer ping-pong — same overlap pattern as CCE and `target_mi.pto`. diff --git a/docs/repro/pipelining/fixtures/current_slow_vmi.pto b/docs/repro/pipelining/fixtures/current_slow_vmi.pto new file mode 100644 index 0000000000..4494a4cce1 --- /dev/null +++ b/docs/repro/pipelining/fixtures/current_slow_vmi.pto @@ -0,0 +1,47 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: Persistent vector tile: stages=1, block_k=512. + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind + // stages = 1 + // block_k = 512 +} { + func.func @persistent_quant_tile_k512( + %x_ub: !pto.ptr, + %y_ub: !pto.ptr, + %scale_slot: !pto.vmi.vreg<1xf32>, + %block_count: index) + attributes {pto.kernel} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c512 = arith.constant 512 : index + + pto.vecscope { + %scale = pto.vmi.vbrc %scale_slot + : !pto.vmi.vreg<1xf32> -> !pto.vmi.vreg<512xf32> + scf.for %i = %c0 to %block_count step %c1 { + %off = arith.muli %i, %c512 : index + %x_bf16 = pto.vmi.vload %x_ub[%off] + : !pto.ptr -> !pto.vmi.vreg<512xbf16> + %x_f32 = pto.vmi.vcvt %x_bf16 + : !pto.vmi.vreg<512xbf16> -> !pto.vmi.vreg<512xf32> + %scaled = pto.vmi.vmul %x_f32, %scale + : !pto.vmi.vreg<512xf32>, !pto.vmi.vreg<512xf32> + -> !pto.vmi.vreg<512xf32> + %y = pto.vmi.vcvt %scaled {saturate = "SAT"} + : !pto.vmi.vreg<512xf32> -> !pto.vmi.vreg<512xf8E4M3FN> + pto.vmi.vstore %y, %y_ub[%off] + : !pto.vmi.vreg<512xf8E4M3FN>, !pto.ptr + } + } + return + } +} diff --git a/docs/repro/pipelining/fixtures/current_slow_vmi.py b/docs/repro/pipelining/fixtures/current_slow_vmi.py new file mode 100644 index 0000000000..2c0f4e96b2 --- /dev/null +++ b/docs/repro/pipelining/fixtures/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/repro/pipelining/fixtures/desired_vmi.pto b/docs/repro/pipelining/fixtures/desired_vmi.pto new file mode 100644 index 0000000000..da24d2cdb3 --- /dev/null +++ b/docs/repro/pipelining/fixtures/desired_vmi.pto @@ -0,0 +1,62 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Idiomatic Persistent VMI: stages=2 and block_k=1024 (does not lower today). + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind, + pto.pipeline.stages = 2 : i32, + pto.block_k = 1024 : i32 +} { + func.func @persistent_quant_tile_k1024( + %x_ub0: !pto.ptr, + %x_ub1: !pto.ptr, + %y_ub0: !pto.ptr, + %y_ub1: !pto.ptr, + %scale_slot: !pto.vmi.vreg<1xf32>, + %block_count: index) + attributes {pto.kernel} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c1024 = arith.constant 1024 : index + + pto.vecscope { + %scale = pto.vmi.vbrc %scale_slot + : !pto.vmi.vreg<1xf32> -> !pto.vmi.vreg<1024xf32> + scf.for %i = %c0 to %block_count step %c1 { + %off = arith.muli %i, %c1024 : index + %stage = arith.remui %i, %c2 : index + %use1 = arith.cmpi eq, %stage, %c1 : index + %x_ub = scf.if %use1 -> (!pto.ptr) { + scf.yield %x_ub1 : !pto.ptr + } else { + scf.yield %x_ub0 : !pto.ptr + } + %y_ub = scf.if %use1 -> (!pto.ptr) { + scf.yield %y_ub1 : !pto.ptr + } else { + scf.yield %y_ub0 : !pto.ptr + } + %x_bf16 = pto.vmi.vload %x_ub[%off] + : !pto.ptr -> !pto.vmi.vreg<1024xbf16> + %x_f32 = pto.vmi.vcvt %x_bf16 + : !pto.vmi.vreg<1024xbf16> -> !pto.vmi.vreg<1024xf32> + %scaled = pto.vmi.vmul %x_f32, %scale + : !pto.vmi.vreg<1024xf32>, !pto.vmi.vreg<1024xf32> + -> !pto.vmi.vreg<1024xf32> + %y = pto.vmi.vcvt %scaled {saturate = "SAT"} + : !pto.vmi.vreg<1024xf32> -> !pto.vmi.vreg<1024xf8E4M3FN> + pto.vmi.vstore %y, %y_ub[%off] + : !pto.vmi.vreg<1024xf8E4M3FN>, !pto.ptr + } + } + return + } +} diff --git a/docs/repro/pipelining/fixtures/desired_vmi.py b/docs/repro/pipelining/fixtures/desired_vmi.py new file mode 100644 index 0000000000..cdee3d6176 --- /dev/null +++ b/docs/repro/pipelining/fixtures/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/repro/pipelining/fixtures/lowered_vpto.pto b/docs/repro/pipelining/fixtures/lowered_vpto.pto new file mode 100644 index 0000000000..7bcff4e15e --- /dev/null +++ b/docs/repro/pipelining/fixtures/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/repro/pipelining/fixtures/reference_asc_cce.asc b/docs/repro/pipelining/fixtures/reference_asc_cce.asc new file mode 100644 index 0000000000..a7989be4b7 --- /dev/null +++ b/docs/repro/pipelining/fixtures/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/repro/pipelining/fixtures/target_mi.pto b/docs/repro/pipelining/fixtures/target_mi.pto new file mode 100644 index 0000000000..39407b79f1 --- /dev/null +++ b/docs/repro/pipelining/fixtures/target_mi.pto @@ -0,0 +1,84 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: stages=2 / block_k=1024 Persistent schedule — dual UB ping-pong, +// pipe-event sync on slot = i % 2, and UNPK / mul / PK4 VF body per tile. +// Mirrors the ASC/CCE reference_asc_cce.asc HardEvent shell. + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind, + pto.pipeline.stages = 2 : i32, + pto.block_k = 1024 : i32 +} { + func.func @persistent_quant_tile_k1024_mi( + %x_ub0: !pto.ptr, + %x_ub1: !pto.ptr, + %y_ub0: !pto.ptr, + %y_ub1: !pto.ptr, + %scale: !pto.vreg<64xf32>, + %block_count: index) attributes {pto.kernel} { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c1024 = arith.constant 1024 : index + + // Prime dual-slot events (ASC: SetFlag MTE3_V / V_MTE2 on slots 0 and 1). + pto.set_flag["PIPE_MTE3", "PIPE_V", "EVENT_ID0"] + pto.set_flag["PIPE_MTE3", "PIPE_V", "EVENT_ID1"] + pto.set_flag["PIPE_V", "PIPE_MTE2", "EVENT_ID0"] + pto.set_flag["PIPE_V", "PIPE_MTE2", "EVENT_ID1"] + + scf.for %i = %c0 to %block_count step %c1 { + %slot = arith.remui %i, %c2 : index + %use1 = arith.cmpi eq, %slot, %c1 : index + %off = arith.muli %i, %c1024 : index + + // Wait producer-ready / consumer-done for this ping-pong slot. + pto.wait_flag_dyn [#pto.pipe, #pto.pipe, %slot] + // (Host MTE2 would fill the selected UB slot here.) + pto.set_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.wait_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.wait_flag_dyn [#pto.pipe, #pto.pipe, %slot] + + %x_ub = scf.if %use1 -> (!pto.ptr) { + scf.yield %x_ub1 : !pto.ptr + } else { + scf.yield %x_ub0 : !pto.ptr + } + %y_ub = scf.if %use1 -> (!pto.ptr) { + scf.yield %y_ub1 : !pto.ptr + } else { + scf.yield %y_ub0 : !pto.ptr + } + + pto.vecscope { + %m32 = pto.pset_b32 "PAT_ALL" : !pto.mask + %x0 = pto.vlds %x_ub[%off] {dist = "UNPK_B16"} + : !pto.ptr -> !pto.vreg<64xf32> + %y0 = pto.vmul %x0, %scale, %m32 + : !pto.vreg<64xf32>, !pto.vreg<64xf32>, !pto.mask + -> !pto.vreg<64xf32> + pto.vsts %y0, %y_ub[%off], %m32 {dist = "PK4_B32"} + : !pto.vreg<64xf32>, !pto.ptr, !pto.mask + } + + pto.set_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.set_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.wait_flag_dyn [#pto.pipe, #pto.pipe, %slot] + pto.set_flag_dyn [#pto.pipe, #pto.pipe, %slot] + } + + // Drain both slots. + pto.wait_flag["PIPE_MTE3", "PIPE_V", "EVENT_ID0"] + pto.wait_flag["PIPE_MTE3", "PIPE_V", "EVENT_ID1"] + pto.wait_flag["PIPE_V", "PIPE_MTE2", "EVENT_ID0"] + pto.wait_flag["PIPE_V", "PIPE_MTE2", "EVENT_ID1"] + return + } +} diff --git a/docs/repro/pipelining/scripts/check_stages.sh b/docs/repro/pipelining/scripts/check_stages.sh new file mode 100755 index 0000000000..b9f3d6b803 --- /dev/null +++ b/docs/repro/pipelining/scripts/check_stages.sh @@ -0,0 +1,194 @@ +#!/usr/bin/env bash +# Compile-check CCE stages=2 vs VMI stages=1/2 Persistent paths; record pass/fail. +set -euo pipefail +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +KERNEL_ROOT="$(cd "${SCRIPT_DIR}/.." && pwd)" +FIXTURES="${KERNEL_ROOT}/fixtures" +OUT="${KERNEL_ROOT}/sim_outputs/check_stages" +mkdir -p "${OUT}" +LOG="${OUT}/compile_results.txt" + +# shellcheck disable=SC1091 +source "${SCRIPT_DIR}/env.sh" + +NPU_ARCH="${ASCEND_NPU_ARCH:-dav-3510}" + +: > "${LOG}" +echo "pipelining check_stages — $(date -u +%Y-%m-%dT%H:%M:%SZ)" | tee -a "${LOG}" +echo "PTOAS_ROOT=${PTOAS_ROOT}" | tee -a "${LOG}" +echo "BISHENG=${BISHENG}" | tee -a "${LOG}" +echo | tee -a "${LOG}" + +# --- reference_asc_cce.asc: CCE stages=2 ping-pong (expect PASS) --- +echo "=== reference_asc_cce.asc (bisheng --cce-aicore-only) ===" | tee -a "${LOG}" +REF="${FIXTURES}/reference_asc_cce.asc" +REF_OBJ="${OUT}/reference_asc_cce.o" +set +e +"${BISHENG}" -O2 -fPIC -std=c++17 --npu-arch="${NPU_ARCH}" --cce-aicore-only -c \ + "${REF}" -o "${REF_OBJ}" \ + -I"${ASCEND}/include" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw/impl" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw/interface" \ + > "${OUT}/reference_asc_cce.log" 2>&1 +ref_rc=$? +set -e +if [ "${ref_rc}" -eq 0 ] && [ -s "${REF_OBJ}" ]; then + echo "PASS: reference_asc_cce.asc -> non-empty .o (CCE stages=2 ping-pong is legal)" | tee -a "${LOG}" +else + echo "FAIL: reference_asc_cce.asc compile exit ${ref_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/reference_asc_cce.log" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +# --- desired_vmi.py: ptodsl frontend + full lower (expect layout reject at lower) --- +echo "=== desired_vmi.py (ptodsl -> MLIR frontend) ===" | tee -a "${LOG}" +set +e +python3 "${FIXTURES}/desired_vmi.py" > "${OUT}/desired_vmi.mlir" 2> "${OUT}/desired_vmi_py.err" +py_rc=$? +set -e +if [ "${py_rc}" -eq 0 ]; then + echo "PASS: desired_vmi.py emitted MLIR (frontend only)" | tee -a "${LOG}" +else + echo "FAIL: desired_vmi.py compile exit ${py_rc}" | tee -a "${LOG}" + tail -30 "${OUT}/desired_vmi_py.err" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +echo "=== desired_vmi.py MLIR (ptoas -> LLVM IR, stages=2 block_k=1024) ===" | tee -a "${LOG}" +DESIRED_LL="${OUT}/desired_vmi.ll" +set +e +if command -v ptoas >/dev/null 2>&1 && [ -f "${OUT}/desired_vmi.mlir" ]; then + ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto-llvm-ir "${OUT}/desired_vmi.mlir" -o "${DESIRED_LL}" \ + > "${OUT}/desired_vmi_lower.log" 2>&1 + desired_lower_rc=$? +else + desired_lower_rc=127 +fi +set -e +if [ "${desired_lower_rc}" -eq 0 ]; then + echo "PASS: desired_vmi full lower (unexpected — gap may be closed)" | tee -a "${LOG}" +else + echo "FAIL: desired_vmi full lower exit ${desired_lower_rc} (layout reject on 1024-wide vmul)" | tee -a "${LOG}" + tail -15 "${OUT}/desired_vmi_lower.log" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +# --- desired_vmi.pto: ptoas --emit-vpto (expect layout / lower fail) --- +echo "=== desired_vmi.pto (ptoas --emit-vpto) ===" | tee -a "${LOG}" +DESIRED_PTO="${FIXTURES}/desired_vmi.pto" +DESIRED_VPTO="${OUT}/desired_vmi.vpto" +set +e +if command -v ptoas >/dev/null 2>&1; then + ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto "${DESIRED_PTO}" -o "${DESIRED_VPTO}" \ + > "${OUT}/desired_vmi_ptoas.log" 2>&1 + desired_pto_rc=$? +else + echo "SKIP: ptoas not on PATH" | tee -a "${LOG}" + desired_pto_rc=127 +fi +set -e +if [ "${desired_pto_rc}" -eq 0 ] && [ -s "${DESIRED_VPTO}" ]; then + echo "PASS: desired_vmi.pto --emit-vpto (unexpected — gap may be closed)" | tee -a "${LOG}" +else + echo "FAIL: desired_vmi.pto --emit-vpto exit ${desired_pto_rc}" | tee -a "${LOG}" + tail -15 "${OUT}/desired_vmi_ptoas.log" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +# --- target_mi.pto: emit-vpto then LLVM IR -> bisheng (expect bisheng crash) --- +echo "=== target_mi.pto (ptoas --emit-vpto) ===" | tee -a "${LOG}" +TARGET_MI="${FIXTURES}/target_mi.pto" +TARGET_VPTO="${OUT}/target_mi.vpto" +LLVM_IR="${OUT}/target_mi.ll" +OBJ="${OUT}/target_mi.o" + +set +e +if command -v ptoas >/dev/null 2>&1; then + ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto "${TARGET_MI}" -o "${TARGET_VPTO}" \ + > "${OUT}/target_mi_emit_vpto.log" 2>&1 + emit_vpto_rc=$? +else + echo "SKIP: ptoas not on PATH" | tee -a "${LOG}" + emit_vpto_rc=127 +fi +set -e + +if [ "${emit_vpto_rc}" -eq 0 ] && [ -s "${TARGET_VPTO}" ]; then + echo "PASS: target_mi.pto --emit-vpto" | tee -a "${LOG}" +else + echo "FAIL: target_mi.pto --emit-vpto exit ${emit_vpto_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/target_mi_emit_vpto.log" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +echo "=== target_mi.pto (ptoas -> LLVM IR -> bisheng object) ===" | tee -a "${LOG}" +set +e +if command -v ptoas >/dev/null 2>&1; then + ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto-llvm-ir "${TARGET_MI}" -o "${LLVM_IR}" \ + > "${OUT}/target_mi_ptoas.log" 2>&1 + pto_rc=$? +else + pto_rc=127 +fi +set -e + +if [ "${pto_rc}" -eq 0 ]; then + echo "PASS: ptoas --emit-vpto-llvm-ir target_mi.pto" | tee -a "${LOG}" + set +e + "${BISHENG}" -O2 -fPIC -std=c++17 --npu-arch="${NPU_ARCH}" -c -x ir \ + "${LLVM_IR}" -o "${OBJ}" > "${OUT}/target_mi_bisheng.log" 2>&1 + b_rc=$? + set -e + if [ "${b_rc}" -eq 0 ] && [ -s "${OBJ}" ]; then + echo "PASS: bisheng object compile (unexpected — gap may be closed)" | tee -a "${LOG}" + else + echo "FAIL: bisheng object compile exit ${b_rc} (expected on vmi-v0.1.3)" | tee -a "${LOG}" + tail -10 "${OUT}/target_mi_bisheng.log" | tee -a "${LOG}" + fi +else + echo "FAIL: ptoas emit-vpto-llvm-ir exit ${pto_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/target_mi_ptoas.log" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +# --- current_slow_vmi stages=1: lowered_vpto.pto is the working MI path --- +echo "=== current_slow_vmi.py (ptodsl -> MLIR frontend, stages=1 block_k=512) ===" | tee -a "${LOG}" +set +e +python3 "${FIXTURES}/current_slow_vmi.py" > "${OUT}/current_slow_vmi.mlir" 2> "${OUT}/current_slow_vmi_py.err" +slow_py_rc=$? +set -e +if [ "${slow_py_rc}" -eq 0 ]; then + echo "PASS: current_slow_vmi.py emitted MLIR (frontend only)" | tee -a "${LOG}" +else + echo "FAIL: current_slow_vmi.py compile exit ${slow_py_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/current_slow_vmi_py.err" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +echo "=== lowered_vpto.pto (stages=1 chunked MI, ptoas -> LLVM IR) ===" | tee -a "${LOG}" +LOWERED_PTO="${FIXTURES}/lowered_vpto.pto" +SLOW_LL="${OUT}/lowered_vpto.ll" +set +e +if command -v ptoas >/dev/null 2>&1; then + ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto-llvm-ir "${LOWERED_PTO}" -o "${SLOW_LL}" \ + > "${OUT}/lowered_vpto_ptoas.log" 2>&1 + lowered_rc=$? +else + lowered_rc=127 +fi +set -e +if [ "${lowered_rc}" -eq 0 ] && [ -s "${SLOW_LL}" ]; then + echo "PASS: lowered_vpto.pto -> LLVM IR (stages=1 working MI path)" | tee -a "${LOG}" +else + echo "FAIL: lowered_vpto.pto emit-vpto-llvm-ir exit ${lowered_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/lowered_vpto_ptoas.log" | tee -a "${LOG}" +fi + +echo | tee -a "${LOG}" +echo "Full log: ${LOG}" | tee -a "${LOG}" diff --git a/docs/repro/pipelining/scripts/env.sh b/docs/repro/pipelining/scripts/env.sh new file mode 100755 index 0000000000..33fff862d4 --- /dev/null +++ b/docs/repro/pipelining/scripts/env.sh @@ -0,0 +1,23 @@ +#!/usr/bin/env bash +# CANN + PTOAS env for pipelining stage compile checks. +set -euo pipefail +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +KERNEL_ROOT="$(cd "${SCRIPT_DIR}/.." && pwd)" +ASCEND_HOME_PATH="${ASCEND_HOME_PATH:-${ASCEND_TOOLKIT_HOME:-/usr/local/Ascend/cann-9.0.0}}" + +# pipelining -> repro -> docs -> repo (three levels up from repro root). +PTOAS_ROOT="${PTOAS_ROOT:-$(cd "${KERNEL_ROOT}/../../.." && pwd)}" + +# shellcheck disable=SC1091 +source "${ASCEND_HOME_PATH}/bin/setenv.bash" +if [ -n "${PTOAS_ROOT}" ] && [ -f "${PTOAS_ROOT}/scripts/ptoas_env.sh" ]; then + export PTOAS_ENV_SKIP_SMOKE_TEST="${PTOAS_ENV_SKIP_SMOKE_TEST:-1}" + # shellcheck disable=SC1091 + source "${PTOAS_ROOT}/scripts/ptoas_env.sh" + [ -d "${PTOAS_ROOT}/build/tools/ptoas" ] && export PATH="${PTOAS_ROOT}/build/tools/ptoas:${PATH}" + [ -d "${PTOAS_ROOT}/install/bin" ] && export PATH="${PTOAS_ROOT}/install/bin:${PATH}" +fi + +export ASCEND_NPU_ARCH="${ASCEND_NPU_ARCH:-dav-3510}" +export BISHENG="${BISHENG:-${ASCEND_HOME_PATH}/tools/bisheng_compiler/bin/bisheng}" +export ASCEND="${ASCEND_HOME_PATH}" From 013ab628d59ee54b30417c2b5c6a181b5bb55b80 Mon Sep 17 00:00:00 2001 From: learning-chip Date: Tue, 28 Jul 2026 19:03:47 +0800 Subject: [PATCH 2/3] Expand pipelining repro with isolate matrix and on-device CCE host Add stages/block_k and cast-back fixtures, extend check_stages.sh, and a minimal AscendC host that runs stages=2 scale on NPU. --- docs/repro/pipelining/.gitignore | 2 + docs/repro/pipelining/README.md | 115 ++++---- docs/repro/pipelining/device/build.sh | 43 +++ docs/repro/pipelining/device/main.cpp | 65 ++++ docs/repro/pipelining/device/run.sh | 12 + .../repro/pipelining/device/scale_stages2.asc | 68 +++++ .../fixtures/cast_back_stages1_k512_vmi.pto | 47 +++ .../fixtures/cast_back_stages1_k512_vmi.py | 32 ++ .../fixtures/cast_back_stages2_k1024_vmi.pto | 62 ++++ .../fixtures/cast_back_stages2_k1024_vmi.py | 41 +++ .../isolate_stages1_blockk1024_vmi.pto | 47 +++ .../isolate_stages1_blockk1024_vmi.py | 35 +++ .../isolate_stages2_blockk512_vmi.pto | 62 ++++ .../fixtures/isolate_stages2_blockk512_vmi.py | 38 +++ .../fixtures/reference_asc_cce_cast_back.asc | 57 ++++ .../fixtures/reference_asc_cce_f8.asc | 59 ++++ docs/repro/pipelining/scripts/check_stages.sh | 277 +++++++++--------- docs/repro/pipelining/scripts/env.sh | 13 +- 18 files changed, 885 insertions(+), 190 deletions(-) create mode 100755 docs/repro/pipelining/device/build.sh create mode 100644 docs/repro/pipelining/device/main.cpp create mode 100755 docs/repro/pipelining/device/run.sh create mode 100644 docs/repro/pipelining/device/scale_stages2.asc create mode 100644 docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.pto create mode 100644 docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.py create mode 100644 docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.pto create mode 100644 docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.py create mode 100644 docs/repro/pipelining/fixtures/isolate_stages1_blockk1024_vmi.pto create mode 100644 docs/repro/pipelining/fixtures/isolate_stages1_blockk1024_vmi.py create mode 100644 docs/repro/pipelining/fixtures/isolate_stages2_blockk512_vmi.pto create mode 100644 docs/repro/pipelining/fixtures/isolate_stages2_blockk512_vmi.py create mode 100644 docs/repro/pipelining/fixtures/reference_asc_cce_cast_back.asc create mode 100644 docs/repro/pipelining/fixtures/reference_asc_cce_f8.asc diff --git a/docs/repro/pipelining/.gitignore b/docs/repro/pipelining/.gitignore index a95b67ec69..4478ddd95b 100644 --- a/docs/repro/pipelining/.gitignore +++ b/docs/repro/pipelining/.gitignore @@ -1,5 +1,7 @@ build/ sim_outputs/ +device/build/ + *.so *.o log_ca/ diff --git a/docs/repro/pipelining/README.md b/docs/repro/pipelining/README.md index 53912f0160..9dd400ea20 100644 --- a/docs/repro/pipelining/README.md +++ b/docs/repro/pipelining/README.md @@ -1,81 +1,90 @@ # pipelining — VMI Persistent stages vs CCE ping-pong repro -Self-contained compile repro for PTOAS gap 05: 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. Ascend CCE can express stages=2 ping-pong -with `block_k=1024`; VMI on tag **vmi-v0.1.3** cannot reach the same schedule -— layout assignment rejects wide tiles, and the hand-written MI path crashes -bisheng even after ptoas emits LLVM IR. +Self-contained compile (+ optional on-device CCE) repro for long-K Persistent +quant / dequant and residual-mix-style scale: production wants stages=2 and +`block_k > 512` so MTE and vector stay overlapped. Ascend CCE can express and +**run** stages=2 ping-pong; VMI on tag **vmi-v0.1.3** cannot reach the same +wide-tile schedule — layout assignment rejects 1024-wide tiles, and the +hand-written MI path crashes bisheng even after ptoas emits LLVM IR. -Fixtures are copied from `docs/feature-gaps/vmi/05_persistent_stages_block_k/`. +## Pin + +| Component | Ref | +|-----------|-----| +| PTOAS | tag **vmi-v0.1.3** (`ptoas` from that build; not PyPI-only) | +| Branch | `zjw/pipelining_issue` | +| CANN | 9.0 / `bisheng` | +| `PTOAS_ROOT` | repo root (`../../..` from this directory) | + +If `install/bin/ptoas` crashes on missing MLIR libs, point tools at a build tree: + +```bash +export PTOAS_TOOLS_ROOT=/path/to/PTOAS-built-at-vmi-v0.1.3 +``` ## What breaks | Path | stages | block_k | Fault mode | |------|-------:|--------:|------------| -| CCE (`fixtures/reference_asc_cce.asc`) | 2 | 1024 | **OK** — bisheng `--cce-aicore-only` produces a non-empty `.o` | -| VMI desired (`fixtures/desired_vmi.*`) | 2 | 1024 | **Layout reject** — `pto.vmi.mulf` on 1024-wide tile: deinterleaved=2 vs contiguous | -| VMI target MI (`fixtures/target_mi.pto`) | 2 | 1024 | **Bisheng crash** — ptoas `--emit-vpto` and LLVM IR OK; object compile exit 70 | -| VMI current slow (`fixtures/current_slow_vmi.*`) | 1 | 512 | **Frontend OK** — ptodsl emits MLIR; idiomatic VMI `.pto` hits same layout issue | -| VMI lowered MI (`fixtures/lowered_vpto.pto`) | 1 | 512 | **OK to LLVM IR** — chunked 128-wide `vlds`/`vmul`/`vsts` path lowers | +| CCE (`fixtures/reference_asc_cce*.asc`, `device/scale_stages2.asc`) | 2 | 512–1024 | **OK** — compile; device host numerical OK | +| VMI stages=1 / 512 (`current_slow_vmi.*`) | 1 | 512 | **emit-vpto OK** | +| VMI stages=2 / 512 (`isolate_stages2_blockk512_vmi.*`) | 2 | 512 | **emit-vpto OK** (dual-buffer alone is fine) | +| VMI stages=1 / 1024 (`isolate_stages1_blockk1024_vmi.*`) | 1 | 1024 | **Layout reject** on 1024-wide vmul | +| VMI stages=2 / 1024 (`desired_vmi.*`) | 2 | 1024 | **Layout reject** (stages+wide) | +| Cast-back VMI (`cast_back_stages*_*.*`) | 1–2 | 512–1024 | **Layout reject** on f8↔bf16 Persistent | +| VMI target MI (`target_mi.pto`) | 2 | 1024 | **Bisheng crash** after LLVM IR emit | -Run `./scripts/check_stages.sh` to refresh the table below. +So the blocker for the desired quant schedule is **wide `block_k=1024` layout**, +not dual-buffer stages alone. Cast-back Persistent hits the same layout wall. -## Check results (vmi-v0.1.3, 2026-07-27) +## Check results (vmi-v0.1.3 tools, 2026-07-28) -Measured by `./scripts/check_stages.sh`: +```bash +source scripts/env.sh +# optional: export PTOAS_TOOLS_ROOT=... +./scripts/check_stages.sh +``` | Step | Result | |------|--------| -| `reference_asc_cce.asc` → bisheng `--cce-aicore-only` | **PASS** | -| `desired_vmi.py` → MLIR frontend | **PASS** (frontend only) | -| `desired_vmi.py` MLIR → LLVM IR | **FAIL** — layout reject on 1024-wide vmul | -| `desired_vmi.pto` → `--emit-vpto` | **FAIL** — same layout reject | -| `target_mi.pto` → `--emit-vpto` | **PASS** | -| `target_mi.pto` → LLVM IR → bisheng `.o` | **FAIL** — bisheng crash (exit 70) | -| `current_slow_vmi.py` → MLIR frontend | **PASS** (frontend only) | -| `lowered_vpto.pto` → LLVM IR | **PASS** — stages=1 chunked MI path | +| `reference_asc_cce.asc` / `_f8` / `_cast_back` | **PASS** bisheng `--cce-aicore-only` | +| `device/scale_stages2.asc` aicore object | **PASS** | +| `current_slow_vmi.pto` `--emit-vpto` | **PASS** | +| `isolate_stages2_blockk512_vmi.pto` | **PASS** | +| `isolate_stages1_blockk1024_vmi.pto` | **FAIL** layout | +| `desired_vmi.pto` | **FAIL** layout | +| `cast_back_stages1_k512_vmi.pto` | **FAIL** layout | +| `cast_back_stages2_k1024_vmi.pto` | **FAIL** layout | +| `lowered_vpto.pto` → LLVM IR | **PASS** | +| `target_mi.pto` → emit + LLVM IR | **PASS**; bisheng `.o` **FAIL** | Full log: `sim_outputs/check_stages/compile_results.txt` -## Layout +## On-device CCE +```bash +cd device +./run.sh # build + launch stages=2 f32 scale, check max_err +MSOPPROF=1 ./run.sh # optional msopprof wrap if installed ``` -pipelining/ - fixtures/ IR dumps and reference ASC from gap 05 - scripts/ env.sh, check_stages.sh -``` - -CCE stages=2 baseline lives in `fixtures/reference_asc_cce.asc` (compile-only -is enough for this repro; on-device overlap profiling is TODO below). -## Prerequisites +Measured 2026-07-28: -- CANN 9.0 (`bisheng`) -- Built PTOAS at tag **vmi-v0.1.3** (branch `zjw/pipelining_issue`) - -Push remote hint: https://github.com/learning-chip/PTOAS.git - -## Run - -From this directory (`docs/repro/pipelining/`): - -```bash -source scripts/env.sh -./scripts/check_stages.sh +``` +scale_stages2: N=2048 mism=0 max_err=0 OK ``` -## Pin - -- PTOAS tag: **vmi-v0.1.3** -- Branch: `zjw/pipelining_issue` -- `PTOAS_ROOT` defaults to the repo root (`../../..` from this repro directory). +CCE stages=2 runs with dual-buffer HardEvent overlap. VMI stages=2 / `block_k=1024` +does not compile through layout — that is the PTOAS gap this package isolates. -## TODO +## Layout -- On-device run and **msopprof** bandwidth study for stages=2 CCE vs stages=1 - VMI on large hidden sizes. -- Guide: https://gitcode.com/Ascend/msopprof/blob/master/docs/en/user_guide/msopprof_user_guide.md +``` +pipelining/ + fixtures/ CCE twins, VMI isolate + cast-back cells, target_mi + device/ AscendC host for stages=2 CCE numerical check + scripts/ env.sh, check_stages.sh +``` ## Done when diff --git a/docs/repro/pipelining/device/build.sh b/docs/repro/pipelining/device/build.sh new file mode 100755 index 0000000000..bc61cc0646 --- /dev/null +++ b/docs/repro/pipelining/device/build.sh @@ -0,0 +1,43 @@ +#!/usr/bin/env bash +# Build launchable CCE stages=2 scale host (relative CANN / PTOAS paths). +set -euo pipefail +DEVICE_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +# shellcheck disable=SC1091 +source "${DEVICE_DIR}/../scripts/env.sh" + +AICORE_ARCH="${PTO_AICORE_ARCH:-dav-c310}" +OUT="${DEVICE_DIR}/build" +mkdir -p "${OUT}" + +echo "bisheng=${BISHENG}" +echo "ASCEND=${ASCEND}" + +"${BISHENG}" -c -fPIC -O2 -std=c++17 -xcce \ + -Xhost-start -Xhost-end \ + -mllvm -cce-aicore-stack-size=0x8000 \ + -mllvm -cce-aicore-function-stack-size=0x8000 \ + -mllvm -cce-aicore-record-overflow=true \ + -mllvm -cce-aicore-addr-transform \ + -mllvm -cce-aicore-dcci-insert-for-scalar=false \ + --cce-aicore-arch="${AICORE_ARCH}" \ + -I"${ASCEND}/include" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw/impl" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw/interface" \ + "${DEVICE_DIR}/scale_stages2.asc" \ + -o "${OUT}/scale_stages2.o" + +"${BISHENG}" -fPIC -shared --cce-fatobj-link \ + -o "${OUT}/libscale_stages2.so" \ + "${OUT}/scale_stages2.o" \ + -Wl,--no-as-needed -L"${ASCEND}/lib64" -lruntime + +"${BISHENG}" -O2 -std=c++17 -xc++ \ + -I"${ASCEND}/include" \ + "${DEVICE_DIR}/main.cpp" \ + -o "${OUT}/scale_stages2_host" \ + -L"${OUT}" -L"${ASCEND}/lib64" \ + -lscale_stages2 -lascendcl -lruntime -ltiling_api -lplatform -lc_sec -ldl -lnnopbase \ + -Wl,-rpath,"${OUT}" -Wl,-rpath,"${ASCEND}/lib64" + +echo "OK: ${OUT}/scale_stages2_host" diff --git a/docs/repro/pipelining/device/main.cpp b/docs/repro/pipelining/device/main.cpp new file mode 100644 index 0000000000..793e8bba6f --- /dev/null +++ b/docs/repro/pipelining/device/main.cpp @@ -0,0 +1,65 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// Minimal host: fill f32 input, launch CCE stages=2 scale, check vs host ref. +#include "acl/acl.h" +#include +#include +#include +#include + +void LaunchScaleStages2(void* x, void* scale, void* out, void* stream); + +int main() { + constexpr int N = 2048; + const size_t nbytes = N * sizeof(float); + const float scale_v = 1.5f; + + aclInit(nullptr); + aclrtSetDevice(0); + aclrtStream stream; + aclrtCreateStream(&stream); + + float *x_h = nullptr, *y_h = nullptr, *scale_h = nullptr; + void *x_d = nullptr, *y_d = nullptr, *scale_d = nullptr; + aclrtMallocHost((void**)&x_h, nbytes); + aclrtMallocHost((void**)&y_h, nbytes); + aclrtMallocHost((void**)&scale_h, sizeof(float)); + aclrtMalloc(&x_d, nbytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&y_d, nbytes, ACL_MEM_MALLOC_HUGE_FIRST); + aclrtMalloc(&scale_d, sizeof(float), ACL_MEM_MALLOC_HUGE_FIRST); + + scale_h[0] = scale_v; + for (int i = 0; i < N; ++i) { + x_h[i] = 0.25f * static_cast(i % 17); + y_h[i] = 0.f; + } + + aclrtMemcpy(x_d, nbytes, x_h, nbytes, ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(scale_d, sizeof(float), scale_h, sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE); + aclrtMemcpy(y_d, nbytes, y_h, nbytes, ACL_MEMCPY_HOST_TO_DEVICE); + + LaunchScaleStages2(x_d, scale_d, y_d, stream); + aclrtSynchronizeStream(stream); + aclrtMemcpy(y_h, nbytes, y_d, nbytes, ACL_MEMCPY_DEVICE_TO_HOST); + + int mism = 0; + float max_err = 0.f; + for (int i = 0; i < N; ++i) { + float ref = x_h[i] * scale_v; + float err = std::fabs(ref - y_h[i]); + if (err > max_err) max_err = err; + if (err > 1e-5f) ++mism; + } + std::printf("scale_stages2: N=%d mism=%d max_err=%.4g %s\n", N, mism, max_err, + mism == 0 ? "OK" : "FAIL"); + + aclrtFree(x_d); + aclrtFree(y_d); + aclrtFree(scale_d); + aclrtFreeHost(x_h); + aclrtFreeHost(y_h); + aclrtFreeHost(scale_h); + aclrtDestroyStream(stream); + aclrtResetDevice(0); + aclFinalize(); + return mism == 0 ? 0 : 1; +} diff --git a/docs/repro/pipelining/device/run.sh b/docs/repro/pipelining/device/run.sh new file mode 100755 index 0000000000..9fd7f074fb --- /dev/null +++ b/docs/repro/pipelining/device/run.sh @@ -0,0 +1,12 @@ +#!/usr/bin/env bash +# Build and run CCE stages=2 on NPU. Optional: MSOPPROF=1 for msopprof wrap. +set -euo pipefail +DEVICE_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +bash "${DEVICE_DIR}/build.sh" +# shellcheck disable=SC1091 +source "${DEVICE_DIR}/../scripts/env.sh" +HOST="${DEVICE_DIR}/build/scale_stages2_host" +if [ "${MSOPPROF:-0}" = "1" ] && command -v msopprof >/dev/null 2>&1; then + msopprof --application="${HOST}" --output="${DEVICE_DIR}/build/msopprof" || true +fi +exec "${HOST}" diff --git a/docs/repro/pipelining/device/scale_stages2.asc b/docs/repro/pipelining/device/scale_stages2.asc new file mode 100644 index 0000000000..49487f548c --- /dev/null +++ b/docs/repro/pipelining/device/scale_stages2.asc @@ -0,0 +1,68 @@ +// Launchable CCE stages=2 scale: GM f32 * scale -> GM f32 (real MTE + VF). +// Proves dual-buffer HardEvent ping-pong runs on device; VMI stages=2 is the gap. +#include "kernel_operator.h" + +__simd_vf__ inline void vf_scale_f32(__ubuf__ float* x_ub, __ubuf__ float* scale_ub, + __ubuf__ float* out_ub) { + vector_bool v32 = pset_b32(PAT_ALL); + for (int i = 0; i < 512; i += 64) { + vector_f32 xv, sv, yv; + vlds(xv, x_ub + i, 0, NORM); + vlds(sv, scale_ub, 0, BRC_B32); + vmul(yv, xv, sv, v32, MODE_ZEROING); + vsts(yv, out_ub + i, 0, NORM_B32, v32); + } +} + +// N = 2048 f32, block_k = 512, stages = 2, 4 tiles. +extern "C" __global__ __vector__ void scale_stages2(__gm__ float* x, __gm__ float* scale, + __gm__ float* out) { + AscendC::InitSocState(); + __ubuf__ uint8_t* ub = (__ubuf__ uint8_t*)0; + + constexpr int kBlockK = 512; + constexpr int kTiles = 4; + constexpr int kBytes = kBlockK * 4; + + __ubuf__ float* sc_ub = (__ubuf__ float*)(ub + 4 * kBytes); + copy_gm_to_ubuf_align_v2((__ubuf__ uint8_t*)sc_ub, (__gm__ uint8_t*)scale, 0, 1, 4, 0, 0, 0, 0, + 32, 32); + AscendC::SetFlag(2); + AscendC::WaitFlag(2); + + AscendC::SetFlag(0); + AscendC::SetFlag(1); + AscendC::SetFlag(0); + AscendC::SetFlag(1); + + for (int w = 0; w < kTiles; ++w) { + int slot = w & 1; + __ubuf__ float* x_ub = (__ubuf__ float*)(ub + slot * kBytes); + __ubuf__ float* o_ub = (__ubuf__ float*)(ub + (2 + slot) * kBytes); + + AscendC::WaitFlag(slot); + copy_gm_to_ubuf_align_v2((__ubuf__ uint8_t*)x_ub, (__gm__ uint8_t*)(x + w * kBlockK), 0, 1, + kBytes, 0, 0, 0, 0, kBytes, kBytes); + AscendC::SetFlag(slot); + AscendC::WaitFlag(slot); + AscendC::WaitFlag(slot); + + vf_scale_f32(x_ub, sc_ub, o_ub); + + AscendC::SetFlag(slot); + AscendC::SetFlag(slot); + AscendC::WaitFlag(slot); + copy_ubuf_to_gm_align_v2((__gm__ void*)(out + w * kBlockK), (__ubuf__ void*)o_ub, 0, 1, kBytes, + 4, kBytes, kBytes); + AscendC::SetFlag(slot); + } + + AscendC::WaitFlag(0); + AscendC::WaitFlag(1); + AscendC::WaitFlag(0); + AscendC::WaitFlag(1); +} + +void LaunchScaleStages2(void* x, void* scale, void* out, void* stream) { + scale_stages2<<<1, nullptr, stream>>>((__gm__ float*)x, (__gm__ float*)scale, (__gm__ float*)out); +} diff --git a/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.pto b/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.pto new file mode 100644 index 0000000000..8712b03b56 --- /dev/null +++ b/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.pto @@ -0,0 +1,47 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Cast-back Persistent VMI: stages=1, tile_k=512 (f8 -> bf16 * scale). + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind, + pto.pipeline.stages = 1 : i32, + pto.block_k = 512 : i32 +} { + func.func @persistent_cast_back_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_f8 = pto.vmi.vload %x_ub[%off] + : !pto.ptr -> !pto.vmi.vreg<512xf8E4M3FN> + %x_f32 = pto.vmi.vcvt %x_f8 + : !pto.vmi.vreg<512xf8E4M3FN> -> !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 + : !pto.vmi.vreg<512xf32> -> !pto.vmi.vreg<512xbf16> + pto.vmi.vstore %y, %y_ub[%off] + : !pto.vmi.vreg<512xbf16>, !pto.ptr + } + } + return + } +} diff --git a/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.py b/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.py new file mode 100644 index 0000000000..faff9e5b67 --- /dev/null +++ b/docs/repro/pipelining/fixtures/cast_back_stages1_k512_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. + +"""Cast-back Persistent: e4m3 -> f32 * scale -> bf16, stages=1, tile_k=512.""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def persistent_cast_back_k512(*, BLOCK_COUNT: pto.const_expr = 4): + x_ub = pto.alloc_tile(shape=[1, 512], dtype=pto.f8e4m3) + y_ub = pto.alloc_tile(shape=[1, 512], dtype=pto.bf16) + 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_f8 = pto.vmi.vload(x_ub.as_ptr(), off, size=512) + x_f32 = pto.vmi.vcvt(x_f8, pto.f32) + scaled = pto.vmi.vmul(x_f32, scale) + y = pto.vmi.vcvt(scaled, pto.bf16) + pto.vmi.vstore(y, y_ub.as_ptr(), off) + + +if __name__ == "__main__": + print(persistent_cast_back_k512.compile().mlir_text()) diff --git a/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.pto b/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.pto new file mode 100644 index 0000000000..ae01c38e38 --- /dev/null +++ b/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_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. + +// Cast-back Persistent VMI: stages=2, tile_k=1024 (f8 -> bf16 * scale). + +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_cast_back_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_f8 = pto.vmi.vload %x_ub[%off] + : !pto.ptr -> !pto.vmi.vreg<1024xf8E4M3FN> + %x_f32 = pto.vmi.vcvt %x_f8 + : !pto.vmi.vreg<1024xf8E4M3FN> -> !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 + : !pto.vmi.vreg<1024xf32> -> !pto.vmi.vreg<1024xbf16> + pto.vmi.vstore %y, %y_ub[%off] + : !pto.vmi.vreg<1024xbf16>, !pto.ptr + } + } + return + } +} diff --git a/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.py b/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.py new file mode 100644 index 0000000000..8546d041e1 --- /dev/null +++ b/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.py @@ -0,0 +1,41 @@ +# 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. + +"""Cast-back Persistent: e4m3 -> f32 * scale -> bf16, stages=2, tile_k=1024.""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def persistent_cast_back_k1024(*, BLOCK_COUNT: pto.const_expr = 4): + x_ub0 = pto.alloc_tile(shape=[1, 1024], dtype=pto.f8e4m3) + x_ub1 = pto.alloc_tile(shape=[1, 1024], dtype=pto.f8e4m3) + y_ub0 = pto.alloc_tile(shape=[1, 1024], dtype=pto.bf16) + y_ub1 = pto.alloc_tile(shape=[1, 1024], dtype=pto.bf16) + 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, y_ptr = x_ub1.as_ptr(), y_ub1.as_ptr() + else: + x_ptr, y_ptr = x_ub0.as_ptr(), y_ub0.as_ptr() + x_f8 = pto.vmi.vload(x_ptr, off, size=1024) + x_f32 = pto.vmi.vcvt(x_f8, pto.f32) + scaled = pto.vmi.vmul(x_f32, scale) + y = pto.vmi.vcvt(scaled, pto.bf16) + pto.vmi.vstore(y, y_ptr, off) + + +if __name__ == "__main__": + try: + print(persistent_cast_back_k1024.compile().mlir_text()) + except Exception as exc: # noqa: BLE001 + print(f"emit/lower may fail (layout / stages): {exc}") diff --git a/docs/repro/pipelining/fixtures/isolate_stages1_blockk1024_vmi.pto b/docs/repro/pipelining/fixtures/isolate_stages1_blockk1024_vmi.pto new file mode 100644 index 0000000000..f7a7d21e96 --- /dev/null +++ b/docs/repro/pipelining/fixtures/isolate_stages1_blockk1024_vmi.pto @@ -0,0 +1,47 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Isolate: stages=1, block_k=1024 — wide tile without dual-buffer stages. + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind, + pto.pipeline.stages = 1 : i32, + pto.block_k = 1024 : i32 +} { + func.func @persistent_quant_stages1_k1024( + %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 + %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 + %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/repro/pipelining/fixtures/isolate_stages1_blockk1024_vmi.py b/docs/repro/pipelining/fixtures/isolate_stages1_blockk1024_vmi.py new file mode 100644 index 0000000000..61d3a341d0 --- /dev/null +++ b/docs/repro/pipelining/fixtures/isolate_stages1_blockk1024_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. + +"""Isolate cell: stages=1, block_k=1024 — wide tile without dual-buffer stages.""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def persistent_quant_stages1_k1024(*, BLOCK_COUNT: pto.const_expr = 4): + x_ub = pto.alloc_tile(shape=[1, 1024], dtype=pto.bf16) + y_ub = 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 + x_bf16 = pto.vmi.vload(x_ub.as_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_ub.as_ptr(), off) + + +if __name__ == "__main__": + try: + print(persistent_quant_stages1_k1024.compile().mlir_text()) + except Exception as exc: # noqa: BLE001 + print(f"emit/lower may fail (layout): {exc}") diff --git a/docs/repro/pipelining/fixtures/isolate_stages2_blockk512_vmi.pto b/docs/repro/pipelining/fixtures/isolate_stages2_blockk512_vmi.pto new file mode 100644 index 0000000000..0f6c918a9b --- /dev/null +++ b/docs/repro/pipelining/fixtures/isolate_stages2_blockk512_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. + +// Isolate: stages=2, block_k=512 — dual-buffer without 1024-wide layout. + +module attributes { + pto.target_arch = "a5", + pto.kernel_kind = #pto.kernel_kind, + pto.pipeline.stages = 2 : i32, + pto.block_k = 512 : i32 +} { + func.func @persistent_quant_stages2_k512( + %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 + %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 + %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<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/repro/pipelining/fixtures/isolate_stages2_blockk512_vmi.py b/docs/repro/pipelining/fixtures/isolate_stages2_blockk512_vmi.py new file mode 100644 index 0000000000..ca7e01e7d3 --- /dev/null +++ b/docs/repro/pipelining/fixtures/isolate_stages2_blockk512_vmi.py @@ -0,0 +1,38 @@ +# 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. + +"""Isolate cell: stages=2, block_k=512 — dual-buffer without wide-tile layout.""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def persistent_quant_stages2_k512(*, BLOCK_COUNT: pto.const_expr = 4): + x_ub0 = pto.alloc_tile(shape=[1, 512], dtype=pto.bf16) + x_ub1 = pto.alloc_tile(shape=[1, 512], dtype=pto.bf16) + y_ub0 = pto.alloc_tile(shape=[1, 512], dtype=pto.f8e4m3) + y_ub1 = 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 + if (i % 2) == 1: + x_ptr, y_ptr = x_ub1.as_ptr(), y_ub1.as_ptr() + else: + x_ptr, y_ptr = x_ub0.as_ptr(), y_ub0.as_ptr() + x_bf16 = pto.vmi.vload(x_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_ptr, off) + + +if __name__ == "__main__": + print(persistent_quant_stages2_k512.compile().mlir_text()) diff --git a/docs/repro/pipelining/fixtures/reference_asc_cce_cast_back.asc b/docs/repro/pipelining/fixtures/reference_asc_cce_cast_back.asc new file mode 100644 index 0000000000..f68534594c --- /dev/null +++ b/docs/repro/pipelining/fixtures/reference_asc_cce_cast_back.asc @@ -0,0 +1,57 @@ +// CCE twin for cast-back Persistent schedule (stages=2 / block_k=1024). +// HardEvent ping-pong shell aligned to quant CCE; ends in bf16 store. +#include "kernel_operator.h" + +__simd_vf__ inline void vf_cast_back_scale(__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_cast_back_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); + 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_cast_back_scale(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/repro/pipelining/fixtures/reference_asc_cce_f8.asc b/docs/repro/pipelining/fixtures/reference_asc_cce_f8.asc new file mode 100644 index 0000000000..0a3d0d480b --- /dev/null +++ b/docs/repro/pipelining/fixtures/reference_asc_cce_f8.asc @@ -0,0 +1,59 @@ +// CCE twin for long-K Persistent quant schedule (stages=2 / block_k=1024). +// Same HardEvent ping-pong + UNPK/vmul/PK body as reference_asc_cce.asc; kernel +// name marks the f8-store consumer. bisheng legality of the stages=2 shell is +// what this file proves (f8 PK4 store shapes live in target_mi.pto). +#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_quant_stages2_f8(__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); + 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/repro/pipelining/scripts/check_stages.sh b/docs/repro/pipelining/scripts/check_stages.sh index b9f3d6b803..08f8e0643b 100755 --- a/docs/repro/pipelining/scripts/check_stages.sh +++ b/docs/repro/pipelining/scripts/check_stages.sh @@ -1,5 +1,5 @@ #!/usr/bin/env bash -# Compile-check CCE stages=2 vs VMI stages=1/2 Persistent paths; record pass/fail. +# Compile-check CCE stages=2 vs VMI stages×block_k isolate + cast-back matrix. set -euo pipefail SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" KERNEL_ROOT="$(cd "${SCRIPT_DIR}/.." && pwd)" @@ -19,175 +19,180 @@ echo "PTOAS_ROOT=${PTOAS_ROOT}" | tee -a "${LOG}" echo "BISHENG=${BISHENG}" | tee -a "${LOG}" echo | tee -a "${LOG}" -# --- reference_asc_cce.asc: CCE stages=2 ping-pong (expect PASS) --- -echo "=== reference_asc_cce.asc (bisheng --cce-aicore-only) ===" | tee -a "${LOG}" -REF="${FIXTURES}/reference_asc_cce.asc" -REF_OBJ="${OUT}/reference_asc_cce.o" +compile_asc() { + local name="$1" + local src="${FIXTURES}/${name}" + local obj="${OUT}/${name%.asc}.o" + local logf="${OUT}/${name%.asc}.log" + echo "=== ${name} (bisheng --cce-aicore-only) ===" | tee -a "${LOG}" + set +e + "${BISHENG}" -O2 -fPIC -std=c++17 --npu-arch="${NPU_ARCH}" --cce-aicore-only -c \ + "${src}" -o "${obj}" \ + -I"${ASCEND}/include" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw/impl" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw/interface" \ + > "${logf}" 2>&1 + local rc=$? + set -e + if [ "${rc}" -eq 0 ] && [ -s "${obj}" ]; then + echo "PASS: ${name} -> non-empty .o" | tee -a "${LOG}" + else + echo "FAIL: ${name} compile exit ${rc}" | tee -a "${LOG}" + tail -20 "${logf}" | tee -a "${LOG}" + fi + echo | tee -a "${LOG}" +} + +emit_vpto() { + local name="$1" + local expect="$2" # pass|fail + local src="${FIXTURES}/${name}" + local outf="${OUT}/${name%.pto}.vpto" + local logf="${OUT}/${name%.pto}_emit.log" + echo "=== ${name} (ptoas --emit-vpto) expect=${expect} ===" | tee -a "${LOG}" + set +e + if command -v ptoas >/dev/null 2>&1; then + ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto "${src}" -o "${outf}" > "${logf}" 2>&1 + local rc=$? + else + echo "SKIP: ptoas not on PATH" | tee -a "${LOG}" + rc=127 + fi + set -e + if [ "${rc}" -eq 0 ] && [ -s "${outf}" ]; then + if [ "${expect}" = "pass" ]; then + echo "PASS: ${name} --emit-vpto" | tee -a "${LOG}" + else + echo "PASS: ${name} --emit-vpto (unexpected — gap may be closed)" | tee -a "${LOG}" + fi + else + if [ "${expect}" = "fail" ]; then + echo "FAIL: ${name} --emit-vpto exit ${rc} (expected layout/lower reject)" | tee -a "${LOG}" + else + echo "FAIL: ${name} --emit-vpto exit ${rc}" | tee -a "${LOG}" + fi + tail -12 "${logf}" | tee -a "${LOG}" + fi + echo | tee -a "${LOG}" +} + +ptodsl_py() { + local name="$1" + local outf="${OUT}/${name%.py}.mlir" + local errf="${OUT}/${name%.py}.err" + echo "=== ${name} (ptodsl frontend) ===" | tee -a "${LOG}" + set +e + python3 "${FIXTURES}/${name}" > "${outf}" 2> "${errf}" + local rc=$? + set -e + if [ "${rc}" -eq 0 ]; then + echo "PASS: ${name} emitted MLIR (frontend only)" | tee -a "${LOG}" + else + echo "FAIL: ${name} exit ${rc}" | tee -a "${LOG}" + tail -20 "${errf}" | tee -a "${LOG}" + fi + echo | tee -a "${LOG}" +} + +# --- CCE baselines --- +compile_asc "reference_asc_cce.asc" +compile_asc "reference_asc_cce_f8.asc" +compile_asc "reference_asc_cce_cast_back.asc" + +echo "=== device/scale_stages2.asc (bisheng --cce-aicore-only) ===" | tee -a "${LOG}" set +e "${BISHENG}" -O2 -fPIC -std=c++17 --npu-arch="${NPU_ARCH}" --cce-aicore-only -c \ - "${REF}" -o "${REF_OBJ}" \ + "${KERNEL_ROOT}/device/scale_stages2.asc" -o "${OUT}/scale_stages2_aicore.o" \ -I"${ASCEND}/include" \ -I"${ASCEND}/compiler/tikcpp/tikcfw" \ -I"${ASCEND}/compiler/tikcpp/tikcfw/impl" \ -I"${ASCEND}/compiler/tikcpp/tikcfw/interface" \ - > "${OUT}/reference_asc_cce.log" 2>&1 -ref_rc=$? + > "${OUT}/scale_stages2_aicore.log" 2>&1 +rc=$? set -e -if [ "${ref_rc}" -eq 0 ] && [ -s "${REF_OBJ}" ]; then - echo "PASS: reference_asc_cce.asc -> non-empty .o (CCE stages=2 ping-pong is legal)" | tee -a "${LOG}" +if [ "${rc}" -eq 0 ] && [ -s "${OUT}/scale_stages2_aicore.o" ]; then + echo "PASS: device/scale_stages2.asc -> non-empty .o" | tee -a "${LOG}" else - echo "FAIL: reference_asc_cce.asc compile exit ${ref_rc}" | tee -a "${LOG}" - tail -20 "${OUT}/reference_asc_cce.log" | tee -a "${LOG}" + echo "FAIL: device/scale_stages2.asc compile exit ${rc}" | tee -a "${LOG}" + tail -20 "${OUT}/scale_stages2_aicore.log" | tee -a "${LOG}" fi echo | tee -a "${LOG}" -# --- desired_vmi.py: ptodsl frontend + full lower (expect layout reject at lower) --- -echo "=== desired_vmi.py (ptodsl -> MLIR frontend) ===" | tee -a "${LOG}" -set +e -python3 "${FIXTURES}/desired_vmi.py" > "${OUT}/desired_vmi.mlir" 2> "${OUT}/desired_vmi_py.err" -py_rc=$? -set -e -if [ "${py_rc}" -eq 0 ]; then - echo "PASS: desired_vmi.py emitted MLIR (frontend only)" | tee -a "${LOG}" -else - echo "FAIL: desired_vmi.py compile exit ${py_rc}" | tee -a "${LOG}" - tail -30 "${OUT}/desired_vmi_py.err" | tee -a "${LOG}" -fi -echo | tee -a "${LOG}" -echo "=== desired_vmi.py MLIR (ptoas -> LLVM IR, stages=2 block_k=1024) ===" | tee -a "${LOG}" -DESIRED_LL="${OUT}/desired_vmi.ll" +# --- VMI matrix: stages × block_k --- +# stages=1 block_k=512 (current_slow) — emit OK on vmi-v0.1.3; lowered_vpto is chunked MI twin +ptodsl_py "current_slow_vmi.py" +emit_vpto "current_slow_vmi.pto" "pass" +echo "=== lowered_vpto.pto (stages=1 chunked MI -> LLVM IR) ===" | tee -a "${LOG}" set +e -if command -v ptoas >/dev/null 2>&1 && [ -f "${OUT}/desired_vmi.mlir" ]; then - ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ - --emit-vpto-llvm-ir "${OUT}/desired_vmi.mlir" -o "${DESIRED_LL}" \ - > "${OUT}/desired_vmi_lower.log" 2>&1 - desired_lower_rc=$? -else - desired_lower_rc=127 -fi +ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto-llvm-ir "${FIXTURES}/lowered_vpto.pto" -o "${OUT}/lowered_vpto.ll" \ + > "${OUT}/lowered_vpto_ptoas.log" 2>&1 +lowered_rc=$? set -e -if [ "${desired_lower_rc}" -eq 0 ]; then - echo "PASS: desired_vmi full lower (unexpected — gap may be closed)" | tee -a "${LOG}" +if [ "${lowered_rc}" -eq 0 ] && [ -s "${OUT}/lowered_vpto.ll" ]; then + echo "PASS: lowered_vpto.pto -> LLVM IR" | tee -a "${LOG}" else - echo "FAIL: desired_vmi full lower exit ${desired_lower_rc} (layout reject on 1024-wide vmul)" | tee -a "${LOG}" - tail -15 "${OUT}/desired_vmi_lower.log" | tee -a "${LOG}" + echo "FAIL: lowered_vpto.pto exit ${lowered_rc}" | tee -a "${LOG}" + tail -15 "${OUT}/lowered_vpto_ptoas.log" | tee -a "${LOG}" fi echo | tee -a "${LOG}" -# --- desired_vmi.pto: ptoas --emit-vpto (expect layout / lower fail) --- -echo "=== desired_vmi.pto (ptoas --emit-vpto) ===" | tee -a "${LOG}" -DESIRED_PTO="${FIXTURES}/desired_vmi.pto" -DESIRED_VPTO="${OUT}/desired_vmi.vpto" -set +e -if command -v ptoas >/dev/null 2>&1; then - ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ - --emit-vpto "${DESIRED_PTO}" -o "${DESIRED_VPTO}" \ - > "${OUT}/desired_vmi_ptoas.log" 2>&1 - desired_pto_rc=$? -else - echo "SKIP: ptoas not on PATH" | tee -a "${LOG}" - desired_pto_rc=127 -fi -set -e -if [ "${desired_pto_rc}" -eq 0 ] && [ -s "${DESIRED_VPTO}" ]; then - echo "PASS: desired_vmi.pto --emit-vpto (unexpected — gap may be closed)" | tee -a "${LOG}" -else - echo "FAIL: desired_vmi.pto --emit-vpto exit ${desired_pto_rc}" | tee -a "${LOG}" - tail -15 "${OUT}/desired_vmi_ptoas.log" | tee -a "${LOG}" -fi -echo | tee -a "${LOG}" +# isolate stages=2 / block_k=512 — dual-buffer OK; separates stages from wide-tile fail +ptodsl_py "isolate_stages2_blockk512_vmi.py" +emit_vpto "isolate_stages2_blockk512_vmi.pto" "pass" + +# isolate stages=1 / block_k=1024 — wide tile layout reject +ptodsl_py "isolate_stages1_blockk1024_vmi.py" +emit_vpto "isolate_stages1_blockk1024_vmi.pto" "fail" -# --- target_mi.pto: emit-vpto then LLVM IR -> bisheng (expect bisheng crash) --- -echo "=== target_mi.pto (ptoas --emit-vpto) ===" | tee -a "${LOG}" -TARGET_MI="${FIXTURES}/target_mi.pto" -TARGET_VPTO="${OUT}/target_mi.vpto" -LLVM_IR="${OUT}/target_mi.ll" -OBJ="${OUT}/target_mi.o" +# desired stages=2 / block_k=1024 — stages+wide (layout reject) +ptodsl_py "desired_vmi.py" +emit_vpto "desired_vmi.pto" "fail" +# cast-back cells (f8->bf16); wide / dual-buffer cells reject today +ptodsl_py "cast_back_stages1_k512_vmi.py" +emit_vpto "cast_back_stages1_k512_vmi.pto" "fail" +ptodsl_py "cast_back_stages2_k1024_vmi.py" +emit_vpto "cast_back_stages2_k1024_vmi.pto" "fail" + + +# target_mi: emit OK, bisheng object often crashes +echo "=== target_mi.pto (emit-vpto + LLVM IR + bisheng .o) ===" | tee -a "${LOG}" set +e -if command -v ptoas >/dev/null 2>&1; then - ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ - --emit-vpto "${TARGET_MI}" -o "${TARGET_VPTO}" \ - > "${OUT}/target_mi_emit_vpto.log" 2>&1 - emit_vpto_rc=$? -else - echo "SKIP: ptoas not on PATH" | tee -a "${LOG}" - emit_vpto_rc=127 -fi +ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto "${FIXTURES}/target_mi.pto" -o "${OUT}/target_mi.vpto" \ + > "${OUT}/target_mi_emit_vpto.log" 2>&1 +emit_rc=$? set -e - -if [ "${emit_vpto_rc}" -eq 0 ] && [ -s "${TARGET_VPTO}" ]; then +if [ "${emit_rc}" -eq 0 ] && [ -s "${OUT}/target_mi.vpto" ]; then echo "PASS: target_mi.pto --emit-vpto" | tee -a "${LOG}" else - echo "FAIL: target_mi.pto --emit-vpto exit ${emit_vpto_rc}" | tee -a "${LOG}" - tail -20 "${OUT}/target_mi_emit_vpto.log" | tee -a "${LOG}" + echo "FAIL: target_mi.pto --emit-vpto exit ${emit_rc}" | tee -a "${LOG}" fi -echo | tee -a "${LOG}" - -echo "=== target_mi.pto (ptoas -> LLVM IR -> bisheng object) ===" | tee -a "${LOG}" set +e -if command -v ptoas >/dev/null 2>&1; then - ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ - --emit-vpto-llvm-ir "${TARGET_MI}" -o "${LLVM_IR}" \ - > "${OUT}/target_mi_ptoas.log" 2>&1 - pto_rc=$? -else - pto_rc=127 -fi +ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ + --emit-vpto-llvm-ir "${FIXTURES}/target_mi.pto" -o "${OUT}/target_mi.ll" \ + > "${OUT}/target_mi_ptoas.log" 2>&1 +ll_rc=$? set -e - -if [ "${pto_rc}" -eq 0 ]; then - echo "PASS: ptoas --emit-vpto-llvm-ir target_mi.pto" | tee -a "${LOG}" +if [ "${ll_rc}" -eq 0 ]; then + echo "PASS: target_mi.pto --emit-vpto-llvm-ir" | tee -a "${LOG}" set +e "${BISHENG}" -O2 -fPIC -std=c++17 --npu-arch="${NPU_ARCH}" -c -x ir \ - "${LLVM_IR}" -o "${OBJ}" > "${OUT}/target_mi_bisheng.log" 2>&1 + "${OUT}/target_mi.ll" -o "${OUT}/target_mi.o" > "${OUT}/target_mi_bisheng.log" 2>&1 b_rc=$? set -e - if [ "${b_rc}" -eq 0 ] && [ -s "${OBJ}" ]; then - echo "PASS: bisheng object compile (unexpected — gap may be closed)" | tee -a "${LOG}" + if [ "${b_rc}" -eq 0 ] && [ -s "${OUT}/target_mi.o" ]; then + echo "PASS: bisheng object (unexpected — gap may be closed)" | tee -a "${LOG}" else - echo "FAIL: bisheng object compile exit ${b_rc} (expected on vmi-v0.1.3)" | tee -a "${LOG}" - tail -10 "${OUT}/target_mi_bisheng.log" | tee -a "${LOG}" + echo "FAIL: bisheng object exit ${b_rc} (expected on vmi-v0.1.3)" | tee -a "${LOG}" + tail -8 "${OUT}/target_mi_bisheng.log" | tee -a "${LOG}" fi else - echo "FAIL: ptoas emit-vpto-llvm-ir exit ${pto_rc}" | tee -a "${LOG}" - tail -20 "${OUT}/target_mi_ptoas.log" | tee -a "${LOG}" -fi -echo | tee -a "${LOG}" - -# --- current_slow_vmi stages=1: lowered_vpto.pto is the working MI path --- -echo "=== current_slow_vmi.py (ptodsl -> MLIR frontend, stages=1 block_k=512) ===" | tee -a "${LOG}" -set +e -python3 "${FIXTURES}/current_slow_vmi.py" > "${OUT}/current_slow_vmi.mlir" 2> "${OUT}/current_slow_vmi_py.err" -slow_py_rc=$? -set -e -if [ "${slow_py_rc}" -eq 0 ]; then - echo "PASS: current_slow_vmi.py emitted MLIR (frontend only)" | tee -a "${LOG}" -else - echo "FAIL: current_slow_vmi.py compile exit ${slow_py_rc}" | tee -a "${LOG}" - tail -20 "${OUT}/current_slow_vmi_py.err" | tee -a "${LOG}" -fi -echo | tee -a "${LOG}" - -echo "=== lowered_vpto.pto (stages=1 chunked MI, ptoas -> LLVM IR) ===" | tee -a "${LOG}" -LOWERED_PTO="${FIXTURES}/lowered_vpto.pto" -SLOW_LL="${OUT}/lowered_vpto.ll" -set +e -if command -v ptoas >/dev/null 2>&1; then - ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ - --emit-vpto-llvm-ir "${LOWERED_PTO}" -o "${SLOW_LL}" \ - > "${OUT}/lowered_vpto_ptoas.log" 2>&1 - lowered_rc=$? -else - lowered_rc=127 -fi -set -e -if [ "${lowered_rc}" -eq 0 ] && [ -s "${SLOW_LL}" ]; then - echo "PASS: lowered_vpto.pto -> LLVM IR (stages=1 working MI path)" | tee -a "${LOG}" -else - echo "FAIL: lowered_vpto.pto emit-vpto-llvm-ir exit ${lowered_rc}" | tee -a "${LOG}" - tail -20 "${OUT}/lowered_vpto_ptoas.log" | tee -a "${LOG}" + echo "FAIL: target_mi llvm-ir exit ${ll_rc}" | tee -a "${LOG}" fi echo | tee -a "${LOG}" diff --git a/docs/repro/pipelining/scripts/env.sh b/docs/repro/pipelining/scripts/env.sh index 33fff862d4..e0b682a9d9 100755 --- a/docs/repro/pipelining/scripts/env.sh +++ b/docs/repro/pipelining/scripts/env.sh @@ -14,8 +14,19 @@ if [ -n "${PTOAS_ROOT}" ] && [ -f "${PTOAS_ROOT}/scripts/ptoas_env.sh" ]; then export PTOAS_ENV_SKIP_SMOKE_TEST="${PTOAS_ENV_SKIP_SMOKE_TEST:-1}" # shellcheck disable=SC1091 source "${PTOAS_ROOT}/scripts/ptoas_env.sh" - [ -d "${PTOAS_ROOT}/build/tools/ptoas" ] && export PATH="${PTOAS_ROOT}/build/tools/ptoas:${PATH}" + # Prefer build/tools wrappers over install/bin (install wrappers often miss MLIR libs). [ -d "${PTOAS_ROOT}/install/bin" ] && export PATH="${PTOAS_ROOT}/install/bin:${PATH}" + [ -d "${PTOAS_ROOT}/build/tools/ptoas" ] && export PATH="${PTOAS_ROOT}/build/tools/ptoas:${PATH}" +fi + +# Optional: point at a known-good ptoas tree (e.g. built at tag vmi-v0.1.3). +if [ -n "${PTOAS_TOOLS_ROOT:-}" ]; then + [ -d "${PTOAS_TOOLS_ROOT}/build/tools/ptoas" ] && export PATH="${PTOAS_TOOLS_ROOT}/build/tools/ptoas:${PATH}" +fi + +# Prefer in-tree ptodsl when present (else site-packages). +if [ -d "${PTOAS_ROOT}/ptodsl" ]; then + export PYTHONPATH="${PTOAS_ROOT}/ptodsl${PYTHONPATH:+:${PYTHONPATH}}" fi export ASCEND_NPU_ARCH="${ASCEND_NPU_ARCH:-dav-3510}" From 8319463d1b65a2873889a5588ad71410eb2bc9ee Mon Sep 17 00:00:00 2001 From: learning-chip Date: Tue, 28 Jul 2026 19:10:33 +0800 Subject: [PATCH 3/3] Narrow pipelining repro to kernel-agnostic stages gap MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Drop cast-back / f8 product fixtures; keep stages×block_k isolate and the minimal CCE stages=2 device check. Point e2e MHC/quant BW study elsewhere. --- docs/repro/pipelining/README.md | 40 ++++++------ .../fixtures/cast_back_stages1_k512_vmi.pto | 47 -------------- .../fixtures/cast_back_stages1_k512_vmi.py | 32 ---------- .../fixtures/cast_back_stages2_k1024_vmi.pto | 62 ------------------- .../fixtures/cast_back_stages2_k1024_vmi.py | 41 ------------ .../fixtures/reference_asc_cce_cast_back.asc | 57 ----------------- .../fixtures/reference_asc_cce_f8.asc | 59 ------------------ docs/repro/pipelining/scripts/check_stages.sh | 19 +----- 8 files changed, 25 insertions(+), 332 deletions(-) delete mode 100644 docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.pto delete mode 100644 docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.py delete mode 100644 docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.pto delete mode 100644 docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.py delete mode 100644 docs/repro/pipelining/fixtures/reference_asc_cce_cast_back.asc delete mode 100644 docs/repro/pipelining/fixtures/reference_asc_cce_f8.asc diff --git a/docs/repro/pipelining/README.md b/docs/repro/pipelining/README.md index 9dd400ea20..dc7c4f800d 100644 --- a/docs/repro/pipelining/README.md +++ b/docs/repro/pipelining/README.md @@ -1,11 +1,14 @@ -# pipelining — VMI Persistent stages vs CCE ping-pong repro +# pipelining — VMI Persistent stages vs CCE ping-pong -Self-contained compile (+ optional on-device CCE) repro for long-K Persistent -quant / dequant and residual-mix-style scale: production wants stages=2 and -`block_k > 512` so MTE and vector stay overlapped. Ascend CCE can express and -**run** stages=2 ping-pong; VMI on tag **vmi-v0.1.3** cannot reach the same -wide-tile schedule — layout assignment rejects 1024-wide tiles, and the -hand-written MI path crashes bisheng even after ptoas emits LLVM IR. +Minimal, **kernel-agnostic** compile (+ optional on-device CCE) repro for the +PTOAS gap: long-K Persistent loops want `stages=2` and `block_k > 512` so MTE +and vector stay overlapped. Ascend CCE can express and **run** stages=2 +ping-pong; VMI on tag **vmi-v0.1.3** cannot reach the same wide-tile schedule +— layout assignment rejects 1024-wide tiles, and the hand-written MI path +crashes bisheng even after ptoas emits LLVM IR. + +Product-shaped MHC / quant patterns and ASC↔CCE **bandwidth** live in the +companion study under `vmi-demo-pipeline/pipeline_study` (separate repo). ## Pin @@ -26,16 +29,16 @@ export PTOAS_TOOLS_ROOT=/path/to/PTOAS-built-at-vmi-v0.1.3 | Path | stages | block_k | Fault mode | |------|-------:|--------:|------------| -| CCE (`fixtures/reference_asc_cce*.asc`, `device/scale_stages2.asc`) | 2 | 512–1024 | **OK** — compile; device host numerical OK | +| CCE (`fixtures/reference_asc_cce.asc`, `device/scale_stages2.asc`) | 2 | 512–1024 | **OK** — compile; device host numerical OK | | VMI stages=1 / 512 (`current_slow_vmi.*`) | 1 | 512 | **emit-vpto OK** | | VMI stages=2 / 512 (`isolate_stages2_blockk512_vmi.*`) | 2 | 512 | **emit-vpto OK** (dual-buffer alone is fine) | | VMI stages=1 / 1024 (`isolate_stages1_blockk1024_vmi.*`) | 1 | 1024 | **Layout reject** on 1024-wide vmul | | VMI stages=2 / 1024 (`desired_vmi.*`) | 2 | 1024 | **Layout reject** (stages+wide) | -| Cast-back VMI (`cast_back_stages*_*.*`) | 1–2 | 512–1024 | **Layout reject** on f8↔bf16 Persistent | | VMI target MI (`target_mi.pto`) | 2 | 1024 | **Bisheng crash** after LLVM IR emit | -So the blocker for the desired quant schedule is **wide `block_k=1024` layout**, -not dual-buffer stages alone. Cast-back Persistent hits the same layout wall. +The blocker for the desired schedule is **wide `block_k=1024` layout**, not +dual-buffer stages alone. Fixtures use a trivial scale body so the failure +stays schedule/layout — not a product kernel. ## Check results (vmi-v0.1.3 tools, 2026-07-28) @@ -47,20 +50,18 @@ source scripts/env.sh | Step | Result | |------|--------| -| `reference_asc_cce.asc` / `_f8` / `_cast_back` | **PASS** bisheng `--cce-aicore-only` | +| `reference_asc_cce.asc` | **PASS** bisheng `--cce-aicore-only` | | `device/scale_stages2.asc` aicore object | **PASS** | | `current_slow_vmi.pto` `--emit-vpto` | **PASS** | | `isolate_stages2_blockk512_vmi.pto` | **PASS** | | `isolate_stages1_blockk1024_vmi.pto` | **FAIL** layout | | `desired_vmi.pto` | **FAIL** layout | -| `cast_back_stages1_k512_vmi.pto` | **FAIL** layout | -| `cast_back_stages2_k1024_vmi.pto` | **FAIL** layout | | `lowered_vpto.pto` → LLVM IR | **PASS** | | `target_mi.pto` → emit + LLVM IR | **PASS**; bisheng `.o` **FAIL** | Full log: `sim_outputs/check_stages/compile_results.txt` -## On-device CCE +## On-device CCE (correctness only) ```bash cd device @@ -74,14 +75,17 @@ Measured 2026-07-28: scale_stages2: N=2048 mism=0 max_err=0 OK ``` -CCE stages=2 runs with dual-buffer HardEvent overlap. VMI stages=2 / `block_k=1024` -does not compile through layout — that is the PTOAS gap this package isolates. +CCE stages=2 runs with dual-buffer HardEvent overlap. VMI stages=2 / +`block_k=1024` does not compile through layout — that is the PTOAS gap. + +For e2e bandwidth (stages=1 vs stages=2, MHC / quant-shaped bodies), see +`pipeline_study` in the vmi-demo-pipeline repo. ## Layout ``` pipelining/ - fixtures/ CCE twins, VMI isolate + cast-back cells, target_mi + fixtures/ CCE shell, VMI stages×block_k isolate, target_mi device/ AscendC host for stages=2 CCE numerical check scripts/ env.sh, check_stages.sh ``` diff --git a/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.pto b/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.pto deleted file mode 100644 index 8712b03b56..0000000000 --- a/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.pto +++ /dev/null @@ -1,47 +0,0 @@ -// 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. - -// Cast-back Persistent VMI: stages=1, tile_k=512 (f8 -> bf16 * scale). - -module attributes { - pto.target_arch = "a5", - pto.kernel_kind = #pto.kernel_kind, - pto.pipeline.stages = 1 : i32, - pto.block_k = 512 : i32 -} { - func.func @persistent_cast_back_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_f8 = pto.vmi.vload %x_ub[%off] - : !pto.ptr -> !pto.vmi.vreg<512xf8E4M3FN> - %x_f32 = pto.vmi.vcvt %x_f8 - : !pto.vmi.vreg<512xf8E4M3FN> -> !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 - : !pto.vmi.vreg<512xf32> -> !pto.vmi.vreg<512xbf16> - pto.vmi.vstore %y, %y_ub[%off] - : !pto.vmi.vreg<512xbf16>, !pto.ptr - } - } - return - } -} diff --git a/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.py b/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.py deleted file mode 100644 index faff9e5b67..0000000000 --- a/docs/repro/pipelining/fixtures/cast_back_stages1_k512_vmi.py +++ /dev/null @@ -1,32 +0,0 @@ -# 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. - -"""Cast-back Persistent: e4m3 -> f32 * scale -> bf16, stages=1, tile_k=512.""" - -from ptodsl import pto - - -@pto.jit(target="a5", backend="vpto", mode="explicit") -def persistent_cast_back_k512(*, BLOCK_COUNT: pto.const_expr = 4): - x_ub = pto.alloc_tile(shape=[1, 512], dtype=pto.f8e4m3) - y_ub = pto.alloc_tile(shape=[1, 512], dtype=pto.bf16) - 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_f8 = pto.vmi.vload(x_ub.as_ptr(), off, size=512) - x_f32 = pto.vmi.vcvt(x_f8, pto.f32) - scaled = pto.vmi.vmul(x_f32, scale) - y = pto.vmi.vcvt(scaled, pto.bf16) - pto.vmi.vstore(y, y_ub.as_ptr(), off) - - -if __name__ == "__main__": - print(persistent_cast_back_k512.compile().mlir_text()) diff --git a/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.pto b/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.pto deleted file mode 100644 index ae01c38e38..0000000000 --- a/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.pto +++ /dev/null @@ -1,62 +0,0 @@ -// 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. - -// Cast-back Persistent VMI: stages=2, tile_k=1024 (f8 -> bf16 * scale). - -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_cast_back_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_f8 = pto.vmi.vload %x_ub[%off] - : !pto.ptr -> !pto.vmi.vreg<1024xf8E4M3FN> - %x_f32 = pto.vmi.vcvt %x_f8 - : !pto.vmi.vreg<1024xf8E4M3FN> -> !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 - : !pto.vmi.vreg<1024xf32> -> !pto.vmi.vreg<1024xbf16> - pto.vmi.vstore %y, %y_ub[%off] - : !pto.vmi.vreg<1024xbf16>, !pto.ptr - } - } - return - } -} diff --git a/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.py b/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.py deleted file mode 100644 index 8546d041e1..0000000000 --- a/docs/repro/pipelining/fixtures/cast_back_stages2_k1024_vmi.py +++ /dev/null @@ -1,41 +0,0 @@ -# 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. - -"""Cast-back Persistent: e4m3 -> f32 * scale -> bf16, stages=2, tile_k=1024.""" - -from ptodsl import pto - - -@pto.jit(target="a5", backend="vpto", mode="explicit") -def persistent_cast_back_k1024(*, BLOCK_COUNT: pto.const_expr = 4): - x_ub0 = pto.alloc_tile(shape=[1, 1024], dtype=pto.f8e4m3) - x_ub1 = pto.alloc_tile(shape=[1, 1024], dtype=pto.f8e4m3) - y_ub0 = pto.alloc_tile(shape=[1, 1024], dtype=pto.bf16) - y_ub1 = pto.alloc_tile(shape=[1, 1024], dtype=pto.bf16) - 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, y_ptr = x_ub1.as_ptr(), y_ub1.as_ptr() - else: - x_ptr, y_ptr = x_ub0.as_ptr(), y_ub0.as_ptr() - x_f8 = pto.vmi.vload(x_ptr, off, size=1024) - x_f32 = pto.vmi.vcvt(x_f8, pto.f32) - scaled = pto.vmi.vmul(x_f32, scale) - y = pto.vmi.vcvt(scaled, pto.bf16) - pto.vmi.vstore(y, y_ptr, off) - - -if __name__ == "__main__": - try: - print(persistent_cast_back_k1024.compile().mlir_text()) - except Exception as exc: # noqa: BLE001 - print(f"emit/lower may fail (layout / stages): {exc}") diff --git a/docs/repro/pipelining/fixtures/reference_asc_cce_cast_back.asc b/docs/repro/pipelining/fixtures/reference_asc_cce_cast_back.asc deleted file mode 100644 index f68534594c..0000000000 --- a/docs/repro/pipelining/fixtures/reference_asc_cce_cast_back.asc +++ /dev/null @@ -1,57 +0,0 @@ -// CCE twin for cast-back Persistent schedule (stages=2 / block_k=1024). -// HardEvent ping-pong shell aligned to quant CCE; ends in bf16 store. -#include "kernel_operator.h" - -__simd_vf__ inline void vf_cast_back_scale(__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_cast_back_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); - 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_cast_back_scale(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/repro/pipelining/fixtures/reference_asc_cce_f8.asc b/docs/repro/pipelining/fixtures/reference_asc_cce_f8.asc deleted file mode 100644 index 0a3d0d480b..0000000000 --- a/docs/repro/pipelining/fixtures/reference_asc_cce_f8.asc +++ /dev/null @@ -1,59 +0,0 @@ -// CCE twin for long-K Persistent quant schedule (stages=2 / block_k=1024). -// Same HardEvent ping-pong + UNPK/vmul/PK body as reference_asc_cce.asc; kernel -// name marks the f8-store consumer. bisheng legality of the stages=2 shell is -// what this file proves (f8 PK4 store shapes live in target_mi.pto). -#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_quant_stages2_f8(__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); - 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/repro/pipelining/scripts/check_stages.sh b/docs/repro/pipelining/scripts/check_stages.sh index 08f8e0643b..33ec55def5 100755 --- a/docs/repro/pipelining/scripts/check_stages.sh +++ b/docs/repro/pipelining/scripts/check_stages.sh @@ -1,5 +1,5 @@ #!/usr/bin/env bash -# Compile-check CCE stages=2 vs VMI stages×block_k isolate + cast-back matrix. +# Compile-check CCE stages=2 vs VMI stages×block_k isolate matrix (kernel-agnostic). set -euo pipefail SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" KERNEL_ROOT="$(cd "${SCRIPT_DIR}/.." && pwd)" @@ -96,10 +96,8 @@ ptodsl_py() { echo | tee -a "${LOG}" } -# --- CCE baselines --- +# --- CCE baseline --- compile_asc "reference_asc_cce.asc" -compile_asc "reference_asc_cce_f8.asc" -compile_asc "reference_asc_cce_cast_back.asc" echo "=== device/scale_stages2.asc (bisheng --cce-aicore-only) ===" | tee -a "${LOG}" set +e @@ -120,11 +118,10 @@ else fi echo | tee -a "${LOG}" - # --- VMI matrix: stages × block_k --- -# stages=1 block_k=512 (current_slow) — emit OK on vmi-v0.1.3; lowered_vpto is chunked MI twin ptodsl_py "current_slow_vmi.py" emit_vpto "current_slow_vmi.pto" "pass" + echo "=== lowered_vpto.pto (stages=1 chunked MI -> LLVM IR) ===" | tee -a "${LOG}" set +e ptoas --cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto \ @@ -140,25 +137,15 @@ else fi echo | tee -a "${LOG}" -# isolate stages=2 / block_k=512 — dual-buffer OK; separates stages from wide-tile fail ptodsl_py "isolate_stages2_blockk512_vmi.py" emit_vpto "isolate_stages2_blockk512_vmi.pto" "pass" -# isolate stages=1 / block_k=1024 — wide tile layout reject ptodsl_py "isolate_stages1_blockk1024_vmi.py" emit_vpto "isolate_stages1_blockk1024_vmi.pto" "fail" -# desired stages=2 / block_k=1024 — stages+wide (layout reject) ptodsl_py "desired_vmi.py" emit_vpto "desired_vmi.pto" "fail" -# cast-back cells (f8->bf16); wide / dual-buffer cells reject today -ptodsl_py "cast_back_stages1_k512_vmi.py" -emit_vpto "cast_back_stages1_k512_vmi.pto" "fail" -ptodsl_py "cast_back_stages2_k1024_vmi.py" -emit_vpto "cast_back_stages2_k1024_vmi.pto" "fail" - - # target_mi: emit OK, bisheng object often crashes echo "=== target_mi.pto (emit-vpto + LLVM IR + bisheng .o) ===" | tee -a "${LOG}" set +e