feat(miopen-provider): support standalone binary pointwise

dnn-providers/integration-tests/src/common/BinaryPointwiseCommon.hpp

@@ -0,0 +1,101 @@
+// Copyright © Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier:  MIT
+
+#pragma once
+
+#include <cstdint>
+#include <ostream>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include <hipdnn_data_sdk/utilities/ShapeUtilities.hpp>
+#include <hipdnn_data_sdk/utilities/StringUtil.hpp>
+#include <hipdnn_flatbuffers_sdk/flatbuffer_utilities/FlatbufferTypeHelpers.hpp>
+
+namespace test_binary_pointwise_common
+{
+
+struct BinaryPointwiseShapeCase
+{
+    std::vector<int64_t> xDims;
+    std::vector<int64_t> xStrides;
+    std::vector<int64_t> yDims;
+    std::vector<int64_t> yStrides;
+    std::vector<int64_t> outDims;
+    std::vector<int64_t> outStrides;
+    std::string note;
+
+    BinaryPointwiseShapeCase(std::vector<int64_t>&& dims, std::string&& caseNote)
+        : xDims(dims)
+        , xStrides(hipdnn_data_sdk::utilities::generateStrides(xDims))
+        , yDims(xDims)
+        , yStrides(xStrides)
+        , outDims(std::move(dims))
+        , outStrides(xStrides)
+        , note(std::move(caseNote))
+    {
+    }
+
+    BinaryPointwiseShapeCase(std::vector<int64_t>&& xShapeDims,
+                             std::vector<int64_t>&& yShapeDims,
+                             std::string&& caseNote)
+        : xDims(std::move(xShapeDims))
+        , xStrides(hipdnn_data_sdk::utilities::generateStrides(xDims))
+        , yDims(std::move(yShapeDims))
+        , yStrides(hipdnn_data_sdk::utilities::generateStrides(yDims))
+        , outDims(xDims)
+        , outStrides(xStrides)
+        , note(std::move(caseNote))
+    {
+    }
+
+    friend std::ostream& operator<<(std::ostream& ss, const BinaryPointwiseShapeCase& tc)
+    {
+        using namespace hipdnn_data_sdk::utilities;
+
+        ss << "(xDims:";
+        vecToStream(ss, tc.xDims);
+        ss << " xStrides:";
+        vecToStream(ss, tc.xStrides);
+        ss << " yDims:";
+        vecToStream(ss, tc.yDims);
+        ss << " yStrides:";
+        vecToStream(ss, tc.yStrides);
+        ss << " outDims:";
+        vecToStream(ss, tc.outDims);
+        ss << " outStrides:";
+        vecToStream(ss, tc.outStrides);
+        if(!tc.note.empty())
+        {
+            ss << " note:" << tc.note;
+        }
+        ss << ")";
+
+        return ss;
+    }
+};
+
+inline std::vector<BinaryPointwiseShapeCase> createBinaryPointwiseShapeCases()
+{
+    std::vector<BinaryPointwiseShapeCase> cases;
+
+    cases.push_back(BinaryPointwiseShapeCase{{2, 4, 8, 8}, "4d"});
+    cases.push_back(BinaryPointwiseShapeCase{{2, 4, 8, 8}, {1, 4, 1, 1}, "Broadcast"});
+    cases.push_back(BinaryPointwiseShapeCase{{2, 4, 8, 8, 2}, "5d"});
+    cases.push_back(BinaryPointwiseShapeCase{{2, 8, 64}, "3d"});
+
+    return cases;
+}
+
+inline std::vector<hipdnn_flatbuffers_sdk::data_objects::PointwiseMode> createBinaryPointwiseModes()
+{
+    using PointwiseMode = hipdnn_flatbuffers_sdk::data_objects::PointwiseMode;
+    return {PointwiseMode::ADD,
+            PointwiseMode::SUB,
+            PointwiseMode::MUL,
+            PointwiseMode::MAX_OP,
+            PointwiseMode::MIN_OP};
+}
+
+} // namespace test_binary_pointwise_common

