VMI bug reproducer: L=8 ui8 to ui16 widen

docs/repro/l8_widen/.gitignore

@@ -0,0 +1 @@
+outputs/

docs/repro/l8_widen/README.md

@@ -0,0 +1,80 @@
+# L=8 ui8→ui16 VMI widen repro
+
+Fused quant→dequant needs a compact **L=8** unsigned widen (`ui8→ui16`) on the
+scale-factor path. At PTOAS tag **vmi-v0.1.3**, that shape lowers through MI
+and CCE, but VMI legalization leaves residual `pto.vmi.extui` and fails. Kernels
+that need this step (for example fused `per_token_cast_and_cast_back`) fall back
+to composing separate cast kernels or using larger L=256 tiles until L=8 VMI
+widen is fixed.
+
+Remote: https://github.com/learning-chip/PTOAS.git (branch `zjw/l8_widen_issue`)
+
+## What breaks
+
+| Fixture | Role | vmi-v0.1.3 result |
+|---------|------|-------------------|
+| `fixtures/current_slow_vmi.pto` | L=256 VMI widen (works today) | PASS |
+| `fixtures/desired_vmi.pto` | L=8 group-slot VMI widen (needed) | **FAIL** — illegal `extui` |
+| `fixtures/target_mi.pto` | Small-L MI shape (`PAT_VL8`, `vcvt` EVEN) | PASS |
+| `fixtures/reference_asc_cce.asc` | Hand-written CCE baseline | PASS |
+
+L=256 widen reference: `test/lit/vmi_new/vmi_to_vpto_integer_casts.pto`.
+
+## Layout
+
+```
+l8_widen/
+  fixtures/     IR and ASC inputs copied from gap 06
+  scripts/      env setup and compile check
+  outputs/      generated by check script (gitignored)
+```
+
+## Prerequisites
+
+- CANN 9.0 (`bisheng`, Ascend headers)
+- PTOAS built at tag **vmi-v0.1.3** on branch `zjw/l8_widen_issue`
+  - needs `pto-test-opt`, `ptoas`, and `ptoas_runtime_staging` in `build/`
+
+## Run
+
+From this directory (`docs/repro/l8_widen/`):
+
+```bash
+source scripts/env.sh
+./scripts/check_l8_widen.sh
+```
+
+The script lowers/compiles each fixture and prints PASS or FAIL. Full log:
+`outputs/check_l8_widen/compile_results.txt`.
+
+## Check results (2026-07-27, vmi-v0.1.3)
+
+```
+=== current_slow_vmi.pto (VMI L=256 ui8→ui16 lower) ===
+PASS: current_slow_vmi.pto lowers to VPTO (L=256 vcvt EVEN/ODD)
+
+=== desired_vmi.pto (VMI L=8 ui8→ui16 lower) ===
+FAIL: desired_vmi.pto lower exit 1 (residual pto.vmi.extui on vmi-v0.1.3)
+  error: failed to legalize operation 'pto.vmi.extui' that was explicitly marked illegal
+  error: VMI-RESIDUAL-OP: failed to convert all VMI ops/types to VPTO
+
+=== target_mi.pto (MI PAT_VL8 vcvt EVEN + ui16 store) ===
+PASS: ptoas --emit-vpto target_mi.pto
+PASS: ptoas --emit-vpto-llvm-ir target_mi.pto
+PASS: ptoas device object compile for target_mi.pto
+
+=== reference_asc_cce.asc (bisheng --cce-aicore-only) ===
+PASS: reference_asc_cce.asc -> non-empty .o (CCE PAT_VL8 widen baseline)
+```
+
+## Done when
+
+VMI L=8 `ui8→ui16` in `fixtures/desired_vmi.pto` lowers to the same MI ops as
+`fixtures/target_mi.pto` (no residual `extui`), so fused quant round-trip can
+match ASC bandwidth without compose fallback or L=256 workarounds.
+
+## Pin
+
+- PTOAS tag: **vmi-v0.1.3**
+- Branch: `zjw/l8_widen_issue`
+- `PTOAS_ROOT` defaults to the repo root (`../../..` from this repro directory).

docs/repro/l8_widen/fixtures/current_slow_vmi.pto

@@ -0,0 +1,20 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// Current slow path: L=256 ui8→ui16 vcvt (see test/lit/vmi_new/vmi_to_vpto_integer_casts.pto).
+
+module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind<vector>} {
+  func.func @ui8_to_ui16_l256(
+      %input: !pto.vmi.vreg<256xui8, #pto.vmi.layout<contiguous>>)
+      -> !pto.vmi.vreg<256xui16, #pto.vmi.layout<deinterleaved = 2>> {
+    %wide = pto.vmi.vcvt %input
+        : !pto.vmi.vreg<256xui8, #pto.vmi.layout<contiguous>>
+        -> !pto.vmi.vreg<256xui16, #pto.vmi.layout<deinterleaved = 2>>
+    return %wide : !pto.vmi.vreg<256xui16, #pto.vmi.layout<deinterleaved = 2>>
+  }
+}

docs/repro/l8_widen/fixtures/current_slow_vmi.py

@@ -0,0 +1,27 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+"""Python DSL form of ``current_slow_vmi.pto`` — L=256 ui8→ui16 ``vcvt``.
+
+Canonical IR: ``current_slow_vmi.pto``.
+"""
+
+from ptodsl import pto
+
+
+@pto.jit(target="a5", backend="vpto", mode="explicit")
+def ui8_to_ui16_l256():
+    ub = pto.alloc_tile(shape=[1, 256], dtype=pto.ui8)
+    off = pto.const(0, dtype=pto.index)
+    inp = pto.vmi.vload(ub.as_ptr(), off, size=256)
+    wide = pto.vmi.vcvt(inp, pto.ui16)
+    _ = wide
+
+
+if __name__ == "__main__":
+    print(ui8_to_ui16_l256.compile().mlir_text())

docs/repro/l8_widen/fixtures/desired_vmi.pto

@@ -0,0 +1,20 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// Idiomatic small-L widen: L=8 ui8→ui16 (needed for fused quant→dequant UE8M0).
+
+module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind<vector>} {
+  func.func @ui8_to_ui16_l8(
+      %input: !pto.vmi.vreg<8xui8, #pto.vmi.layout<num_groups = 8, slots = 1>>)
+      -> !pto.vmi.vreg<8xui16, #pto.vmi.layout<num_groups = 8, slots = 1>> {
+    %wide = pto.vmi.vcvt %input
+        : !pto.vmi.vreg<8xui8, #pto.vmi.layout<num_groups = 8, slots = 1>>
+        -> !pto.vmi.vreg<8xui16, #pto.vmi.layout<num_groups = 8, slots = 1>>
+    return %wide : !pto.vmi.vreg<8xui16, #pto.vmi.layout<num_groups = 8, slots = 1>>
+  }
+}

docs/repro/l8_widen/fixtures/desired_vmi.py

@@ -0,0 +1,31 @@
+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+# CANN Open Software License Agreement Version 2.0 (the "License").
+# Please refer to the License for details. You may not use this file except in compliance with the License.
+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+# See LICENSE in the root of the software repository for the full text of the License.
+
+"""Python DSL form of ``desired_vmi.pto`` — L=8 ui8→ui16 widen.
+
+Canonical IR: ``desired_vmi.pto``. VMI legalization leaves residual ``extui``.
+"""
+
+from ptodsl import pto
+
+
+@pto.jit(target="a5", backend="vpto", mode="explicit")
+def ui8_to_ui16_l8():
+    # 32B-aligned tile; load/convert only L=8 lanes (compact SF widen).
+    ub = pto.alloc_tile(shape=[1, 32], dtype=pto.ui8)
+    off = pto.const(0, dtype=pto.index)
+    inp = pto.vmi.vload(ub.as_ptr(), off, size=8)
+    wide = pto.vmi.vcvt(inp, pto.ui16)
+    _ = wide
+
+
+if __name__ == "__main__":
+    try:
+        print(ui8_to_ui16_l8.compile().mlir_text())
+    except Exception as exc:  # noqa: BLE001 — document the gap
+        print(f"desired emit/lower may fail (L=8 widen): {exc}")

