From f9cf62912a317d6d13e074796c847b47739ecc5d Mon Sep 17 00:00:00 2001 From: mindest <30493312+mindest@users.noreply.github.com> Date: Tue, 12 Oct 2021 19:57:38 +0800 Subject: [PATCH] Add same_shape case for BiasDropout (#9188) * bias dropout improvement * add transform case for same shape case * combine kernel * merge with vectorized kernel * use "has_same_shape_bias" * minor: a "N % 4 != 0" case * add op UT for has_same_shape_bias * address comments; add param case for 1d bias; add param case tests for 1d and same-shape bias * rewrite logic condition Co-authored-by: Peng Wang --- docs/ContribOperators.md | 2 +- .../contrib_ops/cuda/math/bias_dropout.cc | 25 +++-- .../contrib_ops/cuda/math/bias_dropout.h | 3 +- .../cuda/math/bias_dropout_impl.cu | 69 +++++++++--- .../core/graph/contrib_ops/contrib_defs.cc | 2 +- .../core/optimizer/bias_dropout_fusion.cc | 56 ++++++---- .../test/contrib_ops/bias_dropout_op_test.cc | 23 +++- .../test/optimizer/graph_transform_test.cc | 6 ++ .../bias_dropout_fusion_dim_is_param.onnx | Bin 0 -> 314 bytes ..._dropout_residual_fusion_dim_is_param.onnx | Bin 0 -> 404 bytes .../fusion/bias_dropout_residual_gen.py | 99 ++++++++++++++++++ ...as_dropout_residual_same_shape_fusion.onnx | Bin 0 -> 376 bytes ...sidual_same_shape_fusion_dim_is_param.onnx | Bin 0 -> 424 bytes .../bias_dropout_same_shape_fusion.onnx | Bin 0 -> 298 bytes ...ropout_same_shape_fusion_dim_is_param.onnx | Bin 0 -> 334 bytes 15 files changed, 230 insertions(+), 55 deletions(-) create mode 100644 onnxruntime/test/testdata/transform/fusion/bias_dropout_fusion_dim_is_param.onnx create mode 100644 onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_fusion_dim_is_param.onnx create mode 100644 onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_same_shape_fusion.onnx create mode 100644 onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_same_shape_fusion_dim_is_param.onnx create mode 100644 onnxruntime/test/testdata/transform/fusion/bias_dropout_same_shape_fusion.onnx create mode 100644 onnxruntime/test/testdata/transform/fusion/bias_dropout_same_shape_fusion_dim_is_param.onnx diff --git a/docs/ContribOperators.md b/docs/ContribOperators.md index 44d6c8ae4b..4ef770568c 100644 --- a/docs/ContribOperators.md +++ b/docs/ContribOperators.md @@ -359,7 +359,7 @@ This version of the operator has been available since version 1 of the 'com.micr
data : T
The input data as Tensor.
bias : T
-
The bias input, a vector with the same shape as last dim of data
+
The bias input, a vector with the same shape as last dim of data OR same shape with data
residual (optional) : T
The residual input, must have the same shape as data
ratio (optional) : T1
diff --git a/onnxruntime/contrib_ops/cuda/math/bias_dropout.cc b/onnxruntime/contrib_ops/cuda/math/bias_dropout.cc index 950106d879..192c28c86d 100644 --- a/onnxruntime/contrib_ops/cuda/math/bias_dropout.cc +++ b/onnxruntime/contrib_ops/cuda/math/bias_dropout.cc @@ -36,7 +36,8 @@ struct BiasDropoutComputeImpl { const Tensor& bias, const Tensor* residual, Tensor& Y, - bool* mask_data) const { + bool* mask_data, + bool has_same_shape_bias) const { typedef typename ToCudaType::MappedType CudaT; const CudaT* X_data = reinterpret_cast(X.template Data()); @@ -52,7 +53,7 @@ struct BiasDropoutComputeImpl { CudaT* Y_data = reinterpret_cast(Y.template MutableData()); - BiasDropoutKernelImpl(prop, stream, N, fdm_dim, ratio_data, generator, X_data, bias_data, residual_data, Y_data, mask_data); + BiasDropoutKernelImpl(prop, stream, N, fdm_dim, ratio_data, generator, X_data, bias_data, residual_data, Y_data, mask_data, has_same_shape_bias); return Status::OK(); } @@ -70,12 +71,16 @@ Status BiasDropout::ComputeInternal(OpKernelContext* context) const { const Tensor* bias = context->Input(1); if (bias == nullptr) return Status(common::ONNXRUNTIME, common::FAIL, "Bias input of BiasDropout is not available."); const TensorShape& bias_shape = bias->Shape(); - if (bias_shape.NumDimensions() != 1) { - return Status(common::ONNXRUNTIME, common::FAIL, "Bias input is not a 1D tensor."); - } - const int64_t dim = bias_shape[0]; - if (dim != x_shape.GetDims().back()) { - return Status(common::ONNXRUNTIME, common::FAIL, "Bias' dimension doesn't match input's last dimension."); + const int64_t dim = bias_shape.GetDims().back(); + bool has_same_shape_bias = (bias_shape == x_shape); + if (!has_same_shape_bias) { + if (bias_shape.NumDimensions() != 1) { + return Status(common::ONNXRUNTIME, common::FAIL, "Bias input is not a 1D tensor."); + } + + if (dim != x_shape.GetDims().back()) { + return Status(common::ONNXRUNTIME, common::FAIL, "Bias' dimension doesn't match input's last dimension."); + } } //Get residual_data @@ -114,9 +119,9 @@ Status BiasDropout::ComputeInternal(OpKernelContext* context) const { utils::MLTypeCallDispatcher t_disp(X->GetElementType()); return t_disp.InvokeRet( - GetDeviceProp(), Stream(), N, fdm_dim, ratio_data, generator, *X, *bias, residual, *Y, mask_data); + GetDeviceProp(), Stream(), N, fdm_dim, ratio_data, generator, *X, *bias, residual, *Y, mask_data, has_same_shape_bias); } } // namespace cuda } // namespace contrib -} // namespace onnxruntime \ No newline at end of file +} // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cuda/math/bias_dropout.h b/onnxruntime/contrib_ops/cuda/math/bias_dropout.h index 97c294e641..eb931dbd54 100644 --- a/onnxruntime/contrib_ops/cuda/math/bias_dropout.h +++ b/onnxruntime/contrib_ops/cuda/math/bias_dropout.h @@ -26,7 +26,8 @@ void BiasDropoutKernelImpl( const T* bias_data, const T* residual_data, T* Y_data, - bool* mask_data); + bool* mask_data, + bool has_same_shape_bias); class BiasDropout final : public CudaKernel { public: diff --git a/onnxruntime/contrib_ops/cuda/math/bias_dropout_impl.cu b/onnxruntime/contrib_ops/cuda/math/bias_dropout_impl.cu index 27fb09c969..ad2ff05a58 100644 --- a/onnxruntime/contrib_ops/cuda/math/bias_dropout_impl.cu +++ b/onnxruntime/contrib_ops/cuda/math/bias_dropout_impl.cu @@ -28,7 +28,7 @@ namespace cuda { constexpr int UNROLL = 4; -template +template __global__ void BiasDropoutKernel( const int64_t N, const fast_divmod fdm_dim, @@ -58,14 +58,19 @@ __global__ void BiasDropoutKernel( // use of Philox_4x32_10 is to generate a multiple of 4 times number of threads. for (CUDA_LONG id = idx * UNROLL; id < N; id += step_size) { rand = curand_uniform4(&state); - + // actual computation #pragma unroll for (int i = 0; i < UNROLL; i++) { CUDA_LONG li = id + i; if (li < N) { - int offset = fdm_dim.mod(li); - float bias = float(bias_data[offset]); + float bias; + if (has_same_shape_bias) { + bias = float(bias_data[li]); + } else { + int offset = fdm_dim.mod(li); + bias = float(bias_data[offset]); + } mask_data[li] = (&rand.x)[i] < p; float output_data = (float(X_data[li]) + bias) * mask_data[li] * scale; @@ -83,7 +88,7 @@ __global__ void BiasDropoutKernel( } -template +template __global__ void BiasDropoutVectorizedKernel( const int64_t N, const fast_divmod fdm_dim, @@ -105,7 +110,7 @@ __global__ void BiasDropoutVectorizedKernel( float4 rand; - // using vectorized data load/store approach when N % 4 == 0 + // using vectorized data load/store approach when N % 4 == 0 // since this is typical case for input shape size using LoadT = aligned_vector; using MaskLoadT = aligned_vector; @@ -113,8 +118,14 @@ __global__ void BiasDropoutVectorizedKernel( for (CUDA_LONG id = idx * UNROLL; id < N; id += step_size) { rand = curand_uniform4(&state); - + // vectorized load into storage + T bias_vec[UNROLL]; + if (has_same_shape_bias) { + LoadT *value0 = reinterpret_cast(&bias_vec); + *value0 = *reinterpret_cast(&bias_data[id]); + } + T src[UNROLL]; LoadT *value1 = reinterpret_cast(&src); *value1 = *reinterpret_cast(&X_data[id]); @@ -131,8 +142,13 @@ __global__ void BiasDropoutVectorizedKernel( // actual computation #pragma unroll for (int ii = 0; ii < UNROLL; ii++) { - int offset = fdm_dim.mod(id + ii); - float bias = float(bias_data[offset]); + float bias; + if (has_same_shape_bias) { + bias = float(bias_vec[ii]); + } else { + int offset = fdm_dim.mod(id + ii); + bias = float(bias_data[offset]); + } mask[ii] = (&rand.x)[ii] < p; float output_data = (float(src[ii]) + bias) * mask[ii] * scale; @@ -162,7 +178,8 @@ void BiasDropoutKernelImpl( const T* bias_data, const T* residual_data, T* Y_data, - bool* mask_data) { + bool* mask_data, + bool has_same_shape_bias) { const int block_size = 256; const int blocks_per_sm = prop.maxThreadsPerMultiProcessor / block_size; const int grid_size = std::min(prop.multiProcessorCount * blocks_per_sm, static_cast(CeilDiv(N, block_size * UNROLL))); @@ -172,16 +189,32 @@ void BiasDropoutKernelImpl( auto seeds = generator.NextPhiloxSeeds(counter_offset); if (N % UNROLL != 0) { - if (residual_data == nullptr) { - BiasDropoutKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + if (has_same_shape_bias) { + if (residual_data == nullptr) { + BiasDropoutKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + } else { + BiasDropoutKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + } } else { - BiasDropoutKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + if (residual_data == nullptr) { + BiasDropoutKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + } else { + BiasDropoutKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + } } } else { - if (residual_data == nullptr) { - BiasDropoutVectorizedKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + if (has_same_shape_bias) { + if (residual_data == nullptr) { + BiasDropoutVectorizedKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + } else { + BiasDropoutVectorizedKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + } } else { - BiasDropoutVectorizedKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + if (residual_data == nullptr) { + BiasDropoutVectorizedKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + } else { + BiasDropoutVectorizedKernel<<>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data); + } } } } @@ -198,7 +231,9 @@ void BiasDropoutKernelImpl( const T* bias_data, \ const T* residual_data, \ T* Y_data, \ - bool* mask_data); + bool* mask_data, \ + bool has_same_shape_bias); + SPECIALIZED_BIAS_DROPOUT_IMPL(float) SPECIALIZED_BIAS_DROPOUT_IMPL(double) diff --git a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc index dfaf089979..375f45eb5e 100644 --- a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc +++ b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc @@ -2856,7 +2856,7 @@ It's an extension of Gelu. It takes the sum of input A and bias input B as the i .Attr("seed", "(Optional) Seed to the random generator, if not specified we will auto generate one.", AttributeProto::INT, OPTIONAL_VALUE) .AllowUncheckedAttributes() .Input(0, "data", "The input data as Tensor.", "T") - .Input(1, "bias", "The bias input, a vector with the same shape as last dim of data", "T") + .Input(1, "bias", "The bias input, a vector with the same shape as last dim of data OR same shape with data", "T") .Input(2, "residual", "The residual input, must have the same shape as data", "T", OpSchema::Optional) .Input(3, "ratio", "The ratio of random dropout, with value in [0, 1). If this input was not set, " diff --git a/onnxruntime/core/optimizer/bias_dropout_fusion.cc b/onnxruntime/core/optimizer/bias_dropout_fusion.cc index 0095f94ddc..2e662ee570 100644 --- a/onnxruntime/core/optimizer/bias_dropout_fusion.cc +++ b/onnxruntime/core/optimizer/bias_dropout_fusion.cc @@ -10,6 +10,18 @@ using namespace ONNX_NAMESPACE; using namespace ::onnxruntime::common; namespace onnxruntime { +static bool IsSameShape(const TensorShapeProto& shape1, const TensorShapeProto& shape2) { + int rank1 = shape1.dim_size(); + if (rank1 != shape2.dim_size()) { + return false; + } + bool same_shape = true; + for (int i = 0; i < rank1; ++i) { + same_shape &= ONNX_NAMESPACE::operator==(shape1.dim(i), shape2.dim(i)); + } + return same_shape; +} + void FuseResidualAddIfAny(Graph& graph, const Node& dropout_node, std::vector& dropout_input, std::vector& dropout_output, @@ -37,17 +49,12 @@ void FuseResidualAddIfAny(Graph& graph, const Node& dropout_node, if (input1_shape == nullptr || input2_shape == nullptr || input1_shape->dim_size() < 1 || - input2_shape->dim_size() < 1 || - input1_shape->dim_size() != input2_shape->dim_size()) { + input2_shape->dim_size() < 1) { continue; } // Inputs of Residual Add must match in shape - bool match = true; - for (int i = 0; i < input1_shape->dim_size(); ++i) { - match &= ONNX_NAMESPACE::operator==(input1_shape->dim(i), input2_shape->dim(i)); - } - if (!match) { + if (!IsSameShape(*input1_shape, *input2_shape)) { continue; } @@ -107,22 +114,29 @@ Status BiasDropoutFusion::ApplyImpl(Graph& graph, bool& modified, int graph_leve continue; } - int last_dim_shape1 = input1_shape->dim_size() - 1; - int last_dim_shape2 = input2_shape->dim_size() - 1; - if (!utils::HasDimValue(input1_shape->dim(last_dim_shape1)) || - !utils::HasDimValue(input2_shape->dim(last_dim_shape2)) || - input1_shape->dim(last_dim_shape1).dim_value() != input2_shape->dim(last_dim_shape2).dim_value()) { - continue; - } - - if (input1_shape->dim_size() == 1) { - dropout_input.push_back(node.MutableInputDefs()[1]); // dropout input - dropout_input.push_back(node.MutableInputDefs()[0]); // bias - } else if (input2_shape->dim_size() == 1) { + if (IsSameShape(*input1_shape, *input2_shape)) { dropout_input.push_back(node.MutableInputDefs()[0]); // dropout input dropout_input.push_back(node.MutableInputDefs()[1]); // bias } else { - continue; + const int last_dim_shape1 = input1_shape->dim_size() - 1; + const int last_dim_shape2 = input2_shape->dim_size() - 1; + if (!(utils::HasDimValue(input1_shape->dim(last_dim_shape1)) && + utils::HasDimValue(input2_shape->dim(last_dim_shape2)) && + input1_shape->dim(last_dim_shape1).dim_value() == input2_shape->dim(last_dim_shape2).dim_value()) && + !(utils::HasDimParam(input1_shape->dim(last_dim_shape1)) && + utils::HasDimParam(input2_shape->dim(last_dim_shape2)) && + input1_shape->dim(last_dim_shape1).dim_param() == input2_shape->dim(last_dim_shape2).dim_param())) { + continue; // continue if no same DimValue && no same DimParam + } + if (input1_shape->dim_size() == 1) { + dropout_input.push_back(node.MutableInputDefs()[1]); // dropout input + dropout_input.push_back(node.MutableInputDefs()[0]); // bias + } else if (input2_shape->dim_size() == 1) { + dropout_input.push_back(node.MutableInputDefs()[0]); // dropout input + dropout_input.push_back(node.MutableInputDefs()[1]); // bias + } else { + continue; + } } Node& add_node = node; nodes_to_fuse.push_back(add_node); @@ -163,7 +177,7 @@ Status BiasDropoutFusion::ApplyImpl(Graph& graph, bool& modified, int graph_leve const std::string op_type = "BiasDropout"; Node& dropout_add_fusion_node = graph.AddNode(graph.GenerateNodeName(op_type), op_type, - "fused Add and Dropout", + "fused Add-Dropout-(Add) for " + dropout_node.Name(), dropout_input, dropout_output, {}, diff --git a/onnxruntime/test/contrib_ops/bias_dropout_op_test.cc b/onnxruntime/test/contrib_ops/bias_dropout_op_test.cc index b16b3f190c..eaa583a633 100644 --- a/onnxruntime/test/contrib_ops/bias_dropout_op_test.cc +++ b/onnxruntime/test/contrib_ops/bias_dropout_op_test.cc @@ -30,7 +30,8 @@ enum TrainingMode { TrainingFalse, #if defined(USE_CUDA) || defined(USE_ROCM) namespace { void RunBiasDropoutTest(const bool use_mask, const std::vector& input_shape, float ratio = -1.0f, - TrainingMode training_mode = TrainingTrue, bool use_float16_ratio = false, bool has_residual = true) { + TrainingMode training_mode = TrainingTrue, bool use_float16_ratio = false, + bool has_residual = true, bool has_same_shape_bias = false) { OpTester t{"BiasDropout", 1, kMSDomain}; const int64_t seed = 42; t.AddAttribute("seed", seed); @@ -40,8 +41,13 @@ void RunBiasDropoutTest(const bool use_mask, const std::vector& input_s const std::vector input = ValueRange(input_size, 1.0f, 1.0f); t.AddInput("data", input_shape, input); - std::vector bias_shape{input_shape.back()}; - const auto bias_size = input_shape.back(); + std::vector bias_shape; + if (has_same_shape_bias) { + bias_shape = input_shape; + } else { + bias_shape.push_back(input_shape.back()); + } + const auto bias_size = has_same_shape_bias ? input_size : input_shape.back(); const std::vector bias = ValueRange(bias_size, 2.0f, 1.0f); t.AddInput("bias", bias_shape, bias); @@ -143,7 +149,7 @@ TEST(BiasDropoutTest, BasicWithoutResidualAndNotVectorized) { RunBiasDropoutTest(false, {10, 5, 5}, 0.75f, TrainingTrue, false, false); } TEST(BiasDropoutTest, MaskAndNotVectorized) { - RunBiasDropoutTest(true, {3, 5, 100}, 0.25f); + RunBiasDropoutTest(true, {3, 5, 10}, 0.25f); } // N % 4 == 0 @@ -178,6 +184,15 @@ TEST(BiasDropoutTest, RatioLimit) { TEST(BiasDropoutTest, EmptyRatio) { RunBiasDropoutTest(true, {2, 7, 1024}); } + +// has_same_bias_shape == true +TEST(BiasDropoutTest, BasicBiasSameShape) { + RunBiasDropoutTest(false, {10, 10, 10}, 0.75f, TrainingTrue, false, true, true); +} + +TEST(BiasDropoutTest, BasicBiasSameShapeNotVectorized) { + RunBiasDropoutTest(false, {10, 5, 5}, 0.75f, TrainingTrue, false, true, true); +} #endif } // namespace test diff --git a/onnxruntime/test/optimizer/graph_transform_test.cc b/onnxruntime/test/optimizer/graph_transform_test.cc index cebd850cff..876e3caf6c 100644 --- a/onnxruntime/test/optimizer/graph_transform_test.cc +++ b/onnxruntime/test/optimizer/graph_transform_test.cc @@ -3053,6 +3053,12 @@ TEST_F(GraphTransformationTests, BiasDropoutFusionTest) { TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_residual_fusion_mismatch.onnx", *logger_, 1); TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_residual_fusion_multiple_consumers1.onnx", *logger_, 1); TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_residual_fusion_multiple_consumers2.onnx", *logger_, 1); + TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_same_shape_fusion.onnx", *logger_); + TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_residual_same_shape_fusion.onnx", *logger_); + TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_fusion_dim_is_param.onnx", *logger_); + TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_residual_fusion_dim_is_param.onnx", *logger_); + TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_same_shape_fusion_dim_is_param.onnx", *logger_); + TestBiasDropoutFusion(MODEL_FOLDER "fusion/bias_dropout_residual_same_shape_fusion_dim_is_param.onnx", *logger_); } TEST_F(GraphTransformationTests, LayerNormFusionTest) { diff --git a/onnxruntime/test/testdata/transform/fusion/bias_dropout_fusion_dim_is_param.onnx b/onnxruntime/test/testdata/transform/fusion/bias_dropout_fusion_dim_is_param.onnx new file mode 100644 index 0000000000000000000000000000000000000000..f61f761c2a3a3d2d2ba02f39e65b2fc2806ab350 GIT binary patch literal 314 zcmd;J7vjm!%d5~$tw_u*$Vs(Y$i$`0#p)PenwQPR>I9;MI1*D*4C3=kOQcvpEG1^g zloT#EF1RokcTr+VW`2BfeqM127jH>XVrE`uUV40PeoCqkt25XPmfXbRY$=YEqWl7& zK?X|fE)bTGuv2DYaXge4?^arznV+X6D!{13a`udmy%Vz0S|S3RT8xZNs5V4tL7gk4 y$R*FgC?w6rF2o9ScDx~sVFY6sM@hi^#LvYGRLcpMNdlSzaVcR_oLIOR1b6`Og;d%A literal 0 HcmV?d00001 diff --git a/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_fusion_dim_is_param.onnx b/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_fusion_dim_is_param.onnx new file mode 100644 index 0000000000000000000000000000000000000000..47898128f646beadcaf60e4691132c376ff9121e GIT binary patch literal 404 zcmd;J7vjm!%d5~$tw_u*$Vs*O&BUe5#p)PenwQPR>I9;MI1*D*4C3=kOQcvpEG1^g zloT!>F1RokcTr+VW`2BfeqM127jH>XVrE`uUV40PeoCqkcS=!y0Z=bUzYt4qVsW+< z2UN&FiQNSv$)(1Htd5H{2y8K{Gl&A)Y6!MfNZ2Veu{a*8E8eZNI5R&_OH_bSiRJ7W zAA2WchiHiiaB49!I-$BJN(<^aAw@2E4n`qqE_NYSpe6B!FoqF~VH_m^3juyEUZ7e| bgbdIWh!Y8$k_1!>_6HF*I@E;VE;%k_bzF==LX6H}+YG_B2?;x8CKks-RmZ!P7H8(?X^9FjDzThB<74lH>;Nqh z0ZuJOMkiEvL@5D%BP7cu&A})n&c!aoTAG&~ZwO-;aWQi=NJJ@vy-Pq5(4Zg!1|#E}1ag27t22lK+iD25RY=$=GqE@xsw>{Dv^X1nTf^mP+q)SX>n$Lo|dQpqY}&6Gd}iC$i`}k2ykjKGCHA}9;F0yqmV3@GzX)Q pI2XGRYiV9~ydjKX#Kp|fAQ7bu_5^-KNkD^~2^i$W!o?uK0|2(pO9lV{ literal 0 HcmV?d00001 diff --git a/onnxruntime/test/testdata/transform/fusion/bias_dropout_same_shape_fusion_dim_is_param.onnx b/onnxruntime/test/testdata/transform/fusion/bias_dropout_same_shape_fusion_dim_is_param.onnx new file mode 100644 index 0000000000000000000000000000000000000000..7bf25a713bb2b5e60ff6e2491bff5f2ac73d3987 GIT binary patch literal 334 zcmd;J7vjm!%d5~$tw_u*$Vs)@%EYC@#p)PenwQPR2BwUJI1*D*4C3=kOQcvpEG1^g zloT#EF1RokcTr+VW`2BfeqM127jH>XVrE`uUV40PeoCqkt24+fA(q_4;%q67l%o6s zpg{&o>@EGPN+6SX+hm9 uq{t=D!6+on#V*7Ov@+fh#xR00jH9&So+7L`31~3H@r2bnv2Za6@Bjez#aXEU literal 0 HcmV?d00001