Skip to content

Refactor subtile to TileInfo architecture with modular emit files#39

Draft
b-shi wants to merge 3 commits into
subtile_mxfrom
save/subtile_mx_rebase_v1
Draft

Refactor subtile to TileInfo architecture with modular emit files#39
b-shi wants to merge 3 commits into
subtile_mxfrom
save/subtile_mx_rebase_v1

Conversation

@b-shi

@b-shi b-shi commented Apr 28, 2026

Copy link
Copy Markdown
Owner
  • Replace monolithic SubtileBasedKernel.py with modular Subtile/ package
  • Split emit logic into SubtileGREmit.py, SubtileLREmit.py, SubtileScaleEmit.py, SubtileGeometry.py
  • Introduce TileInfo_ frozen geometry class replacing legacy TileInfo
  • Integrate LogicalScheduler from subtile_mx with TileInfo_ architecture
  • Method renames: grSubtileRegIdx -> grRegGroupForSubtileRow, grLoadIndex -> grLoadIndexForSubtile, lrTileIndex -> lrTileIndexForSubtile
  • Add lrSubtileShape, lrLocalSubtileGrid, loadRatioLR fields
  • Update loadRatioGR formula to use subtileSize * blockCount
  • Remove legacy ScheduleInfo, rename allocVgprTileRegisters -> _legacy
  • Add test_tileinfo_compare, test_tileinfo_fields, test_gr_offset

@b-shi b-shi marked this pull request as draft April 28, 2026 04:47
Comment thread projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileGeometry.py Outdated
@b-shi b-shi force-pushed the save/subtile_mx_rebase_v1 branch from 9d3d904 to 51dedf8 Compare April 28, 2026 15:38
- Replace monolithic SubtileBasedKernel.py with modular Subtile/ package
- Split emit logic into SubtileGREmit.py, SubtileLREmit.py, SubtileScaleEmit.py, SubtileGeometry.py
- Introduce TileInfo frozen geometry class replacing legacy TileInfo
- Integrate LogicalScheduler from subtile_mx with TileInfo architecture
- Method renames: grSubtileRegIdx -> grRegGroupForSubtileRow, grLoadIndex -> grLoadIndexForSubtile, lrTileIndex -> lrTileIndexForSubtile
- Add lrSubtileShape, lrLocalSubtileGrid, loadRatioLR fields
- Update loadRatioGR formula to use subtileSize * blockCount
- Remove legacy ScheduleInfo, rename allocVgprTileRegisters -> _legacy
- Add test_tileinfo_compare, test_tileinfo_fields, test_gr_offset

Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>
@b-shi b-shi force-pushed the save/subtile_mx_rebase_v1 branch from 51dedf8 to 92be082 Compare April 29, 2026 18:37
@b-shi b-shi requested review from msujon-AMD and sebvince April 30, 2026 13:28

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should we rename it to InstructionEmitter as well ? (now it's in subtile dir)


# Row-major A/B: GR and LR both contiguous along K
AB_B16 = ABTilePair(
gr=ABGRGeometry(tag=GRTag_1x2(), **_B16, blockShape=(1, 2), loadShape=LoadShape(m=1, k=8)), # 128-bit GR: 8 bf16 along K

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

aren't blockShape and subtileShape the same thing ? I mean why do we need 2 different names ?

2x2 WG: A=(1,2), B=(1,2) — 2 cooperating waves each
1x1 WG: A=(1,2), B=(1,2) — 1 wave each

LR subtile shape is always (1,2) regardless of wave config.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually we currently use (1,1) in the scheduler. This is for current PGR0 codepath I guess ?

@amd-chunxlin amd-chunxlin left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Small suggestion: may consider the copyright below to avoid yearly updates:

// Copyright Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier:  MIT

@memmett

memmett commented Apr 30, 2026

Copy link
Copy Markdown

This looks good to me. I don't have any detailed reviews for now.

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.

5 participants