Fix int32 overflow issues for large tensor support in paddle/phi/kernels/impl

[Copilot]

修复 paddle/phi/kernels/impl 中的 int32 溢出问题以支持大张量

概述

本 PR 解决了 paddle/phi/kernels/impl 目录中的 int32 溢出问题，以改进对大张量（超过 20 亿元素）的支持。这是修复 PaddlePaddle 代码库中 int32 溢出漏洞的系统性工作的第 3 阶段。

背景

之前阶段已完成：

• 阶段 1：paddle/phi/kernels/funcs 目录 ✓
• 阶段 2：paddle/phi/kernels/gpu 目录 ✓

大张量可能在以下常见模式中导致整数溢出：

• tensor.numel() 返回 int64_t，但赋值给 int 时会对超过 20 亿元素的张量进行截断
• CUDA 线程索引：处理大张量时 blockIdx.x * blockDim.x 可能溢出 int32
• 单个张量维度可能超过 INT32_MAX（例如，具有 30 亿元素的 1D 张量的 dims[0] = 30 亿）

已完成的修改

在 30 个文件中修复了 85 处潜在的 int32 溢出，并添加了大张量验证检查：

已修改的文件

1. elementwise_grad_kernel_impl.h - 逐元素梯度操作的 CUDA 内核和 CPU 循环
2. accuracy_check_kernel_impl.h - 精度检查的 CUDA 内核（通用 + 复数特化版本）
3. isclose_kernel_impl.h - isclose 操作的 CUDA 内核（模板 + 4 个特化版本）
4. renorm_impl.h - 重归一化内核的网格大小计算
5. unstack_kernel_impl.h - unstack 操作的元素计数变量
6. kldiv_loss_grad_kernel_impl.h - KL 散度梯度的元素计数变量
7. kldiv_loss_kernel_impl.h - KL 散度前向传播的批次维度
8. svdvals_grad_kernel_impl.h - SVD 梯度的批次计数计算
9. gumbel_softmax_kernel_impl.h - Gumbel-Softmax 前向传播的轴维度
10. gumbel_softmax_grad_kernel_impl.h - Gumbel-Softmax 梯度的轴维度
11. lrn_kernel_impl.h - N、C、H、W 张量维度 + 函数签名 + 大张量验证检查
12. frame_kernel_impl.h - 帧操作的帧数和序列长度
13. frame_grad_kernel_impl.h - 帧梯度的帧数和序列长度
14. stft_kernel_impl.h - STFT 操作的帧数和序列长度
15. stft_grad_kernel_impl.h - STFT 梯度的帧数和序列长度
16. fold_kernel_impl.h - Fold 操作的批次大小和输入平面数
17. fold_grad_kernel_impl.h - Fold 梯度的批次大小和输入平面数
18. unfold_kernel_impl.h - Unfold 操作的批次大小
19. unfold_grad_kernel_impl.h - Unfold 梯度的批次大小
20. lstm_kernel_impl.h - LSTM 的帧大小
21. lstsq_kernel_impl.h - 最小二乘的矩阵维度
22. qr_grad_kernel_impl.h - QR 分解梯度的矩阵维度
23. spectral_norm_grad_kernel_impl.h - 谱归一化梯度的维度
24. warpctc_grad_kernel_impl.h - Warp CTC 梯度的序列维度
25. warprnnt_grad_kernel_impl.h - Warp RNN-T 梯度的维度
26. gru_unit_kernel_impl.h - GRU 单元的批次和帧大小
27. spectral_norm_kernel_impl.h - 谱归一化的高度和宽度
28. svd_grad_kernel_impl.h - SVD 梯度的矩阵维度
29. conv_kernel_impl.h - 卷积操作的批次大小和步长
30. conv_grad_kernel_impl.h - 卷积梯度的批次大小和步长

应用的修复模式

1. CUDA 线程索引

// 之前：当 blockIdx.x * blockDim.x > INT32_MAX 时可能溢出
int tid = threadIdx.x + blockIdx.x * blockDim.x;

