Star876
1776
代码
代码
Issues227
Pull Requests319
Actions
项目讨论
Wiki
分析
Star876
1776
已合并
[CANNBot]SoftMarginLoss算子支持Ascend950 AscendC实现 #3729
wangweidong创建于 4月11日
[CANNBot]SoftMarginLoss算子支持Ascend950 AscendC实现 #3729
已合并
wangweidong创建于 4月11日
从gcw_YBIAEfqJ/ops-nngenerate0411合入到cann/ops-nnmaster
共 13 个文件变更+1520-0
Aexperimental/loss/soft_margin_loss/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
1+# ---------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aexperimental/loss/soft_margin_loss/README.md+95-0
@@ -0,0 +1,95 @@
1+# SoftMarginLoss
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| ---- | :----:|
7+|Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件|√|
8+ 
9+## 功能说明
10+ 
11+- 算子功能：计算二分类场景下的soft margin loss（逻辑回归损失），对输入预测值和目标值逐元素计算损失，并支持归约输出。
12+ 
13+- 计算公式：
14+ 给定预测张量$self$和目标张量$target$，逐元素计算损失：
15+ 
16+ $
17+ L_i = \max(0, -t_i \cdot x_i) + \log(1 + \exp(-|t_i \cdot x_i|))
18+ $
19+ 
20+ 其中$x_i = self_i$，$t_i = target_i$。
21+ 
22+ 归约模式：
23+ - $reduction = 0$（none）：$output = L$，输出与输入同shape
24+ - $reduction = 1$（mean）：$output = \frac{1}{N}\sum_{i}L_i$，输出标量
25+ - $reduction = 2$（sum）：$output = \sum_{i}L_i$，输出标量
26+ 
27+- 示例：
28+ 假设输入张量$self = [1, -1, 2, -2]$，目标张量$target = [1, 1, -1, -1]$，$reduction = 0$（none），那么输出张量$output = [0.3133, 1.3133, 2.1269, 0.1269]$，具体计算过程如下：
29+ $
30+ \begin{aligned}
31+ L_0 &= \max(0, -(1)(1)) + \log(1 + \exp(-|1|)) = 0 + \log(1 + 0.3679) = 0.3133 \\
32+ L_1 &= \max(0, -(1)(-1)) + \log(1 + \exp(-|-1|)) = 1 + \log(1 + 0.3679) = 1.3133 \\
33+ L_2 &= \max(0, -(-1)(2)) + \log(1 + \exp(-|-2|)) = 2 + \log(1 + 0.1353) = 2.1269 \\
34+ L_3 &= \max(0, -(-1)(-2)) + \log(1 + \exp(-|-2|)) = 0 + \log(1 + 0.1353) = 0.1269
35+ \end{aligned}
36+ $
37+ 
38+## 参数说明
39+ 
40+<table style="undefined;table-layout: fixed; width: 1250px"><colgroup>
41+ <col style="width: 60px">
42+ <col style="width: 60px">
43+ <col style="width: 150px">
44+ <col style="width: 150px">
45+ <col style="width: 60px">
46+ </colgroup>
47+ <thead>
48+ <tr>
49+ <th>参数名</th>
50+ <th>输入/输出/属性</th>
51+ <th>描述</th>
52+ <th>数据类型</th>
53+ <th>数据格式</th>
54+ </tr></thead>
55+ <tbody>
56+ <tr>
57+ <td>self</td>
58+ <td>输入</td>
59+ <td>预测值张量，公式中的x。</td>
60+ <td>FLOAT、FLOAT16</td>
61+ <td>ND</td>
62+ </tr>
63+ <tr>
64+ <td>target</td>
65+ <td>输入</td>
66+ <td>目标值张量，shape和dtype与self一致，公式中的t。</td>
67+ <td>FLOAT、FLOAT16</td>
68+ <td>ND</td>
69+ </tr>
70+ <tr>
71+ <td>output</td>
72+ <td>输出</td>
73+ <td>输出张量。reduction=none时与输入同shape；reduction=mean/sum时为标量（0维张量）。</td>
74+ <td>FLOAT、FLOAT16</td>
75+ <td>ND</td>
76+ </tr>
77+ <tr>
78+ <td>reduction</td>
79+ <td>属性</td>
80+ <td><ul><li>归约模式：0=none，1=mean，2=sum。</li><li>默认值为1（mean）。</li></ul></td>
81+ <td>Int</td>
82+ <td>-</td>
83+ </tr>
84+ </tbody></table>
85+ 
86+## 约束说明
87+ 
88+- self和target的shape和dtype必须一致。
89+- float16输入在Kernel内部提升至float32计算，仅最终输出转回float16。
90+ 
91+## 调用说明
92+ 
93+| 调用方式 | 调用样例 | 说明 |
94+|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
95+| aclnn调用 | [test_aclnn_soft_margin_loss.cpp](./examples/test_aclnn_soft_margin_loss.cpp) | 通过aclnn接口方式调用SoftMarginLoss算子。 |
Aexperimental/loss/soft_margin_loss/docs/aclnnSoftMarginLoss.md+178-0
@@ -0,0 +1,178 @@
1+# aclnnSoftMarginLoss
2+ 
3+## 支持的产品型号
4+ 
5+| 产品系列 | 产品型号 |
6+|---------|---------|
7+| Atlas A2 训练系列产品 | Atlas 800T A2、Atlas 800I A2、Atlas 900 A2 PoD、Atlas 200I A2 |
8+ 
9+## 功能描述
10+ 
11+计算输入张量 `self` 与目标张量 `target` 之间的 SoftMarginLoss，即逐元素计算：
12+ 
13+$$\text{loss}_i = \log\!\left(1 + e^{-\text{target}_i \times \text{self}_i}\right)$$
14+ 
15+并根据 `reduction` 参数对结果进行归约：
16+ 
17+- `reduction=0`（none）：不归约，输出 shape 与输入相同。
18+- `reduction=1`（mean）：对所有元素求均值，输出为标量。
19+- `reduction=2`（sum）：对所有元素求和，输出为标量。
20+ 
21+- `self` 与 `target` 须具有相同的 shape 和数据类型。
22+- 支持数据类型：float32、float16。
23+ 
24+## 函数原型
25+ 
26+```cpp
27+aclnnStatus aclnnSoftMarginLossGetWorkspaceSize(
28+ const aclTensor *self,
29+ const aclTensor *target,
30+ int64_t reduction,
31+ const aclTensor *out,
32+ uint64_t *workspaceSize,
33+ aclOpExecutor **executor);
34+ 
35+aclnnStatus aclnnSoftMarginLoss(
36+ void *workspace,
37+ uint64_t workspaceSize,
38+ aclOpExecutor *executor,
39+ aclrtStream stream);
40+```
41+ 
42+## aclnnSoftMarginLossGetWorkspaceSize
43+ 
44+### 参数说明
45+ 
46+| 参数名 | 输入/输出 | 描述 |
47+|-------|---------|------|
48+| self | 输入 | 预测值张量。数据类型：float32、float16。数据格式：ND。支持非连续 tensor。 |
49+| target | 输入 | 目标值张量（通常为 +1 或 -1）。数据类型须与 self 相同。数据格式：ND。shape 须与 self 相同。 |
50+| reduction | 输入 | 归约模式。int64_t 类型，取值：0（none）、1（mean，默认）、2（sum）。 |
51+| out | 输出 | 输出张量。数据类型与 self 相同。reduction=0 时 shape 与 self 相同；reduction=1 或 2 时为标量（0 维 tensor）。 |
52+| workspaceSize | 输出 | 算子执行所需 workspace 大小，单位为 Byte。由本函数返回，调用方须据此分配 workspace 内存。 |
53+| executor | 输出 | 算子执行器，包含算子计算流信息，由本函数返回后传入 aclnnSoftMarginLoss 执行。 |
54+ 
55+### 返回值说明
56+ 
57+返回 `aclnnStatus` 错误码，详见 [aclnn 错误码](#错误码)。
58+ 
59+## aclnnSoftMarginLoss
60+ 
61+### 参数说明
62+ 
63+| 参数名 | 输入/输出 | 描述 |
64+|-------|---------|------|
65+| workspace | 输入 | workspace 内存地址。若 workspaceSize 为 0，可传入 nullptr。 |
66+| workspaceSize | 输入 | workspace 大小，由 aclnnSoftMarginLossGetWorkspaceSize 返回。 |
67+| executor | 输入 | 算子执行器，由 aclnnSoftMarginLossGetWorkspaceSize 返回。 |
68+| stream | 输入 | ACL stream，用于异步调度算子执行。 |
69+ 
70+### 返回值说明
71+ 
72+返回 `aclnnStatus` 错误码，详见 [aclnn 错误码](#错误码)。
73+ 
74+## 错误码
75+ 
76+| 错误码 | 描述 |
77+|-------|------|
78+| ACLNN_SUCCESS（0） | 执行成功。 |
79+| ACLNN_ERR_PARAM_NULLPTR（161001） | 输入/输出 tensor 指针为空。 |
80+| ACLNN_ERR_PARAM_INVALID（161002） | 参数非法，包括：数据类型不支持、self 与 target 数据类型不一致、reduction 值不在 {0,1,2} 范围内、out shape 与预期不一致等。 |
81+| ACLNN_ERR_INNER_CREATE_EXECUTOR | 内部创建算子执行器失败。 |
82+| ACLNN_ERR_INNER_NULLPTR | 内部 tensor 分配失败。 |
83+| ACLNN_ERR_INNER_INFERSHAPE_ERROR | 内部 InferShape 失败。 |
84+ 
85+## 约束说明
86+ 
87+- `self` 与 `target` 须为相同数据类型和相同 shape。
88+- `reduction` 仅支持 0（none）、1（mean）、2（sum），其他值返回 `ACLNN_ERR_PARAM_INVALID`。
89+- `out` 的 shape 须与 reduction 模式匹配：reduction=0 时与 self 相同，reduction=1 或 2 时为 0 维标量（**注意：不是 shape=[1]，而是 dimNum=0 的标量 tensor**）。
90+- 支持空 tensor（元素数为 0），此时 workspaceSize 为 0，直接返回成功。
91+- workspace 须在调用 `aclnnSoftMarginLoss` 之前分配，在 stream 中算子执行完成后方可释放。
92+ 
93+## 调用示例
94+ 
95+以下示例展示了 SoftMarginLoss 算子的完整调用流程（reduction=mean）：
96+ 
97+```cpp
98+#include <iostream>
99+#include <vector>
100+#include <cmath>
101+#include "acl/acl.h"
102+#include "aclnn_soft_margin_loss.h"
103+ 
104+int main() {
105+ // 1. 初始化 ACL 及设备
106+ aclInit(nullptr);
107+ aclrtSetDevice(0);
108+ aclrtStream stream;
109+ aclrtCreateStream(&stream);
110+ 
111+ // 2. 准备输入数据（fp32，shape=[4, 8]）
112+ int64_t shape[] = {4, 8};
113+ int64_t strides[] = {8, 1};
114+ 
115+ float self_host[] = {
116+ 0.5f, 1.0f, -0.3f, 2.0f, 0.1f, -1.5f, 0.8f, -0.2f,
117+ -1.0f, 0.7f, 1.5f, -0.5f, 0.3f, 0.9f, -0.8f, 1.2f,
118+ 0.4f, -0.6f, 1.1f, -1.3f, 0.6f, -0.4f, 1.4f, 0.2f,
119+ -0.9f, 1.3f, -0.1f, 0.0f, 1.6f, -1.1f, 0.5f, -0.7f
120+ };
121+ float target_host[] = {
122+ 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f,
123+ -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f,
124+ 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f,
125+ 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f
126+ };
127+ size_t nbytes = 32 * sizeof(float);
128+ 
129+ void *self_dev = nullptr, *target_dev = nullptr, *out_dev = nullptr;
130+ aclrtMalloc(&self_dev, nbytes, ACL_MEM_MALLOC_NORMAL_ONLY);
131+ aclrtMalloc(&target_dev, nbytes, ACL_MEM_MALLOC_NORMAL_ONLY);
132+ aclrtMalloc(&out_dev, sizeof(float), ACL_MEM_MALLOC_NORMAL_ONLY);
133+ aclrtMemcpy(self_dev, nbytes, self_host, nbytes, ACL_MEMCPY_HOST_TO_DEVICE);
134+ aclrtMemcpy(target_dev, nbytes, target_host, nbytes, ACL_MEMCPY_HOST_TO_DEVICE);
135+ 
136+ // 3. 创建 aclTensor
137+ aclTensor *self_t = aclCreateTensor(shape, 2, ACL_FLOAT, strides, 0,
138+ ACL_FORMAT_ND, shape, 2, self_dev);
139+ aclTensor *target_t = aclCreateTensor(shape, 2, ACL_FLOAT, strides, 0,
140+ ACL_FORMAT_ND, shape, 2, target_dev);
141+ 
142+ // reduction=mean/sum 时 out 为 0 维标量 tensor
143+ aclTensor *out_t = aclCreateTensor(nullptr, 0, ACL_FLOAT, nullptr, 0,
144+ ACL_FORMAT_ND, nullptr, 0, out_dev);
145+ 
146+ // 4. 查询 workspace 大小并分配
147+ int64_t reduction = 1; // mean
148+ uint64_t workspaceSize = 0;
149+ aclOpExecutor *executor = nullptr;
150+ aclnnSoftMarginLossGetWorkspaceSize(self_t, target_t, reduction, out_t,
151+ &workspaceSize, &executor);
152+ 
153+ void *workspace = nullptr;
154+ if (workspaceSize > 0)
155+ aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_NORMAL_ONLY);
156+ 
157+ // 5. 执行算子
158+ aclnnSoftMarginLoss(workspace, workspaceSize, executor, stream);
159+ aclrtSynchronizeStream(stream);
160+ 
161+ // 6. 取回结果
162+ float out_host = 0.0f;
163+ aclrtMemcpy(&out_host, sizeof(float), out_dev, sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
164+ printf("SoftMarginLoss (mean) = %f\n", out_host);
165+ // 期望: 0.781925
166+ 
167+ // 7. 释放资源
168+ if (workspace) aclrtFree(workspace);
169+ aclrtFree(self_dev); aclrtFree(target_dev); aclrtFree(out_dev);
170+ aclDestroyTensor(self_t); aclDestroyTensor(target_t); aclDestroyTensor(out_t);
171+ aclrtDestroyStream(stream);
172+ aclrtResetDevice(0);
173+ aclFinalize();
174+ return 0;
175+}
176+```
177+ 
178+> **reduction=none 示例**：若 `reduction=0`，则 `out` 的 shape 须与 `self` 相同（例如 `[4, 8]`），输出为逐元素的 loss 值。
Aexperimental/loss/soft_margin_loss/examples/test_aclnn_soft_margin_loss.cpp+204-0
@@ -0,0 +1,204 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+#include <iostream>
17+#include <vector>
18+#include <cmath>
19+#include "acl/acl.h"
20+#include "aclnn_soft_margin_loss.h"
21+ 
22+#define CHECK_RET(cond, return_expr) \
23+ do { \
24+ if (!(cond)) { \
25+ return_expr; \
26+ } \
27+ } while (0)
28+ 
29+#define LOG_PRINT(message, ...) \
30+ do { \
31+ printf(message, ##__VA_ARGS__); \
32+ } while (0)
33+ 
34+int64_t GetShapeSize(const std::vector<int64_t>& shape)
35+{
36+ int64_t shapeSize = 1;
37+ for (auto i : shape) {
38+ shapeSize *= i;
39+ }
40+ return shapeSize;
41+}
42+ 
43+void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr, aclDataType dtype)
44+{
45+ auto size = GetShapeSize(shape);
46+ if (dtype == ACL_FLOAT) {
47+ std::vector<float> resultData(size, 0);
48+ aclrtMemcpy(resultData.data(), size * sizeof(float), *deviceAddr, size * sizeof(float),
49+ ACL_MEMCPY_DEVICE_TO_HOST);
50+ for (int64_t i = 0; i < size; i++) {
51+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
52+ }
53+ } else if (dtype == ACL_FLOAT16) {
54+ // fp16: 拷回 uint16_t 再转 float 显示
55+ std::vector<uint16_t> resultData(size, 0);
56+ aclrtMemcpy(resultData.data(), size * sizeof(uint16_t), *deviceAddr, size * sizeof(uint16_t),
57+ ACL_MEMCPY_DEVICE_TO_HOST);
58+ for (int64_t i = 0; i < size; i++) {
59+ LOG_PRINT("result[%ld] is: 0x%04x\n", i, resultData[i]);
60+ }
61+ }
62+}
63+ 
64+int Init(int32_t deviceId, aclrtStream* stream)
65+{
66+ auto ret = aclInit(nullptr);
67+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
68+ ret = aclrtSetDevice(deviceId);
69+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
70+ ret = aclrtCreateStream(stream);
71+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
72+ return 0;
73+}
74+ 
75+template <typename T>
76+int CreateAclTensor(
77+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
78+ aclTensor** tensor)
79+{
80+ auto elemCount = hostData.size();
81+ auto size = elemCount * sizeof(T);
82+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
83+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
84+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
85+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
86+ 
87+ std::vector<int64_t> strides(shape.size(), 1);
88+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; i--) {
89+ strides[i] = shape[i + 1] * strides[i + 1];
90+ }
91+ 
92+ *tensor = aclCreateTensor(
93+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
94+ *deviceAddr);
95+ return 0;
96+}
97+ 
98+int main()
99+{
100+ // 1. 初始化
101+ int32_t deviceId = 0;
102+ aclrtStream stream;
103+ auto ret = Init(deviceId, &stream);
104+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
105+ 
106+ // 2. 构造输入数据
107+ // SoftMarginLoss: loss = log(1 + exp(-target * self))
108+ // reduction: 0=none, 1=mean, 2=sum
109+ std::vector<int64_t> selfShape = {4, 8};
110+ int64_t totalNum = GetShapeSize(selfShape);
111+ 
112+ // self: 输入预测值
113+ std::vector<float> selfHostData = {
114+ 0.5f, 1.0f, -0.3f, 2.0f, 0.1f, -1.5f, 0.8f, -0.2f,
115+ -1.0f, 0.7f, 1.5f, -0.5f, 0.3f, 0.9f, -0.8f, 1.2f,
116+ 0.4f, -0.6f, 1.1f, -1.3f, 0.6f, -0.4f, 1.4f, 0.2f,
117+ -0.9f, 1.3f, -0.1f, 0.0f, 1.6f, -1.1f, 0.5f, -0.7f
118+ };
119+ 
120+ // target: 标签值（+1 或 -1）
121+ std::vector<float> targetHostData = {
122+ 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f,
123+ -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f,
124+ 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f,
125+ 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f
126+ };
127+ 
128+ // 3. 计算 CPU golden (reduction=mean)
129+ int64_t reduction = 1; // mean
130+ {
131+ double totalLoss = 0.0;
132+ LOG_PRINT("=== CPU Golden (reduction=mean) ===\n");
133+ for (int64_t i = 0; i < totalNum; i++) {
134+ double loss = std::log(1.0 + std::exp(-targetHostData[i] * selfHostData[i]));
135+ totalLoss += loss;
136+ }
137+ LOG_PRINT("expected mean loss: %f\n\n", totalLoss / totalNum);
138+ }
139+ 
140+ // 4. 创建 aclTensor
141+ aclTensor* selfTensor = nullptr;
142+ void* selfDeviceAddr = nullptr;
143+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, ACL_FLOAT, &selfTensor);
144+ CHECK_RET(ret == ACL_SUCCESS, return ret);
145+ 
146+ aclTensor* targetTensor = nullptr;
147+ void* targetDeviceAddr = nullptr;
148+ ret = CreateAclTensor(targetHostData, selfShape, &targetDeviceAddr, ACL_FLOAT, &targetTensor);
149+ CHECK_RET(ret == ACL_SUCCESS, return ret);
150+ 
151+ // output shape: reduction=none -> same as input; reduction=mean/sum -> scalar (0-dim)
152+ std::vector<int64_t> outShape = (reduction == 0) ? selfShape : std::vector<int64_t>{};
153+ int64_t outNum = (reduction == 0) ? totalNum : 1;
154+ std::vector<float> outHostData(outNum, 0.0f);
155+ 
156+ aclTensor* outTensor = nullptr;
157+ void* outDeviceAddr = nullptr;
158+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, ACL_FLOAT, &outTensor);
159+ CHECK_RET(ret == ACL_SUCCESS, return ret);
160+ 
161+ // 5. 调用 aclnnSoftMarginLoss
162+ uint64_t workspaceSize = 0;
163+ aclOpExecutor* executor;
164+ 
165+ ret = aclnnSoftMarginLossGetWorkspaceSize(selfTensor, targetTensor, reduction, outTensor,
166+ &workspaceSize, &executor);
167+ LOG_PRINT("aclnnSoftMarginLossGetWorkspaceSize returned %d, workspaceSize=%llu\n",
168+ ret, (unsigned long long)workspaceSize);
169+ CHECK_RET(ret == ACL_SUCCESS,
170+ LOG_PRINT("aclnnSoftMarginLossGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
171+ 
172+ void* workspaceAddr = nullptr;
173+ if (workspaceSize > 0) {
174+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
175+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
176+ }
177+ 
178+ ret = aclnnSoftMarginLoss(workspaceAddr, workspaceSize, executor, stream);
179+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSoftMarginLoss failed. ERROR: %d\n", ret); return ret);
180+ 
181+ ret = aclrtSynchronizeStream(stream);
182+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
183+ 
184+ // 6. 打印结果
185+ LOG_PRINT("=== NPU Output ===\n");
186+ PrintOutResult(outShape, &outDeviceAddr, ACL_FLOAT);
187+ 
188+ // 7. 释放资源
189+ aclDestroyTensor(selfTensor);
190+ aclDestroyTensor(targetTensor);
191+ aclDestroyTensor(outTensor);
192+ 
193+ aclrtFree(selfDeviceAddr);
194+ aclrtFree(targetDeviceAddr);
195+ aclrtFree(outDeviceAddr);
196+ if (workspaceSize > 0) {
197+ aclrtFree(workspaceAddr);
198+ }
199+ aclrtDestroyStream(stream);
200+ aclrtResetDevice(deviceId);
201+ aclFinalize();
202+ 
203+ return 0;
204+}
Aexperimental/loss/soft_margin_loss/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ---------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE soft_margin_loss ACLNNTYPE aclnn)
Aexperimental/loss/soft_margin_loss/op_host/soft_margin_loss_def.cpp+61-0
@@ -0,0 +1,61 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file soft_margin_loss_def.cpp
18+ * \brief SoftMarginLoss operator definition - declares inputs, outputs, attributes, and chip configuration
19+ */
20+#include "register/op_def_registry.h"
21+ 
22+namespace ops {
23+class SoftMarginLoss : public OpDef {
24+public:
25+ explicit SoftMarginLoss(const char* name) : OpDef(name)
26+ {
27+ this->Input("self")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
32+ .AutoContiguous();
33+ this->Input("target")
34+ .ParamType(REQUIRED)
35+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
36+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
37+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
38+ .AutoContiguous();
39+ this->Output("output")
40+ .ParamType(REQUIRED)
41+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
42+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
43+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
44+ .AutoContiguous();
45+ this->Attr("reduction")
46+ .Int(1); // default: mean (0=none, 1=mean, 2=sum)
47+ 
48+ OpAICoreConfig aicoreConfig;
49+ aicoreConfig.DynamicCompileStaticFlag(true)
50+ .DynamicFormatFlag(false)
51+ .DynamicRankSupportFlag(true)
52+ .DynamicShapeSupportFlag(true)
53+ .NeedCheckSupportFlag(false)
54+ .PrecisionReduceFlag(true)
55+ .ExtendCfgInfo("opFile.value", "soft_margin_loss");
56+ this->AICore().AddConfig("ascend910b", aicoreConfig);
57+ this->AICore().AddConfig("ascend950", aicoreConfig);
58+ }
59+};
60+OP_ADD(SoftMarginLoss);
61+} // namespace ops
Aexperimental/loss/soft_margin_loss/op_host/soft_margin_loss_infershape.cpp+66-0
@@ -0,0 +1,66 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file soft_margin_loss_infershape.cpp
18+ * \brief SoftMarginLoss shape inference
19+ *
20+ * - reduction=0 (none): output shape = input shape
21+ * - reduction=1 (mean) or 2 (sum): output shape = scalar (dimNum=0)
22+ */
23+ 
24+#include "register/op_impl_registry.h"
25+#include "exe_graph/runtime/infer_shape_context.h"
26+ 
27+using namespace ge;
28+ 
29+namespace ops {
30+ 
31+static ge::graphStatus InferShape4SoftMarginLoss(gert::InferShapeContext* context)
32+{
33+ const gert::Shape* inputShape = context->GetInputShape(0);
34+ if (inputShape == nullptr) {
35+ return ge::GRAPH_FAILED;
36+ }
37+ 
38+ gert::Shape* outputShape = context->GetOutputShape(0);
39+ if (outputShape == nullptr) {
40+ return ge::GRAPH_FAILED;
41+ }
42+ 
43+ // Get reduction attribute
44+ int64_t reduction = 1; // default: mean
45+ auto attrs = context->GetAttrs();
46+ if (attrs != nullptr) {
47+ const int64_t* reductionPtr = attrs->GetAttrPointer<int64_t>(0);
48+ if (reductionPtr != nullptr) {
49+ reduction = *reductionPtr;
50+ }
51+ }
52+ 
53+ if (reduction == 0) {
54+ // none: output shape = input shape
55+ *outputShape = *inputShape;
56+ } else {
57+ // mean or sum: output is scalar (0 dimensions)
58+ *outputShape = gert::Shape();
59+ }
60+ 
61+ return ge::GRAPH_SUCCESS;
62+}
63+ 
64+IMPL_OP_INFERSHAPE(SoftMarginLoss).InferShape(InferShape4SoftMarginLoss);
65+ 
66+} // namespace ops
Aexperimental/loss/soft_margin_loss/op_host/soft_margin_loss_tiling.cpp+251-0
@@ -0,0 +1,251 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+ 
17+/**
18+ * \file soft_margin_loss_tiling.cpp
19+ * \brief SoftMarginLoss Tiling implementation
20+ *
21+ * Computes tiling parameters for SoftMarginLoss operator:
22+ * - Multi-core splitting: totalNum divided evenly across AI Cores
23+ * - UB splitting: each core processes data in chunks of ubFactor elements
24+ * - TilingKey selection: based on input dtype and reduction mode
25+ *
26+ * Iteration 1: Implements float32 + none path, pre-embeds skeleton for other paths.
27+ */
28+ 
29+#include <cstring>
30+#include <cmath>
31+#include "register/op_def_registry.h"
32+#include "op_common/log/log.h"
33+#include "op_common/op_host/util/math_util.h"
34+#include "op_common/op_host/util/platform_util.h"
35+#include "../op_kernel/soft_margin_loss_tiling_data.h"
36+#include "../op_kernel/soft_margin_loss_tiling_key.h"
37+ 
38+namespace optiling {
39+ 
40+using Ops::Base::CeilDiv;
41+using Ops::Base::FloorDiv;
42+using Ops::Base::FloorAlign;
43+using Ops::Base::GetUbBlockSize;
44+ 
45+constexpr uint32_t WS_SYS_SIZE = 0U;
46+ 
47+// float32 + none: selfQueue(x2) + targetQueue(x2) + outputQueue(x2) + tmpBuf1(x1) + tmpBuf2(x1) = 8 float buffers
48+constexpr int64_t BUFFER_NUM_FP32_NONE = 8;
49+// float32 + reduce: selfQueue(x2) + targetQueue(x2) + tmpBuf1(x1) + tmpBuf2(x1) + tmpBuf3(x1) + reduceTmpBuf(x1) + partialSumBuf(x1) = 9
50+constexpr int64_t BUFFER_NUM_FP32_REDUCE = 9;
51+// float16 + none: selfQueue(x2,half=1eq) + targetQueue(x2,half=1eq) + outputQueue(x2,half=1eq) + tmpBuf1(float=2eq) + tmpBuf2(float=2eq) + tmpBuf3(float=2eq) = 6 float-equiv
52+constexpr int64_t BUFFER_NUM_FP16_NONE = 6;
53+// float16 + reduce: similar with reduce buffers = 7
54+constexpr int64_t BUFFER_NUM_FP16_REDUCE = 7;
55+ 
56+static const gert::Shape g_vec_1_shape = {1};
57+ 
58+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape)
59+{
60+ if (in_shape.GetDimNum() == 0) {
61+ return g_vec_1_shape;
62+ }
63+ return in_shape;
64+}
65+ 
66+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
67+{
68+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
69+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
70+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
71+ coreNum = ascendcPlatform.GetCoreNumAiv();
72+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
73+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
74+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
75+ return ge::GRAPH_SUCCESS;
76+}
77+ 
78+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context, size_t usrWorkspaceSize)
79+{
80+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
81+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
82+ currentWorkspace[0] = WS_SYS_SIZE + usrWorkspaceSize;
83+ return ge::GRAPH_SUCCESS;
84+}
85+ 
86+static ge::graphStatus SoftMarginLossTilingFunc(gert::TilingContext* context)
87+{
88+ // 1. Get platform info
89+ uint64_t ubSize;
90+ int64_t coreNum;
91+ OP_CHECK_IF(
92+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
93+ OP_LOGE(context, "GetPlatformInfo error"),
94+ return ge::GRAPH_FAILED);
95+ 
96+ // 2. Get input shape and dtype
97+ auto inputShape = context->GetInputShape(0);
98+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
99+ auto storageShape = EnsureNotScalar(inputShape->GetStorageShape());
100+ int64_t totalNum = storageShape.GetShapeSize();
101+ 
102+ auto inputDesc = context->GetInputDesc(0);
103+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
104+ ge::DataType dtype = inputDesc->GetDataType();
105+ 
106+ // 3. Get reduction attribute
107+ // NOTE: The aclnn runtime stores Int attrs as strings ("none"/"mean"/"sum") in TilingContext,
108+ // while the UT framework stores them as raw int64. We try int64 first, and if the value is
109+ // out of valid range [0,2], fall back to reading as string.
110+ const auto* attrs = context->GetAttrs();
111+ int64_t reduction = -1; // -1 means not yet determined
112+ if (attrs != nullptr) {
113+ const int64_t* intPtr = attrs->GetInt(0);
114+ if (intPtr != nullptr && *intPtr >= 0 && *intPtr <= 2) {
115+ reduction = *intPtr;
116+ OP_LOGI(context, "SoftMarginLoss: got reduction=%ld from int attr", reduction);
117+ } else {
118+ // Runtime may store as string: try string parsing
119+ const char* reductionStr = attrs->GetStr(0);
120+ if (reductionStr != nullptr) {
121+ if (strcmp(reductionStr, "none") == 0) {
122+ reduction = 0;
123+ } else if (strcmp(reductionStr, "mean") == 0) {
124+ reduction = 1;
125+ } else if (strcmp(reductionStr, "sum") == 0) {
126+ reduction = 2;
127+ } else {
128+ OP_LOGE(context, "SoftMarginLoss: invalid reduction string '%s', expected 0/1/2 or none/mean/sum", reductionStr);
129+ return ge::GRAPH_FAILED;
130+ }
131+ OP_LOGI(context, "SoftMarginLoss: got reduction='%s' -> %ld", reductionStr, reduction);
132+ } else if (intPtr == nullptr || *intPtr < 0 || *intPtr > 2) {
133+ // Neither valid int64 nor valid string found
134+ OP_LOGE(context, "SoftMarginLoss: could not read valid reduction attr");
135+ return ge::GRAPH_FAILED;
136+ }
137+ }
138+ } else {
139+ OP_LOGE(context, "SoftMarginLoss: context->GetAttrs() returned nullptr");
140+ return ge::GRAPH_FAILED;
141+ }
142+ 
143+ // 4. Set TilingData
144+ SoftMarginLossTilingData* tiling = context->GetTilingData<SoftMarginLossTilingData>();
145+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
146+ OP_CHECK_IF(
147+ memset_s(tiling, sizeof(SoftMarginLossTilingData), 0, sizeof(SoftMarginLossTilingData)) != EOK,
148+ OP_LOGE(context, "set tiling data error"),
149+ return ge::GRAPH_FAILED);
150+ 
151+ // Handle empty tensor
152+ if (totalNum == 0) {
153+ tiling->totalNum = 0;
154+ tiling->blockFactor = 0;
155+ tiling->ubFactor = 0;
156+ tiling->reductionMode = static_cast<int32_t>(reduction);
157+ tiling->invNumel = (reduction == 1) ? std::nanf("") : 0.0f;
158+ tiling->usedCoreNum = 1;
159+ context->SetBlockDim(1);
160+ // Select TilingKey based on dtype and reduction mode
161+ if (dtype == ge::DT_FLOAT) {
162+ context->SetTilingKey(GET_TPL_TILING_KEY(reduction == 0 ? SML_TPL_SCH_MODE_0 : SML_TPL_SCH_MODE_1));
163+ } else {
164+ context->SetTilingKey(GET_TPL_TILING_KEY(reduction == 0 ? SML_TPL_SCH_MODE_2 : SML_TPL_SCH_MODE_3));
165+ }
166+ OP_CHECK_IF(
167+ GetWorkspaceSize(context, 0) != ge::GRAPH_SUCCESS,
168+ OP_LOGE(context, "GetWorkspaceSize error"),
169+ return ge::GRAPH_FAILED);
170+ return ge::GRAPH_SUCCESS;
171+ }
172+ 
173+ // 5. Multi-core splitting
174+ tiling->totalNum = totalNum;
175+ bool isReduce = (reduction != 0);
176+ 
177+ // For reduce path: use single core to avoid cross-core SyncAll coordination.
178+ // The reduce path produces a single scalar output; all elements are processed
179+ // in tiles on core 0. This is correct and sufficient for iteration 2.
180+ // Multi-core reduce can be added as a performance optimization in a later iteration.
181+ int64_t usedCoreNum;
182+ if (isReduce) {
183+ tiling->blockFactor = totalNum;
184+ usedCoreNum = 1;
185+ } else {
186+ tiling->blockFactor = CeilDiv(totalNum, coreNum);
187+ usedCoreNum = CeilDiv(totalNum, tiling->blockFactor);
188+ }
189+ tiling->usedCoreNum = usedCoreNum;
190+ 
191+ // 6. Reduction mode
192+ tiling->reductionMode = static_cast<int32_t>(reduction);
193+ tiling->invNumel = (reduction == 1) ? (1.0f / static_cast<float>(totalNum)) : 0.0f;
194+ 
195+ // 7. UB splitting and TilingKey selection
196+ int64_t ubCanUse = static_cast<int64_t>(ubSize);
197+ int64_t ubBlockSize = GetUbBlockSize(context);
198+ // Both paths compute internally in float32, so typeSize = 4
199+ constexpr int64_t typeSize = 4;
200+ size_t usrWorkspaceSize = 0;
201+ 
202+ OP_LOGI(context, "SoftMarginLoss: dtype=%d, reduction=%ld, isReduce=%s, totalNum=%ld, coreNum=%ld, usedCoreNum=%ld",
203+ dtype, reduction, isReduce ? "true" : "false", totalNum, coreNum, usedCoreNum);
204+ 
205+ if (dtype == ge::DT_FLOAT) {
206+ if (!isReduce) {
207+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP32_NONE), ubBlockSize);
208+ auto tilingKey = GET_TPL_TILING_KEY(SML_TPL_SCH_MODE_0);
209+ OP_LOGI(context, "SoftMarginLoss: setting tiling key for FP32_NONE, key=%lu", tilingKey);
210+ context->SetTilingKey(tilingKey);
211+ } else {
212+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP32_REDUCE), ubBlockSize);
213+ auto tilingKey = GET_TPL_TILING_KEY(SML_TPL_SCH_MODE_1);
214+ OP_LOGI(context, "SoftMarginLoss: setting tiling key for FP32_REDUCE, key=%lu", tilingKey);
215+ context->SetTilingKey(tilingKey);
216+ usrWorkspaceSize = static_cast<size_t>(usedCoreNum) * 32; // 32-byte aligned per core
217+ }
218+ } else if (dtype == ge::DT_FLOAT16) {
219+ if (!isReduce) {
220+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP16_NONE), ubBlockSize);
221+ context->SetTilingKey(GET_TPL_TILING_KEY(SML_TPL_SCH_MODE_2));
222+ } else {
223+ tiling->ubFactor = FloorAlign(FloorDiv((ubCanUse / typeSize), BUFFER_NUM_FP16_REDUCE), ubBlockSize);
224+ context->SetTilingKey(GET_TPL_TILING_KEY(SML_TPL_SCH_MODE_3));
225+ usrWorkspaceSize = static_cast<size_t>(usedCoreNum) * 32;
226+ }
227+ } else {
228+ OP_LOGE(context, "SoftMarginLoss: unsupported dtype");
229+ return ge::GRAPH_FAILED;
230+ }
231+ 
232+ // 8. Set workspace
233+ OP_CHECK_IF(
234+ GetWorkspaceSize(context, usrWorkspaceSize) != ge::GRAPH_SUCCESS,
235+ OP_LOGE(context, "GetWorkspaceSize error"),
236+ return ge::GRAPH_FAILED);
237+ 
238+ context->SetBlockDim(usedCoreNum);
239+ return ge::GRAPH_SUCCESS;
240+}
241+ 
242+static ge::graphStatus TilingParseForSoftMarginLoss([[maybe_unused]] gert::TilingParseContext* context)
243+{
244+ return ge::GRAPH_SUCCESS;
245+}
246+ 
247+struct SoftMarginLossCompileInfo {};
248+ 
249+IMPL_OP_OPTILING(SoftMarginLoss).Tiling(SoftMarginLossTilingFunc).TilingParse<SoftMarginLossCompileInfo>(TilingParseForSoftMarginLoss);
250+ 
251+} // namespace optiling
Aexperimental/loss/soft_margin_loss/op_kernel/soft_margin_loss.cpp+54-0
@@ -0,0 +1,54 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file soft_margin_loss_arch32.cpp
18+ * \brief SoftMarginLoss kernel entry point (arch32 architecture - Ascend910B)
19+ *
20+ * Dispatches to template instantiations:
21+ * SoftMarginLossNone<float/half> - elementwise output
22+ * SoftMarginLossReduce<float/half> - scalar reduction output
23+ */
24+ 
25+#include "soft_margin_loss.h"
26+ 
27+template <uint32_t schMode>
28+__global__ __aicore__ void soft_margin_loss(GM_ADDR selfInput, GM_ADDR targetInput,
29+ GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling)
30+{
31+ REGISTER_TILING_DEFAULT(SoftMarginLossTilingData);
32+ GET_TILING_DATA_WITH_STRUCT(SoftMarginLossTilingData, tilingData, tiling);
33+ 
34+ if constexpr (schMode == SML_TPL_SCH_MODE_0) {
35+ NsSoftMarginLoss::SoftMarginLossNone<float> op;
36+ op.Init(selfInput, targetInput, output, &tilingData);
37+ op.Process();
38+ }
39+ if constexpr (schMode == SML_TPL_SCH_MODE_1) {
40+ NsSoftMarginLoss::SoftMarginLossReduce<float> op;
41+ op.Init(selfInput, targetInput, output, workspace, &tilingData);
42+ op.Process();
43+ }
44+ if constexpr (schMode == SML_TPL_SCH_MODE_2) {
45+ NsSoftMarginLoss::SoftMarginLossNone<half> op;
46+ op.Init(selfInput, targetInput, output, &tilingData);
47+ op.Process();
48+ }
49+ if constexpr (schMode == SML_TPL_SCH_MODE_3) {
50+ NsSoftMarginLoss::SoftMarginLossReduce<half> op;
51+ op.Init(selfInput, targetInput, output, workspace, &tilingData);
52+ op.Process();
53+ }
54+}
Aexperimental/loss/soft_margin_loss/op_kernel/soft_margin_loss.h+499-0
@@ -0,0 +1,499 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file soft_margin_loss.h
18+ * \brief SoftMarginLoss kernel class definition (arch32 - Ascend910B)
19+ *
20+ * Computes SoftMarginLoss: L[i] = max(0, -t*x) + log(1 + exp(-|t*x|))
21+ * where x = self[i], t = target[i]
22+ *
23+ * Template-based implementation:
24+ * SoftMarginLossNone<T> - elementwise output (reduction='none')
25+ * SoftMarginLossReduce<T> - scalar output (reduction='mean'/'sum')
26+ * where T = float or half
27+ *
28+ * Uses double buffering (depth=2) for pipeline parallelism.
29+ * Reduce path uses two-phase cross-core reduction via workspace.
30+ */
31+#ifndef SOFT_MARGIN_LOSS_H
32+#define SOFT_MARGIN_LOSS_H
33+ 
34+#include "kernel_operator.h"
35+#include "kernel_tiling/kernel_tiling.h"
36+#include "soft_margin_loss_tiling_data.h"
37+#include "soft_margin_loss_tiling_key.h"
38+ 
39+namespace NsSoftMarginLoss {
40+ 
41+using namespace AscendC;
42+ 
43+constexpr int32_t BUFFER_NUM = 2;
44+ 
45+// ============================================================================
46+// ComputeLossCore - 9-step SoftMarginLoss core computation (float32)
47+// L = max(0, -tx) + log(1 + exp(-|tx|))
48+//
49+// outputF32: result written here
50+// txLocal: temp buffer 1 (reused across steps)
51+// negTxLocal: temp buffer 2 (reused across steps)
52+// selfF32: self values in float32
53+// targetF32: target values in float32
54+// ============================================================================
55+__aicore__ inline void ComputeLossCore(
56+ LocalTensor<float>& outputF32,
57+ LocalTensor<float>& txLocal,
58+ LocalTensor<float>& negTxLocal,
59+ LocalTensor<float>& selfF32,
60+ LocalTensor<float>& targetF32,
61+ int64_t currentNum)
62+{
63+ // Step 1: tx = target * self
64+ AscendC::Mul(txLocal, targetF32, selfF32, currentNum);
65+ // Step 2: neg_tx = -tx
66+ AscendC::Muls(negTxLocal, txLocal, static_cast<float>(-1.0f), currentNum);
67+ // Step 3: abs_tx = |tx| (reuse txLocal)
68+ AscendC::Abs(txLocal, txLocal, currentNum);
69+ // Step 4: -|tx| (reuse txLocal)
70+ AscendC::Muls(txLocal, txLocal, static_cast<float>(-1.0f), currentNum);
71+ // Step 5: exp(-|tx|) (reuse txLocal)
72+ AscendC::Exp(txLocal, txLocal, currentNum);
73+ // Step 6: 1 + exp(-|tx|) (reuse txLocal)
74+ AscendC::Adds(txLocal, txLocal, static_cast<float>(1.0f), currentNum);
75+ // Step 7: log(1 + exp(-|tx|)) (reuse txLocal)
76+ AscendC::Log(txLocal, txLocal, currentNum);
77+ // Step 8: max(0, neg_tx) (reuse negTxLocal)
78+ AscendC::Maxs(negTxLocal, negTxLocal, static_cast<float>(0.0f), currentNum);
79+ // Step 9: L = max(0, -tx) + log(1 + exp(-|tx|))
80+ AscendC::Add(outputF32, negTxLocal, txLocal, currentNum);
81+}
82+ 
83+// ============================================================================
84+// SoftMarginLossNone<T> - Elementwise path (reduction='none')
85+// T = float: direct compute
86+// T = half: Cast to float -> compute -> Cast back to half
87+// ============================================================================
88+template <typename T>
89+class SoftMarginLossNone {
90+public:
91+ __aicore__ inline SoftMarginLossNone() {}
92+ 
93+ __aicore__ inline void Init(GM_ADDR selfGm, GM_ADDR targetGm, GM_ADDR outputGm,
94+ const SoftMarginLossTilingData* tilingData);
95+ __aicore__ inline void Process();
96+ 
97+private:
98+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
99+ __aicore__ inline void Compute(int64_t currentNum);
100+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
101+ 
102+private:
103+ TPipe pipe;
104+ TQue<QuePosition::VECIN, BUFFER_NUM> selfQueue;
105+ TQue<QuePosition::VECIN, BUFFER_NUM> targetQueue;
106+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueue;
107+ TBuf<QuePosition::VECCALC> tmpBuf1; // float: tx / abs / exp / log
108+ TBuf<QuePosition::VECCALC> tmpBuf2; // float: neg_tx / max_val
109+ TBuf<QuePosition::VECCALC> castBuf; // float: Cast workspace (half path only)
110+ 
111+ GlobalTensor<T> selfGM;
112+ GlobalTensor<T> targetGM;
113+ GlobalTensor<T> outputGM;
114+ 
115+ int64_t blockLength_ = 0;
116+ int64_t ubLength_ = 0;
117+};
118+ 
119+// ---- Init ----
120+template <typename T>
121+__aicore__ inline void SoftMarginLossNone<T>::Init(GM_ADDR selfGm, GM_ADDR targetGm, GM_ADDR outputGm,
122+ const SoftMarginLossTilingData* tilingData)
123+{
124+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx();
125+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
126+ if (blockLength_ < 0) {
127+ blockLength_ = 0;
128+ }
129+ ubLength_ = tilingData->ubFactor;
130+ 
131+ if (blockLength_ <= 0 || ubLength_ <= 0) {
132+ return;
133+ }
134+ 
135+ int64_t coreOffset = tilingData->blockFactor * AscendC::GetBlockIdx();
136+ selfGM.SetGlobalBuffer((__gm__ T*)selfGm + coreOffset, blockLength_);
137+ targetGM.SetGlobalBuffer((__gm__ T*)targetGm + coreOffset, blockLength_);
138+ outputGM.SetGlobalBuffer((__gm__ T*)outputGm + coreOffset, blockLength_);
139+ 
140+ pipe.InitBuffer(selfQueue, BUFFER_NUM, ubLength_ * sizeof(T));
141+ pipe.InitBuffer(targetQueue, BUFFER_NUM, ubLength_ * sizeof(T));
142+ pipe.InitBuffer(outputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
143+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(float));
144+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(float));
145+ if constexpr (std::is_same_v<T, half>) {
146+ pipe.InitBuffer(castBuf, ubLength_ * sizeof(float));
147+ }
148+}
149+ 
150+// ---- CopyIn ----
151+template <typename T>
152+__aicore__ inline void SoftMarginLossNone<T>::CopyIn(int64_t progress, int64_t currentNum)
153+{
154+ AscendC::LocalTensor<T> selfLocal = selfQueue.template AllocTensor<T>();
155+ AscendC::DataCopyParams copyParams;
156+ copyParams.blockCount = 1;
157+ copyParams.blockLen = currentNum * sizeof(T);
158+ copyParams.srcStride = 0;
159+ copyParams.dstStride = 0;
160+ AscendC::DataCopyPad(selfLocal, selfGM[progress * ubLength_], copyParams, {false, 0, 0, 0});
161+ selfQueue.EnQue(selfLocal);
162+ 
163+ AscendC::LocalTensor<T> targetLocal = targetQueue.template AllocTensor<T>();
164+ AscendC::DataCopyPad(targetLocal, targetGM[progress * ubLength_], copyParams, {false, 0, 0, 0});
165+ targetQueue.EnQue(targetLocal);
166+}
167+ 
168+// ---- Compute ----
169+template <typename T>
170+__aicore__ inline void SoftMarginLossNone<T>::Compute(int64_t currentNum)
171+{
172+ AscendC::LocalTensor<T> selfLocal = selfQueue.template DeQue<T>();
173+ AscendC::LocalTensor<T> targetLocal = targetQueue.template DeQue<T>();
174+ AscendC::LocalTensor<T> outputLocal = outputQueue.template AllocTensor<T>();
175+ AscendC::LocalTensor<float> txLocal = tmpBuf1.Get<float>();
176+ AscendC::LocalTensor<float> negTxLocal = tmpBuf2.Get<float>();
177+ 
178+ if constexpr (std::is_same_v<T, float>) {
179+ // float path: compute directly
180+ ComputeLossCore(outputLocal, txLocal, negTxLocal, selfLocal, targetLocal, currentNum);
181+ } else {
182+ // half path: Cast to float -> compute -> Cast back
183+ AscendC::LocalTensor<float> castLocal = castBuf.Get<float>();
184+ // Cast self (half -> float) into castLocal
185+ AscendC::Cast(castLocal, selfLocal, AscendC::RoundMode::CAST_NONE, currentNum);
186+ // Cast target (half -> float) into txLocal (temporary reuse)
187+ AscendC::Cast(txLocal, targetLocal, AscendC::RoundMode::CAST_NONE, currentNum);
188+ // Compute in float: result in negTxLocal (reused as output buffer)
189+ ComputeLossCore(negTxLocal, txLocal, negTxLocal, castLocal, txLocal, currentNum);
190+ // Wait - negTxLocal is both output and temp. Need separate flow:
191+ // selfF32 = castLocal, targetF32 = txLocal after cast
192+ // But ComputeLossCore overwrites txLocal in step 1 (Mul into txLocal).
193+ // And negTxLocal is used as both neg_tx buffer and output. This works because
194+ // step 9 writes Add result to outputF32 (=negTxLocal) using negTxLocal and txLocal as inputs.
195+ // Actually let's re-examine: we need selfF32 and targetF32 as separate inputs to step 1.
196+ // After Cast: castLocal=selfF32, txLocal=targetF32
197+ // Step 1: Mul(txLocal, targetF32=txLocal, selfF32=castLocal) -> txLocal = tx (OK, txLocal overwritten)
198+ // This is correct because Mul reads txLocal and castLocal first, then writes txLocal.
199+ // The rest follows the same pattern as float path.
200+ // But outputF32=negTxLocal and negTxLocal param is also negTxLocal - that's fine,
201+ // the output is written in step 9, after negTxLocal is last read in step 9 itself.
202+ // Actually no: Add(output, negTx, tx) reads negTx then writes output=negTx. This is in-place, which is OK.
203+ // Let me use a cleaner approach with castLocal as the third buffer:
204+ // selfF32=castLocal, targetF32=txLocal (after second Cast)
205+ // Then ComputeLossCore(output=negTxLocal, tx_buf=txLocal, neg_buf=castLocal, self=castLocal, target=txLocal)
206+ // Problem: castLocal is both selfF32 input and neg_tx buffer. Step 1 reads selfF32=castLocal,
207+ // step 2 writes neg_tx=castLocal. But step 1 is Mul(txLocal, target, self) which only writes txLocal.
208+ // So castLocal (selfF32) is only read in step 1, then reused as neg_tx from step 2 onward. That works!
209+ 
210+ // Re-do with correct buffer assignment:
211+ // castLocal = selfF32 (from Cast), then reused as neg_tx buffer from step 2
212+ // txLocal = targetF32 (from Cast), then reused as tx buffer from step 1
213+ // negTxLocal = output destination
214+ // But wait, ComputeLossCore signature: (output, txBuf, negTxBuf, selfF32, targetF32)
215+ // We need: output=negTxLocal, txBuf=txLocal, negTxBuf=castLocal, selfF32=castLocal, targetF32=txLocal
216+ // negTxBuf and selfF32 are both castLocal - step 1 reads selfF32, step 2 writes negTxBuf.
217+ // Inside ComputeLossCore step 1: Mul(txLocal, targetF32=txLocal, selfF32=castLocal) - reads both, writes txLocal
218+ // Step 2: Muls(negTxLocal=castLocal, txLocal, -1) - reads txLocal (from step 1), writes castLocal. OK!
219+ // Step 9: Add(output=negTxLocal, negTxLocal=castLocal, txLocal) -> writes negTxLocal. OK!
220+ // This works correctly.
221+ ComputeLossCore(negTxLocal, txLocal, castLocal, castLocal, txLocal, currentNum);
222+ // Cast result (float -> half)
223+ AscendC::Cast(outputLocal, negTxLocal, AscendC::RoundMode::CAST_ROUND, currentNum);
224+ }
225+ 
226+ outputQueue.template EnQue<T>(outputLocal);
227+ selfQueue.FreeTensor(selfLocal);
228+ targetQueue.FreeTensor(targetLocal);
229+}
230+ 
231+// ---- CopyOut ----
232+template <typename T>
233+__aicore__ inline void SoftMarginLossNone<T>::CopyOut(int64_t progress, int64_t currentNum)
234+{
235+ AscendC::LocalTensor<T> outputLocal = outputQueue.template DeQue<T>();
236+ AscendC::DataCopyParams copyParams;
237+ copyParams.blockCount = 1;
238+ copyParams.blockLen = currentNum * sizeof(T);
239+ copyParams.srcStride = 0;
240+ copyParams.dstStride = 0;
241+ AscendC::DataCopyPad(outputGM[progress * ubLength_], outputLocal, copyParams);
242+ outputQueue.FreeTensor(outputLocal);
243+}
244+ 
245+// ---- Process ----
246+template <typename T>
247+__aicore__ inline void SoftMarginLossNone<T>::Process()
248+{
249+ if (blockLength_ <= 0 || ubLength_ <= 0) {
250+ return;
251+ }
252+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
253+ for (int64_t i = 0; i < loopCount; i++) {
254+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
255+ CopyIn(i, currentNum);
256+ Compute(currentNum);
257+ CopyOut(i, currentNum);
258+ }
259+}
260+ 
261+// ============================================================================
262+// SoftMarginLossReduce<T> - Reduction path (reduction='mean'/'sum')
263+// T = float or half
264+//
265+// Two-phase cross-core reduction:
266+// Phase 1: Each core computes elementwise loss + local ReduceSum -> partialSum
267+// Writes partialSum to workspace[coreIdx * 8] (32-byte aligned)
268+// Phase 2: After SyncAll, core 0 reads all partial sums, aggregates,
269+// applies mean division if needed, writes scalar output
270+// ============================================================================
271+template <typename T>
272+class SoftMarginLossReduce {
273+public:
274+ __aicore__ inline SoftMarginLossReduce() {}
275+ 
276+ __aicore__ inline void Init(GM_ADDR selfGm, GM_ADDR targetGm, GM_ADDR outputGm,
277+ GM_ADDR workspaceGm, const SoftMarginLossTilingData* tilingData);
278+ __aicore__ inline void Process();
279+ 
280+private:
281+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
282+ __aicore__ inline void ComputeAndReduce(int64_t currentNum, float& localSum);
283+ __aicore__ inline void WritePartialSum(float localSum);
284+ __aicore__ inline void CrossCoreReduce();
285+ __aicore__ inline void WriteFinalScalar(float value);
286+ 
287+private:
288+ TPipe pipe;
289+ TQue<QuePosition::VECIN, BUFFER_NUM> selfQueue;
290+ TQue<QuePosition::VECIN, BUFFER_NUM> targetQueue;
291+ TBuf<QuePosition::VECCALC> tmpBuf1; // float: tx / abs / exp / log
292+ TBuf<QuePosition::VECCALC> tmpBuf2; // float: neg_tx / max_val
293+ TBuf<QuePosition::VECCALC> tmpBuf3; // float: loss result (ReduceSum input)
294+ TBuf<QuePosition::VECCALC> castBuf; // float: Cast workspace (half path only)
295+ TBuf<QuePosition::VECCALC> reduceTmpBuf; // float: ReduceSum temporary buffer
296+ TBuf<QuePosition::VECCALC> resultBuf; // small buffer for partial sum / final result
297+ 
298+ GlobalTensor<T> selfGM;
299+ GlobalTensor<T> targetGM;
300+ GlobalTensor<T> outputGM;
301+ GlobalTensor<float> workspaceGM;
302+ 
303+ int64_t blockLength_ = 0;
304+ int64_t ubLength_ = 0;
305+ int32_t reductionMode_ = 0;
306+ float invNumel_ = 0.0f;
307+ int64_t usedCoreNum_ = 0;
308+};
309+ 
310+// ---- Init ----
311+template <typename T>
312+__aicore__ inline void SoftMarginLossReduce<T>::Init(GM_ADDR selfGm, GM_ADDR targetGm, GM_ADDR outputGm,
313+ GM_ADDR workspaceGm,
314+ const SoftMarginLossTilingData* tilingData)
315+{
316+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx();
317+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
318+ if (blockLength_ < 0) {
319+ blockLength_ = 0;
320+ }
321+ ubLength_ = tilingData->ubFactor;
322+ reductionMode_ = tilingData->reductionMode;
323+ invNumel_ = tilingData->invNumel;
324+ usedCoreNum_ = tilingData->usedCoreNum;
325+ 
326+ int64_t coreOffset = tilingData->blockFactor * AscendC::GetBlockIdx();
327+ selfGM.SetGlobalBuffer((__gm__ T*)selfGm + coreOffset, (blockLength_ > 0) ? blockLength_ : 1);
328+ targetGM.SetGlobalBuffer((__gm__ T*)targetGm + coreOffset, (blockLength_ > 0) ? blockLength_ : 1);
329+ outputGM.SetGlobalBuffer((__gm__ T*)outputGm, 1);
330+ workspaceGM.SetGlobalBuffer((__gm__ float*)workspaceGm, usedCoreNum_ * 8);
331+ 
332+ if (ubLength_ > 0) {
333+ pipe.InitBuffer(selfQueue, BUFFER_NUM, ubLength_ * sizeof(T));
334+ pipe.InitBuffer(targetQueue, BUFFER_NUM, ubLength_ * sizeof(T));
335+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(float));
336+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(float));
337+ pipe.InitBuffer(tmpBuf3, ubLength_ * sizeof(float));
338+ pipe.InitBuffer(reduceTmpBuf, ubLength_ * sizeof(float));
339+ if constexpr (std::is_same_v<T, half>) {
340+ pipe.InitBuffer(castBuf, ubLength_ * sizeof(float));
341+ }
342+ }
343+ int64_t resultBufSize = (usedCoreNum_ * sizeof(float) + 31) & ~31;
344+ if (resultBufSize < 32) {
345+ resultBufSize = 32;
346+ }
347+ pipe.InitBuffer(resultBuf, resultBufSize);
348+}
349+ 
350+// ---- CopyIn ----
351+template <typename T>
352+__aicore__ inline void SoftMarginLossReduce<T>::CopyIn(int64_t progress, int64_t currentNum)
353+{
354+ AscendC::LocalTensor<T> selfLocal = selfQueue.template AllocTensor<T>();
355+ AscendC::DataCopyParams copyParams;
356+ copyParams.blockCount = 1;
357+ copyParams.blockLen = currentNum * sizeof(T);
358+ copyParams.srcStride = 0;
359+ copyParams.dstStride = 0;
360+ AscendC::DataCopyPad(selfLocal, selfGM[progress * ubLength_], copyParams, {false, 0, 0, 0});
361+ selfQueue.EnQue(selfLocal);
362+ 
363+ AscendC::LocalTensor<T> targetLocal = targetQueue.template AllocTensor<T>();
364+ AscendC::DataCopyPad(targetLocal, targetGM[progress * ubLength_], copyParams, {false, 0, 0, 0});
365+ targetQueue.EnQue(targetLocal);
366+}
367+ 
368+// ---- ComputeAndReduce ----
369+template <typename T>
370+__aicore__ inline void SoftMarginLossReduce<T>::ComputeAndReduce(int64_t currentNum, float& localSum)
371+{
372+ AscendC::LocalTensor<T> selfLocal = selfQueue.template DeQue<T>();
373+ AscendC::LocalTensor<T> targetLocal = targetQueue.template DeQue<T>();
374+ AscendC::LocalTensor<float> txLocal = tmpBuf1.Get<float>();
375+ AscendC::LocalTensor<float> negTxLocal = tmpBuf2.Get<float>();
376+ AscendC::LocalTensor<float> lossLocal = tmpBuf3.Get<float>();
377+ AscendC::LocalTensor<float> reduceTmp = reduceTmpBuf.Get<float>();
378+ 
379+ if constexpr (std::is_same_v<T, float>) {
380+ ComputeLossCore(lossLocal, txLocal, negTxLocal, selfLocal, targetLocal, currentNum);
381+ } else {
382+ AscendC::LocalTensor<float> castLocal = castBuf.Get<float>();
383+ // Cast self (half -> float) into castLocal
384+ AscendC::Cast(castLocal, selfLocal, AscendC::RoundMode::CAST_NONE, currentNum);
385+ // Cast target (half -> float) into txLocal (temporary)
386+ AscendC::Cast(txLocal, targetLocal, AscendC::RoundMode::CAST_NONE, currentNum);
387+ // Compute: output=lossLocal, txBuf=txLocal, negTxBuf=castLocal, selfF32=castLocal, targetF32=txLocal
388+ ComputeLossCore(lossLocal, txLocal, castLocal, castLocal, txLocal, currentNum);
389+ }
390+ 
391+ // ReduceSum over this tile
392+ AscendC::ReduceSum<float>(lossLocal, lossLocal, reduceTmp, currentNum);
393+ localSum += lossLocal.GetValue(0);
394+ 
395+ selfQueue.FreeTensor(selfLocal);
396+ targetQueue.FreeTensor(targetLocal);
397+}
398+ 
399+// ---- WritePartialSum ----
400+template <typename T>
401+__aicore__ inline void SoftMarginLossReduce<T>::WritePartialSum(float localSum)
402+{
403+ AscendC::LocalTensor<float> resultLocal = resultBuf.Get<float>();
404+ resultLocal.SetValue(0, localSum);
405+ 
406+ AscendC::DataCopyParams copyParams;
407+ copyParams.blockCount = 1;
408+ copyParams.blockLen = 32;
409+ copyParams.srcStride = 0;
410+ copyParams.dstStride = 0;
411+ AscendC::DataCopyPad(workspaceGM[AscendC::GetBlockIdx() * 8], resultLocal, copyParams);
412+}
413+ 
414+// ---- WriteFinalScalar ----
415+template <typename T>
416+__aicore__ inline void SoftMarginLossReduce<T>::WriteFinalScalar(float value)
417+{
418+ AscendC::LocalTensor<float> resultLocal = resultBuf.Get<float>();
419+ 
420+ if constexpr (std::is_same_v<T, float>) {
421+ resultLocal.SetValue(0, value);
422+ } else {
423+ // Cast float scalar to half via ReinterpretCast
424+ AscendC::LocalTensor<half> halfResult = resultLocal.ReinterpretCast<half>();
425+ halfResult.SetValue(0, static_cast<half>(value));
426+ }
427+ 
428+ AscendC::DataCopyParams outParams;
429+ outParams.blockCount = 1;
430+ outParams.blockLen = 32;
431+ outParams.srcStride = 0;
432+ outParams.dstStride = 0;
433+ 
434+ if constexpr (std::is_same_v<T, float>) {
435+ AscendC::DataCopyPad(outputGM[0], resultLocal, outParams);
436+ } else {
437+ AscendC::LocalTensor<half> halfResult = resultLocal.ReinterpretCast<half>();
438+ AscendC::DataCopyPad(outputGM[0], halfResult, outParams);
439+ }
440+}
441+ 
442+// ---- CrossCoreReduce ----
443+template <typename T>
444+__aicore__ inline void SoftMarginLossReduce<T>::CrossCoreReduce()
445+{
446+ if (AscendC::GetBlockIdx() != 0) {
447+ return;
448+ }
449+ 
450+ AscendC::LocalTensor<float> resultLocal = resultBuf.Get<float>();
451+ 
452+ float globalSum = 0.0f;
453+ for (int64_t i = 0; i < usedCoreNum_; i++) {
454+ AscendC::DataCopyParams copyParams;
455+ copyParams.blockCount = 1;
456+ copyParams.blockLen = 32;
457+ copyParams.srcStride = 0;
458+ copyParams.dstStride = 0;
459+ AscendC::DataCopyPad(resultLocal, workspaceGM[i * 8], copyParams, {false, 0, 0, 0});
460+ AscendC::PipeBarrier<PIPE_ALL>();
461+ globalSum += resultLocal.GetValue(0);
462+ }
463+ 
464+ if (reductionMode_ == 1) {
S
Ssu-yueming4月13日

[Low] 空 tensor + mean 模式输出 0.0f 而非 NaN（与 PyTorch 不一致）

文件: op_kernel/soft_margin_loss.h:464-468（CrossCoreReduce 函数）

if (reductionMode_ == 1) {
    globalSum *= invNumel_;  // invNumel = 0.0f when totalNum == 0
}
WriteFinalScalar(globalSum);  // 输出 0.0f * 0.0f = 0.0f

问题：当 totalNum == 0 时，Tiling 侧设置 invNumel = 0.0f，最终输出 0.0f。但 PyTorch 的 torch.nn.functional.soft_margin_loss 在空 tensor + mean 模式下返回 nan。

PyTorch 依据：

  • torch.mean() 官方文档：空 tensor 返回 nan
  • GitHub Issue #50224：社区确认所有 loss 函数（含 soft_margin_loss）在 empty + mean 模式下返回 nan
  • 本地验证：soft_margin_loss(torch.tensor([]), torch.tensor([]), reduction='mean') → nan

建议修复：对空 tensor + mean 模式特殊处理，输出 NaN：

if (reductionMode_ == 1) {
    if (tilingData->totalNum == 0) {
        // 0.0f / 0.0f = NaN, 与 PyTorch mean 空 tensor 行为一致
        globalSum = 0.0f / 0.0f;\n    } else {\n        globalSum *= invNumel_;\n    }\n}
likedislike
465+ globalSum *= invNumel_;
466+ }
467+ 
468+ WriteFinalScalar(globalSum);
469+}
470+ 
471+// ---- Process ----
472+template <typename T>
473+__aicore__ inline void SoftMarginLossReduce<T>::Process()
474+{
475+ float localSum = 0.0f;
476+ 
477+ if (blockLength_ > 0 && ubLength_ > 0) {
478+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
479+ for (int64_t i = 0; i < loopCount; i++) {
480+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
481+ CopyIn(i, currentNum);
482+ ComputeAndReduce(currentNum, localSum);
483+ }
484+ }
485+ 
486+ if (usedCoreNum_ == 1) {
487+ if (reductionMode_ == 1) {
488+ localSum *= invNumel_;
489+ }
490+ WriteFinalScalar(localSum);
491+ } else {
492+ WritePartialSum(localSum);
493+ AscendC::SyncAll();
494+ CrossCoreReduce();
495+ }
496+}
497+ 
498+} // namespace NsSoftMarginLoss
499+#endif // SOFT_MARGIN_LOSS_H
Aexperimental/loss/soft_margin_loss/op_kernel/soft_margin_loss_tiling_data.h+33-0
@@ -0,0 +1,33 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+/**
17+ * \file soft_margin_loss_tiling_data.h
18+ * \brief SoftMarginLoss TilingData structure definition
19+ */
20+ 
21+#ifndef _SOFT_MARGIN_LOSS_TILING_DATA_H_
22+#define _SOFT_MARGIN_LOSS_TILING_DATA_H_
23+ 
24+struct SoftMarginLossTilingData {
25+ int64_t totalNum = 0; // Total number of elements
26+ int64_t blockFactor = 0; // Number of elements per AI Core
27+ int64_t ubFactor = 0; // Number of elements per UB loop iteration (in float32 units)
28+ int32_t reductionMode = 0; // Reduction mode: 0=none, 1=mean, 2=sum
29+ float invNumel = 0.0f; // 1.0f / totalNum (only used in mean mode)
30+ int64_t usedCoreNum = 0; // Actual number of cores used (for cross-core reduction)
31+};
32+ 
33+#endif
Aexperimental/loss/soft_margin_loss/op_kernel/soft_margin_loss_tiling_key.h+49-0
@@ -0,0 +1,49 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/**
12+ * NOTE: Portions of this code were AI-generated and have been
13+ * technically reviewed for functional accuracy and security
14+ */
15+ 
16+ 
17+/**
18+ * \file soft_margin_loss_tiling_key.h
19+ * \brief SoftMarginLoss TilingKey definition
20+ *
21+ * TilingKey mapping:
22+ * - SML_TPL_SCH_MODE_0 (0): FLOAT32 + NONE (elementwise output)
23+ * - SML_TPL_SCH_MODE_1 (1): FLOAT32 + REDUCE (mean/sum)
24+ * - SML_TPL_SCH_MODE_2 (2): FLOAT16 + NONE (elementwise output)
25+ * - SML_TPL_SCH_MODE_3 (3): FLOAT16 + REDUCE (mean/sum)
26+ */
27+ 
28+#ifndef __SOFT_MARGIN_LOSS_TILING_KEY_H__
29+#define __SOFT_MARGIN_LOSS_TILING_KEY_H__
30+ 
31+#include "ascendc/host_api/tiling/template_argument.h"
32+ 
33+#define SML_TPL_SCH_MODE_0 0 // FLOAT32 + NONE
34+#define SML_TPL_SCH_MODE_1 1 // FLOAT32 + REDUCE (mean/sum)
35+#define SML_TPL_SCH_MODE_2 2 // FLOAT16 + NONE
36+#define SML_TPL_SCH_MODE_3 3 // FLOAT16 + REDUCE (mean/sum)
37+ 
38+ASCENDC_TPL_ARGS_DECL(
39+ SoftMarginLoss,
40+ ASCENDC_TPL_UINT_DECL(schMode, 2, ASCENDC_TPL_UI_LIST,
41+ SML_TPL_SCH_MODE_0, SML_TPL_SCH_MODE_1,
42+ SML_TPL_SCH_MODE_2, SML_TPL_SCH_MODE_3));
43+ 
44+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
45+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,
46+ SML_TPL_SCH_MODE_0, SML_TPL_SCH_MODE_1,
47+ SML_TPL_SCH_MODE_2, SML_TPL_SCH_MODE_3)));
48+ 
49+#endif
Aexperimental/loss/soft_margin_loss/tests/.gitkeep+0-0
The file is empty