[CANNBOT]scale适配Ascend950 Ascendc实现-ops-math-AtomGit

Mdocs/zh/op_api_list.md+1-1

	@@ -273,7 +273,7 @@
273	\| [aclnnRsqrt&aclnnInplaceRsqrt](../../math/rsqrt/docs/aclnnRsqrt&aclnnInplaceRsqrt.md) \| 求input(Tensor)每个元素的平方根的倒数。 \| 默认确定性实现\| 默认确定性实现\|	273	\| [aclnnRsqrt&aclnnInplaceRsqrt](../../math/rsqrt/docs/aclnnRsqrt&aclnnInplaceRsqrt.md) \| 求input(Tensor)每个元素的平方根的倒数。 \| 默认确定性实现\| 默认确定性实现\|
274	\| [aclnnRsub](../../math/sub/docs/aclnnRsub.md) \| 完成减法计算。 \| 默认确定性实现\| 默认确定性实现 \|	274	\| [aclnnRsub](../../math/sub/docs/aclnnRsub.md) \| 完成减法计算。 \| 默认确定性实现\| 默认确定性实现 \|
275	\| [aclnnRsubs](../../math/sub/docs/aclnnRsubs.md) \| 完成减法计算。 \| 默认确定性实现\| 默认确定性实现 \|	275	\| [aclnnRsubs](../../math/sub/docs/aclnnRsubs.md) \| 完成减法计算。 \| 默认确定性实现\| 默认确定性实现 \|
276	-\| [aclnnScale](../../math/scale/docs/aclnnScale.md) \| 参见算子文档。 \| 默认确定性实现\| - \|	276	+\| [aclnnScale](../../math/scale/docs/aclnnScale.md) \| 对输入Tensor进行scale和bias计算。若不输入bias，则 $y = x \cdot scale$；若输入bias，则 $y = x \cdot scale + bias$。 \| 默认确定性实现\| 默认确定性实现 \|
277	\| [aclnnSearchSorted](../../math/search_sorted/docs/aclnnSearchSorted.md) \| 在一个已排序的张量（sortedSequence）中查找给定tensor值（self）应该插入的位置。 \| 默认确定性实现\| 默认确定性实现 \|	277	\| [aclnnSearchSorted](../../math/search_sorted/docs/aclnnSearchSorted.md) \| 在一个已排序的张量（sortedSequence）中查找给定tensor值（self）应该插入的位置。 \| 默认确定性实现\| 默认确定性实现 \|
278	\| [aclnnSearchSorteds](../../math/search_sorted/docs/aclnnSearchSorteds.md) \| 在一个已排序的一维张量（sortedSequence）中查找给定Scalar值（self）应该插入的位置。 \| 默认确定性实现\| - \|	278	\| [aclnnSearchSorteds](../../math/search_sorted/docs/aclnnSearchSorteds.md) \| 在一个已排序的一维张量（sortedSequence）中查找给定Scalar值（self）应该插入的位置。 \| 默认确定性实现\| - \|
279	\| [aclnnSign](../../math/sign/docs/aclnnSign.md) \| 对输入的tensor逐元素进行Sign符号函数的运算并输出结果tensor。 \| 默认确定性实现\| 默认确定性实现\|	279	\| [aclnnSign](../../math/sign/docs/aclnnSign.md) \| 对输入的tensor逐元素进行Sign符号函数的运算并输出结果tensor。 \| 默认确定性实现\| 默认确定性实现\|

Mdocs/zh/op_list.md+5-5

	@@ -1579,12 +1579,12 @@
1579	<tr>	1579	<tr>
1580	<td>math</td>	1580	<td>math</td>
1581	<td><a href="../../math/scale/README.md">scale</a></td>	1581	<td><a href="../../math/scale/README.md">scale</a></td>
1582	- <td>×</td>	1582	+ <td>√</td>
1583	- <td>×</td>	1583	+ <td>√</td>
1584	- <td>×</td>	1584	+ <td>√</td>
1585	- <td>×</td>	1585	+ <td>√</td>
1586	<td>AI Core</td>	1586	<td>AI Core</td>
1587	- <td>该算子暂无Ascend C代码实现，欢迎开发者补充贡献，贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>	1587	+ <td>对输入x按scale缩放并可选叠加bias：y = x * scale + bias，支持axis/num_axes/scale_from_blob控制的通道广播。</td>
1588	</tr>	1588	</tr>
1589	<tr>	1589	<tr>
1590	<td>math</td>	1590	<td>math</td>

Mmath/acosh_grad/README.md+6-2

	@@ -14,9 +14,13 @@
14	## 功能说明	14	## 功能说明
15		15
16	- 算子功能：计算Acosh（反双曲余弦）算子的反向梯度。	16	- 算子功能：计算Acosh（反双曲余弦）算子的反向梯度。
17	-- 算子公式：$$	17	+- 算子公式：
		18	+
		19	+ $$
18	z_i = dy_i \cdot \dfrac{1}{\sqrt{y_i^2 - 1}}	20	z_i = dy_i \cdot \dfrac{1}{\sqrt{y_i^2 - 1}}
19	- $$其中：	21	+ $$
		22	+
		23	+ 其中：
20	- $y_i$为前向Acosh算子的输入张量，值域期望$\geq 1$；	24	- $y_i$为前向Acosh算子的输入张量，值域期望$\geq 1$；
21	- $dy_i$为上游传入的梯度；	25	- $dy_i$为上游传入的梯度；
22	- $z_i$为对原始输入张量的梯度，等于上游梯度乘以$1/\sqrt{y_i^2 - 1}$。	26	- $z_i$为对原始输入张量的梯度，等于上游梯度乘以$1/\sqrt{y_i^2 - 1}$。

Mmath/scale/CMakeLists.txt+13-17

	@@ -1,20 +1,16 @@
1	-# ----------------------------------------------------------------------------	1	+# ----------------------------------------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3	-# Copyright (c) 2025 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	-# This file is a part of the CANN Open Software.	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6	# Please refer to the License for details. You may not use this file except in compliance with the License.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	# See LICENSE in the root of the software repository for the full text of the License.	8	# See LICENSE in the root of the software repository for the full text of the License.
10	-# ----------------------------------------------------------------------------	9	+# ----------------------------------------------------------------------------------------------------------
		10	+# Generated By CANNBot
11		11
12	-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)	12	+# 设置算子定义时支持的芯片类型
13	-if(NOT ENABLE_TEST AND NOT BENCHMARK)	13	+set(SUPPORT_COMPUTE_UNIT "ascend950")
14	- list(REMOVE_ITEM CURRENT_DIRS tests)	14	+# 设置每种芯片类型对应的 tiling 文件目录
15	-endif()	15	+set(SUPPORT_TILING_DIR "arch35")
16	-foreach(SUB_DIR ${CURRENT_DIRS})	16	+add_all_modules_sources(OPTYPE scale ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP FALSE)
17	- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18	- add_subdirectory(${SUB_DIR})
19	- endif()
20	-endforeach()

Mmath/scale/README.md+120-1

	@@ -1,3 +1,122 @@
1	# Scale	1	# Scale
2		2
3	-本目录仅包含Scale算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。	3	+## 产品支持情况
		4	+
		5	+\| 产品 \| 是否支持 \|
		6	+\| :----------------------------------------------------------- \| :------: \|
		7	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
		8	+\| <term>Atlas A3训练系列产品/Atlas A3推理系列产品</term> \| √ \|
		Y陈 Y yang-jiwei6月16日看到下面算子的kernel代码与op_def的定义只有950相关的，请确认一下接口支持范围 W wangweidong 6月24日评论： 6月24日评论：陈陈佳良6月24日问题: README 文档与实际实现不一致。表格中标记 Atlas A3/A2/推理系列产品为支持（√），但实际代码只有 ascend950 (arch35) 的 Ascend C 实现。文档误导可能导致用户在不支持的芯片上调用算子失败。建议: 根据实际实现情况更新支持列表，或在表格中注明仅 ascend950 有 Ascend C 实现，其他芯片通过其他方式支持 W wangweidong 6月25日评论： 6月25日评论：
		9	+\| <term>Atlas A2训练系列产品/Atlas A2推理系列产品</term> \| √ \|
		10	+\| <term>Atlas 200I/500 A2推理产品</term> \| × \|
		11	+\| <term>Atlas推理系列产品</term> \| √ \|
		12	+\| <term>Atlas训练系列产品</term> \| × \|
		13	+
		14	+## 功能说明
		15	+
		16	+- 算子功能：对输入Tensor进行scale和bias计算。若不输入bias，则 $y = x \cdot scale$；若输入bias，则 $y = x \cdot scale + bias$。
		17	+- 算子公式：
		18	+
		19	+ 若不输入bias，则
		20	+
		21	+ $$
		22	+ y_i = x_i \cdot scale_i
		23	+ $$
		24	+
		25	+ 若输入bias，则
		26	+
		27	+ $$
		28	+ y_i = x_i \cdot scale_i + bias_i
		29	+ $$
		30	+
		31	+ 其中：
		32	+ - $x_i$为输入Tensor；
		33	+ - $scale_i$为缩放因子Tensor，支持与x的broadcast；
		34	+ - $bias_i$为可选的偏置Tensor，shape与scale保持一致；
		35	+ - $y_i$为输出Tensor，shape与x一致。
		36	+
		37	+## 参数说明
		38	+
		39	+<table style="undefined;table-layout: fixed; width: 1576px"><colgroup>
		40	+ <col style="width: 170px">
		41	+ <col style="width: 170px">
		42	+ <col style="width: 310px">
		43	+ <col style="width: 212px">
		44	+ <col style="width: 100px">
		45	+ </colgroup>
		46	+ <thead>
		47	+ <tr>
		48	+ <th>参数名</th>
		49	+ <th>输入/输出/属性</th>
		50	+ <th>描述</th>
		51	+ <th>数据类型</th>
		52	+ <th>数据格式</th>
		53	+ </tr></thead>
		54	+ <tbody>
		55	+ <tr>
		56	+ <td>x</td>
		57	+ <td>输入</td>
		58	+ <td>算子输入的Tensor。支持空Tensor。</td>
		59	+ <td>FLOAT16, FLOAT32, BFLOAT16</td>
		60	+ <td>ND</td>
		61	+ </tr>
		62	+ <tr>
		63	+ <td>scale</td>
		64	+ <td>输入</td>
		65	+ <td>缩放因子Tensor。数据类型需与x一致，shape满足broadcast要求。</td>
		66	+ <td>FLOAT16, FLOAT32, BFLOAT16</td>
		67	+ <td>ND</td>
		68	+ </tr>
		69	+ <tr>
		70	+ <td>bias</td>
		71	+ <td>可选输入</td>
		72	+ <td>偏置Tensor。不为空时数据类型需与scale一致，shape与scale保持一致。</td>
		73	+ <td>FLOAT16, FLOAT32, BFLOAT16</td>
		74	+ <td>ND</td>
		75	+ </tr>
		76	+ <tr>
		77	+ <td>axis</td>
		78	+ <td>属性</td>
		79	+ <td>指定进行scale的起始轴。取值范围[-x_rank, x_rank)。</td>
		80	+ <td>INT64</td>
		81	+ <td>-</td>
		82	+ </tr>
		83	+ <tr>
		84	+ <td>num_axes</td>
		85	+ <td>属性</td>
		86	+ <td>指定进行scale的轴长度。取值范围>=-1，-1表示从axis轴scale到最后一轴。</td>
		87	+ <td>INT64</td>
		88	+ <td>-</td>
		89	+ </tr>
		90	+ <tr>
		91	+ <td>scale_from_blob</td>
		92	+ <td>属性</td>
		93	+ <td>True：使用numAxes + axis推导scale shape；False：从axis开始按scale的rank推导，忽略numAxes。</td>
		94	+ <td>BOOL</td>
		95	+ <td>-</td>
		96	+ </tr>
		97	+ <tr>
		98	+ <td>y</td>
		99	+ <td>输出</td>
		100	+ <td>输出Tensor。shape和数据类型与x一致。</td>
		101	+ <td>FLOAT16, FLOAT32, BFLOAT16</td>
		102	+ <td>ND</td>
		103	+ </tr>
		104	+ </tbody></table>
		105	+
		106	+## 约束说明
		107	+
		108	+- x与y的shape必须完全一致。
		109	+- x、scale、bias、y的dtype必须一致（均为FLOAT16/FLOAT32/BFLOAT16之一）。
		110	+- bias不为空时，bias与scale的shape必须一致。
		111	+- scale的shape需满足broadcast规则（参见功能说明）。
		112	+- x和scale的shape维度不大于8。
		113	+- axis取值范围为[-x_rank, x_rank)。
		114	+- numAxes取值范围>=-1。
		115	+- 仅支持ND格式。
		116	+- 支持[非连续的Tensor](../../docs/zh/context/非连续的Tensor.md)，非连续的Tensor维度不大于8。
		117	+
		118	+## 调用说明
		119	+
		120	+\| 调用方式 \| 调用样例 \| 说明 \|
		121	+\|---------\|---------------------------------------------------\|------------------------------------------------------------------\|
		122	+\| aclnn调用 \| [test_aclnn_scale](./examples/test_aclnn_scale.cpp) \| 通过[aclnnScale](./docs/aclnnScale.md)接口方式调用Scale算子。 \|

Mmath/scale/docs/aclnnScale.md+1-1

	@@ -4,7 +4,7 @@
4		4
5	\| 产品 \| 是否支持 \|	5	\| 产品 \| 是否支持 \|
6	\| :----------------------------------------------------------- \| :------: \|	6	\| :----------------------------------------------------------- \| :------: \|
7	-\| <term>Ascend 950PR/Ascend 950DT</term> \| × \|	7	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
8	\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|	8	\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|
9	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|	9	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|
10	\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|	10	\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|

Amath/scale/examples/arch35/test_aclnn_scale.cpp+174-0

