conv_transpose_op.cc 20.1 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
sneaxiy 已提交
16
#include <memory>
S
Siddharth Goyal 已提交
17 18
#include <string>
#include <vector>
19
#include "paddle/fluid/framework/data_layout.h"
20
#include "paddle/fluid/platform/cudnn_workspace_helper.h"
C
chengduoZH 已提交
21

J
Jacek Czaja 已提交
22 23 24 25
#ifdef PADDLE_WITH_MKLDNN
#include "paddle/fluid/platform/mkldnn_helper.h"
#endif

C
chengduoZH 已提交
26 27 28
namespace paddle {
namespace operators {

29 30
using DataLayout = framework::DataLayout;

C
chengduoZH 已提交
31
void ConvTransposeOp::InferShape(framework::InferShapeContext* ctx) const {
32 33 34 35 36 37
  PADDLE_ENFORCE_EQ(ctx->HasInput("Input"), true,
                    "Input(Input) of ConvTransposeOp should not be null.");
  PADDLE_ENFORCE_EQ(ctx->HasInput("Filter"), true,
                    "Input(Filter) of ConvTransposeOp should not be null.");
  PADDLE_ENFORCE_EQ(ctx->HasOutput("Output"), true,
                    "Output(Output) of ConvTransposeOp should not be null.");
C
chengduoZH 已提交
38 39 40

  auto in_dims = ctx->GetInputDim("Input");
  auto filter_dims = ctx->GetInputDim("Filter");
41 42
  std::vector<int> output_size =
      ctx->Attrs().Get<std::vector<int>>("output_size");
C
chengduoZH 已提交
43 44
  std::vector<int> strides = ctx->Attrs().Get<std::vector<int>>("strides");
  std::vector<int> paddings = ctx->Attrs().Get<std::vector<int>>("paddings");
C
chengduoZH 已提交
45
  std::vector<int> dilations = ctx->Attrs().Get<std::vector<int>>("dilations");
Y
Yibing Liu 已提交
46
  int groups = ctx->Attrs().Get<int>("groups");
47 48 49 50
  std::string padding_algorithm =
      ctx->Attrs().Get<std::string>("padding_algorithm");
  const DataLayout data_layout = framework::StringToDataLayout(
      ctx->Attrs().Get<std::string>("data_format"));
C
chengduoZH 已提交
51

52 53
  PADDLE_ENFORCE_EQ(in_dims.size() == 4 || in_dims.size() == 5, true,
                    "ConvTransposeOp intput should be 4-D or 5-D tensor.");
C
chengduoZH 已提交
54 55 56
  PADDLE_ENFORCE_EQ(in_dims.size(), filter_dims.size(),
                    "ConvTransposeOp input dimension and filter dimension "
                    "should be the same.");
57 58 59 60
  PADDLE_ENFORCE_EQ(
      in_dims.size() - strides.size(), 2U,
      "ConvTransposeOp input dimension and strides dimension should "
      "be consistent.");
61 62 63 64
  if (output_size.size())
    PADDLE_ENFORCE_EQ(output_size.size(), strides.size(),
                      "ConvTransposeOp output_size dimension and strides "
                      "dimension should be the same.");
C
chengduoZH 已提交
65

66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90
  const int64_t C =
      (data_layout == DataLayout::kNCHW ? in_dims[1]
                                        : in_dims[in_dims.size() - 1]);
  PADDLE_ENFORCE_EQ(
      C, filter_dims[0],
      "The number of input channels of Op(ConvTransposeOp) should "
      "be equal to the number of filter's channels.");

  framework::DDim in_data_dims;
  if (data_layout == DataLayout::kNCHW) {
    in_data_dims = framework::slice_ddim(in_dims, 2, in_dims.size());
  } else {
    in_data_dims = framework::slice_ddim(in_dims, 1, in_dims.size() - 1);
  }
  framework::DDim filter_data_dims =
      framework::slice_ddim(filter_dims, 2, filter_dims.size());
  std::vector<int> ksize = framework::vectorize<int>(filter_data_dims);
  UpdatePaddingAndDilation(&paddings, &dilations, padding_algorithm,
                           in_data_dims, strides, ksize);

  std::vector<int64_t> output_shape({in_dims[0]});
  if (data_layout == DataLayout::kNCHW) {
    output_shape.push_back(filter_dims[1] * groups);
  }
  const int offset = (data_layout == DataLayout::kNCHW ? 2 : 1);
C
chengduoZH 已提交
91
  for (size_t i = 0; i < strides.size(); ++i) {
C
chengduoZH 已提交
92
    auto filter_extent = dilations[i] * (filter_dims[i + 2] - 1) + 1;
93 94
    auto infer_shape = (in_dims[i + offset] - 1) * strides[i] -
                       paddings[2 * i] - paddings[2 * i + 1] + filter_extent;
95
    if (output_size.size()) {
96 97 98 99 100
      PADDLE_ENFORCE_EQ((output_size[i] >= infer_shape &&
                         output_size[i] < infer_shape + strides[i]),
                        true,
                        "output_size of Op(ConvTransposeOp) should be "
                        "in appropriate range.");
101 102 103 104
      output_shape.push_back(output_size[i]);
    } else {
      output_shape.push_back(infer_shape);
    }
C
chengduoZH 已提交
105
  }
106 107 108
  if (data_layout == DataLayout::kNHWC) {
    output_shape.push_back(filter_dims[1] * groups);
  }
C
chengduoZH 已提交
109
  ctx->SetOutputDim("Output", framework::make_ddim(output_shape));
C
chengduoZH 已提交
110 111
}

112 113
framework::OpKernelType ConvTransposeOp::GetExpectedKernelType(
    const framework::ExecutionContext& ctx) const {
J
Jacek Czaja 已提交
114
  framework::LibraryType library_{framework::LibraryType::kPlain};
115
  framework::DataLayout layout_ = framework::DataLayout::kAnyLayout;
116
  bool use_cudnn = ctx.Attr<bool>("use_cudnn");
C
chengduoZH 已提交
117
  use_cudnn &= platform::is_gpu_place(ctx.GetPlace());
C
chengduoZH 已提交
118 119 120 121
#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;
J
Jacek Czaja 已提交
122 123 124
    if (use_cudnn) {
      library_ = framework::LibraryType::kCUDNN;
    }
C
chengduoZH 已提交
125 126
  }
#endif
J
Jacek Czaja 已提交
127 128 129 130 131
#ifdef PADDLE_WITH_MKLDNN
  if (library_ == framework::LibraryType::kPlain &&
      platform::CanMKLDNNBeUsed(ctx)) {
    library_ = framework::LibraryType::kMKLDNN;
    layout_ = framework::DataLayout::kMKLDNN;
132
  }
J
Jacek Czaja 已提交
133
#endif
134

Y
Yu Yang 已提交
135 136
  return framework::OpKernelType(ctx.Input<Tensor>("Input")->type(),
                                 ctx.GetPlace(), layout_, library_);
137 138
}

Y
Yu Yang 已提交
139
void Conv2DTransposeOpMaker::Make() {
J
Jacek Czaja 已提交
140 141 142 143
  AddAttr<bool>("is_test",
                "(bool, default false) Set to true for inference only, false "
                "for training. Some layers may run faster when this is true.")
      .SetDefault(false);
144 145 146 147 148
  AddInput("Input",
           "(Tensor) The input tensor of convolution transpose operator. "
           "The format of input tensor is NCHW or NHWC. Where N is batch size, "
           "C is the number of input channels, H is the height of the feature, "
           "and W is the width of the feature.");
C
chengduoZH 已提交
149 150 151 152 153 154 155 156
  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.");
157 158 159 160 161 162
  AddInput("Bias",
           "(Tensor) Bias to be added to each output of filter application."
           "The format of output tensor is X (one-dimensional) of size equal"
           "to the number of output channels. Only used with MKL-DNN.")
      .AsDispensable();

C
chengduoZH 已提交
163
  AddOutput("Output",
C
chengduoZH 已提交
164
            "(Tensor) The output tensor of convolution transpose operator. "
165
            "The format of output tensor is the same as input tensor.");
166 167 168 169
  AddAttr<std::vector<int>>("output_size",
                            "(vector<int> default: []), the "
                            "size of the output tensor")
      .SetDefault({});
Y
Yibing Liu 已提交
170 171 172 173
  AddAttr<int>("groups",
               "(int default:1), the groups number of the convolution "
               "transpose operator. ")
      .SetDefault(1);
C
chengduoZH 已提交
174 175 176 177 178
  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 已提交
179 180
  AddAttr<std::vector<int>>(
      "strides",
C
chengduoZH 已提交
181
      "(vector<int> default:{1, 1}), the strides(h_stride, w_stride) of "
182
      "convolution transpose operator.")
C
chengduoZH 已提交
183
      .SetDefault({1, 1});
C
chengduoZH 已提交
184 185
  AddAttr<std::vector<int>>(
      "paddings",
C
chengduoZH 已提交
186
      "(vector<int> default:{0, 0}), the paddings(h_pad, w_pad) of convolution "
C
chengduoZH 已提交
187
      "transpose operator.")
C
chengduoZH 已提交
188
      .SetDefault({0, 0});
189 190 191 192
  AddAttr<bool>(
      "use_cudnn",
      "(bool, default false) Only used in cudnn kernel, need install cudnn")
      .SetDefault(false);
J
Jacek Czaja 已提交
193 194 195 196 197
  AddAttr<bool>("use_mkldnn",
                "(bool, default false) Only used in mkldnn kernel")
      .SetDefault(false);
  AddAttr<bool>("fuse_relu", "(bool, default false) Only used in mkldnn kernel")
      .SetDefault(false);
198 199 200 201 202 203 204 205
  AddAttr<std::string>("fuse_activation",
                       "(string, default \"\") Only used in mkldnn kernel")
      .SetDefault("");
  AddAttr<float>("fuse_alpha",
                 "(float, default 0.0) Only used in mkldnn kernel")
      .SetDefault(0.0f);
  AddAttr<float>("fuse_beta", "(float, default 0.0) Only used in mkldnn kernel")
      .SetDefault(0.0f);
206 207 208 209
  AddAttr<std::string>(
      "data_format",
      "(string, default NCHW) Only used in "
      "An optional string from: \"NHWC\", \"NCHW\". "
210 211 212 213 214 215 216 217 218
      "Specify that the data format of the input and output data is "
      "channel_first or channel_last.")
      .SetDefault("NCHW");
  AddAttr<std::string>(
      "padding_algorithm",
      "(string, default \"EXPLICIT\") An optional string from: \"EXPLICIT\","
      "\"SAME\",\"VALID\". Set to \"EXPLICIT\" for explicit padding. "
      "Set to \"SAME\" or \"VALID\" for algorithm of padding. ")
      .SetDefault("EXPLICIT");
219 220 221 222 223 224
  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.")
225
      .SetDefault(platform::kDefaultConvWorkspaceSizeLimitMB);
C
chengduoZH 已提交
226
  AddComment(R"DOC(
C
chengduoZH 已提交
227 228
Convolution2D Transpose Operator.

C
chengduoZH 已提交
229
The convolution transpose operation calculates the output based on the input, filter
C
chengduoZH 已提交
230
and dilations, strides, paddings, groups parameters. The size of each dimension of the
C
chengduoZH 已提交
231
parameters is checked in the infer-shape.
232
Input(Input) and output(Output) are in NCHW or NHWC format. Where N is batchsize, C is the
C
chengduoZH 已提交
233 234 235 236 237 238
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 已提交
239
The input(X) size and output(Out) size may be different.
C
chengduoZH 已提交
240

Y
update  
yi.wu 已提交
241
For an example:
C
chengduoZH 已提交
242
  Input:
C
chengduoZH 已提交
243 244
       Input shape: $(N, C_{in}, H_{in}, W_{in})$
       Filter shape: $(C_{in}, C_{out}, H_f, W_f)$
C
chengduoZH 已提交
245
  Output:
C
chengduoZH 已提交
246 247 248
       Output shape: $(N, C_{out}, H_{out}, W_{out})$
  Where
  $$
249 250
       H_{out} = (H_{in} - 1) * strides[0] - pad_height_top - pad_height_bottom  + dilations[0] * (H_f - 1) + 1 \\
       W_{out} = (W_{in} - 1) * strides[1] - pad_width_left  - pad_width_right + dilations[1] * (W_f - 1) + 1
C
chengduoZH 已提交
251
  $$
C
chengduoZH 已提交
252 253 254
)DOC");
}

Y
Yu Yang 已提交
255
void Conv3DTransposeOpMaker::Make() {
256 257 258 259 260 261
  AddInput(
      "Input",
      "(Tensor) The input tensor of convolution transpose operator."
      "The format of input tensor is NCDHW or NDHWC. 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 已提交
262 263
  AddInput("Filter",
           "(Tensor) The filter tensor of convolution transpose operator."
C
chengduoZH 已提交
264 265 266
           "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 已提交
267 268
           "is the depth of the filter, H is the height of the filter, and "
           "W is the width of the filter."
C
chengduoZH 已提交
269
           "We enforce groups number == 1 and padding == 0 in "
C
chengduoZH 已提交
270
           "the convolution3d transpose scenario.");
C
chengduoZH 已提交
271 272
  AddOutput("Output",
            "(Tensor) The output tensor of convolution transpose operator."
273
            "The format of output tensor is the same as input tensor."
C
chengduoZH 已提交
274
            "Where N is batch size, C is "
C
chengduoZH 已提交
275 276
            "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.");
277 278 279 280
  AddAttr<std::vector<int>>("output_size",
                            "(vector<int> default: []), the "
                            "size of the output tensor")
      .SetDefault({});
C
chengduoZH 已提交
281 282 283 284 285 286
  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 已提交
287
  AddAttr<std::vector<int>>("strides",
C
chengduoZH 已提交
288
                            "(vector<int> default:{1, 1, 1}), the "
289
                            "strides{d_stride, h_stride, w_stride} of "
C
chengduoZH 已提交
290
                            "convolution transpose operator.")
C
chengduoZH 已提交
291
      .SetDefault({1, 1, 1});
C
chengduoZH 已提交
292
  AddAttr<std::vector<int>>("paddings",
C
chengduoZH 已提交
293
                            "(vector<int> default:{0, 0, 0}), paddings(d_pad, "
C
chengduoZH 已提交
294
                            "h_pad, w_pad) of convolution transpose operator.")
C
chengduoZH 已提交
295
      .SetDefault({0, 0, 0});
296 297 298 299
  AddAttr<int>("groups",
               "(int default:1), the groups number of the convolution3d "
               "transpose operator. ")
      .SetDefault(1);
300 301 302 303
  AddAttr<bool>(
      "use_cudnn",
      "(bool, default false) Only used in cudnn kernel, need install cudnn")
      .SetDefault(false);
304 305 306
  AddAttr<bool>("use_mkldnn",
                "(bool, default false) Only used in mkldnn kernel")
      .SetDefault(false);
307 308 309 310
  AddAttr<std::string>(
      "data_format",
      "(string, default NCHW) Only used in "
      "An optional string from: \"NHWC\", \"NCHW\". "
311 312 313 314 315 316 317 318 319
      "Specify that the data format of the input and output data is "
      "channel_first or channel_last.")
      .SetDefault("NCHW");
  AddAttr<std::string>(
      "padding_algorithm",
      "(string, default \"EXPLICIT\") An optional string from: \"EXPLICIT\","
      "\"SAME\",\"VALID\". Set to \"EXPLICIT\" for explicit padding. "
      "Set to \"SAME\" or \"VALID\" for algorithm of padding. ")
      .SetDefault("EXPLICIT");
320 321 322 323 324 325
  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.")
326
      .SetDefault(platform::kDefaultConvWorkspaceSizeLimitMB);
C
chengduoZH 已提交
327
  AddComment(R"DOC(
C
chengduoZH 已提交
328 329
Convolution3D Transpose Operator.

C
chengduoZH 已提交
330
The convolution transpose operation calculates the output based on the input, filter
C
chengduoZH 已提交
331
and dilations, strides, paddings, groups parameters. The size of each dimension of the
C
chengduoZH 已提交
332
parameters is checked in the infer-shape.
333
Input(Input) and output(Output) are in NCDHW or NDHWC format. Where N is batch size, C is the
C
chengduoZH 已提交
334 335 336 337 338 339 340
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 已提交
341
The input(X) size and output(Out) size may be different.
C
chengduoZH 已提交
342

343
Example:
C
chengduoZH 已提交
344
  Input:
C
chengduoZH 已提交
345 346
       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 已提交
347
  Output:
C
chengduoZH 已提交
348 349 350
       Output shape: $(N, C_{out}, D_{out}, H_{out}, W_{out})$
  Where
  $$
351 352 353
       D_{out} = (D_{in} - 1) * strides[0] - pad_depth_front - pad_depth_back + dilations[0] * (D_f - 1) + 1 \\
       H_{out} = (H_{in} - 1) * strides[1] - pad_height_top  - pad_height_bottom + dilations[1] * (H_f - 1) + 1 \\
       W_{out} = (W_{in} - 1) * strides[2] - pad_width_left - pad_width_right + dilations[2] * (W_f - 1) + 1
C
chengduoZH 已提交
354
  $$
C
chengduoZH 已提交
355 356 357
)DOC");
}

C
chengduoZH 已提交
358
void ConvTransposeOpGrad::InferShape(framework::InferShapeContext* ctx) const {
C
chengduoZH 已提交
359 360 361 362 363 364 365 366 367 368
  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);
  }
}

