[Bug fixes] support csa_sparse_attn_bwd_cudnn for h-card by danleifeng · Pull Request #1472 · PaddlePaddle/PaddleFleet

danleifeng · 2026-07-16T07:41:42Z

PR Category

Environment Adaptation

PR Types

Bug fixes

Description

support csa_sparse_attn_bwd_cudnn for h-card

是否引起精度变化

否

risemeup1111

已复查，未发现需要阻塞合入的问题。具体建议见 inline review comments。

^{Powered by Nyanpasu with gpt-5.5 xhigh, please check the suggestions carefully.}

risemeup1111 · 2026-07-16T08:07:42Z

+    import paddle
+
+    S = paddle.device.Stream
+    if not hasattr(S, "cuda_stream"):
+        S.cuda_stream = property(lambda self: self.stream_base.cuda_stream)


优先级：P2
这里把 cuda_stream 直接挂到整个 paddle.device.Stream 基类上了，但这个类本身也承载 XPU / custom device stream。现在的 fallback 少了 CUDA 类型检查，非 CUDA stream 只要访问这个属性就会走到 self.stream_base.cuda_stream，和 Paddle 现有实现的行为不一致。
处理要求：请针对该评论进行回复（同意并已修改请回复 Done，不同意请说明理由）。

Suggested change

import paddle

S = paddle.device.Stream

if not hasattr(S, "cuda_stream"):

S.cuda_stream = property(lambda self: self.stream_base.cuda_stream)

import paddle

from paddle.base.libpaddle import core

S = paddle.device.Stream

if not hasattr(S, "cuda_stream"):

def _cuda_stream(self):

assert isinstance(self.stream_base, core.CUDAStream), (

"cuda_stream is only available for CUDA streams"

)

return self.stream_base.cuda_stream

S.cuda_stream = property(_cuda_stream)

risemeup1111 · 2026-07-16T08:08:19Z

+
+        @contextmanager
+        def _nvtx_range(msg: str, *args, **kwargs):
+            nvtx_cls.range_push(msg)


优先级：P3
这个兼容层接受了 *args, **kwargs，但没有像 torch.cuda.nvtx.range / Paddle 现有实现那样先执行 msg.format(*args, **kwargs)。带占位符的 profiling 标签会丢失参数，后续排查 kernel range 时会看到未格式化的原始模板。
处理要求：请针对该评论进行回复（同意并已修改请回复 Done，不同意请说明理由）。

Suggested change

nvtx_cls.range_push(msg)

nvtx_cls.range_push(msg.format(*args, **kwargs))

Paddle-CI-Bot · 2026-07-16T08:48:49Z

PaddleFleet Log Analysis

Run #29569839443 · Attempt 1

日志分析报告

流水线名称	问题标签	修复建议	日志片段
Integration test (H20, multi-card) — GLM4.5 dpo	NCCL unhandled cuda error / MoE EP barrier 失败	检查 `fused_a2a.py` `barrier_ep` 中 DeepEP dispatch 触发的 NCCL cuda error，结合 nvidia-smi timeout 确认是否为 runner GPU 异常，建议 rerun 并配合 NCCL_DEBUG=INFO 复现	报错代码

失败的测试 case:

Integration test (H20, multi-card) / GLM4.5 dpo
  - glm45_dpo.yaml 8卡训练，rank 7 进程 exit code 1
  - sft_exit_code=241，未找到 '***** eval metrics *****'
  - OSError: NCCL error 'unhandled cuda error' @ process_group_nccl.cc:912
    调用栈: moe_layer.py:fusion_moe_forward → token_dispatcher.py:dispatch
            → fused_a2a.py:DeepEPDispatch.forward → barrier_ep → paddle.distributed.barrier

根本原因分析:

失败发生在 GLM4.5 DPO 8卡训练第一个训练步 MoE EP（Expert Parallelism）token dispatch 阶段。fused_a2a.py::fused_dispatch_forward_func 内调用 barrier_ep 做 EP group barrier 时触发 NCCL unhandled cuda error，同时可见 nvidia-smi 查询全部 8 张 GPU 状态也超时 3s（save gpu log failed: timed out），表明底层 GPU 驱动/通信状态在 DPO 任务运行过程中出现异常。本 PR 修改仅涉及 csa_sparse_attn_bwd_cudnn.py 和对应单测，DPO 任务使用的是 csa_sparse_attn_backend: tilelang（未切到 cudnn），因此该失败与本 PR 改动无直接关联，属于 runner GPU 异常导致的 NCCL 通信错误。

