forked from ccf-ai-infra/TileOPs-Metax
65 Commits
| Author | SHA1 | Message | Date |
|---|---|---|---|
ray24777
|
da567b4239 |
refactor(quant_swiglu): drop stage-profiling variants and their scripts
Remove the `stage=` builder parameter, the read_only/read_shared/read probe branches, the _block_sum_reduce macro, and stage_kernel() from the fused SwiGLU + per-channel FP8-cast kernel. These existed only to attribute fused time to pipeline stages for profiling and had no effect on the production path; the unconditional kernel body is byte-identical to the former stage="full". lru_cache returns to default maxsize=32. Delete scripts/profile_quant_swiglu_stages.py and scripts/quant_swiglu_stage_profile.sh, the only consumers. Verified: tests/ops/test_quant_swiglu_channel_cast_transpose.py 12 passed (bit-exact fp8 contract preserved). Co-Authored-By: Claude <noreply@anthropic.com> |
|
|
ray24777
|
2e196447c8 |
optimize(quant_swiglu): patch-mapped read + register absmax (5.9% sum)
Redesign the read pipeline of the fused SwiGLU+FP8-cast kernel after decomposing the read stage (read_only / read_shared / read / full): the HBM read was fine, but the shared-staging absmax was the hidden cost. - Non-transposed read: 4x4 register-patch mapping (each thread owns TILE_K rows x TILE_K cols, 1-2 rows per load instruction vs 8 for the 16-element row-strip) lifts read_only bandwidth ~0.95 -> ~1.27 TB/s at TILE_Y=64 with an unchanged shared footprint, so the fused kernel's occupancy is untouched. - Non-transposed absmax: folded into registers during the read from the patch values, dropping the act_shared read-back the reduction did before (read-pipeline overhead t1 - t_ro: +42% -> +22% on t8064-h4096). - read_only probe: log-depth tree reduction sink instead of a serial per-block sum (a ~17% C500 read-bandwidth penalty). - Tuning-config validation tightened to the patch read's real constraint (tile_x % (4 * thread_shared_step) == 0), which also rejects configs the old check silently let through as empty loops; the obsolete 16-element row-strip machinery (_best_vectorize_size, thread_global_step) is gone. - Stage framework: add read_shared diagnostic (read + staging, no absmax) to split the read pipeline into staging vs reduction cost. Full-kernel sum over the 10 manifest workloads (MetaX C500): 1698.6 -> 1598.2 us (-5.9%), from the non-transposed t8064 workloads (-8.8% to -11.5%); the transposed t4096 workloads are unchanged (their fragment T.copy + reduce_absmax path is not yet touched). 12/12 bit-exact correctness tests pass. Co-Authored-By: Claude <noreply@anthropic.com> |
|
|
ray24777
|
074cdf83ef |
feat: add roofline plot script and C500 profile
tileops/perf/profiles/c500.yaml | 47 +++++++++ Co-Authored-By: Claude <noreply@anthropic.com> |
|
|
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> |
|
|
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>
|
|
|
Cao Ying
|
9bda1ac537
|
[Refactor][ELEMENTWISE] move strategy into kernel config; drop ctor kwarg (#1778)
Closes #1768 ## Summary - Move elementwise `strategy` selection into the kernel `config` dict; delete the ctor kwarg and its pass-through plumbing from all elementwise kernels and Ops. - fp8/bool coercions and the `register_copy` broadcast downgrade stay inside the kernel; same kernel body selected for every (op, dtype, shape) as before. - Migrate strategy tests to config-based construction; add guards asserting no elementwise Op/kernel exposes a `strategy` kwarg. - Fused-gated strategy bench now records a torch baseline and the measured kernel object. - Validator: `"strategy"` removed from `_CTOR_INFRA_PARAMS` — reintroducing the kwarg on any op now fails validation (validator suite 142 passed). ## Test plan - [x] pre-commit passed; 335 passed across the six modified test files - [x] AC-1: no elementwise Op/kernel exposes `strategy` (signature guards, +2 test nodes) - [x] AC-2: register_copy broadcast-downgrade regression preserved under config form - [x] AC-3: elementwise GPU smoke tier green (219 passed, H200) ## Benchmark NVIDIA H200, CUDA 12.8, PyTorch 2.9.1+cu128, TileLang 0.1.11. Fused-gated explicit_parallel (4096, 4096) fp16: | Op | TileOPs (ms) | torch (ms) | Speedup | BW (TB/s) | | --- | ---: | ---: | ---: | ---: | | SiluAndMul | 0.0292 | 0.1130 | 3.87× | 3.45 | | GeluAndMul | 0.0332 | 0.1213 | 3.65× | 3.03 | | GeluTanhAndMul | 0.0294 | 0.1147 | 3.90× | 3.43 | All rows meet or exceed the documented bandwidth basis (3.04/2.72/3.38 TB/s) — perf-neutral or better. ## Regression `test_register_copy_downgrades_on_broadcast` PASSED — config-form downgrade matches PyTorch under broadcast strides. --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
f9e161144c
|
[Refactor][Tooling] table-driven validator checks, shared emit/probe helpers, case-table tests (#1776)
Closes #1767 ## Summary - Table-drive `check_l0`: per-field validators registered in an `_L0_SECTIONS` table; genuinely custom rules stay as small validators. - Extract L3-parity helpers `_probe_out_of_union` / reason-kind dispatch, shared by the combos and no-combos branches. - Add a shared per-op error emitter (`_emit_to`) used across L0/L1/L2/L3/C3 checks; consolidate shape-literal parsing into `_SHAPE_EQ_RE`/`_shape_eq_literals`. - Convert sibling pass/fail test pairs to case tables with shared scaffolding (`_sig`, `_infer_parity`, `_dtype_parity`, `_write_manifest`, `_fake_op_module`, `_strict_op`). - Combined pair shrinks 7917 -> 6571 lines (-1346); test nodes 142 -> 120 (31 removed named tests consolidated into 9 case-table tests, all intents preserved). ## Test plan - [x] Modified files pass unit tests (`pytest -q tests/test_validate_manifest.py`: 120 passed) - [x] Golden-file diff of validator output over the full manifest is empty (default / `--verbose` / `--strict`: exit codes, stdout, stderr byte-identical to main) - [x] Combined pair shrinks by >=1300 lines; all baseline test intents preserved - [x] `TestCompileContractRegistry` remains collected by the compile-contract-gate invocation - [x] pre-commit passed ## Regression Default, `--verbose`, and `--strict` full-manifest validator output and exit codes are byte-identical to main; 5000 randomized `check_l0` inputs show zero old-vs-new mismatches. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
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> |
|
|
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> |
|
|
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>
|
|
|
Cao Ying
|
d621d44e2b
|
[Test][TestInfra] consolidate root tests into validator-owned schema checks (#1756)
Closes #1753 ## Summary - Move manifest schema policies into `scripts/validate_manifest.py`. - Consolidate validator tests into behavior-grouped case tables. - Relocate the unique `Op.autotune()` unit coverage and retain two roofline formula smokes. - Remove duplicate tier-validation, reclaim-action, compile, and manifest tests. ## Test plan - [x] pre-commit passed - [x] `pytest tests --ignore=tests/ops --ignore=tests/kernels --ignore=tests/perf`: 184 passed - [x] `python scripts/validate_manifest.py --strict`: exits 0 - [x] Root `tests/` nodes: 269 → 184 (-85); unique coverage retained or moved to the validator --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
b439c6d34f
|
[Feat][MANIFEST] Add torch_compile_fullgraph L0 check and compile-contract registry (#1749)
Part of #1748 (step 1 of 3). ## Summary - Validate optional `torch_compile_fullgraph`: literal `true` only and implemented ops only. - Add one lazy compile-evidence registry backed directly by existing fullgraph test cases. - Register 64 ops; manifest declarations and the equality gate remain follow-ups in #1748. ## Test plan - [x] 193 CPU validator tests - [x] 87 CUDA compile tests - [x] Manifest validation, pre-commit, and GPU smoke passed ## Test node delta `tests/test_validate_manifest.py`: 226 → 193 (-33); removed same-branch duplicates only. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
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> |
|
|
Ang Gao
|
a90c28c065
|
[Fix][Manifest] Keep UnaryOp dispatch visible to validator (#1708)
Fixes the unary-op manifest validation regression surfaced around #1707, using the reviewer-preferred approach from this PR discussion. ## Summary - Keep `UnaryOp.__init__` calling `dispatch_kernel(...)` directly, so the S13 manifest rule remains simple and statically visible - Replace the previous helper-call validator approach with explicit unary customization hooks: `_build_kernel_instance(...)` and `_resolve_output_dtype(...)` - Let `LogicalNotFwdOp` override those hooks for the bool-storage path while still exposing public `OUTPUT_DTYPE = torch.bool` - Tighten `validate_manifest.py` source lookup so it locates the exact class `__init__` AST node instead of relying on broader helper-call traversal - Update bool-like elementwise dtype coverage to assert comparison/logical kernels expose `torch.bool` as their public output dtype ## Problem #1705 refactored unary elementwise setup so `UnaryOp.__init__` reached `dispatch_kernel(...)` through `_prepare_unary_instance(...)`. That made the runtime behavior work, but it broke the S13 manifest invariant: the validator expects each op constructor to expose dispatch directly from `__init__`. The first version of this PR taught the validator to chase `self.<helper>()` calls. After Ibuki's review, we changed direction because that made the validator more permissive than the contract we actually want. The better fix is to keep dispatch explicit in the op constructor and move only the customization points into helpers. Separately, op tests exposed that bool-like comparison/logical kernels should advertise the public torch dtype (`torch.bool`) even when their internal storage path uses an integer/uint8 representation. ## Fix `UnaryOp.__init__` now performs the dispatch in place again: - `_build_kernel_instance(...)` prepares the kernel instance - `_resolve_output_dtype(...)` resolves the public output dtype - `dispatch_kernel(kernel_map)` remains directly visible in `__init__` `LogicalNotFwdOp` customizes the two hooks it actually needs instead of routing the whole constructor through a dispatch helper. The manifest validator no longer needs to follow arbitrary helper calls for S13, and the regression test now rejects helper-only dispatch. The bool-like elementwise dtype test now documents the public contract: comparison and logical kernels expose `torch.bool`; any uint8/int storage detail is internal to the implementation. ## Test plan - [x] `pytest -q tests/ops/test_elementwise_config_dtype.py::test_bool_like_elementwise_kernels_expose_torch_dtype_output` - [x] `pytest -q tests/test_validate_manifest.py::TestStrictParityC5Dispatch` - [x] `python3 scripts/validate_manifest.py` - [x] `python3 -m ruff check scripts/validate_manifest.py tests/test_validate_manifest.py tests/ops/test_elementwise_config_dtype.py` - [x] `python3 -m py_compile scripts/validate_manifest.py` - [x] `git diff --check` |
|
|
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. |
|
|
Cao Ying
|
1d99d4cda2
|
[CI] retire run-local venv across CI; install via image-baked stack (#1606)
## What
Full RFC §4 cleanup on top of the ephemeral containerized-runner
migration. The runner image bakes tilelang + the runtime/dev stack, so
CI stops building/copying per-run venvs and installs only tileops
(`--no-deps`) via `scripts/ci/install_tileops.sh`.
- **gpu-smoke**: drop `Set up Python`, `Resolve runtime state` (venv
hash/copy/mtime-sync/divergence guard), `Cleanup isolated fork state`;
install via `install_tileops.sh`; run with the image's `python3`;
trust-route PRs by collaborator permission (write/maintain/admin →
resident `nightly` pool; everyone else + lookup failure → `fork` pool,
fail-closed); reclaim the `/ci-cache` layout (no wheels, no tool-cache
prune); per-test `--timeout` + a `timeout-minutes` backstop so a wedged
kernel cannot hold the single runner.
- **nightly**: `setup_nightly_venv.sh` → `install_tileops.sh`; system
`python3`; cache env `/data7` → `/ci-cache`; per-test timeout on the
pytest runs.
- **runner-maintenance**: drop the retired `venv` runs-on label (→
`nightly`).
- **preflight**: pin CPU installs with `constraints.txt`.
- **pyproject / constraints**: tilelang as a compatibility range
(constraints pins its deps, not tilelang itself); pin `pytest-timeout`.
- **reclaim action + verify script**: `/ci-cache` layout, drop
`WHEEL_DIR`; trim stamp on the persistent cache.
- **delete**: `setup_nightly_venv.sh`, `ci_venv_hash.py` and their
obsolete tests.
## In scope vs follow-up
To keep CI dispatching, installing and importing cleanly on the new
stack, this PR DOES adapt the import surface: `gqa_fwd_fp8` guards its
`from tvm import tir` with a sentinel that raises a targeted error only
when the kernel is built, and the fp8-GQA / topk-selector
kernel-building smoke cases are skipped via a `tvm.tir` availability
gate (with a focused test for the gate). So **CI is green at import**,
not red.
**Out of scope (follow-up):** the actual kernel migration off `tvm.tir`
(`tir.call_extern` → `tilelang.language` `T.*`) and any other new-stack
kernel regressions surfaced by gpu-smoke (e.g. a build hang in
`FP8LightingIndexerKernel`). Those are tracked separately; until they
land, the corresponding smoke cases are skipped or fail on the per-test
timeout rather than wedging the runner.
---------
Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
fefcb10a8b
|
[CI] Reproducible multi-stage runner image + single dependency source (#1600)
Closes #1599 ## Summary - Repo-side foundation of the CI dependency / cache redesign — additive and files-only; nothing consumes these files yet (the workflow cutover is a separate follow-up). - `constraints.txt`: single source of exact version pins for the CI/runner stack (torch, triton, apache-tvm-ffi, and tilelang runtime deps). - `scripts/ci/install_tileops.sh`: requires `tilelang` already present, then `pip install -e . --no-deps -c constraints.txt`; fails clearly if tilelang is missing (so pip cannot drift torch / apache-tvm-ffi). - `.github/runner/Dockerfile`: rewritten as a multi-stage build from the public `nvidia/cuda:12.9.1-devel-ubuntu22.04` base (`runtime` builds python3.12 via the deadsnakes PPA → `post-fa3` → `fullstack` → `final`); bakes no TileOPs source and no runner credentials. ## Test plan - [x] AC-1: `constraints.txt`, `scripts/ci/install_tileops.sh`, and the rewritten multi-stage `.github/runner/Dockerfile` exist with the requested structure (stages runtime/post-fa3/fullstack/final, public CUDA 12.9.1 base, constraints-only Docker COPY, tilelang preflight, `--no-deps` install). - [x] AC-2: `scripts/ci/install_tileops.sh` passes `shellcheck` (0.11.0, no diagnostics). - [x] AC-3: `.github/runner/Dockerfile` passes `hadolint --failure-threshold error` (2.12.0; only warning/info findings). - [ ] AC-4: **DEFERRED-TO-MANUAL on a GPU build host.** The image is intentionally NOT built in CI by this issue. Manual validation (image builds; `torch.version.cuda == "12.9"`; `tilelang` imports; cuBLAS matmul/bmm/einsum probe passes; `pytest -m smoke` passes) is performed by a maintainer on a host with an NVIDIA GPU + nvcc. Reviewers should not expect a CI image build here. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
ChongLi
|
bcce329067
|
[Maintain][Manifest] Add max pool specs (#1578)
## Summary - Add spec-only manifest entries for MaxPool1d/2d/3d values-only and indices-returning variants. - Model `return_indices` as fixed-output variant entries: `MaxPool*dFwdOp` covers `return_indices=False`, and `MaxPool*dIndicesFwdOp` covers `return_indices=True` with `output` plus `int64` `indices`. - Keep all MaxPool `source` paths unset as `null` while the entries remain spec-only, so the manifest does not point at nonexistent or mismatched Op/Kernel/Test/Benchmark files. - Add a narrow validator guard so targeted L1 signature checks emit a warning and skip implementation class resolution for `status: spec-only` entries with `source.op: null`. ## Test plan - `python scripts/validate_manifest.py --levels schema --check-op MaxPool1dFwdOp` - `python scripts/validate_manifest.py --levels schema --check-op MaxPool2dFwdOp` - `python scripts/validate_manifest.py --levels schema --check-op MaxPool3dFwdOp` - `python scripts/validate_manifest.py --levels signature --check-op MaxPool1dFwdOp` - `python scripts/validate_manifest.py --levels signature --check-op MaxPool2dFwdOp` - `python scripts/validate_manifest.py --levels signature --check-op MaxPool3dFwdOp` - `python -m pytest tests/test_ops_manifest.py -q` - `python -m pytest tests/test_validate_manifest.py -q` Closes #1577 |
|
|
Ang Gao
|
adbb005f18
|
[Bench][Attention] Migrate benchmarks to manifest workloads (#1565)
## Summary - migrate attention benchmark params to manifest workloads and shared manifest param helpers - enable `source.bench_manifest_driven: true` for validated attention benchmarks - add attention roofline formulas needed by `ManifestBenchmark` / `op.eval_roofline()` - keep unsupported GQA paged decode `page_size=16` cases explicit with skip reasons Closes #1560. ## Validation - `ruff check scripts/validate_manifest.py benchmarks/ops/attention tileops/perf/formulas.py` - `git diff --check` - `python -m pytest --collect-only -q ...attention benchmark files...` -> 117 collected - `python scripts/validate_manifest.py --levels schema,shape,dtype,bench` -> passed ## Nightly Benchmarking Ran with `tileops-runner:nightly-tl019-fullstack-no-tileops-ldfix` on host GPU1. Container-visible GPU UUID matched host GPU1: `GPU-bea8d0b6-e3c4-082c-e524-50f7791c9a1e`. Results: - affected attention run: 109 collected, 59 passed, 50 skipped, 0 failures - FP8 tensor-core GQA benchmark: 8 passed, 0 failures - targeted recheck: 14 passed, 0 failures No clear regression attributable to this change. The suspicious GQA varlen slowdown was reproduced on a clean base worktree at the original commit, so it appears to be environmental/runtime variance rather than this migration. |
|
|
Cao Ying
|
f22b7fe723
|
[Refactor][CI] retarget README badges to spec / bench coverage (#1470)
## Summary
- Badge 1: relabel `implemented` → `spec coverage` (same metric,
108/132).
- Badge 2: swap from internal `kernel_map coverage` to `bench coverage`
= `implemented ∧ source.bench_manifest_driven` (17/108, 16%). Badges
advertise external commitments; `kernel_map` is a schema-internal field.
- Surface the new bench gap in `--format text` / `--format md` and in
the JSON's `conformance_gaps`.
- README badge URL updated to `manifest-benchmark.json`. After merge,
`push: main` regenerates the `stats` orphan branch and drops the old
`manifest-kernel-map.json`.
Design rationale: `TileOpsGov/rfcs/2026-05-15-manifest-stats-ci.md`
§6.2.
## Test plan
- [x] pre-commit passed
- [x] `--format text` / `--format md` / `--badge-output` all show the
new bench gap
- [ ] Post-merge: `push: main` regenerates stats branch; README badges
live within ~2 min
---------
Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
944f64e386
|
[Feat][CI] add manifest stats reporter and publishing workflow (#1462)
## Summary
- Add `scripts/manifest_stats.py`: aggregates `tileops/manifest/` into
status / per-family / spec-coverage / conformance views. Outputs text,
markdown, JSON, plus shields.io endpoint badge payloads
(`--badge-output`) and PR-vs-base deltas (`--diff`). Exit code is always
0 — informational, not a gate.
- Add `.github/workflows/manifest-stats.yml` with three jobs: per-event
Job Summary; sticky PR comment posted only when the diff is non-zero
(via `marocchino/sticky-pull-request-comment@v2`); publish job that
force-pushes JSON + badge payloads to an orphan `stats` branch on `push:
main` and a daily schedule.
- Wire two shields.io endpoint badges into the README header pointing to
`raw.githubusercontent.com/.../stats/*.json` (`implemented X/Y (Z%)` and
`kernel_map coverage`).
- No source changes outside the three files above; `main` history is not
polluted by generated artifacts.
Design rationale lives in the governance RFC
`2026-05-15-manifest-stats-ci.md` (private repo). Key decisions: orphan
`stats` branch (not `gh-pages`, no GitHub Pages required); in-place
double-checkout for diff baseline (no artifact TTL); spec coverage badge
is `count(implemented ∧ has kernel_map) / count(implemented)`; daily
`schedule:` as belt-and-suspenders.
## Test plan
- [x] pre-commit passed
- [x] script runs locally in all three formats and with `--diff` against
a baseline JSON
- [x] `--badge-output` produces valid shields.io endpoint JSON
(`schemaVersion / label / message / color`)
- [x] workflow YAML parses with `yaml.safe_load`
- [ ] First merge to main triggers `publish-stats`, which creates the
orphan `stats` branch and makes README badges live (verifiable
post-merge)
## Additional context
- Surfaced gap (informational, to be tracked separately): 34 implemented
elementwise ops lack `source.kernel_map`. The script reports this; fix
is out of scope for this PR.
- `bench_manifest_driven` coverage is 17% — currently surfaced only in
the markdown/text reports, not in a badge, pending clarification of
rollout status.
---------
Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
bfae9be4e5
|
[Refactor][Manifest] drop generative-op carrier; allow empty signature.inputs (#1444)
Closes #1440 ## Summary - Relax `validate_manifest.py`: `outputs >= 1 AND (inputs >= 1 OR params >= 1)` replaces the old `inputs >= 1` invariant; the `is_generative` parameter and `ref_api == 'none'` heuristic are removed from `check_l1` and L4. - `tileops/manifest/elementwise_generative.yaml`: `AlibiFwdOp` and `SinusoidalFwdOp` now declare `inputs: {}` with `dtype` promoted to a real `params:` entry; the carve-out preamble is gone. - `tileops/ops/elementwise/_base.py`: drop `_register_generative_custom_op`; `alibi.py` / `sinusoidal.py` drop the `self._device_carrier` allocation and call the kernel directly (eager-only — torch.compile graph capture deferred to a follow-up if needed). - Add `test_signature_inputs_may_be_empty_when_params_present` in `tests/test_validate_manifest.py` exercising the relaxed schema. - No `device_carrier` / `_register_generative_custom_op` / `is_generative` / `generative-op carve-out` tokens remain anywhere in `tileops/`, `scripts/`, `tests/`. ## Test plan - [x] pytest tests/test_validate_manifest.py tests/ops/test_special_elementwise.py — 301 passed, 3 warnings, 0 failed - [x] python scripts/validate_manifest.py — exits 0, "All manifest checks passed." - [x] grep `device_carrier|_register_generative_custom_op|is_generative|Generative-op carve-out|generative-op carve-out` on `tileops/ scripts/ tests/` — no matches - [x] grep `signature.inputs` on validator + tests — no `>= 1` invariant remains; relaxed form is in place - [x] Manifest YAML check: both ops have `inputs: {}` and `dtype` as a param; preamble carve-out narrative removed - [x] New test `test_signature_inputs_may_be_empty_when_params_present` present and passing ## Test node delta ``` File Base HEAD Delta -------------------------------------------------------- tests/test_validate_manifest.py 218 220 +2 -------------------------------------------------------- TOTAL 218 220 +2 Growth: +0.9% ``` **Justification:** one new schema test (`test_signature_inputs_may_be_empty_when_params_present`) is added to lock in the relaxed `outputs >= 1 AND (inputs >= 1 OR params >= 1)` invariant required by AC-6, and one ancillary node tracks the symmetric `params`-less rejection path. AC-6 mandates this addition; without it the relaxation is unverified. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
lcy-seso
|
2cfa372629 |
[Maintain][Manifest] trim process-y carve-out narrative from comments
Comments describing the generative-op workaround framing (mentions of
the schema invariant being satisfied by a carrier, references to a
"carve-out" pattern) age out the moment the workaround is removed.
Tightened to just the mechanical condition and the invariant the code
itself preserves.
Co-Authored-By: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
lcy-seso
|
116a50ccfc |
[Fix][Manifest] distinguish forward() introspection failure from zero positional args
`_forward_positional_params(cls)` returns ``None`` when
``inspect.signature(forward)`` raises, and ``[]`` when forward() has zero
positional args. The L1 generative-op carve-out used ``or []`` which
coalesced both cases to falsy, so an introspection failure silently
satisfied ``not positional_forward_params`` and skipped L1 alignment for
any ``ref_api: "none"`` op.
Use explicit ``positional is not None and len(positional) == 0`` at the
L1 site; mirror the explicit ``len == 0`` form at the C4 site so the two
carve-out predicates stay visually identical and an introspection-failed
class can never bypass either check.
Co-Authored-By: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
lcy-seso
|
c2e01ac8d4 |
[Fix][Manifest] adopt documented shape/workload format in new elementwise families
- Convert signature shape values from YAML lists to string form ("[M, N]")
in elementwise_generative.yaml and elementwise_fused_gated.yaml to match
docs/design/manifest.md R8 + the convolution.yaml precedent so the
validator can bind shape symbols.
- Add per-input workload shape keys (device_carrier_shape, x_shape) to
every workload row to satisfy the workload schema documented in
docs/design/manifest.md and used across normalization.yaml etc.
- Restore exception class/message in check_c4_forward_signature_parity's
warning when inspect.signature(forward) raises.
Co-Authored-By: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
lcy-seso
|
855adcf333 |
[Maintain][Manifest] align L1 generative detection with C4 positional filter
The L1 carve-out at check_l1_pyfile_signature used _get_forward_params,
which counts KEYWORD_ONLY params. check_c4_forward_signature_parity
filters to POSITIONAL_ONLY / POSITIONAL_OR_KEYWORD before applying its
ref_api == none + zero-positional-args carve-out. A forward(self, *,
dtype=None) class therefore diverged: L1 emitted a [signature] error
while C4 returned []. Extract _forward_positional_params and use it from
both sites so the two carve-outs stay in lockstep.
Co-Authored-By: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
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>
|
|
|
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> |
|
|
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> |
|
|
Cao Ying
|
2de052e19a
|
[Refactor][Hygiene] strip development-process metadata from shipped source (#1290)
## Summary
Apply `.claude/rules/code-style.md` "no development-process metadata in
shipped source" across `tileops/` `tests/` `benchmarks/` `scripts/`.
Rewrite `pytest.skip`/`xfail` reasons, docstrings, comments to describe
invariants instead of referencing issue/PR numbers and AC labels.
Prose-only — no behavior change.
## Scope (24 files)
- 16 `pytest.skip`/`xfail` reason rewrites (tilelang 0.1.9 batch, rope,
attention, recurrent linear-attention) + shared
`TILELANG_019_SKIP_REASON` in two conftests
- AC-label strips in 9 test files / 1 bench /
`scripts/validate_manifest.py`
- `TODO(#NNNN)` strips in `tileops/kernels/convolution.py` (×2),
`reduction/logsumexp.py`, `reduction/softmax.py`
- Historical PR refs in `tests/test_validate_manifest.py` rewritten to
describe the contract
Excluded: `tileops/manifest/`, `docs/`, `.claude/` (out of rule scope).
`FIXME(staged-rollout)` blocks already compliant. Single-digit "Finding
#N" labels not matched by the rule's `#[0-9]{3,}` regex.
## Test plan
- [x] discovery scan returns 0 matches
- [x] `pytest tests/ --co -q` — 2752 collected
- [x] touched test files: 186 + 53 passed
- [x] pre-commit clean
Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
1e8f091ec5
|
[Refactor][Manifest] drop eager ops import from tileops package (#1219)
Closes #1216 ## Summary - Drop `from . import ops` (and the matching `__all__` entry) from `tileops/__init__.py`. Importing `tileops` no longer pulls in `tileops.ops` or any `tilelang.*` module as a side effect. - Replace the `importlib.spec_from_file_location` workaround in `scripts/validate_manifest.py` with a direct `from tileops.manifest.shape_rules import dim_range_validity, dim_uniqueness, reduced_axes`. Drop the `_SHAPE_RULES_PATH` / `_shape_rules_spec` / `_shape_rules_module` plumbing and the underscore-prefixed rebindings. ## Test plan - [x] AC-1: Modified files pass existing tests. - [x] AC-2: `import tileops` does not load `tileops.ops` (verified via `sys.modules` snapshot). - [x] AC-3: `from tileops.manifest import shape_rules` does not load `tileops.ops` or any `tilelang.*` module (verified via `sys.modules` snapshot). - [x] AC-4: `scripts/validate_manifest.py` uses standard `from tileops.manifest.shape_rules import ...` instead of `importlib.spec_from_file_location`; the `_SHAPE_RULES_PATH` literal is removed. - [x] AC-5: `python scripts/validate_manifest.py` exit 0 on the checked-in manifest, output bit-identical to pre-change baseline. |
|
|
Cao Ying
|
b81001dae7
|
[Feat][Manifest] AST-walk L0 typo defense for shape_rule callables (#1217) (#1218)
Closes #1217 ## Summary - Add `_check_shape_rule_callables` helper in `scripts/validate_manifest.py` that AST-walks each `shape_rules` entry at L0 and rejects any direct `Name(...)` call whose name is not in `_SHAPE_RULE_BUILTINS`. - SyntaxError in a shape_rule now surfaces as a single `[schema]` error at L0 instead of an L2 eval-time warning, so typos are caught without L2 mock-input setup. - Add four `TestSchema` cases covering unknown callable rejection, known-callable acceptance, attribute-call passthrough, and SyntaxError rejection. ## Test plan - [x] AC-1 — Modified files pass existing tests. (`python -m pytest -q tests/test_validate_manifest.py` → 185 passed) - [x] AC-2 — Typoed shape_rule callable names surface as `[schema]` errors at L0, not L2 NameError warnings. - [x] AC-3 — Existing `tileops/manifest/*.yaml` shape_rules pass the new L0 check with no false positives. (`python scripts/validate_manifest.py` → exit 0) - [x] AC-4 — Test coverage for unknown callable, known callables, attribute call, and syntax error cases. - [x] AC-5 — SyntaxError in a shape_rule moves from L2 eval warning to L0 schema-level rejection without L2 mock input setup. ## Test node delta ``` File Base HEAD Delta -------------------------------------------------------- tests/test_validate_manifest.py 181 185 +4 -------------------------------------------------------- TOTAL 181 185 +4 Growth: +2.2% ``` **Justification:** Four new tests are required by AC-4, one per shape_rule callable case (unknown name → reject, known names → pass, attribute call → pass, SyntaxError → reject). Each case exercises a distinct branch of the new AST walk in `_check_shape_rule_callables`; collapsing them would lose coverage of the L0/L2 boundary that this PR establishes. --------- Co-authored-by: Ibuki — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
f851936a0a
|
[Refactor][Manifest] introduce shape_rules helpers as flat shape_rule builtins (#1210)
Closes #1209. Follow-up #1212 migrates `tileops/manifest/reduction.yaml` to reference these helpers. ## Summary Tooling-only PR — no `tileops/manifest/*.yaml` data is modified. - `tileops/manifest/shape_rules.py`: three pure-Python helpers — `dim_range_validity(x, dim) -> bool`, `dim_uniqueness(x, dim) -> bool`, `reduced_axes(x, dim) -> frozenset` — preserving the inline-form malformed-dim semantics verbatim (TypeError on malformed input is classified by the validator as an eval-error warning, bit-identical pre/post). - `scripts/validate_manifest.py`: helpers wired into `_SHAPE_RULE_BUILTINS` alongside `len`, `range`, `broadcast_shapes`, etc. Callable by bare name from any shape_rule body. The dict is built from an explicit `(name, callable)` list so a duplicate-name registration fails the validator at import time instead of silently shadowing. - `.claude/domain-rules/`: drift-surface cleanup — trim reduction-dim rule, dedupe Op-naming rule across three files, strip restatements that the validator/type-system already enforces. - `tests/test_validate_manifest.py`: drop per-case scaffolding tests; fold pre-existing `TestShapeRuleBroadcastBuiltins` per-case property tests into matrix tests. ## Test plan - [x] `pytest tests/test_validate_manifest.py -q` → 181 passed - [x] `python scripts/validate_manifest.py` → exit 0; pre/post diff empty (pinned by `test_sum_rules_helper_inline_classification_parity`) - [x] `tileops.manifest.shape_rules` exposes the three helpers via `__all__` with docstrings - [x] Helpers callable by bare name from shape_rules eval scope (pinned by `test_shape_rules_helpers_callable_by_bare_name`, driven by `__all__`) - [x] `_SHAPE_RULE_BUILTIN_PAIRS` raises on duplicate names (pinned by `test_shape_rule_builtin_pairs_have_unique_names`) - [x] `.claude/domain-rules/manifest-spec.md` reduction-dim rule points at `tileops.manifest.shape_rules` ## Test node delta `tests/test_validate_manifest.py`: 190 → 181 (−9). +6 contract-level tests for the new helper layer; −15 from folding `TestShapeRuleBroadcastBuiltins` per-case tests into matrix tests with no contract loss. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
8f6f1462b1
|
[Feat][Manifest] align where/clamp/masked_fill multi-input ops to pytorch api (#1109)
Closes #1088 ## Summary - Align the three multi-input elementwise op manifest entries (`where`, `clamp`, `masked_fill`) to PyTorch's public API, per the manifest trust model. - Split families per the "No `Optional[Tensor]`" rule: `Clamp` family becomes 4 entries (`ClampFwdOp` primary + `ClampScalarFwdOp` / `ClampMinFwdOp` / `ClampMaxFwdOp` variants); `MaskedFill` family becomes 2 entries (`MaskedFillFwdOp` primary + `MaskedFillScalarFwdOp` variant). - Update `WhereFwdOp.shape_rules` to use `broadcast_shapes(condition.shape, input.shape, other.shape)` matching `torch.where` broadcasting; remove inline BLOCKED notes from all three families. - Add `broadcast_shapes` and `is_broadcastable_to` to `_SHAPE_RULE_BUILTINS` in `scripts/validate_manifest.py` (pure-Python, no torch dependency) so manifest `shape_rules` can express NumPy/PyTorch broadcast semantics. - Add roofline blocks for the new `Clamp` and `MaskedFill` variants, with `vars.N_total = product(broadcast_shapes(...))` for broadcasting variants and `product(input.shape)` for the scalar-bound `ClampScalarFwdOp`. - All new / restructured entries land as `status: spec-only`; kernel and op-class conformance is tracked in follow-up #1107. ## Test plan - [x] AC-1: `WhereFwdOp.shape_rules` uses `broadcast_shapes(...)` and matches `torch.where` broadcasting semantics; inline BLOCKED note removed. - [x] AC-2: Clamp family expressed as 4 entries (`ClampFwdOp` primary + `ClampScalarFwdOp` / `ClampMinFwdOp` / `ClampMaxFwdOp` variants) with `Number | None` for scalar params; inline BLOCKED note removed. - [x] AC-3: MaskedFill family expressed as 2 entries (`MaskedFillFwdOp` primary + `MaskedFillScalarFwdOp` variant) with bidirectional input/mask broadcasting (out-of-place `Tensor.masked_fill` semantics) and `Number` value type; inline BLOCKED note removed. - [x] AC-4: `_SHAPE_RULE_BUILTINS` includes `broadcast_shapes` and `is_broadcastable_to`, with unit tests in `tests/test_validate_manifest.py`. - [x] AC-5: `scripts/validate_manifest.py --levels schema` passes for all affected entries. - [x] AC-6: `where_fwd_roofline` and new variant roofline blocks are consistent with the post-broadcast `N_total` convention. - [x] AC-7: Follow-up issue filed (#1107) tracking kernel + op-class conformance to all new / restructured entries. Verification commands (all green locally): - `pre-commit run --all-files` → all hooks pass - `python scripts/validate_manifest.py --levels schema` → all manifest checks passed - `pytest tests/test_validate_manifest.py` → 190 passed, 3 warnings ## Test node delta ``` File Base HEAD Delta -------------------------------------------------------- tests/test_validate_manifest.py 171 190 +19 -------------------------------------------------------- TOTAL 171 190 +19 Growth: +11.1% ``` **Justification:** The +19 nodes cover the new `_SHAPE_RULE_BUILTINS` helpers (`broadcast_shapes`, `is_broadcastable_to`): direct unit coverage of helper behavior (broadcasting equivalence, broadcast failures, scalar/empty-shape edge cases, unidirectional broadcastability) plus end-to-end coverage via `_eval_shape_rule` to confirm the helpers are actually reachable from manifest `shape_rules` strings. These helpers gate the new PyTorch-aligned `shape_rules` for `where`, `clamp`, and `masked_fill`; without this coverage, the broadcasting semantics of those entries are validator-untested. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
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> |
|
|
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>
|
|
|
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> |
|
|
Cao Ying
|
b913e9cd72
|
[Feat][TOOLING] validator parity for _infer_output_shapes and _validate_dtypes (#1005)
## Summary Add validator checks that ensure a concrete op's `_infer_output_shapes` and `_validate_dtypes` agree with its manifest declaration. Introduces L2 infer-shape parity and L3 dtype parity checks in `scripts/validate_manifest.py`, plus documents the two new rows in the Consistency Enforcement table and the `parity_opt_out` entry field. Closes #994 ## What this PR does - **L2 infer-shape parity**: for each `status: implemented` op, invokes `_infer_output_shapes` with concrete mock inputs and checks the returned shapes against `signature.outputs[*].shape` declarations, `signature.static_dims`, param defaults, and `shape_rules` (incl. R11a comprehensions). Symbolic dims (`B`, `S`, `H`, `D`, …) are bound from input shapes and also rebound from inferred outputs so output-only symbols like conv `L_out` can be checked against rule-derived values. - **L3 dtype parity**: validates `_validate_dtypes` behaviour against `dtype_combos` (when present, exhaustive) or declared `dtype` unions (full Cartesian product), plus out-of-union negative probes and `same_as(ref)` identity probes on both branches. - **Manifest-data validation (runs unconditionally in `check_l3`)**: combo completeness (every combo row must cover every declared input), invalid dtype tokens, unresolved / cyclic `same_as(ref)`, and union values in combo entries are hard L3 errors regardless of whether the class overrides `_validate_dtypes`. - **`parity_opt_out`**: optional manifest entry field (`true` or subset of `[shape_parity, dtype_parity]`) to suppress parity checks for ops whose method genuinely needs GPU execution. Documented in `docs/manifest.md` (Entry Structure) and `docs/ops-design-reference.md` (Consistency Enforcement). - **Security / bounds**: `eval()` for shape rules is restricted to a builtins allowlist, AST-filtered for dunder attrs, and protected from `ctx` overwriting `__builtins__`. Cartesian product iteration is bounded by `_MAX_DTYPE_COMBOS = 4096` with deterministic skip + warning on over-bound ops. - **Policy (tightened over review rounds)**: unexpected body exceptions from `_infer_output_shapes` / `_validate_dtypes` are **hard parity errors by default**; reserve the soft-warning path for entries that explicitly opt out. Signature-bind failures (via `inspect.signature().bind`) remain a separate category. ## Test plan - [x] **AC-1**: Modified files pass unit tests. - Evidence: `python -m pytest tests/test_validate_manifest.py -q` → 170 passed. - [x] **AC-2**: Validator L2 reports error when `_infer_output_shapes` disagrees with `shape_rules` / declared shapes. - Evidence: `TestInferShapeParity` covers incorrect infer, symbolic dim mismatch, R11/R11a helper rules, input-only precondition separation, conv-like output-only symbols, declared-output-shape exact match, `self.attr` / `self.static_dim` mock populations, body-exception-as-hard-error (+ opt-out downgrade). - [x] **AC-3**: Validator L3 reports error when `_validate_dtypes` accepts non-listed or rejects declared combos. - Evidence: `TestValidateDtypesParity` covers union accept/reject, `dtype_combos` accept/reject listed, first-rejected-later-accepted enumeration, out-of-union probes (both branches), `same_as` identity probes, signature-mismatch vs body-exception separation (body errors are hard unless opted out), Cartesian bound skip with warning, exhausts-the-union warning. - [x] **AC-4**: Consistency Enforcement table adds two new rows. - Evidence: `docs/ops-design-reference.md` adds rows for `_infer_output_shapes` vs `shape_rules` (L2) and `_validate_dtypes` vs `dtype_combos` / unions (L3), plus a "Parity check coverage" paragraph documenting missing-override warning + `parity_opt_out`. `docs/manifest.md` Entry Structure documents the `parity_opt_out` field. ## Integration check ``` python scripts/validate_manifest.py → exit 0, 0 ERROR lines, 48 WARNING lines (expected: missing-override warnings for implemented ops awaiting codegen migration + pre-existing bench warnings). "All manifest checks passed." ``` No existing Op required `parity_opt_out` under the tightened body-exception policy. ## Test node-delta (testing-budget rule) ``` tests/test_validate_manifest.py 101 170 +69 (+68.3%) ``` Justification: the delta covers AC-1..AC-4 parity behaviour plus regression coverage from seven review rounds: - signature-bind vs body-exception separation (body errors are hard L2/L3 by default, opted out via `parity_opt_out`) - `same_as` fixpoint resolution (order-independent); pure-cycle and dangling-ref diagnosis produce hard L3 errors with all cycle participants named - out-of-union negative probes on both branches (`dtype_combos` and Cartesian), sourced from `_TORCH_DTYPES − declared` so sentinel engulfment can't create a vacuous pass; full-coverage case emits a named warning - `same_as(ref)` identity probes on both branches - mock `self` via `cls.__new__(cls)` with `signature.params` defaults, `static_dims` resolved values, and primary `dtype` axis populated - symbolic dim binding from `signature.*.shape`; output-only symbol rebinding from inferred outputs (conv-like `L_out` scenarios) - `static_dims` and scalar-int param defaults both pin declared-output-shape positions to exact values (not just rank/consistency) - combo completeness (every row covers every declared input), invalid dtype value rejection, union-in-combo-value rejection as hard L3 manifest errors — run unconditionally in `check_l3`, independent of `_validate_dtypes` override status - Cartesian product bounded by `_MAX_DTYPE_COMBOS = 4096` with deterministic skip + warning - `eval()` sandbox: restricted `__builtins__`, AST dunder-attr filter, `ctx` cannot overwrite `__builtins__` Tests are validator regressions, not broad combinatorics. --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
a3b2af0a79
|
[Refactor][Roofline] remove manifest-level roofline evaluator (#1024)
## Summary
This PR now targets the design direction that target design should not
have a manifest-level evaluator. Instead of keeping a legacy evaluator
boundary, it removes the legacy evaluator surface and migrates current
callers.
- Delete `tileops/manifest_legacy_roofline.py`.
- Remove manifest exports for `_safe_eval`, `eval_roofline`,
`resolve_roofline_vars`, and `has_roofline_vars`.
- Move benchmark roofline consumption to op-local `eval_roofline()` and
`ManifestBenchmark(op_name, op, workload)`.
- Add local `eval_roofline()` implementations for the currently migrated
norm, MoE, reduction, and softmax-family benchmark surfaces.
- Update validator, tests, and roofline docs for the op-local evaluator
contract.
## Test Plan
- `ruff check` on changed Python files
- `python -m compileall tileops benchmarks scripts tests`
- `pytest tests/test_ops_manifest.py tests/test_workloads_to_params.py
tests/test_validate_manifest.py
benchmarks/tests/test_roofline_workload_protocol.py`
- `python scripts/validate_manifest.py`
Note: `validate_manifest.py` passes and still reports the existing
attention benchmark L4 warnings; those are pre-existing warnings outside
this cleanup scope.
## Follow-up
No follow-up issues or suggestions.
---------
Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com>
|
|
|
ChongLi
|
4da2dbba1d
|
[CI] Run nightly inside containerized runner (#1017)
Closes #1016 ## Summary - Move nightly workflow execution to assume it already runs inside the containerized self-hosted runner. - Remove workflow-level Docker startup, workspace ownership fixes, and the local tileops-runner:latest wrapper. - Add a nightly runner verification script for cache env vars, cache writability, and GPU visibility. ## Test plan - [x] pre-commit passed during git commit - [x] bash -n scripts/ci/verify_nightly_runner.sh - [x] python -c "import yaml; yaml.safe_load(open('.github/workflows/nightly.yml')); print('yaml ok')" - [x] git diff --check -- .github/workflows/nightly.yml scripts/ci/verify_nightly_runner.sh - [x] rg -n "docker run|tileops-runner:latest|Fix workspace ownership|/data/ci-cache" .github/workflows/nightly.yml returns no matches - [ ] actionlint .github/workflows/nightly.yml (not run locally; actionlint is not installed) ## Additional context This PR intentionally keeps the existing runner labels unchanged and does not modify runner-maintenance.yml. Cache cleanup policy is left for a separate follow-up, as scoped in #1016. |
|
|
Cao Ying
|
ee28b2926b
|
[Fix][CI] widen fork-PR gpu-smoke cache hits without widening write surface (#1009)
## Summary
Maximize cache hit rate for fork-PR `gpu-smoke` runs without widening
the write-path attack surface.
Three narrowly-scoped fixes:
1. **Narrow venv HASH** (`scripts/ci_venv_hash.py`): parse
`pyproject.toml` with `tomllib` and hash only `[project]` +
`[build-system]` + `[project.optional-dependencies]`. Edits to
`[tool.pytest]` / `[tool.ruff]` no longer bust the venv cache. Falls
back to full-file hash on parse failure, preserving the 16-char hex
format.
2. **Member-fork → trusted runtime path**
(`.github/workflows/gpu-smoke.yml` policy): when head repo differs from
base but `author_association ∈ {OWNER, MEMBER, COLLABORATOR}`, set
`is_fork=false` so the PR traverses the trusted runtime path
(`RUNTIME_ROOT=/home/ci-runner`). External contributors still go the
isolated path.
3. **Skip in-line atomic age-trim**
(`.github/actions/reclaim-runner-disk/action.yml`): add
`skip-atomic-age-trim` input (default `false`); `gpu-smoke.yml` passes
`true` so the destructive autotuner age-trim only runs via
`runner-maintenance.yml` (daily). Sentinel-repair remains always on.
Closes #1007
## Test plan
- **AC-1**: Changing `[tool.pytest]` / `[tool.ruff]` sections of
`pyproject.toml` does NOT change the computed venv HASH — pass. `python
-m pytest -q tests/test_ci_venv_hash.py tests/test_gpu_smoke_policy.py
tests/test_reclaim_action.py`: 26 passed in 1.20s. Direct hash probe:
base=`c2c5efe07bd7e601`, tool-section mutation=`c2c5efe07bd7e601`,
stable=True.
- **AC-2**: Changing `[project.dependencies]` or
`[build-system].requires` DOES change the computed HASH — pass.
Dependency mutation `d435cbfc2b7aca79` changed=True; build-system
requires mutation `93760788d3b9ecd3` changed=True.
- **AC-3**: Member-fork PRs produce `is_fork=false` and traverse the
trusted runtime path — pass. Policy probe with `HEAD_REPO!=BASE_REPO`
and `AUTHOR_ASSOC=MEMBER` produced `is_fork=false`; Resolve-runtime
probe with `IS_FORK=false` exported `RUNTIME_ROOT=/home/ci-runner`.
Negative probe with `AUTHOR_ASSOC=CONTRIBUTOR` produced `is_fork=true`
and isolated `RUNNER_TEMP` runtime.
- **AC-4**: `gpu-smoke` in-line `reclaim-runner-disk` does NOT invoke
atomic age-trim — pass. Composite-action run with
`SKIP_ATOMIC_AGE_TRIM=true` emitted `Skipping atomic age-trim (opted
out)`, still ran sentinel-repair on half-dead atomic subdir, and
preserved the complete atomic subdir.
- **AC-5**: `runner-maintenance.yml` daily job still performs full
atomic trim — pass. YAML probe confirmed reclaim step has
`force-reclaim=true` and `skip-atomic-age-trim` absent/false. `git diff
main...HEAD` does not modify `runner-maintenance.yml`.
- **AC-6**: Unit tests for `reclaim_cache.sh` still pass — pass.
`tests/test_reclaim_action.py` Base 13 nodes, HEAD 13 nodes, Delta 0;
all 13 sentinel-repair tests pass.
---------
Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
a5e3c48f49
|
[Design][Manifest] Redesign init_dims as static_dims (docs + schema + validator) (#982)
## Summary
Complete the `static_dims` redesign (renamed from `init_dims`) for
arbitrary-rank ops in a single internally consistent PR. Scope now
covers:
1. **Docs** — authoritative spec rewrite: `docs/manifest.md` R20,
`docs/ops-design.md` `_cache_key` section,
`docs/ops-design-reference.md` Base Class Protocol + Consistency
Enforcement table.
2. **Manifest yaml** — 16 ops migrated from `init_dims: {N: {from:
"..."}}` to flat `static_dims: {N: "..."}`. LogSumExpFwdOp's multi-axis
form removed (violates new R20 single-axis rule).
3. **Validator** — `scripts/validate_manifest.py` renamed
`manifest_init_dims` → `manifest_static_dims`, reads
`sig.get("static_dims")`, updated error messages. L1 enforces
static_dims keys ⊆ `__init__` parameters.
4. **Status flips** — 13 `status: implemented` ops flipped to `status:
spec-only` with per-op comments explaining the impl/spec mismatch
(LayerNorm family × 5, BatchNorm/InstanceNorm × 3,
Softmax/LogSoftmax/LogSumExp × 3, Argmax/Argmin × 2). This is the
trust-model-prescribed response to impl divergence.
5. **PyTorch API alignment** — `dim` types include `tuple[int, ...]`;
`isinstance(dim, list)` → `(list, tuple)` throughout manifest;
LogSumExpFwdOp `shape_rules` dim-normalized.
6. **Design-first principle** — one-line statement added to `CLAUDE.md`
(always loaded). Docs focus on technical spec.
Closes #984 (absorbed). Closes #976.
#985 remains open for the Op base `_cache_key` runtime +
empty-static_dims warning.
## Test plan
- [x] `python scripts/validate_manifest.py` passes (24 pre-existing
bench-file warnings, 0 errors)
- [x] 91 validator unit tests pass (4 new `static_dims` tests)
- [x] No `init_dims` remnants anywhere: `grep -rn init_dims docs/
tileops/ scripts/ tests/` returns empty
- [x] 16 `static_dims:` entries in manifest yaml (down from 31
`init_dims` originally — 14 dim=None-supporting reduction ops correctly
had theirs removed per R20 applicability, LogSumExp removed per
single-axis rule)
- [x] `docs/manifest.md` R20 uses `static_dims` flat syntax with
single-axis reference rule, two time points + codegen pseudocode,
LinearFwdOp multi-input example, empty-legal + mandatory `_cache_key`
override (SumFwdOp example)
- [x] Design-first principle stated once in `CLAUDE.md` Project Overview
- [x] Consistency Enforcement table lists only live enforcement
(validator + runtime); deferred mechanisms tracked via issues
## Follow-up
Issues:
- #992 — dim-aware ManifestBenchmark + end-to-end roofline.vars
consumption
- #993 — retrofit 13 spec-only ops to the new static_dims / _cache_key
contract
- #994 — validator codegen parity checks for _infer_output_shapes /
_validate_dtypes
---------
Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
bdbd7a8324
|
[Refactor] Rename base modules to <module>_base.py convention (#950)
## Summary Unify base module file naming to `<module>_base.py` convention across four modules: - `workloads/base.py` -> `workloads/workload_base.py` - `tileops/ops/op.py` -> `tileops/ops/op_base.py` - `tileops/kernels/kernel.py` -> `tileops/kernels/kernel_base.py` - `benchmarks/benchmark.py` -> `benchmarks/benchmark_base.py` All imports, string references, and doc paths updated. No logic changes — pure rename + import propagation. Closes #947 ## Test plan - [x] **AC-1**: All four files renamed to <module>_base.py -- pass - [x] **AC-2**: All imports updated, no broken references -- pass - [x] **AC-3**: docs/testing.md path references updated -- pass - [x] **AC-4**: Pre-commit and existing tests pass -- pass **Test results**: 2368/2368 tests passed, 0 failed --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
ef5d248f2f
|
[Refactor][Benchmark] Extract ManifestBenchmark base class and shared helpers (#915)
## Summary Extract a generic `ManifestBenchmark` base class and shared helpers (`roofline_vars`, `workloads_to_params`) into `benchmarks/benchmark.py`, so manifest-driven bench files only specify the op name and baseline function per test. This eliminates ~300 lines of duplication across 5 benchmark files. Closes #904 ## Changes - **benchmarks/benchmark.py** — Add `ManifestBenchmark(op_name)` subclass of `BenchmarkBase` with `roofline_vars(workload)` and `workloads_to_params(op_name)` public helpers - **benchmarks/ops/bench_{softmax,reduce,argreduce,logical_reduce,vector_norm}.py** — Refactor to use `ManifestBenchmark`, replacing 19 per-op benchmark classes - **scripts/validate_manifest.py** — Teach AST validator to recognize `ManifestBenchmark` subclasses as satisfying `eval_roofline` requirement - **tests/test_validate_manifest.py** — Add tests for indirect bench-validator AST paths and a regression test documenting the subclass override limitation ## Test plan - [x] AC-1: All 5 bench files use `ManifestBenchmark` instead of per-op classes - [x] AC-2: `_roofline_vars` and `_workloads_to_params` exist only in `benchmarks/benchmark.py`, not duplicated - [x] AC-3: `pytest --collect-only` passes for all 5 bench files (63 tests collected) - [x] AC-4: `validate_manifest.py --check-op` produces zero `[bench]` warnings for all 19 affected ops - [x] 89 validator tests pass, lint clean ## Follow-up - #917 — Add import alias and attribute access support to bench AST validator - #918 — Type-narrow roofline_vars parameter from WorkloadBase to typed Protocol --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
d81ec1f5ef
|
[Test][Softmax] Add non-aligned edge-case shapes for softmax-family ops (#914)
## Summary Add targeted edge-case test shapes that exercise the M×N non-aligned combination for softmax-family kernels, covering single-tile and multi-tile (large-N) paths. Closes #906 ## Changes **Tests** (`tests/ops/test_softmax.py`): - 3 new shapes per fixture (SoftmaxFixture, LogSoftmaxFixture, LogSumExpFixture): - `(33, 300)` fp32 — both M and N non-aligned, single-tile path - `(33, 33000)` fp16 — both M and N non-aligned, multi-tile with masked loads - `(33, 32768)` fp16 — non-aligned M with large-N tiled path **Scripts** (`scripts/test_node_delta.py`): - Fix base-ref collection: write temp file in original directory instead of `/tmp/` so pytest can resolve project imports **Foundry** (`.foundry/mold/pre-create-checks.md`): - Add test node delta as pre-create check #3 (soft gate) with correct `--base upstream/main` command ## Test plan - [x] AC-1: All 126 tests pass (`pytest tests/ops/test_softmax.py`) - [x] AC-2: New shapes exercise non-aligned M×N paths (single-tile and multi-tile) ## Test node delta ``` File Base HEAD Delta -------------------------------------------------- tests/ops/test_softmax.py 117 126 +9 -------------------------------------------------- TOTAL 117 126 +9 Growth: +7.7% ``` Justification: 9 new `pytest.mark.full` nodes (3 shapes × 3 fixtures). Each shape targets a distinct non-aligned code path not covered by existing aligned shapes. No cross-dtype expansion (existing fixtures already cover fp32/fp16/bf16 on aligned shapes; the non-aligned path shares the same tiling/masked-load logic across fp16 and bf16). ## Follow-up No follow-up issues or suggestions. --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
0ce134bb24
|
[Refactor][Manifest] Enforce status as required and validate kernel_map schema (#888)
## Summary Update the manifest validator to enforce `status` as required and validate `kernel_map` schema. - `status` added to `_REQUIRED_TOP` — entries missing it produce a schema error - `status` must be a string (`"implemented"` or `"spec-only"`); null and non-string values rejected - `source.kernel_map` validated when present: must be `dict[str, str]` - Missing `kernel_map` for `status: implemented` entries produces a **warning** (not error) — full enforcement deferred to #887 which populates kernel_map for all existing ops - `_is_spec_only()` defaults to `True` on missing/non-string status (safe for `--levels` bypass) - Integration test validates against the real checked-in manifest ## Changed files - `scripts/validate_manifest.py` — status required, kernel_map validation, _is_spec_only safe default - `tests/test_validate_manifest.py` — 9 new tests, updated helper, integration test ## Test plan - [x] `pytest tests/test_validate_manifest.py` — 84 passed, 1 xfailed - [x] `python scripts/validate_manifest.py` exits 0 on real manifest (25 kernel_map warnings) - [x] Pre-commit hooks pass Closes #886 ## Follow-up - #887 — Populate kernel_map for all implemented ops (converts warnings to clean pass) No additional follow-up issues or suggestions. --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
c85712ec45
|
[Refactor][Manifest] Rename keys to PascalCase and add ref_api field (#869)
## Summary
Rename all manifest keys from snake_case to PascalCase class names in
`{Name}{Direction}Op` format and add the required `ref_api` field to
every op entry. Validator enforces exact `cls.__name__ == manifest_key`
identity — no heuristic resolution. No op/kernel/benchmark
implementation code is modified.
**Note**: After this PR merges, benchmark callers that use snake_case op
names (e.g. `load_workloads("rmsnorm_fwd")`) will break. PR-B (#866)
fixes all callers as part of the class rename.
Closes #865
## Test plan
- [x] All manifest keys are PascalCase class names in
`{Name}{Direction}Op` format
- [x] `ref_api` field required for all ops (`"none"` when no external
API counterpart)
- [x] Validator L0 schema enforces `ref_api` as required string
- [x] Validator requires exact `cls.__name__ == manifest_key` — no
single-candidate bypass, no heuristic fallback
- [x] No changes to any file under `tileops/ops/`, `tileops/kernels/`,
or `benchmarks/`
- [x] Integration test `xfail` until PR-B renames Op classes
## Changed files
- `tileops/ops_manifest.yaml` — renamed all 52 keys, added `ref_api` to
every entry
- `tileops/manifest.py` — canonical PascalCase keys only (no legacy
alias resolution)
- `scripts/validate_manifest.py` — `ref_api` required in L0 schema,
exact class-name identity enforcement
- `tests/test_ops_manifest.py` — updated test fixtures, canonical key
tests
- `tests/test_validate_manifest.py` — updated validator tests,
exact-match enforcement tests
---------
Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
dbf58a10b1
|
[Doc][Manifest] Clarify same_as(ref) as identity constraint with validator enforcement (#864)
## Summary Clarify \`same_as(ref)\` as a dtype-only identity constraint in manifest spec, and add validator enforcement. **Docs:** - \`docs/manifest.md\`: R3/R3a define \`same_as(ref)\` as exact dtype identity (not "dependent types"). R5/R8 clarify \`same_as\` is dtype-only — shape must be explicitly declared. R13 status default changed to \`spec-only\`. Overall doc reduced by ~27% (470→342 lines). - \`.claude/domain-rules/manifest-spec.md\`: matching updates — \`same_as\` dtype identity, shape explicit, status default \`spec-only\`. **Validator:** - \`_check_dtype_combos_same_as_identity()\`: enforces that \`dtype_combos\` entries assign the same dtype to \`same_as\`-bound tensors. Rejects partial combos where a bound tensor appears without its reference. - \`_is_spec_only()\`: defaults to \`spec-only\` when \`status\` is absent, matching R13. **Tests:** 5 new test cases — dtype_combos pass, mismatch, multi-binding mismatch, partial-combo rejection, missing-status default. Closes #863 ## Changes - \`docs/manifest.md\` — rewritten: concise, no "dependent types", \`same_as\` dtype-only, status default spec-only - \`.claude/domain-rules/manifest-spec.md\` — matching rule updates - \`scripts/validate_manifest.py\` — \`_build_same_as_map()\`, \`_check_dtype_combos_same_as_identity()\`, \`_is_spec_only()\` default fix - \`tests/test_validate_manifest.py\` — 5 new tests (81 total, all pass) ## Test plan - [x] AC-1: \`docs/manifest.md\` R3/R3a explicitly defines \`same_as(ref)\` as dtype identity constraint - [x] AC-2: Validator flags \`dtype_combos\` with mismatched or partial \`same_as\`-bound entries - [x] AC-3: Full manifest validation passes (\`python scripts/validate_manifest.py\` → 0 errors) ## Follow-up No follow-up issues or suggestions. --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
a93643998e
|
[Refactor][Validator] Deduplicate PascalCase resolution loops in _resolve_op_class (#836)
## Summary Consolidates the duplicated resolution heuristics in `_resolve_op_class` into a single strategy loop and adds ambiguity detection to all matching stages (not just the terminal fallback). - Three separate PascalCase loops (stripped_pascal, full_pascal, suffix_fwd_bwd) replaced with single `for`-loop over an ordered strategies list - Suffix matching (`_fwd`/`_bwd`) now emits an ambiguity warning when multiple candidates match, consistent with the pattern established in PR #830 - One new test (`test_suffix_match_ambiguity_emits_warning`) covers the multi-candidate suffix case Closes #811 ## Acceptance Criteria | ID | Criterion | Status | |----|-----------|--------| | AC-1 | Single resolution loop replaces the duplicated PascalCase loops | PASS | | AC-2 | _fwd/_bwd suffix matching emits ambiguity warning when multiple candidates match | PASS | | AC-3 | All existing TestResolveOpClass tests pass unchanged | PASS | | AC-4 | New test covers suffix-match ambiguity (multiple fwd candidates) | PASS | | AC-5 | Full manifest validation passes | PASS | ## Test plan - [x] 25 TestResolveOpClass tests pass (24 pre-existing + 1 new) — 2.93s - [x] Full test suite: 75 passed, 0 failed - [x] `python scripts/validate_manifest.py` — all manifest checks passed - [x] pre-commit lint passes (ruff, codespell, mdformat) ## Follow-up No follow-up issues or suggestions. --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
e814e01bf9
|
[Fix][Validator] Emit warning on ambiguous op class resolution instead of silent fallback (#830)
## Summary When `_resolve_op_class` finds multiple candidate classes but no naming heuristic matches, it previously fell back silently to `candidates[0]` (alphabetical order). This could produce wrong validation results without any diagnostic. This PR replaces the silent fallback with a `UserWarning` that identifies the op name, all candidate class names, and the file path, then returns an unresolved result (`cls=None`) so the caller can handle it explicitly. Closes #810 ## Test plan - [x] AC-1: When all heuristics fail, a warning or error is emitted identifying the ambiguous resolution - [x] AC-2: Full manifest validation passes (no existing ops broken) — `python scripts/validate_manifest.py` -> All checks passed - [x] AC-3: Unit test covers the ambiguous-fallback code path — `test_ambiguous_fallback_returns_none_with_warning` added **Test results**: 79 tests passed, 0 failed ## Follow-up - #811 — Consolidate resolution heuristics + fix `_fwd`/`_bwd` suffix-match silent fallback Suggestions: strengthen ambiguous-fallback test to match full diagnostic fields (op_name, candidate_names, op_file) --------- Co-authored-by: Ibuki 🍃 — a wind born from Claude Opus <Ibuki-wind@users.noreply.github.com> |