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

Add CPU reference computation

Browse Source
This commit is contained in:
Matti Eskelinen
2026-01-26 09:14:16 +00:00
parent 5f1d79b42d
commit 04f8f3ed5d
1 changed files with 109 additions and 46 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -14,30 +14,92 @@
#include "ck_tile/host/kernel_launch.hpp"
template <typename Tuple>
class TestCkTileSinkHorn: public ::testing::Test
class TestCkTileSinkHorn : public ::testing::Test
{
protected:
using XDataType = std::tuple_element_t<0, Tuple>;
using ComputeDataType = std::tuple_element_t<1, Tuple>;
using YDataType = std::tuple_element_t<2, Tuple>;
using BlockWarps_ = std::tuple_element_t<3, Tuple>;
using BlockTile_ = std::tuple_element_t<4, Tuple>;
using WarpTile_ = std::tuple_element_t<5, Tuple>;
using ThreadTile_ = std::tuple_element_t<6, Tuple>;
using XDataType = std::tuple_element_t<0, Tuple>;
using ComputeDataType = std::tuple_element_t<1, Tuple>;
using YDataType = std::tuple_element_t<2, Tuple>;
using BlockWarps_ = std::tuple_element_t<3, Tuple>;
using BlockTile_ = std::tuple_element_t<4, Tuple>;
using WarpTile_ = std::tuple_element_t<5, Tuple>;
using ThreadTile_ = std::tuple_element_t<6, Tuple>;
using TestSinkhornShape =
ck_tile::SinkhornKnoppShape<
BlockWarps_,
BlockTile_,
WarpTile_,
ThreadTile_
>;
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_m = input_shape[0];
auto default_stride = {input_m, 1};
auto input_n = input_shape[0];
auto default_stride = {input_n, 1};
ck_tile::HostTensor<XDataType> h_x(input_shape, default_stride);
ck_tile::HostTensor<YDataType> h_y(input_shape, default_stride);
@@ -48,24 +110,19 @@ class TestCkTileSinkHorn: public ::testing::Test
ck_tile::DeviceMem d_x_mem(h_x.get_element_space_size_in_bytes());
ck_tile::DeviceMem d_y_mem(buffer_size);
ck_tile::SinkhornKnoppArgs args{
static_cast<void*>(d_y_mem.GetDeviceBuffer()),
static_cast<void*>(d_x_mem.GetDeviceBuffer()),
input_m,
max_iterations
};
ck_tile::SinkhornKnoppArgs args{static_cast<void*>(d_y_mem.GetDeviceBuffer()),
static_cast<void*>(d_x_mem.GetDeviceBuffer()),
input_n,
max_iterations};
d_x_mem.ToDevice(h_x.data());
d_y_mem.ToDevice(h_y.data());
using Problem = ck_tile::SinkhornKnoppProblem<XDataType,
YDataType,
TestSinkhornShape,
ComputeDataType
>;
using Kernel = ck_tile::SinkhornKnoppKernelDummyNonStochastic<
Problem,
ck_tile::SinkhornKnoppDefaultPolicy>;
using Problem =
ck_tile::SinkhornKnoppProblem<XDataType, YDataType, TestSinkhornShape, ComputeDataType>;
using Kernel =
ck_tile::SinkhornKnoppKernelDummyNonStochastic<Problem,
ck_tile::SinkhornKnoppDefaultPolicy>;
// Launch configuration
const ck_tile::index_t kBlockSize = Kernel::BlockSize();
@@ -73,26 +130,32 @@ class TestCkTileSinkHorn: public ::testing::Test
ck_tile::index_t kGridSize = 1; // TODO
//TODO
// if(!Kernel::IsSupportedArgument())
// {
// throw std::runtime_error("Wrong! Arguments not supported!\n");
// }
// TODO
// if(!Kernel::IsSupportedArgument())
// {
// throw std::runtime_error("Wrong! Arguments not supported!\n");
// }
ck_tile::launch_kernel(
ck_tile::stream_config{nullptr, false, 0},
ck_tile::make_kernel<kBlockPerCu>(Kernel{},
kGridSize,
kBlockSize,
0,
args));
ck_tile::make_kernel<kBlockPerCu>(Kernel{}, kGridSize, kBlockSize, 0, args));
// Reference computation
// TODO
ck_tile::HostTensor<YDataType> h_y_ref(input_shape, default_stride);
sinkhorn_knopp_ref(h_x, h_y_ref, max_iterations);
// Transfer data from device and check error for each operation
// TODO
// TODO: Test whether or not output is actually doubly stochastic
EXPECT_TRUE(true); // TODO
}
// TODO: Refine tolerances
const float rtol = 1e-7;
const float atol = 1e-8;
// Transfer data from device and check that it matches reference
d_y_mem.FromDevice(h_y.data());
bool result = true;
result &= ck_tile::check_err(
h_y, h_y_ref, "Error: Sinkhorn-Knopp doesn't match CPU reference!", rtol, atol);
EXPECT_TRUE(result);
}
};