 
 // Compute y gradient - initial naive implementation with atomic add.
 template <typename T>
-__global__ void ConvShiftDy(const T *x, const T *dout, T *dy, int x_width,
-                            int y_width, int y_half_width, int batch_size) {
+__global__ void ConvShiftDy(const T *x, const T *dout, int x_width, int y_width,
+                            int y_half_width, int batch_size, T *dy) {
   int i = blockIdx.x * blockDim.x + threadIdx.x;  // x index
   int j = blockIdx.y;                             // y index
   int k = blockIdx.z;                             // batch index

// y is fairly small. For large y, it would probably be more efficient
// to also tile across y.
template <typename T>
__global__ void ConvShiftForward(const T *x, const T *y, T *out, int x_width,
                                 int y_width, int y_half_width,
                                 int batch_size) {
  extern __shared__ T mem[];

  int tx = threadIdx.x;

    int y_width = Y->dims()[1];
    int y_half_width = (y_width - 1) / 2;

    auto &device_ctx = context.cuda_device_context();
    math::SetConstant<platform::GPUPlace, T> zero;

    const int x_per_block = 256;
    int num_x_blocks = DivUp(x_width, x_per_block);