@@ -0,0 +1,174 @@
	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	+/* Generated By CANNBot */
	12	+/**
	13	+ * This program is free software, you can redistribute it and/or modify.
	14	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	15	+ * This file is a part of the CANN Open Software.
	16	+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	17	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	18	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	19	+ * See LICENSE in the root of the software repository for the full text of the License.
	20	+ */
	21	+
	22	+#include <iostream>
	23	+#include <vector>
	24	+#include "acl/acl.h"
	25	+#include "aclnnop/aclnn_scale.h"
	26	+
	27	+#define CHECK_RET(cond, return_expr) \
	28	+ do { \
	29	+ if (!(cond)) { \
	30	+ return_expr; \
	31	+ } \
	32	+ } while (0)
	33	+
	34	+#define LOG_PRINT(message, ...) \
	35	+ do { \
	36	+ printf(message, ##__VA_ARGS__); \
	37	+ } while (0)
	38	+
	39	+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
	40	+ int64_t shapeSize = 1;
	41	+ for (auto i : shape) {
	42	+ shapeSize *= i;
	43	+ }
	44	+ return shapeSize;
	45	+}
	46	+
	47	+int Init(int32_t deviceId, aclrtStream* stream) {
	48	+ // 固定写法，资源初始化
	49	+ auto ret = aclInit(nullptr);
	50	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	51	+ ret = aclrtSetDevice(deviceId);
	52	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	53	+ ret = aclrtCreateStream(stream);
	54	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	55	+ return 0;
	56	+}
	57	+
	58	+template <typename T>
	59	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
	60	+ aclDataType dataType, aclTensor** tensor) {
	61	+ auto size = GetShapeSize(shape) * sizeof(T);
	62	+ // 调用aclrtMalloc申请device侧内存
	63	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	64	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	65	+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
	66	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	67	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	68	+
	69	+ // 计算连续tensor的strides
	70	+ std::vector<int64_t> strides(shape.size(), 1);
	71	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	72	+ strides[i] = shape[i + 1] * strides[i + 1];
	73	+ }
	74	+
	75	+ // 调用aclCreateTensor接口创建aclTensor
	76	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	77	+ shape.data(), shape.size(), *deviceAddr);
	78	+ return 0;
	79	+}
	80	+
	81	+int main() {
	82	+ // 1. （固定写法）device/stream初始化，参考acl API手册
	83	+ // 根据自己的实际device填写deviceId
	84	+ int32_t deviceId = 0;
	85	+ aclrtStream stream;
	86	+ auto ret = Init(deviceId, &stream);
	87	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	88	+
	89	+ // 2. 构造输入与输出，需要根据API的接口自定义构造
	90	+ std::vector<int64_t> selfShape = {4, 2};
	91	+ std::vector<int64_t> tensor1Shape = {4};
	陈陈陈佳良6月24日问题: 示例代码参数配置错误。scale 和 bias 的 shape 定义为 {4}（4个元素），但 hostData 却有 8 个元素，不匹配。根据 axis=0, numAxes=1, scaleFromBlob=true，scale 应为 {2}（匹配 x 的第一维大小）。错误配置会导致示例运行时出错或产生错误结果。建议: 修正为：tensor1Shape = {2}，tensor1HostData = {2, 2}，tensor2Shape = {2}，tensor2HostData = {2, 2} W wangweidong 6月24日评论： 6月24日评论：
	92	+ std::vector<int64_t> tensor2Shape = {4};
	93	+ std::vector<int64_t> outShape = {4, 2};
	94	+ void* selfDeviceAddr = nullptr;
	95	+ void* tensor1DeviceAddr = nullptr;
	96	+ void* tensor2DeviceAddr = nullptr;
	97	+ void* outDeviceAddr = nullptr;
	98	+ aclTensor* self = nullptr;
	99	+ aclTensor* tensor1 = nullptr;
	100	+ aclTensor* tensor2 = nullptr;
	101	+ aclTensor* out = nullptr;
	102	+
	103	+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
	104	+ std::vector<float> tensor1HostData = {2, 2, 2, 2};
	105	+ std::vector<float> tensor2HostData = {2, 2, 2, 2};
	106	+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
	107	+ int64_t axis = 0;
	108	+ int64_t numAxes = 1;
	109	+ bool fromBlob = true;
	110	+
	111	+ // 创建self aclTensor
	112	+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
	113	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	114	+ // 创建tensor1 aclTensor
	115	+ ret = CreateAclTensor(tensor1HostData, tensor1Shape, &tensor1DeviceAddr, aclDataType::ACL_FLOAT, &tensor1);
	116	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	117	+ // 创建tensor2 aclTensor
	118	+ ret = CreateAclTensor(tensor2HostData, tensor2Shape, &tensor2DeviceAddr, aclDataType::ACL_FLOAT, &tensor2);
	119	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	120	+ // 创建out aclTensor
	121	+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
	122	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	123	+
	124	+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
	125	+ uint64_t workspaceSize = 0;
	126	+ aclOpExecutor* executor;
	127	+ // 调用aclnnScale第一段接口
	128	+ ret = aclnnScaleGetWorkspaceSize(self, tensor1, tensor2, axis, numAxes, fromBlob, out, &workspaceSize, &executor);
	129	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnScaleGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
	130	+ // 根据第一段接口计算出的workspaceSize申请device内存
	131	+ void* workspaceAddr = nullptr;
	132	+ if (workspaceSize > 0) {
	133	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	134	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
	135	+ }
	136	+ // 调用aclnnScale第二段接口
	137	+ ret = aclnnScale(workspaceAddr, workspaceSize, executor, stream);
	138	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnScale failed. ERROR: %d\n", ret); return ret);
	139	+
	140	+ // 4. （固定写法）同步等待任务执行结束
	141	+ ret = aclrtSynchronizeStream(stream);
	142	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	143	+
	144	+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
	145	+ auto size = GetShapeSize(outShape);
	146	+ std::vector<float> resultData(size, 0);
	147	+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]),
	148	+ outDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
	149	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy resultData from device to host failed. ERROR: %d\n", ret);
	150	+ return ret);
	151	+ for (int64_t i = 0; i < size; i++) {
	152	+ LOG_PRINT("resultData[%ld] is: %f\n", i, resultData[i]);
	153	+ }
	154	+
	155	+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
	156	+ aclDestroyTensor(self);
	157	+ aclDestroyTensor(tensor1);
	158	+ aclDestroyTensor(tensor2);
	159	+ aclDestroyTensor(out);
	160	+
	161	+ // 7.释放device资源，需要根据具体API的接口定义修改
	162	+ aclrtFree(selfDeviceAddr);
	163	+ aclrtFree(tensor1DeviceAddr);
	164	+ aclrtFree(tensor2DeviceAddr);
	165	+ aclrtFree(outDeviceAddr);
	166	+ if (workspaceSize > 0) {
	167	+ aclrtFree(workspaceAddr);
	168	+ }
	169	+ aclrtDestroyStream(stream);
	170	+ aclrtResetDevice(deviceId);
	171	+ aclFinalize();
	172	+
	173	+ return 0;
	174	+}

Rmath/scale/op_host/op_api/aclnn_scale.cpp→math/scale/op_api/aclnn_scale.cpp+1-0

	@@ -19,6 +19,7 @@
19	#include "opdev/data_type_utils.h"	19	#include "opdev/data_type_utils.h"
20	#include "opdev/format_utils.h"	20	#include "opdev/format_utils.h"
21	#include "opdev/make_op_executor.h"	21	#include "opdev/make_op_executor.h"
		22	+#include "opdev/platform.h"
22	#include "aclnn_kernels/common/op_error_check.h"	23	#include "aclnn_kernels/common/op_error_check.h"
23		24
24	using namespace op;	25	using namespace op;

