     const auto x_dims = x->dims();
     auto *x_data = x->data<T>();
     auto *y_data = y->mutable_data<T>(ctx.GetPlace());
-    auto *mean_data = mean->mutable_data<LayerNormParamType<T>>(ctx.GetPlace());
-    auto *var_data = var->mutable_data<LayerNormParamType<T>>(ctx.GetPlace());
-    auto *scale_data =
-        (scale == nullptr ? nullptr : scale->data<LayerNormParamType<T>>());
-    auto *bias_data =
-        (bias == nullptr ? nullptr : bias->data<LayerNormParamType<T>>());
+    auto *mean_data = mean->mutable_data<U>(ctx.GetPlace());
+    auto *var_data = var->mutable_data<U>(ctx.GetPlace());
+    auto *scale_data = (scale == nullptr ? nullptr : scale->data<U>());
+    auto *bias_data = (bias == nullptr ? nullptr : bias->data<U>());
 
     auto matrix_dim = framework::flatten_to_2d(x_dims, begin_norm_axis);
     int batch_size = static_cast<int>(matrix_dim[0]);


template <typename T>
__inline__ __device__ T rsqrt(const T val) {
  return ::rsqrt(val);
}

template <>

    : public framework::OpKernel<T> {
 public:
  void Compute(const framework::ExecutionContext &ctx) const override {
    using U = LayerNormParamType<T>;
    const float epsilon = ctx.Attr<float>("epsilon");
    auto *scale = ctx.Input<Tensor>("Scale");
    auto *bias = ctx.Input<Tensor>("Bias");