369 370 371
framework::OpKernelType ConvTransposeOpGrad::GetExpectedKernelType(
    const framework::ExecutionContext& ctx) const {
  bool use_cudnn = ctx.Attr<bool>("use_cudnn");
372
  use_cudnn &= platform::is_gpu_place(ctx.GetPlace());
C
chengduoZH 已提交
373 374 375 376 377 378
#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
379 380 381 382 383 384 385
  framework::LibraryType library_;
  if (use_cudnn) {
    library_ = framework::LibraryType::kCUDNN;
  } else {
    library_ = framework::LibraryType::kPlain;
  }

386
  framework::DataLayout layout_ = framework::DataLayout::kAnyLayout;
Y
Yu Yang 已提交
387 388
  return framework::OpKernelType(ctx.Input<Tensor>("Input")->type(),
                                 ctx.GetPlace(), layout_, library_);
389 390
}

S
sneaxiy 已提交
391 392 393 394 395 396 397 398 399 400 401 402 403 404 405 406 407 408 409 410 411 412
class ConvTransposeGradOpDescMaker : public framework::SingleGradOpDescMaker {
 public:
  using framework::SingleGradOpDescMaker::SingleGradOpDescMaker;

 protected:
  std::unique_ptr<framework::OpDesc> Apply() const override {
    std::unique_ptr<framework::OpDesc> op(new framework::OpDesc());
    op->SetType(ForwardOp().Type() + "_grad");
    op->SetInput("Input", Input("Input"));
    op->SetInput("Filter", Input("Filter"));
    op->SetOutput(framework::GradVarName("Input"), InputGrad("Input"));
    op->SetOutput(framework::GradVarName("Filter"), InputGrad("Filter"));
    if (ForwardOp().Inputs().count("Bias") > 0) {
      op->SetInput("Bias", Input("Bias"));
      op->SetOutput(framework::GradVarName("Bias"), InputGrad("Bias"));
    }
    op->SetInput(framework::GradVarName("Output"), OutputGrad("Output"));
    op->SetAttrMap(Attrs());
    return op;
  }
};

