Skip to content

[Enhancement] Support region padding when convert buffer load to buffer region - #342

Merged
LeiWang1999 merged 2 commits into
tile-ai:mainfrom
LeiWang1999:fix_20250406
Apr 6, 2025
Merged

LeiWang1999 merged 2 commits into
tile-ai:mainfrom
LeiWang1999:fix_20250406

Conversation

@LeiWang1999

Copy link
Copy Markdown
Member

Enhance error checking in RegionOp and buffer_load_to_tile_region

- Added detailed error messages to the index size check in `RegionOp` to aid debugging.
- Implemented a check in `buffer_load_to_tile_region` to ensure the length of indices matches extents, with a fallback to expand extents if necessary. This improves robustness in handling buffer loads with mismatched dimensions.

- Added detailed error messages to the index size check in `RegionOp` to aid debugging.
- Implemented a check in `buffer_load_to_tile_region` to ensure the length of indices matches extents, with a fallback to expand extents if necessary. This improves robustness in handling buffer loads with mismatched dimensions.
@LeiWang1999
LeiWang1999 merged commit b0bde55 into tile-ai:main Apr 6, 2025
LeiWang1999 added a commit to LeiWang1999/tilelang that referenced this pull request Jul 18, 2025
…er region (tile-ai#342)

* Enhance error checking in RegionOp and buffer_load_to_tile_region

- Added detailed error messages to the index size check in `RegionOp` to aid debugging.
- Implemented a check in `buffer_load_to_tile_region` to ensure the length of indices matches extents, with a fallback to expand extents if necessary. This improves robustness in handling buffer loads with mismatched dimensions.

* lint fix
LeiWang1999 added a commit to LeiWang1999/tilelang that referenced this pull request Jul 20, 2025
…er region (tile-ai#342)

* Enhance error checking in RegionOp and buffer_load_to_tile_region

- Added detailed error messages to the index size check in `RegionOp` to aid debugging.
- Implemented a check in `buffer_load_to_tile_region` to ensure the length of indices matches extents, with a fallback to expand extents if necessary. This improves robustness in handling buffer loads with mismatched dimensions.

* lint fix
jimmychou0 added a commit to PTO-ISA/tilelang that referenced this pull request Sep 30, 2026
* feat(ascend): migrate mixed GEMM PTO codegen

* fix(ascend): rebind local vars in PTO sections

* fix(ascend): correct mixed GEMM PTO codegen

* feat(pto): remove ssa value workaround

* fix(pto): add back kernel_kind for pto.jit. default is vector, section set None

* feat(pto): GEMM support dual_copy(mixed=true)

* feat(ascend): add PTO blockscaled GEMM codegen

* feat(ascend): lower L0 blockscaled GEMM to PTO

* test(ascend): harden PTO blockscaled GEMM coverage

* refactor(ascend): centralize blockscaled MAD flag handling

* feat(ascend): add FP4 blockscaled GEMM PTO support

* feat(ascend): preserve explicit PTO L1-to-L0 controls

* fix(pto): lower explicit T.set_atomic in PTO codegen

Wire VisitExpr_ for tl.ascend_set_atomic / set_atomic_none on target=pto
so kernels can arm/disarm store-atomic state without annotation-based
T.copy(..., atomic_op=...). AscendC path already existed; no store-atomic
copy lowering changes.

* fix(example_gemm): include initial c0 in atomic-accumulate reference

For acc=True the kernel computes C = c0 + A@B^T by accumulating into the
pre-initialized output via store-atomic. The checking code compared against
ref_program(x, w) which omitted c0, so max_diff equaled the magnitude of the
random c0 (~5-6) and the case falsely FAILed. Pass c=c0 to the reference.
It is a no-op for acc=False (c0 is None).

* feat(ascend): GEMM L0 PTO codegen support

* fix(ascend): allow float16 in PTO GEMM

* feat(example): add target option for annotate unlimit memory

* feat(ascend): support explicit FP4 PTO L1-to-L0 loads

* style(ascend): format PTO mixed-kernel changes

* test(ascend): unify Ascend and PTO GEMM coverage

Add explicit target selection to GEMM examples and performance helpers. Parameterize auto, manual mixed, and blockscaled GEMM tests across the Ascend and PTO backends, including mixed, HF32, atomic, FP8, and FP4 paths.

* fix(pto): address mixed GEMM codegen review findings

* fix(pto): validate mixed section topology

* fix(pto): reject unpaired physical sections

* fix(pto): validate mixed kernel AIV count

---------

Co-authored-by: jimmychou <47636600+jimmychou0@users.noreply.github.com>
Co-authored-by: caojian5 <caojian5@huawei.com>
Co-authored-by: andodo <andodo@Mac.local>
Co-authored-by: andodo <andodo@andododeMacBook-Air.local>
(cherry picked from commit a0ec76e11d68b8c18a345b6161642e9a78fdd291)
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant