 
   if (bid < nbatch) {
     extern __shared__ char buf[];  // Shared memory
-    int tid = threadIdx.x + threadIdx.y * blockDim.x;
     const DType* col_vals = in_data + bid * nchannel;
-    BlockWelfordOnlineSum(tid, nthread, col_vals, nchannel, mean, sigma2, count);
+    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)



template<typename AType, typename DType, typename IType>
__device__ __forceinline__ void BlockWelfordOnlineSum(const int tid,
                                                      const int nthread,
                                                      const DType* __restrict__ col_vals,
                                                      const int nchannel,
                                                      AType& mean,
                                                      AType& sigma2,
                                                      IType& count) {
  // Each thread takes charge of 4 consecutive numbers. This should optimize the loading speed using
  // vectorized types like float4.
  // Also, to minimize branch divergence, we split the for-loop into two parts.

    }
    __syncthreads();
  }
  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;