Rmath/scale/op_host/op_api/aclnn_scale.h→math/scale/op_api/aclnn_scale.h+3-3

	@@ -26,11 +26,11 @@ extern "C" {
26	* 支持非连续的Tensor。	26	* 支持非连续的Tensor。
27	* @param [in] scale: npu device侧的aclTensor, 数据类型支持float, bf16, float16	27	* @param [in] scale: npu device侧的aclTensor, 数据类型支持float, bf16, float16
28	* @param [in] bias: npu device侧的aclTensor，数据类型支持float, bf16, float16	28	* @param [in] bias: npu device侧的aclTensor，数据类型支持float, bf16, float16
29	- * @param [in] axis: host侧的aclScalar，数据类型int64_t	29	+ * @param [in] axis: host侧的int64_t类型参数
30	* @param [in] numAxes: host侧的aclScalar，数据类型int64_t	30	* @param [in] numAxes: host侧的aclScalar，数据类型int64_t
31	- * @param [in] scaleFromBlob: host侧的aclScalar, 数据类型bool	31	+ * @param [in] scaleFromBlob: host侧的bool类型参数
32	* @param [in] y: Scale计算的出参。npu device侧的aclTensor,	32	* @param [in] y: Scale计算的出参。npu device侧的aclTensor,
33	- * 数据类型支持int8, 数据格式支持ND,	33	+ * 数据类型支持float16, bfloat16, float32, 数据格式支持ND,
34	* 支持非连续的Tensor。	34	* 支持非连续的Tensor。
35	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。	35	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
36	* @param [out] executor: 返回op执行器，包含算子计算流程。	36	* @param [out] executor: 返回op执行器，包含算子计算流程。

Rmath/scale/op_host/op_api/scale.cpp→math/scale/op_api/scale.cpp+0-0

文件重命名但无更改。

Rmath/scale/op_host/op_api/scale.h→math/scale/op_api/scale.h+0-0

文件重命名但无更改。

Amath/scale/op_host/arch35/scale_tiling_arch35.cpp+502-0

@@ -0,0 +1,502 @@
	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	+/* Generated By CANNBot */
	12	+/**
	13	+ * Scale Tiling — arch35 实现 (RegBase Broadcast 范式)
	14	+ * 基于 adam_apply_one_assign_tiling_arch35.cpp 模板适配
	15	+ * 目标架构: DAV_3510 (arch35 / ascend950)
	16	+ */
	17	+#include "register/op_def_registry.h"
	18	+#include "op_common/log/log.h"
	19	+#include "op_common/op_host/util/math_util.h"
	20	+#include "op_common/op_host/util/platform_util.h"
	21	+#include "../../op_kernel/arch35/scale_tiling_struct.h"
	22	+#include "../../op_kernel/arch35/scale_struct.h"
	23	+#include "scale_tiling_arch35.h"
	24	+
	25	+#include <algorithm>
	26	+#include <sstream>
	27	+#include <vector>
	28	+
	29	+namespace optiling {
	30	+
	31	+using Ops::Base::CeilDiv;
	32	+
	33	+// ============================================================
	34	+// Tiling 函数模块 — namespace scale
	35	+// ============================================================
	36	+
	37	+namespace scale {
	38	+
	39	+// 5.0 CheckBroadcastShape — 补 1 后逐维校验 broadcast 兼容性：同维非 1 大小必须一致
	40	+// 单组（入参或出参）在第 d 维校验，非 1 大小须与共享 ref 一致；ref 跨入参/出参累积
	41	+static bool CheckGroupDim(const std::vector<std::vector<int64_t>>& group, int64_t d,
	42	+ const char* tag, int64_t& ref)
	43	+{
	44	+ for (size_t i = 0; i < group.size(); i++) {
	45	+ int64_t v = group[i][d];
	46	+ if (v == 1) continue;
	47	+ if (ref == -1) {
	48	+ ref = v;
	49	+ } else if (v != ref) {
	50	+ OP_LOGE("PadAndSqueeze", "dim %d broadcast incompatible: %s[%d] size %d != %d",
	51	+ (int)d, tag, (int)i, (int)v, (int)ref);
	52	+ return false;
	53	+ }
	54	+ }
	55	+ return true;
	56	+}
	57	+
	58	+static bool CheckBroadcastShape(
	59	+ const std::vector<std::vector<int64_t>>& padded_in,
	60	+ const std::vector<std::vector<int64_t>>& padded_out,
	61	+ int64_t max_rank)
	62	+{
	63	+ for (int64_t d = 0; d < max_rank; d++) {
	64	+ int64_t ref = -1;
	65	+ if (!CheckGroupDim(padded_in, d, "input", ref)) return false;
	66	+ if (!CheckGroupDim(padded_out, d, "output", ref)) return false;
	67	+ }
	68	+ return true;
	69	+}
	70	+
	71	+// 补 1 — 低 rank 在 shape 最前面补 1，拉齐到 max_rank
	72	+static std::vector<std::vector<int64_t>> PadFront(
	73	+ const std::vector<std::vector<int64_t>>& shapes, int64_t max_rank)
	74	+{
	75	+ std::vector<std::vector<int64_t>> out(shapes.size());
	76	+ for (size_t i = 0; i < shapes.size(); i++) {
	77	+ out[i].assign(max_rank - (int64_t)shapes[i].size(), 1);
	78	+ out[i].insert(out[i].end(), shapes[i].begin(), shapes[i].end());
	79	+ }
	80	+ return out;
	81	+}
	82	+
	83	+// 单维去 1：所有入参+出参在该维都为 1 则废维 squeeze；否则取 max 并保留归一化分量
	84	+static bool SqueezeDim(
	85	+ const std::vector<std::vector<int64_t>>& padded_in,
	86	+ const std::vector<std::vector<int64_t>>& padded_out, int64_t d,
	87	+ std::vector<int64_t>& maximum_bro_shape,
	88	+ std::vector<std::vector<int64_t>>& normal_input_shapes,
	89	+ std::vector<std::vector<int64_t>>& normal_output_shapes)
	90	+{
	91	+ bool all_one = true;
	92	+ int64_t max_dim = 0;
	93	+ for (auto& s : padded_in) { if (s[d] != 1) all_one = false; max_dim = std::max(max_dim, s[d]); }
	94	+ for (auto& s : padded_out) { if (s[d] != 1) all_one = false; max_dim = std::max(max_dim, s[d]); }
	95	+ if (all_one) return false;
	96	+ maximum_bro_shape.push_back(max_dim);
	97	+ for (size_t i = 0; i < padded_in.size(); i++) normal_input_shapes[i].push_back(padded_in[i][d]);
	98	+ for (size_t i = 0; i < padded_out.size(); i++) normal_output_shapes[i].push_back(padded_out[i][d]);
	99	+ return true;
	100	+}
	101	+
	102	+// 5.1 PadAndSqueeze
	103	+bool PadAndSqueeze(
	104	+ const std::vector<std::vector<int64_t>>& input_shapes,
	105	+ const std::vector<std::vector<int64_t>>& output_shapes,
	106	+ std::vector<int64_t>& maximum_bro_shape,
	107	+ std::vector<std::vector<int64_t>>& normal_input_shapes,
	108	+ std::vector<std::vector<int64_t>>& normal_output_shapes)
	109	+{
	110	+ int64_t num_inputs = (int64_t)input_shapes.size();
	111	+ int64_t num_outputs = (int64_t)output_shapes.size();
	112	+ int64_t max_rank = 0;
	113	+ for (auto& s : input_shapes) max_rank = std::max(max_rank, (int64_t)s.size());
	114	+ for (auto& s : output_shapes) max_rank = std::max(max_rank, (int64_t)s.size());
	115	+ // 补 1 + 校验 broadcast 兼容性（同维非 1 大小必须一致）
	116	+ auto padded_in = PadFront(input_shapes, max_rank);
	117	+ auto padded_out = PadFront(output_shapes, max_rank);
	118	+ if (!CheckBroadcastShape(padded_in, padded_out, max_rank)) return false;
	119	+ // 去 1 — 废维 squeeze，其余轴归一化保留
	120	+ maximum_bro_shape.clear();
	121	+ normal_input_shapes.assign(num_inputs, std::vector<int64_t>());
	122	+ normal_output_shapes.assign(num_outputs, std::vector<int64_t>());
	123	+ for (int64_t d = 0; d < max_rank; d++) {
	124	+ SqueezeDim(padded_in, padded_out, d, maximum_bro_shape, normal_input_shapes, normal_output_shapes);
	125	+ }
	126	+ if (maximum_bro_shape.empty()) { // 全标量归一为 (1,)
	127	+ maximum_bro_shape.push_back(1);
	128	+ for (int64_t i = 0; i < num_inputs; i++) normal_input_shapes[i].push_back(1);
	129	+ for (int64_t i = 0; i < num_outputs; i++) normal_output_shapes[i].push_back(1);
	130	+ }
	131	+ return true;
	132	+}
	133	+
	134	+// 5.2 FindSplitAxis
	135	+// per_buf_elems 统一按 FP32 计算（cast.md §Tile）
	136	+bool FindSplitAxis(const std::vector<int64_t>& max_bro_shape,
	137	+ int64_t /dtype_size/, int64_t ub_per_core, int64_t phys_nodes, SplitResult& out)
	138	+{
	139	+ if (phys_nodes <= 0) { // 除 0 保护：phys_nodes 异常直接判失败
	140	+ OP_LOGE("FindSplitAxis", "phys_nodes(%ld) must be > 0", phys_nodes);
	141	+ return false;
	142	+ }
	143	+ int64_t per_buf_bytes = (ub_per_core / phys_nodes) & ~31LL; // 32B 对齐
	144	+ int64_t per_buf_elems = per_buf_bytes / 4; // 统一按 FP32 计算
	145	+ int64_t rank = (int64_t)max_bro_shape.size();
	146	+ int64_t inner = 1;
	147	+ for (int64_t k = rank - 1; k >= 0; k--) {
	148	+ if (max_bro_shape[k] * inner > per_buf_elems) {
	149	+ out.a_i = per_buf_elems / inner;
	150	+ if (out.a_i < 1) out.a_i = 1; // 除 0 保护：单元素超 buffer 时至少切 1，避免后续 /a_i %a_i 除 0
	151	+ out.a_o = (max_bro_shape[k] + out.a_i - 1) / out.a_i;
	152	+ int64_t rem = max_bro_shape[k] % out.a_i;
	153	+ out.a_i_tail = (rem == 0) ? out.a_i : rem;
	154	+ out.axis = k;
	155	+ return true;
	156	+ }
	157	+ if (k == 0) {
	158	+ out.axis = 0;
	159	+ out.a_i = max_bro_shape[0];
	160	+ out.a_o = 1;
	161	+ out.a_i_tail = max_bro_shape[0];
	162	+ return true;
	163	+ }
	164	+ inner *= max_bro_shape[k];
	165	+ }
	166	+ return true;
	167	+}
	168	+
	169	+// 5.3 MultiCoreSplit
	170	+bool MultiCoreSplit(const std::vector<int64_t>& max_bro_shape,
	171	+ const SplitResult& ub_split, int64_t max_cores, MultiCoreResult& out)
	172	+{
	173	+ int64_t k = ub_split.axis, outer_prod = 1;
	174	+ for (int64_t j = 0; j < k; j++) outer_prod *= max_bro_shape[j];
	175	+ out.total_tiles = outer_prod * ub_split.a_o;
	176	+ out.num_cores = (out.total_tiles < max_cores) ? out.total_tiles : max_cores;
	177	+ out.tiles_main = out.total_tiles / out.num_cores;
	178	+ out.cores_tail = out.total_tiles % out.num_cores;
	179	+ return true;
	180	+}
	181	+
	182	+// 5.5 地址偏移计算
	183	+bool PrecomputeInputStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides) {
	184	+ int64_t rank = (int64_t)s.size();
	185	+ strides.assign(rank, 0);
	186	+ for (int64_t d = rank - 1; d >= 0; d--) {
	187	+ if (s[d] == 1) { strides[d] = 0; continue; }
	188	+ int64_t prod = 1;
	189	+ for (int64_t j = d + 1; j < rank; j++) prod *= s[j];
	190	+ strides[d] = prod;
	191	+ }
	192	+ return true;
	193	+}
	194	+
	195	+bool PrecomputeOutputStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides) {
	196	+ int64_t rank = (int64_t)s.size();
	197	+ strides.assign(rank, 0);
	198	+ for (int64_t d = rank - 1; d >= 0; d--) {
	199	+ if (s[d] == 1) { strides[d] = 0; continue; }
	200	+ int64_t prod = 1;
	201	+ for (int64_t j = d + 1; j < rank; j++) prod *= s[j];
	202	+ strides[d] = prod;
	203	+ }
	204	+ return true;
	205	+}
	206	+
	207	+} // namespace scale
	208	+
	209	+// ============================================================
	210	+// ScaleTiling — CANN 主线 Tiling
	211	+// ============================================================
	212	+// ScaleCompileInfo 定义统一在 scale_tiling_arch35.h，避免重复定义
	213	+
	214	+static std::string Arr2String(const int64_t* arr, int64_t n)
	215	+{
	216	+ std::ostringstream oss;
	217	+ oss << "[";
	218	+ if (n > 0) {
	219	+ for (int64_t i = 0; i < n - 1; ++i) {
	220	+ oss << arr[i] << ",";
	221	+ }
	222	+ oss << arr[n - 1];
	223	+ }
	224	+ oss << "]";
	225	+ return oss.str();
	226	+}
	227	+
	228	+class ScaleTiling {
	229	+public:
	230	+ explicit ScaleTiling(gert::TilingContext* ctx) : ctx_(ctx) {}
	231	+
	232	+ ge::graphStatus RunTiling()
	233	+ {
	234	+ ge::graphStatus ret = GetShapeInfo();
	235	+ if (ret != ge::GRAPH_SUCCESS) return ret;
	236	+
	237	+ // 空 Tensor 防御：输出 shape 任一维为 0 则跳过 Tiling
	238	+ int64_t total_out = 1;
	239	+ for (auto d : raw_output_shapes_[0]) total_out *= d;
	240	+ if (total_out == 0) {
	241	+ ctx_->SetBlockDim(1);
	242	+ return ge::GRAPH_SUCCESS;
	243	+ }
	244	+
	245	+ int64_t mapped = (rank_ <= 4) ? 4 : 8;
	246	+ if (mapped == 4) {
	247	+ ret = DoTilingAndSet<4>();
	248	+ ctx_->SetTilingKey(GET_TPL_TILING_KEY(SCALE_RANK_4));
	249	+ } else {
	250	+ ret = DoTilingAndSet<8>();
	251	+ ctx_->SetTilingKey(GET_TPL_TILING_KEY(SCALE_RANK_8));
	252	+ }
	253	+ return ret;
	254	+ }
	255	+
	256	+private:
	257	+ // 读取平台信息：核数 + UB 大小
	258	+ ge::graphStatus ReadPlatform()
	259	+ {
	260	+ fe::PlatFormInfos* platformInfo = ctx_->GetPlatformInfo();
	261	+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, platformInfo);
	262	+ auto ap = platform_ascendc::PlatformAscendC(platformInfo);
	263	+ coreNum_ = ap.GetCoreNumAiv();
	264	+ ap.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize_);
	265	+ return ge::GRAPH_SUCCESS;
	266	+ }
	267	+
	268	+ // 读取输入 shape：x(0)、scale(1) REQUIRED，bias(2) OPTIONAL；并读取全部输出 shape
	269	+ // bias 是否存在以 OPTIONAL 输入是否实例化判定，与 aclnn bias!=nullptr 对齐，避免按输入个数误判越界
	270	+ ge::graphStatus ReadShapes()
	271	+ {
	272	+ auto readInto = [&](const gert::StorageShape* shape, std::vector<std::vector<int64_t>>& dst) {
	273	+ std::vector<int64_t> dims;
	274	+ gert::Shape s = shape->GetStorageShape();
	275	+ for (size_t d = 0; d < s.GetDimNum(); ++d) dims.push_back(s.GetDim(d));
	276	+ dst.push_back(dims);
	277	+ };
	278	+ auto xShape = ctx_->GetInputShape(0);
	279	+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, xShape);
	280	+ readInto(xShape, raw_input_shapes_);
	281	+ auto scShape = ctx_->GetInputShape(1);
	282	+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, scShape);
	283	+ readInto(scShape, raw_input_shapes_);
	284	+ auto biasShape = ctx_->GetOptionalInputShape(2);
	285	+ has_bias_ = (biasShape != nullptr) ? 1 : 0;
	286	+ if (biasShape != nullptr) readInto(biasShape, raw_input_shapes_);
	287	+ for (size_t i = 0; i < ctx_->GetComputeNodeInfo()->GetOutputsNum(); ++i) {
	288	+ auto shape = ctx_->GetOutputShape(i);
	289	+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, shape);
	290	+ readInto(shape, raw_output_shapes_);
	291	+ }
	292	+ return ge::GRAPH_SUCCESS;
	293	+ }
	294	+
	295	+ ge::graphStatus ReadDtypeSize()
	296	+ {
	297	+ auto inputDesc = ctx_->GetInputDesc(0);
	298	+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, inputDesc);
	299	+ ge::DataType dtype = inputDesc->GetDataType();
	300	+ if (dtype == ge::DT_FLOAT16 \|\| dtype == ge::DT_BF16) dtype_size_ = 2;
	301	+ else if (dtype == ge::DT_FLOAT) dtype_size_ = 4;
	302	+ else {
	303	+ OP_LOGE(ctx_->GetNodeName(), "Unsupported dtype");
	304	+ return ge::GRAPH_FAILED;
	305	+ }
	306	+ return ge::GRAPH_SUCCESS;
	307	+ }
	308	+
	309	+ // 读取 Scale 特有属性：axis, num_axes, scale_from_blob
	310	+ void ReadAttrs()
	311	+ {
	312	+ const auto* attrs = ctx_->GetAttrs();
	313	+ if (attrs == nullptr) return;
	314	+ const int64_t* axisPtr = attrs->GetInt(0);
	315	+ const int64_t* numAxesPtr = attrs->GetInt(1);
	316	+ const bool* sfbPtr = attrs->GetBool(2);
	317	+ axis_ = axisPtr ? *axisPtr : 1;
	318	+ num_axes_ = numAxesPtr ? *numAxesPtr : 1;
	319	+ scale_from_blob_ = sfbPtr ? *sfbPtr : true;
	320	+ }
	321	+
	322	+ // 单个 scale/bias shape 按 axis 位置预补 1 到 xRank（不改变标量与同 rank 情形）
	323	+ void PadVecToAxis(std::vector<int64_t>& vshape, int64_t xRank, int64_t newAxis) const
	324	+ {
	325	+ int64_t vrank = (int64_t)vshape.size();
	326	+ // 标量 (rank=0)：广播全 1，PadAndSqueeze 右补 1 即可，显式跳过预补
	327	+ if (vrank == 0) return;
	328	+ if (!(vrank < xRank && newAxis >= 0 && newAxis < xRank)) return;
	329	+ std::vector<int64_t> padded(xRank, 1);
	330	+ if (scale_from_blob_) {
	331	+ // scaleFromBlob=true: numAxes 个维度从 axis 位置开始（numAxes==0 为标量，保持全 1）
	332	+ int64_t scaleLen = (num_axes_ == -1) ? (xRank - newAxis) : num_axes_;
	333	+ if (scaleLen != 0)
	334	+ for (int64_t i = 0; i < vrank && (newAxis + i) < xRank; i++) padded[newAxis + i] = vshape[i];
	335	+ } else {
	336	+ // scaleFromBlob=false: 右对齐到 x 末尾（标准 broadcast 口径）；
	337	+ // 非法 shape 由后续 CheckBroadcastShape 逐维校验拦截并报错
	338	+ int64_t offset = xRank - vrank;
	339	+ for (int64_t i = 0; i < vrank; i++) padded[offset + i] = vshape[i];
	340	+ }
	341	+ vshape = padded;
	342	+ }
	343	+
	344	+ // Scale 特有：按 axis/numAxes/scaleFromBlob 预补 scale 和 bias shape
	345	+ // 标准 PadAndSqueeze 右对齐补 1，但 Caffe Scale 的 scale/bias 需按 axis 位置对齐
	346	+ // scale(idx=1) 与 bias(idx=2) 与 x 的广播口径相同，必须同时预补，否则 bias 错位
	347	+ void PrePadScaleBias()
	348	+ {
	349	+ if (raw_input_shapes_.size() < 2) return;
	350	+ int64_t xRank = (int64_t)raw_input_shapes_[0].size();
	351	+ // axis<0 归一，与 aclnn newAxis = axis>=0 ? axis : xRank+axis 口径一致
	352	+ int64_t newAxis = axis_ >= 0 ? axis_ : xRank + axis_;
	353	+ PadVecToAxis(raw_input_shapes_[1], xRank, newAxis); // scale
	354	+ if (raw_input_shapes_.size() >= 3) PadVecToAxis(raw_input_shapes_[2], xRank, newAxis); // bias
	355	+ }
	356	+
	357	+ ge::graphStatus GetShapeInfo()
	358	+ {
	359	+ if (ReadPlatform() != ge::GRAPH_SUCCESS) return ge::GRAPH_FAILED;
	360	+ if (ReadShapes() != ge::GRAPH_SUCCESS) return ge::GRAPH_FAILED;
	361	+ if (ReadDtypeSize() != ge::GRAPH_SUCCESS) return ge::GRAPH_FAILED;
	362	+ ReadAttrs();
	363	+ PrePadScaleBias();
	364	+
	365	+ if (!scale::PadAndSqueeze(raw_input_shapes_, raw_output_shapes_,
	366	+ max_bro_shape_, normal_input_shapes_, normal_output_shapes_)) {
	367	+ OP_LOGE(ctx_->GetNodeName(), "scale/bias shape cannot broadcast to x");
	368	+ return ge::GRAPH_FAILED;
	369	+ }
	370	+ rank_ = (int64_t)max_bro_shape_.size();
	371	+
	372	+ // rank>8 拦截：TilingData 数组上限为 kRank=8，超限会越界写，与 aclnn OP_CHECK_MAX_DIM(8) 一致
	373	+ if (rank_ > 8) {
	374	+ OP_LOGE(ctx_->GetNodeName(), "rank(%ld) exceeds max supported dim 8", rank_);
	375	+ return ge::GRAPH_FAILED;
	376	+ }
	377	+
	378	+ OP_LOGI(ctx_->GetNodeName(), "GetShapeInfo done rank %ld dtype %ld ub %lu core %lu has_bias %ld",
	379	+ rank_, dtype_size_, ubSize_, coreNum_, has_bias_);
	380	+
	381	+ return ge::GRAPH_SUCCESS;
	382	+ }
	383	+
	384	+ // 各 slot 填 TilingData：前补 shape=1/stride=0，归一化值右移；未用 slot 全填 1/0
	385	+ template<int64_t R, typename ShapeArr, typename StrideArr>
	386	+ static void FillSlots(ShapeArr shapes, StrideArr strides,
	387	+ const std::vector<std::vector<int64_t>>& norm,
	388	+ const std::vector<std::vector<int64_t>>& norm_strides,
	389	+ int64_t num, int64_t max_slots, int64_t delta)
	390	+ {
	391	+ int64_t rank = R - delta;
	392	+ for (int64_t i = 0; i < num; i++) {
	393	+ for (int64_t d = 0; d < delta; d++) { shapes[i][d] = 1; strides[i][d] = 0; }
	394	+ for (int64_t d = 0; d < rank; d++) {
	395	+ shapes[i][d + delta] = norm[i][d];
	396	+ strides[i][d + delta] = norm_strides[i][d];
	397	+ }
	398	+ }
	399	+ for (int64_t i = num; i < max_slots; i++)
	400	+ for (int64_t d = 0; d < R; d++) { shapes[i][d] = 1; strides[i][d] = 0; }
	401	+ }
	402	+
	403	+ template<int64_t R>
	404	+ ge::graphStatus DoTilingAndSet()
	405	+ {
	406	+ auto* tiling = ctx_->GetTilingData<ScaleTilingData<R>>();
	407	+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, tiling);
	408	+
	409	+ int64_t ub_per_core = (int64_t)ubSize_;
	410	+ int64_t per_buf_bytes = (ub_per_core / kPhysNodes) & ~31LL;
	411	+
	412	+ scale::FindSplitAxis(max_bro_shape_, dtype_size_, ub_per_core, kPhysNodes, tiling->split);
	413	+ scale::MultiCoreSplit(max_bro_shape_, tiling->split, (int64_t)coreNum_, tiling->multicore);
	414	+ tiling->per_buf_bytes = per_buf_bytes;
	415	+ tiling->per_buf_elems = per_buf_bytes / 4; // 统一按 FP32
	416	+
	417	+ int64_t num_in = (int64_t)normal_input_shapes_.size();
	418	+ int64_t num_out = (int64_t)normal_output_shapes_.size();
	419	+ std::vector<std::vector<int64_t>> in_strides(num_in), out_strides(num_out);
	420	+ for (int64_t i = 0; i < num_in; i++) scale::PrecomputeInputStrides(normal_input_shapes_[i], in_strides[i]);
	421	+ for (int64_t i = 0; i < num_out; i++) scale::PrecomputeOutputStrides(normal_output_shapes_[i], out_strides[i]);
	422	+
	423	+ tiling->rank = rank_;
	424	+ tiling->has_bias = has_bias_;
	425	+ int64_t delta = R - rank_; // 前补维数
	426	+
	427	+ // max_bro_shape: 前补 1，实际值右移
	428	+ for (int64_t d = 0; d < delta; d++) tiling->max_bro_shape[d] = 1;
	429	+ for (int64_t d = 0; d < rank_; d++) tiling->max_bro_shape[d + delta] = max_bro_shape_[d];
	430	+
	431	+ // split axis 右平移
	432	+ tiling->split.axis += delta;
	433	+
	434	+ tiling->num_inputs = num_in;
	435	+ tiling->num_outputs = num_out;
	436	+
	437	+ FillSlots<R>(tiling->input_shapes, tiling->input_strides, normal_input_shapes_, in_strides,
	438	+ num_in, kMaxInputSlots, delta);
	439	+ FillSlots<R>(tiling->output_shapes, tiling->output_strides, normal_output_shapes_, out_strides,
	440	+ num_out, kMaxOutputSlots, delta);
	441	+
	442	+ ctx_->SetBlockDim(tiling->multicore.num_cores);
	443	+ LogTilingData<R>(tiling, num_in, num_out);
	444	+ return ge::GRAPH_SUCCESS;
	445	+ }
	446	+
	447	+ template<int64_t R>
	448	+ void LogTilingData(ScaleTilingData<R>* tiling, int64_t num_in, int64_t num_out)
	449	+ {
	450	+ OP_LOGI(ctx_->GetNodeName(), "TilingData: per_buf_bytes=%ld rank=%ld->R=%d "
	451	+ "max_bro_shape=%s "
	452	+ "split(axis=%ld a_i=%ld a_o=%ld a_i_tail=%ld) "
	453	+ "multi(cores=%ld tiles=%ld main=%ld core_tail=%ld) num_in=%ld num_out=%ld has_bias=%ld",
	454	+ tiling->per_buf_bytes, rank_, (int)R,
	455	+ Arr2String(tiling->max_bro_shape, R).c_str(),
	456	+ tiling->split.axis, tiling->split.a_i, tiling->split.a_o, tiling->split.a_i_tail,
	457	+ tiling->multicore.num_cores, tiling->multicore.total_tiles,
	458	+ tiling->multicore.tiles_main, tiling->multicore.cores_tail,
	459	+ num_in, num_out, has_bias_);
	460	+ }
	461	+
	462	+ gert::TilingContext* ctx_;
	463	+ std::vector<std::vector<int64_t>> raw_input_shapes_;
	464	+ std::vector<std::vector<int64_t>> raw_output_shapes_;
	465	+ std::vector<int64_t> max_bro_shape_;
	466	+ std::vector<std::vector<int64_t>> normal_input_shapes_;
	467	+ std::vector<std::vector<int64_t>> normal_output_shapes_;
	468	+ int64_t dtype_size_ = 0;
	469	+ int64_t rank_ = 0;
	470	+ int64_t has_bias_ = 0;
	471	+ int64_t axis_ = 1;
	472	+ int64_t num_axes_ = 1;
	473	+ bool scale_from_blob_ = true;
	474	+ uint64_t coreNum_ = 0;
	475	+ uint64_t ubSize_ = 0;
	476	+};
	477	+
	478	+static ge::graphStatus TilingFuncScale(gert::TilingContext* context)
	479	+{
	480	+ ScaleTiling scaleTiling(context);
	481	+ auto ret = scaleTiling.RunTiling();
	482	+ if (ret != ge::GRAPH_SUCCESS) return ret;
	483	+ size_t* workspaces = context->GetWorkspaceSizes(1);
	484	+ workspaces[0] = 0; // Scale 无 workspace 需求
	485	+ return ge::GRAPH_SUCCESS;
	486	+}
	487	+
	488	+ge::graphStatus TilingPrepareForScale(gert::TilingParseContext* context)
	489	+{
	490	+ fe::PlatFormInfos* platformInfo = context->GetPlatformInfo();
	491	+ auto compileInfo = context->GetCompiledInfo<ScaleCompileInfo>();
	492	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
	493	+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
	494	+ auto ap = platform_ascendc::PlatformAscendC(platformInfo);
	495	+ compileInfo->coreNum = ap.GetCoreNumAiv();
	496	+ ap.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfo->ubSize);
	497	+ return ge::GRAPH_SUCCESS;
	498	+}
	499	+
	500	+IMPL_OP_OPTILING(Scale).Tiling(TilingFuncScale).TilingParse<ScaleCompileInfo>(TilingPrepareForScale);
	501	+
	502	+} // namespace optiling

