Star1325
1440
代码
代码
Issues114
Pull Requests240
Actions
项目讨论
Wiki
分析
Star1325
1440
已合并
提交Ascend C实现的Cast算子 #646
skywang2创建于 1月4日
提交Ascend C实现的Cast算子 #646
已合并
skywang2创建于 1月4日
从skywang2/ops-mathmath-Cast合入到cann/ops-mathmaster
共 13 个文件变更+2387-0
Aexperimental/math/cast/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 cast ACLNNTYPE aclnn_exclude)
Aexperimental/math/cast/README.md+38-0
@@ -0,0 +1,38 @@
1+# Cast
2+## 贡献说明
3+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
4+|--------|------------------|-------|-----------|-----------|
5+| skywang2 | 个人开发者 | Cast | 2025/12/31 | 新增Cast算子 |
6+ 
7+### 算子描述
8+`Cast`算子提供将tensor从源数据类型转换为目标数据类型的功能。
9+ 
10+### 算子规格描述
11+ 
12+<table>
13+<tr><th align="center">算子类型(OpType)</th><th colspan="4" align="center">Cast</th></tr>
14+<tr><td rowspan="2" align="center">算子输入</td><td align="center">name</td><td align="center">type</td><td align="center">data type</td><td align="center">format</td></tr>
15+<tr><td align="center">x</td><td align="center">tensor</td><td align="center">float16, float32, int32, int8, uint8, bool, int64, bfloat16, int16</td><td align="center">ND</td></tr>
16+<tr><td rowspan="1" align="center">算子输出</td><td align="center">out</td><td align="center">tensor</td><td align="center">float16, float32, int32, int8, uint8, bool, int64, bfloat16, int16</td><td align="center">ND</td></tr>
17+<tr><td rowspan="1" align="center">算子属性</td><td align="center">dstType</td><td align="center">attr</td><td align="center">int64</td><td align="center">-</td></tr>
18+<tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">cast</td></td></tr>
19+</table>
20+ 
21+### 支持的产品型号
22+本样例支持如下产品型号：
23+- Atlas A2训练系列产品
24+- Atlas 800I A2推理产品
25+ 
26+ 
27+### 环境要求
28+编译运行此样例前，请参考[《CANN软件安装指南》](https://hiascend.com/document/redirect/CannCommunityInstSoftware)完成开发运行环境的部署。
29+ 
30+### 算子调用
31+测试命令调用方式：[build.sh](../../../docs/zh/invocation/quick_op_invocation.md)
32+<table>
33+ <th>目录</th><th>描述</th>
34+ <tr>
35+ <td><a href="./examples/test_aclnn_cast.cpp"> test_aclnn_cast.cpp</td><td>通过aclnn调用的方式调用Cast算子</td>
36+ </tr>
37+</table>
38+ 
Aexperimental/math/cast/docs/aclnnCast.md+358-0
@@ -0,0 +1,358 @@
1+# aclnnCast
2+ 
3+ 
4+## 产品支持情况
5+ 
6+| 产品 | 是否支持 |
7+| :----------------------------------------------------------- | :------: |
8+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
9+ 
10+## 功能说明
11+ 
12+将输入tensor转换为指定的dtype类型。
13+ 
14+## 函数原型
15+ 
16+每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnCastGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnCast”接口执行计算。
17+```Cpp
18+aclnnStatus aclnnCastGetWorkspaceSize(
19+ const aclTensor *self,
20+ const aclDataType dtype,
21+ aclTensor *out,
22+ uint64_t *workspaceSize,
23+ aclOpExecutor **executor)
24+```
25+```Cpp
26+aclnnStatus aclnnCast(
27+ void *workspace,
28+ uint64_t workspaceSize,
29+ aclOpExecutor *executor,
30+ aclrtStream stream)
31+```
32+## aclnnCastGetWorkspaceSize
33+ 
34+- **参数说明：**
35+ 
36+ <table style="undefined;table-layout: fixed; width: 1495px"><colgroup>
37+ <col style="width: 146px">
38+ <col style="width: 110px">
39+ <col style="width: 301px">
40+ <col style="width: 219px">
41+ <col style="width: 328px">
42+ <col style="width: 101px">
43+ <col style="width: 143px">
44+ <col style="width: 147px">
45+ </colgroup>
46+ <thead>
47+ <tr>
48+ <th>参数名</th>
49+ <th>输入/输出</th>
50+ <th>描述</th>
51+ <th>使用说明</th>
52+ <th>数据类型</th>
53+ <th>数据格式</th>
54+ <th>维度(shape)</th>
55+ <th>非连续Tensor</th>
56+ </tr></thead>
57+ <tbody>
58+ <tr>
59+ <td>self</td>
60+ <td>输入</td>
61+ <td>待进行cast计算的入参，Device侧的aclTensor。</td>
62+ <td>-</td>
63+ <td>FLOAT16、FLOAT、DOUBLE、INT8、UINT8、INT16、UINT16、INT32、UINT32、INT64、UINT64、BOOL、COMPLEX32、COMPLEX64、COMPLEX128、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2</td>
64+ <td>ND</td>
65+ <td>0-8</td>
66+ <td>√</td>
67+ </tr>
68+ <tr>
69+ <td>dtype</td>
70+ <td>属性</td>
71+ <td>输入tensor要转换的目标dtype。</td>
72+ <td>-</td>
73+ <td>const aclDataType</td>
74+ <td>-</td>
75+ <td>-</td>
76+ <td>-</td>
77+ </tr>
78+ <tr>
79+ <td>out</td>
80+ <td>输出</td>
81+ <td>待进行cast计算的出参，Device侧的aclTensor。</td>
82+ <td>shape与self相同。</td>
83+ <td>FLOAT16、FLOAT、DOUBLE、INT8、UINT8、INT16、UINT16、INT32、UINT32、INT64、UINT64、BOOL、COMPLEX32、COMPLEX64、COMPLEX128、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2、INT4（暂不支持非连续Tensor）</td>
84+ <td>ND</td>
85+ <td>0-8</td>
86+ <td>√</td>
87+ </tr>
88+ <tr>
89+ <td>workspaceSize</td>
90+ <td>输出</td>
91+ <td>返回需要在Device侧申请的workspace大小。</td>
92+ <td>-</td>
93+ <td>-</td>
94+ <td>-</td>
95+ <td>-</td>
96+ <td>-</td>
97+ </tr>
98+ <tr>
99+ <td>executor</td>
100+ <td>输出</td>
101+ <td>返回op执行器，包含了算子计算流程。</td>
102+ <td>-</td>
103+ <td>-</td>
104+ <td>-</td>
105+ <td>-</td>
106+ <td>-</td>
107+ </tr>
108+ </tbody></table>
109+
110+ - <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>：不支持COMPLEX32、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2、INT4。
111+ 
112+- **返回值：**
113+ 
114+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
115+ 
116+ 第一段接口会完成入参校验，出现以下场景时报错：
117+
118+ <table style="undefined;table-layout: fixed;width: 1155px"><colgroup>
119+ <col style="width: 319px">
120+ <col style="width: 144px">
121+ <col style="width: 671px">
122+ </colgroup>
123+ <thead>
124+ <tr>
125+ <th>返回码</th>
126+ <th>错误码</th>
127+ <th>描述</th>
128+ </tr>
129+ </thead>
130+ <tbody>
131+ <tr>
132+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
133+ <td>161001</td>
134+ <td>传入的tensor或out是空指针。</td>
135+ </tr>
136+ <tr>
137+ <td rowspan="5">ACLNN_ERR_PARAM_INVALID</td>
138+ <td rowspan="5">161002</td>
139+ <td>self的数据类型和数据格式不在支持的范围之内。</td>
140+ </tr>
141+ <tr>
142+ <td>self的数据格式与out的数据格式不同。</td>
143+ </tr>
144+ <tr>
145+ <td>self的shape与out的shape不同。</td>
146+ </tr>
147+ <tr>
148+ <td>参数dtype不在输出支持的数据格式范围之内。</td>
149+ </tr>
150+ <tr>
151+ <td>out的数据类型为INT4时，self为非连续Tensor。</td>
152+ </tr>
153+ <tr>
154+ <td>ACLNN_ERR_INNER_TILING_ERROR</td>
155+ <td>561002</td>
156+ <td>out的数据类型为INT4时，self的shape尾轴为奇数。</td>
157+ </tr>
158+ </tbody></table>
159+ 
160+## aclnnCast
161+ 
162+- **参数说明：**
163+ <table style="undefined;table-layout: fixed; width: 598px"><colgroup>
164+ <col style="width: 173px">
165+ <col style="width: 173px">
166+ <col style="width: 668px">
167+ </colgroup>
168+ <thead>
169+ <tr>
170+ <th>参数名</th>
171+ <th>输入/输出</th>
172+ <th>描述</th>
173+ </tr></thead>
174+ <tbody>
175+ <tr>
176+ <td>workspace</td>
177+ <td>输入</td>
178+ <td>在Device侧申请的workspace内存地址。</td>
179+ </tr>
180+ <tr>
181+ <td>workspaceSize</td>
182+ <td>输入</td>
183+ <td>在Device侧申请的workspace大小，由第一段接口aclnnCastGetWorkspaceSize获取。</td>
184+ </tr>
185+ <tr>
186+ <td>executor</td>
187+ <td>输入</td>
188+ <td>op执行器，包含了算子计算流程。</td>
189+ </tr>
190+ <tr>
191+ <td>stream</td>
192+ <td>输入</td>
193+ <td>指定执行任务的Stream。</td>
194+ </tr>
195+ </tbody>
196+ </table>
197+ 
198+ 
199+- **返回值：**
200+ 
201+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
202+ 
203+## 约束说明
204+ 
205+- 确定性计算：
206+ - aclnnCast默认确定性实现。
207+ 
208+- 针对数据类型从浮点数转换为整型的场景：
209+ 输入数据中存在nan，则将nan转换为0。
210+ 
211+- 针对输入数据类型为BOOL、COMPLEX32、COMPLEX64、COMPLEX128、FLOAT4_E2M1、FLOAT4_E1M2的场景：
212+ 不支持输入为非连续。
213+ 
214+- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>：
215+ - 针对数据类型从int32转换为int8的场景：
216+ 只能保证输入数据在(-2048, 1920)范围内精度无误差。
217+ - 针对数据类型从float64/complex64/complex128转换为uint8的场景：
218+ 只能保证输入数据为非负数精度无误差。
219+ 
220+ 
221+ 
222+## 调用示例
223+ 
224+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
225+```Cpp
226+#include <iostream>
227+#include <vector>
228+#include "acl/acl.h"
229+#include "aclnnop/aclnn_cast.h"
230+ 
231+#define CHECK_RET(cond, return_expr) \
232+ do { \
233+ if (!(cond)) { \
234+ return_expr; \
235+ } \
236+ } while (0)
237+ 
238+#define LOG_PRINT(message, ...) \
239+ do { \
240+ printf(message, ##__VA_ARGS__); \
241+ } while (0)
242+ 
243+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
244+ int64_t shapeSize = 1;
245+ for (auto i : shape) {
246+ shapeSize *= i;
247+ }
248+ return shapeSize;
249+}
250+ 
251+int Init(int32_t deviceId, aclrtStream* stream) {
252+ // 固定写法，初始化
253+ auto ret = aclInit(nullptr);
254+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
255+ ret = aclrtSetDevice(deviceId);
256+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
257+ ret = aclrtCreateStream(stream);
258+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
259+ return 0;
260+}
261+ 
262+template <typename T>
263+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
264+ aclDataType dataType, aclTensor** tensor) {
265+ auto size = GetShapeSize(shape) * sizeof(T);
266+ // 调用aclrtMalloc申请device侧内存
267+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
268+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
269+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
270+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
271+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
272+ 
273+ // 计算连续tensor的strides
274+ std::vector<int64_t> strides(shape.size(), 1);
275+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
276+ strides[i] = shape[i + 1] * strides[i + 1];
277+ }
278+ 
279+ // 调用aclCreateTensor接口创建aclTensor
280+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
281+ shape.data(), shape.size(), *deviceAddr);
282+ return 0;
283+}
284+ 
285+int main() {
286+ // 1. （固定写法）device/stream初始化，参考acl API文档
287+ // 根据自己的实际device填写deviceId
288+ int32_t deviceId = 0;
289+ aclrtStream stream;
290+ auto ret = Init(deviceId, &stream);
291+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
292+ 
293+ // 2. 构造输入与输出，需要根据API的接口自定义构造
294+ std::vector<int64_t> selfShape = {4, 2};
295+ std::vector<int64_t> outShape = {4, 2};
296+ 
297+ void* selfDeviceAddr = nullptr;
298+ void* outDeviceAddr = nullptr;
299+ aclTensor* self = nullptr;
300+ aclTensor* out = nullptr;
301+ 
302+ std::vector<float> selfHostData = {0.1, 1.1, 2.1, 3.1, 4.1, 5.1, 6.1, 7.1};
303+ std::vector<double> outHostData = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0};
304+ 
305+ // 创建self aclTensor
306+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
307+ CHECK_RET(ret == ACL_SUCCESS, return ret);
308+ // 创建out aclTensor
309+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_DOUBLE, &out);
310+ CHECK_RET(ret == ACL_SUCCESS, return ret);
311+ 
312+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
313+ uint64_t workspaceSize = 0;
314+ aclOpExecutor* executor;
315+ // 调用aclnnCast第一段接口
316+ ret = aclnnCastGetWorkspaceSize(self, aclDataType::ACL_DOUBLE, out, &workspaceSize, &executor);
317+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnCastGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
318+ // 根据第一段接口计算出的workspaceSize申请device内存
319+ void* workspaceAddr = nullptr;
320+ if (workspaceSize > 0) {
321+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
322+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
323+ }
324+ // 调用aclnnCast第二段接口
325+ ret = aclnnCast(workspaceAddr, workspaceSize, executor, stream);
326+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnCast failed. ERROR: %d\n", ret); return ret);
327+ 
328+ // 4. （固定写法）同步等待任务执行结束
329+ ret = aclrtSynchronizeStream(stream);
330+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
331+ 
332+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
333+ auto size = GetShapeSize(outShape);
334+ std::vector<double> resultData(size, 0);
335+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
336+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
337+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
338+ for (int64_t i = 0; i < size; i++) {
339+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
340+ }
341+ 
342+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
343+ aclDestroyTensor(self);
344+ aclDestroyTensor(out);
345+ 
346+ // 7. 释放device 资源
347+ aclrtFree(selfDeviceAddr);
348+ aclrtFree(outDeviceAddr);
349+ if (workspaceSize > 0) {
350+ aclrtFree(workspaceAddr);
351+ }
352+ aclrtDestroyStream(stream);
353+ aclrtResetDevice(deviceId);
354+ aclFinalize();
355+ 
356+ return 0;
357+}
358+```
Aexperimental/math/cast/examples/test_aclnn_cast.cpp+142-0
@@ -0,0 +1,142 @@
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 <vector>
13+#include "acl/acl.h"
14+#include "aclnnop/aclnn_cast.h"
15+ 
16+#define CHECK_RET(cond, return_expr) \
17+ do { \
18+ if (!(cond)) { \
19+ return_expr; \
20+ } \
21+ } while (0)
22+ 
23+#define LOG_PRINT(message, ...) \
24+ do { \
25+ printf(message, ##__VA_ARGS__); \
26+ } while (0)
27+ 
28+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
29+ int64_t shapeSize = 1;
30+ for (auto i : shape) {
31+ shapeSize *= i;
32+ }
33+ return shapeSize;
34+}
35+ 
36+int Init(int32_t deviceId, aclrtStream* stream) {
37+ // 固定写法，初始化
38+ auto ret = aclInit(nullptr);
39+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
40+ ret = aclrtSetDevice(deviceId);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtCreateStream(stream);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
44+ return 0;
45+}
46+ 
47+template <typename T>
48+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
49+ aclDataType dataType, aclTensor** tensor) {
50+ auto size = GetShapeSize(shape) * sizeof(T);
51+ // 调用aclrtMalloc申请device侧内存
52+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
53+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
54+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
55+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
57+ 
58+ // 计算连续tensor的strides
59+ std::vector<int64_t> strides(shape.size(), 1);
60+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
61+ strides[i] = shape[i + 1] * strides[i + 1];
62+ }
63+ 
64+ // 调用aclCreateTensor接口创建aclTensor
65+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
66+ shape.data(), shape.size(), *deviceAddr);
67+ return 0;
68+}
69+ 
70+int main() {
71+ // 1. （固定写法）device/stream初始化，参考acl API文档
72+ // 根据自己的实际device填写deviceId
73+ int32_t deviceId = 0;
74+ aclrtStream stream;
75+ auto ret = Init(deviceId, &stream);
76+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
77+ 
78+ // 2. 构造输入与输出，需要根据API的接口自定义构造
79+ std::vector<int64_t> selfShape = {4, 2};
80+ std::vector<int64_t> outShape = {4, 2};
81+ 
82+ void* selfDeviceAddr = nullptr;
83+ void* outDeviceAddr = nullptr;
84+ aclTensor* self = nullptr;
85+ aclTensor* out = nullptr;
86+ 
87+ std::vector<float> selfHostData = {0.1, 1.1, 2.1, 3.1, 4.1, 5.1, 6.1, 7.1};
88+ std::vector<double> outHostData = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0};
89+ 
90+ // 创建self aclTensor
91+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
92+ CHECK_RET(ret == ACL_SUCCESS, return ret);
93+ // 创建out aclTensor
94+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_DOUBLE, &out);
95+ CHECK_RET(ret == ACL_SUCCESS, return ret);
96+ 
97+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
98+ uint64_t workspaceSize = 0;
99+ aclOpExecutor* executor;
100+ // 调用aclnnCast第一段接口
101+ ret = aclnnCastGetWorkspaceSize(self, aclDataType::ACL_DOUBLE, out, &workspaceSize, &executor);
102+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnCastGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
103+ // 根据第一段接口计算出的workspaceSize申请device内存
104+ void* workspaceAddr = nullptr;
105+ if (workspaceSize > static_cast<uint64_t>(0)) {
106+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
107+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
108+ }
109+ // 调用aclnnCast第二段接口
110+ ret = aclnnCast(workspaceAddr, workspaceSize, executor, stream);
111+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnCast failed. ERROR: %d\n", ret); return ret);
112+ 
113+ // 4. （固定写法）同步等待任务执行结束
114+ ret = aclrtSynchronizeStream(stream);
115+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
116+ 
117+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
118+ auto size = GetShapeSize(outShape);
119+ std::vector<double> resultData(size, 0);
120+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
121+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
122+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
123+ for (int64_t i = 0; i < size; i++) {
124+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
125+ }
126+ 
127+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
128+ aclDestroyTensor(self);
129+ aclDestroyTensor(out);
130+ 
131+ // 7. 释放device 资源
132+ aclrtFree(selfDeviceAddr);
133+ aclrtFree(outDeviceAddr);
134+ if (workspaceSize > static_cast<uint64_t>(0)) {
135+ aclrtFree(workspaceAddr);
136+ }
137+ aclrtDestroyStream(stream);
138+ aclrtResetDevice(deviceId);
139+ aclFinalize();
140+ 
141+ return 0;
142+}
Aexperimental/math/cast/op_api/aclnn_cast.cpp+176-0
@@ -0,0 +1,176 @@
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+#include "aclnn_cast.h"
11+#include "aclnn_kernels/cast.h"
12+#include "aclnn_kernels/contiguous.h"
13+#include "aclnn_kernels/common/op_error_check.h"
14+#include "aclnn/aclnn_base.h"
15+#include "opdev/common_types.h"
16+#include "opdev/shape_utils.h"
17+#include "opdev/data_type_utils.h"
18+#include "opdev/format_utils.h"
19+#include "opdev/op_dfx.h"
20+#include "opdev/op_executor.h"
21+#include "opdev/op_log.h"
22+#include "opdev/tensor_view_utils.h"
23+#include "opdev/platform.h"
24+ 
25+using namespace op;
26+#ifdef __cplusplus
27+extern "C" {
28+#endif
29+ 
30+/* Cast 算子的完整计算流程如下:
31+ * self dtype
32+ * \ /
33+ * Contiguous(workspace_0) /
34+ * \ /
35+ * Cast(workspace_1)
36+ * |
37+ * ViewCopy
38+ * |
39+ * result
40+ */
41+ 
42+static const size_t MAX_DIM = 8;
43+ 
44+// 根据API定义，需要列出所能支持的所有dtype
45+static const std::initializer_list<op::DataType> ASCEND910_DTYPE_SUPPORT_LIST = {
46+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_DOUBLE, op::DataType::DT_INT8,
47+ op::DataType::DT_UINT8, op::DataType::DT_INT16, op::DataType::DT_INT32, op::DataType::DT_INT64,
48+ op::DataType::DT_UINT16, op::DataType::DT_UINT32, op::DataType::DT_UINT64, op::DataType::DT_BOOL,
49+ op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128};
50+static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST_DEFAULT = {
51+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_DOUBLE, op::DataType::DT_INT8,
52+ op::DataType::DT_UINT8, op::DataType::DT_INT16, op::DataType::DT_INT32, op::DataType::DT_INT64,
53+ op::DataType::DT_UINT16, op::DataType::DT_UINT32, op::DataType::DT_UINT64, op::DataType::DT_BOOL,
54+ op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128, op::DataType::DT_BF16};
55+static const std::initializer_list<op::DataType> ASCEND950_DTYPE_SUPPORT_LIST = {
56+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_DOUBLE,
57+ op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_INT16,
58+ op::DataType::DT_INT32, op::DataType::DT_INT64, op::DataType::DT_UINT16,
59+ op::DataType::DT_UINT32, op::DataType::DT_UINT64, op::DataType::DT_BOOL,
60+ op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128, op::DataType::DT_BF16,
61+ op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2, op::DataType::DT_FLOAT8_E4M3FN,
62+ op::DataType::DT_COMPLEX32, op::DataType::DT_FLOAT4_E1M2, op::DataType::DT_FLOAT4_E2M1};
63+static const std::initializer_list<op::DataType> ASCEND950_SELF_DTYPE_SUPPORT_LIST = {
64+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_DOUBLE,
65+ op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_INT16,
66+ op::DataType::DT_INT32, op::DataType::DT_INT64, op::DataType::DT_UINT16,
67+ op::DataType::DT_UINT32, op::DataType::DT_UINT64, op::DataType::DT_BOOL,
68+ op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128, op::DataType::DT_BF16,
69+ op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2, op::DataType::DT_FLOAT8_E4M3FN,
70+ op::DataType::DT_COMPLEX32, op::DataType::DT_FLOAT4_E1M2, op::DataType::DT_FLOAT4_E2M1};
71+ 
72+static bool CheckNotNull(const aclTensor* self, const aclTensor* out)
73+{
74+ OP_CHECK_NULL(self, return false);
75+ OP_CHECK_NULL(out, return false);
76+ return true;
77+}
78+ 
79+static bool CheckDtypeValid(const aclTensor* self, const DataType dtype)
80+{
81+ // 检查self的数据类型是否在算子的支持列表内
82+ bool isASCEND910B = (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201);
83+ bool isASCEND910_93 = (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201);
84+ bool isASCEND950 = (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_3510);
85+ bool isAscend910BC = isASCEND910B || isASCEND910_93;
86+ 
87+ auto supportList = ASCEND910_DTYPE_SUPPORT_LIST;
88+ auto selfSupportList = ASCEND910_DTYPE_SUPPORT_LIST;
89+ if (isAscend910BC) {
90+ supportList = DTYPE_SUPPORT_LIST_DEFAULT;
91+ selfSupportList = DTYPE_SUPPORT_LIST_DEFAULT;
92+ } else if (isASCEND950) {
93+ supportList = ASCEND950_DTYPE_SUPPORT_LIST;
94+ selfSupportList = ASCEND950_SELF_DTYPE_SUPPORT_LIST;
95+ }
96+ OP_CHECK_DTYPE_NOT_SUPPORT(self, selfSupportList, return false);
97+ bool isSupport = CheckType(dtype, supportList);
98+ // 检查参数dtype是否在Cast算子的输出数据类型支持列表内
99+ if (!isSupport) {
100+ OP_LOGE(
101+ ACLNN_ERR_PARAM_INVALID, "The param dtype not implemented for %s, should be in dtype support list %s.",
102+ op::ToString(dtype).GetString(), op::ToString(supportList).GetString());
103+ return false;
104+ }
105+ return true;
106+}
107+ 
108+static inline bool CheckShape(const aclTensor* self, const aclTensor* out)
109+{
110+ OP_CHECK_MAX_DIM(self, MAX_DIM, return false);
111+ OP_CHECK_SHAPE_NOT_EQUAL(out, self, return false);
112+ return true;
113+}
114+ 
115+static inline aclnnStatus CheckParams(const aclTensor* self, const DataType dtype, const aclTensor* out)
116+{
117+ // 1. 检查参数是否为空指针
118+ CHECK_RET(CheckNotNull(self, out), ACLNN_ERR_PARAM_NULLPTR);
119+ 
120+ // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
121+ CHECK_RET(CheckDtypeValid(self, dtype), ACLNN_ERR_PARAM_INVALID);
122+ 
123+ // 3. 检查输入tensor的shape是否为异常，输出和输入的shape是否相同
124+ CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID);
125+ 
126+ return ACLNN_SUCCESS;
127+}
128+ 
129+aclnnStatus aclnnCastGetWorkspaceSize(
130+ const aclTensor* self, const aclDataType dtype, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)
131+{
132+ L2_DFX_PHASE_1(aclnnCast, DFX_IN(self, dtype), DFX_OUT(out));
133+ 
134+ // 固定写法，创建OpExecutor
135+ auto uniqueExecutor = CREATE_EXECUTOR();
136+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
137+ 
138+ // 固定写法，参数检查
139+ auto ret = CheckParams(self, op::ToOpDataType(dtype), out);
140+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
141+ 
142+ // 输入为空tensor时，直接返回dtype类型的空tensor
143+ if (self->IsEmpty()) {
144+ *workspaceSize = 0;
145+ uniqueExecutor.ReleaseTo(executor);
146+ return ACLNN_SUCCESS;
147+ }
148+ 
149+ // 固定写法，将输入self转换成连续的tensor
150+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
151+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
152+ 
153+ // 调用Cast算子kernel,将输入self的数据类型转换成指定的数据类型
154+ auto castOut = l0op::Cast(selfContiguous, op::ToOpDataType(dtype), uniqueExecutor.get());
155+ CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
156+ 
157+ // 固定写法，将计算结果拷贝到输出out上，out可能是非连续的tensor
158+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
159+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
160+ 
161+ // 固定写法，获取计算过程中需要使用的workspace大小
162+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
163+ uniqueExecutor.ReleaseTo(executor);
164+ return ACLNN_SUCCESS;
165+}
166+ 
167+aclnnStatus aclnnCast(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
168+{
169+ L2_DFX_PHASE_2(aclnnCast);
170+ // 固定写法，调用框架能力，完成计算
171+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
172+}
173+ 
174+#ifdef __cplusplus
175+}
176+#endif
Aexperimental/math/cast/op_api/aclnn_cast.h+57-0
@@ -0,0 +1,57 @@
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_CAST_H_
12+#define OP_API_INC_CAST_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+ 
17+#ifdef __cplusplus
18+extern "C" {
19+#endif
20+ 
21+/**
22+ * @brief aclnnCast的第一段接口，根据具体的计算流程，计算workspace大小。
23+ * @domain aclnn_math
24+ *
25+ * 算子功能：将输入tensor转换为指定的dtype类型。
26+ *
27+ * @param [in] self: npu
28+ * device侧的aclTensor，数据类型支持FLOAT16、FLOAT、FlOAT64、INT8、UINT8、INT16、INT32、INT64、BOOL。
29+ * 支持非连续的Tensor，数据格式支持ND。
30+ * @param [in] dtype: host侧的aclDataType，输入tensor要转换的目标dtype。
31+ * @param [in] out: npu
32+ * device侧的aclTensor，数据类型支持FLOAT16、FLOAT、FlOAT64、INT8、UINT8、INT16、INT32、INT64、BOOL、
33+ * COMPLEX64、COMPLEX128。数据类型为dtype，shape与self相同，数据格式支持ND，且数据格式需要与self一致。
34+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
35+ * @param [out] executor: 返回op执行器，包含算子计算流程。
36+ * @return aclnnStatus: 返回状态码。
37+ */
38+ACLNN_API aclnnStatus aclnnCastGetWorkspaceSize(
39+ const aclTensor* self, const aclDataType dtype, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
40+/**
41+ * @brief aclnnCast的第二段接口，用于执行计算。
42+ *
43+ * 算子功能：将输入tensor转换为指定的dtype类型。
44+ *
45+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
46+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnCastGetWorkspaceSize获取。
47+ * @param [in] executor: op执行器，包含了算子计算流程。
48+ * @param [in] stream: acl stream流。
49+ * @return aclnnStatus: 返回状态码。
50+ */
51+ACLNN_API aclnnStatus aclnnCast(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
52+ 
53+#ifdef __cplusplus
54+}
55+#endif
56+ 
57+#endif // OP_API_INC_CAST_H_
Aexperimental/math/cast/op_host/cast_def.cpp+40-0
@@ -0,0 +1,40 @@
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 cast_def.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class Cast : public OpDef {
19+public:
20+ explicit Cast(const char* name) : OpDef(name)
21+ {
22+ this->Input("x")
23+ .ParamType(REQUIRED)
24+ .DataType({ ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16 , ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT , ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32 , ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8 , ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8 , ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL , ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64 , ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16 , ge::DT_INT16, ge::DT_INT16, ge::DT_INT16, ge::DT_INT16, ge::DT_INT16, ge::DT_INT16 })
CANN-robot
CANN-robotCANN-robot1月4日

代码结构与可维护性: 输入数据类型列表存在明显的重复模式，代码可读性差且难以维护。从代码可以看出，同一数据类型（如ge::DT_FLOAT16）在列表中重复出现了7次，ge::DT_FLOAT重复了8次，ge::DT_INT32重复了8次等。这种硬编码的重复列表不仅难以阅读，也容易在修改时出错。应该使用更清晰的数据结构或生成方式来定义这些类型映射关系。

问题类型: 代码结构与可维护性 文件路径: experimental/math/cast/op_host/cast_def.cpp 行号: 25 问题代码:

.DataType({             ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16   , ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT            , ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32            , ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8                    , ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8          , ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL                 , ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64           , ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16              , ge::DT_INT16, ge::DT_INT16, ge::DT_INT16, ge::DT_INT16, ge::DT_INT16, ge::DT_INT16        })

修改建议:

建议重构代码，使用更清晰的数据结构来定义输入输出类型映射。可以考虑：1) 使用数组或向量来存储支持的数据类型；2) 使用循环或模板来生成类型列表；3) 将类型映射关系定义为常量数组，然后通过循环添加。这样既能提高代码可读性，也便于后续维护和扩展。

此评论由代码审查工具自动生成

likedislike
25+ .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND })
CANN-robot
CANN-robotCANN-robot1月4日

代码结构与可维护性: Format和UnknownShapeFormat列表包含大量重复的ge::FORMAT_ND值，这些重复的硬编码值使得代码难以阅读和维护。从代码可以看出，每个数据类型对应一个FORMAT_ND格式，但代码中重复写了数十次相同的值。这种重复不仅增加了代码体积，也容易在修改时出现不一致的情况。

问题类型: 代码结构与可维护性 文件路径: experimental/math/cast/op_host/cast_def.cpp 行号: 26 问题代码:

.Format({               ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND          , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND    , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND    , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND    , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND   , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND   , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND   , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND  , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND  })

修改建议:

建议简化格式定义：1) 如果所有数据类型都使用相同的格式（FORMAT_ND），可以考虑使用一个更简洁的API来设置默认格式；2) 或者使用循环来为每个数据类型设置相同的格式值；3) 如果框架支持，可以使用默认值或省略这些重复的设置。

此评论由代码审查工具自动生成

likedislike
26+ .UnknownShapeFormat({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND });
27+ this->Output("y")
28+ .ParamType(REQUIRED)
29+ .DataType({ ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32, ge::DT_INT16, ge::DT_UINT8, ge::DT_BOOL, ge::DT_BF16 , ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT32, ge::DT_INT64, ge::DT_BOOL, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16 , ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_INT64, ge::DT_BOOL , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT8, ge::DT_BOOL, ge::DT_BF16, ge::DT_INT16, ge::DT_INT64 , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT8, ge::DT_INT16, ge::DT_BF16, ge::DT_INT64 , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8, ge::DT_BF16, ge::DT_INT64 , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32, ge::DT_UINT8, ge::DT_BOOL, ge::DT_INT16, ge::DT_BF16 , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32, ge::DT_UINT8, ge::DT_BOOL , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32, ge::DT_UINT8, ge::DT_INT64 })
CANN-robot
CANN-robotCANN-robot1月4日

