VMI/MI 2-stage double buffer compilation error vs ref CCE

docs/repro/pipelining/.gitignore

@@ -0,0 +1,12 @@
+build/
+sim_outputs/
+device/build/
+
+*.so
+*.o
+log_ca/
+instr.bin
+__pycache__/
+*.pyc
+.pytest_cache/
+*.egg-info/

docs/repro/pipelining/README.md

@@ -0,0 +1,96 @@
+# pipelining — VMI Persistent stages vs CCE ping-pong
+
+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
+
+| 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`, `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) |
+| VMI target MI (`target_mi.pto`) | 2 | 1024 | **Bisheng crash** after LLVM IR emit |
+
+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)
+
+```bash
+source scripts/env.sh
+# optional: export PTOAS_TOOLS_ROOT=...
+./scripts/check_stages.sh
+```
+
+| Step | Result |
+|------|--------|
+| `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 |
+| `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 (correctness only)
+
+```bash
+cd device
+./run.sh                 # build + launch stages=2 f32 scale, check max_err
+MSOPPROF=1 ./run.sh      # optional msopprof wrap if installed
+```
+
+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.
+
+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 shell, VMI stages×block_k isolate, target_mi
+  device/       AscendC host for stages=2 CCE numerical check
+  scripts/      env.sh, check_stages.sh
+```
+
+## 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`.

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"

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 <cmath>
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+
+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<float>(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;
+}

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}"

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<AscendC::HardEvent::MTE2_V>(2);
+  AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(2);
+
+  AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(0);
+  AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(1);
+  AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(0);
+  AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(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<AscendC::HardEvent::V_MTE2>(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<AscendC::HardEvent::MTE2_V>(slot);
+    AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(slot);
+    AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(slot);
+
+    vf_scale_f32(x_ub, sc_ub, o_ub);
+
+    AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(slot);
+    AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(slot);
+    AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(slot);
+    copy_ubuf_to_gm_align_v2((__gm__ void*)(out + w * kBlockK), (__ubuf__ void*)o_ub, 0, 1, kBytes,
+                             4, kBytes, kBytes);
+    AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(slot);
+  }
+
+  AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(0);
+  AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(1);
+  AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(0);
+  AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(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);
+}

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<vector>
+  // stages = 1
+  // block_k = 512
+} {
+  func.func @persistent_quant_tile_k512(
+      %x_ub: !pto.ptr<bf16, ub>,
+      %y_ub: !pto.ptr<f8E4M3FN, ub>,
+      %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<bf16, ub> -> !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<f8E4M3FN, ub>
+      }
+    }
+    return
+  }
+}

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())