 // y is fairly small. For large y, it would probably be more efficient
 // to also tile across y.
 template <typename T>
-__global__ void conv_shift_forward(const T *x, const T *y, T *out, int x_width,
-                                   int y_width, int y_half_width,
-                                   int batch_size) {
+__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;

    auto stream = context.cuda_device_context().stream();

    const int x_per_block = 256;
    int num_x_blocks = div_up(x_width, x_per_block);
    dim3 grid_dim(num_x_blocks, y_width, batch_size);

    if (dX) {
      T *dx_data = dX->mutable_data<T>(context.GetPlace());
      cudaMemsetAsync(dx_data, 0, dX->numel() * sizeof(T), stream);
      conv_shift_dx<T><<<grid_dim, x_per_block, 0, stream>>>(
          dout_data, y_data, dx_data, x_width, y_width, y_half_width,
          batch_size);
    }
    if (dY) {
      T *dy_data = dY->mutable_data<T>(context.GetPlace());
      cudaMemsetAsync(dy_data, 0, dY->numel() * sizeof(T), stream);
      conv_shift_dy<T><<<grid_dim, x_per_block, 0, stream>>>(
          x_data, dout_data, dy_data, x_width, y_width, y_half_width,
          batch_size);
    }

    int y_half_width = (y_width - 1) / 2;

    const int x_per_block = 256;
    int num_x_blocks = DivUp(x_width, x_per_block);
    int mem_per_block = (x_per_block + 2 * y_width) * sizeof(T);

    dim3 grid_dim(num_x_blocks, batch_size);

    auto stream = context.cuda_device_context().stream();

    ConvShiftForward<T><<<grid_dim, x_per_block, mem_per_block, stream>>>(
        x_data, y_data, out_data, x_width, y_width, y_half_width, batch_size);
  }
};
