diff --git a/paddle/phi/kernels/cpu/attention_lstm_kernel.cc b/paddle/phi/kernels/cpu/attention_lstm_kernel.cc
index ecacaedfee885..8a82f83cd8809 100644
--- a/paddle/phi/kernels/cpu/attention_lstm_kernel.cc
+++ b/paddle/phi/kernels/cpu/attention_lstm_kernel.cc
@@ -169,7 +169,9 @@ void AttentionLSTMKernel(const Context& dev_ctx,
       bias_relu<T>(seq_len, cur_atten_x_data, &prev_cell_bias, fc_out_data);
       // 1c. fc scalar
       if (atten_scalar_data) {
-        blas.SCAL(seq_len, *atten_scalar_data, fc_out_data);
+        for (int j = 0; j < seq_len; ++j) {
+          fc_out_data[j] = fc_out_data[j] * (*atten_scalar_data);
+        }
         bias_relu<T>(seq_len, fc_out_data, atten_scalar_bias_data, fc_out_data);
       }
       // 1d. softmax
diff --git a/paddle/phi/kernels/cpu/elementwise_grad.h b/paddle/phi/kernels/cpu/elementwise_grad.h
index 08bfa5fbe6602..14d2d045fb92f 100644
--- a/paddle/phi/kernels/cpu/elementwise_grad.h
+++ b/paddle/phi/kernels/cpu/elementwise_grad.h
@@ -15,8 +15,8 @@ limitations under the License. */
 #pragma once
 
 #include "paddle/phi/backends/cpu/cpu_context.h"
+#include "paddle/phi/common/memory_utils.h"
 #include "paddle/phi/core/dense_tensor.h"
-#include "paddle/phi/kernels/funcs/blas/blas.h"
 #include "paddle/phi/kernels/funcs/eigen/common.h"
 #include "paddle/phi/kernels/funcs/elementwise_grad_base.h"
 
@@ -90,13 +90,20 @@ ElementwiseAddGrad(const CPUContext& dev_ctx,
                    DenseTensor* dx,
                    DenseTensor* dy,
                    int axis = -1) {
-  auto blas = funcs::GetBlas<CPUContext, T>(dev_ctx);
   if (dx) {
-    blas.VCOPY(dout.numel(), dout.data<T>(), dev_ctx.template Alloc<T>(dx));
+    memory_utils::Copy(dev_ctx.GetPlace(),
+                       dev_ctx.template Alloc<T>(dx),
+                       dev_ctx.GetPlace(),
+                       dout.data<T>(),
+                       static_cast<size_t>(dout.numel()) * sizeof(T));
   }
 
   if (dy) {
-    blas.VCOPY(dout.numel(), dout.data<T>(), dev_ctx.template Alloc<T>(dy));
+    memory_utils::Copy(dev_ctx.GetPlace(),
+                       dev_ctx.template Alloc<T>(dy),
+                       dev_ctx.GetPlace(),
+                       dout.data<T>(),
+                       static_cast<size_t>(dout.numel()) * sizeof(T));
   }
 }
 
diff --git a/paddle/phi/kernels/cpu/gelu_grad_kernel.cc b/paddle/phi/kernels/cpu/gelu_grad_kernel.cc
index c58160728928b..7125b425247b4 100644
--- a/paddle/phi/kernels/cpu/gelu_grad_kernel.cc
+++ b/paddle/phi/kernels/cpu/gelu_grad_kernel.cc
@@ -79,17 +79,19 @@ struct GeluGradFunctor {
       funcs::CBlas<T>::AXPY(n, static_cast<T>(M_SQRT1_2), x_data, 1, first, 1);
       funcs::CBlas<T>::VMERF(n, first, first, VML_LA);
       for (int i = 0; i < n; i++) {
-        first[i] += static_cast<T>(1);
+        first[i] = (first[i] + static_cast<T>(1)) * static_cast<T>(0.5);
       }
-      funcs::CBlas<T>::SCAL(n, static_cast<T>(0.5), first, 1);
 
       // second = (0.5 * 2/sqrt(pi) * 1/sqrt(2) * x * exp(-0.5 * x^2))
       funcs::CBlas<T>::VSQUARE(n, x_data, second);
-      funcs::CBlas<T>::SCAL(n, -static_cast<T>(0.5), second, 1);
+      for (int i = 0; i < n; i++) {
+        second[i] = second[i] * -static_cast<T>(0.5);
+      }
       funcs::CBlas<T>::VEXP(n, second, second);
       funcs::CBlas<T>::VMUL(n, x_data, second, second);
-      funcs::CBlas<T>::SCAL(
-          n, static_cast<T>(0.5 * M_2_SQRTPI * M_SQRT1_2), second, 1);
+      for (int i = 0; i < n; i++) {
+        second[i] = second[i] * static_cast<T>(0.5 * M_2_SQRTPI * M_SQRT1_2);
+      }
 
       // dx = dout * (first + second);
       funcs::CBlas<T>::VADD(n, first, second, first);
diff --git a/paddle/phi/kernels/cpu/hsigmoid_loss_grad.h b/paddle/phi/kernels/cpu/hsigmoid_loss_grad.h
index 8381604b74107..aa53d2ed4a109 100644
--- a/paddle/phi/kernels/cpu/hsigmoid_loss_grad.h
+++ b/paddle/phi/kernels/cpu/hsigmoid_loss_grad.h
@@ -81,7 +81,9 @@ void HSigmoidLossGradKernelImpl(const Context& dev_ctx,
   int64_t dim1 = pre_out_grad.dims()[1];
   for (int64_t i = 0; i < dim0; ++i) {
     T tmp = out_grad_data[i];
-    blas.SCAL(dim1, tmp, pre_out_grad_data + i * dim1);
+    for (int64_t j = 0; j < dim1; ++j) {
+      pre_out_grad_data[i * dim1 + j] = pre_out_grad_data[i * dim1 + j] * tmp;
+    }
   }
   // TODO(guosheng): multiply pre_out_grad with subgradient of clipping to
   // be consistent with the clipping in forward.
diff --git a/paddle/phi/kernels/cpu/sparse_weight_embedding_kernel.cc b/paddle/phi/kernels/cpu/sparse_weight_embedding_kernel.cc
index 9fe3cbf20afaf..c025a67ca05e9 100644
--- a/paddle/phi/kernels/cpu/sparse_weight_embedding_kernel.cc
+++ b/paddle/phi/kernels/cpu/sparse_weight_embedding_kernel.cc
@@ -17,7 +17,6 @@
 #include "paddle/phi/core/kernel_registry.h"
 #include "paddle/phi/core/utils/data_type.h"
 #include "paddle/phi/kernels/embedding_kernel.h"
-#include "paddle/phi/kernels/funcs/blas/blas.h"
 #include "paddle/phi/kernels/funcs/embedding_util.h"
 
 namespace phi {
@@ -45,7 +44,6 @@ struct EmbeddingCPUSparseFunctor {
     int64_t row_width = table_t.value().dims()[1];
     const auto* table = table_t.value().template data<T>();
     auto* output = dev_ctx_.template Alloc<T>(output_t);
-    auto input_data_type = table_t.value().dtype();
 
     for (int64_t i = 0; i < ids_numel; ++i) {
       if (padding_idx_ != kNoPadding && ids[i] == padding_idx_) {
@@ -64,15 +62,9 @@ struct EmbeddingCPUSparseFunctor {
             common::errors::InvalidArgument(
                 "the input key should be exists. But received %d.", id_index));
 
-        if (input_data_type == DataType::BFLOAT16) {
-          memcpy(output + i * row_width,
-                 table + id_index * row_width,
-                 row_width * sizeof(T));
-        } else {
-          auto blas = funcs::GetBlas<CPUContext, T>(dev_ctx_);
-          blas.VCOPY(
-              row_width, table + id_index * row_width, output + i * row_width);
-        }
+        memcpy(output + i * row_width,
+               table + id_index * row_width,
+               static_cast<size_t>(row_width) * sizeof(T));
       }
     }
   }
diff --git a/paddle/phi/kernels/funcs/blas/blas.h b/paddle/phi/kernels/funcs/blas/blas.h
index 2ee3def3e30b2..2e75f06e5a8fe 100644
--- a/paddle/phi/kernels/funcs/blas/blas.h
+++ b/paddle/phi/kernels/funcs/blas/blas.h
@@ -287,9 +287,6 @@ class Blas {
   template <typename T>
   void VDIV(int n, const T* x, const T* y, T* z) const;
 
-  template <typename T>
-  void VCOPY(int n, const T* x, T* y) const;
-
   template <typename T>
   void VEXP(int n, const T* x, T* y) const;
 
@@ -316,9 +313,6 @@ class Blas {
   void CUDOT(
       int n, const T* x, int incx, const T* y, int incy, T* result) const;
 
-  template <typename T>
-  void SCAL(int n, const T a, T* x) const;
-
   template <typename T>
   T ASUM(int n, T* x, int inc) const;
 
@@ -575,11 +569,6 @@ class BlasT : private Blas<DeviceContext> {
     Base()->template VDIV<T>(args...);
   }
 
-  template <typename... ARGS>
-  void VCOPY(ARGS... args) const {
-    Base()->template VCOPY<T>(args...);
-  }
-
   template <typename... ARGS>
   void VEXP(ARGS... args) const {
     Base()->template VEXP<T>(args...);
@@ -610,11 +599,6 @@ class BlasT : private Blas<DeviceContext> {
     Base()->template CUDOT<T>(args...);
   }
 
-  template <typename... ARGS>
-  void SCAL(ARGS... args) const {
-    Base()->template SCAL<T>(args...);
-  }
-
   template <typename... ARGS>
   T ASUM(ARGS... args) const {
     return Base()->template ASUM<T>(args...);
diff --git a/paddle/phi/kernels/funcs/blas/blas_impl.cu.h b/paddle/phi/kernels/funcs/blas/blas_impl.cu.h
index 66c86b8d7666c..37bab10cb1f24 100644
--- a/paddle/phi/kernels/funcs/blas/blas_impl.cu.h
+++ b/paddle/phi/kernels/funcs/blas/blas_impl.cu.h
@@ -48,16 +48,6 @@ struct CUBlas<float> {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cublasSaxpy(args...));
   }
 
-  template <typename... ARGS>
-  static void SCAL(ARGS... args) {
-    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cublasSscal(args...));
-  }
-
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cublasScopy(args...));
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cublasSgemv(args...));
@@ -231,16 +221,6 @@ struct CUBlas<double> {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cublasDaxpy(args...));
   }
 
-  template <typename... ARGS>
-  static void SCAL(ARGS... args) {
-    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cublasDscal(args...));
-  }
-
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cublasDcopy(args...));
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cublasDgemv(args...));
@@ -2480,20 +2460,6 @@ inline void Blas<phi::GPUContext>::CUDOT(int n,
   });
 }
 
