Star1291
1404
代码
代码
Issues108
Pull Requests204
Actions
项目讨论
Wiki
分析
Star1291
1404
已合并
SignBitsUnpack算子贡献 #2351
松柏创建于 4月21日
SignBitsUnpack算子贡献 #2351
已合并
松柏创建于 4月21日
从ilovescrapy/ops-mathunpack-master合入到cann/ops-mathmaster
共 22 个文件变更+1761-0
Aexperimental/math/sign_bits_unpack/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_all_modules_sources(OPTYPE sign_bits_unpack ACLNNTYPE aclnn_exclude)
Aexperimental/math/sign_bits_unpack/README.md+80-0
@@ -0,0 +1,80 @@
1+# SignBitsUnpack
2+ 
3+## 产品支持情况
4+ 
5+ 
6+| 产品 | 是否支持 |
7+| ------------------------------------------------------------------ | :------: |
8+| Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件 | √ |
9+ 
10+## 功能说明
11+ 
12+- 算子功能：对输入进行unpack。
13+ 
14+当位置为1时取1.0，位置为0时取0.0
15+ 
16+## 参数说明
17+ 
18+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
19+ <col style="width: 100px">
20+ <col style="width: 150px">
21+ <col style="width: 280px">
22+ <col style="width: 330px">
23+ <col style="width: 120px">
24+ </colgroup>
25+ <thead>
26+ <tr>
27+ <th>参数名</th>
28+ <th>输入/输出/属性</th>
29+ <th>描述</th>
30+ <th>数据类型</th>
31+ <th>数据格式</th>
32+ </tr></thead>
33+ <tbody>
34+ <tr>
35+ <td>self</td>
36+ <td>输入</td>
37+ <td>待进行SignBitsUnpack计算的入参，公式中的x1。</td>
38+ <td>uint8</td>
39+ <td>ND</td>
40+ </tr>
41+ <tr>
42+ <td>size</td>
43+ <td>参数</td>
44+ <td>reshape时输出张量的第一个维度</td>
45+ <td>int64</td>
46+ <td>1</td>
47+ </tr>
48+ <tr>
49+ <td>dtype</td>
50+ <td>参数</td>
51+ <td>决定输出的数据类型</td>
52+ <td>int64</td>
53+ <td>1</td>
54+ </tr>
55+ <tr>
56+ <td>y</td>
57+ <td>输出</td>
58+ <td>待进行SignBitsUnpack计算的出参，公式中的输出。</td>
59+ <td>float16,float</td>
60+ <td>ND</td>
61+ </tr>
62+ </tbody></table>
63+ 
64+## 约束说明
65+ 
66+无
67+ 
68+## 调用说明
69+ 
70+ 
71+| 调用方式 | 调用样例 | 说明 |
72+| --------- | ------------------------------------------------------------------------- | ---------------------------------------------------------------------------------------- |
73+| aclnn调用 | [test_aclnn_sign_bits_unpack](./examples/test_aclnn_sign_bits_unpack.cpp) | 通过[aclnnSignBitsUnpack](./docs/aclnnSignBitsUnpack.md)接口方式调用SignBitsUnpack算子。 |
74+ 
75+## 贡献说明
76+ 
77+ 
78+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
79+| ----------- | ---------- | -------------- | --------- | ---------------------------- |
80+| ilovescrapy | 个人开发者 | SignBitsUnpack | 2026/4/21 | SignBitsUnpack算子适配开源仓 |
Aexperimental/math/sign_bits_unpack/docs/aclnnSignBitsUnpack.md+268-0
@@ -0,0 +1,268 @@
1+# aclnnSignBitsUnpack
2+ 
3+## 产品支持情况
4+ 
5+ 
6+| 产品 | 是否支持 |
7+| :------------------------------------------------------- | :------: |
8+| <term>Ascend 950PR/Ascend 950DT</term> | × |
9+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | x |
10+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
11+| <term>Atlas 200I/500 A2 推理产品</term> | × |
12+| <term>Atlas 推理系列产品</term> | × |
13+| <term>Atlas 训练系列产品</term> | × |
14+ 
15+## 功能说明
16+ 
17+- 算子功能：将uint8类型1位Adam拆包为float32或者float16。
18+ 
19+## 函数原型
20+ 
21+每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnSignBitsUnpackGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnSignBitsUnpack”接口执行计算。
22+ 
23+- `aclnnStatus aclnnSignBitsUnpackGetWorkspaceSize(const aclTensor* self, int64_t size, aclDataType dtype, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)`
24+- `aclnnStatus aclnnSignBitsUnpack(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)`
25+ 
26+## aclnnSignBitsUnpackGetWorkspaceSize
27+ 
28+- **参数说明：**
29+ 
30+ - self（aclTensor*, 计算输入）：表示用于计算的1D张量，Device侧的aclTensor，支持空tensor场景，数据类型支持UINT8，支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。
31+ - size（int64_t, 入参）：表示维度处理，Host侧的整型，reshape时输出张量的第一个维度，数据类型支持INT64。
32+ - dtype（aclDataType, 入参）：表示量化输出Tensor的数据类型，支持ACL_FLOAT16、ACL_FLOAT。
33+ - out（aclTensor*, 计算输出）：Device侧的aclTensor，数据类型支持FLOAT16、FLOAT，由dtype参数决定。[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)。
34+ - workspaceSize（uint64_t*, 出参）：返回需要在Device侧申请的workspace大小。
35+ - executor（aclOpExecutor**, 出参）：返回op执行器，包含了算子计算流程。
36+- **返回值：**
37+ 
38+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
39+ 
40+ 第一段接口完成入参校验，出现以下场景时报错：
41+ 
42+ <table style="undefined;table-layout: fixed; width: 1149px"><colgroup>
43+ <col style="width: 300px">
44+ <col style="width: 136px">
45+ <col style="width: 713px">
46+ </colgroup>
47+ <thead>
48+ <tr>
49+ <th>返回值</th>
50+ <th>错误码</th>
51+ <th>描述</th>
52+ </tr></thead>
53+ <tbody>
54+ <tr>
55+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
56+ <td>161001</td>
57+ <td>传入的self或out是空指针。</td>
58+ </tr>
59+ <tr>
60+ <td rowspan="5">ACLNN_ERR_PARAM_INVALID</td>
61+ <td rowspan="5">161002</td>
62+ <td>self、out的数据类型\数据格式不在支持的范围内。</td>
63+ </tr>
64+ <tr>
65+ <td>size小于等于0或者 （self的元素个数）* 8 % size != 0。</td>
66+ </tr>
67+ <tr>
68+ <td>out的数据类型与dtype不一致。</td>
69+ </tr>
70+ <tr>
71+ <td>self的维度不是1维。</td>
72+ </tr>
73+ <tr>
74+ <td>out的第一维度与size不一致。</td>
75+ </tr>
76+ </tbody>
77+ </table>
78+ 
79+## aclnnSignBitsUnpack
80+ 
81+- **参数说明：**
82+ 
83+ <table style="undefined;table-layout: fixed; width: 1149px"><colgroup>
84+ <col style="width: 167px">
85+ <col style="width: 134px">
86+ <col style="width: 848px">
87+ </colgroup>
88+ <thead>
89+ <tr>
90+ <th>参数名</th>
91+ <th>输入/输出</th>
92+ <th>描述</th>
93+ </tr></thead>
94+ <tbody>
95+ <tr>
96+ <td>workspace</td>
97+ <td>输入</td>
98+ <td>在Device侧申请的workspace内存地址。</td>
99+ </tr>
100+ <tr>
101+ <td>workspaceSize</td>
102+ <td>输入</td>
103+ <td>在Device侧申请的workspace大小，由第一段接口aclnnSignBitsUnpackGetWorkspaceSize获取。</td>
104+ </tr>
105+ <tr>
106+ <td>executor</td>
107+ <td>输入</td>
108+ <td>op执行器，包含了算子计算流程。</td>
109+ </tr>
110+ <tr>
111+ <td>stream</td>
112+ <td>输入</td>
113+ <td>指定执行任务的Stream。</td>
114+ </tr>
115+ </tbody>
116+ </table>
117+- **返回值：**
118+ 
119+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
120+ 
121+## 约束说明
122+ 
123+- 确定性计算：
124+ - aclnnSignBitsUnpack默认确定性实现。
125+ 
126+## 调用示例
127+ 
128+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
129+ 
130+```Cpp
131+#include <memory>
132+#include <iostream>
133+#include <vector>
134+#include "acl/acl.h"
135+#include "aclnnop/aclnn_sign_bits_unpack.h"
136+ 
137+#define CHECK_RET(cond, return_expr) \
138+ do { \
139+ if (!(cond)) { \
140+ return_expr; \
141+ } \
142+ } while (0)
143+ 
144+#define LOG_PRINT(message, ...) \
145+ do { \
146+ printf(message, ##__VA_ARGS__); \
147+ } while (0)
148+ 
149+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
150+ int64_t shapeSize = 1;
151+ for (auto i : shape) {
152+ shapeSize *= i;
153+ }
154+ return shapeSize;
155+}
156+ 
157+int Init(int32_t deviceId, aclrtStream* stream) {
158+ // 固定写法，资源初始化
159+ auto ret = aclInit(nullptr);
160+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
161+ ret = aclrtSetDevice(deviceId);
162+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
163+ ret = aclrtCreateStream(stream);
164+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
165+ return 0;
166+}
167+ 
168+template <typename T>
169+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
170+ aclDataType dataType, aclTensor** tensor) {
171+ auto size = GetShapeSize(shape) * sizeof(T);
172+ // 调用aclrtMalloc申请device侧内存
173+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
174+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
175+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
176+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
177+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
178+ 
179+ // 计算连续tensor的strides
180+ std::vector<int64_t> strides(shape.size(), 1);
181+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
182+ strides[i] = shape[i + 1] * strides[i + 1];
183+ }
184+ 
185+ // 调用aclCreateTensor接口创建aclTensor
186+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
187+ shape.data(), shape.size(), *deviceAddr);
188+ return 0;
189+}
190+ 
191+int main() {
192+ // 1. （固定写法）device/stream初始化，参考acl API手册
193+ // 根据自己的实际device填写deviceId
194+ int32_t deviceId = 0;
195+ aclrtStream stream;
196+ auto ret = Init(deviceId, &stream);
197+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
198+ 
199+ // 2. 构造输入与输出，需要根据API的接口自定义构造
200+ std::vector<int64_t> selfShape = {2};
201+ std::vector<int64_t> outShape = {2, 8};
202+ 
203+ int64_t outsize = 2;
204+ aclDataType dataType = ACL_FLOAT;
205+ 
206+ void* selfDeviceAddr = nullptr;
207+ void* outDeviceAddr = nullptr;
208+
209+ aclTensor* self = nullptr;
210+ aclTensor* out = nullptr;
211+ 
212+ std::vector<uint8_t> selfHostData = {159, 15};
213+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
214+ 
215+ // 创建self aclTensor
216+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_UINT8, &self);
217+ CHECK_RET(ret == ACL_SUCCESS, return ret);
218+ // 创建out aclTensor
219+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
220+ CHECK_RET(ret == ACL_SUCCESS, return ret);
221+ 
222+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
223+ uint64_t workspaceSize = 0;
224+ aclOpExecutor* executor;
225+ // 调用aclnnSignBitsUnpack第一段接口
226+ ret = aclnnSignBitsUnpackGetWorkspaceSize(self, outsize, dataType, out, &workspaceSize, &executor);
227+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsUnpackGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
228+ // 根据第一段接口计算出的workspaceSize申请device内存
229+ void* workspaceAddr = nullptr;
230+ if (workspaceSize > 0) {
231+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
232+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
233+ }
234+ // 调用aclnnSignBitsUnpack第二段接口
235+ ret = aclnnSignBitsUnpack(workspaceAddr, workspaceSize, executor, stream);
236+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsUnpack failed. ERROR: %d\n", ret); return ret);
237+ 
238+ // 4. （固定写法）同步等待任务执行结束
239+ ret = aclrtSynchronizeStream(stream);
240+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
241+ 
242+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
243+ auto size = GetShapeSize(outShape);
244+ std::vector<float> resultData(size, 0);
245+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
246+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
247+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
248+ for (int64_t i = 0; i < size; i++) {
249+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
250+ }
251+ 
252+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
253+ aclDestroyTensor(self);
254+ aclDestroyTensor(out);
255+
256+ // 7. 释放device 资源
257+ aclrtFree(selfDeviceAddr);
258+ aclrtFree(outDeviceAddr);
259+ if (workspaceSize > 0) {
260+ aclrtFree(workspaceAddr);
261+ }
262+ aclrtDestroyStream(stream);
263+ aclrtResetDevice(deviceId);
264+ aclFinalize();
265+
266+ return 0;
267+}
268+```
Aexperimental/math/sign_bits_unpack/examples/test_aclnn_sign_bits_unpack.cpp+137-0
@@ -0,0 +1,137 @@
1+#include <memory>
2+#include <iostream>
3+#include <vector>
4+#include "acl/acl.h"
5+#include "aclnnop/aclnn_sign_bits_unpack.h"
6+ 
7+#define CHECK_RET(cond, return_expr) \
8+ do { \
9+ if (!(cond)) { \
10+ return_expr; \
11+ } \
12+ } while (0)
13+ 
14+#define LOG_PRINT(message, ...) \
15+ do { \
16+ printf(message, ##__VA_ARGS__); \
17+ } while (0)
18+ 
19+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
20+ int64_t shapeSize = 1;
21+ for (auto i : shape) {
22+ shapeSize *= i;
23+ }
24+ return shapeSize;
25+}
26+ 
27+int Init(int32_t deviceId, aclrtStream* stream) {
28+ // 固定写法，资源初始化
29+ auto ret = aclInit(nullptr);
30+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
31+ ret = aclrtSetDevice(deviceId);
32+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
33+ ret = aclrtCreateStream(stream);
34+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
35+ return 0;
36+}
37+ 
38+template <typename T>
39+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
40+ aclDataType dataType, aclTensor** tensor) {
41+ auto size = GetShapeSize(shape) * sizeof(T);
42+ // 调用aclrtMalloc申请device侧内存
43+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
44+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
45+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
46+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
47+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
48+ 
49+ // 计算连续tensor的strides
50+ std::vector<int64_t> strides(shape.size(), 1);
51+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
52+ strides[i] = shape[i + 1] * strides[i + 1];
53+ }
54+ 
55+ // 调用aclCreateTensor接口创建aclTensor
56+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
57+ shape.data(), shape.size(), *deviceAddr);
58+ return 0;
59+}
60+ 
61+int main() {
62+ // 1. （固定写法）device/stream初始化，参考acl API手册
63+ // 根据自己的实际device填写deviceId
64+ int32_t deviceId = 0;
65+ aclrtStream stream;
66+ auto ret = Init(deviceId, &stream);
67+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
68+ 
69+ // 2. 构造输入与输出，需要根据API的接口自定义构造
70+ std::vector<int64_t> selfShape = {2};
71+ std::vector<int64_t> outShape = {2, 8};
72+ 
73+ int64_t outsize = 2;
74+ aclDataType dataType = ACL_FLOAT;
75+ 
76+ void* selfDeviceAddr = nullptr;
77+ void* outDeviceAddr = nullptr;
78+
79+ aclTensor* self = nullptr;
80+ aclTensor* out = nullptr;
81+ 
82+ std::vector<uint8_t> selfHostData = {128, 128};
83+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
84+ 
85+ // 创建self aclTensor
86+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_UINT8, &self);
87+ CHECK_RET(ret == ACL_SUCCESS, return ret);
88+ // 创建out aclTensor
89+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
90+ CHECK_RET(ret == ACL_SUCCESS, return ret);
91+ 
92+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
93+ uint64_t workspaceSize = 0;
94+ aclOpExecutor* executor;
95+ // 调用aclnnSignBitsUnpack第一段接口
96+ ret = aclnnSignBitsUnpackGetWorkspaceSize(self, outsize, dataType, out, &workspaceSize, &executor);
97+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsUnpackGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
98+ // 根据第一段接口计算出的workspaceSize申请device内存
99+ void* workspaceAddr = nullptr;
100+ if (workspaceSize > 0) {
101+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
102+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
103+ }
104+ // 调用aclnnSignBitsUnpack第二段接口
105+ ret = aclnnSignBitsUnpack(workspaceAddr, workspaceSize, executor, stream);
106+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsUnpack failed. ERROR: %d\n", ret); return ret);
107+ 
108+ // 4. （固定写法）同步等待任务执行结束
109+ ret = aclrtSynchronizeStream(stream);
110+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
111+ 
112+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
113+ auto size = GetShapeSize(outShape);
114+ std::vector<float> resultData(size, 0);
115+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
116+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
117+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
118+ for (int64_t i = 0; i < size; i++) {
119+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
120+ }
121+ 
122+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
123+ aclDestroyTensor(self);
124+ aclDestroyTensor(out);
125+
126+ // 7. 释放device 资源
127+ aclrtFree(selfDeviceAddr);
128+ aclrtFree(outDeviceAddr);
129+ if (workspaceSize > 0) {
130+ aclrtFree(workspaceAddr);
131+ }
132+ aclrtDestroyStream(stream);
133+ aclrtResetDevice(deviceId);
134+ aclFinalize();
135+
136+ return 0;
137+}
Aexperimental/math/sign_bits_unpack/op_api/aclnn_sign_bits_unpack.cpp+220-0
@@ -0,0 +1,220 @@
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+#include "aclnn_sign_bits_unpack.h"
12+#include "sign_bits_unpack.h"
13+#include "aclnn_kernels/contiguous.h"
14+#include "aclnn_kernels/common/op_error_check.h"
15+#include "aclnn/aclnn_base.h"
16+#include "opdev/common_types.h"
17+#include "opdev/shape_utils.h"
18+#include "opdev/data_type_utils.h"
19+#include "opdev/format_utils.h"
20+#include "opdev/op_dfx.h"
21+#include "opdev/op_executor.h"
22+#include "opdev/op_log.h"
23+#include "opdev/tensor_view_utils.h"
24+#include "opdev/platform.h"
25+#include "opdev/make_op_executor.h"
26+#include "op_api/aclnn_check.h"
27+ 
28+using namespace op;
29+#ifdef __cplusplus
30+extern "C" {
31+#endif
32+ 
33+/* SignBitsUnpack 算子的完整计算流程如下:
34+ * self dtype size
35+ * \ / /
36+ * Contiguous(workspace_0) / /
37+ * \ / /
38+ * SignBitsUnpack(workspace_1)
39+ * |
40+ * ViewCopy
41+ * |
42+ * result
43+ */
44+ 
45+static const size_t DIM_NUM_1D = 1;
46+static const size_t DIM_NUM_2D = 2;
47+ 
48+static const int64_t PACK_SIZE = 8;
49+ 
50+// 根据API定义，需要列出所能支持的所有dtype
51+static const std::initializer_list<op::DataType> SELF_DTYPE_SUPPORT_LIST = {
52+ op::DataType::DT_UINT8};
53+ 
54+static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST = {
55+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16};
56+ 
57+static bool CheckNotNull(const aclTensor *self, const aclTensor *out) {
58+ // 检查输入和输出是否是空指针
59+ OP_CHECK_NULL(self, return false);
60+ OP_CHECK_NULL(out, return false);
61+ return true;
62+}
63+ 
64+static bool CheckDtypeValid(const aclTensor *self, const aclTensor *out, aclDataType dtype) {
65+ // 检查芯片类型是否支持
66+ if(GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_2201) {
67+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "SignBitsUnpack is not supported on this device.");
68+ return false;
69+ }
70+ 
71+ // 检查self的数据类型是否在算子的支持列表内
72+ OP_CHECK_DTYPE_NOT_SUPPORT(self, SELF_DTYPE_SUPPORT_LIST, return false);
73+ 
74+ // 检查out的数据类型是否在算子的支持列表内
75+ OP_CHECK_DTYPE_NOT_SUPPORT(out, OUT_DTYPE_SUPPORT_LIST, return false);
76+ 
77+ // 检查out与dtype的数据类型是否一致
78+ op::DataType dtypeOP = op::ToOpDataType(dtype);
79+ OP_CHECK_DTYPE_NOT_MATCH(out, dtypeOP, return false);
80+ 
81+ return true;
82+}
83+ 
84+static bool CheckFormat(const aclTensor *self, const aclTensor *out) {
85+ // 检查self的format是否为ND
86+ if(self->GetStorageFormat() != Format::FORMAT_ND) {
87+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "self fromat should be ND. Actual: self is [%s].",
88+ op::ToString(self->GetStorageFormat()).GetString());
89+ return false;
90+ }
91+ 
92+ // 检查out的format是否为ND
93+ if(out->GetStorageFormat() != Format::FORMAT_ND) {
94+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "out fromat should be ND. Actual: out is [%s].",
95+ op::ToString(out->GetStorageFormat()).GetString());
96+ return false;
97+ }
98+ 
99+ // 如果输入格式是私有格式，记录日志，直接报错
100+ if(op::IsPrivateFormat(self->GetStorageFormat())) {
101+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Not support format [%s].",
102+ op::ToString(self->GetStorageFormat()).GetString());
103+ return false;
104+ }
105+ 
106+ return true;
107+}
108+ 
109+static inline bool CheckShape(const aclTensor *self, const aclTensor *out) {
110+ // 检查输入shape
111+ size_t selfDimNum = self->GetViewShape().GetDimNum();
112+ if(selfDimNum != DIM_NUM_1D) {
113+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self dims %zu should be 1.", selfDimNum);
114+ return false;
115+ }
116+ 
117+ // 检查输出shape
118+ size_t outDimNum = out->GetViewShape().GetDimNum();
119+ if(outDimNum != DIM_NUM_2D) {
120+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Out dims %zu should be 2.", outDimNum);
121+ return false;
122+ }
123+ return true;
124+}
125+ 
126+static inline bool CheckValue(const aclTensor *self, int64_t size, const aclTensor *out) {
127+ // 检查size的值
128+ if(size <= 0) {
129+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Size value must bigger than zero.");
130+ return false;
131+ }
132+ 
133+ for(size_t i = 0; i < out->GetViewShape().GetDimNum(); i++) {
134+ if(out->GetViewShape().GetDim(i) < 0) {
135+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Dim Value of out is negative.");
136+ return false;
137+ }
138+ }
139+ 
140+ // size可被uint8的拆包输出整除。输出大小为（self的元素个数) * 8
141+ int64_t selfDim = self->GetViewShape().GetDim(0);
142+ if((selfDim * PACK_SIZE) % size != 0) {
143+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The number of elements in self * 8 cannot be divided by size.");
144+ return false;
145+ }
146+ 
147+ // 检查out第一维度是否等于size
148+ int64_t outDimOneNum = out->GetViewShape().GetDim(0);
149+ if(size != outDimOneNum) {
150+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The value of the first dimension of 'out' is incorrect and should be equal to size.");
151+ return false;
152+ }
153+ return true;
154+}
155+ 
156+static inline aclnnStatus CheckParams(const aclTensor *self, aclDataType dtype, int64_t size, const aclTensor *out) {
157+ // 1. 检查参数是否为空指针
158+ CHECK_RET(CheckNotNull(self, out), ACLNN_ERR_PARAM_NULLPTR);
159+ 
160+ // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
161+ CHECK_RET(CheckDtypeValid(self, out, dtype), ACLNN_ERR_PARAM_INVALID);
162+ 
163+ // 3. 检查数据格式是否支持
164+ CHECK_RET(CheckFormat(self, out), ACLNN_ERR_PARAM_INVALID);
165+ 
166+ // 4. 检查入参数值是否合法
167+ CHECK_RET(CheckValue(self, size, out), ACLNN_ERR_PARAM_INVALID);
168+ 
169+ // 5. 检查输入和输出tensor的shape是否为异常，输入必须1维
170+ CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID);
171+ 
172+ return ACLNN_SUCCESS;
173+}
174+ 
175+aclnnStatus aclnnSignBitsUnpackGetWorkspaceSize(const aclTensor* self, int64_t size, aclDataType dtype, aclTensor* out,
176+ uint64_t* workspaceSize, aclOpExecutor** executor) {
177+ L2_DFX_PHASE_1(aclnnSignBitsUnpack, DFX_IN(self, size, dtype), DFX_OUT(out));
178+ 
179+ // 固定写法，创建OpExecutor
180+ auto uniqueExecutor = CREATE_EXECUTOR();
181+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
182+ 
183+ // 固定写法，参数检查
184+ auto ret = CheckParams(self, dtype, size, out);
185+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
186+ 
187+ // 输入为空tensor时，直接返回dtype类型的空tensor
188+ if (self->IsEmpty() || out->IsEmpty()) {
189+ *workspaceSize = 0;
190+ uniqueExecutor.ReleaseTo(executor);
191+ return ACLNN_SUCCESS;
192+ }
193+ 
194+ // 固定写法，将输入self转换成连续的tensor
195+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
196+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
197+ 
198+ // 调用SignBitsUnpack算子kernel,将输入self的数据类型转换成指定的数据类型
199+ auto castOut = l0op::SignBitsUnpack(selfContiguous, size, op::ToOpDataType(dtype), uniqueExecutor.get());
200+ CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
201+ 
202+ // 固定写法，将计算结果拷贝到输出out上，out可能是非连续的tensor
203+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
204+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
205+ 
206+ // 固定写法，获取计算过程中需要使用的workspace大小
207+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
208+ uniqueExecutor.ReleaseTo(executor);
209+ return ACLNN_SUCCESS;
210+}
211+ 
212+aclnnStatus aclnnSignBitsUnpack(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream) {
213+ L2_DFX_PHASE_2(aclnnSignBitsUnpack);
214+ // 固定写法，调用框架能力，完成计算
215+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
216+}
217+ 
218+#ifdef __cplusplus
219+}
220+#endif
Aexperimental/math/sign_bits_unpack/op_api/aclnn_sign_bits_unpack.h+53-0
@@ -0,0 +1,53 @@
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+#ifndef OP_API_INC_SIGN_BITS_UNPACK_H_
11+#define OP_API_INC_SIGN_BITS_UNPACK_H_
12+ 
13+#include "aclnn/aclnn_base.h"
14+#include "aclnn_util.h"
15+ 
16+#ifdef __cplusplus
17+extern "C" {
18+#endif
19+ 
20+/**
21+ * @brief aclnnSignBitsUnpack的第一段接口，根据具体的计算流程，计算workspace大小。
22+ * @domain aclnnop_ops_train
23+ *
24+ * 算子功能：将uint8类型1位Adam拆包为float32或者float16。
25+ *
26+ * @param [in] self:
27+ * device侧的aclTensor，数据类型支持UINT8，支持空Tensor。支持非连续的Tensor，数据类型支持UINT8，数据格式支持ND。
28+ * @param [in] size: host侧的int64_t，reshape时输出张量的第一个维度。
29+ * @param [in] dtype: host侧的aclDataType，表示量化输出Tensor的数据类型，数据类型支持ACL_FLOAT16、ACL_FLOAT。
30+ * @param [in] out:
31+ * device侧的aclTensor，数据类型支持FLOAT16、FLOAT。数据类型由dtype决定，支持非连续的Tensor，数据格式支持ND。
32+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
33+ * @param [out] executor: 返回op执行器，包含算子计算流程。
34+ * @return aclnnStatus: 返回状态码。
35+ */
36+ACLNN_API aclnnStatus aclnnSignBitsUnpackGetWorkspaceSize(const aclTensor* self, int64_t size, aclDataType dtype, aclTensor* out,
37+ uint64_t* workspaceSize, aclOpExecutor** executor);
38+/**
39+ * @brief aclnnSignBitsUnpack的第二段接口，用于执行计算。
40+ *
41+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
42+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnSignBitsUnpackGetWorkspaceSize获取。
43+ * @param [in] executor: op执行器，包含了算子计算流程。
44+ * @param [in] stream: acl stream流。
45+ * @return aclnnStatus: 返回状态码。
46+ */
47+ACLNN_API aclnnStatus aclnnSignBitsUnpack(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
48+ 
49+#ifdef __cplusplus
50+}
51+#endif
52+ 
53+#endif // OP_API_INC_SIGN_BITS_UNPACK_H_
Aexperimental/math/sign_bits_unpack/op_api/sign_bits_unpack.cpp+52-0
@@ -0,0 +1,52 @@
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+#include "sign_bits_unpack.h"
12+#include "opdev/make_op_executor.h"
13+#include "opdev/op_def.h"
14+#include "opdev/op_dfx.h"
15+#include "opdev/op_executor.h"
16+#include "opdev/shape_utils.h"
17+ 
18+ 
19+using namespace op;
20+ 
21+namespace l0op {
22+OP_TYPE_REGISTER(SignBitsUnpack);
23+static constexpr size_t OUT_DIM = 2;
24+static constexpr size_t OUT_SIZE = 8;
25+ 
26+const aclTensor *SignBitsUnpack(const aclTensor *self, int64_t size, op::DataType dtype, aclOpExecutor *executor) {
27+ L0_DFX(SignBitsUnpack, self, size, dtype);
28+ 
29+ int64_t selfDimOne = self->GetViewShape().GetDim(0);
30+ op::Shape outShape;
31+ outShape.SetDimNum(OUT_DIM);
32+ outShape.SetDim(0, size);
33+ 
34+ if(size <= 0) {
35+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "size is zero or less than zero, cannot perform division.");
36+ return nullptr;
37+ }
38+ outShape.SetDim(1, (selfDimOne * OUT_SIZE) / size);
39+ 
40+ auto out = executor->AllocTensor(outShape, dtype, op::Format::FORMAT_ND);
41+ if(out == nullptr) {
42+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "alloc out tensor failed.");
43+ return nullptr;
44+ }
45+ 
46+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(SignBitsUnpack, OP_INPUT(self), OP_OUTPUT(out), OP_ATTR(size, dtype));
47+ OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "SignBitsUnpack ADD_TO_LAUNCHER_LIST_AICORE failed."),
48+ return nullptr);
49+ 
50+ return out;
51+}
52+} // namespace l0op
Aexperimental/math/sign_bits_unpack/op_api/sign_bits_unpack.h+20-0
@@ -0,0 +1,20 @@
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+#ifndef OP_API_INC_LEVEL0_SIGN_BITS_UNPACK_H_
12+#define OP_API_INC_LEVEL0_SIGN_BITS_UNPACK_H_
13+ 
14+#include "opdev/op_executor.h"
15+ 
16+namespace l0op {
17+const aclTensor *SignBitsUnpack(const aclTensor *self, int64_t size, op::DataType dtype, aclOpExecutor *executor);
18+} // namespace l0op
19+ 
20+#endif // OP_API_INC_LEVEL0_SIGN_BITS_UNPACK_H_
Aexperimental/math/sign_bits_unpack/op_host/sign_bits_unpack_def.cpp+37-0
@@ -0,0 +1,37 @@
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+ * \file sign_bits_unpack_def.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class SignBitsUnpack : public OpDef {
19+public:
20+ explicit SignBitsUnpack(const char* name) : OpDef(name)
21+ {
22+ this->Input("self")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_UINT8, ge::DT_UINT8})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
27+ this->Output("y")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
32+ 
33+ this->AICore().AddConfig("ascend910b"); // 其他的soc版本补充部分配置项
34+ }
35+};
36+OP_ADD(SignBitsUnpack); // 添加算子信息库
37+} // namespace ops
Aexperimental/math/sign_bits_unpack/op_host/sign_bits_unpack_infershape.cpp+43-0
@@ -0,0 +1,43 @@
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+ * \file sign_bits_unpack_infershape.cpp
13+ * \brief
14+*/
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+ 
20+namespace ops {
21+static constexpr int64_t IDX_0 = 0;
22+ 
23+static ge::graphStatus InferShapeSignBitsUnpack(gert::InferShapeContext* context)
24+{
25+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
26+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeSignBitsUnpack");
27+ 
28+ // get input shapes
29+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
30+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
31+ 
32+ // get output shapes
33+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
34+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
35+ 
36+ // 填充输出shape大小
37+ *yShape = *xShape;
38+ OP_LOGD(context->GetNodeName(), "End to do InferShapeSignBitsUnpack");
39+ return GRAPH_SUCCESS;
40+}
41+ 
42+IMPL_OP_INFERSHAPE(SignBitsUnpack).InferShape(InferShapeSignBitsUnpack);
43+}
Aexperimental/math/sign_bits_unpack/op_host/sign_bits_unpack_tiling.cpp+211-0
@@ -0,0 +1,211 @@
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+ * \file sign_bits_unpack_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "util/platform_util.h"
19+#include "op_host/tiling_util.h"
20+#include "op_host/tiling_templates_registry.h"
21+#include "tiling/platform/platform_ascendc.h"
22+#include "register/op_impl_registry.h"
23+#include "../op_kernel/sign_bits_unpack_tiling_data.h"
24+#include "../op_kernel/sign_bits_unpack_tiling_key.h"
25+ 
26+namespace optiling {
27+ 
28+using namespace Ops::Math::OpTiling;
29+ 
30+#define UB_NUM 2U
31+#define BLOCK_SIZE 64U
32+constexpr uint32_t BUFFER_NUM = 2;
33+constexpr uint32_t WS_SYS_SIZE = 0;
34+ 
35+struct SignBitsUnpackCompileInfo {};
36+ 
37+static ge::graphStatus TilingParseForSignBitsUnpack([[maybe_unused]] gert::TilingParseContext* context)
38+{
39+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
40+ return ge::GRAPH_SUCCESS;
41+}
42+ 
43+// 获取平台信息如ubSize, coreNum
44+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
45+{
46+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
47+ // 获取ubsize coreNum
48+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
49+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
50+ coreNum = ascendcPlatform.GetCoreNum();
51+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
52+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
53+ return ge::GRAPH_SUCCESS;
54+}
55+ 
56+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
57+{
58+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
59+ size_t usrSize = 0;
60+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
61+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
62+ size_t* currentWorkspace = context->GetWorkspaceSizes(
63+ 1); // 通过框架获取workspace的指针，GetWorkspaceSizes入参为所需workspace的块数。当前限制使用一块。
64+ currentWorkspace[0] = usrSize + sysWorkspaceSize;
65+ return ge::GRAPH_SUCCESS;
66+}
67+ 
68+static ge::graphStatus GetShapeAttrsInfo(
69+ gert::TilingContext* context, uint64_t ubSize, int64_t coreNum, uint32_t& typeLengthOut, uint64_t& bufferOpen, uint64_t& inputNum, uint64_t& inputBytes, uint64_t& tileBlockNum,
70+ uint64_t& tileDataNum, uint64_t& inputLengthAlgin)
71+{
72+ OP_CHECK_IF(
73+ context == nullptr || context->GetInputShape(0) == nullptr, OP_LOGE(context, "context is nullptr"),
74+ return ge::GRAPH_FAILED);
75+ inputNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
76+ uint32_t typeLength = 0;
77+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
78+ typeLengthOut = 0;
79+ ge::TypeUtils::GetDataTypeLength(context->GetOutputDesc(0)->GetDataType(), typeLengthOut);
80+ uint64_t inputLength = inputNum * typeLength;
81+ if (inputNum == 0) {
82+ OP_LOGE(context, "inputNum is 0");
83+ return ge::GRAPH_FAILED;
84+ }
85+ inputBytes = inputLength / inputNum;
86+ uint64_t ubDataNumber = UB_NUM;
87+ bufferOpen = 1;
88+ inputLengthAlgin = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE);
89+ // 重点
90+ //half uint8 + half * 8
91+ // 1 + 2 * 8 = 17
92+ uint64_t ubSizeTmp = ubSize - (((ubSize + 17 - 1) / 17) * 16);
93+ if (typeLengthOut == 4) {
94+ //额外使用一块half
95+ //1+ 2*8 + 4*8
96+ ubSizeTmp = ubSize - (((ubSize + 49 - 1) / 49) * 48);
97+ }
98+ if (inputLengthAlgin < coreNum * (((ubSizeTmp / BLOCK_SIZE) * BLOCK_SIZE) / 1U)) {
99+ ubDataNumber = 1U;
100+ bufferOpen = 0;
101+ }
102+ 
103+ // double buffer需要重新计算输入所用的UB空间
104+ if (bufferOpen == 1) {
105+ ubSizeTmp = ubSize - (((ubSize + 34 - 1) / 34) * 32);
106+ if (typeLengthOut == 4) {
107+ //1*2 + 2*8 + 4*8*2
108+ ubSizeTmp = ubSize - (((ubSize + 82 - 1) / 82) * 80);
109+ }
110+ }
111+ if (ubDataNumber == 0 || BLOCK_SIZE == 0) {
112+ OP_LOGE(context, "ubDataNumber or BLOCK_SIZE is 0");
113+ return ge::GRAPH_FAILED;
114+ }
115+ tileBlockNum = (ubSizeTmp / BLOCK_SIZE) / ubDataNumber;
116+ if (inputBytes == 0) {
117+ OP_LOGE(context, "inputBytes is 0");
118+ return ge::GRAPH_FAILED;
119+ }
120+ tileDataNum = (tileBlockNum * BLOCK_SIZE) / inputBytes;
121+ return ge::GRAPH_SUCCESS;
122+}
123+ 
124+static ge::graphStatus CalculateCoreBlockNums(
125+ gert::TilingContext* context, uint64_t inputLengthAlgin, int64_t coreNum, uint64_t tileBlockNum, uint64_t inputBytes,
126+ uint64_t tileDataNum, uint64_t& smallCoreDataNum, uint64_t& bigCoreDataNum, uint64_t& smallTailDataNum, uint64_t& bigTailDataNum,
127+ uint64_t& finalSmallTileNum, uint64_t& finalBigTileNum, uint64_t& tailBlockNum)
128+{
129+ if (0 == BLOCK_SIZE || 0 == coreNum || 0 == tileBlockNum || 0 == inputBytes) {
130+ OP_LOGE(context, "BLOCK_SIZE or coreNum or tileBlockNum or inputBytes is 0");
131+ return ge::GRAPH_FAILED;
132+ }
133+ uint64_t everyCoreInputBlockNum = inputLengthAlgin / BLOCK_SIZE / coreNum;
134+ tailBlockNum = (inputLengthAlgin / BLOCK_SIZE) % coreNum;
135+ smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
136+ uint64_t smallTileNum = everyCoreInputBlockNum / tileBlockNum;
137+ finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
138+ smallTailDataNum = smallCoreDataNum - (tileDataNum * smallTileNum);
139+ smallTailDataNum = smallTailDataNum == 0 ? tileDataNum : smallTailDataNum;
140+ 
141+ everyCoreInputBlockNum += 1;
142+ bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
143+ uint64_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
144+ finalBigTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
145+ bigTailDataNum = bigCoreDataNum - tileDataNum * bigTileNum;
146+ bigTailDataNum = bigTailDataNum == 0 ? tileDataNum : bigTailDataNum;
147+ 
148+ return ge::GRAPH_SUCCESS;
149+}
150+ 
151+// tiling 分发入口
152+static ge::graphStatus SignBitsUnpackTilingFunc(gert::TilingContext* context)
153+{
154+ // 1、获取平台运行信息
155+ uint64_t ubSize;
156+ int64_t coreNum;
157+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, coreNum);
158+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
159+ // 2、获取shape、属性信息
160+ uint64_t bufferOpen;
161+ uint32_t typeLengthOut;
162+ uint64_t inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin;
163+ ret = GetShapeAttrsInfo(context, ubSize, coreNum, typeLengthOut, bufferOpen, inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin);
164+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
165+ // 3、获取WorkspaceSize信息
166+ OP_CHECK_IF(
167+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
168+ return ge::GRAPH_FAILED);
169+ // 4、设置tiling信息
170+ SignBitsUnpackTilingData* tiling = context->GetTilingData<SignBitsUnpackTilingData>();
171+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
172+ OP_CHECK_IF(
173+ memset_s(tiling, sizeof(SignBitsUnpackTilingData), 0, sizeof(SignBitsUnpackTilingData)) != EOK,
174+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
175+ 
176+ if (tileDataNum >= inputNum) {
177+ coreNum = 1;
178+ } else {
179+ coreNum = (static_cast<uint64_t>(coreNum) < inputLengthAlgin / BLOCK_SIZE) ? coreNum : inputLengthAlgin / BLOCK_SIZE;
180+ }
181+ // 计算每个core处理的数据块数
182+ uint64_t smallCoreDataNum, bigCoreDataNum, smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum;
183+ ret = CalculateCoreBlockNums(
184+ context, inputLengthAlgin, coreNum, tileBlockNum, inputBytes, tileDataNum, smallCoreDataNum, bigCoreDataNum,
185+ smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum);
186+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "CalculateCoreBlockNums error"), return ge::GRAPH_FAILED);
187+ // 设置tiling数据
188+ tiling->smallCoreDataNum = static_cast<uint64_t>(smallCoreDataNum);
189+ tiling->bigCoreDataNum = static_cast<uint64_t>(bigCoreDataNum);
190+ tiling->tileDataNum = static_cast<uint64_t>(tileDataNum);
191+ tiling->smallTailDataNum = static_cast<uint64_t>(smallTailDataNum);
192+ tiling->bigTailDataNum = static_cast<uint64_t>(bigTailDataNum);
193+ tiling->finalSmallTileNum = static_cast<uint64_t>(finalSmallTileNum);
194+ tiling->finalBigTileNum = static_cast<uint64_t>(finalBigTileNum);
195+ tiling->tailBlockNum = static_cast<uint64_t>(tailBlockNum);
196+ 
197+ tiling->bufferOpen = static_cast<uint64_t>(bufferOpen);
198+ 
199+ context->SetBlockDim(coreNum);
200+ uint64_t tilingKey = 0;
201+ tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0);
202+ if (typeLengthOut == 2) {
203+ tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_1);
204+ }
205+ context->SetTilingKey(tilingKey);
206+ return ge::GRAPH_SUCCESS;
207+}
208+ 
209+// tiling注册入口.
210+IMPL_OP_OPTILING(SignBitsUnpack).Tiling(SignBitsUnpackTilingFunc).TilingParse<SignBitsUnpackCompileInfo>(TilingParseForSignBitsUnpack);
211+} // namespace optiling
Aexperimental/math/sign_bits_unpack/op_kernel/sign_bits_unpack.cpp+34-0
@@ -0,0 +1,34 @@
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+ * \file sign_bits_unpack.cpp
13+ * \brief
14+ */
15+ 
16+#include "sign_bits_unpack.h"
17+ 
18+template <uint32_t schMode>
19+__global__ __aicore__ void sign_bits_unpack(GM_ADDR self, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling)
20+{
21+ REGISTER_TILING_DEFAULT(SignBitsUnpackTilingData);
22+ GET_TILING_DATA_WITH_STRUCT(SignBitsUnpackTilingData, tilingData, tiling);
23+ if (TILING_KEY_IS(1)) {
24+ NsSignBitsUnpack::KernelSignBitsUnpack<half> op; // 算子kernel实例获取
25+ op.Init(self, out, tilingData.smallCoreDataNum, tilingData.bigCoreDataNum, tilingData.finalBigTileNum, tilingData.finalSmallTileNum, tilingData.tileDataNum,
26+ tilingData.smallTailDataNum, tilingData.bigTailDataNum, tilingData.tailBlockNum, tilingData.bufferOpen);
27+ op.Process();
28+ } else {
29+ NsSignBitsUnpack::KernelSignBitsUnpack<float> op; // 算子kernel实例获取
30+ op.Init(self, out, tilingData.smallCoreDataNum, tilingData.bigCoreDataNum, tilingData.finalBigTileNum, tilingData.finalSmallTileNum, tilingData.tileDataNum,
31+ tilingData.smallTailDataNum, tilingData.bigTailDataNum, tilingData.tailBlockNum, tilingData.bufferOpen);
32+ op.Process();
33+ }
34+}
Aexperimental/math/sign_bits_unpack/op_kernel/sign_bits_unpack.h+157-0
@@ -0,0 +1,157 @@
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+ * \file sign_bits_unpack.h
13+ * \brief
14+ */
15+#ifndef SIGN_BITS_UNPACK_H_
16+#define SIGN_BITS_UNPACK_H_
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "sign_bits_unpack_tiling_data.h"
21+#include "sign_bits_unpack_tiling_key.h"
22+ 
23+#include "kernel_operator.h"
24+ 
25+constexpr uint64_t DOUBLE_BUFFER = 2;
26+constexpr uint64_t SINGLE_BUFFER = 1;
27+ 
28+namespace NsSignBitsUnpack {
29+ 
30+using namespace AscendC;
31+ 
32+template <typename T>
33+class KernelSignBitsUnpack {
34+ 
35+public:
36+ __aicore__ inline KernelSignBitsUnpack(){};
37+ 
38+ __aicore__ inline void Init(GM_ADDR self, GM_ADDR out, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
39+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum, uint64_t tailBlockNum, uint64_t bufferOpen);
40+ __aicore__ inline void Process();
41+ 
42+private:
43+ __aicore__ inline void CopyIn(int64_t progress);
44+ __aicore__ inline void CopyOut(int64_t progress);
45+ __aicore__ inline void Compute(int64_t progress);
46+ 
47+private:
48+ AscendC::TPipe pipe;
49+ AscendC::TQue<AscendC::TPosition::VECIN, DOUBLE_BUFFER> inQueueSelf;
50+ AscendC::TQue<AscendC::TPosition::VECOUT, DOUBLE_BUFFER> outQueueOut;
51+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpQueue0;
52+ 
53+ AscendC::GlobalTensor<uint8_t> selfGm;
54+ AscendC::GlobalTensor<T> outGm;
55+ uint64_t coreDataNum = 0;
56+ uint64_t tileNum = 0;
57+ uint64_t tileDataNum = 0;
58+ uint64_t tailDataNum = 0;
59+ uint64_t tileDataNumOut = 0;
60+ uint64_t processDataNumIn = 0;
61+ uint64_t processDataNumOut = 0;
62+};
63+ 
64+template <typename T>
65+__aicore__ inline void KernelSignBitsUnpack<T>::Init(GM_ADDR self, GM_ADDR out, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
66+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum, uint64_t tailBlockNum, uint64_t bufferOpen)
67+{
68+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
69+ uint64_t coreId = AscendC::GetBlockIdx();
70+ uint64_t globalBufferIndex = bigCoreDataNum * coreId;
71+ uint64_t outCoreIndx = bigCoreDataNum * coreId * 8;
72+ this->tileDataNum = tileDataNum;
73+ if (coreId < tailBlockNum) {
74+ this->coreDataNum = bigCoreDataNum;
75+ this->tileNum = finalBigTileNum;
76+ this->tailDataNum = bigTailDataNum;
77+ } else {
78+ this->coreDataNum = smallCoreDataNum;
79+ this->tileNum = finalSmallTileNum;
80+ this->tailDataNum = smallTailDataNum;
81+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (coreId - tailBlockNum);
82+ outCoreIndx -= 8 * (bigCoreDataNum - smallCoreDataNum) * (coreId - tailBlockNum);
83+ }
84+ uint64_t BUFFER_NUM = DOUBLE_BUFFER;
85+ if (bufferOpen == 0) {
86+ BUFFER_NUM = SINGLE_BUFFER;
87+ }
88+ uint64_t outCoreNum = this->coreDataNum * 8;
89+ this->tileDataNumOut = this->tileDataNum * 8;
90+ selfGm.SetGlobalBuffer((__gm__ uint8_t*)self + globalBufferIndex, this->coreDataNum);
91+ outGm.SetGlobalBuffer((__gm__ T*)out + outCoreIndx, outCoreNum);
92+ pipe.InitBuffer(inQueueSelf, BUFFER_NUM, this->tileDataNum * sizeof(uint8_t));
93+ pipe.InitBuffer(outQueueOut, BUFFER_NUM, this->tileDataNumOut * sizeof(T));
94+ if constexpr (std::is_same_v<T, float>) {
95+ pipe.InitBuffer(tmpQueue0, this->tileDataNumOut * sizeof(half));
96+ }
97+}
98+ 
99+template <typename T>
100+__aicore__ inline void KernelSignBitsUnpack<T>::CopyIn(int64_t progress)
101+{
102+ AscendC::LocalTensor<uint8_t> selfLocal = inQueueSelf.template AllocTensor<uint8_t>();
103+ AscendC::DataCopy(selfLocal, selfGm[progress * this->tileDataNum], this->processDataNumIn);
104+ inQueueSelf.EnQue(selfLocal);
105+}
106+ 
107+template <typename T>
108+__aicore__ inline void KernelSignBitsUnpack<T>::CopyOut(int64_t progress)
109+{
110+ AscendC::LocalTensor<T> outLocal = outQueueOut.template DeQue<T>();
111+ AscendC::DataCopy(outGm[progress * this->tileDataNumOut], outLocal, this->processDataNumOut);
112+ outQueueOut.FreeTensor(outLocal);
113+}
114+ 
115+template <typename T>
116+__aicore__ inline void KernelSignBitsUnpack<T>::Compute(int64_t progress)
117+{
118+ if constexpr (std::is_same_v<T, float>) {
119+ AscendC::LocalTensor<uint8_t> selfLocal = inQueueSelf.template DeQue<uint8_t>();
120+ AscendC::LocalTensor<float> outLocal = outQueueOut.template AllocTensor<float>();
121+ AscendC::LocalTensor<half> tmp0Local = tmpQueue0.AllocTensor<half>();
122+ AscendC::Duplicate(outLocal, static_cast<float>(1.0), this->processDataNumOut);
123+ AscendC::Select(outLocal, selfLocal, outLocal, static_cast<float>(-1.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNumOut);
124+ outQueueOut.template EnQue<float>(outLocal);
125+ inQueueSelf.FreeTensor(selfLocal);
126+ } else {
127+ AscendC::LocalTensor<uint8_t> selfLocal = inQueueSelf.template DeQue<uint8_t>();
128+ AscendC::LocalTensor<half> outLocal = outQueueOut.template AllocTensor<half>();
129+ AscendC::Duplicate(outLocal, static_cast<half>(1.0), this->processDataNumOut);
130+ //数据对齐
131+ AscendC::Select(outLocal, selfLocal, outLocal, static_cast<half>(1.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNumOut);
132+ AscendC::Select(outLocal, selfLocal, outLocal, static_cast<half>(-1.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNumOut);
133+ outQueueOut.template EnQue<half>(outLocal);
134+ inQueueSelf.FreeTensor(selfLocal);
135+ }
136+}
137+ 
138+template <typename T>
139+__aicore__ inline void KernelSignBitsUnpack<T>::Process()
140+{
141+ int32_t loopCount = this->tileNum;
142+ this->processDataNumIn = this->tileDataNum;
143+ this->processDataNumOut = this->processDataNumIn * 8;
144+ for (int64_t i = 0; i < loopCount - 1; i++) {
145+ CopyIn(i);
146+ Compute(i);
147+ CopyOut(i);
148+ }
149+ this->processDataNumIn = this->tailDataNum;
150+ this->processDataNumOut = this->processDataNumIn * 8;
151+ CopyIn(loopCount - 1);
152+ Compute(loopCount - 1);
153+ CopyOut(loopCount - 1);
154+}
155+ 
156+} // namespace NsSignBitsUnpack
157+#endif // SIGN_BITS_UNPACK_H
Aexperimental/math/sign_bits_unpack/op_kernel/sign_bits_unpack_tiling_data.h+30-0
@@ -0,0 +1,30 @@
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+ * \file sign_bits_unpack_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef SIGN_BITS_UNPACK_TILING_DATA_H_
17+#define SIGN_BITS_UNPACK_TILING_DATA_H_
18+ 
19+struct SignBitsUnpackTilingData {
20+ uint64_t smallCoreDataNum;
21+ uint64_t bigCoreDataNum;
22+ uint64_t finalBigTileNum;
23+ uint64_t finalSmallTileNum;
24+ uint64_t tileDataNum;
25+ uint64_t smallTailDataNum;
26+ uint64_t bigTailDataNum;
27+ uint64_t tailBlockNum;
28+ uint64_t bufferOpen;
29+};
30+#endif
Aexperimental/math/sign_bits_unpack/op_kernel/sign_bits_unpack_tiling_key.h+31-0
@@ -0,0 +1,31 @@
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+ * \file sign_bits_unpack_tiling_key.h
13+ * \brief sign_bits_unpack tiling key declare
14+ */
15+ 
16+#ifndef __SIGN_BITS_UNPACK_TILING_KEY_H__
17+#define __SIGN_BITS_UNPACK_TILING_KEY_H__
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+#define ELEMENTWISE_TPL_SCH_MODE_0 0
22+#define ELEMENTWISE_TPL_SCH_MODE_1 1
23+ 
24+ASCENDC_TPL_ARGS_DECL(
25+ SignBitsUnpack,
26+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
27+ 
28+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
29+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));
30+ 
31+#endif
Aexperimental/math/sign_bits_unpack/tests/ut/op_host/sign_bits_unpack_tiling.h+29-0
@@ -0,0 +1,29 @@
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+ * \file sign_bits_unpack_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef BITWISE_AND_TILING_H
17+#define BITWISE_AND_TILING_H
18+ 
19+#include "register/tilingdata_base.h"
20+ 
21+namespace optiling {
22+struct SignBitsUnpackCompileInfo {
23+ int32_t totalCoreNum = 0;
24+ int64_t ubSize = 0;
25+ bool isRegbase = false;
26+};
27+} // namespace optiling
28+ 
29+#endif // BITWISE_AND_TILING_H
Aexperimental/math/sign_bits_unpack/tests/ut/op_host/test_sign_bits_unpack_tiling.cpp+51-0
@@ -0,0 +1,51 @@
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+#include <iostream>
12+#include <gtest/gtest.h>
13+#include "sign_bits_unpack_tiling.h"
14+#include "../../../op_kernel/sign_bits_unpack_tiling_data.h"
15+#include "../../../op_kernel/sign_bits_unpack_tiling_key.h"
16+#include "tiling_context_faker.h"
17+#include "tiling_case_executor.h"
18+ 
19+using namespace std;
20+using namespace optiling;
21+ 
22+class SignBitsUnpackTiling : public testing::Test {
23+protected:
24+ static void SetUpTestCase()
25+ {
26+ cout << "SignBitsUnpackTiling SetUp" << endl;
27+ }
28+ 
29+ static void TearDownTestCase()
30+ {
31+ cout << "SignBitsUnpackTiling TearDown " << endl;
32+ }
33+};
34+ 
35+TEST_F(SignBitsUnpackTiling, ascend9101_test_tiling_FLOAT_001)
36+{
37+ optiling::SignBitsUnpackCompileInfo compileInfo = {40, 196608, false};
38+ gert::TilingContextPara tilingContextPara(
39+ "SignBitsUnpack",
40+ {
41+ {{{128}, {128}}, ge::DT_UINT8, ge::FORMAT_ND},
42+ },
43+ {
44+ {{{1024}, {1024}}, ge::DT_FLOAT, ge::FORMAT_ND},
45+ },
46+ &compileInfo);
47+ uint64_t expectTilingKey = 0;
48+ string expectTilingData = "128 192 1 1 5312 128 192 0 0 ";
49+ std::vector<size_t> expectWorkspaces = {16777216};
50+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
51+}
Aexperimental/math/sign_bits_unpack/tests/ut/op_kernel/CMakeLists.txt+29-0
@@ -0,0 +1,29 @@
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+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
13+ # 需要将Tiling依赖的文件添加到CMakeLists.txt中
14+ # set(elewise_common_tiling_files
15+ # ${CANN_ROOT}/ops/built-in/op_tiling/runtime/elewise_tiling.cc
16+ # )
17+ # 算子自己的tiling文件路径
18+ set(sign_bits_unpack_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/sign_bits_unpack_tiling.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/sign_bits_unpack_infershape.cpp
21+ # ${elewise_common_tiling_files}
22+ )
23+ # 使用AddOpTestCase
24+ # param1：算子名称，以kernel方式命名
25+ # param2：soc版本，多个以分号分隔，例如："ascend950pr_9599;AscendB1"
26+ # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔，例如："-DDTYPE_X=float -DDTYPE_Y=float"
27+ # param4：该算子依赖的所有tiling源码文件
28+ AddOpTestCase(sign_bits_unpack "ascend910b" "-DDTYPE_SELF=float_t" "${sign_bits_unpack_tiling_files}")
29+endif()
Aexperimental/math/sign_bits_unpack/tests/ut/op_kernel/sign_bits_unpack_data/compare_data.py+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+import sys
12+import numpy as np
13+import glob
14+import os
15+ 
16+curr_dir = os.path.dirname(os.path.realpath(__file__))
17+ 
18+def compare_data(golden_file_lists, output_file_lists, d_type):
19+ if d_type == "float":
20+ np_dtype = np.float32
21+ elif d_type == "float16":
22+ np_dtype = np.float16
23+ else:
24+ raise ValueError("d_type must be float16 or float")
25+
26+ data_same = True
27+ for gold, out in zip(golden_file_lists, output_file_lists):
28+ tmp_out = np.fromfile(out, np_dtype)
29+ tmp_gold = np.fromfile(gold, np_dtype)
30+ diff_res = np.isclose(tmp_out, tmp_gold, 0, 0, True)
31+ diff_idx = np.where(diff_res != True)[0]
32+ if len(diff_idx) == 0:
33+ print("PASSED!")
34+ else:
35+ print("FAILED!")
36+ for idx in diff_idx[:5]:
37+ print(f"index: {idx}, output: {tmp_out[idx]}, golden: {tmp_gold[idx]}")
38+ data_same = False
39+ return data_same
40+ 
41+def get_file_lists(dtype):
42+ golden_file_lists = sorted(glob.glob(curr_dir + "/*golden*.bin"))
43+ output_file_lists = sorted(glob.glob(curr_dir + "/*output*.bin"))
44+ return golden_file_lists, output_file_lists
45+ 
46+def process(d_type):
47+ golden_file_lists, output_file_lists = get_file_lists(d_type)
48+ result = compare_data(golden_file_lists, output_file_lists, d_type)
49+ print("compare result:", result)
50+ return result
51+ 
52+if __name__ == '__main__':
53+ ret = process(sys.argv[1])
54+ exit(0 if ret else 1)
Aexperimental/math/sign_bits_unpack/tests/ut/op_kernel/sign_bits_unpack_data/gen_data.py+45-0
@@ -0,0 +1,45 @@
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+import sys
12+import os
13+import numpy as np
14+import re
15+ 
16+def parse_str_to_shape_list(shape_str):
17+ shape_str = shape_str.strip('(').strip(')')
18+ shape_list = [int(x) for x in shape_str.split(",")]
19+ return np.array(shape_list)
20+ 
21+ 
22+def gen_data_and_golden(shape_str, d_type="uint8"):
23+ d_type_dict = {
24+ "uint8": np.uint8
25+ }
26+ np_type = d_type_dict[d_type]
27+ shape = parse_str_to_shape_list(shape_str)
28+ size = np.prod(shape)
29+
30+ input_self = np.random.randint(0, 10, shape).astype(np_type)
31+ 
32+ golden = np.unpackbits(input_self, bitorder='little').astype(np.float32)
33+ golden[golden == 0] = -1
34+ golden = golden.astype(np.float32)
35+ input_self.astype(np_type).tofile("uint8_input_self_sign_bits_unpack.bin")
36+ golden.astype(np.float32).tofile("float_golden_t_bitwiese_and.bin")
37+ 
38+ 
39+if __name__ == "__main__":
40+ if len(sys.argv) != 3:
41+ print("Param num must be 3.")
42+ exit(1)
43+ # 清理bin文件
44+ os.system("rm -rf *.bin")
45+ gen_data_and_golden(sys.argv[1], sys.argv[2])
Aexperimental/math/sign_bits_unpack/tests/ut/op_kernel/sign_bits_unpack_tiling.h+59-0
@@ -0,0 +1,59 @@
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+ * \file sign_bits_unpack_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef _SIGN_BITS_UNPACK_TILING_H_
17+#define _SIGN_BITS_UNPACK_TILING_H_
18+ 
19+#include <cstdint>
20+ 
21+#include "../../../op_kernel/sign_bits_unpack_tiling_data.h"
22+#include "kernel_tiling/kernel_tiling.h"
23+#include "../../../op_kernel/sign_bits_unpack.h"
24+ 
25+using namespace NsSignBitsUnpack;
26+ 
27+#define __aicore__
28+#ifdef __NPU_TILING__
29+inline[aicore] void InitTilingData(const __gm__ uint8_t *tiling, SignBitsUnpackTilingData *constData)
30+{
31+ const __gm__ uint32_t *src = (const __gm__ uint32_t *)tiling;
32+ uint32_t *dst = (uint32_t *)constData;
33+ for (size_t i = 0; i < sizeof(SignBitsUnpackTilingData) / 4; i++) {
34+ *(dst + i) = *(src + i);
35+ }
36+}
37+#else
38+inline void InitTilingData(uint8_t *tiling, SignBitsUnpackTilingData *constData)
39+{
40+ memcpy(constData, tiling, sizeof(SignBitsUnpackTilingData));
41+}
42+#endif // __NPU_TILING__
43+ 
44+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
45+ __ubuf__ tilingStruct *tilingDataPointer = \
46+ reinterpret_cast<__ubuf__ tilingStruct *>((__ubuf__ uint8_t *)(tilingPointer));
47+ 
48+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
49+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
50+ 
51+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingArg) \
52+ tilingStruct tilingData; \
53+ InitTilingData(tilingArg, &tilingData)
54+ 
55+#define GET_TILING_DATA(tilingData, tilingArg) \
56+ SignBitsUnpackTilingData tilingData; \
57+ InitTilingData(tilingArg, &tilingData)
58+ 
59+#endif // _SIGN_BITS_UNPACK_TILING_H_
Aexperimental/math/sign_bits_unpack/tests/ut/op_kernel/test_sign_bits_unpack.cpp+110-0
@@ -0,0 +1,110 @@
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+#include <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+ 
15+#include <array>
16+#include <vector>
17+#include <iostream>
18+#include <string>
19+#include <cstdint>
20+#include "gtest/gtest.h"
21+#include "tikicpulib.h"
22+#include "data_utils.h"
23+ 
24+#include "../../../op_kernel/sign_bits_unpack.cpp"
25+ 
26+using namespace std;
27+ 
28+constexpr uint32_t smallCoreDataNum = 128;
29+constexpr uint32_t bigCoreDataNum = 160;
30+constexpr uint32_t tileDataNum = 2048;
31+constexpr uint32_t smallTailDataNum = 128;
32+constexpr uint32_t bigTailDataNum = 160;
33+ 
34+extern "C" __global__ __aicore__ void sign_bits_unpack(GM_ADDR self, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling);
35+ 
36+class SignBitsUnpackTest : public testing::Test {
37+protected:
38+ static void SetUpTestCase()
39+ {
40+ std::cout << "sign_bits_unpack_test SetUp" << std::endl;
41+ const string cmd = "cp -rf " + dataPath + " ./";
42+ system(cmd.c_str());
43+ system("chmod -R 755 ./sign_bits_unpack_data/");
44+ }
45+ static void TearDownTestCase()
46+ {
47+ std::cout << "sign_bits_unpack_test TearDown" << std::endl;
48+ }
49+ 
50+private:
51+ const static std::string rootPath;
52+ const static std::string dataPath;
53+};
54+ 
55+const std::string SignBitsUnpackTest::rootPath = "../../../../experimental/";
56+const std::string SignBitsUnpackTest::dataPath = rootPath + "math/sign_bits_unpack/tests/ut/op_kernel/sign_bits_unpack_data";
57+ 
58+template <typename T1, typename T2>
59+inline T1 CeilAlign(T1 a, T2 b)
60+{
61+ return (a + b - 1) / b * b;
62+}
63+ 
64+TEST_F(SignBitsUnpackTest, test_case_float_1)
65+{
66+ uint32_t blockDim = 1;
67+ system("cd ./sign_bits_unpack_data/ && python3 gen_data.py '(128)' 'uint8'");
68+ uint32_t dataCount = 1024;
69+ uint32_t dataCountIn = 128;
70+ size_t inputByteSize = dataCountIn * sizeof(uint8_t);
71+ 
72+ std::string self_fileName = "./sign_bits_unpack_data/uint8_input_self_sign_bits_unpack.bin";
73+ 
74+ uint8_t* self = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
75+ 
76+ ReadFile(self_fileName, inputByteSize, self, inputByteSize);
77+ 
78+ size_t outputByteSize = dataCount * sizeof(float);
79+ uint8_t* out = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
80+ 
81+ size_t workspaceSize = 32 * 1024 * 1024;
82+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
83+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SignBitsUnpackTilingData));
84+ 
85+ SignBitsUnpackTilingData* tilingData = reinterpret_cast<SignBitsUnpackTilingData*>(tiling);
86+ 
87+ tilingData->smallCoreDataNum = smallCoreDataNum;
88+ tilingData->bigCoreDataNum = bigCoreDataNum;
89+ tilingData->tileDataNum = tileDataNum;
90+ tilingData->smallTailDataNum = smallTailDataNum;
91+ tilingData->bigTailDataNum = bigTailDataNum;
92+ tilingData->finalSmallTileNum = 1;
93+ tilingData->finalBigTileNum = 1;
94+ tilingData->tailBlockNum = 0;
95+ tilingData->bufferOpen = 0;
96+ 
97+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
98+ auto func = sign_bits_unpack<ELEMENTWISE_TPL_SCH_MODE_0>;
99+ ICPU_RUN_KF(func, blockDim, self, out, workspace, (uint8_t*)(tilingData));
100+ 
101+ std::string fileName = "./sign_bits_unpack_data/float_output_t_sign_bits_unpack.bin";
102+ WriteFile(fileName, out, outputByteSize);
103+ 
104+ AscendC::GmFree((void*)(self));
105+ AscendC::GmFree((void*)(out));
106+ AscendC::GmFree((void*)workspace);
107+ AscendC::GmFree((void*)tiling);
108+ 
109+ system("cd ./sign_bits_unpack_data/ && python3 compare_data.py 'float'");
110+}