Amath/scale/op_host/arch35/scale_tiling_arch35.h+25-0

@@ -0,0 +1,25 @@
	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 OPS_MATH_SCALE_OP_HOST_ARCH35_SCALE_TILING_ARCH35_H
	12	+#define OPS_MATH_SCALE_OP_HOST_ARCH35_SCALE_TILING_ARCH35_H
	13	+
	14	+#include "../../op_kernel/arch35/scale_tiling_struct.h"
	15	+
	16	+namespace optiling {
	17	+
	18	+struct ScaleCompileInfo {
	19	+ uint64_t coreNum;
	20	+ uint64_t ubSize;
	21	+};
	22	+
	23	+} // namespace optiling
	24	+
	25	+#endif // OPS_MATH_SCALE_OP_HOST_ARCH35_SCALE_TILING_ARCH35_H

Amath/scale/op_host/scale_def.cpp+60-0

@@ -0,0 +1,60 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+#include "register/op_def_registry.h"
	13	+
	14	+namespace ops {
	15	+class Scale : public OpDef {
	16	+public:
	17	+ explicit Scale(const char* name) : OpDef(name)
	18	+ {
	19	+ this->Input("x")
	20	+ .ParamType(REQUIRED)
	21	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
	22	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	23	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	24	+ .AutoContiguous();
	25	+ this->Input("scale")
	26	+ .ParamType(REQUIRED)
	27	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
	28	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	29	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	30	+ .AutoContiguous();
	31	+ this->Input("bias")
	32	+ .ParamType(OPTIONAL)
	33	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
	34	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	35	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	36	+ .AutoContiguous();
	37	+ this->Output("y")
	38	+ .ParamType(REQUIRED)
	39	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
	40	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	41	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	42	+ .AutoContiguous();
	43	+
	44	+ this->Attr("axis").Int(1L);
	45	+ this->Attr("num_axes").Int(1L);
	46	+ this->Attr("scale_from_blob").Bool(true);
	47	+
	48	+ OpAICoreConfig aiCoreConfig;
	49	+ aiCoreConfig.DynamicCompileStaticFlag(true)
	50	+ .DynamicFormatFlag(false)
	51	+ .DynamicRankSupportFlag(true)
	52	+ .DynamicShapeSupportFlag(true)
	53	+ .NeedCheckSupportFlag(false)
	54	+ .PrecisionReduceFlag(true)
	55	+ .ExtendCfgInfo("opFile.value", "scale_apt");
	56	+ this->AICore().AddConfig("ascend950", aiCoreConfig);
	57	+ }
	58	+};
	59	+OP_ADD(Scale);
	60	+} // namespace ops

Amath/scale/op_host/scale_infershape.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	+/* Generated By CANNBot */
	12	+/**
	13	+ * Scale 算子形状推导实现
	14	+ * y.shape = x.shape（输出形状与输入 x 一致）
	15	+ */
	16	+#include "register/op_impl_registry.h"
	17	+#include "exe_graph/runtime/infer_shape_context.h"
	18	+#include "op_common/log/log.h"
	19	+
	20	+using namespace ge;
	21	+
	22	+namespace ops {
	23	+
	24	+static ge::graphStatus InferShape4Scale(gert::InferShapeContext* context)
	25	+{
	26	+ const gert::Shape* input_shape = context->GetInputShape(0);
	27	+ OP_CHECK_NULL_WITH_CONTEXT(context, input_shape);
	28	+
	29	+ gert::Shape* output_shape = context->GetOutputShape(0);
	30	+ OP_CHECK_NULL_WITH_CONTEXT(context, output_shape);
	31	+
	32	+ // y.shape = x.shape
	33	+ output_shape = input_shape;
	34	+
	35	+ return ge::GRAPH_SUCCESS;
	36	+}
	37	+
	38	+IMPL_OP_INFERSHAPE(Scale).InferShape(InferShape4Scale);
	39	+
	40	+} // namespace ops

Amath/scale/op_kernel/arch35/scale_kernel.h+395-0

