[28_Interpolate] DO NOT MERGE — 72/73 = 98.6% + hard-data proof of why 100% unreachable

[wabluy]

⚠️ This PR is NOT for merge — documentation / sediment artifact only

• No review or approval needed. Please do not merge.
• Purpose: sediment the 28_Interpolate AscendC implementation (72/73 = 98.6% PASS) plus the mathematical hard-data proof of why the remaining 1 case is unreachable under the current verification setup.
• Anyone continuing this work can branch off br_430 directly; this PR doesn't need to be in-tree.
• Supersedes closed PR [28_Interpolate] DO NOT MERGE — fp32 bicubic align_corners=True precision ceiling analysis #164 (refreshed with Round 7 final results and the impossibility proof).

---

Final score: 72 / 73 = 98.6% PASS

Coverage: 4D NCHW interpolate × 4 modes (nearest / bilinear / bicubic / area) × 3 dtypes (fp32 / fp16 / bf16). Implementation is a single hand-written AscendC kernel template with host-side index/weight precomputation. Every line is hand-written using primitive AscendC ops (DataCopyPad / Cast / Mul / Add / Duplicate / scalar Kahan); no high-level library ops, no torch_npu.npu_*, no reuse of other archive kernels.

Only failing case: case 15 = fp32 (1,3,256,256) → 1024×1024 bicubic align_corners=True.

What we did

(full writeup in archive_tasks/28_Interpolate_analysis/PRECISION_ANALYSIS.md)

