[BugFix] Bound a sliced fill by the end of the region, not just its extent - #3352

Refs #2979.

Why

The static bound check in the Fill constructor validates min >= 0 and extent <= shape per dimension. extent <= shape only bounds a region anchored at the start of the buffer, so a slice that starts further in passes both checks and is lowered to an out-of-bounds store:

As = T.alloc_shared((8,), "int32")
T.fill(As[6:12], 7)   # min=6 >= 0, extent=6 <= 8, but 6 + 6 = 12 > 8

T.fill(As[6:12], 7)    compiled, emitted As[(((int)threadIdx.x) + 6)] = 7   # writes 6..11
T.fill(As[0:12], 7)    rejected: extent_imm->value <= shape_imm->value (12 vs. 8)
T.fill(As[2:8], 7)     compiled, in bounds

T.clear shares the same node and had the same gap: T.clear(As[6:12]) compiled and emitted the same store.

What changed

Two commits.

The end of the region is checked against the destination shape as well, which is what has to fit. The new term needs both the start and the extent to be static, so a region anchored at zero and one with a symbolic start are both unaffected.

Scope: the report also names T.copy, which accepts the same slice — T.copy(A[0:6], As[6:12]) compiles and writes past the end. Its regions are not validated in this node at all, so it needs a check where they are built rather than another term in this one; that is left out here rather than guessed at.

Verification

• Four cases added to testing/python/language/test_tilelang_language_clear.py, which already covers both T.fill and T.clear: each spelling against a slice that runs past the end and against one that fits. The two out-of-bounds cases fail on an unmodified tree with DID NOT RAISE Exception; the two in-bounds controls pass in both states.
• test_tilelang_language_clear.py, test_tilelang_transform_lower_tile_op.py, test_tilelang_language_transpose.py, test_tilelang_language_reshape.py: 48 passed.
• testing/python/transform/: 669 passed, 41 skipped, 2 deselected. The two deselected are pre-existing environment limits — test_tiled_ws_multi_versioned_wide_linear_smem_keeps_tma and test_tiled_ws_wide_linear_smem_three_stages fail on Failed to set the allowed dynamic shared memory size to 131072 / 147456.
• The four new cases are host-only: validation happens while lowering, which needs a target context but no device, so they pass with CUDA_VISIBLE_DEVICES unset.
• The kernel cache was cleared before each run.
• pre-commit run --files <changed files>: all 13 hooks pass.

A second commit makes the three checks in that block report a frontend validation failure rather than an internal invariant, since the region comes from the TileLang program under compilation. support/check.h is already included and src/op/reduce.cc uses CHECK(..., ValueError) for the same kind of condition. An out-of-range slice now fails as ValueError: Check failed: (min_imm->value + extent_imm->value <= ... instead of an InternalError, which is what the style guide asks for:

Use user-facing validation errors where the problem can be caused by a TileLang program rather than an internal invariant. Keep internal invariant failures distinguishable from frontend validation failures.

The second commit's regression run: 737 passed, 42 skipped, 2 deselected.

Summary

• Fill now checks that a statically known region’s start plus extent does not exceed the destination shape. This catches out-of-bounds slices such as As[6:12] on a buffer of shape 8.
• The check applies when the start, extent, and destination shape are static. T.copy remains out of scope because its regions require validation elsewhere.
• Added tests for in-bounds and out-of-bounds slices with both T.fill and T.clear.

Validation

The author reports four new cases passing, 48 passing tests across four named test files, and 669 passed, 41 skipped, and 2 deselected in testing/python/transform/. The two deselected tests failed while setting the allowed dynamic shared memory size. The author also reports that all 13 pre-commit hooks passed. No independent test run or current review findings were provided.

C++ style / lint notes

• The PR changes C++ code and touches checks covered by docs/developer_guide/cpp_style.md, including descriptive local names and TVM handle usage. No exported or public declarations changed.
• CI includes the “C++ API Style Audit (warning only)” step. No audit findings for this change were supplied, so warning status is unavailable. Style warnings are separate from correctness, build, and test results.

Verified end to end, together with the other contributions from this batch: VERIFICATION.md