Remove backward op for slow 2d transposed convolution (#70333)
Summary: Pull Request resolved: https://github.com/pytorch/pytorch/pull/70333
Test Plan: Imported from OSS
Reviewed By: bdhirsh
Differential Revision: D33301402
Pulled By: jbschlosser
fbshipit-source-id: 3cfb3165589fe1620f22479b05139676d20dc493
diff --git a/aten/src/ATen/core/aten_interned_strings.h b/aten/src/ATen/core/aten_interned_strings.h
index 3f8af56..7254d14 100644
--- a/aten/src/ATen/core/aten_interned_strings.h
+++ b/aten/src/ATen/core/aten_interned_strings.h
@@ -706,7 +706,6 @@
_(aten, slow_conv_dilated3d) \
_(aten, slow_conv_dilated3d_backward) \
_(aten, slow_conv_transpose2d) \
-_(aten, slow_conv_transpose2d_backward) \
_(aten, slow_conv_transpose3d) \
_(aten, slow_conv_transpose3d_backward) \
_(aten, threshold) \
diff --git a/aten/src/ATen/native/ConvUtils.h b/aten/src/ATen/native/ConvUtils.h
index 6fedcae..5e09e1a 100644
--- a/aten/src/ATen/native/ConvUtils.h
+++ b/aten/src/ATen/native/ConvUtils.h
@@ -31,6 +31,10 @@
const at::Tensor&, const at::Tensor&, const at::Tensor&, at::IntArrayRef, at::IntArrayRef,
at::IntArrayRef, at::IntArrayRef, std::array<bool, 3>);
DECLARE_DISPATCH(slow_conv_dilated2d_backward_fn, slow_conv_dilated2d_backward_stub);
+using slow_conv_transpose2d_backward_fn = std::tuple<at::Tensor,at::Tensor,at::Tensor>(*)(
+ const at::Tensor&, const at::Tensor&, const at::Tensor&, at::IntArrayRef, at::IntArrayRef,
+ at::IntArrayRef, at::IntArrayRef, at::IntArrayRef, std::array<bool,3>);
+DECLARE_DISPATCH(slow_conv_transpose2d_backward_fn, slow_conv_transpose2d_backward_stub);
// NOLINTNEXTLINE(cppcoreguidelines-pro-type-member-init)
struct ConvParams {
diff --git a/aten/src/ATen/native/Convolution.cpp b/aten/src/ATen/native/Convolution.cpp
index 5e39e8a..7c177b1 100644
--- a/aten/src/ATen/native/Convolution.cpp
+++ b/aten/src/ATen/native/Convolution.cpp
@@ -29,6 +29,7 @@
DEFINE_DISPATCH(miopen_convolution_transpose_backward_stub);
DEFINE_DISPATCH(miopen_depthwise_convolution_backward_stub);
DEFINE_DISPATCH(slow_conv_dilated2d_backward_stub);
+DEFINE_DISPATCH(slow_conv_transpose2d_backward_stub);
REGISTER_NO_CPU_DISPATCH(cudnn_convolution_backward_stub, cudnn_convolution_backward_fn);
REGISTER_NO_CPU_DISPATCH(cudnn_convolution_transpose_backward_stub, cudnn_convolution_transpose_backward_fn);
REGISTER_NO_CPU_DISPATCH(miopen_convolution_backward_stub, miopen_convolution_backward_fn);
@@ -1461,9 +1462,9 @@
return at::slow_conv_dilated3d_backward(
grad_output, input, weight, kernel_size, params.stride, params.padding, params.dilation, output_mask);
case ConvBackend::SlowTranspose2d:
- return at::slow_conv_transpose2d_backward(
- grad_output, input, weight, kernel_size, params.stride, params.padding, params.output_padding,
- params.dilation, output_mask);
+ return slow_conv_transpose2d_backward_stub(
+ input.device().type(), grad_output, input, weight, kernel_size, params.stride, params.padding,
+ params.output_padding, params.dilation, output_mask);
case ConvBackend::SlowTranspose3d:
return at::slow_conv_transpose3d_backward(
grad_output, input, weight, kernel_size, params.stride, params.padding, params.output_padding,
diff --git a/aten/src/ATen/native/NaiveConvolutionTranspose2d.cpp b/aten/src/ATen/native/NaiveConvolutionTranspose2d.cpp
index 13532e1..560d021 100644
--- a/aten/src/ATen/native/NaiveConvolutionTranspose2d.cpp
+++ b/aten/src/ATen/native/NaiveConvolutionTranspose2d.cpp
@@ -4,6 +4,7 @@
#include <ATen/TensorUtils.h>
#include <ATen/core/Tensor.h>
+#include <ATen/native/ConvUtils.h>
#include <ATen/native/CPUBlas.h>
#include <ATen/native/im2col.h>
@@ -876,5 +877,7 @@
return std::tuple<Tensor, Tensor, Tensor>(grad_input, grad_weight, grad_bias);
}
+REGISTER_ALL_CPU_DISPATCH(slow_conv_transpose2d_backward_stub, &slow_conv_transpose2d_backward_cpu);
+
} // namespace native
} // namespace at
diff --git a/aten/src/ATen/native/cuda/NaiveConvolutionTranspose2d.cu b/aten/src/ATen/native/cuda/NaiveConvolutionTranspose2d.cu
index 653d2e6..a04d118 100644
--- a/aten/src/ATen/native/cuda/NaiveConvolutionTranspose2d.cu
+++ b/aten/src/ATen/native/cuda/NaiveConvolutionTranspose2d.cu
@@ -8,6 +8,7 @@
#include <ATen/cuda/CUDABlas.h>
#include <ATen/cuda/CUDAContext.h>
+#include <ATen/native/ConvUtils.h>
#include <ATen/native/cuda/im2col.cuh>
namespace at {
@@ -816,5 +817,7 @@
return std::tuple<Tensor, Tensor, Tensor>(grad_input, grad_weight, grad_bias);
}
+REGISTER_CUDA_DISPATCH(slow_conv_transpose2d_backward_stub, &slow_conv_transpose2d_backward_cuda);
+
} // namespace native
} // namespace at
diff --git a/aten/src/ATen/native/native_functions.yaml b/aten/src/ATen/native/native_functions.yaml
index b2f8358..2c3ef0a 100644
--- a/aten/src/ATen/native/native_functions.yaml
+++ b/aten/src/ATen/native/native_functions.yaml
@@ -9961,18 +9961,6 @@
python_module: nn
structured_delegate: slow_conv_transpose2d.out
-- func: slow_conv_transpose2d_backward.grad_output(Tensor grad_output, Tensor self, Tensor weight, int[2] kernel_size, int[2] stride, int[2] padding, int[2] output_padding, int[2] dilation, *, Tensor(a!) grad_input, Tensor(b!) grad_weight, Tensor(c!) grad_bias) -> (Tensor(a!), Tensor(b!), Tensor(c!))
- python_module: nn
- dispatch:
- CPU: slow_conv_transpose2d_backward_out_cpu
- CUDA: slow_conv_transpose2d_backward_out_cuda
-
-- func: slow_conv_transpose2d_backward.output_mask(Tensor grad_output, Tensor self, Tensor weight, int[2] kernel_size, int[2] stride, int[2] padding, int[2] output_padding, int[2] dilation, bool[3] output_mask) -> (Tensor grad_input, Tensor grad_weight, Tensor grad_bias)
- python_module: nn
- dispatch:
- CPU: slow_conv_transpose2d_backward_cpu
- CUDA: slow_conv_transpose2d_backward_cuda
-
- func: slow_conv_transpose3d.out(Tensor self, Tensor weight, int[3] kernel_size, Tensor? bias=None, int[3] stride=1, int[3] padding=0, int[3] output_padding=0, int[3] dilation=1, *, Tensor(a!) out) -> Tensor(a!)
python_module: nn
dispatch:
diff --git a/tools/autograd/derivatives.yaml b/tools/autograd/derivatives.yaml
index 03271ec..b4300cf 100644
--- a/tools/autograd/derivatives.yaml
+++ b/tools/autograd/derivatives.yaml
@@ -2027,10 +2027,7 @@
grad_output, input, weight: _convolution_double_backward(grads[0], grads[1], grads[2], grad_output, weight, input, stride, padding, dilation, false, output_padding, groups, false, false, false, false, grad_input_mask)
- name: slow_conv_transpose2d(Tensor self, Tensor weight, int[2] kernel_size, Tensor? bias=None, int[2] stride=1, int[2] padding=0, int[2] output_padding=0, int[2] dilation=1) -> Tensor
- self, weight, bias: "grad.defined() ? slow_conv_transpose2d_backward(grad, self, weight, kernel_size, stride, padding, output_padding, dilation, grad_input_mask) : std::tuple<Tensor, Tensor, Tensor>()"
-
-- name: slow_conv_transpose2d_backward.output_mask(Tensor grad_output, Tensor self, Tensor weight, int[2] kernel_size, int[2] stride, int[2] padding, int[2] output_padding, int[2] dilation, bool[3] output_mask) -> (Tensor grad_input, Tensor grad_weight, Tensor grad_bias)
- grad_output, self, weight: _convolution_double_backward(grads[0], grads[1], grads[2], grad_output, weight, self, stride, padding, dilation, true, output_padding, 1, false, false, false, false, grad_input_mask)
+ self, weight, bias: "grad.defined() ? convolution_backward(grad, self, weight, bias->sizes(), stride, padding, dilation, true, output_padding, 1, grad_input_mask) : std::tuple<Tensor, Tensor, Tensor>()"
- name: slow_conv_transpose3d(Tensor self, Tensor weight, int[3] kernel_size, Tensor? bias=None, int[3] stride=1, int[3] padding=0, int[3] output_padding=0, int[3] dilation=1) -> Tensor
self, weight, bias: "grad.defined() ? slow_conv_transpose3d_backward(grad, self, weight, kernel_size, stride, padding, output_padding, dilation, grad_input_mask) : std::tuple<Tensor, Tensor, Tensor>()"
