From 83742918f19eccec8460edb1b1f30c18eee26308 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Mon, 2 Nov 2026 15:10:24 +0000 Subject: [PATCH 1/9] feat(miopen-provider): add RAII wrapper for MIOpen pooling descriptors MiopenPoolingDescriptor owns an N-d miopenPoolingDescriptor_t for 2 or 3 spatial dimensions. The index type and the workspace index mode are set after miopenSetNdPoolingDescriptor, which resets both. Indices are always recorded in Image mode, a flattened spatial offset into the unpadded input, which is hipDNN's index convention. If a setter fails after the descriptor is created, the descriptor is destroyed before rethrowing. Co-Authored-By: Claude Opus 5.5 (1M context) --- dnn-providers/miopen-provider/CMakeLists.txt | 1 + .../MiopenPoolingDescriptor.cpp | 100 ++++++++ .../MiopenPoolingDescriptor.hpp | 39 +++ .../miopen-provider/tests/CMakeLists.txt | 1 + .../tests/TestMiopenPoolingDescriptor.cpp | 223 ++++++++++++++++++ 5 files changed, 364 insertions(+) create mode 100644 dnn-providers/miopen-provider/MiopenPoolingDescriptor.cpp create mode 100644 dnn-providers/miopen-provider/MiopenPoolingDescriptor.hpp create mode 100644 dnn-providers/miopen-provider/tests/TestMiopenPoolingDescriptor.cpp diff --git a/dnn-providers/miopen-provider/CMakeLists.txt b/dnn-providers/miopen-provider/CMakeLists.txt index 53470809d58a..f7ff5d9ad982 100644 --- a/dnn-providers/miopen-provider/CMakeLists.txt +++ b/dnn-providers/miopen-provider/CMakeLists.txt @@ -229,6 +229,7 @@ add_library( MiopenContainer.cpp MiopenConvDescriptor.cpp MiopenPluginPublic.cpp + MiopenPoolingDescriptor.cpp MiopenTensor.cpp MiopenUtils.cpp ) diff --git a/dnn-providers/miopen-provider/MiopenPoolingDescriptor.cpp b/dnn-providers/miopen-provider/MiopenPoolingDescriptor.cpp new file mode 100644 index 000000000000..4841acf2d01c --- /dev/null +++ b/dnn-providers/miopen-provider/MiopenPoolingDescriptor.cpp @@ -0,0 +1,100 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include + +#include "MiopenPoolingDescriptor.hpp" +#include "MiopenUtils.hpp" + +namespace miopen_plugin +{ + +MiopenPoolingDescriptor::MiopenPoolingDescriptor(miopenPoolingMode_t mode, + const std::vector& window, + const std::vector& padding, + const std::vector& stride, + miopenIndexType_t indexType) +{ + if(window.size() != padding.size() || window.size() != stride.size()) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_BAD_PARAM, + "MiopenPoolingDescriptor: window, padding and stride must have the same length"); + } + + // MIOpen pooling handles 4D and 5D tensors only, i.e. 2 or 3 spatial dims. + if(window.size() != 2 && window.size() != 3) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_BAD_PARAM, + "MiopenPoolingDescriptor: only 2 or 3 spatial dimensions are supported, got " + + std::to_string(window.size())); + } + + THROW_ON_MIOPEN_FAILURE(miopenCreatePoolingDescriptor(&_descriptor)); + + try + { + THROW_ON_MIOPEN_FAILURE(miopenSetNdPoolingDescriptor(_descriptor, + mode, + static_cast(window.size()), + window.data(), + padding.data(), + stride.data())); + + // Setting the N-d descriptor resets the index type and the workspace index mode, so both + // must be applied afterwards. MIOpen defaults to mask mode (an offset inside the window), + // which its N-d max pooling solver rejects when indices are saved. + THROW_ON_MIOPEN_FAILURE(miopenSetPoolingIndexType(_descriptor, indexType)); + THROW_ON_MIOPEN_FAILURE( + miopenSetPoolingWorkSpaceIndexMode(_descriptor, miopenPoolingWorkspaceIndexImage)); + } + catch(...) + { + LOG_ON_MIOPEN_FAILURE(miopenDestroyPoolingDescriptor(_descriptor)); + _descriptor = nullptr; + throw; + } +} + +MiopenPoolingDescriptor::MiopenPoolingDescriptor(MiopenPoolingDescriptor&& other) noexcept + : _descriptor(other._descriptor) +{ + other._descriptor = nullptr; +} + +MiopenPoolingDescriptor& + MiopenPoolingDescriptor::operator=(MiopenPoolingDescriptor&& other) noexcept +{ + if(this == &other) + { + return *this; + } + + if(_descriptor != nullptr) + { + LOG_ON_MIOPEN_FAILURE(miopenDestroyPoolingDescriptor(_descriptor)); + } + + _descriptor = other._descriptor; + other._descriptor = nullptr; + + return *this; +} + +MiopenPoolingDescriptor::~MiopenPoolingDescriptor() +{ + if(_descriptor != nullptr) + { + LOG_ON_MIOPEN_FAILURE(miopenDestroyPoolingDescriptor(_descriptor)); + } +} + +miopenPoolingDescriptor_t MiopenPoolingDescriptor::poolingDescriptor() const +{ + return _descriptor; +} + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/MiopenPoolingDescriptor.hpp b/dnn-providers/miopen-provider/MiopenPoolingDescriptor.hpp new file mode 100644 index 000000000000..d76541e10a45 --- /dev/null +++ b/dnn-providers/miopen-provider/MiopenPoolingDescriptor.hpp @@ -0,0 +1,39 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include + +#include "MiopenApi.hpp" + +namespace miopen_plugin +{ + +// RAII owner of an N-d miopenPoolingDescriptor_t. Indices, when requested, are always recorded +// in image mode: a flattened spatial offset into the unpadded input, which is hipDNN's index +// convention. +class MiopenPoolingDescriptor +{ +public: + MiopenPoolingDescriptor(miopenPoolingMode_t mode, + const std::vector& window, + const std::vector& padding, + const std::vector& stride, + miopenIndexType_t indexType); + + MiopenPoolingDescriptor(const MiopenPoolingDescriptor&) = delete; + MiopenPoolingDescriptor& operator=(const MiopenPoolingDescriptor&) = delete; + + MiopenPoolingDescriptor(MiopenPoolingDescriptor&& other) noexcept; + MiopenPoolingDescriptor& operator=(MiopenPoolingDescriptor&& other) noexcept; + + ~MiopenPoolingDescriptor(); + + miopenPoolingDescriptor_t poolingDescriptor() const; + +private: + miopenPoolingDescriptor_t _descriptor{nullptr}; +}; + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/tests/CMakeLists.txt b/dnn-providers/miopen-provider/tests/CMakeLists.txt index 1c75613d4609..0ac575832516 100644 --- a/dnn-providers/miopen-provider/tests/CMakeLists.txt +++ b/dnn-providers/miopen-provider/tests/CMakeLists.txt @@ -45,6 +45,7 @@ add_executable( TestMiopenHipdnnEnginePluginHandle.cpp TestMiopenHipdnnPluginExecutionContext.cpp TestMiopenPluginPublic.cpp + TestMiopenPoolingDescriptor.cpp TestMiopenTensor.cpp TestMiopenUtils.cpp ) diff --git a/dnn-providers/miopen-provider/tests/TestMiopenPoolingDescriptor.cpp b/dnn-providers/miopen-provider/tests/TestMiopenPoolingDescriptor.cpp new file mode 100644 index 000000000000..0672be919ed3 --- /dev/null +++ b/dnn-providers/miopen-provider/tests/TestMiopenPoolingDescriptor.cpp @@ -0,0 +1,223 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include + +#include "MiopenApi.hpp" +#include +#include + +#include "MiopenPoolingDescriptor.hpp" + +using namespace miopen_plugin; + +struct PoolingDescriptorCase +{ + std::string name; + miopenPoolingMode_t mode; + std::vector window; + std::vector padding; + std::vector stride; + miopenIndexType_t indexType; +}; + +namespace +{ + +std::vector getRoundTripCases() +{ + return { + {"Max2dUint8", miopenPoolingMax, {2, 3}, {1, 0}, {2, 1}, miopenIndexUint8}, + {"Max2dUint32", miopenPoolingMax, {3, 2}, {0, 1}, {1, 2}, miopenIndexUint32}, + {"Average2d", miopenPoolingAverage, {3, 3}, {1, 1}, {2, 2}, miopenIndexUint8}, + {"AverageInclusive2d", + miopenPoolingAverageInclusive, + {2, 2}, + {0, 0}, + {2, 2}, + miopenIndexUint16}, + {"Max3dUint32", miopenPoolingMax, {2, 3, 4}, {1, 0, 2}, {2, 1, 3}, miopenIndexUint32}, + {"Max3dUint64", miopenPoolingMax, {1, 1, 1}, {0, 0, 0}, {1, 1, 1}, miopenIndexUint64}, + {"Average3d", miopenPoolingAverage, {3, 3, 3}, {1, 1, 1}, {1, 2, 3}, miopenIndexUint8}, + {"AverageInclusive3d", + miopenPoolingAverageInclusive, + {2, 4, 2}, + {1, 2, 1}, + {2, 2, 2}, + miopenIndexUint32}, + }; +} + +void expectBadParam(const std::vector& window, + const std::vector& padding, + const std::vector& stride) +{ + try + { + const MiopenPoolingDescriptor descriptor( + miopenPoolingMax, window, padding, stride, miopenIndexUint8); + ADD_FAILURE() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_BAD_PARAM); + } +} + +} // namespace + +class TestMiopenPoolingDescriptorRoundTrip : public ::testing::TestWithParam +{ +}; + +TEST_P(TestMiopenPoolingDescriptorRoundTrip, DescriptorReportsConfiguredValues) +{ + const auto& param = GetParam(); + const MiopenPoolingDescriptor descriptor( + param.mode, param.window, param.padding, param.stride, param.indexType); + ASSERT_NE(descriptor.poolingDescriptor(), nullptr); + + const auto spatialRank = static_cast(param.window.size()); + miopenPoolingMode_t mode{}; + int nbDims = 0; + std::vector window(param.window.size(), -1); + std::vector padding(param.window.size(), -1); + std::vector stride(param.window.size(), -1); + ASSERT_EQ(miopenGetNdPoolingDescriptor(descriptor.poolingDescriptor(), + spatialRank, + &mode, + &nbDims, + window.data(), + padding.data(), + stride.data()), + miopenStatusSuccess); + + EXPECT_EQ(mode, param.mode); + EXPECT_EQ(nbDims, spatialRank); + EXPECT_EQ(window, param.window); + EXPECT_EQ(padding, param.padding); + EXPECT_EQ(stride, param.stride); + + // miopenSetNdPoolingDescriptor resets both of these, so they only survive if the descriptor + // applies them afterwards. + miopenIndexType_t indexType{}; + ASSERT_EQ(miopenGetPoolingIndexType(descriptor.poolingDescriptor(), &indexType), + miopenStatusSuccess); + EXPECT_EQ(indexType, param.indexType); + + // Image mode for 2D too, where MIOpen's own default is mask mode. + miopenPoolingWorkspaceIndexMode_t indexMode{}; + ASSERT_EQ(miopenGetPoolingWorkSpaceIndexMode(descriptor.poolingDescriptor(), &indexMode), + miopenStatusSuccess); + EXPECT_EQ(indexMode, miopenPoolingWorkspaceIndexImage); +} + +INSTANTIATE_TEST_SUITE_P(AllCases, + TestMiopenPoolingDescriptorRoundTrip, + ::testing::ValuesIn(getRoundTripCases()), + [](const ::testing::TestParamInfo& info) { + return info.param.name; + }); + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForPaddingLengthMismatch) +{ + expectBadParam({2, 2}, {0}, {2, 2}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForStrideLengthMismatch) +{ + expectBadParam({2, 2}, {0, 0}, {2, 2, 2}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForWindowLengthMismatch) +{ + expectBadParam({2, 2, 2}, {0, 0}, {2, 2}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForNoSpatialDims) +{ + expectBadParam({}, {}, {}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForOneSpatialDim) +{ + expectBadParam({2}, {0}, {2}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForFourSpatialDims) +{ + expectBadParam({2, 2, 2, 2}, {0, 0, 0, 0}, {1, 1, 1, 1}); +} + +TEST(TestMiopenPoolingDescriptor, MoveConstructorTransfersOwnershipAndNullsSource) +{ + MiopenPoolingDescriptor source(miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint32); + auto* const handle = source.poolingDescriptor(); + ASSERT_NE(handle, nullptr); + + const MiopenPoolingDescriptor destination(std::move(source)); + + EXPECT_EQ(destination.poolingDescriptor(), handle); + // NOLINTNEXTLINE(bugprone-use-after-move) + EXPECT_EQ(source.poolingDescriptor(), nullptr); + + miopenIndexType_t indexType{}; + ASSERT_EQ(miopenGetPoolingIndexType(destination.poolingDescriptor(), &indexType), + miopenStatusSuccess); + EXPECT_EQ(indexType, miopenIndexUint32); +} + +TEST(TestMiopenPoolingDescriptor, MoveAssignmentTransfersOwnershipAndNullsSource) +{ + MiopenPoolingDescriptor source( + miopenPoolingAverage, {3, 3, 3}, {1, 1, 1}, {1, 1, 1}, miopenIndexUint8); + MiopenPoolingDescriptor destination( + miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint32); + auto* const handle = source.poolingDescriptor(); + ASSERT_NE(handle, nullptr); + ASSERT_NE(destination.poolingDescriptor(), handle); + + destination = std::move(source); + + EXPECT_EQ(destination.poolingDescriptor(), handle); + // NOLINTNEXTLINE(bugprone-use-after-move) + EXPECT_EQ(source.poolingDescriptor(), nullptr); + + miopenPoolingMode_t mode{}; + int nbDims = 0; + ASSERT_EQ(miopenGetNdPoolingDescriptor( + destination.poolingDescriptor(), 0, &mode, &nbDims, nullptr, nullptr, nullptr), + miopenStatusSuccess); + EXPECT_EQ(mode, miopenPoolingAverage); + EXPECT_EQ(nbDims, 3); +} + +TEST(TestMiopenPoolingDescriptor, MoveAssignmentFromMovedFromSourceLeavesBothNull) +{ + MiopenPoolingDescriptor first(miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint8); + const MiopenPoolingDescriptor keeper(std::move(first)); + MiopenPoolingDescriptor second(miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint8); + + // NOLINTNEXTLINE(bugprone-use-after-move) + second = std::move(first); + + EXPECT_EQ(second.poolingDescriptor(), nullptr); + // NOLINTNEXTLINE(bugprone-use-after-move) + EXPECT_EQ(first.poolingDescriptor(), nullptr); + EXPECT_NE(keeper.poolingDescriptor(), nullptr); +} + +TEST(TestMiopenPoolingDescriptor, SelfMoveAssignmentKeepsDescriptor) +{ + MiopenPoolingDescriptor descriptor(miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint8); + auto* const handle = descriptor.poolingDescriptor(); + + // Through a reference so the compiler does not reject the self-move outright. + auto& alias = descriptor; + descriptor = std::move(alias); + + // NOLINTNEXTLINE(bugprone-use-after-move) + EXPECT_EQ(descriptor.poolingDescriptor(), handle); +} From 0707ec044f4b368353d4a7d865790068b1a35ed9 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Mon, 2 Nov 2026 19:12:46 +0000 Subject: [PATCH 2/9] test(hipdnn): add a resample forward graph spec to the test SDK Add ResampleFwdTensorSpec, ResampleFwdGraphSpec and createResampleFwdGraph() to the test SDK (FlatbufferGraphTestUtils.hpp), next to createValidResampleFwdGraph(). They describe a configurable single-node ResampleFwd graph. The defaults are a valid fp32 max pool; packed() builds a pooling graph with a floor-mode Y and the padding mode that matches the resample mode, optionally with an INT8 or INT32 index; makeChannelsLast() gives X and Y channels-last strides. Each tensor can be made virtual, pass-by-value, ragged or strideless, or left out of the graph, so provider tests can build the graphs their checks decline. createValidResampleFwdGraph() stays; moving its callers to the spec is left for a follow-up. TestFlatbufferGraphTestUtils.cpp covers packed(), makeChannelsLast() and each branch of createResampleFwdGraph(). Co-Authored-By: Claude Opus 5.5 (1M context) --- .../utilities/FlatbufferGraphTestUtils.hpp | 212 ++++++++++++++++++ projects/hipdnn/test_sdk/tests/CMakeLists.txt | 1 + .../TestFlatbufferGraphTestUtils.cpp | 196 ++++++++++++++++ 3 files changed, 409 insertions(+) create mode 100644 projects/hipdnn/test_sdk/tests/utilities/TestFlatbufferGraphTestUtils.cpp diff --git a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp index 129776cb4e40..23f28882472d 100644 --- a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp +++ b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp @@ -3,7 +3,11 @@ #pragma once +#include #include +#include +#include +#include #include #include @@ -3427,6 +3431,214 @@ inline flatbuffers::FlatBufferBuilder return builder; } +// A tensor of a ResampleFwdGraphSpec. +struct ResampleFwdTensorSpec +{ + // A dense, non-virtual tensor that is part of the graph. + static ResampleFwdTensorSpec dense(int64_t uid, + std::vector dims, + std::vector strides, + hipdnn_flatbuffers_sdk::data_objects::DataType dataType + = hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT) + { + ResampleFwdTensorSpec tensor; + tensor.uid = uid; + tensor.dims = std::move(dims); + tensor.strides = std::move(strides); + tensor.dataType = dataType; + return tensor; + } + + int64_t uid = 0; + std::vector dims; + // nullopt leaves the strides field out of the flatbuffer. + std::optional> strides; + hipdnn_flatbuffers_sdk::data_objects::DataType dataType + = hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT; + bool isVirtual = false; + bool isRuntimePassByValue = false; + std::optional raggedOffsetTensorUid; + // false leaves the tensor out of the graph, so the node's reference to `uid` dangles. + bool inGraph = true; +}; + +// Configuration for a single-node ResampleFwdAttributes graph. The defaults describe a valid graph: +// fp32 max pooling of a packed NCHW 2x3x6x6 tensor with a 3x3 window, stride 2 and padding 1 +// (NEG_INF_PAD) into 2x3x3x3, without an index. +struct ResampleFwdGraphSpec +{ + static constexpr int64_t X_UID = 1; + static constexpr int64_t Y_UID = 2; + static constexpr int64_t INDEX_UID = 3; + + // Pools `xDims` with `window`, `stride` and the same `padding` before and after every spatial + // dimension: packed channels-first X and Y of `dataType`, Y sized by the floor-mode formula, and + // NEG_INF_PAD for MAXPOOL or ZERO_PAD for the average modes. With `indexType`, the node + // generates an index into a packed channels-first tensor shaped like Y. + static ResampleFwdGraphSpec + packed(hipdnn_flatbuffers_sdk::data_objects::ResampleMode mode, + const std::vector& xDims, + const std::vector& window, + const std::vector& stride, + const std::vector& padding, + hipdnn_flatbuffers_sdk::data_objects::DataType dataType + = hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT, + std::optional indexType + = std::nullopt) + { + using hipdnn_data_sdk::utilities::generateStrides; + + std::vector yDims{xDims.at(0), xDims.at(1)}; + for(size_t i = 0; i + 2 < xDims.size(); ++i) + { + yDims.push_back((xDims[i + 2] + 2 * padding.at(i) - window.at(i)) / stride.at(i) + 1); + } + + ResampleFwdGraphSpec spec; + spec.mode = mode; + spec.paddingMode = mode == hipdnn_flatbuffers_sdk::data_objects::ResampleMode::MAXPOOL + ? hipdnn_flatbuffers_sdk::data_objects::PaddingMode::NEG_INF_PAD + : hipdnn_flatbuffers_sdk::data_objects::PaddingMode::ZERO_PAD; + spec.x = ResampleFwdTensorSpec::dense(X_UID, xDims, generateStrides(xDims), dataType); + spec.y = ResampleFwdTensorSpec::dense(Y_UID, yDims, generateStrides(yDims), dataType); + if(indexType.has_value()) + { + spec.index = ResampleFwdTensorSpec::dense( + INDEX_UID, yDims, generateStrides(yDims), indexType.value()); + spec.generateIndex = true; + } + spec.window = window; + spec.stride = stride; + spec.setPadding(padding); + return spec; + } + + // Gives X and Y packed channels-last strides (NHWC, NDHWC). The index keeps its strides. + void makeChannelsLast() + { + using hipdnn_data_sdk::utilities::generateStrides; + using hipdnn_data_sdk::utilities::strideOrderNhwc; + x.strides = generateStrides(x.dims, strideOrderNhwc(x.dims.size())); + y.strides = generateStrides(y.dims, strideOrderNhwc(y.dims.size())); + } + + // Sets the same padding before and after every spatial dimension. + void setPadding(const std::vector& padding) + { + prePadding = padding; + postPadding = padding; + } + + hipdnn_flatbuffers_sdk::data_objects::ResampleMode mode + = hipdnn_flatbuffers_sdk::data_objects::ResampleMode::MAXPOOL; + hipdnn_flatbuffers_sdk::data_objects::PaddingMode paddingMode + = hipdnn_flatbuffers_sdk::data_objects::PaddingMode::NEG_INF_PAD; + ResampleFwdTensorSpec x = ResampleFwdTensorSpec::dense(X_UID, {2, 3, 6, 6}, {108, 36, 6, 1}); + ResampleFwdTensorSpec y = ResampleFwdTensorSpec::dense(Y_UID, {2, 3, 3, 3}, {27, 9, 3, 1}); + // Sets the node's index_tensor_uid; nullopt leaves it unset. + std::optional index; + // nullopt leaves the vector out of the node. + std::optional> window = std::vector{3, 3}; + std::optional> stride = std::vector{2, 2}; + std::optional> prePadding = std::vector{1, 1}; + std::optional> postPadding = std::vector{1, 1}; + std::optional generateIndex; + hipdnn_flatbuffers_sdk::data_objects::DataType computeDataType + = hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT; + bool overrideShapeEnabled = false; + // Number of identical resample nodes in the graph. + uint32_t nodeCount = 1; +}; + +// Builds a ResampleFwdAttributes graph from `spec`. Graph tensor uids are unique, so a tensor +// whose uid an earlier one already has is emitted once, and the node references it twice. +inline flatbuffers::FlatBufferBuilder createResampleFwdGraph(const ResampleFwdGraphSpec& spec) +{ + namespace data_objects = hipdnn_flatbuffers_sdk::data_objects; + + const auto toFlatbuffers = [](const auto& value) { + using T = typename std::decay_t::value_type; + return value.has_value() ? ::flatbuffers::Optional(value.value()) + : ::flatbuffers::Optional(::flatbuffers::nullopt); + }; + + flatbuffers::FlatBufferBuilder builder; + std::vector> tensorAttributes; + std::vector emittedUids; + + const auto addTensor = [&](const ResampleFwdTensorSpec& tensor, const char* name) { + if(!tensor.inGraph + || std::find(emittedUids.begin(), emittedUids.end(), tensor.uid) != emittedUids.end()) + { + return; + } + emittedUids.push_back(tensor.uid); + tensorAttributes.push_back(data_objects::CreateTensorAttributesDirect( + builder, + tensor.uid, + name, + tensor.dataType, + tensor.strides.has_value() ? &tensor.strides.value() : nullptr, + &tensor.dims, + tensor.isVirtual, + data_objects::TensorValue::NONE, + 0, + tensor.isRuntimePassByValue, + toFlatbuffers(tensor.raggedOffsetTensorUid))); + }; + + addTensor(spec.x, "x"); + addTensor(spec.y, "y"); + if(spec.index.has_value()) + { + addTensor(spec.index.value(), "index"); + } + + const auto optionalVector = [](const std::optional>& values) { + return values.has_value() ? &values.value() : nullptr; + }; + const auto indexUid = spec.index.has_value() + ? ::flatbuffers::Optional(spec.index->uid) + : ::flatbuffers::Optional(::flatbuffers::nullopt); + + std::vector> nodes; + for(uint32_t i = 0; i < spec.nodeCount; ++i) + { + const auto attributes + = data_objects::CreateResampleFwdAttributesDirect(builder, + spec.x.uid, + spec.y.uid, + indexUid, + optionalVector(spec.prePadding), + optionalVector(spec.postPadding), + optionalVector(spec.stride), + optionalVector(spec.window), + spec.mode, + spec.paddingMode, + toFlatbuffers(spec.generateIndex)); + + const auto name = "resample_fwd_" + std::to_string(i); + nodes.push_back( + data_objects::CreateNodeDirect(builder, + name.c_str(), + spec.computeDataType, + data_objects::NodeAttributes::ResampleFwdAttributes, + attributes.Union())); + } + + const auto graph = data_objects::CreateGraphDirect(builder, + "test", + data_objects::DataType::FLOAT, + data_objects::DataType::FLOAT, + data_objects::DataType::FLOAT, + &tensorAttributes, + &nodes, + flatbuffers::nullopt, + spec.overrideShapeEnabled); + builder.Finish(graph); + return builder; +} + inline flatbuffers::FlatBufferBuilder createValidResampleBwdGraph(bool useIndex = true, hipdnn_flatbuffers_sdk::data_objects::ResampleMode resampleMode diff --git a/projects/hipdnn/test_sdk/tests/CMakeLists.txt b/projects/hipdnn/test_sdk/tests/CMakeLists.txt index 3f52a9711dc9..1e2551a9249d 100644 --- a/projects/hipdnn/test_sdk/tests/CMakeLists.txt +++ b/projects/hipdnn/test_sdk/tests/CMakeLists.txt @@ -50,6 +50,7 @@ add_executable( utilities/TestDynamicTolerancesSdpa.cpp utilities/TestGraphTestHelpers.cpp utilities/TestFlatbufferDatatypeMapping.cpp + utilities/TestFlatbufferGraphTestUtils.cpp utilities/TestFlatbufferTensorAttributesUtils.cpp utilities/TestLoadGraphAndTensors.cpp utilities/TestLoggingUtils.cpp diff --git a/projects/hipdnn/test_sdk/tests/utilities/TestFlatbufferGraphTestUtils.cpp b/projects/hipdnn/test_sdk/tests/utilities/TestFlatbufferGraphTestUtils.cpp new file mode 100644 index 000000000000..93cee0b88e6c --- /dev/null +++ b/projects/hipdnn/test_sdk/tests/utilities/TestFlatbufferGraphTestUtils.cpp @@ -0,0 +1,196 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include + +#include + +#include +#include + +using hipdnn_flatbuffers_sdk::data_objects::DataType; +using hipdnn_flatbuffers_sdk::data_objects::GetGraph; +using hipdnn_flatbuffers_sdk::data_objects::PaddingMode; +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_test_sdk::utilities::createResampleFwdGraph; +using hipdnn_test_sdk::utilities::ResampleFwdGraphSpec; + +namespace +{ + +std::vector toVector(const flatbuffers::Vector& values) +{ + return {values.begin(), values.end()}; +} + +} // namespace + +// ============================================================================ +// ResampleFwdGraphSpec and createResampleFwdGraph +// ============================================================================ + +TEST(TestResampleFwdGraphSpec, PackedSizesFloorModeOutputAndMatchesPaddingMode) +{ + // Y: H = (7 + 2 - 3) / 2 + 1 = 4, W = (5 - 2) / 1 + 1 = 4. + const auto spec = ResampleFwdGraphSpec::packed(ResampleMode::AVGPOOL_INCLUDE_PADDING, + {1, 2, 7, 5}, + {3, 2}, + {2, 1}, + {1, 0}, + DataType::HALF); + + EXPECT_EQ(spec.mode, ResampleMode::AVGPOOL_INCLUDE_PADDING); + EXPECT_EQ(spec.paddingMode, PaddingMode::ZERO_PAD); + EXPECT_EQ(spec.x.dims, (std::vector{1, 2, 7, 5})); + EXPECT_EQ(spec.x.strides, (std::vector{70, 35, 5, 1})); + EXPECT_EQ(spec.x.dataType, DataType::HALF); + EXPECT_EQ(spec.y.dims, (std::vector{1, 2, 4, 4})); + EXPECT_EQ(spec.y.strides, (std::vector{32, 16, 4, 1})); + EXPECT_EQ(spec.y.dataType, DataType::HALF); + EXPECT_EQ(spec.window, (std::vector{3, 2})); + EXPECT_EQ(spec.stride, (std::vector{2, 1})); + EXPECT_EQ(spec.prePadding, (std::vector{1, 0})); + EXPECT_EQ(spec.postPadding, (std::vector{1, 0})); + EXPECT_FALSE(spec.index.has_value()); + EXPECT_FALSE(spec.generateIndex.has_value()); +} + +TEST(TestResampleFwdGraphSpec, PackedWithIndexTypeAddsChannelsFirstIndexShapedLikeY) +{ + const auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 6, 6}, + {3, 3}, + {2, 2}, + {1, 1}, + DataType::FLOAT, + DataType::INT8); + + EXPECT_EQ(spec.paddingMode, PaddingMode::NEG_INF_PAD); + ASSERT_TRUE(spec.index.has_value()); + EXPECT_EQ(spec.index->uid, ResampleFwdGraphSpec::INDEX_UID); + EXPECT_EQ(spec.index->dims, spec.y.dims); + EXPECT_EQ(spec.index->strides, (std::vector{27, 9, 3, 1})); + EXPECT_EQ(spec.index->dataType, DataType::INT8); + EXPECT_EQ(spec.generateIndex, std::optional(true)); +} + +TEST(TestResampleFwdGraphSpec, MakeChannelsLastRestridesXAndYButNotTheIndex) +{ + auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 6, 6}, + {3, 3}, + {2, 2}, + {1, 1}, + DataType::FLOAT, + DataType::INT32); + spec.makeChannelsLast(); + + EXPECT_EQ(spec.x.strides, (std::vector{108, 1, 18, 3})); + EXPECT_EQ(spec.y.strides, (std::vector{27, 1, 9, 3})); + ASSERT_TRUE(spec.index.has_value()); + EXPECT_EQ(spec.index->strides, (std::vector{27, 9, 3, 1})); +} + +TEST(TestCreateResampleFwdGraph, WritesTensorsAndNodeAttributes) +{ + auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 6, 6}, + {3, 3}, + {2, 2}, + {1, 1}, + DataType::FLOAT, + DataType::INT32); + spec.computeDataType = DataType::HALF; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + ASSERT_EQ(graph->tensors()->size(), 3u); + EXPECT_EQ(graph->tensors()->Get(2)->uid(), ResampleFwdGraphSpec::INDEX_UID); + EXPECT_EQ(graph->tensors()->Get(2)->data_type(), DataType::INT32); + ASSERT_EQ(graph->nodes()->size(), 1u); + EXPECT_EQ(graph->nodes()->Get(0)->compute_data_type(), DataType::HALF); + EXPECT_FALSE(graph->is_override_shape_enabled()); + + const auto* attributes = graph->nodes()->Get(0)->attributes_as_ResampleFwdAttributes(); + ASSERT_NE(attributes, nullptr); + EXPECT_EQ(attributes->x_tensor_uid(), ResampleFwdGraphSpec::X_UID); + EXPECT_EQ(attributes->y_tensor_uid(), ResampleFwdGraphSpec::Y_UID); + ASSERT_TRUE(attributes->index_tensor_uid().has_value()); + EXPECT_EQ(attributes->index_tensor_uid().value(), ResampleFwdGraphSpec::INDEX_UID); + EXPECT_EQ(toVector(*attributes->window()), (std::vector{3, 3})); + EXPECT_EQ(toVector(*attributes->stride()), (std::vector{2, 2})); + EXPECT_EQ(toVector(*attributes->pre_padding()), (std::vector{1, 1})); + EXPECT_EQ(toVector(*attributes->post_padding()), (std::vector{1, 1})); + EXPECT_EQ(attributes->resample_mode(), ResampleMode::MAXPOOL); + EXPECT_EQ(attributes->padding_mode(), PaddingMode::NEG_INF_PAD); + ASSERT_TRUE(attributes->generate_index().has_value()); + EXPECT_TRUE(attributes->generate_index().value()); +} + +TEST(TestCreateResampleFwdGraph, LeavesOutTensorNotInGraph) +{ + ResampleFwdGraphSpec spec; + spec.y.inGraph = false; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + ASSERT_EQ(graph->tensors()->size(), 1u); + EXPECT_EQ(graph->tensors()->Get(0)->uid(), ResampleFwdGraphSpec::X_UID); + // The node still names y, so its output uid dangles. + const auto* attributes = graph->nodes()->Get(0)->attributes_as_ResampleFwdAttributes(); + ASSERT_NE(attributes, nullptr); + EXPECT_EQ(attributes->y_tensor_uid(), ResampleFwdGraphSpec::Y_UID); + EXPECT_FALSE(attributes->index_tensor_uid().has_value()); +} + +TEST(TestCreateResampleFwdGraph, WritesNoStridesOrPoolingVectorsWhenUnset) +{ + ResampleFwdGraphSpec spec; + spec.x.strides = std::nullopt; + spec.window = std::nullopt; + spec.generateIndex = false; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + ASSERT_EQ(graph->tensors()->size(), 2u); + EXPECT_EQ(graph->tensors()->Get(0)->strides(), nullptr); + ASSERT_NE(graph->tensors()->Get(1)->strides(), nullptr); + const auto* attributes = graph->nodes()->Get(0)->attributes_as_ResampleFwdAttributes(); + ASSERT_NE(attributes, nullptr); + EXPECT_EQ(attributes->window(), nullptr); + ASSERT_NE(attributes->stride(), nullptr); + ASSERT_TRUE(attributes->generate_index().has_value()); + EXPECT_FALSE(attributes->generate_index().value()); +} + +TEST(TestCreateResampleFwdGraph, EmitsAliasedTensorOnce) +{ + ResampleFwdGraphSpec spec; + spec.y.uid = spec.x.uid; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + ASSERT_EQ(graph->tensors()->size(), 1u); + const auto* attributes = graph->nodes()->Get(0)->attributes_as_ResampleFwdAttributes(); + ASSERT_NE(attributes, nullptr); + EXPECT_EQ(attributes->x_tensor_uid(), attributes->y_tensor_uid()); +} + +TEST(TestCreateResampleFwdGraph, WritesNodeCountAndOverrideShapeFlag) +{ + ResampleFwdGraphSpec spec; + spec.nodeCount = 2; + spec.overrideShapeEnabled = true; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + EXPECT_EQ(graph->nodes()->size(), 2u); + EXPECT_TRUE(graph->is_override_shape_enabled()); +} From 68db273f05cc09c09ac5537692b4e37cf61f8e60 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Mon, 2 Nov 2026 19:19:06 +0000 Subject: [PATCH 3/9] feat(miopen-provider): add applicability checks for pooling forward resample_fwd_applicability::resolve() maps a single-node ResampleFwd graph to the MIOpen pooling configuration that the plan executes. If the graph is not supported it throws BAD_PARAM naming the reason; isSupported() logs that reason and declines. Accepted graphs: - MAXPOOL, AVGPOOL_EXCLUDE_PADDING or AVGPOOL_INCLUDE_PADDING, with FP32 compute; - FLOAT, HALF or BFLOAT16 X and Y of one type, rank 4 or 5; - symmetric padding smaller than the window, and a floor-mode Y; - NEG_INF_PAD for MAXPOOL and ZERO_PAD for average pooling; any other padding mode only when all padding is zero; - an INT8 or INT32 image-mode index, for max pooling in NCHW/NCDHW only. Average pooling with generate_index set but no index tensor runs without one: average pooling produces no index. X and Y are handed to MIOpen with the canonical packed strides of the layout they share. MIOpen infers each tensor's layout from its strides and picks native or transposing kernels per tensor, but sizes the transpose workspace from Y alone, so the two must never be classified differently. A spatial dimension with a single output gets the smallest stride that still gives one output. This changes nothing numerically, but keeps MIOpen's workspace and tile estimates sane for global pooling. Channels-last graphs are declined when their per-lane tile exceeds the scratch limit of MIOpen's tiled kernels, which have no naive fallback for these layouts. The unit tests check each rule once, matching its decline message, plus the boundaries and transformations the checks compute themselves: padding against the window, the padding-mode rule, the INT8 index range, the channels-last tile limit, stride normalization, canonical strides and the layout choice. They build their graphs with the test SDK's ResampleFwdGraphSpec. Co-Authored-By: Claude Opus 5.5 (1M context) --- dnn-providers/miopen-provider/CMakeLists.txt | 1 + .../engines/plans/MiopenResampleFwdChecks.cpp | 468 ++++++++++++ .../engines/plans/MiopenResampleFwdChecks.hpp | 58 ++ .../miopen-provider/tests/CMakeLists.txt | 1 + .../plans/TestMiopenResampleFwdChecks.cpp | 684 ++++++++++++++++++ 5 files changed, 1212 insertions(+) create mode 100644 dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.cpp create mode 100644 dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.hpp create mode 100644 dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdChecks.cpp diff --git a/dnn-providers/miopen-provider/CMakeLists.txt b/dnn-providers/miopen-provider/CMakeLists.txt index f7ff5d9ad982..c6ba057ede9c 100644 --- a/dnn-providers/miopen-provider/CMakeLists.txt +++ b/dnn-providers/miopen-provider/CMakeLists.txt @@ -221,6 +221,7 @@ add_library( engines/plans/MiopenConvPlanBuilder.cpp engines/plans/MiopenConvWrwPlan.cpp engines/plans/MiopenPointwiseTensorChecks.cpp + engines/plans/MiopenResampleFwdChecks.cpp engines/plans/MiopenUnaryActivationChecks.cpp engines/plans/MiopenUnaryActivationPlan.cpp engines/plans/MiopenUnaryActivationPlanBuilder.cpp diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.cpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.cpp new file mode 100644 index 000000000000..e9f9a623ae93 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.cpp @@ -0,0 +1,468 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include "engines/plans/MiopenResampleFwdChecks.hpp" + +namespace miopen_plugin::resample_fwd_applicability +{ + +using hipdnn_flatbuffers_sdk::data_objects::DataType; +using hipdnn_flatbuffers_sdk::data_objects::NodeAttributes; +using hipdnn_flatbuffers_sdk::data_objects::PaddingMode; +using hipdnn_flatbuffers_sdk::data_objects::ResampleFwdAttributes; +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_flatbuffers_sdk::data_objects::TensorAttributes; +using hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph; + +namespace +{ + +// The prefix isSupported() logs. Kept out of every throw below so it never doubles. +constexpr auto LOG_PREFIX = "Resample forward plan builder: "; + +[[noreturn]] void decline(const std::string& reason) +{ + throw hipdnn_plugin_sdk::HipdnnPluginException(HIPDNN_PLUGIN_STATUS_BAD_PARAM, reason); +} + +enum class Layout +{ + CHANNELS_FIRST, // NCHW / NCDHW + CHANNELS_LAST, // NHWC / NDHWC +}; + +struct TensorInfo +{ + const TensorAttributes* attr; + int64_t uid; + std::vector dims; + std::vector strides; +}; + +TensorInfo describeTensor(const IGraph& opGraph, int64_t uid, const std::string& name) +{ + const auto& tensorMap = opGraph.getTensorMap(); + const auto it = tensorMap.find(uid); + if(it == tensorMap.end() || it->second == nullptr) + { + decline(name + " tensor (uid " + std::to_string(uid) + + ") does not resolve in the graph's tensor map"); + } + + const auto* attr = it->second; + if(attr->virtual_()) + { + decline(name + " tensor must not be virtual"); + } + if(hipdnn_flatbuffers_sdk::utilities::isPassByValueTensor(attr)) + { + decline(name + " tensor must not be pass-by-value"); + } + if(attr->ragged_offset_tensor_uid().has_value()) + { + decline(name + " tensor must not be ragged"); + } + if(attr->dims() == nullptr || attr->strides() == nullptr + || attr->dims()->size() != attr->strides()->size()) + { + decline(name + " tensor must have one stride per dimension"); + } + + TensorInfo info{attr, + uid, + std::vector(attr->dims()->begin(), attr->dims()->end()), + std::vector(attr->strides()->begin(), attr->strides()->end())}; + + // MIOpen's pooling kernels address tensors with 32-bit offsets. + int64_t elementCount = 1; + for(const auto dim : info.dims) + { + if(dim <= 0) + { + decline(name + " tensor has a non-positive dimension"); + } + if(elementCount > std::numeric_limits::max() / dim) + { + decline(name + " tensor element count exceeds INT32_MAX"); + } + elementCount *= dim; + } + + return info; +} + +std::vector packedStrides(const std::vector& dims, Layout layout) +{ + using namespace hipdnn_data_sdk::utilities; + return layout == Layout::CHANNELS_FIRST ? generateStrides(dims) + : generateStrides(dims, strideOrderNhwc(dims.size())); +} + +// True when the tensor puts every element where a packed tensor of `layout` would. Length-1 +// axes are never stepped along, so their strides do not matter. +bool addressesLike(const TensorInfo& tensor, Layout layout) +{ + const auto expected = packedStrides(tensor.dims, layout); + for(size_t i = 0; i < tensor.dims.size(); ++i) + { + if(tensor.dims[i] > 1 && tensor.strides[i] != expected[i]) + { + return false; + } + } + return true; +} + +std::vector toSpatialVector(const flatbuffers::Vector* values, + size_t spatialRank, + const std::string& name, + int64_t minValue) +{ + if(values == nullptr || values->size() != spatialRank) + { + decline(name + " must have one entry per spatial dimension (" + std::to_string(spatialRank) + + ")"); + } + + std::vector result; + result.reserve(spatialRank); + for(const auto value : *values) + { + if(value < minValue || value > std::numeric_limits::max()) + { + decline(name + " values must be in [" + std::to_string(minValue) + ", INT_MAX], got " + + std::to_string(value)); + } + result.push_back(static_cast(value)); + } + return result; +} + +miopenPoolingMode_t toMiopenPoolingMode(ResampleMode mode) +{ + switch(mode) + { + case ResampleMode::MAXPOOL: + return miopenPoolingMax; + case ResampleMode::AVGPOOL_EXCLUDE_PADDING: + return miopenPoolingAverage; + case ResampleMode::AVGPOOL_INCLUDE_PADDING: + return miopenPoolingAverageInclusive; + default: + decline("unsupported resample mode " + + std::string(hipdnn_flatbuffers_sdk::data_objects::EnumNameResampleMode(mode))); + } +} + +miopenIndexType_t toMiopenIndexType(DataType dataType) +{ + // hipDNN's index is signed, MIOpen's unsigned; the range check in resolve() keeps every + // index MIOpen can write representable in the signed type. + switch(dataType) + { + case DataType::INT8: + return miopenIndexUint8; + case DataType::INT32: + return miopenIndexUint32; + default: + decline("index tensor data type must be INT8 or INT32, got " + + std::string(hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType))); + } +} + +// MIOpen pools channels-last tensors only by transposing them around its tiled NCHW kernels +// (TransposedPoolingFwd2d/Nd); unlike channels-first there is no naive fallback. A tiled kernel is +// applicable only while its per-lane tile fits in scratch, sized as sizeof_private_memory() in +// MIOpen's src/solver/pooling/forward2d.cpp and forwardNd.cpp; past that the forward call fails +// with "No solver found". The limit assumes wave64, the smaller of MIOpen's two. +bool channelsLastTileFits(const ResampleFwdConfig& config) +{ + constexpr uint64_t K_MAX_WAVE_SCRATCH_SIZE = uint64_t{1024} * 8191; + constexpr uint64_t K_LANE_SCRATCH_LIMIT = K_MAX_WAVE_SCRATCH_SIZE / 64; + + auto roundUpTo16 = [](uint64_t bytes) { return (bytes + 15) / 16 * 16; }; + + // Every factor is at least 1, so stopping as soon as the tile outgrows the limit also keeps + // the product from overflowing. + uint64_t tileBytes = config.dataType == DataType::FLOAT ? 4 : 2; + auto grow = [&tileBytes](uint64_t factor) { + if(factor > K_LANE_SCRATCH_LIMIT || tileBytes * factor > K_LANE_SCRATCH_LIMIT) + { + return false; + } + tileBytes *= factor; + return true; + }; + + const auto inputsSpanned = [&config](size_t dim, uint64_t outputsPerLane) { + return (outputsPerLane - 1) * static_cast(config.stride[dim]) + + static_cast(config.window[dim]); + }; + + if(config.window.size() == 2) + { + // 16 output rows of one output column per lane, plus their 16 fp32 accumulators. + return grow(inputsSpanned(0, 16)) && grow(inputsSpanned(1, 1)) + && roundUpTo16(tileBytes) + roundUpTo16(sizeof(float) * 16) <= K_LANE_SCRATCH_LIMIT; + } + + // A 4 x 4 x 4 block of outputs per lane. + return grow(inputsSpanned(0, 4)) && grow(inputsSpanned(1, 4)) && grow(inputsSpanned(2, 4)) + && roundUpTo16(tileBytes) <= K_LANE_SCRATCH_LIMIT; +} + +} // namespace + +ResampleFwdConfig resolve(const IGraph& opGraph) +{ + if(opGraph.nodeCount() != 1) + { + decline("applicable only for single-node graphs. Graph has " + + std::to_string(opGraph.nodeCount()) + " nodes"); + } + + if(!opGraph.hasOnlySupportedAttributes( + std::set{NodeAttributes::ResampleFwdAttributes})) + { + decline("graph contains unsupported node attributes"); + } + + if(opGraph.getNode(0).compute_data_type() != DataType::FLOAT) + { + decline("only supports nodes with an fp32 compute_data_type"); + } + + const auto& attrs = opGraph.getNodeWrapper(0).attributesAs(); + + ResampleFwdConfig config; + config.mode = toMiopenPoolingMode(attrs.resample_mode()); + const bool isMaxPool = attrs.resample_mode() == ResampleMode::MAXPOOL; + + if(attrs.x_tensor_uid() == attrs.y_tensor_uid()) + { + decline("does not support in-place execution; X and Y must be distinct tensors"); + } + + const auto x = describeTensor(opGraph, attrs.x_tensor_uid(), "X"); + const auto y = describeTensor(opGraph, attrs.y_tensor_uid(), "Y"); + + const auto rank = x.dims.size(); + if(rank != 4 && rank != 5) + { + decline("X must be 4D or 5D, got rank " + std::to_string(rank) + + "; MIOpen pooling has no 1D path"); + } + if(y.dims.size() != rank) + { + decline("X and Y must have the same rank"); + } + + const auto dataType = x.attr->data_type(); + if(dataType != DataType::FLOAT && dataType != DataType::HALF && dataType != DataType::BFLOAT16) + { + decline("X data type must be FLOAT, HALF or BFLOAT16, got " + + std::string(hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType))); + } + if(y.attr->data_type() != dataType) + { + decline("Y data type must match X"); + } + + const auto spatialRank = rank - 2; + config.window = toSpatialVector(attrs.window(), spatialRank, "window", 1); + config.stride = toSpatialVector(attrs.stride(), spatialRank, "stride", 1); + config.padding = toSpatialVector(attrs.pre_padding(), spatialRank, "pre_padding", 0); + + if(toSpatialVector(attrs.post_padding(), spatialRank, "post_padding", 0) != config.padding) + { + decline("asymmetric padding is not supported; MIOpen pools with one symmetric pad per " + "spatial dimension, so pre_padding must equal post_padding"); + } + + for(size_t i = 0; i < spatialRank; ++i) + { + if(config.padding[i] >= config.window[i]) + { + // A window that lies entirely in the padding is where hipDNN and MIOpen disagree: + // hipDNN's index sentinel is -1, MIOpen writes 0. + decline("padding must be smaller than the window in every spatial dimension, got " + "padding " + + std::to_string(config.padding[i]) + " and window " + + std::to_string(config.window[i]) + " at spatial dimension " + + std::to_string(i)); + } + } + + // MIOpen's max pooling can only select a padded position when no input in the window is above + // the lowest finite value of the type (all -inf, NaN or lowest), and its average pooling counts + // padded positions as zero. Any other padding_mode, PADDING_NOT_SET included, only agrees with + // that when there are no padded positions. + const auto implicitPaddingMode = isMaxPool ? PaddingMode::NEG_INF_PAD : PaddingMode::ZERO_PAD; + const bool hasPadding + = std::any_of(config.padding.begin(), config.padding.end(), [](int p) { return p > 0; }); + if(attrs.padding_mode() != implicitPaddingMode && hasPadding) + { + decline("padding_mode " + + std::string( + hipdnn_flatbuffers_sdk::data_objects::EnumNamePaddingMode(attrs.padding_mode())) + + " with resample mode " + + hipdnn_flatbuffers_sdk::data_objects::EnumNameResampleMode(attrs.resample_mode()) + + " requires all padding to be zero; MIOpen pads this mode as " + + hipdnn_flatbuffers_sdk::data_objects::EnumNamePaddingMode(implicitPaddingMode)); + } + + if(x.dims[0] != y.dims[0] || x.dims[1] != y.dims[1]) + { + decline("X and Y must have the same batch and channel dimensions"); + } + + for(size_t i = 0; i < spatialRank; ++i) + { + const auto paddedInput = x.dims[i + 2] + 2 * static_cast(config.padding[i]); + if(paddedInput < config.window[i]) + { + decline("the window must fit inside the padded input at spatial dimension " + + std::to_string(i)); + } + + // MIOpen does not check Y against its own (floor mode) output size formula. + const auto expected = (paddedInput - config.window[i]) / config.stride[i] + 1; + if(y.dims[i + 2] != expected) + { + decline("Y spatial dimension " + std::to_string(i) + " is " + + std::to_string(y.dims[i + 2]) + ", expected " + std::to_string(expected) + + " = floor((x + pre_padding + post_padding - window) / stride) + 1"); + } + + // With a single output the stride never steps, so the smallest stride that still gives + // one output changes nothing. MIOpen's estimates do depend on it: V2 sizes channels-last + // scratch for an input of y * stride - 2 * pad + window - 1, and the tiled kernels size + // their tile from the stride. Global pooling would otherwise overshoot both. + if(expected == 1) + { + config.stride[i] = static_cast(paddedInput - config.window[i] + 1); + } + } + + Layout layout = Layout::CHANNELS_FIRST; + if(addressesLike(x, Layout::CHANNELS_FIRST) && addressesLike(y, Layout::CHANNELS_FIRST)) + { + layout = Layout::CHANNELS_FIRST; + } + else if(addressesLike(x, Layout::CHANNELS_LAST) && addressesLike(y, Layout::CHANNELS_LAST)) + { + layout = Layout::CHANNELS_LAST; + } + else + { + decline("X and Y must both be packed channels-first (NCHW/NCDHW) or both packed " + "channels-last (NHWC/NDHWC)"); + } + + config.dataType = dataType; + if(layout == Layout::CHANNELS_LAST && !channelsLastTileFits(config)) + { + decline("the window and stride are too large for MIOpen's channels-last (NHWC/NDHWC) " + "pooling kernels"); + } + + config.xUid = x.uid; + config.xDims = x.dims; + config.xStrides = packedStrides(x.dims, layout); + config.yUid = y.uid; + config.yDims = y.dims; + config.yStrides = packedStrides(y.dims, layout); + + const bool generateIndex = attrs.generate_index().value_or(false); + const auto indexUid = attrs.index_tensor_uid(); + // Average pooling produces no index, so for it generate_index without an index tensor asks + // for nothing and is ignored; the frontend builds such graphs. + if(generateIndex && !indexUid.has_value() && isMaxPool) + { + decline("generate_index is set but the node has no index tensor"); + } + if(!generateIndex && indexUid.has_value()) + { + decline("the node has an index tensor but generate_index is not set"); + } + + if(indexUid.has_value()) + { + if(!isMaxPool) + { + decline("only max pooling produces indices"); + } + + // MIOpen's channels-last path transposes X and Y around a channels-first kernel but + // leaves the indices in packed NC(D)HW order. + if(layout != Layout::CHANNELS_FIRST) + { + decline("max pooling with generate_index requires channels-first (NCHW/NCDHW) X and " + "Y"); + } + + if(*indexUid == x.uid || *indexUid == y.uid) + { + decline("the index tensor must be distinct from X and Y"); + } + + const auto index = describeTensor(opGraph, *indexUid, "index"); + config.indexType = toMiopenIndexType(index.attr->data_type()); + + if(index.dims != y.dims) + { + decline("index tensor dimensions must equal Y dimensions"); + } + if(!addressesLike(index, Layout::CHANNELS_FIRST)) + { + decline("index tensor must be packed channels-first, the order MIOpen writes indices " + "in"); + } + + // Image-mode indices are flattened spatial offsets into X, from 0 to the X spatial size + // minus one. X's element count is already known to fit in INT32, so only an INT8 index + // can overflow. + int64_t spatialSize = 1; + for(size_t i = 2; i < rank; ++i) + { + spatialSize *= x.dims[i]; + } + if(index.attr->data_type() == DataType::INT8 + && spatialSize - 1 > std::numeric_limits::max()) + { + decline("X spatial size " + std::to_string(spatialSize) + + " exceeds the range of an INT8 index"); + } + + config.indexUid = indexUid; + } + + return config; +} + +bool isSupported(const IGraph& opGraph) +{ + try + { + resolve(opGraph); + return true; + } + catch(const std::exception& e) + { + HIPDNN_PLUGIN_LOG_INFO(LOG_PREFIX << e.what()); + return false; + } +} + +} // namespace miopen_plugin::resample_fwd_applicability diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.hpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.hpp new file mode 100644 index 000000000000..01baca10f752 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.hpp @@ -0,0 +1,58 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include + +#include +#include + +#include "MiopenApi.hpp" + +namespace miopen_plugin::resample_fwd_applicability +{ + +// Everything MiopenResampleFwdPlan needs, resolved from a graph the checks accepted. +// +// The X and Y strides are the canonical packed strides of the one layout both tensors share, +// not the graph's own strides. They address the same elements (only the strides of length-1 +// axes can differ), but MIOpen infers a tensor's layout from its strides and picks its native or +// transposing pooling path per tensor, while it sizes the transpose workspace from Y alone. +// Canonical strides make MIOpen classify X and Y identically. +struct ResampleFwdConfig +{ + miopenPoolingMode_t mode{miopenPoolingMax}; + std::vector window; + std::vector padding; // pre_padding == post_padding, MIOpen takes one symmetric pad + std::vector stride; + + hipdnn_flatbuffers_sdk::data_objects::DataType dataType{ + hipdnn_flatbuffers_sdk::data_objects::DataType::UNSET}; + + int64_t xUid{0}; + std::vector xDims; + std::vector xStrides; + + int64_t yUid{0}; + std::vector yDims; + std::vector yStrides; + + // Set when the node has generate_index. MIOpen writes the indices, packed in NC(D)HW order, + // to the front of its workspace; the plan copies them into this tensor. + std::optional indexUid; + miopenIndexType_t indexType{miopenIndexUint8}; +}; + +// Resolves a single-node ResampleFwd graph to a MIOpen pooling forward configuration. Throws +// HipdnnPluginException(HIPDNN_PLUGIN_STATUS_BAD_PARAM) naming the reason when the graph is not +// supported. Does not look at is_override_shape_enabled; the plan builder declines that first. +ResampleFwdConfig resolve(const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph); + +// Returns whether resolve() accepts the graph, logging the reason at INFO with prefix +// "Resample forward plan builder: " if it does not. +bool isSupported(const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph); + +} // namespace miopen_plugin::resample_fwd_applicability diff --git a/dnn-providers/miopen-provider/tests/CMakeLists.txt b/dnn-providers/miopen-provider/tests/CMakeLists.txt index 0ac575832516..35a0025def00 100644 --- a/dnn-providers/miopen-provider/tests/CMakeLists.txt +++ b/dnn-providers/miopen-provider/tests/CMakeLists.txt @@ -36,6 +36,7 @@ add_executable( engines/plans/TestMiopenConvPlanBuilder.cpp engines/plans/TestMiopenConvWrwPlan.cpp engines/plans/TestMiopenPointwiseTensorChecks.cpp + engines/plans/TestMiopenResampleFwdChecks.cpp engines/plans/TestMiopenUnaryActivationPlanBuilder.cpp engines/TestMiopenEngine.cpp main.cpp diff --git a/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdChecks.cpp b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdChecks.cpp new file mode 100644 index 000000000000..8b1a4d39ecfa --- /dev/null +++ b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdChecks.cpp @@ -0,0 +1,684 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// Applicability rules of resample_fwd_applicability::resolve(), one test per rule, each matching +// its decline message, plus the boundaries and transformations the provider computes itself. None +// of these tests needs a GPU: resolve() only reads the graph. + +#include +#include +#include +#include +#include + +#include +#include + +#include +#include +#include + +#include "engines/plans/MiopenResampleFwdChecks.hpp" + +using namespace miopen_plugin; + +using hipdnn_flatbuffers_sdk::data_objects::DataType; +using hipdnn_flatbuffers_sdk::data_objects::PaddingMode; +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; +using hipdnn_test_sdk::utilities::createResampleFwdGraph; +using hipdnn_test_sdk::utilities::ResampleFwdGraphSpec; +using resample_fwd_applicability::ResampleFwdConfig; + +namespace +{ + +constexpr int64_t INT_MAX_VALUE = std::numeric_limits::max(); + +// 2x3x6x6 -> 2x3x3x3 with a 3x3 window, stride 2 and padding 1, so the padding rules have +// something to bite on. +ResampleFwdGraphSpec defaultSpec(ResampleMode mode = ResampleMode::MAXPOOL, + std::optional indexType = std::nullopt) +{ + return ResampleFwdGraphSpec::packed( + mode, {2, 3, 6, 6}, {3, 3}, {2, 2}, {1, 1}, DataType::FLOAT, indexType); +} + +ResampleFwdConfig resolveSpec(const ResampleFwdGraphSpec& spec) +{ + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + return resample_fwd_applicability::resolve(graph); +} + +void expectAccepted(const ResampleFwdGraphSpec& spec) +{ + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + ASSERT_TRUE(graph.isValid()); + + EXPECT_TRUE(resample_fwd_applicability::isSupported(graph)); + try + { + resample_fwd_applicability::resolve(graph); + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + ADD_FAILURE() << "resolve() declined: " << e.what(); + } +} + +// Declined with BAD_PARAM, for the reason the message fragment names. The fragment guards against +// a decline test passing because some other rule happened to fire first. +void expectDeclined(const flatbuffers::FlatBufferBuilder& fbb, const std::string& reasonFragment) +{ + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + ASSERT_TRUE(graph.isValid()); + + EXPECT_FALSE(resample_fwd_applicability::isSupported(graph)); + try + { + resample_fwd_applicability::resolve(graph); + ADD_FAILURE() << "resolve() accepted a graph it should decline (" << reasonFragment << ")"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_BAD_PARAM); + EXPECT_THAT(e.what(), ::testing::HasSubstr(reasonFragment)); + } +} + +void expectDeclined(const ResampleFwdGraphSpec& spec, const std::string& reasonFragment) +{ + expectDeclined(createResampleFwdGraph(spec), reasonFragment); +} + +// Max pooling with a 1x1 window and stride, so Y matches X and only the X spatial size matters. +ResampleFwdGraphSpec identityWindowSpec(const std::vector& xDims, DataType indexType) +{ + const std::vector ones(xDims.size() - 2, 1); + const std::vector zeros(xDims.size() - 2, 0); + return ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, xDims, ones, ones, zeros, DataType::FLOAT, indexType); +} + +} // namespace + +// ============================================================================ +// Graph-level rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMultiNodeGraph) +{ + auto spec = defaultSpec(); + spec.nodeCount = 2; + expectDeclined(spec, "single-node"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForOtherNodeType) +{ + expectDeclined(hipdnn_test_sdk::utilities::createValidResampleBwdGraph(), + "unsupported node attributes"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForNonFloatComputeType) +{ + auto spec = defaultSpec(); + spec.computeDataType = DataType::HALF; + expectDeclined(spec, "compute_data_type"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForModeNotSet) +{ + auto spec = defaultSpec(); + spec.mode = ResampleMode::NOT_SET; + expectDeclined(spec, "unsupported resample mode"); +} + +// ============================================================================ +// X and Y tensor rules. Every per-tensor rule is checked on X; Y goes through the same check. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForYAliasingX) +{ + auto spec = defaultSpec(); + spec.y.uid = spec.x.uid; + expectDeclined(spec, "in-place"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnresolvableX) +{ + auto spec = defaultSpec(); + spec.x.inGraph = false; + expectDeclined(spec, "X tensor (uid 1) does not resolve"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnresolvableY) +{ + auto spec = defaultSpec(); + spec.y.inGraph = false; + expectDeclined(spec, "Y tensor (uid 2) does not resolve"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForVirtualX) +{ + auto spec = defaultSpec(); + spec.x.isVirtual = true; + expectDeclined(spec, "X tensor must not be virtual"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForPassByValueX) +{ + auto spec = defaultSpec(); + spec.x.isRuntimePassByValue = true; + expectDeclined(spec, "X tensor must not be pass-by-value"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForRaggedX) +{ + auto spec = defaultSpec(); + spec.x.raggedOffsetTensorUid = 7; + expectDeclined(spec, "X tensor must not be ragged"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMissingXStrides) +{ + auto spec = defaultSpec(); + spec.x.strides = std::nullopt; + expectDeclined(spec, "X tensor must have one stride per dimension"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForZeroDimInX) +{ + auto spec = defaultSpec(); + spec.x.dims[0] = 0; + expectDeclined(spec, "X tensor has a non-positive dimension"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForXElementCountAboveInt32Max) +{ + // 46341^2 > INT32_MAX; the stride collapses Y to a single element. + expectDeclined(ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 1, 46341, 46341}, {1, 1}, {46341, 46341}, {0, 0}), + "X tensor element count exceeds INT32_MAX"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForThreeDimensionalTensors) +{ + expectDeclined(ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {2, 3, 6}, {2}, {2}, {0}), + "X must be 4D or 5D"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForRankMismatch) +{ + auto spec = defaultSpec(); + spec.y.dims = {2, 3, 3, 3, 1}; + spec.y.strides = std::vector{27, 9, 3, 1, 1}; + expectDeclined(spec, "X and Y must have the same rank"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnsupportedDataType) +{ + auto spec = defaultSpec(); + spec.x.dataType = DataType::DOUBLE; + spec.y.dataType = DataType::DOUBLE; + expectDeclined(spec, "X data type must be FLOAT, HALF or BFLOAT16"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMismatchedYDataType) +{ + auto spec = defaultSpec(); + spec.y.dataType = DataType::HALF; + expectDeclined(spec, "Y data type must match X"); +} + +// ============================================================================ +// Pooling parameter rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForWindowLengthMismatch) +{ + auto spec = defaultSpec(); + spec.window = std::vector{3, 3, 3}; + expectDeclined(spec, "window must have one entry per spatial dimension"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForZeroStride) +{ + auto spec = defaultSpec(); + spec.stride = std::vector{0, 2}; + expectDeclined(spec, "stride values must be in [1, INT_MAX]"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForWindowAboveIntMax) +{ + auto spec = defaultSpec(); + spec.window = std::vector{INT_MAX_VALUE + 1, 3}; + expectDeclined(spec, "window values must be in [1, INT_MAX]"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForAsymmetricPadding) +{ + auto spec = defaultSpec(); + spec.postPadding = std::vector{0, 0}; + expectDeclined(spec, "asymmetric padding"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForPaddingEqualToWindow) +{ + expectDeclined( + ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {2, 3, 6, 6}, {2, 2}, {2, 2}, {0, 2}), + "padding must be smaller than the window"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForPaddingOneBelowWindow) +{ + expectAccepted( + ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {1, 1, 4, 4}, {3, 3}, {1, 1}, {2, 2})); +} + +// MIOpen pads max pooling with -inf and average pooling with zeros. Any other padding_mode is only +// equivalent when no position is padded. +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMismatchedPaddingModeWithPadding) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL); + spec.paddingMode = PaddingMode::ZERO_PAD; + expectDeclined(spec, "requires all padding to be zero"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForMismatchedPaddingModeWithoutPadding) +{ + auto spec = ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_INCLUDE_PADDING, {2, 3, 6, 6}, {3, 3}, {2, 2}, {0, 0}); + spec.paddingMode = PaddingMode::NEG_INF_PAD; + expectAccepted(spec); +} + +// ============================================================================ +// Shape rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForBatchMismatch) +{ + auto spec = defaultSpec(); + spec.y.dims = {1, 3, 3, 3}; + spec.y.strides = std::vector{27, 9, 3, 1}; + expectDeclined(spec, "same batch and channel dimensions"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForWindowLargerThanPaddedInput) +{ + // 2 + 2 * 1 < 5 in both spatial dimensions. + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {1, 1, 4, 4}, {5, 5}, {1, 1}, {1, 1}); + spec.x.dims = {1, 1, 2, 2}; + spec.x.strides = std::vector{4, 4, 2, 1}; + spec.y.dims = {1, 1, 1, 1}; + spec.y.strides = std::vector{1, 1, 1, 1}; + expectDeclined(spec, "window must fit inside the padded input"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForWindowCoveringPaddedInputExactly) +{ + expectAccepted(ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_INCLUDE_PADDING, {1, 1, 2, 2}, {4, 4}, {1, 1}, {1, 1})); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForOutputOfOtherThanFloorModeSize) +{ + auto spec = defaultSpec(); + spec.y.dims = {2, 3, 4, 3}; + spec.y.strides = std::vector{36, 12, 3, 1}; + expectDeclined(spec, "Y spatial dimension 0 is 4, expected 3"); +} + +// ============================================================================ +// Layout rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForChannelsFirstXWithChannelsLastY) +{ + auto spec = defaultSpec(); + spec.y.strides = std::vector{27, 1, 9, 3}; + expectDeclined(spec, "both be packed channels-first"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForArbitraryStridesOnLengthOneAxes) +{ + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {1, 1, 6, 6}, {3, 3}, {2, 2}, {1, 1}); + spec.x.strides = std::vector{7, 999, 6, 1}; + spec.y.strides = std::vector{0, 1, 3, 1}; + expectAccepted(spec); +} + +// ============================================================================ +// Index rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMaxPoolGenerateIndexWithoutIndexTensor) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL); + spec.generateIndex = true; + expectDeclined(spec, "generate_index is set but the node has no index tensor"); +} + +// Average pooling produces no index: the flag without an index tensor asks for nothing and is +// ignored. Graph::resample() builds this graph for average pooling. +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForAveragePoolingGenerateIndexWithoutIndexTensor) +{ + auto spec = defaultSpec(ResampleMode::AVGPOOL_EXCLUDE_PADDING); + spec.generateIndex = true; + expectAccepted(spec); + EXPECT_FALSE(resolveSpec(spec).indexUid.has_value()); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexTensorWithoutGenerateIndex) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.generateIndex = std::nullopt; + expectDeclined(spec, "generate_index is not set"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexWithAveragePooling) +{ + expectDeclined(defaultSpec(ResampleMode::AVGPOOL_EXCLUDE_PADDING, DataType::INT32), + "only max pooling produces indices"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexWithChannelsLastTensors) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.makeChannelsLast(); + expectDeclined(spec, "requires channels-first"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexAliasingX) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.index->uid = spec.x.uid; + expectDeclined(spec, "index tensor must be distinct from X and Y"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnresolvableIndex) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.index->inGraph = false; + expectDeclined(spec, "index tensor (uid 3) does not resolve"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnsupportedIndexDataType) +{ + expectDeclined(defaultSpec(ResampleMode::MAXPOOL, DataType::INT64), + "index tensor data type must be INT8 or INT32"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexDimsMismatch) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.index->dims = {2, 3, 3, 4}; + spec.index->strides = std::vector{36, 12, 4, 1}; + expectDeclined(spec, "index tensor dimensions must equal Y dimensions"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForChannelsLastIndex) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.index->strides = std::vector{27, 1, 9, 3}; + expectDeclined(spec, "index tensor must be packed channels-first"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForInt8IndexAtSpatialSize128) +{ + // Indices run up to 8 * 16 - 1 = 127, the largest INT8 value. N and C do not count. + expectAccepted(identityWindowSpec({1, 2, 8, 16}, DataType::INT8)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForInt8IndexAtSpatialSize129) +{ + expectDeclined(identityWindowSpec({1, 2, 3, 43}, DataType::INT8), + "X spatial size 129 exceeds the range of"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForInt32IndexAtSpatialSize129) +{ + expectAccepted(identityWindowSpec({1, 2, 3, 43}, DataType::INT32)); +} + +// ============================================================================ +// resolve() output. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, ResolveMapsResampleModesToMiopenModes) +{ + EXPECT_EQ(resolveSpec(defaultSpec(ResampleMode::MAXPOOL)).mode, miopenPoolingMax); + EXPECT_EQ(resolveSpec(defaultSpec(ResampleMode::AVGPOOL_EXCLUDE_PADDING)).mode, + miopenPoolingAverage); + EXPECT_EQ(resolveSpec(defaultSpec(ResampleMode::AVGPOOL_INCLUDE_PADDING)).mode, + miopenPoolingAverageInclusive); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReportsPoolingParameters) +{ + // Y: H = (7 + 2 - 3) / 2 + 1 = 4, W = (5 + 0 - 2) / 1 + 1 = 4. + const auto config = resolveSpec(ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_EXCLUDE_PADDING, {1, 2, 7, 5}, {3, 2}, {2, 1}, {1, 0})); + EXPECT_EQ(config.window, (std::vector{3, 2})); + EXPECT_EQ(config.padding, (std::vector{1, 0})); + EXPECT_EQ(config.stride, (std::vector{2, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReportsDataTypeUidsAndDims) +{ + auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 6, 6}, + {3, 3}, + {2, 2}, + {1, 1}, + DataType::BFLOAT16, + DataType::INT32); + spec.x.uid = 10; + spec.y.uid = 20; + spec.index->uid = 30; + + const auto config = resolveSpec(spec); + EXPECT_EQ(config.dataType, DataType::BFLOAT16); + EXPECT_EQ(config.xUid, 10); + EXPECT_EQ(config.yUid, 20); + EXPECT_EQ(config.indexUid, std::optional(30)); + EXPECT_EQ(config.xDims, (std::vector{2, 3, 6, 6})); + EXPECT_EQ(config.yDims, (std::vector{2, 3, 3, 3})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReportsIndexTypes) +{ + const auto int32Config = resolveSpec(defaultSpec(ResampleMode::MAXPOOL, DataType::INT32)); + EXPECT_EQ(int32Config.indexUid, std::optional(ResampleFwdGraphSpec::INDEX_UID)); + EXPECT_EQ(int32Config.indexType, miopenIndexUint32); + + const auto int8Config = resolveSpec(defaultSpec(ResampleMode::MAXPOOL, DataType::INT8)); + EXPECT_EQ(int8Config.indexType, miopenIndexUint8); + + EXPECT_FALSE(resolveSpec(defaultSpec(ResampleMode::MAXPOOL)).indexUid.has_value()); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReportsChannelsLastStrides) +{ + auto spec = defaultSpec(); + spec.makeChannelsLast(); + const auto config = resolveSpec(spec); + EXPECT_EQ(config.xStrides, (std::vector{108, 1, 18, 3})); + EXPECT_EQ(config.yStrides, (std::vector{27, 1, 9, 3})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReplacesStridesOfLengthOneAxes) +{ + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {1, 1, 6, 6}, {3, 3}, {2, 2}, {1, 1}); + spec.x.strides = std::vector{7, 999, 6, 1}; + spec.y.strides = std::vector{0, 1, 3, 1}; + + const auto config = resolveSpec(spec); + EXPECT_EQ(config.xStrides, (std::vector{36, 36, 6, 1})); + EXPECT_EQ(config.yStrides, (std::vector{9, 9, 3, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolvePrefersChannelsFirstWhenBothLayoutsMatch) +{ + // With C == 1, NHWC strides address like NCHW for both tensors. + auto spec = ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_EXCLUDE_PADDING, {2, 1, 4, 4}, {2, 2}, {2, 2}, {0, 0}); + spec.makeChannelsLast(); + + const auto config = resolveSpec(spec); + EXPECT_EQ(config.xStrides, (std::vector{16, 16, 4, 1})); + EXPECT_EQ(config.yStrides, (std::vector{4, 4, 2, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveFollowsChannelsLastInputForAmbiguousOutput) +{ + // 2x4x3x3 -> 2x4x1x1 with channels-last X. Y has only length-1 spatial axes, so its + // channels-first-style strides match both layouts; X settles it as channels-last. + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {2, 4, 3, 3}, {3, 3}, {1, 1}, {0, 0}); + spec.makeChannelsLast(); + spec.y.strides = std::vector{4, 1, 1, 1}; + + const auto config = resolveSpec(spec); + EXPECT_EQ(config.xStrides, (std::vector{36, 1, 12, 4})); + EXPECT_EQ(config.yStrides, (std::vector{4, 1, 4, 4})); +} + +// ============================================================================ +// Stride normalization: a spatial dimension with a single output never steps its window, so +// resolve() replaces its stride by the smallest one that still yields one output, +// x + 2 * pad - window + 1. Dimensions with more outputs keep their stride. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, ResolveNormalizesStrideOfGlobalPooling) +{ + const auto config = resolveSpec(ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 2, 64, 64}, {64, 64}, {64, 64}, {0, 0})); + EXPECT_EQ(config.yDims, (std::vector{1, 2, 1, 1})); + EXPECT_EQ(config.stride, (std::vector{1, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveNormalizesStrideAgainstPaddedInput) +{ + // 5 + 2 * 1 - 3 + 1 = 5; the padding itself is left alone. + const auto config = resolveSpec(ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 1, 5, 5}, {3, 3}, {10, 10}, {1, 1})); + EXPECT_EQ(config.stride, (std::vector{5, 5})); + EXPECT_EQ(config.padding, (std::vector{1, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveNormalizesOnlySingleOutputDims) +{ + // H: one output, 7 - 3 + 1 = 5. W: (12 - 3) / 4 + 1 = 3 outputs, stride kept. + const auto config = resolveSpec(ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 1, 7, 12}, {3, 3}, {100, 4}, {0, 0})); + EXPECT_EQ(config.yDims, (std::vector{1, 1, 1, 3})); + EXPECT_EQ(config.stride, (std::vector{5, 4})); +} + +// ============================================================================ +// Channels-last tile limit. MIOpen pools NHWC/NDHWC only through its tiled Forward2d/ForwardNd +// kernels, which are applicable while their per-lane tile fits 1024 * 8191 / 64 = 131056 bytes: +// 2D: roundUp16(esize * ((16 - 1) * stride_h + window_h) * window_w) + 64 +// 3D: roundUp16(esize * prod((4 - 1) * stride + window)) +// with the strides after normalization. Channels-first graphs fall back to the naive kernel. +// ============================================================================ + +namespace +{ + +constexpr const char* TILE_LIMIT_REASON = "too large for MIOpen's channels-last"; + +ResampleFwdGraphSpec tileSpec(const std::vector& xDims, + const std::vector& window, + const std::vector& stride, + bool channelsLast, + DataType dataType, + const std::optional>& padding = std::nullopt) +{ + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + xDims, + window, + stride, + padding.value_or(std::vector(window.size(), 0)), + dataType); + if(channelsLast) + { + spec.makeChannelsLast(); + } + return spec; +} + +} // namespace + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForPlanarFloatTileAtLimit) +{ + // 4 * (15 * 545 + 12) * 4 = 130992, + 64 = 131056. H keeps two outputs so its stride is not + // normalized. + expectAccepted(tileSpec({1, 2, 557, 4}, {12, 4}, {545, 1}, true, DataType::FLOAT)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForPlanarFloatTileOverLimit) +{ + // One window row more: 4 * 8188 * 4 = 131008, + 64 = 131072. + expectDeclined(tileSpec({1, 2, 558, 4}, {13, 4}, {545, 1}, true, DataType::FLOAT), + TILE_LIMIT_REASON); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForPlanarFloatTileOverLimitInChannelsFirst) +{ + expectAccepted(tileSpec({1, 2, 558, 4}, {13, 4}, {545, 1}, false, DataType::FLOAT)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForPlanar16BitTileAtLimit) +{ + // 2 * (15 * 545 + 12) * 8 = 130992, + 64 = 131056. + expectAccepted(tileSpec({1, 2, 557, 8}, {12, 8}, {545, 1}, true, DataType::HALF)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForFloatTileThatFitsAs16Bit) +{ + // The 16-bit at-limit shape needs twice the bytes in fp32. + expectDeclined(tileSpec({1, 2, 557, 8}, {12, 8}, {545, 1}, true, DataType::FLOAT), + TILE_LIMIT_REASON); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForVolumetricFloatTileUnderLimit) +{ + // Window 7, stride 8: 4 * 31^3 = 119164. + expectAccepted(tileSpec({1, 2, 16, 16, 16}, {7, 7, 7}, {8, 8, 8}, true, DataType::FLOAT)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForVolumetricFloatTileOverLimit) +{ + // Window 8, stride 8: 4 * 32^3 = 131072. + expectDeclined(tileSpec({1, 2, 16, 16, 16}, {8, 8, 8}, {8, 8, 8}, true, DataType::FLOAT), + TILE_LIMIT_REASON); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForChannelsLastGlobalPoolingThanksToNormalization) +{ + // Unnormalized, stride 64 would need 4 * (15 * 64 + 64) * 64 = 262144 bytes; normalized to 1, + // 4 * (15 + 64) * 64 = 20224. + expectAccepted(tileSpec({1, 2, 64, 64}, {64, 64}, {64, 64}, true, DataType::FLOAT)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForChannelsLastIntMaxWindowWithoutOverflow) +{ + // Padding INT_MAX - 1 lets an INT_MAX window and stride produce two outputs, so nothing is + // normalized and the tile estimate would overflow without the early stop. + const auto big = INT_MAX_VALUE; + expectDeclined(tileSpec({1, 2, 2, 2}, + {big, 1}, + {big, 1}, + true, + DataType::FLOAT, + std::vector{big - 1, 0}), + TILE_LIMIT_REASON); +} From a5576a2dba0220f5eefbb88e10393790938252ee Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Mon, 2 Nov 2026 19:34:19 +0000 Subject: [PATCH 4/9] feat(miopen-provider): add pooling forward plan backed by miopenPoolingForward MiopenResampleFwdPlan runs a resolved ResampleFwd configuration with miopenPoolingForward, using alpha 1 and beta 0. Its workspace is MIOpen's own, sized by miopenPoolingGetWorkSpaceSizeV2. That covers the indices and, for channels-last tensors, the transposed copies of X and Y. When the node generates an index, MIOpen leaves it in the first yElements * sizeof(index) bytes of the workspace, packed in NC(D)HW order. The plan copies it into the index tensor on MIOpen's stream. MIOpen never checks the size of the workspace it is given and knows nothing of the index tensor, so the plan throws INVALID_VALUE when a required workspace or the index buffer is missing. The tests run the plan on the GPU with values worked out by hand. Each fails when the plan's own code is wrong: the descriptor and tensors it builds, the workspace it reports and checks, alpha and beta, the index it copies out, or the canonical strides it hands MIOpen. Co-Authored-By: Claude Opus 5.5 (1M context) --- dnn-providers/miopen-provider/CMakeLists.txt | 1 + .../engines/plans/MiopenResampleFwdPlan.cpp | 142 ++++++ .../engines/plans/MiopenResampleFwdPlan.hpp | 56 +++ .../miopen-provider/tests/CMakeLists.txt | 1 + .../plans/TestMiopenResampleFwdPlan.cpp | 452 ++++++++++++++++++ 5 files changed, 652 insertions(+) create mode 100644 dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.cpp create mode 100644 dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.hpp create mode 100644 dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlan.cpp diff --git a/dnn-providers/miopen-provider/CMakeLists.txt b/dnn-providers/miopen-provider/CMakeLists.txt index c6ba057ede9c..a1b4ff2f4b81 100644 --- a/dnn-providers/miopen-provider/CMakeLists.txt +++ b/dnn-providers/miopen-provider/CMakeLists.txt @@ -222,6 +222,7 @@ add_library( engines/plans/MiopenConvWrwPlan.cpp engines/plans/MiopenPointwiseTensorChecks.cpp engines/plans/MiopenResampleFwdChecks.cpp + engines/plans/MiopenResampleFwdPlan.cpp engines/plans/MiopenUnaryActivationChecks.cpp engines/plans/MiopenUnaryActivationPlan.cpp engines/plans/MiopenUnaryActivationPlanBuilder.cpp diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.cpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.cpp new file mode 100644 index 000000000000..94c93746a734 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.cpp @@ -0,0 +1,142 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include + +#include + +#include "MiopenUtils.hpp" +#include "engines/plans/MiopenResampleFwdPlan.hpp" + +namespace +{ + +size_t indexTypeSize(miopenIndexType_t indexType) +{ + switch(indexType) + { + case miopenIndexUint8: + return sizeof(uint8_t); + case miopenIndexUint16: + return sizeof(uint16_t); + case miopenIndexUint32: + return sizeof(uint32_t); + case miopenIndexUint64: + return sizeof(uint64_t); + default: + throw hipdnn_plugin_sdk::HipdnnPluginException(HIPDNN_PLUGIN_STATUS_INTERNAL_ERROR, + "MiopenResampleFwdPlan: unknown index type"); + } +} + +size_t elementCount(const std::vector& dims) +{ + return static_cast( + std::accumulate(dims.begin(), dims.end(), int64_t{1}, std::multiplies<>())); +} + +} // namespace + +namespace miopen_plugin +{ + +MiopenResampleFwdPlan::MiopenResampleFwdPlan( + const resample_fwd_applicability::ResampleFwdConfig& config) + : _x(config.xUid, config.dataType, config.xDims, config.xStrides) + , _y(config.yUid, config.dataType, config.yDims, config.yStrides) + , _pooling(config.mode, config.window, config.padding, config.stride, config.indexType) + , _indexUid(config.indexUid) +{ + if(_indexUid.has_value()) + { + _indexBytes = elementCount(config.yDims) * indexTypeSize(config.indexType); + } + + THROW_ON_MIOPEN_FAILURE(miopenPoolingGetWorkSpaceSizeV2( + _pooling.poolingDescriptor(), _y.tensorDescriptor(), &_workspaceSize)); + + if(_workspaceSize < _indexBytes) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_INTERNAL_ERROR, + "MiopenResampleFwdPlan: MIOpen's pooling workspace (" + std::to_string(_workspaceSize) + + " bytes) cannot hold the " + std::to_string(_indexBytes) + " bytes of indices"); + } +} + +size_t + MiopenResampleFwdPlan::getWorkspaceSize([[maybe_unused]] const HipdnnMiopenHandle& handle) const +{ + return _workspaceSize; +} + +void MiopenResampleFwdPlan::execute(const HipdnnMiopenHandle& handle, + const hipdnnPluginDeviceBuffer_t* deviceBuffers, + uint32_t numDeviceBuffers, + void* workspace) const +{ + const auto xBuf + = hipdnn_plugin_sdk::findDeviceBuffer(_x.uid(), deviceBuffers, numDeviceBuffers); + const auto yBuf + = hipdnn_plugin_sdk::findDeviceBuffer(_y.uid(), deviceBuffers, numDeviceBuffers); + + void* indexPtr = nullptr; + if(_indexUid.has_value()) + { + indexPtr + = hipdnn_plugin_sdk::findDeviceBuffer(*_indexUid, deviceBuffers, numDeviceBuffers).ptr; + if(indexPtr == nullptr) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_INVALID_VALUE, + "MiopenResampleFwdPlan: the index tensor has no device buffer"); + } + } + + // MIOpen never checks the size of the workspace it is given, and rejects a null one only for + // max pooling that saves indices: a missing or short workspace is otherwise an out-of-bounds + // write on the device, not an error. + if(_workspaceSize > 0 && workspace == nullptr) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_INVALID_VALUE, + "MiopenResampleFwdPlan: requires " + std::to_string(_workspaceSize) + + " workspace bytes but none was provided"); + } + + const float alpha = 1.0f; + const float beta = 0.0f; + + THROW_ON_MIOPEN_FAILURE(miopenPoolingForward(handle.miopenHandle, + _pooling.poolingDescriptor(), + &alpha, + _x.tensorDescriptor(), + xBuf.ptr, + &beta, + _y.tensorDescriptor(), + yBuf.ptr, + _indexUid.has_value(), + workspace, + _workspaceSize)); + + if(_indexUid.has_value()) + { + // Enqueue on MIOpen's stream so the copy is ordered after the pooling kernel. + hipStream_t stream = nullptr; + THROW_ON_MIOPEN_FAILURE(miopenGetStream(handle.miopenHandle, &stream)); + + const auto status + = hipMemcpyAsync(indexPtr, workspace, _indexBytes, hipMemcpyDeviceToDevice, stream); + if(status != hipSuccess) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_INTERNAL_ERROR, + std::string("MiopenResampleFwdPlan: copying the pooling indices failed: ") + + hipGetErrorString(status)); + } + } +} + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.hpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.hpp new file mode 100644 index 000000000000..a9991bc27f82 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.hpp @@ -0,0 +1,56 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include + +#include + +#include "HipdnnMiopenHandle.hpp" +#include "MiopenPoolingDescriptor.hpp" +#include "MiopenTensor.hpp" +#include "engines/plans/MiopenResampleFwdChecks.hpp" + +namespace miopen_plugin +{ + +// Execution plan for a standalone resample forward (pooling) node, backed by +// miopenPoolingForward. +// +// The workspace is MIOpen's own (miopenPoolingGetWorkSpaceSizeV2): room for the indices plus, +// for channels-last tensors, the transposed copies of X and Y. When the node generates an +// index, MIOpen leaves it in the first yElements * sizeof(index) bytes of the workspace and the +// plan copies it into the index tensor. +class MiopenResampleFwdPlan : public hipdnn_plugin_sdk::IPlan +{ +public: + explicit MiopenResampleFwdPlan(const resample_fwd_applicability::ResampleFwdConfig& config); + + MiopenResampleFwdPlan(const MiopenResampleFwdPlan&) = delete; + MiopenResampleFwdPlan& operator=(const MiopenResampleFwdPlan&) = delete; + + MiopenResampleFwdPlan(MiopenResampleFwdPlan&&) = delete; + MiopenResampleFwdPlan& operator=(MiopenResampleFwdPlan&&) = delete; + + ~MiopenResampleFwdPlan() override = default; + + size_t getWorkspaceSize(const HipdnnMiopenHandle& handle) const override; + + void execute(const HipdnnMiopenHandle& handle, + const hipdnnPluginDeviceBuffer_t* deviceBuffers, + uint32_t numDeviceBuffers, + void* workspace = nullptr) const override; + +private: + MiopenTensor _x; + MiopenTensor _y; + MiopenPoolingDescriptor _pooling; + std::optional _indexUid; + size_t _indexBytes{0}; + size_t _workspaceSize{0}; +}; + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/tests/CMakeLists.txt b/dnn-providers/miopen-provider/tests/CMakeLists.txt index 35a0025def00..f45f291476ae 100644 --- a/dnn-providers/miopen-provider/tests/CMakeLists.txt +++ b/dnn-providers/miopen-provider/tests/CMakeLists.txt @@ -37,6 +37,7 @@ add_executable( engines/plans/TestMiopenConvWrwPlan.cpp engines/plans/TestMiopenPointwiseTensorChecks.cpp engines/plans/TestMiopenResampleFwdChecks.cpp + engines/plans/TestMiopenResampleFwdPlan.cpp engines/plans/TestMiopenUnaryActivationPlanBuilder.cpp engines/TestMiopenEngine.cpp main.cpp diff --git a/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlan.cpp b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlan.cpp new file mode 100644 index 000000000000..e4316440d75c --- /dev/null +++ b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlan.cpp @@ -0,0 +1,452 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// Executes MiopenResampleFwdPlan on small inputs with values worked out by hand. Each test fails +// when the plan's own code is wrong: the descriptor or tensors it builds from the configuration, +// the workspace it reports and checks, the alpha and beta it passes, or the index it copies out. +// How MIOpen pools is MIOpen's to test. + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include "MiopenApi.hpp" +#include +#include +#include +#include +#include +#include +#include +#include + +#include "HipdnnMiopenHandle.hpp" +#include "engines/plans/MiopenResampleFwdChecks.hpp" +#include "engines/plans/MiopenResampleFwdPlan.hpp" + +using namespace miopen_plugin; + +using hipdnn_data_sdk::types::half; +using hipdnn_data_sdk::utilities::iterateAlongDimensions; +using hipdnn_data_sdk::utilities::Tensor; +using hipdnn_flatbuffers_sdk::data_objects::DataType; +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; +using hipdnn_test_sdk::utilities::createResampleFwdGraph; +using hipdnn_test_sdk::utilities::ResampleFwdGraphSpec; + +namespace +{ + +constexpr float Y_SENTINEL = -12345.0f; +constexpr int INDEX_SENTINEL = -7; + +std::vector ascending(int64_t count, float first = 0.0f) +{ + std::vector values(static_cast(count)); + std::iota(values.begin(), values.end(), first); + return values; +} + +// Y and the indices in row-major logical order over the Y dims; `index` is empty when the node +// has no index tensor. +struct PoolingResult +{ + std::vector y; + std::vector index; +}; + +void expectYNear(const std::vector& actual, + const std::vector& expected, + double relativeTolerance) +{ + ASSERT_EQ(actual.size(), expected.size()); + for(size_t i = 0; i < actual.size(); ++i) + { + const auto want = static_cast(expected[i]); + EXPECT_NEAR( + static_cast(actual[i]), want, relativeTolerance * std::max(1.0, std::abs(want))) + << "Y element " << i; + } +} + +void expectMatches(const PoolingResult& actual, + const PoolingResult& expected, + double relativeTolerance = 1e-5) +{ + expectYNear(actual.y, expected.y, relativeTolerance); + EXPECT_EQ(actual.index, expected.index); +} + +// Host/device tensors for one execution: X filled from logical (row-major) values, Y and the +// index pre-filled with sentinels so unwritten elements show up. +template +class PoolingTensors +{ +public: + PoolingTensors(const ResampleFwdGraphSpec& spec, const std::vector& xValues) + : _spec(spec) + , _x(spec.x.dims, spec.x.strides.value()) + , _y(spec.y.dims, spec.y.strides.value()) + { + size_t next = 0; + iterateAlongDimensions(_spec.x.dims, [&](const std::vector& index) { + _x.setHostValue(static_cast(xValues.at(next++)), index); + }); + _x.markHostModified(); + + _y.fillWithValue(static_cast(Y_SENTINEL)); + + if(_spec.index.has_value()) + { + _index.emplace(_spec.index->dims, _spec.index->strides.value()); + _index->fillWithValue(static_cast(INDEX_SENTINEL)); + } + } + + std::vector deviceBuffers(bool includeIndex = true) + { + std::vector buffers{{_spec.x.uid, _x.rawDeviceData()}, + {_spec.y.uid, _y.rawDeviceData()}}; + if(includeIndex && _index.has_value()) + { + buffers.push_back({_spec.index->uid, _index->rawDeviceData()}); + } + return buffers; + } + + // Reads Y and the index back in logical order. The device work must have completed. + PoolingResult read() + { + PoolingResult result; + + _y.markDeviceModified(); + _y.rawHostData(); + iterateAlongDimensions(_spec.y.dims, [&](const std::vector& index) { + result.y.push_back(static_cast(_y.getHostValue(index))); + }); + + if(_index.has_value()) + { + _index->markDeviceModified(); + _index->rawHostData(); + iterateAlongDimensions(_spec.index->dims, [&](const std::vector& index) { + result.index.push_back(static_cast(_index->getHostValue(index))); + }); + } + return result; + } + +private: + ResampleFwdGraphSpec _spec; + Tensor _x; + Tensor _y; + std::optional> _index; +}; + +std::unique_ptr buildPlan(const ResampleFwdGraphSpec& spec) +{ + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + return std::make_unique(resample_fwd_applicability::resolve(graph)); +} + +// 4x4 input pooled 2x2 / 2 into 2x2. Window maxima (row, col): (0,1)=5, (1,2)=8, (3,1)=11, +// (2,3)=7. +const std::vector& planarMaxInput() +{ + // clang-format off + static const std::vector s_values{ + 1, 5, 2, 0, + 3, 4, 8, 6, + 9, 2, 1, 7, + 0, 11, 3, 2, + }; + // clang-format on + return s_values; +} + +const PoolingResult& planarMaxExpected() +{ + static const PoolingResult s_expected{{5, 8, 11, 7}, {1, 6, 13, 11}}; + return s_expected; +} + +ResampleFwdGraphSpec planarMaxSpec(std::optional indexType, + DataType dataType = DataType::FLOAT) +{ + return ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 1, 4, 4}, {2, 2}, {2, 2}, {0, 0}, dataType, indexType); +} + +// `channels` slices of 3x3, 3x3 window, stride 1, padding 1: corner windows see 4 valid inputs, +// edge windows 6 and the centre window all 9. +ResampleFwdGraphSpec paddedAverageSpec(int64_t channels = 1) +{ + return ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_EXCLUDE_PADDING, {1, channels, 3, 3}, {3, 3}, {1, 1}, {1, 1}); +} + +// paddedAverageSpec over 1..9: window sums divided by 4 (corners), 6 (edges) and 9 (centre). +const std::vector& paddedAverageExpected() +{ + static const std::vector s_expected{ + 3.0f, 3.5f, 4.0f, 4.5f, 5.0f, 5.5f, 6.0f, 6.5f, 7.0f}; + return s_expected; +} + +class ResampleFwdPlanFixture : public ::testing::Test +{ +protected: + void SetUp() override + { + SKIP_IF_NO_DEVICES(); + _handle = std::make_unique(); + } + + void synchronize() const + { + hipStream_t stream = nullptr; + ASSERT_EQ(miopenGetStream(_handle->miopenHandle, &stream), miopenStatusSuccess); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + } + + // Builds the plan for `spec`, executes it with a workspace of the size the plan asks for and + // returns Y (and the index) in logical order. + template + PoolingResult run(const ResampleFwdGraphSpec& spec, const std::vector& xValues) + { + const auto plan = buildPlan(spec); + PoolingTensors tensors(spec, xValues); + auto buffers = tensors.deviceBuffers(); + + const hipdnn_data_sdk::utilities::Workspace workspace(plan->getWorkspaceSize(*_handle)); + plan->execute( + *_handle, buffers.data(), static_cast(buffers.size()), workspace.get()); + synchronize(); + + return tensors.read(); + } + + std::unique_ptr _handle; +}; + +} // namespace + +class TestGpuMiopenResampleFwdPlan : public ResampleFwdPlanFixture +{ +}; + +// ============================================================================ +// Workspace: MIOpen's size for the index type the plan passes. +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlan, WorkspaceHoldsInt32Indices) +{ + const auto plan = buildPlan(planarMaxSpec(DataType::INT32)); + EXPECT_EQ(plan->getWorkspaceSize(*_handle), 4 * sizeof(int32_t)); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, WorkspaceHoldsInt8Indices) +{ + const auto plan = buildPlan(planarMaxSpec(DataType::INT8)); + EXPECT_EQ(plan->getWorkspaceSize(*_handle), 4 * sizeof(int8_t)); +} + +// ============================================================================ +// Execution. Y and the index start as sentinels, so a wrong alpha or beta, or an index that is +// not copied out, shows in the values. +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlan, MaxPoolWritesValuesAndInt32Indices) +{ + expectMatches(run(planarMaxSpec(DataType::INT32), planarMaxInput()), planarMaxExpected(), 0.0); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, MaxPoolWritesInt8Indices) +{ + expectMatches(run(planarMaxSpec(DataType::INT8), planarMaxInput()), + planarMaxExpected(), + 0.0); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, MaxPoolWritesHalfValues) +{ + // The inputs are small integers, exact in half precision. + expectMatches( + run(planarMaxSpec(DataType::INT32, DataType::HALF), planarMaxInput()), + planarMaxExpected(), + 0.0); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, MaxPoolWithoutIndexWritesValues) +{ + const auto result = run(planarMaxSpec(std::nullopt), planarMaxInput()); + expectYNear(result.y, planarMaxExpected().y, 0.0); + EXPECT_TRUE(result.index.empty()); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, AvgExcludePaddingDividesByValidCount) +{ + const auto result = run(paddedAverageSpec(), ascending(9, 1.0f)); + expectYNear(result.y, paddedAverageExpected(), 1e-5); + EXPECT_TRUE(result.index.empty()); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, AvgPoolingNhwcWritesValues) +{ + // Two channels: 1..9 and 10..18, whose averages are those of channel 0 plus 9. The plan hands + // MIOpen channels-last strides, so MIOpen transposes X and Y through the workspace. + auto spec = paddedAverageSpec(/*channels=*/2); + spec.makeChannelsLast(); + ASSERT_GT(buildPlan(spec)->getWorkspaceSize(*_handle), 0u); + + auto expected = paddedAverageExpected(); + for(const auto value : paddedAverageExpected()) + { + expected.push_back(value + 9.0f); + } + expectYNear(run(spec, ascending(18, 1.0f)).y, expected, 1e-5); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, PoolingChannelsLastStyleInputIntoNchwOutput) +{ + // X {2, 3, 1, 1} carries NHWC strides {3, 1, 3, 3}, which address like packed NCHW; Y is + // packed NCHW. Taken at face value MIOpen would read X as NHWC and Y as NCHW; the plan hands + // it canonical NCHW strides for both. Every 2x2 window holds the single valid pixel. + auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 1, 1}, + {2, 2}, + {1, 1}, + {1, 1}, + DataType::FLOAT, + DataType::INT32); + spec.x.strides = std::vector{3, 1, 3, 3}; + ASSERT_EQ(spec.y.strides.value(), (std::vector{12, 4, 2, 1})); + + const auto x = ascending(6, 1.0f); + PoolingResult expected; + for(const auto value : x) + { + expected.y.insert(expected.y.end(), 4, value); + expected.index.insert(expected.index.end(), 4, 0); + } + expectMatches(run(spec, x), expected, 0.0); +} + +// ============================================================================ +// Execution errors +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteThrowsInvalidValueWithoutRequiredWorkspace) +{ + const auto spec = planarMaxSpec(DataType::INT32); + const auto plan = buildPlan(spec); + ASSERT_GT(plan->getWorkspaceSize(*_handle), 0u); + + PoolingTensors tensors(spec, planarMaxInput()); + auto buffers = tensors.deviceBuffers(); + try + { + plan->execute(*_handle, buffers.data(), static_cast(buffers.size()), nullptr); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_INVALID_VALUE); + } +} + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteThrowsInvalidValueWithoutWorkspaceForNhwcAveragePooling) +{ + // No index, but channels-last tensors still need MIOpen's transpose buffers. + auto spec = ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_INCLUDE_PADDING, {2, 3, 5, 5}, {3, 3}, {2, 2}, {1, 1}); + spec.makeChannelsLast(); + const auto plan = buildPlan(spec); + ASSERT_GT(plan->getWorkspaceSize(*_handle), 0u); + + PoolingTensors tensors(spec, ascending(150)); + auto buffers = tensors.deviceBuffers(); + try + { + plan->execute(*_handle, buffers.data(), static_cast(buffers.size()), nullptr); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_INVALID_VALUE); + } +} + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteSucceedsWithoutWorkspaceWhenNoneIsRequired) +{ + const auto spec = paddedAverageSpec(); + const auto plan = buildPlan(spec); + ASSERT_EQ(plan->getWorkspaceSize(*_handle), 0u); + + PoolingTensors tensors(spec, ascending(9, 1.0f)); + auto buffers = tensors.deviceBuffers(); + ASSERT_NO_THROW( + plan->execute(*_handle, buffers.data(), static_cast(buffers.size()), nullptr)); + synchronize(); + + expectYNear(tensors.read().y, paddedAverageExpected(), 1e-5); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteThrowsInvalidValueForMissingIndexBuffer) +{ + const auto spec = planarMaxSpec(DataType::INT32); + const auto plan = buildPlan(spec); + + PoolingTensors tensors(spec, planarMaxInput()); + auto buffers = tensors.deviceBuffers(/*includeIndex=*/false); + const hipdnn_data_sdk::utilities::Workspace workspace(plan->getWorkspaceSize(*_handle)); + try + { + plan->execute( + *_handle, buffers.data(), static_cast(buffers.size()), workspace.get()); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_INVALID_VALUE); + } +} + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteThrowsInvalidValueForNullIndexPointer) +{ + const auto spec = planarMaxSpec(DataType::INT32); + const auto plan = buildPlan(spec); + + PoolingTensors tensors(spec, planarMaxInput()); + auto buffers = tensors.deviceBuffers(); + ASSERT_EQ(buffers.size(), 3u); + ASSERT_EQ(buffers[2].uid, ResampleFwdGraphSpec::INDEX_UID); + buffers[2].ptr = nullptr; + + const hipdnn_data_sdk::utilities::Workspace workspace(plan->getWorkspaceSize(*_handle)); + try + { + plan->execute( + *_handle, buffers.data(), static_cast(buffers.size()), workspace.get()); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_INVALID_VALUE); + } + synchronize(); + + // The plan threw before MIOpen ran: Y and the index still hold their sentinels. + const auto result = tensors.read(); + EXPECT_EQ(result.y, std::vector(4, Y_SENTINEL)); + EXPECT_EQ(result.index, std::vector(4, INDEX_SENTINEL)); +} From e227ba3951391e07971e8989a0641a7d8faa5767 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Mon, 2 Nov 2026 19:42:01 +0000 Subject: [PATCH 5/9] feat(miopen-provider): add plan builder for pooling forward on both MIOpen engines MiopenResampleFwdPlanBuilder routes a standalone ResampleFwd node to MiopenResampleFwdPlan. isApplicable declines graphs with override shapes and otherwise defers to the ResampleFwd checks. getMaxWorkspaceSize and buildPlan resolve the configuration through the same checks, so the workspace the builder reports is the one the plan asks for, and an unsupported graph throws BAD_PARAM from both. The builder has no execution settings and no custom knobs. The builder is registered on MIOPEN_ENGINE and MIOPEN_ENGINE_DETERMINISTIC. MIOpen's pooling forward kernels compute each output in one thread without atomics, so their results do not vary from run to run. The tests cover the override-shape decline, the deferral to the checks, a channels-last max-pooling build whose plan workspace equals the reported maximum, the BAD_PARAM errors, and both engines accepting a ResampleFwd graph from the real container. Co-Authored-By: Claude Opus 5.5 (1M context) --- dnn-providers/miopen-provider/CMakeLists.txt | 1 + .../miopen-provider/MiopenContainer.cpp | 3 + .../plans/MiopenResampleFwdPlanBuilder.cpp | 76 +++++++++ .../plans/MiopenResampleFwdPlanBuilder.hpp | 55 ++++++ .../miopen-provider/tests/CMakeLists.txt | 1 + .../tests/TestMiopenEngineManager.cpp | 36 ++++ .../TestMiopenResampleFwdPlanBuilder.cpp | 160 ++++++++++++++++++ 7 files changed, 332 insertions(+) create mode 100644 dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.cpp create mode 100644 dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.hpp create mode 100644 dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp diff --git a/dnn-providers/miopen-provider/CMakeLists.txt b/dnn-providers/miopen-provider/CMakeLists.txt index a1b4ff2f4b81..394fb2aa4487 100644 --- a/dnn-providers/miopen-provider/CMakeLists.txt +++ b/dnn-providers/miopen-provider/CMakeLists.txt @@ -223,6 +223,7 @@ add_library( engines/plans/MiopenPointwiseTensorChecks.cpp engines/plans/MiopenResampleFwdChecks.cpp engines/plans/MiopenResampleFwdPlan.cpp + engines/plans/MiopenResampleFwdPlanBuilder.cpp engines/plans/MiopenUnaryActivationChecks.cpp engines/plans/MiopenUnaryActivationPlan.cpp engines/plans/MiopenUnaryActivationPlanBuilder.cpp diff --git a/dnn-providers/miopen-provider/MiopenContainer.cpp b/dnn-providers/miopen-provider/MiopenContainer.cpp index 75679ca1c3c2..cc06d6710628 100644 --- a/dnn-providers/miopen-provider/MiopenContainer.cpp +++ b/dnn-providers/miopen-provider/MiopenContainer.cpp @@ -8,6 +8,7 @@ #include "engines/plans/MiopenBinaryPointwisePlanBuilder.hpp" #include "engines/plans/MiopenConvFwdBiasActivPlanBuilder.hpp" #include "engines/plans/MiopenConvPlanBuilder.hpp" +#include "engines/plans/MiopenResampleFwdPlanBuilder.hpp" #include "engines/plans/MiopenUnaryActivationPlanBuilder.hpp" #include @@ -52,6 +53,7 @@ const std::vector& MiopenContainer::getEngine engine->addPlanBuilder(std::make_unique(false)); engine->addPlanBuilder(std::make_unique()); engine->addPlanBuilder(std::make_unique()); + engine->addPlanBuilder(std::make_unique()); return engine; }}, @@ -68,6 +70,7 @@ const std::vector& MiopenContainer::getEngine engine->addPlanBuilder(std::make_unique(true)); engine->addPlanBuilder(std::make_unique()); engine->addPlanBuilder(std::make_unique()); + engine->addPlanBuilder(std::make_unique()); return engine; }} diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.cpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.cpp new file mode 100644 index 000000000000..678d09070094 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.cpp @@ -0,0 +1,76 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include +#include + +#include "engines/plans/MiopenResampleFwdChecks.hpp" +#include "engines/plans/MiopenResampleFwdPlan.hpp" +#include "engines/plans/MiopenResampleFwdPlanBuilder.hpp" + +namespace miopen_plugin +{ + +bool MiopenResampleFwdPlanBuilder::isApplicable( + [[maybe_unused]] const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph) const +{ + if(opGraph.getGraph().is_override_shape_enabled()) + { + HIPDNN_PLUGIN_LOG_INFO("Resample forward plan builder does not support override shapes"); + return false; + } + + return resample_fwd_applicability::isSupported(opGraph); +} + +size_t MiopenResampleFwdPlanBuilder::getMaxWorkspaceSize( + const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + [[maybe_unused]] const HipdnnMiopenSettings& executionSettings) const +{ + const MiopenResampleFwdPlan plan(resample_fwd_applicability::resolve(opGraph)); + return plan.getWorkspaceSize(handle); +} + +void MiopenResampleFwdPlanBuilder::initializeExecutionSettings( + [[maybe_unused]] const HipdnnMiopenHandle& handle, + [[maybe_unused]] const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + [[maybe_unused]] const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IEngineConfig& + engineConfig, + [[maybe_unused]] HipdnnMiopenSettings& executionSettings) const +{ +} + +void MiopenResampleFwdPlanBuilder::buildPlan( + [[maybe_unused]] const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + [[maybe_unused]] const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IEngineConfig& + engineConfig, + HipdnnMiopenContext& executionContext) const +{ + const auto& nodeWrapper = opGraph.getNodeWrapper(0); + const auto& attrs + = nodeWrapper.attributesAs(); + + HIPDNN_PLUGIN_LOG_INFO( + "Building resample forward " + << hipdnn_flatbuffers_sdk::data_objects::EnumNameResampleMode(attrs.resample_mode()) + << " plan for node: " << nodeWrapper.name()); + + auto plan + = std::make_unique(resample_fwd_applicability::resolve(opGraph)); + executionContext.setPlan(std::move(plan)); +} + +std::vector + MiopenResampleFwdPlanBuilder::getCustomKnobs( + [[maybe_unused]] const HipdnnMiopenHandle& handle, + [[maybe_unused]] const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph) const +{ + return {}; +} + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.hpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.hpp new file mode 100644 index 000000000000..bec442d7a82b --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.hpp @@ -0,0 +1,55 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include + +#include +#include + +#include "HipdnnMiopenContext.hpp" +#include "HipdnnMiopenHandle.hpp" +#include "HipdnnMiopenSettings.hpp" + +namespace miopen_plugin +{ + +// PlanBuilder for a standalone resample forward (pooling) node: MAXPOOL, +// AVGPOOL_EXCLUDE_PADDING and AVGPOOL_INCLUDE_PADDING, routed to miopenPoolingForward. +class MiopenResampleFwdPlanBuilder : public hipdnn_plugin_sdk::IPlanBuilder +{ +public: + MiopenResampleFwdPlanBuilder() = default; + ~MiopenResampleFwdPlanBuilder() override = default; + + MiopenResampleFwdPlanBuilder(const MiopenResampleFwdPlanBuilder&) = delete; + MiopenResampleFwdPlanBuilder& operator=(const MiopenResampleFwdPlanBuilder&) = delete; + + bool isApplicable( + const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph) const override; + + size_t getMaxWorkspaceSize(const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + const HipdnnMiopenSettings& executionSettings) const override; + + void initializeExecutionSettings( + const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IEngineConfig& engineConfig, + HipdnnMiopenSettings& executionSettings) const override; + + void buildPlan(const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IEngineConfig& engineConfig, + HipdnnMiopenContext& executionContext) const override; + + std::vector getCustomKnobs( + const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph) const override; +}; + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/tests/CMakeLists.txt b/dnn-providers/miopen-provider/tests/CMakeLists.txt index f45f291476ae..1aa0d3b103b5 100644 --- a/dnn-providers/miopen-provider/tests/CMakeLists.txt +++ b/dnn-providers/miopen-provider/tests/CMakeLists.txt @@ -38,6 +38,7 @@ add_executable( engines/plans/TestMiopenPointwiseTensorChecks.cpp engines/plans/TestMiopenResampleFwdChecks.cpp engines/plans/TestMiopenResampleFwdPlan.cpp + engines/plans/TestMiopenResampleFwdPlanBuilder.cpp engines/plans/TestMiopenUnaryActivationPlanBuilder.cpp engines/TestMiopenEngine.cpp main.cpp diff --git a/dnn-providers/miopen-provider/tests/TestMiopenEngineManager.cpp b/dnn-providers/miopen-provider/tests/TestMiopenEngineManager.cpp index 4d2960c7cd41..c48f7fed4df4 100644 --- a/dnn-providers/miopen-provider/tests/TestMiopenEngineManager.cpp +++ b/dnn-providers/miopen-provider/tests/TestMiopenEngineManager.cpp @@ -4,16 +4,21 @@ #include #include #include +#include #include +#include +#include #include #include +#include #include #include #include #include "HipdnnMiopenContext.hpp" #include "HipdnnMiopenHandle.hpp" +#include "MiopenContainer.hpp" #include "mocks/MockEngine.hpp" #include "mocks/MockHipdnnMiopenContext.hpp" @@ -243,3 +248,34 @@ TEST(TestMiopenEngineManager, InitializeExecutionContextThrowsOnInvalidEngineId) manager.initializeExecutionContext(dummyHandle, mockGraph, mockEngineConfig, execCtx), hipdnn_plugin_sdk::HipdnnPluginException); } + +// The real engine set: MIOpen's pooling forward is bit-reproducible (one thread per output, no +// atomics), so both the default and the deterministic engine take a ResampleFwd graph. +TEST(TestGpuMiopenContainer, DefaultAndDeterministicEnginesAcceptResampleFwd) +{ + SKIP_IF_NO_DEVICES(); + + using hipdnn_data_sdk::utilities::MIOPEN_ENGINE_DETERMINISTIC_ID; + using hipdnn_data_sdk::utilities::MIOPEN_ENGINE_ID; + using hipdnn_flatbuffers_sdk::data_objects::PaddingMode; + using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; + using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; + + MiopenContainer container; + HipdnnMiopenHandle handle; + auto& manager = container.getEngineManager(); + + for(const auto mode : {ResampleMode::MAXPOOL, ResampleMode::AVGPOOL_EXCLUDE_PADDING}) + { + SCOPED_TRACE(hipdnn_flatbuffers_sdk::data_objects::EnumNameResampleMode(mode)); + auto fbb = createValidResampleFwdGraph( + mode == ResampleMode::MAXPOOL ? std::optional(true) : std::nullopt, + mode, + mode == ResampleMode::MAXPOOL ? PaddingMode::NEG_INF_PAD : PaddingMode::ZERO_PAD); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + + const auto applicable = manager.getApplicableEngineIds(handle, graph); + EXPECT_THAT(applicable, ::testing::Contains(MIOPEN_ENGINE_ID)); + EXPECT_THAT(applicable, ::testing::Contains(MIOPEN_ENGINE_DETERMINISTIC_ID)); + } +} diff --git a/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp new file mode 100644 index 000000000000..b0b241cb51c4 --- /dev/null +++ b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp @@ -0,0 +1,160 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +#include "HipdnnMiopenContext.hpp" +#include "HipdnnMiopenHandle.hpp" +#include "HipdnnMiopenSettings.hpp" +#include "engines/plans/MiopenResampleFwdPlanBuilder.hpp" + +using namespace miopen_plugin; + +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; +using hipdnn_test_sdk::utilities::createResampleFwdGraph; +using hipdnn_test_sdk::utilities::MockEngineConfig; +using hipdnn_test_sdk::utilities::ResampleFwdGraphSpec; + +namespace +{ + +// A spec the checks decline: asymmetric padding. +ResampleFwdGraphSpec unsupportedSpec() +{ + ResampleFwdGraphSpec spec; + spec.postPadding = std::vector{0, 0}; + return spec; +} + +} // namespace + +class TestGpuMiopenResampleFwdPlanBuilder : public ::testing::Test +{ +protected: + void SetUp() override + { + SKIP_IF_NO_DEVICES(); + _handle = std::make_unique(); + } + + bool isApplicable(const ResampleFwdGraphSpec& spec) const + { + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + return _planBuilder.isApplicable(*_handle, graph); + } + + size_t maxWorkspaceSize(const ResampleFwdGraphSpec& spec) const + { + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + const HipdnnMiopenSettings settings; + return _planBuilder.getMaxWorkspaceSize(*_handle, graph, settings); + } + + MiopenResampleFwdPlanBuilder _planBuilder; + MockEngineConfig _engineConfig; + std::unique_ptr _handle; +}; + +// ============================================================================ +// isApplicable +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, IsApplicableDefersToTheChecks) +{ + EXPECT_TRUE(isApplicable(ResampleFwdGraphSpec{})); + EXPECT_FALSE(isApplicable(unsupportedSpec())); +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, IsApplicableReturnsFalseForOverrideShapeEnabledGraph) +{ + ResampleFwdGraphSpec spec; + spec.overrideShapeEnabled = true; + EXPECT_FALSE(isApplicable(spec)); +} + +// ============================================================================ +// getMaxWorkspaceSize and buildPlan +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, BuildPlanWorkspaceMatchesMaxWorkspaceSize) +{ + // Channels-last max pooling: MIOpen's workspace holds the index region and the transposed X + // and Y, so the builder and the plan must size it from the same resolved configuration. + ResampleFwdGraphSpec spec; + spec.makeChannelsLast(); + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + const auto expected = maxWorkspaceSize(spec); + ASSERT_GT(expected, 0u); + + HipdnnMiopenContext context; + ASSERT_NO_THROW(_planBuilder.buildPlan(*_handle, graph, _engineConfig, context)); + ASSERT_TRUE(context.hasValidPlan()); + EXPECT_EQ(context.plan().getWorkspaceSize(*_handle), expected); +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, MaxWorkspaceSizeThrowsBadParamForUnsupportedGraph) +{ + try + { + maxWorkspaceSize(unsupportedSpec()); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_BAD_PARAM); + } +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, BuildPlanThrowsBadParamForUnsupportedGraph) +{ + auto fbb = createResampleFwdGraph(unsupportedSpec()); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + + HipdnnMiopenContext context; + try + { + _planBuilder.buildPlan(*_handle, graph, _engineConfig, context); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_BAD_PARAM); + } + EXPECT_FALSE(context.hasValidPlan()); +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, InitializeExecutionSettingsLeavesSettingsUntouched) +{ + auto fbb = createResampleFwdGraph(ResampleFwdGraphSpec{}); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + + HipdnnMiopenSettings settings; + EXPECT_NO_THROW( + _planBuilder.initializeExecutionSettings(*_handle, graph, _engineConfig, settings)); + EXPECT_FALSE(settings.workspaceSizeLimit().has_value()); + EXPECT_FALSE(settings.defaultWorkspaceSize().has_value()); + EXPECT_FALSE(settings.benchmarkingEnabled()); +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, GetCustomKnobsReturnsEmpty) +{ + auto fbb = createResampleFwdGraph(ResampleFwdGraphSpec{}); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + + EXPECT_TRUE(_planBuilder.getCustomKnobs(*_handle, graph).empty()); +} From 3bf68b7afeee8da73c837343b3efbdcd953cd2a6 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Mon, 2 Nov 2026 20:00:35 +0000 Subject: [PATCH 6/9] docs(miopen-provider): document standalone resample forward The MIOpen provider's OperationSupport.md, hipDNN's OperationSupport.md, miopen-plugin.rst and plugins.rst list standalone resample forward on the MIOpen provider. The provider notes give what the applicability checks accept: the modes, ranks, layouts and datatypes, the padding rules, the floor-mode output size and the index rules. The deterministic engine section now includes resample forward. Co-Authored-By: Claude Opus 5.5 (1M context) --- .../miopen-provider/docs/OperationSupport.md | 13 ++++++++++++- projects/hipdnn/docs/OperationSupport.md | 1 + .../docs/user-guides/reference/miopen-plugin.rst | 7 ++++++- .../hipdnn/docs/user-guides/reference/plugins.rst | 1 + 4 files changed, 20 insertions(+), 2 deletions(-) diff --git a/dnn-providers/miopen-provider/docs/OperationSupport.md b/dnn-providers/miopen-provider/docs/OperationSupport.md index e61df76971ab..76e8e1c3367b 100644 --- a/dnn-providers/miopen-provider/docs/OperationSupport.md +++ b/dnn-providers/miopen-provider/docs/OperationSupport.md @@ -24,6 +24,7 @@ The following table lists all operations currently supported in hipDNN: | Convolution Wgrad | FP16, BFP16, FP32 | NCL, NLC, NCHW, NHWC, NCDHW, NDHWC | Cross-correlation only2,6, Deterministic5 | | Pointwise Activation (standalone) | FP16, FP32 | NCHW, NHWC | Single-node graph7 | | Pointwise Binary Elementwise (standalone) | FP16, FP32 | NCHW, NCDHW | Single-node graph8 | +| Resample Forward (pooling, standalone) | FP16, BFP16, FP32 | NCHW, NHWC, NCDHW, NDHWC | Single-node graph9, Deterministic5 | ¹ See Batchnorm Operations note below ² See Convolution Operations note below @@ -33,6 +34,7 @@ The following table lists all operations currently supported in hipDNN: ⁶ 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 +⁹ See Standalone Resample Forward note below ## Detailed Requirements @@ -86,6 +88,15 @@ The following table lists all operations currently supported in hipDNN: > must share one datatype. SUB is executed internally as Add with `in_1` scaled by −1. In-place > execution is not supported. +> [!NOTE] +> **Standalone Resample Forward:** A single-node graph supports MAXPOOL, AVGPOOL_EXCLUDE_PADDING +> and AVGPOOL_INCLUDE_PADDING across ranks 4-5, with packed X and Y sharing one datatype and +> layout, and an FP32 compute type. Padding must be symmetric and smaller than the window; unless +> it is all zero, MAXPOOL needs NEG_INF_PAD and the average modes need ZERO_PAD. Y must have the +> floor-mode output size. The index output (`generate_index`) is supported for MAXPOOL in +> NCHW/NCDHW only, as INT8 or INT32; with INT8 the input spatial size may not exceed 128. Very +> large windows on channels-last tensors are not supported. Max pooling ignores NaN inputs. + > [!NOTE] > **Sparse Support:** All operations currently work with dense tensors only. @@ -94,7 +105,7 @@ The following table lists all operations currently supported in hipDNN: ## Deterministic Engine Support -The MIOpen provider offers a deterministic execution engine (`MIOPEN_ENGINE_DETERMINISTIC`) for convolution, unary activation, and binary pointwise operations. This engine guarantees bit-reproducible results across multiple executions with the same inputs. +The MIOpen provider offers a deterministic execution engine (`MIOPEN_ENGINE_DETERMINISTIC`) for convolution, unary activation, binary pointwise and resample forward (pooling) operations. This engine guarantees bit-reproducible results across multiple executions with the same inputs. To use the deterministic engine, set it as the preferred engine on your graph before building: diff --git a/projects/hipdnn/docs/OperationSupport.md b/projects/hipdnn/docs/OperationSupport.md index afe2019fa33b..befaa000bbb9 100644 --- a/projects/hipdnn/docs/OperationSupport.md +++ b/projects/hipdnn/docs/OperationSupport.md @@ -21,6 +21,7 @@ hipDNN operations are implemented through plugins. Each plugin provides its own - Fused operation graphs - Standalone pointwise activation - Standalone pointwise binary elementwise (ADD, SUB, MUL, MAX, MIN) + - Standalone resample forward - **[hipBLASLt Provider Plugin](../../../dnn-providers/hipblaslt-provider/docs/OperationSupport.md)** - Integration with AMD's hipBLASLt library that provides optimized GEMM operations. - Stand-alone Matmul (FP32/FP16/BF16 inputs and outputs, with optional fused bias and activation) diff --git a/projects/hipdnn/docs/user-guides/reference/miopen-plugin.rst b/projects/hipdnn/docs/user-guides/reference/miopen-plugin.rst index c821a3512798..00a87ee49a6a 100644 --- a/projects/hipdnn/docs/user-guides/reference/miopen-plugin.rst +++ b/projects/hipdnn/docs/user-guides/reference/miopen-plugin.rst @@ -85,10 +85,14 @@ This table lists all operations supported in hipDNN: - ``FP16``, ``FP32`` - NCHW, NCDHW - Single-node graph⁷ + * - Resample Forward (pooling, standalone) + - ``FP16``, ``BFP16``, ``FP32`` + - NCHW, NHWC, NCDHW, NDHWC + - Single-node graph⁸ .. note:: - - For annotations ¹ through ⁴, ⁶ and ⁷, see :ref:`operations`. + - For annotations ¹ through ⁴ and ⁶ through ⁸, see :ref:`operations`. - For annotation ⁵, see :ref:`detailed` for more information. .. _detailed: @@ -146,6 +150,7 @@ Operation notes - **Activation Functions (fused)**: Supports ReLU, Clipped ReLU (with configurable upper clip), and CLAMP (with configurable lower/upper clips). Leaky ReLU, Sigmoid and Tanh are not supported when fused onto a producer op; use a standalone activation node for those. - **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. - ⁷ **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. In-place execution (``out_0`` aliasing an input) isn't supported. +- ⁸ **Standalone Resample Forward**: A single-node graph supports ``MAXPOOL``, ``AVGPOOL_EXCLUDE_PADDING`` and ``AVGPOOL_INCLUDE_PADDING`` across ranks 4-5, with packed X and Y sharing one datatype and layout, and an ``FP32`` compute type. Padding must be symmetric and smaller than the window; unless it is all zero, ``MAXPOOL`` needs ``NEG_INF_PAD`` and the average modes need ``ZERO_PAD``. Y must have the floor-mode output size. The index output (``generate_index``) is supported for ``MAXPOOL`` in NCHW/NCDHW only, as ``INT8`` or ``INT32``; with ``INT8`` the input spatial size may not exceed 128. Very large windows on channels-last tensors aren't supported. Max pooling ignores NaN inputs. - **Sparse Support**: All operations only work with dense tensors. Knobs diff --git a/projects/hipdnn/docs/user-guides/reference/plugins.rst b/projects/hipdnn/docs/user-guides/reference/plugins.rst index ebadc73ef636..c1771fa2432f 100644 --- a/projects/hipdnn/docs/user-guides/reference/plugins.rst +++ b/projects/hipdnn/docs/user-guides/reference/plugins.rst @@ -19,6 +19,7 @@ Plugins - Batchnorm operations (Training, Backward, Inference) - Fused operation graphs - Standalone pointwise activation and binary elementwise operations + - Standalone resample forward - :ref:`hipblaslt`: Provides integration with AMD's `hipBLASLt library `_ that provides optimized GEMM operations. From d9450a4f1fffa6c36ed06fe9466950f42bb9b0db Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Tue, 3 Nov 2026 13:17:08 +0000 Subject: [PATCH 7/9] test(miopen-provider): claim ResampleFwd bundles on MIOPEN_ENGINE MIOPEN_ENGINE can run few of the existing ResampleFwd cases: Variant2 uses asymmetric padding, and Default is max pooling with an index. The new quick bundle ResampleFwd/SymmetricPadding adds 47 cases covering all three modes, four layouts and three datatypes. MIOPEN_ENGINE claims those cases and the Default cases it runs in every tier, on gfx90a, gfx942 and gfx950 (Linux). One quick Default case joins ffm-quick. SUPPORT_MATRIX_OVERVIEW.md is regenerated. Co-Authored-By: Claude Opus 5.5 --- .../SUPPORT_MATRIX_OVERVIEW.md | 9 +- .../ResampleFwd/Default/support.json | 20 + .../full/ResampleFwd/Default/support.json | 19 + .../quick/ResampleFwd/Default/support.json | 20 + .../SymmetricPadding/graph.template.json | 52 + .../ResampleFwd/SymmetricPadding/support.json | 69 + .../ResampleFwd/SymmetricPadding/sweep.json | 3394 +++++++++++++++++ .../standard/ResampleFwd/Default/support.json | 18 + .../test_categories_integration.yaml | 3 + 9 files changed, 3600 insertions(+), 4 deletions(-) create mode 100644 dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/graph.template.json create mode 100644 dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/support.json create mode 100644 dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/sweep.json diff --git a/dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md b/dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md index a3e0b4937690..88a6be90638c 100644 --- a/dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md +++ b/dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md @@ -2,7 +2,7 @@ -Generated from the committed support claims (11174 graphs, 486 catalogued options). Percentages are shares of each op's API options, not of its graphs. Options count singly, not in combination. +Generated from the committed support claims (11221 graphs, 486 catalogued options). Percentages are shares of each op's API options, not of its graphs. Options count singly, not in combination. **Jump to:** [RDNA (consumer)](#rdna-consumer) · [CDNA (data center)](#cdna-data-center) · [By engine](#by-engine) · [Op groups](#op-groups) · [Options per op](#options-per-op) · [GPU targets](#gpu-targets) @@ -43,7 +43,7 @@ Generated from the committed support claims (11174 graphs, 486 catalogued option | **[Matmul](#group-matmul)** (35) | 🟡 37% | 🟡 37% | 🟡 37% | | **[Pointwise](#group-pointwise)** (58) | 🟡 20% | 🟡 20% | 🟡 20% | | **[Reduction](#group-reduction)** (26) | ❌ | ❌ | ❌ | -| **[Resample](#group-resample)** (34) | 🟡 47% | 🟡 47% | 🟡 47% | +| **[Resample](#group-resample)** (34) | 🟡 52% | 🟡 52% | 🟡 52% | | **[RMSNorm](#group-rmsnorm)** (38) | 🟡 86% | 🟡 86% | 🟡 86% | ## By engine @@ -95,6 +95,7 @@ An op appears under an engine that claims at least one of its graphs somewhere. | **[ConvolutionFwd](#op-convolutionfwd)** (17) | 🟡 76% | 🟡 76% | 🟡 76% | 🟡 76% | ❌ | 🟡 76% | ❌ | 🟡 76% | 🟡 76% | 🟡 76% | ❌ | 🟡 76% | | **[ConvolutionWrw](#op-convolutionwrw)** (17) | 🟡 76% | 🟡 76% | 🟡 76% | 🟡 76% | ❌ | 🟡 76% | ❌ | 🟡 76% | 🟡 76% | 🟡 76% | ❌ | 🟡 76% | | **[Pointwise](#op-pointwise)** (58) | 🟡 13% | 🟡 17% | 🟡 17% | 🟡 17% | ❌ | 🟡 17% | ❌ | 🟡 17% | 🟡 17% | 🟡 17% | ❌ | 🟡 17% | +| **[ResampleFwd](#op-resamplefwd)** (19) | 🟡 84% | 🟡 84% | 🟡 84% | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ### MIOPEN_ENGINE_DETERMINISTIC @@ -127,7 +128,7 @@ The API node types (`Attributes` in the flatbuffers schema) that make up eac | **Matmul** | 459 | 35 | [Matmul](#op-matmul) (16), [MoeGroupedMatmul](#op-moegroupedmatmul) (13, **no test graph**), [MoeGroupedMatmulBwd](#op-moegroupedmatmulbwd) (6, **no test graph**) | | **Pointwise** | 6162 | 58 | [Pointwise](#op-pointwise) (58) | | **Reduction** | 114 | 26 | [Reduction](#op-reduction) (26) | -| **Resample** | 48 | 34 | [ResampleBwd](#op-resamplebwd) (15, **no test graph**), [ResampleFwd](#op-resamplefwd) (19) | +| **Resample** | 95 | 34 | [ResampleBwd](#op-resamplebwd) (15, **no test graph**), [ResampleFwd](#op-resamplefwd) (19) | | **RMSNorm** | 5008 | 38 | [RMSNorm](#op-rmsnorm) (21), [RMSNormBackward](#op-rmsnormbackward) (17) | ## Options per op @@ -156,7 +157,7 @@ Every API option the percentages count, catalogued per op in `dnn-providers/inte | **Pointwise** (58) | `mode:gelu_approx_tanh_fwd`, `mode:relu_fwd`, `mode:sigmoid_fwd`, `mode:swish_fwd`, `mode:tanh_fwd`, `mode:add`, `mode:max_op`, `mode:min_op`, `mode:mul`, `mode:relu_bwd`, `mode:sub`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `precision:mixed`, `broadcast:in`, `relu:lower_clip`, `relu:upper_clip`, `relu:leaky` | `mode:abs`, `mode:ceil`, `mode:elu_fwd`, `mode:erf`, `mode:exp`, `mode:floor`, `mode:gelu_fwd`, `mode:gen_index`, `mode:identity`, `mode:log`, `mode:logical_not`, `mode:neg`, `mode:reciprocal`, `mode:rsqrt`, `mode:sin`, `mode:softplus_fwd`, `mode:sqrt`, `mode:tan`, `mode:add_square`, `mode:cmp_eq`, `mode:cmp_ge`, `mode:cmp_gt`, `mode:cmp_le`, `mode:cmp_lt`, `mode:cmp_neq`, `mode:div`, `mode:elu_bwd`, `mode:gelu_approx_tanh_bwd`, `mode:gelu_bwd`, `mode:logical_and`, `mode:logical_or`, `mode:sigmoid_bwd`, `mode:softplus_bwd`, `mode:swish_bwd`, `mode:tanh_bwd`, `mode:binary_select`, `swish:beta`, `elu:alpha`, `softplus:beta` | | **Reduction** (26) | `mode:add`, `mode:mul`, `mode:min`, `mode:max`, `mode:amax`, `mode:avg`, `mode:norm1`, `mode:norm2`, `mode:mul_no_zeros`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `dims:4`, `reduce:all`, `reduce:trailing`, `reduce:leading`, `layout:nchw`, `layout:nhwc`, `deterministic:off` | `precision:mixed_output`, `dims:1`, `dims:2`, `dims:3`, `dims:5`, `reduce:interleaved`, `deterministic:on` | | **ResampleBwd** (15) | — | `mode:maxpool`, `mode:avgpool_exclude_padding`, `mode:avgpool_include_padding`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `spatial:2d`, `spatial:3d`, `layout:nchw`, `layout:nhwc`, `padding:symmetric`, `padding:asymmetric`, `stride:strided`, `window:overlapping`, `window:gapped` | -| **ResampleFwd** (19) | `mode:maxpool`, `mode:avgpool_exclude_padding`, `mode:avgpool_include_padding`, `padding_mode:zero_pad`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `spatial:2d`, `spatial:3d`, `layout:nchw`, `layout:nhwc`, `padding:symmetric`, `padding:asymmetric`, `stride:strided`, `window:overlapping`, `index:generated` | `padding_mode:neg_inf_pad`, `window:gapped`, `index:none` | +| **ResampleFwd** (19) | `mode:maxpool`, `mode:avgpool_exclude_padding`, `mode:avgpool_include_padding`, `padding_mode:neg_inf_pad`, `padding_mode:zero_pad`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `spatial:2d`, `spatial:3d`, `layout:nchw`, `layout:nhwc`, `padding:symmetric`, `padding:asymmetric`, `stride:strided`, `window:overlapping`, `index:generated`, `index:none` | `window:gapped` | | **RMSNorm** (21) | `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `precision:mixed_io`, `precision:mixed_affine`, `precision:mixed_stats`, `dims:4`, `dims:5`, `layout:nchw`, `layout:nhwc`, `norm:last_dim`, `norm:multi_dim`, `norm:all_but_batch`, `bias:none`, `bias:bound`, `phase:inference`, `phase:training`, `epsilon:fixed`, `epsilon:runtime` | `dims:2`, `dims:3` | | **RMSNormBackward** (17) | `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `precision:mixed_io`, `precision:mixed_affine`, `precision:mixed_stats`, `dims:4`, `dims:5`, `layout:nchw`, `layout:nhwc`, `norm:last_dim`, `norm:multi_dim`, `norm:all_but_batch`, `dbias:none`, `dbias:bound` | `dims:2`, `dims:3` | | **Sdpa** (36) | `dtype:bf16`, `dtype:fp16`, `dtype:fp8e4m3`, `head_dim:64`, `head_dim:128`, `head_dim:192x128`, `heads:mha`, `heads:gqa`, `heads:mqa`, `layout:bhsd`, `layout:bshd`, `layout:ragged`, `mask:none`, `mask:causal_top_left`, `mask:causal_bottom_right`, `seq:q_eq_kv`, `seq:q_ne_kv`, `varlen:seq_lens`, `scale:fixed`, `scale:runtime`, `out:stats`, `precision:mixed_output` | `dtype:fp8e4m3_fnuz`, `head_dim:192`, `mask:sliding_window`, `seq:decode`, `scale:default`, `out:max_sum_exp`, `out:rng_dump`, `fp8:amax`, `bias:additive`, `bias:alibi`, `dropout:on`, `kv:paged`, `input:block_mask`, `input:sink_token` | diff --git a/dnn-providers/integration-tests/integration-test-bundles/comprehensive/ResampleFwd/Default/support.json b/dnn-providers/integration-tests/integration-test-bundles/comprehensive/ResampleFwd/Default/support.json index d4e6a9881ce7..d8c6c18cfa8b 100644 --- a/dnn-providers/integration-tests/integration-test-bundles/comprehensive/ResampleFwd/Default/support.json +++ b/dnn-providers/integration-tests/integration-test-bundles/comprehensive/ResampleFwd/Default/support.json @@ -38,6 +38,26 @@ ] } } + ], + "MIOPEN_ENGINE": [ + { + "cases": [ + "1_2_3_2_2_bfp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_4_3_bfp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "2_3_3_2_fp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0" + ], + "support": { + "gfx90a": [ + "linux" + ], + "gfx942": [ + "linux" + ], + "gfx950": [ + "linux" + ] + } + } ] }, "version": 1 diff --git a/dnn-providers/integration-tests/integration-test-bundles/full/ResampleFwd/Default/support.json b/dnn-providers/integration-tests/integration-test-bundles/full/ResampleFwd/Default/support.json index babb41a2cd77..0a9868c231a5 100644 --- a/dnn-providers/integration-tests/integration-test-bundles/full/ResampleFwd/Default/support.json +++ b/dnn-providers/integration-tests/integration-test-bundles/full/ResampleFwd/Default/support.json @@ -40,6 +40,25 @@ ] } } + ], + "MIOPEN_ENGINE": [ + { + "cases": [ + "1_2_3_2_2_fp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "2_3_3_2_bfp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0" + ], + "support": { + "gfx90a": [ + "linux" + ], + "gfx942": [ + "linux" + ], + "gfx950": [ + "linux" + ] + } + } ] }, "version": 1 diff --git a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/Default/support.json b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/Default/support.json index c80f6f85403b..c7b7cffeb657 100644 --- a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/Default/support.json +++ b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/Default/support.json @@ -38,6 +38,26 @@ ] } } + ], + "MIOPEN_ENGINE": [ + { + "cases": [ + "1_2_3_2_2_fp32_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_4_3_fp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "1_2_4_3_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0" + ], + "support": { + "gfx90a": [ + "linux" + ], + "gfx942": [ + "linux" + ], + "gfx950": [ + "linux" + ] + } + } ] }, "version": 1 diff --git a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/graph.template.json b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/graph.template.json new file mode 100644 index 000000000000..2b94e8877255 --- /dev/null +++ b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/graph.template.json @@ -0,0 +1,52 @@ +{ + "compute_data_type": "float", + "intermediate_data_type": "float", + "io_data_type": "${case.io_data_type}", + "is_override_shape_enabled": false, + "min_required_engine_api_version": { + "major": 1, + "minor": 0, + "patch": 0 + }, + "name": "ResampleFwdTest", + "nodes": [ + { + "compute_data_type": "float", + "generate_index": false, + "name": "ResampleFwd_0", + "padding_mode": "${case.attributes.padding_mode}", + "post_padding": "${case.attributes.post_padding}", + "pre_padding": "${case.attributes.pre_padding}", + "resample_mode": "${case.attributes.resample_mode}", + "stride": "${case.attributes.stride}", + "type": "ResampleFwdAttributes", + "window": "${case.attributes.window}", + "x_tensor_uid": 1, + "y_tensor_uid": 0 + } + ], + "tensors": [ + { + "alignment": 16, + "data_type": "${case.data_type}", + "dims": "${case.dims}", + "is_runtime_pass_by_value": false, + "name": "ResampleFwd_0::Y", + "ragged_offset_multiplier": 1, + "strides": "${case.strides}", + "uid": 0, + "virtual": false + }, + { + "alignment": 16, + "data_type": "${case.data_type}", + "dims": "${case.dims}", + "is_runtime_pass_by_value": false, + "name": "X", + "ragged_offset_multiplier": 1, + "strides": "${case.strides}", + "uid": 1, + "virtual": false + } + ] +} diff --git a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/support.json b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/support.json new file mode 100644 index 000000000000..3792c3e454c7 --- /dev/null +++ b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/support.json @@ -0,0 +1,69 @@ +{ + "claims": { + "MIOPEN_ENGINE": [ + { + "cases": [ + "1_2_16_16_16_fp32_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_2_3_2_bfp16_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "1_2_2_3_2_bfp16_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "1_2_2_3_2_fp16_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "1_2_2_3_2_fp16_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "1_2_2_3_2_fp32_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "1_2_2_3_2_fp32_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "1_2_3_2_2_bfp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_3_2_2_bfp16_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_3_2_2_fp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_3_2_2_fp16_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_3_2_2_fp32_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_3_2_2_fp32_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "1_2_5_2_4_bfp16_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "1_2_5_2_4_bfp16_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "1_2_5_2_4_fp16_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "1_2_5_2_4_fp16_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "1_2_5_2_4_fp32_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "1_2_5_2_4_fp32_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "1_2_64_64_fp32_nhwc_padding_modezero_pad_postpad0x0_prepad0x0", + "2_3_3_2_bfp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "2_3_3_2_bfp16_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "2_3_3_2_fp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "2_3_3_2_fp16_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "2_3_3_2_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "2_3_3_2_fp32_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "2_3_4_3_bfp16_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "2_3_4_3_bfp16_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "2_3_4_3_bfp16_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "2_3_4_3_bfp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "2_3_4_3_fp16_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "2_3_4_3_fp16_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "2_3_4_3_fp16_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "2_3_4_3_fp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "2_3_4_3_fp32_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "2_3_4_3_fp32_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "2_3_4_3_fp32_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "2_3_4_3_fp32_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "2_3_4_4_bfp16_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "2_3_4_4_bfp16_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "2_3_4_4_fp16_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "2_3_4_4_fp16_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "2_3_4_4_fp32_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "2_3_4_4_fp32_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "2_8_17_16_bfp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "2_8_17_16_fp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "2_8_17_16_fp32_nhwc_padding_modezero_pad_postpad1x1_prepad1x1" + ], + "support": { + "gfx90a": [ + "linux" + ], + "gfx942": [ + "linux" + ], + "gfx950": [ + "linux" + ] + } + } + ] + }, + "version": 1 +} diff --git a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/sweep.json b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/sweep.json new file mode 100644 index 000000000000..98a4f501f214 --- /dev/null +++ b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/sweep.json @@ -0,0 +1,3394 @@ +{ + "version": 1, + "cases": [ + { + "id": "2_3_4_3_bfp16_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_8_17_16_bfp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 8, + 17, + 16 + ], + "strides": [ + 2176, + 1, + 128, + 8 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 8, + 33, + 31 + ], + "strides": [ + 8184, + 1, + 248, + 8 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_2d_nhwc_large_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_2d_nhwc_large_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_bfp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_bfp16_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 12, + 6, + 2, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_3d_ncdhw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_3d_ncdhw_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_bfp16_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 1, + 12, + 4, + 2 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_3d_ndhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_3d_ndhwc_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_4_bfp16_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 16, + 4, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_2d_nchw_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_2d_nchw_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "2_3_4_4_bfp16_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 1, + 12, + 3 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_2d_nhwc_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_2d_nhwc_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "1_2_5_2_4_bfp16_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 40, + 8, + 4, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "1_2_5_2_4_bfp16_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 1, + 16, + 8, + 2 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_bfp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_bfp16_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_bfp16_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_3_2_2_bfp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 12, + 4, + 2, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_3d_ncdhw_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_3d_ncdhw_pad0_win2_stride1x2x1" + } + }, + { + "id": "1_2_3_2_2_bfp16_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 1, + 8, + 4, + 2 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_3d_ndhwc_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_3d_ndhwc_pad0_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_bfp16_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_zero_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_zero_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_fp16_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_8_17_16_fp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 8, + 17, + 16 + ], + "strides": [ + 2176, + 1, + 128, + 8 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 8, + 33, + 31 + ], + "strides": [ + 8184, + 1, + 248, + 8 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_2d_nhwc_large_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_2d_nhwc_large_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_fp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_fp16_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 12, + 6, + 2, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_3d_ncdhw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_3d_ncdhw_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_fp16_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 1, + 12, + 4, + 2 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_3d_ndhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_3d_ndhwc_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_4_fp16_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 16, + 4, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_2d_nchw_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_2d_nchw_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "2_3_4_4_fp16_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 1, + 12, + 3 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_2d_nhwc_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_2d_nhwc_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "1_2_5_2_4_fp16_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 40, + 8, + 4, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "1_2_5_2_4_fp16_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 1, + 16, + 8, + 2 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_fp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_fp16_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_fp16_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_3_2_2_fp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 12, + 4, + 2, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_3d_ncdhw_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_3d_ncdhw_pad0_win2_stride1x2x1" + } + }, + { + "id": "1_2_3_2_2_fp16_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 1, + 8, + 4, + 2 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_3d_ndhwc_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_3d_ndhwc_pad0_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_fp16_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_zero_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_zero_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_fp32_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_8_17_16_fp32_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 8, + 17, + 16 + ], + "strides": [ + 2176, + 1, + 128, + 8 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 8, + 33, + 31 + ], + "strides": [ + 8184, + 1, + 248, + 8 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_2d_nhwc_large_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_2d_nhwc_large_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_fp32_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_fp32_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 12, + 6, + 2, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_3d_ncdhw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_3d_ncdhw_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_fp32_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 1, + 12, + 4, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_3d_ndhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_3d_ndhwc_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_4_fp32_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 16, + 4, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_2d_nchw_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_2d_nchw_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "2_3_4_4_fp32_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 1, + 12, + 3 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_2d_nhwc_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_2d_nhwc_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "1_2_5_2_4_fp32_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 40, + 8, + 4, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "1_2_5_2_4_fp32_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 1, + 16, + 8, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_fp32_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_fp32_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_3_2_2_fp32_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 12, + 4, + 2, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_3d_ncdhw_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_3d_ncdhw_pad0_win2_stride1x2x1" + } + }, + { + "id": "1_2_3_2_2_fp32_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 1, + 8, + 4, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_3d_ndhwc_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_3d_ndhwc_pad0_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_fp32_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_zero_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_zero_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "1_2_64_64_fp32_nhwc_padding_modezero_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 1, + 1 + ], + "strides": [ + 2, + 1, + 2, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 64, + 64 + ], + "strides": [ + 8192, + 1, + 128, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 64, + 64 + ], + "window": [ + 64, + 64 + ] + } + }, + "metadata": { + "format_version": 1, + "reference_source": "graph logged via HIPDNN_LOG_GRAPH_DIR from miopen-provider Smoke/IntegrationGpuResampleFwdPositiveControls.BuildsGraph/NhwcGlobalAvgPool", + "seed": 1, + "generator": "import_graph.py" + } + }, + { + "id": "1_2_16_16_16_fp32_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 1, + 1, + 1 + ], + "strides": [ + 2, + 1, + 2, + 2, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 16, + 16, + 16 + ], + "strides": [ + 8192, + 1, + 512, + 32, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 16, + 16, + 16 + ], + "window": [ + 16, + 16, + 16 + ] + } + }, + "metadata": { + "format_version": 1, + "reference_source": "graph logged via HIPDNN_LOG_GRAPH_DIR from miopen-provider Smoke/IntegrationGpuResampleFwdPositiveControls.BuildsGraph/NdhwcGlobalMaxPool", + "seed": 1, + "generator": "import_graph.py" + } + } + ] +} diff --git a/dnn-providers/integration-tests/integration-test-bundles/standard/ResampleFwd/Default/support.json b/dnn-providers/integration-tests/integration-test-bundles/standard/ResampleFwd/Default/support.json index 8c8ad3b55983..95be95573405 100644 --- a/dnn-providers/integration-tests/integration-test-bundles/standard/ResampleFwd/Default/support.json +++ b/dnn-providers/integration-tests/integration-test-bundles/standard/ResampleFwd/Default/support.json @@ -36,6 +36,24 @@ ] } } + ], + "MIOPEN_ENGINE": [ + { + "cases": [ + "2_3_3_2_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0" + ], + "support": { + "gfx90a": [ + "linux" + ], + "gfx942": [ + "linux" + ], + "gfx950": [ + "linux" + ] + } + } ] }, "version": 1 diff --git a/dnn-providers/miopen-provider/test_categories_integration.yaml b/dnn-providers/miopen-provider/test_categories_integration.yaml index d65955adb101..f423744c8eba 100644 --- a/dnn-providers/miopen-provider/test_categories_integration.yaml +++ b/dnn-providers/miopen-provider/test_categories_integration.yaml @@ -107,6 +107,9 @@ test_categories: # signal; this is the smallest quick-tier layernorm bundle case. - "quick_Layernorm_Default.2_2_3_2_bfp16_nhwc_normalized_dim_count1_6d6330" - "quick_SdpaFwd_bhsd_bf16_hd128_nomask_batch_Small.Small" + # ResampleFwd: max pooling that also writes an index, so the run goes + # through the index copy out of MIOpen's workspace. + - "quick_ResampleFwd_Default.1_2_4_3_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0" exclude: *common_excludes labels: - "ffm-quick" From effd758df345913711d145d3db5a87abb58d70c3 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Tue, 3 Nov 2026 15:28:05 +0000 Subject: [PATCH 8/9] fix(hipdnn): pass optionals straight through in createResampleFwdGraph flatbuffers::Optional is std::optional here, so the toFlatbuffers lambda unwrapped and rewrapped the same type. clang-tidy rejects that (bugprone-optional-value-conversion) and fails the build. Co-Authored-By: Claude Opus 5.5 txt Co-Authored-By: Claude Opus 5.5 --- .../utilities/FlatbufferGraphTestUtils.hpp | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) diff --git a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp index 682b72b8efc0..b037c1f4fa1b 100644 --- a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp +++ b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp @@ -7,7 +7,6 @@ #include #include #include -#include #include #include @@ -3626,12 +3625,6 @@ inline flatbuffers::FlatBufferBuilder createResampleFwdGraph(const ResampleFwdGr { namespace data_objects = hipdnn_flatbuffers_sdk::data_objects; - const auto toFlatbuffers = [](const auto& value) { - using T = typename std::decay_t::value_type; - return value.has_value() ? ::flatbuffers::Optional(value.value()) - : ::flatbuffers::Optional(::flatbuffers::nullopt); - }; - flatbuffers::FlatBufferBuilder builder; std::vector> tensorAttributes; std::vector emittedUids; @@ -3654,7 +3647,7 @@ inline flatbuffers::FlatBufferBuilder createResampleFwdGraph(const ResampleFwdGr data_objects::TensorValue::NONE, 0, tensor.isRuntimePassByValue, - toFlatbuffers(tensor.raggedOffsetTensorUid))); + tensor.raggedOffsetTensorUid)); }; addTensor(spec.x, "x"); @@ -3685,7 +3678,7 @@ inline flatbuffers::FlatBufferBuilder createResampleFwdGraph(const ResampleFwdGr optionalVector(spec.window), spec.mode, spec.paddingMode, - toFlatbuffers(spec.generateIndex)); + spec.generateIndex); const auto name = "resample_fwd_" + std::to_string(i); nodes.push_back( From 997b9114f21b5f04a66d41fc78f651fb30b2a864 Mon Sep 17 00:00:00 2001 From: Evgenii Averin <86725875+averinevg@users.noreply.github.com> Date: Tue, 3 Nov 2026 19:34:44 +0000 Subject: [PATCH 9/9] fix(miopen-provider): drop unused ResampleMode using in plan builder test clang-tidy misc-unused-using-decls flagged it as an error in CI. Co-Authored-By: Claude Opus 5.5 --- .../tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp index b0b241cb51c4..85e4c3b33819 100644 --- a/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp +++ b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp @@ -21,7 +21,6 @@ using namespace miopen_plugin; -using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; using hipdnn_test_sdk::utilities::createResampleFwdGraph; using hipdnn_test_sdk::utilities::MockEngineConfig;