[Operator Mechanism] Fix big tensor int64 to int32 narrowing bugs by feixi139 · Pull Request #79246 · PaddlePaddle/Paddle

feixi139 · 2026-06-04T11:29:02Z

PR Category

Operator Mechanism

PR Types

Bug fixes

Description

本 PR 修复了在扫描 Paddle 源码时发现的多类 int64_t / size_t / long / uint64_t 到 int / int32_t / uint32_t 的真实截断风险，重点解决大张量场景下可能出现的int64 -> int32 溢出、截断，以及 int 乘法在提升到 int64_t 前已经溢出的问题。

扫描和处理目录

本次主要扫描并处理了以下 Paddle 源码目录：

paddle/common
paddle/fluid/pybind
paddle/phi/backends
paddle/phi/core
paddle/phi/kernels/cpu
paddle/phi/kernels/funcs
paddle/phi/kernels/fusion/gpu
paddle/phi/kernels/gpu
paddle/phi/kernels/gpudnn
paddle/phi/kernels/impl
paddle/phi/kernels/primitive

公共检查能力

本 PR 在 paddle/common/enforce.h 中新增了 PADDLE_ENFORCE_LE_UINT32_MAX，用于处理 uint64_t / size_t / int64_t 等宽类型进入 uint32_t 边界前的上界检查，典型场景包括 CUDA dim3 / kernel launch 维度。

PADDLE_ENFORCE_LE_INT_MAX 用于 int / int32_t 边界
PADDLE_ENFORCE_LE_UINT32_MAX 用于 uint32_t 边界
两类宏分别对应 signed int 边界和 CUDA launch unsigned dimension 边界，避免混用

本 PR 主要处理了以下 case：

CUDA launch 维度截断
- 处理 grid / block / dim3 等 launch dimension 中由 tensor size、shape product、numel 等计算得到的值
- 使用 PADDLE_ENFORCE_LE_UINT32_MAX 做上界检查
- 将 launch dimension 显式转换为 uint32_t
- 保持 launch 调用处变量名不变，避免引入额外重构
下游 int API 边界
- 处理 BLAS、cuBLAS、cuDNN、CUB、JIT、自定义 kernel 等只接受 int 参数的边界
- 使用 PADDLE_ENFORCE_LE_INT_MAX 做上界检查
- 再通过 static_cast<int> 显式转换
tensor size 相关截断
- 处理 tensor dim、shape dim、numel、stride、offset、shape product 等来源的窄化转换
- 对没有必要转为 int 的内部计算，尽量保持为 int64_t
写入 INT32 tensor / int 参数结构
- 处理 size_t / int64_t 写入 int* buffer、INT32 tensor 或 int 参数结构前的截断风险
- 在写入前增加上界检查并显式 cast
int 乘法溢出
- 处理两个 int / int32_t 相乘后再赋值给 int64_t 的场景
- 避免乘法先在 32-bit 范围内溢出，再被提升为 int64_t

本 PR 跳过了以下不属于真实大张量截断风险或语义边界不清晰的 case：

axis / rank / ndim / device id 等 metadata-only 转换
operand / result / attribute / block arg 数量等 IR 元信息
已经被真实上界覆盖的派生值，例如由 total_numel < INT_MAX 保护的 offset 或 shard length
HOSTDEVICE 路径中不适合直接加入普通 runtime enforce 的位置
variadic forwarding wrapper 等无法可靠判断语义边界的位置

是否引起精度变化

否

agent check step 2

PaddlePaddle-bot · 2026-06-04T20:41:40Z

🤖 Paddle-CI-Agent | ci_status_monitor | 2026-06-10 20:03:48 UTC+08:00

CI报告基于以下代码生成（30分钟更新一次）:
PR commit: 948c677 | Merge base: 82c1f90 (branch: develop)

1 Required任务 : 43/48 通过

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

任务	错误类型	置信度	日志
`Fleet Unit test (single card)`	PR问题	中	Job
`Check approval`	需要 Approval	高	Job

2 失败详情

🔴 Fleet Unit test (single card) — PR问题（置信度: 中）

分析器: 通用分析(fallback)

失败用例:

用例	错误摘要
`tests/single_card_tests/custom_ops/test_sparse_mqa_flash_mla.py`, `test_tilelang_csa_indexer.py`, `test_tilelang_csa_indexer_cp.py`	单卡 pytest 进程 Aborted(core dumped)，CI 退出 1
`tests/single_card_tests/transformer/test_dsv4_hybrid_attention.py`	同一 Job 汇总为失败，compact 日志未给出完整 traceback

关键日志:

