Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into fix_bug_for_lstmp
xuezhong
6 years ago
commit
4028943125
@ -0,0 +1,80 @@
|
||||
// Copyright (c) 2019 PaddlePaddle 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 "paddle/fluid/framework/ir/identity_scale_op_clean_pass.h"
|
||||
#include <string>
|
||||
#include "paddle/fluid/framework/ir/graph_pattern_detector.h"
|
||||
|
||||
namespace paddle {
|
||||
namespace framework {
|
||||
namespace ir {
|
||||
|
||||
std::unique_ptr<ir::Graph> IdentityScaleOpCleanPass::ApplyImpl(
|
||||
std::unique_ptr<ir::Graph> graph) const {
|
||||
FusePassBase::Init("identity_scale_op_clean", graph.get());
|
||||
|
||||
// pre_op -> scale_in -> scale_op -> scale_out
|
||||
// ->
|
||||
// pre_op -> scale_out
|
||||
GraphPatternDetector detector;
|
||||
auto pre_op = detector.mutable_pattern()->NewNode("pre_op")->assert_is_op();
|
||||
auto scale_in = detector.mutable_pattern()
|
||||
->NewNode("scale_in")
|
||||
->assert_is_op_input("scale")
|
||||
->AsIntermediate();
|
||||
auto scale_op = detector.mutable_pattern()
|
||||
->NewNode("scale_fuse")
|
||||
->assert_is_op("scale")
|
||||
->assert_op_attr<float>("scale", 1.)
|
||||
->assert_op_attr<float>("bias", 0.);
|
||||
auto scale_out = detector.mutable_pattern()
|
||||
->NewNode("scale_out")
|
||||
->assert_is_op_output("scale");
|
||||
|
||||
pre_op->LinksTo({scale_in});
|
||||
scale_op->LinksFrom({scale_in}).LinksTo({scale_out});
|
||||
|
||||
GraphPatternDetector::handle_t handler = [&](
|
||||
const GraphPatternDetector::subgraph_t& subgraph, Graph* graph) {
|
||||
Node* scale_op_var = subgraph.at(scale_op);
|
||||
Node* scale_in_var = subgraph.at(scale_in);
|
||||
Node* scale_out_var = subgraph.at(scale_out);
|
||||
Node* pre_op_var = subgraph.at(pre_op);
|
||||
// Link pre_op directly to scale_out
|
||||
const std::string scale_in_name = scale_in_var->Name();
|
||||
const std::string scale_out_name = scale_out_var->Name();
|
||||
// Remove links in graph
|
||||
GraphSafeRemoveNodes(graph, {scale_in_var, scale_op_var});
|
||||
// Modify proto message
|
||||
auto* pre_op_desc = pre_op_var->Op();
|
||||
for (auto& parameter : *pre_op_desc->Proto()->mutable_outputs()) {
|
||||
auto* arguments = parameter.mutable_arguments();
|
||||
auto it = std::find(arguments->begin(), arguments->end(), scale_in_name);
|
||||
PADDLE_ENFORCE(it != arguments->end());
|
||||
*it = scale_out_name;
|
||||
}
|
||||
|
||||
IR_NODE_LINK_TO(pre_op_var, scale_out_var);
|
||||
};
|
||||
|
||||
detector(graph.get(), handler);
|
||||
return graph;
|
||||
}
|
||||
|
||||
} // namespace ir
|
||||
} // namespace framework
|
||||
} // namespace paddle
|
||||
|
||||
REGISTER_PASS(identity_scale_op_clean_pass,
|
||||
paddle::framework::ir::IdentityScaleOpCleanPass);
|
||||
@ -0,0 +1,33 @@
|
||||
// Copyright (c) 2019 PaddlePaddle 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.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "paddle/fluid/framework/ir/fuse_pass_base.h"
|
||||
|
||||
namespace paddle {
|
||||
namespace framework {
|
||||
namespace ir {
|
||||
|
||||
class IdentityScaleOpCleanPass : public FusePassBase {
|
||||
protected:
|
||||
std::unique_ptr<ir::Graph> ApplyImpl(std::unique_ptr<ir::Graph> graph) const;
|
||||
|
||||
private:
|
||||
virtual ~IdentityScaleOpCleanPass() = default;
|
||||
};
|
||||
|
||||
} // namespace ir
|
||||
} // namespace framework
|
||||
} // namespace paddle
|
||||
@ -1,4 +1,4 @@
|
||||
cc_library(benchmark SRCS benchmark.cc DEPS enforce)
|
||||
cc_test(test_benchmark SRCS benchmark_tester.cc DEPS benchmark)
|
||||
#cc_binary(visualizer SRCS visualizer.cc DEPS analysis
|
||||
# paddle_pass_builder ir_pass_manager pass graph_viz_pass analysis_passes)
|
||||
cc_binary(visualizer SRCS visualizer.cc DEPS analysis
|
||||
paddle_pass_builder ir_pass_manager pass graph_viz_pass analysis_passes)
|
||||
|
||||
@ -0,0 +1,86 @@
|
||||
/* Copyright (c) 2018 PaddlePaddle 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 "paddle/fluid/operators/detection/box_clip_op.h"
|
||||
#include "paddle/fluid/framework/op_registry.h"
|
||||
|
||||
namespace paddle {
|
||||
namespace operators {
|
||||
|
||||
class BoxClipOp : public framework::OperatorWithKernel {
|
||||
public:
|
||||
using framework::OperatorWithKernel::OperatorWithKernel;
|
||||
|
||||
protected:
|
||||
void InferShape(framework::InferShapeContext* ctx) const override {
|
||||
PADDLE_ENFORCE(ctx->HasInput("Input"),
|
||||
"Input(Input) of BoxClipOp should not be null.");
|
||||
PADDLE_ENFORCE(ctx->HasInput("ImInfo"),
|
||||
"Input(ImInfo) of BoxClipOp should not be null.");
|
||||
|
||||
auto input_box_dims = ctx->GetInputDim("Input");
|
||||
auto im_info_dims = ctx->GetInputDim("ImInfo");
|
||||
|
||||
if (ctx->IsRuntime()) {
|
||||
auto input_box_size = input_box_dims.size();
|
||||
PADDLE_ENFORCE_EQ(input_box_dims[input_box_size - 1], 4,
|
||||
"The last dimension of Input must be 4");
|
||||
PADDLE_ENFORCE_EQ(im_info_dims.size(), 2,
|
||||
"The rank of Input(Input) in BoxClipOp must be 2");
|
||||
PADDLE_ENFORCE_EQ(im_info_dims[1], 3,
|
||||
"The last dimension of ImInfo must be 3");
|
||||
}
|
||||
ctx->ShareDim("Input", /*->*/ "Output");
|
||||
ctx->ShareLoD("Input", /*->*/ "Output");
|
||||
}
|
||||
};
|
||||
|
||||
class BoxClipOpMaker : public framework::OpProtoAndCheckerMaker {
|
||||
public:
|
||||
void Make() override {
|
||||
AddInput("Input",
|
||||
"(LoDTensor) "
|
||||
"Input is a LoDTensor with shape [..., 4] holds 4 points"
|
||||
"in last dimension in format [xmin, ymin, xmax, ymax]");
|
||||
AddInput("ImInfo",
|
||||
"(Tensor) Information for image reshape is in shape (N, 3), "
|
||||
"in format (height, width, im_scale)");
|
||||
AddOutput("Output",
|
||||
"(LoDTensor) "
|
||||
"Output is a LoDTensor with the same shape as Input"
|
||||
"and it is the result after clip");
|
||||
AddComment(R"DOC(
|
||||
This operator clips input boxes to original input images.
|
||||
|
||||
For each input box, The formula is given as follows:
|
||||
|
||||
$$xmin = \max(\min(xmin, im_w - 1), 0)$$
|
||||
$$ymin = \max(\min(ymin, im_h - 1), 0)$$
|
||||
$$xmax = \max(\min(xmax, im_w - 1), 0)$$
|
||||
$$ymax = \max(\min(ymax, im_h - 1), 0)$$
|
||||
|
||||
where im_w and im_h are computed from ImInfo, the formula is given as follows:
|
||||
|
||||
$$im_w = \round(width / im_scale)$$
|
||||
$$im_h = \round(height / im_scale)$$
|
||||
)DOC");
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace operators
|
||||
} // namespace paddle
|
||||
|
||||
namespace ops = paddle::operators;
|
||||
REGISTER_OPERATOR(box_clip, ops::BoxClipOp, ops::BoxClipOpMaker,
|
||||
paddle::framework::EmptyGradOpMaker);
|
||||
REGISTER_OP_CPU_KERNEL(
|
||||
box_clip, ops::BoxClipKernel<paddle::platform::CPUDeviceContext, float>,
|
||||
ops::BoxClipKernel<paddle::platform::CPUDeviceContext, double>);
|
||||
@ -0,0 +1,74 @@
|
||||
/* Copyright (c) 2018 PaddlePaddle 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 <algorithm>
|
||||
#include "paddle/fluid/framework/op_registry.h"
|
||||
#include "paddle/fluid/operators/detection/box_clip_op.h"
|
||||
#include "paddle/fluid/operators/math/math_function.h"
|
||||
#include "paddle/fluid/platform/cuda_primitives.h"
|
||||
#include "paddle/fluid/platform/hostdevice.h"
|
||||
|
||||
namespace paddle {
|
||||
namespace operators {
|
||||
|
||||
using Tensor = framework::Tensor;
|
||||
using LoDTenso = framework::LoDTensor;
|
||||
|
||||
static constexpr int ImInfoSize = 3;
|
||||
|
||||
template <typename T, int BlockSize>
|
||||
static __global__ void GPUBoxClip(const T *input, const size_t *lod,
|
||||
const size_t width, const T *im_info,
|
||||
T *output) {
|
||||
T im_w = round(im_info[blockIdx.x * ImInfoSize + 1] /
|
||||
im_info[blockIdx.x * ImInfoSize + 2]);
|
||||
T im_h = round(im_info[blockIdx.x * ImInfoSize] /
|
||||
im_info[blockIdx.x * ImInfoSize + 2]);
|
||||
for (int i = threadIdx.x; i < (lod[blockIdx.x + 1] - lod[blockIdx.x]) * width;
|
||||
i += BlockSize) {
|
||||
int idx = lod[blockIdx.x] * width + i;
|
||||
T im_size = (idx % 2 == 0) ? im_w : im_h;
|
||||
output[idx] = max(min(input[idx], im_size - 1), T(0.));
|
||||
}
|
||||
}
|
||||
|
||||
template <typename DeviceContext, typename T>
|
||||
class GPUBoxClipKernel : public framework::OpKernel<T> {
|
||||
public:
|
||||
void Compute(const framework::ExecutionContext &context) const override {
|
||||
PADDLE_ENFORCE(platform::is_gpu_place(context.GetPlace()),
|
||||
"This kernel only runs on GPU device.");
|
||||
auto *input = context.Input<LoDTensor>("Input");
|
||||
auto *im_info = context.Input<Tensor>("ImInfo");
|
||||
auto *output = context.Output<LoDTensor>("Output");
|
||||
const int64_t num = input->dims()[0];
|
||||
const int64_t bbox_width = input->numel() / num;
|
||||
auto lod = input->lod();
|
||||
framework::LoD abs_offset_lod = framework::ToAbsOffset(lod);
|
||||
auto &dev_ctx = context.template device_context<DeviceContext>();
|
||||
auto stream = dev_ctx.stream();
|
||||
const size_t batch_size = lod.back().size() - 1;
|
||||
T *output_data = output->mutable_data<T>(dev_ctx.GetPlace());
|
||||
GPUBoxClip<T, 512><<<batch_size, 512, 0, stream>>>(
|
||||
input->data<T>(), abs_offset_lod[0].CUDAMutableData(dev_ctx.GetPlace()),
|
||||
bbox_width, im_info->data<T>(), output_data);
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace operators
|
||||
} // namespace paddle
|
||||
|
||||
namespace ops = paddle::operators;
|
||||
REGISTER_OP_CUDA_KERNEL(
|
||||
box_clip, ops::GPUBoxClipKernel<paddle::platform::CUDADeviceContext, float>,
|
||||
ops::GPUBoxClipKernel<paddle::platform::CUDADeviceContext, double>);
|
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue