Star1328
1442
代码
代码
Issues109
Pull Requests237
Actions
项目讨论
Wiki
分析
Star1328
1442
已合并
【社区任务】AscendC实现SignBitsPack算子贡献 #2494
infinity_infi创建于 4月26日
【社区任务】AscendC实现SignBitsPack算子贡献 #2494
已合并
infinity_infi创建于 4月26日
从infinity_infi/ops-mathmaster合入到cann/ops-mathmaster
共 23 个文件变更+1894-0
Aexperimental/math/sign_bits_pack/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_pack ACLNNTYPE aclnn_exclude)
Aexperimental/math/sign_bits_pack/README.md+72-0
@@ -0,0 +1,72 @@
1+# SignBitsPack
2+ 
3+## 产品支持情况
4+ 
5+ 
6+| 产品 | 是否支持 |
7+| ------------------------------------------------------------------ | :------: |
8+| Atlas A2 训练系列产品/Atlas 800I A2 推理产品 | √ |
9+ 
10+## 功能说明
11+ 
12+- 算子功能：将float16类型或者float32类型的输入的符号位打包为uint8。
13+ 
14+## 算子原型
15+ 
16+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
17+ <col style="width: 100px">
18+ <col style="width: 150px">
19+ <col style="width: 280px">
20+ <col style="width: 330px">
21+ <col style="width: 120px">
22+ </colgroup>
23+ <thead>
24+ <tr>
25+ <th>参数名</th>
26+ <th>输入/输出/属性</th>
27+ <th>描述</th>
28+ <th>数据类型</th>
29+ <th>数据格式</th>
30+ </tr></thead>
31+ <tbody>
32+ <tr>
33+ <td>x</td>
34+ <td>输入</td>
35+ <td>入参x,计算的1D张量</td>
36+ <td>float16,float</td>
37+ <td>ND</td>
38+ </tr>
39+ <tr>
40+ <td>size</td>
41+ <td>属性</td>
42+ <td>表示处理维度，reshape时输出张量的第一个维度</td>
43+ <td>int</td>
44+ <td>scale</td>
45+ </tr>
46+ <tr>
47+ <tr>
48+ <td>y</td>
49+ <td>输出</td>
50+ <td>出参y, 只支持二维, x元素个数不被8整除时为（x元素个数 // 8） + 1，在被8整除时为（x元素个数/8）</td>
51+ <td>uint8</td>
52+ <td>ND</td>
53+ </tr>
54+ </tbody></table>
55+ 
56+## 约束说明
57+ 
58+无
59+ 
60+## 调用说明
61+ 
62+ 
63+| 调用方式 | 调用样例 | 说明 |
64+| --------- | ------------------------------------------------------------- | ------------------------------------------------------------------------------------------- |
65+| aclnn调用 | [test_aclnn_sign_bits_pack](./examples/test_aclnn_sign_bits_pack.cpp) | 通过[aclnnSignBitsPack](./docs/aclnnSignBitsPack.md)接口方式调用SignBitsPack算子。 |
66+ 
67+## 贡献说明
68+ 
69+ 
70+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
71+| ----------- | ---------- | -------- | --------- | ---------------------- |
72+| infinity | 个人开发者 | SignBitsPack | 2026/4/26 | SignBitsPack算子适配开源仓 |
Aexperimental/math/sign_bits_pack/docs/aclnnSignBitsPack.md+349-0
@@ -0,0 +1,349 @@
1+# aclnnSignBitsPack
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
8+ 
9+## 功能说明
10+ 
11+将float16类型或者float32类型的1位Adam打包为uint8。
12+ 
13+## 函数原型
14+ 
15+每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnSignBitsPackGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnSignBitsPack”接口执行计算。
16+ 
17+- `aclnnStatus aclnnSignBitsPackGetWorkspaceSize(const aclTensor* self, int64_t size, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)`
18+ 
19+- `aclnnStatus aclnnSignBitsPack(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)`
20+ 
21+## aclnnSignBitsPackGetWorkspaceSize
22+ 
23+- **参数说明**：
24+ 
25+ <table style="undefined;table-layout: fixed; width: 1550px"><colgroup>
26+ <col style="width: 170px">
27+ <col style="width: 120px">
28+ <col style="width: 271px">
29+ <col style="width: 330px">
30+ <col style="width: 223px">
31+ <col style="width: 101px">
32+ <col style="width: 190px">
33+ <col style="width: 145px">
34+ </colgroup>
35+ <thead>
36+ <tr>
37+ <th>参数名</th>
38+ <th>输入/输出</th>
39+ <th>描述</th>
40+ <th>使用说明</th>
41+ <th>数据类型</th>
42+ <th>数据格式</th>
43+ <th>维度(shape)</th>
44+ <th>非连续Tensor</th>
45+ </tr></thead>
46+ <tbody>
47+ <tr>
48+ <td>self</td>
49+ <td>输入</td>
50+ <td>表示用于计算的1D张量。</td>
51+ <td><ul><li>不支持空Tensor。</li><li>只支持1D输入，不支持Nan值输入。</li></ul></td>
52+ <td>FLOAT32、FLOAT16</td>
53+ <td>ND</td>
54+ <td>1</td>
55+ <td>√</td>
56+ </tr>
57+ <tr>
58+ <td>size</td>
59+ <td>输入</td>
60+ <td>表示处理维度，reshape时输出张量的第一个维度。</td>
61+ <td>-</td>
62+ <td>int64_t</td>
63+ <td>-</td>
64+ <td>-</td>
65+ <td>-</td>
66+ </tr>
67+ <tr>
68+ <td>out</td>
69+ <td>输出</td>
70+ <td>表示计算的输出张量。对应公式中的`out`。</td>
71+ <td><ul><li>不支持空Tensor。</li><li>只支持二维。</li><li>out总长度在self元素个数不被8整除时为（self元素个数 // 8） + 1，在被8整除时为（self元素个数/8）。</li></ul></td>
72+ <td>UINT8</td>
73+ <td>ND</td>
74+ <td>2</td>
75+ <td>√</td>
76+ </tr>
77+ <tr>
78+ <td>workspaceSize</td>
79+ <td>输出</td>
80+ <td>返回需要在Device侧申请的workspace大小。</td>
81+ <td>-</td>
82+ <td>-</td>
83+ <td>-</td>
84+ <td>-</td>
85+ <td>-</td>
86+ </tr>
87+ <tr>
88+ <td>executor</td>
89+ <td>输出</td>
90+ <td>返回op执行器，包含了算子计算流程。</td>
91+ <td>-</td>
92+ <td>-</td>
93+ <td>-</td>
94+ <td>-</td>
95+ <td>-</td>
96+ </tr>
97+ </tbody>
98+ </table>
99+ 
100+- **返回值**：
101+ 
102+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
103+ 
104+ 第一段接口完成入参校验，出现以下场景时报错：
105+ 
106+ <table style="undefined;table-layout: fixed; width: 1148px"><colgroup>
107+ <col style="width: 300px">
108+ <col style="width: 136px">
109+ <col style="width: 712px">
110+ </colgroup>
111+ <thead>
112+ <tr>
113+ <th>返回值</th>
114+ <th>错误码</th>
115+ <th>描述</th>
116+ </tr></thead>
117+ <tbody>
118+ <tr>
119+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
120+ <td>161001</td>
121+ <td>传入的self或out是空指针。</td>
122+ </tr>
123+ <tr>
124+ <td rowspan="4">ACLNN_ERR_PARAM_INVALID</td>
125+ <td rowspan="4">161002</td>
126+ <td>传入的self或out的数据类型/数据格式不在支持的范围之内。</td>
127+ </tr>
128+ <tr>
129+ <td>self的维度不是一维。</td>
130+ </tr>
131+ <tr>
132+ <td>size小于0。</td>
133+ </tr>
134+ <tr>
135+ <td>out输出总长度无法整除size。</td>
136+ </tr>
137+ </tbody>
138+ </table>
139+ 
140+## aclnnSignBitsPack
141+ 
142+- **参数说明**：
143+ 
144+ <table style="undefined;table-layout: fixed; width: 1149px"><colgroup>
145+ <col style="width: 167px">
146+ <col style="width: 134px">
147+ <col style="width: 848px">
148+ </colgroup>
149+ <thead>
150+ <tr>
151+ <th>参数名</th>
152+ <th>输入/输出</th>
153+ <th>描述</th>
154+ </tr></thead>
155+ <tbody>
156+ <tr>
157+ <td>workspace</td>
158+ <td>输入</td>
159+ <td>在Device侧申请的workspace内存地址。</td>
160+ </tr>
161+ <tr>
162+ <td>workspaceSize</td>
163+ <td>输入</td>
164+ <td>在Device侧申请的workspace大小，由第一段接口aclnnSignBitsPackGetWorkspaceSize获取。</td>
165+ </tr>
166+ <tr>
167+ <td>executor</td>
168+ <td>输入</td>
169+ <td>op执行器，包含了算子计算流程。</td>
170+ </tr>
171+ <tr>
172+ <td>stream</td>
173+ <td>输入</td>
174+ <td>指定执行任务的Stream。</td>
175+ </tr>
176+ </tbody>
177+ </table>
178+ 
179+- **返回值：**
180+ 
181+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
182+ 
183+## 约束说明
184+ 
185+- 确定性计算：
186+ - aclnnSignBitsPack默认确定性实现。
187+ - size的值不能大于out输出总长度。
188+ 
189+## 调用示例
190+ 
191+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
192+ 
193+```Cpp
194+#include <iostream>
195+#include <memory>
196+#include <vector>
197+#include "acl/acl.h"
198+#include "aclnnop/aclnn_sign_bits_pack.h"
199+ 
200+#define CHECK_RET(cond, return_expr) \
201+ do { \
202+ if (!(cond)) { \
203+ return_expr; \
204+ } \
205+ } while (0)
206+ 
207+#define CHECK_FREE_RET(cond, return_expr) \
208+ do { \
209+ if (!(cond)) { \
210+ Finalize(deviceId, stream); \
211+ return_expr; \
212+ } \
213+ } while (0)
214+ 
215+#define LOG_PRINT(message, ...) \
216+ do { \
217+ printf(message, ##__VA_ARGS__); \
218+ } while (0)
219+ 
220+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
221+ int64_t shapeSize = 1;
222+ for (auto i : shape) {
223+ shapeSize *= i;
224+ }
225+ return shapeSize;
226+}
227+ 
228+int Init(int32_t deviceId, aclrtStream* stream) {
229+ // 固定写法，资源初始化
230+ auto ret = aclInit(nullptr);
231+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
232+ ret = aclrtSetDevice(deviceId);
233+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
234+ ret = aclrtCreateStream(stream);
235+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
236+ return 0;
237+}
238+ 
239+template <typename T>
240+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
241+ aclDataType dataType, aclTensor** tensor) {
242+ auto size = GetShapeSize(shape) * sizeof(T);
243+ // 调用aclrtMalloc申请device侧内存
244+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
245+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
246+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
247+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
248+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
249+ 
250+ // 计算连续tensor的strides
251+ std::vector<int64_t> strides(shape.size(), 1);
252+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
253+ strides[i] = shape[i + 1] * strides[i + 1];
254+ }
255+ 
256+ // 调用aclCreateTensor接口创建aclTensor
257+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
258+ shape.data(), shape.size(), *deviceAddr);
259+ return 0;
260+}
261+ 
262+void Finalize(int32_t deviceId, aclrtStream& stream)
263+{
264+ aclrtDestroyStream(stream);
265+ aclrtResetDevice(deviceId);
266+ aclFinalize();
267+}
268+ 
269+int aclnnSignBitsPackTest(int32_t deviceId, aclrtStream& stream) {
270+ auto ret = Init(deviceId, &stream);
271+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
272+ 
273+ // 2. 构造输入与输出，需要根据API的接口自定义构造
274+ int64_t outsize = 2;
275+
276+ std::vector<int64_t> selfShape = {14};
277+ std::vector<int64_t> outShape = {2,1};
278+ 
279+ void* selfDeviceAddr = nullptr;
280+ void* outDeviceAddr = nullptr;
281+ 
282+ aclTensor* self = nullptr;
283+ aclTensor* out = nullptr;
284+ 
285+ std::vector<float> selfHostData{5, 4, 3, 2, 0, -1, -2, 4, 3, 2, 1, 0, -1, -2};
286+ std::vector<uint8_t> outHostData{0, 0};
287+ 
288+ // 创建self aclTensor
289+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
290+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor *)> selfTensorPtr(self, aclDestroyTensor);
291+ std::unique_ptr<void, aclError (*)(void *)> selfDeviceAddrPtr(selfDeviceAddr, aclrtFree);
292+ CHECK_FREE_RET(ret == ACL_SUCCESS, return ret);
293+ 
294+ // 创建out aclTensor
295+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_UINT8, &out);
296+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor *)> outTensorPtr(out, aclDestroyTensor);
297+ std::unique_ptr<void, aclError (*)(void *)> outDeviceAddrPtr(outDeviceAddr, aclrtFree);
298+ CHECK_FREE_RET(ret == ACL_SUCCESS, return ret);
299+ 
300+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
301+ uint64_t workspaceSize = 0;
302+ aclOpExecutor* executor;
303+ // 调用aclnnSignBitsPack第一段接口
304+ ret = aclnnSignBitsPackGetWorkspaceSize(self,
305+ outsize,
306+ out,
307+ &workspaceSize,
308+ &executor);
309+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsPackGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
310+ // 根据第一段接口计算出的workspaceSize申请device内存
311+ void* workspaceAddr = nullptr;
312+ std::unique_ptr<void, aclError (*)(void *)> workspaceAddrPtr(nullptr, aclrtFree);
313+ if (workspaceSize > 0) {
314+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
315+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
316+ workspaceAddrPtr.reset(workspaceAddr);
317+ }
318+ // 调用aclnnSignBitsPack第二段接口
319+ ret = aclnnSignBitsPack(workspaceAddr, workspaceSize, executor, stream);
320+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsPack failed. ERROR: %d\n", ret); return ret);
321+ 
322+ // 4. （固定写法）同步等待任务执行结束
323+ ret = aclrtSynchronizeStream(stream);
324+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
325+ 
326+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
327+ auto size = GetShapeSize(outShape);
328+ std::vector<uint8_t> outData(size, 0);
329+ ret = aclrtMemcpy(outData.data(), outData.size() * sizeof(outData[0]), outDeviceAddr,
330+ size * sizeof(outData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
331+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
332+ for (int64_t i = 0; i < size; i++) {
333+ LOG_PRINT("out result[%ld] is: %u\n", i, outData[i]);
334+ }
335+ return ACL_SUCCESS;
336+}
337+ 
338+int main() {
339+ // 1. （固定写法）device/stream初始化，参考acl API手册
340+ // 根据自己的实际device填写deviceId
341+ int32_t deviceId = 0;
342+ aclrtStream stream;
343+ auto ret = aclnnSignBitsPackTest(deviceId, stream);
344+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsPackTest failed. ERROR: %d\n", ret); return ret);
345+ 
346+ Finalize(deviceId, stream);
347+ return 0;
348+}
349+```
Aexperimental/math/sign_bits_pack/examples/test_aclnn_sign_bits_pack.cpp+165-0
@@ -0,0 +1,165 @@
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 <memory>
13+#include <vector>
14+#include "acl/acl.h"
15+#include "aclnnop/aclnn_sign_bits_pack.h"
16+ 
17+#define CHECK_RET(cond, return_expr) \
18+ do { \
19+ if (!(cond)) { \
20+ return_expr; \
21+ } \
22+ } while (0)
23+ 
24+#define CHECK_FREE_RET(cond, return_expr) \
25+ do { \
26+ if (!(cond)) { \
27+ Finalize(deviceId, stream); \
28+ return_expr; \
29+ } \
30+ } while (0)
31+ 
32+#define LOG_PRINT(message, ...) \
33+ do { \
34+ printf(message, ##__VA_ARGS__); \
35+ } while (0)
36+ 
37+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
38+ int64_t shapeSize = 1;
39+ for (auto i : shape) {
40+ shapeSize *= i;
41+ }
42+ return shapeSize;
43+}
44+ 
45+int Init(int32_t deviceId, aclrtStream* stream) {
46+ // 固定写法，资源初始化
47+ auto ret = aclInit(nullptr);
48+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
49+ ret = aclrtSetDevice(deviceId);
50+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
51+ ret = aclrtCreateStream(stream);
52+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
53+ return 0;
54+}
55+ 
56+template <typename T>
57+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
58+ aclDataType dataType, aclTensor** tensor) {
59+ auto size = GetShapeSize(shape) * sizeof(T);
60+ // 调用aclrtMalloc申请device侧内存
61+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
62+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
63+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
64+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
65+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
66+ 
67+ // 计算连续tensor的strides
68+ std::vector<int64_t> strides(shape.size(), 1);
69+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
70+ strides[i] = shape[i + 1] * strides[i + 1];
71+ }
72+ 
73+ // 调用aclCreateTensor接口创建aclTensor
74+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
75+ shape.data(), shape.size(), *deviceAddr);
76+ return 0;
77+}
78+ 
79+void Finalize(int32_t deviceId, aclrtStream& stream)
80+{
81+ aclrtDestroyStream(stream);
82+ aclrtResetDevice(deviceId);
83+ aclFinalize();
84+}
85+ 
86+int aclnnSignBitsPackTest(int32_t deviceId, aclrtStream& stream) {
87+ auto ret = Init(deviceId, &stream);
88+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
89+ 
90+ // 2. 构造输入与输出，需要根据API的接口自定义构造
91+ int64_t outsize = 2;
92+
93+ std::vector<int64_t> selfShape = {14};
94+ std::vector<int64_t> outShape = {2,1};
95+ 
96+ void* selfDeviceAddr = nullptr;
97+ void* outDeviceAddr = nullptr;
98+ 
99+ aclTensor* self = nullptr;
100+ aclTensor* out = nullptr;
101+ 
102+ std::vector<float> selfHostData{5, 4, 3, 2, 0, -1, -2, 4, 3, 2, 1, 0, -1, -2};
103+ std::vector<uint8_t> outHostData{0, 0};
104+ 
105+ // 创建self aclTensor
106+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
107+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor *)> selfTensorPtr(self, aclDestroyTensor);
108+ std::unique_ptr<void, aclError (*)(void *)> selfDeviceAddrPtr(selfDeviceAddr, aclrtFree);
109+ CHECK_FREE_RET(ret == ACL_SUCCESS, return ret);
110+ 
111+ // 创建out aclTensor
112+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_UINT8, &out);
113+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor *)> outTensorPtr(out, aclDestroyTensor);
114+ std::unique_ptr<void, aclError (*)(void *)> outDeviceAddrPtr(outDeviceAddr, aclrtFree);
115+ CHECK_FREE_RET(ret == ACL_SUCCESS, return ret);
116+ 
117+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
118+ uint64_t workspaceSize = 0;
119+ aclOpExecutor* executor;
120+ // 调用aclnnSignBitsPack第一段接口
121+ ret = aclnnSignBitsPackGetWorkspaceSize(self,
122+ outsize,
123+ out,
124+ &workspaceSize,
125+ &executor);
126+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsPackGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
127+ // 根据第一段接口计算出的workspaceSize申请device内存
128+ void* workspaceAddr = nullptr;
129+ std::unique_ptr<void, aclError (*)(void *)> workspaceAddrPtr(nullptr, aclrtFree);
130+ if (workspaceSize > 0) {
131+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
132+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
133+ workspaceAddrPtr.reset(workspaceAddr);
134+ }
135+ // 调用aclnnSignBitsPack第二段接口
136+ ret = aclnnSignBitsPack(workspaceAddr, workspaceSize, executor, stream);
137+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsPack failed. ERROR: %d\n", ret); return ret);
138+ 
139+ // 4. （固定写法）同步等待任务执行结束
140+ ret = aclrtSynchronizeStream(stream);
141+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
142+ 
143+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
144+ auto size = GetShapeSize(outShape);
145+ std::vector<uint8_t> outData(size, 0);
146+ ret = aclrtMemcpy(outData.data(), outData.size() * sizeof(outData[0]), outDeviceAddr,
147+ size * sizeof(outData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
148+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
149+ for (int64_t i = 0; i < size; i++) {
150+ LOG_PRINT("out result[%ld] is: %u\n", i, outData[i]);
151+ }
152+ return ACL_SUCCESS;
153+}
154+ 
155+int main() {
156+ // 1. （固定写法）device/stream初始化，参考acl API手册
157+ // 根据自己的实际device填写deviceId
158+ int32_t deviceId = 0;
159+ aclrtStream stream;
160+ auto ret = aclnnSignBitsPackTest(deviceId, stream);
161+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSignBitsPackTest failed. ERROR: %d\n", ret); return ret);
162+ 
163+ Finalize(deviceId, stream);
164+ return 0;
165+}
Aexperimental/math/sign_bits_pack/op_api/aclnn_sign_bits_pack.cpp+205-0
@@ -0,0 +1,205 @@
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_pack.h"
12+#include "sign_bits_pack.h"
13+#include "opdev/common_types.h"
14+#include "opdev/data_type_utils.h"
15+#include "opdev/format_utils.h"
16+#include "opdev/op_dfx.h"
17+#include "opdev/op_executor.h"
18+#include "opdev/op_log.h"
19+#include "opdev/shape_utils.h"
20+#include "opdev/tensor_view_utils.h"
21+#include "opdev/framework_op.h"
22+#include "op_api/aclnn_check.h"
23+#include "op_api/op_api_def.h"
24+#include "aclnn_kernels/common/op_error_check.h"
25+#include "opdev/make_op_executor.h"
26+#include "aclnn_kernels/contiguous.h"
27+ 
28+using namespace op;
29+#ifdef __cplusplus
30+extern "C" {
31+#endif
32+ 
33+static const int64_t DIM_NUM_1D = 1;
34+static const int64_t DIM_NUM_2D = 2;
35+static const int64_t PACK_SIZE = 8;
36+ 
37+static const std::initializer_list<op::DataType> SELF_DTYPE_SUPPORT_LIST = {
38+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16};
39+ 
40+static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST = {
41+ op::DataType::DT_UINT8};
42+ 
43+static bool CheckNotNullPtr(const aclTensor *self, aclTensor *out)
44+{
45+ OP_CHECK_NULL(self, return false);
46+ OP_CHECK_NULL(out, return false);
47+ return true;
48+}
49+ 
50+static inline bool CheckSocVersionIsSupport(void) {
51+ return GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910B ||
52+ GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND310P ||
R
Rrxtfeng4月27日

支持的硬件类型和aclnn文档中内容不一致，代码中支持3种芯片类型，文档只写了一种

likedislike
I
infinity_infi
4月28日 评论：
F
fulltower
5月9日 评论：
53+ GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_93;
54+}
55+ 
56+static bool CheckDtypeValid(const aclTensor *self, const aclTensor *out)
57+{
58+ bool isSupport = CheckSocVersionIsSupport();
59+ if(!isSupport) {
60+ auto socVersion = GetCurrentPlatformInfo().GetSocVersion();
61+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
62+ "SignBitsPack is unsupported by the current SOC version [%s]",
63+ op::ToString(socVersion).GetString());
64+ return false;
65+ }
66+ // 检查self的数据类型是否在支持列表内
67+ OP_CHECK_DTYPE_NOT_SUPPORT(self, SELF_DTYPE_SUPPORT_LIST, return false);
68+ // 检查out的数据类型是否在支持列表内
69+ OP_CHECK_DTYPE_NOT_SUPPORT(out, OUT_DTYPE_SUPPORT_LIST, return false);
70+ 
71+ return true;
72+}
73+ 
74+static bool CheckFormat(const aclTensor *self, const aclTensor *out)
75+{
76+ // 检查self的format是否为ND
77+ if (self->GetStorageFormat() != Format::FORMAT_ND) {
78+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "self format should be ND. Actual: self is [%s].",
79+ op::ToString(self->GetStorageFormat()).GetString());
80+ return false;
81+ }
82+ 
83+ // 检查out的format是否为ND
84+ if (out->GetStorageFormat() != Format::FORMAT_ND) {
85+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "out format should be ND. Actual: out is [%s].",
86+ op::ToString(out->GetStorageFormat()).GetString());
87+ return false;
88+ }
89+ 
90+ // 如果输入格式是私有格式，记录日志，直接报错
L
Lloov4月27日