@@ -0,0 +1,395 @@
	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	+/* Generated By CANNBot */
	12	+
	13	+/**
	14	+ * Scale Kernel — ScaleKernel<T, RANK>
	15	+ * RegBase Broadcast 范式，基于 DESIGN.md §3.5
	16	+ * 目标架构: DAV_3510 (arch35)
	17	+ */
	18	+#pragma once
	19	+#include "kernel_operator.h"
	20	+#include "scale_tiling_struct.h"
	21	+#include "scale_struct.h"
	22	+
	23	+// ============================================================
	24	+// VF 函数前置声明 — ScaleMulAddVF: dst = src0 * src1 + dst
	25	+// (DESIGN.md §3.6, Reg::MulAddDst 包装；FP32 直算, FP16/BF16 以 float 实例化)
	26	+// ============================================================
	27	+template <typename T>
	28	+__simd_vf__ inline void ScaleMulAddVF(
	29	+ __ubuf__ T* dstAddr, __ubuf__ T* src0Addr, __ubuf__ T* src1Addr,
	30	+ uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes);
	31	+
	32	+// ============================================================
	33	+// Kernel 侧辅助函数 (int64_t* 版本, 无 std::vector)
	34	+// ============================================================
	35	+
	36	+__aicore__ inline void GetCoreRange(int64_t core_id, int64_t tiles_main, int64_t cores_tail,
	37	+ int64_t& start, int64_t& end)
	38	+{
	39	+ if (core_id < cores_tail) {
	40	+ start = core_id * (tiles_main + 1);
	41	+ end = start + tiles_main + 1;
	42	+ } else {
	43	+ start = cores_tail * (tiles_main + 1) + (core_id - cores_tail) * tiles_main;
	44	+ end = start + tiles_main;
	45	+ }
	46	+}
	47	+
	48	+__aicore__ inline int64_t GetUBSplitRange(
	49	+ int64_t a_o_off, int64_t a_o, int64_t a_i, int64_t a_i_tail)
	50	+{
	51	+ return (a_o_off == a_o - 1) ? a_i_tail : a_i;
	52	+}
	53	+
	54	+__aicore__ inline bool FlatToEffectiveCoord(int64_t flat, const int64_t* max_bro_shape,
	陈陈陈佳良6月24日问题: 坐标转换逻辑需验证。当 split_axis 不是最后一维时，坐标转换逻辑（从 split_axis-1 向下遍历）是否正确？建议增加 UT 测试验证多轴切分场景。建议: 增加多轴切分的 UT 测试用例，验证坐标转换的正确性 W wangweidong 6月25日评论： 6月25日评论：
	55	+ int64_t rank, int64_t split_axis, int64_t a_i, int64_t a_o, int64_t* eff_coord)
	56	+{
	57	+ for (int64_t d = 0; d < rank; d++)
	58	+ eff_coord[d] = 0;
	59	+ if (a_o <= 0) return false; // 除 0 保护：a_o 由 Tiling 保证 ≥1，异常时安全退出
	陈陈陈佳良6月24日问题: 边界条件保护依赖外部保证。代码注释提到 a_o 由 Tiling 保证 ≥1，但 Kernel 中仍添加了保护。需要验证 Tiling 代码是否真的保证了这个条件，否则保护可能失效。建议: 在 Tiling 文档或代码注释中明确说明保证条件，或增加更多防御性检查 W wangweidong 6月25日评论： 6月25日评论：
	60	+ int64_t a_o_off = flat % a_o;
	61	+ int64_t outer = flat / a_o;
	62	+ for (int64_t d = split_axis - 1; d >= 0; d--) {
	63	+ eff_coord[d] = outer % max_bro_shape[d];
	64	+ outer /= max_bro_shape[d];
	65	+ }
	66	+ eff_coord[split_axis] = a_o_off * a_i;
	67	+ return true;
	68	+}
	69	+
	70	+__aicore__ inline int64_t CalcInputOffset(
	71	+ const int64_t* eff_coord, const int64_t* strides, int64_t rank)
	72	+{
	73	+ int64_t offset = 0;
	74	+ for (int64_t d = 0; d < rank; d++)
	75	+ offset += eff_coord[d] * strides[d];
	76	+ return offset;
	77	+}
	78	+
	79	+__aicore__ inline int64_t CalcOutputOffset(
	80	+ const int64_t* eff_coord, const int64_t* strides, int64_t rank)
	81	+{
	82	+ int64_t offset = 0;
	83	+ for (int64_t d = 0; d < rank; d++)
	84	+ offset += eff_coord[d] * strides[d];
	85	+ return offset;
	86	+}
	87	+
	88	+// ============================================================
	89	+// ScaleKernel<T, RANK> — 模板化 Kernel 类
	90	+// ============================================================
	91	+
	92	+template <typename T, int64_t RANK>
	93	+class ScaleKernel {
	94	+ // NDDMA 维度数 (最大5), RANK>5 时外层走 Flat loop
	95	+ static constexpr int64_t ND = (RANK <= 5) ? RANK : 5;
	96	+ static constexpr bool NEED_CAST = !std::is_same_v<T, float>;
	97	+ // VF 以 float 实例化（FP32 直算, FP16/BF16 Cast 后均为 FP32）
	98	+ static constexpr uint32_t VL_F32 = AscendC::GetVecLen() / sizeof(float);
	99	+
	100	+ AscendC::TPipe pipe_;
	101	+ const ScaleTilingData<RANK>* td_;
	102	+ AscendC::GlobalTensor<T> gmIn_[kMaxInputSlots];
	103	+ AscendC::GlobalTensor<T> gmOut_[kMaxOutputSlots];
	104	+ AscendC::TBuf<AscendC::TPosition::VECCALC> buf_[kPhysNodes];
	105	+ AscendC::MultiCopyParams<T, ND> nddmaParams_[kMaxInputSlots];
	106	+ int64_t nddmaOuterIters_[kMaxInputSlots];
	107	+ int64_t nddma_dims_;
	108	+
	109	+public:
	110	+ __aicore__ inline void Init(GM_ADDR inputs[kMaxInputSlots], GM_ADDR outputs[kMaxOutputSlots],
	111	+ const ScaleTilingData<RANK>* td)
	112	+ {
	113	+ td_ = td;
	114	+ for (int i = 0; i < kMaxInputSlots; i++)
	115	+ gmIn_[i].SetGlobalBuffer((__gm__ T*)inputs[i]);
	116	+ for (int i = 0; i < kMaxOutputSlots; i++)
	117	+ gmOut_[i].SetGlobalBuffer((__gm__ T*)outputs[i]);
	118	+ for (int i = 0; i < kPhysNodes; i++)
	119	+ pipe_.InitBuffer(buf_[i], td_->per_buf_bytes);
	120	+
	121	+ // NDDMA 参数预计算
	122	+ const int64_t* dstShape = td_->max_bro_shape;
	123	+ int64_t k = td_->split.axis;
	124	+ nddma_dims_ = (RANK - k <= ND) ? (RANK - k) : ND;
	125	+ for (int inp = 0; inp < kMaxInputSlots; inp++) {
	126	+ int64_t inner = 1;
	127	+ int64_t nd = 0;
	128	+ for (int64_t d = RANK - 1; d >= k && nd < ND; d--) {
	129	+ nddmaParams_[inp].loopInfo.loopSize[nd] = (d == k) ? 0 : dstShape[d];
	130	+ nddmaParams_[inp].loopInfo.loopSrcStride[nd] = td_->input_strides[inp][d];
	131	+ nddmaParams_[inp].loopInfo.loopDstStride[nd] = inner;
	132	+ nddmaParams_[inp].loopInfo.loopLpSize[nd] = 0;
	133	+ nddmaParams_[inp].loopInfo.loopRpSize[nd] = 0;
	134	+ inner *= (d == k) ? td_->split.a_i : dstShape[d];
	135	+ nd++;
	136	+ }
	137	+ for (; nd < ND; nd++) {
	138	+ nddmaParams_[inp].loopInfo.loopSize[nd] = 1;
	139	+ nddmaParams_[inp].loopInfo.loopSrcStride[nd] = 0;
	140	+ nddmaParams_[inp].loopInfo.loopDstStride[nd] = inner;
	141	+ nddmaParams_[inp].loopInfo.loopLpSize[nd] = 0;
	142	+ nddmaParams_[inp].loopInfo.loopRpSize[nd] = 0;
	143	+ }
	144	+ nddmaOuterIters_[inp] = 1;
	145	+ for (int64_t d = k; d < RANK - nddma_dims_; d++)
	146	+ nddmaOuterIters_[inp] *= (d == k) ? td_->split.a_i : dstShape[d];
	147	+ }
	148	+ }
	149	+
	150	+ __aicore__ inline void Process()
	151	+ {
	152	+ if constexpr (NEED_CAST) {
	153	+ ProcessWithCast();
	154	+ } else {
	155	+ ProcessFP32();
	156	+ }
	157	+ }
	158	+
	159	+private:
	160	+ struct PipeEvents { int32_t mte2toV; int32_t vtoMte2; int32_t vtoMte3; int32_t mte3toMte2; };
	161	+
	162	+ __aicore__ inline PipeEvents FetchEvents()
	163	+ {
	164	+ return {
	165	+ static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_V)),
	166	+ static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_MTE2)),
	167	+ static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_MTE3)),
	168	+ static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE3_MTE2)),
	169	+ };
	170	+ }
	171	+
	172	+ __aicore__ inline int64_t ComputeInnerCount()
	173	+ {
	174	+ int64_t inner_count = 1;
	175	+ for (int64_t d = td_->split.axis + 1; d < RANK; d++) inner_count *= td_->max_bro_shape[d];
	176	+ return inner_count;
	177	+ }
	178	+
	179	+ // ============================================================
	180	+ // FP32 Process — 直接计算，无 Cast
	181	+ // ============================================================
	182	+ __aicore__ inline void ProcessFP32()
	183	+ {
	184	+ PipeEvents ev = FetchEvents();
	185	+ int64_t start, end;
	186	+ GetCoreRange(AscendC::GetBlockIdx(), td_->multicore.tiles_main,
	187	+ td_->multicore.cores_tail, start, end);
	188	+ int64_t inner_count = ComputeInnerCount();
	189	+ int64_t coord[8] = {};
	190	+ for (int64_t flat = start; flat < end; flat++) {
	191	+ int64_t a_i_seg = GetUBSplitRange(flat % td_->split.a_o, td_->split.a_o,
	192	+ td_->split.a_i, td_->split.a_i_tail);
	193	+ int64_t count = a_i_seg * inner_count;
	194	+ FlatToEffectiveCoord(flat, td_->max_bro_shape, RANK,
	195	+ td_->split.axis, td_->split.a_i, td_->split.a_o, coord);
	196	+ // 上轮 CopyOut(MTE3) 结束 → 本轮 CopyIn(MTE2) 可以开始
	197	+ if (flat != start) AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(ev.mte3toMte2);
	198	+ TileFP32(coord, count, a_i_seg, ev, flat != end - 1);
	199	+ }
	200	+ }
	201	+
	202	+ __aicore__ inline void TileFP32(const int64_t* coord, int64_t count, int64_t a_i_seg,
	203	+ const PipeEvents& ev, bool not_last)
	204	+ {
	205	+ constexpr int B0 = 0, B1 = 1, B2 = 2, B3 = 3;
	206	+ constexpr int IN_X = 0, IN_SCALE = 1, IN_BIAS = 2, OUT_Y = 0;
	207	+ // S1a: CopyInBrc x → B0
	208	+ CopyInBrc(coord, IN_X, B0, a_i_seg);
	209	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	210	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	211	+ // S1b: CopyInBrc scale → B1
	212	+ CopyInBrc(coord, IN_SCALE, B1, a_i_seg);
	213	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	214	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	215	+ int outBuf = B3;
	216	+ if (td_->has_bias) {
	217	+ // S1c: CopyInBrc bias → B2；S1d: MulAddDst(B2 ← B0·B1 + B2) in-place, 省 B3 中间读写
	218	+ CopyInBrc(coord, IN_BIAS, B2, a_i_seg);
	219	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	220	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	221	+ uint16_t rep = AscendC::CeilDivision(count, VL_F32);
	222	+ asc_vf_call<ScaleMulAddVF<float>>(
	223	+ (__ubuf__ float*)buf_[B2].Get<float>().GetPhyAddr(),
	陈陈陈佳良6月24日问题: Buffer 地址类型转换不安全。使用 reinterpret_cast 或直接强制转换 UB buffer 物理地址为 ubuf 指针，这是不安全的做法，可能导致内存访问错误或未定义行为。VF 函数需要 ubuf 类型指针，但 Ascend C API 提供的是 LocalTensor 的物理地址。建议: 使用 Ascend C 提供的安全 API：buf_[B2].Get().reinterpret_cast() 或查看官方 VF 示例中推荐的获取 ubuf 指针的方法 W wangweidong 6月25日评论： 6月25日评论：
	224	+ (__ubuf__ float*)buf_[B0].Get<float>().GetPhyAddr(),
	225	+ (__ubuf__ float*)buf_[B1].Get<float>().GetPhyAddr(),
	226	+ count, VL_F32, rep);
	227	+ outBuf = B2;
	228	+ } else {
	229	+ // S1c: Mul(B0, B1 → B3) 峰值 P_FP32=3
	230	+ AscendC::Mul(buf_[B3].Get<float>(), buf_[B0].Get<float>(), buf_[B1].Get<float>(), count);
	231	+ }
	232	+ // V→MTE3: V 写完 → MTE3 可搬；CopyOut → GM
	233	+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(ev.vtoMte3);
	234	+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(ev.vtoMte3);
	235	+ CopyOutOne(coord, OUT_Y, outBuf, a_i_seg);
	236	+ if (not_last) AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(ev.mte3toMte2);
	237	+ }
	238	+
	239	+ // ============================================================
	240	+ // FP16/BF16 Process — Cast→FP32→计算→Cast 回原始 dtype
	241	+ // ============================================================
	242	+ __aicore__ inline void ProcessWithCast()
	243	+ {
	244	+ PipeEvents ev = FetchEvents();
	245	+ int64_t start, end;
	246	+ GetCoreRange(AscendC::GetBlockIdx(), td_->multicore.tiles_main,
	247	+ td_->multicore.cores_tail, start, end);
	248	+ int64_t inner_count = ComputeInnerCount();
	249	+ int64_t coord[8] = {};
	250	+ for (int64_t flat = start; flat < end; flat++) {
	251	+ int64_t a_i_seg = GetUBSplitRange(flat % td_->split.a_o, td_->split.a_o,
	252	+ td_->split.a_i, td_->split.a_i_tail);
	253	+ int64_t count = a_i_seg * inner_count;
	254	+ FlatToEffectiveCoord(flat, td_->max_bro_shape, RANK,
	255	+ td_->split.axis, td_->split.a_i, td_->split.a_o, coord);
	256	+ if (flat != start) AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(ev.mte3toMte2);
	257	+ TileCast(coord, count, a_i_seg, ev, flat != end - 1);
	258	+ }
	259	+ }
	260	+
	261	+ __aicore__ inline void TileCast(const int64_t* coord, int64_t count, int64_t a_i_seg,
	262	+ const PipeEvents& ev, bool not_last)
	263	+ {
	264	+ constexpr int B0 = 0, B1 = 1, B2 = 2, B3 = 3;
	265	+ constexpr int IN_X = 0, IN_SCALE = 1, IN_BIAS = 2, OUT_Y = 0;
	266	+ // S1a: CopyInBrc x(T)→B0; S1b: Cast B0→B1 (x_fp32, CAST_NONE)
	267	+ CopyInBrc(coord, IN_X, B0, a_i_seg);
	268	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	269	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	270	+ AscendC::Cast(buf_[B1].template Get<float>(), buf_[B0].template Get<T>(),
	271	+ AscendC::RoundMode::CAST_NONE, count);
	272	+ // WAR: 等 Cast 读完 B0 再让 MTE2 覆写 B0(scale)；S1c CopyIn scale→B0; S1d Cast B0→B2
	273	+ AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(ev.vtoMte2);
	274	+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(ev.vtoMte2);
	275	+ CopyInBrc(coord, IN_SCALE, B0, a_i_seg);
	276	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	277	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	278	+ AscendC::Cast(buf_[B2].template Get<float>(), buf_[B0].template Get<T>(),
	279	+ AscendC::RoundMode::CAST_NONE, count);
	280	+ if (td_->has_bias) {
	281	+ // WAR 后 S2a CopyIn bias→B0; S2b Cast B0→B3(累加器); S2c MulAddDst(B3 ← B1·B2 + B3)
	282	+ AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(ev.vtoMte2);
	283	+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(ev.vtoMte2);
	284	+ CopyInBrc(coord, IN_BIAS, B0, a_i_seg);
	285	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	286	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(ev.mte2toV);
	287	+ AscendC::Cast(buf_[B3].template Get<float>(), buf_[B0].template Get<T>(),
	288	+ AscendC::RoundMode::CAST_NONE, count);
	289	+ uint16_t rep = AscendC::CeilDivision(count, VL_F32);
	290	+ asc_vf_call<ScaleMulAddVF<float>>(
	291	+ (__ubuf__ float*)buf_[B3].template Get<float>().GetPhyAddr(),
	292	+ (__ubuf__ float*)buf_[B1].template Get<float>().GetPhyAddr(),
	293	+ (__ubuf__ float*)buf_[B2].template Get<float>().GetPhyAddr(),
	294	+ count, VL_F32, rep);
	295	+ } else {
	296	+ // S1e: Mul(B1, B2 → B3) 峰值 P=3
	297	+ AscendC::Mul(buf_[B3].template Get<float>(), buf_[B1].template Get<float>(),
	298	+ buf_[B2].template Get<float>(), count);
	299	+ }
	300	+ // S3a: Cast(B3→B0) float→T (CAST_RINT)；V→MTE3 后 CopyOut B0→GM
	301	+ AscendC::Cast(buf_[B0].template Get<T>(), buf_[B3].template Get<float>(),
	302	+ AscendC::RoundMode::CAST_RINT, count);
	303	+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(ev.vtoMte3);
	304	+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(ev.vtoMte3);
	305	+ CopyOutOne(coord, OUT_Y, B0, a_i_seg);
	306	+ if (not_last) AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(ev.mte3toMte2);
	307	+ }
	308	+
	309	+ // ============================================================
	310	+ // CopyInBrc — NDDMA 多维搬运，利用 stride 实现随路 broadcast
	311	+ // ============================================================
	312	+ __aicore__ inline void CopyInBrc(
	313	+ const int64_t* coord, int inputIdx, int slot, int64_t a_i_seg)
	314	+ {
	315	+ int64_t k = td_->split.axis;
	316	+ int64_t off = CalcInputOffset(coord, td_->input_strides[inputIdx], RANK);
	317	+ const int64_t* dstShape = td_->max_bro_shape;
	318	+
	319	+ auto params = nddmaParams_[inputIdx];
	320	+ int64_t k_nd = RANK - 1 - k;
	321	+ int64_t inner = 1;
	322	+ for (int64_t nd = 0; nd < ND; nd++) {
	323	+ if (nd == k_nd) params.loopInfo.loopSize[nd] = a_i_seg;
	324	+ params.loopInfo.loopDstStride[nd] = inner;
	325	+ inner *= params.loopInfo.loopSize[nd];
	326	+ }
	327	+
	328	+ static constexpr AscendC::NdDmaConfig cfg = { false, AscendC::NdDmaConfig::unsetPad,
	329	+ AscendC::NdDmaConfig::unsetPad, false };
	330	+
	331	+ if constexpr (RANK <= 5) {
	332	+ AscendC::DataCopy<T, ND, cfg>(
	333	+ buf_[slot].template Get<T>(), gmIn_[inputIdx][off], params);
	334	+ } else {
	335	+ AscendC::LocalTensor<T> buf = buf_[slot].template Get<T>();
	336	+ int64_t elem_base = off;
	337	+ for (int64_t oi = 0; oi < nddmaOuterIters_[inputIdx]; oi++) {
	338	+ int64_t elem_adj = 0, tmp = oi;
	339	+ for (int64_t d = RANK - nddma_dims_ - 1; d >= k; d--) {
	340	+ int64_t sz = (d == k) ? a_i_seg : dstShape[d];
	341	+ elem_adj += (tmp % sz) * td_->input_strides[inputIdx][d];
	342	+ tmp /= sz;
	343	+ }
	344	+ AscendC::DataCopy<T, ND, cfg>(
	345	+ buf[oi * inner], gmIn_[inputIdx][elem_base + elem_adj], params);
	346	+ }
	347	+ }
	348	+ }
	349	+
	350	+ // ============================================================
	351	+ // CopyOutOne — DataCopyPad 将结果从 UB 写回 GM
	352	+ // ============================================================
	353	+ __aicore__ inline void CopyOutOne(
	354	+ const int64_t* coord, int outputIdx, int slot, int64_t a_i_seg)
	355	+ {
	356	+ int64_t off = CalcOutputOffset(coord, td_->output_strides[outputIdx], RANK);
	357	+ // 输出 transfer count: split 轴上 a_i_seg（y.shape = x.shape，无 broadcast）
	358	+ int64_t split_elems = a_i_seg;
	359	+ int64_t inner_elems = 1;
	360	+ for (int64_t d = td_->split.axis + 1; d < RANK; d++)
	361	+ inner_elems *= td_->output_shapes[outputIdx][d];
	362	+ int64_t cnt = split_elems * inner_elems;
	363	+
	364	+ AscendC::DataCopyExtParams extParams;
	365	+ extParams.blockCount = 1;
	366	+ extParams.blockLen = cnt * sizeof(T); // DataCopyPad 不要求 32B 对齐
	S S sunday6月24日 CopyOutOne 使用 DataCopyExtParams，当 T=half/bf16 且尾块元素数 cnt 较小时（如 cnt=1 时 blockLen=2字节），blockLen < 32字节违反 AscendC DMA 对 GM 写入的 32B 对齐要求。建议对尾块改用 isPad=true 的 DataCopyPad（配合 DataCopyPadParams 补齐至 32B），或在 tiling 中保证最小输出元素数 >= 16（half）/ >= 8（float）。
	367	+ extParams.srcStride = 0;
	368	+ extParams.dstStride = 0;
	369	+ AscendC::DataCopyPad(gmOut_[outputIdx][off], buf_[slot].template Get<T>(), extParams);
	370	+ }
	371	+};
	372	+
	373	+// ============================================================
	374	+// ScaleMulAddVF — VF 函数: dst = src0 * src1 + dst (Reg::MulAddDst 包装)
	375	+// DESIGN.md §3.6; FP32 直算 / FP16·BF16 Cast 后均以 float 实例化
	376	+// ============================================================
	377	+template <typename T>
	378	+__simd_vf__ inline void ScaleMulAddVF(
	Y Y yang-jiwei6月16日这个函数跟28行的vf啥关系，为啥有两个命名一样且参数一样的vf函数 W wangweidong 6月23日评论： 6月23日评论：
	379	+ __ubuf__ T* dstAddr, __ubuf__ T* src0Addr, __ubuf__ T* src1Addr,
	380	+ uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes)
	381	+{
	382	+ AscendC::Reg::RegTensor<T> srcReg0, srcReg1, dstReg;
	383	+ AscendC::Reg::MaskReg mask;
	384	+ AscendC::Reg::AddrReg aReg;
	385	+ for (uint16_t i = 0; i < repeatTimes; ++i) {
	386	+ aReg = AscendC::Reg::CreateAddrReg<T>(i, oneRepeatSize);
	387	+ uint32_t remain = count - (uint32_t)i * oneRepeatSize;
	388	+ mask = AscendC::Reg::UpdateMask<T>(remain);
	389	+ AscendC::Reg::LoadAlign(srcReg0, src0Addr, aReg);
	390	+ AscendC::Reg::LoadAlign(srcReg1, src1Addr, aReg);
	391	+ AscendC::Reg::LoadAlign(dstReg, dstAddr, aReg);
	392	+ AscendC::Reg::MulAddDst(dstReg, srcReg0, srcReg1, mask);
	393	+ AscendC::Reg::StoreAlign(dstAddr, dstReg, aReg, mask);
	394	+ }
	395	+}

