[Maintain][Manifest] Add adaptive 2D pooling ops

[RMLYC]

Part of #1799

Summary

• Add spec-only manifest entries for adaptive average pooling, adaptive max pooling, and the max-pooling indices variant.
• Define FP16/BF16 signatures, PyTorch-aligned output_size rules, representative workloads, and exact adaptive-bin roofline formulas.
• Register future kernel, Op, test, and benchmark paths while keeping manifest-driven benchmarks disabled until implementation lands.

Test plan

• python scripts/validate_manifest.py
• Pre-commit CI passed
• Roofline synthesis succeeds for all three entries
• Test node delta: 0 (no tests changed)

[gemini-code-assist]

Caution

The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased.

[gemini-code-assist]

Caution

The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased.

[superAngGao]

Thanks for adding the adaptive 2D pooling specs. I tested this in the TileOpsGov runner image requested for CI parity:

ghcr.io/tile-ai/tileops-runner:65dbc98-torch2.10

with the source and cache mounts, without installing or upgrading tilelang, torch, or any dependencies. Full manifest validation passes:

python scripts/validate_manifest.py -> All manifest checks passed.

I think there are two manifest/spec issues to fix before landing:

1. The adaptive pooling roofline FLOP expression undercounts non-divisible output grids.

The new entries use:

kH = (H_in + out_H - 1) // out_H
kW = (W_in + out_W - 1) // out_W
flops = N * C * out_H * out_W * kH * kW

For PyTorch adaptive pooling, each output bin is based on floor/ceil bin boundaries, so bin sizes vary and may overlap. For the PR's own nondiv-7x7 workload [2, 64, 55, 57] -> [7, 7], I get:

• manifest formula: 451584
• adaptive-bin scan count: 491904
• ratio: 0.918033

So the formula is about 8.2% low for this representative workload. This affects all three new ops: AdaptiveAvgPool2dFwdOp, AdaptiveMaxPool2dFwdOp, and AdaptiveMaxPool2dIndicesFwdOp.

Could we either switch this to an exact adaptive-bin formula/helper, or avoid pairing a non-divisible adaptive workload with this approximate roofline expression?

2. The output_size shape rules accept tuple/list forms that PyTorch rejects.

The current rules treat a one-element tuple/list like a square output, but PyTorch only does that for a bare int. In the same runner image:

• adaptive_avg_pool2d(x, 6) -> OK, output [N, C, 6, 6]
• adaptive_avg_pool2d(x, [6]) -> RuntimeError: output_size must be 2
• adaptive_avg_pool2d(x, (6,)) -> RuntimeError: output_size must be 2
• adaptive_avg_pool2d(x, [None, 3]) -> OK

So I think the manifest should require tuple/list output_size to have length 2, and reserve the square-output behavior for scalar int only.

One non-blocking note: targeted validate_manifest.py --check-op AdaptiveAvgPool2dFwdOp fails because these are status: spec-only entries and the corresponding Op classes do not exist yet. That matches the manifest design doc behavior for targeted L0-L4 checks, so I am not treating that as a blocker here.

[zhen8838]

Request changes on head ceed9f7.

1. The three adaptive entries still describe only [N, C, H_in, W_in], while the referenced PyTorch API also accepts unbatched [C, H_in, W_in]. The spec must either include the CHW form or explicitly narrow the supported API before promotion.
2. output_size is incomplete and permissive: the parameter type omits scalar None, and the len(...) > 1 fallback in the shape rules accepts a one-element tuple/list as a square output. Require the exact accepted forms and a two-element tuple/list, matching the runtime contract.
3. The roofline's ceil(H_in/H_out) and ceil(W_in/W_out) treat all adaptive bins as uniform. Adaptive bins vary and can overlap for non-divisible and expanded outputs, so the current FLOPs are not authoritative. Define the work from the actual bin boundaries before marking these entries implemented.

[RMLYC]

@superAngGao Both points verified locally with the same numbers (451584 vs 491904, ratio 0.918). Addressed in 0972fde:

1. Roofline switched to the exact adaptive-bin scan: scan_H = H_in + sum(1 for j in range(1, out_H) if (j * H_in) % out_H != 0) (same for W), flops = N * C * scan_H * scan_W. This is exact by construction — it sums the true per-bin sizes (separable over H/W), counting the one-row/col overlap of adjacent bins on non-divisible grids. It evaluates to 491904 for nondiv-7x7, matches the previous values on divisible workloads, and also fixes expanded outputs, where the ceil approximation was ~44% low.

2. Shape rules now require isinstance(output_size, int) or (isinstance(output_size, (tuple, list)) and len(output_size) == 2), and the len > 1 fallback is dropped — matching the runtime contract ([6] / (6,) → RuntimeError). The Op implementation in [Feat][Pool] Implement adaptive 2D pooling ops #1809 already rejects non-length-2 tuples, so the two PRs stay consistent.

[RMLYC]

@zhen8838 Addressed in 0972fde:

1. The entries intentionally declare the fixed-rank batched NCHW form, as every existing pool entry does (F.avg_pool2d / F.max_pool2d also accept unbatched CHW). The Op implementation in [Feat][Pool] Implement adaptive 2D pooling ops #1809 accepts CHW by unsqueezing to NCHW. All three entries now carry an explicit comment documenting this narrowing.

2. Scalar None is rejected by PyTorch itself: F.adaptive_avg_pool2d(x, None) → TypeError: object of type 'NoneType' has no len()' (torch 2.10). The declared type int | tuple[int | None, int | None] already matches the runtime contract, so we are not adding scalar None. The length-2 requirement for tuple/list is added in the same commit.

3. Roofline now defines work from the actual bin boundaries (exact scan formula, see the thread above), verified against a per-bin enumeration.

[RMLYC]

@zhen8838 Follow-up on point 2 — you were right about the documented contract: the torch 2.13 reference pages declare output_size (int | None | tuple[int | None, int | None]) for both nn.AdaptiveAvgPool2d and nn.AdaptiveMaxPool2d. The manifest type now includes scalar None (115007f), resolving to (None, None) — output equals the input spatial size.

One nuance worth recording: the torch 2.10 runtime still raises TypeError for scalar None in all three call paths (F.adaptive_avg_pool2d, nn.AdaptiveAvgPool2d, nn.AdaptiveMaxPool2d), so the Op in #1809 implements the documented semantics rather than the current runtime behavior, and its tests map the PyTorch reference to (None, None) until the runtime catches up.

[zhen8838]

Re-reviewed current head 115007f.

The previous manifest blockers are resolved: output_size now covers scalar None and rejects non-2-element tuple/list forms, and the roofline uses exact adaptive-bin scan counts rather than a uniform ceil-divided window. The manifest remains correctly status: spec-only, with the NCHW pool-family convention and CHW Op-layer support explicitly documented. Manifest, compile-contract, benchmark-contract, pre-commit, actionlint, and gitleaks checks pass; GPU smoke is appropriately skipped for this spec-only diff.

[superAngGao]

Thanks for the updates. The adaptive-bin math looks correct now: in the official runner image, the new formula evaluates the PR's nondiv-7x7 workload to 491904, matching a direct per-bin enumeration, so the previous undercount is fixed.

I think there is still one blocking issue before this manifest can land:

The inline roofline expressions are not synthesizeable by the current TileOps roofline codegen. The new vars entries reference shape symbols such as N, C, H_in, and W_in directly, e.g. out_H: "H_in if oH is None else oH". However, inline roofline codegen only binds declared tensor inputs, params, elem_bytes, helpers, and earlier vars; it does not automatically materialize shape dimension names as Python locals. In the official runner image, this reproduces with:

from tileops.manifest import load_manifest
from tileops.ops._roofline_codegen import synthesize_eval_roofline
entry = load_manifest()["AdaptiveAvgPool2dFwdOp"]
synthesize_eval_roofline(
    "AdaptiveAvgPool2dFwdOp",
    roofline=entry["roofline"],
    signature=entry["signature"],
)