-template <>
-template <typename T>
-void Blas<phi::GPUContext>::SCAL(int n, const T alpha, T *x) const {
-  dev_ctx_.CublasCall(
-      [&](cublasHandle_t handle) { CUBlas<T>::SCAL(handle, n, &alpha, x, 1); });
-}
-
-template <>
-template <typename T>
-void Blas<phi::GPUContext>::VCOPY(int n, const T *x, T *y) const {
-  dev_ctx_.CublasCall(
-      [&](cublasHandle_t handle) { CUBlas<T>::VCOPY(handle, n, x, 1, y, 1); });
-}
-
 template <>
 template <typename T>
 void Blas<phi::GPUContext>::GEMV(bool trans_a,
diff --git a/paddle/phi/kernels/funcs/blas/blas_impl.h b/paddle/phi/kernels/funcs/blas/blas_impl.h
index 67389cbb1acff..fdb8eb7341ff5 100644
--- a/paddle/phi/kernels/funcs/blas/blas_impl.h
+++ b/paddle/phi/kernels/funcs/blas/blas_impl.h
@@ -19,6 +19,7 @@
 
 #include <algorithm>
 #include <cmath>
+#include <cstring>
 #include <limits>
 #include <vector>
 
@@ -44,22 +45,10 @@ template <typename T>
 struct CBlas;
 
 template <>
-struct CBlas<int8_t> {
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    PADDLE_THROW(common::errors::Unimplemented(
-        "Blas VCOPY do not supported on CPU, please check your code"));
-  }
-};
+struct CBlas<int8_t> {};
 
 template <>