Amath/scale/op_kernel/arch35/scale_struct.h+33-0

@@ -0,0 +1,33 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+// Scale TilingKey 模板参数
	14	+// 位置：operators/scale/op_kernel/arch35/scale_struct.h
	15	+#ifndef SCALE_STRUCT_H_
	16	+#define SCALE_STRUCT_H_
	17	+
	18	+#include "ascendc/host_api/tiling/template_argument.h"
	19	+
	20	+#define SCALE_RANK_4 4
	21	+#define SCALE_RANK_8 8
	22	+
	23	+ASCENDC_TPL_ARGS_DECL(Scale,
	24	+ ASCENDC_TPL_UINT_DECL(RANK, 8, ASCENDC_TPL_UI_LIST,
	25	+ SCALE_RANK_4, SCALE_RANK_8)
	26	+);
	27	+
	28	+ASCENDC_TPL_SEL(
	29	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, SCALE_RANK_4)),
	30	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, SCALE_RANK_8))
	31	+);
	32	+
	33	+#endif // SCALE_STRUCT_H_

Amath/scale/op_kernel/arch35/scale_tiling_struct.h+50-0

@@ -0,0 +1,50 @@
	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	+/* Generated By CANNBot */
	12	+// Scale TilingData — 按 rank 模板化，体积分两档
	13	+#pragma once
	14	+#include <cstdint>
	15	+
	16	+// === 算子特定常量 ===
	17	+constexpr int64_t kMaxInputSlots = 3; // x, scale, bias（可选）
	18	+constexpr int64_t kMaxOutputSlots = 1; // y
	19	+constexpr int64_t kPhysNodes = 4; // 物理存活节点 P（DESIGN.md §3.4 P trace 结论）
	20	+
	21	+struct SplitResult {
	22	+ int64_t axis;
	23	+ int64_t a_i;
	24	+ int64_t a_o;
	25	+ int64_t a_i_tail;
	26	+};
	27	+
	28	+struct MultiCoreResult {
	29	+ int64_t num_cores;
	30	+ int64_t total_tiles;
	31	+ int64_t tiles_main;
	32	+ int64_t cores_tail;
	33	+};
	34	+
	35	+template<int64_t kRank>
	36	+struct ScaleTilingData {
	37	+ SplitResult split;
	38	+ MultiCoreResult multicore;
	39	+ int64_t rank; // 实际 rank (1~8)，Kernel 运行期读取
	40	+ int64_t per_buf_bytes; // UB/P 向下对齐 32B，Kernel 用此初始化 TBuf
	41	+ int64_t per_buf_elems; // per_buf_bytes / 4（统一按 FP32 计算）
	42	+ int64_t max_bro_shape[kRank];
	43	+ int64_t num_inputs;
	44	+ int64_t num_outputs;
	45	+ int64_t has_bias; // 1=有 bias, 0=无 bias
	46	+ int64_t input_shapes [kMaxInputSlots][kRank];
	47	+ int64_t input_strides[kMaxInputSlots][kRank];
	48	+ int64_t output_shapes[kMaxOutputSlots][kRank];
	49	+ int64_t output_strides[kMaxOutputSlots][kRank];
	50	+};

Amath/scale/op_kernel/scale_apt.cpp+47-0

