ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-12 09:16:52 +00:00
Code Issues Packages Projects Releases Wiki Activity

Move reference to host/reference

Browse Source
This commit is contained in:
Matti Eskelinen
2026-01-26 09:18:52 +00:00
parent 04f8f3ed5d
commit e9485e0ecb
3 changed files with 82 additions and 68 deletions
Download Patch File Download Diff File Expand all files Collapse all files

68
test/ck_tile/sinkhorn/test_sinkhorn_impl.hpp
View File

@@ -28,74 +28,6 @@ class TestCkTileSinkHorn : public ::testing::Test
using TestSinkhornShape =
ck_tile::SinkhornKnoppShape<BlockWarps_, BlockTile_, WarpTile_, ThreadTile_>;
// template <typename XDataType, typename ComputeDataType, typename YDataType>
void sinkhorn_knopp_ref_single_iter(ck_tile::HostTensor<ComputeDataType>& c_n_n,
ck_tile::HostTensor<ComputeDataType>& acc_n)
{
const ck_tile::index_t input_n = acc_n.get_length(0);
// Sum and scale rowwise
for(ck_tile::index_t i = 0; i < input_n; ++i)
{
acc_n(i) = 0;
for(ck_tile::index_t j = 0; j < input_n; ++j)
{
acc_n(i) += c_n_n(i, j);
}
for(ck_tile::index_t j = 0; j < input_n; ++j)
{
c_n_n(i, j) /= acc_n(i);
}
}
// Repeat columnwise
for(ck_tile::index_t i = 0; i < input_n; ++i)
{
acc_n(i) = 0;
for(ck_tile::index_t j = 0; j < input_n; ++j)
{
acc_n(i) += c_n_n(j, i);
}
for(ck_tile::index_t j = 0; j < input_n; ++j)
{
c_n_n(j, i) /= acc_n(i);
}
}
}
void sinkhorn_knopp_ref(const ck_tile::HostTensor<XDataType>& x_n_n,
ck_tile::HostTensor<YDataType>& y_n_n,
const int n_iter)
{
const ck_tile::index_t input_n = x_n_n.get_length(0);
ck_tile::HostTensor<ComputeDataType> c_n_n({input_n, input_n}, {1, input_n});
ck_tile::HostTensor<ComputeDataType> acc_n({input_n}, {1});
// First apply exp to make input nonnegative
for(ck_tile::index_t i = 0; i < input_n; ++i)
{
for(ck_tile::index_t j = 0; j < input_n; ++j)
{
c_n_n(i, j) = exp(ck_tile::type_convert<ComputeDataType>(x_n_n(i, j)));
}
}
// Iterate normalization on rows and columns
for(auto it = 0; it < n_iter; ++it)
{
sinkhorn_knopp_ref_single_iter(c_n_n, c_n_n);
}
// Copy and cast to output type
for(ck_tile::index_t i = 0; i < input_n; ++i)
{
for(ck_tile::index_t j = 0; j < input_n; ++j)
{
y_n_n(i, j) = ck_tile::type_convert<YDataType>(c_n_n(i, j));
}
}
}
void RunGenericTest(const std::vector<ck_tile::index_t>& input_shape, const int max_iterations)
{
auto input_n = input_shape[0];