已合并
修复MoeInitRoutingQuantV2第三个输出结果精度问题 #8354
wuxiyuan创建于 7月7日
修复MoeInitRoutingQuantV2第三个输出结果精度问题 #8354
已合并
Pull Request已成功合入, 合并人@CANN-robot
（感谢 wuxiyuan 的贡献）
atomgit-bot
7月7日 评论：
7月7日 评论：
变更摘要
该 PR 修复了 MoeInitRoutingQuantV2 算子中第三个输出结果的精度问题。问题根因在于 CumsumComputeSimt SIMT kernel 内通过 UBUF (lastExpertIdCunsumAddr_) 跨线程传递最后一个元素的 expertId 和 cumsum 索引时存在竞争条件，导致读取到的值可能不正确。修复方案是移除 UBUF 中转机制，改为在 Process 方法中直接通过 GlobalTensor (expandedExpertIdxGmGT_) 从 GM 读取最后一个元素的 expertId，并以 totalLength_ 直接作为 cumsum 计数，彻底消除了线程间同步带来的精度隐患。
主要改动
- 移除 UBUF 缓冲区和相关成员变量：删除了
expertIdCunsumTBuf_（VECCALC 缓冲区）、lastExpertIdCunsum_（LocalTensor）和lastExpertIdCunsumAddr_（UBUF 指针），不再使用 UBUF 作为跨线程数据传递的中介。 - 新增
expandedExpertIdxGmGT_成员：引入GlobalTensor<int32_t>类型的expandedExpertIdxGmGT_，在Init中通过SetGlobalBuffer将其指向与expandedExpertIdxGm_相同的 workspace 内存区域，便于后续直接从 GM 读取数据。 - 简化
CumsumComputeSimt函数签名与逻辑：移除lastExpertIdCunsumAddr参数，并删除函数末尾向 UBUF 写入curExpertId和curIndex + 1的代码（原先由threadIdx.y == 0的线程执行），消除了潜在的多线程写入竞争。 - 重写
Process末尾的 expertId/cunsumNum 获取方式：将原来从lastExpertIdCunsumAddr_读取改为通过expandedExpertIdxGmGT_.GetValue(this->totalLength_ - 1)获取 expertId，cunsumNum直接取static_cast<int32_t>(this->totalLength_)，确保数据来源可靠。
7月7日 添加了label：cann-cla/yes
CANN-robot
7月7日 评论：
7月7日 评论：
CANN-robot
7月7日 评论：
7月7日 评论：
Thanks for your pull-request.
The full list of commands accepted by me can be found at here。
You can get sig-info at here
PR Approval Progress
✅ Congratulations! All modules have met the lgtm and approve requirements.
Module Approval Details
| module | lgtm status | approve status |
|---|---|---|
| ** | ✅ 蒋礼锐, gcw_kUomxQ2l (2/2) | ✅ gcw_kUomxQ2l (1/1) |
💡 Tip:
- Committer can comment
/approveor/lgtm- Commenting
/approveimplies both code review (lgtm) and intent to merge (approve)
7月7日 将monologue815,jiang-lirui,Allan_Yu,liudan12,mabing1118,gcw_kUomxQ2l,Andy-lb设为评审人
7月7日 将gcw_kUomxQ2l,Andy-lb设为审查人
7月7日 添加了label：ci-pipeline-running
CANN-robot
7月7日 评论：
7月7日 评论：
流水线任务触发成功
任务链接 [424176c7bf98467cbc13c30b9f18131a][流水线指导]
| 任务名称 | 状态 | 日志 | 下载链接 |
|---|---|---|---|
| Compile_Ascend_X86 | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_Ascend_ARM | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_Ascend_X86_mobile_station | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_Ascend_X86_experimental | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_Ascend_ARM_experimental | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_Pre | ✅ SUCCESS | >>>>> | |
| pre_comment | ✅ SUCCESS | >>>>> | |
| Compile_X86_monitor_910b | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_X86_monitor_910c | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_X86_monitor_950 | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_Ascend_X86_experimental_950 | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_Ascend_ARM_experimental_950 | ✅ SUCCESS | >>>>> | >>>>> |
| Compile_harmony-infer-chs-transformer | ✅ SUCCESS | >>>>> | |
| UT_Test | ✅ SUCCESS | ||
| PreSmoke_A900 | ✅ SUCCESS | >>>>> | |
| PreSmoke_ATK_Test_A2 | ✅ SUCCESS | >>>>> |
[2026-07-07 10:19:45] CI执行结束
CANN-robot
7月7日 评论：
7月7日 评论：
流水线任务触发成功
任务链接 [18fd3a4071d24e13bcdf930b7272fcfa][流水线指导]
| 任务名称 | 状态 | 日志 | 下载链接 |
|---|---|---|---|
| codecheck | ✅ SUCCESS | >>>>> | |
| SCA | ✅ SUCCESS | >>>>> | |
| antipoison | ✅ SUCCESS | >>>>> | |
| codecheck_style | ✅ SUCCESS | >>>>> | |
| StaticCheck_codespell_check | ✅ SUCCESS | ||
| StaticCheck_link_validity_check | ✅ SUCCESS | ||
| StaticCheck_resource_existence_check | ✅ SUCCESS | ||
| StaticCheck_tag_closed_check | ✅ SUCCESS | ||
| StaticCheck_markdownlint | ✅ SUCCESS | ||
| codecheck_checkpr | ✅ SUCCESS | ||
| codecheck_precommit | ⚠️ WARNING | >>>>> | >>>>> |
[2026-07-07 10:04:55] CI执行结束
7月7日 删除了label：ci-pipeline-running
7月7日 添加了label：ci-pipeline-passed
liulun0308
7月7日 评论：
7月7日 评论：
/approve
7月7日 添加了label：approved
jiang-lirui
7月7日 评论：
7月7日 评论：
/lgtm
7月7日 添加了label：lgtm
7月7日 合入了pull request
描述
在 CumsumComputeSimt 中，当处理到最后一个元素（curIndex == totalLength - 1）时，原代码直接 return 跳过了该元素的 cumsum 写入，仅通过 lastExpertIdCunsumAddr 缓存最后一条 expert 的 ID 和 cumsum 值到 ubuf。随后在 Process() 中从 ubuf 读取该缓存值来填充尾部 expert 的 cumsum 输出。
这种依赖 ubuf 跨线程传递缓存值的方式存在可靠性问题，导致第三个输出结果精度不正确。
具体变更
关联的Issue
https://gitcode.com/cann/ops-transformer/issues/3598
测试
已完成自验证 + 二级冒烟 + ST验证
文档更新
无
类型标签