       mean = mean_buf[0];
       sigma2 = sigma2_buf[0] / nchannel;
     }
-    if(bid == 0 && threadIdx.x == 0 && threadIdx.y == 0) printf("mean=%g, var=%g\n", mean, sigma2);
     // Calculate the out_data: gamma * (x - mean) / sqrt(var + eps) + beta
     DType std_eps = sqrt(sigma2 + static_cast<DType>(eps));
     DType invstd_eps = static_cast<DType>(1) / std_eps;

  int nbatch = data_shape[0];
  int nchannel = data_shape[1];
  float eps = param.eps;
  std::cout << "nbatch = " << nbatch << ", nchannel = " << nchannel << ", eps = " << eps << std::endl;
  int ngrid_x = (nbatch > kMaxGridDim) ? (nbatch + kBaseGridNum - 1) / kBaseGridNum : nbatch;
  int ngrid_y = (nbatch > kMaxGridDim) ? kBaseGridNum : 1;
  std::cout << "ngrid = " << ngrid_x << " " << ngrid_y << std::endl;
  int nthread_y = 0;
  const dim3 dimGrid(ngrid_x, ngrid_y, 1);
  if(nchannel <= 32) {

  extern __shared__ char buf[];  // Shared memory size

  if (bid < nbatch) {
    int tid = threadIdx.x + threadIdx.y * blockDim.x;
    const DType* col_vals = in_data + bid * nchannel;
    // Each thread takes charge of 4 consecutive numbers