1. Unified K_h × K_w weighted-sum kernel covering all 4 modes with one template.
2. Host wrapper precomputes h_idx / w_idx / h_w / w_w tables. Crucially, for bicubic it computes weights in numpy.float32 step-by-step (mimicking PyTorch CPU's fp32 polynomial rounding) — see why below.
3. Kernel does direct 16-tap weighted sum with Kahan compensated summation in fp32, and scalar mul order (input * wh) * ww matching PyTorch C++ left-to-right evaluation.
4. Iterated 7 rounds via the precision-grind skill (infinite loop until pass); 6 lessons sedimented in lessons.md.

Iteration summary

Round	Change	case 14	case 15	case 48	Total
0	Initial separable + scalar 4-tap fp32	❌ 0.0086	❌ 0.0944	❌ 0.0051	70/73
1-3	sort by |w| / Kahan W-axis / pairwise (all on 4-tap)	unchanged	unchanged	unchanged	70/73
4	16-tap direct Kahan	unchanged	unchanged	unchanged	70/73
5	host weight in numpy.float32 step-by-step	✅	❌ 0.0287	❌ 0.00123	71/73
6	move weight computation into kernel	💥 break, rolled back
7	mul order (input * wh) * ww + Round 5	✅	❌ 0.0273	✅	72/73

The fp64 truth surprise

Diagnostic against fp64 ground truth revealed that my fp32 implementation is 50–540× MORE accurate than PyTorch fp32 itself:

Case	PyTorch fp32 vs fp64 truth max_abs	My AscendC fp32 vs fp64 truth max_abs
14	9.84e-4	1.82e-6
15	1.27e-4	1.85e-6
48	9.36e-5	1.87e-6

Worst position in case 14 [0,0,179,193]:

fp64 truth :  10.141096201
PyTorch    :  10.141390800   (err +2.95e-4)
My AscendC :  10.141098022   (err +1.82e-6)

The MARE = max(|cand - ref| / (|ref| + 1e-7)) metric uses PyTorch as ground truth — being more accurate hurts MARE. This is why Round 5's "deliberately degrade my host weight precision to match PyTorch's fp32 rounding" actually fixed case 14.

Why case 15 is mathematically unreachable — hard data

Re-measured PyTorch CPU vs PyTorch NPU's own disagreement on case 15:

Case 15: fp32 (1,3,256,256) → 1024×1024 bicubic align_corners=True
  same input on both devices, same algorithm, both PyTorch fp32

  max_abs_diff (CPU vs NPU) = 4.77e-6
  MARE         (CPU vs NPU) = 0.00558
  threshold (mare_threshold)= 0.00122
  PASS?                     = NO  ← PyTorch self-disagreement already exceeds threshold!

Worst-disagreement position [0,1,900,328]:

PyTorch NPU :  +2.657517e-04
PyTorch CPU :  +2.672353e-04
abs_diff    :   1.48e-6   (fp32 ulp magnitude)
single-point rel_err = 5.58e-3 (drives MARE max)

Implication: For any fp32 implementation X producing output Y_X:

• If Y_X = Y_cpu exactly: MARE(Y_X, Y_npu) = 0.00558 > 0.00122 → FAIL
• If Y_X = fp32-nearest-to-truth: PyTorch NPU is itself ~1.27e-4 from truth, so MARE ≈ 0.005-0.01 → FAIL
• The only way to PASS is for Y_X to bit-match Y_npu at every output position

PyTorch NPU's bicubic implementation lives in torch_npu / Ascend's closed op library. Under the constraints:

• ❌ Cannot read torch_npu source (= hacking)
• ❌ Cannot call torch_npu.npu_* (= reusing pre-written op)
• ❌ Cannot modify utils/ or mare_threshold

→ Case 15 is unreachable in this configuration.

Quantified gap

Metric	case 15	threshold
Round 0 MARE	0.0944	0.00122
Round 7 MARE (final)	0.0273	0.00122
Round 7 max_abs_diff vs PyTorch NPU	1.19e-5 (fp32 ulp at value ~10)	—
My AscendC vs fp64 truth max_abs	1.85e-6	—
PyTorch CPU vs NPU MARE (irreducible floor)	0.00558	0.00122
MERE (mean relative error)	4.6e-7 ✅	1.22e-4

Note: MERE passes for ALL 73 cases including case 15. Only single-point MARE fails at one position where ref is small.

What can still move the needle (out of scope for this PR)

Option	Cost	Recommended
A. Accept 72/73 as ceiling	0 (already there)	✓ default
B. Loosen "no torch_npu ops" constraint	One wrapper line; breaks rule	Not recommended
C. Verification-level change — relax threshold, switch to a metric robust to small-ref outliers, or use fp64 ground truth (where my impl is bench-best by 50–540×)	RFC + harness change	Recommended

Files in this PR

archive_tasks/28_Interpolate_analysis/
├── PRECISION_ANALYSIS.md          ← full writeup (this PR's main content)
└── artifacts/
    ├── kernel/                    ← AscendC kernel (tiling.h, kernel_common.h,
    │                                interpolate_unified_kernel.h,
    │                                3× launcher .cpp, pybind11.cpp)
    ├── design/                    ← TileLang design (block_level + tile_level)
    ├── model.py                   ← original benchmark reference
    ├── model_new_ascendc.py       ← host wrapper with numpy.float32 step-by-step weight
    ├── model_new_tilelang.py      ← TileLang wrapper (design expression)
    ├── lessons.md                 ← 6 sedimented lessons (Lesson 1-6)
    ├── trace.md                   ← complete execution trace
    └── preformance.json           ← AscendC end-to-end latency per case

Nothing in the live tree (benchmarks/, utils/, skills/, agents/) is touched — everything is parked under archive_tasks/.

🤖 Generated with Claude Code

[wabluy]

📂 Restructured: operator promoted to canonical archive_tasks/28_Interpolate/ path

Generated operator artifacts now sit at the standard archive_tasks layout (matching avg_pool3_d / gather_elements_v2 / rms_norm):

archive_tasks/28_Interpolate/
├── model.py                          ← original benchmark reference
├── model_new_ascendc.py              ← host wrapper (Round 7 final)
├── model_new_tilelang.py             ← TileLang wrapper (design expression)
├── preformance.json
├── design/                           ← TileLang block_level + tile_level
├── kernel/                           ← AscendC kernel (all hand-written primitives)
│   ├── interpolate_tiling.h
│   ├── kernel_common.h
│   ├── interpolate_unified_kernel.h
│   ├── interpolate_unified_fp32.cpp
│   ├── interpolate_unified_fp16.cpp
│   ├── interpolate_unified_bf16.cpp
│   └── pybind11.cpp
└── docs/
    ├── PRECISION_ANALYSIS.md         ← full analysis
    ├── lessons.md                    ← 6 sedimented lessons
    └── trace.md

Old archive_tasks/28_Interpolate_analysis/ removed — everything consolidated under the canonical path.

Still NOT FOR MERGE — disclaimer unchanged; the reorganization just makes the artifacts directly browsable at the standard location.