-struct CBlas<int16_t> {
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    PADDLE_THROW(common::errors::Unimplemented(
-        "Blas VCOPY do not supported on CPU, please check your code"));
-  }
-};
+struct CBlas<int16_t> {};
 
 template <>
 struct CBlas<phi::bfloat16> {
@@ -68,13 +57,6 @@ struct CBlas<phi::bfloat16> {
     detail::axpy(args...);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args UNUSED) {
-    PADDLE_THROW(common::errors::Unimplemented(
-        "Blas VCOPY do not supported on CPU with bfloat16,"
-        " please check your code"));
-  }
-
   template <typename... ARGS>
   static void VADD(int n,
                    const phi::bfloat16 *x,
@@ -146,11 +128,6 @@ struct CBlas<float> {
     phi::dynload::cblas_saxpy(args...);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    phi::dynload::cblas_scopy(args...);
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     phi::dynload::cblas_sgemv(args...);
@@ -161,11 +138,6 @@ struct CBlas<float> {
     return phi::dynload::cblas_sdot(args...);
   }
 
-  template <typename... ARGS>
-  static void SCAL(ARGS... args) {
-    phi::dynload::cblas_sscal(args...);
-  }
-
   template <typename... ARGS>
   static float ASUM(ARGS... args) {
     return phi::dynload::cblas_sasum(args...);
@@ -272,11 +244,6 @@ struct CBlas<double> {
     phi::dynload::cblas_daxpy(args...);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    phi::dynload::cblas_dcopy(args...);
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     phi::dynload::cblas_dgemv(args...);
@@ -287,11 +254,6 @@ struct CBlas<double> {
     return phi::dynload::cblas_ddot(args...);
   }
 
-  template <typename... ARGS>
-  static void SCAL(ARGS... args) {
-    phi::dynload::cblas_dscal(args...);
-  }
-
   template <typename... ARGS>
   static double ASUM(ARGS... args) {
     return phi::dynload::cblas_dasum(args...);
@@ -371,11 +333,6 @@ struct CBlas<phi::complex64> {
     phi::dynload::cblas_caxpy(n, &alpha, X, incX, Y, incY);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    phi::dynload::cblas_ccopy(args...);
-  }
-
   // the libmklml_intel.so paddle used has no vcAdd, vcSub,
   // vcMul, vcDiv apis before rebuild from source
   // so replace with the raw operator methods
@@ -569,11 +526,6 @@ struct CBlas<phi::complex128> {
     phi::dynload::cblas_zaxpy(n, &alpha, X, incX, Y, incY);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    phi::dynload::cblas_zcopy(args...);
-  }
-
   // the libmklml_intel.so paddle used has no vzAdd, vzSub,
   // vzMul, vzDiv apis before rebuild from source
   // so replace with the raw operator methods
@@ -768,11 +720,6 @@ struct CBlas<float> {
     phi::dynload::cblas_saxpy(args...);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    phi::dynload::cblas_scopy(args...);
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     phi::dynload::cblas_sgemv(args...);
@@ -783,11 +730,6 @@ struct CBlas<float> {
     return phi::dynload::cblas_sdot(args...);
   }
 
-  template <typename... ARGS>
-  static void SCAL(ARGS... args) {
-    phi::dynload::cblas_sscal(args...);
-  }
-
   template <typename... ARGS>
   static float ASUM(ARGS... args) {
     return phi::dynload::cblas_sasum(args...);
@@ -856,11 +798,6 @@ struct CBlas<double> {
     phi::dynload::cblas_daxpy(args...);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    phi::dynload::cblas_dcopy(args...);
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     phi::dynload::cblas_dgemv(args...);
@@ -871,11 +808,6 @@ struct CBlas<double> {
     return phi::dynload::cblas_ddot(args...);
   }
 
-  template <typename... ARGS>
-  static void SCAL(ARGS... args) {
-    phi::dynload::cblas_dscal(args...);
-  }
-
   template <typename... ARGS>
   static double ASUM(ARGS... args) {
     return phi::dynload::cblas_dasum(args...);
@@ -944,11 +876,6 @@ struct CBlas<phi::complex64> {
     phi::dynload::cblas_caxpy(n, &alpha, X, incX, Y, incY);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    phi::dynload::cblas_ccopy(args...);
-  }
-
   template <typename... ARGS>
   static void VADD(int n,
                    const phi::complex64 *a,
@@ -1117,11 +1044,6 @@ struct CBlas<phi::complex128> {
     phi::dynload::cblas_zaxpy(n, &alpha, X, incX, Y, incY);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    phi::dynload::cblas_zcopy(args...);
-  }
-
   template <typename... ARGS>
   static void VADD(int n,
                    const phi::complex128 *a,
@@ -1292,11 +1214,6 @@ struct CBlas<float> {
     cblas_saxpy(args...);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    cblas_scopy(args...);
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     cblas_sgemv(args...);
@@ -1320,11 +1237,6 @@ struct CBlas<double> {
     cblas_daxpy(args...);
   }
 
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    cblas_dcopy(args...);
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     cblas_dgemv(args...);
@@ -1338,11 +1250,6 @@ struct CBlas<double> {
 
 template <>
 struct CBlas<phi::complex64> {
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    cblas_ccopy(args...);
-  }
-
   template <typename... ARGS>
   static void AXPY(int n,
                    const phi::complex64 alpha,
@@ -1406,11 +1313,6 @@ struct CBlas<phi::complex64> {
 
 template <>
 struct CBlas<phi::complex128> {
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    cblas_zcopy(args...);
-  }
-
   template <typename... ARGS>
   static void AXPY(int n,
                    const phi::complex128 alpha,
@@ -1505,10 +1407,6 @@ struct CBlas<phi::float16> {
     PADDLE_THROW(common::errors::Unimplemented(
         "float16 DOT not supported on CPU, please check your code"));
   };
-  static void SCAL(...) {
-    PADDLE_THROW(common::errors::Unimplemented(
-        "float16 SCAL not supported on CPU, please check your code"));
-  };
   static void ASUM(...) {
     PADDLE_THROW(common::errors::Unimplemented(
         "float16 ASUM not supported on CPU, please check your code"));
@@ -1786,12 +1684,6 @@ void Blas<CPUContext>::AXPY(int n, T alpha, const T *x, T *y) const {
   CBlas<T>::AXPY(n, alpha, x, 1, y, 1);
 }
 
-template <>
-template <typename T>
-void Blas<CPUContext>::VCOPY(int n, const T *x, T *y) const {
-  CBlas<T>::VCOPY(n, x, 1, y, 1);
-}
-
 template <>
 template <typename T>
 void Blas<CPUContext>::VADD(int n, const T *x, const T *y, T *z) const {
@@ -1801,7 +1693,9 @@ void Blas<CPUContext>::VADD(int n, const T *x, const T *y, T *z) const {
   if (x == z) {
     this->template AXPY<T>(n, (T)(1.), y, z);
   } else {
-    this->template VCOPY<T>(n, y, z);
+    if (y != z) {
+      std::memcpy(z, y, n * sizeof(T));
+    }
     this->template AXPY<T>(n, (T)(1.), x, z);
   }
 #endif
@@ -1898,19 +1792,6 @@ T Blas<CPUContext>::DOT(int n, const T *x, const T *y) const {
 #endif
 }
 
-template <>
-template <typename T>
-void Blas<CPUContext>::SCAL(int n, const T a, T *x) const {
-#if defined(PADDLE_WITH_MKLML) || defined(PADDLE_WITH_HML)
-  CBlas<T>::SCAL(n, a, x, 1);
-#else
-  // try to find if openblas support cblas_scal
-  for (int i = 0; i < n; ++i) {
-    x[i] = a * x[i];
-  }
-#endif
-}
-
 template <>
 template <typename T>
 T Blas<CPUContext>::ASUM(int n, T *x, int inc) const {
diff --git a/paddle/phi/kernels/funcs/blas/blas_impl.hip.h b/paddle/phi/kernels/funcs/blas/blas_impl.hip.h
index 72054ff0d9e99..b00988d872c56 100644
--- a/paddle/phi/kernels/funcs/blas/blas_impl.hip.h
+++ b/paddle/phi/kernels/funcs/blas/blas_impl.hip.h
@@ -43,16 +43,6 @@ struct CUBlas<float> {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_saxpy(args...));
   }
 
-  template <typename... ARGS>
-  static void SCAL(ARGS... args) {
-    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_sscal(args...));
-  }
-
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_scopy(args...));
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_sgemv(args...));
@@ -118,16 +108,6 @@ struct CUBlas<double> {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_daxpy(args...));
   }
 
-  template <typename... ARGS>
-  static void SCAL(ARGS... args) {
-    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_dscal(args...));
-  }
-
-  template <typename... ARGS>
-  static void VCOPY(ARGS... args) {
-    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_dcopy(args...));
-  }
-
   template <typename... ARGS>
   static void GEMV(ARGS... args) {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_dgemv(args...));
@@ -1273,20 +1253,6 @@ void Blas<GPUContext>::AXPY(int n, T alpha, const T *x, T *y) const {
   });
 }
 
-template <>
-template <typename T>
-void Blas<GPUContext>::SCAL(int n, const T alpha, T *x) const {
-  dev_ctx_.CublasCall(
-      [&](rocblas_handle handle) { CUBlas<T>::SCAL(handle, n, &alpha, x, 1); });
-}
-
-template <>
-template <typename T>
-void Blas<GPUContext>::VCOPY(int n, const T *x, T *y) const {
-  dev_ctx_.CublasCall(
-      [&](rocblas_handle handle) { CUBlas<T>::VCOPY(handle, n, x, 1, y, 1); });
-}
-
 template <>
 template <typename T>
 void Blas<GPUContext>::GEMV(bool trans_a,
diff --git a/paddle/phi/kernels/funcs/selected_rows_functor.cc b/paddle/phi/kernels/funcs/selected_rows_functor.cc
index 2cf67ed95fbd6..b7ea95ba95766 100644
--- a/paddle/phi/kernels/funcs/selected_rows_functor.cc
+++ b/paddle/phi/kernels/funcs/selected_rows_functor.cc
@@ -272,13 +272,16 @@ struct SelectedRowsSumTo<CPUContext, T> {
 
     auto* in2_value = input2->mutable_value();
     auto* in2_data = in2_value->data<T>();
-    auto blas = phi::funcs::GetBlas<CPUContext, T>(dev_ctx);
     size_t offset = 0u;
     for (size_t i = 0u; i != input1.size(); ++i) {
       auto& in_value = input1[i]->value();
       const auto* in_data = in_value.data<T>();
       offset += input2_offsets[i];
-      blas.VCOPY(in_value.numel(), in_data, in2_data + offset);
+      memory_utils::Copy(input2->place(),
+                         in2_data + offset,
+                         input1[i]->place(),
+                         in_data,
+                         static_cast<size_t>(in_value.numel()) * sizeof(T));
     }
   }
 };
diff --git a/paddle/phi/kernels/funcs/sequence_pooling.cc b/paddle/phi/kernels/funcs/sequence_pooling.cc
index 39e1ae6400899..29de06b819a77 100644
--- a/paddle/phi/kernels/funcs/sequence_pooling.cc
+++ b/paddle/phi/kernels/funcs/sequence_pooling.cc
@@ -18,7 +18,6 @@ limitations under the License. */
 
 #include "paddle/common/enforce.h"
 
-#include "paddle/phi/kernels/funcs/blas/blas.h"
 #include "paddle/phi/kernels/funcs/eigen/common.h"
 #include "paddle/phi/kernels/funcs/jit/kernels.h"
 #include "paddle/phi/kernels/funcs/math_function.h"
@@ -323,7 +322,6 @@ class SumSeqPoolGradFunctor {
                           out_w));
     const T* out_g_data = out_grad.data<T>();
     T* in_g_data = dev_ctx.template Alloc<T>(in_grad);
-    auto blas = funcs::GetBlas<CPUContext, T>(dev_ctx);
     for (int i = 0; i < static_cast<int>(lod.size()) - 1; ++i) {
       int64_t h = static_cast<int64_t>(lod[i + 1] - lod[i]);
       if (h == 0) continue;
@@ -331,7 +329,8 @@ class SumSeqPoolGradFunctor {
       const T* out_pos = out_g_data + i * out_w;
       T* in_pos = in_g_data + in_offset;
       for (int r = 0; r != h; ++r) {
-        blas.VCOPY(in_w, out_pos, in_pos + r * in_w);
+        std::memcpy(
+            in_pos + r * in_w, out_pos, static_cast<size_t>(in_w) * sizeof(T));
       }
     }
   }
diff --git a/paddle/phi/kernels/fusion/cpu/fusion_lstm_kernel.cc b/paddle/phi/kernels/fusion/cpu/fusion_lstm_kernel.cc
index 424a6186557bb..aa14d9670a2cf 100644
--- a/paddle/phi/kernels/fusion/cpu/fusion_lstm_kernel.cc
+++ b/paddle/phi/kernels/fusion/cpu/fusion_lstm_kernel.cc
@@ -12,6 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
+#include <cstring>
 #include <string>
 
 #include "paddle/phi/core/kernel_registry.h"
@@ -279,10 +280,10 @@ void BatchCompute(const Context &dev_ctx,
     const T *c0_data = c0->data<T>();
     prev_h_data = reordered_h0_data;
     prev_c_data = reordered_c0_data;
-    size_t sz = D;
+    size_t sz_bytes = sizeof(T) * D;
     for (int i = 0; i < max_bs; ++i) {
-      blas.VCOPY(sz, h0_data + seq_order[i] * D, reordered_h0_data);
-      blas.VCOPY(sz, c0_data + seq_order[i] * D, reordered_c0_data);
+      std::memcpy(reordered_h0_data, h0_data + seq_order[i] * D, sz_bytes);
+      std::memcpy(reordered_c0_data, c0_data + seq_order[i] * D, sz_bytes);
       reordered_h0_data += D;
       reordered_c0_data += D;
     }
diff --git a/paddle/phi/kernels/fusion/cpu/fusion_seqconv_eltadd_relu_kernel.cc b/paddle/phi/kernels/fusion/cpu/fusion_seqconv_eltadd_relu_kernel.cc
index a5f0b8d174c87..f0eeb670a80b5 100644
--- a/paddle/phi/kernels/fusion/cpu/fusion_seqconv_eltadd_relu_kernel.cc
+++ b/paddle/phi/kernels/fusion/cpu/fusion_seqconv_eltadd_relu_kernel.cc
@@ -77,7 +77,6 @@ void FusionSeqConvEltAddReluKernel(const Context& dev_ctx,
       dst_data = dst_data + up_pad * src_mat_w;
       int copy_size = col_mat_w_sz - up_pad * src_mat_w_sz;
       for (int j = 0; j < up_pad; ++j) {
-        // blas.VCOPY?
         std::memcpy(dst_data, src_data, copy_size);
         dst_data += (col_mat_w - src_mat_w);
         copy_size += src_mat_w_sz;
diff --git a/paddle/phi/kernels/impl/addmm_grad_kernel_impl.h b/paddle/phi/kernels/impl/addmm_grad_kernel_impl.h
index 5d5908372cc0a..47391f6ba0dbd 100644
--- a/paddle/phi/kernels/impl/addmm_grad_kernel_impl.h
+++ b/paddle/phi/kernels/impl/addmm_grad_kernel_impl.h
@@ -105,7 +105,6 @@ void AddmmGradKernel(const Context& dev_ctx,
   }
 
   auto blas = funcs::GetBlas<Context, T>(dev_ctx);
-  auto mt_blas = funcs::GetBlas<Context, MPType>(dev_ctx);
   if (input_grad) {
     dev_ctx.template Alloc<T>(input_grad);
     total_elems = in_dims[0] * in_dims[1];
@@ -151,29 +150,18 @@ void AddmmGradKernel(const Context& dev_ctx,
                                          .reshape(eigen_dinput_shape)
                                          .template cast<T>();
       }
-    } else {
-      // The VCOPY does not support the float16, bfloat16
-      if (!is_float16_or_bfloat16 && !is_big_tensor) {
-        mt_blas.VCOPY(
-            total_elems, out_grad.data<MPType>(), input_grad->data<MPType>());
-      } else {
-        funcs::ForRange<Context> for_range(dev_ctx, total_elems);
-        CopyOrScaleFunctor<T> functor(
-            1, out_grad.data<T>(), input_grad->data<T>(), total_elems);
-        for_range(functor);
-      }
-    }
-
-    // The SCAL does not support the float16, bfloat16
-    if (!is_float16_or_bfloat16 && !is_big_tensor) {
-      mt_blas.SCAL(total_elems, beta, input_grad->data<MPType>());
     } else {
       funcs::ForRange<Context> for_range(dev_ctx, total_elems);
       CopyOrScaleFunctor<T> functor(
-          beta, input_grad->data<T>(), input_grad->data<T>(), total_elems);
+          1, out_grad.data<T>(), input_grad->data<T>(), total_elems);
       for_range(functor);
     }
 
+    funcs::ForRange<Context> for_range(dev_ctx, total_elems);
+    CopyOrScaleFunctor<T> functor(
+        beta, input_grad->data<T>(), input_grad->data<T>(), total_elems);
+    for_range(functor);
+
     if (input.dims().size() == 1) {
       input_grad->Resize(input.dims());
     }
@@ -193,28 +181,20 @@ void AddmmGradKernel(const Context& dev_ctx,
     total_elems = x.dims()[0] * x.dims()[1];
     // x_grad = out_grad * y'. x_grad: M x K, out_grad : M x N, y : K x N
     blas.MatMul(out_grad, false, y, true, x_grad);
-    if (!is_float16_or_bfloat16 && !is_big_tensor) {
-      mt_blas.SCAL(total_elems, alpha, x_grad->data<MPType>());
-    } else {
-      funcs::ForRange<Context> for_range(dev_ctx, total_elems);
-      CopyOrScaleFunctor<T> functor(
-          alpha, x_grad->data<T>(), x_grad->data<T>(), total_elems);
-      for_range(functor);
-    }
+    funcs::ForRange<Context> for_range(dev_ctx, total_elems);
+    CopyOrScaleFunctor<T> functor(
+        alpha, x_grad->data<T>(), x_grad->data<T>(), total_elems);
+    for_range(functor);
   }
   if (y_grad) {
     dev_ctx.template Alloc<T>(y_grad);
     total_elems = x.dims()[1] * y.dims()[1];
     // y_grad = x' * out_grad. y_grad K x N, out_grad : M x N, x : M x K
     blas.MatMul(x, true, out_grad, false, y_grad);
-    if (!is_float16_or_bfloat16 && !is_big_tensor) {
-      mt_blas.SCAL(total_elems, alpha, y_grad->data<MPType>());
-    } else {
-      funcs::ForRange<Context> for_range(dev_ctx, total_elems);
-      CopyOrScaleFunctor<T> functor(
-          alpha, y_grad->data<T>(), y_grad->data<T>(), total_elems);
-      for_range(functor);
-    }
+    funcs::ForRange<Context> for_range(dev_ctx, total_elems);
+    CopyOrScaleFunctor<T> functor(
+        alpha, y_grad->data<T>(), y_grad->data<T>(), total_elems);
+    for_range(functor);
   }
 }
 
diff --git a/paddle/phi/kernels/impl/baddbmm_grad_kernel_impl.h b/paddle/phi/kernels/impl/baddbmm_grad_kernel_impl.h
index 67bc4e35c8876..b32a5824f4abb 100644
--- a/paddle/phi/kernels/impl/baddbmm_grad_kernel_impl.h
+++ b/paddle/phi/kernels/impl/baddbmm_grad_kernel_impl.h
@@ -89,7 +89,6 @@ void BaddbmmGradKernel(const Context& dev_ctx,
   }
 
   auto blas = funcs::GetBlas<Context, T>(dev_ctx);
-  auto mt_blas = funcs::GetBlas<Context, MPType>(dev_ctx);
   if (input_grad) {
     dev_ctx.template Alloc<T>(input_grad);
     total_elems = in_dims[0] * in_dims[1] * in_dims[2];
@@ -180,28 +179,17 @@ void BaddbmmGradKernel(const Context& dev_ctx,
                                          .reshape(eigen_dinput_shape)
                                          .template cast<T>();
       }
-    } else {
-      // The VCOPY does not support the float16, bfloat16
-      if (!is_float16_or_bfloat16) {
-        mt_blas.VCOPY(
-            total_elems, out_grad.data<MPType>(), input_grad->data<MPType>());
-      } else {
-        funcs::ForRange<Context> for_range(dev_ctx, total_elems);
-        BCopyOrScaleFunctor<T> functor(
-            1, out_grad.data<T>(), input_grad->data<T>(), total_elems);
-        for_range(functor);
-      }
-    }
-
-    // The SCAL does not support the float16, bfloat16
-    if (!is_float16_or_bfloat16) {
-      mt_blas.SCAL(total_elems, beta, input_grad->data<MPType>());
     } else {
       funcs::ForRange<Context> for_range(dev_ctx, total_elems);
       BCopyOrScaleFunctor<T> functor(
-          beta, input_grad->data<T>(), input_grad->data<T>(), total_elems);
+          1, out_grad.data<T>(), input_grad->data<T>(), total_elems);
       for_range(functor);
     }
+
+    funcs::ForRange<Context> for_range(dev_ctx, total_elems);
+    BCopyOrScaleFunctor<T> functor(
+        beta, input_grad->data<T>(), input_grad->data<T>(), total_elems);
+    for_range(functor);
   }
   if (x_grad) {
     dev_ctx.template Alloc<T>(x_grad);
@@ -248,14 +236,10 @@ void BaddbmmGradKernel(const Context& dev_ctx,
                        K_dim * N_dim);
     }
     if (FLAGS_use_accuracy_compatible_kernel) {
-      if (!is_float16_or_bfloat16) {
-        mt_blas.SCAL(total_elems, alpha, x_grad->data<MPType>());
-      } else {
-        funcs::ForRange<Context> for_range(dev_ctx, total_elems);
-        BCopyOrScaleFunctor<T> functor(
-            alpha, x_grad->data<T>(), x_grad->data<T>(), total_elems);
-        for_range(functor);
-      }
+      funcs::ForRange<Context> for_range(dev_ctx, total_elems);
+      BCopyOrScaleFunctor<T> functor(
+          alpha, x_grad->data<T>(), x_grad->data<T>(), total_elems);
+      for_range(functor);
     }
   }
   if (y_grad) {
@@ -303,14 +287,10 @@ void BaddbmmGradKernel(const Context& dev_ctx,
                        M_dim * N_dim);
     }
     if (FLAGS_use_accuracy_compatible_kernel) {
-      if (!is_float16_or_bfloat16) {
-        mt_blas.SCAL(total_elems, alpha, y_grad->data<MPType>());
-      } else {
-        funcs::ForRange<Context> for_range(dev_ctx, total_elems);
-        BCopyOrScaleFunctor<T> functor(
-            alpha, y_grad->data<T>(), y_grad->data<T>(), total_elems);
-        for_range(functor);
-      }
+      funcs::ForRange<Context> for_range(dev_ctx, total_elems);
+      BCopyOrScaleFunctor<T> functor(
+          alpha, y_grad->data<T>(), y_grad->data<T>(), total_elems);
+      for_range(functor);
     }
   }
 }
diff --git a/paddle/phi/kernels/selected_rows/cpu/lookup_table_kernel.cc b/paddle/phi/kernels/selected_rows/cpu/lookup_table_kernel.cc
index 74fc5d496519c..81366745b37ce 100644
--- a/paddle/phi/kernels/selected_rows/cpu/lookup_table_kernel.cc
+++ b/paddle/phi/kernels/selected_rows/cpu/lookup_table_kernel.cc
@@ -57,7 +57,6 @@ void LookupTableKernel(const Context &dev_ctx,
   int64_t row_width = table_t.value().dims()[1];
   const auto *table = table_t.value().data<T>();
   auto *output = dev_ctx.template Alloc<T>(output_t);
-  auto input_data_type = table_t.value().dtype();
   for (int64_t i = 0; i < ids_numel; ++i) {
     if (padding_idx != kNoPadding && ids[i] == padding_idx) {
       memset(output + i * row_width, 0, row_width * sizeof(T));
@@ -72,18 +71,9 @@ void LookupTableKernel(const Context &dev_ctx,
         auto id_index = table_t.GetIndexFromId(ids[i]);
 
         if (id_index != -1) {
-          if (input_data_type == phi::DataType::INT8 ||
-              input_data_type == phi::DataType::INT16 ||
-              input_data_type == phi::DataType::BFLOAT16) {
-            memcpy(output + i * row_width,
-                   table + id_index * row_width,
-                   row_width * sizeof(T));
-          } else {
-            auto blas = funcs::GetBlas<CPUContext, T>(dev_ctx);
-            blas.VCOPY(row_width,
-                       table + id_index * row_width,
-                       output + i * row_width);
-          }
+          memcpy(output + i * row_width,
+                 table + id_index * row_width,
+                 static_cast<size_t>(row_width) * sizeof(T));
         } else {
           memset(output + i * row_width, 0, row_width * sizeof(T));
         }
@@ -101,17 +91,9 @@ void LookupTableKernel(const Context &dev_ctx,
             common::errors::InvalidArgument(
                 "the input key should be exists. But received %d.", id_index));
 
-        if (input_data_type == phi::DataType::INT8 ||
-            input_data_type == phi::DataType::INT16 ||
-            input_data_type == phi::DataType::BFLOAT16) {
-          memcpy(output + i * row_width,
-                 table + id_index * row_width,
-                 row_width * sizeof(T));
-        } else {
-          auto blas = funcs::GetBlas<CPUContext, T>(dev_ctx);
-          blas.VCOPY(
-              row_width, table + id_index * row_width, output + i * row_width);
-        }
+        memcpy(output + i * row_width,
+               table + id_index * row_width,
+               static_cast<size_t>(row_width) * sizeof(T));
       }
     }
   }
diff --git a/paddle/phi/kernels/sparse/gpu/addmm_grad_kernel.cu b/paddle/phi/kernels/sparse/gpu/addmm_grad_kernel.cu
index 9c915800de209..055fe381fe55c 100644
--- a/paddle/phi/kernels/sparse/gpu/addmm_grad_kernel.cu
+++ b/paddle/phi/kernels/sparse/gpu/addmm_grad_kernel.cu
@@ -15,14 +15,29 @@ limitations under the License. */
 #include "paddle/phi/kernels/sparse/addmm_grad_kernel.h"
 
 #include "paddle/phi/backends/gpu/gpu_context.h"
+#include "paddle/phi/common/memory_utils.h"
 #include "paddle/phi/core/kernel_registry.h"
 #include "paddle/phi/kernels/empty_kernel.h"
 #include "paddle/phi/kernels/funcs/blas/blas.h"
+#include "paddle/phi/kernels/funcs/for_range.h"
 #include "paddle/phi/kernels/sparse/matmul_grad_kernel.h"
 
 namespace phi {
 namespace sparse {
 
+template <typename T>
+struct ScaleFunctor {
+  ScaleFunctor(const T scale, T* data) : scale_(scale), data_(data) {}
+
+  HOSTDEVICE void operator()(int64_t idx) const {
+    data_[idx] = data_[idx] * scale_;
+  }
+
+ private:
+  const T scale_;
+  T* data_;
+};
+
 template <typename T, typename Context>
 void AddmmCooDenseGradKernel(const Context& dev_ctx,
                              const DenseTensor& input,
@@ -34,17 +49,30 @@ void AddmmCooDenseGradKernel(const Context& dev_ctx,
                              DenseTensor* dinput,
                              SparseCooTensor* dx,
                              DenseTensor* dy) {
-  auto blas = funcs::GetBlas<Context, T>(dev_ctx);
   if (dinput) {
     dinput->Resize(input.dims());
     dev_ctx.template Alloc<T>(dinput);
 
-    blas.VCOPY(input.numel(), dout.data<T>(), dinput->data<T>());
-    blas.SCAL(input.numel(), beta, dinput->data<T>());
+    memory_utils::Copy(dev_ctx.GetPlace(),
+                       dinput->data<T>(),
+                       dev_ctx.GetPlace(),
+                       dout.data<T>(),
+                       static_cast<size_t>(input.numel()) * sizeof(T),
+                       dev_ctx.stream());
+    funcs::ForRange<Context> for_range(dev_ctx, input.numel());
+    ScaleFunctor<T> functor(static_cast<T>(beta), dinput->data<T>());
+    for_range(functor);
   }
   DenseTensor dout_scale = EmptyLike<T, Context>(dev_ctx, dout);
-  blas.VCOPY(dout.numel(), dout.data<T>(), dout_scale.data<T>());
-  blas.SCAL(dout.numel(), alpha, dout_scale.data<T>());
+  memory_utils::Copy(dev_ctx.GetPlace(),
+                     dout_scale.data<T>(),
+                     dev_ctx.GetPlace(),
+                     dout.data<T>(),
+                     static_cast<size_t>(dout.numel()) * sizeof(T),
+                     dev_ctx.stream());
+  funcs::ForRange<Context> for_range(dev_ctx, dout.numel());
+  ScaleFunctor<T> functor(static_cast<T>(alpha), dout_scale.data<T>());
+  for_range(functor);
   MatmulCooDenseGradKernel<T, Context>(dev_ctx, x, y, dout_scale, dx, dy);
 }
 
@@ -60,17 +88,30 @@ void AddmmCsrDenseGradKernel(const Context& dev_ctx,
                              DenseTensor* dinput,
                              SparseCsrTensor* dx,
                              DenseTensor* dy) {
-  auto blas = funcs::GetBlas<Context, T>(dev_ctx);
   if (dinput) {
     dinput->Resize(input.dims());
     dev_ctx.template Alloc<T>(dinput);
 
-    blas.VCOPY(input.numel(), dout.data<T>(), dinput->data<T>());
-    blas.SCAL(input.numel(), beta, dinput->data<T>());
+    memory_utils::Copy(dev_ctx.GetPlace(),
+                       dinput->data<T>(),
+                       dev_ctx.GetPlace(),
+                       dout.data<T>(),
+                       static_cast<size_t>(input.numel()) * sizeof(T),
+                       dev_ctx.stream());
+    funcs::ForRange<Context> for_range(dev_ctx, input.numel());
+    ScaleFunctor<T> functor(static_cast<T>(beta), dinput->data<T>());
+    for_range(functor);
   }
   DenseTensor dout_scale = EmptyLike<T, Context>(dev_ctx, dout);
-  blas.VCOPY(dout.numel(), dout.data<T>(), dout_scale.data<T>());
-  blas.SCAL(dout.numel(), alpha, dout_scale.data<T>());
+  memory_utils::Copy(dev_ctx.GetPlace(),
+                     dout_scale.data<T>(),
+                     dev_ctx.GetPlace(),
+                     dout.data<T>(),
+                     static_cast<size_t>(dout.numel()) * sizeof(T),
+                     dev_ctx.stream());
+  funcs::ForRange<Context> for_range(dev_ctx, dout.numel());
+  ScaleFunctor<T> functor(static_cast<T>(alpha), dout_scale.data<T>());
+  for_range(functor);
   MatmulCsrDenseGradKernel<T, Context>(dev_ctx, x, y, dout_scale, dx, dy);
 }