ci/single_card_test.sh: line 43: 1690 Aborted (core dumped) pytest "$test_file"
Test FAILED: tests/single_card_tests/custom_ops/test_sparse_mqa_flash_mla.py, see log for details...
Some single card tests failed:
- tests/single_card_tests/custom_ops/test_sparse_mqa_flash_mla.py
- tests/single_card_tests/custom_ops/test_tilelang_csa_indexer.py
- tests/single_card_tests/custom_ops/test_tilelang_csa_indexer_cp.py
- tests/single_card_tests/transformer/test_dsv4_hybrid_attention.py
Process completed with exit code 1.

根因摘要: attention/custom op 单卡用例在 PR wheel 上 abort
PR 修改范围包含 flash_attn_utils.h、masked_multihead_attention_kernel.cu、qkv_unpack_mha_kernel.cu、fused_rope* / fused_partial_rope* 等 GPU attention/rope 路径，并在这些路径新增 int/uint32 上界检查和 cast。失败用例正覆盖 sparse MQA/MLA、TileLang CSA indexer、DSV4 hybrid attention，和上述功能路径相关；该 Job 已重跑 5 次仍失败，暂不符合“重跑后通过”的 flaky 证据。当前 compact 日志没有完整 C++ traceback，因此无法锁定到具体某个 kernel 行，置信度标为中。

修复建议:

在 PaddleFleet 单卡环境复现 tests/single_card_tests/custom_ops/test_sparse_mqa_flash_mla.py，优先检查 flash_attn_utils.h:204 的 max_seqlen_* 上界检查、qkv_unpack_mha_kernel.cu:357 的 grid cast、masked_multihead_attention_kernel.cu:740 的 shared memory cast 是否改变该用例运行路径。
对 test_tilelang_csa_indexer* 和 test_dsv4_hybrid_attention.py 分别补充完整 traceback 后，再判断是否由同一 attention/rope kernel 修改触发；如果是合法大尺寸输入，应避免在非 int 下游边界前过早抛错。

关联变更: paddle/phi/kernels/gpu/flash_attn_utils.h:204, paddle/phi/kernels/fusion/gpu/qkv_unpack_mha_kernel.cu:357, paddle/phi/kernels/fusion/gpu/masked_multihead_attention_kernel.cu:740, paddle/phi/kernels/fusion/gpu/fused_partial_rope_kernel.cu:112, paddle/phi/kernels/fusion/gpu/fused_rope_grad_kernel.cu:70

🔴 Check approval — 需要 Approval（置信度: 高）

该 Job 需要人工 Approval，完成审批后 CI 才会继续执行。

根因摘要: 需要人工审批
修复建议: 请通过人工审批
关联变更: 无

