Limited optimization to paddle_with_cuda.

Author: limin2021

paddle/fluid/operators/layer_norm_kernel.cu.h

@@ -172,6 +172,7 @@ __inline__ __device__ half rsqrt_(const half val) {
 }
 #endif
 
+#ifdef PADDLE_WITH_CUDA
 template <typename T, typename U, typename ScaleT = U, int VecSize = 8,
  int WARPS_M = 4, int WARPS_N = 1, int BYTES_PER_LDG = 16,
  int ELTS_PER_ROW = 1024, int THREADS_PER_WARP = 32,
@@ -281,6 +282,7 @@ __global__ __launch_bounds__(THREADS_PER_CTA) void ln_fwd_1024_kernel(
  }
  }
 }
+#endif
 
 template <typename T, typename U, bool ScaleBiasWithSameTypeX>
 using LayerNormScaleBiasT =

paddle/fluid/operators/layer_norm_op.cu

@@ -112,11 +112,11 @@ class LayerNormKernel<platform::CUDADeviceContext, T>
  } \
  } while (0)
 
+#ifdef PADDLE_WITH_CUDA
  bool can_call_1024_kernel = false;
  if (feature_size == 1024 && scale != nullptr && bias != nullptr) {
  can_call_1024_kernel = true;
  }
-
  if (can_call_1024_kernel) {
  const int WARPS_M = 4;
  const int WARPS_N = 1;
@@ -145,12 +145,15 @@ class LayerNormKernel<platform::CUDADeviceContext, T>
  y_data);
  }
  } else {
+#endif
  if (is_scale_bias_same_dtype_with_x) {
  PADDLE_LAUNCH_LAYERNORM_FWD(T, true);
  } else {
  PADDLE_LAUNCH_LAYERNORM_FWD(U, false);
  }
+#ifdef PADDLE_WITH_CUDA
  }
+#endif
 
 #undef PADDLE_LAUNCH_LAYERNORM_FWD
  }