From 05110a65586a86f5b671a850b93dc0434589af43 Mon Sep 17 00:00:00 2001 From: Du Li Date: Fri, 31 May 2019 11:48:43 -0700 Subject: [PATCH] Adding custom op ConvTransposeWithDynamicPads. (#638) * Adding custom op ConvTransposeWithDynamicPads. * Adding custom op ConvTransposeWithDynamicPads. * adding cuda kernels * fix a bug * fix build issue. * Integrate PR comments. --- cmake/onnxruntime_providers.cmake | 8 ++- .../onnxruntime/core/framework/op_kernel.h | 7 +++ onnxruntime/contrib_ops/contrib_kernels.cc | 3 + onnxruntime/contrib_ops/contrib_kernels.h | 5 +- .../cpu/conv_transpose_with_dynamic_pads.cc | 16 +++++ .../cpu/conv_transpose_with_dynamic_pads.h | 20 ++++++ .../cuda/conv_transpose_with_dynamic_pads.cc | 19 ++++++ .../cuda/conv_transpose_with_dynamic_pads.h | 23 +++++++ .../core/graph/contrib_ops/contrib_defs.cc | 57 ++++++++++++++++- .../core/providers/cpu/nn/conv_transpose.cc | 61 ++++++++++++------- .../core/providers/cpu/nn/conv_transpose.h | 5 +- .../providers/cuda/contrib_kernels_cuda.cc | 17 ++++++ .../providers/cuda/cuda_execution_provider.cc | 3 + .../core/providers/cuda/nn/conv_transpose.cc | 11 +++- .../core/providers/cuda/nn/conv_transpose.h | 1 + .../conv_transpose_with_dynamic_pads_test.cc | 23 +++++++ 16 files changed, 248 insertions(+), 31 deletions(-) create mode 100644 onnxruntime/contrib_ops/cpu/conv_transpose_with_dynamic_pads.cc create mode 100644 onnxruntime/contrib_ops/cpu/conv_transpose_with_dynamic_pads.h create mode 100644 onnxruntime/contrib_ops/cuda/conv_transpose_with_dynamic_pads.cc create mode 100644 onnxruntime/contrib_ops/cuda/conv_transpose_with_dynamic_pads.h create mode 100644 onnxruntime/core/providers/cuda/contrib_kernels_cuda.cc create mode 100644 onnxruntime/test/contrib_ops/conv_transpose_with_dynamic_pads_test.cc diff --git a/cmake/onnxruntime_providers.cmake b/cmake/onnxruntime_providers.cmake index 3b12b78bcd..319514d203 100644 --- a/cmake/onnxruntime_providers.cmake +++ b/cmake/onnxruntime_providers.cmake @@ -8,9 +8,11 @@ file(GLOB_RECURSE onnxruntime_providers_srcs CONFIGURE_DEPENDS file(GLOB_RECURSE onnxruntime_contrib_ops_srcs CONFIGURE_DEPENDS "${ONNXRUNTIME_ROOT}/contrib_ops/*.h" - "${ONNXRUNTIME_ROOT}/contrib_ops/*.cc" + "${ONNXRUNTIME_ROOT}/contrib_ops/contrib_kernels.cc" "${ONNXRUNTIME_ROOT}/contrib_ops/cpu/*.h" "${ONNXRUNTIME_ROOT}/contrib_ops/cpu/*.cc" + "${ONNXRUNTIME_ROOT}/contrib_ops/cpu/attnlstm/*.h" + "${ONNXRUNTIME_ROOT}/contrib_ops/cpu/attnlstm/*.cc" ) file(GLOB onnxruntime_providers_common_srcs CONFIGURE_DEPENDS @@ -58,12 +60,14 @@ if (onnxruntime_USE_CUDA) file(GLOB_RECURSE onnxruntime_providers_cuda_cc_srcs CONFIGURE_DEPENDS "${ONNXRUNTIME_ROOT}/core/providers/cuda/*.h" "${ONNXRUNTIME_ROOT}/core/providers/cuda/*.cc" + "${ONNXRUNTIME_ROOT}/contrib_ops/cuda/*.h" + "${ONNXRUNTIME_ROOT}/contrib_ops/cuda/*.cc" ) file(GLOB_RECURSE onnxruntime_providers_cuda_cu_srcs CONFIGURE_DEPENDS "${ONNXRUNTIME_ROOT}/core/providers/cuda/*.cu" "${ONNXRUNTIME_ROOT}/core/providers/cuda/*.cuh" ) - source_group(TREE ${ONNXRUNTIME_ROOT}/core FILES ${onnxruntime_providers_cuda_cc_srcs} ${onnxruntime_providers_cuda_cu_srcs}) + source_group(TREE ${ONNXRUNTIME_ROOT} FILES ${onnxruntime_providers_cuda_cc_srcs} ${onnxruntime_providers_cuda_cu_srcs}) add_library(onnxruntime_providers_cuda ${onnxruntime_providers_cuda_cc_srcs} ${onnxruntime_providers_cuda_cu_srcs}) if (UNIX) target_compile_options(onnxruntime_providers_cuda PRIVATE "$<$:SHELL:-Xcompiler -Wno-reorder>" diff --git a/include/onnxruntime/core/framework/op_kernel.h b/include/onnxruntime/core/framework/op_kernel.h index ccfe13d79c..e41896adec 100644 --- a/include/onnxruntime/core/framework/op_kernel.h +++ b/include/onnxruntime/core/framework/op_kernel.h @@ -221,6 +221,13 @@ template KernelCreateInfo BuildKernelCreateInfo(); } // namespace contrib +namespace contrib { +namespace cuda { +template +KernelCreateInfo BuildKernelCreateInfo(); +} // namespace cuda +} // namespace contrib + using BuildKernelCreateInfoFn = KernelCreateInfo (*)(); // Naming convention for operator kernel classes diff --git a/onnxruntime/contrib_ops/contrib_kernels.cc b/onnxruntime/contrib_ops/contrib_kernels.cc index f190c3344c..ea40cde388 100644 --- a/onnxruntime/contrib_ops/contrib_kernels.cc +++ b/onnxruntime/contrib_ops/contrib_kernels.cc @@ -20,6 +20,7 @@ class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, Murmu class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, float, MaxpoolWithMask); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, Pad); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, Unique); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, float, ConvTransposeWithDynamicPads); // This section includes all opkernel declarations for former experimental ops which have now been removed from onnx. // To maintain backward compatibility these are added as contrib ops. @@ -64,6 +65,8 @@ void RegisterContribKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + // These ops were experimental ops in onnx domain which have been removed now. We add them here as // contrib ops to main backward compatibility BuildKernelCreateInfo, diff --git a/onnxruntime/contrib_ops/contrib_kernels.h b/onnxruntime/contrib_ops/contrib_kernels.h index 48a7d53614..d53b11ba6e 100644 --- a/onnxruntime/contrib_ops/contrib_kernels.h +++ b/onnxruntime/contrib_ops/contrib_kernels.h @@ -9,5 +9,8 @@ namespace onnxruntime { namespace contrib { void RegisterContribKernels(KernelRegistry& kernel_registry); -} // namespace contrib +namespace cuda { +void RegisterCudaContribKernels(KernelRegistry& kernel_registry); +} +} // namespace contrib } // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cpu/conv_transpose_with_dynamic_pads.cc b/onnxruntime/contrib_ops/cpu/conv_transpose_with_dynamic_pads.cc new file mode 100644 index 0000000000..dc603f4778 --- /dev/null +++ b/onnxruntime/contrib_ops/cpu/conv_transpose_with_dynamic_pads.cc @@ -0,0 +1,16 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "conv_transpose_with_dynamic_pads.h" + +namespace onnxruntime { +namespace contrib { +ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( + ConvTransposeWithDynamicPads, + 1, + float, + KernelDefBuilder() + .TypeConstraint("T", DataTypeImpl::GetTensorType()), + ConvTransposeWithDynamicPads); +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cpu/conv_transpose_with_dynamic_pads.h b/onnxruntime/contrib_ops/cpu/conv_transpose_with_dynamic_pads.h new file mode 100644 index 0000000000..4bcb5fba93 --- /dev/null +++ b/onnxruntime/contrib_ops/cpu/conv_transpose_with_dynamic_pads.h @@ -0,0 +1,20 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include "core/providers/cpu/nn/conv_transpose.h" + +namespace onnxruntime { +namespace contrib { +template +class ConvTransposeWithDynamicPads : public ConvTranspose { + public: + ConvTransposeWithDynamicPads(const OpKernelInfo& info) : ConvTranspose(info) {} + + Status Compute(OpKernelContext* context) const override { + return ConvTranspose::DoConvTranspose(context, true); + } +}; +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cuda/conv_transpose_with_dynamic_pads.cc b/onnxruntime/contrib_ops/cuda/conv_transpose_with_dynamic_pads.cc new file mode 100644 index 0000000000..e526c23e64 --- /dev/null +++ b/onnxruntime/contrib_ops/cuda/conv_transpose_with_dynamic_pads.cc @@ -0,0 +1,19 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "contrib_ops/cuda/conv_transpose_with_dynamic_pads.h" + +namespace onnxruntime { +namespace contrib { +namespace cuda { +ONNX_OPERATOR_TYPED_KERNEL_EX( + ConvTransposeWithDynamicPads, + kMSDomain, + 1, + float, + kCudaExecutionProvider, + KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()).InputMemoryType(2), + ConvTransposeWithDynamicPads); +} // namespace cuda +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cuda/conv_transpose_with_dynamic_pads.h b/onnxruntime/contrib_ops/cuda/conv_transpose_with_dynamic_pads.h new file mode 100644 index 0000000000..6f7a04d059 --- /dev/null +++ b/onnxruntime/contrib_ops/cuda/conv_transpose_with_dynamic_pads.h @@ -0,0 +1,23 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include "core/providers/cuda/nn/conv_transpose.h" + +namespace onnxruntime { +namespace contrib { +namespace cuda { + +template +class ConvTransposeWithDynamicPads : public ::onnxruntime::cuda::ConvTranspose { + public: + ConvTransposeWithDynamicPads(const OpKernelInfo& info) : ::onnxruntime::cuda::ConvTranspose(info) {} + + Status ComputeInternal(OpKernelContext* context) const override { + return ::onnxruntime::cuda::ConvTranspose::DoConvTranspose(context, true); + } +}; +} // namespace cuda +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc index 65f2cb4725..13e4cea71b 100644 --- a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc +++ b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc @@ -530,6 +530,61 @@ Sample echo operator.)DOC"); ONNX_NAMESPACE::convPoolShapeInference(ctx, false, true, 0, 1); }); + ONNX_CONTRIB_OPERATOR_SCHEMA(ConvTransposeWithDynamicPads) + .SetDomain(kMSDomain) + .SinceVersion(1) + .SetDoc(R"DOC()DOC") + .Attr( + "kernel_shape", + "", + AttributeProto::INTS, + OPTIONAL) + .Attr("output_padding", + "", + AttributeProto::INTS, + OPTIONAL) + .Attr( + "dilations", + "", + AttributeProto::INTS, + OPTIONAL) + .Attr( + "strides", + "", + AttributeProto::INTS, + OPTIONAL) + .Attr( + "auto_pad", + "", + AttributeProto::STRING, + std::string("NOTSET")) + .Attr( + "group", + "", + AttributeProto::INT, + static_cast(1)) + .Input( + 0, + "X", + "", + "T") + .Input( + 1, + "W", + "", + "T") + .Input(2, "Pads", "", "tensor(int64)", OpSchema::Optional) + .Input(3, "B", "", "T", OpSchema::Optional) + .Output( + 0, + "Y", + "", + "T") + .TypeConstraint("T", {"tensor(float16)", "tensor(float)", "tensor(double)"}, "Constrain input and output types to float tensors") + .TypeAndShapeInferenceFunction([](ONNX_NAMESPACE::InferenceContext& ctx) { + propagateElemTypeFromInputToOutput(ctx, 0, 0); + }); + ONNX_CONTRIB_OPERATOR_SCHEMA(FusedConv) .SetDomain(kMSDomain) .SinceVersion(1) @@ -1192,6 +1247,6 @@ Example 4: // register internal ops RegisterInternalSchemas(); #endif -} +} // namespace contrib } // namespace contrib } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cpu/nn/conv_transpose.cc b/onnxruntime/core/providers/cpu/nn/conv_transpose.cc index 0cf531efbb..14f13ccd20 100644 --- a/onnxruntime/core/providers/cpu/nn/conv_transpose.cc +++ b/onnxruntime/core/providers/cpu/nn/conv_transpose.cc @@ -53,30 +53,31 @@ inline void ComputeTransposePadAndOutputShape( } return; } - if (pad_type != AutoPadType::NOTSET) { - switch (pad_type) { - // We handle cases of AutoPadType::VALID and AutoPadType::SAME_UPPER/LOWER, - // the same way - case AutoPadType::VALID: - case AutoPadType::SAME_UPPER: - case AutoPadType::SAME_LOWER: - *pad_head = 0; - *pad_tail = 0; - *out_size = (in_size - 1) * stride + kernel + dilation - 1 + adj; - break; - default: - throw NotImplementedException("pad type not supported"); - } - } else { - *out_size = - (in_size - 1) * stride + kernel + dilation - 1 + adj - *pad_head - *pad_tail; + if (pad_type != AutoPadType::NOTSET) { + switch (pad_type) { + // We handle cases of AutoPadType::VALID and AutoPadType::SAME_UPPER/LOWER, + // the same way + case AutoPadType::VALID: + case AutoPadType::SAME_UPPER: + case AutoPadType::SAME_LOWER: + *pad_head = 0; + *pad_tail = 0; + *out_size = (in_size - 1) * stride + kernel + dilation - 1 + adj; + break; + default: + throw NotImplementedException("pad type not supported"); } + } else { + *out_size = + (in_size - 1) * stride + kernel + dilation - 1 + adj - *pad_head - *pad_tail; + } } -Status ConvTransposeBase::PrepareForCompute(OpKernelContext* context, bool has_bias, ConvTransposeBase::Prepare& p) const { - const auto* X = context->Input(0); - const auto* F = context->Input(1); - const Tensor* B = has_bias ? context->Input(2) : nullptr; +Status ConvTransposeBase::PrepareForCompute(OpKernelContext* context, bool has_bias, ConvTransposeBase::Prepare& p, bool dynamic_padding) const { + const Tensor* X = context->Input(0); + const Tensor* F = context->Input(1); + const Tensor* Pads = dynamic_padding ? context->Input(2) : nullptr; + const Tensor* B = has_bias ? (dynamic_padding ? context->Input(3) : context->Input(2)) : nullptr; const TensorShape& input_shape = X->Shape(); // input validations @@ -129,7 +130,15 @@ Status ConvTransposeBase::PrepareForCompute(OpKernelContext* context, bool has_b if (output_padding.empty()) { output_padding.resize(kernel_shape.size(), 0); } - std::vector pads(pads_); + std::vector pads; + pads.reserve(2 * (input_shape.NumDimensions() - 2)); + if (dynamic_padding) { + for (int64_t i = 0; i < Pads->Shape().SizeFromDimension(0); ++i) { + pads.push_back(Pads->Data()[i]); + } + } else { + pads.assign(pads_.begin(), pads_.end()); + } if (pads.empty()) { pads.resize(kernel_shape.size() * 2, 0); } @@ -214,9 +223,15 @@ void ConvTransposeBase::ComputePadsAndOutputShape( template Status ConvTranspose::Compute(OpKernelContext* context) const { + return ConvTranspose::DoConvTranspose(context, false); +} + +template +Status ConvTranspose::DoConvTranspose(OpKernelContext* context, bool dynamic_padding) const { size_t num_inputs = OpKernel::Node().InputDefs().size(); Prepare p; - ORT_RETURN_IF_ERROR(PrepareForCompute(context, num_inputs == 3, p)); + bool has_bias = dynamic_padding ? num_inputs == 4 : num_inputs == 3; + ORT_RETURN_IF_ERROR(PrepareForCompute(context, has_bias, p, dynamic_padding)); const int64_t input_image_size = p.H * p.W; const int64_t X_offset = p.num_input_channels / group_ * input_image_size; diff --git a/onnxruntime/core/providers/cpu/nn/conv_transpose.h b/onnxruntime/core/providers/cpu/nn/conv_transpose.h index 35dcdc08bc..19e32f1512 100644 --- a/onnxruntime/core/providers/cpu/nn/conv_transpose.h +++ b/onnxruntime/core/providers/cpu/nn/conv_transpose.h @@ -45,7 +45,7 @@ class ConvTransposeBase : public ConvBase { std::vector strides; }; - Status PrepareForCompute(OpKernelContext* context, bool has_bias, Prepare& p) const; + Status PrepareForCompute(OpKernelContext* context, bool has_bias, Prepare& p, bool dynamic_padding = false) const; void ComputePadsAndOutputShape(TensorShape input_shape, int64_t output_channel, const std::vector& kernel_shape, const std::vector& strides, @@ -62,6 +62,9 @@ class ConvTranspose : public OpKernel, public ConvTransposeBase { ConvTranspose(const OpKernelInfo& info) : OpKernel(info), ConvTransposeBase(info) {} Status Compute(OpKernelContext* context) const override; + + protected: + Status DoConvTranspose(OpKernelContext* context, bool dynamic_padding) const; }; } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/contrib_kernels_cuda.cc b/onnxruntime/core/providers/cuda/contrib_kernels_cuda.cc new file mode 100644 index 0000000000..923d4aca8e --- /dev/null +++ b/onnxruntime/core/providers/cuda/contrib_kernels_cuda.cc @@ -0,0 +1,17 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "contrib_ops/contrib_kernels.h" +#include "core/graph/constants.h" + +namespace onnxruntime { +namespace contrib { +namespace cuda { +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, ConvTransposeWithDynamicPads); + +void RegisterCudaContribKernels(KernelRegistry& kernel_registry) { + kernel_registry.Register(BuildKernelCreateInfo()); +} +} // namespace cuda +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index a430e385c9..d28777af1c 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -8,6 +8,7 @@ #include "cuda_allocator.h" #include "core/framework/kernel_registry.h" #include "core/framework/compute_capability.h" +#include "contrib_ops/contrib_kernels.h" using namespace onnxruntime::common; @@ -881,6 +882,8 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) { std::shared_ptr GetCudaKernelRegistry() { std::shared_ptr kernel_registry = std::make_shared(); RegisterCudaKernels(*kernel_registry); + ::onnxruntime::contrib::cuda::RegisterCudaContribKernels(*kernel_registry); + return kernel_registry; } diff --git a/onnxruntime/core/providers/cuda/nn/conv_transpose.cc b/onnxruntime/core/providers/cuda/nn/conv_transpose.cc index c15bd2978c..9c5865c78f 100644 --- a/onnxruntime/core/providers/cuda/nn/conv_transpose.cc +++ b/onnxruntime/core/providers/cuda/nn/conv_transpose.cc @@ -22,6 +22,11 @@ REGISTER_KERNEL_TYPED(MLFloat16) template Status ConvTranspose::ComputeInternal(OpKernelContext* context) const { + return DoConvTranspose(context, false); +} + +template +Status ConvTranspose::DoConvTranspose(OpKernelContext* context, bool dynamic_padding) const { typedef typename ToCudaType::MappedType CudaT; const Tensor* X = context->Input(0); @@ -35,7 +40,7 @@ Status ConvTranspose::ComputeInternal(OpKernelContext* context) const { auto w_data = reinterpret_cast(W->template Data()); size_t num_inputs = OpKernel::Node().InputDefs().size(); - bool has_bias = (num_inputs == 3); + bool has_bias = dynamic_padding ? num_inputs == 4 : num_inputs == 3; CudaT* y_data = nullptr; @@ -54,7 +59,7 @@ Status ConvTranspose::ComputeInternal(OpKernelContext* context) const { } Prepare p; - ORT_RETURN_IF_ERROR(PrepareForCompute(context, has_bias, p)); + ORT_RETURN_IF_ERROR(PrepareForCompute(context, has_bias, p, dynamic_padding)); const auto& y_dims = p.Y->Shape().GetDims(); s_.y_dims = y_dims; @@ -143,7 +148,7 @@ Status ConvTranspose::ComputeInternal(OpKernelContext* context) const { y_data)); if (has_bias) { - const Tensor* B = context->Input(2); + const Tensor* B = dynamic_padding ? context->Input(3) : context->Input(2); auto b_data = reinterpret_cast(B->template Data()); CUDNN_RETURN_IF_ERROR(cudnnAddTensor(CudnnHandle(), &alpha, s_.b_tensor, b_data, &alpha, s_.y_tensor, y_data)); } diff --git a/onnxruntime/core/providers/cuda/nn/conv_transpose.h b/onnxruntime/core/providers/cuda/nn/conv_transpose.h index 9b6abd46f9..9030e8f856 100644 --- a/onnxruntime/core/providers/cuda/nn/conv_transpose.h +++ b/onnxruntime/core/providers/cuda/nn/conv_transpose.h @@ -15,6 +15,7 @@ class ConvTranspose : public CudaKernel, public ConvTransposeBase { public: ConvTranspose(const OpKernelInfo& info) : CudaKernel(info), ConvTransposeBase(info){}; Status ComputeInternal(OpKernelContext* context) const override; + Status DoConvTranspose(OpKernelContext* context, bool dynamic_padding) const; private: mutable CudnnConvState s_; diff --git a/onnxruntime/test/contrib_ops/conv_transpose_with_dynamic_pads_test.cc b/onnxruntime/test/contrib_ops/conv_transpose_with_dynamic_pads_test.cc new file mode 100644 index 0000000000..092d07cc0e --- /dev/null +++ b/onnxruntime/test/contrib_ops/conv_transpose_with_dynamic_pads_test.cc @@ -0,0 +1,23 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "gtest/gtest.h" +#include "test/providers/provider_test_utils.h" + +namespace onnxruntime { +namespace test { +TEST(ContribOpTest, ConvTransposeWithDynamicPads) { + OpTester test("ConvTransposeWithDynamicPads", 1, onnxruntime::kMSDomain); + test.AddAttribute("kernel_shape", std::vector{3, 3}); + test.AddAttribute("output_padding", std::vector{1, 1}); + test.AddAttribute("strides", std::vector{2, 2}); + test.AddAttribute("dilations", std::vector{1, 1}); + + test.AddInput("X", {1, 1, 3, 3}, std::vector{0.16857791f, -0.15161794f, 0.08540368f, 0.1820628f, -0.21746576f, 0.08245695f, 0.1431433f, -0.43156421f, 0.30591947f}); + test.AddInput("W", {1, 1, 3, 3}, std::vector{-0.06230065f, 0.37932432f, -0.25388849f, 0.33878803f, 0.43709868f, -0.22477469f, 0.04118127f, -0.44696793f, 0.06373066f}); + test.AddInput("Pads", {4}, std::vector{1, 1, 1, 1}); + test.AddOutput("Y", {1, 1, 6, 6}, std::vector{0.07368518f, -0.08925839f, -0.06627201f, 0.06301362f, 0.03732984f, -0.01919658f, -0.00628807f, -0.02817563f, -0.01472169f, 0.04392925f, -0.00689478f, -0.01549204f, 0.07957941f, -0.11459791f, -0.09505399f, 0.07681622f, 0.03604182f, -0.01853423f, -0.0270785f, -0.00680824f, -0.06650258f, 0.08004665f, 0.07918708f, -0.0724144f, 0.06256775f, -0.17838378f, -0.18863615f, 0.20064656f, 0.133717f, -0.06876295f, -0.06398046f, -0.00864975f, 0.19289537f, -0.01490572f, -0.13673618f, 0.01949645f}); + test.Run(); +} +} // namespace test +} // namespace onnxruntime