VMI pad-8 vs CCE ONEPT reduce store issue

.gitignore

@@ -99,3 +99,12 @@ test/samples/**/npu_validation/*
 .planning/
 .work/
 .local/
+
+# local cannsim / build junk under docs/repro
+docs/repro/**/build/
+docs/repro/**/sim_outputs/
+docs/repro/**/*.so
+docs/repro/**/traces/
+docs/repro/**/*.zip
+docs/repro/**/route_table.txt
+docs/repro/**/rtb_debug.txt

docs/repro/store_pad8/.gitignore

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

docs/repro/store_pad8/README.md

@@ -0,0 +1,113 @@
+# store_pad8 — VMI pad-8 reduce store repro
+
+Self-contained cannsim microbench for PTOAS gap 03: after a cross-lane
+`vcadd`, CCE stores one f32 scalar per group with `vsts ... ONEPT_B32`.
+Product VMI kernels often broadcast the reduced value to 8 lanes (`vbrc`)
+and store under an 8-wide mask.
+
+This repro isolates that epilogue and compares **kernel formulations** on a
+built PTOAS pin (branch `zjw/store_pad_issue`, tag lineage **vmi-v0.1.3**).
+
+## What to compare
+
+| Path | Store shape | Notes |
+|------|-------------|-------|
+| CCE (`cce/`) | ONEPT / 1 scalar | Baseline — see `fixtures/reference_asc_cce.asc` |
+| VMI pad-8 (`vmi/store_pad8_vmi.py`) | pad-8 + mask8 | Product workaround |
+| VMI mask1 (`vmi/store_mask1_vmi.py`) | compact `[N]` + mask1 | Masked 1-lane store |
+| VMI **group1** (`vmi/store_group1_vmi.py`) | `vcadd(group=1)` + `vstore(group=1)` | Compact score → point store / `1PT_B32` |
+| VMI desired (`fixtures/desired_vmi.*`) | mask-1 after `vcadd` | Compile-only single-tile fixture |
+| VMI target MI (`fixtures/target_mi.pto`) | `dist = "1PT_B32"` | Hand-written MI; bisheng may still crash |
+
+**Best VMI path:** `group1`. It lowers to `vsts ... {dist = "1PT_B32"}` and
+matches CCE’s `vcadd` + scalar store shape in cannsim (see table below).
+
+## Layout
+
+```
+store_pad8/
+  cce/          CCE ONEPT kernel + ctypes launcher
+  vmi/          VMI pad-8 / mask1 / group1 kernels (ptodsl)
+  common/       torch runtime, golden ref, build helper, top launcher
+  fixtures/     IR dumps and reference ASC from gap 03
+  test/         correctness test
+  scripts/      env, cannsim runners, check_desired, run_three_way
+```
+
+## Prerequisites
+
+- CANN 9.0 simulator (`cannsim`, `bisheng`)
+- Built PTOAS from this repo (`./quick_install.sh` with
+  `-DPTOAS_ENABLE_WERROR=OFF` if needed for older gcc). LLVM/MLIR comes from
+  the docker image. Install the produced wheel so `ptoas` on `PATH` is the
+  wheel launcher (not only `build/tools/ptoas`).
+- `torch_npu` for cannsim host launch
+
+Push remote hint: https://github.com/learning-chip/PTOAS.git
+
+## Run
+
+From this directory (`docs/repro/store_pad8/`):
+
+```bash
+source scripts/env.sh
+
+# Four-way: CCE + VMI pad8 / mask1 / group1
+STORE_PAD8_CASE=large ./scripts/run_three_way.sh
+STORE_PAD8_CASE=small ./scripts/run_three_way.sh
+
+# Or one backend
+STORE_PAD8_CASE=large ./scripts/run_cannsim.sh cce
+STORE_PAD8_VARIANT=group1 STORE_PAD8_CASE=large \
+  ./scripts/run_sim.sh test/test_store_pad8.py sim_outputs/store_pad8_large_vmi_group1
+
+./scripts/check_desired.sh
+```
+
+Correctness only (no cannsim):
+
+```bash
+source scripts/env.sh
+TLVF_VMI_BACKEND=cce STORE_PAD8_CASE=large python3 test/test_store_pad8.py
+TLVF_VMI_BACKEND=vmi STORE_PAD8_VARIANT=group1 STORE_PAD8_CASE=large python3 test/test_store_pad8.py
+```
+
+## Measured RVEC (Ascend950 cannsim, 2026-07-29)
+
+Built PTOAS from this branch (wheel `ptoas 0.52`). All paths PASS
+(`maxDiff≈1e-6`).
+
+### N_ACC=20 (LARGE)
+
+| Backend | RVEC span | vs CCE | Top ops (approx.) |
+|---------|----------:|-------:|-------------------|
+| CCE ONEPT | **63** | 1.00× | `RV_VLDI=20, RV_VCADD=20, RV_VSTI=20, RV_PSET=1` |
+| VMI **group1** | **59** | **0.94×** | `RV_VCADD=20, RV_VSTI=20, RV_VLDI=16, RV_VLDS=4, RV_PSET=3` |
+| VMI mask1 | **66** | 1.05× | `RV_VCADD=20, RV_VADD=20, RV_VSTS=17, RV_VLDI=16, RV_VLDS=4` |
+| VMI pad8 | **99** | 1.57× | `RV_VCADD=20, RV_VADD=20, RV_VDUP=20, RV_VSTI=20, RV_VLDI=16` |
+
+### N_ACC=4 (SMALL)
+
+| Backend | RVEC span | vs CCE | Top ops (approx.) |
+|---------|----------:|-------:|-------------------|
+| CCE ONEPT | **38** | 1.00× | `RV_VLDI=4, RV_VCADD=4, RV_VSTI=4, RV_PSET=1` |
+| VMI **group1** | **43** | 1.13× | `RV_VLDI=4, RV_VCADD=4, RV_VSTI=4, RV_PSET=3, RV_VDUP=1` |
+| VMI mask1 | **43** | 1.13× | `RV_VLDI=4, RV_VCADD=4, RV_VADD=4, RV_PSET=3, RV_VSTS=3` |
+| VMI pad8 | **53** | 1.39× | `RV_PSET=4, RV_VLDI=4, RV_VCADD=4, RV_VADD=4, RV_VDUP=4, RV_VSTI=4` |
+
+**Read:** `vcadd(group=1)` + `vstore(group=1, stride=1)` is the VMI form that
+gets ONEPT/`1PT_B32` lowering. At N_ACC=20 it matches (slightly beats) CCE
+RVEC. Pad-8 and mask1 remain useful as slower baselines. Older PyPI
+`ptoas 0.51` without a built tree overstated the pad-8 tax (~281 RVEC).
+
+## Done when
+
+Product kernels use the group1 / compact-score store (or equivalent ONEPT
+lowering) instead of pad-8. This repro already shows the fast path on the
+built pin above.
+
+## Pin
+
+- PTOAS: this branch `zjw/store_pad_issue` (from tag **vmi-v0.1.3**), built +
+  wheel-installed (`ptoas 0.52` in the measured run)
+- `PTOAS_ROOT` defaults to the repo root (`../../..` from this repro directory)

