           &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;
   }
 }
 

  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,

  std::shared_ptr<void> ws_ptr = state_->pool.allocate(ws_size);

  // Performs forward operation.
  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();
  const float alpha = 1.f;
  const float beta = 0.f;
  const float *x_ptr = CDATA(x);
  const float *w_ptr = CDATA(w);
  float *y_ptr = MDATA(y);
  for (std::uint32_t bn = 0; bn < w_shape.batch(); ++bn) {
    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,
          &beta, y_desc.get(), y_ptr));
    x_ptr += x_shift;
    w_ptr += w_shift;
    y_ptr += y_shift;
  }
}