wanghuancoder · 2026-06-05T07:31:14Z

        PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize());
      }
      if (FLAGS_alloc_fill_value >= 0) {
+        PADDLE_ENFORCE_LE_INT_MAX(FLAGS_alloc_fill_value, "memset fill value");


这个可以约束FLAGS_alloc_fill_value的取值范围是0~255

wanghuancoder · 2026-06-05T07:41:22Z

  if constexpr (dispatch == GradDispatchTag::MulInputGrad) {
    COMPUTE_OFFSET_SINGLE_OUTPUT(replace_index, 1, tid, 2)
-    atomicMax(aux_buffer + replace_index, tid);
+    atomicMax(aux_buffer + replace_index, static_cast<int>(tid));


这个再斟酌斟酌有风险吗？

在gpu侧无法使用PADDLE_ENFORCE_LE_INT_MAX，已经在对应的host侧添加了对应的PADDLE_ENFORCE_LE_INT_MAX限制

codecov-commenter · 2026-06-08T07:32:42Z

Codecov Report

❌ Patch coverage is 97.50000% with 1 line in your changes missing coverage. Please review.
⚠️ Please upload report for BASE (develop@d5cd4b9). Learn more about missing BASE report.

Files with missing lines	Patch %	Lines
paddle/phi/kernels/funcs/gather_scatter_functor.cc	85.71%	1 Missing ⚠️

Additional details and impacted files

@@            Coverage Diff             @@
##             develop   #79246   +/-   ##
==========================================
  Coverage           ?   97.50%           
==========================================
  Files              ?        6           
  Lines              ?       40           
  Branches           ?        0           
==========================================
  Hits               ?       39           
  Misses             ?        1           
  Partials           ?        0

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

A-nnonymous · 2026-06-11T03:25:37Z

-  uint32_t main_offset = (numel / (read_lens * threads)) * read_lens * threads;
-  uint32_t tail_tid = numel % (read_lens * threads);
+  const int64_t block_len = static_cast<int64_t>(read_lens) * threads;
+  const int64_t main_offset_64 = (numel_64 / block_len) * block_len;


block_len（read_lens）有没有取0的可能性？全链路检查一下

已进行全链路检查，paddle/phi/kernels/funcs/broadcast_function.h:839 已经在outs[0]->numel() == 0 时直接 return，因此不会继续进入 LaunchBroadcastKernel。

A-nnonymous · 2026-06-11T03:30:24Z

      constexpr int _block = 512;
-      int64_t grid = (self_size + _block - 1) / _block;
+      const uint32_t grid = GetCudaLaunchGrid(
+          self_size, _block, "CastDivKernel CUDA launch grid size");


我理解这里的操作是一个对输入空间的一个收窄，之前的超uint32的grid是否就完全无法支持了？

CUDA 中 dim3 的 x/y/z 成员类型是 unsigned int。因此超过 uint32_t 范围的 grid 本身就无法被正确表达和支持。旧代码虽然用 int64_t grid 计算，但传入 kernel launch 时仍会隐式转换到 dim3，存在截断/非法 launch 风险；本次修改只是将这种隐式风险改为通过 PADDLE_ENFORCE 显式报错，并不是收窄了原本可正确支持的输入空间。

A-nnonymous · 2026-06-11T03:30:53Z

-  int64_t ndim = index_dims.size();
  int64_t n = inner_dim_size * select_dim_size * outer_dim_size;
-  int64_t grid = (n + block - 1) / block;
+  const uint32_t grid = GetCudaLaunchGrid(


同“输入空间收窄”

CUDA 中 dim3 的 x/y/z 成员类型是 unsigned int。因此超过 uint32_t 范围的 grid 本身就无法被正确表达和支持。旧代码虽然用 int64_t grid 计算，但传入 kernel launch 时仍会隐式转换到 dim3，存在截断/非法 launch 风险；本次修改只是将这种隐式风险改为通过 PADDLE_ENFORCE 显式报错，并不是收窄了原本可正确支持的输入空间。

A-nnonymous · 2026-06-11T03:34:36Z

  const int threads = 256;
  int64_t x_numel = x->numel();
  int64_t fsize = H * W * D;
+  PADDLE_ENFORCE_LE_INT_MAX(fsize, "sync_batch_norm grad feature size");


假设过强了吧？HWD不是很容易超过uint32吗？是否这种大shape之后就不支持了？

进一步审查发现，在完成enforce检查后，在代码中还会调用int32与int64两种数据类型。因此已将int32调用处的相关函数参数类型提升至int64

A-nnonymous · 2026-06-11T03:40:58Z

  auto dims = weight.dims();
  const int rank = dims.size();
-  std::vector<int> real_dims;
+  std::vector<int64_t> real_dims;


这里real dims里面push了autotype的普通dims，安全吗

dims 的类型是 phi::DDim，dims[dim] 返回的是维度值，类型是 int64_t 。此处修复了原本的int64到int32的截断问题

A-nnonymous · 2026-06-12T06:08:58Z

+    const int64_t axis_dim) {
  int64_t batch_size = prob->dims()[0];
  int64_t class_num = prob->dims()[1];
+  PADDLE_ENFORCE_EQ(


这里是新加的约束条件？我理解这是对输入的约束增强，是必须加入的吗？请解释一下。如果两变量一定相等的话，就说明语义存在冗余，变量数应该进行缩减

在Infermeta处会对于不想等的情况进行拦截，因此此处的判定并不是必须的正确性问题判定，已删除

PaddlePaddle-bot

🤖 Paddle-CI-Agent | pr_review | 2026-06-12 16:10:37

📋 Review 摘要

PR 概述：修复多处 Paddle C++/CUDA kernel 中 64 位计数向 32 位参数传递前的截断风险。

变更范围：paddle/fluid/pybind/、paddle/phi/backends/、paddle/phi/core/、paddle/phi/kernels/

影响面 Tag：[Operator Mechanism] [Custom Device] [Performance Optimization]

问题

级别	文件	概述
🟡 建议	PR 级别	本 PR 变更量较大（124 文件 / 4249 行），建议按模块拆分以降低审查和回归风险。

本轮达到时间上限，未形成新的可阻塞发现；未覆盖部分待后续深挖。

历史 Findings 修复情况

Finding	问题	状态
F1	`sync_batch_norm_kernel.cu` 中 `fsize` 仍基于已收窄到 `int` 的 `H/W/D`	⚠️ 仍存在

📝 PR 规范检查

PR 标题格式合规（[Operator Mechanism]），四个必填 section（PR Category / PR Types / Description / 是否引起精度变化）均已填写，符合模板结构要求。

总体评价

⚠️ 本 PR 变更量较大（124 文件 / 4249 行），建议拆分以降低审查难度和合入风险。

建议拆分方案：

PR 1: 公共边界检查与基础设施 — paddle/common/enforce.h、paddle/phi/backends/gpu/gpu_launch_config.h、paddle/phi/core/memory/allocation/allocator.h
PR 2: BLAS/cuDNN/OneDNN 参数边界 — paddle/phi/kernels/funcs/blas/*、paddle/phi/kernels/funcs/cudnn_rnn_cache.h、paddle/phi/kernels/gpudnn/*
PR 3: 通用 funcs 与 CPU kernel 计数修复 — paddle/phi/kernels/funcs/*、paddle/phi/kernels/cpu/*
PR 4: GPU/fusion kernel launch 与索引修复 — paddle/phi/kernels/gpu/*、paddle/phi/kernels/fusion/gpu/*
PR 5: pybind 与 impl 层窄化修复 — paddle/fluid/pybind/*、paddle/phi/kernels/impl/*

本轮按风险优先审查了公共宏可用性、GPU launch 维度、sync_batch_norm 历史问题、broadcast/conv/cuDNN RNN/BLASLt 等关键语义变更点，未全量覆盖其余机械迁移文件。

A-nnonymous · 2026-06-16T02:39:13Z

LGTM，合入后建议等一段时间CICE反馈

wanghuancoder

LGTM

paddle-bot · 2026-07-13T12:23:24Z

你的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.

Backport the SplitFunctor<CPUContext> int->int64_t narrowing fix from PaddlePaddle#79246 so unbind/split on CPU no longer overflows when a per-output column count exceeds INT_MAX (e.g. unbind([2, 2^31], axis=0) via the non-strided path). Pairs with PaddlePaddle#79737 which removed the numel<=INT_MAX guard.

fix int64 = int32

de2d332

agent check step 2

This comment was marked as outdated.

Sign in to view

paddle-bot Bot added the contributor External developers label Jun 4, 2026

fix bugs

2ea6462

This comment was marked as outdated.

Sign in to view

feixi139 added 2 commits June 5, 2026 05:18

fix bugs

5482ce0

fix bugs

65a9a1e

This comment was marked as outdated.

Sign in to view

fix bugs

5f11698

This comment was marked as outdated.

Sign in to view

fix bugs

db7dc4e

This comment was marked as outdated.

Sign in to view

wanghuancoder reviewed Jun 5, 2026

View reviewed changes

Address review comments

4794f02

This comment was marked as outdated.

Sign in to view

fix bugs

67f33cd

feixi139 added 2 commits June 8, 2026 07:49

fix bugs

37ca03a

fix bugs

9af173a

feixi139 force-pushed the fix_int_convert branch from 1efbb45 to 9af173a Compare June 8, 2026 13:04

feixi139 added 2 commits June 9, 2026 03:05

fix codestyle

f6a021f

revert for_range.h

948c677

A-nnonymous reviewed Jun 11, 2026

View reviewed changes

fix premature int narrowing in gpu kernels

30c28c3

This comment was marked as outdated.

Sign in to view

Merge branch 'develop' into fix_int_convert

ca8bc6c

This comment was marked as outdated.

Sign in to view

Update generate_proposals_kernel.cu

d5765d5

This comment was marked as outdated.

Sign in to view

A-nnonymous reviewed Jun 12, 2026

View reviewed changes

Update cross_entropy.cu

0e09fca

PaddlePaddle-bot reviewed Jun 12, 2026

View reviewed changes

A-nnonymous approved these changes Jun 16, 2026

View reviewed changes

wanghuancoder approved these changes Jun 16, 2026

View reviewed changes

wanghuancoder merged commit a8b4238 into PaddlePaddle:develop Jun 16, 2026
120 of 125 checks passed

Conversation

feixi139 commented Jun 4, 2026

PR Category

PR Types

Description

扫描和处理目录

公共检查能力

是否引起精度变化

Uh oh!

This comment was marked as outdated.

Uh oh!

PaddlePaddle-bot commented Jun 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

1 Required任务 : 43/48 通过

2 失败详情

Uh oh!

This comment was marked as outdated.

Uh oh!

This comment was marked as outdated.

Uh oh!

This comment was marked as outdated.

Uh oh!

This comment was marked as outdated.

Uh oh!

This comment was marked as outdated.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

feixi139 Jun 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment was marked as outdated.

Uh oh!

codecov-commenter commented Jun 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

This comment was marked as outdated.

Uh oh!

This comment was marked as outdated.

Uh oh!

This comment was marked as outdated.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

feixi139 Jun 12, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

PaddlePaddle-bot commented Jun 4, 2026 •

edited

Loading

feixi139 Jun 8, 2026 •

edited

Loading

codecov-commenter commented Jun 8, 2026 •

edited

Loading

feixi139 Jun 12, 2026 •

edited

Loading