dnn-providers/integration-tests/src/integration-tests/pointwise/CMakeLists.txt

@@ -3,4 +3,5 @@
 
 add_cpp_graph_test_sources(
     IntegrationGpuActivationForward.cpp
+    IntegrationGpuPointwiseBinary.cpp
 )

dnn-providers/integration-tests/src/integration-tests/pointwise/IntegrationGpuPointwiseBinary.cpp

@@ -0,0 +1,121 @@
+// Copyright © Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier:  MIT
+
+#include <hip/hip_runtime.h>
+
+#include <tuple>
+
+#include <hipdnn_data_sdk/utilities/PlatformUtils.hpp>
+#include <hipdnn_test_sdk/utilities/CpuFpReferenceValidation.hpp>
+#include <hipdnn_test_sdk/utilities/SdkFrontendTypeConversions.hpp>
+#include <hipdnn_test_sdk/utilities/Seeds.hpp>
+#include <hipdnn_test_sdk/utilities/TestUtilities.hpp>
+
+#include "common/BinaryPointwiseCommon.hpp"
+#include "harness/IntegrationGraphVerificationHarness.hpp"
+
+using namespace hipdnn_frontend;
+using namespace hipdnn_data_sdk::utilities;
+using namespace hipdnn_test_sdk::utilities;
+using namespace hipdnn_integration_tests;
+using namespace test_binary_pointwise_common;
+
+namespace
+{
+
+using SdkPointwiseMode = hipdnn_flatbuffers_sdk::data_objects::PointwiseMode;
+using PointwiseBinaryTestCase = std::tuple<BinaryPointwiseShapeCase, SdkPointwiseMode>;
+
+template <typename DataType>
+class PointwiseBinary
+    : public IntegrationGraphVerificationHarness<DataType, PointwiseBinaryTestCase>
+{
+public:
+    struct GraphOutputs
+    {
+        std::shared_ptr<graph::TensorAttributes> out;
+    };
+
+    static std::pair<graph::Graph, GraphOutputs> buildGraph(hipdnnHandle_t handle,
+                                                            const PointwiseBinaryTestCase& tc)
+    {
+        const auto& [shape, mode] = tc;
+
+        graph::Graph graphObj;
+        graphObj.set_name("PointwiseBinaryTest");
+
+        auto dataType = getDataTypeEnumFromType<DataType>();
+        graphObj.set_intermediate_data_type(dataType)
+            .set_compute_data_type(hipdnn_frontend::DataType::FLOAT)
+            .set_io_data_type(dataType);
+
+        auto xAttr = graph::makeTensorAttributes("x", shape.xDims, shape.xStrides);
+        auto xTensorAttr = std::make_shared<graph::TensorAttributes>(std::move(xAttr));
+
+        auto yAttr = graph::makeTensorAttributes("y", shape.yDims, shape.yStrides);
+        auto yTensorAttr = std::make_shared<graph::TensorAttributes>(std::move(yAttr));
+
+        graph::PointwiseAttributes pwAttrs;
+        pwAttrs.set_mode(sdkToFrontendPointwiseMode(mode));
+
+        auto outTensorAttr = graphObj.pointwise(xTensorAttr, yTensorAttr, pwAttrs);
+        outTensorAttr->set_output(true);
+
+        auto validateResult = graphObj.validate();
+        if(validateResult.is_bad())
+        {
+            throw std::runtime_error("Failed to validate graph: " + validateResult.get_message());
+        }
+
+        auto buildResult = graphObj.build_operation_graph(handle);
+        if(buildResult.is_bad())
+        {
+            throw std::runtime_error("Failed to build operation graph: "
+                                     + buildResult.get_message());
+        }
+
+        return std::make_pair(std::move(graphObj), GraphOutputs{outTensorAttr});
+    }
+
+protected:
+    void runGraphTest() override
+    {
+        const auto& testCase = this->GetParam();
+        const auto& shape = std::get<BinaryPointwiseShapeCase>(testCase);
+
+        auto [graphObj, outputs] = buildGraph(getSharedHandle(), testCase);
+
+        this->registerValidator(outputs.out, this->getTolerance(graphObj, outputs.out));
+
+        this->setTestCaseNote(shape.note);
+        this->inputFillRecipes().setGlobalSeed(hipdnn_test_sdk::utilities::getGlobalTestSeed());
+        this->verifyGraph(graphObj);
+    }
+};
+
+using IntegrationGpuPointwiseBinaryFp32 = PointwiseBinary<float>;
+using IntegrationGpuPointwiseBinaryFp16 = PointwiseBinary<half>;
+
+} // namespace
+
+GTEST_ALLOW_UNINSTANTIATED_PARAMETERIZED_TEST(IntegrationGpuPointwiseBinaryFp32);
+TEST_P(IntegrationGpuPointwiseBinaryFp32, Correctness)
+{
+    runGraphTest();
+}
+
+GTEST_ALLOW_UNINSTANTIATED_PARAMETERIZED_TEST(IntegrationGpuPointwiseBinaryFp16);
+TEST_P(IntegrationGpuPointwiseBinaryFp16, Correctness)
+{
+    runGraphTest();
+}
+
+INSTANTIATE_TEST_SUITE_P(Smoke,
+                         IntegrationGpuPointwiseBinaryFp32,
+                         testing::Combine(testing::ValuesIn(createBinaryPointwiseShapeCases()),
+                                          testing::ValuesIn(createBinaryPointwiseModes())));
+
+INSTANTIATE_TEST_SUITE_P(Smoke,
+                         IntegrationGpuPointwiseBinaryFp16,
+                         testing::Combine(testing::ValuesIn(createBinaryPointwiseShapeCases()),
+                                          testing::ValuesIn(createBinaryPointwiseModes())));

