     const int num, const int height, const int width, const int channels,
     const int size, const T negative_beta, const T cache_ratio,
     T* bottom_diff) {
-  int index = threadIdx.x + blockIdx.x * blockDim.x;
-  if (index < nthreads) {
+  CUDA_1D_KERNEL_LOOP(index, nthreads) {
     // find out the local channel offset
     int c = index % channels;
     int pre_pad = size / 2;

    const int num, const int channels, const int height,
    const int width, const int size, const T alpha_over_size,
    const T bias, T* scale) {
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index < nthreads) {
    // find out the local offset
    int w = index % width;
    int h = (index / width) % height;

      scale[(head - post_pad) * step] = bias + accum_scale * alpha_over_size;
      ++head;
    }
    // recover the pointers for the next loop.
    in -= offset;
    scale -= offset;
  }
}