修复建议:

先 rerun 此 job，确认是否为偶发 GPU 异常（同 runner 多次失败则上报 CI 维护人员检查 paddle-10-1567435 机器 nvidia-container-runtime 状态）。
若需进一步定位，在 glm45_dpo.sh 中增加 NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=ALL 环境变量以获取详细 NCCL 错误信息。
本 PR 代码改动（csa_sparse_attn_bwd_cudnn）不需要修改；DPO 任务走 tilelang 后端，与 cudnn bwd 路径无交集。

🔍 准确性记录：请点击评论底部 😊 图标，选择 👍（准确）或 👎（有误），将自动记录到 CI 监控系统

_{🔄 每次 Re-run 后自动更新}

risemeup1111

上一轮两条兼容性建议已经按回复和新提交处理。复查新增 UT 后，我补充了一条测试收集阶段的非阻塞建议，细节见 inline comment；当前没有发现需要阻塞合入的代码问题。

^{Powered by Nyanpasu with gpt-5.5 xhigh, please check the suggestions carefully.}

risemeup1111 · 2026-07-16T11:49:54Z

+# Load module under test from local source tree
+_BWD_CUDNN_PATH = os.path.join(
+    os.path.dirname(__file__),
+    "../../../../src/paddlefleet/cudnn_ops/attn/csa_sparse_attn_bwd_cudnn.py",
+)
+_BWD_CUDNN_PATH = os.path.abspath(_BWD_CUDNN_PATH)
+
+_spec = importlib.util.spec_from_file_location("_bwd_cudnn_ut", _BWD_CUDNN_PATH)
+_bwd_mod = importlib.util.module_from_spec(_spec)
+_spec.loader.exec_module(_bwd_mod)
+
+_patch_cutlass_nvgpu = _bwd_mod._patch_cutlass_nvgpu
+_patch_paddle_nvtx_range = _bwd_mod._patch_paddle_nvtx_range
+_patch_paddle_stream_cuda_stream = _bwd_mod._patch_paddle_stream_cuda_stream


优先级：P2
这里把 _bwd_cudnn 模块在文件导入阶段直接 exec_module 了。这样一来，只要本地/CI 环境缺少 paddlefleet_ops、cudnn 或相关扩展，这个测试文件在 collection 阶段就会直接报错；它上面已经有按依赖可用性跳过的写法，这里绕开后会把原本可跳过的测试变成硬失败。
处理要求：请针对该评论进行回复（同意并已修改请回复 Done，不同意请说明理由）。

Suggested change

# Load module under test from local source tree

_BWD_CUDNN_PATH = os.path.join(

os.path.dirname(__file__),

"../../../../src/paddlefleet/cudnn_ops/attn/csa_sparse_attn_bwd_cudnn.py",

)

_BWD_CUDNN_PATH = os.path.abspath(_BWD_CUDNN_PATH)

_spec = importlib.util.spec_from_file_location("_bwd_cudnn_ut", _BWD_CUDNN_PATH)

_bwd_mod = importlib.util.module_from_spec(_spec)

_spec.loader.exec_module(_bwd_mod)

_patch_cutlass_nvgpu = _bwd_mod._patch_cutlass_nvgpu

_patch_paddle_nvtx_range = _bwd_mod._patch_paddle_nvtx_range

_patch_paddle_stream_cuda_stream = _bwd_mod._patch_paddle_stream_cuda_stream

try:

_spec = importlib.util.spec_from_file_location("_bwd_cudnn_ut", _BWD_CUDNN_PATH)

_bwd_mod = importlib.util.module_from_spec(_spec)

_spec.loader.exec_module(_bwd_mod)

except (ImportError, RuntimeError, AttributeError) as exc:

raise unittest.SkipTest(

f"csa_sparse_attn_bwd_cudnn import unavailable: {exc}"

)

_patch_cutlass_nvgpu = _bwd_mod._patch_cutlass_nvgpu

_patch_paddle_nvtx_range = _bwd_mod._patch_paddle_nvtx_range

_patch_paddle_stream_cuda_stream = _bwd_mod._patch_paddle_stream_cuda_stream