dnn-providers/miopen-provider/CMakeLists.txt

@@ -203,29 +203,33 @@ set_target_properties(${MIOPEN_PROVIDER_MIOPEN_LIB} PROPERTIES SYSTEM TRUE)
 add_library(
     miopen_plugin_impl OBJECT
     engines/MiopenEngine.cpp
-    engines/plans/MiopenUnaryActivationPlan.cpp
-    engines/plans/MiopenUnaryActivationChecks.cpp
-    engines/plans/MiopenUnaryActivationPlanBuilder.cpp
     engines/plans/MiopenBatchnormApplicabilityChecks.cpp
-    engines/plans/MiopenBatchnormPlanBuilder.cpp
     engines/plans/MiopenBatchnormBwdPlan.cpp
     engines/plans/MiopenBatchnormFwdInferencePlan.cpp
     engines/plans/MiopenBatchnormFwdInferenceWithVariancePlan.cpp
     engines/plans/MiopenBatchnormFwdTrainingPlan.cpp
     engines/plans/MiopenBatchnormFwdTrainingPlanBuilder.cpp
+    engines/plans/MiopenBatchnormPlanBuilder.cpp
+    engines/plans/MiopenBinaryPointwiseChecks.cpp
+    engines/plans/MiopenBinaryPointwisePlan.cpp
+    engines/plans/MiopenBinaryPointwisePlanBuilder.cpp
     engines/plans/MiopenConvBwdPlan.cpp
     engines/plans/MiopenConvFwdBiasActivPlan.cpp
     engines/plans/MiopenConvFwdBiasActivPlanBuilder.cpp
     engines/plans/MiopenConvFwdPlan.cpp
     engines/plans/MiopenConvPlanBuilder.cpp
     engines/plans/MiopenConvWrwPlan.cpp
+    engines/plans/MiopenPointwiseTensorChecks.cpp
+    engines/plans/MiopenUnaryActivationChecks.cpp
+    engines/plans/MiopenUnaryActivationPlan.cpp
+    engines/plans/MiopenUnaryActivationPlanBuilder.cpp
     HipdnnMiopenHandle.cpp
     MiopenActivationDescriptor.cpp
     MiopenContainer.cpp
     MiopenConvDescriptor.cpp
-    MiopenUtils.cpp
-    MiopenTensor.cpp
     MiopenPluginPublic.cpp
+    MiopenTensor.cpp
+    MiopenUtils.cpp
 )
 
 target_include_directories(miopen_plugin_impl PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})

