ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-01 12:17:00 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add compile check for assumed row-mjor layout.

Browse Source
This commit is contained in:
Ville Pietilä
2025-09-26 08:39:39 +00:00
parent 1764c77fb2
commit 354dd5039c
1 changed files with 8 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp
View File

@@ -219,6 +219,13 @@ struct CShuffleEpilogue
template <typename Problem>
CK_TILE_HOST_DEVICE static constexpr auto MakeLdsBlockDescriptor()
{
if constexpr(NumGroupsToMerge > 1)
{
// We haven't yet tested the ColumnMajor case.
static_assert(std::is_same_v<ELayout, tensor_layout::gemm::RowMajor>,
"Currently, the CShuffle Epilogue with NumGroupsToMerge > 1 only supports the Row Major layout");
}
// N is contiguous dimension
if constexpr(std::is_same_v<ELayout, tensor_layout::gemm::RowMajor>)
{
@@ -366,7 +373,7 @@ struct CShuffleEpilogue
sequence<0, 3, 0, 3>>{};
constexpr auto dram_tile_distribution = make_static_tile_distribution(dram_tile_encoding);
// The LDS data has the following 4D layout
// The LDS data has the following 4D layout in the row-major case.
// linear_index = c + Gs * m + Gs * MPerGroup * n + Gs * MPerGroup * NPerGroup * r
// for 4D coordinates (r,c,m,n) where (r,c) is the group index and (m,n) is the index within the group.
// We pick-up only the diagonal blocks where r == c.