前面已经检查了，只支持ND格式，这里还有必要再检查吗？

likedislike
I
infinity_infi
4月28日 评论：
91+ if (op::IsPrivateFormat(self->GetStorageFormat())) {
92+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Not support format [%s].", op::ToString(self->GetStorageFormat()).GetString());
93+ return false;
94+ }
95+ 
96+ return true;
97+}
98+ 
99+static bool CheckShape(const aclTensor *self, const aclTensor *out)
100+{
101+ size_t selfdimNum = self->GetViewShape().GetDimNum();
102+ if (selfdimNum != DIM_NUM_1D) {
103+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self dims %zu should be 1.", selfdimNum);
104+ return false;
105+ }
106+ 
107+ size_t outdimNum = out->GetViewShape().GetDimNum();
108+ if (outdimNum != DIM_NUM_2D) {
109+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Out dims %zu should be 2.", outdimNum);
110+ return false;
111+ }
112+ return true;
113+}
114+ 
115+static bool CheckValue(const aclTensor *self, int64_t size, const aclTensor *out)
116+{
117+ for (size_t i = 0; i < out->GetViewShape().GetDimNum(); i++) {
118+ if (out->GetViewShape().GetDim(i) < 0) {
119+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Dim value of out is negative.");
120+ return false;
121+ }
122+ }
123+ 
124+ size_t selfdim = self->GetViewShape().GetDim(0);
125+ auto ysize = (selfdim + 7) / 8;
126+ if(size <= 0)
127+ {
128+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "size value must bigger zero.");
129+ return false;
130+ }
131+ 
132+ if(size != 0 && ysize % size != 0){
L
Lloov4月27日

前面的检查已经保证size !=0了，这里可以不用检查size !=0

likedislike
I
infinity_infi
4月28日 评论：
133+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "all must need be divisible by size");
134+ return false;
135+ }
136+ return true;
137+}
138+ 
139+static aclnnStatus CheckParams(const aclTensor *self, int64_t size, aclTensor *out)
140+{
141+ // 检查参数是否为空指针
142+ CHECK_RET(CheckNotNullPtr(self, out), ACLNN_ERR_PARAM_NULLPTR);
143+ 
144+ // 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
145+ CHECK_RET(CheckDtypeValid(self, out), ACLNN_ERR_PARAM_INVALID);
146+ 
147+ // 检查数据格式是否支持
148+ CHECK_RET(CheckFormat(self, out), ACLNN_ERR_PARAM_INVALID);
149+ 
150+ // 检查参数值是否合法
151+ CHECK_RET(CheckValue(self, size, out), ACLNN_ERR_PARAM_INVALID);
152+ 
153+ // 检查数据维度是否合法
154+ CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID);
155+ 
156+ return ACLNN_SUCCESS;
157+}
158+ 
159+aclnnStatus aclnnSignBitsPackGetWorkspaceSize(const aclTensor *self, int64_t size, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor)
160+{
161+ L2_DFX_PHASE_1(aclnnSignBitsPack, DFX_IN(self, size), DFX_OUT(out));
162+ 
163+ auto uniqueExecutor = CREATE_EXECUTOR();
164+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
165+ 
166+ // 固定写法，参数检查
167+ auto ret = CheckParams(self, size, out);
168+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
169+ 
170+ // 如果是空tensor，直接返回
171+ if (self->IsEmpty() || out->IsEmpty()) {
L
Lloov4月27日

建议空tensot的检查扩在参数检查的前面

likedislike
I
infinity_infi
4月28日 评论：
172+ *workspaceSize = 0;
173+ uniqueExecutor.ReleaseTo(executor);
174+ return ACLNN_SUCCESS;
175+ }
176+ 
177+ // 固定写法，将输入self转换成连续的tensor
178+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
179+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
180+ 
181+ const aclTensor *calcOut = nullptr;
182+ calcOut = l0op::SignBitsPack(selfContiguous, size, uniqueExecutor.get());
183+ CHECK_RET(calcOut != nullptr, ACLNN_ERR_PARAM_NULLPTR);
184+ 
185+ auto viewCopyResult = l0op::ViewCopy(calcOut, out, uniqueExecutor.get());
186+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
187+ 
188+ // 固定写法，获取计算过程中需要使用的workspace大小
189+ 
190+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
191+ uniqueExecutor.ReleaseTo(executor); // 需要把 uniqueExecutor持有executor转移给executor
192+ 
193+ return ACLNN_SUCCESS;
194+}
195+ 
196+aclnnStatus aclnnSignBitsPack(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)
197+{
198+ L2_DFX_PHASE_2(aclnnSignBitsPack);
199+ // 固定写法，调用框架能力，完成计算
200+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
201+}
202+ 
203+#ifdef __cplusplus
204+}
205+#endif
Aexperimental/math/sign_bits_pack/op_api/aclnn_sign_bits_pack.h+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+#ifndef OP_API_INC_SIGNBITSPACK_H_
11+#define OP_API_INC_SIGNBITSPACK_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 aclnnSignBitsPack的第一段接口，根据具体的计算流程，计算workspace大小。
22+ * @domain aclnnop_ops_train
23+ *
24+ * 算子功能：将float16类型或者float32类型的1位Adam打包为uint8。
25+ *
26+ * @param [in] self:
27+ * device侧的aclTensor，数据类型支持FLOAT、FLOAT16L。数据格式支持ND。
28+ * 支持非连续的Tensor。
29+ * @param [in] size: reshape时输出张量的第一个维度， 数据类型支持INT64。
30+ * @param [in] out: device侧的aclTensor，数据类型UINT8。数据格式支持ND.
31+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
32+ * @param [out] executor: 返回op执行器，包含算子计算流程。
33+ * @return aclnnStatus: 返回状态码。
34+ */
35+ACLNN_API aclnnStatus aclnnSignBitsPackGetWorkspaceSize(const aclTensor* self, int64_t size,
36+ aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
37+ 
38+/**
39+ * @brief aclnnSignBitsPack的第二段接口，用于执行计算。
40+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
41+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnAminGetWorkspaceSize获取。
42+ * @param [in] executor: op执行器，包含了算子计算流程。
43+ * @param [in] stream: acl stream流。
44+ * @return aclnnStatus: 返回状态码。
45+ */
46+ACLNN_API aclnnStatus aclnnSignBitsPack(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
47+ 
48+#ifdef __cplusplus
49+}
50+#endif
51+ 
52+#endif // OP_API_INC_SIGNBITSPACK_H_
Aexperimental/math/sign_bits_pack/op_api/sign_bits_pack.cpp+46-0
@@ -0,0 +1,46 @@
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_pack.h"
12+#include "opdev/op_dfx.h"
13+#include "opdev/shape_utils.h"
14+#include "opdev/make_op_executor.h"
15+ 
16+using namespace op;
17+namespace l0op {
18+OP_TYPE_REGISTER(SignBitsPack);
19+static constexpr size_t OUT_DIM = 2;
20+ 
21+const aclTensor* SignBitsPack(const aclTensor* self, int64_t size,aclOpExecutor* executor) {
22+ L0_DFX(SignBitsPack, self, size);
23+
24+ int64_t selfDimOne = self->GetViewShape().GetDim(0);
25+ int64_t outDimTwo = 0;
26+ 
27+ auto ysize = (selfDimOne + 7) / 8;
28+ if(size != 0)
29+ {
30+ outDimTwo = ysize / size;
31+ }
32+
33+ op::Shape outShape;
34+ outShape.SetDimNum(OUT_DIM);
35+ outShape.SetDim(0, size);
36+ outShape.SetDim(1, outDimTwo);
37+ 
38+ auto out = executor->AllocTensor(outShape, op::DataType::DT_UINT8, op::Format::FORMAT_ND);
39+ CHECK_RET(out != nullptr, nullptr);
40+ 
41+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(SignBitsPack, OP_INPUT(self), OP_OUTPUT(out), OP_ATTR(size));
42+ OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "SignBitsPackAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),
43+ return nullptr);
44+ return out;
45+}
46+} // namespace l0op
Aexperimental/math/sign_bits_pack/op_api/sign_bits_pack.h+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+#ifndef OP_API_INC_LEVEL0_SIGNBITSPACK_H_OP
11+#define OP_API_INC_LEVEL0_SIGNBITSPACK_H_OP
12+ 
13+#include "opdev/op_executor.h"
14+ 
15+namespace l0op {
16+const aclTensor *SignBitsPack(const aclTensor *self, int64_t size, aclOpExecutor *executor);
17+}
18+ 
19+#endif // OP_API_INC_LEVEL0_SIGNBITSPACK_H_OP
Aexperimental/math/sign_bits_pack/op_host/sign_bits_pack_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_pack_def.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class SignBitsPack : public OpDef {
19+public:
20+ explicit SignBitsPack(const char* name) : OpDef(name)
21+ {
22+ this->Input("x")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
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_UINT8, ge::DT_UINT8})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
32+ this->Attr("size").AttrType(REQUIRED).Int(1);
33+ this->AICore().AddConfig("ascend910b"); // 其他的soc版本补充部分配置项
R
Rrxtfeng4月27日