dnn-providers/miopen-provider/MiopenContainer.cpp

@@ -5,6 +5,7 @@
 #include "engines/MiopenEngine.hpp"
 #include "engines/plans/MiopenBatchnormFwdTrainingPlanBuilder.hpp"
 #include "engines/plans/MiopenBatchnormPlanBuilder.hpp"
+#include "engines/plans/MiopenBinaryPointwisePlanBuilder.hpp"
 #include "engines/plans/MiopenConvFwdBiasActivPlanBuilder.hpp"
 #include "engines/plans/MiopenConvPlanBuilder.hpp"
 #include "engines/plans/MiopenUnaryActivationPlanBuilder.hpp"
@@ -50,6 +51,7 @@ const std::vector<MiopenContainer::EngineDefinition>& MiopenContainer::getEngine
              engine->addPlanBuilder(std::make_unique<MiopenConvPlanBuilder>(false));
              engine->addPlanBuilder(std::make_unique<MiopenConvFwdBiasActivPlanBuilder>(false));
              engine->addPlanBuilder(std::make_unique<MiopenUnaryActivationPlanBuilder>());
+             engine->addPlanBuilder(std::make_unique<MiopenBinaryPointwisePlanBuilder>());
 
              return engine;
          }},

dnn-providers/miopen-provider/docs/OperationSupport.md

@@ -23,6 +23,7 @@ The following table lists all operations currently supported in hipDNN:
 | Convolution Forward + (Bias) + Activation<sup>4</sup> | FP16, BFP16, FP32 | NCL, NLC, NCHW, NHWC, NCDHW, NDHWC | Fused graph<sup>2,3,6</sup>, Deterministic<sup>5</sup> |
 | Convolution Wgrad   | FP16, BFP16, FP32 | NCL, NLC, NCHW, NHWC, NCDHW, NDHWC | Cross-correlation only<sup>2,6</sup>, Deterministic<sup>5</sup> |
 | Pointwise Activation (standalone) | FP16, FP32 | NCHW, NHWC | Single-node graph<sup>7</sup> |
+| Pointwise Binary Elementwise (standalone) | FP16, FP32 | NCHW, NCDHW | Single-node graph<sup>8</sup> |
 
 ¹ See Batchnorm Operations note below
 ² See Convolution Operations note below
@@ -31,6 +32,7 @@ The following table lists all operations currently supported in hipDNN:
 ⁵ See Deterministic Engine Support section
 ⁶ 3D tensors are internally padded to 4D for MIOpen compatibility. For convolution, the padding also extends the padding, stride and dilation vectors with a trailing spatial dimension of length 1.
 ⁷ See Standalone Activation note below
+⁸ See Standalone Binary Elementwise note below
 
 ## Detailed Requirements
 
@@ -77,6 +79,13 @@ The following table lists all operations currently supported in hipDNN:
 > [!NOTE]
 > **Standalone Activation:** A single-node pointwise graph supports ReLU, Clipped ReLU (configurable upper clip), CLAMP (configurable lower/upper clips), Leaky ReLU (configurable negative slope), Sigmoid and Tanh.
 
+> [!NOTE]
+> **Standalone Binary Elementwise:** A single-node pointwise graph supports ADD, SUB, MUL, MAX and MIN
+> across ranks 3-5, with packed channels-first strides and FP16/FP32 datatypes only. `in_0` must be
+> the full-shape operand; only `in_1` may broadcast per-axis at matching rank, and all three tensors
+> must share one datatype. SUB is executed internally as Add with `in_1` scaled by −1. In-place
+> execution is not supported.
+
 > [!NOTE]
 > **Sparse Support:** All operations currently work with dense tensors only.