     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) {
+  CUDA_1D_KERNEL_LOOP(index, nthreads) {
     // find out the local offset
     int w = index % width;
     int h = (index / width) % height;

          bottom_data[(head - post_pad) * step] * accum_ratio;
      ++head;
    }
  }
}


    const int width, const int size, const T negative_beta,
    const T cache_ratio,
    T* bottom_diff) {
  CUDA_1D_KERNEL_LOOP(index, nthreads) {
    // find out the local offset
    int w = index % width;
    int h = (index / width) % height;
