saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-28 20:11:22 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Fix FusedConv for ROCm (#15460)

Browse source 
1. Fix undesired runtime optimization for non-Relu activation.
3. Fix false positive runtime error log due to fusion failure.
This commit is contained in:
cloudhan 2023-04-17 11:41:00 +08:00 committed by GitHub
parent ac6ceffb2c
commit 7ed3bfde51
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
4 changed files with 72 additions and 80 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

26
onnxruntime/contrib_ops/cuda/fused_conv.cc
View file

@ -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<T, false> {
using Base = onnxruntime::cuda::Conv<T, false>;
FusedConv(const OpKernelInfo& info) : onnxruntime::cuda::Conv<T, false>(info) {
std::string activation;
if (info.GetAttr<std::string>("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<double>::max());
}
ORT_THROW_IF_ERROR(info.GetAttr<std::string>("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<double>::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<OrtMutex> 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<T, false> {
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;
};

93
onnxruntime/contrib_ops/rocm/fused_conv.cc
View file

@ -3,6 +3,7 @@
#include <unordered_set>
#include <unordered_map>
#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<int> spatial_dims;
std::vector<int> pads;
std::vector<int> strides;
std::vector<int> dilations;
constexpr const int kMaxSpatialDim = 5;
std::vector<int> pads{kMaxSpatialDim};
std::vector<int> strides{kMaxSpatialDim};
std::vector<int> 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<kMaxSpatialDim; i++) {
if (miopenStatusSuccess == miopenGetConvolutionNdDescriptor(
cdesc, i, &spatial_dim, pads.data(), strides.data(), dilations.data(), &mode)) {
spatial_dim_guessed = true;
break;
}
}
ORT_ENFORCE(spatial_dim_guessed, "Failed to guess the actual spatial dimension");
// Remove the extra dimension
pads.resize(spatial_dim);
strides.resize(spatial_dim);
dilations.resize(spatial_dim);
(*this) << spatial_dim;
(*this) << pads;
(*this) << strides;
(*this) << dilations;
#endif
}
@ -129,27 +117,22 @@ class FusedConv : public onnxruntime::rocm::Conv<T, false> {
using Base = onnxruntime::rocm::Conv<T, false>;
FusedConv(const OpKernelInfo& info) : onnxruntime::rocm::Conv<T, false>(info) {
std::string activation;
if (info.GetAttr<std::string>("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<std::string>("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<OrtMutex> lock(Base::s_.mutex);
ORT_RETURN_IF_ERROR(Base::UpdateState(context, true));
@ -160,7 +143,7 @@ class FusedConv : public onnxruntime::rocm::Conv<T, false> {
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<T, false> {
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<T, false> {
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,

8
onnxruntime/core/optimizer/conv_activation_fusion.cc
View file

@ -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;
}
}

25
onnxruntime/test/optimizer/graph_transform_test.cc
View file

@ -1221,9 +1221,6 @@ TEST_F(GraphTransformationTests, FuseCpuConvAdd) {
#if !defined(DISABLE_CONTRIB_OPS)
TEST_F(GraphTransformationTests, FuseConvActivation) {
#ifdef USE_CUDA
std::unordered_map<PathString, std::string> model_to_op_name{{ORT_TSTR("fusion/conv_relu.onnx"), "Relu"}};
#else
std::unordered_map<PathString, std::string> 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<std::string, int> op_to_count = CountOpsInGraph(graph);
ASSERT_TRUE(op_to_count[model.second] >= 1);
std::map<std::string, int> 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<ConvActivationFusion>(), 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<std::string, int> op_to_count_after_fusion = CountOpsInGraph(graph);
#if defined(USE_CUDA) || defined(USE_ROCM)
std::set<std::string> 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
}
}