@@ -0,0 +1,47 @@
	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	+/* Generated By CANNBot */
	12	+/**
	13	+ * Scale Kernel 入口 — scale_apt.cpp
	14	+ * RANK 来自 TilingKey, DTYPE 来自 CANN 框架 (def 注册的 Input("x") 类型)
	15	+ * 目标架构: DAV_3510 (arch35 / ascend950)
	16	+ */
	17	+#include "kernel_operator.h"
	18	+#include "arch35/scale_kernel.h"
	19	+#include "arch35/scale_tiling_struct.h"
	20	+
	21	+using TilingData4 = ScaleTilingData<4>; // RANK≤4: 数组维度 [4]
	22	+using TilingData8 = ScaleTilingData<8>; // RANK>4: 数组维度 [8]
	23	+
	24	+template<int RANK>
	25	+__global__ __aicore__ void scale(
	26	+ GM_ADDR x, GM_ADDR scale_in, GM_ADDR bias,
	27	+ GM_ADDR y,
	28	+ GM_ADDR workspace, GM_ADDR tiling)
	29	+{
	30	+ GM_ADDR ins[3] = {x, scale_in, bias};
	31	+ GM_ADDR outs[1] = {y};
	32	+
	33	+ REGISTER_NONE_TILING;
	34	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	35	+
	36	+ if constexpr (RANK == 4) {
	37	+ GET_TILING_DATA_WITH_STRUCT(TilingData4, td, tiling);
	38	+ ScaleKernel<DTYPE_X, 4> kernel;
	39	+ kernel.Init(ins, outs, &td);
	40	+ kernel.Process();
	41	+ } else {
	42	+ GET_TILING_DATA_WITH_STRUCT(TilingData8, td, tiling);
	43	+ ScaleKernel<DTYPE_X, 8> kernel;
	44	+ kernel.Init(ins, outs, &td);
	45	+ kernel.Process();
	46	+ }
	47	+}

Mmath/scale/tests/ut/op_host/CMakeLists.txt+2-1

	@@ -16,5 +16,6 @@ foreach(SUB_DIR ${CURRENT_DIRS})