代码结构与可维护性: 输出数据类型列表的组织方式不清晰，难以理解输入类型和输出类型之间的映射关系。代码中输出类型列表似乎是按照某种模式组织的，但没有明确的注释说明这种映射关系。这种隐式的映射关系使得代码难以理解和维护，特别是当需要添加新的类型支持时。

问题类型: 代码结构与可维护性 文件路径: experimental/math/cast/op_host/cast_def.cpp 行号: 30 问题代码:

.DataType({             ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32, ge::DT_INT16, ge::DT_UINT8, ge::DT_BOOL, ge::DT_BF16                    , ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT32, ge::DT_INT64, ge::DT_BOOL, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16             , ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_INT64, ge::DT_BOOL             , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_UINT8, ge::DT_BOOL, ge::DT_BF16, ge::DT_INT16, ge::DT_INT64            , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT8, ge::DT_INT16, ge::DT_BF16, ge::DT_INT64          , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8, ge::DT_BF16, ge::DT_INT64          , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32, ge::DT_UINT8, ge::DT_BOOL, ge::DT_INT16, ge::DT_BF16            , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32, ge::DT_UINT8, ge::DT_BOOL        , ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_INT32, ge::DT_UINT8, ge::DT_INT64       })

修改建议:

建议：1) 添加详细的注释说明输入类型和输出类型之间的映射关系；2) 使用更清晰的数据结构来定义类型转换矩阵，例如使用std::map或std::unordered_map来存储源类型到目标类型的映射；3) 考虑将类型映射关系提取到单独的配置文件中，提高可维护性。

此评论由代码审查工具自动生成

likedislike
30+ .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND })
31+ .UnknownShapeFormat({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND , ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND });
32+ this->Attr("dst_type")
33+ .AttrType(REQUIRED)
34+ .Int();
35+ 
36+ this->AICore().AddConfig("ascend910b");
CANN-robot
CANN-robotCANN-robot1月4日

代码结构与可维护性: 硬编码了AI Core配置为"ascend910b"，这限制了代码的可移植性和灵活性。如果将来需要支持其他AI Core型号，需要修改代码。这种硬编码的配置值应该定义为常量或从配置中读取。

问题类型: 代码结构与可维护性 文件路径: experimental/math/cast/op_host/cast_def.cpp 行号: 37 问题代码:

this->AICore().AddConfig("ascend910b");

修改建议:

建议：1) 将"ascend910b"定义为常量或配置项；2) 考虑支持多种AI Core配置；3) 如果可能，从配置文件或环境变量中读取AI Core配置。

此评论由代码审查工具自动生成

