[CANNBot]Applyadadelta算子Ascend950适配-ops-nn-AtomGit

Aexperimental/optim/apply_adadelta/CMakeLists.txt+14-0

@@ -0,0 +1,14 @@
	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	+
	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,
	W W wangyongguang4月20日补充PR描述和Issue信息
	8	+# 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.
	10	+# ----------------------------------------------------------------------------------------------------------
	11	+# NOTE: Portions of this code were AI-generated and have been
	12	+# technically reviewed for functional accuracy and security
	13	+
	14	+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE apply_adadelta ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)

Aexperimental/optim/apply_adadelta/README.md+145-0

@@ -0,0 +1,145 @@
	1	+# ApplyAdadelta
	2	+
	3	+## 算子说明
	4	+
	5	+执行 Adadelta 优化器的单步参数更新。根据当前梯度、梯度平方累积和更新量平方累积，计算参数更新量并原地更新权重参数（inplace 语义）。对标 TensorFlow `tf.raw_ops.ApplyAdadelta` 接口和 PyTorch `torch.optim.Adadelta` 的计算语义。
	6	+
	7	+### 计算公式
	8	+
	9	+$$
	10	+\begin{aligned}
	11	+accum_{t} &= \rho \cdot accum_{t-1} + (1 - \rho) \cdot grad^2 \\
	12	+update &= \frac{\sqrt{accum\_update_{t-1} + \epsilon}}{\sqrt{accum_{t} + \epsilon}} \cdot grad \\
	13	+var_{t} &= var_{t-1} - lr \cdot update \\
	14	+accum\_update_{t} &= \rho \cdot accum\_update_{t-1} + (1 - \rho) \cdot update^2
	15	+\end{aligned}
	16	+$$
	17	+
	chenfeng4月20日 op_api_list.md op_list.md 交付件缺少，以及缺少ut Nerddddddddddd 4月24日评论： 4月24日评论：
	18	+其中：
	19	+- `var`：待更新的权重参数
	20	+- `accum`：梯度平方的指数移动平均（$E[g^2]$）
	21	+- `accum_update`：参数更新量平方的指数移动平均（$E[\Delta x^2]$）
	22	+- `lr`：学习率
	23	+- `rho`：衰减系数，取值范围 [0, 1)
	24	+- `epsilon`：数值稳定常数，必须 > 0
	25	+- `grad`：当前梯度
	26	+
	27	+## 产品支持情况
	28	+
	29	+\| 产品 \| 是否支持 \|
	30	+\|------\|:------:\|
	31	+\| Ascend 950PR / Ascend 950DT（Atlas A5 推理系列，DAV_3510） \| √ \|
	32	+\| Atlas A3 训练系列 / Atlas A3 推理系列 \| × \|
	33	+\| Atlas A2 训练系列 / Atlas A2 推理系列 \| × \|
	34	+\| Atlas 200I/500 A2 推理产品 \| × \|
	35	+\| Atlas 推理系列 \| × \|
	36	+\| Atlas 训练系列 \| × \|
	37	+
	38	+## 调用方式支持
	39	+
	40	+\| 调用方式 \| 是否支持 \|
	41	+\|---------\|:--------:\|
	42	+\| ACLNN 单算子调用 \| √ \|
	43	+\| torch_npu 单算子调用 \| × \|
	44	+\| torch.compile 入图 \| × \|
	45	+\| GE 图模式（静态 shape） \| × \|
	46	+\| GE 图模式（动态 shape） \| × \|
	47	+
	48	+> 说明：本算子仅支持 ACLNN 单算子调用，暂不支持 GE 图模式、torch_npu 单算子及 torch.compile 入图。因此仅提供 ACLNN 调用示例 `test_aclnn_apply_adadelta.cpp`，不提供 GE IR 示例。
	49	+
	50	+## 数据类型
	51	+
	52	+\| 参数 \| 支持 dtype \|
	53	+\|-----\|-----------\|
	54	+\| var / accum / accumUpdate / grad \| FLOAT16、FLOAT \|
	55	+\| lr / rho / epsilon (aclScalar) \| FLOAT16、FLOAT（需与 Tensor dtype 一致） \|
	56	+
	57	+## 参数说明
	58	+
	59	+\| 参数名 \| 输入/输出 \| 说明 \|
	60	+\|-------\|----------\|------\|
	61	+\| var \| 输入+输出 (inplace) \| 权重参数 Tensor，dtype ∈ {FLOAT16, FLOAT}，1-8 维 ND 格式 \|
	62	+\| accum \| 输入+输出 (inplace) \| 梯度平方累积 Tensor，shape/dtype 必须与 var 一致 \|
	63	+\| accumUpdate \| 输入+输出 (inplace) \| 更新量平方累积 Tensor，shape/dtype 必须与 var 一致 \|
	64	+\| lr \| 输入 \| 学习率标量，dtype 必须与 Tensor 一致 \|
	65	+\| rho \| 输入 \| 衰减系数标量，取值范围 [0, 1) \|
	66	+\| epsilon \| 输入 \| 数值稳定常数标量，必须 > 0 \|
	67	+\| grad \| 输入 \| 梯度 Tensor，shape/dtype 必须与 var 一致 \|
	68	+\| varOut \| 输出 \| 与 var 共享 Device 内存 \|
	69	+\| accumOut \| 输出 \| 与 accum 共享 Device 内存 \|
	70	+\| accumUpdateOut \| 输出 \| 与 accumUpdate 共享 Device 内存 \|
	71	+
	72	+详细接口参数与返回码定义请参见 [aclnnApplyAdadelta 接口文档](docs/aclnnApplyAdadelta.md)。
	73	+
	74	+## 约束说明
	75	+
	76	+- 本算子仅支持 Ascend 950PR / Ascend 950DT 产品（Atlas A5 推理系列，DAV_3510 架构）。
	77	+- `var`、`accum`、`accumUpdate`、`grad` 的 shape 必须完全相同，不支持广播。
	78	+- `var`、`accum`、`accumUpdate`、`grad` 的 dtype 必须完全一致，`lr`、`rho`、`epsilon` 的 dtype 也需与之一致。
	79	+- 不支持 0 维标量 Tensor 作为 var / accum / accumUpdate / grad。
	80	+- 参数值域：`epsilon > 0`，`rho ∈ [0, 1)`。
	81	+- Inplace 语义：`varOut / accumOut / accumUpdateOut` 分别与 `var / accum / accumUpdate` 共享 GM 地址，输出直接写回输入内存。调用方需自行保存调用前的原始值（如需）。
	82	+- 空 Tensor 处理：元素数为 0 时直接返回成功，不执行计算。
	83	+- FP16 精度路径：Kernel 内部提升为 FP32 计算，输出时转回 FP16，以保证精度。
	84	+- 确定性：本算子为逐元素运算，无 Reduce，相同输入产出相同输出。
	85	+
	86	+## 目录结构
	87	+
	88	+```
	89	+apply_adadelta/
	90	+├── CMakeLists.txt
	91	+├── README.md
	92	+├── build.sh
	93	+├── docs/
	94	+│ ├── aclnnApplyAdadelta.md # aclnn 接口说明
	95	+│ ├── REQUIREMENTS.md # 需求文档
	96	+│ ├── DESIGN.md # 详细设计
	97	+│ ├── TEST.md # 测试设计
	98	+│ ├── precision-report.md # 最终精度验收报告
	99	+│ └── performance-report.md # 性能验收报告
	100	+├── examples/
	101	+│ ├── CMakeLists.txt
	102	+│ └── test_aclnn_apply_adadelta.cpp # aclnn 调用示例
	103	+├── op_host/
	104	+├── op_kernel/
	105	+├── op_api/
	106	+└── tests/
	107	+ ├── ut/
	108	+ └── st/
	109	+```
	110	+
	111	+## 调用示例
	112	+
	113	+本算子提供 ACLNN 调用示例，样例代码位于 [`examples/test_aclnn_apply_adadelta.cpp`](examples/test_aclnn_apply_adadelta.cpp)。
	114	+
	115	+### 编译运行
	116	+
	117	+前置条件：已安装 CANN Toolkit，并已编译安装本算子自定义包（生成 `opp/vendors/apply_adadelta_custom/`）。
	118	+
	119	+```bash
	120	+# 1. 设置环境
	121	+source ${ASCEND_HOME_PATH}/bin/setenv.bash
	122	+
	123	+# 2. 编译示例
	124	+cd examples
	125	+mkdir -p build && cd build
	126	+cmake ..
	127	+make
	128	+
	129	+# 3. 运行示例
	130	+./test_aclnn_apply_adadelta
	131	+```
	132	+
	133	+### 示例输出
	134	+
	135	+示例以 shape=[4,2] 的 fp32 张量、`lr=0.01, rho=0.9, epsilon=1e-6` 执行一次 Adadelta 更新，典型输出：
	136	+
	137	+```
	138	+var_out[0]=0.998526 accum_out[0]=0.115000 accum_update_out[0]=0.011174
	139	+var_out[1]=2.000976 accum_out[1]=0.189000 accum_update_out[1]=0.018952
	140	+...
	141	+```
	142	+
	143	+> 不支持 GE 图模式：如前述调用方式支持矩阵所述，本算子不支持 GE 图模式，因此不提供 `test_geir_apply_adadelta.cpp` 示例。
	144	+
	145	+具体编译和运行过程可参考 [编译与运行样例](../../../docs/context/编译与运行样例.md)。

Aexperimental/optim/apply_adadelta/examples/arch35/test_aclnn_apply_adadelta.cpp+218-0

