ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 11:16:59 +00:00
Code Issues Packages Projects Releases Wiki Activity

can pass; but no logic

Browse Source
This commit is contained in:
joye
2025-04-24 19:05:47 -05:00
parent 90a4501869
commit df9769afba
2 changed files with 13 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
include/ck_tile/ops/transpose/block_transpose.hpp
View File

@@ -8,6 +8,9 @@
namespace ck_tile {
template <typename T>
struct Debug;
template <typename Layout_, index_t kRow, index_t kCol>
struct TransposeTraits
{
@@ -118,7 +121,7 @@ struct BlockTranspose
auto input_tile_window =
make_tile_window(input_window, Policy::template MakeInputDistribution<Problem>());
auto output_tile_window =
make_tile_window(out_window, Policy::template MakeOutputDistribution<Problem>());
make_tile_window(out_window, Policy::template MakeLdsLoadTileDistribution<Problem>());
DataType* p_lds_ptr = static_cast<DataType*>(p_smem);
constexpr auto in_lds_block_desc = Policy::template MakeLdsStoreBlockDescriptor<Problem>();
@@ -145,13 +148,15 @@ struct BlockTranspose
block_sync_lds();
auto y = load_tile(load_from_lds_window);
//Debug<remove_cvref_t<decltype(y)>> cccc;
// auto load_from_lds_window =
// make_tile_window(output_lds_block,
// make_tuple(number<kSecondSizePerBlock>{},
// number<kLeadSizePerBlock>{}), {0, 0}, Policy::template
// MakeLdsLoadTileDistribution<Problem>());
// auto y = load_tile_transpose(load_from_lds_window);
//auto y = load_tile_transpose(load_from_lds_window);
store_tile(output_tile_window, y);
}
};

9
include/ck_tile/ops/transpose/transpose_policy.hpp
View File

@@ -157,6 +157,7 @@ struct TransposePolicy
kSecond / QuartTransposeTraits<typename Problem::DataType>::ksecondDim;
using xdllevel_dstr_encoding = typename QuartTransposeTraits<
typename Problem::DataType>::template TileDistribution<kSecondDimstr, kLeadDimstr>;
constexpr index_t kLeadIterPerWarp = Problem::kLeadXdlNumPerWarp;
constexpr index_t kSecondIterPerWarp = Problem::kSecondXdlNumPerWarp;
constexpr index_t kLeadNumWarps = Problem::kLeadNumWarps;
@@ -165,12 +166,14 @@ struct TransposePolicy
tile_distribution_encoding<sequence<>,
tuple<sequence<kSecondIterPerWarp, kSecondNumWarps>,
sequence<kLeadIterPerWarp, kLeadNumWarps>>,
tuple<sequence<1>, sequence<2>>,
tuple<sequence<1>, sequence<1>>,
tuple<sequence<1, 2>>,
tuple<sequence<1, 1>>,
sequence<1, 2>,
sequence<0, 0>>{};
return detail::make_embed_tile_distribution_encoding(block_outer_dst_encoding,
constexpr auto blk_distr_encode = detail::make_embed_tile_distribution_encoding(block_outer_dst_encoding,
xdllevel_dstr_encoding{});
constexpr auto block_dstr = make_static_tile_distribution(blk_distr_encode);
return block_dstr;
}
};