docs/repro/l8_widen/fixtures/lowered_vpto.pto

@@ -0,0 +1,9 @@
+module attributes {pto.kernel_kind = #pto.kernel_kind<vector>, pto.target_arch = "a5"} {
+  func.func @ui8_to_ui16_l256(%arg0: !pto.vreg<256xui8>) -> (!pto.vreg<128xui16>, !pto.vreg<128xui16>) {
+    %0 = pto.pset_b8 "PAT_ALL" : !pto.mask<b8>
+    %1 = pto.vcvt %arg0, %0 {part = "EVEN"} : !pto.vreg<256xui8>, !pto.mask<b8> -> !pto.vreg<128xui16>
+    %2 = pto.vcvt %arg0, %0 {part = "ODD"} : !pto.vreg<256xui8>, !pto.mask<b8> -> !pto.vreg<128xui16>
+    return %1, %2 : !pto.vreg<128xui16>, !pto.vreg<128xui16>
+  }
+}
+

docs/repro/l8_widen/fixtures/reference_asc_cce.asc

@@ -0,0 +1,21 @@
+// ASC/CCE working baseline for gap 06: small-L ui8→ui16 widen (PART_EVEN).
+// Compiles with bisheng --cce-aicore-only (see README).
+#include "kernel_operator.h"
+
+__simd_vf__ inline void vf_ui8_widen_l8(__ubuf__ uint8_t* src, __ubuf__ uint16_t* dst) {
+  // L=8-shaped mask: the compact SF widen VMI wants on the fused path.
+  vector_bool v = pset_b16(PAT_VL8);
+  vector_u8 x;
+  vlds(x, src, 0, NORM);
+  vector_u16 y;
+  vcvt(y, x, v, PART_EVEN, MODE_ZEROING);
+  vsts(y, dst, 0, NORM_B16, v);
+}
+
+extern "C" __global__ __vector__ void ref_gap06_ui8_widen(__gm__ uint8_t* /*gm_in*/,
+                                                         __gm__ uint16_t* /*gm_out*/) {
+  AscendC::InitSocState();
+  __ubuf__ uint8_t* ub8 = (__ubuf__ uint8_t*)0;
+  __ubuf__ uint16_t* ub16 = (__ubuf__ uint16_t*)(ub8 + 256);
+  vf_ui8_widen_l8(ub8, ub16);
+}

docs/repro/l8_widen/fixtures/target_mi.pto

@@ -0,0 +1,29 @@
+// Copyright (c) 2026 Huawei Technologies Co., Ltd.
+// This program is free software, you can redistribute it and/or modify it under the terms and conditions of
+// CANN Open Software License Agreement Version 2.0 (the "License").
+// Please refer to the License for details. You may not use this file except in compliance with the License.
+// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
+// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
+// See LICENSE in the root of the software repository for the full text of the License.
+
+// Desired MI: small-L (L=8) ui8→ui16 widen via vcvt PART_EVEN, then ui16 store.
+// Fused quant→dequant needs this shape; VMI currently illegalizes to extui.
+
+module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind<vector>} {
+  func.func @ui8_to_ui16_widen_l8_mi(
+      %src: !pto.ptr<ui8, ub>,
+      %dst: !pto.ptr<ui16, ub>,
+      %off: index) attributes {pto.kernel} {
+    pto.vecscope {
+      %m8 = pto.pset_b8 "PAT_VL8" : !pto.mask<b8>
+      %in = pto.vlds %src[%off] {dist = "NORM"}
+          : !pto.ptr<ui8, ub> -> !pto.vreg<256xui8>
+      %even = pto.vcvt %in, %m8 {part = "EVEN"}
+          : !pto.vreg<256xui8>, !pto.mask<b8> -> !pto.vreg<128xui16>
+      %m16 = pto.pset_b16 "PAT_VL8" : !pto.mask<b16>
+      pto.vsts %even, %dst[%off], %m16
+          : !pto.vreg<128xui16>, !pto.ptr<ui16, ub>, !pto.mask<b16>
+    }
+    return
+  }
+}