docs/repro/store_pad8/cce/csrc/inc/common.h

@@ -0,0 +1,22 @@
+#pragma once
+
+#include "acl/acl.h"
+#include "kernel_operator.h"
+#include "simt_api/asc_bf16.h"
+#include "simt_api/asc_fp16.h"
+
+#include "numeric_limits.h"
+#include "simd_inst.h"
+
+// Reduce N_ACC live f32 accumulator vectors (VL=64 lanes each) to N_ACC
+// scalars on one core. N_ACC=20 (LARGE) matches a typical residual-mix
+// epilogue. N_ACC=4 (SMALL) is a quick smoke case.
+constexpr int32_t kVL = 64;
+constexpr int32_t kMaxNAcc = 20;
+
+// UB layout (bytes):
+//   acc[N_ACC * kVL] f32      @0
+//   reduced[N_ACC] f32        @kMaxNAcc*kVL*4  (compact ONEPT store target)
+constexpr int32_t kAccUb = 0;
+constexpr int32_t kReducedUb = kMaxNAcc * kVL * 4;
+constexpr int32_t kUbBytes = kReducedUb + kMaxNAcc * 4;

docs/repro/store_pad8/cce/csrc/inc/simd_inst.h

@@ -0,0 +1,39 @@
+#pragma once
+
+// Minimal CCE SIMD helpers for store_pad8_vf_sim_kernel.cpp.
+// Only float vlds / vcadd / vsts — not a full TileLang simd_inst dump.
+
+#include "kernel_operator.h"
+
+namespace simd_inst {
+
+template <typename T> struct vec {};
+
+template <> struct vec<float> {
+  using type = vector_f32;
+};
+
+template <typename T> using vec_t = typename vec<T>::type;
+
+template <typename T, typename Offset, typename Dist>
+__simd_callee__ inline vec_t<T> vlds(__ubuf__ T *src, Offset offset,
+                                     Dist dist) {
+  vec_t<T> dst;
+  ::vlds(dst, src, offset, dist);
+  return dst;
+}
+
+template <typename SrcVec, typename Mode>
+__simd_callee__ inline SrcVec vcadd(SrcVec src, vector_bool mask, Mode mode) {
+  SrcVec dst;
+  ::vcadd(dst, src, mask, mode);
+  return dst;
+}
+
+template <typename T, typename Dist>
+__simd_callee__ inline void vsts(vec_t<T> data, __ubuf__ T *base,
+                                 int32_t offset, Dist dist, vector_bool mask) {
+  ::vsts(data, base, offset, dist, mask);
+}
+
+} // namespace simd_inst

