     int y_width = Y->dims()[1];
     int y_half_width = (y_width - 1) / 2;
 
-    auto stream = context.cuda_device_context().stream();
+    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);


// Compute x gradient - initial naive implementation with atomic add.
template <typename T>
__global__ void ConvShiftGradX(const T *dout, const T *y, T *dx, int x_width,
                               int y_width, int y_half_width, int batch_size) {
  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, int x_width,
                                 int y_width, int y_half_width, int batch_size,
                                 T *out) {
  extern __shared__ T mem[];

  int tx = threadIdx.x;
