ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 19:44:39 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add unit tests for split-K autodeduction.

Browse Source
This commit is contained in:
Ville Pietilä
2025-07-10 14:11:55 +00:00
parent 08ebf5ecc2
commit 3ad1ec2981
4 changed files with 36 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
profiler/include/profiler/profile_grouped_conv_bwd_weight_impl.hpp
View File

@@ -262,7 +262,7 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
using AccDataType =
std::conditional_t<std::is_same_v<ComputeType, int8_t>, int32_t, float>;
const index_t num_accums = output.GetElementSize() / conv_param.K_;
const index_t num_accums_split_k = split_k_list[split_k_id];
const index_t num_accums_split_k = split_k_value;
// Calculate thresholds
auto rtol =
ck::utils::get_relative_threshold<ComputeType, WeiDataType, AccDataType>(

2
test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp
View File

@@ -30,7 +30,7 @@ class TestGroupedConvndBwdWeight : public ::testing::Test
using NDimSpatial = std::tuple_element_t<6, Tuple>;
std::vector<ck::utils::conv::ConvParam> conv_params;
std::vector<ck::index_t> split_ks{1, 2};
std::vector<ck::index_t> split_ks{-1, 1, 2};
bool skip_case(const ck::index_t split_k)
{

44
test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_interface_xdl.cpp
View File

@@ -52,7 +52,7 @@ class TestGroupedConvndBwdWeight : public ::testing::Test
// clang-format on
ck::utils::conv::ConvParam conv_param;
ck::index_t split_k{2};
std::vector<ck::index_t> split_ks{-1, 2};
template <ck::index_t NDimSpatial>
bool Run()
@@ -96,24 +96,30 @@ class TestGroupedConvndBwdWeight : public ::testing::Test
auto conv = GroupedConvBwdWeightDeviceInstance{};
auto argument = conv.MakeArgument(nullptr,
nullptr,
nullptr,
input_lengths,
input_strides,
filter_lengths,
weights_strides,
output_lengths,
output_strides,
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads,
PassThrough{},
PassThrough{},
PassThrough{},
split_k);
return conv.IsSupportedArgument(argument);
bool is_supported = true;
for (const auto split_k : split_ks)
{
auto argument = conv.MakeArgument(nullptr,
nullptr,
nullptr,
input_lengths,
input_strides,
filter_lengths,
weights_strides,
output_lengths,
output_strides,
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads,
PassThrough{},
PassThrough{},
PassThrough{},
split_k);
is_supported &=conv.IsSupportedArgument(argument);
}
return is_supported;
}
};

12
test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_v3_interface_xdl.cpp
View File

@@ -52,7 +52,7 @@ class TestGroupedConvndBwdWeight : public ::testing::Test
// clang-format on
ck::utils::conv::ConvParam conv_param;
ck::index_t split_k{2};
std::vector<ck::index_t> split_ks{-1, 2};
template <ck::index_t NDimSpatial>
bool Run()
@@ -96,7 +96,11 @@ class TestGroupedConvndBwdWeight : public ::testing::Test
auto conv = GroupedConvBwdWeightDeviceInstance{};
auto argument = conv.MakeArgument(nullptr,
bool is_supported = true;
for (const auto split_k : split_ks)
{
auto argument = conv.MakeArgument(nullptr,
nullptr,
nullptr,
input_lengths,
@@ -113,7 +117,9 @@ class TestGroupedConvndBwdWeight : public ::testing::Test
PassThrough{},
PassThrough{},
split_k);
return conv.IsSupportedArgument(argument);
is_supported &=conv.IsSupportedArgument(argument);
}
return is_supported;
}
};