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 0000000000..f61f761c2a Binary files /dev/null and b/onnxruntime/test/testdata/transform/fusion/bias_dropout_fusion_dim_is_param.onnx differ 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 0000000000..47898128f6 Binary files /dev/null and b/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_fusion_dim_is_param.onnx differ diff --git a/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_gen.py b/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_gen.py index b0642194c6..aa010fb409 100644 --- a/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_gen.py +++ b/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_gen.py @@ -133,3 +133,102 @@ graph = helper.make_graph( model = helper.make_model(graph, producer_name='onnx-example', **kwargs) onnx.save(model, 'bias_dropout_residual_fusion_multiple_consumers2.onnx') + + +# Create the model (ModelProto) +A2 = helper.make_tensor_value_info('A2', TensorProto.FLOAT, ['unk_1', 'unk_2', 3072]) + +bias = helper.make_node("Add", ["A", "A2"], ["add0_out"], "add0") +dropout_12 = helper.make_node("Dropout", ["add0_out", "ratio_const", "training_mode"], ["C", "mask"], "dropout0") + +graph = helper.make_graph( + [bias, dropout_12], + "Bias_Dropout_Fusion", #name + [A, A2], + [C], + [ratio, training_mode]) + +model = helper.make_model(graph, producer_name='onnx-example', **kwargs) +onnx.save(model, 'bias_dropout_same_shape_fusion.onnx') + +# Create the model (ModelProto) +bias = helper.make_node("Add", ["A", "A2"], ["add0_out"], "add0") +dropout_12 = helper.make_node("Dropout", ["add0_out", "ratio_const", "training_mode"], ["dropout_out", "mask"], "dropout0") +residual = helper.make_node("Add", ["dropout_out", "R"], ["C"], "add1") + +graph = helper.make_graph( + [bias, dropout_12, residual], + "Bias_Dropout_Fusion", #name + [A, A2, R], + [C], + [ratio, training_mode]) + +model = helper.make_model(graph, producer_name='onnx-example', **kwargs) +onnx.save(model, 'bias_dropout_residual_same_shape_fusion.onnx') + + +# Create the model (ModelProto) +A_unk = helper.make_tensor_value_info('A_unk', TensorProto.FLOAT, ['unk_1', 'unk_2', 'unk_3']) +B_unk = helper.make_tensor_value_info('B_unk', TensorProto.FLOAT, ['unk_3']) +C_unk = helper.make_tensor_value_info('C_unk', TensorProto.FLOAT, ['unk_1', 'unk_2', 'unk_3']) + +bias = helper.make_node("Add", ["A_unk", "B_unk"], ["add0_out"], "add0") +dropout_12 = helper.make_node("Dropout", ["add0_out", "ratio_const", "training_mode"], ["C_unk", "mask"], "dropout0") + +graph = helper.make_graph( + [bias, dropout_12], + "Bias_Dropout_Fusion", #name + [A_unk, B_unk], + [C_unk], + [ratio, training_mode]) + +model = helper.make_model(graph, producer_name='onnx-example', **kwargs) +onnx.save(model, 'bias_dropout_fusion_dim_is_param.onnx') + +# Create the model (ModelProto) +R_unk = helper.make_tensor_value_info('R_unk', TensorProto.FLOAT, ['unk_1', 'unk_2', 'unk_3']) + +bias = helper.make_node("Add", ["A_unk", "B_unk"], ["add0_out"], "add0") +dropout_12 = helper.make_node("Dropout", ["add0_out", "ratio_const", "training_mode"], ["dropout_out", "mask"], "dropout0") +residual = helper.make_node("Add", ["dropout_out", "R_unk"], ["C_unk"], "add1") + +graph = helper.make_graph( + [bias, dropout_12, residual], + "Bias_Dropout_Fusion", #name + [A_unk, B_unk, R_unk], + [C_unk], + [ratio, training_mode]) + +model = helper.make_model(graph, producer_name='onnx-example', **kwargs) +onnx.save(model, 'bias_dropout_residual_fusion_dim_is_param.onnx') + +# Create the model (ModelProto) +A_unk2 = helper.make_tensor_value_info('A_unk2', TensorProto.FLOAT, ['unk_1', 'unk_2', 'unk_3']) + +bias = helper.make_node("Add", ["A_unk", "A_unk2"], ["add0_out"], "add0") +dropout_12 = helper.make_node("Dropout", ["add0_out", "ratio_const", "training_mode"], ["C_unk", "mask"], "dropout0") + +graph = helper.make_graph( + [bias, dropout_12], + "Bias_Dropout_Fusion", #name + [A_unk, A_unk2], + [C_unk], + [ratio, training_mode]) + +model = helper.make_model(graph, producer_name='onnx-example', **kwargs) +onnx.save(model, 'bias_dropout_same_shape_fusion_dim_is_param.onnx') + +# Create the model (ModelProto) +bias = helper.make_node("Add", ["A_unk", "A_unk2"], ["add0_out"], "add0") +dropout_12 = helper.make_node("Dropout", ["add0_out", "ratio_const", "training_mode"], ["dropout_out", "mask"], "dropout0") +residual = helper.make_node("Add", ["dropout_out", "R_unk"], ["C_unk"], "add1") + +graph = helper.make_graph( + [bias, dropout_12, residual], + "Bias_Dropout_Fusion", #name + [A_unk, A_unk2, R_unk], + [C_unk], + [ratio, training_mode]) + +model = helper.make_model(graph, producer_name='onnx-example', **kwargs) +onnx.save(model, 'bias_dropout_residual_same_shape_fusion_dim_is_param.onnx') diff --git a/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_same_shape_fusion.onnx b/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_same_shape_fusion.onnx new file mode 100644 index 0000000000..56488c30a1 Binary files /dev/null and b/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_same_shape_fusion.onnx differ diff --git a/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_same_shape_fusion_dim_is_param.onnx b/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_same_shape_fusion_dim_is_param.onnx new file mode 100644 index 0000000000..caaedef826 Binary files /dev/null and b/onnxruntime/test/testdata/transform/fusion/bias_dropout_residual_same_shape_fusion_dim_is_param.onnx differ diff --git a/onnxruntime/test/testdata/transform/fusion/bias_dropout_same_shape_fusion.onnx b/onnxruntime/test/testdata/transform/fusion/bias_dropout_same_shape_fusion.onnx new file mode 100644 index 0000000000..51a0b64b98 Binary files /dev/null and b/onnxruntime/test/testdata/transform/fusion/bias_dropout_same_shape_fusion.onnx differ 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 0000000000..7bf25a713b Binary files /dev/null and b/onnxruntime/test/testdata/transform/fusion/bias_dropout_same_shape_fusion_dim_is_param.onnx differ