Open
feat(loss): add cross_entropy_forward for Ascend 950 #63
feat(loss): add cross_entropy_forward for Ascend 950 #63
Open
yangqiangAdd on 5 days ago
yangqiang
5 days ago

改动

新增loss算子类别及cross_entropy_forward公共接口,提供硬标签交叉熵前向计算。

  • 支持FP32、FP16、BF16输入,int64标签,FP32输出。
  • 支持none、sum、mean、ignore_index、空输入、全忽略及非重叠行间距。
  • 接口负责校验、分配和选路,内核使用Ascend SIMD与SIMT。
  • 提供torch.ops.cann_ops_tilelang.cross_entropy_forward自定义入口及fake实现,可在torch.no_grad()下使用torch.compile进行无梯度前向;不提供backward。
  • 提供FP64参考、分级精度与路径测试、72档完整词表benchmark及接口文档docs/cross_entropy.md。

验证

在Ascend 950、CANN 9.3.0 weekly 20261021、PyTorch与torch_npu 2.10.0、TileLang提交f2b8866e上验证。

  • 全仓ruff检查与格式检查通过。
  • 本地CE回归727项通过,含41项torch接入、fake与编译回放测试。
  • 本次提交CI结果以检查页为准。

性能证据

完整词表V=129280,共72档设备耗时、带宽利用率图、完整数值表及CSV见性能数据与图表。测试环境、逻辑流量、实测搬运参考及统计方式均在资料页中说明。

性能基线JSONL和显存画像暂不随本PR提交。有效标签范围和有限输入由调用者保证,输出由算子独立分配。

likedislike
Merge blocked
Yyangqiang
5 days ago Created pull request, commit bde7d375
CANN-robot
CANN-robotCollaborators
5 days ago Comments:

Hi @gcw_wOhhD7pg, welcome to submitting your first PR to CANN!

PR Merge Steps

1. CLA Signing

If the current PR label includes cann-cla/yes, it means you have signed the CLA and can proceed to the next step. If the label includes cann-cla/no, please sign the CLA first. If you have any questions, please refer to the FAQ.

2. CI Check

Please comment /compile to trigger the CI pipeline check. If the CI run is successful, the PR will be tagged with ci-pipeline-passed and you can proceed to the next step. If the CI run fails, the PR will be tagged with ci-pipeline-failed, please check the CI logs to fix the issues in the PR. If you have any questions, please refer to the FAQ.

3. Code Review

After CI passes, please refer to the PR Approval Progress and proactively @ the committers in the table to review the code. After approval, committers will comment /lgtm and /approve. Once the lgtm and approved labels are successfully added, the PR will be merged automatically.

likedislike
CANN-robotCANN-robotCollaborators
5 days ago Added label: cann-cla/no
CANN-robotCANN-robotCollaborators
5 days ago Added label: stat/needs-squash
CANN-robot
CANN-robotCollaborators
5 days ago Comments:

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.
You can self-configure the PR merge rules for this repository. For more details, please refer to Here.
For more, you also can visit HICANN.


PR Approval Progress

⚠️ This PR does not yet meet the following requirements:lgtm (requires ≥ 2 person(s) per module)、approve (requires ≥ 1 person(s) per module)

Module Approval Details

module lgtm status approve status
repo-cann/ops-tilelang ❌ (0/2)(You can also ask: 俞郑中, 黄俊健, 刘丹, kevin365, 马兵) ❌ (0/1)(You can also ask: kevin365, 蒋礼锐, 刘丹, 黄俊健, ChaoyangJi)

💡 Tip:

  • Committer can comment /approve or /lgtm
  • Commenting /approve implies both code review (lgtm) and intent to merge (approve)

CLA Signature Pass

gcw_wOhhD7pg, thanks for your pull request. All authors of the commits have signed the CLA. 👍

likedislike
283 messages are collapsed here。 See more
CANN-robotCANN-robotCollaborators
13 hours ago Removed label: ci-pipeline-running
CANN-robotCANN-robotCollaborators
13 hours ago Added label: ci-pipeline-passed
Yyangqiang
12 hours ago Draft state has been cancelled
Yyangqiang
12 hours ago Change title to "feat(loss): add cross_entropy_forward for Ascend 950", originally titled "[WIP] feat(loss): add cross_entropy_forward for Ascend 950".
gong-zhihaoCollaborators12 hours agoStart code review1
docs/cross_entropy.md
@@ -0,0 +1,0 @@
1+# Cross entropy 前向
gong-zhihao12 hours agoComments:

md 文档改一下里面的描述吧,看着ai味道有点重

likedislike