     }
     __syncthreads();
   }
-  for(int offset = blockDim.y/2;  offset > 0;  offset /= 2) {
+  for(int offset = blockDim.y/2;  offset > 0;  offset >>= 1) {
     if(threadIdx.y < offset) {
       int idx1 = threadIdx.y * blockDim.x + threadIdx.x;
       int idx2 = (threadIdx.y + offset) * blockDim.x + threadIdx.x;

    __syncthreads();
    // Begin for inter-warp reduce
    if(npart > 1) {
      for(int offset = blockDim.y/2; offset > 0; offset /= 2) {
        if(threadIdx.y < offset) {
          int idx1 = tid;
          int idx2 = tid + offset * blockDim.x;


  if (bid < nbatch) {
    extern __shared__ char buf[];  // Shared memory
    const DType* col_vals = in_data + bid * nchannel;
    BlockWelfordOnlineSum(col_vals, nchannel, mean, sigma2, count);

    // Merge the mean/sigma2 within a warp
    // Use the Chan's Parallel Algorithm to merge all (mean, sigma2, counts)