docs/repro/store_pad8/cce/csrc/store_pad8_vf_sim_kernel.cpp

@@ -0,0 +1,61 @@
+/*===========================================================================
+ * store_pad8 microbench — CCE ONEPT path
+ *
+ * Per accumulator: vcadd then a single scalar vsts ... ONEPT_B32 store.
+ * No padding, no broadcast — the CCE baseline the VMI pad-8 port
+ * (../vmi/store_pad8_vmi.py) is compared against.
+ *===========================================================================*/
+
+#include "inc/common.h"
+
+namespace {
+
+template <int32_t N_ACC>
+__simd_vf__ inline void simd_vf_store_pad8_cce(__ubuf__ uint8_t *ub) {
+  auto *acc = (__ubuf__ float *)(ub + kAccUb);
+  auto *reduced = (__ubuf__ float *)(ub + kReducedUb);
+#pragma unroll
+  for (int32_t i = 0; i < N_ACC; ++i) {
+    vector_bool v = pset_b32(PAT_ALL);
+    auto acc_v = simd_inst::vlds<float>(acc + i * kVL, 0, NORM);
+    auto red = simd_inst::vcadd(acc_v, v, MODE_ZEROING);
+    simd_inst::vsts(red, reduced + i, 0, ONEPT_B32, v);
+  }
+}
+
+}  // namespace
+
+template <int32_t N_ACC>
+__global__ __vector__ void store_pad8_cce_kernel(__gm__ float *acc_g,
+                                                  __gm__ float *reduced_g) {
+  AscendC::InitSocState();
+  __ubuf__ uint8_t *ub = (__ubuf__ uint8_t *)0x0;
+
+  copy_gm_to_ubuf_align_v2(ub + kAccUb, (__gm__ uint8_t *)acc_g, 0, 1,
+                           N_ACC * kVL * 4, 0, 0, 0, 0, N_ACC * kVL * 4,
+                           N_ACC * kVL * 4);
+  AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(0);
+  AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(0);
+
+  simd_vf_store_pad8_cce<N_ACC>(ub);
+
+  AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(0);
+  AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(0);
+  copy_ubuf_to_gm_align_v2((__gm__ void *)reduced_g,
+                           (__ubuf__ void *)(ub + kReducedUb), 0, 1, N_ACC * 4,
+                           4, N_ACC * 4, N_ACC * 4);
+}
+
+extern "C" {
+
+void call_store_pad8_cce_large(void *stream, void *acc, void *reduced) {
+  store_pad8_cce_kernel<20><<<1, kUbBytes, stream>>>(
+      (__gm__ float *)acc, (__gm__ float *)reduced);
+}
+
+void call_store_pad8_cce_small(void *stream, void *acc, void *reduced) {
+  store_pad8_cce_kernel<4><<<1, kUbBytes, stream>>>(
+      (__gm__ float *)acc, (__gm__ float *)reduced);
+}
+
+}  // extern "C"

docs/repro/store_pad8/cce/launcher.py

@@ -0,0 +1,50 @@
+"""ctypes launcher for the store_pad8 CCE VF sim kernel."""
+
+from __future__ import annotations
+
+import ctypes
+from pathlib import Path
+import sys
+
+import torch
+
+_REPRO = Path(__file__).resolve().parent.parent
+if str(_REPRO) not in sys.path:
+    sys.path.insert(0, str(_REPRO))
+
+from common.torch_runtime import empty_npu, stream_ptr, sync  # noqa: E402
+
+_VL = 64
+
+_LIB: ctypes.CDLL | None = None
+
+
+def _lib() -> ctypes.CDLL:
+    global _LIB
+    if _LIB is None:
+        from common.cce_vf_build import build_cce_root
+
+        path = build_cce_root(Path(__file__).resolve().parent)
+        lib = ctypes.CDLL(str(path))
+        for name in ("call_store_pad8_cce_large", "call_store_pad8_cce_small"):
+            fn = getattr(lib, name)
+            fn.argtypes = [ctypes.c_void_p] * 3
+            fn.restype = None
+        _LIB = lib
+    return _LIB
+
+
+def _vp(t: torch.Tensor) -> ctypes.c_void_p:
+    return ctypes.c_void_p(t.data_ptr())
+
+
+def launch(acc: torch.Tensor, n_acc: int) -> torch.Tensor:
+    """acc: [n_acc, VL] f32 on npu -> [n_acc] f32 (compact ONEPT reduce)."""
+    lib = _lib()
+    name = "call_store_pad8_cce_large" if n_acc == 20 else "call_store_pad8_cce_small"
+    if n_acc not in (4, 20):
+        raise ValueError(f"unsupported n_acc={n_acc} (build only has large=20 / small=4)")
+    reduced = empty_npu((n_acc,), torch.float32)
+    getattr(lib, name)(stream_ptr(), _vp(acc), _vp(reduced))
+    sync()
+    return reduced