Skip to content

[CK Tile] Add transposed tile load implementation, and tests for load_and_convert_tile - #5510

Merged
SamiAario-AMD merged 60 commits into
developfrom
users/samaario/ck/test-load-tile-transpose
Jun 15, 2026
Merged

[CK Tile] Add transposed tile load implementation, and tests for load_and_convert_tile#5510
SamiAario-AMD merged 60 commits into
developfrom
users/samaario/ck/test-load-tile-transpose

Conversation

@SamiAario-AMD

@SamiAario-AMD SamiAario-AMD commented Mar 17, 2026

Copy link
Copy Markdown
Contributor

Motivation

Mixed precision b/fp16 x fp8 requires a transposed tile load implementation that supports mixed precision using these types. Implement this, use it in load_and_convert_tile, and add a unit test for load_and_convert_tile which covers this functionality.

Technical Details

Test Plan

Test Result

Submission Checklist

@SamiAario-AMD
SamiAario-AMD requested a review from a team as a code owner March 17, 2026 11:47
@SamiAario-AMD
SamiAario-AMD marked this pull request as draft March 17, 2026 11:47
@SamiAario-AMD
SamiAario-AMD force-pushed the users/samaario/ck/test-load-tile-transpose branch from b381791 to 459d1d9 Compare March 17, 2026 12:00
@SamiAario-AMD
SamiAario-AMD force-pushed the users/samaario/ck/test-load-tile-transpose branch from 459d1d9 to 6694ced Compare March 17, 2026 13:20
@SamiAario-AMD
SamiAario-AMD marked this pull request as ready for review March 17, 2026 13:33
@SamiAario-AMD SamiAario-AMD changed the title Users/samaario/ck/test load tile transpose [CK Tile] test load_and_convert_tile Mar 18, 2026
@SamiAario-AMD
SamiAario-AMD force-pushed the users/samaario/ck/test-load-tile-transpose branch 6 times, most recently from 236c248 to 713dd57 Compare March 19, 2026 13:54
@SamiAario-AMD
SamiAario-AMD force-pushed the users/samaario/ck/test-load-tile-transpose branch from 713dd57 to a6309fa Compare March 20, 2026 06:49
@SamiAario-AMD SamiAario-AMD changed the title [CK Tile] test load_and_convert_tile [CK Tile] Add transposed tile load implementation, and tests for load_and_convert_tile Mar 23, 2026

@jakpiase jakpiase left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

I have left one minor request, apart from that it does look great.

Comment thread projects/composablekernel/include/ck_tile/host/fill.hpp Outdated
Comment thread projects/composablekernel/test/ck_tile/load_and_convert_tile/CMakeLists.txt Outdated

@aosewski aosewski left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Re-reviewed the changes since my 2026-04-27 round, focused on whether the blocking findings landed. They did:

  • Layering inversion — fixed. load_tile_transpose.hpp no longer pulls in the ops/elementwise header; the converter is injected as a PassThroughPack_ template arg, with PassThroughPackSelector living in the ops/common layer. Core stays op-free.
  • print_matrix.hpp — fixed. Moved under test/, #pragma once, wrapped in namespace ck_tile, real lens[0]/lens[1], ASCII only, and out of the public host.hpp aggregate.
  • LDS WAR hazard — fixed. block_sync_lds() now sits between the LDS read and the next iteration's write.
  • gfx950 guard — your rebuttal holds: amd_transpose_load_to_vgpr stays behind #if defined(__gfx950__) in buffer_view.hpp, the convert path is CMake-gated to gfx950, and reverting the static_assert makes sense given the new gfx1250 branch.
  • Transpose test coverage — on a closer read I was wrong here. The test is an identity round-trip through a transposed LDS view, and the fill is independent-random (not symmetric), so a broken transpose distribution would scramble values and fail the check. It does exercise the layout.

Warnings are largely handled too (vectorized non-transpose path, gfx950 down-cast builtins, integer_divide_ceil tail handling, PassThroughPackSelector generalization, target gating + rationale). What's left is minor/stylistic — dead LoadTranspose arg in the pk branch, inner Problem shadowing, single test shape — fine to leave or fold in later.

LGTM, approving.

@aosewski
aosewski dismissed cgmillette’s stale review June 9, 2026 13:35

Comments addressed.

SamiAario-AMD and others added 2 commits June 9, 2026 16:44
Use row-major (N, 1) strides that match HostTensor for global and LDS views.

Co-authored-by: Cursor <cursoragent@cursor.com>
@SamiAario-AMD
SamiAario-AMD force-pushed the users/samaario/ck/test-load-tile-transpose branch from 19186dd to 1521356 Compare June 10, 2026 14:51
SamiAario-AMD and others added 3 commits June 11, 2026 13:26
Always launch kBlockSize threads so the tile distribution matches blockDim on wave32,
and compile the gfx11 and gfx12 test targets with CK_TILE_USE_LAUNCH_BOUNDS=0 to
avoid HIP launch failures from default launch bounds on large blocks.

Co-authored-by: Cursor <cursoragent@cursor.com>
@SamiAario-AMD
SamiAario-AMD force-pushed the users/samaario/ck/test-load-tile-transpose branch from 68d0e5a to 523db71 Compare June 12, 2026 12:45
@SamiAario-AMD
SamiAario-AMD merged commit 8415c8c into develop Jun 15, 2026
40 checks passed
@SamiAario-AMD
SamiAario-AMD deleted the users/samaario/ck/test-load-tile-transpose branch June 15, 2026 06:41
assistant-librarian Bot pushed a commit to ROCm/composable_kernel that referenced this pull request Jun 15, 2026
[CK Tile] Add transposed tile load implementation, and tests
 for load_and_convert_tile (#5510)

## Motivation

Mixed precision b/fp16 x fp8 requires a transposed tile load
implementation that supports mixed precision using these types.
Implement this, use it in `load_and_convert_tile`, and add a unit test
for `load_and_convert_tile` which covers this functionality.

## Technical Details

<!-- Explain the changes along with any relevant GitHub links. -->

## Test Plan

<!-- Explain any relevant testing done to verify this PR. -->

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
samremes added a commit that referenced this pull request Jun 15, 2026
Merge develop (a2d2d13; +84 commits incl MX GEMM correctness PR #6663 and
transposed tile-load PR #5510). Two conflicts in the MX test suite, both
"independent fixtures added on each side" -- resolved by keeping both:
- test_mx_gemm_config.hpp: ours = 32x32x64 transpose configs (MxGemmConfig32,
  MXfp4/8_GemmConfig32); develop = MXfp8_GemmConfig16_PadMN.
- test_mx_gemm_async.cpp: ours = TestMxGemmTranspose (8 layout cases);
  develop = preshuffle split-K + fp8/fp4 regression + PadMN fixtures (#6663).
The shared Run() now takes k_batch=1 default (develop), compatible with our
3-arg transpose calls. Verified after auto-merge: dispatcher 32x32x64 f8f6f4
spec, kDefaultScale unity, and load_tile_transpose all intact.

Co-authored-by: Cursor <cursoragent@cursor.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants