From 311eea2258e1b064e59b41c5b7d31a69bc21e9b1 Mon Sep 17 00:00:00 2001 From: hhhfccz Date: Mon, 5 Dec 2022 13:09:27 +0000 Subject: [PATCH 1/5] add fused get target_offsets --- oneflow/core/functional/functional_api.yaml | 4 + oneflow/core/functional/impl/math_functor.cpp | 26 ++++ oneflow/ir/include/OneFlow/OneFlowUserOps.td | 21 +++- .../fused_get_target_offsets_kernel.cu | 115 ++++++++++++++++++ .../user/ops/fused_get_target_offsets_op.cpp | 68 +++++++++++ .../modules/test_fused_get_target_offsets.py | 72 +++++++++++ 6 files changed, 304 insertions(+), 2 deletions(-) create mode 100644 oneflow/user/kernels/fused_get_target_offsets_kernel.cu create mode 100644 oneflow/user/ops/fused_get_target_offsets_op.cpp create mode 100644 python/oneflow/test/modules/test_fused_get_target_offsets.py diff --git a/oneflow/core/functional/functional_api.yaml b/oneflow/core/functional/functional_api.yaml index 41c87ad5777..a9a4e64fbee 100644 --- a/oneflow/core/functional/functional_api.yaml +++ b/oneflow/core/functional/functional_api.yaml @@ -2506,6 +2506,10 @@ signature: "TensorTuple (Tensor c2_diff, Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Float eps) => FusedGetConvexDiagonalSquaredGrad" bind_python: False +- name: "fused_get_target_offsets" + signature: "Tensor (Tensor gxy, Tensor gxi, Float g) => FusedGetTargetOffsets" + bind_python: True + - name: "grouped_matmul_bias" signature: "TensorTuple (TensorTuple xs, TensorTuple weights, TensorTuple biases) => GroupedMatmulBias" bind_python: True diff --git a/oneflow/core/functional/impl/math_functor.cpp b/oneflow/core/functional/impl/math_functor.cpp index eff9f311492..d37d4b3155f 100644 --- a/oneflow/core/functional/impl/math_functor.cpp +++ b/oneflow/core/functional/impl/math_functor.cpp @@ -3630,6 +3630,31 @@ class FusedGetConvexDiagonalSquaredGradFunctor { std::shared_ptr op_; }; +class FusedGetTargetOffsetsFunctor { + public: + FusedGetTargetOffsetsFunctor() { + op_ = CHECK_JUST(one::OpBuilder("fused_get_target_offsets") + .Input("gxy") + .Input("gxi") + .Output("j") + .Build()); + } + + Maybe operator()(const std::shared_ptr& gxy, + const std::shared_ptr& gxi, + const float g) const { + auto& attrs = THREAD_CACHED_MUTABLE_ATTR_MAP("g"); + attrs.SetAllAttrs(g); + auto shape = gxy->shape(); + CHECK_EQ_OR_RETURN(shape->NumAxes(), 2); + CHECK_EQ_OR_RETURN(shape->At(1), 2); + return OpInterpUtil::Dispatch(*op_, {gxy, gxi}, attrs); + } + + private: + std::shared_ptr op_; +}; + } // namespace impl using namespace impl; @@ -3761,6 +3786,7 @@ ONEFLOW_FUNCTION_LIBRARY(m) { m.add_functor("FusedGetConvexDiagonalSquared"); m.add_functor( "FusedGetConvexDiagonalSquaredGrad"); + m.add_functor("FusedGetTargetOffsets"); }; } // namespace functional diff --git a/oneflow/ir/include/OneFlow/OneFlowUserOps.td b/oneflow/ir/include/OneFlow/OneFlowUserOps.td index e766e50f799..ef4c48cbcac 100644 --- a/oneflow/ir/include/OneFlow/OneFlowUserOps.td +++ b/oneflow/ir/include/OneFlow/OneFlowUserOps.td @@ -2264,8 +2264,8 @@ def OneFlow_EagerSymmetricSToPOp : OneFlow_BaseOp<"eager_symmetric_s_to_p", [NoS #endif // GET_ONEFLOW_EAGER_OP_DEFINITIONS // Group: FUSED -// cudnn_fused_normalization_add_relu, cudnn_fused_normalization_add_relu_grad, fused_bias_add_gelu, fused_bias_add_gelu_grad, fused_bias_add_mask_scale, fused_cast_scale, fused_scale_mask_softmax, fused_scale_mask_softmax_dropout, fused_scale_mask_softmax_dropout_grad, fused_scale_mask_softmax_grad, fused_scale_tril, fused_self_attention_query_mul_key_and_value, fused_self_attention_query_mul_key_and_value_grad, fused_tril_scale_softmax_mask_scale, fused_tril_scale_softmax_mask_scale_grad, normalization_add_relu_grad, fused_dot_feature_interaction, fused_dot_feature_interaction_grad, fused_cross_feature_interaction, fused_cross_feature_interaction_grad_v1, fused_cross_feature_interaction_grad_v2, fused_lstm_cell, fused_lstm_cell_grad, fused_gru_cell, fused_gru_cell_grad, fused_multi_head_attention_inference, fused_fast_gelu_mul, fused_fast_gelu_mul_grad, fused_get_boundding_boxes_coord, fused_get_boundding_boxes_coord_grad, fused_get_center_dist, fused_get_center_dist_grad, fused_get_ciou_result, fused_get_ciou_result_grad, fused_get_ciou_diagonal_angle, fused_get_ciou_diagonal_angle_grad, fused_get_intersection_area, fused_get_intersection_area_grad -// Total: 42 +// cudnn_fused_normalization_add_relu, cudnn_fused_normalization_add_relu_grad, fused_bias_add_gelu, fused_bias_add_gelu_grad, fused_bias_add_mask_scale, fused_cast_scale, fused_scale_mask_softmax, fused_scale_mask_softmax_dropout, fused_scale_mask_softmax_dropout_grad, fused_scale_mask_softmax_grad, fused_scale_tril, fused_self_attention_query_mul_key_and_value, fused_self_attention_query_mul_key_and_value_grad, fused_tril_scale_softmax_mask_scale, fused_tril_scale_softmax_mask_scale_grad, normalization_add_relu_grad, fused_dot_feature_interaction, fused_dot_feature_interaction_grad, fused_cross_feature_interaction, fused_cross_feature_interaction_grad_v1, fused_cross_feature_interaction_grad_v2, fused_lstm_cell, fused_lstm_cell_grad, fused_gru_cell, fused_gru_cell_grad, fused_multi_head_attention_inference, fused_fast_gelu_mul, fused_fast_gelu_mul_grad, fused_get_boundding_boxes_coord, fused_get_boundding_boxes_coord_grad, fused_get_center_dist, fused_get_center_dist_grad, fused_get_ciou_result, fused_get_ciou_result_grad, fused_get_ciou_diagonal_angle, fused_get_ciou_diagonal_angle_grad, fused_get_intersection_area, fused_get_intersection_area_grad, fused_get_target_offsets +// Total: 45 #ifdef GET_ONEFLOW_FUSED_OP_DEFINITIONS @@ -3211,6 +3211,23 @@ def OneFlow_FusedGetConvexDiagonalSquaredGradOp : OneFlow_BaseOp<"fused_get_conv let has_data_type_infer_fn = 1; } +def OneFlow_FusedGetTargetOffsetsOp : OneFlow_BaseOp<"fused_get_target_offsets", [NoSideEffect, DeclareOpInterfaceMethods]> { + let input = (ins + OneFlow_Tensor:$gxy, + OneFlow_Tensor:$gxi + ); + let output = (outs + OneFlow_Tensor:$j + ); + let attrs = (ins + DefaultValuedAttr:$g + ); + let has_logical_tensor_desc_infer_fn = 1; + let has_physical_tensor_desc_infer_fn = 1; + let has_get_sbp_fn = 1; + let has_data_type_infer_fn = 1; +} + #endif // GET_ONEFLOW_FUSED_OP_DEFINITIONS // Group: IDEMPOTENT diff --git a/oneflow/user/kernels/fused_get_target_offsets_kernel.cu b/oneflow/user/kernels/fused_get_target_offsets_kernel.cu new file mode 100644 index 00000000000..7a597bc8105 --- /dev/null +++ b/oneflow/user/kernels/fused_get_target_offsets_kernel.cu @@ -0,0 +1,115 @@ +/* +Copyright 2020 The OneFlow Authors. All rights reserved. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +*/ +#include "oneflow/core/common/data_type.pb.h" +#include "oneflow/core/device/cuda_util.h" +#include "oneflow/core/framework/framework.h" +#include "oneflow/core/framework/op_kernel.h" +#include "oneflow/core/kernel/new_kernel_util.h" +#include "oneflow/core/ep/common/primitive/binary_functor.h" + +namespace oneflow { + +namespace { + +template +struct ModFunctor { + __device__ T Compute(T input_tensor) const { + return fmod(input_tensor, static_cast(1.0)); + } +}; + +template<> +struct ModFunctor { + ModFunctor float_functor; + __device__ half Compute(half input_tensor) const { + return __float2half(float_functor.Compute(__half2float(input_tensor))); + } +}; + +template +struct GetStatsFunctor { + __device__ bool Compute(T input_tensor, T input_tensor_mod_1, float g) const { + return (input_tensor_mod_1 < static_cast(g)) && (input_tensor > static_cast(1.0)); + } +}; + +template +__global__ void FusedGetTargetOffsetsForward(MOD_FUNCTOR mod_functor, + GET_STATUS_FUNCTOR get_stats_functor, + const int n, const T* gxy, const T* gxi, + const float g, bool* output_tensor, + const int64_t rows) { + CUDA_1D_KERNEL_LOOP(i, n) { + const T gxy_i = gxy[i]; + const T gxi_i = gxi[i]; + const T gxy_mod_1 = mod_functor.Compute(gxy_i); + const T gxi_mod_1 = mod_functor.Compute(gxi_i); + const bool stats_1 = get_stats_functor.Compute(gxy_i, gxy_mod_1, g); + const bool stats_2 = get_stats_functor.Compute(gxi_i, gxi_mod_1, g); + if (i % 2 == 0) { + const int64_t extra_cols = i / 2; + output_tensor[i - extra_cols + rows] = stats_1; + output_tensor[i + n - extra_cols + rows] = stats_2; + } else { + const int64_t extra_cols = (i + n - 1) / 2; + output_tensor[extra_cols + rows] = stats_1; + output_tensor[n + extra_cols + rows] = stats_2; + } + } +} + +} // namespace + +template +class FusedGetTargetOffsetsKernel final : public user_op::OpKernel { + public: + FusedGetTargetOffsetsKernel() = default; + ~FusedGetTargetOffsetsKernel() = default; + + private: + using user_op::OpKernel::Compute; + void Compute(user_op::KernelComputeContext* ctx) const override { + const user_op::Tensor* gxy = ctx->Tensor4ArgNameAndIndex("gxy", 0); + const user_op::Tensor* gxi = ctx->Tensor4ArgNameAndIndex("gxi", 0); + const float g = ctx->Attr("g"); + + user_op::Tensor* j = ctx->Tensor4ArgNameAndIndex("j", 0); + + const int64_t elem_cnt = gxy->shape_view().elem_cnt(); + const int64_t rows = gxy->shape_view().At(0); + + ModFunctor mod_functor{}; + GetStatsFunctor get_stats_functor{}; + + RUN_CUDA_KERNEL((FusedGetTargetOffsetsForward), ctx->stream(), elem_cnt, + mod_functor, get_stats_functor, elem_cnt, + gxy->dptr(), gxi->dptr(), g, j->mut_dptr(), rows); + } + bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } +}; + +#define REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(dtype) \ + REGISTER_USER_KERNEL("fused_get_target_offsets") \ + .SetCreateFn>() \ + .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ + && (user_op::HobDataType("j", 0) == DataType::kBool) \ + && (user_op::HobDataType("gxy", 0) == GetDataType::value)); + +REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(float) +REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(double) +REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(half) + +} // namespace oneflow diff --git a/oneflow/user/ops/fused_get_target_offsets_op.cpp b/oneflow/user/ops/fused_get_target_offsets_op.cpp new file mode 100644 index 00000000000..6d2372ed2be --- /dev/null +++ b/oneflow/user/ops/fused_get_target_offsets_op.cpp @@ -0,0 +1,68 @@ +/* +Copyright 2020 The OneFlow Authors. All rights reserved. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +*/ +#include "oneflow/core/framework/framework.h" +#include "oneflow/core/framework/op_generated.h" + +namespace oneflow { + +Maybe FusedGetTargetOffsetsOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { + const user_op::TensorDesc& gxy = ctx->InputTensorDesc("gxy", 0); + const user_op::TensorDesc& gxi = ctx->InputTensorDesc("gxi", 0); + + const Shape& gxy_shape = gxy.shape(); + + CHECK_EQ_OR_RETURN(gxy.shape().NumAxes(), 2); + CHECK_EQ_OR_RETURN(gxy.shape().At(1), 2); + CHECK_EQ_OR_RETURN(gxy.shape(), gxi.shape()); + + user_op::TensorDesc* j = ctx->MutOutputTensorDesc("j", 0); + j->set_is_dynamic(gxy.is_dynamic()); + j->set_shape(Shape({gxy_shape.At(1) * 2 + 1, gxy_shape.At(0)})); + + return Maybe::Ok(); +} + +Maybe FusedGetTargetOffsetsOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedGetTargetOffsetsOp::InferLogicalTensorDesc(ctx); +} + +Maybe FusedGetTargetOffsetsOp::InferDataType(user_op::InferContext* ctx) { + const user_op::TensorDesc& gxy = ctx->InputTensorDesc("gxy", 0); + const user_op::TensorDesc& gxi = ctx->InputTensorDesc("gxi", 0); + + CHECK_EQ_OR_RETURN(gxy.data_type(), gxi.data_type()); + + user_op::TensorDesc* j = ctx->MutOutputTensorDesc("j", 0); + j->set_data_type(DataType::kBool); + // j->set_data_type(gxy.data_type()); + return Maybe::Ok(); +} + +Maybe FusedGetTargetOffsetsOp::GetSbp(user_op::SbpContext* ctx) { + const user_op::TensorDesc& gxy = ctx->LogicalTensorDesc4InputArgNameAndIndex("gxy", 0); + FOR_RANGE(int64_t, i, 0, gxy.shape().NumAxes()) { + if (i != 1) { + ctx->NewBuilder() + .Split(user_op::OpArg("gxy", 0), i) + .Split(user_op::OpArg("gxi", 0), i) + .Split(user_op::OpArg("j", 0), i) + .Build(); + } + } + return Maybe::Ok(); +} + +} // namespace oneflow diff --git a/python/oneflow/test/modules/test_fused_get_target_offsets.py b/python/oneflow/test/modules/test_fused_get_target_offsets.py new file mode 100644 index 00000000000..a84376066fe --- /dev/null +++ b/python/oneflow/test/modules/test_fused_get_target_offsets.py @@ -0,0 +1,72 @@ +""" +Copyright 2020 The OneFlow Authors. All rights reserved. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +""" +import unittest +from collections import OrderedDict + +import numpy as np +import torch +from oneflow.test_utils.test_util import GenArgList + +import oneflow as flow +import oneflow.unittest + + +def torch_get_target_offsets(gxy, gxi, g): + j, k = ((gxy % 1 < g) & (gxy > 1)).T + l, m = ((gxi % 1 < g) & (gxi > 1)).T + return torch.stack((j, k, l, m, torch.ones_like(j))) + + +def _test_fused_get_target_offsets_impl(test_case, device, shape, g): + x = [] + torch_x = [] + for _ in range(3): + tmp = flow.tensor( + np.random.randn(*shape), + dtype=flow.float32, + device=flow.device(device), + requires_grad=False, + ) + x.append(tmp) + torch_x.append( + torch.tensor( + tmp.numpy(), + dtype=torch.float32, + device=torch.device(device), + requires_grad=False, + ) + ) + gxy, gxi = x[0] + x[1], x[0] + x[2] + torch_gxy, torch_gxi = torch_x[0] + torch_x[1], torch_x[0] + torch_x[2] + j = flow._C.fused_get_target_offsets(gxy, gxi, g) + torch_j = torch_get_target_offsets(torch_gxy, torch_gxi, g) + test_case.assertTrue((j.cpu().numpy() == torch_j.cpu().numpy()).any()) + + +@flow.unittest.skip_unless_1n1d() +class TestGetTargetOffsetsModule(flow.unittest.TestCase): + def test_fused_get_target_offsets_area(test_case): + arg_dict = OrderedDict() + arg_dict["test_fun"] = [_test_fused_get_target_offsets_impl] + arg_dict["device"] = ["cuda"] + arg_dict["shape"] = [(1234, 2), (583, 2), (128, 2)] + arg_dict["g"] = [0.1, 0.5, 0.9] + for arg in GenArgList(arg_dict): + arg[0](test_case, *arg[1:]) + + +if __name__ == "__main__": + unittest.main() From 42702477dfb41d77aa8b2283d75d8469840df6af Mon Sep 17 00:00:00 2001 From: hhhfccz Date: Tue, 6 Dec 2022 02:27:55 +0000 Subject: [PATCH 2/5] fix: rename to yolov5 --- .../autograd/gradient_funcs/fused_center.cpp | 2 +- .../fused_get_boundding_boxes_coord.cpp | 2 +- .../fused_get_ciou_diagonal_angle.cpp | 2 +- .../gradient_funcs/fused_get_ciou_result.cpp | 2 +- .../fused_get_convex_diagonal_squared.cpp | 2 +- .../fused_get_intersection_area.cpp | 2 +- .../autograd/gradient_funcs/fused_get_iou.cpp | 2 +- oneflow/core/functional/functional_api.yaml | 30 ++++++++--------- oneflow/core/functional/impl/math_functor.cpp | 30 ++++++++--------- oneflow/ir/include/OneFlow/OneFlowUserOps.td | 32 +++++++++---------- oneflow/user/kernels/fused_center_kernel.cu | 4 +-- .../fused_get_bounding_boxes_coord_kernel.cu | 4 +-- .../fused_get_ciou_diagonal_angle_kernel.cu | 4 +-- .../kernels/fused_get_ciou_result_kernel.cu | 4 +-- ...used_get_convex_diagonal_squared_kernel.cu | 4 +-- .../fused_get_intersection_area_kernel.cu | 4 +-- oneflow/user/kernels/fused_get_iou_kernel.cu | 4 +-- .../fused_get_target_offsets_kernel.cu | 2 +- .../test_fused_get_boundding_boxes_coord.py | 2 +- ...enter.py => test_fused_get_center_dist.py} | 2 +- .../test_fused_get_ciou_diagonal_angle.py | 2 +- .../modules/test_fused_get_ciou_result.py | 2 +- .../test_fused_get_convex_diagonal_squared.py | 2 +- .../test_fused_get_intersection_area.py | 2 +- .../test/modules/test_fused_get_iou.py | 2 +- .../modules/test_fused_get_target_offsets.py | 2 +- 26 files changed, 76 insertions(+), 76 deletions(-) rename python/oneflow/test/modules/{test_fused_center.py => test_fused_get_center_dist.py} (98%) diff --git a/oneflow/core/autograd/gradient_funcs/fused_center.cpp b/oneflow/core/autograd/gradient_funcs/fused_center.cpp index 02c65930182..e31609b6e00 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_center.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_center.cpp @@ -66,7 +66,7 @@ class FusedCenterGrad : public OpExprGradFunction { } }; -REGISTER_OP_EXPR_GRAD_FUNCTION("fused_get_center_dist", FusedCenterGrad); +REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_center_dist", FusedCenterGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp b/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp index 7555efdb146..6580e2eee10 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp @@ -65,7 +65,7 @@ class FusedGetBounddingBoxesCoordGrad } }; -REGISTER_OP_EXPR_GRAD_FUNCTION("fused_get_boundding_boxes_coord", FusedGetBounddingBoxesCoordGrad); +REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_boundding_boxes_coord", FusedGetBounddingBoxesCoordGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp b/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp index 61f45e8b555..dd1b6348ce6 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp @@ -67,7 +67,7 @@ class FusedGetCiouDiagonalAngleGrad : public OpExprGradFunction { AttrMap base_attrs_; }; -REGISTER_OP_EXPR_GRAD_FUNCTION("fused_get_iou", FusedGetIouGrad); +REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_iou", FusedGetIouGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/functional/functional_api.yaml b/oneflow/core/functional/functional_api.yaml index a9a4e64fbee..2a79e76924e 100644 --- a/oneflow/core/functional/functional_api.yaml +++ b/oneflow/core/functional/functional_api.yaml @@ -1662,27 +1662,27 @@ Bool align_corners=False, Int64List[3] output_size=None, String data_format="channels_first") => UpsampleTrilinear3DGrad' bind_python: False -- name: "fused_get_boundding_boxes_coord" +- name: "fused_yolov5_get_boundding_boxes_coord" signature: "TensorTuple (Tensor x1, Tensor y1, Tensor w1, Tensor h1, Tensor x2, Tensor y2, Tensor w2, Tensor h2) => FusedGetBounddingBoxesCoord" bind_python: True -- name: "fused_get_boundding_boxes_coord_grad" +- name: "fused_yolov5_get_boundding_boxes_coord_grad" signature: "TensorTuple (Tensor b1_x1_diff, Tensor b1_x2_diff, Tensor b1_y1_diff, Tensor b1_y2_diff, Tensor b2_x1_diff, Tensor b2_x2_diff, Tensor b2_y1_diff, Tensor b2_y2_diff) => FusedGetBounddingBoxesCoordGrad" bind_python: False -- name: "fused_get_ciou_result" +- name: "fused_yolov5_get_ciou_result" signature: "TensorTuple (Tensor v, Tensor iou, Tensor rho2, Tensor c2, Float eps) => FusedGetCiouResult" bind_python: True -- name: "fused_get_ciou_result_grad" +- name: "fused_yolov5_get_ciou_result_grad" signature: "TensorTuple (Tensor dy ,Tensor alpha, Tensor rho2, Tensor c2) => FusedGetCiouResultGrad" bind_python: False -- name: "fused_get_iou" +- name: "fused_yolov5_get_iou" signature: "Tensor (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor inter, Float eps) => FusedGetIou" bind_python: True -- name: "fused_get_iou_grad" +- name: "fused_yolov5_get_iou_grad" signature: "TensorTuple (Tensor diou, Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor inter, Float eps) => FusedGetIouGrad" bind_python: False @@ -2474,39 +2474,39 @@ signature: "Tensor (Tensor query, Tensor key, Tensor value, Int64 num_heads, Bool causal=False, Int64 query_hidden_slice_start=0, Int64 query_hidden_slice_end=-1, Int64 key_hidden_slice_start=0, Int64 key_hidden_slice_end=-1, Int64 value_hidden_slice_start=0, Int64 value_hidden_slice_end=-1) => FusedMultiHeadAttentionInference" bind_python: True -- name: "fused_get_center_dist" +- name: "fused_yolov5_get_center_dist" signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2) => FusedCenter" bind_python: True -- name: "fused_get_center_dist_grad" +- name: "fused_yolov5_get_center_dist_grad" signature: "TensorTuple (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Tensor rho2_diff) => FusedCenterGrad" bind_python: False -- name: "fused_get_intersection_area" +- name: "fused_yolov5_get_intersection_area" signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2) => FusedGetIntersectionArea" bind_python: True -- name: "fused_get_intersection_area_grad" +- name: "fused_yolov5_get_intersection_area_grad" signature: "TensorTuple (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Tensor inter_diff) => FusedGetIntersectionAreaGrad" bind_python: False -- name: "fused_get_ciou_diagonal_angle" +- name: "fused_yolov5_get_ciou_diagonal_angle" signature: "Tensor (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Float eps) => FusedGetCiouDiagonalAngle" bind_python: True -- name: "fused_get_ciou_diagonal_angle_grad" +- name: "fused_yolov5_get_ciou_diagonal_angle_grad" signature: "TensorTuple (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor v_diff, Float eps) => FusedGetCiouDiagonalAngleGrad" bind_python: False -- name: "fused_get_convex_diagonal_squared" +- name: "fused_yolov5_get_convex_diagonal_squared" signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Float eps) => FusedGetConvexDiagonalSquared" bind_python: True -- name: "fused_get_convex_diagonal_squared_grad" +- name: "fused_yolov5_get_convex_diagonal_squared_grad" signature: "TensorTuple (Tensor c2_diff, Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Float eps) => FusedGetConvexDiagonalSquaredGrad" bind_python: False -- name: "fused_get_target_offsets" +- name: "fused_yolov5_get_target_offsets" signature: "Tensor (Tensor gxy, Tensor gxi, Float g) => FusedGetTargetOffsets" bind_python: True diff --git a/oneflow/core/functional/impl/math_functor.cpp b/oneflow/core/functional/impl/math_functor.cpp index d37d4b3155f..f7fe43ac5e7 100644 --- a/oneflow/core/functional/impl/math_functor.cpp +++ b/oneflow/core/functional/impl/math_functor.cpp @@ -3184,7 +3184,7 @@ class FusedWeightedSumFunctor { class FusedCenterFunctor { public: FusedCenterFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_center_dist") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_center_dist") .Input("b1_x1") .Input("b1_x2") .Input("b2_x1") @@ -3213,7 +3213,7 @@ class FusedCenterFunctor { class FusedCenterGradFunctor { public: FusedCenterGradFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_center_dist_grad") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_center_dist_grad") .Input("b1_x1") .Input("b1_x2") .Input("b2_x1") @@ -3251,7 +3251,7 @@ class FusedCenterGradFunctor { class FusedGetIntersectionAreaFunctor { public: FusedGetIntersectionAreaFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_intersection_area") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_intersection_area") .Input("b1_x1") .Input("b1_x2") .Input("b2_x1") @@ -3280,7 +3280,7 @@ class FusedGetIntersectionAreaFunctor { class FusedGetIntersectionAreaGradFunctor { public: FusedGetIntersectionAreaGradFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_intersection_area_grad") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_intersection_area_grad") .Input("b1_x1") .Input("b1_x2") .Input("b2_x1") @@ -3318,7 +3318,7 @@ class FusedGetIntersectionAreaGradFunctor { class FusedGetBounddingBoxesCoordFunctor { public: FusedGetBounddingBoxesCoordFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_boundding_boxes_coord") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_boundding_boxes_coord") .Input("x1") .Input("y1") .Input("w1") @@ -3353,7 +3353,7 @@ class FusedGetBounddingBoxesCoordFunctor { class FusedGetBounddingBoxesCoordGradFunctor { public: FusedGetBounddingBoxesCoordGradFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_boundding_boxes_coord_grad") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_boundding_boxes_coord_grad") .Input("b1_x1_diff") .Input("b1_x2_diff") .Input("b1_y1_diff") @@ -3394,7 +3394,7 @@ class FusedGetBounddingBoxesCoordGradFunctor { class FusedGetCiouDiagonalAngleFunctor { public: FusedGetCiouDiagonalAngleFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_ciou_diagonal_angle") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_ciou_diagonal_angle") .Input("w1") .Input("h1") .Input("w2") @@ -3419,7 +3419,7 @@ class FusedGetCiouDiagonalAngleFunctor { class FusedGetCiouResultFunctor { public: FusedGetCiouResultFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_ciou_result") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_ciou_result") .Input("v") .Input("iou") .Input("rho2") @@ -3445,7 +3445,7 @@ class FusedGetCiouResultFunctor { class FusedGetCiouDiagonalAngleGradFunctor { public: FusedGetCiouDiagonalAngleGradFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_ciou_diagonal_angle_grad") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_ciou_diagonal_angle_grad") .Input("w1") .Input("h1") .Input("w2") @@ -3475,7 +3475,7 @@ class FusedGetCiouDiagonalAngleGradFunctor { class FusedGetCiouResultGradFunctor { public: FusedGetCiouResultGradFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_ciou_result_grad") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_ciou_result_grad") .Input("dy") .Input("alpha") .Input("rho2") @@ -3501,7 +3501,7 @@ class FusedGetCiouResultGradFunctor { class FusedGetIouFunctor { public: FusedGetIouFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_iou") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_iou") .Input("w1") .Input("h1") .Input("w2") @@ -3528,7 +3528,7 @@ class FusedGetIouFunctor { class FusedGetIouGradFunctor { public: FusedGetIouGradFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_iou_grad") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_iou_grad") .Input("diou") .Input("w1") .Input("h1") @@ -3559,7 +3559,7 @@ class FusedGetIouGradFunctor { class FusedGetConvexDiagonalSquaredFunctor { public: FusedGetConvexDiagonalSquaredFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_convex_diagonal_squared") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_convex_diagonal_squared") .Input("b1_x1") .Input("b1_x2") .Input("b2_x1") @@ -3593,7 +3593,7 @@ class FusedGetConvexDiagonalSquaredFunctor { class FusedGetConvexDiagonalSquaredGradFunctor { public: FusedGetConvexDiagonalSquaredGradFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_convex_diagonal_squared_grad") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_convex_diagonal_squared_grad") .Input("c2_diff") .Input("b1_x1") .Input("b1_x2") @@ -3633,7 +3633,7 @@ class FusedGetConvexDiagonalSquaredGradFunctor { class FusedGetTargetOffsetsFunctor { public: FusedGetTargetOffsetsFunctor() { - op_ = CHECK_JUST(one::OpBuilder("fused_get_target_offsets") + op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_target_offsets") .Input("gxy") .Input("gxi") .Output("j") diff --git a/oneflow/ir/include/OneFlow/OneFlowUserOps.td b/oneflow/ir/include/OneFlow/OneFlowUserOps.td index ef4c48cbcac..81fd9c31066 100644 --- a/oneflow/ir/include/OneFlow/OneFlowUserOps.td +++ b/oneflow/ir/include/OneFlow/OneFlowUserOps.td @@ -2264,7 +2264,7 @@ def OneFlow_EagerSymmetricSToPOp : OneFlow_BaseOp<"eager_symmetric_s_to_p", [NoS #endif // GET_ONEFLOW_EAGER_OP_DEFINITIONS // Group: FUSED -// cudnn_fused_normalization_add_relu, cudnn_fused_normalization_add_relu_grad, fused_bias_add_gelu, fused_bias_add_gelu_grad, fused_bias_add_mask_scale, fused_cast_scale, fused_scale_mask_softmax, fused_scale_mask_softmax_dropout, fused_scale_mask_softmax_dropout_grad, fused_scale_mask_softmax_grad, fused_scale_tril, fused_self_attention_query_mul_key_and_value, fused_self_attention_query_mul_key_and_value_grad, fused_tril_scale_softmax_mask_scale, fused_tril_scale_softmax_mask_scale_grad, normalization_add_relu_grad, fused_dot_feature_interaction, fused_dot_feature_interaction_grad, fused_cross_feature_interaction, fused_cross_feature_interaction_grad_v1, fused_cross_feature_interaction_grad_v2, fused_lstm_cell, fused_lstm_cell_grad, fused_gru_cell, fused_gru_cell_grad, fused_multi_head_attention_inference, fused_fast_gelu_mul, fused_fast_gelu_mul_grad, fused_get_boundding_boxes_coord, fused_get_boundding_boxes_coord_grad, fused_get_center_dist, fused_get_center_dist_grad, fused_get_ciou_result, fused_get_ciou_result_grad, fused_get_ciou_diagonal_angle, fused_get_ciou_diagonal_angle_grad, fused_get_intersection_area, fused_get_intersection_area_grad, fused_get_target_offsets +// cudnn_fused_normalization_add_relu, cudnn_fused_normalization_add_relu_grad, fused_bias_add_gelu, fused_bias_add_gelu_grad, fused_bias_add_mask_scale, fused_cast_scale, fused_scale_mask_softmax, fused_scale_mask_softmax_dropout, fused_scale_mask_softmax_dropout_grad, fused_scale_mask_softmax_grad, fused_scale_tril, fused_self_attention_query_mul_key_and_value, fused_self_attention_query_mul_key_and_value_grad, fused_tril_scale_softmax_mask_scale, fused_tril_scale_softmax_mask_scale_grad, normalization_add_relu_grad, fused_dot_feature_interaction, fused_dot_feature_interaction_grad, fused_cross_feature_interaction, fused_cross_feature_interaction_grad_v1, fused_cross_feature_interaction_grad_v2, fused_lstm_cell, fused_lstm_cell_grad, fused_gru_cell, fused_gru_cell_grad, fused_multi_head_attention_inference, fused_fast_gelu_mul, fused_fast_gelu_mul_grad, fused_yolov5_get_boundding_boxes_coord, fused_yolov5_get_boundding_boxes_coord_grad, fused_yolov5_get_center_dist, fused_yolov5_get_center_dist_grad, fused_yolov5_get_ciou_result, fused_yolov5_get_ciou_result_grad, fused_yolov5_get_ciou_diagonal_angle, fused_yolov5_get_ciou_diagonal_angle_grad, fused_yolov5_get_intersection_area, fused_yolov5_get_intersection_area_grad, fused_yolov5_get_target_offsets // Total: 45 #ifdef GET_ONEFLOW_FUSED_OP_DEFINITIONS @@ -2880,7 +2880,7 @@ def OneFlow_FusedFastGeluMulGradOp : OneFlow_BaseOp<"fused_fast_gelu_mul_grad", let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetBounddingBoxesCoordOp : OneFlow_BaseOp<"fused_get_boundding_boxes_coord", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetBounddingBoxesCoordOp : OneFlow_BaseOp<"fused_yolov5_get_boundding_boxes_coord", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$x1, OneFlow_Tensor:$y1, @@ -2907,7 +2907,7 @@ def OneFlow_FusedGetBounddingBoxesCoordOp : OneFlow_BaseOp<"fused_get_boundding_ let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetBounddingBoxesCoordGradOp : OneFlow_BaseOp<"fused_get_boundding_boxes_coord_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetBounddingBoxesCoordGradOp : OneFlow_BaseOp<"fused_yolov5_get_boundding_boxes_coord_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1_diff, OneFlow_Tensor:$b1_x2_diff, @@ -2934,7 +2934,7 @@ def OneFlow_FusedGetBounddingBoxesCoordGradOp : OneFlow_BaseOp<"fused_get_boundd let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetCiouResultOp : OneFlow_BaseOp<"fused_get_ciou_result", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetCiouResultOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_result", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$v, OneFlow_Tensor:$iou, @@ -2954,7 +2954,7 @@ def OneFlow_FusedGetCiouResultOp : OneFlow_BaseOp<"fused_get_ciou_result", [NoSi let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetCiouResultGradOp : OneFlow_BaseOp<"fused_get_ciou_result_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetCiouResultGradOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_result_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$dy, OneFlow_Tensor:$alpha, @@ -2973,7 +2973,7 @@ def OneFlow_FusedGetCiouResultGradOp : OneFlow_BaseOp<"fused_get_ciou_result_gra let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetIouOp : OneFlow_BaseOp<"fused_get_iou", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetIouOp : OneFlow_BaseOp<"fused_yolov5_get_iou", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$w1, OneFlow_Tensor:$h1, @@ -2993,7 +2993,7 @@ def OneFlow_FusedGetIouOp : OneFlow_BaseOp<"fused_get_iou", [NoSideEffect, Decla let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetIouGradOp : OneFlow_BaseOp<"fused_get_iou_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetIouGradOp : OneFlow_BaseOp<"fused_yolov5_get_iou_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$diou, OneFlow_Tensor:$w1, @@ -3019,7 +3019,7 @@ def OneFlow_FusedGetIouGradOp : OneFlow_BaseOp<"fused_get_iou_grad", [NoSideEffe let has_data_type_infer_fn = 1; } -def OneFlow_FusedCenterOp : OneFlow_BaseOp<"fused_get_center_dist", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedCenterOp : OneFlow_BaseOp<"fused_yolov5_get_center_dist", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3039,7 +3039,7 @@ def OneFlow_FusedCenterOp : OneFlow_BaseOp<"fused_get_center_dist", [NoSideEffec let has_data_type_infer_fn = 1; } -def OneFlow_FusedCenterGradOp : OneFlow_BaseOp<"fused_get_center_dist_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedCenterGradOp : OneFlow_BaseOp<"fused_yolov5_get_center_dist_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3067,7 +3067,7 @@ def OneFlow_FusedCenterGradOp : OneFlow_BaseOp<"fused_get_center_dist_grad", [No let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetCiouDiagonalAngleOp : OneFlow_BaseOp<"fused_get_ciou_diagonal_angle", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetCiouDiagonalAngleOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_diagonal_angle", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$w1, OneFlow_Tensor:$h1, @@ -3086,7 +3086,7 @@ def OneFlow_FusedGetCiouDiagonalAngleOp : OneFlow_BaseOp<"fused_get_ciou_diagona let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetCiouDiagonalAngleGradOp : OneFlow_BaseOp<"fused_get_ciou_diagonal_angle_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetCiouDiagonalAngleGradOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_diagonal_angle_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$w1, OneFlow_Tensor:$h1, @@ -3109,7 +3109,7 @@ def OneFlow_FusedGetCiouDiagonalAngleGradOp : OneFlow_BaseOp<"fused_get_ciou_dia let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetIntersectionAreaOp : OneFlow_BaseOp<"fused_get_intersection_area", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetIntersectionAreaOp : OneFlow_BaseOp<"fused_yolov5_get_intersection_area", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3129,7 +3129,7 @@ def OneFlow_FusedGetIntersectionAreaOp : OneFlow_BaseOp<"fused_get_intersection_ let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetIntersectionAreaGradOp : OneFlow_BaseOp<"fused_get_intersection_area_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetIntersectionAreaGradOp : OneFlow_BaseOp<"fused_yolov5_get_intersection_area_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3157,7 +3157,7 @@ def OneFlow_FusedGetIntersectionAreaGradOp : OneFlow_BaseOp<"fused_get_intersect let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetConvexDiagonalSquaredOp : OneFlow_BaseOp<"fused_get_convex_diagonal_squared", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetConvexDiagonalSquaredOp : OneFlow_BaseOp<"fused_yolov5_get_convex_diagonal_squared", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3180,7 +3180,7 @@ def OneFlow_FusedGetConvexDiagonalSquaredOp : OneFlow_BaseOp<"fused_get_convex_d let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetConvexDiagonalSquaredGradOp : OneFlow_BaseOp<"fused_get_convex_diagonal_squared_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetConvexDiagonalSquaredGradOp : OneFlow_BaseOp<"fused_yolov5_get_convex_diagonal_squared_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$c2_diff, OneFlow_Tensor:$b1_x1, @@ -3211,7 +3211,7 @@ def OneFlow_FusedGetConvexDiagonalSquaredGradOp : OneFlow_BaseOp<"fused_get_conv let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetTargetOffsetsOp : OneFlow_BaseOp<"fused_get_target_offsets", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedGetTargetOffsetsOp : OneFlow_BaseOp<"fused_yolov5_get_target_offsets", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$gxy, OneFlow_Tensor:$gxi diff --git a/oneflow/user/kernels/fused_center_kernel.cu b/oneflow/user/kernels/fused_center_kernel.cu index 7a93bbaf8b8..d31a12c87c3 100644 --- a/oneflow/user/kernels/fused_center_kernel.cu +++ b/oneflow/user/kernels/fused_center_kernel.cu @@ -106,7 +106,7 @@ class FusedCenterKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CENTER_DIST_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_center_dist") \ + REGISTER_USER_KERNEL("fused_yolov5_get_center_dist") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("rho2", 0) == GetDataType::value)); @@ -156,7 +156,7 @@ class FusedCenterGradKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CENTER_DIST_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_center_dist_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_center_dist_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu b/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu index 538c8d60c03..aa6c65bd2f9 100644 --- a/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu +++ b/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu @@ -107,7 +107,7 @@ class FusedGetBounddingBoxesCoordGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_boundding_boxes_coord") \ + REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); @@ -156,7 +156,7 @@ class FusedGetBounddingBoxesCoordGradGpuKernel final : public user_op::OpKernel }; #define REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_boundding_boxes_coord_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1_diff", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu b/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu index 64d843e5fac..f7b2b90d497 100644 --- a/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu +++ b/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu @@ -197,7 +197,7 @@ class FusedGetCiouDiagonalAngleKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_ciou_diagonal_angle") \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("v", 0) == GetDataType::value)); @@ -243,7 +243,7 @@ class FusedGetCiouDiagonalAngleGradKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_ciou_diagonal_angle_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("w1_diff", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_ciou_result_kernel.cu b/oneflow/user/kernels/fused_get_ciou_result_kernel.cu index 7511ce48290..78adc7ca1ee 100644 --- a/oneflow/user/kernels/fused_get_ciou_result_kernel.cu +++ b/oneflow/user/kernels/fused_get_ciou_result_kernel.cu @@ -73,7 +73,7 @@ class FusedGetCiouResultGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CIOU_RESULT_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_ciou_result") \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_result") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("v", 0) == GetDataType::value)); @@ -111,7 +111,7 @@ class FusedGetCiouResultGradGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_ciou_result_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_result_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("dy", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu b/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu index c7e6778db11..9cf5a3318f6 100644 --- a/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu +++ b/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu @@ -91,7 +91,7 @@ class FusedGetConvexDiagonalSquaredKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_convex_diagonal_squared") \ + REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); @@ -143,7 +143,7 @@ class FusedGetConvexDiagonalSquaredGradKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_convex_diagonal_squared_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_intersection_area_kernel.cu b/oneflow/user/kernels/fused_get_intersection_area_kernel.cu index 3a242dd9327..fb52d0a3a3f 100644 --- a/oneflow/user/kernels/fused_get_intersection_area_kernel.cu +++ b/oneflow/user/kernels/fused_get_intersection_area_kernel.cu @@ -125,7 +125,7 @@ class FusedGetIntersectionAreaKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_INTERSECTION_AREA_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_intersection_area") \ + REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("inter", 0) == GetDataType::value)); @@ -179,7 +179,7 @@ class FusedGetIntersectionAreaGradKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_intersection_area_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1_diff", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_iou_kernel.cu b/oneflow/user/kernels/fused_get_iou_kernel.cu index 83d16c40f07..35e2e763c94 100644 --- a/oneflow/user/kernels/fused_get_iou_kernel.cu +++ b/oneflow/user/kernels/fused_get_iou_kernel.cu @@ -74,7 +74,7 @@ class FusedGetIouGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_IOU_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_iou") \ + REGISTER_USER_KERNEL("fused_yolov5_get_iou") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("iou", 0) == GetDataType::value)); @@ -115,7 +115,7 @@ class FusedGetIouGradGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_IOU_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_iou_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_iou_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("diou", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_target_offsets_kernel.cu b/oneflow/user/kernels/fused_get_target_offsets_kernel.cu index 7a597bc8105..45e83eba51a 100644 --- a/oneflow/user/kernels/fused_get_target_offsets_kernel.cu +++ b/oneflow/user/kernels/fused_get_target_offsets_kernel.cu @@ -102,7 +102,7 @@ class FusedGetTargetOffsetsKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_get_target_offsets") \ + REGISTER_USER_KERNEL("fused_yolov5_get_target_offsets") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("j", 0) == DataType::kBool) \ diff --git a/python/oneflow/test/modules/test_fused_get_boundding_boxes_coord.py b/python/oneflow/test/modules/test_fused_get_boundding_boxes_coord.py index f9257423bab..3c4c53bb127 100644 --- a/python/oneflow/test/modules/test_fused_get_boundding_boxes_coord.py +++ b/python/oneflow/test/modules/test_fused_get_boundding_boxes_coord.py @@ -54,7 +54,7 @@ def _test_get_boundding_boxes_coord_impl(test_case, device, shape): b2_x2, b2_y1, b2_y2, - ) = flow._C.fused_get_boundding_boxes_coord(x1, y1, w1, h1, x2, y2, w2, h2) + ) = flow._C.fused_yolov5_get_boundding_boxes_coord(x1, y1, w1, h1, x2, y2, w2, h2) torch_x1, torch_y1, torch_w1, torch_h1, torch_x2, torch_y2, torch_w2, torch_h2 = ( torch_x[0], torch_x[1], diff --git a/python/oneflow/test/modules/test_fused_center.py b/python/oneflow/test/modules/test_fused_get_center_dist.py similarity index 98% rename from python/oneflow/test/modules/test_fused_center.py rename to python/oneflow/test/modules/test_fused_get_center_dist.py index 524733dfeb4..d746597f446 100644 --- a/python/oneflow/test/modules/test_fused_center.py +++ b/python/oneflow/test/modules/test_fused_get_center_dist.py @@ -87,7 +87,7 @@ def compare(a, b, rtol=1e-5, atol=1e-5): torch_x[6], torch_x[7], ) - rho2 = flow._C.fused_get_center_dist( + rho2 = flow._C.fused_yolov5_get_center_dist( b1_x1, b1_x2, b2_x1, b2_x2, b1_y1, b1_y2, b2_y1, b2_y2 ) torch_rho2 = torch_center( diff --git a/python/oneflow/test/modules/test_fused_get_ciou_diagonal_angle.py b/python/oneflow/test/modules/test_fused_get_ciou_diagonal_angle.py index 20d343a5ca9..38647a9788b 100644 --- a/python/oneflow/test/modules/test_fused_get_ciou_diagonal_angle.py +++ b/python/oneflow/test/modules/test_fused_get_ciou_diagonal_angle.py @@ -71,7 +71,7 @@ def compare(a, b, rtol=1e-5, atol=1e-5): torch_x[2], torch_x[3], ) - v = flow._C.fused_get_ciou_diagonal_angle(w1, h1, w2, h2, eps=1e-8) + v = flow._C.fused_yolov5_get_ciou_diagonal_angle(w1, h1, w2, h2, eps=1e-8) torch_v = torch_get_ciou_diagonal_angle(torch_w1, torch_h1, torch_w2, torch_h2,) compare(v, torch_v) diff --git a/python/oneflow/test/modules/test_fused_get_ciou_result.py b/python/oneflow/test/modules/test_fused_get_ciou_result.py index 36594d59fa1..47ebf041fd4 100644 --- a/python/oneflow/test/modules/test_fused_get_ciou_result.py +++ b/python/oneflow/test/modules/test_fused_get_ciou_result.py @@ -46,7 +46,7 @@ def _test_get_ciou_result_impl(test_case, device, shape): ) ) v, iou, rho2, c2 = x[0], x[1], x[2], x[3] - y = flow._C.fused_get_ciou_result(v, iou, rho2, c2, eps)[0] + y = flow._C.fused_yolov5_get_ciou_result(v, iou, rho2, c2, eps)[0] torch_v, torch_iou, torch_rho2, torch_c2 = ( torch_x[0], torch_x[1], diff --git a/python/oneflow/test/modules/test_fused_get_convex_diagonal_squared.py b/python/oneflow/test/modules/test_fused_get_convex_diagonal_squared.py index a0d935d8604..415bd876d00 100644 --- a/python/oneflow/test/modules/test_fused_get_convex_diagonal_squared.py +++ b/python/oneflow/test/modules/test_fused_get_convex_diagonal_squared.py @@ -91,7 +91,7 @@ def compare(a, b, rtol=1e-5, atol=1e-5): torch_x[6], torch_x[7], ) - c2 = flow._C.fused_get_convex_diagonal_squared( + c2 = flow._C.fused_yolov5_get_convex_diagonal_squared( b1_x1, b1_x2, b2_x1, b2_x2, b1_y1, b1_y2, b2_y1, b2_y2, eps ) torch_c2 = torch_fused_get_convex_diagonal_squared( diff --git a/python/oneflow/test/modules/test_fused_get_intersection_area.py b/python/oneflow/test/modules/test_fused_get_intersection_area.py index 19fcebdc546..e77280b6287 100644 --- a/python/oneflow/test/modules/test_fused_get_intersection_area.py +++ b/python/oneflow/test/modules/test_fused_get_intersection_area.py @@ -86,7 +86,7 @@ def compare(a, b, rtol=1e-5, atol=1e-5): torch_x[6], torch_x[7], ) - inter = flow._C.fused_get_intersection_area( + inter = flow._C.fused_yolov5_get_intersection_area( b1_x1, b1_x2, b2_x1, b2_x2, b1_y1, b1_y2, b2_y1, b2_y2 ) torch_inter = torch_get_intersection_area( diff --git a/python/oneflow/test/modules/test_fused_get_iou.py b/python/oneflow/test/modules/test_fused_get_iou.py index 9e5a08f4913..c009294edf6 100644 --- a/python/oneflow/test/modules/test_fused_get_iou.py +++ b/python/oneflow/test/modules/test_fused_get_iou.py @@ -46,7 +46,7 @@ def _test_get_iou_impl(test_case, device, shape): ) ) w1, h1, w2, h2, inter = x[0], x[1], x[2], x[3], x[4] - iou = flow._C.fused_get_iou(w1, h1, w2, h2, inter, eps) + iou = flow._C.fused_yolov5_get_iou(w1, h1, w2, h2, inter, eps) torch_w1, torch_h1, torch_w2, torch_h2, torch_inter = ( torch_x[0], torch_x[1], diff --git a/python/oneflow/test/modules/test_fused_get_target_offsets.py b/python/oneflow/test/modules/test_fused_get_target_offsets.py index a84376066fe..27f3330a0c0 100644 --- a/python/oneflow/test/modules/test_fused_get_target_offsets.py +++ b/python/oneflow/test/modules/test_fused_get_target_offsets.py @@ -51,7 +51,7 @@ def _test_fused_get_target_offsets_impl(test_case, device, shape, g): ) gxy, gxi = x[0] + x[1], x[0] + x[2] torch_gxy, torch_gxi = torch_x[0] + torch_x[1], torch_x[0] + torch_x[2] - j = flow._C.fused_get_target_offsets(gxy, gxi, g) + j = flow._C.fused_yolov5_get_target_offsets(gxy, gxi, g) torch_j = torch_get_target_offsets(torch_gxy, torch_gxi, g) test_case.assertTrue((j.cpu().numpy() == torch_j.cpu().numpy()).any()) From 8eb2e9db5b71f09ec5a01a983e72991635e31b38 Mon Sep 17 00:00:00 2001 From: hhhfccz Date: Tue, 6 Dec 2022 02:29:35 +0000 Subject: [PATCH 3/5] format --- .../fused_get_boundding_boxes_coord.cpp | 3 ++- .../fused_get_ciou_diagonal_angle.cpp | 3 ++- oneflow/core/functional/impl/math_functor.cpp | 3 +-- oneflow/user/kernels/fused_center_kernel.cu | 4 +-- .../fused_get_bounding_boxes_coord_kernel.cu | 4 +-- .../fused_get_ciou_diagonal_angle_kernel.cu | 4 +-- .../kernels/fused_get_ciou_result_kernel.cu | 4 +-- ...used_get_convex_diagonal_squared_kernel.cu | 4 +-- .../fused_get_intersection_area_kernel.cu | 4 +-- oneflow/user/kernels/fused_get_iou_kernel.cu | 4 +-- .../fused_get_target_offsets_kernel.cu | 26 +++++++++---------- 11 files changed, 31 insertions(+), 32 deletions(-) diff --git a/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp b/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp index 6580e2eee10..392f6e59006 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp @@ -65,7 +65,8 @@ class FusedGetBounddingBoxesCoordGrad } }; -REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_boundding_boxes_coord", FusedGetBounddingBoxesCoordGrad); +REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_boundding_boxes_coord", + FusedGetBounddingBoxesCoordGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp b/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp index dd1b6348ce6..30d0dc3a169 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp @@ -67,7 +67,8 @@ class FusedGetCiouDiagonalAngleGrad : public OpExprGradFunction operator()(const std::shared_ptr& gxy, - const std::shared_ptr& gxi, - const float g) const { + const std::shared_ptr& gxi, const float g) const { auto& attrs = THREAD_CACHED_MUTABLE_ATTR_MAP("g"); attrs.SetAllAttrs(g); auto shape = gxy->shape(); diff --git a/oneflow/user/kernels/fused_center_kernel.cu b/oneflow/user/kernels/fused_center_kernel.cu index d31a12c87c3..5de62719283 100644 --- a/oneflow/user/kernels/fused_center_kernel.cu +++ b/oneflow/user/kernels/fused_center_kernel.cu @@ -106,7 +106,7 @@ class FusedCenterKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CENTER_DIST_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_center_dist") \ + REGISTER_USER_KERNEL("fused_yolov5_get_center_dist") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("rho2", 0) == GetDataType::value)); @@ -156,7 +156,7 @@ class FusedCenterGradKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CENTER_DIST_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_center_dist_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_center_dist_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu b/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu index aa6c65bd2f9..0e0b00f09a3 100644 --- a/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu +++ b/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu @@ -107,7 +107,7 @@ class FusedGetBounddingBoxesCoordGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord") \ + REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); @@ -156,7 +156,7 @@ class FusedGetBounddingBoxesCoordGradGpuKernel final : public user_op::OpKernel }; #define REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1_diff", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu b/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu index f7b2b90d497..89fb5f857de 100644 --- a/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu +++ b/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu @@ -197,7 +197,7 @@ class FusedGetCiouDiagonalAngleKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle") \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("v", 0) == GetDataType::value)); @@ -243,7 +243,7 @@ class FusedGetCiouDiagonalAngleGradKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("w1_diff", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_ciou_result_kernel.cu b/oneflow/user/kernels/fused_get_ciou_result_kernel.cu index 78adc7ca1ee..b540c994d52 100644 --- a/oneflow/user/kernels/fused_get_ciou_result_kernel.cu +++ b/oneflow/user/kernels/fused_get_ciou_result_kernel.cu @@ -73,7 +73,7 @@ class FusedGetCiouResultGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CIOU_RESULT_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_ciou_result") \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_result") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("v", 0) == GetDataType::value)); @@ -111,7 +111,7 @@ class FusedGetCiouResultGradGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_ciou_result_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_result_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("dy", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu b/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu index 9cf5a3318f6..18619dede4c 100644 --- a/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu +++ b/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu @@ -91,7 +91,7 @@ class FusedGetConvexDiagonalSquaredKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared") \ + REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); @@ -143,7 +143,7 @@ class FusedGetConvexDiagonalSquaredGradKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_intersection_area_kernel.cu b/oneflow/user/kernels/fused_get_intersection_area_kernel.cu index fb52d0a3a3f..5bd4cd79aba 100644 --- a/oneflow/user/kernels/fused_get_intersection_area_kernel.cu +++ b/oneflow/user/kernels/fused_get_intersection_area_kernel.cu @@ -125,7 +125,7 @@ class FusedGetIntersectionAreaKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_INTERSECTION_AREA_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area") \ + REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("inter", 0) == GetDataType::value)); @@ -179,7 +179,7 @@ class FusedGetIntersectionAreaGradKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1_diff", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_iou_kernel.cu b/oneflow/user/kernels/fused_get_iou_kernel.cu index 35e2e763c94..977ecffc467 100644 --- a/oneflow/user/kernels/fused_get_iou_kernel.cu +++ b/oneflow/user/kernels/fused_get_iou_kernel.cu @@ -74,7 +74,7 @@ class FusedGetIouGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_IOU_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_iou") \ + REGISTER_USER_KERNEL("fused_yolov5_get_iou") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("iou", 0) == GetDataType::value)); @@ -115,7 +115,7 @@ class FusedGetIouGradGpuKernel final : public user_op::OpKernel { }; #define REGISTER_FUSED_GET_IOU_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_iou_grad") \ + REGISTER_USER_KERNEL("fused_yolov5_get_iou_grad") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("diou", 0) == GetDataType::value)); diff --git a/oneflow/user/kernels/fused_get_target_offsets_kernel.cu b/oneflow/user/kernels/fused_get_target_offsets_kernel.cu index 45e83eba51a..6af9dc97d58 100644 --- a/oneflow/user/kernels/fused_get_target_offsets_kernel.cu +++ b/oneflow/user/kernels/fused_get_target_offsets_kernel.cu @@ -26,9 +26,7 @@ namespace { template struct ModFunctor { - __device__ T Compute(T input_tensor) const { - return fmod(input_tensor, static_cast(1.0)); - } + __device__ T Compute(T input_tensor) const { return fmod(input_tensor, static_cast(1.0)); } }; template<> @@ -47,11 +45,10 @@ struct GetStatsFunctor { }; template -__global__ void FusedGetTargetOffsetsForward(MOD_FUNCTOR mod_functor, - GET_STATUS_FUNCTOR get_stats_functor, - const int n, const T* gxy, const T* gxi, - const float g, bool* output_tensor, - const int64_t rows) { +__global__ void FusedGetTargetOffsetsForward(MOD_FUNCTOR mod_functor, + GET_STATUS_FUNCTOR get_stats_functor, const int n, + const T* gxy, const T* gxi, const float g, + bool* output_tensor, const int64_t rows) { CUDA_1D_KERNEL_LOOP(i, n) { const T gxy_i = gxy[i]; const T gxi_i = gxi[i]; @@ -59,8 +56,8 @@ __global__ void FusedGetTargetOffsetsForward(MOD_FUNCTOR mod_functor, const T gxi_mod_1 = mod_functor.Compute(gxi_i); const bool stats_1 = get_stats_functor.Compute(gxy_i, gxy_mod_1, g); const bool stats_2 = get_stats_functor.Compute(gxi_i, gxi_mod_1, g); - if (i % 2 == 0) { - const int64_t extra_cols = i / 2; + if (i % 2 == 0) { + const int64_t extra_cols = i / 2; output_tensor[i - extra_cols + rows] = stats_1; output_tensor[i + n - extra_cols + rows] = stats_2; } else { @@ -94,15 +91,16 @@ class FusedGetTargetOffsetsKernel final : public user_op::OpKernel { ModFunctor mod_functor{}; GetStatsFunctor get_stats_functor{}; - RUN_CUDA_KERNEL((FusedGetTargetOffsetsForward), ctx->stream(), elem_cnt, - mod_functor, get_stats_functor, elem_cnt, - gxy->dptr(), gxi->dptr(), g, j->mut_dptr(), rows); + RUN_CUDA_KERNEL( + (FusedGetTargetOffsetsForward), + ctx->stream(), elem_cnt, mod_functor, get_stats_functor, elem_cnt, gxy->dptr(), + gxi->dptr(), g, j->mut_dptr(), rows); } bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; #define REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_target_offsets") \ + REGISTER_USER_KERNEL("fused_yolov5_get_target_offsets") \ .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("j", 0) == DataType::kBool) \ From dab7fc6f3ecf8042a4ded734a9e16da05da396bf Mon Sep 17 00:00:00 2001 From: hhhfccz Date: Thu, 15 Dec 2022 05:47:53 +0000 Subject: [PATCH 4/5] rename to yolov5 --- ...used_yolov5_get_boundding_boxes_coord.cpp} | 6 +- ...r.cpp => fused_yolov5_get_center_dist.cpp} | 8 +- ... fused_yolov5_get_ciou_diagonal_angle.cpp} | 7 +- ...t.cpp => fused_yolov5_get_ciou_result.cpp} | 6 +- ...ed_yolov5_get_convex_diagonal_squared.cpp} | 6 +- ...=> fused_yolov5_get_intersection_area.cpp} | 9 +- ...d_get_iou.cpp => fused_yolov5_get_iou.cpp} | 6 +- oneflow/core/functional/functional_api.yaml | 66 ++++++------- oneflow/core/functional/impl/math_functor.cpp | 97 ++++++++++--------- oneflow/ir/include/OneFlow/OneFlowUserOps.td | 30 +++--- ...yolov5_get_bounding_boxes_coord_kernel.cu} | 40 ++++---- ...=> fused_yolov5_get_center_dist_kernel.cu} | 32 +++--- ..._yolov5_get_ciou_diagonal_angle_kernel.cu} | 40 ++++---- ...=> fused_yolov5_get_ciou_result_kernel.cu} | 32 +++--- ...ov5_get_convex_diagonal_squared_kernel.cu} | 40 ++++---- ...ed_yolov5_get_intersection_area_kernel.cu} | 36 +++---- ...rnel.cu => fused_yolov5_get_iou_kernel.cu} | 32 +++--- ...fused_yolov5_get_target_offsets_kernel.cu} | 16 +-- ...d_yolov5_get_boundding_boxes_coord_op.cpp} | 24 +++-- ...pp => fused_yolov5_get_center_dist_op.cpp} | 20 ++-- ...sed_yolov5_get_ciou_diagonal_angle_op.cpp} | 22 +++-- ...pp => fused_yolov5_get_ciou_result_op.cpp} | 20 ++-- ...yolov5_get_convex_diagonal_squared_op.cpp} | 22 +++-- ...fused_yolov5_get_intersection_area_op.cpp} | 22 +++-- ...iou_op.cpp => fused_yolov5_get_iou_op.cpp} | 20 ++-- ...=> fused_yolov5_get_target_offsets_op.cpp} | 11 +-- ...fused_yolov5_get_boundding_boxes_coord.py} | 0 ...y => test_fused_yolov5_get_center_dist.py} | 0 ...t_fused_yolov5_get_ciou_diagonal_angle.py} | 0 ...y => test_fused_yolov5_get_ciou_result.py} | 0 ...sed_yolov5_get_convex_diagonal_squared.py} | 0 ...est_fused_yolov5_get_intersection_area.py} | 0 ...et_iou.py => test_fused_yolov5_get_iou.py} | 0 ...> test_fused_yolov5_get_target_offsets.py} | 0 34 files changed, 341 insertions(+), 329 deletions(-) rename oneflow/core/autograd/gradient_funcs/{fused_get_boundding_boxes_coord.cpp => fused_yolov5_get_boundding_boxes_coord.cpp} (93%) rename oneflow/core/autograd/gradient_funcs/{fused_center.cpp => fused_yolov5_get_center_dist.cpp} (88%) rename oneflow/core/autograd/gradient_funcs/{fused_get_ciou_diagonal_angle.cpp => fused_yolov5_get_ciou_diagonal_angle.cpp} (89%) rename oneflow/core/autograd/gradient_funcs/{fused_get_ciou_result.cpp => fused_yolov5_get_ciou_result.cpp} (92%) rename oneflow/core/autograd/gradient_funcs/{fused_get_convex_diagonal_squared.cpp => fused_yolov5_get_convex_diagonal_squared.cpp} (94%) rename oneflow/core/autograd/gradient_funcs/{fused_get_intersection_area.cpp => fused_yolov5_get_intersection_area.cpp} (90%) rename oneflow/core/autograd/gradient_funcs/{fused_get_iou.cpp => fused_yolov5_get_iou.cpp} (91%) rename oneflow/user/kernels/{fused_get_bounding_boxes_coord_kernel.cu => fused_yolov5_get_bounding_boxes_coord_kernel.cu} (86%) rename oneflow/user/kernels/{fused_center_kernel.cu => fused_yolov5_get_center_dist_kernel.cu} (87%) rename oneflow/user/kernels/{fused_get_ciou_diagonal_angle_kernel.cu => fused_yolov5_get_ciou_diagonal_angle_kernel.cu} (89%) rename oneflow/user/kernels/{fused_get_ciou_result_kernel.cu => fused_yolov5_get_ciou_result_kernel.cu} (81%) rename oneflow/user/kernels/{fused_get_convex_diagonal_squared_kernel.cu => fused_yolov5_get_convex_diagonal_squared_kernel.cu} (84%) rename oneflow/user/kernels/{fused_get_intersection_area_kernel.cu => fused_yolov5_get_intersection_area_kernel.cu} (88%) rename oneflow/user/kernels/{fused_get_iou_kernel.cu => fused_yolov5_get_iou_kernel.cu} (84%) rename oneflow/user/kernels/{fused_get_target_offsets_kernel.cu => fused_yolov5_get_target_offsets_kernel.cu} (88%) rename oneflow/user/ops/{fused_get_boundding_boxes_coord_op.cpp => fused_yolov5_get_boundding_boxes_coord_op.cpp} (88%) rename oneflow/user/ops/{fused_center_op.cpp => fused_yolov5_get_center_dist_op.cpp} (91%) rename oneflow/user/ops/{fused_get_ciou_diagonal_angle_op.cpp => fused_yolov5_get_ciou_diagonal_angle_op.cpp} (85%) rename oneflow/user/ops/{fused_get_ciou_result_op.cpp => fused_yolov5_get_ciou_result_op.cpp} (81%) rename oneflow/user/ops/{fused_get_convex_diagonal_squared_op.cpp => fused_yolov5_get_convex_diagonal_squared_op.cpp} (85%) rename oneflow/user/ops/{fused_get_intersection_area_op.cpp => fused_yolov5_get_intersection_area_op.cpp} (91%) rename oneflow/user/ops/{fused_get_iou_op.cpp => fused_yolov5_get_iou_op.cpp} (81%) rename oneflow/user/ops/{fused_get_target_offsets_op.cpp => fused_yolov5_get_target_offsets_op.cpp} (81%) rename python/oneflow/test/modules/{test_fused_get_boundding_boxes_coord.py => test_fused_yolov5_get_boundding_boxes_coord.py} (100%) rename python/oneflow/test/modules/{test_fused_get_center_dist.py => test_fused_yolov5_get_center_dist.py} (100%) rename python/oneflow/test/modules/{test_fused_get_ciou_diagonal_angle.py => test_fused_yolov5_get_ciou_diagonal_angle.py} (100%) rename python/oneflow/test/modules/{test_fused_get_ciou_result.py => test_fused_yolov5_get_ciou_result.py} (100%) rename python/oneflow/test/modules/{test_fused_get_convex_diagonal_squared.py => test_fused_yolov5_get_convex_diagonal_squared.py} (100%) rename python/oneflow/test/modules/{test_fused_get_intersection_area.py => test_fused_yolov5_get_intersection_area.py} (100%) rename python/oneflow/test/modules/{test_fused_get_iou.py => test_fused_yolov5_get_iou.py} (100%) rename python/oneflow/test/modules/{test_fused_get_target_offsets.py => test_fused_yolov5_get_target_offsets.py} (100%) diff --git a/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_boundding_boxes_coord.cpp similarity index 93% rename from oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp rename to oneflow/core/autograd/gradient_funcs/fused_yolov5_get_boundding_boxes_coord.cpp index 392f6e59006..7c1c10863ec 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_get_boundding_boxes_coord.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_boundding_boxes_coord.cpp @@ -25,7 +25,7 @@ struct FusedGetBounddingBoxesCoordGradCaptureState : public AutoGradCaptureState std::vector requires_grad; }; -class FusedGetBounddingBoxesCoordGrad +class FusedYolov5GetBounddingBoxesCoordGrad : public OpExprGradFunction { public: Maybe Init(const OpExpr& op) override { return Maybe::Ok(); } @@ -54,7 +54,7 @@ class FusedGetBounddingBoxesCoordGrad const auto& b2_y2_diff = out_grads.at(7); in_grads->resize(8); - auto result = JUST(functional::FusedGetBounddingBoxesCoordGrad( + auto result = JUST(functional::FusedYolov5GetBounddingBoxesCoordGrad( b1_x1_diff, b1_x2_diff, b1_y1_diff, b1_y2_diff, b2_x1_diff, b2_x2_diff, b2_y1_diff, b2_y2_diff)); CHECK_EQ_OR_RETURN(result->size(), INPUT_LEN); @@ -66,7 +66,7 @@ class FusedGetBounddingBoxesCoordGrad }; REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_boundding_boxes_coord", - FusedGetBounddingBoxesCoordGrad); + FusedYolov5GetBounddingBoxesCoordGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/autograd/gradient_funcs/fused_center.cpp b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_center_dist.cpp similarity index 88% rename from oneflow/core/autograd/gradient_funcs/fused_center.cpp rename to oneflow/core/autograd/gradient_funcs/fused_yolov5_get_center_dist.cpp index e31609b6e00..3802cc76920 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_center.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_center_dist.cpp @@ -24,7 +24,7 @@ struct FusedCenterCaptureState : public AutoGradCaptureState { std::vector requires_grad; }; -class FusedCenterGrad : public OpExprGradFunction { +class FusedYolov5GetCenterDistGrad : public OpExprGradFunction { public: Maybe Init(const OpExpr& op) override { return Maybe::Ok(); } @@ -55,8 +55,8 @@ class FusedCenterGrad : public OpExprGradFunction { const auto& b2_y2 = ctx->SavedTensors().at(7); in_grads->resize(INPUT_LEN); - auto result = JUST(functional::FusedCenterGrad(b1_x1, b1_x2, b2_x1, b2_x2, b1_y1, b1_y2, b2_y1, - b2_y2, rho2_diff)); + auto result = JUST(functional::FusedYolov5GetCenterDistGrad(b1_x1, b1_x2, b2_x1, b2_x2, b1_y1, + b1_y2, b2_y1, b2_y2, rho2_diff)); CHECK_EQ_OR_RETURN(result->size(), INPUT_LEN); for (int i = 0; i < INPUT_LEN; i++) { @@ -66,7 +66,7 @@ class FusedCenterGrad : public OpExprGradFunction { } }; -REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_center_dist", FusedCenterGrad); +REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_center_dist", FusedYolov5GetCenterDistGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_ciou_diagonal_angle.cpp similarity index 89% rename from oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp rename to oneflow/core/autograd/gradient_funcs/fused_yolov5_get_ciou_diagonal_angle.cpp index 30d0dc3a169..d919d9b6047 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_get_ciou_diagonal_angle.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_ciou_diagonal_angle.cpp @@ -26,7 +26,7 @@ struct FusedCiouAngleCaptureState : public AutoGradCaptureState { float eps = 1e-8; }; -class FusedGetCiouDiagonalAngleGrad : public OpExprGradFunction { +class FusedYolov5GetCiouDiagonalAngleGrad : public OpExprGradFunction { public: Maybe Init(const OpExpr& op) override { return Maybe::Ok(); } @@ -56,7 +56,8 @@ class FusedGetCiouDiagonalAngleGrad : public OpExprGradFunctionSavedTensors().at(2); const auto& h2 = ctx->SavedTensors().at(3); - auto result = JUST(functional::FusedGetCiouDiagonalAngleGrad(w1, h1, w2, h2, v_diff, ctx->eps)); + auto result = + JUST(functional::FusedYolov5GetCiouDiagonalAngleGrad(w1, h1, w2, h2, v_diff, ctx->eps)); CHECK_EQ_OR_RETURN(result->size(), INPUT_LEN); in_grads->resize(INPUT_LEN); @@ -68,7 +69,7 @@ class FusedGetCiouDiagonalAngleGrad : public OpExprGradFunction { +class FusedYolov5GetCiouResultGrad : public OpExprGradFunction { public: Maybe Init(const OpExpr& op) override { return Maybe::Ok(); } @@ -60,7 +60,7 @@ class FusedGetCiouResultGrad : public OpExprGradFunctionresize(4); - auto result = JUST(functional::FusedGetCiouResultGrad(dy, alpha, rho2, c2)); + auto result = JUST(functional::FusedYolov5GetCiouResultGrad(dy, alpha, rho2, c2)); CHECK_EQ_OR_RETURN(result->size(), 4); if (ctx->v_requires_grad && ctx->iou_requires_grad && ctx->rho2_requires_grad && ctx->c2_requires_grad) { @@ -76,7 +76,7 @@ class FusedGetCiouResultGrad : public OpExprGradFunction { public: Maybe Init(const OpExpr& op) override { @@ -64,7 +64,7 @@ class FusedGetConvexDiagonalSquaredGrad const auto& b2_y2 = ctx->SavedTensors().at(7); in_grads->resize(INPUT_LEN); - auto result = JUST(functional::FusedGetConvexDiagonalSquaredGrad( + auto result = JUST(functional::FusedYolov5GetConvexDiagonalSquaredGrad( c2_diff, b1_x1, b1_x2, b2_x1, b2_x2, b1_y1, b1_y2, b2_y1, b2_y2, ctx->eps)); CHECK_EQ_OR_RETURN(result->size(), INPUT_LEN); @@ -79,7 +79,7 @@ class FusedGetConvexDiagonalSquaredGrad }; REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_convex_diagonal_squared", - FusedGetConvexDiagonalSquaredGrad); + FusedYolov5GetConvexDiagonalSquaredGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/autograd/gradient_funcs/fused_get_intersection_area.cpp b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_intersection_area.cpp similarity index 90% rename from oneflow/core/autograd/gradient_funcs/fused_get_intersection_area.cpp rename to oneflow/core/autograd/gradient_funcs/fused_yolov5_get_intersection_area.cpp index 6e8c1d7d004..504f41ce078 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_get_intersection_area.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_intersection_area.cpp @@ -24,7 +24,7 @@ struct FusedGetIntersectionAreaCaptureState : public AutoGradCaptureState { std::vector requires_grad; }; -class FusedGetIntersectionAreaGrad +class FusedYolov5GetIntersectionAreaGrad : public OpExprGradFunction { public: Maybe Init(const OpExpr& op) override { return Maybe::Ok(); } @@ -56,8 +56,8 @@ class FusedGetIntersectionAreaGrad const auto& b2_y2 = ctx->SavedTensors().at(7); in_grads->resize(INPUT_LEN); - auto result = JUST(functional::FusedGetIntersectionAreaGrad(b1_x1, b1_x2, b2_x1, b2_x2, b1_y1, - b1_y2, b2_y1, b2_y2, rho2_diff)); + auto result = JUST(functional::FusedYolov5GetIntersectionAreaGrad( + b1_x1, b1_x2, b2_x1, b2_x2, b1_y1, b1_y2, b2_y1, b2_y2, rho2_diff)); CHECK_EQ_OR_RETURN(result->size(), INPUT_LEN); for (int i = 0; i < INPUT_LEN; i++) { @@ -67,7 +67,8 @@ class FusedGetIntersectionAreaGrad } }; -REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_intersection_area", FusedGetIntersectionAreaGrad); +REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_intersection_area", + FusedYolov5GetIntersectionAreaGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/autograd/gradient_funcs/fused_get_iou.cpp b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_iou.cpp similarity index 91% rename from oneflow/core/autograd/gradient_funcs/fused_get_iou.cpp rename to oneflow/core/autograd/gradient_funcs/fused_yolov5_get_iou.cpp index b44a4b4d95e..4c836dd5577 100644 --- a/oneflow/core/autograd/gradient_funcs/fused_get_iou.cpp +++ b/oneflow/core/autograd/gradient_funcs/fused_yolov5_get_iou.cpp @@ -26,7 +26,7 @@ struct FusedGetIouGradCaptureState : public AutoGradCaptureState { float eps = 1e-8; }; -class FusedGetIouGrad : public OpExprGradFunction { +class FusedYolov5GetIouGrad : public OpExprGradFunction { public: Maybe Init(const OpExpr& op) override { const UserOpExpr* fw_op_expr = dynamic_cast(&op); @@ -67,7 +67,7 @@ class FusedGetIouGrad : public OpExprGradFunction { const auto& inter = saved_tensors.at(4); in_grads->resize(5); - auto result = JUST(functional::FusedGetIouGrad(diou, w1, h1, w2, h2, inter, ctx->eps)); + auto result = JUST(functional::FusedYolov5GetIouGrad(diou, w1, h1, w2, h2, inter, ctx->eps)); CHECK_EQ_OR_RETURN(result->size(), 3); if (ctx->requires_grad) { in_grads->at(0) = result->at(0); @@ -81,7 +81,7 @@ class FusedGetIouGrad : public OpExprGradFunction { AttrMap base_attrs_; }; -REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_iou", FusedGetIouGrad); +REGISTER_OP_EXPR_GRAD_FUNCTION("fused_yolov5_get_iou", FusedYolov5GetIouGrad); } // namespace one } // namespace oneflow diff --git a/oneflow/core/functional/functional_api.yaml b/oneflow/core/functional/functional_api.yaml index 2a79e76924e..169e2ca44a5 100644 --- a/oneflow/core/functional/functional_api.yaml +++ b/oneflow/core/functional/functional_api.yaml @@ -1662,30 +1662,6 @@ Bool align_corners=False, Int64List[3] output_size=None, String data_format="channels_first") => UpsampleTrilinear3DGrad' bind_python: False -- name: "fused_yolov5_get_boundding_boxes_coord" - signature: "TensorTuple (Tensor x1, Tensor y1, Tensor w1, Tensor h1, Tensor x2, Tensor y2, Tensor w2, Tensor h2) => FusedGetBounddingBoxesCoord" - bind_python: True - -- name: "fused_yolov5_get_boundding_boxes_coord_grad" - signature: "TensorTuple (Tensor b1_x1_diff, Tensor b1_x2_diff, Tensor b1_y1_diff, Tensor b1_y2_diff, Tensor b2_x1_diff, Tensor b2_x2_diff, Tensor b2_y1_diff, Tensor b2_y2_diff) => FusedGetBounddingBoxesCoordGrad" - bind_python: False - -- name: "fused_yolov5_get_ciou_result" - signature: "TensorTuple (Tensor v, Tensor iou, Tensor rho2, Tensor c2, Float eps) => FusedGetCiouResult" - bind_python: True - -- name: "fused_yolov5_get_ciou_result_grad" - signature: "TensorTuple (Tensor dy ,Tensor alpha, Tensor rho2, Tensor c2) => FusedGetCiouResultGrad" - bind_python: False - -- name: "fused_yolov5_get_iou" - signature: "Tensor (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor inter, Float eps) => FusedGetIou" - bind_python: True - -- name: "fused_yolov5_get_iou_grad" - signature: "TensorTuple (Tensor diou, Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor inter, Float eps) => FusedGetIouGrad" - bind_python: False - - name: "abs" signature: "Tensor (Tensor x) => Abs" bind_python: True @@ -2474,40 +2450,64 @@ signature: "Tensor (Tensor query, Tensor key, Tensor value, Int64 num_heads, Bool causal=False, Int64 query_hidden_slice_start=0, Int64 query_hidden_slice_end=-1, Int64 key_hidden_slice_start=0, Int64 key_hidden_slice_end=-1, Int64 value_hidden_slice_start=0, Int64 value_hidden_slice_end=-1) => FusedMultiHeadAttentionInference" bind_python: True +- name: "fused_yolov5_get_boundding_boxes_coord" + signature: "TensorTuple (Tensor x1, Tensor y1, Tensor w1, Tensor h1, Tensor x2, Tensor y2, Tensor w2, Tensor h2) => FusedYolov5GetBounddingBoxesCoord" + bind_python: True + +- name: "fused_yolov5_get_boundding_boxes_coord_grad" + signature: "TensorTuple (Tensor b1_x1_diff, Tensor b1_x2_diff, Tensor b1_y1_diff, Tensor b1_y2_diff, Tensor b2_x1_diff, Tensor b2_x2_diff, Tensor b2_y1_diff, Tensor b2_y2_diff) => FusedYolov5GetBounddingBoxesCoordGrad" + bind_python: False + +- name: "fused_yolov5_get_ciou_result" + signature: "TensorTuple (Tensor v, Tensor iou, Tensor rho2, Tensor c2, Float eps) => FusedYolov5GetCiouResult" + bind_python: True + +- name: "fused_yolov5_get_ciou_result_grad" + signature: "TensorTuple (Tensor dy ,Tensor alpha, Tensor rho2, Tensor c2) => FusedYolov5GetCiouResultGrad" + bind_python: False + +- name: "fused_yolov5_get_iou" + signature: "Tensor (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor inter, Float eps) => FusedYolov5GetIou" + bind_python: True + +- name: "fused_yolov5_get_iou_grad" + signature: "TensorTuple (Tensor diou, Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor inter, Float eps) => FusedYolov5GetIouGrad" + bind_python: False + - name: "fused_yolov5_get_center_dist" - signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2) => FusedCenter" + signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2) => FusedYolov5GetCenterDist" bind_python: True - name: "fused_yolov5_get_center_dist_grad" - signature: "TensorTuple (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Tensor rho2_diff) => FusedCenterGrad" + signature: "TensorTuple (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Tensor rho2_diff) => FusedYolov5GetCenterDistGrad" bind_python: False - name: "fused_yolov5_get_intersection_area" - signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2) => FusedGetIntersectionArea" + signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2) => FusedYolov5GetIntersectionArea" bind_python: True - name: "fused_yolov5_get_intersection_area_grad" - signature: "TensorTuple (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Tensor inter_diff) => FusedGetIntersectionAreaGrad" + signature: "TensorTuple (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Tensor inter_diff) => FusedYolov5GetIntersectionAreaGrad" bind_python: False - name: "fused_yolov5_get_ciou_diagonal_angle" - signature: "Tensor (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Float eps) => FusedGetCiouDiagonalAngle" + signature: "Tensor (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Float eps) => FusedYolov5GetCiouDiagonalAngle" bind_python: True - name: "fused_yolov5_get_ciou_diagonal_angle_grad" - signature: "TensorTuple (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor v_diff, Float eps) => FusedGetCiouDiagonalAngleGrad" + signature: "TensorTuple (Tensor w1, Tensor h1, Tensor w2, Tensor h2, Tensor v_diff, Float eps) => FusedYolov5GetCiouDiagonalAngleGrad" bind_python: False - name: "fused_yolov5_get_convex_diagonal_squared" - signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Float eps) => FusedGetConvexDiagonalSquared" + signature: "Tensor (Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Float eps) => FusedYolov5GetConvexDiagonalSquared" bind_python: True - name: "fused_yolov5_get_convex_diagonal_squared_grad" - signature: "TensorTuple (Tensor c2_diff, Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Float eps) => FusedGetConvexDiagonalSquaredGrad" + signature: "TensorTuple (Tensor c2_diff, Tensor b1_x1, Tensor b1_x2, Tensor b2_x1, Tensor b2_x2, Tensor b1_y1, Tensor b1_y2, Tensor b2_y1, Tensor b2_y2, Float eps) => FusedYolov5GetConvexDiagonalSquaredGrad" bind_python: False - name: "fused_yolov5_get_target_offsets" - signature: "Tensor (Tensor gxy, Tensor gxi, Float g) => FusedGetTargetOffsets" + signature: "Tensor (Tensor gxy, Tensor gxi, Float g) => FusedYolov5GetTargetOffsets" bind_python: True - name: "grouped_matmul_bias" diff --git a/oneflow/core/functional/impl/math_functor.cpp b/oneflow/core/functional/impl/math_functor.cpp index 32b9d54c0fd..8c1ee9d7dc1 100644 --- a/oneflow/core/functional/impl/math_functor.cpp +++ b/oneflow/core/functional/impl/math_functor.cpp @@ -3181,9 +3181,9 @@ class FusedWeightedSumFunctor { std::vector> op_; }; -class FusedCenterFunctor { +class FusedYolov5GetCenterDistFunctor { public: - FusedCenterFunctor() { + FusedYolov5GetCenterDistFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_center_dist") .Input("b1_x1") .Input("b1_x2") @@ -3210,9 +3210,9 @@ class FusedCenterFunctor { std::shared_ptr op_; }; -class FusedCenterGradFunctor { +class FusedYolov5GetCenterDistGradFunctor { public: - FusedCenterGradFunctor() { + FusedYolov5GetCenterDistGradFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_center_dist_grad") .Input("b1_x1") .Input("b1_x2") @@ -3248,9 +3248,9 @@ class FusedCenterGradFunctor { std::shared_ptr op_; }; -class FusedGetIntersectionAreaFunctor { +class FusedYolov5GetIntersectionAreaFunctor { public: - FusedGetIntersectionAreaFunctor() { + FusedYolov5GetIntersectionAreaFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_intersection_area") .Input("b1_x1") .Input("b1_x2") @@ -3277,9 +3277,9 @@ class FusedGetIntersectionAreaFunctor { std::shared_ptr op_; }; -class FusedGetIntersectionAreaGradFunctor { +class FusedYolov5GetIntersectionAreaGradFunctor { public: - FusedGetIntersectionAreaGradFunctor() { + FusedYolov5GetIntersectionAreaGradFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_intersection_area_grad") .Input("b1_x1") .Input("b1_x2") @@ -3315,9 +3315,9 @@ class FusedGetIntersectionAreaGradFunctor { std::shared_ptr op_; }; -class FusedGetBounddingBoxesCoordFunctor { +class FusedYolov5GetBounddingBoxesCoordFunctor { public: - FusedGetBounddingBoxesCoordFunctor() { + FusedYolov5GetBounddingBoxesCoordFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_boundding_boxes_coord") .Input("x1") .Input("y1") @@ -3350,9 +3350,9 @@ class FusedGetBounddingBoxesCoordFunctor { std::shared_ptr op_; }; -class FusedGetBounddingBoxesCoordGradFunctor { +class FusedYolov5GetBounddingBoxesCoordGradFunctor { public: - FusedGetBounddingBoxesCoordGradFunctor() { + FusedYolov5GetBounddingBoxesCoordGradFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_boundding_boxes_coord_grad") .Input("b1_x1_diff") .Input("b1_x2_diff") @@ -3391,9 +3391,9 @@ class FusedGetBounddingBoxesCoordGradFunctor { std::shared_ptr op_; }; -class FusedGetCiouDiagonalAngleFunctor { +class FusedYolov5GetCiouDiagonalAngleFunctor { public: - FusedGetCiouDiagonalAngleFunctor() { + FusedYolov5GetCiouDiagonalAngleFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_ciou_diagonal_angle") .Input("w1") .Input("h1") @@ -3416,9 +3416,9 @@ class FusedGetCiouDiagonalAngleFunctor { std::shared_ptr op_; }; -class FusedGetCiouResultFunctor { +class FusedYolov5GetCiouResultFunctor { public: - FusedGetCiouResultFunctor() { + FusedYolov5GetCiouResultFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_ciou_result") .Input("v") .Input("iou") @@ -3442,9 +3442,9 @@ class FusedGetCiouResultFunctor { std::shared_ptr op_; }; -class FusedGetCiouDiagonalAngleGradFunctor { +class FusedYolov5GetCiouDiagonalAngleGradFunctor { public: - FusedGetCiouDiagonalAngleGradFunctor() { + FusedYolov5GetCiouDiagonalAngleGradFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_ciou_diagonal_angle_grad") .Input("w1") .Input("h1") @@ -3472,9 +3472,9 @@ class FusedGetCiouDiagonalAngleGradFunctor { std::shared_ptr op_; }; -class FusedGetCiouResultGradFunctor { +class FusedYolov5GetCiouResultGradFunctor { public: - FusedGetCiouResultGradFunctor() { + FusedYolov5GetCiouResultGradFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_ciou_result_grad") .Input("dy") .Input("alpha") @@ -3498,9 +3498,9 @@ class FusedGetCiouResultGradFunctor { std::shared_ptr op_; }; -class FusedGetIouFunctor { +class FusedYolov5GetIouFunctor { public: - FusedGetIouFunctor() { + FusedYolov5GetIouFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_iou") .Input("w1") .Input("h1") @@ -3525,9 +3525,9 @@ class FusedGetIouFunctor { std::shared_ptr op_; }; -class FusedGetIouGradFunctor { +class FusedYolov5GetIouGradFunctor { public: - FusedGetIouGradFunctor() { + FusedYolov5GetIouGradFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_iou_grad") .Input("diou") .Input("w1") @@ -3556,9 +3556,9 @@ class FusedGetIouGradFunctor { std::shared_ptr op_; }; -class FusedGetConvexDiagonalSquaredFunctor { +class FusedYolov5GetConvexDiagonalSquaredFunctor { public: - FusedGetConvexDiagonalSquaredFunctor() { + FusedYolov5GetConvexDiagonalSquaredFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_convex_diagonal_squared") .Input("b1_x1") .Input("b1_x2") @@ -3590,9 +3590,9 @@ class FusedGetConvexDiagonalSquaredFunctor { std::shared_ptr op_; }; -class FusedGetConvexDiagonalSquaredGradFunctor { +class FusedYolov5GetConvexDiagonalSquaredGradFunctor { public: - FusedGetConvexDiagonalSquaredGradFunctor() { + FusedYolov5GetConvexDiagonalSquaredGradFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_convex_diagonal_squared_grad") .Input("c2_diff") .Input("b1_x1") @@ -3630,9 +3630,9 @@ class FusedGetConvexDiagonalSquaredGradFunctor { std::shared_ptr op_; }; -class FusedGetTargetOffsetsFunctor { +class FusedYolov5GetTargetOffsetsFunctor { public: - FusedGetTargetOffsetsFunctor() { + FusedYolov5GetTargetOffsetsFunctor() { op_ = CHECK_JUST(one::OpBuilder("fused_yolov5_get_target_offsets") .Input("gxy") .Input("gxi") @@ -3768,24 +3768,25 @@ ONEFLOW_FUNCTION_LIBRARY(m) { m.add_functor("EinSum"); m.add_functor("Inv"); m.add_functor("GeluWithApproximate"); - m.add_functor("Trunc"); - m.add_functor("FusedWeightedSum"); - m.add_functor("FusedCenter"); - m.add_functor("FusedCenterGrad"); - m.add_functor("FusedGetBounddingBoxesCoord"); - m.add_functor("FusedGetBounddingBoxesCoordGrad"); - m.add_functor("FusedGetCiouDiagonalAngle"); - m.add_functor("FusedGetCiouDiagonalAngleGrad"); - m.add_functor("FusedGetCiouResult"); - m.add_functor("FusedGetCiouResultGrad"); - m.add_functor("FusedGetIntersectionArea"); - m.add_functor("FusedGetIntersectionAreaGrad"); - m.add_functor("FusedGetIou"); - m.add_functor("FusedGetIouGrad"); - m.add_functor("FusedGetConvexDiagonalSquared"); - m.add_functor( - "FusedGetConvexDiagonalSquaredGrad"); - m.add_functor("FusedGetTargetOffsets"); + m.add_functor("Trunc"); + m.add_functor("FusedWeightedSum"); + m.add_functor("FusedYolov5GetCenterDist"); + m.add_functor("FusedYolov5GetCenterDistGrad"); + m.add_functor("FusedYolov5GetBounddingBoxesCoord"); + m.add_functor( + "FusedYolov5GetBounddingBoxesCoordGrad"); + m.add_functor("FusedYolov5GetCiouDiagonalAngle"); + m.add_functor("FusedYolov5GetCiouDiagonalAngleGrad"); + m.add_functor("FusedYolov5GetCiouResult"); + m.add_functor("FusedYolov5GetCiouResultGrad"); + m.add_functor("FusedYolov5GetIntersectionArea"); + m.add_functor("FusedYolov5GetIntersectionAreaGrad"); + m.add_functor("FusedYolov5GetIou"); + m.add_functor("FusedYolov5GetIouGrad"); + m.add_functor("FusedYolov5GetConvexDiagonalSquared"); + m.add_functor( + "FusedYolov5GetConvexDiagonalSquaredGrad"); + m.add_functor("FusedYolov5GetTargetOffsets"); }; } // namespace functional diff --git a/oneflow/ir/include/OneFlow/OneFlowUserOps.td b/oneflow/ir/include/OneFlow/OneFlowUserOps.td index 81fd9c31066..12ba63e8896 100644 --- a/oneflow/ir/include/OneFlow/OneFlowUserOps.td +++ b/oneflow/ir/include/OneFlow/OneFlowUserOps.td @@ -2880,7 +2880,7 @@ def OneFlow_FusedFastGeluMulGradOp : OneFlow_BaseOp<"fused_fast_gelu_mul_grad", let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetBounddingBoxesCoordOp : OneFlow_BaseOp<"fused_yolov5_get_boundding_boxes_coord", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetBounddingBoxesCoordOp : OneFlow_BaseOp<"fused_yolov5_get_boundding_boxes_coord", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$x1, OneFlow_Tensor:$y1, @@ -2907,7 +2907,7 @@ def OneFlow_FusedGetBounddingBoxesCoordOp : OneFlow_BaseOp<"fused_yolov5_get_bou let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetBounddingBoxesCoordGradOp : OneFlow_BaseOp<"fused_yolov5_get_boundding_boxes_coord_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetBounddingBoxesCoordGradOp : OneFlow_BaseOp<"fused_yolov5_get_boundding_boxes_coord_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1_diff, OneFlow_Tensor:$b1_x2_diff, @@ -2934,7 +2934,7 @@ def OneFlow_FusedGetBounddingBoxesCoordGradOp : OneFlow_BaseOp<"fused_yolov5_get let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetCiouResultOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_result", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetCiouResultOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_result", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$v, OneFlow_Tensor:$iou, @@ -2954,7 +2954,7 @@ def OneFlow_FusedGetCiouResultOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_result" let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetCiouResultGradOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_result_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetCiouResultGradOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_result_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$dy, OneFlow_Tensor:$alpha, @@ -2973,7 +2973,7 @@ def OneFlow_FusedGetCiouResultGradOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_res let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetIouOp : OneFlow_BaseOp<"fused_yolov5_get_iou", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetIouOp : OneFlow_BaseOp<"fused_yolov5_get_iou", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$w1, OneFlow_Tensor:$h1, @@ -2993,7 +2993,7 @@ def OneFlow_FusedGetIouOp : OneFlow_BaseOp<"fused_yolov5_get_iou", [NoSideEffect let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetIouGradOp : OneFlow_BaseOp<"fused_yolov5_get_iou_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetIouGradOp : OneFlow_BaseOp<"fused_yolov5_get_iou_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$diou, OneFlow_Tensor:$w1, @@ -3019,7 +3019,7 @@ def OneFlow_FusedGetIouGradOp : OneFlow_BaseOp<"fused_yolov5_get_iou_grad", [NoS let has_data_type_infer_fn = 1; } -def OneFlow_FusedCenterOp : OneFlow_BaseOp<"fused_yolov5_get_center_dist", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetCenterDistOp : OneFlow_BaseOp<"fused_yolov5_get_center_dist", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3039,7 +3039,7 @@ def OneFlow_FusedCenterOp : OneFlow_BaseOp<"fused_yolov5_get_center_dist", [NoSi let has_data_type_infer_fn = 1; } -def OneFlow_FusedCenterGradOp : OneFlow_BaseOp<"fused_yolov5_get_center_dist_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetCenterDistGradOp : OneFlow_BaseOp<"fused_yolov5_get_center_dist_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3067,7 +3067,7 @@ def OneFlow_FusedCenterGradOp : OneFlow_BaseOp<"fused_yolov5_get_center_dist_gra let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetCiouDiagonalAngleOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_diagonal_angle", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetCiouDiagonalAngleOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_diagonal_angle", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$w1, OneFlow_Tensor:$h1, @@ -3086,7 +3086,7 @@ def OneFlow_FusedGetCiouDiagonalAngleOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_ let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetCiouDiagonalAngleGradOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_diagonal_angle_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetCiouDiagonalAngleGradOp : OneFlow_BaseOp<"fused_yolov5_get_ciou_diagonal_angle_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$w1, OneFlow_Tensor:$h1, @@ -3109,7 +3109,7 @@ def OneFlow_FusedGetCiouDiagonalAngleGradOp : OneFlow_BaseOp<"fused_yolov5_get_c let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetIntersectionAreaOp : OneFlow_BaseOp<"fused_yolov5_get_intersection_area", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetIntersectionAreaOp : OneFlow_BaseOp<"fused_yolov5_get_intersection_area", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3129,7 +3129,7 @@ def OneFlow_FusedGetIntersectionAreaOp : OneFlow_BaseOp<"fused_yolov5_get_inters let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetIntersectionAreaGradOp : OneFlow_BaseOp<"fused_yolov5_get_intersection_area_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetIntersectionAreaGradOp : OneFlow_BaseOp<"fused_yolov5_get_intersection_area_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3157,7 +3157,7 @@ def OneFlow_FusedGetIntersectionAreaGradOp : OneFlow_BaseOp<"fused_yolov5_get_in let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetConvexDiagonalSquaredOp : OneFlow_BaseOp<"fused_yolov5_get_convex_diagonal_squared", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetConvexDiagonalSquaredOp : OneFlow_BaseOp<"fused_yolov5_get_convex_diagonal_squared", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$b1_x1, OneFlow_Tensor:$b1_x2, @@ -3180,7 +3180,7 @@ def OneFlow_FusedGetConvexDiagonalSquaredOp : OneFlow_BaseOp<"fused_yolov5_get_c let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetConvexDiagonalSquaredGradOp : OneFlow_BaseOp<"fused_yolov5_get_convex_diagonal_squared_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetConvexDiagonalSquaredGradOp : OneFlow_BaseOp<"fused_yolov5_get_convex_diagonal_squared_grad", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$c2_diff, OneFlow_Tensor:$b1_x1, @@ -3211,7 +3211,7 @@ def OneFlow_FusedGetConvexDiagonalSquaredGradOp : OneFlow_BaseOp<"fused_yolov5_g let has_data_type_infer_fn = 1; } -def OneFlow_FusedGetTargetOffsetsOp : OneFlow_BaseOp<"fused_yolov5_get_target_offsets", [NoSideEffect, DeclareOpInterfaceMethods]> { +def OneFlow_FusedYolov5GetTargetOffsetsOp : OneFlow_BaseOp<"fused_yolov5_get_target_offsets", [NoSideEffect, DeclareOpInterfaceMethods]> { let input = (ins OneFlow_Tensor:$gxy, OneFlow_Tensor:$gxi diff --git a/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_bounding_boxes_coord_kernel.cu similarity index 86% rename from oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu rename to oneflow/user/kernels/fused_yolov5_get_bounding_boxes_coord_kernel.cu index 0e0b00f09a3..7e27edc7894 100644 --- a/oneflow/user/kernels/fused_get_bounding_boxes_coord_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_bounding_boxes_coord_kernel.cu @@ -69,10 +69,10 @@ __global__ void FusedGetBounddingBoxesCoordBackward( }; // namespace template -class FusedGetBounddingBoxesCoordGpuKernel final : public user_op::OpKernel { +class FusedYolov5GetBounddingBoxesCoordGpuKernel final : public user_op::OpKernel { public: - FusedGetBounddingBoxesCoordGpuKernel() = default; - ~FusedGetBounddingBoxesCoordGpuKernel() = default; + FusedYolov5GetBounddingBoxesCoordGpuKernel() = default; + ~FusedYolov5GetBounddingBoxesCoordGpuKernel() = default; private: using user_op::OpKernel::Compute; @@ -106,21 +106,21 @@ class FusedGetBounddingBoxesCoordGpuKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord") \ - .SetCreateFn>() \ - .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ +#define REGISTER_FUSED_YOLOV5_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(dtype) \ + REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord") \ + .SetCreateFn>() \ + .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); -REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(float) -REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(half) -REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_BOUNDDING_BOXES_COORD_CUDA_KERNEL(double) template -class FusedGetBounddingBoxesCoordGradGpuKernel final : public user_op::OpKernel { +class FusedYolov5GetBounddingBoxesCoordGradGpuKernel final : public user_op::OpKernel { public: - FusedGetBounddingBoxesCoordGradGpuKernel() = default; - ~FusedGetBounddingBoxesCoordGradGpuKernel() = default; + FusedYolov5GetBounddingBoxesCoordGradGpuKernel() = default; + ~FusedYolov5GetBounddingBoxesCoordGradGpuKernel() = default; private: using user_op::OpKernel::Compute; @@ -155,14 +155,14 @@ class FusedGetBounddingBoxesCoordGradGpuKernel final : public user_op::OpKernel bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord_grad") \ - .SetCreateFn>() \ - .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ +#define REGISTER_FUSED_YOLOV5_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(dtype) \ + REGISTER_USER_KERNEL("fused_yolov5_get_boundding_boxes_coord_grad") \ + .SetCreateFn>() \ + .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1_diff", 0) == GetDataType::value)); -REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(float) -REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(half) -REGISTER_FUSED_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_BOUNDDING_BOXES_COORD_GRAD_CUDA_KERNEL(double) } // namespace oneflow diff --git a/oneflow/user/kernels/fused_center_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_center_dist_kernel.cu similarity index 87% rename from oneflow/user/kernels/fused_center_kernel.cu rename to oneflow/user/kernels/fused_yolov5_get_center_dist_kernel.cu index 5de62719283..85721a3b7d7 100644 --- a/oneflow/user/kernels/fused_center_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_center_dist_kernel.cu @@ -74,10 +74,10 @@ __global__ void FusedCenterBackward(const int n, const T* b1_x1, const T* b1_x2, } // namespace template -class FusedCenterKernel final : public user_op::OpKernel { +class FusedYolov5GetCenterDistKernel final : public user_op::OpKernel { public: - FusedCenterKernel() = default; - ~FusedCenterKernel() = default; + FusedYolov5GetCenterDistKernel() = default; + ~FusedYolov5GetCenterDistKernel() = default; private: using user_op::OpKernel::Compute; @@ -105,21 +105,21 @@ class FusedCenterKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_CENTER_DIST_CUDA_KERNEL(dtype) \ +#define REGISTER_FUSED_YOLOV5_GET_CENTER_DIST_CUDA_KERNEL(dtype) \ REGISTER_USER_KERNEL("fused_yolov5_get_center_dist") \ - .SetCreateFn>() \ + .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("rho2", 0) == GetDataType::value)); -REGISTER_FUSED_GET_CENTER_DIST_CUDA_KERNEL(float) -REGISTER_FUSED_GET_CENTER_DIST_CUDA_KERNEL(double) -REGISTER_FUSED_GET_CENTER_DIST_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_CENTER_DIST_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_CENTER_DIST_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_CENTER_DIST_CUDA_KERNEL(half) template -class FusedCenterGradKernel final : public user_op::OpKernel { +class FusedYolov5GetCenterDistGradKernel final : public user_op::OpKernel { public: - FusedCenterGradKernel() = default; - ~FusedCenterGradKernel() = default; + FusedYolov5GetCenterDistGradKernel() = default; + ~FusedYolov5GetCenterDistGradKernel() = default; private: using user_op::OpKernel::Compute; @@ -155,14 +155,14 @@ class FusedCenterGradKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_CENTER_DIST_GRAD_CUDA_KERNEL(dtype) \ +#define REGISTER_FUSED_YOLOV5_GET_CENTER_DIST_GRAD_CUDA_KERNEL(dtype) \ REGISTER_USER_KERNEL("fused_yolov5_get_center_dist_grad") \ - .SetCreateFn>() \ + .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); -REGISTER_FUSED_GET_CENTER_DIST_GRAD_CUDA_KERNEL(float) -REGISTER_FUSED_GET_CENTER_DIST_GRAD_CUDA_KERNEL(double) -REGISTER_FUSED_GET_CENTER_DIST_GRAD_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_CENTER_DIST_GRAD_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_CENTER_DIST_GRAD_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_CENTER_DIST_GRAD_CUDA_KERNEL(half) } // namespace oneflow diff --git a/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_ciou_diagonal_angle_kernel.cu similarity index 89% rename from oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu rename to oneflow/user/kernels/fused_yolov5_get_ciou_diagonal_angle_kernel.cu index 89fb5f857de..2546e70d27a 100644 --- a/oneflow/user/kernels/fused_get_ciou_diagonal_angle_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_ciou_diagonal_angle_kernel.cu @@ -168,10 +168,10 @@ __global__ void FusedCiouAngleBackward(FUNCTOR_BACKWARD functor_backward, } // namespace template -class FusedGetCiouDiagonalAngleKernel final : public user_op::OpKernel { +class FusedYolov5GetCiouDiagonalAngleKernel final : public user_op::OpKernel { public: - FusedGetCiouDiagonalAngleKernel() = default; - ~FusedGetCiouDiagonalAngleKernel() = default; + FusedYolov5GetCiouDiagonalAngleKernel() = default; + ~FusedYolov5GetCiouDiagonalAngleKernel() = default; private: using user_op::OpKernel::Compute; @@ -196,21 +196,21 @@ class FusedGetCiouDiagonalAngleKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle") \ - .SetCreateFn>() \ - .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ +#define REGISTER_FUSED_YOLOV5_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(dtype) \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle") \ + .SetCreateFn>() \ + .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("v", 0) == GetDataType::value)); -REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(float) -REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(double) -REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_CIOU_DIAGONAL_ANGLE_CUDA_KERNEL(half) template -class FusedGetCiouDiagonalAngleGradKernel final : public user_op::OpKernel { +class FusedYolov5GetCiouDiagonalAngleGradKernel final : public user_op::OpKernel { public: - FusedGetCiouDiagonalAngleGradKernel() = default; - ~FusedGetCiouDiagonalAngleGradKernel() = default; + FusedYolov5GetCiouDiagonalAngleGradKernel() = default; + ~FusedYolov5GetCiouDiagonalAngleGradKernel() = default; private: using user_op::OpKernel::Compute; @@ -242,14 +242,14 @@ class FusedGetCiouDiagonalAngleGradKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle_grad") \ - .SetCreateFn>() \ - .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ +#define REGISTER_FUSED_YOLOV5_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(dtype) \ + REGISTER_USER_KERNEL("fused_yolov5_get_ciou_diagonal_angle_grad") \ + .SetCreateFn>() \ + .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("w1_diff", 0) == GetDataType::value)); -REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(float) -REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(double) -REGISTER_FUSED_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_CIOU_DIAGONAL_ANGLE_GRAD_CUDA_KERNEL(half) } // namespace oneflow diff --git a/oneflow/user/kernels/fused_get_ciou_result_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_ciou_result_kernel.cu similarity index 81% rename from oneflow/user/kernels/fused_get_ciou_result_kernel.cu rename to oneflow/user/kernels/fused_yolov5_get_ciou_result_kernel.cu index b540c994d52..682e0b77569 100644 --- a/oneflow/user/kernels/fused_get_ciou_result_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_ciou_result_kernel.cu @@ -46,10 +46,10 @@ __global__ void FusedGetCiouResultBackward(const int n, const T* dy, const T* al }; // namespace template -class FusedGetCiouResultGpuKernel final : public user_op::OpKernel { +class FusedYolov5GetCiouResultGpuKernel final : public user_op::OpKernel { public: - FusedGetCiouResultGpuKernel() = default; - ~FusedGetCiouResultGpuKernel() = default; + FusedYolov5GetCiouResultGpuKernel() = default; + ~FusedYolov5GetCiouResultGpuKernel() = default; private: using user_op::OpKernel::Compute; @@ -72,21 +72,21 @@ class FusedGetCiouResultGpuKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_CIOU_RESULT_CUDA_KERNEL(dtype) \ +#define REGISTER_FUSED_YOLOV5_GET_CIOU_RESULT_CUDA_KERNEL(dtype) \ REGISTER_USER_KERNEL("fused_yolov5_get_ciou_result") \ - .SetCreateFn>() \ + .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("v", 0) == GetDataType::value)); -REGISTER_FUSED_GET_CIOU_RESULT_CUDA_KERNEL(float) -REGISTER_FUSED_GET_CIOU_RESULT_CUDA_KERNEL(half) -REGISTER_FUSED_GET_CIOU_RESULT_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_CIOU_RESULT_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_CIOU_RESULT_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_CIOU_RESULT_CUDA_KERNEL(double) template -class FusedGetCiouResultGradGpuKernel final : public user_op::OpKernel { +class FusedYolov5GetCiouResultGradGpuKernel final : public user_op::OpKernel { public: - FusedGetCiouResultGradGpuKernel() = default; - ~FusedGetCiouResultGradGpuKernel() = default; + FusedYolov5GetCiouResultGradGpuKernel() = default; + ~FusedYolov5GetCiouResultGradGpuKernel() = default; private: using user_op::OpKernel::Compute; @@ -110,14 +110,14 @@ class FusedGetCiouResultGradGpuKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(dtype) \ +#define REGISTER_FUSED_YOLOV5_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(dtype) \ REGISTER_USER_KERNEL("fused_yolov5_get_ciou_result_grad") \ - .SetCreateFn>() \ + .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("dy", 0) == GetDataType::value)); -REGISTER_FUSED_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(float) -REGISTER_FUSED_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(half) -REGISTER_FUSED_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_CIOU_RESULT_GRAD_CUDA_KERNEL(double) } // namespace oneflow diff --git a/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_convex_diagonal_squared_kernel.cu similarity index 84% rename from oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu rename to oneflow/user/kernels/fused_yolov5_get_convex_diagonal_squared_kernel.cu index 18619dede4c..c7328bcb2fa 100644 --- a/oneflow/user/kernels/fused_get_convex_diagonal_squared_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_convex_diagonal_squared_kernel.cu @@ -60,10 +60,10 @@ __global__ void FusedGetConvexDiagonalSquaredBackward( } // namespace template -class FusedGetConvexDiagonalSquaredKernel final : public user_op::OpKernel { +class FusedYolov5GetConvexDiagonalSquaredKernel final : public user_op::OpKernel { public: - FusedGetConvexDiagonalSquaredKernel() = default; - ~FusedGetConvexDiagonalSquaredKernel() = default; + FusedYolov5GetConvexDiagonalSquaredKernel() = default; + ~FusedYolov5GetConvexDiagonalSquaredKernel() = default; private: using user_op::OpKernel::Compute; @@ -90,21 +90,21 @@ class FusedGetConvexDiagonalSquaredKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared") \ - .SetCreateFn>() \ - .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ +#define REGISTER_FUSED_YOLOV5_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(dtype) \ + REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared") \ + .SetCreateFn>() \ + .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); -REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(float) -REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(double) -REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_CONVEX_DIAGOAL_SQUARED_CUDA_KERNEL(half) template -class FusedGetConvexDiagonalSquaredGradKernel final : public user_op::OpKernel { +class FusedYolov5GetConvexDiagonalSquaredGradKernel final : public user_op::OpKernel { public: - FusedGetConvexDiagonalSquaredGradKernel() = default; - ~FusedGetConvexDiagonalSquaredGradKernel() = default; + FusedYolov5GetConvexDiagonalSquaredGradKernel() = default; + ~FusedYolov5GetConvexDiagonalSquaredGradKernel() = default; private: using user_op::OpKernel::Compute; @@ -142,14 +142,14 @@ class FusedGetConvexDiagonalSquaredGradKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared_grad") \ - .SetCreateFn>() \ - .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ +#define REGISTER_FUSED_YOLOV5_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(dtype) \ + REGISTER_USER_KERNEL("fused_yolov5_get_convex_diagonal_squared_grad") \ + .SetCreateFn>() \ + .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1", 0) == GetDataType::value)); -REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(float) -REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(double) -REGISTER_FUSED_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_CONVEX_DIAGOAL_SQUARED_GRAD_CUDA_KERNEL(half) } // namespace oneflow diff --git a/oneflow/user/kernels/fused_get_intersection_area_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_intersection_area_kernel.cu similarity index 88% rename from oneflow/user/kernels/fused_get_intersection_area_kernel.cu rename to oneflow/user/kernels/fused_yolov5_get_intersection_area_kernel.cu index 5bd4cd79aba..dd4fcd93dc9 100644 --- a/oneflow/user/kernels/fused_get_intersection_area_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_intersection_area_kernel.cu @@ -93,10 +93,10 @@ __global__ void FusedGetIntersectionAreaForward(FUNCTOR functor, const int n, co } // namespace template -class FusedGetIntersectionAreaKernel final : public user_op::OpKernel { +class FusedYolov5GetIntersectionAreaKernel final : public user_op::OpKernel { public: - FusedGetIntersectionAreaKernel() = default; - ~FusedGetIntersectionAreaKernel() = default; + FusedYolov5GetIntersectionAreaKernel() = default; + ~FusedYolov5GetIntersectionAreaKernel() = default; private: using user_op::OpKernel::Compute; @@ -124,21 +124,21 @@ class FusedGetIntersectionAreaKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_INTERSECTION_AREA_CUDA_KERNEL(dtype) \ +#define REGISTER_FUSED_YOLOV5_GET_INTERSECTION_AREA_CUDA_KERNEL(dtype) \ REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area") \ - .SetCreateFn>() \ + .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("inter", 0) == GetDataType::value)); -REGISTER_FUSED_GET_INTERSECTION_AREA_CUDA_KERNEL(float) -REGISTER_FUSED_GET_INTERSECTION_AREA_CUDA_KERNEL(double) -REGISTER_FUSED_GET_INTERSECTION_AREA_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_INTERSECTION_AREA_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_INTERSECTION_AREA_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_INTERSECTION_AREA_CUDA_KERNEL(half) template -class FusedGetIntersectionAreaGradKernel final : public user_op::OpKernel { +class FusedYolov5GetIntersectionAreaGradKernel final : public user_op::OpKernel { public: - FusedGetIntersectionAreaGradKernel() = default; - ~FusedGetIntersectionAreaGradKernel() = default; + FusedYolov5GetIntersectionAreaGradKernel() = default; + ~FusedYolov5GetIntersectionAreaGradKernel() = default; private: using user_op::OpKernel::Compute; @@ -178,14 +178,14 @@ class FusedGetIntersectionAreaGradKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(dtype) \ - REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area_grad") \ - .SetCreateFn>() \ - .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ +#define REGISTER_FUSED_YOLOV5_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(dtype) \ + REGISTER_USER_KERNEL("fused_yolov5_get_intersection_area_grad") \ + .SetCreateFn>() \ + .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("b1_x1_diff", 0) == GetDataType::value)); -REGISTER_FUSED_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(float) -REGISTER_FUSED_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(double) -REGISTER_FUSED_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_INTERSECTION_AREA_GRAD_CUDA_KERNEL(half) } // namespace oneflow diff --git a/oneflow/user/kernels/fused_get_iou_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_iou_kernel.cu similarity index 84% rename from oneflow/user/kernels/fused_get_iou_kernel.cu rename to oneflow/user/kernels/fused_yolov5_get_iou_kernel.cu index 977ecffc467..cf3f75ad78e 100644 --- a/oneflow/user/kernels/fused_get_iou_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_iou_kernel.cu @@ -47,10 +47,10 @@ __global__ void FusedGetIouBackward(const int n, const T* diou, const T* w1, con }; // namespace template -class FusedGetIouGpuKernel final : public user_op::OpKernel { +class FusedYolov5GetIouGpuKernel final : public user_op::OpKernel { public: - FusedGetIouGpuKernel() = default; - ~FusedGetIouGpuKernel() = default; + FusedYolov5GetIouGpuKernel() = default; + ~FusedYolov5GetIouGpuKernel() = default; private: using user_op::OpKernel::Compute; @@ -73,21 +73,21 @@ class FusedGetIouGpuKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_IOU_CUDA_KERNEL(dtype) \ +#define REGISTER_FUSED_YOLOV5_GET_IOU_CUDA_KERNEL(dtype) \ REGISTER_USER_KERNEL("fused_yolov5_get_iou") \ - .SetCreateFn>() \ + .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("iou", 0) == GetDataType::value)); -REGISTER_FUSED_GET_IOU_CUDA_KERNEL(float) -REGISTER_FUSED_GET_IOU_CUDA_KERNEL(half) -REGISTER_FUSED_GET_IOU_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_IOU_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_IOU_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_IOU_CUDA_KERNEL(double) template -class FusedGetIouGradGpuKernel final : public user_op::OpKernel { +class FusedYolov5GetIouGradGpuKernel final : public user_op::OpKernel { public: - FusedGetIouGradGpuKernel() = default; - ~FusedGetIouGradGpuKernel() = default; + FusedYolov5GetIouGradGpuKernel() = default; + ~FusedYolov5GetIouGradGpuKernel() = default; private: using user_op::OpKernel::Compute; @@ -114,14 +114,14 @@ class FusedGetIouGradGpuKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_IOU_GRAD_CUDA_KERNEL(dtype) \ +#define REGISTER_FUSED_YOLOV5_GET_IOU_GRAD_CUDA_KERNEL(dtype) \ REGISTER_USER_KERNEL("fused_yolov5_get_iou_grad") \ - .SetCreateFn>() \ + .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("diou", 0) == GetDataType::value)); -REGISTER_FUSED_GET_IOU_GRAD_CUDA_KERNEL(float) -REGISTER_FUSED_GET_IOU_GRAD_CUDA_KERNEL(half) -REGISTER_FUSED_GET_IOU_GRAD_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_IOU_GRAD_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_IOU_GRAD_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_IOU_GRAD_CUDA_KERNEL(double) } // namespace oneflow diff --git a/oneflow/user/kernels/fused_get_target_offsets_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_target_offsets_kernel.cu similarity index 88% rename from oneflow/user/kernels/fused_get_target_offsets_kernel.cu rename to oneflow/user/kernels/fused_yolov5_get_target_offsets_kernel.cu index 6af9dc97d58..f3f953f31c3 100644 --- a/oneflow/user/kernels/fused_get_target_offsets_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_target_offsets_kernel.cu @@ -71,10 +71,10 @@ __global__ void FusedGetTargetOffsetsForward(MOD_FUNCTOR mod_functor, } // namespace template -class FusedGetTargetOffsetsKernel final : public user_op::OpKernel { +class FusedYolov5GetTargetOffsetsKernel final : public user_op::OpKernel { public: - FusedGetTargetOffsetsKernel() = default; - ~FusedGetTargetOffsetsKernel() = default; + FusedYolov5GetTargetOffsetsKernel() = default; + ~FusedYolov5GetTargetOffsetsKernel() = default; private: using user_op::OpKernel::Compute; @@ -99,15 +99,15 @@ class FusedGetTargetOffsetsKernel final : public user_op::OpKernel { bool AlwaysComputeWhenAllOutputsEmpty() const override { return false; } }; -#define REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(dtype) \ +#define REGISTER_FUSED_YOLOV5_GET_TARGET_OFFSETS_CUDA_KERNEL(dtype) \ REGISTER_USER_KERNEL("fused_yolov5_get_target_offsets") \ - .SetCreateFn>() \ + .SetCreateFn>() \ .SetIsMatchedHob((user_op::HobDeviceType() == DeviceType::kCUDA) \ && (user_op::HobDataType("j", 0) == DataType::kBool) \ && (user_op::HobDataType("gxy", 0) == GetDataType::value)); -REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(float) -REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(double) -REGISTER_FUSED_GET_TARGET_OFFSETS_CUDA_KERNEL(half) +REGISTER_FUSED_YOLOV5_GET_TARGET_OFFSETS_CUDA_KERNEL(float) +REGISTER_FUSED_YOLOV5_GET_TARGET_OFFSETS_CUDA_KERNEL(double) +REGISTER_FUSED_YOLOV5_GET_TARGET_OFFSETS_CUDA_KERNEL(half) } // namespace oneflow diff --git a/oneflow/user/ops/fused_get_boundding_boxes_coord_op.cpp b/oneflow/user/ops/fused_yolov5_get_boundding_boxes_coord_op.cpp similarity index 88% rename from oneflow/user/ops/fused_get_boundding_boxes_coord_op.cpp rename to oneflow/user/ops/fused_yolov5_get_boundding_boxes_coord_op.cpp index dcb4187f2c7..66752aa1644 100644 --- a/oneflow/user/ops/fused_get_boundding_boxes_coord_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_boundding_boxes_coord_op.cpp @@ -18,7 +18,8 @@ limitations under the License. namespace oneflow { -Maybe FusedGetBounddingBoxesCoordOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetBounddingBoxesCoordOp::InferLogicalTensorDesc( + user_op::InferContext* ctx) { const user_op::TensorDesc& x1 = ctx->InputTensorDesc("x1", 0); Shape x1_shape = x1.shape(); @@ -57,11 +58,12 @@ Maybe FusedGetBounddingBoxesCoordOp::InferLogicalTensorDesc(user_op::Infer return Maybe::Ok(); } -Maybe FusedGetBounddingBoxesCoordOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetBounddingBoxesCoordOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetBounddingBoxesCoordOp::InferPhysicalTensorDesc( + user_op::InferContext* ctx) { + return FusedYolov5GetBounddingBoxesCoordOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetBounddingBoxesCoordOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetBounddingBoxesCoordOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& x1 = ctx->InputTensorDesc("x1", 0); user_op::TensorDesc* b1_x1 = ctx->MutOutputTensorDesc("b1_x1", 0); @@ -91,7 +93,7 @@ Maybe FusedGetBounddingBoxesCoordOp::InferDataType(user_op::InferContext* return Maybe::Ok(); } -Maybe FusedGetBounddingBoxesCoordOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetBounddingBoxesCoordOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& x1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("x1", 0); FOR_RANGE(int64_t, i, 0, x1.shape().NumAxes()) { ctx->NewBuilder() @@ -116,7 +118,8 @@ Maybe FusedGetBounddingBoxesCoordOp::GetSbp(user_op::SbpContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetBounddingBoxesCoordGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetBounddingBoxesCoordGradOp::InferLogicalTensorDesc( + user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1_diff = ctx->InputTensorDesc("b1_x1_diff", 0); user_op::TensorDesc* x1_diff = ctx->MutOutputTensorDesc("x1_diff", 0); @@ -154,11 +157,12 @@ Maybe FusedGetBounddingBoxesCoordGradOp::InferLogicalTensorDesc(user_op::I return Maybe::Ok(); } -Maybe FusedGetBounddingBoxesCoordGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetBounddingBoxesCoordGradOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetBounddingBoxesCoordGradOp::InferPhysicalTensorDesc( + user_op::InferContext* ctx) { + return FusedYolov5GetBounddingBoxesCoordGradOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetBounddingBoxesCoordGradOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetBounddingBoxesCoordGradOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1_diff = ctx->InputTensorDesc("b1_x1_diff", 0); user_op::TensorDesc* x1_diff = ctx->MutOutputTensorDesc("x1_diff", 0); @@ -188,7 +192,7 @@ Maybe FusedGetBounddingBoxesCoordGradOp::InferDataType(user_op::InferConte return Maybe::Ok(); } -Maybe FusedGetBounddingBoxesCoordGradOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetBounddingBoxesCoordGradOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& b1_x1_diff = ctx->LogicalTensorDesc4InputArgNameAndIndex("b1_x1_diff", 0); FOR_RANGE(int64_t, i, 0, b1_x1_diff.shape().NumAxes()) { diff --git a/oneflow/user/ops/fused_center_op.cpp b/oneflow/user/ops/fused_yolov5_get_center_dist_op.cpp similarity index 91% rename from oneflow/user/ops/fused_center_op.cpp rename to oneflow/user/ops/fused_yolov5_get_center_dist_op.cpp index 3cc9a5cc102..f914c1ef598 100644 --- a/oneflow/user/ops/fused_center_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_center_dist_op.cpp @@ -18,7 +18,7 @@ limitations under the License. namespace oneflow { -Maybe FusedCenterOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCenterDistOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); const user_op::TensorDesc& b1_x2 = ctx->InputTensorDesc("b1_x2", 0); const user_op::TensorDesc& b1_y1 = ctx->InputTensorDesc("b1_y1", 0); @@ -43,11 +43,11 @@ Maybe FusedCenterOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { return Maybe::Ok(); } -Maybe FusedCenterOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedCenterOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetCenterDistOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetCenterDistOp::InferLogicalTensorDesc(ctx); } -Maybe FusedCenterOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCenterDistOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); const user_op::TensorDesc& b1_x2 = ctx->InputTensorDesc("b1_x2", 0); const user_op::TensorDesc& b1_y1 = ctx->InputTensorDesc("b1_y1", 0); @@ -70,7 +70,7 @@ Maybe FusedCenterOp::InferDataType(user_op::InferContext* ctx) { return Maybe::Ok(); } -Maybe FusedCenterOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetCenterDistOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("b1_x1", 0); FOR_RANGE(int64_t, i, 0, b1_x1.shape().NumAxes()) { ctx->NewBuilder() @@ -88,7 +88,7 @@ Maybe FusedCenterOp::GetSbp(user_op::SbpContext* ctx) { return Maybe::Ok(); } -Maybe FusedCenterGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCenterDistGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); const user_op::TensorDesc& b1_x2 = ctx->InputTensorDesc("b1_x2", 0); const user_op::TensorDesc& b1_y1 = ctx->InputTensorDesc("b1_y1", 0); @@ -143,11 +143,11 @@ Maybe FusedCenterGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx return Maybe::Ok(); } -Maybe FusedCenterGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedCenterGradOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetCenterDistGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetCenterDistGradOp::InferLogicalTensorDesc(ctx); } -Maybe FusedCenterGradOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCenterDistGradOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); const user_op::TensorDesc& b1_x2 = ctx->InputTensorDesc("b1_x2", 0); const user_op::TensorDesc& b1_y1 = ctx->InputTensorDesc("b1_y1", 0); @@ -194,7 +194,7 @@ Maybe FusedCenterGradOp::InferDataType(user_op::InferContext* ctx) { return Maybe::Ok(); } -Maybe FusedCenterGradOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetCenterDistGradOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("b1_x1", 0); FOR_RANGE(int64_t, i, 0, b1_x1.shape().NumAxes()) { ctx->NewBuilder() diff --git a/oneflow/user/ops/fused_get_ciou_diagonal_angle_op.cpp b/oneflow/user/ops/fused_yolov5_get_ciou_diagonal_angle_op.cpp similarity index 85% rename from oneflow/user/ops/fused_get_ciou_diagonal_angle_op.cpp rename to oneflow/user/ops/fused_yolov5_get_ciou_diagonal_angle_op.cpp index c3244eba9ed..6478209ec4c 100644 --- a/oneflow/user/ops/fused_get_ciou_diagonal_angle_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_ciou_diagonal_angle_op.cpp @@ -18,7 +18,7 @@ limitations under the License. namespace oneflow { -Maybe FusedGetCiouDiagonalAngleOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCiouDiagonalAngleOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& w1 = ctx->InputTensorDesc("w1", 0); const user_op::TensorDesc& h1 = ctx->InputTensorDesc("h1", 0); const user_op::TensorDesc& w2 = ctx->InputTensorDesc("w2", 0); @@ -35,11 +35,11 @@ Maybe FusedGetCiouDiagonalAngleOp::InferLogicalTensorDesc(user_op::InferCo return Maybe::Ok(); } -Maybe FusedGetCiouDiagonalAngleOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetCiouDiagonalAngleOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetCiouDiagonalAngleOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetCiouDiagonalAngleOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetCiouDiagonalAngleOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCiouDiagonalAngleOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& w1 = ctx->InputTensorDesc("w1", 0); const user_op::TensorDesc& h1 = ctx->InputTensorDesc("h1", 0); const user_op::TensorDesc& w2 = ctx->InputTensorDesc("w2", 0); @@ -54,7 +54,7 @@ Maybe FusedGetCiouDiagonalAngleOp::InferDataType(user_op::InferContext* ct return Maybe::Ok(); } -Maybe FusedGetCiouDiagonalAngleOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetCiouDiagonalAngleOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("w1", 0); FOR_RANGE(int64_t, i, 0, b1_x1.shape().NumAxes()) { ctx->NewBuilder() @@ -68,7 +68,8 @@ Maybe FusedGetCiouDiagonalAngleOp::GetSbp(user_op::SbpContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetCiouDiagonalAngleGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCiouDiagonalAngleGradOp::InferLogicalTensorDesc( + user_op::InferContext* ctx) { const user_op::TensorDesc& w1 = ctx->InputTensorDesc("w1", 0); const user_op::TensorDesc& h1 = ctx->InputTensorDesc("h1", 0); const user_op::TensorDesc& w2 = ctx->InputTensorDesc("w2", 0); @@ -99,11 +100,12 @@ Maybe FusedGetCiouDiagonalAngleGradOp::InferLogicalTensorDesc(user_op::Inf return Maybe::Ok(); } -Maybe FusedGetCiouDiagonalAngleGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetCiouDiagonalAngleGradOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetCiouDiagonalAngleGradOp::InferPhysicalTensorDesc( + user_op::InferContext* ctx) { + return FusedYolov5GetCiouDiagonalAngleGradOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetCiouDiagonalAngleGradOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCiouDiagonalAngleGradOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& w1 = ctx->InputTensorDesc("w1", 0); const user_op::TensorDesc& h1 = ctx->InputTensorDesc("h1", 0); const user_op::TensorDesc& w2 = ctx->InputTensorDesc("w2", 0); @@ -134,7 +136,7 @@ Maybe FusedGetCiouDiagonalAngleGradOp::InferDataType(user_op::InferContext return Maybe::Ok(); } -Maybe FusedGetCiouDiagonalAngleGradOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetCiouDiagonalAngleGradOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& w1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("w1", 0); FOR_RANGE(int64_t, i, 0, w1.shape().NumAxes()) { ctx->NewBuilder() diff --git a/oneflow/user/ops/fused_get_ciou_result_op.cpp b/oneflow/user/ops/fused_yolov5_get_ciou_result_op.cpp similarity index 81% rename from oneflow/user/ops/fused_get_ciou_result_op.cpp rename to oneflow/user/ops/fused_yolov5_get_ciou_result_op.cpp index 55b224925f3..7a0b7443bbb 100644 --- a/oneflow/user/ops/fused_get_ciou_result_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_ciou_result_op.cpp @@ -18,7 +18,7 @@ limitations under the License. namespace oneflow { -Maybe FusedGetCiouResultOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCiouResultOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& v = ctx->InputTensorDesc("v", 0); user_op::TensorDesc* y = ctx->MutOutputTensorDesc("y", 0); @@ -32,11 +32,11 @@ Maybe FusedGetCiouResultOp::InferLogicalTensorDesc(user_op::InferContext* return Maybe::Ok(); } -Maybe FusedGetCiouResultOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetCiouResultOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetCiouResultOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetCiouResultOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetCiouResultOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCiouResultOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& v = ctx->InputTensorDesc("v", 0); user_op::TensorDesc* y = ctx->MutOutputTensorDesc("y", 0); @@ -48,7 +48,7 @@ Maybe FusedGetCiouResultOp::InferDataType(user_op::InferContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetCiouResultOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetCiouResultOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& v = ctx->LogicalTensorDesc4InputArgNameAndIndex("v", 0); FOR_RANGE(int64_t, i, 0, v.shape().NumAxes()) { ctx->NewBuilder() @@ -63,7 +63,7 @@ Maybe FusedGetCiouResultOp::GetSbp(user_op::SbpContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetCiouResultGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCiouResultGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& dy = ctx->InputTensorDesc("dy", 0); user_op::TensorDesc* dv = ctx->MutOutputTensorDesc("dv", 0); @@ -85,11 +85,11 @@ Maybe FusedGetCiouResultGradOp::InferLogicalTensorDesc(user_op::InferConte return Maybe::Ok(); } -Maybe FusedGetCiouResultGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetCiouResultGradOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetCiouResultGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetCiouResultGradOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetCiouResultGradOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetCiouResultGradOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& dy = ctx->InputTensorDesc("dy", 0); user_op::TensorDesc* dv = ctx->MutOutputTensorDesc("dv", 0); @@ -107,7 +107,7 @@ Maybe FusedGetCiouResultGradOp::InferDataType(user_op::InferContext* ctx) return Maybe::Ok(); } -Maybe FusedGetCiouResultGradOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetCiouResultGradOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& dy = ctx->LogicalTensorDesc4InputArgNameAndIndex("dy", 0); FOR_RANGE(int64_t, i, 0, dy.shape().NumAxes()) { ctx->NewBuilder() diff --git a/oneflow/user/ops/fused_get_convex_diagonal_squared_op.cpp b/oneflow/user/ops/fused_yolov5_get_convex_diagonal_squared_op.cpp similarity index 85% rename from oneflow/user/ops/fused_get_convex_diagonal_squared_op.cpp rename to oneflow/user/ops/fused_yolov5_get_convex_diagonal_squared_op.cpp index f6962b60ba2..1eee7bf6867 100644 --- a/oneflow/user/ops/fused_get_convex_diagonal_squared_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_convex_diagonal_squared_op.cpp @@ -18,7 +18,8 @@ limitations under the License. namespace oneflow { -Maybe FusedGetConvexDiagonalSquaredOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetConvexDiagonalSquaredOp::InferLogicalTensorDesc( + user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); user_op::TensorDesc* c2 = ctx->MutOutputTensorDesc("c2", 0); @@ -28,11 +29,12 @@ Maybe FusedGetConvexDiagonalSquaredOp::InferLogicalTensorDesc(user_op::Inf return Maybe::Ok(); } -Maybe FusedGetConvexDiagonalSquaredOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetConvexDiagonalSquaredOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetConvexDiagonalSquaredOp::InferPhysicalTensorDesc( + user_op::InferContext* ctx) { + return FusedYolov5GetConvexDiagonalSquaredOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetConvexDiagonalSquaredOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetConvexDiagonalSquaredOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); user_op::TensorDesc* c2 = ctx->MutOutputTensorDesc("c2", 0); @@ -40,7 +42,7 @@ Maybe FusedGetConvexDiagonalSquaredOp::InferDataType(user_op::InferContext return Maybe::Ok(); } -Maybe FusedGetConvexDiagonalSquaredOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetConvexDiagonalSquaredOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("b1_x1", 0); FOR_RANGE(int64_t, i, 0, b1_x1.shape().NumAxes()) { ctx->NewBuilder() @@ -58,7 +60,7 @@ Maybe FusedGetConvexDiagonalSquaredOp::GetSbp(user_op::SbpContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetConvexDiagonalSquaredGradOp::InferLogicalTensorDesc( +Maybe FusedYolov5GetConvexDiagonalSquaredGradOp::InferLogicalTensorDesc( user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); @@ -97,12 +99,12 @@ Maybe FusedGetConvexDiagonalSquaredGradOp::InferLogicalTensorDesc( return Maybe::Ok(); } -Maybe FusedGetConvexDiagonalSquaredGradOp::InferPhysicalTensorDesc( +Maybe FusedYolov5GetConvexDiagonalSquaredGradOp::InferPhysicalTensorDesc( user_op::InferContext* ctx) { - return FusedGetConvexDiagonalSquaredGradOp::InferLogicalTensorDesc(ctx); + return FusedYolov5GetConvexDiagonalSquaredGradOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetConvexDiagonalSquaredGradOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetConvexDiagonalSquaredGradOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); user_op::TensorDesc* b1_x1_diff = ctx->MutOutputTensorDesc("b1_x1_diff", 0); @@ -132,7 +134,7 @@ Maybe FusedGetConvexDiagonalSquaredGradOp::InferDataType(user_op::InferCon return Maybe::Ok(); } -Maybe FusedGetConvexDiagonalSquaredGradOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetConvexDiagonalSquaredGradOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("b1_x1", 0); FOR_RANGE(int64_t, i, 0, b1_x1.shape().NumAxes()) { ctx->NewBuilder() diff --git a/oneflow/user/ops/fused_get_intersection_area_op.cpp b/oneflow/user/ops/fused_yolov5_get_intersection_area_op.cpp similarity index 91% rename from oneflow/user/ops/fused_get_intersection_area_op.cpp rename to oneflow/user/ops/fused_yolov5_get_intersection_area_op.cpp index 8bfc42b7863..ca1df673281 100644 --- a/oneflow/user/ops/fused_get_intersection_area_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_intersection_area_op.cpp @@ -18,7 +18,7 @@ limitations under the License. namespace oneflow { -Maybe FusedGetIntersectionAreaOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetIntersectionAreaOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); const user_op::TensorDesc& b1_x2 = ctx->InputTensorDesc("b1_x2", 0); const user_op::TensorDesc& b1_y1 = ctx->InputTensorDesc("b1_y1", 0); @@ -43,11 +43,11 @@ Maybe FusedGetIntersectionAreaOp::InferLogicalTensorDesc(user_op::InferCon return Maybe::Ok(); } -Maybe FusedGetIntersectionAreaOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetIntersectionAreaOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetIntersectionAreaOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetIntersectionAreaOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetIntersectionAreaOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetIntersectionAreaOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); const user_op::TensorDesc& b1_x2 = ctx->InputTensorDesc("b1_x2", 0); const user_op::TensorDesc& b1_y1 = ctx->InputTensorDesc("b1_y1", 0); @@ -70,7 +70,7 @@ Maybe FusedGetIntersectionAreaOp::InferDataType(user_op::InferContext* ctx return Maybe::Ok(); } -Maybe FusedGetIntersectionAreaOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetIntersectionAreaOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("b1_x1", 0); FOR_RANGE(int64_t, i, 0, b1_x1.shape().NumAxes()) { ctx->NewBuilder() @@ -88,7 +88,8 @@ Maybe FusedGetIntersectionAreaOp::GetSbp(user_op::SbpContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetIntersectionAreaGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetIntersectionAreaGradOp::InferLogicalTensorDesc( + user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); const user_op::TensorDesc& b1_x2 = ctx->InputTensorDesc("b1_x2", 0); const user_op::TensorDesc& b1_y1 = ctx->InputTensorDesc("b1_y1", 0); @@ -143,11 +144,12 @@ Maybe FusedGetIntersectionAreaGradOp::InferLogicalTensorDesc(user_op::Infe return Maybe::Ok(); } -Maybe FusedGetIntersectionAreaGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetIntersectionAreaGradOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetIntersectionAreaGradOp::InferPhysicalTensorDesc( + user_op::InferContext* ctx) { + return FusedYolov5GetIntersectionAreaGradOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetIntersectionAreaGradOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetIntersectionAreaGradOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->InputTensorDesc("b1_x1", 0); const user_op::TensorDesc& b1_x2 = ctx->InputTensorDesc("b1_x2", 0); const user_op::TensorDesc& b1_y1 = ctx->InputTensorDesc("b1_y1", 0); @@ -194,7 +196,7 @@ Maybe FusedGetIntersectionAreaGradOp::InferDataType(user_op::InferContext* return Maybe::Ok(); } -Maybe FusedGetIntersectionAreaGradOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetIntersectionAreaGradOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& b1_x1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("b1_x1", 0); FOR_RANGE(int64_t, i, 0, b1_x1.shape().NumAxes()) { ctx->NewBuilder() diff --git a/oneflow/user/ops/fused_get_iou_op.cpp b/oneflow/user/ops/fused_yolov5_get_iou_op.cpp similarity index 81% rename from oneflow/user/ops/fused_get_iou_op.cpp rename to oneflow/user/ops/fused_yolov5_get_iou_op.cpp index 70eee4ff4b3..1aa00f39ff8 100644 --- a/oneflow/user/ops/fused_get_iou_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_iou_op.cpp @@ -18,7 +18,7 @@ limitations under the License. namespace oneflow { -Maybe FusedGetIouOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetIouOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& w1 = ctx->InputTensorDesc("w1", 0); user_op::TensorDesc* iou = ctx->MutOutputTensorDesc("iou", 0); @@ -28,11 +28,11 @@ Maybe FusedGetIouOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetIouOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetIouOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetIouOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetIouOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetIouOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetIouOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& w1 = ctx->InputTensorDesc("w1", 0); user_op::TensorDesc* iou = ctx->MutOutputTensorDesc("iou", 0); @@ -41,7 +41,7 @@ Maybe FusedGetIouOp::InferDataType(user_op::InferContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetIouOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetIouOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& w1 = ctx->LogicalTensorDesc4InputArgNameAndIndex("w1", 0); FOR_RANGE(int64_t, i, 0, w1.shape().NumAxes()) { ctx->NewBuilder() @@ -56,7 +56,7 @@ Maybe FusedGetIouOp::GetSbp(user_op::SbpContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetIouGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetIouGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& diou = ctx->InputTensorDesc("diou", 0); user_op::TensorDesc* dw1 = ctx->MutOutputTensorDesc("dw1", 0); @@ -74,11 +74,11 @@ Maybe FusedGetIouGradOp::InferLogicalTensorDesc(user_op::InferContext* ctx return Maybe::Ok(); } -Maybe FusedGetIouGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetIouGradOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetIouGradOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetIouGradOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetIouGradOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetIouGradOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& diou = ctx->InputTensorDesc("diou", 0); user_op::TensorDesc* dw1 = ctx->MutOutputTensorDesc("dw1", 0); @@ -93,7 +93,7 @@ Maybe FusedGetIouGradOp::InferDataType(user_op::InferContext* ctx) { return Maybe::Ok(); } -Maybe FusedGetIouGradOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetIouGradOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& dy = ctx->LogicalTensorDesc4InputArgNameAndIndex("dy", 0); FOR_RANGE(int64_t, i, 0, dy.shape().NumAxes()) { ctx->NewBuilder() diff --git a/oneflow/user/ops/fused_get_target_offsets_op.cpp b/oneflow/user/ops/fused_yolov5_get_target_offsets_op.cpp similarity index 81% rename from oneflow/user/ops/fused_get_target_offsets_op.cpp rename to oneflow/user/ops/fused_yolov5_get_target_offsets_op.cpp index 6d2372ed2be..e2f053f088f 100644 --- a/oneflow/user/ops/fused_get_target_offsets_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_target_offsets_op.cpp @@ -18,7 +18,7 @@ limitations under the License. namespace oneflow { -Maybe FusedGetTargetOffsetsOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { +Maybe FusedYolov5GetTargetOffsetsOp::InferLogicalTensorDesc(user_op::InferContext* ctx) { const user_op::TensorDesc& gxy = ctx->InputTensorDesc("gxy", 0); const user_op::TensorDesc& gxi = ctx->InputTensorDesc("gxi", 0); @@ -35,11 +35,11 @@ Maybe FusedGetTargetOffsetsOp::InferLogicalTensorDesc(user_op::InferContex return Maybe::Ok(); } -Maybe FusedGetTargetOffsetsOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { - return FusedGetTargetOffsetsOp::InferLogicalTensorDesc(ctx); +Maybe FusedYolov5GetTargetOffsetsOp::InferPhysicalTensorDesc(user_op::InferContext* ctx) { + return FusedYolov5GetTargetOffsetsOp::InferLogicalTensorDesc(ctx); } -Maybe FusedGetTargetOffsetsOp::InferDataType(user_op::InferContext* ctx) { +Maybe FusedYolov5GetTargetOffsetsOp::InferDataType(user_op::InferContext* ctx) { const user_op::TensorDesc& gxy = ctx->InputTensorDesc("gxy", 0); const user_op::TensorDesc& gxi = ctx->InputTensorDesc("gxi", 0); @@ -47,11 +47,10 @@ Maybe FusedGetTargetOffsetsOp::InferDataType(user_op::InferContext* ctx) { user_op::TensorDesc* j = ctx->MutOutputTensorDesc("j", 0); j->set_data_type(DataType::kBool); - // j->set_data_type(gxy.data_type()); return Maybe::Ok(); } -Maybe FusedGetTargetOffsetsOp::GetSbp(user_op::SbpContext* ctx) { +Maybe FusedYolov5GetTargetOffsetsOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& gxy = ctx->LogicalTensorDesc4InputArgNameAndIndex("gxy", 0); FOR_RANGE(int64_t, i, 0, gxy.shape().NumAxes()) { if (i != 1) { diff --git a/python/oneflow/test/modules/test_fused_get_boundding_boxes_coord.py b/python/oneflow/test/modules/test_fused_yolov5_get_boundding_boxes_coord.py similarity index 100% rename from python/oneflow/test/modules/test_fused_get_boundding_boxes_coord.py rename to python/oneflow/test/modules/test_fused_yolov5_get_boundding_boxes_coord.py diff --git a/python/oneflow/test/modules/test_fused_get_center_dist.py b/python/oneflow/test/modules/test_fused_yolov5_get_center_dist.py similarity index 100% rename from python/oneflow/test/modules/test_fused_get_center_dist.py rename to python/oneflow/test/modules/test_fused_yolov5_get_center_dist.py diff --git a/python/oneflow/test/modules/test_fused_get_ciou_diagonal_angle.py b/python/oneflow/test/modules/test_fused_yolov5_get_ciou_diagonal_angle.py similarity index 100% rename from python/oneflow/test/modules/test_fused_get_ciou_diagonal_angle.py rename to python/oneflow/test/modules/test_fused_yolov5_get_ciou_diagonal_angle.py diff --git a/python/oneflow/test/modules/test_fused_get_ciou_result.py b/python/oneflow/test/modules/test_fused_yolov5_get_ciou_result.py similarity index 100% rename from python/oneflow/test/modules/test_fused_get_ciou_result.py rename to python/oneflow/test/modules/test_fused_yolov5_get_ciou_result.py diff --git a/python/oneflow/test/modules/test_fused_get_convex_diagonal_squared.py b/python/oneflow/test/modules/test_fused_yolov5_get_convex_diagonal_squared.py similarity index 100% rename from python/oneflow/test/modules/test_fused_get_convex_diagonal_squared.py rename to python/oneflow/test/modules/test_fused_yolov5_get_convex_diagonal_squared.py diff --git a/python/oneflow/test/modules/test_fused_get_intersection_area.py b/python/oneflow/test/modules/test_fused_yolov5_get_intersection_area.py similarity index 100% rename from python/oneflow/test/modules/test_fused_get_intersection_area.py rename to python/oneflow/test/modules/test_fused_yolov5_get_intersection_area.py diff --git a/python/oneflow/test/modules/test_fused_get_iou.py b/python/oneflow/test/modules/test_fused_yolov5_get_iou.py similarity index 100% rename from python/oneflow/test/modules/test_fused_get_iou.py rename to python/oneflow/test/modules/test_fused_yolov5_get_iou.py diff --git a/python/oneflow/test/modules/test_fused_get_target_offsets.py b/python/oneflow/test/modules/test_fused_yolov5_get_target_offsets.py similarity index 100% rename from python/oneflow/test/modules/test_fused_get_target_offsets.py rename to python/oneflow/test/modules/test_fused_yolov5_get_target_offsets.py From f59f6dacbeca9416dc1d62c3cffaff3a74482087 Mon Sep 17 00:00:00 2001 From: hhhfccz Date: Fri, 16 Dec 2022 06:07:31 +0000 Subject: [PATCH 5/5] fix: init for first row --- .../fused_yolov5_get_target_offsets_kernel.cu | 8 ++++++++ .../ops/fused_yolov5_get_target_offsets_op.cpp | 14 ++++++-------- .../test_fused_yolov5_get_target_offsets.py | 6 +++--- 3 files changed, 17 insertions(+), 11 deletions(-) diff --git a/oneflow/user/kernels/fused_yolov5_get_target_offsets_kernel.cu b/oneflow/user/kernels/fused_yolov5_get_target_offsets_kernel.cu index f3f953f31c3..5bc63b67e03 100644 --- a/oneflow/user/kernels/fused_yolov5_get_target_offsets_kernel.cu +++ b/oneflow/user/kernels/fused_yolov5_get_target_offsets_kernel.cu @@ -29,6 +29,13 @@ struct ModFunctor { __device__ T Compute(T input_tensor) const { return fmod(input_tensor, static_cast(1.0)); } }; +template<> +struct ModFunctor { + __device__ float Compute(float input_tensor) const { + return fmodf(input_tensor, static_cast(1.0)); + } +}; + template<> struct ModFunctor { ModFunctor float_functor; @@ -56,6 +63,7 @@ __global__ void FusedGetTargetOffsetsForward(MOD_FUNCTOR mod_functor, const T gxi_mod_1 = mod_functor.Compute(gxi_i); const bool stats_1 = get_stats_functor.Compute(gxy_i, gxy_mod_1, g); const bool stats_2 = get_stats_functor.Compute(gxi_i, gxi_mod_1, g); + if (i < rows) { output_tensor[i] = true; } if (i % 2 == 0) { const int64_t extra_cols = i / 2; output_tensor[i - extra_cols + rows] = stats_1; diff --git a/oneflow/user/ops/fused_yolov5_get_target_offsets_op.cpp b/oneflow/user/ops/fused_yolov5_get_target_offsets_op.cpp index e2f053f088f..94547ddd862 100644 --- a/oneflow/user/ops/fused_yolov5_get_target_offsets_op.cpp +++ b/oneflow/user/ops/fused_yolov5_get_target_offsets_op.cpp @@ -30,7 +30,7 @@ Maybe FusedYolov5GetTargetOffsetsOp::InferLogicalTensorDesc(user_op::Infer user_op::TensorDesc* j = ctx->MutOutputTensorDesc("j", 0); j->set_is_dynamic(gxy.is_dynamic()); - j->set_shape(Shape({gxy_shape.At(1) * 2 + 1, gxy_shape.At(0)})); + j->set_shape(Shape({5, gxy_shape.At(0)})); return Maybe::Ok(); } @@ -53,13 +53,11 @@ Maybe FusedYolov5GetTargetOffsetsOp::InferDataType(user_op::InferContext* Maybe FusedYolov5GetTargetOffsetsOp::GetSbp(user_op::SbpContext* ctx) { const user_op::TensorDesc& gxy = ctx->LogicalTensorDesc4InputArgNameAndIndex("gxy", 0); FOR_RANGE(int64_t, i, 0, gxy.shape().NumAxes()) { - if (i != 1) { - ctx->NewBuilder() - .Split(user_op::OpArg("gxy", 0), i) - .Split(user_op::OpArg("gxi", 0), i) - .Split(user_op::OpArg("j", 0), i) - .Build(); - } + ctx->NewBuilder() + .Broadcast(user_op::OpArg("gxy", 0)) + .Broadcast(user_op::OpArg("gxi", 0)) + .Broadcast(user_op::OpArg("j", 0)) + .Build(); } return Maybe::Ok(); } diff --git a/python/oneflow/test/modules/test_fused_yolov5_get_target_offsets.py b/python/oneflow/test/modules/test_fused_yolov5_get_target_offsets.py index 27f3330a0c0..b1a015ef08b 100644 --- a/python/oneflow/test/modules/test_fused_yolov5_get_target_offsets.py +++ b/python/oneflow/test/modules/test_fused_yolov5_get_target_offsets.py @@ -27,7 +27,7 @@ def torch_get_target_offsets(gxy, gxi, g): j, k = ((gxy % 1 < g) & (gxy > 1)).T l, m = ((gxi % 1 < g) & (gxi > 1)).T - return torch.stack((j, k, l, m, torch.ones_like(j))) + return torch.stack((torch.ones_like(j), j, k, l, m)) def _test_fused_get_target_offsets_impl(test_case, device, shape, g): @@ -53,12 +53,12 @@ def _test_fused_get_target_offsets_impl(test_case, device, shape, g): torch_gxy, torch_gxi = torch_x[0] + torch_x[1], torch_x[0] + torch_x[2] j = flow._C.fused_yolov5_get_target_offsets(gxy, gxi, g) torch_j = torch_get_target_offsets(torch_gxy, torch_gxi, g) - test_case.assertTrue((j.cpu().numpy() == torch_j.cpu().numpy()).any()) + test_case.assertTrue((j.cpu().numpy() == torch_j.cpu().numpy()).all()) @flow.unittest.skip_unless_1n1d() class TestGetTargetOffsetsModule(flow.unittest.TestCase): - def test_fused_get_target_offsets_area(test_case): + def test_fused_get_target_offsets(test_case): arg_dict = OrderedDict() arg_dict["test_fun"] = [_test_fused_get_target_offsets_impl] arg_dict["device"] = ["cuda"]