     auto stream = context.cuda_device_context().stream();
 
     const int x_per_block = 256;
-    int num_x_blocks = div_up(x_width, x_per_block);
+    int num_x_blocks = DivUp(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>>>(
+      ConvShiftGradX<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>>>(
+      ConvShiftDy<T><<<grid_dim, x_per_block, 0, stream>>>(
           x_data, dout_data, dy_data, x_width, y_width, y_half_width,
           batch_size);
     }


namespace {

inline int div_up(int x, int y) { return (x + y - 1) / y; }

// Some notes on the design:
//

// 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;
