ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-11 08:50:17 +00:00
Code Issues Packages Projects Releases Wiki Activity

Get the test to compile

Browse Source
This commit is contained in:
Damien Lejeune
2026-01-21 04:57:01 -05:00
parent e3efa236ec
commit 557a8d3f21
3 changed files with 28 additions and 19 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
include/ck_tile/ops/sinkhorn_knopp/pipeline/sinkhorn_knopp_problem.hpp
View File

@@ -29,7 +29,7 @@ template <typename BlockWarps, // num warps along seq<M, N>
typename BlockTile, // block size, seq<M, N>
typename WarpTile, // warp size, seq<M, N>
typename ThreadTile> // contiguous pixels(vector size) along seq<M, N>
struct SinkHornKnoppShape
struct SinkhornKnoppShape
{
static constexpr index_t Block_M = BlockTile::at(number<0>{});
static constexpr index_t Block_N = BlockTile::at(number<1>{});

22
include/ck_tile/ops/sinkhorn_knopp/sinkhorn_knopp_kernel.hpp
View File

@@ -53,6 +53,14 @@ struct SinkhornKnoppKernelReduce
template <typename Problem, typename Policy>
struct SinkhornKnoppKernelDummyNonStochastic
{
static constexpr index_t kBlockSize = Problem::BlockShape::BlockSize;
CK_TILE_HOST static constexpr auto BlockSize()
{
return is_wave32() ? kBlockSize / 2 : kBlockSize;
}
template <typename XDistributedTensor_>
CK_TILE_DEVICE static auto MakeYBlockTile()
{
@@ -67,16 +75,16 @@ struct SinkhornKnoppKernelDummyNonStochastic
return tensor;
}
CK_TILE_DEVICE void operator()(const SinkhornKnoppArgs& args) const
CK_TILE_DEVICE void operator()([[maybe_unused]]const SinkhornKnoppArgs& args) const
{
using S = Problem::BlockShape;
// using S = Problem::BlockShape;
static_assert(S::Block_M == S::Block_N, "Input must be a square matrix!");
// static_assert(S::Block_M == S::Block_N, "Input must be a square matrix!");
const auto x_desc = make_naive_tensor_descriptor(make_tuple(args.input_m, args.input_m),
make_tuple(args.input_m, 1),
number<4>{}, // TODO: Hardcoded vectorization, we should calculate it!
number<1>{});
// const auto x_desc = make_naive_tensor_descriptor(make_tuple(args.input_m, args.input_m),
// make_tuple(args.input_m, 1),
// number<4>{}, // TODO: Hardcoded vectorization, we should calculate it!
// number<1>{});
// auto buffer_view = make_buffer_view<address_space_enum::global>(
// args.p_x, desc.get_element_space_size(), number<0>{});