// 之后：安全转换防止溢出
int64_t tid = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;

2. 张量元素计数

// 之前：numel() 返回 int64_t 但赋值给 int
int numel = input_grad->numel();
int grid = (numel + block - 1) / block;

// 之后：使用 int64_t 保留完整范围
int64_t numel = input_grad->numel();
int64_t grid = (numel + block - 1) / block;

3. 维度访问

// 之前：单个维度可能超过 INT32_MAX
int axis_dim = x.dims()[axis];
int N = x_dims[0];
int batch_size = static_cast<int>(x.dims()[0]);

// 之后：对维度值使用 int64_t
int64_t axis_dim = x.dims()[axis];
int64_t N = x_dims[0];
int64_t batch_size = x.dims()[0];

4. 函数签名

// 之前：参数类型不匹配
struct LRNFunctor {
  void operator()(..., int N, int C, int H, int W, ...);
};

// 之后：更新以匹配 int64_t 调用者
struct LRNFunctor {
  void operator()(..., int64_t N, int64_t C, int64_t H, int64_t W, ...);
};

5. 大张量验证检查

// 对于底层实现仍使用 int 的操作，添加验证
// TODO(large-tensor): LRN GPU kernel implementation still uses int for dimensions.
// Need to update GPU kernel to support dimensions > INT32_MAX.
PADDLE_ENFORCE_LE(
    N * C * H * W,
    std::numeric_limits<int>::max(),
    common::errors::InvalidArgument(
        "The total number of elements (N*C*H*W = %ld) exceeds the maximum "
        "value that int can represent (%d). LRN operation does not support "
        "such large tensors yet.",
        N * C * H * W,
        std::numeric_limits<int>::max()));

TODO 列表

impl 目录（剩余 12 个文件，约 62 处修复）

• 复杂文件（需要仔细审查）：

  • llm_int8_matmul_kernel_impl.h（10 处模式）- 整个 API 使用 int32_t
  • matmul_kernel_impl.h（25 处模式）- BLAS API 兼容性考虑
  • conv_transpose_kernel_impl.h, conv_transpose_grad_kernel_impl.h
  • deformable_conv_grad_kernel_impl.h

• 简单文件（仅 dims[] 模式）：

  • bilinear_grad_kernel_impl.h
  • cross_entropy2_kernel_impl.h
  • eigh_grad_kernel_impl.h
  • slogdeterminant_grad_kernel_impl.h
  • softmax_kernel_impl.h, softmax_grad_kernel_impl.h

cpu 目录（约 361 处修复）

• 模式 2（numel）：50 处匹配
• 模式 3（offset）：13 处匹配
• 模式 4（dims）：298 处匹配

其他子目录（约 868 处修复）

• fusion/ 目录（462 处匹配）
• sparse/ 目录（165 处匹配）
• legacy/ 目录（116 处匹配）
• primitive/ 目录（54 处匹配）
• gpudnn/ 目录（29 处匹配）
• stride/ 目录（25 处匹配）
• selected_rows/ 目录（11 处匹配）
• strings/ 目录（6 处匹配）

测试

• ✅ 代码审查通过（无问题）
• ✅ CodeQL 安全扫描通过（无漏洞）
• ✅ 所有更改遵循之前成功 PR 的既定模式
• ✅ 更改最小且精确

范围

本 PR 完成了 impl 目录的 71%（42 个文件中的 30 个）。剩余工作包括：

• paddle/phi/kernels/impl 中的 12 个文件（约 62 处修复）
• paddle/phi/kernels/cpu 目录（约 361 处修复）
• 其他子目录：fusion/、sparse/、legacy/ 等（约 868 处修复）

项目总范围：约 200 个文件中的约 1,376 处潜在修复

影响

这些更改使 PaddlePaddle 能够安全处理超过 20 亿元素的张量，而不会出现整数溢出错误或静默数据损坏。这些修复向后兼容，在扩展对大张量工作负载的支持的同时保持现有的 API 契约。

