Skip to content

[CK_TILE] Fix CShuffleEpilogue test to use correct GEMM accumulator distribution - #4260

Closed
assistant-librarian[bot] wants to merge 9 commits into
developfrom
import/develop/ROCm_composable_kernel/pr-3701
Closed

[CK_TILE] Fix CShuffleEpilogue test to use correct GEMM accumulator distribution#4260
assistant-librarian[bot] wants to merge 9 commits into
developfrom
import/develop/ROCm_composable_kernel/pr-3701

Conversation

@assistant-librarian

@assistant-librarian assistant-librarian Bot commented Feb 3, 2026

Copy link
Copy Markdown
Contributor

Summary

The test was using LDS distribution to create the accumulator tile, but CShuffleEpilogue expects the GEMM accumulator distribution that BlockGemm produces. This mismatch caused incorrect data permutation.

Changes

  • Use WarpGemmDispatcher to get correct accumulator distribution encoding
  • Load test input from host-initialized global memory for deterministic verification
  • Shard tests by data type (FP16, FP8) with gfx950-specific FP8 tests
  • Extract scale tests into separate target for better organization
  • Implement exact permutation verification (all unique values appear once)
  • Reduce tile size from 256x256 to 128x128 to fit in unique fp16 range
  • Add parameterized test configurations for various warp layouts and MFMA types

Test plan

  • Run new cshuffle epilogue tests

🤖 Generated with Claude Code

Co-Authored-By: Claude noreply@anthropic.com

tenpercent and others added 2 commits February 3, 2026 16:30
…and improved verification

This commit refactors and expands the CShuffleEpilogue test suite with the following improvements:

- Shard tests by data type (FP16, FP8) and add gfx950-specific FP8 tests
- Extract scale tests into a separate target for better organization
- Implement robust verification using set-based comparison to ensure all expected values are present
- Add distribution uniformity checks to verify proper data shuffling
- Strengthen verification to require kBlockSize unique values and validate all rows independently
- Add support for parameterized test configurations covering various warp layouts and MFMA types
- Improve code organization by separating concerns and extracting common utilities
- Fix memory leaks and improve code hygiene
- Add CMake support for OCP FP8 compilation flag

The refactored test suite provides better coverage, more rigorous verification, and improved
maintainability for the CShuffleEpilogue functionality.

Co-Authored-By: Claude <noreply@anthropic.com>
@assistant-librarian assistant-librarian Bot added the external contribution Code contribution from users community.. label Feb 3, 2026
@ThomasNing

Copy link
Copy Markdown

Comment:
VGPR and AGPR are basically the same when we are loading out to the dram. In this test, we should still initialize it in FP32 in the host memory and load it to the VGPR and then cast the epilogue operation. For 32x32 and 16x16 matrix, we should have the same result of output as the host input. In this way, we could call the test passed.

@tenpercent
tenpercent force-pushed the import/develop/ROCm_composable_kernel/pr-3701 branch from 3a5dfae to 4b3e611 Compare February 4, 2026 20:52
@tenpercent tenpercent changed the title [CK-Tile] Improve cshuffle epilogue mfma tile coverage [CK_TILE] Fix CShuffleEpilogue test to use correct GEMM accumulator distribution Feb 4, 2026
The test was using LDS distribution to create the accumulator tile, but
CShuffleEpilogue expects the GEMM accumulator distribution that BlockGemm
produces. This mismatch caused incorrect data permutation.

Changes:
- Use WarpGemmDispatcher to get correct accumulator distribution encoding
- Load test input from host-initialized global memory instead of generating
  values per-thread in kernel (ensures deterministic, verifiable input)
- Generate unique fp16 bit patterns for exact permutation verification
- Simplify convert_fp16_bits to return float directly (was misleading template)
- Update verification to check exact permutation (all values appear once)
- Reduce tile size from 256x256 to 128x128 to fit in unique fp16 range
- Use HostTensor consistently for input/output management
- Add file organization documentation to test_cshuffle_epilogue_common.hpp

Co-Authored-By: Claude <noreply@anthropic.com>
@tenpercent
tenpercent force-pushed the import/develop/ROCm_composable_kernel/pr-3701 branch from 4b3e611 to 552568d Compare February 4, 2026 21:28
@tenpercent
tenpercent marked this pull request as ready for review February 5, 2026 01:31
@tenpercent
tenpercent requested a review from a team as a code owner February 5, 2026 01:31
@ThomasNing
ThomasNing enabled auto-merge (squash) February 9, 2026 19:36
@ThomasNing

Copy link
Copy Markdown

@tenpercent Clang format error.

tenpercent and others added 4 commits February 10, 2026 10:48
Run clang-format-18 on modified epilogue test files to align with project
coding standards. Changes include alignment of variable assignments and
wrapping of long lines/comments.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
@tenpercent

Copy link
Copy Markdown
Contributor

Since Math CI never runs in this PR copied it over to #4518

@tenpercent tenpercent closed this Feb 11, 2026
auto-merge was automatically disabled February 11, 2026 19:12

Pull request was closed

@tenpercent
tenpercent deleted the import/develop/ROCm_composable_kernel/pr-3701 branch March 6, 2026 23:29
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants