   for (; l + 3 < nchannel; l += 4 * nthread) {
 #pragma unroll
     for (int i = 0; i < 4; ++i) {
-      welford_online_sum_step(col_vals[l + i], mean, sigma2, count);
+      StepWelfordOnlineSum(col_vals[l + i], mean, sigma2, count);
     }
   }
   for(; l < nchannel; ++l) {
-    welford_online_sum_step(col_vals[l], mean, sigma2, count);
+    StepWelfordOnlineSum(col_vals[l], mean, sigma2, count);
   }
 }
 

        }
        __syncthreads();
        if (threadIdx.x == 0 && threadIdx.y < offset) {
          chan_merge_partition(mean_buf[threadIdx.y],
                               sigma2_buf[threadIdx.y],
                               count_buf[threadIdx.y], mean, sigma2, count);
        }
        __syncthreads();
      }



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.
