Commit Graph

117 Commits

Author SHA1 Message Date
Beckylu 13e9aee4a4 docs: simplify operator claim rules 2026-07-31 11:30:31 +08:00
FrRay 9586e0164f docs: Update TileKernels-MACA operator migration list 2026-07-31 11:17:47 +08:00
FrRay 00dea71686 docs: add TileKernels-MACA operator migration list 2026-07-31 10:32:31 +08:00
Beckylu 79d4882172 docs: refine summer camp operator selection scope 2026-07-30 17:28:50 +08:00
Beckylu 3eb4f8f92b docs: add summer camp project naming guidance 2026-07-30 14:10:13 +08:00
Beckylu 8f9d78123f Merge pull request '新增证书说明' (#13) from docs/add-task-awards-overview into summer-camp-2026 2026-07-30 13:55:07 +08:00
Beckylu 86f9d33b99 docs: add certificate wording and online participation plan 2026-07-30 13:49:39 +08:00
FrRay f58405e124 docs: remove unnecessary line breaks from READMEs 2026-07-30 12:53:19 +08:00
Beckylu 94a471453a Merge latest summer-camp-2026 documentation updates 2026-07-30 11:22:20 +08:00
Beckylu 0daed2916c docs: clarify operator claim priorities 2026-07-30 11:19:38 +08:00
FrRay 19d16f9760 docs: clarify PR B rebasing after PR A merge 2026-07-30 11:16:22 +08:00
Beckylu ac6512ed86 docs: use one branch for both PR stages 2026-07-29 17:04:18 +08:00
Beckylu 8444e34a7f docs: add Manifest PR templates 2026-07-29 16:09:14 +08:00
Beckylu 45d01cbd03 docs: recommend summer camp MetaX image 2026-07-29 10:48:29 +08:00
wawahejun ce1b15c28f docs: add MetaX C500 summer-camp guide on top of upstream dev
Source tree is identical to MetaX-MACA/TileOPs-Metax dev at f02d3d8; this commit
carries only the summer-camp documentation and PR templates. Content verified by
running everything on a real MetaX C500 (MACA 3.7.1.5, torch
2.8.0+metax3.7.1.3, tilelang 0.1.10+cuda.gitf549117c, sGPU slice 16000 MiB).

Installation (high severity). The documented `make install`,
`pip install tileops`, and bare `python3 -m venv .venv` steps destroy a working
MACA environment. The container's TileLang is an in-place source build imported
via PYTHONPATH, so pip reports it as absent and resolves the official CUDA wheel
over it; a venv without --system-site-packages cuts off the MetaX PyTorch build
and the ABI-coupled apache-tvm-ffi. Replace those steps with the PYTHONPATH
setup, document that tileops needs no install at all, and note that --no-deps is
the only safe install form (as scripts/ci/install_tileops.sh already does).
Flag -c constraints.txt as CUDA-CI-only for the same ABI reason. Add TileLang
provenance and backend checks to the verification list, which previously covered
mx-smi, torch and einops but not the component most likely to be wrong.

Quick start. GemmOp(M, N, K, dtype=...) does not match the implementation --
GemmOp is input-inferred and takes only trans_a/trans_b. Fix the signature and
document the trans_b default, keeping the original M,N,K of 1024,1024,512, which
passes on C500 via the MACA dispatch path.

New sections 1.2 and 1.3 in the migration guide. Document the is_maca() dispatch
to *_maca.py kernels, and that a gated kernel does not imply an unusable Op:
GemmKernel declares [89, 90] and is gated on C500, yet GemmOp works because it
dispatches to gemm_maca.py ([80, 86, 89, 90]). Availability must be judged from
what the Op layer dispatches to, not from one kernel's supported_archs. List the
20 declarations that exclude 80 as unsuitable migration targets, and note that
adding a *_maca.py kernel plus dispatch is a good target instead. Record that
get_sm_version() reuses NVIDIA's encoding, so C500 reports 80 while sharing
nothing with Ampere, and that the raw "architecture 80" message names no device.
Document that a usable Op still has shape limits: SoftmaxFwdOp fails above a
1024-wide reduction dimension (mcErrorInvalidValue), independent of row count.
Document that a parent process which has imported tilelang will see any
subprocess that imports it again SIGKILLed with no output, which aborts
tests/test_validate_manifest.py at exit 137, and give the deselect workaround.

Roofline. Record the sGPU slice quota and state whether peaks are whole-card or
slice-scaled; dividing a slice measurement by a whole-card peak yields an
unexplainable efficiency.

Verified on C500 against this tree: validate_manifest.py exit 0; 29 passed
across test_ops_manifest.py, test_kernel_map_install.py and benchmarks/tests;
GemmOp passes at 1024x1024x512, 1024^3 and 4096^3; the documented quick-start
snippet and every self-check command run as written. pre-commit and ruff are
unavailable in this container (installing them would invoke pip dependency
resolution), so formatting was checked via git diff --check and end-of-file
newlines instead.

Squashed documentation commits by Beckylu <648245013@qq.com> and
FrRay <1077376663@qq.com> covering the summer-camp guide, PR templates, and
README translations.

Co-Authored-By: Beckylu <648245013@qq.com>
Co-Authored-By: FrRay <1077376663@qq.com>
2026-07-28 17:55:09 +00:00
Beckylu 09e6b0021a Update README.zh-CN.md 2026-07-27 14:46:41 +08:00
Beckylu 8e3160ea8a docs: add 2026 summer camp contribution workflow 2026-07-27 14:24:21 +08:00
Cao Ying c986df5407
[Refactor][POOL] strip over-design and scaffolding from the merged cleanup PRs (#1779)
## Summary

- Remove over-engineered pool hooks and private snapshot tests; make
indexed max-pool forwards explicit.
- Deduplicate GLA, Mamba, and formula tests; remove dead helpers,
commented benchmark rows, and process metadata.
- Consolidate benchmarks onto manifest workloads and
`ManifestBenchmark`; delete obsolete benchmark modules and factor
redundant sweep axes.
- Require implemented ops to declare `kernel_map` and manifest-driven
benchmark coverage, filling the corresponding manifest gaps.
- Reduce the repository by 2,415 net lines without changing runtime
operator behavior.

## Test plan

- [x] pre-commit passed
- [x] Pool tests passed: 246
- [x] Perf/formula and validator tests passed: 128
- [x] Changed benchmark modules collected: 538 nodes
- [x] Test node delta: 476 → 391 (-85)
- [x] `python scripts/validate_manifest.py --strict` passed

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-07-27 10:05:38 +08:00
Cao Ying c9097efb64
[Refactor][POOL] collapse pool op class bodies into ndim-generic family bases (#1777)
Closes #1766

## Summary

- Collapse the nine pool op class bodies into two ndim-parametrized
generic bases (`_AvgPoolFwdOpBase`, `_MaxPoolFwdOpBase`);
`tileops/ops/pool.py` shrinks 1431 -> 918 lines (-513).
- All nine public classes stay as thin subclasses with unchanged names,
rank-specific ctor/forward signatures, and per-op `kernel_map`.
- Add a cross-family contract snapshot suite: ctor signatures,
`kernel_map` precedence, kernel-cache-key separation, `register_fake`
shapes/dtypes, error messages, compile behavior.
- Document the dimension-parametrized family-base + `ndim` protocol in
`docs/design/ops-design.md`.
- Zero `tileops/manifest/` edits; `scripts/validate_manifest.py` passes
unchanged.

## Test plan

- [x] Modified files pass unit tests
- [x] Contract snapshot suite green before and after each collapse step
- [x] `tileops/ops/pool.py` shrinks by >=500 lines with all nine public
classes intact
- [x] `python scripts/validate_manifest.py` passes with zero manifest
edits
- [x] CUDA eager and `torch.compile(fullgraph=True)` smoke parity for
all nine ops
- [x] `docs/design/ops-design.md` documents the generic-base + `ndim`
protocol

## Regression

- Contract snapshots freeze the public surface: `inspect.signature` for
all nine ctors is byte-identical to the pre-collapse baseline;
fake-output shapes/dtypes (incl. int64 MaxPool indices) and error
messages are snapshotted.
- Avg-pool 1d/3d-vs-2d fast-path asymmetry and the explicit kernel-cache
keys are preserved as-is.

## Test node delta

```
File                      Base    HEAD    Delta
-----------------------------------------------
tests/ops/test_pool.py     194     244      +50
-----------------------------------------------
TOTAL                      194     244      +50

Growth: +25.8%
```

**Justification:** the added nodes are the contract snapshot suite that
guards the collapse — each case pins a distinct public class or
preservation path (ctor signature, cache-key separation, fake
registration, error message, fullgraph compile), so the refactor cannot
silently change the public surface.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-07-26 20:06:55 +08:00
Cao Ying 0a9bf1e40a
[Chore][Cleanup] repo slimming: dead code, duplicated tests/benches, compat shims, file defragmentation (#1764)
Closes #1763

## Summary

- Remove dead helpers, redundant elementwise benchmark drivers,
decorative banners, and stale compatibility aliases.
- Consolidate duplicated tests and private benchmark wrappers onto
shared parametrized fixtures and `ManifestBenchmark`.
- Merge fragmented workload, reduction-op, MHC, and normalization
benchmark modules; retarget imports and manifest source paths.
- Preserve canonical runtime behavior while reducing the repository by
roughly 4.3k net lines.

## Test plan

- [x] pre-commit passed
- [x] `python scripts/validate_manifest.py --strict` passed
- [x] Repository-wide pytest collection completed: 5,178 tests, 0
collection errors
- [x] Touched test and benchmark modules collect cleanly

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-07-26 12:59:48 +08:00
Cao Ying a72e78b939
[Refactor][OPS] narrow manifest key-format rule; move variant words before direction suffix (#1765)
Closes #1758

## Summary

- Narrow the key-format rule in `docs/design/manifest.md`: direction
suffix is required only where a direction pair (or bwd sibling) exists;
variant words always precede `FwdOp`/`BwdOp`.
- Rename `GroupNormFwdOpNoAffine` → `GroupNormNoAffineFwdOp` and
`InstanceNormFwdOpNoAffine` → `InstanceNormNoAffineFwdOp` (class +
manifest key + exports + call sites).
- Add the narrowed format check to `check_l0` with one guard per reject
branch; the shipped manifest passes.

The 21 single-direction keys without a direction suffix are legitimized
by the narrowed rule, not an omission. Manifest+code rename atomicity is
human-driven, enforced by the `cls.__name__ == key` validator check.

## Test plan

- [x] `pytest tests/test_validate_manifest.py` — 142 passed, including
one test per new reject branch
- [x] `pytest tests/ops/test_group_norm.py
tests/ops/test_instance_norm.py` — 73 passed (GPU)
- [x] `python scripts/validate_manifest.py --levels
schema,signature,shape,dtype,bench --strict` — shipped manifest passes

## Test node delta

```
File                               Base    HEAD    Delta
--------------------------------------------------------
tests/test_validate_manifest.py     140     142       +2
--------------------------------------------------------
```

**Justification:** one test per new `check_l0` reject branch (variant
word after direction suffix; missing direction suffix with a bwd
sibling). The norm test files are pure renames with no node growth
(base-side collection fails against the renamed package, so the script
reports them as new).

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-07-26 11:52:07 +08:00
Cao Ying 139e60ee3d
[Feat][OPS] compile dispatch boundary; pool declarations; design-doc audit (#1757)
## Summary

- New contract: a traced `Op.forward` never constructs Kernels. Shared
registry (`tileops/ops/compile_boundary.py`) + opaque custom-op
dispatch; auto-registration via `Op.dispatch_kernel`. Recorded in
ops-design.md.
- Pool adopts it: all 9 ops compile cold under `fullgraph=True`; kernel
files untouched; `pool.yaml` declares them. BatchNorm migrates to the
shared registry.
- Design-doc audit (all 7 docs): stale facts fixed, implementation
snapshots and non-decisions removed; MoE class tour moved to its module
docstring; scaffold example made fully fictional. testing.md needed
nothing.

## Test plan

- [x] 9 pool ops cold-compile on H200, match `F.*_pool*`; eager
regression clean
- [x] Equality gate 73 == 73; validator + 189 structural tests pass;
pool smoke 86, norm 5
- [x] Node delta +3 (new avg_pool compile cases); pre-commit clean

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-07-26 10:36:34 +08:00
Cao Ying a3641e25f9
[Feat][Manifest] enforce torch_compile_fullgraph declarations with always-on gate (#1752)
## Summary

Part of #1748 (step 3/3, final).

- One equality test: manifest `torch_compile_fullgraph` declarations ==
registered compile-test evidence. The three interim registry tests it
subsumes are removed.
- New always-on preflight job `compile-contract-gate` (no path
condition) runs it on every PR.
- `architecture.md` L2 claim narrowed to the declared set; one design
bullet in `ops-design.md`.

## Test plan

- [x] `pytest tests/test_validate_manifest.py` — 189 passed (gate live:
64 == 64)
- [x] `compile-contract-gate` green on this PR
- [x] actionlint + pre-commit clean

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-07-26 07:09:05 +08:00
Cao Ying 7892d02d6b
[Feat][Manifest] declare torch_compile_fullgraph on evidence-backed ops (#1751)
Part of #1748 (step 2 of 3).

## Summary

- Declare `torch_compile_fullgraph: true` on the 64 ops backed by
registered cold fullgraph tests.
- Document the field semantics and Manifest-stage ownership.
- Leave the equality gate and CI enforcement to step 3; ops without
qualifying evidence remain undeclared.

## Test plan

- [x] Declared set equals registry: 64 == 64
- [x] Validator and 193 validator tests passed
- [x] Pre-commit passed

## Test node delta

No test files changed.

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-07-26 00:13:45 +08:00
Cao Ying 24886ce6d0
[Fix][Benchmark] stop losing results on native failures; derive workload keys from signature (#1745)
Closes #1740. Closes #1746.

## Summary

- **One failure no longer destroys the session's results.** tilelang's
fd-level output suppression corrupted pytest's capture stream; when the
07-24 nightly stalled >900s, the timeout handler crashed and all 1328
completed results were lost. Suppression now applies only when
stdout/stderr are the real fds 1/2. The stall itself did not reproduce
(7 attempts, latency parity with the last green nightly); mechanism
unidentified. Process isolation: #1744.
- **Measurement never degrades silently.** CUPTI projection failure
raises a dedicated sentinel — genuine CUDA errors/OOM propagate; the
CUDA-events fallback and the >1 GiB clone-skip are logged and marked in
results.
- **Workload keys derive from the manifest signature.** The hardcoded
shape-key allowlist (which let #1734 break dropout/fft collection) is
gone; the contract lives once in `tileops.manifest`, consumed by both
`workloads_to_params` and validator rule R21. 26 manifest entries
aligned to their signatures (`x_shape`→`input_shape`;
`DropoutOp`/`FFTC2COp` input `x`→`input` incl. `forward()`). Unknown
signature/entry/workload keys are rejected; the validator stays total
over malformed YAML.
- **Cleanup.** "harness" jargon removed repo-wide; benchmark_base tests
folded into one file; 24 redundant/tombstone validator tests deleted;
`init_dims`/`parity_opt_out` special cases replaced by generic
unknown-key rejection. Net −350 lines.

## Test plan

- `validate_manifest.py` full pass; 207 validator + 17 benchmark unit
tests; `--collect-only benchmarks/ops` 1638, zero errors
- GPU (H200): dropout/fft tests 37 passed, benches 6 passed;
`bench_moe_grouped_gemm_nopad` 4/4; fallback + error-propagation
contract tests

## Benchmark

No kernel change and the timing protocol (CUPTI kernel-only,
warmup/repeats/trials, L2 flush) is untouched; this PR changes
suppression selection and fallback/error semantics around it. Latency
parity: `bench_moe_grouped_gemm_nopad` decode-down 2.3061 ms vs 07-17's
2.3088 ms.

## Regression

- Bench files keep the `(shape, dtype[, extra])` parametrize contract
- Post-merge check: next nightly progresses past
`bench_moe_grouped_gemm_nopad.py` with a non-empty `bench_results.xml`

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-07-25 22:11:21 +08:00
Ang Gao c2e98526de
[Feat][Ops] Infer shape metadata from op inputs (#1651)
## Summary

Closes #1648.

This updates the initial #1648 scope so derivable shape metadata is
inferred from op inputs instead of being required as constructor
arguments:

- Norm family: `FusedAddLayerNormFwdOp`, `FusedAddRMSNormFwdOp`,
`AdaLayerNormFwdOp`, `AdaLayerNormZeroFwdOp`
- Scan family: `CumsumFwdOp`, `CumprodFwdOp`
- MoE leaf ops: `FusedTopKOp`, `MoePermuteNopadFwdOp`

The old explicit constructor arguments remain supported as
compatibility/strict-validation commitments, while the preferred API
derives `M`, `N`, token counts, top-k shape, hidden size, and dtype from
the input tensors at `forward()` time. Kernels are built lazily and
cached by the inferred specialization.

## Manifest / callers

- Removed input-derivable `static_dims` from the affected norm and scan
manifest entries.
- Updated `MoePermuteNopadFwdOp` manifest params/workloads/roofline to
treat `hidden_states` and `topk_ids` shapes as the source of truth,
keeping `num_experts` explicit.
- Updated direct benchmarks and composite call sites to use the
preferred input-inferred APIs.
- Added compatibility and cache-reuse/changing-shape test coverage.

## Testing

Using the TileOpsGov runner image
`ghcr.io/tile-ai/tileops-runner:65dbc98-torch2.10` on H200 GPU:

- `python -m pytest -q tests/ops/test_cumulative.py
tests/ops/test_fused_add_layer_norm.py
tests/ops/test_fused_add_rms_norm.py tests/ops/test_ada_layer_norm.py
tests/ops/test_ada_layer_norm_zero.py tests/ops/test_moe_fused_topk.py
tests/ops/test_moe_permute_nopad.py -m smoke --tb=short --timeout=900
--timeout-method=thread`
  - `67 passed, 52 deselected, 14 warnings`
- `python -m pytest -q tests/ops/test_cumulative.py
tests/ops/test_fused_add_layer_norm.py
tests/ops/test_fused_add_rms_norm.py tests/ops/test_ada_layer_norm.py
tests/ops/test_ada_layer_norm_zero.py tests/ops/test_moe_fused_topk.py
tests/ops/test_moe_permute_nopad.py --tb=short --timeout=900
--timeout-method=thread`
  - `119 passed, 14 warnings`
- `python scripts/validate_manifest.py`
  - `All manifest checks passed`

Note: the validator reports an advisory warning that
`MoePermuteNopadFwdOp.num_experts` has an `__init__` default of `None`
while the manifest has no default. This is intentional to preserve old
positional construction compatibility; `forward()` still requires
`num_experts` to be provided.
2026-07-06 17:51:47 +08:00
Zheng QiHang 2bcb708413
[Doc][Trace] Doc-ready docstrings + in-kernel timeline tutorial (#1640)
## What

Makes the `tileops.trace` module documentation-ready and adds a usage
tutorial,
so the trace tool renders cleanly on the docs site (mkdocstrings +
Material).

- **Docstrings: Sphinx roles → plain code spans.** Across the whole
`tileops/trace/` package (`__init__`, `api`, `decode`, `markers`,
`passes`,
`record`, `state`, `ui`), replace `:func:` / `:class:` / `:mod:` /
`:data:`
  cross-reference roles with double-backtick code spans. The project is
Google-docstring-only; mkdocstrings does not resolve Sphinx roles, so
they
  rendered as literal `:func:`...`` text on the site.
- **`ui.py` timeline legend fix.** The horizontal legend overlapped the
x-axis
  title. Give the x-axis title a `standoff`, anchor the legend below it
  (`y: -0.30`, `yanchor: top`), and grow the bottom margin (120 → 170).
- **New tutorial: `docs/perf/trace-timeline.md`.** End-to-end
walkthrough —
  write a traced warp-specialized GEMM, run it, enable tracing, read the
timeline. Each code block uses Material code annotations (the `+`
markers)
  linked to the auto-generated API reference.
- **`docs/perf/README.md`**: add a "Tools & Guides" entry pointing at
the tutorial.

## Why

The trace tool shipped in #1638 with no rendered API docs or guide; this
makes
its public surface (`tileops.trace.api`) render correctly and gives
users a
single walkthrough.

## Notes

- No behavior change: docstring text and one Plotly layout dict only.
- The site-side wiring (mkdocstrings page, nav, embedded timeline) is a
separate
  PR against `tile-ai/TileOPs.github.io`.

Closes #1639
2026-06-30 13:27:07 +08:00
Maybe_Jayden 48852858ba
[Refactor][MoE] Unify FusedMoEExperts hierarchy with 3WG persistent kernel (#1514)
Closes #1513

### Summary

Refactor the MoE experts implementation to align with vLLM's
`FusedMoEExpertsModular` interface while integrating the 3WG persistent
grouped GEMM kernel as the default backend. Eliminates 299 lines of
duplicated code and establishes a clean inheritance hierarchy.

### Architecture

```
Op → FusedMoEExperts (ABC) → FusedMoEExpertsModular (ABC)
  ├── FusedMoEExpertsNopadPersistent3WGFwdOp (nopad, 3WG default)
  └── FusedMoEExpertsPaddedFwdOp (padded baseline)
```

### Changes

| Category | Detail |
|----------|--------|
| Rename ABCs | `MoEExperts` → `FusedMoEExperts`, `MoEExpertsModular` →
`FusedMoEExpertsModular`, `MoEPrepareAndFinalize` →
`FusedMoEPrepareAndFinalize` |
| Inherit Op | `FusedMoEExperts` now inherits from `Op`; unify on
`forward()` |
| Default GEMM | `GroupedGemmPersistent3WGKernel` with auto-fallback for
unaligned dims |
| Delete duplicate | `fused_moe_experts.py` removed (-299 lines) |
| Rename files | `test_moe_experts_nopad.py` →
`test_fused_moe_experts.py`, `bench_moe_experts_nopad.py` →
`bench_fused_moe_experts.py` |
| Migrate callers | `FusedMoe`, `prepare_finalize/no_dp_ep.py`,
manifest, docs |
| Tests | Replace deleted-op reference with PyTorch per-expert ground
truth |

### Performance (H200, bf16, H=7168, F=2048, K=8)

| Workload | TileOPs 3WG | vLLM Triton | Ratio |
|----------|-------------|-------------|-------|
| T=4096, E=128 (prefill) | 7.0ms / 410 TF | 7.7ms / 375 TF | **109%** |
| T=4096, E=256 (prefill) | 9.2ms / 314 TF | 9.5ms / 304 TF | **104%** |
| T=512, E=128 (decode) | 3.2ms / 113 TF | 2.8ms / 128 TF | 88% |
| T=512, E=256 (decode) | 6.1ms / 59 TF | 5.5ms / 66 TF | 89% |

Prefill: matches or exceeds vLLM Triton. Decode: ~10% behind (persistent
scheduling overhead at small batch).

### Testing

- `test_fused_moe_experts.py`: 15/15 passed
- `test_moe_fused_moe.py` (integration): smoke 11/11 passed
- ruff: all checks passed

### How to verify

```bash
# Unit tests
TILELANG_CLEANUP_TEMP_FILES=1 conda run -n tileops-tl9 \
    python -m pytest tests/ops/test_fused_moe_experts.py -vv

# Integration
TILELANG_CLEANUP_TEMP_FILES=1 conda run -n tileops-tl9 \
    python -m pytest tests/ops/test_moe_fused_moe.py -k smoke -v
```
2026-05-29 17:36:45 +08:00
Cao Ying 7064bc323f [Chore][MANIFEST] enumerate source.test/source.bench in status-flip carve-out (#1493)
Closes #1492

## Summary

- Enumerate `source.test` and `source.bench` explicitly in the manifest
carve-out (`.claude/rules/manifest-trust-model.md`) as **in-scope** for
status-flip PRs, with a `Why:` line citing validator behavior
(presence-only + AST check, no path canonicality).
- Re-align `MoeGroupedGemmNopadFwdOp` in `tileops/manifest/moe.yaml`:
- `source.test`: `tests/ops/test_moe_fused_moe.py` →
`tests/ops/test_moe_grouped_gemm_nopad.py`
- `source.bench`: `benchmarks/ops/bench_moe_fused_moe.py` →
`benchmarks/ops/bench_moe_grouped_gemm_nopad.py`
- Audit of all 132 ops: 0 missing files, 1 actionable umbrella mismatch
(the one fixed here), 115 umbrella-by-design (per-family shared
harness), 16 fully aligned.

## Decision rationale (Option A)

`source.test` / `source.bench` are **discoverability pointers**, not
trust-bearing fields:

- `docs/design/trust-model.md` §Manifest §OWNS lists signatures, dtypes,
workload shapes, roofline formulas, status, and `kernel_map`.
`source.test` / `source.bench` are not OWNS items.
- `scripts/validate_manifest.py` only requires their **presence**
(`_REQUIRED_SOURCE`); `check_l4_benchmark` opens the bench file to
AST-check for `load_workloads` + `eval_roofline`, but does not enforce
path canonicality.

Therefore: realigning these pointers onto the per-op test/bench file
that an implementation PR has just authored does not weaken any
trust-bearing field, and belongs in the carve-out alongside `status` and
`source.kernel_map`. Every other `source.*` key (`source.kernel`,
`source.op`, `source.bench_manifest_driven`) remains out of scope and
continues to require a separate manifest-only PR.

## Audit (132 ops scanned)

| Category | Count |
| --- | --- |
| `test` or `bench` file missing | 0 |
| Actionable umbrella mismatch | 1 (`MoeGroupedGemmNopadFwdOp`, fixed in
this PR) |
| Umbrella by design (per-family shared harness) | 115 |
| Fully aligned | 16 |

Umbrella-by-design breakdown: attention 7, convolution 2, elementwise
71, moe 6, normalization 8, reduction 19, scan 2. These reflect
deliberate per-family harness choices (e.g. `test_unary_math.py` covers
17 unary math ops); splitting them is per-family design work, out of
scope here.

Full audit artifact:
`.foundry/runs/issue-1492/pipeline/audit-source-paths.md` (run-local,
not shipped).

## Test plan

- [x] **AC-1**: `.claude/rules/manifest-trust-model.md` carve-out
explicitly names `source.test` and `source.bench` with no ambiguity
(line 12 lists them in-scope with `Why:` rationale; line 15 enumerates
excluded `source.*` keys).
- [x] **AC-2**: Audit enumerates every manifest entry whose
`source.test` / `source.bench` is stale or umbrella (132 ops scanned;
classified as missing / actionable-umbrella / umbrella-by-design with
explicit per-op canonical-path derivation rule).
- [x] **AC-3**: `MoeGroupedGemmNopadFwdOp.source.test =
tests/ops/test_moe_grouped_gemm_nopad.py` and `source.bench =
benchmarks/ops/bench_moe_grouped_gemm_nopad.py` (verified via
`ruamel.yaml` parse of `tileops/manifest/moe.yaml`).
- [x] **AC-4**: `python scripts/validate_manifest.py --strict` exits 0.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-17 10:12:56 +08:00
Cao Ying 7e8deff27b [Feat][Codegen] generate eval_roofline bodies from manifest roofline (#1471)
Closes #1459

## Summary

- New `tileops/ops/_roofline_codegen.py`: synthesizes each op's
`eval_roofline()` from its manifest `roofline` block. Two modes — `func`
(dotted path eagerly resolved at class build) and `inline` (vars + flops
+ bytes, AST-validated against the §4.4.3/§4.4.4 namespace, emitted as
plain Python with no runtime `eval`).
- `Op.__init_subclass__` installs both `_validate_dtypes` (from #1466)
and `eval_roofline` (new) on every `status: implemented` subclass;
skipped on spec-only ops and on classes / bases that supply their own
override (MRO-checked).
- Inline mode reads op state under a 2-tier contract per
`signature.inputs`/`params` names *referenced* by the roofline: tensor
inputs as `self.<input>` (with both `.shape` and `.ndim`) or
`self.<input>_shape` (shape tuple/list); params as `self.<param>`.
Unreferenced manifest names are not bound. Missing required bindings
raise `ValueError(op, input, conventions)`.
- Vars-layer AST validator is scope-aware: each comprehension
generator's iterable is visited before its target binds (matching Python
semantics); `.shape` / `.ndim` are accepted only directly on a declared
tensor-input Name — chained (`x.shape.ndim`), subscripted
(`x.shape[0].shape`), and local-operand (`N.shape`) forms reject; tensor
inputs may not appear as bare values (no `sum(x)`, no `x[0]`); calls
restricted to the §4.4.4 helper table; vars keys colliding with
inputs/params/helpers/`elem_bytes`/earlier vars are rejected at
synthesis.
- Arithmetic-layer validator is a positive AST allowlist (BinOp,
UnaryOp, BoolOp, IfExp, Compare, Call to allowed helper, Constant, Name)
— collection literals (`Tuple`/`List`/`Set`/`Dict`) and any other node
kind fail synthesis.
- `PreluFwdOp` / `NanToNumFwdOp` (the two real inline-mode ops) declare
bindings explicitly: `self.input_shape`, `self.weight_shape`,
`self.input_shape`.
- `docs/design/roofline.md` §4.4.3 names the referenced-only binding
decision and links to the rule.
- `.claude/domain-rules/ops-design.md` adds the concrete attribute
contract for op authors.

## Test plan

- [x] pre-commit passed
- [x] `python -m pytest tests/test_roofline_codegen.py` — 36 passed
- [x] `python -m pytest tests/test_validate_manifest.py` — 221 passed
- [x] All 108 `status: implemented` ops synthesize without error against
the live manifest
- [x] `python scripts/validate_manifest.py --strict`: this PR is *scoped
to the C7 class* (`eval_roofline is the Op base stub`), and that count
is **0** on this branch (was **41** on `testbed`). The script still
exits non-zero with **58 `kw_only mismatch` errors in reduction ops**
that are pre-existing on `testbed` (`testbed` itself reports 99 errors =
41 C7 + 58 kw_only) and are out of scope for this PR — they are tracked
separately.

## Test node delta

`python scripts/test_node_delta.py --base upstream/testbed`:

```
File                              Base    HEAD    Delta
-------------------------------------------------------
tests/test_roofline_codegen.py     new      36    (new)
-------------------------------------------------------
TOTAL                                0      36      +36
```

**Justification:** new codegen module; tests cover both synthesis modes,
the install-hook condition matrix (spec-only / explicit override /
inherited override / missing metadata), the 2-tier binding contract
(positive + tier-1 shape/ndim fallthrough + missing-binding error),
vars-layer comprehension scope (target binds after iter), attribute
whitelist + chained/subscripted/local-operand rejection, tensor-input
bare-reference rejection, vars-key collision rejection, arithmetic-layer
collection-literal rejection, param exposure contract, and the
live-manifest C7 parity invariant.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-16 08:41:27 +08:00
Cao Ying 8e4248dad5
[Refactor][Reduction] replace pytorch scalar fallback with native 0-D path (#1453)
Closes #1448

## Summary

- Replace the PyTorch runtime scalar fallback in reduction ops with a
TileOps-native 0-D path covering
Sum/Mean/Amax/Amin/Prod/All/Any/CountNonzero/Std/Var/VarMean (and
Welford variants).
- Update `docs/design/trust-model.md` and
`.claude/domain-rules/manifest-spec.md` so PyTorch is documented as the
spec oracle, not the default runtime implementation; allowed-fallback
conditions are narrowed and named.
- Extend `tests/ops/test_reduce_scalar_conformance.py` and
`tests/ops/test_reduction_scalar_input.py` to cover scalar inputs across
all reduction ops and the documented exception boundary (invalid DOF).

## Test plan

- [x] AC-1: Modified files pass unit tests. — `CUDA_VISIBLE_DEVICES=3
python -m pytest -q tests/ops/test_reduction*.py
tests/ops/test_reduce*.py tests/ops/test_*reduce*.py` → 1054 passed.
- [x] AC-2: Scalar reduction paths no longer call PyTorch reference
helpers at runtime where TileOps can implement the behavior directly. —
Helper-blocking probe patched
`torch.sum/mean/amax/amin/prod/std/var/var_mean/all/any/count_nonzero`
to raise; scalar ops still returned correct results.
- [x] AC-3: Docs/specs clearly state when a host fallback is allowed and
when TileOps must provide its own implementation. —
`docs/design/trust-model.md` and `.claude/domain-rules/manifest-spec.md`
updated.
- [x] AC-4: Existing and new reduction tests cover scalar inputs and any
documented exception boundary. — Scalar conformance and reduction
scalar-input suites extended (+14 nodes).

## Test node delta

```
File                                           Base    HEAD    Delta
--------------------------------------------------------------------
tests/ops/test_reduce_scalar_conformance.py      96     108      +12
tests/ops/test_reduction_scalar_input.py         50      52       +2
--------------------------------------------------------------------
TOTAL                                           146     160      +14
```

**Justification:** Each new node pins a previously-uncovered scalar
contract: prod / welford (var, std, var_mean) / logical (all, any) /
count_nonzero scalar fast paths, plus the invalid-DOF exception boundary
required by AC-4. No redundant coverage; the additions are the minimum
needed to lock the new native 0-D path.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-14 17:57:40 +08:00
Cao Ying 2eb6f7de26
[Doc] introduce docs/known-issues.md with logical reduction perf entry (#1452)
## Summary

- Add `docs/known-issues.md` as a project-level register of acknowledged
gaps where the spec/manifest is correct but the implementation has a
non-trivial performance, precision, or specialization deficit.
- First entry: logical reductions (`AllFwdOp`, `AnyFwdOp`,
`CountNonzeroFwdOp`) — after #1451 aligned manifest signatures with
PyTorch (`int32`, `int64`, `bool`, `complex64`, `complex128`), the
TileLang kernel still only ingests `float16/bfloat16/float32`; the op
layer bridges via `to_logical_float32`, materializing a full-sized
`float32` intermediate before dispatch.
- Entry tabulates per-dtype DRAM costs, notes the redundant in-kernel
`!= 0` predicate on the pre-converted path, and the
non-dtype-specialized `source.kernel_map`.
- Points at where a future fix lands: dtype-specialized TileLang kernels
that consume `int{32,64}`, `bool`, `complex{64,128}` directly and
predicate on first load.

## Test plan

- [x] pre-commit passed (mdformat applied)
- [x] no code changes — docs-only

## Additional context

This is a register, not a bug tracker. Entries here are *acknowledged*
deviations the project chooses not to fix immediately. Per-issue
tracking still goes through GitHub issues.

Future entries should follow the same template: affected ops, spec
status, gap, concrete costs, why not fixed now, pointer for a fix.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-14 17:19:50 +08:00
Cao Ying 275b18a50f
[Fix][Reduction] support 0-D scalar input on status-implemented reductions (#1446)
Closes #1442

## Summary

- Add 0-D scalar fast path to the seven `status: implemented` reduction
families (Sum/Mean/Amax/Amin/All/Any/CountNonzero) and the Welford trio
(Var/Std/VarMean) so `dim ∈ {None, 0, -1, (), []}` matches `torch.<op>`
on scalar input.
- Welford scalar path preserves PyTorch's `nan` (default `correction=1`)
and `UserWarning` emission; multi-element kernel path is untouched.
- Subclass guard limits the fast path to the covered reduce ops; other
subclasses fall through to the pre-existing rank-guard `ValueError`
unchanged.
- New conformance test `tests/ops/test_reduction_scalar_input.py`
parametrizes 10 ops × 5 dim forms against PyTorch reference (50 cases).
- Tighten reviewer rules so `New-path coverage` and `Critical-path
floor` are anchored to "output-distinguishing input" rather than literal
control-flow branches, ruling out scaffold-style negative tests on no-op
aliases. Touches `.claude/review-checklists/pre-review.md`,
`.claude/review-checklists/approval-gate.md`,
`docs/design/trust-model.md`.
- Manifest and kernels are not modified — any `shape_rules` update is a
separate manifest-only PR per the trust-model carve-out.

## Test plan

- [x] AC-1: `pytest tests/ops/test_reduction*.py
tests/ops/test_logical_reduce.py tests/ops/test_reduce*.py` green on
CUDA
- [x] AC-2: `python scripts/validate_manifest.py` exits 0; `pytest
tests/test_validate_manifest.py` → 218 passed
- [x] AC-3: Sum/Mean/Amax/Amin/All/Any/CountNonzero scalar input matches
`torch.<op>` for `dim ∈ {None, 0, -1, (), []}`
- [x] AC-4: Var/Std/VarMean scalar input matches `torch.<op>` value,
shape `()`, and `UserWarning`
- [x] AC-5: `tests/ops/test_reduction_scalar_input.py` → 50 passed
(cartesian op × dim form)
- [x] AC-6: rank-≥1 reduction suites — no regression
- [x] AC-7: `git diff upstream/main -- tileops/manifest/
tileops/kernels/` is empty
- [x] pre-commit passed

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-14 11:05:49 +08:00
lcy-seso 267bd50e84 [Maintain][Docs] drop generative-op carve-out from manifest design doc
The deleted section described an implementation workaround (carrier
scalar to satisfy the >=1-input invariant) as if it were a design
decision. The hack and its mechanism live in validator code + manifest
YAML comments where they belong; design docs should not document
transient impl details.

Co-Authored-By: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-13 12:34:49 +08:00
lcy-seso ff32723d16 [Maintain][Manifest] add generative-op carve-out + restore alibi/sinusoidal
L1 and C4 forward()-arity checks now skip when ``ref_api: "none"`` and
``forward()`` takes zero positional args. The five elementwise ops
without manifest entries -- AlibiFwdOp, SinusoidalFwdOp, plus the three
fused-gated entries already present -- are all covered by the manifest;
the two generative ops carry a single ``device_carrier`` scalar input to
satisfy ``test_every_signature_has_inputs_and_outputs`` while the new
carve-out unblocks the forward()-order check.

Co-Authored-By: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-13 12:34:49 +08:00
Cao Ying 7a153cfb30 [Refactor][MANIFEST] unify per-element flop convention for activation/clamp (#1389)
Closes #1379

## Summary

- Document a single per-element FLOP convention in
`docs/design/roofline.md` (Convention subsection, dated) covering
arithmetic, transcendental, compare-and-select, and clamp ops.
- Apply the convention to activation helpers (`relu`, `leaky_relu`,
`prelu`, `gelu`, `silu`, `swish`, `tanh`, `sigmoid`, `hardtanh`) and the
clamp/min-max family (`clamp_scalar`, `clamp_min`, `clamp_max`,
`maximum`, `minimum`) in `tileops/perf/formulas.py` and the
corresponding manifest entries.
- Each helper now carries a one-line `# FLOPs: ...` derivation comment;
the clamp/min-max family returns the same per-element FLOP count.
- Add `scripts/perf/flop_convention_delta.py` plus checked-in artifacts
under `docs/perf/` so the before/after delta is reproducible without a
GPU.

## Test plan

- [x] AC-1: `pytest tests/perf/ tests/test_validate_manifest.py` — 288
passed.
- [x] AC-2: `docs/design/roofline.md` Convention subsection present and
dated.
- [x] AC-3: each activation helper carries a `# FLOPs: ...` derivation
comment.
- [x] AC-4: `hardtanh` / `clamp_scalar` / `clamp_min` / `clamp_max` /
`maximum` / `minimum` return identical per-element FLOP counts.
- [x] AC-5: before/after roofline-utilization table for ≥3
representative benches included below.
- [x] AC-6: external tracker flip — out of scope.

## Benchmark

## Per-element FLOP convention — before/after delta

Formula-only evaluation (no GPU required). Reproduce with:

```bash
python scripts/perf/flop_convention_delta.py \
  --out docs/perf/flop_convention_delta.csv
```

| family | op | label | shape | dtype | flops before | flops after |
flops delta | bytes before | bytes after | bytes delta |
| ---------- | ------------- | -------------------- | --------- |
------- | -----------: | ----------: | ----------: | -----------: |
----------: | ----------: |
| activation | ReluFwdOp | hidden-state-prefill | 2048x4096 | float16 |
16,777,216 | 8,388,608 | -8,388,608 | 33,554,432 | 33,554,432 | 0 |
| clamp | HardtanhFwdOp | hidden-state-prefill | 2048x4096 | float16 |
33,554,432 | 8,388,608 | -25,165,824 | 33,554,432 | 33,554,432 | 0 |
| min-max | ClampFwdOp | elementwise-16M | 4096x4096 | float16 |
33,554,432 | 16,777,216 | -16,777,216 | 134,217,728 | 134,217,728 | 0 |

`flops before` reflects the coefficients on `upstream/testbed`
immediately before the convention commit. `flops after` is evaluated
from the manifest formulas (`ReluFwdOp` / `HardtanhFwdOp`) or
`clamp_fwd_roofline` (`ClampFwdOp`) on the current checkout. Byte counts
are unchanged by the convention. For these elementwise workloads
`memory_time` already dominates, so the FLOP-coefficient reduction
shifts each workload further into the memory-bound regime without
changing predicted achievable bandwidth.

See
[docs/perf/flop_convention_delta.md](docs/perf/flop_convention_delta.md)
and
[docs/perf/flop_convention_delta.csv](docs/perf/flop_convention_delta.csv).

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-12 09:57:31 +08:00
Cao Ying 70d30f0ed1
[Design] Trust-model rules become review lens; retire strict-mode auto-reject (#1424)
## Summary

- Remove strict / principled mode split in
`.claude/review-checklists/pre-review.md`.
- Trust-model rules apply uniformly as a review lens — no auto-reject
keyed on directory layout.
- Provenance labels (`automated` / `needs-review` / `nightshift`) mark
origin only, not rule semantics.
- State the substantive criteria reviewers cite (oracle origin, coverage
set, new-path coverage).

## Motivation

The retired rule treated a semantic property (joint authorship of oracle
+ impl enables construction-by-design cheating) as a syntactic check
(whether `tests/` and `tileops/` appear in the same diff). Observed
cost:

- **False positives** block legitimate contract-consistency fixes — see
PR #1408, #1409.
- **Induced evasion** — dev agent splits joint changes with "tests
follow in sibling PR"; sibling never lands. PR #1410 hit the
unsatisfiable variant (new behavior branch with no prior coverage;
kernel-only rejected by reviewer, joint diff rejected by lexical rule).
- **Residual false negatives** — agent-fabricated oracles in test-only
diffs are invisible to the lexical check anyway.

Trust-model semantics in `docs/design/trust-model.md` are intentionally
unchanged — the design intent is preserved as a review lens.
Semantic-level enforcement (oracle-origin classifier, coverage-delta
lint) is future work, not part of this PR.

## Test plan

- [ ] Documentation-only change; no runtime impact.
- [ ] Spot-check that `pre-review.md` still loads cleanly as a
review-skill input (no broken cross-refs).

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-11 18:16:25 +08:00
Cao Ying 0451c058a9
[Chore][Release] merge testbed into main: op-family alignment + new ops (#1380)
## Summary
Merge 27 commits accumulated on `testbed` into `main`. Covers
manifest-spec alignment for the elementwise (unary/binary/multi-input),
reduction, and normalization families; new ops (LerpTensorFwd,
Reciprocal, fused tanh-GELU, InstanceNormFwdOpNoAffine,
GroupNorm/InstanceNorm no-affine variants); shared dispatch/refactor
cleanup; and supporting manifest DSL (`promote_int_to_float`), tests,
and process docs.

## Highlights
- **Family alignment to manifest spec**: elementwise_unary_activation
(#1211), elementwise_binary (#1222), elementwise_multi_input (#1229),
reduction (#1235), normalization (#1240).
- **New ops / variants**: LerpTensorFwd (#1264), Reciprocal +
`promote_int_to_float` DSL (#1252), fused tanh-GELU (#1262),
GroupNorm/InstanceNorm no-affine (#1274, #1284),
InstanceNormFwdOpNoAffine running-stats (#1375).
- **Refactors**: activation shared bases (#1230), dispatch_kernel
routing (#1260, #1368), normalization hotfix cleanup (#1256),
elementwise_binary torch-fallback removal (#1248).
- **Tests / docs**: manifest-driven L1 parity for elementwise_binary
(#1370), drop manifest-mirror tests (#1374), FLIP_STATUS carve-out
(#1258), spec-only -> implemented promotion rule (#1271).

## Test plan
- [ ] CI green on the PR branch
- [ ] \`make test\` passes locally
- [ ] Manifest validator and strict-parity gates pass

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-09 00:03:37 +08:00
Cao Ying 56977a02db
[Design][Ops] forbid promoting per-op exceptions to base-class shared mechanisms (#1283)
Closes #1273

## Summary

- Add normative bullet to `.claude/domain-rules/ops-design.md`: per-op
workarounds MUST NOT be promoted to a base-class shared mechanism within
the same op-family migration PR; promotion requires a separate design PR
proving genuine family invariant.
- Drop the stale "no parity opt-out" trailing clause in
`docs/design/ops-design.md` Step 5 Validation paragraph
(`parity_opt_out` field was retired in PR #1292; defensive reminder no
longer relevant).
- Qualify the same paragraph: validator runs L2/L3 only on `status:
implemented` entries; `spec-only` entries skip. Parity failures route to
`strict_errors`, downgraded to warnings in advisory mode and blocking
under `--strict` / `MANIFEST_STRICT_BLOCKING=1`.

## Test plan

- [x] pre-commit passed
- [x] manual: `python scripts/validate_manifest.py` runs unchanged (this
PR is docs-only)

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-08 19:42:50 +08:00
Cao Ying a127738c3e
[Feat][Manifest] strict parity CI for status: implemented ops (#1292)
Refs #1291 · backlog + flip-to-blocking in #1372

Adds C1–C7 strict parity gates to `scripts/validate_manifest.py` for
`status: implemented` ops; `spec-only` skip. Removes the
`parity_opt_out` escape hatch.

| Check | Catches |
| --- | --- |
| C1 / C2 | `_infer_output_shapes` / `_validate_dtypes` disagrees with
manifest |
| C3 | ctor param defaults + kw-only flag mismatch |
| C4 | forward positional names mismatch |
| C5 | Slot S12+S13: `__init__` lacks `kernel_map` kwarg, or body never
calls `self.dispatch_kernel(...)` / `super().__init__(...)` (static
inspect + AST; no construction) |
| C6 / C7 | `_validate_dtypes` / `eval_roofline` is the base stub |

Default **advisory** — failures → warnings, exit 0. `--strict` /
`MANIFEST_STRICT_BLOCKING=1` flips to blocking. `--levels` gates:
C3/C4/C5 ↦ `signature`, C6 ↦ `dtype`, C7 ↦ `bench`.

## Test plan

- [x] pre-commit
- [x] `pytest tests/test_validate_manifest.py -q` — 205 passed (+19
nodes)
- [x] `python scripts/validate_manifest.py` — exit 0 (advisory)
- [x] `python scripts/validate_manifest.py --strict` — exit 1 (83
errors: 22 ctor + 46 + 15 stub; tracked in #1372)

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-08 17:33:05 +08:00
Cao Ying 3bd33cf069
[Chore][Pipeline] eliminate dev-pipeline cross-stage contradictions (#1208)
Closes #1204

## Summary

Audit-driven cleanup of dev-pipeline cross-stage contradictions surfaced
by the nightshift run on PR #1203 / issue #1192. Four root causes
addressed:

- **Knowledge-loading asymmetry** (A.1) — `.foundry/config.json` now
populates `knowledge.reviewer[]` (10 paths) and `knowledge.gatekeeper[]`
(3 paths). `knowledge.developer[]` stays `null` because `CLAUDE.md`
auto-load + the on-demand domain-rules table already cover the
developer's stage. Foundry codex reviewer (Layer 1, dev rounds) and
`review-tileops` (Layer 2, post-PR) now converge on a single rule set.
- **Duplicate trust-model documents** (A.2) —
`.claude/rules/manifest-trust-model.md` rule 1 (which restated the
design doc's Manifest stage boundary) is now a one-line pointer to
`docs/design/trust-model.md §Manifest`; rules 2–4 stay verbatim as
operational rules.
- **Manifest YAML hygiene** (A.3) — 18 process-metadata comments removed
across five manifest files. New rule lives in `manifest-spec.md`:
technical content permitted, project-management metadata bound to a
specific issue / PR / commit / round forbidden. Test phrased so it
survives renaming.
- **`MAY READ` whitelist confusion** (B.1) — removed the `MAY READ` axis
from every stage in `trust-model.md`. Reads are not policed; the trust
model controls writes (`OWNS`, `MUST NOT WRITE`) and structural
couplings (`MUST NOT`, optional).

C.1 (reviewer architecture) is not changed in this PR — the two-layer
split is preserved. Revisit only if subsequent runs show persistent
disagreement.

## Test plan

- [x] AC-1: Modified files pass unit tests (`pytest -q
tests/test_ops_manifest.py
tests/test_validate_manifest.py::TestIntegration` → 23 passed;
pre-commit clean).
- [x] AC-2: `.foundry/config.json` populated correctly —
`knowledge.reviewer[]` 10 paths, `knowledge.gatekeeper[]` 3 paths,
`knowledge.developer[]` null. Every listed path resolves.
- [x] AC-3: `manifest-trust-model.md` rule 1 replaced by pointer; rules
2/3/4 retained verbatim.
- [x] AC-4: `grep -rE 'follow.?up|fix in.*PR|drift'
tileops/manifest/*.yaml` returns 0 matches.
- [x] AC-5: `code-style.md` no longer contains the `status:
spec-only`-line issue-number allowance.
- [x] AC-6: `manifest-spec.md` carries the no-process-metadata rule with
a renaming-survival test.
- [x] AC-7: `trust-model.md` has zero `MAY READ`; every stage has `OWNS`
+ `MUST NOT WRITE`; `MUST NOT` permitted only for non-write coupling
(currently only Benchmark); page intro documents the taxonomy.
- [x] AC-8: Every `status: spec-only` op has a corresponding open
`follow-up` issue. New: #1205 (MoE × 11), #1206 (Conv1dBiasFwdOp), #1207
(Conv1dFwdOp); relabelled: #1192, #1196, #1198, #1194, #1200, #1142,
#879, #402.
- [x] AC-9: Regression test — re-run `/foundry:pipeline 1192
--nightshift` on a fresh `.foundry/runs/issue-1192/` after merge;
expected zero T010-equivalent findings and zero diff-correctness
violations from `review-tileops` tied to the four root causes.
- [x] AC-10: `pre-commit run --all-files` passes.

## Follow-up

- #1209 — lift the reduction-dim `shape_rules` patterns into a shared
helper module so manifest YAMLs reference helpers by name instead of
pasting Python strings.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-05 22:26:58 +08:00
Cao Ying cf6dc5b0fc
[Doc][Skills] regroup tileops-skills.md by scope and admit workflow skills (#1141)
Closes #1136

## Summary

- Restructured `docs/tileops-skills.md` "Skills in detail" into four H3
sections: per op, per op family, manifest, workflow.
- Admitted workflow skills (`review-tileops`, `resolve-tileops`,
`follow-up`) as a concise table; op/manifest skills retain their full
per-block detail.
- Updated the at-a-glance matrix and intent table to match the new scope
ordering, and dropped the workflow-out-of-scope carve-out from
Maintenance.

## Test plan

- [x] AC-1: `docs/tileops-skills.md` "Skills in detail" has exactly four
H3 sections (per-op, per-op-family, manifest, workflow) and no other
H3s.
- [x] AC-2: Every directory under `.claude/skills/` is documented by
exactly one bold-led skill block in the doc.
- [x] AC-3: At-a-glance matrix row labels read 'per op' / 'per op
family' / 'manifest' / 'workflow', in that order.
- [x] AC-4: Maintenance section no longer carves out 'process / workflow
skills' as out of scope.
- [x] AC-5: mdformat passes; no Codex/Claude/Anthropic/OpenAI/GPT tool
names in `docs/**/*.md` or
`.claude/skills/*/{SKILL.md,README.md,*.yaml}` outside path references.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-02 17:23:44 +08:00
Cao Ying 1fa9f7b8f6
[Chore][Skill] move follow-up skill (formerly create-follow-up-issue) under TileOPs (#1131)
## Summary

- Move `create-follow-up-issue` skill from `~/.claude/skills/` (the
personal ibuki config) into this repo at
`.claude/skills/create-follow-up-issue/SKILL.md`.
- Skill is generic but in practice only used here; versioning it with
the project means edits go through normal review.

## Test plan

- [x] pre-commit passed (mdformat, codespell, gitleaks)
- [ ] N/A — pure relocation, no code paths changed

## Additional context

Skill content is unchanged from the ibuki copy. Removal commit on the
ibuki side stays local until a separate `/push-ibuki`.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-05-01 08:49:54 +08:00
Cao Ying 6e9fdddff7
[Refactor][Skill] review-tileops prompt: trim redundancy, fold verify into Task (#1127)
## Summary

- **Move iteration block into `## Task`.** Drop the standalone `##
Iteration round (round N of M)` wrapper in round 2+. Round number is
already in the prompt title, and the verify-prior-blockers /
no-new-problems content is the actual task instruction — fold it under
`## Task` as a one-line round-2+ prefix.
- **Trim redundancy with `procedure.md`.** procedure.md is in Codex's
session memory after round 1. Drop everything in `compose_prompt` that
restates it: the `## Project-specific regression guards` subtitle
(handled by procedure.md step 2), the `(inbox, one-shot)` parenthetical,
and the `## Task` body lines "Read the diff … Apply the loaded
checklists. Submit ONE atomic review …" (steps 1/3/7).
- **Doc fix:** point the `TileOPs.github.io` link in `docs/README.md` at
the actual Pages site (`tile-ai.github.io/TileOPs.github.io/`) rather
than the source repo on GitHub.com.

## Test plan

- [x] pre-commit passed
- [x] `bash -n loop.sh` syntax check

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-04-30 23:44:14 +08:00
Cao Ying a4d4b27c41
[Doc] move tileops-skills.md out of docs/design/ (#1108)
## Summary

- Move `docs/design/tileops-skills.md` → `docs/tileops-skills.md`. It is
a developer user guide for picking the right repo skill, not a design
document, so it does not belong under `docs/design/`.
- Update cross-references in `CLAUDE.md` and `docs/README.md`.
- Rebase the moved file's relative links to `.claude/skills/...` (one
fewer `../`).

## Test plan

- [x] pre-commit passed
- [x] no code changes; doc-only move

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-04-29 11:01:38 +08:00
Maybe_Jayden e3ac5c28d1
[Feat][MoE] Add modular MoE interface: ABCs + FusedMoe orchestration (#1087)
## Summary

- Introduce strategy-pattern interface for MoE forward pass — pluggable
expert GEMM and EP/quantization backends without changing core routing.
- New ABCs in `tileops/ops/moe/abc.py`: `MoEExperts` /
`MoEExpertsModular`, `MoEPrepareAndFinalize`, `WeightedReduce` /
`WeightedReduceNoOp`, and `PrepareResult` dataclass carrying `hidden_q /
scale / topk_weights / topk_ids` from `prepare()` to `apply()`.
- Concrete implementations: `MoEPrepareAndFinalizeNoDPEP` (pass-through,
no EP/quantization), `MoEExpertsNopadFwdOp` (wraps existing nopad
pipeline), `MoEExpertsPaddedFwdOp` (wraps existing padded pipeline).
- `FusedMoe` refactor (`fused_moe.py`): accepts optional
`prepare_finalize` and `experts` strategy args, defaults preserve
existing behavior; `forward()` delegates via the ABC protocol; fixes
silent drop of `expert_map` (now stored and forwarded to both
`prepare()` and `apply()`).
- Manifest (`tileops/manifest/moe.yaml`): add `MoEExpertsNopadFwdOp`,
`MoEExpertsPaddedFwdOp` (replace prior `FusedMoeExperts*FwdOp` modular
entries), `FusedMoeFwdOp` (softmax scoring), `FusedMoeFwdCbFwdOp`
(sigmoid + correction_bias). Add `fused_moe_fwd_bytes` roofline formula
in `tileops/perf/formulas.py`. Both experts entries set
`bench_manifest_driven: false` — strategy objects are not `Op`
subclasses and cannot implement `eval_roofline()`.
- Tests (`tests/ops/test_moe_experts_nopad.py`): 15 smoke tests covering
ABC contracts and the three concrete strategies, parametrized over
`bfloat16` and `float16`; correctness verified against
`FusedMoeExperts*FwdOp` references.
- Benchmark (`benchmarks/ops/bench_moe_experts_nopad.py`):
manifest-driven workloads, benches nopad and padded side-by-side;
baseline vLLM `fused_experts` when available, torch `index_add_` loop
fallback.
- Docs (`docs/design/ops-design.md`): add "MoE Modular Protocol" section
— ABC strategy pattern, component table, `FusedMoe.forward()` data flow,
EP/quantization extension points.

## Test plan

- [x] `pre-commit run` passed (SKIP=gitleaks)
- [x] `pytest tests/ops/test_moe_experts_nopad.py` — 15 passed (bfloat16
+ float16)
- [x] `pytest tests/ops/test_moe_fused_moe.py` — 20 passed (no
`FusedMoe` regression)
- [x] Merge conflict with `main` resolved (`ops_manifest.yaml` →
`tileops/manifest/moe.yaml` migration)
- [ ] Full CI smoke suite

## Regression

- `FusedMoe` defaults to `MoEPrepareAndFinalizeNoDPEP` +
`MoEExpertsNopadFwdOp` / `MoEExpertsPaddedFwdOp`, so existing call sites
remain backwards-compatible.
- `expert_map` was previously silently ignored; this PR wires it
through. Confirmed via existing `test_moe_fused_moe.py` (20 passed).

## Additional context

- Lays groundwork for future EP dispatch and per-expert quantization
backends without touching `FusedMoe.forward()`.
- Both new experts ops carry `bench_manifest_driven: false` because
strategy objects do not subclass `Op`; manifest-driven benchmarking will
be revisited when the modular path graduates to first-class `Op` status.

---------

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
2026-04-28 20:51:01 +08:00
Cao Ying 46d0918a79
[Refactor][Manifest] split elementwise.yaml into per-sub-family shards (#1095)
Closes #1093

## Summary

- Split `tileops/manifest/elementwise.yaml` (62 ops, ~2050 lines) into 4
sub-family shards by signature shape:
- `elementwise_unary_math.yaml` — 24 ops (~746 lines): unary math
(single-input single-output, math kernels)
- `elementwise_unary_activation.yaml` — 13 ops (~485 lines): unary
activations (single-input single-output, NN activations)
- `elementwise_binary.yaml` — 23 ops (~739 lines): binary (two-input
single-output)
- `elementwise_multi_input.yaml` — 2 ops (~84 lines): 3+ input or
multi-output
- Removed the monolithic `tileops/manifest/elementwise.yaml`; entries
moved verbatim via `ruamel.yaml` round-trip, preserving `family:
elementwise` on every entry.
- Updated `tileops/manifest/__init__.py` module docstring and
`docs/manifest.md` / `CLAUDE.md` to describe the sharded layout.

No op-code or kernel-code changes — pure manifest file partitioning.
Total op count stays at 114; `load_manifest()` / `load_workloads()`
behavior is dict-equal to upstream/main.

## Test plan

- [x] AC-1: `tileops/manifest/elementwise.yaml` removed; ops live in 4
sub-files, every shard < 1000 lines and ≤ 30 ops (largest: 746 lines /
24 ops).
- [x] AC-2: `load_manifest()` returns the same 114 op names with
dict-level content equal to upstream/main (verified via Python
comparison: `name_sets_equal=True`, `content_equal=True`).
- [x] AC-3: `python scripts/validate_manifest.py` exits 0 — "All
manifest checks passed" (58 warnings, unchanged from baseline).
- [x] AC-4: `pytest tests/test_ops_manifest.py
tests/test_validate_manifest.py tests/test_workloads_to_params.py` → 198
passed, matching upstream/main baseline (198 passed) at branch point.
- [x] pre-commit passed on all changed files.

## Additional context

Follows the same trust-model rules as PR #1092 (which sharded the
manifest into one-file-per-family): manifest-only PR, no op or kernel
changes bundled. The elementwise family was the only family still over
the ~1000-line / ~30-op soft cap after #1092 — this PR completes the
per-sub-family layout.

---------

Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com>
2026-04-28 19:30:40 +08:00
Cao Ying 9448edbfba
[Refactor][Manifest] flatten yaml top-level ops key (#1094)
## Summary

- Each family yaml previously wrapped its entries under a top-level
`ops:` key, costing one indent level on every entry without serving any
purpose — family scope is already conveyed by the filename, and the
loader just unwrapped the key.
- Flatten the format: each family file is now a direct mapping of op
name → entry. Updated 8 yaml files (4546 lines), the `tileops.manifest`
loader, the `--manifest-path` branch in `scripts/validate_manifest.py`,
the `tests/test_validate_manifest.py` fixtures, and the example in
`docs/manifest.md`.
- Pre-release project, so no compat shim — the old format is no longer
accepted.

## Test plan

- [x] pre-commit passed (full hook suite on all changed files)
- [x] `python scripts/validate_manifest.py` → `All manifest checks
passed.`
- [x] `pytest tests/test_validate_manifest.py
tests/test_gpu_smoke_policy.py` → 176 passed
- [x] `from tileops.manifest import load_manifest` smoke check → 114 ops
loaded

## Regression

- Loader migration is mechanical (`data.get("ops") or {}` →
`yaml.safe_load(...) or {}`). All 114 ops load identically; downstream
consumers (`load_workloads`, `validate_manifest`) untouched.
- `_load_manifest` private alias removed (no remaining callers in repo).

## Additional context

There is no near-term plan for family-level metadata. If that need
arises later, a wrapper can be reintroduced — YAGNI for now.

---------

Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com>
2026-04-28 16:58:35 +08:00
Cao Ying 19ee4c2b18
[Refactor][Manifest] split monolithic yaml into per-family package (#1092)
## Summary

- Split `tileops/ops_manifest.yaml` (monolithic, 4564 lines) into one
YAML file per op family under `tileops/manifest/` — `elementwise.yaml`
(62), `reduction.yaml` (19), `attention.yaml` (12), `normalization.yaml`
(10), `moe.yaml` (7), `scan.yaml` (2), `convolution.yaml` (2).
- Convert `tileops/manifest.py` into the `tileops/manifest/` package;
`__init__.py` scans the directory, merges all `ops` dicts, errors on
duplicate keys, caches with `lru_cache(1)`. Public API
(`load_workloads`) is unchanged; new helpers `load_manifest()` and
`manifest_files()` exposed.
- Update direct consumers: `scripts/validate_manifest.py` (default load
via package; `manifest_path=` kept for synthetic fixtures),
`tests/test_ops_manifest.py` (fixtures via `load_manifest()`),
`pyproject.toml` package-data, `.github/workflows/preflight.yml` path
globs.
- Synchronize all documentation: 22 `.md` files (CLAUDE.md, docs/,
.claude/skills/, .claude/domain-rules/, .claude/rules/) plus docstrings
in 5 benchmark files and `tileops/perf/formulas.py`. CLAUDE.md "Reading
the ops manifest" section and docs/manifest.md "Layout" section
rewritten for the new structure. `add-manifest` / `fix-manifest` skills
note family-file routing.

## Test plan

- [x] `python scripts/validate_manifest.py` passes
- [x] `pytest tests/test_ops_manifest.py tests/test_validate_manifest.py
tests/test_workloads_to_params.py` — 197 passed
- [x] pre-commit passed
- [x] `python3 -c "from tileops.manifest import load_manifest; assert
len(load_manifest()) == 114"` — round-trip preserves all 114 ops, no
duplicates

## Additional context

The split keeps the conceptual "single source of truth" (the merged
`ops` dict) while addressing the practical pain of one 4500+ line file.
Largest family file is `elementwise.yaml` at 2030 lines (62 ops); a
follow-up PR can sub-split elementwise into unary/binary/multi-input if
desired. No op or kernel code is touched, so this PR does not conflict
with the trust-model rule prohibiting joint op+manifest changes.

## Structural Readiness

All checks passed.

## Follow-up

Issues:
- #1093 — sub-split elementwise.yaml (2030 lines / 62 ops) into
unary/binary/multi-input sub-files.

---------

Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
2026-04-28 13:41:10 +08:00
Cao Ying aa349659eb
[Refactor][Skills] redesign add-manifest input + redivide labour with fix-manifest (#1066)
## Summary

- `add-manifest`: input is `op_name` + `ref_url`. Idempotent,
one-entry-per-invocation. Class-based source lookup. Required fields
BLOCK with `evidence_needed` instead of falling through.
- `fix-manifest`: scope reduced to `source.kernel_map` +
`signature.static_dims` (on-disk-derivable only). Reference-derivable
fields → `add-manifest`.
- `docs/tileops-skills.md`: trust-model redivided by source-of-truth.

## Test plan

- [x] `pre-commit` passes (mdformat, codespell, gitleaks).
- [x] Skill files have valid frontmatter (`name`, `description`).
- [x] Cross-links in `docs/tileops-skills.md` resolve
(`../.claude/skills/<name>/SKILL.md`).
- [ ] Smoke test: re-align one normalization op via `/add-manifest
<op_name> <pytorch_url>` in a follow-up PR (first real exercise of the
new contract).
- [ ] Smoke test: invoke `/fix-manifest <op_name> --field=kernel_map` on
an op missing only `kernel_map` (verifies the narrowed scope still works
for the most common gap).

---------

Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com>
2026-04-27 14:18:23 +08:00