@@ -0,0 +1,218 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ *
	15	+ * @file test_aclnn_apply_adadelta.cpp
	16	+ * @brief aclnnApplyAdadelta 算子调用示例。
	17	+ *
	18	+ * 功能：Adadelta 优化器单步参数更新（inplace）
	19	+ * 公式：
	20	+ * accum_new = rho * accum + (1 - rho) * grad^2
	21	+ * update = sqrt(accum_update + eps) / sqrt(accum_new + eps) * grad
	22	+ * var_new = var - lr * update
	23	+ * accum_update_new = rho * accum_update + (1 - rho) * update^2
	24	+ *
	25	+ * 注意：本算子为 inplace 语义，varOut/accumOut/accumUpdateOut 分别与
	26	+ * var/accum/accumUpdate 共享 Device 内存。
	27	+ */
	28	+#include <iostream>
	29	+#include <vector>
	30	+#include "acl/acl.h"
	31	+#include "aclnn_apply_adadelta.h"
	32	+
	33	+#define CHECK_RET(cond, return_expr) \
	34	+ do { \
	35	+ if (!(cond)) { \
	36	+ return_expr; \
	37	+ } \
	38	+ } while (0)
	39	+
	40	+#define LOG_PRINT(message, ...) \
	41	+ do { \
	42	+ printf(message, ##__VA_ARGS__); \
	43	+ } while (0)
	44	+
	45	+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
	46	+ int64_t shape_size = 1;
	47	+ for (auto i : shape) {
	48	+ shape_size *= i;
	49	+ }
	50	+ return shape_size;
	51	+}
	52	+
	53	+int Init(int32_t deviceId, aclrtStream* stream) {
	54	+ // 固定写法，资源初始化
	55	+ auto ret = aclInit(nullptr);
	56	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	57	+ ret = aclrtSetDevice(deviceId);
	58	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	59	+ ret = aclrtCreateStream(stream);
	60	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	61	+ return 0;
	62	+}
	63	+
	64	+template <typename T>
	65	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
	66	+ aclDataType dataType, aclTensor** tensor) {
	67	+ auto size = GetShapeSize(shape) * sizeof(T);
	68	+ // 调用aclrtMalloc申请Device侧内存
	69	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	70	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	71	+ // 调用aclrtMemcpy将Host侧数据拷贝到Device侧内存上
	72	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	73	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	74	+ // 计算连续tensor的strides
	75	+ std::vector<int64_t> strides(shape.size(), 1);
	76	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	77	+ strides[i] = shape[i + 1] * strides[i + 1];
	78	+ }
	79	+ // 调用aclCreateTensor接口创建aclTensor
	80	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	81	+ shape.data(), shape.size(), *deviceAddr);
	82	+ return 0;
	83	+}
	84	+
	85	+int main() {
	86	+ // 1. （固定写法）device/stream 初始化
	87	+ int32_t deviceId = 0;
	88	+ aclrtStream stream;
	89	+ auto ret = Init(deviceId, &stream);
	90	+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	91	+
	92	+ // 2. 构造输入与输出
	93	+ // var/accum/accumUpdate/grad shape 必须完全一致，dtype 必须一致（fp16 或 fp32）。
	94	+ // lr/rho/epsilon 为 aclScalar，dtype 需与 Tensor 一致。
	95	+ std::vector<int64_t> shape = {4, 2};
	96	+ void* varDeviceAddr = nullptr;
	97	+ void* accumDeviceAddr = nullptr;
	98	+ void* accumUpdateDeviceAddr = nullptr;
	99	+ void* gradDeviceAddr = nullptr;
	100	+ aclTensor* var = nullptr;
	101	+ aclTensor* accum = nullptr;
	102	+ aclTensor* accumUpdate = nullptr;
	103	+ aclTensor* grad = nullptr;
	104	+ aclTensor* varOut = nullptr;
	105	+ aclTensor* accumOut = nullptr;
	106	+ aclTensor* accumUpdateOut = nullptr;
	107	+ aclScalar* lr = nullptr;
	108	+ aclScalar* rho = nullptr;
	109	+ aclScalar* epsilon = nullptr;
	110	+
	111	+ std::vector<float> varHostData = {1.0f, 2.0f, -1.0f, 0.5f, 0.0f, -0.5f, 1.5f, -2.0f};
	112	+ std::vector<float> accumHostData = {0.1f, 0.2f, 0.3f, 0.4f, 0.5f, 0.1f, 0.2f, 0.3f};
	113	+ std::vector<float> accumUpdateHostData = {0.01f, 0.02f, 0.03f, 0.04f, 0.05f, 0.01f, 0.02f, 0.03f};
	114	+ std::vector<float> gradHostData = {0.5f, -0.3f, 0.1f, -0.2f, 0.4f, -0.1f, 0.3f, -0.4f};
	115	+ float lrValue = 0.01f;
	116	+ float rhoValue = 0.9f;
	117	+ float epsilonValue = 1e-6f;
	118	+
	119	+ // 创建 var/accum/accumUpdate/grad aclTensor
	120	+ ret = CreateAclTensor(varHostData, shape, &varDeviceAddr, aclDataType::ACL_FLOAT, &var);
	121	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	122	+ ret = CreateAclTensor(accumHostData, shape, &accumDeviceAddr, aclDataType::ACL_FLOAT, &accum);
	123	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	124	+ ret = CreateAclTensor(accumUpdateHostData, shape, &accumUpdateDeviceAddr, aclDataType::ACL_FLOAT, &accumUpdate);
	125	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	126	+ ret = CreateAclTensor(gradHostData, shape, &gradDeviceAddr, aclDataType::ACL_FLOAT, &grad);
	127	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	128	+
	129	+ // inplace 语义：输出 tensor 复用输入 Device 地址
	130	+ std::vector<int64_t> strides(shape.size(), 1);
	131	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	132	+ strides[i] = shape[i + 1] * strides[i + 1];
	133	+ }
	134	+ varOut = aclCreateTensor(shape.data(), shape.size(), aclDataType::ACL_FLOAT, strides.data(), 0,
	135	+ aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), varDeviceAddr);
	136	+ accumOut = aclCreateTensor(shape.data(), shape.size(), aclDataType::ACL_FLOAT, strides.data(), 0,
	137	+ aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), accumDeviceAddr);
	138	+ accumUpdateOut = aclCreateTensor(shape.data(), shape.size(), aclDataType::ACL_FLOAT, strides.data(), 0,
	139	+ aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), accumUpdateDeviceAddr);
	140	+
	141	+ // 创建 Scalar 参数（dtype 需与 Tensor 保持一致）
	142	+ lr = aclCreateScalar(&lrValue, aclDataType::ACL_FLOAT);
	143	+ rho = aclCreateScalar(&rhoValue, aclDataType::ACL_FLOAT);
	144	+ epsilon = aclCreateScalar(&epsilonValue, aclDataType::ACL_FLOAT);
	145	+
	146	+ // 3. 调用 CANN 算子库 API（两段式接口）
	147	+ uint64_t workspaceSize = 0;
	148	+ aclOpExecutor* executor = nullptr;
	149	+
	150	+ // 调用 aclnnApplyAdadelta 第一段接口获取 workspace 大小
	151	+ ret = aclnnApplyAdadeltaGetWorkspaceSize(var, accum, accumUpdate, lr, rho, epsilon, grad,
	152	+ varOut, accumOut, accumUpdateOut,
	153	+ &workspaceSize, &executor);
	154	+ CHECK_RET(ret == ACL_SUCCESS,
	155	+ LOG_PRINT("aclnnApplyAdadeltaGetWorkspaceSize failed. ERROR: %d\n", ret);
	156	+ return ret);
	157	+
	158	+ // 根据第一段接口计算出的 workspaceSize 申请 Device 内存
	159	+ void* workspaceAddr = nullptr;
	160	+ if (workspaceSize > 0) {
	161	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	162	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
	163	+ }
	164	+
	165	+ // 调用 aclnnApplyAdadelta 第二段接口执行计算
	166	+ ret = aclnnApplyAdadelta(workspaceAddr, workspaceSize, executor, stream);
	167	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnApplyAdadelta failed. ERROR: %d\n", ret); return ret);
	168	+
	169	+ // 4. （固定写法）同步等待任务执行结束
	170	+ ret = aclrtSynchronizeStream(stream);
	171	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	172	+
	173	+ // 5. 获取输出（inplace：从输入 Device 地址读回）
	174	+ auto size = GetShapeSize(shape);
	175	+ std::vector<float> varResult(size, 0);
	176	+ std::vector<float> accumResult(size, 0);
	177	+ std::vector<float> accumUpdateResult(size, 0);
	178	+
	179	+ ret = aclrtMemcpy(varResult.data(), varResult.size() * sizeof(float),
	180	+ varDeviceAddr, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
	181	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy var result failed. ERROR: %d\n", ret); return ret);
	182	+ ret = aclrtMemcpy(accumResult.data(), accumResult.size() * sizeof(float),
	183	+ accumDeviceAddr, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
	184	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy accum result failed. ERROR: %d\n", ret); return ret);
	185	+ ret = aclrtMemcpy(accumUpdateResult.data(), accumUpdateResult.size() * sizeof(float),
	186	+ accumUpdateDeviceAddr, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
	187	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy accum_update result failed. ERROR: %d\n", ret); return ret);
	188	+
	189	+ for (int64_t i = 0; i < size; i++) {
	190	+ LOG_PRINT("var_out[%ld]=%.6f accum_out[%ld]=%.6f accum_update_out[%ld]=%.6f\n",
	191	+ i, varResult[i], i, accumResult[i], i, accumUpdateResult[i]);
	192	+ }
	193	+
	194	+ // 6. 释放 aclTensor / aclScalar
	195	+ aclDestroyTensor(var);
	196	+ aclDestroyTensor(accum);
	197	+ aclDestroyTensor(accumUpdate);
	198	+ aclDestroyTensor(grad);
	199	+ aclDestroyTensor(varOut);
	200	+ aclDestroyTensor(accumOut);
	201	+ aclDestroyTensor(accumUpdateOut);
	202	+ aclDestroyScalar(lr);
	203	+ aclDestroyScalar(rho);
	204	+ aclDestroyScalar(epsilon);
	205	+
	206	+ // 7. 释放 Device 资源
	207	+ aclrtFree(varDeviceAddr);
	208	+ aclrtFree(accumDeviceAddr);
	209	+ aclrtFree(accumUpdateDeviceAddr);
	210	+ aclrtFree(gradDeviceAddr);
	211	+ if (workspaceSize > 0) {
	212	+ aclrtFree(workspaceAddr);
	213	+ }
	214	+ aclrtDestroyStream(stream);
	215	+ aclrtResetDevice(deviceId);
	216	+ aclFinalize();
	217	+ return 0;
	218	+}

Aexperimental/optim/apply_adadelta/op_api/aclnn_apply_adadelta.cpp+256-0

@@ -0,0 +1,256 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/**
	16	+ * @file aclnn_apply_adadelta.cpp
	17	+ * @brief ACLNN L2 API implementation for ApplyAdadelta
	18	+ *
	19	+ * Two-stage interface:
	20	+ * 1. aclnnApplyAdadeltaGetWorkspaceSize - parameter checking, Contiguous, L0 dispatch
	21	+ * 2. aclnnApplyAdadelta - execute computation
	22	+ *
	23	+ * Scalar parameters (lr, rho, epsilon) are extracted from aclScalar and passed
	24	+ * as attrs to the Tiling function via the executor.
	25	+ */
	26	+
	27	+#include "aclnn_apply_adadelta.h"
	28	+#include "apply_adadelta.h"
	29	+#include "aclnn_kernels/contiguous.h"
	30	+#include "aclnn_kernels/common/op_error_check.h"
	31	+#include "opdev/op_log.h"
	32	+#include "opdev/op_dfx.h"
	33	+#include "opdev/common_types.h"
	34	+#include "opdev/data_type_utils.h"
	35	+#include "opdev/make_op_executor.h"
	36	+#include "opdev/platform.h"
	37	+
	38	+using namespace op;
	39	+
	40	+#define ACLNN_MAX_SHAPE_RANK 8
	41	+
	42	+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
	43	+ DataType::DT_FLOAT, DataType::DT_FLOAT16
	44	+};
	45	+
	46	+static bool CheckNotNull(
	47	+ const aclTensor* var, const aclTensor* accum, const aclTensor* accumUpdate,
	48	+ const aclScalar* lr, const aclScalar* rho, const aclScalar* epsilon,
	49	+ const aclTensor* grad, const aclTensor* varOut,
	50	+ const aclTensor* accumOut, const aclTensor* accumUpdateOut)
	51	+{
	52	+ OP_CHECK_NULL(var, return false);
	53	+ OP_CHECK_NULL(accum, return false);
	54	+ OP_CHECK_NULL(accumUpdate, return false);
	55	+ OP_CHECK_NULL(lr, return false);
	56	+ OP_CHECK_NULL(rho, return false);
	57	+ OP_CHECK_NULL(epsilon, return false);
	58	+ OP_CHECK_NULL(grad, return false);
	59	+ OP_CHECK_NULL(varOut, return false);
	60	+ OP_CHECK_NULL(accumOut, return false);
	61	+ OP_CHECK_NULL(accumUpdateOut, return false);
	62	+ return true;
	63	+}
	64	+
	65	+static bool CheckDtypeValid(
	66	+ const aclTensor* var, const aclTensor* accum,
	67	+ const aclTensor* accumUpdate, const aclTensor* grad)
	68	+{
	69	+ auto dtype = var->GetDataType();
	70	+ if (!CheckType(dtype, AICORE_DTYPE_SUPPORT_LIST)) {
	71	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	72	+ "Unsupported dtype: %d. Only FLOAT and FLOAT16 are supported.",
	73	+ static_cast<int>(dtype));
	74	+ return false;
	75	+ }
	76	+ // All tensor dtypes must match
	77	+ OP_CHECK_DTYPE_NOT_MATCH(accum, dtype, return false);
	78	+ OP_CHECK_DTYPE_NOT_MATCH(accumUpdate, dtype, return false);
	79	+ OP_CHECK_DTYPE_NOT_MATCH(grad, dtype, return false);
	80	+ return true;
	81	+}
	82	+
	83	+static bool CheckShapeConsistent(
	84	+ const aclTensor* var, const aclTensor* accum,
	85	+ const aclTensor* accumUpdate, const aclTensor* grad)
	86	+{
	87	+ OP_CHECK_MAX_DIM(var, ACLNN_MAX_SHAPE_RANK, return false);
	88	+
	89	+ auto varShape = var->GetViewShape();
	90	+ auto accumShape = accum->GetViewShape();
	91	+ auto auShape = accumUpdate->GetViewShape();
	92	+ auto gradShape = grad->GetViewShape();
	93	+
	94	+ // Reject 0-dim scalar tensors (REQUIREMENTS.md §5.3: supported shape dim 1~8).
	95	+ // See docs/aclnnApplyAdadelta.md "不支持 0 维标量 Tensor".
	96	+ if (varShape.GetDimNum() == 0 \|\| accumShape.GetDimNum() == 0 \|\|
	97	+ auShape.GetDimNum() == 0 \|\| gradShape.GetDimNum() == 0) {
	98	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	99	+ "0-dim scalar Tensor is not supported; supported dim range is 1~%d",
	100	+ ACLNN_MAX_SHAPE_RANK);
	101	+ return false;
	102	+ }
	103	+
	104	+ if (varShape.GetDimNum() != accumShape.GetDimNum() \|\|
	105	+ varShape.GetDimNum() != auShape.GetDimNum() \|\|
	106	+ varShape.GetDimNum() != gradShape.GetDimNum()) {
	107	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	108	+ "Shape dims mismatch: var=%zu, accum=%zu, accumUpdate=%zu, grad=%zu",
	109	+ varShape.GetDimNum(), accumShape.GetDimNum(),
	110	+ auShape.GetDimNum(), gradShape.GetDimNum());
	111	+ return false;
	112	+ }
	113	+
	114	+ for (size_t i = 0; i < varShape.GetDimNum(); i++) {
	115	+ if (varShape.GetDim(i) != accumShape.GetDim(i) \|\|
	116	+ varShape.GetDim(i) != auShape.GetDim(i) \|\|
	117	+ varShape.GetDim(i) != gradShape.GetDim(i)) {
	118	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	119	+ "Shape mismatch at dim %zu: var=%ld, accum=%ld, accumUpdate=%ld, grad=%ld",
	120	+ i, varShape.GetDim(i), accumShape.GetDim(i),
	121	+ auShape.GetDim(i), gradShape.GetDim(i));
	122	+ return false;
	123	+ }
	124	+ }
	125	+ return true;
	126	+}
	127	+
	128	+// Validate scalar values:
	129	+// - epsilon must be > 0 (numerical stability constant)
	130	+// - rho must be in [0, 1) (decay coefficient, per REQUIREMENTS.md)
	131	+// - lr must be >= 0 (learning rate; standard Adadelta requires non-negative lr.
	132	+// Note: Under negative lr, NPU observations show var_out is left unchanged
	133	+// while accum/accum_update are correctly updated, i.e. the computation is
	134	+// ill-defined. The current hypothesis is related to the CANN ACLNN tiling
	135	+// cache key missing scalar attrs; see docs/precision-report.md
	136	+// "CANN ACLNN tiling 缓存键缺失 attr" section for background. We therefore
	137	+// reject lr < 0 at the ACLNN layer as a conservative guard.)
	138	+static bool CheckScalarValues(const aclScalar* lr, const aclScalar* rho, const aclScalar* epsilon)
	chenfeng4月20日【建议级】lr<0 拒绝原因注释需澄清注释提到 "CANN ACLNN tiling 缓存键缺失 attr" 问题，但未引用具体 Issue ID。建议：如果是已知 CANN 问题，应引用具体 Issue/PR 链接如果是推测，应标注为 "待验证" 考虑是否可以规避此问题（如将 lr/rho/epsilon 加入 tiling cache key） Nerddddddddddd 4月24日评论： 4月24日评论：
	139	+{
	140	+ float lrVal = lr->ToFloat();
	141	+ float rhoVal = rho->ToFloat();
	142	+ float epsVal = epsilon->ToFloat();
	143	+ if (lrVal < 0.0f) {
	144	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	145	+ "lr must be >= 0, got %f", lrVal);
	146	+ return false;
	147	+ }
	chenfeng4月20日【建议级】rho 范围检查需对标 TensorFlow 当前检查 `rho in [0, 1)`（不包括 rho=1），但需要确认是否对标 TensorFlow `tf.raw_ops.ApplyAdadelta`。参考： TensorFlow 源码中 rho 通常允许 [0, 1] 范围 rho=1 时退化为 RMSProp-like 行为 PyTorch `torch.optim.Adadelta` 默认 rho=0.9，但未显式限制 rho=1 建议：查阅 TensorFlow 源码确认 rho 的实际限制范围。 Nerddddddddddd 4月24日评论： 4月24日评论：
	148	+ if (epsVal <= 0.0f) {
	149	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	150	+ "epsilon must be > 0, got %f", epsVal);
	151	+ return false;
	152	+ }
	153	+ if (rhoVal < 0.0f \|\| rhoVal >= 1.0f) {
	154	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	155	+ "rho must be in [0, 1), got %f", rhoVal);
	156	+ return false;
	157	+ }
	158	+ return true;
	159	+}
	160	+
	161	+static aclnnStatus CheckParams(
	162	+ const aclTensor* var, const aclTensor* accum, const aclTensor* accumUpdate,
	163	+ const aclScalar* lr, const aclScalar* rho, const aclScalar* epsilon,
	164	+ const aclTensor* grad, const aclTensor* varOut,
	165	+ const aclTensor* accumOut, const aclTensor* accumUpdateOut)
	166	+{
	167	+ if (!CheckNotNull(var, accum, accumUpdate, lr, rho, epsilon, grad, varOut, accumOut, accumUpdateOut)) {
	168	+ OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "Null pointer in input parameters");
	169	+ return ACLNN_ERR_PARAM_NULLPTR;
	170	+ }
	171	+ if (!CheckDtypeValid(var, accum, accumUpdate, grad)) {
	172	+ return ACLNN_ERR_PARAM_INVALID;
	173	+ }
	174	+ if (!CheckShapeConsistent(var, accum, accumUpdate, grad)) {
	175	+ return ACLNN_ERR_PARAM_INVALID;
	176	+ }
	177	+ if (!CheckScalarValues(lr, rho, epsilon)) {
	178	+ return ACLNN_ERR_PARAM_INVALID;
	179	+ }
	180	+ return ACLNN_SUCCESS;
	181	+}
	182	+
	183	+extern "C" aclnnStatus aclnnApplyAdadeltaGetWorkspaceSize(
	184	+ const aclTensor* var,
	185	+ const aclTensor* accum,
	186	+ const aclTensor* accumUpdate,
	187	+ const aclScalar* lr,
	188	+ const aclScalar* rho,
	189	+ const aclScalar* epsilon,
	190	+ const aclTensor* grad,
	191	+ aclTensor* varOut,
	192	+ aclTensor* accumOut,
	193	+ aclTensor* accumUpdateOut,
	194	+ uint64_t* workspaceSize,
	195	+ aclOpExecutor** executor)
	196	+{
	197	+ L2_DFX_PHASE_1(aclnnApplyAdadelta,
	198	+ DFX_IN(var, accum, accumUpdate, grad),
	199	+ DFX_OUT(varOut, accumOut, accumUpdateOut));
	200	+
	201	+ auto uniqueExecutor = CREATE_EXECUTOR();
	202	+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
	203	+
	204	+ auto ret = CheckParams(var, accum, accumUpdate, lr, rho, epsilon, grad, varOut, accumOut, accumUpdateOut);
	205	+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
	206	+
	207	+ if (var->IsEmpty()) {
	208	+ *workspaceSize = 0;
	209	+ uniqueExecutor.ReleaseTo(executor);
	210	+ return ACLNN_SUCCESS;
	211	+ }
	212	+
	213	+ // Make inputs contiguous
	214	+ auto varContiguous = l0op::Contiguous(var, uniqueExecutor.get());
	215	+ CHECK_RET(varContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	216	+ auto accumContiguous = l0op::Contiguous(accum, uniqueExecutor.get());
	217	+ CHECK_RET(accumContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	218	+ auto accumUpdateContiguous = l0op::Contiguous(accumUpdate, uniqueExecutor.get());
	219	+ CHECK_RET(accumUpdateContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	220	+ auto gradContiguous = l0op::Contiguous(grad, uniqueExecutor.get());
	221	+ CHECK_RET(gradContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	222	+
	223	+ // Extract scalar values and pass to L0 API
	224	+ float lrVal = lr->ToFloat();
	225	+ float rhoVal = rho->ToFloat();
	226	+ float epsVal = epsilon->ToFloat();
	227	+
	228	+ // Call L0 API
	229	+ auto opResults = l0op::ApplyAdadelta(
	230	+ varContiguous, accumContiguous, accumUpdateContiguous, gradContiguous,
	231	+ lrVal, rhoVal, epsVal,
	232	+ uniqueExecutor.get());
	233	+ CHECK_RET(opResults.varOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
	234	+
	235	+ // ViewCopy for all inplace outputs
	236	+ auto viewCopy1 = l0op::ViewCopy(opResults.varOut, varOut, uniqueExecutor.get());
	237	+ CHECK_RET(viewCopy1 != nullptr, ACLNN_ERR_INNER_NULLPTR);
	238	+ auto viewCopy2 = l0op::ViewCopy(opResults.accumOut, accumOut, uniqueExecutor.get());
	239	+ CHECK_RET(viewCopy2 != nullptr, ACLNN_ERR_INNER_NULLPTR);
	240	+ auto viewCopy3 = l0op::ViewCopy(opResults.accumUpdateOut, accumUpdateOut, uniqueExecutor.get());
	241	+ CHECK_RET(viewCopy3 != nullptr, ACLNN_ERR_INNER_NULLPTR);
	242	+
	243	+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
	244	+ uniqueExecutor.ReleaseTo(executor);
	245	+ return ACLNN_SUCCESS;
	246	+}
	247	+
	248	+extern "C" aclnnStatus aclnnApplyAdadelta(
	249	+ void* workspace,
	250	+ uint64_t workspaceSize,
	251	+ aclOpExecutor* executor,
	252	+ aclrtStream stream)
	253	+{
	254	+ L2_DFX_PHASE_2(aclnnApplyAdadelta);
	255	+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
	256	+}

Aexperimental/optim/apply_adadelta/op_api/aclnn_apply_adadelta.h+61-0

@@ -0,0 +1,61 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/**
	16	+ * @file aclnn_apply_adadelta.h
	17	+ * @brief ACLNN L2 API declaration for ApplyAdadelta
	18	+ *
	19	+ * Two-stage interface:
	20	+ * - aclnnApplyAdadeltaGetWorkspaceSize: compute workspace size, create executor
	21	+ * - aclnnApplyAdadelta: execute computation
	22	+ */
	23	+
	24	+#ifndef ACLNN_APPLY_ADADELTA_H_
	25	+#define ACLNN_APPLY_ADADELTA_H_
	26	+
	27	+#include "aclnn/aclnn_base.h"
	28	+
	29	+#ifndef ACLNN_API
	30	+#define ACLNN_API __attribute__((visibility("default")))
	31	+#endif
	32	+
	33	+#ifdef __cplusplus
	34	+extern "C" {
	35	+#endif
	36	+
	37	+ACLNN_API aclnnStatus aclnnApplyAdadeltaGetWorkspaceSize(
	38	+ const aclTensor *var,
	39	+ const aclTensor *accum,
	40	+ const aclTensor *accumUpdate,
	41	+ const aclScalar *lr,
	42	+ const aclScalar *rho,
	43	+ const aclScalar *epsilon,
	44	+ const aclTensor *grad,
	45	+ aclTensor *varOut,
	46	+ aclTensor *accumOut,
	47	+ aclTensor *accumUpdateOut,
	48	+ uint64_t *workspaceSize,
	49	+ aclOpExecutor **executor);
	50	+
	51	+ACLNN_API aclnnStatus aclnnApplyAdadelta(
	52	+ void *workspace,
	53	+ uint64_t workspaceSize,
	54	+ aclOpExecutor *executor,
	55	+ aclrtStream stream);
	56	+
	57	+#ifdef __cplusplus
	58	+}
	59	+#endif
	60	+
	61	+#endif // ACLNN_APPLY_ADADELTA_H_

Aexperimental/optim/apply_adadelta/op_api/apply_adadelta.cpp+109-0

@@ -0,0 +1,109 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/**
	16	+ * @file apply_adadelta.cpp
	17	+ * @brief ACLNN L0 API implementation for ApplyAdadelta
	18	+ *
	19	+ * L0 API: InferShape, IsAiCoreSupport, AllocTensor, ApplyAdadeltaAiCore
	20	+ *
	21	+ * This is an inplace operator. The Kernel receives both input and output GM addresses,
	22	+ * where output addresses overlap with input addresses (var/accum/accumUpdate).
	23	+ */
	24	+
	25	+#include "apply_adadelta.h"
	26	+#include "opdev/op_log.h"
	27	+#include "opdev/op_dfx.h"
	28	+#include "opdev/shape_utils.h"
	29	+#include "opdev/make_op_executor.h"
	30	+
	31	+using namespace op;
	32	+
	33	+namespace l0op {
	34	+
	35	+OP_TYPE_REGISTER(ApplyAdadelta);
	36	+
	37	+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
	38	+ DataType::DT_FLOAT, DataType::DT_FLOAT16
	39	+};
	40	+
	41	+static bool IsAiCoreSupport(const aclTensor* var)
	42	+{
	43	+ return CheckType(var->GetDataType(), AICORE_DTYPE_SUPPORT_LIST);
	44	+}
	45	+
	46	+static const aclTensor* ApplyAdadeltaAiCore(
	47	+ const aclTensor* var,
	48	+ const aclTensor* accum,
	49	+ const aclTensor* accumUpdate,
	50	+ const aclTensor* grad,
	51	+ const aclTensor* varOut,
	52	+ const aclTensor* accumOut,
	53	+ const aclTensor* accumUpdateOut,
	54	+ float lr, float rho, float epsilon,
	55	+ aclOpExecutor* executor)
	56	+{
	57	+ L0_DFX(ApplyAdadeltaAiCore, var, accum, accumUpdate, grad, varOut, accumOut, accumUpdateOut);
	58	+
	59	+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(ApplyAdadelta,
	60	+ OP_INPUT(var, accum, accumUpdate, grad),
	61	+ OP_OUTPUT(varOut, accumOut, accumUpdateOut),
	62	+ OP_ATTR(lr, rho, epsilon));
	63	+ OP_CHECK(
	64	+ ret == ACLNN_SUCCESS,
	65	+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "ApplyAdadeltaAiCore failed."),
	66	+ return nullptr);
	67	+ return varOut;
	68	+}
	69	+
	70	+ApplyAdadeltaOutputs ApplyAdadelta(
	71	+ const aclTensor* var,
	72	+ const aclTensor* accum,
	73	+ const aclTensor* accumUpdate,
	74	+ const aclTensor* grad,
	75	+ float lr,
	76	+ float rho,
	77	+ float epsilon,
	78	+ aclOpExecutor* executor)
	79	+{
	80	+ ApplyAdadeltaOutputs emptyResult = {nullptr, nullptr, nullptr};
	81	+
	82	+ if (!IsAiCoreSupport(var)) {
	83	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	84	+ "ApplyAdadelta not supported: dtype=%d.",
	85	+ static_cast<int>(var->GetDataType()));
	86	+ return emptyResult;
	87	+ }
	88	+
	89	+ // Output shape = input shape (inplace)
	90	+ auto varShape = var->GetViewShape();
	91	+
	92	+ // Allocate output tensors (inplace: framework handles address aliasing;
	93	+ // the L2 API completes aliasing / write-back to the original var/accum/
	94	+ // accumUpdate buffers via ViewCopy after this L0 call returns.)
	95	+ const aclTensor* varOut = executor->AllocTensor(varShape, var->GetDataType());
	96	+ const aclTensor* accumOut = executor->AllocTensor(varShape, var->GetDataType());
	97	+ const aclTensor* accumUpdateOut = executor->AllocTensor(varShape, var->GetDataType());
	98	+
	99	+ const aclTensor* result = ApplyAdadeltaAiCore(var, accum, accumUpdate, grad,
	100	+ varOut, accumOut, accumUpdateOut,
	101	+ lr, rho, epsilon, executor);
	102	+ if (result == nullptr) {
	103	+ return emptyResult;
	104	+ }
	105	+
	106	+ return {varOut, accumOut, accumUpdateOut};
	107	+}
	108	+
	109	+} // namespace l0op

Aexperimental/optim/apply_adadelta/op_api/apply_adadelta.h+49-0

@@ -0,0 +1,49 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/**
	16	+ * @file apply_adadelta.h
	17	+ * @brief ACLNN L0 API declaration for ApplyAdadelta
	18	+ */
	19	+
	20	+#ifndef OP_API_INC_LEVEL0_APPLY_ADADELTA_H_
	21	+#define OP_API_INC_LEVEL0_APPLY_ADADELTA_H_
	22	+
	23	+#include "opdev/op_executor.h"
	24	+
	25	+namespace l0op {
	26	+
	27	+/**
	28	+ * ApplyAdadelta L0 API - returns 3 output tensors via a struct
	29	+ * (var_out, accum_out, accum_update_out)
	30	+ */
	31	+struct ApplyAdadeltaOutputs {
	32	+ const aclTensor* varOut;
	33	+ const aclTensor* accumOut;
	34	+ const aclTensor* accumUpdateOut;
	35	+};
	36	+
	37	+ApplyAdadeltaOutputs ApplyAdadelta(
	38	+ const aclTensor* var,
	39	+ const aclTensor* accum,
	40	+ const aclTensor* accumUpdate,
	41	+ const aclTensor* grad,
	42	+ float lr,
	43	+ float rho,
	44	+ float epsilon,
	45	+ aclOpExecutor* executor);
	46	+
	47	+} // namespace l0op
	48	+
	49	+#endif // OP_API_INC_LEVEL0_APPLY_ADADELTA_H_

Aexperimental/optim/apply_adadelta/op_host/apply_adadelta_def.cpp+80-0

@@ -0,0 +1,80 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file apply_adadelta_def.cpp
	17	+ * \brief ApplyAdadelta operator definition
	18	+ *
	19	+ * Inputs: var, accum, accumUpdate, grad (Tensor) + lr, rho, epsilon (Scalar via TilingData)
	20	+ * Outputs: varOut, accumOut, accumUpdateOut (Tensor, inplace with inputs)
	21	+ * Target: ascend950 (arch35)
	22	+ */
	23	+#include "register/op_def_registry.h"
	24	+
	25	+namespace ops {
	26	+
	27	+class ApplyAdadelta : public OpDef {
	28	+public:
	29	+ explicit ApplyAdadelta(const char* name) : OpDef(name)
	30	+ {
	31	+ // Helper: configure a Tensor Input with standard (FP32/FP16, ND) schema.
	32	+ auto cfgInput = [this](const char* n) -> void {
	33	+ this->Input(n)
	34	+ .ParamType(REQUIRED)
	35	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
	36	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	37	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	38	+ .AutoContiguous();
	39	+ };
	40	+ auto cfgOutput = [this](const char* n) -> void {
	41	+ this->Output(n)
	42	+ .ParamType(REQUIRED)
	43	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
	44	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	45	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	46	+ .AutoContiguous();
	47	+ };
	48	+
	49	+ // 4 Tensor inputs: var, accum, accumUpdate, grad
	50	+ cfgInput("var");
	51	+ cfgInput("accum");
	52	+ cfgInput("accumUpdate");
	53	+ cfgInput("grad");
	54	+
	55	+ // 3 Tensor outputs: varOut, accumOut, accumUpdateOut (inplace)
	56	+ cfgOutput("varOut");
	57	+ cfgOutput("accumOut");
	58	+ cfgOutput("accumUpdateOut");
	59	+
	60	+ // Scalar attrs: lr, rho, epsilon (set by ACLNN L0 API, accessed by index in Tiling)
	61	+ this->Attr("lr").AttrType(REQUIRED).Float();
	62	+ this->Attr("rho").AttrType(REQUIRED).Float();
	63	+ this->Attr("epsilon").AttrType(REQUIRED).Float();
	64	+
	65	+ // ascend950 (arch35) configuration
	66	+ OpAICoreConfig aiCoreConfig;
	67	+ aiCoreConfig.DynamicCompileStaticFlag(true)
	68	+ .DynamicFormatFlag(false)
	69	+ .DynamicRankSupportFlag(true)
	70	+ .DynamicShapeSupportFlag(true)
	71	+ .NeedCheckSupportFlag(false)
	72	+ .PrecisionReduceFlag(true)
	73	+ .ExtendCfgInfo("opFile.value", "apply_adadelta");
	74	+ this->AICore().AddConfig("ascend950", aiCoreConfig);
	75	+ }
	76	+};
	77	+
	78	+OP_ADD(ApplyAdadelta);
	79	+
	80	+} // namespace ops

Aexperimental/optim/apply_adadelta/op_host/apply_adadelta_infershape.cpp+52-0

@@ -0,0 +1,52 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file apply_adadelta_infershape.cpp
	17	+ * \brief ApplyAdadelta shape inference
	18	+ *
	19	+ * All inputs share the same shape. Outputs = input shape (inplace).
	20	+ * 3 outputs: varOut, accumOut, accumUpdateOut
	21	+ */
	22	+
	23	+#include "register/op_impl_registry.h"
	24	+#include "exe_graph/runtime/infer_shape_context.h"
	25	+
	26	+using namespace ge;
	27	+
	28	+namespace ops {
	29	+
	30	+static ge::graphStatus InferShape4ApplyAdadelta(gert::InferShapeContext* context)
	31	+{
	32	+ // Input(0) = var
	33	+ const gert::Shape* inputShape = context->GetInputShape(0);
	34	+ if (inputShape == nullptr) {
	35	+ return ge::GRAPH_FAILED;
	chenfeng4月20日 OP_CHECK_NULL_WITH_CONTEXT判断即可 Nerddddddddddd 4月24日评论： 4月24日评论：
	36	+ }
	37	+
	38	+ // 3 outputs: varOut(0), accumOut(1), accumUpdateOut(2)
	39	+ for (size_t i = 0; i < 3; i++) {
	40	+ gert::Shape* outputShape = context->GetOutputShape(i);
	41	+ if (outputShape == nullptr) {
	chenfeng4月20日 OP_CHECK_NULL_WITH_CONTEXT判断即可 Nerddddddddddd 4月24日评论： 4月24日评论：
	42	+ return ge::GRAPH_FAILED;
	43	+ }
	44	+ outputShape = inputShape;
	45	+ }
	46	+
	47	+ return ge::GRAPH_SUCCESS;
	48	+}
	49	+
	50	+IMPL_OP_INFERSHAPE(ApplyAdadelta).InferShape(InferShape4ApplyAdadelta);
	51	+
	52	+} // namespace ops

Aexperimental/optim/apply_adadelta/op_host/apply_adadelta_tiling.cpp+220-0

@@ -0,0 +1,220 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file apply_adadelta_tiling.cpp
	17	+ * \brief ApplyAdadelta tiling implementation (arch35)
	18	+ *
	19	+ * TilingKey selection:
	20	+ * TilingKey_0: D_T_X=C_DT_FLOAT, BUFFER_MODE=0 (FP32, single buffer)
	chenfeng4月20日【建议级】MIN_SPLIT_THRESHOLD 阈值偏小当前阈值 `1024` 导致小 shape（如 1024 元素）也启用双缓冲，但实际上对于小数据量，双缓冲的队列初始化开销可能超过收益。建议：调整为 `8192` 或更高（参考其他算子实践）或通过性能测试确定最优阈值 Nerddddddddddd 4月24日评论： 4月24日评论：
	21	+ * TilingKey_1: D_T_X=C_DT_FLOAT, BUFFER_MODE=1 (FP32, double buffer)
	22	+ * TilingKey_2: D_T_X=C_DT_FLOAT16, BUFFER_MODE=0 (FP16, single buffer)
	23	+ * TilingKey_3: D_T_X=C_DT_FLOAT16, BUFFER_MODE=1 (FP16, double buffer)
	24	+ *
	25	+ * Scalar parameters (lr, rho, epsilon) are passed via TilingData.
	26	+ * The ACLNN layer reads aclScalar values and writes them into attrs
	27	+ * named "lr", "rho", "epsilon" for the Tiling function to extract.
	28	+ */
	29	+
	30	+#include "register/op_def_registry.h"
	31	+#include "op_common/log/log.h"
	32	+#include "op_common/op_host/util/math_util.h"
	33	+#include "op_common/op_host/util/platform_util.h"
	34	+#include "../op_kernel/apply_adadelta_tiling_data.h"
	35	+#include "../op_kernel/apply_adadelta_tiling_key.h"
	chenfeng4月20日【建议级】GetTensorSlots 注释表述易混淆注释说 `"FP16 single: 38 bytes/elem"`，但实际计算公式应为： `FP16 single: 7 Que x 1 (half) + 6 fp32 TBuf = 7 * ubFactor * 2B + 6 * ubFactor * 4B = ubFactor * (14 + 24) B = ubFactor * 38 B` 建议重新表述为： `// FP16 single buffer: 7 half-sized TQue (2B) + 6 float-sized TBuf (4B) // Total: 7ubFactor2B + 6ubFactor4B = 38*ubFactor bytes // Equivalent float slots = 38/4 = 9.5 -> round up to 10` Nerddddddddddd 4月24日评论： 4月24日评论：
	36	+
	37	+namespace optiling {
	38	+
	39	+using Ops::Base::CeilDiv;
	40	+using Ops::Base::CeilAlign;
	41	+using Ops::Base::FloorDiv;
	42	+using Ops::Base::FloorAlign;
	43	+using Ops::Base::GetUbBlockSize;
	44	+
	45	+constexpr int64_t MIN_SPLIT_THRESHOLD = 1024;
	46	+
	47	+// UB tensor slot count per TilingKey (see DESIGN.md 3.4.7/3.5.6/3.6.6/3.7.6)
	48	+//
	49	+// Per-element UB byte cost (in fp32-equivalent slots; 1 slot == 4 B == sizeof(float)):
	50	+// FP32 single buffer: 7 TQue (fp32, 4B) x 1 + 2 TBuf (fp32, 4B)
	51	+// = (7 * 4 + 2 * 4) B/elem = 36 B/elem = 9 slots
	52	+// FP32 double buffer: 7 TQue (fp32, 4B) x 2 + 2 TBuf (fp32, 4B)
	53	+// = (7 * 4 * 2 + 2 * 4) B/elem = 64 B/elem = 16 slots
	54	+// FP16 single buffer: 7 TQue (half, 2B) x 1 + 6 TBuf (fp32, 4B)
	55	+// = (7 * 2 + 6 * 4) B/elem = 38 B/elem
	56	+// => ceil(38 / 4) = 10 slots
	57	+// FP16 double buffer: 7 TQue (half, 2B) x 2 + 6 TBuf (fp32, 4B)
	58	+// = (7 * 2 * 2 + 6 * 4) B/elem = 52 B/elem = 13 slots
	59	+//
	60	+// Note: BUFFER_MODE only affects TQue (which may be double-buffered via
	61	+// BUFFER_NUM=2); TBuf is NEVER double-buffered regardless of BUFFER_MODE.
	62	+// The "+2 TBuf" / "+6 fp32 TBuf" terms above stay constant across both
	63	+// single-buffer and double-buffer paths.
	64	+static int64_t GetTensorSlots(ge::DataType dataType, uint64_t bufferMode)
	65	+{
	66	+ if (dataType == ge::DT_FLOAT) {
	67	+ // FP32: see header comment for derivation
	68	+ return bufferMode ? 16 : 9;
	69	+ } else {
	chenfeng4月20日魔鬼数字 Nerddddddddddd 4月23日评论： 4月23日评论：
	70	+ // FP16: see header comment for derivation
	71	+ return bufferMode ? 13 : 10;
	72	+ }
	73	+}
	74	+
	75	+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
	76	+{
	77	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	78	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	79	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	80	+ coreNum = ascendcPlatform.GetCoreNumAiv();
	81	+ OP_CHECK_IF(coreNum <= 0, OP_LOGE(context, "coreNum must be > 0, got %ld", coreNum), return ge::GRAPH_FAILED);
	82	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
	83	+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
	84	+ return ge::GRAPH_SUCCESS;
	85	+}
	86	+
	87	+// Attr order matches op_def: lr(0), rho(1), epsilon(2).
	88	+// epsilon == 0 would trigger Div-by-zero in sqrt(accum+eps) denominator.
	89	+static ge::graphStatus ParseScalarAttrs(gert::TilingContext* context, ApplyAdadeltaTilingData* tiling)
	90	+{
	91	+ auto* attrs = context->GetAttrs();
	92	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
	93	+ const float* lrPtr = attrs->GetFloat(0);
	94	+ OP_CHECK_NULL_WITH_CONTEXT(context, lrPtr);
	95	+ const float* rhoPtr = attrs->GetFloat(1);
	96	+ OP_CHECK_NULL_WITH_CONTEXT(context, rhoPtr);
	97	+ const float* epsPtr = attrs->GetFloat(2);
	98	+ OP_CHECK_NULL_WITH_CONTEXT(context, epsPtr);
	99	+ OP_CHECK_IF(
	100	+ *epsPtr <= 0.0f,
	101	+ OP_LOGE(context, "epsilon must be > 0, got %f", *epsPtr),
	chenfeng4月20日日志有最新要求的整改，请按照最新日志要求整改日志 Nerddddddddddd 4月24日评论： 4月24日评论：
	102	+ return ge::GRAPH_FAILED);
	103	+ tiling->lr = *lrPtr;
	104	+ tiling->rho = *rhoPtr;
	105	+ tiling->oneMinusRho = 1.0f - *rhoPtr;
	106	+ tiling->epsilon = *epsPtr;
	107	+ return ge::GRAPH_SUCCESS;
	108	+}
	109	+
	110	+static int64_t ComputeSplit(ApplyAdadeltaTilingData* tiling, ge::DataType dataType,
	111	+ uint64_t ubSize, int64_t coreNum, int64_t ubBlockSize, uint64_t& bufferMode)
	112	+{
	113	+ tiling->blockFactor = CeilAlign(CeilDiv(tiling->totalNum, coreNum), ubBlockSize);
	114	+ int64_t usedCoreNum = CeilDiv(tiling->totalNum, tiling->blockFactor);
	115	+ bufferMode = (tiling->totalNum > MIN_SPLIT_THRESHOLD) ? 1 : 0;
	116	+ int64_t slots = GetTensorSlots(dataType, bufferMode);
	117	+ constexpr int64_t elemBytes = 4; // Internal computation in fp32
	118	+ tiling->ubFactor = FloorAlign(
	119	+ FloorDiv(static_cast<int64_t>(ubSize) / elemBytes, slots), ubBlockSize);
	120	+ return usedCoreNum;
	121	+}
	122	+
	123	+static ge::graphStatus GetShapeAndDataType(gert::TilingContext* context,
	124	+ int64_t& totalNum, ge::DataType& dataType)
	125	+{
	126	+ auto varShape = context->GetInputShape(0);
	127	+ OP_CHECK_NULL_WITH_CONTEXT(context, varShape);
	128	+ totalNum = varShape->GetStorageShape().GetShapeSize();
	129	+ auto varDesc = context->GetInputDesc(0);
	130	+ OP_CHECK_NULL_WITH_CONTEXT(context, varDesc);
	131	+ dataType = varDesc->GetDataType();
	132	+ return ge::GRAPH_SUCCESS;
	133	+}
	134	+
	135	+static ge::graphStatus SetWorkspace(gert::TilingContext* context)
	136	+{
	137	+ size_t* ws = context->GetWorkspaceSizes(1);
	138	+ OP_CHECK_NULL_WITH_CONTEXT(context, ws);
	139	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	140	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	141	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	142	+ // User workspace is 0 (operator needs no extra workspace), system workspace
	143	+ // is obtained from AscendC platform interface to stay consistent with framework
	144	+ // expectations across SoC versions.
	145	+ ws[0] = static_cast<size_t>(ascendcPlatform.GetLibApiWorkSpaceSize());
	146	+ return ge::GRAPH_SUCCESS;
	147	+}
	148	+
	149	+static ge::graphStatus HandleEmptyTensor(gert::TilingContext* context, uint32_t dTypeX)
	150	+{
	151	+ context->SetBlockDim(1);
	152	+ uint32_t bufMode = 0;
	153	+ ASCENDC_TPL_SEL_PARAM(context, dTypeX, bufMode);
	154	+ return ge::GRAPH_SUCCESS;
	155	+}
	156	+
	157	+static ge::graphStatus InitTilingData(gert::TilingContext* context,
	158	+ int64_t totalNum, ApplyAdadeltaTilingData*& tiling)
	159	+{
	160	+ tiling = context->GetTilingData<ApplyAdadeltaTilingData>();
	161	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	162	+ OP_CHECK_IF(
	163	+ memset_s(tiling, sizeof(tiling), 0, sizeof(tiling)) != EOK,
	164	+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
	165	+ tiling->totalNum = totalNum;
	166	+ return ge::GRAPH_SUCCESS;
	167	+}
	168	+
	169	+static ge::graphStatus ApplyAdadeltaTilingFunc(gert::TilingContext* context)
	170	+{
	171	+ uint64_t ubSize;
	172	+ int64_t coreNum;
	173	+ OP_CHECK_IF(
	174	+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
	175	+ OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
	176	+
	177	+ int64_t totalNum = 0;
	178	+ ge::DataType dataType = ge::DT_FLOAT;
	179	+ OP_CHECK_IF(
	180	+ GetShapeAndDataType(context, totalNum, dataType) != ge::GRAPH_SUCCESS,
	181	+ OP_LOGE(context, "GetShapeAndDataType error"), return ge::GRAPH_FAILED);
	182	+
	183	+ OP_CHECK_IF(
	184	+ SetWorkspace(context) != ge::GRAPH_SUCCESS,
	185	+ OP_LOGE(context, "SetWorkspace error"), return ge::GRAPH_FAILED);
	186	+
	chenfeng4月20日 WS_SYS_SIZE 改为用接口获取的ascendc的workspace，不要用自己写的0或者16m或者32b这种数字，asendc接口是ascendcPlatform.GetLibApiWorkSpaceSize(); Nerddddddddddd 4月24日评论： 4月24日评论：
	187	+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
	188	+ if (totalNum == 0) {
	189	+ return HandleEmptyTensor(context, dTypeX);
	190	+ }
	191	+
	192	+ ApplyAdadeltaTilingData* tiling = nullptr;
	193	+ OP_CHECK_IF(
	194	+ InitTilingData(context, totalNum, tiling) != ge::GRAPH_SUCCESS,
	195	+ OP_LOGE(context, "InitTilingData error"), return ge::GRAPH_FAILED);
	196	+
	197	+ uint64_t bufferMode = 0;
	198	+ int64_t usedCoreNum = ComputeSplit(tiling, dataType, ubSize, coreNum, GetUbBlockSize(context), bufferMode);
	199	+
	200	+ OP_CHECK_IF(
	201	+ ParseScalarAttrs(context, tiling) != ge::GRAPH_SUCCESS,
	202	+ OP_LOGE(context, "ParseScalarAttrs error"), return ge::GRAPH_FAILED);
	203	+
	204	+ context->SetBlockDim(usedCoreNum);
	205	+ ASCENDC_TPL_SEL_PARAM(context, dTypeX, static_cast<uint32_t>(bufferMode));
	206	+ return ge::GRAPH_SUCCESS;
	207	+}
	208	+
	209	+static ge::graphStatus TilingParseForApplyAdadelta([[maybe_unused]] gert::TilingParseContext* context)
	210	+{
	211	+ return ge::GRAPH_SUCCESS;
	212	+}
	213	+
	214	+struct ApplyAdadeltaCompileInfo {};
	215	+
	216	+IMPL_OP_OPTILING(ApplyAdadelta)
	217	+ .Tiling(ApplyAdadeltaTilingFunc)
	218	+ .TilingParse<ApplyAdadeltaCompileInfo>(TilingParseForApplyAdadelta);
	219	+
	220	+} // namespace optiling

Aexperimental/optim/apply_adadelta/op_kernel/apply_adadelta.cpp+37-0

@@ -0,0 +1,37 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file apply_adadelta.cpp
	17	+ * \brief ApplyAdadelta kernel entry (arch35)
	18	+ *
	19	+ * Template parameters (matching apply_adadelta_tiling_key.h):
	20	+ * - D_T_X: Data type, from ASCENDC_TPL_DATATYPE_DECL
	21	+ * - BUFFER_MODE: Buffer mode (0=single, 1=double), from ASCENDC_TPL_UINT_DECL
	22	+ */
	23	+
	24	+#include "apply_adadelta.h"
	25	+
	26	+template <typename D_T_X, int BUFFER_MODE>
	27	+__global__ __aicore__ void apply_adadelta(
	28	+ GM_ADDR var, GM_ADDR accum, GM_ADDR accumUpdate, GM_ADDR grad,
	29	+ GM_ADDR varOut, GM_ADDR accumOut, GM_ADDR accumUpdateOut,
	30	+ GM_ADDR workspace, GM_ADDR tiling)
	31	+{
	32	+ REGISTER_TILING_DEFAULT(ApplyAdadeltaTilingData);
	33	+ GET_TILING_DATA_WITH_STRUCT(ApplyAdadeltaTilingData, tilingData, tiling);
	34	+ NsApplyAdadelta::ApplyAdadelta<D_T_X, BUFFER_MODE> op;
	35	+ op.Init(var, accum, accumUpdate, grad, varOut, accumOut, accumUpdateOut, &tilingData);
	36	+ op.Process();
	37	+}

Aexperimental/optim/apply_adadelta/op_kernel/apply_adadelta.h+306-0

@@ -0,0 +1,306 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file apply_adadelta.h
	17	+ * \brief ApplyAdadelta kernel class definition (arch35)
	18	+ *
	19	+ * Template parameters:
	20	+ * - T: Data type (float / half)
	21	+ * - BUFFER_MODE: Buffer mode (0=single buffer, 1=double buffer)
	22	+ *
	23	+ * Adadelta update formula (4 steps):
	24	+ * Step1: accum_new = rho * accum + (1-rho) * grad^2
	25	+ * Step2: update = sqrt(accum_update + eps) / sqrt(accum_new + eps) * grad
	26	+ * Step3: var_new = var - lr * update
	27	+ * Step4: accum_update_new = rho * accum_update + (1-rho) * update^2
	28	+ */
	29	+#ifndef APPLY_ADADELTA_H
	30	+#define APPLY_ADADELTA_H
	31	+
	32	+#include "kernel_operator.h"
	33	+#include "kernel_tiling/kernel_tiling.h"
	34	+#include "apply_adadelta_tiling_data.h"
	35	+#include "apply_adadelta_tiling_key.h"
	36	+
	37	+namespace NsApplyAdadelta {
	38	+
	39	+using namespace AscendC;
	40	+
	41	+template <typename T, int BUFFER_MODE>
	42	+class ApplyAdadelta {
	43	+ static constexpr int32_t BUFFER_NUM = BUFFER_MODE ? 2 : 1;
	44	+ static constexpr bool IS_FP16 = std::is_same_v<T, half>;
	45	+
	46	+public:
	47	+ __aicore__ inline ApplyAdadelta() {}
	48	+
	49	+ __aicore__ inline void Init(
	50	+ GM_ADDR var, GM_ADDR accum, GM_ADDR accumUpdate, GM_ADDR grad,
	51	+ GM_ADDR varOut, GM_ADDR accumOut, GM_ADDR accumUpdateOut,
	52	+ const ApplyAdadeltaTilingData* tilingData);
	53	+
	54	+ __aicore__ inline void Process();
	55	+
	56	+private:
	57	+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
	58	+ __aicore__ inline void Compute(int64_t currentNum);
	59	+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
	60	+
	61	+private:
	62	+ TPipe pipe;
	63	+ // 4 input queues
	64	+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueVar;
	65	+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueAccum;
	66	+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueAccumUpdate;
	67	+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueGrad;
	68	+ // 3 output queues
	69	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueVar;
	70	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueAccum;
	71	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueAccumUpdate;
	72	+ // FP32 path: 2 temporary buffers for intermediate computation
	73	+ TBuf<QuePosition::VECCALC> tmpBuf1; // tmp1: grad^2 / den / updateSq etc.
	74	+ TBuf<QuePosition::VECCALC> tmpBuf2; // tmp2: num / update / lr*update etc.
	75	+ // FP16 path: 6 fp32 TBuf for Cast workspace + computation
	76	+ // (4 for Cast-up workspace: var/accum/accumUpdate/grad fp32 copies,
	77	+ // 2 for intermediate computation: tmp1/tmp2)
	78	+ TBuf<QuePosition::VECCALC> fp32VarBuf;
	79	+ TBuf<QuePosition::VECCALC> fp32AccumBuf;
	80	+ TBuf<QuePosition::VECCALC> fp32AccumUpdateBuf;
	81	+ TBuf<QuePosition::VECCALC> fp32GradBuf;
	82	+
	83	+ GlobalTensor<T> varGm, accumGm, accumUpdateGm, gradGm;
	84	+ GlobalTensor<T> varOutGm, accumOutGm, accumUpdateOutGm;
	85	+
	86	+ int64_t blockLength_ = 0;
	87	+ int64_t ubLength_ = 0;
	88	+ float lr_ = 0.0f;
	89	+ float rho_ = 0.0f;
	90	+ float oneMinusRho_ = 1.0f;
	91	+ float eps_ = 1e-6f;
	92	+};
	93	+
	94	+template <typename T, int BUFFER_MODE>
	95	+__aicore__ inline void ApplyAdadelta<T, BUFFER_MODE>::Init(
	96	+ GM_ADDR var, GM_ADDR accum, GM_ADDR accumUpdate, GM_ADDR grad,
	97	+ GM_ADDR varOut, GM_ADDR accumOut, GM_ADDR accumUpdateOut,
	98	+ const ApplyAdadeltaTilingData* tilingData)
	99	+{
	100	+ // Compute per-core element range
	101	+ int64_t remain = tilingData->totalNum - tilingData->blockFactor * GetBlockIdx();
	102	+ blockLength_ = (remain > tilingData->blockFactor) ? tilingData->blockFactor : remain;
	103	+ ubLength_ = tilingData->ubFactor;
	104	+
	105	+ // Read scalar parameters from TilingData
	106	+ lr_ = tilingData->lr;
	107	+ rho_ = tilingData->rho;
	108	+ eps_ = tilingData->epsilon;
	109	+ oneMinusRho_ = tilingData->oneMinusRho;
	110	+
	111	+ // Set up GM tensors with per-core offset
	112	+ int64_t off = tilingData->blockFactor * GetBlockIdx();
	113	+ varGm.SetGlobalBuffer((__gm__ T*)var + off, blockLength_);
	114	+ accumGm.SetGlobalBuffer((__gm__ T*)accum + off, blockLength_);
	115	+ accumUpdateGm.SetGlobalBuffer((__gm__ T*)accumUpdate + off, blockLength_);
	116	+ gradGm.SetGlobalBuffer((__gm__ T*)grad + off, blockLength_);
	117	+ varOutGm.SetGlobalBuffer((__gm__ T*)varOut + off, blockLength_);
	118	+ accumOutGm.SetGlobalBuffer((__gm__ T*)accumOut + off, blockLength_);
	119	+ accumUpdateOutGm.SetGlobalBuffer((__gm__ T*)accumUpdateOut + off, blockLength_);
	120	+
	121	+ // Initialize queues
	122	+ pipe.InitBuffer(inQueVar, BUFFER_NUM, ubLength_ * sizeof(T));
	123	+ pipe.InitBuffer(inQueAccum, BUFFER_NUM, ubLength_ * sizeof(T));
	124	+ pipe.InitBuffer(inQueAccumUpdate, BUFFER_NUM, ubLength_ * sizeof(T));
	125	+ pipe.InitBuffer(inQueGrad, BUFFER_NUM, ubLength_ * sizeof(T));
	126	+ pipe.InitBuffer(outQueVar, BUFFER_NUM, ubLength_ * sizeof(T));
	127	+ pipe.InitBuffer(outQueAccum, BUFFER_NUM, ubLength_ * sizeof(T));
	128	+ pipe.InitBuffer(outQueAccumUpdate, BUFFER_NUM, ubLength_ * sizeof(T));
	129	+
	130	+ // Initialize temporary buffers based on dtype path
	131	+ if constexpr (IS_FP16) {
	132	+ // FP16 path: 4 fp32 Cast workspace + 2 fp32 tmp = 6 fp32 TBuf
	133	+ pipe.InitBuffer(fp32VarBuf, ubLength_ * sizeof(float));
	134	+ pipe.InitBuffer(fp32AccumBuf, ubLength_ * sizeof(float));
	135	+ pipe.InitBuffer(fp32AccumUpdateBuf, ubLength_ * sizeof(float));
	136	+ pipe.InitBuffer(fp32GradBuf, ubLength_ * sizeof(float));
	137	+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(float));
	138	+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(float));
	139	+ } else {
	140	+ // FP32 path: 2 fp32 TBuf for intermediate computation
	141	+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(float));
	142	+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(float));
	143	+ }
	144	+}
	145	+
	146	+template <typename T, int BUFFER_MODE>
	147	+__aicore__ inline void ApplyAdadelta<T, BUFFER_MODE>::CopyIn(int64_t progress, int64_t currentNum)
	148	+{
	149	+ LocalTensor<T> v = inQueVar.template AllocTensor<T>();
	150	+ LocalTensor<T> a = inQueAccum.template AllocTensor<T>();
	151	+ LocalTensor<T> au = inQueAccumUpdate.template AllocTensor<T>();
	152	+ LocalTensor<T> g = inQueGrad.template AllocTensor<T>();
	153	+
	154	+ DataCopyExtParams cp;
	155	+ cp.blockCount = 1;
	156	+ cp.blockLen = currentNum * sizeof(T);
	157	+ cp.srcStride = 0;
	158	+ cp.dstStride = 0;
	159	+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
	160	+ int64_t gmOff = progress * ubLength_;
	161	+ DataCopyPad(v, varGm[gmOff], cp, padParams);
	162	+ DataCopyPad(a, accumGm[gmOff], cp, padParams);
	163	+ DataCopyPad(au, accumUpdateGm[gmOff], cp, padParams);
	164	+ DataCopyPad(g, gradGm[gmOff], cp, padParams);
	165	+
	166	+ inQueVar.EnQue(v);
	167	+ inQueAccum.EnQue(a);
	168	+ inQueAccumUpdate.EnQue(au);
	169	+ inQueGrad.EnQue(g);
	170	+}
	171	+
	172	+template <typename T, int BUFFER_MODE>
	173	+__aicore__ inline void ApplyAdadelta<T, BUFFER_MODE>::Compute(int64_t currentNum)
	174	+{
	175	+ LocalTensor<T> v = inQueVar.template DeQue<T>();
	176	+ LocalTensor<T> a = inQueAccum.template DeQue<T>();
	177	+ LocalTensor<T> au = inQueAccumUpdate.template DeQue<T>();
	178	+ LocalTensor<T> g = inQueGrad.template DeQue<T>();
	179	+
	180	+ LocalTensor<T> vOut = outQueVar.template AllocTensor<T>();
	181	+ LocalTensor<T> aOut = outQueAccum.template AllocTensor<T>();
	182	+ LocalTensor<T> auOut = outQueAccumUpdate.template AllocTensor<T>();
	183	+
	184	+ // Get tmp buffers
	185	+ LocalTensor<float> tmp1 = tmpBuf1.Get<float>();
	186	+ LocalTensor<float> tmp2 = tmpBuf2.Get<float>();
	187	+
	188	+ if constexpr (IS_FP16) {
	chenfeng4月20日你只有这些操作可以用atovoss模板去做，不用自己实现，你现在全部用memory base实现，性能是否可以达标，编译器是否可以做到vf融合？，如果改为atvoss实现，记得ascend-config文件开启vf融合编译选项 Nerddddddddddd 4月21日评论： 4月21日评论： Nerddddddddddd 4月24日评论： 4月24日评论：
	189	+ // FP16 path: Cast up to fp32 -> compute -> Cast down to fp16
	190	+ // Get fp32 workspace buffers
	191	+ LocalTensor<float> vF = fp32VarBuf.Get<float>();
	192	+ LocalTensor<float> aF = fp32AccumBuf.Get<float>();
	193	+ LocalTensor<float> auF = fp32AccumUpdateBuf.Get<float>();
	194	+ LocalTensor<float> gF = fp32GradBuf.Get<float>();
	195	+
	196	+ // Cast up: half -> fp32
	197	+ Cast(vF, v, RoundMode::CAST_NONE, currentNum);
	198	+ Cast(aF, a, RoundMode::CAST_NONE, currentNum);
	199	+ Cast(auF, au, RoundMode::CAST_NONE, currentNum);
	200	+ Cast(gF, g, RoundMode::CAST_NONE, currentNum);
	201	+
	202	+ // Step1: accum_new = accum * rho + grad^2 * (1 - rho)
	203	+ Mul(tmp1, gF, gF, currentNum); // tmp1 = grad^2
	204	+ Muls(tmp1, tmp1, oneMinusRho_, currentNum); // tmp1 = grad^2 * (1-rho)
	205	+ Muls(aF, aF, rho_, currentNum); // aF = accum * rho
	206	+ Add(aF, aF, tmp1, currentNum); // aF = accum_new
	207	+
	208	+ // Step2: update = sqrt(au_old + eps) / sqrt(accum_new + eps) * grad
	209	+ Adds(tmp1, auF, eps_, currentNum); // tmp1 = au_old + eps
	210	+ Sqrt(tmp1, tmp1, currentNum); // tmp1 = sqrt(au_old + eps) = numerator
	211	+ Adds(tmp2, aF, eps_, currentNum); // tmp2 = accum_new + eps
	212	+ Sqrt(tmp2, tmp2, currentNum); // tmp2 = sqrt(accum_new + eps) = denominator
	213	+ Div(tmp1, tmp1, tmp2, currentNum); // tmp1 = ratio
	214	+ Mul(tmp1, tmp1, gF, currentNum); // tmp1 = update
	215	+
	216	+ // Step3: var_new = var - lr * update
	217	+ Muls(tmp2, tmp1, lr_, currentNum); // tmp2 = lr * update
	218	+ Sub(vF, vF, tmp2, currentNum); // vF = var_new
	219	+
	220	+ // Step4: accum_update_new = au_old * rho + update^2 * (1 - rho)
	221	+ Mul(tmp2, tmp1, tmp1, currentNum); // tmp2 = update^2
	222	+ Muls(tmp2, tmp2, oneMinusRho_, currentNum); // tmp2 = update^2 * (1-rho)
	223	+ Muls(auF, auF, rho_, currentNum); // auF = au_old * rho
	224	+ Add(auF, auF, tmp2, currentNum); // auF = accum_update_new
	225	+
	226	+ // Cast down: fp32 -> half
	227	+ Cast(vOut, vF, RoundMode::CAST_RINT, currentNum);
	228	+ Cast(aOut, aF, RoundMode::CAST_RINT, currentNum);
	229	+ Cast(auOut, auF, RoundMode::CAST_RINT, currentNum);
	230	+ } else {
	231	+ // FP32 path: T == float, so v/a/au/g/vOut/aOut/auOut are already
	232	+ // LocalTensor<float>. No Cast needed; operate directly on them.
	233	+
	234	+ // Step1: accum_new = rho * accum + (1-rho) * grad^2
	235	+ Mul(tmp1, g, g, currentNum); // tmp1 = grad^2
	236	+ Muls(tmp1, tmp1, oneMinusRho_, currentNum); // tmp1 = grad^2 * (1-rho)
	237	+ Muls(aOut, a, rho_, currentNum); // aOut = accum * rho
	238	+ Add(aOut, aOut, tmp1, currentNum); // aOut = accumrho + grad^2(1-rho)
	239	+
	240	+ // Step2: update = sqrt(accum_update + eps) / sqrt(accum_new + eps) * grad
	241	+ Adds(tmp1, au, eps_, currentNum); // tmp1 = accum_update + eps
	242	+ Sqrt(tmp1, tmp1, currentNum); // tmp1 = sqrt(accum_update + eps) = numerator
	243	+ Adds(tmp2, aOut, eps_, currentNum); // tmp2 = accum_new + eps
	244	+ Sqrt(tmp2, tmp2, currentNum); // tmp2 = sqrt(accum_new + eps) = denominator
	245	+ Div(tmp1, tmp1, tmp2, currentNum); // tmp1 = ratio = num / den
	246	+ Mul(tmp1, tmp1, g, currentNum); // tmp1 = update = ratio * grad
	247	+
	248	+ // Step3: var_new = var - lr * update
	249	+ Muls(tmp2, tmp1, lr_, currentNum); // tmp2 = lr * update
	250	+ Sub(vOut, v, tmp2, currentNum); // vOut = var - lr * update
	251	+
	252	+ // Step4: accum_update_new = rho * accum_update + (1-rho) * update^2
	253	+ Mul(tmp2, tmp1, tmp1, currentNum); // tmp2 = update^2
	254	+ Muls(tmp2, tmp2, oneMinusRho_, currentNum); // tmp2 = update^2 * (1-rho)
	255	+ Muls(auOut, au, rho_, currentNum); // auOut = accum_update * rho
	256	+ Add(auOut, auOut, tmp2, currentNum); // auOut = aurho + update^2(1-rho)
	257	+ }
	258	+
	259	+ // Free input tensors
	260	+ inQueVar.FreeTensor(v);
	261	+ inQueAccum.FreeTensor(a);
	262	+ inQueAccumUpdate.FreeTensor(au);
	263	+ inQueGrad.FreeTensor(g);
	264	+
	265	+ // Enqueue output tensors
	266	+ outQueVar.EnQue(vOut);
	267	+ outQueAccum.EnQue(aOut);
	268	+ outQueAccumUpdate.EnQue(auOut);
	269	+}
	270	+
	271	+template <typename T, int BUFFER_MODE>
	272	+__aicore__ inline void ApplyAdadelta<T, BUFFER_MODE>::CopyOut(int64_t progress, int64_t currentNum)
	273	+{
	274	+ LocalTensor<T> vOut = outQueVar.template DeQue<T>();
	275	+ LocalTensor<T> aOut = outQueAccum.template DeQue<T>();
	276	+ LocalTensor<T> auOut = outQueAccumUpdate.template DeQue<T>();
	277	+
	278	+ DataCopyExtParams cp;
	279	+ cp.blockCount = 1;
	280	+ cp.blockLen = currentNum * sizeof(T);
	281	+ cp.srcStride = 0;
	282	+ cp.dstStride = 0;
	chenfeng4月20日【阻塞级】Process() 未实现真正的双缓冲流水线当前代码串行执行 `CopyIn -> Compute -> CopyOut`，未发挥 BUFFER_MODE=2 的性能优势。建议改为 3-stage 流水线实现： if (loopCount >= 3) { CopyIn(0, ubLength_); CopyIn(1, ubLength_); Compute(0, ubLength_); for (int64_t i = 0; i < loopCount - 2; i++) { int64_t currentNum = (i == loopCount - 3) ? (blockLength_ - ubLength_ * (i + 2)) : ubLength_; CopyOut(i, ubLength_); CopyIn(i + 2, currentNum); Compute(i + 1, ubLength_); } CopyOut(loopCount - 2, ubLength_); Compute(loopCount - 1, currentNum); CopyOut(loopCount - 1, currentNum); } else { // fallback: serial execution for (int64_t i = 0; i < loopCount; i++) { ... } } 参考：`asc-devkit/examples/00_introduction/01_add/basic_api_memory_allocator_add/` 的双缓冲流水线示例。 Nerddddddddddd 4月24日评论： 4月24日评论：
	283	+ int64_t gmOff = progress * ubLength_;
	284	+ DataCopyPad(varOutGm[gmOff], vOut, cp);
	285	+ DataCopyPad(accumOutGm[gmOff], aOut, cp);
	286	+ DataCopyPad(accumUpdateOutGm[gmOff], auOut, cp);
	287	+
	288	+ outQueVar.FreeTensor(vOut);
	289	+ outQueAccum.FreeTensor(aOut);
	290	+ outQueAccumUpdate.FreeTensor(auOut);
	291	+}
	292	+
	293	+template <typename T, int BUFFER_MODE>
	294	+__aicore__ inline void ApplyAdadelta<T, BUFFER_MODE>::Process()
	295	+{
	296	+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
	297	+ for (int64_t i = 0; i < loopCount; i++) {
	298	+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
	299	+ CopyIn(i, currentNum);
	300	+ Compute(currentNum);
	301	+ CopyOut(i, currentNum);
	302	+ }
	303	+}
	304	+
	305	+} // namespace NsApplyAdadelta
	306	+#endif // APPLY_ADADELTA_H

Aexperimental/optim/apply_adadelta/op_kernel/apply_adadelta_tiling_data.h+33-0

@@ -0,0 +1,33 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file apply_adadelta_tiling_data.h
	17	+ * \brief ApplyAdadelta TilingData structure
	18	+ */
	19	+
	20	+#ifndef _APPLY_ADADELTA_TILING_DATA_H_
	21	+#define _APPLY_ADADELTA_TILING_DATA_H_
	22	+
	23	+struct ApplyAdadeltaTilingData {
	24	+ int64_t totalNum = 0; // Total number of elements
	25	+ int64_t blockFactor = 0; // Number of elements per core (aligned to UB block)
	26	+ int64_t ubFactor = 0; // Number of elements per UB loop iteration (aligned to UB block)
	27	+ float lr = 0.0f; // Learning rate (from aclScalar, stored as fp32)
	28	+ float rho = 0.0f; // Decay coefficient
	29	+ float epsilon = 0.0f; // Numerical stability constant
	30	+ float oneMinusRho = 0.0f; // 1 - rho, pre-computed to avoid Kernel redundant computation
	31	+};
	32	+
	33	+#endif

Aexperimental/optim/apply_adadelta/op_kernel/apply_adadelta_tiling_key.h+45-0

@@ -0,0 +1,45 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file apply_adadelta_tiling_key.h
	17	+ * \brief ApplyAdadelta TilingKey template parameter declaration
	18	+ *
	19	+ * Template parameters:
	20	+ * - D_T_X: Data type (C_DT_FLOAT, C_DT_FLOAT16)
	21	+ * - BUFFER_MODE: Buffer mode (0=single buffer, 1=double buffer)
	22	+ */
	23	+
	24	+#ifndef __APPLY_ADADELTA_TILING_KEY_H__
	25	+#define __APPLY_ADADELTA_TILING_KEY_H__
	26	+
	27	+#include "ascendc/host_api/tiling/template_argument.h"
	28	+
	29	+ASCENDC_TPL_ARGS_DECL(ApplyAdadelta,
	30	+ ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, ASCENDC_TPL_INPUT(0)),
	31	+ ASCENDC_TPL_UINT_DECL(BUFFER_MODE, 8, ASCENDC_TPL_UI_LIST, 0, 1)
	32	+);
	33	+
	34	+ASCENDC_TPL_SEL(
	35	+ ASCENDC_TPL_ARGS_SEL(
	36	+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT),
	37	+ ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)
	38	+ ),
	39	+ ASCENDC_TPL_ARGS_SEL(
	40	+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16),
	41	+ ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)
	42	+ )
	43	+);
	44	+
	45	+#endif

Aexperimental/optim/apply_adadelta/tests/.gitkeep+0-0

The file is empty