16	endif()	16	endif()
17	endforeach()	17	endforeach()
18		18
19	-if(OP_API_UT OR (UT_TEST_ALL AND NOT AICPU_ONLY))	19	+if(TILING_UT OR PROTO_UT OR OP_API_UT OR (UT_TEST_ALL AND NOT AICPU_ONLY))
		20	+ add_modules_llt_sources(HOSTNAME ${OPTEST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
20	endif()	21	endif()

Rmath/scale/op_host/CMakeLists.txt→math/scale/tests/ut/op_host/arch35/CMakeLists.txt+3-1

	@@ -9,4 +9,6 @@
9	# See LICENSE in the root of the software repository for the full text of the License.	9	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	10	# ----------------------------------------------------------------------------
11		11
12	-add_modules_sources(OPTYPE scale ACLNNTYPE aclnn_exclude)	12	+if(TILING_UT OR (UT_TEST_ALL AND NOT AICPU_ONLY))
		13	+ add_modules_llt_sources(HOSTNAME ${OPTEST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
		14	+endif()

Amath/scale/tests/ut/op_host/arch35/test_scale_tiling_arch35.cpp+279-0

@@ -0,0 +1,279 @@
	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 <gtest/gtest.h>
	14	+#include "tiling_context_faker.h"
	15	+#include "tiling_case_executor.h"
	16	+#include "../../../../op_host/arch35/scale_tiling_arch35.h"
	17	+
	18	+using namespace std;
	19	+
	20	+class ScaleTilingTest : public testing::Test {
	21	+ protected:
	22	+ static void SetUpTestCase() {
	23	+ std::cout << "ScaleTilingTest SetUp" << std::endl;
	24	+ }
	25	+
	26	+ static void TearDownTestCase() {
	27	+ std::cout << "ScaleTilingTest TearDown" << std::endl;
	28	+ }
	29	+};
	30	+
	31	+TEST_F(ScaleTilingTest, scale_tiling_float_with_bias)
	32	+{
	33	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	34	+ gert::TilingContextPara tilingContextPara(
	35	+ "Scale",
	36	+ {
	37	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	38	+ {{{5}, {5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	39	+ {{{5}, {5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	40	+ },
	41	+ {
	42	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	43	+ },
	44	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	45	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	46	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	47	+ {1, 1, 1}, {1},
	48	+ &compileInfo);
	49	+ TilingInfo tilingInfo;
	50	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	51	+ EXPECT_TRUE(result);
	52	+ EXPECT_EQ(tilingInfo.blockNum, 1);
	53	+}
	54	+
	55	+TEST_F(ScaleTilingTest, scale_tiling_float_no_bias)
	56	+{
	57	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	58	+ gert::TilingContextPara tilingContextPara(
	59	+ "Scale",
	60	+ {
	61	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	62	+ {{{5}, {5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	63	+ },
	64	+ {
	65	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	66	+ },
	67	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	68	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	69	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	70	+ {1, 1, 0}, {1},
	71	+ &compileInfo);
	72	+ TilingInfo tilingInfo;
	73	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	74	+ EXPECT_TRUE(result);
	75	+ EXPECT_EQ(tilingInfo.blockNum, 1);
	76	+}
	77	+
	78	+TEST_F(ScaleTilingTest, scale_tiling_float16_with_bias)
	79	+{
	80	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	81	+ gert::TilingContextPara tilingContextPara(
	82	+ "Scale",
	83	+ {
	84	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	85	+ {{{5}, {5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	86	+ {{{5}, {5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	87	+ },
	88	+ {
	89	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	90	+ },
	91	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	92	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	93	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	94	+ {1, 1, 1}, {1},
	95	+ &compileInfo);
	96	+ TilingInfo tilingInfo;
	97	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	98	+ EXPECT_TRUE(result);
	99	+}
	100	+
	101	+TEST_F(ScaleTilingTest, scale_tiling_bf16)
	102	+{
	103	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	104	+ gert::TilingContextPara tilingContextPara(
	105	+ "Scale",
	106	+ {
	107	+ {{{4, 8}, {4, 8}}, ge::DT_BF16, ge::FORMAT_ND},
	108	+ {{{8}, {8}}, ge::DT_BF16, ge::FORMAT_ND},
	109	+ {{{8}, {8}}, ge::DT_BF16, ge::FORMAT_ND},
	110	+ },
	111	+ {
	112	+ {{{4, 8}, {4, 8}}, ge::DT_BF16, ge::FORMAT_ND},
	113	+ },
	114	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	115	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	116	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	117	+ {1, 1, 1}, {1},
	118	+ &compileInfo);
	119	+ TilingInfo tilingInfo;
	120	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	121	+ EXPECT_TRUE(result);
	122	+}
	123	+
	124	+TEST_F(ScaleTilingTest, scale_tiling_scalar_scale)
	125	+{
	126	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	127	+ gert::TilingContextPara tilingContextPara(
	128	+ "Scale",
	129	+ {
	130	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	131	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	132	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	133	+ },
	134	+ {
	135	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	136	+ },
	137	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(0)),
	138	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(0)),
	139	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	140	+ {1, 1, 1}, {1},
	141	+ &compileInfo);
	142	+ TilingInfo tilingInfo;
	143	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	144	+ EXPECT_TRUE(result);
	145	+}
	146	+
	147	+TEST_F(ScaleTilingTest, scale_tiling_1d)
	148	+{
	149	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	150	+ gert::TilingContextPara tilingContextPara(
	151	+ "Scale",
	152	+ {
	153	+ {{{128}, {128}}, ge::DT_FLOAT, ge::FORMAT_ND},
	154	+ {{{128}, {128}}, ge::DT_FLOAT, ge::FORMAT_ND},
	155	+ },
	156	+ {
	157	+ {{{128}, {128}}, ge::DT_FLOAT, ge::FORMAT_ND},
	158	+ },
	159	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(0)),
	160	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	161	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	162	+ {1, 1, 0}, {1},
	163	+ &compileInfo);
	164	+ TilingInfo tilingInfo;
	165	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	166	+ EXPECT_TRUE(result);
	167	+}
	168	+
	169	+TEST_F(ScaleTilingTest, scale_tiling_4d)
	170	+{
	171	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	172	+ gert::TilingContextPara tilingContextPara(
	173	+ "Scale",
	174	+ {
	175	+ {{{2, 3, 4, 5}, {2, 3, 4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	176	+ {{{4, 5}, {4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	177	+ {{{4, 5}, {4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	178	+ },
	179	+ {
	180	+ {{{2, 3, 4, 5}, {2, 3, 4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	181	+ },
	182	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(2)),
	183	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(2)),
	184	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	185	+ {1, 1, 1}, {1},
	186	+ &compileInfo);
	187	+ TilingInfo tilingInfo;
	188	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	189	+ EXPECT_TRUE(result);
	190	+}
	191	+
	192	+TEST_F(ScaleTilingTest, scale_tiling_empty_tensor)
	193	+{
	194	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	195	+ gert::TilingContextPara tilingContextPara(
	196	+ "Scale",
	197	+ {
	198	+ {{{0, 5}, {0, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	199	+ {{{0, 5}, {0, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	200	+ },
	201	+ {
	202	+ {{{0, 5}, {0, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	203	+ },
	204	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(0)),
	205	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(2)),
	206	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	207	+ {1, 1, 0}, {1},
	208	+ &compileInfo);
	209	+ TilingInfo tilingInfo;
	210	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	211	+ EXPECT_TRUE(result);
	212	+}
	213	+
	214	+TEST_F(ScaleTilingTest, scale_tiling_unsupported_dtype)
	215	+{
	216	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	217	+ gert::TilingContextPara tilingContextPara(
	218	+ "Scale",
	219	+ {
	220	+ {{{3, 5}, {3, 5}}, ge::DT_INT32, ge::FORMAT_ND},
	221	+ {{{5}, {5}}, ge::DT_INT32, ge::FORMAT_ND},
	222	+ },
	223	+ {
	224	+ {{{3, 5}, {3, 5}}, ge::DT_INT32, ge::FORMAT_ND},
	225	+ },
	226	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	227	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	228	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	229	+ {1, 1, 0}, {1},
	230	+ &compileInfo);
	231	+ TilingInfo tilingInfo;
	232	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	233	+ EXPECT_FALSE(result);
	234	+}
	235	+
	236	+TEST_F(ScaleTilingTest, scale_tiling_scale_from_blob_false)
	237	+{
	238	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	239	+ gert::TilingContextPara tilingContextPara(
	240	+ "Scale",
	241	+ {
	242	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	243	+ {{{5}, {5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	244	+ {{{5}, {5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	245	+ },
	246	+ {
	247	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	248	+ },
	249	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	250	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	251	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(false))},
	252	+ {1, 1, 1}, {1},
	253	+ &compileInfo);
	254	+ TilingInfo tilingInfo;
	255	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	256	+ EXPECT_TRUE(result);
	257	+}
	258	+
	259	+TEST_F(ScaleTilingTest, scale_tiling_negative_axis)
	260	+{
	261	+ optiling::ScaleCompileInfo compileInfo = {64, 262144};
	262	+ gert::TilingContextPara tilingContextPara(
	263	+ "Scale",
	264	+ {
	265	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	266	+ {{{5}, {5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	267	+ },
	268	+ {
	269	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	270	+ },
	271	+ {gert::TilingContextPara::OpAttr("axis", Ops::Math::AnyValue::CreateFrom<int64_t>(-1)),
	272	+ gert::TilingContextPara::OpAttr("num_axes", Ops::Math::AnyValue::CreateFrom<int64_t>(1)),
	273	+ gert::TilingContextPara::OpAttr("scale_from_blob", Ops::Math::AnyValue::CreateFrom<bool>(true))},
	274	+ {1, 1, 0}, {1},
	275	+ &compileInfo);
	276	+ TilingInfo tilingInfo;
	277	+ bool result = ExecuteTiling(tilingContextPara, tilingInfo);
	278	+ EXPECT_TRUE(result);
	279	+}

Amath/scale/tests/ut/op_host/test_scale_infershape.cpp+113-0

@@ -0,0 +1,113 @@
	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 <gtest/gtest.h>
	12	+#include <iostream>
	13	+#include "infershape_context_faker.h"
	14	+#include "infershape_case_executor.h"
	15	+
	16	+class ScaleInferShapeTest : public testing::Test {
	17	+ protected:
	18	+ static void SetUpTestCase() {
	19	+ std::cout << "ScaleInferShapeTest SetUp" << std::endl;
	20	+ }
	21	+
	22	+ static void TearDownTestCase() {
	23	+ std::cout << "ScaleInferShapeTest TearDown" << std::endl;
	24	+ }
	25	+};
	26	+
	27	+TEST_F(ScaleInferShapeTest, infershape_float_2d)
	28	+{
	29	+ gert::InfershapeContextPara infershapeContextPara(
	30	+ "Scale",
	31	+ {
	32	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	33	+ {{{5}, {5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	34	+ {{{5}, {5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	35	+ },
	36	+ {
	37	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	38	+ });
	39	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	40	+ {3, 5},
	41	+ };
	42	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	43	+}
	44	+
	45	+TEST_F(ScaleInferShapeTest, infershape_float16_3d)
	46	+{
	47	+ gert::InfershapeContextPara infershapeContextPara(
	48	+ "Scale",
	49	+ {
	50	+ {{{2, 3, 4}, {2, 3, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	51	+ {{{4}, {4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	52	+ },
	53	+ {
	54	+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	55	+ });
	56	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	57	+ {2, 3, 4},
	58	+ };
	59	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	60	+}
	61	+
	62	+TEST_F(ScaleInferShapeTest, infershape_bf16_4d)
	63	+{
	64	+ gert::InfershapeContextPara infershapeContextPara(
	65	+ "Scale",
	66	+ {
	67	+ {{{2, 3, 4, 5}, {2, 3, 4, 5}}, ge::DT_BF16, ge::FORMAT_ND},
	68	+ {{{4, 5}, {4, 5}}, ge::DT_BF16, ge::FORMAT_ND},
	69	+ {{{4, 5}, {4, 5}}, ge::DT_BF16, ge::FORMAT_ND},
	70	+ },
	71	+ {
	72	+ {{{}, {}}, ge::DT_BF16, ge::FORMAT_ND},
	73	+ });
	74	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	75	+ {2, 3, 4, 5},
	76	+ };
	77	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	78	+}
	79	+
	80	+TEST_F(ScaleInferShapeTest, infershape_1d)
	81	+{
	82	+ gert::InfershapeContextPara infershapeContextPara(
	83	+ "Scale",
	84	+ {
	85	+ {{{128}, {128}}, ge::DT_FLOAT, ge::FORMAT_ND},
	86	+ {{{128}, {128}}, ge::DT_FLOAT, ge::FORMAT_ND},
	87	+ },
	88	+ {
	89	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	90	+ });
	91	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	92	+ {128},
	93	+ };
	94	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	95	+}
	96	+
	97	+TEST_F(ScaleInferShapeTest, infershape_scalar_scale)
	98	+{
	99	+ gert::InfershapeContextPara infershapeContextPara(
	100	+ "Scale",
	101	+ {
	102	+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	103	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	104	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	105	+ },
	106	+ {
	107	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	108	+ });
	109	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	110	+ {3, 5},
	111	+ };
	112	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	113	+}

Amath/scale/tests/ut/op_kernel/CMakeLists.txt+17-0

@@ -0,0 +1,17 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify it.
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This file is a part of the CANN Open Software.
	5	+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
	8	+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+# See LICENSE in the root of the software repository for the full text of the License.
	10	+# ----------------------------------------------------------------------------
	11	+
	12	+if (UT_TEST_ALL OR OP_KERNEL_UT)
	13	+ set(scale_tiling_files
	14	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/scale_tiling_arch35.cpp
	15	+ )
	16	+ AddOpTestCase(scale "ascend950" "" "${scale_tiling_files}")
	17	+endif()

Amath/scale/tests/ut/op_kernel/test_scale.cpp+190-0

@@ -0,0 +1,190 @@
	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 <vector>
	12	+#include <iostream>
	13	+#include <cstdint>
	14	+#include <cstring>
	15	+#include "gtest/gtest.h"
	16	+#include "tikicpulib.h"
	17	+
	18	+#include "../../../op_kernel/arch35/scale_tiling_struct.h"
	19	+#include "../../../op_kernel/arch35/scale_kernel.h"
	20	+
	21	+#ifndef GET_TILING_DATA_WITH_STRUCT
	22	+#define REGISTER_TILINGDATA_SIZE(tiling_struct, counter)
	23	+#if defined(ASCENDC_CPU_DEBUG)
	24	+template <class T>
	25	+inline __aicore__ void InitTilingData(const __gm__ uint8_t p, T td)
	26	+{
	27	+ constexpr uint64_t sz = sizeof(T);
	28	+ constexpr uint32_t judge = sz > 15 ? sz - 15 : 0;
	29	+ uint32_t i = 0;
	30	+ if (judge > 0) {
	31	+ for (; i < judge; i += 16) {
	32	+ ((uint64_t)((uint8_t)td + i)) = ((const __gm__ uint64_t)((const __gm__ uint8_t )p + i));
	33	+ ((uint64_t)((uint8_t)td + i + 8)) = ((const __gm__ uint64_t)((const __gm__ uint8_t )p + i + 8));
	34	+ }
	35	+ }
	36	+ if (sz & 0x08) { ((uint64_t)((uint8_t)td + i)) = ((const __gm__ uint64_t)((const __gm__ uint8_t)p + i)); i += 8; }
	37	+ if (sz & 0x04) { ((uint32_t)((uint8_t)td + i)) = ((const __gm__ uint32_t)((const __gm__ uint8_t)p + i)); i += 4; }
	38	+ if (sz & 0x02) { ((uint16_t)((uint8_t)td + i)) = ((const __gm__ uint16_t)((const __gm__ uint8_t)p + i)); i += 2; }
	39	+ if (sz & 0x01) { ((uint8_t)((uint8_t)td + i)) = ((const __gm__ uint8_t)((const __gm__ uint8_t)p + i)); }
	40	+}
	41	+#endif
	42	+#define GET_TILING_DATA_WITH_STRUCT(tiling_struct, tiling_data, tiling_arg) \
	43	+ REGISTER_TILINGDATA_SIZE(tiling_struct, __COUNTER__); \
	44	+ tiling_struct tiling_data; \
	45	+ InitTilingData<tiling_struct>(tiling_arg, &tiling_data);
	46	+#endif
	47	+
	48	+void scale_float_rank4_no_bias(GM_ADDR x, GM_ADDR scale_in, GM_ADDR bias,
	49	+ GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
	50	+{
	51	+ GET_TILING_DATA_WITH_STRUCT(ScaleTilingData<4>, td, tiling);
	52	+ GM_ADDR ins[3] = {x, scale_in, bias};
	53	+ GM_ADDR outs[1] = {y};
	54	+ ScaleKernel<float, 4> kernel;
	55	+ kernel.Init(ins, outs, &td);
	56	+ kernel.Process();
	57	+}
	58	+
	59	+void scale_float_rank4_with_bias(GM_ADDR x, GM_ADDR scale_in, GM_ADDR bias,
	60	+ GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
	61	+{
	62	+ GET_TILING_DATA_WITH_STRUCT(ScaleTilingData<4>, td, tiling);
	63	+ GM_ADDR ins[3] = {x, scale_in, bias};
	64	+ GM_ADDR outs[1] = {y};
	65	+ ScaleKernel<float, 4> kernel;
	66	+ kernel.Init(ins, outs, &td);
	67	+ kernel.Process();
	68	+}
	69	+
	70	+class ScaleKernelTest : public testing::Test {
	71	+protected:
	72	+ static void SetUpTestCase()
	73	+ {
	74	+ std::cout << "ScaleKernelTest SetUp" << std::endl;
	75	+ }
	76	+ static void TearDownTestCase()
	77	+ {
	78	+ std::cout << "ScaleKernelTest TearDown" << std::endl;
	79	+ }
	80	+};
	81	+
	82	+static void FillScaleTilingData4(ScaleTilingData<4>* td, bool hasBias)
	83	+{
	84	+ memset(td, 0, sizeof(ScaleTilingData<4>));
	85	+ td->split = {2, 3, 1, 3};
	86	+ td->multicore = {1, 1, 1, 0};
	87	+ td->rank = 2;
	88	+ td->per_buf_bytes = 15744;
	89	+ td->per_buf_elems = 3936;
	90	+ td->max_bro_shape[0] = 1;
	91	+ td->max_bro_shape[1] = 1;
	92	+ td->max_bro_shape[2] = 3;
	93	+ td->max_bro_shape[3] = 5;
	94	+ td->num_inputs = hasBias ? 3 : 2;
	95	+ td->num_outputs = 1;
	96	+ td->has_bias = hasBias ? 1 : 0;
	97	+ td->input_shapes[0][0] = 1; td->input_shapes[0][1] = 1; td->input_shapes[0][2] = 3; td->input_shapes[0][3] = 5;
	98	+ td->input_strides[0][0] = 0; td->input_strides[0][1] = 0; td->input_strides[0][2] = 5; td->input_strides[0][3] = 1;
	99	+ td->input_shapes[1][0] = 1; td->input_shapes[1][1] = 1; td->input_shapes[1][2] = 1; td->input_shapes[1][3] = 5;
	100	+ td->input_strides[1][0] = 0; td->input_strides[1][1] = 0; td->input_strides[1][2] = 0; td->input_strides[1][3] = 1;
	101	+ if (hasBias) {
	102	+ td->input_shapes[2][0] = 1; td->input_shapes[2][1] = 1; td->input_shapes[2][2] = 1; td->input_shapes[2][3] = 5;
	103	+ td->input_strides[2][0] = 0; td->input_strides[2][1] = 0; td->input_strides[2][2] = 0; td->input_strides[2][3] = 1;
	104	+ }
	105	+ td->output_shapes[0][0] = 1; td->output_shapes[0][1] = 1; td->output_shapes[0][2] = 3; td->output_shapes[0][3] = 5;
	106	+ td->output_strides[0][0] = 0; td->output_strides[0][1] = 0; td->output_strides[0][2] = 5; td->output_strides[0][3] = 1;
	107	+}
	108	+
	109	+TEST_F(ScaleKernelTest, test_float_no_bias_rank4)
	陈陈陈佳良6月24日问题: Kernel UT 测试覆盖不完整。只有 rank4 + float 的 2 个测试用例，缺少关键场景测试：1) RANK=8 的测试；2) FP16/BF16 的 Cast 流程测试；3) 多核切分场景测试；4) 大数据量测试。测试覆盖不足可能导致边缘场景 bug 未被发现。建议: 补充测试用例：test_float16_with_bias_rank8、test_bf16_no_bias_rank4、test_multicore_split、test_large_data
	110	+{
	111	+ constexpr int64_t N = 15;
	112	+ constexpr size_t ELEM_SIZE = sizeof(float);
	113	+
	114	+ uint8_t* x = (uint8_t)AscendC::GmAlloc(N ELEM_SIZE);
	115	+ uint8_t* scale_in = (uint8_t)AscendC::GmAlloc(N ELEM_SIZE);
	116	+ uint8_t* bias = (uint8_t)AscendC::GmAlloc(N ELEM_SIZE);
	117	+ uint8_t* y = (uint8_t)AscendC::GmAlloc(N ELEM_SIZE);
	118	+
	119	+ float* xF = reinterpret_cast<float*>(x);
	120	+ float* scF = reinterpret_cast<float*>(scale_in);
	121	+ float* yF = reinterpret_cast<float*>(y);
	122	+ for (int i = 0; i < N; i++) xF[i] = static_cast<float>(i + 1);
	123	+ for (int i = 0; i < 5; i++) scF[i] = 2.0f;
	124	+ memset(bias, 0, N * ELEM_SIZE);
	125	+ memset(y, 0, N * ELEM_SIZE);
	126	+
	127	+ uint8_t* workspace = (uint8_t)AscendC::GmAlloc(16 1024 * 1024);
	128	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(ScaleTilingData<4>));
	129	+
	130	+ ScaleTilingData<4>* td = reinterpret_cast<ScaleTilingData<4>*>(tiling);
	131	+ FillScaleTilingData4(td, false);
	132	+
	133	+ ICPU_SET_TILING_KEY(4);
	134	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	135	+ ICPU_RUN_KF(scale_float_rank4_no_bias, 1, x, scale_in, bias, y, workspace, tiling);
	136	+
	137	+ for (int i = 0; i < N; i++) {
	138	+ float expected = static_cast<float>(i + 1) * 2.0f;
	139	+ EXPECT_FLOAT_EQ(yF[i], expected);
	140	+ }
	141	+
	142	+ AscendC::GmFree(x);
	143	+ AscendC::GmFree(scale_in);
	144	+ AscendC::GmFree(bias);
	145	+ AscendC::GmFree(y);
	146	+ AscendC::GmFree(workspace);
	147	+ AscendC::GmFree(tiling);
	148	+}
	149	+
	150	+TEST_F(ScaleKernelTest, test_float_with_bias_rank4)
	151	+{
	152	+ constexpr int64_t N = 15;
	153	+ constexpr size_t ELEM_SIZE = sizeof(float);
	154	+
	155	+ uint8_t* x = (uint8_t)AscendC::GmAlloc(N ELEM_SIZE);
	156	+ uint8_t* scale_in = (uint8_t)AscendC::GmAlloc(N ELEM_SIZE);
	157	+ uint8_t* bias = (uint8_t)AscendC::GmAlloc(N ELEM_SIZE);
	158	+ uint8_t* y = (uint8_t)AscendC::GmAlloc(N ELEM_SIZE);
	159	+
	160	+ float* xF = reinterpret_cast<float*>(x);
	161	+ float* scF = reinterpret_cast<float*>(scale_in);
	162	+ float* biF = reinterpret_cast<float*>(bias);
	163	+ float* yF = reinterpret_cast<float*>(y);
	164	+ for (int i = 0; i < N; i++) xF[i] = static_cast<float>(i + 1);
	165	+ for (int i = 0; i < 5; i++) scF[i] = 2.0f;
	166	+ for (int i = 0; i < 5; i++) biF[i] = 1.0f;
	167	+ memset(y, 0, N * ELEM_SIZE);
	168	+
	169	+ uint8_t* workspace = (uint8_t)AscendC::GmAlloc(16 1024 * 1024);
	170	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(ScaleTilingData<4>));
	171	+
	172	+ ScaleTilingData<4>* td = reinterpret_cast<ScaleTilingData<4>*>(tiling);
	173	+ FillScaleTilingData4(td, true);
	174	+
	175	+ ICPU_SET_TILING_KEY(4);
	176	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	177	+ ICPU_RUN_KF(scale_float_rank4_with_bias, 1, x, scale_in, bias, y, workspace, tiling);
	178	+
	179	+ for (int i = 0; i < N; i++) {
	180	+ float expected = static_cast<float>(i + 1) * 2.0f + 1.0f;
	181	+ EXPECT_FLOAT_EQ(yF[i], expected);
	182	+ }
	183	+
	184	+ AscendC::GmFree(x);
	185	+ AscendC::GmFree(scale_in);
	186	+ AscendC::GmFree(bias);
	187	+ AscendC::GmFree(y);
	188	+ AscendC::GmFree(workspace);
	189	+ AscendC::GmFree(tiling);
	190	+}

	@@ -4,7 +4,7 @@
4		4
5	\| 产品 \| 是否支持 \|	5	\| 产品 \| 是否支持 \|
6	\| :----------------------------------------------------------- \| :------: \|	6	\| :----------------------------------------------------------- \| :------: \|
7	-\| <term>Ascend 950PR/Ascend 950DT</term> \| × \|	7	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
8	\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|	8	\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|
9	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|	9	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|
10	\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|	10	\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|