C
chengduoZH 已提交
413 414 415 416
}  // namespace operators
}  // namespace paddle

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

418
// conv2d_transpose
Y
Yang Yang 已提交
419 420
REGISTER_OPERATOR(conv2d_transpose, ops::ConvTransposeOp,
                  ops::Conv2DTransposeOpMaker,
S
sneaxiy 已提交
421
                  ops::ConvTransposeGradOpDescMaker);
422
REGISTER_OPERATOR(conv2d_transpose_grad, ops::ConvTransposeOpGrad);
C
chengduoZH 已提交
423 424

REGISTER_OP_CPU_KERNEL(
C
chengduoZH 已提交
425
    conv2d_transpose,
Q
QI JUN 已提交
426 427
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
C
chengduoZH 已提交
428
REGISTER_OP_CPU_KERNEL(
C
chengduoZH 已提交
429
    conv2d_transpose_grad,
Q
QI JUN 已提交
430 431 432
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
                                     double>);
C
chengduoZH 已提交
433

434
// conv3d_transpose
Y
Yang Yang 已提交
435 436
REGISTER_OPERATOR(conv3d_transpose, ops::ConvTransposeOp,
                  ops::Conv3DTransposeOpMaker,
S
sneaxiy 已提交
437
                  ops::ConvTransposeGradOpDescMaker);
438
REGISTER_OPERATOR(conv3d_transpose_grad, ops::ConvTransposeOpGrad);
C
chengduoZH 已提交
439 440

REGISTER_OP_CPU_KERNEL(
C
chengduoZH 已提交
441
    conv3d_transpose,
Q
QI JUN 已提交
442 443
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
C
chengduoZH 已提交
444
REGISTER_OP_CPU_KERNEL(
C
chengduoZH 已提交
445
    conv3d_transpose_grad,
Q
QI JUN 已提交
446 447 448
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
                                     double>);
449 450 451 452

// depthwise conv2d_transpose
REGISTER_OPERATOR(depthwise_conv2d_transpose, ops::ConvTransposeOp,
                  ops::Conv2DTransposeOpMaker,
S
sneaxiy 已提交
453
                  ops::ConvTransposeGradOpDescMaker);
454 455 456 457 458 459 460 461 462 463 464
REGISTER_OPERATOR(depthwise_conv2d_transpose_grad, ops::ConvTransposeOpGrad);

REGISTER_OP_CPU_KERNEL(
    depthwise_conv2d_transpose,
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL(
    depthwise_conv2d_transpose_grad,
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
    ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
                                     double>);