From a285e176dd9d50ef18dd6c7810ad0f86e3ff5ded Mon Sep 17 00:00:00 2001
From: DarylHawkinsAMD <Daryl.Hawkins@amd.com>
Date: Wed, 4 Mar 2026 16:28:42 -0700
Subject: [PATCH] [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.
---
 .../factory/helpers/ck/conv_elementwise_op.hpp       | 12 ++++++++++++
 .../builder/test/unit_conv_elementwise_op.cpp        | 12 ++++++++++++
 2 files changed, 24 insertions(+)

diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp
index 0cc43fc679..bfc6fb40c3 100644
--- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp
+++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp
@@ -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>
 {
diff --git a/experimental/builder/test/unit_conv_elementwise_op.cpp b/experimental/builder/test/unit_conv_elementwise_op.cpp
index 610edd281e..3b66179d70 100644
--- a/experimental/builder/test/unit_conv_elementwise_op.cpp
+++ b/experimental/builder/test/unit_conv_elementwise_op.cpp
@@ -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;