     auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
     Tensor temp_var;
     temp_var.mutable_data<T>(var->dims(), ctx.GetPlace());
-
-    set_zero(dev_ctx, mean, static_cast<T>(0));
-    set_zero(dev_ctx, &temp_var, static_cast<T>(0));
-
     auto* x_data = x->data<T>();
     auto* y_data = y->data<T>();
     auto* mean_data = mean->data<T>();

    ds_sum += dy[index] * x[index];
    db_sum += dy[index];
  }
  CudaAtomicAddWithWarp(&ds[nc], ds_sum);
  CudaAtomicAddWithWarp(&db[nc], db_sum);
}

template <typename T>

  int i = blockIdx.x;
  AccT x_mean = static_cast<AccT>(0);
  AccT x_var = static_cast<AccT>(0);
  x += i * size;
  const int input_offset = ((uint64_t)x) % ALIGN_BYTES / sizeof(T);
  phi::Array<const T*, 1> ins;
  ins[0] = x;
  ThreadReduce<T, AccT, VecSize, 1>(ins, size, input_offset, &x_mean, &x_var);
  ReduceMeanAndVar<AccT>(mean, var, x_mean, x_var, size);
}

template <typename T, int flags>
