mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-03-20 23:27:39 +00:00
147210ac72
[CK_BUILDER] Add BILINEAR and ADD_CLAMP elementwise operation mappings to CK builder (#5026) ## Motivation The CK kernels that MIOpen consumes use the BILINEAR and ADD_CLAMP operations. The operation mappings in the CK Builder API need to be in place to be able to instantiate those kernels using the builder. ## Technical Details Add the BILINEAR and ADD_CLAMP operation mappings to the builder ## Test Plan * Added builder tests for new helpers ## Test Result * New tests pass locally, waiting for test run ## Submission Checklist - [X] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
58 lines
1.8 KiB
C++
58 lines
1.8 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#include <gtest/gtest.h>
|
|
#include <type_traits>
|
|
|
|
#include "ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp"
|
|
|
|
namespace {
|
|
|
|
using ::ck_tile::builder::ElementwiseOperation;
|
|
using ::ck_tile::builder::factory::internal::ElementwiseOpToCK;
|
|
|
|
TEST(ConvElementwiseOp, AssignsOpsForPassThrough)
|
|
{
|
|
using Op = ElementwiseOpToCK<ElementwiseOperation::PASS_THROUGH>::Op;
|
|
EXPECT_TRUE((std::is_same_v<Op, ck::tensor_operation::element_wise::PassThrough>));
|
|
}
|
|
|
|
TEST(ConvElementwiseOp, AssignsOpsForScale)
|
|
{
|
|
using Op = ElementwiseOpToCK<ElementwiseOperation::SCALE>::Op;
|
|
EXPECT_TRUE((std::is_same_v<Op, ck::tensor_operation::element_wise::Scale>));
|
|
}
|
|
|
|
TEST(ConvElementwiseOp, AssignsOpsForBilinear)
|
|
{
|
|
using Op = ElementwiseOpToCK<ElementwiseOperation::BILINEAR>::Op;
|
|
EXPECT_TRUE((std::is_same_v<Op, ck::tensor_operation::element_wise::Bilinear>));
|
|
}
|
|
|
|
TEST(ConvElementwiseOp, AssignsOpsForClamp)
|
|
{
|
|
using Op = ElementwiseOpToCK<ElementwiseOperation::CLAMP>::Op;
|
|
EXPECT_TRUE((std::is_same_v<Op, ck::tensor_operation::element_wise::Clamp>));
|
|
}
|
|
|
|
TEST(ConvElementwiseOp, AssignsOpsForScaleAddScaleAddRelu)
|
|
{
|
|
using Op = ElementwiseOpToCK<ElementwiseOperation::SCALEADD_SCALEADD_RELU>::Op;
|
|
EXPECT_TRUE((std::is_same_v<Op, ck::tensor_operation::element_wise::ScaleAddScaleAddRelu>));
|
|
}
|
|
|
|
TEST(ConvElementwiseOp, AssignsOpsForAddClamp)
|
|
{
|
|
using Op = ElementwiseOpToCK<ElementwiseOperation::ADD_CLAMP>::Op;
|
|
EXPECT_TRUE((std::is_same_v<Op, ck::tensor_operation::element_wise::AddClamp>));
|
|
}
|
|
|
|
TEST(ConvElementwiseOp, AssignsOpsForBiasNormClamp)
|
|
{
|
|
using Op = ElementwiseOpToCK<ElementwiseOperation::BIAS_BNORM_CLAMP>::Op;
|
|
EXPECT_TRUE(
|
|
(std::is_same_v<Op, ck::tensor_operation::element_wise::BiasNormalizeInInferClamp>));
|
|
}
|
|
|
|
} // namespace
|