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

ck-builder: ck factory grouped conv fwd scaleadd scaleadd relu

Browse Source
This commit is contained in:
Robin Voetter
2025-10-27 15:21:33 +01:00
parent 2a309d7534
commit 2cb1d61ec6
2 changed files with 106 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
experimental/builder/test/CMakeLists.txt
View File

@@ -37,3 +37,4 @@ add_ck_factory_test(test_ck_factory_grouped_convolution_forward_convscale test_c
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_scale test_ck_factory_grouped_convolution_forward_scale.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_scaleadd_ab test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_bias_clamp test_ck_factory_grouped_convolution_forward_bias_clamp.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp)

105
experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp Normal file
View File

@@ -0,0 +1,105 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scaleadd_scaleadd_relu.hpp>
#include "testing_utils.hpp"
using ck_tile::test::InstanceSet;
using ck_tile::test::InstancesMatch;
namespace {
constexpr static auto NumDimSpatial = 3;
using InLayout = ck::tensor_layout::convolution::NDHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWGK;
using ck::tensor_layout::convolution::G_K;
using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD;
using ck::tensor_operation::element_wise::PassThrough;
using ck::tensor_operation::element_wise::ScaleAddScaleAddRelu;
template <typename T, typename U = T>
using DeviceOp = DeviceGroupedConvFwdMultipleABD<NumDimSpatial,
InLayout,
WeiLayout,
ck::Tuple<OutLayout, G_K>, // DsLayout
OutLayout,
T, // InDataType
T, // WeiDataType
ck::Tuple<U, U>, // DsDataType
T, // OutDataType
PassThrough,
PassThrough,
ScaleAddScaleAddRelu,
T>; // ComputeType
} // namespace
template <typename Case>
struct CkFactoryTestConvFwd : public testing::Test
{
static auto get_actual_instances()
{
return InstanceSet::from_factory<typename Case::DeviceOp>();
}
static auto get_expected_instances() { return InstanceSet(Case::expected); }
};
struct F32
{
using DeviceOp = ::DeviceOp<float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F16
{
using DeviceOp = ::DeviceOp<ck::half_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct BF16
{
using DeviceOp = ::DeviceOp<ck::bhalf_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct S8
{
using DeviceOp = ::DeviceOp<int8_t, float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
using TestTypes = ::testing::Types<F32, F16, BF16, S8>;
TYPED_TEST_SUITE(CkFactoryTestConvFwd, TestTypes);
TYPED_TEST(CkFactoryTestConvFwd, TestInstances)
{
auto actual = TestFixture::get_actual_instances();
auto expected = TestFixture::get_expected_instances();
EXPECT_THAT(actual, InstancesMatch(expected));
}