 }
 
 void CUDA::conv2d_bw_impl(
-    const Tensor &x, const Tensor &w, const Tensor &y, const Tensor &gy,
+    const Tensor &x, const Tensor &w, const Tensor &, const Tensor &gy,
     Tensor &gx, Tensor &gw) {
-  THROW_NOT_IMPLEMENTED;
+  const Shape x_shape = x.shape();
+  const Shape w_shape = w.shape();
+  const Shape y_shape = gy.shape();
+
+  // Sets target device.
+  CUDA_CALL(::cudaSetDevice(dev_id_));
+
+  // Makes descriptors.
+  const cuda::CuDNNTensorDescriptor x_desc(
+      w_shape.has_batch() ? 1 : x_shape.batch(),
+      x_shape[2], x_shape[1], x_shape[0]);
+  const cuda::CuDNNTensorDescriptor y_desc(
+      w_shape.has_batch() ? 1 : y_shape.batch(),
+      y_shape[2], y_shape[1], y_shape[0]);
+  const cuda::CuDNNFilterDescriptor w_desc(
+      w_shape[3], w_shape[2], w_shape[1], w_shape[0]);
+  const cuda::CuDNNConvolutionDescriptor conv_desc(0, 0, 1, 1, 1, 1);
+
+  // Obtains the most efficient algorithms.
+  ::cudnnConvolutionBwdDataAlgo_t x_algo;
+  ::cudnnConvolutionBwdFilterAlgo_t w_algo;
+  CUDNN_CALL(::cudnnGetConvolutionBackwardDataAlgorithm(
+        state_->cudnn.get(),
+        w_desc.get(), y_desc.get(), conv_desc.get(), x_desc.get(),
+        CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST, 0, &x_algo));
+  CUDNN_CALL(::cudnnGetConvolutionBackwardFilterAlgorithm(
+        state_->cudnn.get(),
+        x_desc.get(), y_desc.get(), conv_desc.get(), w_desc.get(),
+        CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST, 0, &w_algo));
+
+  // Obtains workspace memory.
+  std::size_t x_ws_size, w_ws_size;
+  CUDNN_CALL(::cudnnGetConvolutionBackwardDataWorkspaceSize(
+        state_->cudnn.get(),
+        w_desc.get(), y_desc.get(), conv_desc.get(), x_desc.get(),
+        x_algo, &x_ws_size));
+  CUDNN_CALL(::cudnnGetConvolutionBackwardFilterWorkspaceSize(
+        state_->cudnn.get(),
+        x_desc.get(), y_desc.get(), conv_desc.get(), w_desc.get(),
+        w_algo, &w_ws_size));
+  const std::size_t ws_size = std::max(x_ws_size, w_ws_size);
+  std::shared_ptr<void> ws_ptr = state_->pool.allocate(ws_size);
+
+  // Performs backward operations.
+  float alpha = 1.f;
+  float beta = 1.f;
+  const float *x_ptr = CDATA(x);
+  const float *w_ptr = CDATA(w);
+  const float *gy_ptr = CDATA(gy);
+  float *gx_ptr = MDATA(gx);
+  float *gw_ptr = MDATA(gw);
+  if (w_shape.has_batch()) {
+    const std::size_t x_shift = x_shape.has_batch() * x_shape.volume();
+    const std::size_t w_shift = w_shape.volume();
+    const std::size_t y_shift = y_shape.volume();
+    for (std::uint32_t bn = 0; bn < w_shape.batch(); ++bn) {
+      CUDNN_CALL(::cudnnConvolutionBackwardData(
+            state_->cudnn.get(),
+            &alpha, w_desc.get(), w_ptr, y_desc.get(), gy_ptr,
+            conv_desc.get(), x_algo, ws_ptr.get(), ws_size,
+            &beta, x_desc.get(), gx_ptr));
+      CUDNN_CALL(::cudnnConvolutionBackwardFilter(
+            state_->cudnn.get(),
+            &alpha, x_desc.get(), x_ptr, y_desc.get(), gy_ptr,
+            conv_desc.get(), w_algo, ws_ptr.get(), ws_size,
+            &beta, w_desc.get(), gw_ptr));
+      x_ptr += x_shift;
+      w_ptr += w_shift;
+      gy_ptr += y_shift;
+      gx_ptr += x_shift;
+      gw_ptr += w_shift;
+    }
+  } else {
+    CUDNN_CALL(::cudnnConvolutionBackwardData(
+          state_->cudnn.get(),
+          &alpha, w_desc.get(), w_ptr, y_desc.get(), gy_ptr,
+          conv_desc.get(), x_algo, ws_ptr.get(), ws_size,
+          &beta, x_desc.get(), gx_ptr));
+    CUDNN_CALL(::cudnnConvolutionBackwardFilter(
+          state_->cudnn.get(),
+          &alpha, x_desc.get(), x_ptr, y_desc.get(), gy_ptr,
+          conv_desc.get(), w_algo, ws_ptr.get(), ws_size,
+          &beta, w_desc.get(), gw_ptr));
+  }
 }
 
 void CUDA::inplace_multiply_const_impl(float k, Tensor &x) {

  const Shape w_shape = w.shape();
  const Shape y_shape = y.shape();

  // Sets device.
  CUDA_CALL(::cudaSetDevice(dev_id_));

  // Makes descriptors.

      y_ptr += y_shift;
    }
  } else {
    CUDNN_CALL(::cudnnConvolutionForward(
          state_->cudnn.get(),
          &alpha, x_desc.get(), x_ptr, w_desc.get(), w_ptr,
          conv_desc.get(), algo, ws_ptr.get(), ws_size,