which raises:

ValueError: AdaptiveAvgPool2dFwdOp: roofline.vars['out_H'] references unknown name 'H_in'

This affects all three new adaptive pooling entries. Full manifest validation does not catch it because status: spec-only only gets L0 schema validation, and the roofline design doc explicitly says codegen is the authoritative gate for name/form correctness.

A minimal fix is to define the shape-derived locals explicitly in roofline.vars before the existing adaptive-bin variables:

N: "input.shape[0]"
C: "input.shape[1]"
H_in: "input.shape[2]"
W_in: "input.shape[3]"

I tested that shape-binding approach locally in the same runner image with the existing formula body, and synthesize_eval_roofline() then succeeds and returns (491904, 815104) for the nondiv-7x7 avg-pool workload.

Non-blocking note: I also verified the scalar None behavior. The latest PyTorch docs describe scalar None, but the current runner's torch 2.10.0+cu129 still rejects F.adaptive_avg_pool2d(x, None) / F.adaptive_max_pool2d(x, None). I am not blocking on that here if the intended TileOps contract is to follow the newer documented semantics, but the follow-up correctness/reference path should account for that runtime mismatch.

[RMLYC]

@superAngGao Confirmed and fixed in 62c1a8b.

Reproduced the codegen failure locally with your snippet (roofline.vars['out_H'] references unknown name 'H_in'). Worth noting: the same failure reproduces on the existing AvgPool2dFwdOp / MaxPool2dIndicesFwdOp entries, so the unsynthesizeable-vars pattern is pre-existing family debt, not introduced by this PR — but the new entries should not propagate it.

Fix: all three adaptive entries now bind the shape dims explicitly at the top of roofline.vars:

N: "input.shape[0]"
C: "input.shape[1]"
H_in: "input.shape[2]"
W_in: "input.shape[3]"

Verification (tileops-dev, torch 2.10.0+cu129): synthesize_eval_roofline() succeeds for all three entries, and evaluating the synthesized eval_roofline over every declared workload reproduces the exact-scan numbers — nondiv-7x7 avg → (491904, 815104), matching your runner-image result. python scripts/validate_manifest.py still passes, and #1809 is updated in lockstep (its hand-written eval_roofline already computes the same exact-scan values, so manifest and implementation now agree at both layers).

On the scalar-None note: agreed — that's the trade we made. The manifest follows the 2.13 documented contract (int | None | tuple[int | None, int | None]), while the Op documents that the 2.10 runtime rejects scalar None and its tests map the PyTorch reference to (None, None) until the runtime catches up (see the follow-up in @zhen8838's thread).

[superAngGao]

Re-reviewed latest head 62c1a8b. The previous inline roofline codegen blocker is fixed: all three adaptive pool entries now bind N/C/H_in/W_in explicitly, synthesize_eval_roofline() succeeds, and the exact adaptive-bin scan matches direct enumeration for the declared workloads including nondiv-7x7. The spec-only manifest shape and bench_manifest_driven=false usage are consistent with the current repository conventions. Looks good to me.

[Ibuki-wind]

Overall

Approval is blocked by stale PR metadata and review-reply cleanup required by the repository approval gate.

Cross-cutting concerns

• Rewrite the PR body to record the current final state: it still says the roofline uses one ceil-divided bin size and must be replaced later, but head 62c1a8b already uses the exact adaptive-bin scan. Keep the body concise and add the required pre-commit and test-node-delta verification facts.
• Reduce the four author replies at issue comments 5126310781, 5126311041, 5127183876, and 5127907638 to outcome-only forms such as Done in <sha>. or Won’t-fix: <one-line reason>.; the current root-cause and design restatements violate the approval-gate reply discipline.

Withdrawn so the PR metadata can be cleaned up directly.

[Ibuki-wind]

Clean — no issues.