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

[rocm-libraries] ROCm/rocm-libraries#5026 (commit fd5325b)

Browse Source 
[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.
This commit is contained in:
DarylHawkinsAMD
2026-03-04 23:29:35 +00:00
committed by assistant-librarian[bot]
parent ae4e632c7d
commit 147210ac72
2 changed files with 24 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp
View File

@@ -29,6 +29,12 @@ struct ElementwiseOpToCK<ElementwiseOperation::SCALE>
using Op = ck::tensor_operation::element_wise::Scale;
};
template <>
struct ElementwiseOpToCK<ElementwiseOperation::BILINEAR>
{
using Op = ck::tensor_operation::element_wise::Bilinear;
};
template <>
struct ElementwiseOpToCK<ElementwiseOperation::CLAMP>
{
@@ -41,6 +47,12 @@ struct ElementwiseOpToCK<ElementwiseOperation::SCALEADD_SCALEADD_RELU>
using Op = ck::tensor_operation::element_wise::ScaleAddScaleAddRelu;
};
template <>
struct ElementwiseOpToCK<ElementwiseOperation::ADD_CLAMP>
{
using Op = ck::tensor_operation::element_wise::AddClamp;
};
template <>
struct ElementwiseOpToCK<ElementwiseOperation::BIAS_BNORM_CLAMP>
{

12
experimental/builder/test/unit_conv_elementwise_op.cpp
View File

@@ -23,6 +23,12 @@ TEST(ConvElementwiseOp, AssignsOpsForScale)
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;
@@ -35,6 +41,12 @@ TEST(ConvElementwiseOp, AssignsOpsForScaleAddScaleAddRelu)
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;