  *
  */
 template<typename DType>
-__device__ warp_merge_mean_sigma2(DType mean, DType sigma2, DType count) {
+__device__ void warp_merge_mean_sigma2(DType mean, DType sigma2, DType count) {
   for (int l = 0; l <= 4; ++l) {
     int src_lane = (threadIdx.x + (1<<l)) & 31;
     DType meanB = WARP_SHFL(mean, src_lane);

    nthread_y = 4;
  }
  const dim3 dimBlock(32, nthread_y, 1);
  int nshared = nthread_y > 1 ? nthread_y * sizeof(DType) + (nthread_y / 2) * sizeof(DType) : 0;
  cudaStream_t stream = Stream<gpu>::GetStream(ctx.get_stream<gpu>());
  CheckLaunchParam(dimGrid, dimBlock);
  LayerNormFusedForwardKernelContig<<<dimBlock, dimGrid, nshared, stream>>>
    (nbatch, nchannel, eps, in_data.dptr_, gamma.dptr_, beta.dptr_, out_data.dptr_, std_data.dptr_);
  MSHADOW_CUDA_POST_KERNEL_CHECK(LayerNormFusedForwardKernelContig);
}

template<>

 *
 */
template<typename DType>
__device__ void welford_online_sum_step(const DType curr,
                                   DType& mean,
                                   DType& sigma2,
                                   DType& count) {
