Skip to content

[Operator Mechanism] Replace VCOPY/SCAL with direct copy and ForRange scale in CPU/GPU kernels #79468

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
wanghuancoder merged 3 commits into from
Jul 16, 2026
Merged

[Operator Mechanism] Replace VCOPY/SCAL with direct copy and ForRange scale in CPU/GPU kernels #79468

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
3e1d1bd
Replace VCOPY/SCAL with direct copy and ForRange scale in CPU/GPU ker…
feixi139 Jul 13, 2026
4920be8
fix bugs
feixi139 Jul 14, 2026
e188ee1
Merge branch 'PaddlePaddle:develop' into Paddle_fix_VCOPY
feixi139 Jul 16, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion paddle/phi/kernels/cpu/attention_lstm_kernel.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {

@A-nnonymous A-nnonymous Jul 15, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这里的int j = 0作为for-loop索引 不会越界吗？

@feixi139 feixi139 Jul 15, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

seq_len的定义为 int seq_len = static_cast(x_lod[0][i + 1] - x_lod[0][i]); 所以这里的j并不会越界，对于seq_len的强制转换为int，暂未分析，后续再看一下

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
Expand Down
15 changes: 11 additions & 4 deletions paddle/phi/kernels/cpu/elementwise_grad.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"

Expand Down Expand Up @@ -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(),

@A-nnonymous A-nnonymous Jul 15, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这里需要保留 inplace 别名安全：elementwise_add_grad 注册了 Out@GRAD -> X@GRAD 复用，Alloc(dx) 可能与 dout.data() 相同；CPU memory_utils::Copy 最终调用 std::memcpy，会形成 memcpy(p, p, n) 的未定义行为。建议像 GPU 路径一样先比较指针，仅在目标与 dout 不同时复制（dy 同理）。

@feixi139 feixi139 Jul 15, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

因为过CI需要一定时间，所以此PR暂不修改，后续PR将会修复

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));
}
}

Expand Down
12 changes: 7 additions & 5 deletions paddle/phi/kernels/cpu/gelu_grad_kernel.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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++) {

@A-nnonymous A-nnonymous Jul 15, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

同int 索引

@feixi139 feixi139 Jul 15, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

int n = std::min(x.size(), dx.size());
这里同上，n为int，所以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);
Expand Down
4 changes: 3 additions & 1 deletion paddle/phi/kernels/cpu/hsigmoid_loss_grad.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
14 changes: 3 additions & 11 deletions paddle/phi/kernels/cpu/sparse_weight_embedding_kernel.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 {
Expand Down Expand Up @@ -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_) {
Expand All @@ -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));
}
}
}
Expand Down
16 changes: 0 additions & 16 deletions paddle/phi/kernels/funcs/blas/blas.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;

Expand All @@ -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;

Expand Down Expand Up @@ -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...);
Expand Down Expand Up @@ -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...);
Expand Down
34 changes: 0 additions & 34 deletions paddle/phi/kernels/funcs/blas/blas_impl.cu.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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...));
Expand Down Expand Up @@ -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...));
Expand Down Expand Up @@ -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,
Expand Down
Loading
Loading