docs/repro/l8_widen/scripts/check_l8_widen.sh

@@ -0,0 +1,131 @@
+#!/usr/bin/env bash
+# Compile/lower L=8 ui8→ui16 widen fixtures; print PASS/FAIL for each path.
+set -euo pipefail
+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+REPRO_ROOT="$(cd "${SCRIPT_DIR}/.." && pwd)"
+FIXTURES="${REPRO_ROOT}/fixtures"
+OUT="${REPRO_ROOT}/outputs/check_l8_widen"
+LOG="${OUT}/compile_results.txt"
+VMI_PASSES=(-vmi-lower-unified-to-legacy -vmi-to-vpto)
+
+# shellcheck disable=SC1091
+source "${SCRIPT_DIR}/env.sh"
+
+mkdir -p "${OUT}"
+: > "${LOG}"
+
+BISHENG="${BISHENG:-${ASCEND_HOME_PATH}/tools/bisheng_compiler/bin/bisheng}"
+ASCEND="${ASCEND_HOME_PATH}"
+NPU_ARCH="${ASCEND_NPU_ARCH:-dav-3510}"
+PTOAS_FLAGS=(--cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto)
+
+if ! command -v pto-test-opt >/dev/null 2>&1; then
+  echo "ERROR: pto-test-opt not found; build PTOAS at tag vmi-v0.1.3 first." | tee -a "${LOG}"
+  exit 1
+fi
+if ! command -v ptoas >/dev/null 2>&1; then
+  echo "ERROR: ptoas not found; build PTOAS at tag vmi-v0.1.3 first." | tee -a "${LOG}"
+  exit 1
+fi
+
+echo "l8_widen compile check — $(date -u +%Y-%m-%dT%H:%M:%SZ)" | tee -a "${LOG}"
+echo "PTOAS_ROOT=${PTOAS_ROOT}" | tee -a "${LOG}"
+echo "pto-test-opt=$(command -v pto-test-opt)" | tee -a "${LOG}"
+echo "ptoas=$(command -v ptoas)" | tee -a "${LOG}"
+echo | tee -a "${LOG}"
+
+# --- current_slow_vmi.pto: L=256 ui8→ui16 VMI (known-good widen shape) ---
+echo "=== current_slow_vmi.pto (VMI L=256 ui8→ui16 lower) ===" | tee -a "${LOG}"
+CURRENT_OUT="${OUT}/current_slow_vmi_vpto.pto"
+set +e
+pto-test-opt "${FIXTURES}/current_slow_vmi.pto" "${VMI_PASSES[@]}" -o "${CURRENT_OUT}" \
+  > "${OUT}/current_slow_vmi.log" 2>&1
+rc=$?
+set -e
+if [ "${rc}" -eq 0 ] && [ -s "${CURRENT_OUT}" ]; then
+  echo "PASS: current_slow_vmi.pto lowers to VPTO (L=256 vcvt EVEN/ODD)" | tee -a "${LOG}"
+else
+  echo "FAIL: current_slow_vmi.pto lower exit ${rc}" | tee -a "${LOG}"
+  tail -20 "${OUT}/current_slow_vmi.log" | tee -a "${LOG}"
+fi
+echo | tee -a "${LOG}"
+
+# --- desired_vmi.pto: L=8 ui8→ui16 VMI (illegalizes to extui on vmi-v0.1.3) ---
+echo "=== desired_vmi.pto (VMI L=8 ui8→ui16 lower) ===" | tee -a "${LOG}"
+DESIRED_OUT="${OUT}/desired_vmi_vpto.pto"
+set +e
+pto-test-opt "${FIXTURES}/desired_vmi.pto" "${VMI_PASSES[@]}" -o "${DESIRED_OUT}" \
+  > "${OUT}/desired_vmi.log" 2>&1
+rc=$?
+set -e
+if [ "${rc}" -eq 0 ] && [ -s "${DESIRED_OUT}" ]; then
+  echo "PASS: desired_vmi.pto lowers to VPTO (L=8 widen gap may be closed)" | tee -a "${LOG}"
+else
+  echo "FAIL: desired_vmi.pto lower exit ${rc} (residual pto.vmi.extui on vmi-v0.1.3)" | tee -a "${LOG}"
+  rg -n "extui|illegal|VMI-RESIDUAL" "${OUT}/desired_vmi.log" | tail -5 | tee -a "${LOG}" || tail -15 "${OUT}/desired_vmi.log" | tee -a "${LOG}"
+fi
+echo | tee -a "${LOG}"
+
+# --- target_mi.pto: small-L MI shape (ptoas lower + device object) ---
+echo "=== target_mi.pto (MI PAT_VL8 vcvt EVEN + ui16 store) ===" | tee -a "${LOG}"
+TARGET_VPTO="${OUT}/target_mi_vpto.pto"
+TARGET_LL="${OUT}/target_mi.ll"
+TARGET_OBJ="${OUT}/target_mi.o"
+
+set +e
+ptoas "${PTOAS_FLAGS[@]}" --emit-vpto "${FIXTURES}/target_mi.pto" -o "${TARGET_VPTO}" \
+  > "${OUT}/target_mi_emit_vpto.log" 2>&1
+vpto_rc=$?
+set -e
+if [ "${vpto_rc}" -ne 0 ]; then
+  echo "FAIL: ptoas --emit-vpto target_mi.pto exit ${vpto_rc}" | tee -a "${LOG}"
+  tail -20 "${OUT}/target_mi_emit_vpto.log" | tee -a "${LOG}"
+else
+  echo "PASS: ptoas --emit-vpto target_mi.pto" | tee -a "${LOG}"
+  set +e
+  ptoas "${PTOAS_FLAGS[@]}" --emit-vpto-llvm-ir "${FIXTURES}/target_mi.pto" -o "${TARGET_LL}" \
+    > "${OUT}/target_mi_emit_llvm.log" 2>&1
+  llvm_rc=$?
+  set -e
+  if [ "${llvm_rc}" -ne 0 ]; then
+    echo "FAIL: ptoas --emit-vpto-llvm-ir target_mi.pto exit ${llvm_rc}" | tee -a "${LOG}"
+    tail -20 "${OUT}/target_mi_emit_llvm.log" | tee -a "${LOG}"
+  else
+    echo "PASS: ptoas --emit-vpto-llvm-ir target_mi.pto" | tee -a "${LOG}"
+    set +e
+    ptoas "${PTOAS_FLAGS[@]}" "${FIXTURES}/target_mi.pto" -o "${TARGET_OBJ}" \
+      > "${OUT}/target_mi_ptoas.log" 2>&1
+    obj_rc=$?
+    set -e
+    if [ "${obj_rc}" -eq 0 ] && [ -s "${TARGET_OBJ}" ]; then
+      echo "PASS: ptoas device object compile for target_mi.pto" | tee -a "${LOG}"
+    else
+      echo "FAIL: ptoas device object compile exit ${obj_rc}" | tee -a "${LOG}"
+      tail -20 "${OUT}/target_mi_ptoas.log" | tee -a "${LOG}"
+    fi
+  fi
+fi
+echo | tee -a "${LOG}"
+
+# --- reference_asc_cce.asc: hand-written CCE baseline ---
+echo "=== reference_asc_cce.asc (bisheng --cce-aicore-only) ===" | tee -a "${LOG}"
+REF_OBJ="${OUT}/reference_asc_cce.o"
+set +e
+"${BISHENG}" -O2 -fPIC -std=c++17 --npu-arch="${NPU_ARCH}" --cce-aicore-only -c \
+  "${FIXTURES}/reference_asc_cce.asc" -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 PAT_VL8 widen baseline)" | 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}"
+echo "Full log: outputs/check_l8_widen/compile_results.txt" | tee -a "${LOG}"