likedislike
37+ }
38+};
39+OP_ADD(Cast); // 添加算子信息库
40+} // namespace ops
Aexperimental/math/cast/op_host/cast_infershape.cpp+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 cast_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+ 
22+static ge::graphStatus InferShapeCast(gert::InferShapeContext* context)
23+{
24+ const gert::Shape* x1_shape = context->GetInputShape(0);
25+ gert::Shape* y_shape = context->GetOutputShape(0);
26+ *y_shape = *x1_shape;
27+ return GRAPH_SUCCESS;
28+}
29+ 
30+IMPL_OP_INFERSHAPE(Cast).InferShape(InferShapeCast);
31+} // namespace ops
Aexperimental/math/cast/op_host/cast_tiling.cpp+325-0
@@ -0,0 +1,325 @@
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 cast_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "op_host/tiling_util.h"
19+#include "op_host/tiling_templates_registry.h"
20+#include "../op_kernel/cast_tiling_data.h"
21+#include "../op_kernel/cast_tiling_key.h"
22+ 
23+namespace optiling {
24+ 
25+const uint64_t BLOCK_SIZE = 32;
26+const uint64_t BUFFER_NUM = 2;
27+uint64_t ubDataNumMap[40][40] = {};
28+uint64_t tilingKeyMap[40][40] = {};
29+uint64_t minDataTypeLengthMap[40][40] = {};
30+void UbDataNumMapInit()
31+{
32+ // = InputBytes * BUFFER_NUM + OutputBytes * BUFFER_NUM + AlltmpBytes
33+ ubDataNumMap[ge::DT_FLOAT16][ge::DT_FLOAT] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
34+ ubDataNumMap[ge::DT_FLOAT16][ge::DT_INT32] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
35+ ubDataNumMap[ge::DT_FLOAT16][ge::DT_INT8] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 4;
36+ ubDataNumMap[ge::DT_FLOAT16][ge::DT_UINT8] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 4;
37+ ubDataNumMap[ge::DT_FLOAT16][ge::DT_BOOL] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
38+ ubDataNumMap[ge::DT_FLOAT16][ge::DT_INT16] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
39+ ubDataNumMap[ge::DT_FLOAT16][ge::DT_BF16] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
40+ ubDataNumMap[ge::DT_FLOAT][ge::DT_FLOAT16] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
41+ ubDataNumMap[ge::DT_FLOAT][ge::DT_BF16] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
42+ ubDataNumMap[ge::DT_FLOAT][ge::DT_INT32] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
43+ ubDataNumMap[ge::DT_FLOAT][ge::DT_INT64] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
44+ ubDataNumMap[ge::DT_FLOAT][ge::DT_BOOL] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 4;
45+ ubDataNumMap[ge::DT_FLOAT][ge::DT_INT8] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 6;
46+ ubDataNumMap[ge::DT_FLOAT][ge::DT_UINT8] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 6;
47+ ubDataNumMap[ge::DT_FLOAT][ge::DT_INT16] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
48+ ubDataNumMap[ge::DT_INT32][ge::DT_FLOAT] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
49+ ubDataNumMap[ge::DT_INT32][ge::DT_FLOAT16] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
50+ ubDataNumMap[ge::DT_INT32][ge::DT_BF16] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
51+ ubDataNumMap[ge::DT_INT32][ge::DT_INT8] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 6;
52+ ubDataNumMap[ge::DT_INT32][ge::DT_UINT8] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 6;
53+ ubDataNumMap[ge::DT_INT32][ge::DT_INT16] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
54+ ubDataNumMap[ge::DT_INT32][ge::DT_INT64] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
55+ ubDataNumMap[ge::DT_INT32][ge::DT_BOOL] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 4;
56+ ubDataNumMap[ge::DT_INT8][ge::DT_FLOAT16] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
57+ ubDataNumMap[ge::DT_INT8][ge::DT_FLOAT] = 1 * BUFFER_NUM + 4 * BUFFER_NUM + 2;
58+ ubDataNumMap[ge::DT_INT8][ge::DT_INT32] = 1 * BUFFER_NUM + 4 * BUFFER_NUM + 2;
59+ ubDataNumMap[ge::DT_INT8][ge::DT_UINT8] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
60+ ubDataNumMap[ge::DT_INT8][ge::DT_BOOL] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
61+ ubDataNumMap[ge::DT_INT8][ge::DT_INT16] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 2;
62+ ubDataNumMap[ge::DT_INT8][ge::DT_INT64] = 1 * BUFFER_NUM + 8 * BUFFER_NUM + 6;
63+ ubDataNumMap[ge::DT_INT8][ge::DT_BF16] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 6;
64+ ubDataNumMap[ge::DT_UINT8][ge::DT_FLOAT16] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
65+ ubDataNumMap[ge::DT_UINT8][ge::DT_FLOAT] = 1 * BUFFER_NUM + 4 * BUFFER_NUM + 2;
66+ ubDataNumMap[ge::DT_UINT8][ge::DT_INT32] = 1 * BUFFER_NUM + 4 * BUFFER_NUM + 2;
67+ ubDataNumMap[ge::DT_UINT8][ge::DT_INT8] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 6;
68+ ubDataNumMap[ge::DT_UINT8][ge::DT_INT16] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 2;
69+ ubDataNumMap[ge::DT_UINT8][ge::DT_INT64] = 1 * BUFFER_NUM + 8 * BUFFER_NUM + 6;
70+ ubDataNumMap[ge::DT_UINT8][ge::DT_BF16] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 6;
71+ ubDataNumMap[ge::DT_BOOL][ge::DT_FLOAT16] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
72+ ubDataNumMap[ge::DT_BOOL][ge::DT_FLOAT] = 1 * BUFFER_NUM + 4 * BUFFER_NUM + 2;
73+ ubDataNumMap[ge::DT_BOOL][ge::DT_INT32] = 1 * BUFFER_NUM + 4 * BUFFER_NUM + 2;
74+ ubDataNumMap[ge::DT_BOOL][ge::DT_UINT8] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
75+ ubDataNumMap[ge::DT_BOOL][ge::DT_INT8] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
76+ ubDataNumMap[ge::DT_BOOL][ge::DT_INT64] = 1 * BUFFER_NUM + 8 * BUFFER_NUM + 6;
77+ ubDataNumMap[ge::DT_BOOL][ge::DT_BF16] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 6;
78+ ubDataNumMap[ge::DT_INT64][ge::DT_FLOAT16] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
79+ ubDataNumMap[ge::DT_INT64][ge::DT_FLOAT] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
80+ ubDataNumMap[ge::DT_INT64][ge::DT_INT32] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
81+ ubDataNumMap[ge::DT_INT64][ge::DT_UINT8] = 8 * BUFFER_NUM + 1 * BUFFER_NUM + 10;
82+ ubDataNumMap[ge::DT_INT64][ge::DT_INT8] = 8 * BUFFER_NUM + 1 * BUFFER_NUM + 10;
83+ ubDataNumMap[ge::DT_INT64][ge::DT_BOOL] = 8 * BUFFER_NUM + 1 * BUFFER_NUM + 6;
84+ ubDataNumMap[ge::DT_INT64][ge::DT_BF16] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
85+ ubDataNumMap[ge::DT_INT64][ge::DT_INT16] = 4 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
86+ ubDataNumMap[ge::DT_BF16][ge::DT_FLOAT16] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 2;
87+ ubDataNumMap[ge::DT_BF16][ge::DT_FLOAT] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
88+ ubDataNumMap[ge::DT_BF16][ge::DT_INT32] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
89+ ubDataNumMap[ge::DT_BF16][ge::DT_INT8] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 10;
90+ ubDataNumMap[ge::DT_BF16][ge::DT_UINT8] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 10;
91+ ubDataNumMap[ge::DT_BF16][ge::DT_BOOL] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 6;
92+ ubDataNumMap[ge::DT_INT16][ge::DT_FLOAT16] = 1 * BUFFER_NUM + 1 * BUFFER_NUM + 0;
93+ ubDataNumMap[ge::DT_INT16][ge::DT_FLOAT] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 0;
94+ ubDataNumMap[ge::DT_INT16][ge::DT_INT32] = 1 * BUFFER_NUM + 2 * BUFFER_NUM + 2;
95+ ubDataNumMap[ge::DT_INT16][ge::DT_INT8] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 4;
96+ ubDataNumMap[ge::DT_INT16][ge::DT_UINT8] = 2 * BUFFER_NUM + 1 * BUFFER_NUM + 4;
97+ ubDataNumMap[ge::DT_INT16][ge::DT_INT64] = 1 * BUFFER_NUM + 4 * BUFFER_NUM + 2;
98+}
99+void TilingKeyMapInit()
100+{
101+ // = Tiling Key
102+ tilingKeyMap[ge::DT_FLOAT16][ge::DT_FLOAT] = 1;
103+ tilingKeyMap[ge::DT_FLOAT16][ge::DT_INT32] = 1;
104+ tilingKeyMap[ge::DT_FLOAT16][ge::DT_INT8] = 3;
105+ tilingKeyMap[ge::DT_FLOAT16][ge::DT_UINT8] = 3;
106+ tilingKeyMap[ge::DT_FLOAT16][ge::DT_BOOL] = 1;
107+ tilingKeyMap[ge::DT_FLOAT16][ge::DT_INT16] = 1;
108+ tilingKeyMap[ge::DT_FLOAT16][ge::DT_BF16] = 2;
109+ tilingKeyMap[ge::DT_FLOAT][ge::DT_FLOAT16] = 1;
110+ tilingKeyMap[ge::DT_FLOAT][ge::DT_BF16] = 1;
111+ tilingKeyMap[ge::DT_FLOAT][ge::DT_INT32] = 1;
112+ tilingKeyMap[ge::DT_FLOAT][ge::DT_INT64] = 1;
113+ tilingKeyMap[ge::DT_FLOAT][ge::DT_BOOL] = 3;
114+ tilingKeyMap[ge::DT_FLOAT][ge::DT_INT8] = 4;
115+ tilingKeyMap[ge::DT_FLOAT][ge::DT_UINT8] = 4;
116+ tilingKeyMap[ge::DT_FLOAT][ge::DT_INT16] = 1;
117+ tilingKeyMap[ge::DT_INT32][ge::DT_FLOAT] = 1;
118+ tilingKeyMap[ge::DT_INT32][ge::DT_FLOAT16] = 2;
119+ tilingKeyMap[ge::DT_INT32][ge::DT_BF16] = 2;
120+ tilingKeyMap[ge::DT_INT32][ge::DT_INT8] = 4;
121+ tilingKeyMap[ge::DT_INT32][ge::DT_UINT8] = 4;
122+ tilingKeyMap[ge::DT_INT32][ge::DT_INT16] = 1;
123+ tilingKeyMap[ge::DT_INT32][ge::DT_INT64] = 1;
124+ tilingKeyMap[ge::DT_INT32][ge::DT_BOOL] = 3;
125+ tilingKeyMap[ge::DT_INT8][ge::DT_FLOAT16] = 1;
126+ tilingKeyMap[ge::DT_INT8][ge::DT_FLOAT] = 5;
127+ tilingKeyMap[ge::DT_INT8][ge::DT_INT32] = 5;
128+ tilingKeyMap[ge::DT_INT8][ge::DT_UINT8] = 8;
129+ tilingKeyMap[ge::DT_INT8][ge::DT_BOOL] = 5;
130+ tilingKeyMap[ge::DT_INT8][ge::DT_INT16] = 5;
131+ tilingKeyMap[ge::DT_INT8][ge::DT_INT64] = 6;
132+ tilingKeyMap[ge::DT_INT8][ge::DT_BF16] = 6;
133+ tilingKeyMap[ge::DT_UINT8][ge::DT_FLOAT16] = 1;
134+ tilingKeyMap[ge::DT_UINT8][ge::DT_FLOAT] = 5;
135+ tilingKeyMap[ge::DT_UINT8][ge::DT_INT32] = 5;
136+ tilingKeyMap[ge::DT_UINT8][ge::DT_INT8] = 8;
137+ tilingKeyMap[ge::DT_UINT8][ge::DT_INT16] = 5;
138+ tilingKeyMap[ge::DT_UINT8][ge::DT_INT64] = 6;
139+ tilingKeyMap[ge::DT_UINT8][ge::DT_BF16] = 6;
140+ tilingKeyMap[ge::DT_BOOL][ge::DT_FLOAT16] = 1;
141+ tilingKeyMap[ge::DT_BOOL][ge::DT_FLOAT] = 5;
142+ tilingKeyMap[ge::DT_BOOL][ge::DT_INT32] = 5;
143+ tilingKeyMap[ge::DT_BOOL][ge::DT_UINT8] = 8;
144+ tilingKeyMap[ge::DT_BOOL][ge::DT_INT8] = 8;
145+ tilingKeyMap[ge::DT_BOOL][ge::DT_INT64] = 6;
146+ tilingKeyMap[ge::DT_BOOL][ge::DT_BF16] = 6;
147+ tilingKeyMap[ge::DT_INT64][ge::DT_FLOAT16] = 2;
148+ tilingKeyMap[ge::DT_INT64][ge::DT_FLOAT] = 1;
149+ tilingKeyMap[ge::DT_INT64][ge::DT_INT32] = 1;
150+ tilingKeyMap[ge::DT_INT64][ge::DT_UINT8] = 7;
151+ tilingKeyMap[ge::DT_INT64][ge::DT_INT8] = 7;
152+ tilingKeyMap[ge::DT_INT64][ge::DT_BOOL] = 6;
153+ tilingKeyMap[ge::DT_INT64][ge::DT_BF16] = 2;
154+ tilingKeyMap[ge::DT_INT64][ge::DT_INT16] = 2;
155+ tilingKeyMap[ge::DT_BF16][ge::DT_FLOAT16] = 2;
156+ tilingKeyMap[ge::DT_BF16][ge::DT_FLOAT] = 1;
157+ tilingKeyMap[ge::DT_BF16][ge::DT_INT32] = 1;
158+ tilingKeyMap[ge::DT_BF16][ge::DT_INT8] = 7;
159+ tilingKeyMap[ge::DT_BF16][ge::DT_UINT8] = 7;
160+ tilingKeyMap[ge::DT_BF16][ge::DT_BOOL] = 6;
161+ tilingKeyMap[ge::DT_INT16][ge::DT_FLOAT16] = 1;
162+ tilingKeyMap[ge::DT_INT16][ge::DT_FLOAT] = 1;
163+ tilingKeyMap[ge::DT_INT16][ge::DT_INT32] = 2;
164+ tilingKeyMap[ge::DT_INT16][ge::DT_INT8] = 3;
165+ tilingKeyMap[ge::DT_INT16][ge::DT_UINT8] = 3;
166+ tilingKeyMap[ge::DT_INT16][ge::DT_INT64] = 2;
167+}
168+void MinDataTypeLengthMapInit()
169+{
170+ minDataTypeLengthMap[ge::DT_FLOAT16][ge::DT_FLOAT] = 2;
171+ minDataTypeLengthMap[ge::DT_FLOAT16][ge::DT_INT32] = 2;
172+ minDataTypeLengthMap[ge::DT_FLOAT16][ge::DT_INT8] = 1;
173+ minDataTypeLengthMap[ge::DT_FLOAT16][ge::DT_UINT8] = 1;
174+ minDataTypeLengthMap[ge::DT_FLOAT16][ge::DT_BOOL] = 1;
175+ minDataTypeLengthMap[ge::DT_FLOAT16][ge::DT_INT16] = 2;
176+ minDataTypeLengthMap[ge::DT_FLOAT16][ge::DT_BF16] = 2;
177+ minDataTypeLengthMap[ge::DT_FLOAT][ge::DT_FLOAT16] = 2;
178+ minDataTypeLengthMap[ge::DT_FLOAT][ge::DT_BF16] = 2;
179+ minDataTypeLengthMap[ge::DT_FLOAT][ge::DT_INT32] = 4;
180+ minDataTypeLengthMap[ge::DT_FLOAT][ge::DT_INT64] = 4;
181+ minDataTypeLengthMap[ge::DT_FLOAT][ge::DT_BOOL] = 1;
182+ minDataTypeLengthMap[ge::DT_FLOAT][ge::DT_INT8] = 1;
183+ minDataTypeLengthMap[ge::DT_FLOAT][ge::DT_UINT8] = 1;
184+ minDataTypeLengthMap[ge::DT_FLOAT][ge::DT_INT16] = 2;
185+ minDataTypeLengthMap[ge::DT_INT32][ge::DT_FLOAT] = 4;
186+ minDataTypeLengthMap[ge::DT_INT32][ge::DT_FLOAT16] = 2;
187+ minDataTypeLengthMap[ge::DT_INT32][ge::DT_BF16] = 2;
188+ minDataTypeLengthMap[ge::DT_INT32][ge::DT_INT8] = 1;
189+ minDataTypeLengthMap[ge::DT_INT32][ge::DT_UINT8] = 1;
190+ minDataTypeLengthMap[ge::DT_INT32][ge::DT_INT16] = 2;
191+ minDataTypeLengthMap[ge::DT_INT32][ge::DT_INT64] = 4;
192+ minDataTypeLengthMap[ge::DT_INT32][ge::DT_BOOL] = 1;
193+ minDataTypeLengthMap[ge::DT_INT8][ge::DT_FLOAT16] = 1;
194+ minDataTypeLengthMap[ge::DT_INT8][ge::DT_FLOAT] = 1;
195+ minDataTypeLengthMap[ge::DT_INT8][ge::DT_INT32] = 1;
196+ minDataTypeLengthMap[ge::DT_INT8][ge::DT_UINT8] = 1;
197+ minDataTypeLengthMap[ge::DT_INT8][ge::DT_BOOL] = 1;
198+ minDataTypeLengthMap[ge::DT_INT8][ge::DT_INT16] = 1;
199+ minDataTypeLengthMap[ge::DT_INT8][ge::DT_INT64] = 1;
200+ minDataTypeLengthMap[ge::DT_INT8][ge::DT_BF16] = 1;
201+ minDataTypeLengthMap[ge::DT_UINT8][ge::DT_FLOAT16] = 1;
202+ minDataTypeLengthMap[ge::DT_UINT8][ge::DT_FLOAT] = 1;
203+ minDataTypeLengthMap[ge::DT_UINT8][ge::DT_INT32] = 1;
204+ minDataTypeLengthMap[ge::DT_UINT8][ge::DT_INT8] = 1;
205+ minDataTypeLengthMap[ge::DT_UINT8][ge::DT_INT16] = 1;
206+ minDataTypeLengthMap[ge::DT_UINT8][ge::DT_INT64] = 1;
207+ minDataTypeLengthMap[ge::DT_UINT8][ge::DT_BF16] = 1;
208+ minDataTypeLengthMap[ge::DT_BOOL][ge::DT_FLOAT16] = 1;
209+ minDataTypeLengthMap[ge::DT_BOOL][ge::DT_FLOAT] = 1;
210+ minDataTypeLengthMap[ge::DT_BOOL][ge::DT_INT32] = 1;
211+ minDataTypeLengthMap[ge::DT_BOOL][ge::DT_UINT8] = 1;
212+ minDataTypeLengthMap[ge::DT_BOOL][ge::DT_INT8] = 1;
213+ minDataTypeLengthMap[ge::DT_BOOL][ge::DT_INT64] = 1;
214+ minDataTypeLengthMap[ge::DT_BOOL][ge::DT_BF16] = 1;
215+ minDataTypeLengthMap[ge::DT_INT64][ge::DT_FLOAT16] = 2;
216+ minDataTypeLengthMap[ge::DT_INT64][ge::DT_FLOAT] = 4;
217+ minDataTypeLengthMap[ge::DT_INT64][ge::DT_INT32] = 4;
218+ minDataTypeLengthMap[ge::DT_INT64][ge::DT_UINT8] = 1;
219+ minDataTypeLengthMap[ge::DT_INT64][ge::DT_INT8] = 1;
220+ minDataTypeLengthMap[ge::DT_INT64][ge::DT_BOOL] = 1;
221+ minDataTypeLengthMap[ge::DT_INT64][ge::DT_BF16] = 2;
222+ minDataTypeLengthMap[ge::DT_INT64][ge::DT_INT16] = 2;
223+ minDataTypeLengthMap[ge::DT_BF16][ge::DT_FLOAT16] = 2;
224+ minDataTypeLengthMap[ge::DT_BF16][ge::DT_FLOAT] = 2;
225+ minDataTypeLengthMap[ge::DT_BF16][ge::DT_INT32] = 2;
226+ minDataTypeLengthMap[ge::DT_BF16][ge::DT_INT8] = 1;
227+ minDataTypeLengthMap[ge::DT_BF16][ge::DT_UINT8] = 1;
228+ minDataTypeLengthMap[ge::DT_BF16][ge::DT_BOOL] = 1;
229+ minDataTypeLengthMap[ge::DT_INT16][ge::DT_FLOAT16] = 2;
230+ minDataTypeLengthMap[ge::DT_INT16][ge::DT_FLOAT] = 2;
231+ minDataTypeLengthMap[ge::DT_INT16][ge::DT_INT32] = 2;
232+ minDataTypeLengthMap[ge::DT_INT16][ge::DT_INT8] = 1;
233+ minDataTypeLengthMap[ge::DT_INT16][ge::DT_UINT8] = 1;
234+ minDataTypeLengthMap[ge::DT_INT16][ge::DT_INT64] = 2;
235+}
236+ 
237+// tiling 分发入口
238+static ge::graphStatus CastTilingFunc(gert::TilingContext* context)
239+{
240+ CastTilingData* tiling = context->GetTilingData<CastTilingData>();
241+ uint64_t ubSize;
242+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
243+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
244+ auto coreNum = ascendcPlatform.GetCoreNum();
245+ auto socVersion = ascendcPlatform.GetSocVersion();
246+ if (socVersion != platform_ascendc::SocVersion::ASCEND910B && socVersion != platform_ascendc::SocVersion::ASCEND310P && context->GetInputDesc(0)->GetDataType() == ge::DT_BF16) {
247+ OP_LOGE(context, "socVersion is invalid");
248+ return ge::GRAPH_FAILED;
F
Ffulltower4月8日

异常返回添加日志打印

likedislike
skywang2
skywang2
4月8日 评论：
249+ }
250+ const gert::RuntimeAttrs *attrs = context->GetAttrs();
251+ const int32_t *dst_type = attrs->GetAttrPointer<int32_t>(0);
252+ if (*dst_type != context->GetOutputDesc(0)->GetDataType()) {
253+ OP_LOGE(context, "dst_type is invalid");
254+ return ge::GRAPH_FAILED;
F
Ffulltower4月8日

异常返回添加日志打印

likedislike
skywang2
skywang2
4月8日 评论：
255+ }
256+ 
257+ uint64_t inputNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
258+ if (inputNum == 0) {
F
Ffulltower4月8日

异常返回添加日志打印

likedislike
skywang2
skywang2
4月8日 评论：
259+ OP_LOGE(context, "inputNum is 0");
260+ return ge::GRAPH_FAILED;
261+ }
262+ UbDataNumMapInit();
263+ TilingKeyMapInit();
264+ MinDataTypeLengthMapInit();
265+ auto inputDatatype = context->GetInputDesc(0)->GetDataType();
266+ auto outputDatatype = context->GetOutputDesc(0)->GetDataType();
267+ uint64_t tilingKey = tilingKeyMap[inputDatatype][outputDatatype];
268+ uint64_t typeLength = minDataTypeLengthMap[inputDatatype][outputDatatype];
269+ uint64_t inputLength = inputNum * typeLength;
270+ uint64_t inputBytes = inputLength / inputNum;
271+
272+ uint64_t ubDataNumber = ubDataNumMap[inputDatatype][outputDatatype];
273+ uint64_t tileBlockNum = (ubSize / BLOCK_SIZE) / ubDataNumber;
274+ uint64_t tileDataNum = (tileBlockNum * BLOCK_SIZE) / inputBytes;
275+
276+ uint64_t inputLengthAlgin32 = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE);
277+ 
278+ if (inputNum <= tileDataNum)
279+ {
280+ coreNum = 1;
281+ }
282+ else
283+ {
284+ coreNum = (coreNum < inputLengthAlgin32 / BLOCK_SIZE) ? coreNum : inputLengthAlgin32 / BLOCK_SIZE;
285+ }
286+ if (coreNum == 0 || BLOCK_SIZE == 0)
287+ {
288+ OP_LOGE(context, "coreNum or BLOCK_SIZE is 0");
289+ return ge::GRAPH_FAILED;
290+ }
291+ uint64_t everyCoreInputBlockNum = inputLengthAlgin32 / BLOCK_SIZE / coreNum;
292+ uint64_t tailBlockNum = (inputLengthAlgin32 / BLOCK_SIZE) % coreNum;
293+
294+ uint64_t smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
295+ uint64_t smallTileNum = everyCoreInputBlockNum / tileBlockNum;
296+ uint64_t finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
297+ uint64_t smallTailDataNum = smallCoreDataNum - (tileDataNum * smallTileNum);
298+ smallTailDataNum = smallTailDataNum == 0 ? tileDataNum : smallTailDataNum;
299+
300+ everyCoreInputBlockNum += 1;
301+ uint64_t bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
302+ uint64_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
303+ uint64_t finalBigTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
304+ uint64_t bigTailDataNum = bigCoreDataNum - tileDataNum * bigTileNum;
305+ bigTailDataNum = bigTailDataNum == 0 ? tileDataNum : bigTailDataNum;
306+ 
307+ tiling->smallCoreDataNum = smallCoreDataNum;
308+ tiling->bigCoreDataNum = bigCoreDataNum;
309+ tiling->tileDataNum = tileDataNum;
310+ tiling->smallTailDataNum = smallTailDataNum;
311+ tiling->bigTailDataNum = bigTailDataNum;
312+ tiling->finalSmallTileNum = finalSmallTileNum;
313+ tiling->finalBigTileNum = finalBigTileNum;
314+ tiling->tailBlockNum = tailBlockNum;
315+ 
316+ context->SetBlockDim(coreNum);
317+ context->SetTilingKey(tilingKey);
318+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
319+ currentWorkspace[0] = 0;
320+ return ge::GRAPH_SUCCESS;
321+}
322+ 
323+// tiling注册入口.
324+IMPL_OP_OPTILING(Cast).Tiling(CastTilingFunc);
325+} // namespace optiling
Aexperimental/math/cast/op_kernel/cast.cpp+189-0
@@ -0,0 +1,189 @@
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 cast.cpp
13+ * \brief
14+ */
15+ 
16+#include "cast.h"
17+ 
18+using namespace NsCast;
19+ 
20+template <uint32_t schMode>
21+__global__ __aicore__ void cast(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
22+{
23+ REGISTER_TILING_DEFAULT(CastTilingData);
24+ GET_TILING_DATA_WITH_STRUCT(CastTilingData, tiling_data, tiling);
25+ 
26+ TPipe pipe;
27+ if (TILING_KEY_IS(1))
28+ {
29+ if constexpr (std::is_same_v<DTYPE_X, bool>)
30+ {
31+ KernelCast0TBuf<int8_t, DTYPE_Y> op;
32+ op.Init(x, y, tiling_data.smallCoreDataNum,
33+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
34+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
35+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
36+ tiling_data.tailBlockNum, &pipe);
37+ op.Process();
38+ }
39+ else if constexpr (std::is_same_v<DTYPE_Y, bool>)
40+ {
41+ KernelCast0TBuf<DTYPE_X, int8_t> op;
42+ op.Init(x, y, tiling_data.smallCoreDataNum,
43+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
44+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
45+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
46+ tiling_data.tailBlockNum, &pipe);
47+ op.Process();
48+ }
49+ else
50+ {
51+ KernelCast0TBuf<DTYPE_X, DTYPE_Y> op;
52+ op.Init(x, y, tiling_data.smallCoreDataNum,
53+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
54+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
55+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
56+ tiling_data.tailBlockNum, &pipe);
57+ op.Process();
58+ }
59+ }
60+ else if (TILING_KEY_IS(2))
61+ {
62+ KernelCast1TBuf4B<DTYPE_X, DTYPE_Y> op;
63+ op.Init(x, y, tiling_data.smallCoreDataNum,
64+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
65+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
66+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
67+ tiling_data.tailBlockNum, &pipe);
68+ op.Process();
69+ }
70+ else if (TILING_KEY_IS(3))
71+ {
72+ if constexpr (std::is_same_v<DTYPE_Y, bool>)
73+ {
74+ KernelCast2TBuf2B<DTYPE_X, int8_t> op;
75+ op.Init(x, y, tiling_data.smallCoreDataNum,
76+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
77+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
78+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
79+ tiling_data.tailBlockNum, &pipe);
80+ op.Process();
81+ }
82+ else
83+ {
84+ KernelCast2TBuf2B<DTYPE_X, DTYPE_Y> op;
85+ op.Init(x, y, tiling_data.smallCoreDataNum,
86+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
87+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
88+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
89+ tiling_data.tailBlockNum, &pipe);
90+ op.Process();
91+ }
92+ }
93+ else if (TILING_KEY_IS(4))
94+ {
95+ KernelCast3TBuf2B<DTYPE_X, DTYPE_Y> op;
96+ op.Init(x, y, tiling_data.smallCoreDataNum,
97+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
98+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
99+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
100+ tiling_data.tailBlockNum, &pipe);
101+ op.Process();
102+ }
103+ else if (TILING_KEY_IS(5))
104+ {
105+ if constexpr (std::is_same_v<DTYPE_X, bool>)
106+ {
107+ KernelCast1TBuf2B<int8_t, DTYPE_Y> op;
108+ op.Init(x, y, tiling_data.smallCoreDataNum,
109+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
110+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
111+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
112+ tiling_data.tailBlockNum, &pipe);
113+ op.Process();
114+ }
115+ else if constexpr (std::is_same_v<DTYPE_Y, bool>)
116+ {
117+ KernelCast1TBuf2B<DTYPE_X, int8_t> op;
118+ op.Init(x, y, tiling_data.smallCoreDataNum,
119+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
120+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
121+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
122+ tiling_data.tailBlockNum, &pipe);
123+ op.Process();
124+ }
125+ else
126+ {
127+ KernelCast1TBuf2B<DTYPE_X, DTYPE_Y> op;
128+ op.Init(x, y, tiling_data.smallCoreDataNum,
129+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
130+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
131+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
132+ tiling_data.tailBlockNum, &pipe);
133+ op.Process();
134+ }
135+ }
136+ else if (TILING_KEY_IS(6))
137+ {
138+ if constexpr (std::is_same_v<DTYPE_X, bool>)
139+ {
140+ KernelCast1TBuf2B1TBuf4B<int8_t, DTYPE_Y> op;
141+ op.Init(x, y, tiling_data.smallCoreDataNum,
142+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
143+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
144+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
145+ tiling_data.tailBlockNum, &pipe);
146+ op.Process();
147+ }
148+ else if constexpr (std::is_same_v<DTYPE_Y, bool>)
149+ {
150+ KernelCast1TBuf2B1TBuf4B<DTYPE_X, int8_t> op;
151+ op.Init(x, y, tiling_data.smallCoreDataNum,
152+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
153+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
154+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
155+ tiling_data.tailBlockNum, &pipe);
156+ op.Process();
157+ }
158+ else
159+ {
160+ KernelCast1TBuf2B1TBuf4B<DTYPE_X, DTYPE_Y> op;
161+ op.Init(x, y, tiling_data.smallCoreDataNum,
162+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
163+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
164+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
165+ tiling_data.tailBlockNum, &pipe);
166+ op.Process();
167+ }
168+ }
169+ else if (TILING_KEY_IS(7))
170+ {
171+ KernelCast3TBuf2B1TBuf4B<DTYPE_X, DTYPE_Y> op;
172+ op.Init(x, y, tiling_data.smallCoreDataNum,
173+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
174+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
175+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
176+ tiling_data.tailBlockNum, &pipe);
177+ op.Process();
178+ }
179+ else if (TILING_KEY_IS(8))
180+ {
181+ KernelCastTQueBind op;
182+ op.Init(x, y, tiling_data.smallCoreDataNum,
183+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
184+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
185+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
186+ tiling_data.tailBlockNum, &pipe);
187+ op.Process();
188+ }
189+}
Aexperimental/math/cast/op_kernel/cast.h+956-0
@@ -0,0 +1,956 @@
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 cast.h
13+ * \brief
14+ */
15+#ifndef __CAST_H__
16+#define __CAST_H__
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "cast_tiling_data.h"
21+#include "cast_tiling_key.h"
22+ 
23+namespace NsCast {
24+ 
25+using namespace AscendC;
26+ 
27+constexpr int32_t BUFFER_NUM = 2;
28+constexpr int16_t CONST_128 = 128;
29+constexpr int16_t CONST_NE128 = -128;
30+constexpr int16_t CONST_255 = 255;
31+constexpr int16_t CONST_1 = 1;
32+constexpr half HALF_ONE = 1.0;
33+ 
34+template <typename TYPE_X, typename TYPE_Y>
35+class BaseKernelCast
36+{
37+public:
38+ __aicore__ inline BaseKernelCast() {}
39+ 
40+protected:
41+ __aicore__ inline void CopyIn(int32_t progress)
42+ {
43+ LocalTensor<TYPE_X> xLocal = inQueueX.AllocTensor<TYPE_X>();
44+ DataCopy(xLocal, xGm[progress * this->tileDataNum], this->processDataNum);
45+ inQueueX.EnQue(xLocal);
46+ }
47+ __aicore__ inline void CopyOut(int32_t progress)
48+ {
49+ LocalTensor<TYPE_Y> yLocal = outQueueY.DeQue<TYPE_Y>();
50+ DataCopy(yGm[progress * this->tileDataNum], yLocal, this->processDataNum);
51+ outQueueY.FreeTensor(yLocal);
52+ }
53+ 
54+protected:
55+ TPipe *pipe;
56+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
57+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
58+ GlobalTensor<TYPE_X> xGm;
59+ GlobalTensor<TYPE_Y> yGm;
60+ uint32_t coreDataNum = 0;
61+ uint32_t tileNum = 0;
62+ uint32_t tileDataNum = 0;
63+ uint32_t tailDataNum = 0;
64+ uint32_t processDataNum = 0;
65+};
66+ 
67+template <typename TYPE_X, typename TYPE_Y>
68+class KernelCast0TBuf : public BaseKernelCast<TYPE_X, TYPE_Y>
69+{
70+ /*
71+ 无临时变量
72+ half -> float
73+ half -> int32 (TRUNC)
74+ half -> bool (Abs)
75+ half -> int16 (TRUNC)
76+ float -> half
77+ float -> bfloat16 (RINT)
78+ float -> int32 (TRUNC)
79+ float -> int64 (TRUNC)
80+ float -> int16 (TRUNC)
81+ int32 -> float
82+ int32 -> int64
83+ int32 -> int16
84+ int8 -> half
85+ uint8 -> half
86+ bool -> half
87+ int64 -> float (ROUND)
88+ int64 -> int32
89+ bfloat16 -> float
90+ bfloat16 -> int32 (TRUNC)
91+ int16 -> float
92+ int16 -> half
93+ */
94+public:
95+ __aicore__ inline KernelCast0TBuf() {}
96+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum,
97+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
98+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
99+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
100+ uint32_t tailBlockNum, TPipe *pipeIn)
101+ {
102+ this->pipe = pipeIn;
103+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
104+ uint32_t coreNum = GetBlockIdx();
105+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
106+ this->tileDataNum = tileDataNum;
107+ if (coreNum < tailBlockNum)
108+ {
109+ this->coreDataNum = bigCoreDataNum;
110+ this->tileNum = finalBigTileNum;
111+ this->tailDataNum = bigTailDataNum;
112+ }
113+ else
114+ {
115+ this->coreDataNum = smallCoreDataNum;
116+ this->tileNum = finalSmallTileNum;
117+ this->tailDataNum = smallTailDataNum;
118+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
119+ }
120+ this->xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
121+ this->yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
122+ BufferInit();
123+ }
124+ __aicore__ inline void Process()
125+ {
126+ int32_t loopCount = this->tileNum;
127+ this->processDataNum = this->tileDataNum;
128+ for (int32_t i = 0; i < loopCount - 1; i++)
129+ {
130+ this->CopyIn(i);
131+ Compute(i);
132+ this->CopyOut(i);
133+ }
134+ this->processDataNum = this->tailDataNum;
135+ this->CopyIn(loopCount - 1);
136+ Compute(loopCount - 1);
137+ this->CopyOut(loopCount - 1);
138+ }
139+ 
140+private:
141+ __aicore__ inline void BufferInit()
142+ {
143+ this->pipe->InitBuffer(this->inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
144+ this->pipe->InitBuffer(this->outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
145+ }
146+ __aicore__ inline void Compute(int32_t progress)
147+ {
148+ LocalTensor<TYPE_X> xLocal = this->inQueueX.template DeQue<TYPE_X>();
149+ LocalTensor<TYPE_Y> yLocal = this->outQueueY.template AllocTensor<TYPE_Y>();
150+ if constexpr ((std::is_same_v<DTYPE_X, half> && std::is_same_v<DTYPE_Y, int32_t>) ||
151+ (std::is_same_v<DTYPE_X, half> && std::is_same_v<DTYPE_Y, int16_t>) ||
152+ (std::is_same_v<DTYPE_X, float> && std::is_same_v<DTYPE_Y, int32_t>) ||
153+ (std::is_same_v<DTYPE_X, float> && std::is_same_v<DTYPE_Y, int16_t>) ||
154+ (std::is_same_v<DTYPE_X, float> && std::is_same_v<DTYPE_Y, int64_t>) ||
155+ (std::is_same_v<DTYPE_X, bfloat16_t> && std::is_same_v<DTYPE_Y, int32_t>))
156+ {
157+ Cast(yLocal, xLocal, RoundMode::CAST_TRUNC, this->processDataNum);
158+ }
159+ else if constexpr ((std::is_same_v<DTYPE_X, int64_t> && std::is_same_v<DTYPE_Y, float>))
160+ {
161+ Cast(yLocal, xLocal, RoundMode::CAST_ROUND, this->processDataNum);
162+ }
163+ else if constexpr ((std::is_same_v<DTYPE_X, half> && std::is_same_v<DTYPE_Y, bool>))
164+ {
165+ Abs(xLocal, xLocal, this->processDataNum);
166+ Mins(xLocal, xLocal, HALF_ONE, this->processDataNum);
167+ Cast(yLocal, xLocal, RoundMode::CAST_CEIL, this->processDataNum);
168+ }
169+ else if constexpr ((std::is_same_v<DTYPE_X, half> && std::is_same_v<DTYPE_Y, float>) ||
170+ (std::is_same_v<DTYPE_X, float> && std::is_same_v<DTYPE_Y, half>) ||
171+ (std::is_same_v<DTYPE_X, int32_t> && std::is_same_v<DTYPE_Y, int16_t>) ||
172+ (std::is_same_v<DTYPE_X, int32_t> && std::is_same_v<DTYPE_Y, int64_t>) ||
173+ (std::is_same_v<DTYPE_X, int32_t> && std::is_same_v<DTYPE_Y, float>) ||
174+ (std::is_same_v<DTYPE_X, int8_t> && std::is_same_v<DTYPE_Y, half>) ||
175+ (std::is_same_v<DTYPE_X, uint8_t> && std::is_same_v<DTYPE_Y, half>) ||
176+ (std::is_same_v<DTYPE_X, bool> && std::is_same_v<DTYPE_Y, half>) ||
177+ (std::is_same_v<DTYPE_X, int64_t> && std::is_same_v<DTYPE_Y, int32_t>) ||
178+ (std::is_same_v<DTYPE_X, bfloat16_t> && std::is_same_v<DTYPE_Y, float>) ||
179+ (std::is_same_v<DTYPE_X, int16_t> && std::is_same_v<DTYPE_Y, half>) ||
180+ (std::is_same_v<DTYPE_X, int16_t> && std::is_same_v<DTYPE_Y, float>))
181+ {
182+ Cast(yLocal, xLocal, RoundMode::CAST_NONE, this->processDataNum);
183+ }
184+ else if constexpr ((std::is_same_v<DTYPE_X, float> && std::is_same_v<DTYPE_Y, bfloat16_t>))
185+ {
186+ Cast(yLocal, xLocal, RoundMode::CAST_RINT, this->processDataNum);
187+ }
188+ this->outQueueY.template EnQue<TYPE_Y>(yLocal);
189+ this->inQueueX.template FreeTensor(xLocal);
190+ }
191+};
192+ 
193+template <typename TYPE_X, typename TYPE_Y>
194+class KernelCast1TBuf4B : public BaseKernelCast<TYPE_X, TYPE_Y>
195+{
196+ /*
197+ 1个4Bytes的临时变量
198+ half -> float ->(RINT) bfloat16
199+ int32 -> float ->(RINT) bfloat16
200+ int32 -> float -> half
201+ int64 ->(ROUND) float -> half
202+ int64 ->(ROUND) float ->(RINT) bfloat16
203+ int64 -> int32 -> int16
204+ bfloat16 -> float -> half
205+ int16 -> float ->(ROUND) int32
206+ int16 -> float ->(ROUND) int64
207+ */
208+public:
209+ __aicore__ inline KernelCast1TBuf4B() {}
210+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum,
211+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
212+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
213+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
214+ uint32_t tailBlockNum, TPipe *pipeIn)
215+ {
216+ this->pipe = pipeIn;
217+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
218+ uint32_t coreNum = GetBlockIdx();
219+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
220+ this->tileDataNum = tileDataNum;
221+ if (coreNum < tailBlockNum)
222+ {
223+ this->coreDataNum = bigCoreDataNum;
224+ this->tileNum = finalBigTileNum;
225+ this->tailDataNum = bigTailDataNum;
226+ }
227+ else
228+ {
229+ this->coreDataNum = smallCoreDataNum;
230+ this->tileNum = finalSmallTileNum;
231+ this->tailDataNum = smallTailDataNum;
232+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
233+ }
234+ this->xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
235+ this->yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
236+ BufferInit();
237+ }
238+ __aicore__ inline void Process()
239+ {
240+ int32_t loopCount = this->tileNum;
241+ this->processDataNum = this->tileDataNum;
242+ for (int32_t i = 0; i < loopCount - 1; i++)
243+ {
244+ this->CopyIn(i);
245+ Compute(i);
246+ this->CopyOut(i);
247+ }
248+ this->processDataNum = this->tailDataNum;
249+ this->CopyIn(loopCount - 1);
250+ Compute(loopCount - 1);
251+ this->CopyOut(loopCount - 1);
252+ }
253+ 
254+private:
255+ __aicore__ inline void BufferInit()
256+ {
257+ this->pipe->InitBuffer(this->inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
258+ this->pipe->InitBuffer(this->outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
259+ this->pipe->InitBuffer(tmp4Bytes1, this->tileDataNum * sizeof(float));
260+ }
261+ __aicore__ inline void Compute(int32_t progress)
262+ {
263+ LocalTensor<TYPE_X> xLocal = this->inQueueX.template DeQue<TYPE_X>();
264+ LocalTensor<TYPE_Y> yLocal = this->outQueueY.template AllocTensor<TYPE_Y>();
265+ if constexpr ((std::is_same_v<DTYPE_X, half> && std::is_same_v<DTYPE_Y, bfloat16_t>) ||
266+ (std::is_same_v<DTYPE_X, int32_t> && std::is_same_v<DTYPE_Y, bfloat16_t>) ||
267+ (std::is_same_v<DTYPE_X, int32_t> && std::is_same_v<DTYPE_Y, half>) ||
268+ (std::is_same_v<DTYPE_X, bfloat16_t> && std::is_same_v<DTYPE_Y, half>))
269+ {
270+ LocalTensor<float> tmp1 = tmp4Bytes1.Get<float>();
271+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
272+ if constexpr (std::is_same_v<DTYPE_Y, half>)
273+ {
274+ Cast(yLocal, tmp1, RoundMode::CAST_NONE, this->processDataNum);
275+ }
276+ else if constexpr (std::is_same_v<DTYPE_Y, bfloat16_t>)
277+ {
278+ Cast(yLocal, tmp1, RoundMode::CAST_RINT, this->processDataNum);
279+ }
280+ }
281+ else if constexpr ((std::is_same_v<DTYPE_X, int64_t> && std::is_same_v<DTYPE_Y, half>) ||
282+ (std::is_same_v<DTYPE_X, int64_t> && std::is_same_v<DTYPE_Y, bfloat16_t>))
283+ {
284+ LocalTensor<float> tmp1 = tmp4Bytes1.Get<float>();
285+ Cast(tmp1, xLocal, RoundMode::CAST_ROUND, this->processDataNum);
286+ if constexpr (std::is_same_v<DTYPE_Y, half>)
287+ {
288+ Cast(yLocal, tmp1, RoundMode::CAST_NONE, this->processDataNum);
289+ }
290+ else if constexpr (std::is_same_v<DTYPE_Y, bfloat16_t>)
291+ {
292+ Cast(yLocal, tmp1, RoundMode::CAST_RINT, this->processDataNum);
293+ }
294+ }
295+ else if constexpr ((std::is_same_v<DTYPE_X, int64_t> && std::is_same_v<DTYPE_Y, int16_t>))
296+ {
297+ LocalTensor<int32_t> tmp1 = tmp4Bytes1.Get<int32_t>();
298+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
299+ Cast(yLocal, tmp1, RoundMode::CAST_NONE, this->processDataNum);
300+ }
301+ else if constexpr ((std::is_same_v<DTYPE_X, int16_t> && std::is_same_v<DTYPE_Y, int32_t>) ||
302+ (std::is_same_v<DTYPE_X, int16_t> && std::is_same_v<DTYPE_Y, int64_t>))
303+ {
304+ LocalTensor<float> tmp1 = tmp4Bytes1.Get<float>();
305+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
306+ Cast(yLocal, tmp1, RoundMode::CAST_ROUND, this->processDataNum);
307+ }
308+ this->outQueueY.template EnQue<TYPE_Y>(yLocal);
309+ this->inQueueX.template FreeTensor(xLocal);
310+ }
311+ 
312+private:
313+ TBuf<QuePosition::VECCALC> tmp4Bytes1;
314+};
315+ 
316+template <typename TYPE_X, typename TYPE_Y>
317+class KernelCast2TBuf2B : public BaseKernelCast<TYPE_X, TYPE_Y>
318+{
319+ /*
320+ 2个2Bytes的临时变量
321+ half -> int8
322+ half -> uint8
323+ float -> bool
324+ int32 -> bool
325+ int16 -> int8
326+ int16 -> uint8
327+ */
328+public:
329+ __aicore__ inline KernelCast2TBuf2B() {}
330+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum,
331+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
332+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
333+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
334+ uint32_t tailBlockNum, TPipe *pipeIn)
335+ {
336+ this->pipe = pipeIn;
337+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
338+ uint32_t coreNum = GetBlockIdx();
339+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
340+ this->tileDataNum = tileDataNum;
341+ if (coreNum < tailBlockNum)
342+ {
343+ this->coreDataNum = bigCoreDataNum;
344+ this->tileNum = finalBigTileNum;
345+ this->tailDataNum = bigTailDataNum;
346+ }
347+ else
348+ {
349+ this->coreDataNum = smallCoreDataNum;
350+ this->tileNum = finalSmallTileNum;
351+ this->tailDataNum = smallTailDataNum;
352+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
353+ }
354+ this->xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
355+ this->yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
356+ BufferInit();
357+ }
358+ __aicore__ inline void Process()
359+ {
360+ int32_t loopCount = this->tileNum;
361+ this->processDataNum = this->tileDataNum;
362+ for (int32_t i = 0; i < loopCount - 1; i++)
363+ {
364+ this->CopyIn(i);
365+ Compute(i);
366+ this->CopyOut(i);
367+ }
368+ this->processDataNum = this->tailDataNum;
369+ this->CopyIn(loopCount - 1);
370+ Compute(loopCount - 1);
371+ this->CopyOut(loopCount - 1);
372+ }
373+ 
374+private:
375+ __aicore__ inline void BufferInit()
376+ {
377+ this->pipe->InitBuffer(this->inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
378+ this->pipe->InitBuffer(this->outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
379+ this->pipe->InitBuffer(tmp2Bytes1, this->tileDataNum * sizeof(half));
380+ this->pipe->InitBuffer(tmp2Bytes2, this->tileDataNum * sizeof(half));
381+ }
382+ __aicore__ inline void Compute(int32_t progress)
383+ {
384+ LocalTensor<TYPE_X> xLocal = this->inQueueX.template DeQue<TYPE_X>();
385+ LocalTensor<TYPE_Y> yLocal = this->outQueueY.template AllocTensor<TYPE_Y>();
386+ if constexpr ((std::is_same_v<DTYPE_X, half> && std::is_same_v<DTYPE_Y, int8_t>) ||
387+ (std::is_same_v<DTYPE_X, half> && std::is_same_v<DTYPE_Y, uint8_t>))
388+ {
389+ LocalTensor<int16_t> tmp1 = tmp2Bytes1.Get<int16_t>();
390+ LocalTensor<int16_t> tmp2 = tmp2Bytes2.Get<int16_t>();
391+ Cast(tmp1, xLocal, RoundMode::CAST_TRUNC, this->processDataNum);
392+ Duplicate(tmp2, CONST_255, this->processDataNum);
393+ And(tmp1, tmp1, tmp2, this->processDataNum);
394+ if constexpr (std::is_same_v<DTYPE_Y, int8_t>)
395+ {
396+ Adds(tmp1, tmp1, CONST_128, this->processDataNum);
397+ And(tmp1, tmp1, tmp2, this->processDataNum);
398+ Adds(tmp1, tmp1, CONST_NE128, this->processDataNum);
399+ }
400+ Cast(xLocal, tmp1, RoundMode::CAST_NONE, this->processDataNum);
401+ Cast(yLocal, xLocal, RoundMode::CAST_NONE, this->processDataNum);
402+ }
403+ else if constexpr ((std::is_same_v<DTYPE_X, float> && std::is_same_v<DTYPE_Y, bool>))
404+ {
405+ LocalTensor<int16_t> tmp1 = tmp2Bytes1.Get<int16_t>();
406+ LocalTensor<half> tmp2 = tmp2Bytes2.Get<half>();
407+ Abs(xLocal, xLocal, this->processDataNum);
408+ Cast(tmp1, xLocal, RoundMode::CAST_CEIL, this->processDataNum);
409+ Mins(tmp1, tmp1, CONST_1, this->processDataNum);
410+ Cast(tmp2, tmp1, RoundMode::CAST_NONE, this->processDataNum);
411+ Cast(yLocal, tmp2, RoundMode::CAST_NONE, this->processDataNum);
412+ }
413+ else if constexpr ((std::is_same_v<DTYPE_X, int32_t> && std::is_same_v<DTYPE_Y, bool>))
414+ {
415+ LocalTensor<int16_t> tmp1 = tmp2Bytes1.Get<int16_t>();
416+ LocalTensor<half> tmp2 = tmp2Bytes2.Get<half>();
417+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
418+ Cast(tmp2, tmp1, RoundMode::CAST_NONE, this->processDataNum);
419+ Abs(tmp2, tmp2, this->processDataNum);
420+ Mins(tmp2, tmp2, HALF_ONE, this->processDataNum);
421+ Cast(yLocal, tmp2, RoundMode::CAST_NONE, this->processDataNum);
422+ }
423+ else if constexpr ((std::is_same_v<DTYPE_X, int16_t> && std::is_same_v<DTYPE_Y, int8_t>) ||
424+ (std::is_same_v<DTYPE_X, int16_t> && std::is_same_v<DTYPE_Y, uint8_t>))
425+ {
426+ LocalTensor<int16_t> tmp1 = tmp2Bytes1.Get<int16_t>();
427+ LocalTensor<half> tmp2 = tmp2Bytes2.Get<half>();
428+ Duplicate(tmp1, CONST_255, this->processDataNum);
429+ And(xLocal, xLocal, tmp1, this->processDataNum);
430+ if constexpr (std::is_same_v<DTYPE_Y, int8_t>)
431+ {
432+ Adds(xLocal, xLocal, CONST_128, this->processDataNum);
433+ And(xLocal, xLocal, tmp1, this->processDataNum);
434+ Adds(xLocal, xLocal, CONST_NE128, this->processDataNum);
435+ }
436+ Cast(tmp2, xLocal, RoundMode::CAST_NONE, this->processDataNum);
437+ Cast(yLocal, tmp2, RoundMode::CAST_NONE, this->processDataNum);
438+ }
439+ this->outQueueY.template EnQue<TYPE_Y>(yLocal);
440+ this->inQueueX.template FreeTensor(xLocal);
441+ }
442+ 
443+private:
444+ TBuf<QuePosition::VECCALC> tmp2Bytes1;
445+ TBuf<QuePosition::VECCALC> tmp2Bytes2;
446+};
447+ 
448+template <typename TYPE_X, typename TYPE_Y>
449+class KernelCast3TBuf2B : public BaseKernelCast<TYPE_X, TYPE_Y>
450+{
451+ /*
452+ 3个2Bytes的临时变量
453+ float -> int8
454+ float -> uint8
455+ int32 -> int8
456+ int32 -> uint8
457+ */
458+public:
459+ __aicore__ inline KernelCast3TBuf2B() {}
460+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum,
461+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
462+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
463+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
464+ uint32_t tailBlockNum, TPipe *pipeIn)
465+ {
466+ this->pipe = pipeIn;
467+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
468+ uint32_t coreNum = GetBlockIdx();
469+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
470+ this->tileDataNum = tileDataNum;
471+ if (coreNum < tailBlockNum)
472+ {
473+ this->coreDataNum = bigCoreDataNum;
474+ this->tileNum = finalBigTileNum;
475+ this->tailDataNum = bigTailDataNum;
476+ }
477+ else
478+ {
479+ this->coreDataNum = smallCoreDataNum;
480+ this->tileNum = finalSmallTileNum;
481+ this->tailDataNum = smallTailDataNum;
482+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
483+ }
484+ this->xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
485+ this->yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
486+ BufferInit();
487+ }
488+ __aicore__ inline void Process()
489+ {
490+ int32_t loopCount = this->tileNum;
491+ this->processDataNum = this->tileDataNum;
492+ for (int32_t i = 0; i < loopCount - 1; i++)
493+ {
494+ this->CopyIn(i);
495+ Compute(i);
496+ this->CopyOut(i);
497+ }
498+ this->processDataNum = this->tailDataNum;
499+ this->CopyIn(loopCount - 1);
500+ Compute(loopCount - 1);
501+ this->CopyOut(loopCount - 1);
502+ }
503+ 
504+private:
505+ __aicore__ inline void BufferInit()
506+ {
507+ this->pipe->InitBuffer(this->inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
508+ this->pipe->InitBuffer(this->outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
509+ this->pipe->InitBuffer(tmp2Bytes1, this->tileDataNum * sizeof(half));
510+ this->pipe->InitBuffer(tmp2Bytes2, this->tileDataNum * sizeof(half));
511+ this->pipe->InitBuffer(tmp2Bytes3, this->tileDataNum * sizeof(half));
512+ }
513+ __aicore__ inline void Compute(int32_t progress)
514+ {
515+ LocalTensor<TYPE_X> xLocal = this->inQueueX.template DeQue<TYPE_X>();
516+ LocalTensor<TYPE_Y> yLocal = this->outQueueY.template AllocTensor<TYPE_Y>();
517+ if constexpr ((std::is_same_v<DTYPE_X, float> && std::is_same_v<DTYPE_Y, int8_t>) ||
518+ (std::is_same_v<DTYPE_X, float> && std::is_same_v<DTYPE_Y, uint8_t>) ||
519+ (std::is_same_v<DTYPE_X, int32_t> && std::is_same_v<DTYPE_Y, int8_t>) ||
520+ (std::is_same_v<DTYPE_X, int32_t> && std::is_same_v<DTYPE_Y, uint8_t>))
521+ {
522+ LocalTensor<int16_t> tmp1 = tmp2Bytes1.Get<int16_t>();
523+ LocalTensor<int16_t> tmp2 = tmp2Bytes2.Get<int16_t>();
524+ LocalTensor<half> tmp3 = tmp2Bytes3.Get<half>();
525+ if constexpr (std::is_same_v<DTYPE_X, float>)
526+ {
527+ Cast(tmp1, xLocal, RoundMode::CAST_TRUNC, this->processDataNum);
528+ }
529+ else if constexpr (std::is_same_v<DTYPE_X, int32_t>)
530+ {
531+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
532+ }
533+ Duplicate(tmp2, CONST_255, this->processDataNum);
534+ And(tmp1, tmp1, tmp2, this->processDataNum);
535+ if constexpr (std::is_same_v<DTYPE_Y, int8_t>)
536+ {
537+ Adds(tmp1, tmp1, CONST_128, this->processDataNum);
538+ And(tmp1, tmp1, tmp2, this->processDataNum);
539+ Adds(tmp1, tmp1, CONST_NE128, this->processDataNum);
540+ }
541+ Cast(tmp3, tmp1, RoundMode::CAST_NONE, this->processDataNum);
542+ Cast(yLocal, tmp3, RoundMode::CAST_NONE, this->processDataNum);
543+ }
544+ this->outQueueY.template EnQue<TYPE_Y>(yLocal);
545+ this->inQueueX.template FreeTensor(xLocal);
546+ }
547+ 
548+private:
549+ TBuf<QuePosition::VECCALC> tmp2Bytes1;
550+ TBuf<QuePosition::VECCALC> tmp2Bytes2;
551+ TBuf<QuePosition::VECCALC> tmp2Bytes3;
552+};
553+ 
554+template <typename TYPE_X, typename TYPE_Y>
555+class KernelCast1TBuf2B : public BaseKernelCast<TYPE_X, TYPE_Y>
556+{
557+ /*
558+ 1个2Bytes的临时变量
559+ int8/uint8/bool -> float
560+ int8/uint8/bool -> int32
561+ int8/uint8 -> int16
562+ int8 -> bool
563+ */
564+public:
565+ __aicore__ inline KernelCast1TBuf2B() {}
566+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum,
567+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
568+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
569+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
570+ uint32_t tailBlockNum, TPipe *pipeIn)
571+ {
572+ this->pipe = pipeIn;
573+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
574+ uint32_t coreNum = GetBlockIdx();
575+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
576+ this->tileDataNum = tileDataNum;
577+ if (coreNum < tailBlockNum)
578+ {
579+ this->coreDataNum = bigCoreDataNum;
580+ this->tileNum = finalBigTileNum;
581+ this->tailDataNum = bigTailDataNum;
582+ }
583+ else
584+ {
585+ this->coreDataNum = smallCoreDataNum;
586+ this->tileNum = finalSmallTileNum;
587+ this->tailDataNum = smallTailDataNum;
588+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
589+ }
590+ this->xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
591+ this->yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
592+ BufferInit();
593+ }
594+ __aicore__ inline void Process()
595+ {
596+ int32_t loopCount = this->tileNum;
597+ this->processDataNum = this->tileDataNum;
598+ for (int32_t i = 0; i < loopCount - 1; i++)
599+ {
600+ this->CopyIn(i);
601+ Compute(i);
602+ this->CopyOut(i);
603+ }
604+ this->processDataNum = this->tailDataNum;
605+ this->CopyIn(loopCount - 1);
606+ Compute(loopCount - 1);
607+ this->CopyOut(loopCount - 1);
608+ }
609+ 
610+private:
611+ __aicore__ inline void BufferInit()
612+ {
613+ this->pipe->InitBuffer(this->inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
614+ this->pipe->InitBuffer(this->outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
615+ this->pipe->InitBuffer(tmp2Bytes1, this->tileDataNum * sizeof(half));
616+ }
617+ __aicore__ inline void Compute(int32_t progress)
618+ {
619+ LocalTensor<TYPE_X> xLocal = this->inQueueX.template DeQue<TYPE_X>();
620+ LocalTensor<TYPE_Y> yLocal = this->outQueueY.template AllocTensor<TYPE_Y>();
621+ if constexpr ((std::is_same_v<DTYPE_X, int8_t> && std::is_same_v<DTYPE_Y, float>) ||
622+ (std::is_same_v<DTYPE_X, uint8_t> && std::is_same_v<DTYPE_Y, float>) ||
623+ (std::is_same_v<DTYPE_X, bool> && std::is_same_v<DTYPE_Y, float>) ||
624+ (std::is_same_v<DTYPE_X, int8_t> && std::is_same_v<DTYPE_Y, int32_t>) ||
625+ (std::is_same_v<DTYPE_X, uint8_t> && std::is_same_v<DTYPE_Y, int32_t>) ||
626+ (std::is_same_v<DTYPE_X, bool> && std::is_same_v<DTYPE_Y, int32_t>) ||
627+ (std::is_same_v<DTYPE_X, int8_t> && std::is_same_v<DTYPE_Y, int16_t>) ||
628+ (std::is_same_v<DTYPE_X, uint8_t> && std::is_same_v<DTYPE_Y, int16_t>) ||
629+ (std::is_same_v<DTYPE_X, bool> && std::is_same_v<DTYPE_Y, uint8_t>))
630+ {
631+ LocalTensor<half> tmp1 = tmp2Bytes1.Get<half>();
632+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
633+ if constexpr (std::is_same_v<DTYPE_Y, int32_t> || std::is_same_v<DTYPE_Y, int16_t> || std::is_same_v<DTYPE_Y, uint8_t>)
634+ {
635+ Cast(yLocal, tmp1, RoundMode::CAST_TRUNC, this->processDataNum);
636+ }
637+ else
638+ {
639+ Cast(yLocal, tmp1, RoundMode::CAST_NONE, this->processDataNum);
640+ }
641+ }
642+ else if constexpr ((std::is_same_v<DTYPE_X, int8_t> && std::is_same_v<DTYPE_Y, bool>))
643+ {
644+ LocalTensor<half> tmp1 = tmp2Bytes1.Get<half>();
645+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
646+ Abs(tmp1, tmp1, this->processDataNum);
647+ Mins(tmp1, tmp1, HALF_ONE, this->processDataNum);
648+ Cast(yLocal, tmp1, RoundMode::CAST_NONE, this->processDataNum);
649+ }
650+ this->outQueueY.template EnQue<TYPE_Y>(yLocal);
651+ this->inQueueX.template FreeTensor(xLocal);
652+ }
653+ 
654+private:
655+ TBuf<QuePosition::VECCALC> tmp2Bytes1;
656+};
657+ 
658+template <typename TYPE_X, typename TYPE_Y>
659+class KernelCast1TBuf2B1TBuf4B : public BaseKernelCast<TYPE_X, TYPE_Y>
660+{
661+ /*
662+ 1个2Bytes,1个4Bytes的临时变量
663+ int8/uint8/bool -> half ->(TRUNC) int32 -> int64
664+ int8/uint8/bool -> half -> float -> bfloat16
665+ int64 -> bool
666+ bfloat16 -> bool
667+ */
668+public:
669+ __aicore__ inline KernelCast1TBuf2B1TBuf4B() {}
670+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum,
671+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
672+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
673+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
674+ uint32_t tailBlockNum, TPipe *pipeIn)
675+ {
676+ this->pipe = pipeIn;
677+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
678+ uint32_t coreNum = GetBlockIdx();
679+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
680+ this->tileDataNum = tileDataNum;
681+ if (coreNum < tailBlockNum)
682+ {
683+ this->coreDataNum = bigCoreDataNum;
684+ this->tileNum = finalBigTileNum;
685+ this->tailDataNum = bigTailDataNum;
686+ }
687+ else
688+ {
689+ this->coreDataNum = smallCoreDataNum;
690+ this->tileNum = finalSmallTileNum;
691+ this->tailDataNum = smallTailDataNum;
692+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
693+ }
694+ this->xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
695+ this->yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
696+ BufferInit();
697+ }
698+ __aicore__ inline void Process()
699+ {
700+ int32_t loopCount = this->tileNum;
701+ this->processDataNum = this->tileDataNum;
702+ for (int32_t i = 0; i < loopCount - 1; i++)
703+ {
704+ this->CopyIn(i);
705+ Compute(i);
706+ this->CopyOut(i);
707+ }
708+ this->processDataNum = this->tailDataNum;
709+ this->CopyIn(loopCount - 1);
710+ Compute(loopCount - 1);
711+ this->CopyOut(loopCount - 1);
712+ }
713+ 
714+private:
715+ __aicore__ inline void BufferInit()
716+ {
717+ this->pipe->InitBuffer(this->inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
718+ this->pipe->InitBuffer(this->outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
719+ this->pipe->InitBuffer(tmp2Bytes1, this->tileDataNum * sizeof(half));
720+ this->pipe->InitBuffer(tmp4Bytes1, this->tileDataNum * sizeof(float));
721+ }
722+ __aicore__ inline void Compute(int32_t progress)
723+ {
724+ LocalTensor<TYPE_X> xLocal = this->inQueueX.template DeQue<TYPE_X>();
725+ LocalTensor<TYPE_Y> yLocal = this->outQueueY.template AllocTensor<TYPE_Y>();
726+ if constexpr ((std::is_same_v<DTYPE_X, int8_t> && std::is_same_v<DTYPE_Y, int64_t>) ||
727+ (std::is_same_v<DTYPE_X, uint8_t> && std::is_same_v<DTYPE_Y, int64_t>) ||
728+ (std::is_same_v<DTYPE_X, bool> && std::is_same_v<DTYPE_Y, int64_t>))
729+ {
730+ LocalTensor<half> tmp1 = tmp2Bytes1.Get<half>();
731+ LocalTensor<int32_t> tmp2 = tmp4Bytes1.Get<int32_t>();
732+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
733+ Cast(tmp2, tmp1, RoundMode::CAST_TRUNC, this->processDataNum);
734+ Cast(yLocal, tmp2, RoundMode::CAST_NONE, this->processDataNum);
735+ }
736+ else if constexpr ((std::is_same_v<DTYPE_X, int8_t> && std::is_same_v<DTYPE_Y, bfloat16_t>) ||
737+ (std::is_same_v<DTYPE_X, uint8_t> && std::is_same_v<DTYPE_Y, bfloat16_t>) ||
738+ (std::is_same_v<DTYPE_X, bool> && std::is_same_v<DTYPE_Y, bfloat16_t>))
739+ {
740+ LocalTensor<half> tmp1 = tmp2Bytes1.Get<half>();
741+ LocalTensor<float> tmp2 = tmp4Bytes1.Get<float>();
742+ Cast(tmp1, xLocal, RoundMode::CAST_NONE, this->processDataNum);
743+ Cast(tmp2, tmp1, RoundMode::CAST_NONE, this->processDataNum);
744+ Cast(yLocal, tmp2, RoundMode::CAST_RINT, this->processDataNum);
745+ }
746+ else if constexpr ((std::is_same_v<DTYPE_X, int64_t> && std::is_same_v<DTYPE_Y, bool>) ||
747+ (std::is_same_v<DTYPE_X, bfloat16_t> && std::is_same_v<DTYPE_Y, bool>))
748+ {
749+ LocalTensor<half> tmp1 = tmp2Bytes1.Get<half>();
750+ LocalTensor<float> tmp2 = tmp4Bytes1.Get<float>();
751+ if constexpr (std::is_same_v<DTYPE_X, int64_t>)
752+ {
753+ Cast(tmp2, xLocal, RoundMode::CAST_ROUND, this->processDataNum);
754+ }
755+ else if constexpr (std::is_same_v<DTYPE_X, bfloat16_t>)
756+ {
757+ Cast(tmp2, xLocal, RoundMode::CAST_NONE, this->processDataNum);
758+ }
759+ Cast(tmp1, tmp2, RoundMode::CAST_CEIL, this->processDataNum);
760+ Abs(tmp1, tmp1, this->processDataNum);
761+ Mins(tmp1, tmp1, HALF_ONE, this->processDataNum);
762+ Cast(yLocal, tmp1, RoundMode::CAST_CEIL, this->processDataNum);
763+ }
764+ this->outQueueY.template EnQue<TYPE_Y>(yLocal);
765+ this->inQueueX.template FreeTensor(xLocal);
766+ }
767+ 
768+private:
769+ TBuf<QuePosition::VECCALC> tmp2Bytes1;
770+ TBuf<QuePosition::VECCALC> tmp4Bytes1;
771+};
772+ 
773+template <typename TYPE_X, typename TYPE_Y>
774+class KernelCast3TBuf2B1TBuf4B : public BaseKernelCast<TYPE_X, TYPE_Y>
775+{
776+ /*
777+ 3个2Bytes,1个4Bytes的临时变量
778+ int64 -> int8
779+ int64 -> uint8
780+ bfloat16 -> int8
781+ bfloat16 -> uint8
782+ */
783+public:
784+ __aicore__ inline KernelCast3TBuf2B1TBuf4B() {}
785+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum,
786+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
787+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
788+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
789+ uint32_t tailBlockNum, TPipe *pipeIn)
790+ {
791+ this->pipe = pipeIn;
792+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
793+ uint32_t coreNum = GetBlockIdx();
794+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
795+ this->tileDataNum = tileDataNum;
796+ if (coreNum < tailBlockNum)
797+ {
798+ this->coreDataNum = bigCoreDataNum;
799+ this->tileNum = finalBigTileNum;
800+ this->tailDataNum = bigTailDataNum;
801+ }
802+ else
803+ {
804+ this->coreDataNum = smallCoreDataNum;
805+ this->tileNum = finalSmallTileNum;
806+ this->tailDataNum = smallTailDataNum;
807+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
808+ }
809+ this->xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
810+ this->yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
811+ BufferInit();
812+ }
813+ __aicore__ inline void Process()
814+ {
815+ int32_t loopCount = this->tileNum;
816+ this->processDataNum = this->tileDataNum;
817+ for (int32_t i = 0; i < loopCount - 1; i++)
818+ {
819+ this->CopyIn(i);
820+ Compute(i);
821+ this->CopyOut(i);
822+ }
823+ this->processDataNum = this->tailDataNum;
824+ this->CopyIn(loopCount - 1);
825+ Compute(loopCount - 1);
826+ this->CopyOut(loopCount - 1);
827+ }
828+ 
829+private:
830+ __aicore__ inline void BufferInit()
831+ {
832+ this->pipe->InitBuffer(this->inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
833+ this->pipe->InitBuffer(this->outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
834+ this->pipe->InitBuffer(tmp2Bytes1, this->tileDataNum * sizeof(half));
835+ this->pipe->InitBuffer(tmp2Bytes2, this->tileDataNum * sizeof(half));
836+ this->pipe->InitBuffer(tmp2Bytes3, this->tileDataNum * sizeof(half));
837+ this->pipe->InitBuffer(tmp4Bytes1, this->tileDataNum * sizeof(float));
838+ }
839+ __aicore__ inline void Compute(int32_t progress)
840+ {
841+ LocalTensor<TYPE_X> xLocal = this->inQueueX.template DeQue<TYPE_X>();
842+ LocalTensor<TYPE_Y> yLocal = this->outQueueY.template AllocTensor<TYPE_Y>();
843+ if constexpr ((std::is_same_v<DTYPE_X, int64_t> && std::is_same_v<DTYPE_Y, int8_t>) ||
844+ (std::is_same_v<DTYPE_X, int64_t> && std::is_same_v<DTYPE_Y, uint8_t>) ||
845+ (std::is_same_v<DTYPE_X, bfloat16_t> && std::is_same_v<DTYPE_Y, int8_t>) ||
846+ (std::is_same_v<DTYPE_X, bfloat16_t> && std::is_same_v<DTYPE_Y, uint8_t>))
847+ {
848+ LocalTensor<int16_t> tmp1 = tmp2Bytes1.Get<int16_t>();
849+ LocalTensor<int16_t> tmp2 = tmp2Bytes2.Get<int16_t>();
850+ LocalTensor<half> tmp3 = tmp2Bytes3.Get<half>();
851+ LocalTensor<int32_t> tmp4 = tmp4Bytes1.Get<int32_t>();
852+ if constexpr (std::is_same_v<DTYPE_X, int64_t>)
853+ {
854+ Cast(tmp4, xLocal, RoundMode::CAST_NONE, this->processDataNum);
855+ }
856+ else if constexpr (std::is_same_v<DTYPE_X, bfloat16_t>)
857+ {
858+ Cast(tmp4, xLocal, RoundMode::CAST_TRUNC, this->processDataNum);
859+ }
860+ Cast(tmp1, tmp4, RoundMode::CAST_NONE, this->processDataNum);
861+ Duplicate(tmp2, CONST_255, this->processDataNum);
862+ And(tmp1, tmp1, tmp2, this->processDataNum);
863+ if constexpr (std::is_same_v<DTYPE_Y, int8_t>)
864+ {
865+ Adds(tmp1, tmp1, CONST_128, this->processDataNum);
866+ And(tmp1, tmp1, tmp2, this->processDataNum);
867+ Adds(tmp1, tmp1, CONST_NE128, this->processDataNum);
868+ }
869+ Cast(tmp3, tmp1, RoundMode::CAST_NONE, this->processDataNum);
870+ Cast(yLocal, tmp3, RoundMode::CAST_NONE, this->processDataNum);
871+ }
872+ this->outQueueY.template EnQue<TYPE_Y>(yLocal);
873+ this->inQueueX.template FreeTensor(xLocal);
874+ }
875+ 
876+private:
877+ TBuf<QuePosition::VECCALC> tmp2Bytes1;
878+ TBuf<QuePosition::VECCALC> tmp2Bytes2;
879+ TBuf<QuePosition::VECCALC> tmp2Bytes3;
880+ TBuf<QuePosition::VECCALC> tmp4Bytes1;
881+};
882+ 
883+class KernelCastTQueBind
884+{
885+ /*
886+ 使用TQueBind直接传输8bit的数据类型
887+ bool -> int8/uint8
888+ int8 -> uint8
889+ uint8 -> int8
890+ */
891+public:
892+ __aicore__ inline KernelCastTQueBind() {}
893+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum,
894+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
895+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
896+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
897+ uint32_t tailBlockNum, TPipe *pipeIn)
898+ {
899+ pipe = pipeIn;
900+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
901+ uint32_t coreNum = GetBlockIdx();
902+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
903+ this->tileDataNum = tileDataNum;
904+ if (coreNum < tailBlockNum)
905+ {
906+ this->coreDataNum = bigCoreDataNum;
907+ this->tileNum = finalBigTileNum;
908+ this->tailDataNum = bigTailDataNum;
909+ }
910+ else
911+ {
912+ this->coreDataNum = smallCoreDataNum;
913+ this->tileNum = finalSmallTileNum;
914+ this->tailDataNum = smallTailDataNum;
915+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
916+ }
917+ xGm.SetGlobalBuffer((__gm__ uint8_t *)x + globalBufferIndex, this->coreDataNum);
918+ yGm.SetGlobalBuffer((__gm__ uint8_t *)y + globalBufferIndex, this->coreDataNum);
919+ pipe->InitBuffer(queBind, BUFFER_NUM, this->tileDataNum * sizeof(uint8_t));
920+ }
921+ __aicore__ inline void Process()
922+ {
923+ int32_t loopCount = this->tileNum;
924+ this->processDataNum = this->tileDataNum;
925+ for (int32_t i = 0; i < loopCount - 1; i++)
926+ {
927+ auto bindLocal = queBind.AllocTensor<uint8_t>();
928+ DataCopy(bindLocal, xGm[i * this->tileDataNum], this->processDataNum);
929+ queBind.EnQue(bindLocal);
930+ bindLocal = queBind.DeQue<uint8_t>();
931+ DataCopy(yGm[i * this->tileDataNum], bindLocal, this->processDataNum);
932+ queBind.FreeTensor(bindLocal);
933+ }
934+ this->processDataNum = this->tailDataNum;
935+ auto bindLocal = queBind.AllocTensor<uint8_t>();
936+ DataCopy(bindLocal, xGm[(loopCount - 1) * this->tileDataNum], this->processDataNum);
937+ queBind.EnQue(bindLocal);
938+ bindLocal = queBind.DeQue<uint8_t>();
939+ DataCopy(yGm[(loopCount - 1) * this->tileDataNum], bindLocal, this->processDataNum);
940+ queBind.FreeTensor(bindLocal);
941+ }
942+ 
943+private:
944+ TPipe *pipe;
945+ TQueBind<TPosition::VECIN, TPosition::VECOUT, BUFFER_NUM> queBind;
946+ GlobalTensor<uint8_t> xGm;
947+ GlobalTensor<uint8_t> yGm;
948+ uint32_t coreDataNum = 0;
949+ uint32_t tileNum = 0;
950+ uint32_t tileDataNum = 0;
951+ uint32_t tailDataNum = 0;
952+ uint32_t processDataNum = 0;
953+};
954+ 
955+} // namespace NsCast
956+#endif // CAST_H
Aexperimental/math/cast/op_kernel/cast_tiling_data.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 cast_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef __CAST_TILLING_DATA_H__
17+#define __CAST_TILLING_DATA_H__
18+ 
19+struct CastTilingData {
20+ uint32_t smallCoreDataNum;
21+ uint32_t bigCoreDataNum;
22+ uint32_t finalBigTileNum;
23+ uint32_t finalSmallTileNum;
24+ uint32_t tileDataNum;
25+ uint32_t smallTailDataNum;
26+ uint32_t bigTailDataNum;
27+ uint32_t tailBlockNum;
28+};
29+#endif
Aexperimental/math/cast/op_kernel/cast_tiling_key.h+35-0
@@ -0,0 +1,35 @@
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 cast_tiling_key.h
13+ * \brief cast tiling key declare
14+ */
15+ 
16+#ifndef __CAST_TILING_KEY_H__
17+#define __CAST_TILING_KEY_H__
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+/* Mode场景定义 */
22+#define ELEMENTWISE_TPL_SCH_MODE_0 0
23+#define ELEMENTWISE_TPL_SCH_MODE_1 1
24+/* 继续定义其他Mode场景... */
25+ 
26+/* 模板参数 */
27+ASCENDC_TPL_ARGS_DECL(
28+ Cast,
29+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
30+ 
31+/* 模板参数组合 */
32+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
33+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));
34+ 
35+#endif