[Cherry-Pick][Operator Mechanism] fix Broadcast support for big tensor(#79439) by feixi139 · Pull Request #79440 · PaddlePaddle/Paddle

feixi139 · 2026-07-08T07:42:15Z

PR Category

Operator Mechanism

PR Types

Bug fixes

Description

BroadcastConfig 中使用 FastDivMod 对输出线性索引进行维度坐标拆解。当 tensor 某一维度大小超过 INT_MAX（2^31 - 1）时，int 类型发生溢出，导致索引计算错误，广播结果产生静默错误。

修复

将 BroadcastConfig 中的 FastDivMod 替换为 FastDivMod<int64_t>，移除 PADDLE_ENFORCE_LE_INT_MAX 的限制，使广播操作支持单维度超过 2^31 的大 tensor。

devPR:#79439

是否引起精度变化

否

fix bugs

PaddlePaddle-bot · 2026-07-09T09:36:52Z

🤖 Paddle-CI-Agent | ci_status_monitor | 2026-07-16 08:30:14 UTC+08:00

CI报告基于以下代码生成（30分钟更新一次）:
PR commit: f65527e | Merge base: eaa369b (branch: release/3.4)

1 Required任务 : 47/48 通过

总执行（rerun次数）	总任务	✅ 通过	❌ 失败	⏳ 运行中	⏸️ 等待中	跳过
164(82)	82	80	1	1	0	0

任务	错误类型	置信度	日志
`Api-Benchmark / Performance data storage`	未知	中	Job

2 失败详情

🔴 Api-Benchmark / Performance data storage — 未知（置信度: 中）

分析器: 通用分析(fallback)

失败用例: (按根因聚类合并)

用例	错误摘要
`paddle.Tensor.ceil_`	API Benchmark 判定动态图 CPU 前向调度性能相对 baseline 下降，进入 `doubt` 区间
`Tensor_reciprocal__0`、`numel_0`、`dropout2d_0` 等 9 个复核 case	日志列为 `double_check_case`；最终 `worse` 为空，`doubt` 仅包含 `paddle.Tensor.ceil_`

关键日志:

double_check_case are: ['Tensor_reciprocal__0', 'numel_0', 'Tensor_ceil__0', 'equal_all_0', 'isinf_0', 'Tensor_exp__0', 'dropout3d_0', 'Tensor_sqrt__0', 'dropout2d_0']
以下为pr79440引入之后，api调度性能相对于baseline的变化。worse表示性能下降超过30%的api，doubt表示性能下降为15%~30%之间的api
{'worse': [], 'doubt': [{'paddle.Tensor.ceil_': -1.159176230008546}]}
详情差异请点击以下链接查询: http://paddletest.baidu-int.com:8081/#/paddle/benchmark/apiBenchmark/report/56470&56445
Exception: 该pr会导致动态图cpu前向调度性能下降，请修复！！！

根因摘要: API benchmark 疑似 CPU 调度退化
日志明确是性能门禁失败，不是编译错误或单元测试断言失败。本 PR 仅修改 paddle/phi/kernels/funcs/broadcast_function.h 和 paddle/phi/kernels/primitive/datamover_primitives.h 中 GPU/XPU broadcast/datamover 的索引类型与启动逻辑；未发现对 paddle.Tensor.ceil_ 或 dropout2d CPU 调度路径的直接修改，因此目前无法确认是 PR 代码直接导致，按未知处理。

修复建议:

先用日志给出的 baseline wheel 与 PR wheel 复现 paddle.Tensor.ceil_ 的 CPU 前向调度差异。
若复现稳定，再排查本 PR 新增的公共头文件/模板变更是否间接影响 wheel 或 dispatcher；若无法复现，建议 rerun 或请 API benchmark 维护者确认基线波动。

关联变更: paddle/phi/kernels/funcs/broadcast_function.h、paddle/phi/kernels/primitive/datamover_primitives.h；未发现直接触碰失败 API 的 CPU 实现。

paddle-bot · 2026-07-14T14:00:11Z

你的PR提交成功，感谢你对开源项目的贡献!
请关注后续CI自动化测试结果，详情请参考Paddle-CI手册。
Your PR has been submitted. Thanks for your contribution!
Please wait for the result of CI firstly. See Paddle CI Manual for details.

A-nnonymous · 2026-07-15T09:43:11Z

 */
 struct BroadcastConfig {
-  funcs::FastDivMod<int> divmoders[DDim::kMaxRank];
+  funcs::FastDivMod<int64_t> divmoders[DDim::kMaxRank];


这里当前无条件使用 FastDivMod<int64_t>，会让普通小 tensor 广播也走 64 位 divmod；但本 PR 的 launcher 仍专门选择 uint32_t 快路径。源 PR #79439 最新 head 已将 BroadcastConfig 模板化为 BroadcastConfig，小 tensor 保留 32 位、仅大 tensor 用 64 位。建议同步该最新修正，避免 release 分支引入全量 broadcast 热路径的性能回退。

done。已修改

wanghuancoder · 2026-07-15T09:43:49Z

  static __device__ __forceinline__ void Apply(const Array &ins,
                                               ArgsT *args,
-                                               uint32_t index_bc[][VecSize]) {
+                                               uint64_t index_bc[][VecSize]) {


这里可以用IndexT吗？

已修改为IndexT

wanghuancoder · 2026-07-15T09:44:40Z

 #else
  if (LoadType == kBroadcast) {
-    uint32_t index_bc[Arity][VecSize] = {0};
+    uint64_t index_bc[Arity][VecSize] = {0};


这里可以用IndexT吗？

已修改为IndexT

wanghuancoder · 2026-07-15T09:44:52Z

 #pragma unroll
    for (int k = 0; k < VecSize; ++k) {
-      uint32_t idx = thread_offset + k;
+      uint64_t idx = thread_offset + k;


这里可以用IndexT吗？

已修改为IndexT

wanghuancoder · 2026-07-15T09:45:16Z

+  uint64_t block_offset =
+      static_cast<uint64_t>(BLOCK_ID_X) * BLOCK_NUM_X * read_lens;
+  uint64_t stride = static_cast<uint64_t>(BLOCK_NUM_X) * GRID_NUM_X * read_lens;
+  for (; block_offset < static_cast<uint64_t>(main_offset);
+       block_offset += stride) {


这里可以用IndexT吗？

wanghuancoder · 2026-07-15T09:45:30Z

+  int64_t num =
+      static_cast<int64_t>(numel) - static_cast<int64_t>(block_offset);


这里可以用IndexT吗？

wanghuancoder · 2026-07-15T09:46:14Z

                                            configs,
-                                            tail_tid,
-                                            block_offset,
+                                            static_cast<uint32_t>(num),


这里为啥用int32？

尾部 scalarize 的元素数受单 block 工作量约束，恒在 uint32 内，Impl 尾部形参也是 uint32_t，

wanghuancoder · 2026-07-15T09:47:19Z

                              NumOuts,
                              VecSize,
-                              kElementwise>
+                              false>


这里为啥要改成false？

我知道了，你还是需要XPU和GPU分开

此处为git显示为问题

wanghuancoder · 2026-07-15T09:48:56Z

+    const IndexT numel = static_cast<IndexT>(classifier.numel);
+    const int threads = 64;
+    const int blocks = 8;
+    int read_lens = classifier.configs[0].buf_len;
+    auto stream = dev_ctx.x_context()->xpu_stream;
+    const IndexT block_len = static_cast<IndexT>(read_lens) * threads;
+    const IndexT main_offset = (numel / block_len) * block_len;
+    const IndexT tail_tid = numel % block_len;


这里在GPU Kernel外部，可以用int64

此处 host 侧的 IndexT 是用来选择 device 侧模板特化的标签，而不是参与大数运算。这里将决定实例化哪一个kernel模版

wanghuancoder · 2026-07-15T09:50:21Z

 struct BroadcastConfig {
-  funcs::FastDivMod<int> divmoders[DDim::kMaxRank];
+  funcs::FastDivMod<int64_t> divmoders[DDim::kMaxRank];


这个Config在Broadcast Kernel里使用。因此是GPU运算使用的代码。需要用IndexT区分

wanghuancoder · 2026-07-15T09:51:57Z

+  launch_with_index_t(uint32_t{0});
+#else
+  const uint64_t index_limit = std::numeric_limits<uint32_t>::max();
+  const uint64_t grid_stride =
+      static_cast<uint64_t>(gpu_config.block_per_grid.x) *
+      static_cast<uint64_t>(gpu_config.GetBlockSize()) * VecSize;
+  if (static_cast<uint64_t>(classifier.numel) <= index_limit &&
+      grid_stride <= index_limit) {
+    launch_with_index_t(uint32_t{0});
  } else {
-    VectorizedBroadcastKernel<Functor, OutT, Arity, NumOuts, VecSize, kMixed>
-        <<<blocks, threads, 0, stream>>>(classifier.ins_data,
-                                         classifier.outs_data,
-                                         classifier.use_broadcast,
-                                         numel,
-                                         classifier.configs,
-                                         main_offset,
-                                         tail_tid,
-                                         VecSize,
-                                         func);
+    launch_with_index_t(uint64_t{0});


XPU的逻辑可以不用动。
GPU的逻辑区分std::numeric_limits<uint32_t>::max()即可。
结合实际情况看看launch_with_index_t是否必要，如果必要就封装，不必要可以平铺写。

lambda内需要根据uint32、uint64构造多个kernel变体，封装代码更为简洁。

Make the GPU BroadcastConfig a template on IndexT so normal-sized tensors use 32-bit FastDivMod again instead of always paying for 64-bit divmod. Build the typed configs at launch time (once IndexT is chosen) from the simplified dims stored on the classifier, and tighten the uint32 fast-path gate to INT_MAX for 32-bit FastDivMod ctor safety.

wanghuancoder

LGTM

PaddlePaddle-bot

🤖 Paddle-CI-Agent | pr_review | 2026-07-15 19:17:37

📋 Review 摘要

PR 概述：修复 Broadcast 在大 tensor 场景下的 GPU/XPU 索引溢出问题。
变更范围：paddle/phi/kernels/funcs/broadcast_function.h、paddle/phi/kernels/primitive/datamover_primitives.h
影响面 Tag：Operator Mechanism

问题

级别	文件	概述
🟡 建议	`paddle/phi/kernels/funcs/broadcast_function.h:657`	大 tensor 的 64 位 broadcast 路径缺少回归覆盖

📝 PR 规范检查

符合规范。历史标题问题已修复：当前标题已使用 [Cherry-Pick][Operator Mechanism] ...(#79439) 格式，描述包含 develop 原 PR 链接并填写了精度变化。

总体评价

实现方向与问题描述一致，索引类型和尾块 underflow 风险已有针对性处理。建议补上触发 64 位 broadcast 分支的回归测试，避免这类静默错误后续回归。

PaddlePaddle-bot · 2026-07-15T11:19:47Z

-                                         tail_tid,
-                                         VecSize,
-                                         func);
+    launch_with_index_t(uint64_t{0});


🟡 建议 这里新增了 uint64_t 广播索引路径，但本 PR 没有同步增加覆盖该路径的回归测试。

这条路径只有 classifier.numel 或 grid_stride 超过 INT_MAX 时才会执行，普通小 tensor 测试仍会走 uint32_t 分支，无法防止 FastDivMod<int64_t>、BroadcastConfig<uint64_t> 或尾块 grid-stride 逻辑回退。原问题是静默错误，缺测试后续很难发现。

建议修复方式：
补充一个 GPU 回归用例，显式构造会触发 launch_with_index_t(uint64_t{0}) 的 broadcast 场景并校验抽样结果；如果完整分配超过 2^31 元素成本过高，可以新增更底层的 BroadcastConfig<uint64_t> / kernel smoke test，至少覆盖大于 INT_MAX 的输出维度和 tail block。

done，后续会补测

fix bugs

b0efc76

fix bugs

feixi139 changed the title ~~[release/3.4][Operator Mechanism] fix Broadcast do not support big_tensor~~ [release/3.4][Operator Mechanism] fix BroadcastConfig do not support big_tensor Jul 8, 2026