conv_transpose_op.cc 14.6 KB
Newer Older
1
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
C
chengduoZH 已提交
2

L
Luo Tao 已提交
3 4 5
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
C
chengduoZH 已提交
6

L
Luo Tao 已提交
7
    http://www.apache.org/licenses/LICENSE-2.0
C
chengduoZH 已提交
8

L
Luo Tao 已提交
9 10 11 12 13
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. */
C
chengduoZH 已提交
14

Y
Yi Wang 已提交
15
#include "paddle/fluid/operators/conv_transpose_op.h"
S
Siddharth Goyal 已提交
16 17
#include <string>
#include <vector>
C
chengduoZH 已提交
18 19 20 21

namespace paddle {
namespace operators {

C
chengduoZH 已提交
22
void ConvTransposeOp::InferShape(framework::InferShapeContext* ctx) const {
C
chengduoZH 已提交
23
  PADDLE_ENFORCE(ctx->HasInput("Input"),
C
chengduoZH 已提交
24
                 "Input(Input) of ConvTransposeOp should not be null.");
C
chengduoZH 已提交
25
  PADDLE_ENFORCE(ctx->HasInput("Filter"),
C
chengduoZH 已提交
26
                 "Input(Filter) of ConvTransposeOp should not be null.");
C
chengduoZH 已提交
27
  PADDLE_ENFORCE(ctx->HasOutput("Output"),
C
chengduoZH 已提交
28
                 "Output(Output) of ConvTransposeOp should not be null.");
C
chengduoZH 已提交
29 30 31 32 33

  auto in_dims = ctx->GetInputDim("Input");
  auto filter_dims = ctx->GetInputDim("Filter");
  std::vector<int> strides = ctx->Attrs().Get<std::vector<int>>("strides");
  std::vector<int> paddings = ctx->Attrs().Get<std::vector<int>>("paddings");
C
chengduoZH 已提交
34
  std::vector<int> dilations = ctx->Attrs().Get<std::vector<int>>("dilations");
Y
Yibing Liu 已提交
35
  int groups = ctx->Attrs().Get<int>("groups");
C
chengduoZH 已提交
36

C
chengduoZH 已提交
37 38 39 40 41 42 43 44 45
  PADDLE_ENFORCE(in_dims.size() == 4 || in_dims.size() == 5,
                 "ConvTransposeOp intput should be 4-D or 5-D tensor.");
  PADDLE_ENFORCE_EQ(in_dims.size(), filter_dims.size(),
                    "ConvTransposeOp input dimension and filter dimension "
                    "should be the same.");
  PADDLE_ENFORCE(in_dims.size() - strides.size() == 2U,
                 "ConvTransposeOp input dimension and strides dimension should "
                 "be consistent.");
  PADDLE_ENFORCE_EQ(paddings.size(), strides.size(),
C
chengduoZH 已提交
46
                    "ConvTransposeOp paddings dimension and strides "
C
chengduoZH 已提交
47
                    "dimension should be the same.");
C
chengduoZH 已提交
48 49 50
  PADDLE_ENFORCE_EQ(paddings.size(), dilations.size(),
                    "ConvTransposeOp paddings dimension and dilations "
                    "dimension should be the same.");
C
chengduoZH 已提交
51
  PADDLE_ENFORCE_EQ(in_dims[1], filter_dims[0],
Y
Yibing Liu 已提交
52
                    "In ConvTransposeOp, The number of input channels should "
53
                    "be equal to the number of filter's channels.");
C
chengduoZH 已提交
54

Y
Yibing Liu 已提交
55
  std::vector<int64_t> output_shape({in_dims[0], filter_dims[1] * groups});
C
chengduoZH 已提交
56
  for (size_t i = 0; i < strides.size(); ++i) {
C
chengduoZH 已提交
57
    auto filter_extent = dilations[i] * (filter_dims[i + 2] - 1) + 1;
C
chengduoZH 已提交
58
    output_shape.push_back((in_dims[i + 2] - 1) * strides[i] - 2 * paddings[i] +
C
chengduoZH 已提交
59
                           filter_extent);
C
chengduoZH 已提交
60
  }
C
chengduoZH 已提交
61
  ctx->SetOutputDim("Output", framework::make_ddim(output_shape));
C
chengduoZH 已提交
62 63
}

64 65 66
framework::OpKernelType ConvTransposeOp::GetExpectedKernelType(
    const framework::ExecutionContext& ctx) const {
  bool use_cudnn = ctx.Attr<bool>("use_cudnn");
C
chengduoZH 已提交
67
  use_cudnn &= platform::is_gpu_place(ctx.GetPlace());
C
chengduoZH 已提交
68 69 70 71 72 73
#ifdef PADDLE_WITH_CUDA
  if (platform::is_gpu_place(ctx.GetPlace())) {
    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
    use_cudnn &= dev_ctx.cudnn_handle() != nullptr;
  }
#endif
74 75 76 77 78 79 80 81 82 83 84 85 86 87
  framework::LibraryType library_;
  if (use_cudnn) {
    library_ = framework::LibraryType::kCUDNN;
  } else {
    library_ = framework::LibraryType::kPlain;
  }

  std::string data_format = ctx.Attr<std::string>("data_format");
  framework::DataLayout layout_ = framework::StringToDataLayout(data_format);
  return framework::OpKernelType(
      framework::ToDataType(ctx.Input<Tensor>("Input")->type()), ctx.GetPlace(),
      layout_, library_);
}

Y
Yu Yang 已提交
88
void Conv2DTransposeOpMaker::Make() {
C
chengduoZH 已提交
89 90 91 92
  AddInput(
      "Input",
      "(Tensor) The input tensor of convolution transpose operator. "
      "The format of input tensor is NCHW. Where N is batch size, C is the "
C
chengduoZH 已提交
93 94
      "number of input channels, H is the height of the feature, and "
      "W is the width of the feature.");
C
chengduoZH 已提交
95 96 97 98 99 100 101 102
  AddInput(
      "Filter",
      "(Tensor) The filter tensor of convolution transpose operator. "
      "The format of the filter tensor is MCHW, where M is the number of "
      "input feature channels, C is the number of "
      "output feature channels,"
      "H is the height of the filter, and W is the width of the filter. "
      "We enforce groups number == 1 in the convolution transpose scenario.");
C
chengduoZH 已提交
103
  AddOutput("Output",
C
chengduoZH 已提交
104
            "(Tensor) The output tensor of convolution transpose operator. "
C
chengduoZH 已提交
105
            "The format of output tensor is also NCHW.");
Y
Yibing Liu 已提交
106 107 108 109
  AddAttr<int>("groups",
               "(int default:1), the groups number of the convolution "
               "transpose operator. ")
      .SetDefault(1);
C
chengduoZH 已提交
110 111 112 113 114
  AddAttr<std::vector<int>>("dilations",
                            "(vector<int> default:{1, 1}), the "
                            "dilations(h_dilation, w_dilation) of convolution "
                            "transpose operator.")
      .SetDefault({1, 1});
C
chengduoZH 已提交
115 116
  AddAttr<std::vector<int>>(
      "strides",
C
chengduoZH 已提交
117
      "(vector<int> default:{1, 1}), the strides(h_stride, w_stride) of "
118
      "convolution transpose operator.")
C
chengduoZH 已提交
119
      .SetDefault({1, 1});
C
chengduoZH 已提交
120 121
  AddAttr<std::vector<int>>(
      "paddings",
C
chengduoZH 已提交
122
      "(vector<int> default:{0, 0}), the paddings(h_pad, w_pad) of convolution "
C
chengduoZH 已提交
123
      "transpose operator.")
C
chengduoZH 已提交
124
      .SetDefault({0, 0});
125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143
  AddAttr<bool>(
      "use_cudnn",
      "(bool, default false) Only used in cudnn kernel, need install cudnn")
      .SetDefault(false);
  AddAttr<std::string>(
      "data_format",
      "(string, default NCHW) Only used in "
      "An optional string from: \"NHWC\", \"NCHW\". "
      "Defaults to \"NHWC\". Specify the data format of the output data, "
      "the input will be transformed automatically. ")
      .SetDefault("AnyLayout");
  // TODO(dzhwinter): need to registered layout transform function
  AddAttr<int>("workspace_size_MB",
               "Used in cudnn kernel only. workspace size for cudnn, in MB, "
               "workspace is a section of GPU memory which will be "
               "allocated/freed each time the operator runs, larger "
               "workspace size can increase performance but also requires "
               "better hardward. This size should be carefully setted.")
      .SetDefault(4096);
C
chengduoZH 已提交
144
  AddComment(R"DOC(
C
chengduoZH 已提交
145 146
Convolution2D Transpose Operator.

C
chengduoZH 已提交
147
The convolution transpose operation calculates the output based on the input, filter
C
chengduoZH 已提交
148
and dilations, strides, paddings, groups parameters. The size of each dimension of the
C
chengduoZH 已提交
149
parameters is checked in the infer-shape.
C
chengduoZH 已提交
150 151 152 153 154 155 156
Input(Input) and output(Output) are in NCHW format. Where N is batchsize, C is the
number of channels, H is the height of the feature, and W is the width of the feature.
Filter(Input) is in MCHW format. Where M is the number of input feature channels,
C is the number of output feature channels, H is the height of the filter,
and W is the width of the filter.
Parameters(strides, paddings) are two elements. These two elements represent height
and width, respectively.
C
chengduoZH 已提交
157
The input(X) size and output(Out) size may be different.
C
chengduoZH 已提交
158

C
chengduoZH 已提交
159 160
Example:
  Input:
C
chengduoZH 已提交
161 162
       Input shape: $(N, C_{in}, H_{in}, W_{in})$
       Filter shape: $(C_{in}, C_{out}, H_f, W_f)$
C
chengduoZH 已提交
163
  Output:
C
chengduoZH 已提交
164 165 166
       Output shape: $(N, C_{out}, H_{out}, W_{out})$
  Where
  $$
167 168
       H_{out} = (H_{in} - 1) * strides[0] - 2 * paddings[0] + dilations[0] * (H_f - 1) + 1 \\
       W_{out} = (W_{in} - 1) * strides[1] - 2 * paddings[1] + dilations[1] * (W_f - 1) + 1
C
chengduoZH 已提交
169
  $$
C
chengduoZH 已提交
170 171 172
)DOC");
}

Y
Yu Yang 已提交
173
void Conv3DTransposeOpMaker::Make() {
C
chengduoZH 已提交
174 175 176 177 178 179
  AddInput("Input",
           "(Tensor) The input tensor of convolution transpose operator."
           "The format of input tensor is NCDHW. Where N is batch size, C is "
           "the number of channels, D is the depth of the feature, H is the "
           "height of the feature, and "
           "W is the width of the feature.");
C
chengduoZH 已提交
180 181
  AddInput("Filter",
           "(Tensor) The filter tensor of convolution transpose operator."
C
chengduoZH 已提交
182 183 184
           "The format of the filter tensor is MCDHW, where M is the number of "
           "input feature channels, C is the number of "
           "output feature channels, D "
C
chengduoZH 已提交
185 186
           "is the depth of the filter, H is the height of the filter, and "
           "W is the width of the filter."
C
chengduoZH 已提交
187
           "We enforce groups number == 1 and padding == 0 in "
C
chengduoZH 已提交
188
           "the convolution3d transpose scenario.");
C
chengduoZH 已提交
189 190 191 192
  AddOutput("Output",
            "(Tensor) The output tensor of convolution transpose operator."
            "The format of output tensor is also NCDHW."
            "Where N is batch size, C is "
C
chengduoZH 已提交
193 194
            "the number of channels, D is the depth of the feature, H is the "
            "height of the feature, and W is the width of the feature.");
C
chengduoZH 已提交
195 196 197 198 199 200 201

  AddAttr<std::vector<int>>(
      "dilations",
      "(vector<int> default:{1, 1, 1}), the "
      "dilations(d_dilation,h_dilation, w_dilation) of convolution "
      "transpose operator.")
      .SetDefault({1, 1, 1});
C
chengduoZH 已提交
202
  AddAttr<std::vector<int>>("strides",
C
chengduoZH 已提交
203
                            "(vector<int> default:{1, 1, 1}), the "
204
                            "strides{d_stride, h_stride, w_stride} of "
C
chengduoZH 已提交
205
                            "convolution transpose operator.")
C
chengduoZH 已提交
206
      .SetDefault({1, 1, 1});
C
chengduoZH 已提交
207
  AddAttr<std::vector<int>>("paddings",
C
chengduoZH 已提交
208
                            "(vector<int> default:{0, 0, 0}), paddings(d_pad, "
C
chengduoZH 已提交
209
                            "h_pad, w_pad) of convolution transpose operator.")
C
chengduoZH 已提交
210
      .SetDefault({0, 0, 0});
211 212 213 214
  AddAttr<int>("groups",
               "(int default:1), the groups number of the convolution3d "
               "transpose operator. ")
      .SetDefault(1);
215 216 217 218 219 220 221 222 223 224 225 226 227 228 229 230 231 232 233
  AddAttr<bool>(
      "use_cudnn",
      "(bool, default false) Only used in cudnn kernel, need install cudnn")
      .SetDefault(false);
  AddAttr<std::string>(
      "data_format",
      "(string, default NCHW) Only used in "
      "An optional string from: \"NHWC\", \"NCHW\". "
      "Defaults to \"NHWC\". Specify the data format of the output data, "
      "the input will be transformed automatically. ")
      .SetDefault("AnyLayout");
  // TODO(dzhwinter): need to registered layout transform function
  AddAttr<int>("workspace_size_MB",
               "Used in cudnn kernel only. workspace size for cudnn, in MB, "
               "workspace is a section of GPU memory which will be "
               "allocated/freed each time the operator runs, larger "
               "workspace size can increase performance but also requires "
               "better hardward. This size should be carefully setted.")
      .SetDefault(4096);
C
chengduoZH 已提交
234
  AddComment(R"DOC(
C
chengduoZH 已提交
235 236
Convolution3D Transpose Operator.

C
chengduoZH 已提交
237
The convolution transpose operation calculates the output based on the input, filter
C
chengduoZH 已提交
238
and dilations, strides, paddings, groups parameters. The size of each dimension of the
C
chengduoZH 已提交
239
parameters is checked in the infer-shape.
C
chengduoZH 已提交
240 241 242 243 244 245 246 247
Input(Input) and output(Output) are in NCDHW format. Where N is batch size, C is the
number of channels, D is the depth of the feature, H is the height of the feature,
and W is the width of the feature.
Filter(Input) is in MCDHW format. Where M is the number of input feature channels,
C is the number of output feature channels, D is the depth of the filter,H is the
height of the filter, and W is the width of the filter.
Parameters(strides, paddings) are three elements. These three elements represent
depth, height and width, respectively.
C
chengduoZH 已提交
248
The input(X) size and output(Out) size may be different.
C
chengduoZH 已提交
249 250

Example:   
C
chengduoZH 已提交
251
  Input:
C
chengduoZH 已提交
252 253
       Input shape: $(N, C_{in}, D_{in}, H_{in}, W_{in})$
       Filter shape: $(C_{in}, C_{out}, D_f, H_f, W_f)$
C
chengduoZH 已提交
254
  Output:
C
chengduoZH 已提交
255 256 257
       Output shape: $(N, C_{out}, D_{out}, H_{out}, W_{out})$
  Where
  $$
258 259 260
       D_{out} = (D_{in} - 1) * strides[0] - 2 * paddings[0] + dilations[0] * (D_f - 1) + 1 \\
       H_{out} = (H_{in} - 1) * strides[1] - 2 * paddings[1] + dilations[1] * (H_f - 1) + 1 \\
       W_{out} = (W_{in} - 1) * strides[2] - 2 * paddings[2] + dilations[2] * (W_f - 1) + 1
C
chengduoZH 已提交
261
  $$
C
chengduoZH 已提交
262 263 264
)DOC");
}

C
chengduoZH 已提交
265
void ConvTransposeOpGrad::InferShape(framework::InferShapeContext* ctx) const {
C
chengduoZH 已提交
266 267 268 269 270 271 272 273 274 275
  auto in_dims = ctx->GetInputDim("Input");
  auto filter_dims = ctx->GetInputDim("Filter");
  if (ctx->HasOutput(framework::GradVarName("Input"))) {
    ctx->SetOutputDim(framework::GradVarName("Input"), in_dims);
  }
  if (ctx->HasOutput(framework::GradVarName("Filter"))) {
    ctx->SetOutputDim(framework::GradVarName("Filter"), filter_dims);
  }
}

276 277 278
framework::OpKernelType ConvTransposeOpGrad::GetExpectedKernelType(
    const framework::ExecutionContext& ctx) const {
  bool use_cudnn = ctx.Attr<bool>("use_cudnn");
279
  use_cudnn &= platform::is_gpu_place(ctx.GetPlace());
C
chengduoZH 已提交
280 281 282 283 284 285
#ifdef PADDLE_WITH_CUDA
  if (platform::is_gpu_place(ctx.GetPlace())) {
    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
    use_cudnn &= dev_ctx.cudnn_handle() != nullptr;
  }
#endif
286 287 288 289 290 291 292 293 294 295 296 297 298 299
  framework::LibraryType library_;
  if (use_cudnn) {
    library_ = framework::LibraryType::kCUDNN;
  } else {
    library_ = framework::LibraryType::kPlain;
  }

  std::string data_format = ctx.Attr<std::string>("data_format");
  framework::DataLayout layout_ = framework::StringToDataLayout(data_format);
  return framework::OpKernelType(
      framework::ToDataType(ctx.Input<Tensor>("Input")->type()), ctx.GetPlace(),
      layout_, library_);
}

C
chengduoZH 已提交
300 301 302 303
}  // namespace operators
}  // namespace paddle

