From 7ed3bfde513a7cad7ba8cc15e10ff82c0b030c94 Mon Sep 17 00:00:00 2001 From: cloudhan Date: Mon, 17 Apr 2023 11:41:00 +0800 Subject: [PATCH] Fix FusedConv for ROCm (#15460) 1. Fix undesired runtime optimization for non-Relu activation. 3. Fix false positive runtime error log due to fusion failure. --- onnxruntime/contrib_ops/cuda/fused_conv.cc | 26 +++--- onnxruntime/contrib_ops/rocm/fused_conv.cc | 93 ++++++++----------- .../core/optimizer/conv_activation_fusion.cc | 8 +- .../test/optimizer/graph_transform_test.cc | 25 +++-- 4 files changed, 72 insertions(+), 80 deletions(-) diff --git a/onnxruntime/contrib_ops/cuda/fused_conv.cc b/onnxruntime/contrib_ops/cuda/fused_conv.cc index 48881ddca4..e126f8bcb3 100644 --- a/onnxruntime/contrib_ops/cuda/fused_conv.cc +++ b/onnxruntime/contrib_ops/cuda/fused_conv.cc @@ -1,6 +1,7 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "core/common/status.h" #include "core/providers/cuda/nn/conv.h" #include "core/providers/cuda/cuda_common.h" @@ -14,28 +15,24 @@ class FusedConv : public onnxruntime::cuda::Conv { using Base = onnxruntime::cuda::Conv; FusedConv(const OpKernelInfo& info) : onnxruntime::cuda::Conv(info) { std::string activation; - if (info.GetAttr("activation", &activation) == Status::OK() && - MapMode(activation) == Status::OK() && - cudnnCreateActivationDescriptor(&activation_desc_) == CUDNN_STATUS_SUCCESS) { - status_ = cudnnSetActivationDescriptor(activation_desc_, - activation_mode_, - cudnnNanPropagation_t::CUDNN_NOT_PROPAGATE_NAN, - std::numeric_limits::max()); - } + ORT_THROW_IF_ERROR(info.GetAttr("activation", &activation)); + ORT_THROW_IF_ERROR(MapMode(activation)); + CUDNN_CALL_THROW(cudnnCreateActivationDescriptor(&activation_desc_)); + CUDNN_CALL_THROW(cudnnSetActivationDescriptor( + activation_desc_, activation_mode_, cudnnNanPropagation_t::CUDNN_NOT_PROPAGATE_NAN, + std::numeric_limits::max())); } ORT_DISALLOW_COPY_AND_ASSIGNMENT(FusedConv); ~FusedConv() { if (activation_desc_) { - cudnnDestroyActivationDescriptor(activation_desc_); - status_ = CUDNN_STATUS_NOT_INITIALIZED; + CUDNN_CALL_THROW(cudnnDestroyActivationDescriptor(activation_desc_)); activation_desc_ = nullptr; } } Status ComputeInternal(OpKernelContext* context) const override { - CUDNN_RETURN_IF_ERROR(status_); std::lock_guard lock(Base::s_.mutex); auto cudnnHandle = this->GetCudnnHandle(context); ORT_RETURN_IF_ERROR(Base::UpdateState(context, true)); @@ -104,13 +101,12 @@ class FusedConv : public onnxruntime::cuda::Conv { if (activaton_mode == "Relu") { activation_mode_ = cudnnActivationMode_t::CUDNN_ACTIVATION_RELU; } else { - return Status(common::StatusCategory::ONNXRUNTIME, - common::StatusCode::INVALID_ARGUMENT, - "unsupported conv activation mode"); + return ORT_MAKE_STATUS( + StatusCategory::ONNXRUNTIME, StatusCode::INVALID_ARGUMENT, + "unsupported conv activation mode \"", activaton_mode, "\""); } return Status::OK(); } - cudnnStatus_t status_ = CUDNN_STATUS_NOT_INITIALIZED; cudnnActivationMode_t activation_mode_; cudnnActivationDescriptor_t activation_desc_ = nullptr; }; diff --git a/onnxruntime/contrib_ops/rocm/fused_conv.cc b/onnxruntime/contrib_ops/rocm/fused_conv.cc index 13b9574780..00ecd2b4af 100644 --- a/onnxruntime/contrib_ops/rocm/fused_conv.cc +++ b/onnxruntime/contrib_ops/rocm/fused_conv.cc @@ -3,6 +3,7 @@ #include #include +#include "core/common/status.h" #include "core/providers/rocm/nn/conv.h" #include "core/providers/rocm/rocm_common.h" @@ -81,41 +82,28 @@ struct FNVHash { // miopenConvolutionDescriptor_t, so we have to guess. // This algorithm is based on a specific behavior of miopenGetConvolutionNdDescriptor, // which fails when requestedSpatialDim > the convolution's spatial dimension - std::vector spatial_dims; - std::vector pads; - std::vector strides; - std::vector dilations; + constexpr const int kMaxSpatialDim = 5; + std::vector pads{kMaxSpatialDim}; + std::vector strides{kMaxSpatialDim}; + std::vector dilations{kMaxSpatialDim}; miopenConvolutionMode_t mode; - while (true) { - spatial_dims.resize(spatial_dim); - pads.resize(spatial_dim); - strides.resize(spatial_dim); - dilations.resize(spatial_dim); - - if (miopenStatusSuccess != miopenGetConvolutionNdDescriptor(cdesc, - spatial_dim, - spatial_dims.data(), - pads.data(), - strides.data(), - dilations.data(), - &mode)) { - // Remove the extra dimension - spatial_dims.resize(spatial_dim - 1); - pads.resize(spatial_dim - 1); - strides.resize(spatial_dim - 1); - dilations.resize(spatial_dim - 1); - (*this) << spatial_dim; - (*this) << spatial_dims; - (*this) << pads; - (*this) << strides; - (*this) << dilations; - break; - } - spatial_dim += 1; - ORT_ENFORCE(spatial_dim < 10, - "miopenGetConvolutionNdDescriptor is supposed to fail before spatial_dim gets to ", - spatial_dim); + bool spatial_dim_guessed = false; + for (int i=0; i { using Base = onnxruntime::rocm::Conv; FusedConv(const OpKernelInfo& info) : onnxruntime::rocm::Conv(info) { std::string activation; - if (info.GetAttr("activation", &activation) == Status::OK() && - MapMode(activation) == Status::OK() && - miopenCreateActivationDescriptor(&activation_desc_) == miopenStatusSuccess) { - status_ = miopenSetActivationDescriptor(activation_desc_, - activation_mode_, - 0.0, 0.0, 0.0); - } + ORT_THROW_IF_ERROR(info.GetAttr("activation", &activation)); + ORT_THROW_IF_ERROR(MapMode(activation)); + MIOPEN_CALL_THROW(miopenCreateActivationDescriptor(&activation_desc_)); + MIOPEN_CALL_THROW(miopenSetActivationDescriptor(activation_desc_, activation_mode_, 0.0, 0.0, 0.0)); } ORT_DISALLOW_COPY_AND_ASSIGNMENT(FusedConv); ~FusedConv() { if (activation_desc_) { - miopenDestroyActivationDescriptor(activation_desc_); - status_ = miopenStatusNotInitialized; + MIOPEN_CALL_THROW(miopenDestroyActivationDescriptor(activation_desc_)); activation_desc_ = nullptr; } } Status ComputeInternal(OpKernelContext* context) const override { - MIOPEN_RETURN_IF_ERROR(status_); std::lock_guard lock(Base::s_.mutex); ORT_RETURN_IF_ERROR(Base::UpdateState(context, true)); @@ -160,7 +143,7 @@ class FusedConv : public onnxruntime::rocm::Conv { bool has_z = nullptr != Base::s_.z_data; bool has_b = nullptr != Base::s_.b_data; auto factory = [this](FusedConvFusionData& fusion) { - return this->DoCreateFusionDesc(fusion); + return this->DoCreateFusionDesc(this->Node().Name(), fusion); }; auto& cached_item = plan_cache_.FindOrCreateFusionPlanCache(Hash(), factory); @@ -266,13 +249,12 @@ class FusedConv : public onnxruntime::rocm::Conv { if (activaton_mode == "Relu") { activation_mode_ = miopenActivationMode_t::miopenActivationRELU; } else { - return Status(common::StatusCategory::ONNXRUNTIME, - common::StatusCode::INVALID_ARGUMENT, - "unsupported conv activation mode"); + return ORT_MAKE_STATUS( + StatusCategory::ONNXRUNTIME, StatusCode::INVALID_ARGUMENT, + "unsupported conv activation mode \"", activaton_mode, "\""); } return Status::OK(); } - miopenStatus_t status_ = miopenStatusNotInitialized; miopenActivationMode_t activation_mode_; miopenActivationDescriptor_t activation_desc_ = nullptr; @@ -389,17 +371,22 @@ class FusedConv : public onnxruntime::rocm::Conv { static FusionPlanCache plan_cache_; - Status DoCreateFusionDesc(FusedConvFusionData& fusion) const { + Status DoCreateFusionDesc(const std::string& node_name, FusedConvFusionData& fusion) const { bool has_z = nullptr != Base::s_.z_data; bool has_b = nullptr != Base::s_.b_data; MIOPEN_RETURN_IF_ERROR(miopenCreateFusionPlan(&fusion.plan, miopenVerticalFusion, Base::s_.x_tensor)); MIOPEN_RETURN_IF_ERROR(miopenCreateOperatorArgs(&fusion.fusion_args)); - MIOPEN_RETURN_IF_ERROR(miopenCreateOpConvForward(fusion.plan, - &fusion.conv_op, - Base::s_.conv_desc, - Base::s_.w_desc)); + auto status = miopenCreateOpConvForward(fusion.plan, &fusion.conv_op, Base::s_.conv_desc, Base::s_.w_desc); + if (status == miopenStatusUnsupportedOp) { + auto msg = MakeString("MIOpen does not support the conv fusion for node \"", + node_name, "\", fallback to unfused implementation."); + LOGS_DEFAULT(WARNING) << msg; + return ORT_MAKE_STATUS(ONNXRUNTIME, NOT_IMPLEMENTED, msg); + } + MIOPEN_RETURN_IF_ERROR(status); + if (has_z) { MIOPEN_RETURN_IF_ERROR(miopenCreateOpBiasForward(fusion.plan, &fusion.bias_z_op, diff --git a/onnxruntime/core/optimizer/conv_activation_fusion.cc b/onnxruntime/core/optimizer/conv_activation_fusion.cc index e87e3b19c1..54fe7c4087 100644 --- a/onnxruntime/core/optimizer/conv_activation_fusion.cc +++ b/onnxruntime/core/optimizer/conv_activation_fusion.cc @@ -70,7 +70,7 @@ class ConvActivation : public NodeSelector { return std::nullopt; } - auto is_supported_non_cuda_ep_activation = [&graph_viewer](const Node& activation_node) { + auto is_supported_non_cuda_rocm_ep_activation = [&graph_viewer](const Node& activation_node) { if (graph_utils::IsSupportedOptypeVersionAndDomain(activation_node, "Relu", {6, 13, 14}) || graph_utils::IsSupportedOptypeVersionAndDomain(activation_node, "Sigmoid", {6, 13}) || graph_utils::IsSupportedOptypeVersionAndDomain(activation_node, "Tanh", {6, 13}) || @@ -94,17 +94,17 @@ class ConvActivation : public NodeSelector { } // check EP type and activation - if (node_ep == kCudaExecutionProvider) { + if (node_ep == kCudaExecutionProvider || node_ep == kRocmExecutionProvider) { if (!graph_utils::IsSupportedOptypeVersionAndDomain(*next_node, "Relu", {6, 13, 14})) { return std::nullopt; } } else if (node_ep.empty() || node_ep == kCpuExecutionProvider) { - if (!is_supported_non_cuda_ep_activation(*next_node) && + if (!is_supported_non_cuda_rocm_ep_activation(*next_node) && !graph_utils::IsSupportedOptypeVersionAndDomain(*next_node, "HardSigmoid", {6})) { return std::nullopt; } } else { - if (!is_supported_non_cuda_ep_activation(*next_node)) { + if (!is_supported_non_cuda_rocm_ep_activation(*next_node)) { return std::nullopt; } } diff --git a/onnxruntime/test/optimizer/graph_transform_test.cc b/onnxruntime/test/optimizer/graph_transform_test.cc index 96ce3d5491..bd2b733225 100755 --- a/onnxruntime/test/optimizer/graph_transform_test.cc +++ b/onnxruntime/test/optimizer/graph_transform_test.cc @@ -1221,9 +1221,6 @@ TEST_F(GraphTransformationTests, FuseCpuConvAdd) { #if !defined(DISABLE_CONTRIB_OPS) TEST_F(GraphTransformationTests, FuseConvActivation) { -#ifdef USE_CUDA - std::unordered_map model_to_op_name{{ORT_TSTR("fusion/conv_relu.onnx"), "Relu"}}; -#else std::unordered_map model_to_op_name{{ORT_TSTR("fusion/conv_relu.onnx"), "Relu"}, {ORT_TSTR("fusion/conv_relu_opset12.onnx"), "Relu"}, {ORT_TSTR("fusion/conv_clip.onnx"), "Clip"}, @@ -1231,7 +1228,6 @@ TEST_F(GraphTransformationTests, FuseConvActivation) { {ORT_TSTR("fusion/conv_tanh.onnx"), "Tanh"}, {ORT_TSTR("fusion/conv_leakyrelu.onnx"), "LeakyRelu"}, {ORT_TSTR("fusion/conv_hardsigmoid.onnx"), "HardSigmoid"}}; -#endif for (const auto& model : model_to_op_name) { PathString model_uri = PathString(MODEL_FOLDER) + model.first; SCOPED_TRACE(ORT_TSTR("model file: ") + model_uri); @@ -1242,17 +1238,30 @@ TEST_F(GraphTransformationTests, FuseConvActivation) { for (auto& node : p_model->MainGraph().Nodes()) { node.SetExecutionProviderType(kCudaExecutionProvider); } +#elif defined(USE_ROCM) + for (auto& node : p_model->MainGraph().Nodes()) { + node.SetExecutionProviderType(kCudaExecutionProvider); + } #endif - std::map op_to_count = CountOpsInGraph(graph); - ASSERT_TRUE(op_to_count[model.second] >= 1); + std::map op_to_count_before_fusion = CountOpsInGraph(graph); + ASSERT_TRUE(op_to_count_before_fusion[model.second] >= 1); // Apply transformer onnxruntime::GraphTransformerManager graph_transformation_mgr{5}; ASSERT_STATUS_OK(graph_transformation_mgr.Register(std::make_unique(), TransformerLevel::Level2)); ASSERT_STATUS_OK(graph_transformation_mgr.ApplyTransformers(graph, TransformerLevel::Level2, *logger_)); - op_to_count = CountOpsInGraph(graph); - ASSERT_TRUE(op_to_count[model.second] == 0); + std::map op_to_count_after_fusion = CountOpsInGraph(graph); +#if defined(USE_CUDA) || defined(USE_ROCM) + std::set cuda_rocm_supported = {"Relu"}; + if (cuda_rocm_supported.find(model.second) == cuda_rocm_supported.end()) { + ASSERT_EQ(op_to_count_before_fusion[model.second], op_to_count_after_fusion[model.second]); + } else { + ASSERT_TRUE(op_to_count_after_fusion[model.second] == 0); + } +#else + ASSERT_TRUE(op_to_count_after_fusion[model.second] == 0); +#endif } }