forked from ccf-ai-infra/TileOPs-Metax
19 Commits
| Author | SHA1 | Message | Date |
|---|---|---|---|
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
|
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> |
|
|
Ang Gao
|
d3e26f7455
|
[Perf][Elementwise] vectorize bool bitwise storage path (#1719)
Refs #1702. Does not close it; #1702 remains the broader broadcast elementwise umbrella. ## Summary - Add uint8-backed bool-storage kernels for `bitwise_and`, `bitwise_or`, and `bitwise_xor`. - Route public `torch.bool` bitwise binary inputs through the uint8 storage path, then view outputs back to `torch.bool`. - Keep integer bitwise paths unchanged. - Add manifest benchmark metadata (`input_shape`, `other_shape`, `output_shape`, `broadcast_kind`) so broadcast and same-shape cases are visible in reports. - Add eager and `torch.compile(fullgraph=True)` coverage for bool bitwise ops. ## Notes from #1705 / #1708 - Avoids exposing raw bool tensors to TileLang vectorized kernels, so we do not hit the `boolx<N>` lowering issue from #1705. - Keeps the normal `BinaryOp.__init__` initialization path; the bool-storage path only customizes `_build_kernel_instance` and `_eager_forward`. - Adds a class-level `_bool_storage = False` default to avoid the ordering-coupling issue called out in #1708. ## Benchmark H200, official TileOpsGov runner image `ghcr.io/tile-ai/tileops-runner:65dbc98-torch2.10`, no tilelang/torch/dependency changes. Baseline is latest `origin/main` at `1a3866c`. | op | workload | before | after | delta | after vs torch | | --- | --- | --- | --- | --- | --- | | BitwiseAndFwdOp bool | same_shape | 0.0290 ms / 0.8690 TB/s | 0.0091 ms / 2.7596 TB/s | -68.6% | torch 0.0101 ms | | BitwiseOrFwdOp bool | same_shape | 0.0288 ms / 0.8726 TB/s | 0.0090 ms / 2.7860 TB/s | -68.8% | torch 0.0099 ms | | BitwiseXorFwdOp bool | same_shape | 0.0288 ms / 0.8724 TB/s | 0.0090 ms / 2.7919 TB/s | -68.8% | torch 0.0101 ms | | BitwiseAndFwdOp bool | channel_broadcast | 0.0421 ms / 0.6108 TB/s | 0.0089 ms / 2.8994 TB/s | -78.9% | torch 0.0464 ms | | BitwiseOrFwdOp bool | channel_broadcast | 0.0417 ms / 0.6158 TB/s | 0.0088 ms / 2.9122 TB/s | -78.9% | torch 0.0455 ms | | BitwiseXorFwdOp bool | channel_broadcast | 0.0420 ms / 0.6112 TB/s | 0.0091 ms / 2.8128 TB/s | -78.3% | torch 0.0464 ms | ## Test plan All test runs used `ghcr.io/tile-ai/tileops-runner:65dbc98-torch2.10` unless noted otherwise. - [x] `python3 -m ruff check tileops/kernels/elementwise.py tileops/ops/elementwise/bitwise.py tests/ops/test_bitwise.py tests/ops/test_elementwise_compile.py benchmarks/ops/bench_elementwise_manifest.py` - [x] `python3 scripts/validate_manifest.py` (advisory warnings only; manifest checks passed) - [x] `git diff --check` - [x] `pytest -q tests/ops/test_bitwise.py tests/ops/test_elementwise_compile.py::test_bitwise_binary_compile tests/ops/test_elementwise_compile.py::test_bool_bitwise_binary_compile --tb=short`: 41 passed - [x] `pytest -q tests/ops/test_elementwise_binary_broadcast.py tests/ops/test_logical.py tests/ops/test_comparison.py --tb=short`: 115 passed - [x] `pytest -q benchmarks/ops/bench_elementwise_manifest.py -k 'bitwise_and_manifest_bench or bitwise_or_manifest_bench or bitwise_xor_manifest_bench' --tb=short`: 18 passed - [x] `pre-commit run --files benchmarks/ops/bench_elementwise_manifest.py tests/ops/test_bitwise.py tests/ops/test_elementwise_compile.py tileops/kernels/elementwise.py tileops/ops/elementwise/bitwise.py` |
|
|
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
|
86c0d1025e
|
[Perf][Elementwise] vectorize bool storage paths (#1705)
Closes #1700. ## Summary - Remove the hot-path `int8 -> torch.bool` PyTorch cast from binary comparison/logical elementwise ops by making the TileLang kernels produce bool outputs directly. - Add uint8-backed bool-storage kernels for bool input paths, so public `torch.bool` tensors are viewed as `torch.uint8` in the Op layer and processed with vectorized-friendly storage kernels instead of TileLang `boolx<N>` loads/stores. - Move float predicate/logical unary bool-output paths from scalar `direct` to `explicit_parallel`, while keeping raw bool input protected from TileLang vectorized bool lowering. - Wire bool-storage dispatch for `logical_and`, `logical_or`, `logical_not`, and bool-input comparison ops without changing public API or output dtype. ## Performance tuning conclusions - Binary bool-output comparison/logical ops were dominated by a separate PyTorch `result.to(torch.bool)` cast kernel; removing that cast turns the hot path into a single TileLang kernel. - TileLang 0.1.11 cannot lower vectorized bool storage such as `boolx8`, so bool input paths need to avoid exposing bool tensors to vectorized TileLang kernels. - Viewing `torch.bool` tensors as `torch.uint8` preserves the public bool API while letting storage kernels use byte-oriented register-copy / explicit-parallel paths. - Predicate and logical unary float inputs can use `explicit_parallel`; only actual `torch.bool` inputs need to remain guarded from bool-vector lowering. ## Benchmark H200, TileLang 0.1.11, TileOpsGov local docker benchmark flow, effective HBM target = 4.0704 TB/s. | kernel | workload | before | after | SOL before | SOL after | |--------|----------|--------|-------|------------|-----------| | eq | hidden-state-prefill fp16 | 0.0367ms / 1.14 TB/s | 0.0147ms / 2.86 TB/s | 28.0% | 70.3% | | eq | hidden-state-prefill fp32 | 0.0462ms / 1.63 TB/s | 0.0231ms / 3.27 TB/s | 40.0% | 80.3% | | logical_and | hidden-state-prefill bool | 0.0514ms / 0.49 TB/s | 0.0088ms / 2.86 TB/s | 12.0% | 70.3% | | logical_and | cnn-feat-broadcast bool | 0.0752ms / 0.34 TB/s | 0.0089ms / 2.90 TB/s | 8.4% | 71.2% | | logical_or | hidden-state-prefill bool | 0.0514ms / 0.49 TB/s | 0.0088ms / 2.85 TB/s | 12.0% | 70.0% | | logical_not | 16M bool | 0.0507ms / 0.66 TB/s | 0.0108ms / 3.11 TB/s | 16.2% | 76.4% | | logical_not | 256M bool | 0.7809ms / 0.69 TB/s | 0.1273ms / 4.22 TB/s | 17.0% | 103.7% | | isnan | 16M fp16 | 0.0549ms / 0.92 TB/s | 0.0212ms / 2.38 TB/s | 22.6% | 58.5% | | isnan | 256M fp16 | 0.8275ms / 0.97 TB/s | 0.2732ms / 2.95 TB/s | 23.8% | 72.5% | ## Test plan - [x] `python3 -m pytest -q benchmarks/ops/bench_unary_elementwise.py -k 'logical_not' --junit-xml=/workspace/TileOPs/results/elementwise_bool_storage_logical_not.xml --tb=short --timeout=900 --timeout-method=thread`: 4 passed - [x] `python3 -m pytest -q benchmarks/ops/bench_unary_elementwise.py::test_logical_not_bench benchmarks/ops/bench_elementwise_manifest.py -k 'logical_and_manifest or logical_or_manifest' --junit-xml=/workspace/TileOPs/results/elementwise_bool_storage_targeted.xml --tb=short --timeout=900 --timeout-method=thread`: 16 passed - [x] Full official elementwise sweep: `bench_unary_elementwise.py`, `bench_elementwise_manifest.py`, `bench_binary_elementwise.py`, `bench_independent_elementwise.py`: 552 passed, 24 skipped |
|
|
Ang Gao
|
af57636ccc
|
[Feat][Manifest] Promote special elementwise ops (#1656)
## Summary Closes #1659. Refs #1650. Promotes implemented special elementwise ops from `spec-only` to `implemented`: - `SiluAndMulFwdOp`, `GeluAndMulFwdOp`, `GeluTanhAndMulFwdOp` - `AlibiFwdOp`, `SinusoidalFwdOp` This is a code/manifest cleanup PR, not a release-plan doc update. #1659 is tracked as a sub-issue of the broader #1650 release-facing manifest coverage tracker. ## Changes - Allow fused gated ops to bind `M`, `N`, and `dtype` from the runtime input tensor, while preserving the old explicit ctor path. - Keep lazy fused gated ops reusable across runtime shapes/dtypes by separating explicit ctor contracts from the last runtime-bound kernel specialization. - Ensure `_eager_forward` validates the runtime input and calls `_ensure_kernel(M, N, x.dtype)` before launch. - Add fused gated `eval_roofline()` support and the correct GELU-tanh FLOPs coefficient. - Add zero-input `_validate_dtypes()`, `_infer_output_shapes()`, `total_memory`, and `eval_roofline()` for `AlibiFwdOp` and `SinusoidalFwdOp`. - Keep generative op `dtype` explicit in the manifest because these ops have no input tensor to infer dtype from. - Add smoke tests for manifest-style fused gated lazy binding and shape rebinding. ## Validation Ran in `ghcr.io/tile-ai/tileops-runner:65dbc98-torch2.10`: ```bash python -m pytest tests/test_ops_manifest.py -q # 21 passed python -m pytest tests/test_ops_manifest.py::TestOpSchema::test_workloads_include_all_required_params -q # 1 passed for op in SiluAndMulFwdOp GeluAndMulFwdOp GeluTanhAndMulFwdOp AlibiFwdOp SinusoidalFwdOp; do python scripts/validate_manifest.py --check-op "$op" --strict done # all targeted checks passed python -m pytest \ tests/ops/test_fused_gated.py::test_silu_and_mul_infers_manifest_shape_contract \ tests/ops/test_fused_gated.py::test_silu_and_mul_lazy_op_rebinds_shape \ tests/ops/test_fused_gated.py::test_fused_gated_rejects_integer_dtype -q # 3 passed python -m ruff check \ tileops/ops/elementwise/_base.py \ tileops/ops/elementwise/fused_gated.py \ tileops/ops/elementwise/alibi.py \ tileops/ops/elementwise/sinusoidal.py \ tests/ops/test_fused_gated.py # All checks passed ``` GitHub CI is green, including `validate-manifest`, `benchmark-contract-tests`, and `gpu-smoke`. |
|
|
lcy-seso
|
12056d153d |
[Chore] strip scaffolding comments and validator-narration docstrings
Tighten agent-facing doc + source/test comments accumulated across the
testbed nightshift cycle:
- .claude/rules/manifest-trust-model.md: replace 40-word validator-AST
Why-clause on the source.test/source.bench bullet with a one-line
principle (discoverability pointers, not contractual fields).
- tileops/ops/moe/experts/{nopad,padded}.py forward(): drop docstrings
that narrated "the validator's Op-class resolution path".
- tileops/ops/op_base.py: shrink the composite-op carve-out comment
to one line.
- tileops/ops/elementwise/masked_fill.py: tighten the bool->uint8
routing comment (both occurrences).
- tileops/kernels/elementwise.py: fold the MaskedFillTensorValue
bool/uint8 paragraph into the existing class-level docstring.
- tileops/perf/formulas.py: shorten the fused-moe roofline docstring;
keep the func-mode rationale, drop the Args/Returns narration.
- tests/ops/test_moe_fused_moe.py: trim test_fused_moe_fwd_op_identity
docstring to the contract under test.
- tests/test_op_base.py: tighten TestCompositeKernelMapOverride class
docstring; drop a redundant method docstring that restated its assert.
No behavior change. Validator --strict and the touched test files pass.
Co-Authored-By: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
cb85703a23 |
[Feat][Elementwise] cover full dtype union for MaskedFillTensorValueFwdKernel; flip MaskedFillFwdOp to implemented (#1499)
Closes #1498 ## Summary - Extend `MaskedFillTensorValueFwdKernel.SUPPORTED_DTYPES` to the full manifest signature union — `float16 | bfloat16 | float32 | bool | uint8 | int8 | int16 | int32 | int64`. - Route `bool` through `uint8` storage at the Op layer, mirroring `MaskedFillScalarFwdOp`. - Flip `MaskedFillFwdOp.status: spec-only -> implemented` in `tileops/manifest/elementwise_binary.yaml` (status-only edit; carve-out compliant). - Expand `tests/ops/test_special_elementwise_conformance.py::test_masked_fill_tensor_value` parametrization to all 9 manifest dtypes × 4 shapes. ## Test plan - [x] `pytest tests/ops/test_special_elementwise.py tests/ops/test_special_elementwise_conformance.py -k masked_fill` — 67 passed - [x] `pytest tests/ops/test_special_elementwise_conformance.py::test_masked_fill_tensor_value` — 36 passed (9 dtypes × 4 shapes) - [x] `python scripts/validate_manifest.py --strict --check-op MaskedFillFwdOp` — `All manifest checks passed` - [x] `python scripts/validate_manifest.py --strict` — `All manifest checks passed` - [x] `pytest benchmarks/ops/bench_masked_fill.py -v` — 5 params pass ## Test node delta ``` File Base HEAD Delta -------------------------------------------------------------------------- tests/ops/test_special_elementwise_conformance.py 49 81 +32 -------------------------------------------------------------------------- TOTAL 49 81 +32 ``` Parametrization expanded to cover the full 9-dtype manifest signature × 4 shape variants on `test_masked_fill_tensor_value`, matching the dtype union the kernel now constructs for. Co-authored-by: Ibuki — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
7e8deff27b |
[Feat][Codegen] generate eval_roofline bodies from manifest roofline (#1471)
Closes #1459 ## Summary - New `tileops/ops/_roofline_codegen.py`: synthesizes each op's `eval_roofline()` from its manifest `roofline` block. Two modes — `func` (dotted path eagerly resolved at class build) and `inline` (vars + flops + bytes, AST-validated against the §4.4.3/§4.4.4 namespace, emitted as plain Python with no runtime `eval`). - `Op.__init_subclass__` installs both `_validate_dtypes` (from #1466) and `eval_roofline` (new) on every `status: implemented` subclass; skipped on spec-only ops and on classes / bases that supply their own override (MRO-checked). - Inline mode reads op state under a 2-tier contract per `signature.inputs`/`params` names *referenced* by the roofline: tensor inputs as `self.<input>` (with both `.shape` and `.ndim`) or `self.<input>_shape` (shape tuple/list); params as `self.<param>`. Unreferenced manifest names are not bound. Missing required bindings raise `ValueError(op, input, conventions)`. - Vars-layer AST validator is scope-aware: each comprehension generator's iterable is visited before its target binds (matching Python semantics); `.shape` / `.ndim` are accepted only directly on a declared tensor-input Name — chained (`x.shape.ndim`), subscripted (`x.shape[0].shape`), and local-operand (`N.shape`) forms reject; tensor inputs may not appear as bare values (no `sum(x)`, no `x[0]`); calls restricted to the §4.4.4 helper table; vars keys colliding with inputs/params/helpers/`elem_bytes`/earlier vars are rejected at synthesis. - Arithmetic-layer validator is a positive AST allowlist (BinOp, UnaryOp, BoolOp, IfExp, Compare, Call to allowed helper, Constant, Name) — collection literals (`Tuple`/`List`/`Set`/`Dict`) and any other node kind fail synthesis. - `PreluFwdOp` / `NanToNumFwdOp` (the two real inline-mode ops) declare bindings explicitly: `self.input_shape`, `self.weight_shape`, `self.input_shape`. - `docs/design/roofline.md` §4.4.3 names the referenced-only binding decision and links to the rule. - `.claude/domain-rules/ops-design.md` adds the concrete attribute contract for op authors. ## Test plan - [x] pre-commit passed - [x] `python -m pytest tests/test_roofline_codegen.py` — 36 passed - [x] `python -m pytest tests/test_validate_manifest.py` — 221 passed - [x] All 108 `status: implemented` ops synthesize without error against the live manifest - [x] `python scripts/validate_manifest.py --strict`: this PR is *scoped to the C7 class* (`eval_roofline is the Op base stub`), and that count is **0** on this branch (was **41** on `testbed`). The script still exits non-zero with **58 `kw_only mismatch` errors in reduction ops** that are pre-existing on `testbed` (`testbed` itself reports 99 errors = 41 C7 + 58 kw_only) and are out of scope for this PR — they are tracked separately. ## Test node delta `python scripts/test_node_delta.py --base upstream/testbed`: ``` File Base HEAD Delta ------------------------------------------------------- tests/test_roofline_codegen.py new 36 (new) ------------------------------------------------------- TOTAL 0 36 +36 ``` **Justification:** new codegen module; tests cover both synthesis modes, the install-hook condition matrix (spec-only / explicit override / inherited override / missing metadata), the 2-tier binding contract (positive + tier-1 shape/ndim fallthrough + missing-binding error), vars-layer comprehension scope (target binds after iter), attribute whitelist + chained/subscripted/local-operand rejection, tensor-input bare-reference rejection, vars-key collision rejection, arithmetic-layer collection-literal rejection, param exposure contract, and the live-manifest C7 parity invariant. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
4a776deeb9 |
[BugFix][MANIFEST] align kw_only on 34 implemented ops to match impl (#1483)
Closes #1456 ## Summary Aligns `kw_only` between manifest and code `__init__` for the 58 C3 parity errors, with per-`(op, param)` direction matching the live PyTorch API (`torch.<op>` / `torch.nn.functional.<op>`). Two routing outcomes are mixed in this single PR: - **Code → positional-or-keyword (manifest stays `false`)** — applied where PyTorch keeps the param positional. Touches `tileops/ops/elementwise/{_base.py,activations.py}` (in-place activation `inplace`), `tileops/ops/norm/batch_norm.py` (`training`/`momentum`/`eps`; also swaps `*spatial: int` varargs for `spatial: Tuple[int, ...] = ()` so `__init__` can declare those three as positional-or-keyword while matching `torch.nn.functional.batch_norm`), `tileops/ops/reduction/{_softmax_base,all_op,any_op,argmax,argmin,count_nonzero,cumulative_base,inf_norm,l1_norm,l2_norm,log_softmax,logsumexp,reduce,softmax}.py` (`dim`/`keepdim`/`ord`). - **Manifest → `kw_only: true` (code stays kw-only)** — applied only to `tileops/manifest/elementwise_binary.yaml` for `AddFwdOp.alpha`, `SubFwdOp.alpha`, `DivFwdOp.rounding_mode`; PyTorch places these after `*`. Test/bench callers (`tests/ops/{test_batch_norm.py, test_elementwise_unary_activation_alignment.py, test_norm_ops.py}`, `benchmarks/ops/bench_batch_norm.py`) are updated to pass the now-positional params correctly (notably `BatchNormFwdOp((N, C), (H, W), eps=..., training=...)` with the `spatial` tuple). ## Test plan - [x] `python scripts/validate_manifest.py --strict 2>&1 | grep -c "kw_only mismatch"` → `0` (down from `58` on testbed) - [x] `python -m pytest -q tests/test_validate_manifest.py tests/test_ops_manifest.py` → 242 passed - [x] Affected per-family tests: `python -m pytest -q tests/ops/{activation,norm,reduction,scan}/ tests/ops/test_norm_ops.py tests/ops/test_batch_norm.py tests/ops/test_elementwise_unary_activation_alignment.py` → 1105 passed, 2 skipped - [x] `pre-commit run --all-files` → clean - [x] AC-1 (manifest validator strict): 0 `[ctor] … kw_only mismatch` errors. Pre-existing `[stub] eval_roofline` errors (41) are unchanged on baseline — out of scope. - [x] AC-2 (no demotion): no op was demoted from `status: implemented` to `status: spec-only`. - [x] AC-3 (all 34 ops addressed): every `[ctor]` C3 site enumerated by the validator on testbed is now resolved. - [x] AC-5 (per-family direction documented): see Summary above and the table in #1456. ## Test node delta ``` File Base HEAD Delta ------------------------------------------------------------------------------ tests/ops/test_batch_norm.py N+0 N+0 0 tests/ops/test_elementwise_unary_activation_alignment.py M+0 M+0 0 tests/ops/test_norm_ops.py K+0 K+0 0 ------------------------------------------------------------------------------ TOTAL 0 ``` **Justification:** Pure signature-alignment PR — every `__init__` and call-site edit preserves the existing test surface. No new test nodes introduced; existing fixtures simply pass the now-positional params correctly. --------- Co-authored-by: Ibuki <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> |
|
|
Cao Ying
|
dfebdd432d |
[Feat][Kernels] int/bool kernel coverage for MaskedFillScalar (#1411)
Closes #1246 - Extend `MaskedFillFwdKernel.SUPPORTED_DTYPES` to cover the int/bool dtypes the manifest declares for `MaskedFillScalarFwdOp` (bool, uint8, int8, int16, int32, int64, float16, bfloat16, float32). - Generalize `_clamp_to_dtype_range` and `_validate_scalar_param_repr` to match PyTorch's `Tensor.masked_fill` scalar coercion: - **Signed int dtypes**: range-check the real float against `iinfo.min/max`, then truncate toward zero (`int32.fill(1.5) -> 1`, `int8.fill(127.5)` raises because `127.5 > 127`). - **`torch.uint8`**: additionally accepts Python ints in `[-255, 0)` and wraps via `value & 0xFF` (`fill(-1) -> 255`). Float negatives stay rejected (matches PyTorch). - **fp16 / bf16 / fp32**: `+/-Inf` and `NaN` pass through untouched. `fp8_e5m2` preserves `Inf`; only `fp8_e4m3fn` (no Inf representation) saturates to `finfo` extrema. - **bool**: any truthy/falsy int or float reduces to `{0, 1}`. - Validator gains an `allow_nonfinite_float` opt-in so masked_fill accepts `Inf`/`NaN` while `elu` / `softplus` / `clamp` keep their finite-only contract. - Dispatch `bool` at the Op layer by viewing input as `uint8` and re-viewing the result as `bool`; TileLang does not vectorize bool storage. Kernel + Op changes only. Manifest entry is **byte-identical** in this PR. - [x] AC-1: `pytest tests/ops/test_special_elementwise.py tests/ops/test_elementwise_independent_fp8.py tests/ops/test_elementwise_caching_autotune.py` -> 131 passed - [x] AC-2: All 9 manifest dtypes (bool, uint8, int8, int16, int32, int64, float16, bfloat16, float32) construct, compile, run, and match `torch.Tensor.masked_fill` reference - [x] AC-3: Scalar coercion validator matches PyTorch exactly: rejects out-of-range integers, `Inf`/`NaN` into int dtypes, and uint8 float negatives; accepts uint8 int wraparound and float `Inf`/`NaN` for floating-point dtypes - [x] AC-4: Per-dtype correctness tests added to `tests/ops/test_special_elementwise.py`: - signed int min/max + uint8 negative wrap (`test_masked_fill_int_dtypes`) - bool truthy/falsy fill values via uint8 storage view (`test_masked_fill_bool`) - `+/-Inf` and `NaN` through fp16/bf16/fp32 (`test_masked_fill_float_nonfinite`) - validator rejection set (`test_masked_fill_rejects_out_of_range`) - [x] pre-commit passed `python scripts/test_node_delta.py --base upstream/testbed` ``` File Base HEAD Delta -------------------------------------------------------------- tests/ops/test_special_elementwise.py 63 81 +18 -------------------------------------------------------------- TOTAL 63 81 +18 Growth: +28.6% ``` --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
731744d0d5 |
[Feat][Kernels] rounding-mode div kernel path for trunc/floor (#1410)
Closes #1245 ## Summary Adds `rounding_mode='trunc' | 'floor'` support to `DivFwdOp`. Previously only `rounding_mode=None` (true division) worked; non-None raised `NotImplementedError`. - `tileops/kernels/elementwise.py`: new `DivTruncFwdKernel` (fp32 promotion → divide → `T.trunc` → cast back). `'floor'` reuses existing `FloorDivideFwdKernel` (byte-identical to what an early draft `DivFloorFwdKernel` would have been; that draft class was removed during review). - `tileops/ops/elementwise/arithmetic.py`: `_DIV_KERNEL_BY_ROUNDING_MODE` dispatch (`None → DivFwdKernel`, `'trunc' → DivTruncFwdKernel`, `'floor' → FloorDivideFwdKernel`). Forward-time `NotImplementedError` guard removed; invalid modes raise `ValueError` at construction. - `tests/ops/test_binary_arith.py`: `test_div_rounding_mode_eager` (6 cells: trunc/floor × fp16/bf16/fp32, mixed-sign quotients vs `torch.div`) + `test_div_rounding_mode_dispatch` (CPU-only dispatch table + invalid-mode rejection). - `tests/ops/test_elementwise_compile.py`: `test_div_rounding_mode_compile` (6 cells under `torch.compile(fullgraph=True)`, same oracle). ## Test plan - [x] `pytest tests/ops/test_binary_arith.py tests/ops/test_elementwise_compile.py -m smoke` — 89 passed. - [x] Manual parity vs `torch.div(..., rounding_mode=...)` on fp16/bf16/fp32; tolerances follow the existing `FloorDivideFwdKernel` convention (`atol=1.0, rtol=0` for half precision rounding-mode kernels, `1e-5` for fp32). - [x] `scripts/validate_manifest.py` clean on `DivFwdOp`. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
4f4ccfb90f |
[Feat][Kernels] alpha-scaled add/sub kernel path for non-default alpha (#1409)
Closes #1244 ## Summary - Bake the scalar `alpha` multiplier into `AddFwdKernel` / `SubFwdKernel` at construction time so `AddFwdOp(alpha=k)` and `SubFwdOp(alpha=k)` for `k != 1` dispatch to a real kernel instead of raising `NotImplementedError`. - Scalar multiply runs in fp32 (per code-style: no narrow-type literal casts) and casts back to storage dtype at the boundary; `alpha == 1` keeps the original `op_func` to preserve the integer-dtype fast path byte-identically. - Removes the `NotImplementedError` raise in `tileops/ops/elementwise/arithmetic.py` for the `alpha != 1` path. ## Scope (kernel-only PR) Per [`docs/design/trust-model.md`](docs/design/trust-model.md), this PR touches only `tileops/kernels/` and `tileops/ops/elementwise/`. Test and benchmark additions for the new `alpha != 1` path are deferred to a sibling PR owned by the `tests/` and `benchmarks/` layers respectively. - AC-1 / AC-2 (kernel correctness): verified in this PR. - AC-3 (per-dtype correctness tests for `alpha != 1`): deferred to a sibling test-layer PR. - AC-4 (benchmark entry at LLaMA workload): deferred to a sibling benchmark-layer PR. ## Test plan - [x] AC-1: `pytest tests/ops/test_binary_arith.py` — 120 passed (existing `alpha=1` fast path, no regression). - [x] AC-1: `pytest tests/test_validate_manifest.py` — 218 passed. - [x] AC-2: `AddFwdOp` / `SubFwdOp` × {fp32, fp16, bf16} × `alpha ∈ {1, 2, -1, 0.5}` — 24/24 cases match `torch.add` / `torch.sub` with `max_err = 0.0` (kernel compile log confirms TileLang dispatch per `(op, dtype, alpha)` specialization, no `NotImplementedError`, no torch passthrough). - [ ] AC-3: per-dtype correctness tests for `alpha != 1` — deferred to sibling tests-layer PR. - [ ] AC-4: benchmark entry at LLaMA workload for `alpha != 1` — deferred to sibling benchmarks-layer PR. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
7a153cfb30 |
[Refactor][MANIFEST] unify per-element flop convention for activation/clamp (#1389)
Closes #1379 ## Summary - Document a single per-element FLOP convention in `docs/design/roofline.md` (Convention subsection, dated) covering arithmetic, transcendental, compare-and-select, and clamp ops. - Apply the convention to activation helpers (`relu`, `leaky_relu`, `prelu`, `gelu`, `silu`, `swish`, `tanh`, `sigmoid`, `hardtanh`) and the clamp/min-max family (`clamp_scalar`, `clamp_min`, `clamp_max`, `maximum`, `minimum`) in `tileops/perf/formulas.py` and the corresponding manifest entries. - Each helper now carries a one-line `# FLOPs: ...` derivation comment; the clamp/min-max family returns the same per-element FLOP count. - Add `scripts/perf/flop_convention_delta.py` plus checked-in artifacts under `docs/perf/` so the before/after delta is reproducible without a GPU. ## Test plan - [x] AC-1: `pytest tests/perf/ tests/test_validate_manifest.py` — 288 passed. - [x] AC-2: `docs/design/roofline.md` Convention subsection present and dated. - [x] AC-3: each activation helper carries a `# FLOPs: ...` derivation comment. - [x] AC-4: `hardtanh` / `clamp_scalar` / `clamp_min` / `clamp_max` / `maximum` / `minimum` return identical per-element FLOP counts. - [x] AC-5: before/after roofline-utilization table for ≥3 representative benches included below. - [x] AC-6: external tracker flip — out of scope. ## Benchmark ## Per-element FLOP convention — before/after delta Formula-only evaluation (no GPU required). Reproduce with: ```bash python scripts/perf/flop_convention_delta.py \ --out docs/perf/flop_convention_delta.csv ``` | family | op | label | shape | dtype | flops before | flops after | flops delta | bytes before | bytes after | bytes delta | | ---------- | ------------- | -------------------- | --------- | ------- | -----------: | ----------: | ----------: | -----------: | ----------: | ----------: | | activation | ReluFwdOp | hidden-state-prefill | 2048x4096 | float16 | 16,777,216 | 8,388,608 | -8,388,608 | 33,554,432 | 33,554,432 | 0 | | clamp | HardtanhFwdOp | hidden-state-prefill | 2048x4096 | float16 | 33,554,432 | 8,388,608 | -25,165,824 | 33,554,432 | 33,554,432 | 0 | | min-max | ClampFwdOp | elementwise-16M | 4096x4096 | float16 | 33,554,432 | 16,777,216 | -16,777,216 | 134,217,728 | 134,217,728 | 0 | `flops before` reflects the coefficients on `upstream/testbed` immediately before the convention commit. `flops after` is evaluated from the manifest formulas (`ReluFwdOp` / `HardtanhFwdOp`) or `clamp_fwd_roofline` (`ClampFwdOp`) on the current checkout. Byte counts are unchanged by the convention. For these elementwise workloads `memory_time` already dominates, so the FLOP-coefficient reduction shifts each workload further into the memory-bound regime without changing predicted achievable bandwidth. See [docs/perf/flop_convention_delta.md](docs/perf/flop_convention_delta.md) and [docs/perf/flop_convention_delta.csv](docs/perf/flop_convention_delta.csv). --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
27a4e532de |
[Refactor][Ops] align W0 ops with manifest parity check (#1393)
Closes #1376 ## Summary - Align 8 W0 ops with the manifest parity check (default `validate_manifest.py` mode reports zero `[shape]/[dtype]` warnings for all 8). - `WhereFwdOp`: add mechanical `_infer_output_shapes` and `_validate_dtypes` overrides matching the manifest broadcast/dtype contract. - Demote `MaskedFillFwdOp`, `MaskedFillScalarFwdOp`, and `RMSNormFwdOp` to `status: spec-only` with documented per-op reasons in their manifest entries. - Clamp family (`ClampFwdOp`, `ClampScalarFwdOp`, `ClampMinFwdOp`, `ClampMaxFwdOp`) was already aligned — no code change needed. ## Per-op approach | Op | Approach | Notes | | --- | --- | --- | | `ClampFwdOp` | already aligned | overrides match manifest; no change | | `ClampScalarFwdOp` | already aligned | overrides match manifest; no change | | `ClampMinFwdOp` | already aligned | overrides match manifest; no change | | `ClampMaxFwdOp` | already aligned | overrides match manifest; no change | | `WhereFwdOp` | mechanical override added | `_infer_output_shapes` + `_validate_dtypes` aligned with manifest broadcast/dtype contract | | `MaskedFillFwdOp` | demoted to `spec-only` | implementation deviates from spec; flip in follow-up impl PR | | `MaskedFillScalarFwdOp` | demoted to `spec-only` | implementation deviates from spec; flip in follow-up impl PR | | `RMSNormFwdOp` | demoted to `spec-only` | implementation deviates from spec; flip in follow-up impl PR | ## Test plan - [x] Modified files pass unit tests (pytest tests/ops/) — 318 passed across the changed-op test files (`test_elementwise_compile.py`, `test_special_elementwise.py`, `test_special_elementwise_conformance.py`, `test_elementwise_independent_fp8.py`, `test_elementwise_unary_activation_alignment.py`, `test_rms_norm.py`, `test_normalization_alignment.py`). - [x] `python scripts/validate_manifest.py` emits zero parity warnings for the 8 W0 ops (default mode). - [x] `tests/test_validate_manifest.py` passes — 218 passed. - [x] Each op is either covered by both overrides aligned with manifest, or demoted to `status: spec-only` with reason in this description. - [x] All 8 W0 ops are explicitly addressed in the per-op approach table above. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
Cao Ying
|
e01a88eea8 |
[Refactor][OPS] split tileops/ops/elementwise.py into per-cluster package (#1388)
Closes #1378 ## Summary - Replaced 3415-line monolith `tileops/ops/elementwise.py` with a `tileops/ops/elementwise/` package, one module per cluster (16 cluster modules + `_base.py`). - Umbrella classes `UnaryOp`, `BinaryOp`, `FusedGatedOp` moved to `_base.py`; `__init__.py` re-exports every previously-public symbol via explicit `__all__`. - Pure file-layout move: no class signature, docstring, or kernel-binding changes; all 295 elementwise tests pass; no file-level lint suppressions introduced. - Pyright shows pre-existing typing issues in the relocated code that were carried over verbatim from the original file — out of scope to fix in this refactor. ## Test plan - [x] AC-1: Modified files pass unit tests (pytest tests/ops/). - [x] AC-2: tileops/ops/elementwise.py no longer exists; tileops/ops/elementwise/ is a package with __init__.py. - [x] AC-3: python -c "from tileops.ops.elementwise import UnaryOp, BinaryOp, FusedGatedOp" succeeds. - [x] AC-4: python -c "from tileops.ops.elementwise import SiluAndMulFwdOp, ReluFwdOp, AddFwdOp, ExpFwdOp, DivFwdOp" succeeds. - [x] AC-5: tests/ops/test_elementwise_fp8.py passes without modification to its imports. - [x] AC-6: grep -nE '^# ruff:|^# flake8:' tileops/ops/elementwise/ returns nothing. - [x] AC-7: External tracker flip — out of scope. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |