forked from ccf-ai-infra/TileOPs-Metax
12 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
|
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`. |
|
|
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
|
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
|
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> |