Star945
1861
代码
代码
Issues223
Pull Requests333
Actions
项目讨论
Wiki
分析
Star945
1861
已合并
将 experimental 下误用的 arch32 修正为 arch22 #8954
clinglai0517创建于 25 天前
将 experimental 下误用的 arch32 修正为 arch22 #8954
已合并
clinglai0517创建于 25 天前
从clinglai0517/ops-nnarch32-fix-pooling合入到cann/ops-nnmaster
共 36 个文件变更+72-70
Mexperimental/activation/celu_v3/README.md+4-4
@@ -46,7 +46,7 @@ aclnnStatus aclnnCeluV3(
46 46 
47| 芯片型号 | 架构 | 编译宏 |47| 芯片型号 | 架构 | 编译宏 |
48|---------|------|--------|48|---------|------|--------|
49-| Ascend 910B | A2 (arch32) | DAV_2201 |49+| Ascend 910B | A2 (arch22) | DAV_2201 |
50 50 
51## 精度标准51## 精度标准
52 52 
@@ -80,12 +80,12 @@ ops/celu_v3/
80│ ├── CMakeLists.txt # Host 侧构建80│ ├── CMakeLists.txt # Host 侧构建
81│ ├── celu_v3_def.cpp # 算子定义（IR、Attr、AICore 配置）81│ ├── celu_v3_def.cpp # 算子定义（IR、Attr、AICore 配置）
82│ ├── celu_v3_infershape.cpp # Shape 推导82│ ├── celu_v3_infershape.cpp # Shape 推导
83-│ └── arch32/83+│ └── arch22/
84│ └── celu_v3_tiling.cpp # Tiling 计算（多核切分 + UB 切分）84│ └── celu_v3_tiling.cpp # Tiling 计算（多核切分 + UB 切分）
85├── op_kernel/85├── op_kernel/
86│ ├── CMakeLists.txt # Kernel 侧构建86│ ├── CMakeLists.txt # Kernel 侧构建
87-│ ├── celu_v3_arch32.cpp # Kernel 入口87+│ ├── celu_v3_arch22.cpp # Kernel 入口
88-│ └── arch32/88+│ └── arch22/
89│ ├── celu_v3.h # Kernel 实现（CopyIn/Compute/CopyOut）89│ ├── celu_v3.h # Kernel 实现（CopyIn/Compute/CopyOut）
90│ ├── celu_v3_tiling_data.h # TilingData 结构定义90│ ├── celu_v3_tiling_data.h # TilingData 结构定义
91│ └── celu_v3_tiling_key.h # TilingKey 模板参数定义91│ └── celu_v3_tiling_key.h # TilingKey 模板参数定义
Mexperimental/activation/celu_v3/op_host/celu_v3_def.cpp+1-1
@@ -22,7 +22,7 @@
22 * Input: self (Tensor, float32/float16/bfloat16)22 * Input: self (Tensor, float32/float16/bfloat16)
23 * Attr: alpha (Float, default=1.0)23 * Attr: alpha (Float, default=1.0)
24 * Output: out (Tensor, same dtype/shape as self)24 * Output: out (Tensor, same dtype/shape as self)
25- * Target: Ascend910B (arch32)25+ * Target: Ascend910B (arch22)
26 */26 */
27#include "register/op_def_registry.h"27#include "register/op_def_registry.h"
28 28 
Mexperimental/activation/celu_v3/op_host/celu_v3_tiling.cpp+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file celu_v3_tiling.cpp17 * \file celu_v3_tiling.cpp
18- * \brief CeluV3 tiling implementation (arch32)18+ * \brief CeluV3 tiling implementation (arch22)
19 *19 *
20 * Tiling strategy:20 * Tiling strategy:
21 * 1. Multi-core: divide total elements evenly across AI Cores21 * 1. Multi-core: divide total elements evenly across AI Cores
Mexperimental/activation/celu_v3/op_kernel/celu_v3.cpp+2-2
@@ -14,8 +14,8 @@
14 */14 */
15 15 
16/**16/**
17- * \file celu_v3_arch32.cpp17+ * \file celu_v3_arch22.cpp
18- * \brief CeluV3 kernel entry point (arch32)18+ * \brief CeluV3 kernel entry point (arch22)
19 *19 *
20 * Template parameter D_T_X maps to data type:20 * Template parameter D_T_X maps to data type:
21 * - float: TilingKey 0 (direct fp32 computation)21 * - float: TilingKey 0 (direct fp32 computation)
Mexperimental/activation/celu_v3/op_kernel/celu_v3.h+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file celu_v3.h17 * \file celu_v3.h
18- * \brief CeluV3 kernel class definition (arch32)18+ * \brief CeluV3 kernel class definition (arch22)
19 *19 *
20 * CELU(x) = max(0, x) + min(0, alpha * (exp(x/alpha) - 1))20 * CELU(x) = max(0, x) + min(0, alpha * (exp(x/alpha) - 1))
21 *21 *
Mexperimental/activation/hard_shrink_grad/README.md+5-5
@@ -23,7 +23,7 @@ output_i = grad_output_i, if |self_i| > lambd
23| 算子类别 | Elementwise（逐元素反向梯度算子） |23| 算子类别 | Elementwise（逐元素反向梯度算子） |
24| 调用方式 | ACLNN (`aclnnHardShrinkGrad`) |24| 调用方式 | ACLNN (`aclnnHardShrinkGrad`) |
25| 目标芯片 | Ascend910B (910B3) |25| 目标芯片 | Ascend910B (910B3) |
26-| 目标架构 | arch32 (DAV_2201) |26+| 目标架构 | arch22 (DAV_2201) |
27 27 
28### 输入28### 输入
29 29 
@@ -122,12 +122,12 @@ hard_shrink_grad/
122│ ├── CMakeLists.txt # Host 构建配置122│ ├── CMakeLists.txt # Host 构建配置
123│ ├── hard_shrink_grad_def.cpp # 算子原型注册123│ ├── hard_shrink_grad_def.cpp # 算子原型注册
124│ ├── hard_shrink_grad_infershape.cpp # Shape 推导（output shape = input shape）124│ ├── hard_shrink_grad_infershape.cpp # Shape 推导（output shape = input shape）
125-│ └── arch32/ # Ascend910B (DAV_2201)125+│ └── arch22/ # Ascend910B (DAV_2201)
126│ └── hard_shrink_grad_tiling.cpp # Tiling 实现（多核切分、UB 切分、TilingKey 选择）126│ └── hard_shrink_grad_tiling.cpp # Tiling 实现（多核切分、UB 切分、TilingKey 选择）
127├── op_kernel/ # Kernel 侧实现127├── op_kernel/ # Kernel 侧实现
128│ ├── CMakeLists.txt # Kernel 构建配置128│ ├── CMakeLists.txt # Kernel 构建配置
129-│ ├── hard_shrink_grad_arch32.cpp # Kernel 入口（模板实例化，双类分发）129+│ ├── hard_shrink_grad_arch22.cpp # Kernel 入口（模板实例化，双类分发）
130-│ └── arch32/ # Ascend910B (DAV_2201)130+│ └── arch22/ # Ascend910B (DAV_2201)
131│ ├── hard_shrink_grad.h # Kernel 类定义与实现131│ ├── hard_shrink_grad.h # Kernel 类定义与实现
132│ ├── hard_shrink_grad_tiling_data.h # TilingData 结构体132│ ├── hard_shrink_grad_tiling_data.h # TilingData 结构体
133│ └── hard_shrink_grad_tiling_key.h # TilingKey 定义（6 种模板组合）133│ └── hard_shrink_grad_tiling_key.h # TilingKey 定义（6 种模板组合）
@@ -176,7 +176,7 @@ hard_shrink_grad/
176 176 
177### 关键技术决策177### 关键技术决策
178 178 
179-1. **arch32 上 Compare API 对 fp16/bf16 不可靠**：统一 Cast 到 fp32 计算，确保精度达标179+1. **arch22 上 Compare API 对 fp16/bf16 不可靠**：统一 Cast 到 fp32 计算，确保精度达标
1802. **bf16 不支持 Abs/Compare/Select API**：复用 fp16 的 CastFp32 路径1802. **bf16 不支持 Abs/Compare/Select API**：复用 fp16 的 CastFp32 路径
1813. **使用 Compare(tensor-tensor) 替代 CompareScalar**：lambd 通过 Duplicate 广播到 tensor，更通用可靠1813. **使用 Compare(tensor-tensor) 替代 CompareScalar**：lambd 通过 Duplicate 广播到 tensor，更通用可靠
182 182 
Mexperimental/activation/hard_shrink_grad/op_host/hard_shrink_grad_tiling.cpp+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file hard_shrink_grad_tiling.cpp17 * \file hard_shrink_grad_tiling.cpp
18- * \brief HardShrinkGrad Tiling implementation for arch32 (Ascend910B)18+ * \brief HardShrinkGrad Tiling implementation for arch22 (Ascend910B)
19 *19 *
20 * Tiling flow:20 * Tiling flow:
21 * 1. Get platform info (coreNum, ubSize)21 * 1. Get platform info (coreNum, ubSize)
Mexperimental/activation/hard_shrink_grad/op_kernel/hard_shrink_grad.cpp+2-2
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file hard_shrink_grad.cpp17 * \file hard_shrink_grad.cpp
18- * \brief HardShrinkGrad kernel entry point (arch32 architecture)18+ * \brief HardShrinkGrad kernel entry point (arch22 architecture)
19 *19 *
20 * Template parameters (matching ASCENDC_TPL_ARGS_DECL in hard_shrink_grad_tiling_key.h):20 * Template parameters (matching ASCENDC_TPL_ARGS_DECL in hard_shrink_grad_tiling_key.h):
21 * - D_T: Data type, from ASCENDC_TPL_DATATYPE_DECL21 * - D_T: Data type, from ASCENDC_TPL_DATATYPE_DECL
@@ -36,7 +36,7 @@ __global__ __aicore__ void hard_shrink_grad(GM_ADDR grad_output, GM_ADDR self, G
36 GET_TILING_DATA_WITH_STRUCT(HardShrinkGradTilingData, tilingData, tiling);36 GET_TILING_DATA_WITH_STRUCT(HardShrinkGradTilingData, tilingData, tiling);
37 37 
38 if constexpr (sizeof(D_T) == 2) {38 if constexpr (sizeof(D_T) == 2) {
39- // fp16/bf16: Cast to fp32 for compute to avoid incorrect Compare on arch3239+ // fp16/bf16: Cast to fp32 for compute to avoid incorrect Compare on arch22
40 NsHardShrinkGrad::HardShrinkGradCastFp32<D_T, BUFFER_MODE> op;40 NsHardShrinkGrad::HardShrinkGradCastFp32<D_T, BUFFER_MODE> op;
41 op.Init(grad_output, self, output, &tilingData);41 op.Init(grad_output, self, output, &tilingData);
42 op.Process();42 op.Process();
Mexperimental/activation/hard_shrink_grad/op_kernel/hard_shrink_grad.h+3-3
@@ -15,14 +15,14 @@
15 15 
16/**16/**
17 * \file hard_shrink_grad.h17 * \file hard_shrink_grad.h
18- * \brief HardShrinkGrad kernel class definitions (arch32 architecture)18+ * \brief HardShrinkGrad kernel class definitions (arch22 architecture)
19 *19 *
20 * Computation: output_i = grad_output_i if |self_i| > lambd, else 020 * Computation: output_i = grad_output_i if |self_i| > lambd, else 0
21 *21 *
22 * Two kernel implementations:22 * Two kernel implementations:
23 * - HardShrinkGradDirect<T, BUFFER_MODE>: for T=float, computes directly in T23 * - HardShrinkGradDirect<T, BUFFER_MODE>: for T=float, computes directly in T
24 * - HardShrinkGradCastFp32<T, BUFFER_MODE>: for T=half/bfloat16_t, casts to fp3224 * - HardShrinkGradCastFp32<T, BUFFER_MODE>: for T=half/bfloat16_t, casts to fp32
25- * for Abs/Compare/Select to avoid incorrect Compare results on arch3225+ * for Abs/Compare/Select to avoid incorrect Compare results on arch22
26 *26 *
27 * Data flow (Direct, fp32):27 * Data flow (Direct, fp32):
28 * CopyIn: GM(grad_output, self) -> UB28 * CopyIn: GM(grad_output, self) -> UB
@@ -200,7 +200,7 @@ __aicore__ inline void HardShrinkGradDirect<T, BUFFER_MODE>::Process()
200 200 
201// ============================================================================201// ============================================================================
202// HardShrinkGradCastFp32: Cast to fp32 for compute (for fp16/bf16)202// HardShrinkGradCastFp32: Cast to fp32 for compute (for fp16/bf16)
203-// On arch32, Compare API produces incorrect results on half type.203+// On arch22, Compare API produces incorrect results on half type.
204// Solution: Cast fp16->fp32, compute in fp32, Cast fp32->fp16.204// Solution: Cast fp16->fp32, compute in fp32, Cast fp32->fp16.
205// ============================================================================205// ============================================================================
206template <typename T, int BUFFER_MODE>206template <typename T, int BUFFER_MODE>
Mexperimental/activation/hard_sigmoid_grad_v2/README.md+4-4
@@ -31,7 +31,7 @@ grad_input = 0 其他情况
31## 目标芯片31## 目标芯片
32 32 
33- **芯片型号**：Ascend910B33- **芯片型号**：Ascend910B
34-- **架构**：arch3234+- **架构**：arch22
35- **CANN 版本**：9.0.035- **CANN 版本**：9.0.0
36 36 
37## 编译与安装37## 编译与安装
@@ -115,8 +115,8 @@ hard_sigmoid_grad_v2/
115├── README.md # 本文件115├── README.md # 本文件
116├── op_kernel/116├── op_kernel/
117│ ├── CMakeLists.txt # Kernel 编译配置117│ ├── CMakeLists.txt # Kernel 编译配置
118-│ ├── hard_sigmoid_grad_v2_arch32.cpp # Kernel 入口（模板实例化）118+│ ├── hard_sigmoid_grad_v2_arch22.cpp # Kernel 入口（模板实例化）
119-│ └── arch32/119+│ └── arch22/
120│ ├── hard_sigmoid_grad_v2.h # Kernel 核心实现120│ ├── hard_sigmoid_grad_v2.h # Kernel 核心实现
121│ ├── hard_sigmoid_grad_v2_tiling_data.h # Tiling 数据结构定义121│ ├── hard_sigmoid_grad_v2_tiling_data.h # Tiling 数据结构定义
122│ └── hard_sigmoid_grad_v2_tiling_key.h # TilingKey 定义（6 个 Key）122│ └── hard_sigmoid_grad_v2_tiling_key.h # TilingKey 定义（6 个 Key）
@@ -124,7 +124,7 @@ hard_sigmoid_grad_v2/
124│ ├── CMakeLists.txt # Host 编译配置124│ ├── CMakeLists.txt # Host 编译配置
125│ ├── hard_sigmoid_grad_v2_def.cpp # 算子注册与 dtype/format 定义125│ ├── hard_sigmoid_grad_v2_def.cpp # 算子注册与 dtype/format 定义
126│ ├── hard_sigmoid_grad_v2_infershape.cpp # Shape 推导（elementwise）126│ ├── hard_sigmoid_grad_v2_infershape.cpp # Shape 推导（elementwise）
127-│ └── arch32/127+│ └── arch22/
128│ └── hard_sigmoid_grad_v2_tiling.cpp # Host 端 Tiling 策略实现128│ └── hard_sigmoid_grad_v2_tiling.cpp # Host 端 Tiling 策略实现
129├── docs/129├── docs/
130│ ├── DEVELOPMENT_LOG.md # 开发日志130│ ├── DEVELOPMENT_LOG.md # 开发日志
Mexperimental/activation/hard_sigmoid_grad_v2/op_host/hard_sigmoid_grad_v2_def.cpp+1-1
@@ -23,7 +23,7 @@
23 * Output:23 * Output:
24 * - grad_input: computed input gradient (same dtype as input)24 * - grad_input: computed input gradient (same dtype as input)
25 *25 *
26- * Target: Ascend910B (arch32, DAV_2201)26+ * Target: Ascend910B (arch22, DAV_2201)
27 */27 */
28#include "register/op_def_registry.h"28#include "register/op_def_registry.h"
29 29 
Mexperimental/activation/hard_sigmoid_grad_v2/op_host/hard_sigmoid_grad_v2_tiling.cpp+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file hard_sigmoid_grad_v2_tiling.cpp17 * \file hard_sigmoid_grad_v2_tiling.cpp
18- * \brief HardSigmoidGradV2 tiling implementation (arch32, Ascend910B)18+ * \brief HardSigmoidGradV2 tiling implementation (arch22, Ascend910B)
19 *19 *
20 * Tiling strategy:20 * Tiling strategy:
21 * - Multi-core: split total elements evenly across AI Cores21 * - Multi-core: split total elements evenly across AI Cores
Mexperimental/activation/hard_sigmoid_grad_v2/op_kernel/hard_sigmoid_grad_v2.cpp+2-2
@@ -14,8 +14,8 @@
14 */14 */
15 15 
16/**16/**
17- * \file hard_sigmoid_grad_v2_arch32.cpp17+ * \file hard_sigmoid_grad_v2_arch22.cpp
18- * \brief HardSigmoidGradV2 kernel entry (arch32)18+ * \brief HardSigmoidGradV2 kernel entry (arch22)
19 *19 *
20 * Template parameters (matching hard_sigmoid_grad_v2_tiling_key.h):20 * Template parameters (matching hard_sigmoid_grad_v2_tiling_key.h):
21 * - D_T_X: Data type, from ASCENDC_TPL_DATATYPE_DECL21 * - D_T_X: Data type, from ASCENDC_TPL_DATATYPE_DECL
Mexperimental/activation/hard_sigmoid_grad_v2/op_kernel/hard_sigmoid_grad_v2.h+2-2
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file hard_sigmoid_grad_v2.h17 * \file hard_sigmoid_grad_v2.h
18- * \brief HardSigmoidGradV2 kernel class definition (arch32)18+ * \brief HardSigmoidGradV2 kernel class definition (arch22)
19 *19 *
20 * Template parameters:20 * Template parameters:
21 * - T: Data type (half/float/bfloat16_t)21 * - T: Data type (half/float/bfloat16_t)
@@ -25,7 +25,7 @@
25 * grad_input = grad_output * ((self > -3) & (self < 3)) * (1/6)25 * grad_input = grad_output * ((self > -3) & (self < 3)) * (1/6)
26 *26 *
27 * For half/float: direct Compares + And + Muls + Select27 * For half/float: direct Compares + And + Muls + Select
28- * For bfloat16_t: Cast to float, compute, Cast back (arch32 does not support bf16 in Compare/Muls/Select)28+ * For bfloat16_t: Cast to float, compute, Cast back (arch22 does not support bf16 in Compare/Muls/Select)
29 */29 */
30#ifndef HARD_SIGMOID_GRAD_V2_H30#ifndef HARD_SIGMOID_GRAD_V2_H
31#define HARD_SIGMOID_GRAD_V2_H31#define HARD_SIGMOID_GRAD_V2_H
Mexperimental/activation/hard_sigmoid_grad_v3/examples/test_aclnn_hard_sigmoid_grad_v3.cpp+9-7
@@ -170,13 +170,15 @@ int main()
170 goto CLEANUP;170 goto CLEANUP;
171 }171 }
172 172 
173- const std::vector<float> expected = {0.f, 2.f / 6.f, 3.f / 6.f, 4.f / 6.f, 5.f / 6.f, 1.f, 0.f, 0.f};173+ {
174- for (size_t i = 0; i < expected.size(); ++i) {174+ const std::vector<float> expected = {0.f, 2.f / 6.f, 3.f / 6.f, 4.f / 6.f, 5.f / 6.f, 1.f, 0.f, 0.f};
175- std::printf("out[%zu] = %.6f\n", i, out_host[i]);175+ for (size_t i = 0; i < expected.size(); ++i) {
176- if (std::abs(out_host[i] - expected[i]) > 1e-5f) {176+ std::printf("out[%zu] = %.6f\n", i, out_host[i]);
177- std::fprintf(stderr, "mismatch at %zu: got %.6f expected %.6f\n", i, out_host[i], expected[i]);177+ if (std::abs(out_host[i] - expected[i]) > 1e-5f) {
178- exitCode = 1;178+ std::fprintf(stderr, "mismatch at %zu: got %.6f expected %.6f\n", i, out_host[i], expected[i]);
179- goto CLEANUP;179+ exitCode = 1;
180+ goto CLEANUP;
181+ }
180 }182 }
181 }183 }
182 184 
Mexperimental/activation/hard_sigmoid_grad_v3/op_host/hard_sigmoid_grad_v3_def.cpp+1-1
@@ -23,7 +23,7 @@
23 * Output:23 * Output:
24 * - grad_input: computed input gradient (same dtype as input)24 * - grad_input: computed input gradient (same dtype as input)
25 *25 *
26- * Target: Ascend910B (arch32, DAV_2201)26+ * Target: Ascend910B (arch22, DAV_2201)
27 */27 */
28#include "register/op_def_registry.h"28#include "register/op_def_registry.h"
29 29 
Mexperimental/activation/hard_sigmoid_grad_v3/op_host/hard_sigmoid_grad_v3_tiling.cpp+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file hard_sigmoid_grad_v3_tiling.cpp17 * \file hard_sigmoid_grad_v3_tiling.cpp
18- * \brief HardSigmoidGradV3 tiling implementation (arch32, Ascend910B)18+ * \brief HardSigmoidGradV3 tiling implementation (arch22, Ascend910B)
19 *19 *
20 * Tiling strategy:20 * Tiling strategy:
21 * - Multi-core: split total elements evenly across AI Cores21 * - Multi-core: split total elements evenly across AI Cores
Mexperimental/activation/hard_sigmoid_grad_v3/op_kernel/hard_sigmoid_grad_v3.cpp+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file hard_sigmoid_grad_v3.cpp17 * \file hard_sigmoid_grad_v3.cpp
18- * \brief HardSigmoidGradV3 kernel entry (arch32)18+ * \brief HardSigmoidGradV3 kernel entry (arch22)
19 *19 *
20 * Template parameters (matching hard_sigmoid_grad_v3_tiling_key.h):20 * Template parameters (matching hard_sigmoid_grad_v3_tiling_key.h):
21 * - D_T_X: Data type, from ASCENDC_TPL_DATATYPE_DECL21 * - D_T_X: Data type, from ASCENDC_TPL_DATATYPE_DECL
Mexperimental/activation/hard_sigmoid_grad_v3/op_kernel/hard_sigmoid_grad_v3.h+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file hard_sigmoid_grad_v3.h17 * \file hard_sigmoid_grad_v3.h
18- * \brief HardSigmoidGradV3 kernel class definition (arch32)18+ * \brief HardSigmoidGradV3 kernel class definition (arch22)
19 */19 */
20#ifndef HARD_SIGMOID_GRAD_V3_H20#ifndef HARD_SIGMOID_GRAD_V3_H
21#define HARD_SIGMOID_GRAD_V3_H21#define HARD_SIGMOID_GRAD_V3_H
Mexperimental/activation/selu/op_host/selu_def.cpp+1-1
@@ -22,7 +22,7 @@
22 * Input: x (Tensor, float32/float16/bfloat16/int32/int8)22 * Input: x (Tensor, float32/float16/bfloat16/int32/int8)
23 * Output: y (Tensor, same dtype/shape as x)23 * Output: y (Tensor, same dtype/shape as x)
24 * No Attr parameters (alpha and scale are fixed constants)24 * No Attr parameters (alpha and scale are fixed constants)
25- * Target: Ascend910B (arch32)25+ * Target: Ascend910B (arch22)
26 */26 */
27#include "register/op_def_registry.h"27#include "register/op_def_registry.h"
28 28 
Mexperimental/activation/selu/op_host/selu_tiling.cpp+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file selu_tiling.cpp17 * \file selu_tiling.cpp
18- * \brief Selu tiling implementation (arch32)18+ * \brief Selu tiling implementation (arch22)
19 *19 *
20 * Tiling strategy:20 * Tiling strategy:
21 * 1. Multi-core: divide total elements evenly across AI Cores21 * 1. Multi-core: divide total elements evenly across AI Cores
Mexperimental/activation/selu/op_kernel/selu.cpp+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file selu.cpp17 * \file selu.cpp
18- * \brief Selu kernel entry point (arch32)18+ * \brief Selu kernel entry point (arch22)
19 *19 *
20 * Template parameter D_T_X maps to data type:20 * Template parameter D_T_X maps to data type:
21 * - float: TilingKey 0 (direct fp32 computation)21 * - float: TilingKey 0 (direct fp32 computation)
Mexperimental/activation/selu/op_kernel/selu.h+3-3
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file selu.h17 * \file selu.h
18- * \brief Selu kernel class definition (arch32)18+ * \brief Selu kernel class definition (arch22)
19 *19 *
20 * SELU(x) = scale * [max(0, x) + min(0, alpha * (exp(x) - 1))]20 * SELU(x) = scale * [max(0, x) + min(0, alpha * (exp(x) - 1))]
21 *21 *
@@ -263,7 +263,7 @@ __aicore__ inline void Selu<T>::ComputeCastFp32(LocalTensor<SrcT>& xLocal, Local
263 LocalTensor<float> tmp2 = tmpBuf2_.template Get<float>();263 LocalTensor<float> tmp2 = tmpBuf2_.template Get<float>();
264 264 
265 // Cast input to fp32265 // Cast input to fp32
266- // Hardware-supported Cast paths on arch32:266+ // Hardware-supported Cast paths on arch22:
267 // int8 -> half -> float (2 steps, no direct int8->float)267 // int8 -> half -> float (2 steps, no direct int8->float)
268 // bfloat16 -> float (direct)268 // bfloat16 -> float (direct)
269 // int32 -> float (direct)269 // int32 -> float (direct)
@@ -293,7 +293,7 @@ __aicore__ inline void Selu<T>::ComputeCastFp32(LocalTensor<SrcT>& xLocal, Local
293 Muls(tmp1, tmp1, SCALE_F32, alignedNum);293 Muls(tmp1, tmp1, SCALE_F32, alignedNum);
294 294 
295 // Cast back to original type295 // Cast back to original type
296- // Hardware-supported output Cast paths on arch32:296+ // Hardware-supported output Cast paths on arch22:
297 // float -> int32 (direct, CAST_TRUNC for truncation toward zero)297 // float -> int32 (direct, CAST_TRUNC for truncation toward zero)
298 // float -> bfloat16 (direct, CAST_ROUND)298 // float -> bfloat16 (direct, CAST_ROUND)
299 // float -> half -> int8 (2 steps, no direct float->int8 or int32->int8)299 // float -> half -> int8 (2 steps, no direct float->int8 or int32->int8)
Mexperimental/activation/softplus_v2/op_kernel/arch35/softplus_v2.h+1-1
@@ -15,7 +15,7 @@
15 15 
16/*!16/*!
17 * \file softplus_v2.h17 * \file softplus_v2.h
18- * \brief SoftplusV2 kernel implementation (arch32)18+ * \brief SoftplusV2 kernel implementation (arch22)
19 *19 *
20 * Iteration 3: Full coverage (FP32 + FP16 + BF16 + edge cases)20 * Iteration 3: Full coverage (FP32 + FP16 + BF16 + edge cases)
21 *21 *
Mexperimental/foreach/foreach_div_scalar/README.md+4-4
@@ -13,7 +13,7 @@ $$y_i = \frac{x_i}{\text{scalar}} \quad (i = 0, 1, \ldots, n-1)$$
13## 支持平台13## 支持平台
14 14 
15- 芯片：Ascend910B（910B1/910B2/910B3/910B4）15- 芯片：Ascend910B（910B1/910B2/910B3/910B4）
16-- 架构：arch32 / DAV_220116+- 架构：arch22 / DAV_2201
17- CANN：9.0.017- CANN：9.0.0
18 18 
19## ACLNN 接口19## ACLNN 接口
@@ -33,15 +33,15 @@ aclnnStatus aclnnForeachDivScalar(
33```text33```text
34ops/foreach_div_scalar/34ops/foreach_div_scalar/
35├── op_kernel/ # Kernel 实现35├── op_kernel/ # Kernel 实现
36-│ ├── foreach_div_scalar_arch32.cpp # 核函数入口36+│ ├── foreach_div_scalar_arch22.cpp # 核函数入口
37-│ └── arch32/37+│ └── arch22/
38│ ├── foreach_div_scalar.h # 核心算子逻辑38│ ├── foreach_div_scalar.h # 核心算子逻辑
39│ ├── foreach_div_scalar_tiling_data.h # TilingData 结构39│ ├── foreach_div_scalar_tiling_data.h # TilingData 结构
40│ └── foreach_div_scalar_tiling_key.h # TilingKey 定义40│ └── foreach_div_scalar_tiling_key.h # TilingKey 定义
41├── op_host/ # Host 侧逻辑41├── op_host/ # Host 侧逻辑
42│ ├── foreach_div_scalar_def.cpp # 算子定义42│ ├── foreach_div_scalar_def.cpp # 算子定义
43│ ├── foreach_div_scalar_infershape.cpp # 形状推导43│ ├── foreach_div_scalar_infershape.cpp # 形状推导
44-│ └── arch32/44+│ └── arch22/
45│ └── foreach_div_scalar_tiling.cpp # Tiling 实现45│ └── foreach_div_scalar_tiling.cpp # Tiling 实现
46├── tests/46├── tests/
47│ ├── ut/ # 单元测试（27 用例）47│ ├── ut/ # 单元测试（27 用例）
Mexperimental/foreach/foreach_div_scalar/op_host/foreach_div_scalar_tiling.cpp+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file foreach_div_scalar_tiling.cpp17 * \file foreach_div_scalar_tiling.cpp
18- * \brief ForeachDivScalar Tiling implementation (arch32 / Ascend910B)18+ * \brief ForeachDivScalar Tiling implementation (arch22 / Ascend910B)
19 *19 *
20 * x is IR index 0 (DYNAMIC), scalar is IR index 1 (REQUIRED).20 * x is IR index 0 (DYNAMIC), scalar is IR index 1 (REQUIRED).
21 * Scalar value is NOT read in tiling; kernel reads it from GM directly.21 * Scalar value is NOT read in tiling; kernel reads it from GM directly.
Mexperimental/foreach/foreach_div_scalar/op_kernel/foreach_div_scalar.cpp+2-2
@@ -14,8 +14,8 @@
14 */14 */
15 15 
16/**16/**
17- * \file foreach_div_scalar_arch32.cpp17+ * \file foreach_div_scalar_arch22.cpp
18- * \brief ForeachDivScalar kernel entry (arch32 / Ascend910B)18+ * \brief ForeachDivScalar kernel entry (arch22 / Ascend910B)
19 *19 *
20 * Template parameter:20 * Template parameter:
21 * - D_T_X: Data type, mapped from ASCENDC_TPL_DATATYPE_DECL21 * - D_T_X: Data type, mapped from ASCENDC_TPL_DATATYPE_DECL
Mexperimental/foreach/foreach_div_scalar/op_kernel/foreach_div_scalar.h+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file foreach_div_scalar.h17 * \file foreach_div_scalar.h
18- * \brief ForeachDivScalar kernel class definition (arch32 / Ascend910B)18+ * \brief ForeachDivScalar kernel class definition (arch22 / Ascend910B)
19 *19 *
20 * Processes a TensorList by dividing each tensor's elements by a scalar value.20 * Processes a TensorList by dividing each tensor's elements by a scalar value.
21 * Uses Muls(x, 1/scalar) for performance.21 * Uses Muls(x, 1/scalar) for performance.
Mexperimental/loss/soft_margin_loss/op_kernel/soft_margin_loss.cpp+2-2
@@ -14,8 +14,8 @@
14 */14 */
15 15 
16/**16/**
17- * \file soft_margin_loss_arch32.cpp17+ * \file soft_margin_loss_arch22.cpp
18- * \brief SoftMarginLoss kernel entry point (arch32 architecture - Ascend910B)18+ * \brief SoftMarginLoss kernel entry point (arch22 architecture - Ascend910B)
19 *19 *
20 * Dispatches to template instantiations:20 * Dispatches to template instantiations:
21 * SoftMarginLossNone<float/half> - elementwise output21 * SoftMarginLossNone<float/half> - elementwise output
Mexperimental/loss/soft_margin_loss/op_kernel/soft_margin_loss.h+1-1
@@ -15,7 +15,7 @@
15 15 
16/**16/**
17 * \file soft_margin_loss.h17 * \file soft_margin_loss.h
18- * \brief SoftMarginLoss kernel class definition (arch32 - Ascend910B)18+ * \brief SoftMarginLoss kernel class definition (arch22 - Ascend910B)
19 *19 *
20 * Computes SoftMarginLoss: L[i] = max(0, -t*x) + log(1 + exp(-|t*x|))20 * Computes SoftMarginLoss: L[i] = max(0, -t*x) + log(1 + exp(-|t*x|))
21 * where x = self[i], t = target[i]21 * where x = self[i], t = target[i]
Mexperimental/norm/bn_infer_grad/CMakeLists.txt+3-3
@@ -2,7 +2,7 @@
2 # Copyright (c) 2026 Huawei Technologies Co., Ltd.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 of3 # 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").4 # CANN Open Software License Agreement Version 2.0 (the "License").
5- 5+ 
6 # Please refer to the License for details. You may not use this file except in compliance with 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,7 # THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8 # INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.8 # INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
@@ -14,5 +14,5 @@
14# 设置算子定义时支持的芯片类型14# 设置算子定义时支持的芯片类型
15set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b")15set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b")
16# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译16# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
17-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32")17+set(SUPPORT_TILING_DIR "arch22" "arch22" "arch22")
18-add_modules_sources(DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE bn_infer_grad ACLNNTYPE aclnn COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)18+add_modules_sources(DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE bn_infer_grad ACLNNTYPE aclnn COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Rexperimental/norm/bn_infer_grad/op_host/arch32/bn_infer_grad_tiling.cpp→experimental/norm/bn_infer_grad/op_host/arch22/bn_infer_grad_tiling.cpp+2-2
@@ -15,10 +15,10 @@
15 15 
16/*!16/*!
17 * \file bn_infer_grad_tiling.cpp17 * \file bn_infer_grad_tiling.cpp
18- * \brief BnInferGrad Tiling 实现（通用，支持 arch32/arch35）18+ * \brief BnInferGrad Tiling 实现（通用，支持 arch22/arch35）
19 *19 *
20 * Tiling 通过 platform API 动态获取平台参数，20 * Tiling 通过 platform API 动态获取平台参数，
21- * 因此同一份代码可适配 arch32 (Ascend910B) 和 arch35 (Ascend950) 等架构。21+ * 因此同一份代码可适配 arch22 (Ascend910B) 和 arch35 (Ascend950) 等架构。
22 *22 *
23 * 迭代三：支持 CONTIGUOUS(TilingKey=0) + NC1HWC0(TilingKey=1)，23 * 迭代三：支持 CONTIGUOUS(TilingKey=0) + NC1HWC0(TilingKey=1)，
24 * 多核切分，NCHW/NHWC/NC1HWC0 三种格式，24 * 多核切分，NCHW/NHWC/NC1HWC0 三种格式，
Mexperimental/norm/bn_infer_grad/op_kernel/bn_infer_grad.cpp+2-2
@@ -14,8 +14,8 @@
14 */14 */
15 15 
16/*!16/*!
17- * \file bn_infer_grad_arch32.cpp17+ * \file bn_infer_grad_arch22.cpp
18- * \brief BnInferGrad Kernel 入口（arch32 架构）18+ * \brief BnInferGrad Kernel 入口（arch22 架构）
19 *19 *
20 * 模板参数说明（与 bn_infer_grad_tiling_key.h 中 ASCENDC_TPL_ARGS_DECL 定义对应）：20 * 模板参数说明（与 bn_infer_grad_tiling_key.h 中 ASCENDC_TPL_ARGS_DECL 定义对应）：
21 * - D_T_X: 数据类型，由 ASCENDC_TPL_DATATYPE_DECL 定义21 * - D_T_X: 数据类型，由 ASCENDC_TPL_DATATYPE_DECL 定义
Mexperimental/pooling/max_pool_with_argmax_v3/op_host/max_pool_with_argmax_v3_def.cpp+1-1
@@ -50,7 +50,7 @@ public:
50 this->Attr("dilations").AttrType(REQUIRED).ListInt();50 this->Attr("dilations").AttrType(REQUIRED).ListInt();
51 this->Attr("ceil_mode").AttrType(OPTIONAL).Bool(false);51 this->Attr("ceil_mode").AttrType(OPTIONAL).Bool(false);
52 52 
53- // Ascend910B 配置 (arch32)53+ // Ascend910B 配置 (arch22)
54 OpAICoreConfig aicoreConfig910B;54 OpAICoreConfig aicoreConfig910B;
55 aicoreConfig910B.DynamicCompileStaticFlag(true)55 aicoreConfig910B.DynamicCompileStaticFlag(true)
56 .DynamicFormatFlag(false)56 .DynamicFormatFlag(false)
Mexperimental/pooling/max_pool_with_argmax_v3/op_host/max_pool_with_argmax_v3_tiling.cpp+2-2
@@ -163,7 +163,7 @@ static ge::graphStatus MaxPoolWithArgmaxV3TilingFunc(gert::TilingContext* contex
163 } else if (dataType == ge::DT_BF16) {163 } else if (dataType == ge::DT_BF16) {
164 context->SetTilingKey(GET_TPL_TILING_KEY(MAXPOOL_TPL_SCH_MODE_BF16));164 context->SetTilingKey(GET_TPL_TILING_KEY(MAXPOOL_TPL_SCH_MODE_BF16));
165 } else {165 } else {
166- OP_LOGE(context, "unsupported dtype on arch32");166+ OP_LOGE(context, "unsupported dtype on arch22");
167 return ge::GRAPH_FAILED;167 return ge::GRAPH_FAILED;
168 }168 }
169 return ge::GRAPH_SUCCESS;169 return ge::GRAPH_SUCCESS;
@@ -251,7 +251,7 @@ static ge::graphStatus MaxPoolWithArgmaxV3TilingFunc(gert::TilingContext* contex
251 } else if (dataType == ge::DT_BF16) {251 } else if (dataType == ge::DT_BF16) {
252 context->SetTilingKey(GET_TPL_TILING_KEY(MAXPOOL_TPL_SCH_MODE_BF16));252 context->SetTilingKey(GET_TPL_TILING_KEY(MAXPOOL_TPL_SCH_MODE_BF16));
253 } else {253 } else {
254- OP_LOGE(context, "unsupported dtype on arch32");254+ OP_LOGE(context, "unsupported dtype on arch22");
255 return ge::GRAPH_FAILED;255 return ge::GRAPH_FAILED;
256 }256 }
257 257 
Mexperimental/pooling/max_pooling_grad/op_host/max_pooling_grad_def.cpp+1-1
@@ -50,7 +50,7 @@ public:
50 .Format({ge::FORMAT_ND, ge::FORMAT_ND})50 .Format({ge::FORMAT_ND, ge::FORMAT_ND})
51 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});51 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
52 52 
53- // Ascend910B (arch32) 配置53+ // Ascend910B (arch22) 配置
54 OpAICoreConfig aicoreConfig910B;54 OpAICoreConfig aicoreConfig910B;
55 aicoreConfig910B.DynamicCompileStaticFlag(true)55 aicoreConfig910B.DynamicCompileStaticFlag(true)
56 .DynamicFormatFlag(false)56 .DynamicFormatFlag(false)