op_def中只支持910b，和aclnn代码中支持的芯片类型存在差异

likedislike
I
infinity_infi
4月28日 评论：
34+ }
35+};
36+OP_ADD(SignBitsPack); // 添加算子信息库
37+} // namespace ops
Aexperimental/math/sign_bits_pack/op_host/sign_bits_pack_infershape.cpp+60-0
@@ -0,0 +1,60 @@
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_pack_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+static constexpr int64_t PACKING_FACTOR = 8;
23+ 
24+static ge::graphStatus InferShapeSignBitsPack(gert::InferShapeContext* context)
25+{
26+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
27+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeSignBitsPack");
28+ 
29+ // get input shapes
30+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
31+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
32+ 
33+ // get output shapes
34+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
35+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
36+ 
37+ const size_t inputSize = xShape->GetDimNum() > 0 ? static_cast<size_t>(xShape->GetDim(0)) : 0;
38+
39+ if (inputSize == 0) {
40+ OP_LOGE(context->GetNodeName(), "Input shape has 0 elements");
41+ return ge::GRAPH_FAILED;
42+ }
43+ // 计算输出元素个数：每8个输入元素的符号位打包为1个字节
44+ // const size_t outputSize = (inputSize + PACKING_FACTOR - 1) / PACKING_FACTOR;
45+ size_t outputSize = (inputSize + PACKING_FACTOR - 1) / PACKING_FACTOR * PACKING_FACTOR;
46+
47+ // 设置输出形状为一维
48+ std::vector<int64_t> outputShapeVec = {static_cast<int64_t>(outputSize)};
49+ gert::Shape outputShape;
50+ outputShape.SetDimNum(1);
51+ outputShape.SetDim(0, static_cast<int64_t>(outputSize));
52+ *yShape = outputShape;
53+ OP_LOGD(context->GetNodeName(), "Input shape size: %zu, Output shape size: %zu",
54+ inputSize, outputSize);
55+ OP_LOGD(context->GetNodeName(), "End to do InferShapeSignBitsPack");
56+ return GRAPH_SUCCESS;
57+}
58+ 
59+IMPL_OP_INFERSHAPE(SignBitsPack).InferShape(InferShapeSignBitsPack);
60+}
Aexperimental/math/sign_bits_pack/op_host/sign_bits_pack_tiling.cpp+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_pack_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_pack_tiling_data.h"
24+#include "../op_kernel/sign_bits_pack_tiling_key.h"
25+ 
26+namespace optiling {
27+ 
28+using namespace Ops::Math::OpTiling;
29+ 
30+constexpr uint32_t BUFFER_NUM = 2;
31+constexpr uint32_t UB_ALIGN = 32;
32+constexpr uint32_t REPEAT_ALIGN = 256;
33+constexpr uint32_t GM_ALIGN = 512;
34+constexpr uint32_t RESERVED_UB_SIZE = 0; // 有些api需要预留ub空间
35+constexpr uint32_t MAX_TILEDATA = 18 * 1024; // 最大可以到
36+ 
37+struct SignBitsPackCompileInfo {};
38+ 
39+static ge::graphStatus TilingParseForSignBitsPack([[maybe_unused]] gert::TilingParseContext* context)
40+{
41+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
42+ return ge::GRAPH_SUCCESS;
43+}
44+ 
45+// 获取平台信息如ubSize, coreNum
46+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
47+{
48+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
49+ // 获取ubsize coreNum
50+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
51+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
52+ coreNum = ascendcPlatform.GetCoreNum();
53+ 
54+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
55+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
56+ return ge::GRAPH_SUCCESS;
57+}
58+ 
59+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
60+{
61+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
62+ size_t usrSize = 0;
63+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
64+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
65+ size_t* currentWorkspace = context->GetWorkspaceSizes(
66+ 1); // 通过框架获取workspace的指针，GetWorkspaceSizes入参为所需workspace的块数。当前限制使用一块。
67+ currentWorkspace[0] = usrSize + sysWorkspaceSize;
68+ return ge::GRAPH_SUCCESS;
69+}
70+ 
71+static ge::graphStatus CalcAndSetTilingData(uint64_t inputNum, uint32_t typeLength, int64_t realCoreNum, SignBitsPackTilingData* tiling, gert::TilingContext* context)
72+{
73+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
74+ OP_CHECK_IF(typeLength == 0, OP_LOGE(context, "typeLength is zero"), return ge::GRAPH_FAILED);
75+ uint64_t elemsPerGmBlock = (GM_ALIGN / typeLength);
76+ uint64_t inputLengthAlgin512 = (inputNum + elemsPerGmBlock - 1) / elemsPerGmBlock * elemsPerGmBlock;
77+
78+ uint64_t tileDataNum = MAX_TILEDATA;
79+ if(typeLength == 2)
80+ {
81+ tileDataNum = 2 * MAX_TILEDATA;
82+ }
83+ int64_t needCoreNum = (inputLengthAlgin512 + tileDataNum * BUFFER_NUM - 1) / (tileDataNum * BUFFER_NUM);
84+ int64_t coreNum = ((realCoreNum) < needCoreNum) ? realCoreNum : needCoreNum;
85+ uint64_t needCoreDataNum = ((inputLengthAlgin512 + coreNum - 1) / coreNum);
86+ if ((coreNum < realCoreNum / 4) && (needCoreDataNum > MAX_TILEDATA / 4)) {
87+ coreNum = coreNum * 2;
88+ needCoreDataNum = ((inputLengthAlgin512 + coreNum - 1) / coreNum);
89+ }
90+ uint32_t bufferNum = BUFFER_NUM;
91+ uint32_t usedDb = 1;
92+ if (needCoreDataNum < (MAX_TILEDATA / 2)) {
93+ bufferNum = 1;
94+ usedDb = 0;
95+ }
96+ uint64_t needTileDataNum = (needCoreDataNum + bufferNum - 1) / bufferNum;
97+ needTileDataNum = (needTileDataNum + elemsPerGmBlock - 1) / elemsPerGmBlock * elemsPerGmBlock;
98+ tileDataNum = (tileDataNum < needTileDataNum) ? tileDataNum : needTileDataNum;
99+ uint64_t everyCoreInputBlockNum = inputLengthAlgin512 / elemsPerGmBlock / coreNum;
100+ uint64_t tailBlockNum = (inputLengthAlgin512 / elemsPerGmBlock) % coreNum;
101+ uint64_t smallCoreDataNum = everyCoreInputBlockNum * elemsPerGmBlock;
102+ uint64_t finalSmallTileNum = (smallCoreDataNum + tileDataNum - 1) / tileDataNum;
103+ uint64_t smallTailDataNum = smallCoreDataNum - (finalSmallTileNum - 1) * tileDataNum;
104+ uint64_t bigCoreDataNum = smallCoreDataNum + elemsPerGmBlock;
105+ uint64_t finalBigTileNum = (bigCoreDataNum + tileDataNum - 1) / tileDataNum;
106+ uint64_t bigTailDataNum = bigCoreDataNum - (finalBigTileNum - 1) * tileDataNum;
107+ 
108+ uint64_t realLastPackLenth = smallTailDataNum - (inputLengthAlgin512 - inputNum);
109+ uint64_t elemsPerRepeat = (REPEAT_ALIGN / typeLength);
110+ uint32_t lastCalcLength = (realLastPackLenth + elemsPerRepeat - 1) / elemsPerRepeat * elemsPerRepeat;
111+ uint64_t elemsPerUbBlock = (UB_ALIGN / typeLength);
112+ uint32_t lastCopyLengthAlign = (realLastPackLenth + elemsPerUbBlock - 1) / elemsPerUbBlock * elemsPerUbBlock;
113+ uint32_t lastCopyLength = realLastPackLenth;
114+ uint32_t rightPaddingElemNums = lastCopyLengthAlign - realLastPackLenth;
115+ 
116+ tiling->smallCoreDataNum = smallCoreDataNum;
117+ tiling->bigCoreDataNum = bigCoreDataNum;
118+ tiling->finalBigTileNum = finalBigTileNum;
119+ tiling->finalSmallTileNum = finalSmallTileNum;
120+ tiling->tileDataNum = tileDataNum;
121+ tiling->smallTailDataNum = smallTailDataNum;
122+ tiling->bigTailDataNum = bigTailDataNum;
123+ tiling->tailBlockNum = tailBlockNum;
124+ tiling->usedDb = usedDb;
125+ tiling->lastCopyLength = lastCopyLength;
126+ tiling->rightPaddingElemNums = rightPaddingElemNums;
127+ tiling->lastCalcLength = lastCalcLength;
128+ context->SetBlockDim(coreNum);
129+ return ge::GRAPH_SUCCESS;
130+}
131+ 
132+// tiling 分发入口
133+static ge::graphStatus SignBitsPackTilingFunc(gert::TilingContext* context)
134+{
135+ // 1、获取平台运行信息
136+ uint64_t ubSize;
137+ int64_t realCoreNum;
138+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, realCoreNum);
139+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
140+ SignBitsPackTilingData* tiling = context->GetTilingData<SignBitsPackTilingData>();
141+ uint64_t inputNum;
142+ inputNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
143+ uint32_t typeLength = 0;
144+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
145+ CalcAndSetTilingData(inputNum, typeLength, realCoreNum, tiling, context);
146+ ret = GetWorkspaceSize(context);
147+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
148+ return ge::GRAPH_FAILED);
149+ uint64_t tilingKey = 0;
150+ tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0);
151+ context->SetTilingKey(tilingKey);
152+ return ge::GRAPH_SUCCESS;
153+}
154+ 
155+// tiling注册入口.
156+IMPL_OP_OPTILING(SignBitsPack).Tiling(SignBitsPackTilingFunc).TilingParse<SignBitsPackCompileInfo>(TilingParseForSignBitsPack);
157+} // namespace optiling
Aexperimental/math/sign_bits_pack/op_kernel/sign_bits_pack.cpp+28-0
@@ -0,0 +1,28 @@
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_pack.cpp
13+ * \brief
14+ */
15+ 
16+#include "sign_bits_pack.h"
17+ 
18+template <uint32_t schMode>
19+__global__ __aicore__ void sign_bits_pack(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
20+{
21+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
22+ REGISTER_TILING_DEFAULT(SignBitsPackTilingData);
23+ GET_TILING_DATA_WITH_STRUCT(SignBitsPackTilingData, tilingData, tiling);
24+ AscendC::TPipe pipe;
25+ NsSignBitsPack::KernelSignBitsPack<DTYPE_X> op; // 算子kernel实例获取
26+ op.Init(x, y, &tilingData, &pipe);
27+ op.Process();
28+}
Aexperimental/math/sign_bits_pack/op_kernel/sign_bits_pack.h+173-0
@@ -0,0 +1,173 @@
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_pack.h
13+ * \brief
14+ */
15+#ifndef SIGN_BITS_PACK_H_
16+#define SIGN_BITS_PACK_H_
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "sign_bits_pack_tiling_data.h"
21+#include "sign_bits_pack_tiling_key.h"
22+ 
23+namespace NsSignBitsPack {
24+ 
25+using namespace AscendC;
26+ 
27+template <typename T>
28+class KernelSignBitsPack {
29+public:
30+ __aicore__ inline KernelSignBitsPack(){};
31+ 
32+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const SignBitsPackTilingData* tilingData, TPipe* pipeIn);
33+ __aicore__ inline void Process();
34+ 
35+private:
36+ __aicore__ inline void CopyIn(int32_t progress);
37+ __aicore__ inline void CopyOut(int32_t progress);
38+ __aicore__ inline void Compute(int32_t progress);
39+ __aicore__ inline void CopyInLast(int32_t progress);
40+ __aicore__ inline void CopyOutLast(int32_t progress);
41+ 
42+private:
43+
44+ TQue<QuePosition::VECIN, 1> inQueueIN;
45+ TQue<QuePosition::VECOUT, 1> outQueueOUT;
46+ GlobalTensor<T> xGm;
47+ GlobalTensor<uint8_t> yGm;
48+ uint64_t coreDataNum = 0;
49+ uint64_t tileNum = 0;
50+ uint64_t tileDataNum = 0;
51+ uint64_t tailDataNum = 0;
52+ uint64_t processDataNum = 0;
53+ int32_t bufferNum = 2;
54+ uint32_t rightPaddingElemNums = 0;
55+ uint32_t lastCopyLength = 0;
56+ uint32_t lastCalcLength = 0;
57+};
58+ 
59+template <typename T>
60+__aicore__ inline void KernelSignBitsPack<T>::Init(GM_ADDR x, GM_ADDR y, const SignBitsPackTilingData* tilingData, TPipe* pipeIn)
61+{
62+ uint64_t coreId = AscendC::GetBlockIdx();
63+ uint64_t globalBufferIndex = tilingData->bigCoreDataNum * AscendC::GetBlockIdx();
64+ this->tileDataNum = tilingData->tileDataNum;
65+ if (coreId < tilingData->tailBlockNum) {
66+ this->coreDataNum = tilingData->bigCoreDataNum;
67+ this->tileNum = tilingData->finalBigTileNum;
68+ this->tailDataNum = tilingData->bigTailDataNum;
69+ } else {
70+ this->coreDataNum = tilingData->smallCoreDataNum;
71+ this->tileNum = tilingData->finalSmallTileNum;
72+ this->tailDataNum = tilingData->smallTailDataNum;
73+ globalBufferIndex -= (tilingData->bigCoreDataNum - tilingData->smallCoreDataNum) * (AscendC::GetBlockIdx() - tilingData->tailBlockNum);
74+ }
75+
76+ this->rightPaddingElemNums = tilingData->rightPaddingElemNums;
77+ this->lastCopyLength = tilingData->lastCopyLength;
78+ this->lastCalcLength = tilingData->lastCalcLength;
79+ this->bufferNum = 1;
80+ if (static_cast<int32_t>(tilingData->usedDb) == 1) {
81+ this->bufferNum = 2;
82+ }
83+ xGm.SetGlobalBuffer((__gm__ T*)x + globalBufferIndex, this->coreDataNum);
84+ yGm.SetGlobalBuffer((__gm__ uint8_t*)y + globalBufferIndex / 8, this->coreDataNum / 8);
85+ pipeIn->InitBuffer(inQueueIN, this->bufferNum, this->tileDataNum * sizeof(T));
86+ pipeIn->InitBuffer(outQueueOUT, this->bufferNum, this->tileDataNum / 4 * sizeof(uint8_t));
87+}
88+ 
89+template <typename T>
90+__aicore__ inline void KernelSignBitsPack<T>::CopyIn(int32_t progress)
91+{
92+ AscendC::LocalTensor<T> xLocal = inQueueIN.AllocTensor<T>();
93+ 
94+ AscendC::DataCopy(xLocal, xGm[progress * this->tileDataNum], this->processDataNum);
95+ 
96+ inQueueIN.EnQue(xLocal);
97+}
98+ 
99+template <typename T> // DatacopyPad
100+__aicore__ inline void KernelSignBitsPack<T>::CopyInLast(int32_t progress)
101+{
102+ AscendC::LocalTensor<T> xLocal = inQueueIN.AllocTensor<T>();
103+ 
104+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(this->lastCopyLength * sizeof(T)), 0, 0, 0};
105+ AscendC::DataCopyPadExtParams<T> padParams{true, 0, static_cast<uint8_t>(this->rightPaddingElemNums), (T)(-1)};
106+ AscendC::DataCopyPad(xLocal, xGm[progress * this->tileDataNum], copyParams, padParams);
107+ 
108+ inQueueIN.EnQue(xLocal);
109+}
110+ 
111+ 
112+template <typename T>
113+__aicore__ inline void KernelSignBitsPack<T>::CopyOut(int32_t progress)
114+{
115+ AscendC::LocalTensor<uint8_t> yLocal = outQueueOUT.DeQue<uint8_t>();
116+ 
117+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>((this->processDataNum / 8) * sizeof(uint8_t)), 0, 0, 0};
118+ AscendC::DataCopyPad(yGm[progress * this->tileDataNum / 8], yLocal, copyParams);
119+ 
120+ outQueueOUT.FreeTensor(yLocal);
121+}
122+ 
123+template <typename T>
124+__aicore__ inline void KernelSignBitsPack<T>::CopyOutLast(int32_t progress)
125+{
126+ AscendC::LocalTensor<uint8_t> yLocal = outQueueOUT.DeQue<uint8_t>();
127+ 
128+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(this->lastCopyLength / 8* sizeof(uint8_t)), 0, 0, 0};
129+ AscendC::DataCopyPad(yGm[progress * this->tileDataNum / 8], yLocal, copyParams);
130+ 
131+ outQueueOUT.FreeTensor(yLocal);
132+}
133+ 
134+template <typename T>
135+__aicore__ inline void KernelSignBitsPack<T>::Compute(int32_t progress)
136+{
137+ LocalTensor<T> xLocal = inQueueIN.DeQue<T>();
138+ LocalTensor<uint8_t> yLocal = outQueueOUT.AllocTensor<uint8_t>();
139+
140+ CompareScalar(yLocal, xLocal, (T)(0), CMPMODE::GE, this->processDataNum);
141+ 
142+ outQueueOUT.EnQue(yLocal);
143+ inQueueIN.FreeTensor(xLocal);
144+}
145+ 
146+template <typename T>
147+__aicore__ inline void KernelSignBitsPack<T>::Process()
148+{
149+ int32_t loopCount = this->tileNum;
150+ this->processDataNum = this->tileDataNum;
151+ 
152+ for (int32_t i = 0; i < loopCount - 1; i++) {
153+ CopyIn(i);
154+ Compute(i);
155+ CopyOut(i);
156+ }
157+ if (AscendC::GetBlockIdx() == (AscendC::GetBlockNum() - 1)) // 最后一个核
158+ {
159+ this->processDataNum = this->lastCalcLength;
160+ CopyInLast(loopCount - 1);
161+ }
162+ else
163+ {
164+ this->processDataNum = this->tailDataNum;
165+ CopyIn(loopCount - 1);
166+ }
167+
168+ Compute(loopCount - 1);
169+ CopyOut(loopCount - 1);
170+}
171+ 
172+} // namespace NsSignBitsPack
173+#endif // SIGN_BITS_PACK_H
Aexperimental/math/sign_bits_pack/op_kernel/sign_bits_pack_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+ * \file sign_bits_pack_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef SIGN_BITS_PACK_TILING_DATA_H_
17+#define SIGN_BITS_PACK_TILING_DATA_H_
18+ 
19+struct SignBitsPackTilingData {
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+ uint32_t usedDb;
29+ uint32_t lastCopyLength;
30+ uint32_t rightPaddingElemNums;
31+ uint32_t lastCalcLength;
32+};
33+#endif
Aexperimental/math/sign_bits_pack/op_kernel/sign_bits_pack_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_pack_tiling_key.h
13+ * \brief sign_bits_pack tiling key declare
14+ */
15+ 
16+#ifndef __SIGN_BITS_PACK_TILING_KEY_H__
17+#define __SIGN_BITS_PACK_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+ SignBitsPack,
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_pack/tests/ut/op_api/test_aclnn_sign_bits_pack.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+/*!
12+ * \file test_aclnn_sign_bits_pack.cpp
13+ * \brief
14+*/
15+#include <vector>
16+#include <array>
17+#include "gtest/gtest.h"
18+ 
19+#include "opdev/platform.h"
20+#include "../../../op_api/aclnn_sign_bits_pack.h"
21+#include "op_api_ut_common/tensor_desc.h"
22+#include "op_api_ut_common/scalar_desc.h"
23+#include "op_api_ut_common/op_api_ut.h"
24+ 
25+class l2_sign_bits_pack_test : public testing::Test {
26+protected:
27+ static void SetUpTestCase()
28+ {
29+ cout << "l2_sign_bits_pack_test SetUp" << endl;
30+ }
31+ 
32+ static void TearDownTestCase()
33+ {
34+ cout << "l2_sign_bits_pack_test TearDown" << endl;
35+ }
36+};
37+ 
38+TEST_F(l2_sign_bits_pack_test, case_01_float)
39+{
40+ op::SetPlatformSocVersion(op::SocVersion::ASCEND910B);
41+ 
42+ auto selfDesc = TensorDesc({14}, ACL_FLOAT, ACL_FORMAT_ND).Value(vector<float>{5, 4, 3, 2, 0, -1, -2, 4, 3, 2, 1, 0, -1, -2});
43+ int64_t size = 2;
44+ auto outDesc = TensorDesc({2, 1}, ACL_UINT8, ACL_FORMAT_ND);
45+ 
46+ auto ut = OP_API_UT(aclnnSignBitsPack, INPUT(selfDesc, size), OUTPUT(outDesc));
47+ 
48+ uint64_t workspace_size = 0;
49+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
50+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
51+}
Aexperimental/math/sign_bits_pack/tests/ut/op_host/sign_bits_pack_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_pack_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef SIGN_BITS_PACK_TILING_H
17+#define SIGN_BITS_PACK_TILING_H
18+ 
19+#include "register/tilingdata_base.h"
20+ 
21+namespace optiling {
22+struct SignBitsPackCompileInfo {
23+ int32_t totalCoreNum = 0;
24+ int64_t ubSize = 0;
25+ bool isRegbase = false;
26+};
27+} // namespace optiling
28+ 
29+#endif // SIGN_BITS_PACK_TILING_H
Aexperimental/math/sign_bits_pack/tests/ut/op_host/test_sign_bits_pack_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_pack_tiling.h"
14+#include "../../../op_kernel/sign_bits_pack_tiling_data.h"
15+#include "../../../op_kernel/sign_bits_pack_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 SignBitsPackTiling : public testing::Test {
23+protected:
24+ static void SetUpTestCase()
25+ {
26+ cout << "SignBitsPackTiling SetUp" << endl;
27+ }
28+ 
29+ static void TearDownTestCase()
30+ {
31+ cout << "SignBitsPackTiling TearDown " << endl;
32+ }
33+};
34+ 
35+TEST_F(SignBitsPackTiling, ascend910_test_tiling_FLOAT_001)
36+{
37+ optiling::SignBitsPackCompileInfo compileInfo = {40, 196608, false};
38+ gert::TilingContextPara tilingContextPara(
39+ "SignBitsPack",
40+ {
41+ {{{14}, {14}}, ge::DT_FLOAT, ge::FORMAT_ND},
42+ },
43+ {
44+ {{{2, 1}, {2, 1}}, ge::DT_UINT8, ge::FORMAT_ND},
45+ },
46+ &compileInfo);
47+ uint64_t expectTilingKey = 0;
48+ string expectTilingData = "128 256 2 1 128 128 128 0 60129542144 274877906946 ";
49+ std::vector<size_t> expectWorkspaces = {16777216};
50+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
51+}
Aexperimental/math/sign_bits_pack/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_pack_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/sign_bits_pack_tiling.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/sign_bits_pack_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_pack "ascend910b" "-DDTYPE_X=float" "${sign_bits_pack_tiling_files}")
29+endif()
Aexperimental/math/sign_bits_pack/tests/ut/op_kernel/sign_bits_pack_data/compare_data.py+65-0
@@ -0,0 +1,65 @@
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 == "uint8":
20+ np_dtype = np.uint8
21+ else:
22+ raise ValueError("d_type must be uint8")
23+ data_same = True
24+
25+ for gold, out in zip(golden_file_lists, output_file_lists):
26+ tmp_out = np.fromfile(out, np.uint8)
27+ tmp_gold = np.fromfile(gold, np.uint8)
28+
29+ # 检查形状是否一致
30+ if tmp_out.shape != tmp_gold.shape:
31+ print(f"FAILED! 形状不匹配: output形状{tmp_out.shape}, golden形状{tmp_gold.shape}")
32+ data_same = False
33+ continue
34+
35+ # 方法1：直接比较（推荐，uint8应该完全相等）
36+ if np.array_equal(tmp_out, tmp_gold):
37+ print(f"PASSED! 文件: {out}")
38+ else:
39+ print(f"FAILED! 文件: {out}")
40+ # 找出不相同的索引
41+ diff_idx = np.where(tmp_out != tmp_gold)[0]
42+ print(f" 不相同的数据数量: {len(diff_idx)}")
43+
44+ # 显示前几个不同的值
45+ for i, idx in enumerate(diff_idx[:10]):
46+ print(f" 索引 {idx}: output={tmp_out[idx]}(0x{tmp_out[idx]:02x}), "
47+ f"golden={tmp_gold[idx]}(0x{tmp_gold[idx]:02x}), "
48+ f"差异={int(tmp_out[idx]) - int(tmp_gold[idx])}")
49+ data_same = False
50+ return data_same
51+ 
52+def get_file_lists(dtype):
53+ golden_file_lists = sorted(glob.glob(curr_dir + "/*golden*.bin"))
54+ output_file_lists = sorted(glob.glob(curr_dir + "/*output*.bin"))
55+ return golden_file_lists, output_file_lists
56+ 
57+def process(d_type):
58+ golden_file_lists, output_file_lists = get_file_lists(d_type)
59+ result = compare_data(golden_file_lists, output_file_lists, d_type)
60+ print("compare result:", result)
61+ return result
62+ 
63+if __name__ == '__main__':
64+ ret = process(sys.argv[1])
65+ exit(0 if ret else 1)
Aexperimental/math/sign_bits_pack/tests/ut/op_kernel/sign_bits_pack_data/gen_data.py+63-0
@@ -0,0 +1,63 @@
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="float"):
23+ d_type_dict = {
24+ "float16": np.float16,
25+ "float": np.float32
26+ }
27+ np_type = d_type_dict[d_type]
28+ shape = parse_str_to_shape_list(shape_str)
29+ size = np.prod(shape)
30+ input_x = np.random.uniform(-10, 10, shape).astype(np_type)
31+ 
32+ signs = (input_x >= 0).astype(np.uint8)
33+ n_elements = len(input_x)
34+ # 计算每个压缩单元包含多少元素
35+ n_packs = (n_elements + 7) // 8 # 每个uint8包含8个符号位
36+ 
37+ # 位压缩
38+ packed = []
39+ for i in range(n_packs):
40+ packed_val = 0
41+ start_idx = i * 8
42+ end_idx = min((i + 1) * 8, n_elements)
43+
44+ for j in range(start_idx, end_idx):
45+ bit_pos = j - start_idx
46+ if signs[j]:
47+ packed_val |= (1 << bit_pos)
48+ packed.append(packed_val)
49+ # 填充到total_packs长度
50+ # golden = np.sign_bits_pack(input_x).astype(np_type)
51+ golden = np.array(packed, dtype=np.uint8)
52+ print(golden.shape)
53+ input_x.astype(np_type).tofile(f"{d_type}_input_t_sign_bits_pack.bin")
54+ golden.astype(np.uint8).tofile(f"uint8_golden_t_sign_bits_pack.bin")
55+ 
56+ 
57+if __name__ == "__main__":
58+ if len(sys.argv) != 3:
59+ print("Param num must be 3.")
60+ exit(1)
61+ # 清理bin文件
62+ os.system("rm -rf *.bin")
63+ gen_data_and_golden(sys.argv[1], sys.argv[2])
Aexperimental/math/sign_bits_pack/tests/ut/op_kernel/sign_bits_pack_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_pack_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef _SIGN_BITS_PACK_TILING_H_
17+#define _SIGN_BITS_PACK_TILING_H_
18+ 
19+#include <cstdint>
20+ 
21+#include "../../../op_kernel/sign_bits_pack_tiling_data.h"
22+#include "kernel_tiling/kernel_tiling.h"
23+#include "../../../op_kernel/sign_bits_pack.h"
24+ 
25+using namespace NsSignBitsPack;
26+ 
27+#define __aicore__
28+#ifdef __NPU_TILING__
29+inline[aicore] void InitTilingData(const __gm__ uint8_t *tiling, SignBitsPackTilingData *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(SignBitsPackTilingData) / 4; i++) {
34+ *(dst + i) = *(src + i);
35+ }
36+}
37+#else
38+inline void InitTilingData(uint8_t *tiling, SignBitsPackTilingData *constData)
39+{
40+ memcpy(constData, tiling, sizeof(SignBitsPackTilingData));
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+ SignBitsPackTilingData tilingData; \
57+ InitTilingData(tilingArg, &tilingData)
58+ 
59+#endif // _SIGN_BITS_PACK_TILING_H_
Aexperimental/math/sign_bits_pack/tests/ut/op_kernel/test_sign_bits_pack.cpp+109-0
@@ -0,0 +1,109 @@
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_pack.cpp"
25+ 
26+using namespace std;
27+ 
28+constexpr uint32_t smallCoreDataNum = 128;
29+constexpr uint32_t bigCoreDataNum = 256;
30+constexpr uint32_t tileDataNum = 128;
31+constexpr uint32_t smallTailDataNum = 1024;
32+constexpr uint32_t bigTailDataNum = 1040;
33+ 
34+extern "C" __global__ __aicore__ void sign_bits_pack(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
35+ 
36+class SignBitsPackTest : public testing::Test {
37+protected:
38+ static void SetUpTestCase()
39+ {
40+ std::cout << "sign_bits_pack_test SetUp" << std::endl;
41+ const string cmd = "cp -rf " + dataPath + " ./";
42+ system(cmd.c_str());
43+ system("chmod -R 755 ./sign_bits_pack_data/");
44+ }
45+ static void TearDownTestCase()
46+ {
47+ std::cout << "sign_bits_pack_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 SignBitsPackTest::rootPath = "../../../../experimental/";
56+const std::string SignBitsPackTest::dataPath = rootPath + "math/sign_bits_pack/tests/ut/op_kernel/sign_bits_pack_data";
57+ 
58+template <typename T1, typename T2>
59+inline T1 CeilAlign(T1 a, T2 b)
60+{
61+ if ( b == 0 )
62+ return 0;
63+ return (a + b - 1) / b * b;
64+}
65+ 
66+TEST_F(SignBitsPackTest, test_case_float_1)
67+{
68+ size_t inputXByteSize = 14 * sizeof(float);
69+ size_t outputYByteSize = 2 * sizeof(uint8_t);
70+ size_t tilingDataSize = sizeof(SignBitsPackTilingData);
71+ 
72+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputXByteSize);
73+ 
74+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(outputYByteSize);
75+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(1024 * 1024 * 1024);
76+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);
77+ uint32_t blockDim = 1;
78+ 
79+ char* path_ = get_current_dir_name();
80+ string path(path_);
81+ 
82+ SignBitsPackTilingData* tilingDatafromBin = reinterpret_cast<SignBitsPackTilingData*>(tiling);
83+ tilingDatafromBin->smallCoreDataNum = 128;
84+ tilingDatafromBin->bigCoreDataNum = 256;
85+ tilingDatafromBin->finalBigTileNum = 2;
86+ tilingDatafromBin->finalSmallTileNum = 1;
87+ tilingDatafromBin->tileDataNum = 128;
88+ tilingDatafromBin->smallTailDataNum = 128;
89+ tilingDatafromBin->bigTailDataNum = 128;
90+ tilingDatafromBin->tailBlockNum = 0;
91+ tilingDatafromBin->usedDb = 0;
92+ tilingDatafromBin->lastCopyLength = 14;
93+ tilingDatafromBin->rightPaddingElemNums = 2;
94+ tilingDatafromBin->lastCalcLength = 64;
95+ 
96+ auto KernelSignBitsPack = [](GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
97+ ::sign_bits_pack<0>(x, y, workspace, tiling);
98+ };
99+ 
100+ ICPU_SET_TILING_KEY(0);
101+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
102+ ICPU_RUN_KF(sign_bits_pack<0>, blockDim, x, y, workspace, (uint8_t*)(tilingDatafromBin));
103+ 
104+ AscendC::GmFree(x);
105+ AscendC::GmFree(y);
106+ AscendC::GmFree(workspace);
107+ AscendC::GmFree(tiling);
108+ free(path_);
109+}