优先级：P1
我重新确认后，这里仍然会在测试收集阶段无条件加载 _bwd_cudnn 模块；当 paddlefleet_ops / cudnn / 相关扩展缺失时，pytest 会在 collection 直接失败，不能只靠 done 视为已解决。CI 的 review 也给出了同一结论。
处理要求：请针对该评论修复并提交新的 commit。
建议保持原来的跳过语义，把加载包在 try/except 里，失败时转成 unittest.SkipTest。

try: _spec = importlib.util.spec_from_file_location(...) _bwd_mod = importlib.util.module_from_spec(_spec) _spec.loader.exec_module(_bwd_mod) except (ImportError, ModuleNotFoundError, RuntimeError, AttributeError) as exc: raise unittest.SkipTest(...)

risemeup1111

已复查，上一轮两处兼容性问题已随新提交处理。当前测试文件在 collection 阶段无条件加载可选依赖的问题仍需要提交修复，细节见已有 inline review thread；本轮没有新增重复 inline comment。

^{Powered by Nyanpasu with gpt-5.5 xhigh, please check the suggestions carefully.}

PaddlePaddle-bot

🤖 Paddle-CI-Agent | pr_review | 2026-07-16 20:17:30

📋 Review 摘要

PR 概述：为 csa_sparse_attn_bwd_cudnn 增加 H-card 兼容 patch，并补充导入期 patch 单测。
变更范围：src/paddlefleet/cudnn_ops/attn/、tests/single_card_tests/ai_edited_test/fusions/
影响面 Tag：OP Tests

问题

级别	文件	概述
🟡 建议	`tests/single_card_tests/ai_edited_test/fusions/test_csa_sparse_attn_backends.py:210`	新增 patch 单测位于 `unittest.main()` 之后，直接执行测试文件时不会被收集
🟡 建议	`src/paddlefleet/cudnn_ops/attn/csa_sparse_attn_bwd_cudnn.py:137`	`cuda_stream` 运行期类型校验依赖 `assert`，优化模式下会失效

历史 Findings 修复情况

Finding	问题	状态
F1	测试收集阶段直接导入依赖可选/构建产物的 `csa_sparse_attn_bwd_cudnn.py`	✅ 已修复

📝 PR 规范检查

符合规范；标题已包含官方 [Bug fixes] Tag，描述包含 checklist §D2 要求的必填章节。

总体评价

H-card 兼容 patch 的方向与 backward 入口调用链匹配，但新增测试的位置会削弱直接运行时的覆盖；兼容属性里的运行期 guard 也建议改为显式异常，避免优化模式下行为漂移。

PaddlePaddle-bot · 2026-07-16T12:19:29Z

+)
+_BWD_CUDNN_PATH = os.path.abspath(_BWD_CUDNN_PATH)
+
+try:


🟡 建议 新增的 patch 单测定义在现有 unittest.main() 之后，直接执行该文件时不会被收集。

这个文件保留了 if __name__ == "__main__": unittest.main()，Python 直接运行到这里就会启动测试收集；本 PR 新增的 _BWD_CUDNN_PATH 加载逻辑和后续 TestMemoryFormatPatch / TestPatchPaddle* 类都还没有定义，因此 python tests/.../test_csa_sparse_attn_backends.py 只会跑旧的 backend 对齐用例，新增 patch 回归测试被漏掉。

建议修复方式：把整个 if __name__ == "__main__": unittest.main() 块移动到文件末尾，确保所有测试类定义完成后再进入 unittest。

PaddlePaddle-bot · 2026-07-16T12:19:29Z

