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 <pengwa@microsoft.com>
This commit is contained in:
mindest 2021-10-12 19:57:38 +08:00 committed by GitHub
parent 2f1204a5d5
commit f9cf62912a
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
15 changed files with 230 additions and 55 deletions

View file

@ -359,7 +359,7 @@ This version of the operator has been available since version 1 of the 'com.micr
<dt><tt>data</tt> : T</dt>
<dd>The input data as Tensor.</dd>
<dt><tt>bias</tt> : T</dt>
<dd>The bias input, a vector with the same shape as last dim of data</dd>
<dd>The bias input, a vector with the same shape as last dim of data OR same shape with data</dd>
<dt><tt>residual</tt> (optional) : T</dt>
<dd>The residual input, must have the same shape as data</dd>
<dt><tt>ratio</tt> (optional) : T1</dt>

View file

@ -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<T>::MappedType CudaT;
const CudaT* X_data = reinterpret_cast<const CudaT*>(X.template Data<T>());
@ -52,7 +53,7 @@ struct BiasDropoutComputeImpl {
CudaT* Y_data = reinterpret_cast<CudaT*>(Y.template MutableData<T>());
BiasDropoutKernelImpl<CudaT>(prop, stream, N, fdm_dim, ratio_data, generator, X_data, bias_data, residual_data, Y_data, mask_data);
BiasDropoutKernelImpl<CudaT>(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<Tensor>(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<ALL_IEEE_FLOAT_DATA_TYPES> t_disp(X->GetElementType());
return t_disp.InvokeRet<Status, BiasDropoutComputeImpl>(
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
} // namespace onnxruntime

View file

@ -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:

View file

@ -28,7 +28,7 @@ namespace cuda {
constexpr int UNROLL = 4;
template <typename T, bool has_residual>
template <typename T, bool has_same_shape_bias, bool has_residual>
__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 <typename T, bool has_residual>
template <typename T, bool has_same_shape_bias, bool has_residual>
__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<T, UNROLL>;
using MaskLoadT = aligned_vector<bool, UNROLL>;
@ -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<LoadT*>(&bias_vec);
*value0 = *reinterpret_cast<const LoadT*>(&bias_data[id]);
}
T src[UNROLL];
LoadT *value1 = reinterpret_cast<LoadT*>(&src);
*value1 = *reinterpret_cast<const LoadT*>(&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<int>(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<T, false><<<grid_size, block_size, 0, stream>>>(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<T, true, false><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
} else {
BiasDropoutKernel<T, true, true><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
}
} else {
BiasDropoutKernel<T, true><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
if (residual_data == nullptr) {
BiasDropoutKernel<T, false, false><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
} else {
BiasDropoutKernel<T, false, true><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
}
}
} else {
if (residual_data == nullptr) {
BiasDropoutVectorizedKernel<T, false><<<grid_size, block_size, 0, stream>>>(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<T, true, false><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
} else {
BiasDropoutVectorizedKernel<T, true, true><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
}
} else {
BiasDropoutVectorizedKernel<T, true><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
if (residual_data == nullptr) {
BiasDropoutVectorizedKernel<T, false, false><<<grid_size, block_size, 0, stream>>>(N, fdm_dim, ratio, seeds, X_data, bias_data, residual_data, Y_data, mask_data);
} else {
BiasDropoutVectorizedKernel<T, false, true><<<grid_size, block_size, 0, stream>>>(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)

View file

@ -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, "

View file

@ -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<NodeArg*>& dropout_input,
std::vector<NodeArg*>& 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,
{},

View file

@ -30,7 +30,8 @@ enum TrainingMode { TrainingFalse,
#if defined(USE_CUDA) || defined(USE_ROCM)
namespace {
void RunBiasDropoutTest(const bool use_mask, const std::vector<int64_t>& 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<int64_t>& input_s
const std::vector<float> input = ValueRange(input_size, 1.0f, 1.0f);
t.AddInput("data", input_shape, input);
std::vector<int64_t> bias_shape{input_shape.back()};
const auto bias_size = input_shape.back();
std::vector<int64_t> 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<float> 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

View file

@ -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) {

View file

@ -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')