namespace ops = paddle::operators;
C
chengduoZH 已提交
304

Y
Yang Yang 已提交
305 306
REGISTER_OPERATOR(conv2d_transpose, ops::ConvTransposeOp,
                  ops::Conv2DTransposeOpMaker,
307 308
                  paddle::framework::DefaultGradOpDescMaker<true>);
REGISTER_OPERATOR(conv2d_transpose_grad, ops::ConvTransposeOpGrad);
C
chengduoZH 已提交
309 310

REGISTER_OP_CPU_KERNEL(
C
chengduoZH 已提交
311
    conv2d_transpose,
Q
QI JUN 已提交
312 313
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
C
chengduoZH 已提交
314
REGISTER_OP_CPU_KERNEL(
C
chengduoZH 已提交
315
    conv2d_transpose_grad,
Q
QI JUN 已提交
316 317 318
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
                                     double>);
C
chengduoZH 已提交
319

Y
Yang Yang 已提交
320 321
REGISTER_OPERATOR(conv3d_transpose, ops::ConvTransposeOp,
                  ops::Conv3DTransposeOpMaker,
322 323
                  paddle::framework::DefaultGradOpDescMaker<true>);
REGISTER_OPERATOR(conv3d_transpose_grad, ops::ConvTransposeOpGrad);
C
chengduoZH 已提交
324 325

REGISTER_OP_CPU_KERNEL(
C
chengduoZH 已提交
326
    conv3d_transpose,
Q
QI JUN 已提交
327 328
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
C
chengduoZH 已提交
329
REGISTER_OP_CPU_KERNEL(
C
chengduoZH 已提交
330
    conv3d_transpose_grad,
Q
QI JUN 已提交
331 332 333
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
                                     double>);