+    if not hasattr(S, "cuda_stream"):
+
+        def _cuda_stream(self):
+            assert isinstance(self.stream_base, core.CUDAStream), (


🟡 建议 这里用 assert 做运行期 stream 类型校验，python -O 下会被移除。

cuda_stream 是导入模块时注入到 paddle.device.Stream 的兼容属性，后续 cuDNN resolve_stream() 会按普通属性访问它；如果遇到非 CUDA stream，当前 guard 在优化模式下不会执行，代码会继续读 self.stream_base.cuda_stream，错误类型和信息都不再可控。测试里也把 AssertionError 固化为期望行为。

建议修复方式：改成显式异常，并同步测试期望，例如：

if not isinstance(self.stream_base, core.CUDAStream): raise RuntimeError("cuda_stream is only available for CUDA streams")

risemeup1111

已复查，上一轮测试收集阶段直接加载可选依赖的问题已经通过后续提交修复；当前没有发现需要阻塞合入的代码问题。仍在运行的 CI 以最终结果为准。

^{Powered by Nyanpasu with gpt-5.5 xhigh, please check the suggestions carefully.}

codecov-commenter · 2026-07-17T11:30:16Z

Codecov Report

✅ All modified and coverable lines are covered by tests.
⚠️ Please upload report for BASE (develop@3415e86). Learn more about missing BASE report.

Additional details and impacted files

@@             Coverage Diff             @@
##             develop     #1472   +/-   ##
===========================================
  Coverage           ?   100.00%           
===========================================
  Files              ?         1           
  Lines              ?        52           
  Branches           ?         6           
===========================================
  Hits               ?        52           
  Misses             ?         0           
  Partials           ?         0

Flag	Coverage Δ
coverage_combine	`100.00% <100.00%> (?)`

Flags with carried forward coverage won't be shown. Click here to find out more.

Files with missing lines	Coverage Δ
...efleet/cudnn_ops/attn/csa_sparse_attn_bwd_cudnn.py	`100.00% <100.00%> (ø)`

🚀 New features to boost your workflow:

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

risemeup1111 · 2026-07-20T12:51:03Z

✅ Cherry-pick successful! Created PR: #1486

risemeup1111 · 2026-07-20T12:51:14Z

✅ Cherry-pick successful! Created PR: #1487

support csa_sparse_attn_bwd_cudnn for h-card

917829d

This comment was marked as outdated.

Sign in to view

danleifeng changed the title ~~support csa_sparse_attn_bwd_cudnn for h-card~~ [Bug fixes] support csa_sparse_attn_bwd_cudnn for h-card Jul 16, 2026

risemeup1111 reviewed Jul 16, 2026

View reviewed changes

add ut

f505a48

risemeup1111 reviewed Jul 16, 2026

View reviewed changes

This comment was marked as outdated.

Sign in to view

danleifeng added 2 commits July 16, 2026 19:56

add ut

de159b9

add ut

7c05c66

risemeup1111 suggested changes Jul 16, 2026

View reviewed changes

PaddlePaddle-bot reviewed Jul 16, 2026

View reviewed changes

danleifeng added 2 commits July 17, 2026 16:23

add ut

703b254

add ut

38b3fd6

risemeup1111 approved these changes Jul 17, 2026

View reviewed changes

zhangbo9674 approved these changes Jul 20, 2026

View reviewed changes

danleifeng merged commit 063870b into PaddlePaddle:develop Jul 20, 2026
26 of 28 checks passed

zhangbo9674 added cherry-pick: release/0.3 cherry-pick: release/0.4 labels Jul 20, 2026

risemeup1111 mentioned this pull request Jul 20, 2026

[release/0.3][Bug fixes] support csa_sparse_attn_bwd_cudnn for h-card #1486

Merged

risemeup1111 mentioned this pull request Jul 20, 2026

[release/0.4][Bug fixes] support csa_sparse_attn_bwd_cudnn for h-card #1487

Merged

github-actions Bot removed cherry-pick: release/0.3 cherry-pick: release/0.4 labels Jul 20, 2026

	nvtx_cls.range_push(msg)
	nvtx_cls.range_push(msg.format(args, *kwargs))

Uh oh!

Conversation

danleifeng commented Jul 16, 2026

PR Category

PR Types

Description

是否引起精度变化

Uh oh!

This comment was marked as outdated.

Uh oh!

risemeup1111 left a comment

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!

Paddle-CI-Bot commented Jul 16, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

PaddleFleet Log Analysis

日志分析报告

Uh oh!

risemeup1111 left a comment

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!

risemeup1111 left a comment

Choose a reason for hiding this comment

Uh oh!

PaddlePaddle-bot left a comment

Choose a reason for hiding this comment

📋 Review 摘要

问题

历史 Findings 修复情况

📝 PR 规范检查

总体评价

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

risemeup1111 left a comment

Choose a reason for hiding this comment

Uh oh!

codecov-commenter commented Jul 17, 2026

Codecov Report

Uh oh!

Uh oh!

risemeup1111 commented Jul 20, 2026

Uh oh!

risemeup1111 commented Jul 20, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

Paddle-CI-Bot commented Jul 16, 2026 •

edited

Loading