         static_cast<BatchNormParamType<T>>(x[i]) - mean[c];
     BatchNormParamType<T> inv_var = 1 / sqrt(variance[c] + epsilon);
     y[i] = static_cast<T>(scale[c] * x_sub_mean * inv_var + bias[c]);
+  }
+}
+
+template <typename T>
+static __global__ void InverseVariance(const BatchNormParamType<T> *variance,
+                                       const double epsilon,
+                                       const int C,
+                                       BatchNormParamType<T> *inv_variance) {
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
+  if (tid < C) {
+    inv_variance[tid] = 1 / sqrt(variance[tid] + epsilon);
+  }
+}
+
+template <typename T, phi::DataLayout layout>
+static __global__ void BN1DForwardInference(
+    const T *x,
+    const BatchNormParamType<T> *mean,
+    const BatchNormParamType<T> *inv_variance,
+    const BatchNormParamType<T> *scale,
+    const BatchNormParamType<T> *bias,
+    const int C,
+    const int N,
+    const int HxW,
+    const double epsilon,
+    T *y) {
+  int gid = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  int num = N * C * HxW;
+  for (int i = gid; i < num; i += stride) {
+    const int c = layout == phi::DataLayout::kNCHW ? i / HxW % C : i % C;
+    BatchNormParamType<T> x_sub_mean =
+        static_cast<BatchNormParamType<T>>(x[i]) - mean[c];
+    y[i] = static_cast<T>(scale[c] * x_sub_mean * inv_variance[c] + bias[c]);
   }
 }
 

                epsilon,
                transformed_y.template data<T>());
      } else {
        BNForwardInference<T, DataLayout::kNHWC>
            <<<grid_size, block_size, 0, ctx.stream()>>>(
                transformed_x.template data<T>(),
                est_mean->template data<BatchNormParamType<T>>(),
                est_var->template data<BatchNormParamType<T>>(),
                scale.template data<BatchNormParamType<T>>(),
                bias.template data<BatchNormParamType<T>>(),
                C,
                N,
                H * W * D,
                epsilon,
                transformed_y.template data<T>());
      }
    } else {
      PADDLE_ENFORCE_GPU_SUCCESS(

//         epsilon));
#else
    const bool use_native_kernel =
        (x_dims.size() == 2 ||
         (x_dims.size() == 3 && N >= CUDNN_SPATIAL_THRESHOLD));
    if (use_native_kernel) {
      const int block_size = 256;
