     int nshared = nthread_y > 1 ? nthread_y * sizeof(DType) + (nthread_y / 2) * sizeof(DType) : 0;
     CheckLaunchParam(dimGrid, dimBlock);
     LayerNormFusedForwardKernelContig<<<dimGrid, dimBlock, nshared, stream>>>
-     (nbatch, nchannel, eps,
+     (nbatch, nchannel, static_cast<DType>(eps),
       in_data.dptr<DType>(), gamma.dptr<DType>(), beta.dptr<DType>(),
       out_data.dptr<DType>(), mean_data.dptr<DType>(), std_data.dptr<DType>());
     MSHADOW_CUDA_POST_KERNEL_CHECK(LayerNormFusedForwardKernelContig);

template<typename DType>
__global__ void LayerNormFusedForwardKernelContig(const int nbatch,
                                                  const int nchannel,
                                                  const float eps,
                                                  const DType* __restrict__ in_data,
                                                  const DType* __restrict__ gamma,
                                                  const DType* __restrict__ beta,

 *  TODO(sxjscience) Explore the possibility of int lhs_count and rhs_count
 */
template<typename DType>
__device__ __inline__ void chan_merge_partition(const DType lhs_mean,
                                                const DType lhs_sigma2,
                                                const DType lhs_count,
                                                DType& rhs_mean,
                                                DType& rhs_sigma2,
                                                DType& rhs_count) {
  DType delta = rhs_mean - lhs_mean;
  DType nA = lhs_count;
  DType nB = rhs_count;