大张量支持限制

某些操作的底层实现（如 LRN GPU 内核）仍然使用 int 参数。在这些情况下，已添加运行时验证检查，如果张量大小超过 INT32_MAX，将抛出清晰的错误消息。这些检查标记为 TODO(large-tensor)，表示需要在未来更新底层实现以完全支持大张量。

Original prompt

大Tensor目前通过“基于业务配置篡改成大Tensor配置方式，以测试case驱动大Tensor问题修复”完成了314个API的前反向问题的修复。
这种模式可能存在遗漏，比如：

• case不够丰富没能覆盖Kernel所有代码路径
• 部分Kernel没有构造出合法case，没能测试
• 有些Fused Kernel没有官方API，未能测试到

修复中，我们发现，出现int32问题的位置往往发生在如下位置：numel、stride、n = shape[1]、index、offset、block.x*threadIdx.x->static_cast<int64_t>(blockIdx.x)threadIdx.x。
因此，未来，我们可以通过关键字模糊检索的方式进一步排雷：
int .threadIdx.|int .blockDim.|int .blockIdx.|int32_t .threadIdx.|int32_t .blockDim.|int32_t .blockIdx.
int .=.numel|int32_t .=.numel
int .=.offset|int32_t .=.offset
int .=.dims[|int32_t .=.dims[|int .=.dims()[|int32_t .=.*dims()[

int .*=.strides[|int32_t .=.strides[|int .=.strides()[|int32_t .=.*strides()[

exclude：cinn,fluid,test,xpu

预计还能发现几百甚至上千处值得修改的代码行。
虽仍不能100%保证大Tensor的安全，但能进一步提升质量，降低业务出问题的概率。Torch也有一些API发生CUDA700或大Tensor精度异常的问题。

大tensor情况：

1. 复杂修改：内核中没有进行兜底的循环，gird和block是三维结构等 需要进行大量的代码修改 correlation支持大tensor PaddlePaddle/Paddle#75383
2. 简单修改：简单的将int32修改为int64 对gird和block的size大小进行限制
3. 无需修改：在执行前已经对数据大小进行了限制，或者数据不会出现较大的情况，所以不需要进行修改

修改进程：

• Paddle/paddle/phi/kernels/funcs 文件的 int .threadIdx.|int .blockDim.|int .blockIdx.|int32_t .threadIdx.|int32_t .blockDim.|int32_t .blockIdx.
• Paddle/paddle/phi/kernels/funcs 文件的 int .*=.numel|int32_t .=.*numel
• Paddle/paddle/phi/kernels/funcs 文件的 int .*=.offset|int32_t .=.*offset
• Paddle/paddle/phi/kernels/funcs 文件的 int .*=.dims[|int32_t .=.dims[|int .=.dims()[|int32_t .=.*dims()[
• Paddle/paddle/phi/kernels/funcs 文件的 int .*=.strides[|int32_t .=.strides[|int .=.strides()[|int32_t .=.*strides()[
• Paddle/paddle/phi/kernels/gpu/ 文件的 int .threadIdx.|int .blockDim.|int .blockIdx.|int32_t .threadIdx.|int32_t .blockDim.|int32_t .blockIdx.
• Paddle/paddle/phi/kernels/gpu/ 文件的 int .*=.numel|int32_t .=.*numel
• Paddle/paddle/phi/kernels/gpu/ 文件的 int .*=.offset|int32_t .=.*offset
• Paddle/paddle/phi/kernels/gpu/ 文件的 int .*=.dims[|int32_t .=.dims[|int .=.dims()[|int32_t .=.*dims()[
• Paddle/paddle/phi/kernels/gpu/ 文件的 int .*=.strides[|int32_t .=.strides[|int .=.strides()[|int32_t .=.*strides()[ 。大tensor情况：
  1. 复杂修改：内核中没有进行兜底的循环，gird和block是三维结构等 需要进行大量的代码修改 参考：correlation支持大tensor PaddlePaddle/Paddle#75383
  2. 简单修改：简单的将int32修改为int64 对gird和block的size大小进行限制
  3. 无需修改：在执行前已经对数据大小进行了限制，或者数据不会出现较大的情况，所以不需要进行修改

修复计划：对phi目录分别使用五个正则式进行检索，分别进行修复再生成pr
1. [进行中] Paddle/paddle/phi/kernels文件夹是最主要的文件夹 其下几个文件大量含有检索结果

i. [已完成] Paddle/paddle/phi/kernels/funcs 目录
ii. [已完成] Paddle/paddle/phi/kernels/gpu 目录
iii. [未开始] Paddle/paddle/phi/kernels/impl 目录
iv. [未开始] Paddle/paddle/phi/kernels/cpu 目录
v. [未开始] Paddle/paddle/phi/kernels目录其他文件
目前完成 funcs 和 gpu 目录 在kernels文件夹中占比约20% - 30%。 现在，需要你首先确定哪些位置还需要修改，然后按步骤依次修改他们。

---

💡 You can make Copilot smarter by setting up custom instructions, customizing its development environment and configuring Model Context Protocol (MCP) servers. Learn more Copilot coding agent tips in the docs.

[zrr1999]

@copilot 描述使用中文，并且加入TODO列表。对于部分没有条件支持bit tensor的地方，添加报错信息指明这里不支持。

[Copilot]

@copilot 描述使用中文，并且加入TODO列表。对于部分没有条件支持bit tensor的地方，添加报错信息指明这里不支持。

已完成以下修改（提交 f1d21ba）：

1. ✅ 已将 PR 描述更新为中文
2. ✅ 已添加详细的 TODO 列表，列出剩余需要修复的文件和目录
3. ✅ 已在 LRN 内核中添加大张量验证检查

对于底层实现仍使用 int 的操作（如 LRN GPU 内核），添加了运行时检查。当张量元素总数超过 INT32_MAX 时，会抛出清晰的错误消息，说明该操作暂不支持如此大的张量。这些位置标记为 TODO(large-tensor)，表示需要在未来更新底层实现。

示例验证检查：

PADDLE_ENFORCE_LE(
    N * C * H * W,
    std::numeric_limits<int>::max(),
    common::errors::InvalidArgument(
        "The total number of elements (N*C*H*W = %ld) exceeds the maximum "
        "value that int can represent (%d). LRN operation does not support "
        "such large tensors yet.",
        N * C * H * W,
        std::numeric_limits<int>::max()));

[zrr1999]

@copilot 继续完成impl文件夹下的内容

[Copilot]

@copilot 继续完成impl文件夹下的内容

正在继续完成 impl 文件夹的修复工作。

已完成的进展（提交 91b6c2f, ef6cb1a, 80e2535）：

已修复 27 个文件（占 impl 目录的 64%），共 72 处 int32 溢出问题：

最新修复的文件：

• stft_kernel_impl.h, stft_grad_kernel_impl.h（帧数和序列长度）
• fold_kernel_impl.h, fold_grad_kernel_impl.h（批次大小和输入平面数）
• unfold_kernel_impl.h, unfold_grad_kernel_impl.h（批次大小）
• lstm_kernel_impl.h（帧大小）
• lstsq_kernel_impl.h, qr_grad_kernel_impl.h（矩阵维度）
• spectral_norm_kernel_impl.h, spectral_norm_grad_kernel_impl.h（高度和宽度）
• warpctc_grad_kernel_impl.h, warprnnt_grad_kernel_impl.h（序列维度）
• gru_unit_kernel_impl.h（批次和帧大小）

剩余工作（约 15 个文件，75 处修复）：

• 复杂文件：matmul（25 处）、llm_int8_matmul（10 处）、conv 相关（约 15 处）、svd_grad（6 处）
• 简单文件：bilinear_grad, cross_entropy2, eigh_grad, slogdeterminant_grad, softmax 相关

继续修复中...