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>
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>
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>
## 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>
## 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>
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>
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>
## 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.
## 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
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>
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>
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>
## 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>
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>
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>
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>
## 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>
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>
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>
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>
## 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>
## 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>
## 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>
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>
## 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>
## 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>
## 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>