  *
  */
 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);

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

                        const OpContext& ctx, const std::vector<TBlob>& inputs,
                        const std::vector<OpReqType>& req,
                        const std::vector<TBlob>& outputs) {
  using namespace mshadow;
  CHECK_EQ(inputs.size(), 3U);
  mxnet::TShape data_shape(2);
  mxnet::TShape mean_shape(1);
