Star1341
1454
代码
代码
Issues120
Pull Requests260
Actions
项目讨论
Wiki
分析
Star1341
1454
已合并
【社区任务】02-03 Tile算子开发任务贡献 #2417
刘十一创建于 4月24日
【社区任务】02-03 Tile算子开发任务贡献 #2417
已合并
刘十一创建于 4月24日
从ElevenLiu/ops-math0203合入到cann/ops-mathmaster
共 25 个文件变更+5360-0
Aexperimental/math/tile/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+ 
9+add_all_modules_sources(OPTYPE tile ACLNNTYPE aclnn_exclude DEPENDENCIES broadcast_to unsqueeze)
Aexperimental/math/tile/README.md+81-0
@@ -0,0 +1,81 @@
1+# Tile
2+ 
3+## 贡献说明
4+ 
5+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
6+|--------|--------|----------|----------|----------|
7+| ElevenLiu | 智子芯元(深圳)科技有限责任公司 | Tile | 2026/04 | 新增Tile算子适配开源仓 |
8+ 
9+## 支持的产品型号
10+ 
11+- Atlas A2系列产品
12+ 
13+## 算子描述
14+ 
15+- 功能描述
16+ 
17+ `Tile`算子将输入张量按指定倍数在各维度重复拼接，等价于 `numpy.tile(input, multiples)` 或 PyTorch `tensor.repeat(multiples)`。
18+ 
19+- 计算公式
20+ 
21+ $output[i_0, i_1, ..., i_N] = input[i_0 \bmod s_0, i_1 \bmod s_1, ..., i_N \bmod s_N]$
22+ 
23+ 其中 $s_0, s_1, ..., s_N$ 为输入各维度大小。
24+ 
25+- 原型信息
26+ 
27+ <table>
28+ <tr><th align="center">算子类型(OpType)</th><th colspan="4" align="center">Tile</th></tr>
29+ <tr><td align="center"> </td><td align="center">name</td><td align="center">Type</td><td align="center">data type</td><td align="center">format</td></tr>
30+ <tr><td rowspan="2" align="center">算子输入</td>
31+ <td align="center">x</td><td align="center">tensor</td><td align="center">float32, float16, bfloat16, int32, int16, int8, uint8, uint16, uint32, uint64, bool, complex64</td><td align="center">ND</td></tr>
32+ <tr><td align="center">multiples</td><td align="center">tensor</td><td align="center">int32, int64</td><td align="center">ND</td></tr>
33+ <tr><td rowspan="1" align="center">算子输出</td>
34+ <td align="center">y</td><td align="center">tensor</td><td align="center">与x一致</td><td align="center">ND</td></tr>
35+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">tile</td></tr>
36+ </table>
37+ 
38+## 约束与限制
39+ 
40+- 输入张量维度范围 1-8，仅支持 ND 格式
41+- int64、double 数据类型暂不支持
R
Rrxtfeng4月27日

这里写的int64、double数据类型不支持，和aclnn_repeat.cpp文件中DTYPE_SUPPORT_内容不一致

likedislike
刘十一
刘十一
4月27日 评论：
42+- multiples 各维度值须为正整数 (≥1)
43+- 输入/输出总元素数须在 int32 范围内
44+ 
45+## 算子使用
46+ 
47+使用该算子前，请参考[社区版CANN开发套件包安装文档](../../../docs/zh/invocation/quick_op_invocation.md)完成开发运行环境的部署。
48+ 
49+### 编译部署
50+ 
51+ - 进入到仓库目录
52+ 
53+ ```bash
54+ cd ${git_clone_path}/ops-math
55+ ```
56+ 
57+ - 执行编译
58+ 
59+ ```bash
60+ bash build.sh --pkg --experimental --soc=ascend910b --ops=tile
61+ ```
62+ 
63+ - 部署算子包
64+ 
65+ ```bash
66+ ./build_out/cann-ops-<vendor_name>-linux.<arch>.run
67+ ```
68+ 
69+## 调用说明
70+ 
71+Tile 算子对应 aclnn 层接口为 `aclnnRepeat`（已内置于 CANN 包中），通过 `tensor.repeat(multiples)` 方式调用。
72+ 
73+| 调用方式 | 调用样例 | 说明 |
74+|----------|----------|------|
75+| aclnn调用 | [test_aclnn_tile](./examples/test_aclnn_tile.cpp) | 通过[aclnnRepeat](./docs/aclnnRepeat.md)接口方式调用Tile算子 |
76+ 
77+### 执行调用
78+ 
79+```bash
80+bash build.sh --experimental --run_example tile eager cust --vendor_name=custom
81+```
Aexperimental/math/tile/docs/aclnnRepeat.md+87-0
@@ -0,0 +1,87 @@
1+# aclnnRepeat
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :--- | :---: |
7+| Atlas A2系列产品 | √ |
8+ 
9+## 功能描述
10+ 
11+- 算子功能：将输入张量按指定倍数在各维度重复拼接，等价于 `numpy.tile(input, multiples)` 或 PyTorch `tensor.repeat(multiples)`。
12+- 计算公式：
13+ 
14+ $$out[i_0, i_1, ..., i_N] = self[i_0 \bmod s_0, i_1 \bmod s_1, ..., i_N \bmod s_N]$$
15+ 
16+ 其中 $s_0, s_1, ..., s_N$ 为输入各维度大小。
17+ 
18+## 实现原理
19+ 
20+调用 Ascend C 的 DataCopy 接口实现纯数据搬运，通过 Host 侧维度合并优化和多核并行来提升性能。
21+ 
22+## 函数原型
23+ 
24+Tile 算子对应内置 aclnnRepeat 接口（[两段式接口](../../../../docs/zh/context/两段式接口.md)）：
25+ 
26+```cpp
27+aclnnStatus aclnnRepeatGetWorkspaceSize(
28+ const aclTensor *self,
29+ const aclIntArray *repeats,
30+ aclTensor *out,
31+ uint64_t *workspaceSize,
32+ aclOpExecutor **executor)
33+```
34+ 
35+```cpp
36+aclnnStatus aclnnRepeat(
37+ void *workspace,
38+ uint64_t workspaceSize,
39+ aclOpExecutor *executor,
40+ const aclrtStream stream)
41+```
42+ 
43+### aclnnRepeatGetWorkspaceSize
44+ 
45+| 参数名 | 输入/输出 | 描述 | 数据类型 | 是否必选 |
46+|--------|----------|------|----------|----------|
47+| self | 输入 | 输入张量 | aclTensor* | 是 |
48+| repeats | 输入 | 各维度重复倍数 | aclIntArray* | 是 |
49+| out | 输出 | 输出张量 | aclTensor* | 是 |
50+| workspaceSize | 输出 | workspace 大小 | uint64_t* | 是 |
51+| executor | 输出 | 算子执行器 | aclOpExecutor** | 是 |
52+ 
53+### aclnnRepeat
54+ 
55+| 参数名 | 输入/输出 | 描述 | 数据类型 | 是否必选 |
56+|--------|----------|------|----------|----------|
57+| workspace | 输入 | workspace 内存地址 | void* | 是 |
58+| workspaceSize | 输入 | workspace 大小 | uint64_t | 是 |
59+| executor | 输入 | 算子执行器 | aclOpExecutor* | 是 |
60+| stream | 输入 | 计算流 | aclrtStream | 是 |
61+ 
62+## 支持的数据类型
63+ 
64+| self 数据类型 | out 数据类型 |
65+|-------------|------------|
66+| FLOAT | FLOAT |
67+| FLOAT16 | FLOAT16 |
68+| BF16 | BF16 |
69+| INT32 | INT32 |
70+| INT16 | INT16 |
71+| INT8 | INT8 |
72+| UINT8 | UINT8 |
73+| UINT16 | UINT16 |
74+| UINT32 | UINT32 |
75+| UINT64 | UINT64 |
76+| BOOL | BOOL |
77+| COMPLEX64 | COMPLEX64 |
78+ 
79+## 约束说明
80+ 
81+- 输入张量维度范围 1-8，仅支持 ND 格式
82+- repeats 各维度值须为非负整数（≥0），出现 0 时输出对应维度为 0
83+- int64、double 数据类型暂不支持
R
Rrxtfeng4月27日

这里写的int64、double数据类型不支持，和aclnn_repeat.cpp文件中DTYPE_SUPPORT_内容不一致

likedislike
刘十一
刘十一
4月27日 评论：
84+ 
85+## 调用示例
86+ 
87+请参考 [test_aclnn_tile.cpp](../examples/test_aclnn_tile.cpp)。
Aexperimental/math/tile/docs/design_doc.md+128-0
@@ -0,0 +1,128 @@
1+# Tile 算子设计文档
2+ 
3+## 一、需求背景
4+ 
5+### 1.1 需求来源
6+ 
7+参考版本内置算子的 TBE 实现，在昇腾 NPU 上使用 Ascend C 编程语言实现相同功能的 Tile 算子，提交到算子开源仓。
8+ 
9+### 1.2 TBE 算子源码参考
10+ 
11+- kernel 实现：`/usr/local/Ascend/cann-8.5.0/opp/built-in/op_impl/ai_core/tbe/impl/ops_legacy/dynamic/tile.py`
12+- 算子信息库：`/usr/local/Ascend/cann-8.5.0/opp/built-in/op_impl/ai_core/tbe/config/ascend910b/aic-ascend910b-ops-info-legacy.json`（Tile 条目）
13+- aclnn 接口：`/usr/local/Ascend/cann-8.5.0/aarch64-linux/include/aclnnop/level2/aclnn_repeat.h`（aclnnRepeat）
14+ 
15+### 1.3 TBE 实现概述
16+ 
17+TBE Tile 算子核心逻辑：
18+ 
19+1. 输入验证：检查 x 数据类型、multiples 类型和维度
20+2. Shape 适配：维度对齐，为每个维度拆分为 (repeat_count, original_size) 两层
21+3. 计算核心：通过 `tbe.broadcast` 完成广播运算
22+4. int8/uint8 类型先 cast 到 float16 再 broadcast 再 cast 回原类型
23+ 
24+AscendC 实现与 TBE 语义完全一致，但采用直接数据搬运方式，无需类型转换。
25+ 
26+## 二、算子实现设计
27+ 
28+### 2.1 工程结构
29+ 
30+```
31+tile/
32+├── op_host/
33+│ ├── tile_def.cpp # 算子信息库 (OP_ADD 注册)
34+│ ├── tile_infershape.cpp # InferShape 实现 (IMPL_OP_INFERSHAPE)
35+│ └── tile_tiling.cpp # Tiling 实现 (IMPL_OP_OPTILING)
36+├── op_api/
37+│ ├── aclnn_repeat.cpp # aclnnRepeat 接口实现
38+│ ├── aclnn_repeat.h # aclnnRepeat 接口声明
39+│ ├── tile.cpp # l0op::Tile 实现
40+│ └── tile.h # l0op::Tile 声明
41+├── op_kernel/
42+│ ├── tile.cpp # Kernel 入口（模板函数 + REGISTER_TILING_DEFAULT）
43+│ ├── tile.h # Kernel 实现（TileOpImpl 模板类）
44+│ ├── tile_tiling_data.h # TilingData 结构体
45+│ └── tile_tiling_key.h # TilingKey 模板注册
46+├── examples/
47+│ └── test_aclnn_tile.cpp # aclnnRepeat 调用示例
48+├── tests/
49+│ └── ut/
50+│ ├── op_host/ # InferShape 单元测试
51+│ └── op_kernel/ # Kernel 测试数据生成/比对脚本
52+├── docs/
53+│ ├── design_doc.md # 算子设计文档
54+│ └── aclnnRepeat.md # aclnn 接口文档
55+├── CMakeLists.txt
56+└── README.md
57+```
58+ 
59+### 2.2 Host 侧设计
60+ 
61+#### 2.2.1 算子定义 (tile_def.cpp)
62+ 
63+通过 `OP_ADD(Tile)` 注册算子信息库：
64+- 输入 x：支持 12 种数据类型 (float32/float16/bfloat16/int32/int16/int8/uint8/uint16/uint32/uint64/bool/complex64)
65+- 输入 multiples：int32 类型 1-D tensor
66+- 输出 y：与 x 类型一致
67+ 
68+#### 2.2.2 Tiling 策略 (tile_tiling.cpp)
69+ 
70+通过 `IMPL_OP_OPTILING(Tile)` 注册，核心逻辑：
71+ 
72+1. 获取平台信息（核数、UB 大小）
73+2. 读取输入 shape 和 multiples tensor 数据
74+3. 维度合并优化：合并 mult=1 的连续维度以扩大 innerDim
75+4. 计算 TileTilingData 并通过 `GET_TPL_TILING_KEY` 设置 tilingKey
76+5. 设置 workspace 和 blockDim
77+ 
78+#### 2.2.3 InferShape (tile_infershape.cpp)
79+ 
80+通过 `IMPL_OP_INFERSHAPE(Tile)` 注册，根据输入 shape 和 multiples 计算输出 shape：`output_shape[i] = input_shape[i] * multiples[i]`。当输入维度与 multiples 长度不一致时，自动左侧补 1 对齐到较大维度数。
81+ 
82+### 2.3 Device 侧设计
83+ 
84+#### 2.3.1 Kernel 入口 (tile.cpp)
85+ 
86+使用单一模板函数 `tile<schMode>()`，通过 `DTYPE_X` 推导数据类型，通过 `REGISTER_TILING_DEFAULT` + `GET_TILING_DATA_WITH_STRUCT` 获取 tiling 数据。
87+ 
88+#### 2.3.2 Kernel 实现 (tile.h)
89+ 
90+模板类 `TileOpImpl<T>` 支持所有 12 种数据类型。根据运行时数据特征动态选择 5 条优化路径：
91+ 
92+| 条件 | 路径 | 说明 |
93+|------|------|------|
94+| innerDim 超过 UB 容量 | ProcessLargeInner | 分块搬运大 innerDim |
95+| outerCount < blockDim 且 innerMult > 1 | ProcessSplitMult | 按 mult 维度细粒度分核 |
96+| innerDim 对齐且 innerMult ≥ 4 | ProcessDoubling | UB 内倍增写出 |
97+| outputInnerDim 可放入 UB | ProcessBuild | UB 内构建完整输出行 |
98+| 以上均不满足 | ProcessPerRow | 逐行处理（通用兜底） |
99+ 
100+核心优化技术：
101+- **UB 内倍增**：小块翻倍复制后大块写出，减少 GM 写次数
102+- **多核分 mult**：outerCount 不足时按 innerMult 维度切分工作
103+- **维度合并**：Host 端合并 mult=1 维度扩大 innerDim
104+ 
105+### 2.4 aclnn 接口
106+ 
107+Tile 算子对应已有内置 aclnn 接口 `aclnnRepeat`，CMakeLists 配置 `ACLNNTYPE aclnn_exclude` 表示不生成新的 aclnn 接口，复用已有接口。
108+ 
109+## 三、精度与性能
110+ 
111+### 3.1 精度
112+ 
113+Tile 为纯数据搬运算子，精度误差恒为 0　12 种数据类型 × 13 种 shape = 156 组，全部通过。
114+ 
115+### 3.2 性能
116+ 
117+测试环境：Atlas A2 系列, blockDim=24, warmup=50, repeat=200
118+ 
119+| 数据类型 | (1024²)×(2,2) | (256²×4)×(2³) | (1×1024)×(1024,1) | (128)×(8192) |
120+|----------|-----------|----------|------------|----------|
121+| float32 | 111% | 102% | 241% | 238% |
122+| float16 | 161% | 104% | 241% | 238% |
123+| int32 | 115% | 103% | 242% | 236% |
124+| int16 | 165% | 106% | 235% | 234% |
125+| int8 | 196% | 423% | 233% | 226% |
126+| uint8 | 184% | 412% | 227% | 232% |
127+ 
128+24 组全部 ≥ 95%，最低 102%，最高 423%。
Aexperimental/math/tile/examples/test_aclnn_tile.cpp+139-0
@@ -0,0 +1,139 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <vector>
13+#include "acl/acl.h"
14+#include "aclnnop/level2/aclnn_repeat.h"
15+ 
16+#define CHECK_RET(cond, return_expr) \
17+ do { \
18+ if (!(cond)) { \
19+ return_expr; \
20+ } \
21+ } while (0)
22+ 
23+#define LOG_PRINT(message, ...) \
24+ do { \
25+ printf(message, ##__VA_ARGS__); \
26+ } while (0)
27+ 
28+int64_t GetShapeSize(const std::vector<int64_t>& shape)
29+{
30+ int64_t shapeSize = 1;
31+ for (auto i : shape) {
32+ shapeSize *= i;
33+ }
34+ return shapeSize;
35+}
36+ 
37+int Init(int32_t deviceId, aclrtStream* stream)
38+{
39+ auto ret = aclInit(nullptr);
40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
41+ ret = aclrtSetDevice(deviceId);
42+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
43+ ret = aclrtCreateStream(stream);
44+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
45+ return 0;
46+}
47+ 
48+template <typename T>
49+int CreateAclTensor(
50+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
51+ aclTensor** tensor)
52+{
53+ auto size = GetShapeSize(shape) * sizeof(T);
54+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
55+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
56+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
57+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
58+ 
59+ std::vector<int64_t> strides(shape.size(), 1);
60+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
61+ strides[i] = shape[i + 1] * strides[i + 1];
62+ }
63+ *tensor = aclCreateTensor(
64+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
65+ *deviceAddr);
66+ return 0;
67+}
68+ 
69+int main()
70+{
71+ int32_t deviceId = 0;
72+ aclrtStream stream;
73+ auto ret = Init(deviceId, &stream);
74+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
75+ 
76+ // input: shape [2, 3], repeats [3, 2] -> output shape [6, 6]
77+ std::vector<int64_t> selfShape = {2, 3};
78+ std::vector<int64_t> outShape = {6, 6};
79+ std::vector<float> selfHostData = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f};
80+ std::vector<float> outHostData(GetShapeSize(outShape), 0.0f);
81+ std::vector<int64_t> repeatsData = {3, 2};
82+ 
83+ void* selfDeviceAddr = nullptr;
84+ void* outDeviceAddr = nullptr;
85+ aclTensor* self = nullptr;
86+ aclTensor* out = nullptr;
87+ 
88+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
89+ CHECK_RET(ret == ACL_SUCCESS, return ret);
90+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
91+ CHECK_RET(ret == ACL_SUCCESS, return ret);
92+ 
93+ aclIntArray* repeats = aclCreateIntArray(repeatsData.data(), repeatsData.size());
94+ 
95+ uint64_t workspaceSize = 0;
96+ aclOpExecutor* executor;
97+ ret = aclnnRepeatGetWorkspaceSize(self, repeats, out, &workspaceSize, &executor);
98+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnRepeatGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
99+ 
100+ void* workspaceAddr = nullptr;
101+ if (workspaceSize > 0) {
102+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
103+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
104+ }
105+ 
106+ ret = aclnnRepeat(workspaceAddr, workspaceSize, executor, stream);
107+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnRepeat failed. ERROR: %d\n", ret); return ret);
108+ 
109+ ret = aclrtSynchronizeStream(stream);
110+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
111+ 
112+ auto size = GetShapeSize(outShape);
113+ std::vector<float> resultData(size, 0);
114+ ret = aclrtMemcpy(
115+ resultData.data(), resultData.size() * sizeof(float), outDeviceAddr, size * sizeof(float),
116+ ACL_MEMCPY_DEVICE_TO_HOST);
117+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result failed. ERROR: %d\n", ret); return ret);
118+ 
119+ LOG_PRINT("aclnnRepeat (Tile) result:\n");
120+ for (int64_t i = 0; i < outShape[0]; i++) {
121+ for (int64_t j = 0; j < outShape[1]; j++) {
122+ LOG_PRINT("%6.1f ", resultData[i * outShape[1] + j]);
123+ }
124+ LOG_PRINT("\n");
125+ }
126+ 
127+ aclDestroyTensor(self);
128+ aclDestroyTensor(out);
129+ aclDestroyIntArray(repeats);
130+ aclrtFree(selfDeviceAddr);
131+ aclrtFree(outDeviceAddr);
132+ if (workspaceSize > 0) {
133+ aclrtFree(workspaceAddr);
134+ }
135+ aclrtDestroyStream(stream);
136+ aclrtResetDevice(deviceId);
137+ aclFinalize();
138+ return 0;
139+}
Aexperimental/math/tile/op_api/aclnn_repeat.cpp+249-0
@@ -0,0 +1,249 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "aclnn_repeat.h"
12+ 
13+#include "op_api/op_api_def.h"
14+#include "op_api/aclnn_check.h"
15+#include "aclnn_kernels/common/op_error_check.h"
16+#include "opdev/common_types.h"
17+#include "opdev/data_type_utils.h"
18+#include "opdev/format_utils.h"
19+#include "opdev/make_op_executor.h"
20+#include "opdev/op_dfx.h"
21+#include "opdev/op_executor.h"
22+#include "opdev/op_log.h"
23+#include "opdev/platform.h"
24+#include "opdev/shape_utils.h"
25+#include "opdev/tensor_view_utils.h"
26+ 
27+#include "conversion/broadcast_to/op_api/broadcast_to.h"
28+#include "aclnn_kernels/cast.h"
29+#include "aclnn_kernels/contiguous.h"
30+#include "aclnn_kernels/reshape.h"
31+#include "aclnn_kernels/transpose.h"
32+#include "tile.h"
33+#include "conversion/unsqueeze/op_host/op_api/unsqueeze.h"
34+ 
35+using namespace op;
36+ 
37+#ifdef __cplusplus
38+extern "C" {
39+#endif
40+ 
41+static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {
42+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_COMPLEX64, op::DataType::DT_UINT8,
43+ op::DataType::DT_INT8, op::DataType::DT_INT16, op::DataType::DT_INT32, op::DataType::DT_BOOL,
44+ op::DataType::DT_BF16, op::DataType::DT_UINT32, op::DataType::DT_UINT64};
45+ 
46+static const std::initializer_list<op::DataType> DTYPE_SUPPORT_910B310P_LIST = {
47+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_COMPLEX64, op::DataType::DT_UINT8,
48+ op::DataType::DT_INT8, op::DataType::DT_INT16, op::DataType::DT_INT32, op::DataType::DT_BOOL,
49+ op::DataType::DT_BF16, op::DataType::DT_UINT32, op::DataType::DT_UINT64, op::DataType::DT_UINT16};
50+ 
51+static const std::initializer_list<op::DataType> DTYPE_SUPPORT_950_LIST = {
52+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_COMPLEX64, op::DataType::DT_UINT8,
53+ op::DataType::DT_INT8, op::DataType::DT_INT16, op::DataType::DT_INT32, op::DataType::DT_BOOL,
54+ op::DataType::DT_BF16, op::DataType::DT_UINT32, op::DataType::DT_UINT64, op::DataType::DT_UINT16,
55+ op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2, op::DataType::DT_FLOAT8_E4M3FN};
56+ 
57+// 检查tensor是否为nullptr
58+static inline bool CheckNotNull(const aclTensor* self, const aclIntArray* repeats, const aclTensor* out)
59+{
60+ OP_CHECK_NULL(self, return false);
61+ OP_CHECK_NULL(repeats, return false);
62+ OP_CHECK_NULL(out, return false);
63+ return true;
64+}
65+ 
66+static bool CheckDtypeValid(const aclTensor* self, const aclTensor* out)
67+{
68+ if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910B ||
69+ GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_93 ||
70+ GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND310P) {
71+ OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_910B310P_LIST, return false);
72+ } else if (IsRegBase()) {
73+ OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_950_LIST, return false);
74+ } else {
75+ OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST, return false);
76+ }
77+ 
78+ OP_CHECK_DTYPE_NOT_SAME(self, out, return false);
79+ return true;
80+}
81+ 
82+// tensor维度数不能超过8维
83+static inline bool CheckTensorDimSize(const aclTensor* self, const aclIntArray* repeats)
84+{
85+ int64_t tensorDimSize = self->GetViewShape().GetDimNum();
86+ int64_t repeatsSize = repeats->Size();
87+ OP_CHECK_MAX_DIM(self, MAX_SUPPORT_DIMS_NUMS, return false);
88+ if (repeatsSize > static_cast<int64_t>(MAX_SUPPORT_DIMS_NUMS)) {
89+ OP_LOGE(
90+ ACLNN_ERR_PARAM_INVALID, "repeats size should not be larger than 8. self %ld, repeats %ld", tensorDimSize,
91+ repeatsSize);
92+ return false;
93+ }
94+ return true;
95+}
96+ 
97+static inline bool CheckRepeatsSize(const aclTensor* self, const aclIntArray* repeats)
98+{
99+ if (self->GetViewShape().GetDimNum() > repeats->Size()) {
100+ OP_LOGE(
101+ ACLNN_ERR_PARAM_INVALID,
102+ "Number of dimensions of repeat dims can not be smaller than number of dimensions of input tensor, self "
103+ "%lu, repeats %lu.",
104+ self->GetViewShape().GetDimNum(), repeats->Size());
105+ return false;
106+ }
107+ return true;
108+}
109+ 
110+// 校验repeat中的值均大于等于0
111+static inline aclnnStatus CheckRepeatsValue(const aclIntArray* repeats)
112+{
113+ for (size_t i = 0; i < repeats->Size(); ++i) {
114+ if ((*repeats)[i] < 0) {
R
Rrxtfeng4月27日

仅判断了(*repeats)[i] < 0，但是错误日志repeats expected %luth value > 0，两者不一致

likedislike
刘十一
刘十一
4月27日 评论：
115+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "repeats expected %luth value >= 0, but get %ld", i + 1, (*repeats)[i]);
116+ return false;
117+ }
118+ }
119+ return true;
120+}
121+ 
122+// 校验repeat中的值存在0
123+static inline aclnnStatus CheckRepeatsZero(const aclIntArray* repeats)
124+{
125+ for (size_t i = 0; i < repeats->Size(); ++i) {
126+ if ((*repeats)[i] == 0) {
127+ return true;
128+ }
129+ }
130+ return false;
131+}
132+ 
133+// 校验repeat算子计算后生成的tensor与out的shape一致
134+static inline bool CheckRepeatOutShape(const aclTensor* repeatRes, const aclTensor* out)
135+{
136+ OP_CHECK_SHAPE_NOT_EQUAL(repeatRes, out, return false);
137+ 
138+ return true;
139+}
140+ 
141+// 返回为[1]的intarray shape
142+static inline aclIntArray* GetBaseShape(aclOpExecutor* executor)
143+{
144+ int64_t tensorShape[1] = {};
145+ tensorShape[0] = 1;
146+ auto res = executor->AllocIntArray(tensorShape, 1);
147+ return res;
148+}
149+ 
150+// 如果self为0维tensor，那么转换为1维tensor。其余情况转成连续tensor
151+static inline const aclTensor* InitializeTensor(const aclTensor* x, aclOpExecutor* executor)
152+{
153+ auto xContiguous = l0op::Contiguous(x, executor);
154+ // 如果tensor为0维，则转换为1维tensor
155+ if (xContiguous->GetViewShape().GetDimNum() == 0) {
156+ auto baseShape = GetBaseShape(executor);
157+ xContiguous = l0op::BroadcastTo(xContiguous, baseShape, executor);
158+ }
159+ return xContiguous;
160+}
161+ 
162+static const aclTensor* ViewToRepeatsSize(const aclTensor* self, const aclIntArray* repeats, aclOpExecutor* executor)
163+{
164+ auto Dims = (int64_t)repeats->Size() - (int64_t)self->GetViewShape().GetDimNum();
165+ std::vector<int64_t> unsqueezeDim(Dims);
166+ for (int64_t idx = 0; idx < Dims; idx++) {
167+ unsqueezeDim[idx] = idx;
168+ }
169+ aclIntArray* dim = executor->AllocIntArray(unsqueezeDim.data(), Dims);
170+ auto unsqeezeSelf = l0op::UnsqueezeNd(self, dim, executor);
171+ return unsqeezeSelf;
172+}
173+ 
174+// 校验开始
175+static aclnnStatus CheckParams(const aclTensor* self, const aclIntArray* repeats, const aclTensor* out)
176+{
177+ CHECK_RET(CheckNotNull(self, repeats, out), ACLNN_ERR_PARAM_NULLPTR);
178+ 
179+ CHECK_RET(CheckDtypeValid(self, out), ACLNN_ERR_PARAM_INVALID);
180+ 
181+ CHECK_RET(CheckRepeatsSize(self, repeats), ACLNN_ERR_PARAM_INVALID);
182+ 
183+ CHECK_RET(CheckRepeatsValue(repeats), ACLNN_ERR_PARAM_INVALID);
184+ 
185+ CHECK_RET(CheckTensorDimSize(self, repeats), ACLNN_ERR_PARAM_INVALID);
186+ 
187+ return ACLNN_SUCCESS;
188+}
189+ 
190+aclnnStatus aclnnRepeatGetWorkspaceSize(
191+ const aclTensor* self, const aclIntArray* repeats, aclTensor* out, uint64_t* workspaceSize,
192+ aclOpExecutor** executor)
193+{
194+ L2_DFX_PHASE_1(aclnnRepeat, DFX_IN(self, repeats), DFX_OUT(out));
195+ 
196+ auto uniqueExecutor = CREATE_EXECUTOR();
197+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
198+ 
199+ auto ret = CheckParams(self, repeats, out);
200+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
201+ 
202+ if (self->IsEmpty() || out->IsEmpty() || CheckRepeatsZero(repeats)) {
203+ *workspaceSize = 0;
204+ uniqueExecutor.ReleaseTo(executor);
205+ return ACLNN_SUCCESS;
206+ }
207+ if (self->GetStorageFormat() != Format::FORMAT_ND) {
208+ OP_LOGW("Format only support ND");
209+ }
210+ if (repeats->Size() == 0) {
211+ auto viewCopyOut = l0op::ViewCopy(self, out, uniqueExecutor.get());
212+ CHECK_RET(viewCopyOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
213+ 
214+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
215+ uniqueExecutor.ReleaseTo(executor);
216+ return ACLNN_SUCCESS;
217+ }
218+ 
219+ auto selfContiguous = InitializeTensor(self, uniqueExecutor.get());
220+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
221+ auto selfUnsqueeze = selfContiguous;
222+ 
223+ if (selfContiguous->GetViewShape().GetDimNum() < repeats->Size()) {
224+ selfUnsqueeze = ViewToRepeatsSize(selfContiguous, repeats, uniqueExecutor.get());
225+ CHECK_RET(selfUnsqueeze != nullptr, ACLNN_ERR_PARAM_INVALID);
226+ }
227+ 
228+ auto repeatOut = l0op::Tile(selfUnsqueeze, repeats, uniqueExecutor.get());
229+ CHECK_RET(repeatOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
230+ CHECK_RET(CheckRepeatOutShape(repeatOut, out), ACLNN_ERR_PARAM_INVALID);
231+ 
232+ auto viewCopyOut = l0op::ViewCopy(repeatOut, out, uniqueExecutor.get());
233+ CHECK_RET(viewCopyOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
234+ 
235+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
236+ uniqueExecutor.ReleaseTo(executor);
237+ return ACLNN_SUCCESS;
238+}
239+ 
240+aclnnStatus aclnnRepeat(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
241+{
242+ L2_DFX_PHASE_2(aclnnRepeat);
243+ 
244+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
245+}
246+ 
247+#ifdef __cplusplus
248+}
249+#endif
Aexperimental/math/tile/op_api/aclnn_repeat.h+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+#ifndef OP_API_INC_LEVEL2_ACLNN_REPEAT_H_
12+#define OP_API_INC_LEVEL2_ACLNN_REPEAT_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+#ifdef __cplusplus
17+extern "C" {
18+#endif
19+ 
20+/**
21+ * @brief aclnnRepeat的第一段接口，根据具体的计算流程，计算workspace大小。
22+ * @domain aclnn_ops_infer
23+ */
24+ACLNN_API aclnnStatus aclnnRepeatGetWorkspaceSize(
25+ const aclTensor* self, const aclIntArray* repeats, aclTensor* out, uint64_t* workspaceSize,
26+ aclOpExecutor** executor);
27+ 
28+/**
29+ * @brief aclnnRepeat的第二段接口，用于执行计算。
30+ */
31+ACLNN_API aclnnStatus aclnnRepeat(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
32+ 
33+#ifdef __cplusplus
34+}
35+#endif
36+ 
37+#endif
Aexperimental/math/tile/op_api/tile.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+ * \file tile.cpp
13+ * \brief
14+ */
15+ 
16+#include "tile.h"
17+#include "op_api/aclnn_check.h"
18+#include "opdev/aicpu/aicpu_task.h"
19+#include "opdev/make_op_executor.h"
20+#include "opdev/op_dfx.h"
21+#include "opdev/op_def.h"
22+#include "opdev/op_executor.h"
23+#include "opdev/op_log.h"
24+#include "opdev/shape_utils.h"
25+#include "opdev/platform.h"
26+ 
27+using namespace op;
28+namespace l0op {
29+OP_TYPE_REGISTER(Tile);
30+ 
31+static inline const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
32+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32,
33+ op::DataType::DT_INT64, op::DataType::DT_BOOL, op::DataType::DT_BF16};
34+ 
35+static inline const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_310P_LIST = {
36+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32, op::DataType::DT_INT64,
37+ op::DataType::DT_BOOL, op::DataType::DT_BF16, op::DataType::DT_INT8, op::DataType::DT_UINT8,
38+ op::DataType::DT_INT16, op::DataType::DT_UINT16, op::DataType::DT_UINT32, op::DataType::DT_UINT64};
39+ 
40+static inline const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_910B_LIST = {
41+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32, op::DataType::DT_INT64,
42+ op::DataType::DT_BOOL, op::DataType::DT_BF16, op::DataType::DT_INT8, op::DataType::DT_UINT8,
43+ op::DataType::DT_INT16, op::DataType::DT_UINT16, op::DataType::DT_UINT32, op::DataType::DT_UINT64,
44+ op::DataType::DT_COMPLEX64};
45+ 
46+static inline const std::initializer_list<op::DataType> ASCEND610LITE_DTYPE_SUPPORT_LIST = {
47+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32, op::DataType::DT_BOOL};
48+ 
49+static const std::initializer_list<DataType> ASCEND950_AICORE_DTYPE_SUPPORT_LIST = {
50+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32,
51+ op::DataType::DT_INT64, op::DataType::DT_BOOL, op::DataType::DT_BF16,
52+ op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_INT16,
53+ op::DataType::DT_UINT16, op::DataType::DT_UINT32, op::DataType::DT_UINT64,
54+ op::DataType::DT_COMPLEX64, op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2,
55+ op::DataType::DT_FLOAT8_E4M3FN};
56+ 
57+static inline bool IsAiCoreSupport(const aclTensor* self)
58+{
59+ if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910B) {
60+ return CheckType(self->GetDataType(), AICORE_DTYPE_SUPPORT_910B_LIST);
61+ }
62+ return CheckType(self->GetDataType(), AICORE_DTYPE_SUPPORT_LIST);
63+}
64+ 
65+static inline const aclTensor* TileAiCore(
66+ const aclTensor* self, const aclTensor* repeats, aclTensor* out, aclOpExecutor* executor)
67+{
68+ L0_DFX(TileAiCore, self, repeats, out);
69+ ADD_TO_LAUNCHER_LIST_AICORE(Tile, OP_INPUT(self, repeats), OP_OUTPUT(out));
70+ return out;
71+}
72+ 
73+static inline const aclTensor* TileAiCpu(
74+ const aclTensor* self, const aclTensor* repeats, aclTensor* out, aclOpExecutor* executor)
75+{
76+ L0_DFX(TileAiCpu, self, repeats, out);
77+ static internal::AicpuTaskSpace space("Tile", ge::DEPEND_IN_SHAPE, true);
78+ op::DataType dtype = repeats->GetDataType();
79+ auto ret = ADD_TO_LAUNCHER_LIST_AICPU(
80+ Tile, OP_ATTR_NAMES({"Tmultiples"}), OP_INPUT(self, repeats), OP_OUTPUT(out), OP_ATTR(dtype));
81+ CHECK_RET(ret == ACLNN_SUCCESS, nullptr);
82+ return out;
83+}
84+ 
85+const aclTensor* Tile(const aclTensor* self, const aclIntArray* repeats, aclOpExecutor* executor)
86+{
87+ int64_t expandDim = static_cast<int64_t>(repeats->Size() - self->GetViewShape().GetDimNum());
88+ op::Shape broadcastShape;
89+ if (expandDim < 0) {
90+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "repeats size should not be smaller than input tensor dim");
91+ return nullptr;
92+ } else {
93+ for (int64_t i = 0; i < expandDim; i++) {
94+ broadcastShape.AppendDim((*repeats)[i]);
95+ }
96+ for (size_t i = 0; i < self->GetViewShape().GetDimNum(); i++) {
97+ broadcastShape.AppendDim(self->GetViewShape()[i] * (*repeats)[i + expandDim]);
98+ }
99+ }
100+ auto out = executor->AllocTensor(broadcastShape, self->GetDataType());
101+ auto repeatTensor = executor->ConvertToTensor(repeats, op::DataType::DT_INT64);
102+ 
103+ if (IsAiCoreSupport(self)) {
104+ return TileAiCore(self, repeatTensor, out, executor);
105+ } else {
106+ return TileAiCpu(self, repeatTensor, out, executor);
107+ }
108+}
109+} // namespace l0op
Aexperimental/math/tile/op_api/tile.h+19-0
@@ -0,0 +1,19 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+#ifndef OP_API_INC_LEVEL0_TILE_H_
11+#define OP_API_INC_LEVEL0_TILE_H_
12+ 
13+#include "opdev/op_executor.h"
14+ 
15+namespace l0op {
16+const aclTensor* Tile(const aclTensor* self, const aclIntArray* repeats, aclOpExecutor* executor);
17+} // namespace l0op
18+ 
19+#endif
Aexperimental/math/tile/op_host/tile_def.cpp+64-0
@@ -0,0 +1,64 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_def.cpp
13+ * \brief
14+ */
15+ 
16+#include <vector>
17+ 
18+#include "register/op_def_registry.h"
19+ 
20+namespace ops {
21+ 
22+static const std::vector<ge::DataType> kXDataTypes = {
23+ ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16,
24+ 
25+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT16, ge::DT_INT16, ge::DT_INT8, ge::DT_INT8,
26+ 
27+ ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT16, ge::DT_UINT16, ge::DT_UINT32, ge::DT_UINT32,
28+ ge::DT_UINT64, ge::DT_UINT64,
29+ 
30+ ge::DT_BOOL, ge::DT_BOOL, ge::DT_COMPLEX64, ge::DT_COMPLEX64};
31+ 
32+static const std::vector<ge::DataType> kMultiplesDataTypes = {
33+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
34+ 
35+ ge::DT_INT32, ge::DT_INT64, ge::DT_INT32, ge::DT_INT64, ge::DT_INT32, ge::DT_INT64,
36+ 
37+ ge::DT_INT32, ge::DT_INT64, ge::DT_INT32, ge::DT_INT64, ge::DT_INT32, ge::DT_INT64, ge::DT_INT32, ge::DT_INT64,
38+ 
39+ ge::DT_INT32, ge::DT_INT64, ge::DT_INT32, ge::DT_INT64};
40+ 
41+static const std::vector<ge::Format> kNdFormats(24, ge::FORMAT_ND);
42+ 
43+class Tile : public OpDef {
44+public:
45+ explicit Tile(const char* name) : OpDef(name)
46+ {
47+ this->Input("x").ParamType(REQUIRED).DataType(kXDataTypes).Format(kNdFormats).UnknownShapeFormat(kNdFormats);
48+ 
49+ this->Input("multiples")
50+ .ParamType(REQUIRED)
51+ .DataType(kMultiplesDataTypes)
52+ .Format(kNdFormats)
53+ .UnknownShapeFormat(kNdFormats)
54+ .ValueDepend(REQUIRED);
55+ 
56+ this->Output("y").ParamType(REQUIRED).DataType(kXDataTypes).Format(kNdFormats).UnknownShapeFormat(kNdFormats);
57+ 
58+ this->AICore().AddConfig("ascend910b");
59+ }
60+};
61+ 
62+OP_ADD(Tile);
63+ 
64+} // namespace ops
Aexperimental/math/tile/op_host/tile_infershape.cpp+119-0
@@ -0,0 +1,119 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_infershape.cpp
13+ * \brief
14+ */
15+#include <vector>
16+ 
17+#include "log/log.h"
18+#include "register/op_impl_registry.h"
19+ 
20+using namespace ge;
21+ 
22+namespace ops {
23+static constexpr size_t IDX_X = 0;
24+static constexpr size_t IDX_MULTIPLES = 1;
25+static constexpr size_t IDX_Y = 0;
26+ 
27+static ge::graphStatus SetUnknownShape(gert::InferShapeContext* context, gert::Shape* yShape, size_t outputDims)
28+{
29+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
30+ yShape->SetDimNum(outputDims);
31+ for (size_t idx = 0; idx < outputDims; ++idx) {
32+ yShape->SetDim(idx, -1);
33+ }
34+ return GRAPH_SUCCESS;
35+}
36+ 
37+template <typename T>
38+static ge::graphStatus GetMultiples(
39+ gert::InferShapeContext* context, const gert::Tensor* multiplesTensor, size_t multiplesSize,
40+ std::vector<int64_t>& multiples)
41+{
42+ const T* multiplesData = multiplesTensor->GetData<T>();
43+ OP_CHECK_NULL_WITH_CONTEXT(context, multiplesData);
44+ multiples.reserve(multiplesSize);
45+ for (size_t idx = 0; idx < multiplesSize; ++idx) {
46+ multiples.push_back(static_cast<int64_t>(multiplesData[idx]));
47+ }
48+ return GRAPH_SUCCESS;
49+}
50+ 
51+static ge::graphStatus InferShapeTile(gert::InferShapeContext* context)
52+{
53+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
54+ 
55+ const gert::Shape* xShape = context->GetInputShape(IDX_X);
56+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
57+ gert::Shape* yShape = context->GetOutputShape(IDX_Y);
58+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
59+ 
60+ const gert::Shape* multiplesShape = context->GetInputShape(IDX_MULTIPLES);
61+ OP_CHECK_NULL_WITH_CONTEXT(context, multiplesShape);
62+ const gert::Tensor* multiplesTensor = context->GetInputTensor(IDX_MULTIPLES);
63+ 
64+ size_t xDims = xShape->GetDimNum();
65+ size_t multiplesSize = 0;
66+ if (multiplesShape->GetDimNum() > 0) {
67+ multiplesSize = static_cast<size_t>(multiplesShape->GetDim(0));
68+ }
69+ if (multiplesSize == 0) {
70+ *yShape = *xShape;
71+ return GRAPH_SUCCESS;
72+ }
73+ size_t outputDims = (xDims > multiplesSize) ? xDims : multiplesSize;
74+ if (multiplesTensor == nullptr) {
75+ return SetUnknownShape(context, yShape, outputDims);
76+ }
77+ 
78+ ge::DataType multiplesDtype = multiplesTensor->GetDataType();
79+ bool hasData = false;
80+ if (multiplesDtype == ge::DT_INT32) {
81+ hasData = (multiplesTensor->GetData<int32_t>() != nullptr);
82+ } else if (multiplesDtype == ge::DT_INT64) {
83+ hasData = (multiplesTensor->GetData<int64_t>() != nullptr);
84+ } else {
85+ OP_LOGE(context, "multiples dtype must be int32 or int64");
86+ return GRAPH_FAILED;
87+ }
88+ if (!hasData) {
89+ return SetUnknownShape(context, yShape, outputDims);
90+ }
91+ 
92+ std::vector<int64_t> multiples;
93+ if (multiplesDtype == ge::DT_INT32) {
94+ OP_CHECK_IF(
95+ GetMultiples<int32_t>(context, multiplesTensor, multiplesSize, multiples) != GRAPH_SUCCESS,
96+ OP_LOGE(context, "get int32 multiples failed"), return GRAPH_FAILED);
97+ } else {
98+ OP_CHECK_IF(
99+ GetMultiples<int64_t>(context, multiplesTensor, multiplesSize, multiples) != GRAPH_SUCCESS,
100+ OP_LOGE(context, "get int64 multiples failed"), return GRAPH_FAILED);
101+ }
102+ 
103+ size_t xOffset = outputDims - xDims;
104+ size_t multiplesOffset = outputDims - multiples.size();
105+ yShape->SetDimNum(outputDims);
106+ for (size_t idx = 0; idx < outputDims; ++idx) {
107+ int64_t xDim = (idx < xOffset) ? 1 : xShape->GetDim(idx - xOffset);
108+ int64_t repeat = (idx < multiplesOffset) ? 1 : multiples[idx - multiplesOffset];
109+ OP_CHECK_IF(
110+ repeat < 0, OP_LOGE(context, "multiples[%zu] must be >= 0, but got %ld", idx - multiplesOffset, repeat),
111+ return GRAPH_FAILED);
112+ yShape->SetDim(idx, xDim * repeat);
113+ }
114+ 
115+ return GRAPH_SUCCESS;
116+}
117+ 
118+IMPL_OP_INFERSHAPE(Tile).InferShape(InferShapeTile).InputsDataDependency({IDX_MULTIPLES});
119+} // namespace ops
Aexperimental/math/tile/op_host/tile_tiling.cpp+774-0
@@ -0,0 +1,774 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_tiling.cpp
13+ * \brief
14+ */
15+#include "log/log.h"
16+#include "tiling/platform/platform_ascendc.h"
17+#include "register/op_impl_registry.h"
18+#include <graph/utils/type_utils.h>
19+#include "../op_kernel/tile_tiling_data.h"
20+#include "../op_kernel/tile_tiling_key.h"
21+#include <vector>
22+#include <cstring>
23+ 
24+namespace optiling {
25+ 
26+// ========== Hardware / platform constants ==========
27+// UB is split into two buffers (input + output) for ping-pong operation
28+constexpr int32_t TILE_BUFFER_NUM = 2;
F
Ffulltower5月11日

🟡 建议：魔鬼数字较多，缺少注释说明

大量constexpr定义的阈值常量，虽然已定义为常量，但缺少注释说明选择依据。

示例：

  • SMALL_INNER_DIM_THRESHOLD = 24
  • LARGE_INNER_DIM_THRESHOLD = 64
  • MIN_PERIOD_RATIO_FOR_BENEFIT = 4

建议：补充注释说明这些阈值的来源和选择依据（如性能测试结果、硬件特性等），提高代码可维护性。

likedislike
刘十一
刘十一
5月12日 评论：
29+// Ascend DMA minimum transfer granularity: 32 bytes (one data block)
30+constexpr int32_t DATA_BLOCK_BYTES = 32;
31+// Gather offset table is uint32-aligned, grouped in units of 8 elements
32+constexpr int32_t BYTE_ALIGN_SIZE = 8;
33+// Vector register width is 256B: holds 128 fp16 elements or 64 fp32 elements
34+constexpr int32_t VREG_WIDTH_FP16 = 128;
35+constexpr int32_t VREG_WIDTH_FP32 = 64;
36+// Each Gather offset table entry occupies 4 bytes (uint32)
37+constexpr int32_t OFFSET_UINT32_SIZE = 4;
38+ 
39+// ========== Data type byte sizes ==========
40+constexpr int32_t ELEM_BYTES_INT8 = 1;
41+constexpr int32_t ELEM_BYTES_FP16 = 2;
42+constexpr int32_t ELEM_BYTES_FP32 = 4;
43+ 
44+// ========== Scheduling mode selection thresholds ==========
45+// The following thresholds are derived from performance profiling on Atlas A2 (Ascend 910B)
46+// and are used to choose the most efficient processing path at runtime.
47+ 
48+// innerDim element-count thresholds that control whether the ScalarBuild path is enabled:
49+// <=24: enabled for 1/2-byte types (SetValue loop overhead is acceptable)
50+// <=32: relaxed for 1/2-byte types with even innerDim (uint32 packing optimization applies)
51+// <=64: relaxed for 4-byte types (per-element SetValue cost is relatively lower)
52+constexpr int32_t SMALL_INNER_DIM_THRESHOLD = 24;
53+constexpr int32_t MODERATE_INNER_DIM_THRESHOLD = 32;
54+constexpr int32_t LARGE_INNER_DIM_THRESHOLD = 64;
55+ 
56+// Minimum workload for the ScalarBuild path: total input elements or per-core rows must reach 256;
57+// prevents using the scalar path on very small tensors where extra logic overhead dominates
58+constexpr int32_t MIN_INPUT_ELEMS_FOR_SCALAR = 256;
59+constexpr int32_t MIN_ROWS_PER_CORE_FOR_SCALAR = 256;
60+ 
61+// Minimum amplification ratio (repeatPeriod / repeatInputPeriod) for seed-amplify to pay off:
62+// >=4: general paths (one seed write reused >=4 times justifies the extra read-back)
63+// >=3: VecGather path has a lower bar (Gather instruction overhead is high, needs less reuse)
64+constexpr int32_t MIN_PERIOD_RATIO_FOR_BENEFIT = 4;
65+constexpr int32_t MIN_PERIOD_RATIO_FOR_VECGATHER = 3;
66+ 
67+// VecGather path constraints:
68+// seed rows <=40 (offset table must stay in UB; too many rows squeeze data space)
69+// output row width >=125 elements (short rows make Gather startup cost disproportionate)
70+// only 2/4-byte types supported (Gather instruction requires element size >= 2 bytes)
71+constexpr int32_t MAX_PERIOD_FOR_VECGATHER = 40;
72+constexpr int32_t MIN_OUTPUT_INNER_DIM_FOR_VECGATHER = 125;
73+constexpr int32_t MIN_ELEM_BYTES_FOR_VECGATHER = 2;
74+constexpr int32_t MAX_ELEM_BYTES_FOR_VECGATHER = 4;
75+ 
76+// Readback path requires innerDim >= 12 elements so that each DMA transfer covers
77+// several data blocks; below this threshold DMA startup cost exceeds scalar build
78+constexpr int32_t MIN_INNER_DIM_FOR_READBACK = 12;
79+ 
80+// DmaBuild path requires innerDim >= 11 and seed rows >= 50:
81+// too-small innerDim makes per-row DMA writes inefficient
82+// too few seed rows means read-back amplification cannot recoup the extra GM I/O
83+constexpr int32_t MIN_INNER_DIM_FOR_DMABUILD = 11;
84+constexpr int32_t MIN_REP_INPUT_PERIOD_FOR_DMABUILD = 50;
85+ 
86+// ========== Multi-core grouping / period reuse thresholds ==========
87+ 
88+// Amplification >= 3 needed for source-grouping (periodsPerSource) to outweigh management overhead
89+constexpr int32_t MIN_AMPLIFICATION_RATIO = 3;
90+// When amplification <= 2, try to find a dimension with higher amplification
91+constexpr int32_t MIN_AMP_THRESHOLD = 2;
92+ 
93+// Unique-source count thresholds for enabling source-grouping optimization:
94+// >= 8: standard bar — enough sources to distribute evenly across cores
95+// >= 4: relaxed for int8 / small-element types (small per-transfer volume needs more parallelism)
96+constexpr int32_t MIN_NSRC_THRESHOLD = 8;
97+constexpr int32_t MIN_NSRC_SMALL = 4;
98+ 
99+// Periods-per-source (pps) range: too large causes excessive serial work per core
100+constexpr int32_t MAX_PPS_THRESHOLD = 20;
101+// int8 types require pps >= 8 before small-source grouping is worthwhile
102+constexpr int32_t MIN_PPS_THRESHOLD = 8;
103+ 
104+// Source grouping only shows clear benefit when seed rows (repeatInputPeriod) >= 100
105+constexpr int32_t MIN_RIP_FOR_BENEFIT = 100;
106+ 
107+// ========== Output size thresholds ==========
108+ 
109+// Total output <= 2 KB: force single core (multi-core scheduling overhead exceeds computation)
110+constexpr int32_t SMALL_OUTPUT_BYTES_THRESHOLD = 2048;
111+// Total output <= 4 KB: cap core count in splitByMult mode to avoid over-partitioning
112+constexpr int32_t MEDIUM_OUTPUT_BYTES_THRESHOLD = 4096;
113+ 
114+// Seed template <= 256 B is considered too small; cap blockDim to AIC core count
115+// to avoid repeated scheduling overhead for tiny templates
116+constexpr int32_t SMALL_TEMPLATE_BYTES = 256;
117+// Total output >= 64 KB: Readback path DMA read-back amplification has enough payoff
118+constexpr int32_t LARGE_OUTPUT_BYTES_THRESHOLD = 65536;
119+ 
120+// In BuildOnce mode, merge cores only when period count >= 10 to ensure >= 2 periods per core
121+constexpr int32_t MIN_PERIODS_FOR_MERGE = 10;
122+ 
123+struct TileCompileInfo {};
124+ 
125+static ge::graphStatus TilingParseForTile([[maybe_unused]] gert::TilingParseContext* context)
126+{
127+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
128+ return ge::GRAPH_SUCCESS;
129+}
130+ 
131+static ge::graphStatus GetPlatformInfo(
132+ gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum, int64_t& aicCoreNum)
133+{
134+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
135+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
136+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
137+ coreNum = ascendcPlatform.GetCoreNum();
138+ aicCoreNum = ascendcPlatform.GetCoreNumAic();
139+ if (aicCoreNum <= 0) {
140+ aicCoreNum = coreNum;
141+ }
142+ OP_CHECK_IF(coreNum <= 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
143+ OP_CHECK_IF(ubSize <= 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
144+ return ge::GRAPH_SUCCESS;
145+}
146+ 
147+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
148+{
149+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
150+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
151+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
152+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
153+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
154+ currentWorkspace[0] = sysWorkspaceSize;
155+ return ge::GRAPH_SUCCESS;
156+}
157+ 
158+static void MergeConsecutiveDims(std::vector<int32_t>& shape, std::vector<int32_t>& multiples)
159+{
160+ bool changed = true;
161+ while (changed) {
162+ changed = false;
163+ for (int32_t idx = static_cast<int32_t>(shape.size()) - 1; idx >= 1; idx--) {
164+ if (multiples[idx] == 1) {
165+ shape[idx - 1] *= shape[idx];
166+ shape.erase(shape.begin() + idx);
167+ multiples.erase(multiples.begin() + idx);
168+ changed = true;
169+ break;
170+ }
171+ }
172+ }
173+}
174+ 
175+static void AlignLastDimForDMA(std::vector<int32_t>& shape, std::vector<int32_t>& multiples, int32_t elemBytes)
176+{
177+ while (shape.size() >= 2) {
L
Lloov4月25日

建议不要用魔鬼数学

likedislike
刘十一
刘十一
4月25日 评论：
178+ int32_t lastIdx = static_cast<int32_t>(shape.size()) - 1;
179+ if (shape[lastIdx] * elemBytes % DATA_BLOCK_BYTES == 0) {
180+ break;
181+ }
182+ int32_t secLast = lastIdx - 1;
183+ if (multiples[secLast] != 1) {
184+ break;
185+ }
186+ shape[lastIdx] = shape[secLast] * shape[lastIdx];
187+ shape.erase(shape.begin() + secLast);
188+ multiples.erase(multiples.begin() + secLast);
189+ }
190+}
191+ 
192+static void RemoveRedundantDims(std::vector<int32_t>& shape, std::vector<int32_t>& multiples)
193+{
194+ std::vector<int32_t> newShape;
195+ std::vector<int32_t> newMult;
196+ for (size_t idx = 0; idx < shape.size(); idx++) {
197+ if (shape[idx] == 1 && multiples[idx] == 1) {
198+ continue;
199+ }
200+ if (!newShape.empty() && newShape.back() == 1 && shape[idx] == 1) {
201+ newMult.back() *= multiples[idx];
202+ } else {
203+ newShape.push_back(shape[idx]);
204+ newMult.push_back(multiples[idx]);
205+ }
206+ }
207+ if (newShape.empty()) {
208+ newShape.push_back(1);
209+ newMult.push_back(1);
210+ }
211+ shape = newShape;
212+ multiples = newMult;
213+}
214+ 
215+static void MergeDims(std::vector<int32_t>& shape, std::vector<int32_t>& multiples, int32_t elemBytes)
216+{
217+ MergeConsecutiveDims(shape, multiples);
218+ AlignLastDimForDMA(shape, multiples, elemBytes);
219+ RemoveRedundantDims(shape, multiples);
220+}
221+ 
222+static void ComputeTilingParams(
223+ TileTilingData& tilingData, const std::vector<int32_t>& origShape, const std::vector<int32_t>& origMult,
224+ int32_t elemBytes, int32_t blockDim, int32_t ubSize, int32_t aicBlockDim)
225+{
226+ (void)aicBlockDim;
227+ std::vector<int32_t> shape = origShape;
228+ std::vector<int32_t> multiples = origMult;
229+ MergeDims(shape, multiples, elemBytes);
230+ 
231+ int32_t ndim = static_cast<int32_t>(shape.size());
232+ tilingData.numDims = ndim;
233+ tilingData.elemBytes = elemBytes;
234+ tilingData.blockDim = blockDim;
235+ tilingData.ubSize = ubSize;
236+ tilingData.totalInputElems = 1;
237+ tilingData.totalOutputElems = 1;
238+ for (int32_t idx = 0; idx < ndim; idx++) {
239+ tilingData.inputShape[idx] = shape[idx];
240+ tilingData.multiples[idx] = multiples[idx];
241+ tilingData.outputShape[idx] = shape[idx] * multiples[idx];
242+ tilingData.totalInputElems *= shape[idx];
243+ tilingData.totalOutputElems *= tilingData.outputShape[idx];
244+ }
245+ tilingData.inputStrides[ndim - 1] = 1;
246+ tilingData.outputStrides[ndim - 1] = 1;
247+ for (int32_t idx = ndim - 2; idx >= 0; idx--) {
248+ tilingData.inputStrides[idx] = tilingData.inputStrides[idx + 1] * tilingData.inputShape[idx + 1];
249+ tilingData.outputStrides[idx] = tilingData.outputStrides[idx + 1] * tilingData.outputShape[idx + 1];
250+ }
251+ 
252+ int32_t innerDim = tilingData.inputShape[ndim - 1];
253+ int32_t innerMult = tilingData.multiples[ndim - 1];
254+ int32_t outputInnerDim = innerDim * innerMult;
255+ int32_t outerCount = (outputInnerDim > 0) ? (tilingData.totalOutputElems / outputInnerDim) : 0;
256+ int32_t alignElems = (elemBytes > 0) ? (DATA_BLOCK_BYTES / elemBytes) : 1;
257+ if (alignElems < 1) {
258+ alignElems = 1;
259+ }
260+ int32_t innerDimAligned = ((innerDim + alignElems - 1) / alignElems) * alignElems;
261+ int32_t bufElems = (elemBytes > 0) ? (ubSize / elemBytes) : 0;
262+ bool splitByMult = (outerCount < blockDim && innerMult > 1 && innerDimAligned <= bufElems);
263+ int32_t totalWork = splitByMult ? (outerCount * innerMult) : outerCount;
264+ if (!splitByMult && innerDimAligned > bufElems && outerCount < blockDim) {
265+ int32_t chunkSize = (bufElems / alignElems) * alignElems;
266+ if (chunkSize > 0) {
267+ int32_t numChunks = (innerDim + chunkSize - 1) / chunkSize;
268+ totalWork = outerCount * numChunks;
269+ }
270+ }
271+ if (totalWork < blockDim) {
272+ tilingData.blockDim = (totalWork > 0) ? totalWork : 1;
273+ }
274+ if (!splitByMult) {
275+ int64_t outputBytes = static_cast<int64_t>(tilingData.totalOutputElems) * elemBytes;
276+ if (outputBytes <= SMALL_OUTPUT_BYTES_THRESHOLD && tilingData.blockDim > 1) {
277+ tilingData.blockDim = 1;
278+ }
279+ }
280+ if (splitByMult && tilingData.blockDim > outerCount && outerCount >= 2) {
281+ int64_t outputBytes = static_cast<int64_t>(tilingData.totalOutputElems) * elemBytes;
282+ if (outputBytes <= MEDIUM_OUTPUT_BYTES_THRESHOLD ||
283+ (innerDim * elemBytes >= DATA_BLOCK_BYTES &&
284+ static_cast<int64_t>(outerCount) * outputInnerDim * elemBytes <= ubSize)) {
285+ tilingData.blockDim = outerCount;
286+ }
287+ }
288+ 
289+ tilingData.repeatPeriod = 0;
290+ tilingData.repeatInputPeriod = 0;
291+ tilingData.periodsPerSource = 0;
292+ tilingData.nUniqueSources = 0;
293+ if (!splitByMult && ndim >= 2) {
294+ int32_t rowsPerCore = (outerCount + tilingData.blockDim - 1) / tilingData.blockDim;
295+ int32_t bestIPeriod = rowsPerCore + 1;
296+ int32_t bestRepeatDim = -1;
297+ for (int32_t d = ndim - 2; d >= 0; d--) {
298+ if (multiples[d] <= 1) {
299+ continue;
300+ }
301+ int32_t stride_d = 1;
302+ for (int32_t dd = d + 1; dd < ndim - 1; dd++) {
303+ stride_d *= tilingData.outputShape[dd];
304+ }
305+ int32_t iPeriod = shape[d] * stride_d;
306+ int32_t oPeriod = tilingData.outputShape[d] * stride_d;
307+ if (iPeriod <= 0 || iPeriod > rowsPerCore) {
308+ continue;
309+ }
310+ if (oPeriod <= iPeriod) {
311+ continue;
312+ }
313+ if (iPeriod < bestIPeriod) {
314+ bestIPeriod = iPeriod;
315+ tilingData.repeatPeriod = oPeriod;
316+ tilingData.repeatInputPeriod = iPeriod;
317+ bestRepeatDim = d;
318+ }
319+ }
320+ if (bestRepeatDim >= 0) {
321+ int32_t bestAmp = tilingData.repeatPeriod / tilingData.repeatInputPeriod;
322+ if (bestAmp <= MIN_AMP_THRESHOLD) {
323+ for (int32_t d = ndim - 2; d >= 0; d--) {
324+ if (multiples[d] <= 1) {
325+ continue;
326+ }
327+ int32_t stride_d = 1;
328+ for (int32_t dd = d + 1; dd < ndim - 1; dd++) {
329+ stride_d *= tilingData.outputShape[dd];
330+ }
331+ int32_t iPeriod = shape[d] * stride_d;
332+ int32_t oPeriod = tilingData.outputShape[d] * stride_d;
333+ if (iPeriod <= 0 || iPeriod > rowsPerCore || oPeriod <= iPeriod) {
334+ continue;
335+ }
336+ int32_t amp = oPeriod / iPeriod;
337+ if (amp <= bestAmp) {
338+ continue;
339+ }
340+ int32_t totalP = outerCount / oPeriod;
341+ if (totalP < tilingData.blockDim / MIN_NSRC_THRESHOLD) {
342+ continue;
343+ }
344+ tilingData.repeatPeriod = oPeriod;
345+ tilingData.repeatInputPeriod = iPeriod;
346+ bestRepeatDim = d;
347+ break;
348+ }
349+ }
350+ int32_t pps = 1;
351+ bool ppsValid = true;
352+ for (int32_t dd = 0; dd < bestRepeatDim; dd++) {
353+ if (shape[dd] == 1) {
354+ pps *= multiples[dd];
355+ } else {
356+ ppsValid = false;
357+ break;
358+ }
359+ }
360+ if (!ppsValid) {
361+ pps = 1;
362+ int32_t candidatePPS = 1;
363+ for (int32_t dd = 0; dd < bestRepeatDim; dd++) {
364+ candidatePPS *= multiples[dd];
365+ }
366+ int32_t nUnique = 1;
367+ for (int32_t dd = 0; dd < bestRepeatDim; dd++) {
368+ nUnique *= shape[dd];
369+ }
370+ int32_t rP_local = tilingData.repeatPeriod;
371+ int32_t totalP_local = (rP_local > 0) ? (outerCount / rP_local) : 0;
372+ int32_t candidateNSrc = (candidatePPS > 0) ? (totalP_local / candidatePPS) : 0;
373+ if (candidateNSrc > 1 && candidateNSrc == nUnique && totalP_local == candidateNSrc * candidatePPS) {
374+ int32_t checkIdx = candidateNSrc * rP_local;
375+ int32_t srcOff1 = 0;
376+ int32_t tmpVal = checkIdx;
377+ for (int32_t d2 = ndim - 2; d2 >= 0; d2--) {
378+ int32_t oc = tmpVal % tilingData.outputShape[d2];
379+ tmpVal /= tilingData.outputShape[d2];
380+ srcOff1 += (oc % tilingData.inputShape[d2]) * tilingData.inputStrides[d2];
381+ }
382+ if (srcOff1 == 0) {
383+ pps = candidatePPS;
384+ }
385+ }
386+ }
387+ int32_t riP = tilingData.repeatInputPeriod;
388+ int32_t rP = tilingData.repeatPeriod;
389+ int32_t totalP = (rP > 0) ? (outerCount / rP) : 0;
390+ int32_t nSrc = (pps > 0 && totalP > 0 && totalP % pps == 0) ? totalP / pps : 0;
391+ if ((nSrc >= MIN_NSRC_THRESHOLD && nSrc <= tilingData.blockDim && riP >= MIN_RIP_FOR_BENEFIT &&
392+ pps <= MAX_PPS_THRESHOLD) ||
393+ (nSrc >= MIN_NSRC_THRESHOLD && nSrc <= tilingData.blockDim && elemBytes == ELEM_BYTES_INT8) ||
394+ (nSrc >= MIN_NSRC_SMALL && nSrc <= tilingData.blockDim && elemBytes == ELEM_BYTES_INT8 &&
395+ pps >= MIN_PPS_THRESHOLD) ||
396+ (nSrc >= MIN_NSRC_SMALL && nSrc <= tilingData.blockDim && pps > MIN_NSRC_SMALL &&
397+ (rP / riP) >= MIN_AMPLIFICATION_RATIO && static_cast<int64_t>(rP) * outputInnerDim <= bufElems) ||
398+ (nSrc >= MIN_NSRC_THRESHOLD && nSrc <= tilingData.blockDim && pps <= MAX_PPS_THRESHOLD &&
399+ totalP >= tilingData.blockDim * 2) ||
400+ (nSrc > tilingData.blockDim && elemBytes <= ELEM_BYTES_FP16 && pps <= MAX_PPS_THRESHOLD &&
401+ (nSrc + tilingData.blockDim - 1) / tilingData.blockDim <= pps)) {
402+ tilingData.periodsPerSource = pps;
403+ }
404+ if (bestRepeatDim >= 0) {
405+ int32_t nUniqueCalc = 1;
406+ for (int32_t dd = 0; dd < bestRepeatDim; dd++) {
407+ nUniqueCalc *= shape[dd];
408+ }
409+ if (nUniqueCalc > 1 && nUniqueCalc < totalP) {
410+ tilingData.nUniqueSources = nUniqueCalc;
411+ }
412+ }
413+ }
414+ if (bestRepeatDim < 0 && innerDimAligned <= bufElems) {
415+ bool smallID =
416+ (innerDim <= SMALL_INNER_DIM_THRESHOLD ||
417+ (elemBytes >= ELEM_BYTES_FP32 && innerDim <= LARGE_INNER_DIM_THRESHOLD));
418+ int32_t candidateIP = outerCount + 1;
419+ for (int32_t d = ndim - 2; d >= 0; d--) {
420+ if (multiples[d] <= 1) {
421+ continue;
422+ }
423+ int32_t stride_d = 1;
424+ for (int32_t dd = d + 1; dd < ndim - 1; dd++) {
425+ stride_d *= tilingData.outputShape[dd];
426+ }
427+ int32_t iPeriod = shape[d] * stride_d;
428+ int32_t oPeriod = tilingData.outputShape[d] * stride_d;
429+ if (iPeriod <= 0 || oPeriod <= iPeriod) {
430+ continue;
431+ }
432+ int32_t amp = oPeriod / iPeriod;
433+ if (amp < MIN_AMPLIFICATION_RATIO) {
434+ continue;
435+ }
436+ int32_t neededBD = outerCount / iPeriod;
437+ if (neededBD < 1)
438+ neededBD = 1;
439+ if (neededBD >= tilingData.blockDim) {
440+ continue;
441+ }
442+ if (!smallID) {
443+ continue;
444+ }
445+ if (iPeriod < candidateIP) {
446+ candidateIP = iPeriod;
447+ tilingData.repeatPeriod = oPeriod;
448+ tilingData.repeatInputPeriod = iPeriod;
449+ bestRepeatDim = d;
450+ tilingData.blockDim = neededBD;
451+ }
452+ }
453+ }
454+ }
455+ if (!splitByMult && tilingData.blockDim > 1 && outerCount > 0 && innerDimAligned <= bufElems) {
456+ int32_t rPC = (outerCount + tilingData.blockDim - 1) / tilingData.blockDim;
457+ int32_t rowBytes = outputInnerDim * elemBytes;
458+ if (rowBytes > 0 && rowBytes % DATA_BLOCK_BYTES != 0 && rPC > 1) {
459+ int32_t g = rowBytes, bv = DATA_BLOCK_BYTES;
460+ while (bv) {
461+ int32_t t = bv;
462+ bv = g % bv;
463+ g = t;
464+ }
465+ int32_t aStep = DATA_BLOCK_BYTES / g;
466+ if (aStep > 1 && rPC > aStep) {
467+ rPC = ((rPC + aStep - 1) / aStep) * aStep;
468+ }
469+ }
470+ int32_t effBD = (rPC > 0) ? ((outerCount + rPC - 1) / rPC) : 1;
471+ if (effBD > 0 && effBD < tilingData.blockDim) {
472+ tilingData.blockDim = effBD;
473+ }
474+ }
475+}
476+ 
477+static ge::graphStatus TileTilingFunc(gert::TilingContext* context)
478+{
479+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
480+ 
481+ uint64_t ubSize;
482+ int64_t coreNum;
483+ int64_t aicCoreNum;
484+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, coreNum, aicCoreNum);
485+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
486+ 
487+ const gert::StorageShape* inputShape = context->GetInputShape(0);
488+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
489+ 
490+ auto inputDesc = context->GetInputDesc(0);
491+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
492+ ge::DataType dtype = inputDesc->GetDataType();
493+ uint32_t typeLength = 0;
494+ ge::TypeUtils::GetDataTypeLength(dtype, typeLength);
495+ int32_t elemBytes = static_cast<int32_t>(typeLength);
496+ OP_CHECK_IF(elemBytes <= 0, OP_LOGE(context, "unsupported data type, elemBytes <= 0"), return ge::GRAPH_FAILED);
497+ 
498+ int32_t ndim = static_cast<int32_t>(inputShape->GetStorageShape().GetDimNum());
499+ std::vector<int32_t> shape(ndim);
500+ for (int32_t idx = 0; idx < ndim; idx++) {
501+ shape[idx] = static_cast<int32_t>(inputShape->GetStorageShape().GetDim(idx));
502+ }
503+ 
504+ const gert::Tensor* multiplesTensor = context->GetInputTensor(1);
505+ OP_CHECK_NULL_WITH_CONTEXT(context, multiplesTensor);
506+ const void* multiplesAddr = multiplesTensor->GetAddr();
507+ OP_CHECK_NULL_WITH_CONTEXT(context, multiplesAddr);
508+ 
509+ auto multiplesDesc = context->GetInputDesc(1);
510+ OP_CHECK_NULL_WITH_CONTEXT(context, multiplesDesc);
511+ ge::DataType multiplesDtype = multiplesDesc->GetDataType();
512+ 
513+ const gert::StorageShape* multiplesShape = context->GetInputShape(1);
514+ OP_CHECK_NULL_WITH_CONTEXT(context, multiplesShape);
515+ int32_t multiplesLen = static_cast<int32_t>(multiplesShape->GetStorageShape().GetDim(0));
516+ std::vector<int32_t> mult(multiplesLen);
517+ if (multiplesDtype == ge::DT_INT64) {
518+ const int64_t* multiplesData64 = reinterpret_cast<const int64_t*>(multiplesAddr);
519+ for (int32_t idx = 0; idx < multiplesLen; idx++) {
520+ mult[idx] = static_cast<int32_t>(multiplesData64[idx]);
521+ }
522+ } else {
523+ const int32_t* multiplesData32 = reinterpret_cast<const int32_t*>(multiplesAddr);
524+ for (int32_t idx = 0; idx < multiplesLen; idx++) {
525+ mult[idx] = multiplesData32[idx];
526+ }
527+ }
528+ 
529+ while (static_cast<int32_t>(shape.size()) < multiplesLen) {
530+ shape.insert(shape.begin(), 1);
531+ }
532+ while (static_cast<int32_t>(mult.size()) < static_cast<int32_t>(shape.size())) {
533+ mult.insert(mult.begin(), 1);
534+ }
535+ 
536+ int32_t blockDim = static_cast<int32_t>(coreNum);
537+ int32_t aicBlockDim = static_cast<int32_t>(aicCoreNum);
538+ int32_t bufferSize = static_cast<int32_t>(ubSize) / TILE_BUFFER_NUM;
539+ 
540+ TileTilingData* tiling = context->GetTilingData<TileTilingData>();
541+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
542+ OP_CHECK_IF(
543+ memset_s(tiling, sizeof(TileTilingData), 0, sizeof(TileTilingData)) != EOK,
544+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
545+ 
546+ ComputeTilingParams(*tiling, shape, mult, elemBytes, blockDim, bufferSize, aicBlockDim);
547+ 
548+ OP_CHECK_IF(
549+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
550+ return ge::GRAPH_FAILED);
551+ 
552+ int32_t ae = (elemBytes > 0) ? (DATA_BLOCK_BYTES / elemBytes) : 1;
553+ if (ae < 1) {
554+ ae = 1;
555+ }
556+ int32_t fID = tiling->inputShape[tiling->numDims - 1];
557+ int32_t fIM = tiling->multiples[tiling->numDims - 1];
558+ int32_t fIDA = ((fID + ae - 1) / ae) * ae;
559+ int32_t fOD = fID * fIM;
560+ bool fAligned = (fID == fIDA);
561+ int32_t fOC = (fOD > 0) ? (tiling->totalOutputElems / fOD) : 0;
562+ bool fSplit = (fOC < tiling->blockDim && fIM > 1 && fIDA <= bufferSize / elemBytes);
563+ int32_t fRowsPerCore = (tiling->blockDim > 0) ? ((fOC + tiling->blockDim - 1) / tiling->blockDim) : 0;
564+ int32_t rPre = tiling->repeatPeriod;
565+ int32_t riPre = tiling->repeatInputPeriod;
566+ bool hasRepeatBenefit =
567+ (rPre > riPre && riPre > 0 && (rPre / riPre) >= MIN_PERIOD_RATIO_FOR_BENEFIT && tiling->totalInputElems > 1);
568+ bool useScalarMode =
569+ !fSplit && !fAligned && fIM > 1 && fOD <= (bufferSize / elemBytes) &&
570+ (tiling->totalInputElems > MIN_INPUT_ELEMS_FOR_SCALAR || fRowsPerCore >= MIN_ROWS_PER_CORE_FOR_SCALAR ||
571+ hasRepeatBenefit) &&
572+ (fID <= SMALL_INNER_DIM_THRESHOLD || (elemBytes >= ELEM_BYTES_FP32 && fID <= LARGE_INNER_DIM_THRESHOLD));
573+ 
574+ int32_t rP = tiling->repeatPeriod;
575+ int32_t riP = tiling->repeatInputPeriod;
576+ int32_t fBufElems = (elemBytes > 0) ? (bufferSize / elemBytes) : 0;
577+ bool canBuildOnce =
578+ rP > riP && riP > 1 && (rP / riP) > 1 && riP * fOD <= fBufElems && (riP - 1) * fID + fIDA <= fBufElems;
579+ if (!useScalarMode && canBuildOnce && !fSplit && !fAligned && fIM > 1 && fOD <= (bufferSize / elemBytes) &&
580+ (tiling->totalInputElems > MIN_INPUT_ELEMS_FOR_SCALAR || fRowsPerCore >= MIN_ROWS_PER_CORE_FOR_SCALAR ||
581+ hasRepeatBenefit) &&
582+ elemBytes <= ELEM_BYTES_FP16 && fID <= MODERATE_INNER_DIM_THRESHOLD && fID % 2 == 0) {
583+ useScalarMode = true;
584+ }
585+ if (!useScalarMode && canBuildOnce && !fSplit && !fAligned && fIM > 1 && fOD <= (bufferSize / elemBytes) &&
586+ hasRepeatBenefit && elemBytes == ELEM_BYTES_INT8 && fID <= LARGE_INNER_DIM_THRESHOLD) {
587+ useScalarMode = true;
588+ }
589+ bool useBuildOnce = useScalarMode && canBuildOnce;
590+ 
591+ bool useVecGather = false;
592+ if (useBuildOnce && elemBytes >= MIN_ELEM_BYTES_FOR_VECGATHER && elemBytes <= MAX_ELEM_BYTES_FOR_VECGATHER &&
593+ fIM > 1 && !fAligned && (rP / riP) >= MIN_PERIOD_RATIO_FOR_VECGATHER && riP <= MAX_PERIOD_FOR_VECGATHER &&
594+ fID * fIM >= MIN_OUTPUT_INNER_DIM_FOR_VECGATHER) {
595+ int32_t vreg = (elemBytes == ELEM_BYTES_FP16) ? VREG_WIDTH_FP16 : VREG_WIDTH_FP32;
596+ int32_t gRow = (fOD > vreg) ? (((fOD + ae - 1) / ae) * ae) : vreg;
597+ int32_t offU = ((fOD + BYTE_ALIGN_SIZE - 1) / BYTE_ALIGN_SIZE) * BYTE_ALIGN_SIZE;
598+ int32_t offT = (offU * OFFSET_UINT32_SIZE + elemBytes - 1) / elemBytes;
599+ int32_t offA = ((offT + ae - 1) / ae) * ae;
600+ int32_t total = offA + riP * gRow;
601+ if (total <= fBufElems && riP * fIDA <= fBufElems) {
602+ useVecGather = true;
603+ }
604+ }
605+ 
606+ bool useDmaBuild = useBuildOnce && elemBytes >= ELEM_BYTES_FP32 && (fID % 2) == 1 &&
607+ fID >= MIN_INNER_DIM_FOR_DMABUILD && riP >= MIN_REP_INPUT_PERIOD_FOR_DMABUILD &&
608+ riP * fIDA <= fBufElems;
609+ 
610+ bool readbackOverride = false;
611+ if (useDmaBuild && rP > riP && riP > 0 && (rP / riP) >= MIN_PERIOD_RATIO_FOR_BENEFIT && riP * fOD <= fBufElems &&
612+ fOD <= fBufElems) {
613+ readbackOverride = true;
614+ }
615+ 
616+ int32_t gRP = rPre, gRiP = riPre;
617+ if (gRP <= gRiP && tiling->numDims >= 2) {
618+ int32_t bestIP = INT32_MAX;
619+ for (int32_t d = tiling->numDims - 2; d >= 0; d--) {
620+ if (tiling->multiples[d] <= 1)
621+ continue;
622+ int32_t sd = 1;
623+ for (int32_t dd = d + 1; dd < tiling->numDims - 1; dd++)
624+ sd *= tiling->outputShape[dd];
625+ int32_t ip = tiling->inputShape[d] * sd;
626+ int32_t op = tiling->outputShape[d] * sd;
627+ if (ip <= 0 || op <= ip)
628+ continue;
629+ if (ip < bestIP) {
630+ bestIP = ip;
631+ gRP = op;
632+ gRiP = ip;
633+ }
634+ }
635+ }
636+ 
637+ bool useReadbackAmplify = !useScalarMode && !fSplit && !fAligned && fIM > 1 && gRP > gRiP && gRiP > 0 &&
638+ (gRP / gRiP) >= MIN_PERIOD_RATIO_FOR_BENEFIT && fOD <= (bufferSize / elemBytes) &&
639+ fID > MIN_INNER_DIM_FOR_READBACK && static_cast<int64_t>(fIDA) * fIM <= fBufElems &&
640+ (static_cast<int64_t>(fID) * elemBytes <= SMALL_TEMPLATE_BYTES ||
641+ static_cast<int64_t>(fOC) * fOD * elemBytes >= LARGE_OUTPUT_BYTES_THRESHOLD);
642+ 
643+ bool useInnerOneFill = fSplit && fID == 1 && fIM > ae;
644+ 
645+ bool buildOnceToReadback = false;
646+ if (useBuildOnce && gRP > gRiP && gRiP > 0) {
647+ int32_t totalPeriods_bo = (gRP > 0) ? (fOC / gRP) : 0;
648+ int32_t totalSeeds_bo = totalPeriods_bo * gRiP;
649+ int32_t seedsPerCore_bo =
650+ (tiling->blockDim > 0) ? (totalSeeds_bo + tiling->blockDim - 1) / tiling->blockDim : 0;
651+ if (totalPeriods_bo > 0 && totalPeriods_bo < tiling->blockDim / 2 && seedsPerCore_bo >= MIN_NSRC_THRESHOLD &&
652+ fOD <= fBufElems && static_cast<int64_t>(fIDA) * fIM <= fBufElems) {
653+ useBuildOnce = false;
654+ buildOnceToReadback = true;
655+ }
656+ }
657+ 
658+ if (useInnerOneFill) {
659+ tiling->blockDim = 1;
660+ }
661+ 
662+ uint64_t tilingKey;
663+ if (readbackOverride) {
664+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_READBACK);
665+ } else if (useDmaBuild) {
666+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_DMABUILD);
667+ } else if (useVecGather) {
668+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_VECGATHER);
669+ } else if (useBuildOnce) {
670+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_BUILDONCE);
671+ } else if (buildOnceToReadback) {
672+ if (gRP != rPre || gRiP != riPre) {
673+ tiling->repeatPeriod = gRP;
674+ tiling->repeatInputPeriod = gRiP;
675+ }
676+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_READBACK);
677+ if (tiling->blockDim * MIN_AMPLIFICATION_RATIO < static_cast<int32_t>(coreNum)) {
678+ tiling->blockDim = static_cast<int32_t>(coreNum);
679+ }
680+ } else if (useInnerOneFill) {
681+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_DMABUILD);
682+ } else if (useReadbackAmplify) {
683+ if (gRP != rPre || gRiP != riPre) {
684+ tiling->repeatPeriod = gRP;
685+ tiling->repeatInputPeriod = gRiP;
686+ }
687+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_READBACK);
688+ } else {
689+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_DEFAULT);
690+ if (!fSplit && !fAligned && fIM > 1 && fOD <= fBufElems && tiling->numDims >= 2 &&
691+ (fID <= MIN_INNER_DIM_FOR_READBACK || (elemBytes >= ELEM_BYTES_FP32 && fID <= LARGE_INNER_DIM_THRESHOLD))) {
692+ int32_t bestIP2 = INT32_MAX;
693+ int32_t foundRP = 0, foundRiP = 0;
694+ for (int32_t d = tiling->numDims - 2; d >= 0; d--) {
695+ if (tiling->multiples[d] <= 1)
696+ continue;
697+ int32_t sd = 1;
698+ for (int32_t dd = d + 1; dd < tiling->numDims - 1; dd++)
699+ sd *= tiling->outputShape[dd];
700+ int32_t ip = tiling->inputShape[d] * sd;
701+ int32_t op = tiling->outputShape[d] * sd;
702+ if (ip <= 0 || op <= ip)
703+ continue;
704+ if ((op / ip) < MIN_PERIOD_RATIO_FOR_BENEFIT)
705+ continue;
706+ bool bo = ip > 1 && ip * fOD <= fBufElems && (ip - 1) * fID + fIDA <= fBufElems;
707+ if (!bo)
708+ continue;
709+ int32_t tp = fOC / op;
710+ if (tp < 1)
711+ continue;
712+ if (ip < bestIP2) {
713+ bestIP2 = ip;
714+ foundRP = op;
715+ foundRiP = ip;
716+ }
717+ }
718+ if (foundRP > 0 && foundRiP > 0) {
719+ tiling->repeatPeriod = foundRP;
720+ tiling->repeatInputPeriod = foundRiP;
721+ tilingKey = GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_BUILDONCE);
722+ }
723+ }
724+ }
725+ context->SetTilingKey(tilingKey);
726+ 
727+ bool isPeriodBased = tilingKey == GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_BUILDONCE) ||
728+ tilingKey == GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_DMABUILD) ||
729+ tilingKey == GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_VECGATHER);
730+ if (isPeriodBased && tiling->repeatPeriod > 0) {
731+ int32_t totalPeriods = fOC / tiling->repeatPeriod;
732+ int32_t remainder = fOC - totalPeriods * tiling->repeatPeriod;
733+ int32_t pps = tiling->periodsPerSource;
734+ int32_t neededCores = 0;
735+ if (pps > 0 && totalPeriods > 0 && totalPeriods % pps == 0) {
736+ int32_t nSrc = totalPeriods / pps;
737+ if (nSrc > 0 && nSrc < tiling->blockDim) {
738+ neededCores = nSrc + (remainder > 0 ? 1 : 0);
739+ }
740+ } else if (pps == 0) {
741+ neededCores = totalPeriods + (remainder > 0 ? 1 : 0);
742+ }
743+ if (neededCores > 0 && neededCores < tiling->blockDim) {
744+ int32_t periodsPerCore = (totalPeriods + neededCores - 1) / neededCores;
745+ if (periodsPerCore < 2 && tilingKey == GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_BUILDONCE) &&
746+ totalPeriods > MIN_PERIODS_FOR_MERGE) {
747+ neededCores = (totalPeriods + 1) / 2;
748+ }
749+ if (neededCores > 0 && neededCores < tiling->blockDim) {
750+ tiling->blockDim = neededCores;
751+ }
752+ }
753+ }
754+ 
755+ if (isPeriodBased && tilingKey == GET_TPL_TILING_KEY(TILE_TPL_SCH_MODE_BUILDONCE) && tiling->repeatPeriod > 0 &&
756+ tiling->repeatInputPeriod > 0) {
757+ int32_t boRiP = tiling->repeatInputPeriod;
758+ int32_t boOID = tiling->inputShape[tiling->numDims - 1] * tiling->multiples[tiling->numDims - 1];
759+ int64_t templateBytes = static_cast<int64_t>(boRiP) * boOID * elemBytes;
760+ if (templateBytes < SMALL_TEMPLATE_BYTES && tiling->blockDim > aicBlockDim && aicBlockDim > 0) {
761+ tiling->blockDim = aicBlockDim;
762+ }
763+ }
764+ 
765+ context->SetBlockDim(static_cast<uint32_t>(tiling->blockDim));
766+ 
767+ return ge::GRAPH_SUCCESS;
768+}
769+ 
770+IMPL_OP_OPTILING(Tile)
771+ .Tiling(TileTilingFunc)
772+ .TilingParse<TileCompileInfo>(TilingParseForTile)
773+ .InputsDataDependency({1});
774+} // namespace optiling
Aexperimental/math/tile/op_kernel/tile.cpp+27-0
@@ -0,0 +1,27 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile.cpp
13+ * \brief
14+ */
15+ 
16+#include "tile.h"
17+ 
18+template <uint32_t schMode>
19+__global__ __aicore__ void tile(GM_ADDR x, GM_ADDR multiples, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
20+{
21+ REGISTER_TILING_DEFAULT(TileTilingData);
22+ GET_TILING_DATA_WITH_STRUCT(TileTilingData, tilingData, tiling);
23+ AscendC::TPipe pipe;
24+ TileKernel::TileOpImpl<DTYPE_X> op;
25+ op.template Init<schMode>(x, y, &tilingData, &pipe);
26+ op.template Process<schMode>();
27+}
Aexperimental/math/tile/op_kernel/tile_tiling_data.h+40-0
@@ -0,0 +1,40 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+#ifndef TILE_TILING_DATA_H_
16+#define TILE_TILING_DATA_H_
17+ 
18+#include <cstdint>
19+ 
20+constexpr int32_t TILE_MAX_DIM = 8;
21+ 
22+struct TileTilingData {
23+ int32_t numDims;
24+ int32_t inputShape[TILE_MAX_DIM];
25+ int32_t multiples[TILE_MAX_DIM];
26+ int32_t outputShape[TILE_MAX_DIM];
27+ int32_t inputStrides[TILE_MAX_DIM];
28+ int32_t outputStrides[TILE_MAX_DIM];
29+ int32_t totalInputElems;
30+ int32_t totalOutputElems;
31+ int32_t elemBytes;
32+ int32_t blockDim;
33+ int32_t ubSize;
34+ int32_t repeatPeriod;
35+ int32_t repeatInputPeriod;
36+ int32_t periodsPerSource;
37+ int32_t nUniqueSources;
38+};
39+ 
40+#endif
Aexperimental/math/tile/op_kernel/tile_tiling_key.h+36-0
@@ -0,0 +1,36 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_tiling_key.h
13+ * \brief tile tiling key declare
14+ */
15+#ifndef TILE_TILING_KEY_H_
16+#define TILE_TILING_KEY_H_
17+ 
18+#include "ascendc/host_api/tiling/template_argument.h"
19+ 
20+#define TILE_TPL_SCH_MODE_DEFAULT 0
21+#define TILE_TPL_SCH_MODE_BUILDONCE 2
22+#define TILE_TPL_SCH_MODE_DMABUILD 3
23+#define TILE_TPL_SCH_MODE_READBACK 4
24+#define TILE_TPL_SCH_MODE_VECGATHER 5
25+ 
26+ASCENDC_TPL_ARGS_DECL(
27+ Tile, ASCENDC_TPL_UINT_DECL(
28+ schMode, 3, ASCENDC_TPL_UI_LIST, TILE_TPL_SCH_MODE_DEFAULT, TILE_TPL_SCH_MODE_BUILDONCE,
29+ TILE_TPL_SCH_MODE_DMABUILD, TILE_TPL_SCH_MODE_READBACK, TILE_TPL_SCH_MODE_VECGATHER));
30+ 
31+ASCENDC_TPL_SEL(
32+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(
33+ schMode, ASCENDC_TPL_UI_LIST, TILE_TPL_SCH_MODE_DEFAULT, TILE_TPL_SCH_MODE_BUILDONCE,
34+ TILE_TPL_SCH_MODE_DMABUILD, TILE_TPL_SCH_MODE_READBACK, TILE_TPL_SCH_MODE_VECGATHER)), );
35+ 
36+#endif
Aexperimental/math/tile/tests/ut/op_api/test_aclnn_repeat.cpp+67-0
@@ -0,0 +1,67 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file test_aclnn_repeat.cpp
13+ * \brief
14+ */
15+#include <vector>
16+#include <array>
17+#include "gtest/gtest.h"
18+ 
19+#include "opdev/platform.h"
20+#include "../../../op_api/aclnn_repeat.h"
21+#include "op_api_ut_common/tensor_desc.h"
22+#include "op_api_ut_common/scalar_desc.h"
23+#include "op_api_ut_common/array_desc.h"
24+#include "op_api_ut_common/op_api_ut.h"
25+ 
26+class l2_repeat_test : public testing::Test {
27+protected:
28+ static void SetUpTestCase()
29+ {
30+ cout << "l2_repeat_test SetUp" << endl;
31+ }
32+ 
33+ static void TearDownTestCase()
34+ {
35+ cout << "l2_repeat_test TearDown" << endl;
36+ }
37+};
38+ 
39+TEST_F(l2_repeat_test, case_01_float_2d)
40+{
41+ op::SetPlatformSocVersion(op::SocVersion::ASCEND910B);
42+ 
43+ auto selfDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND).Value(vector<float>{1.0, 2.0, 3.0, 4.0, 5.0, 6.0});
44+ auto repeatsDesc = IntArrayDesc(vector<int64_t>{3, 2});
45+ auto outDesc = TensorDesc({6, 6}, ACL_FLOAT, ACL_FORMAT_ND);
46+ 
47+ auto ut = OP_API_UT(aclnnRepeat, INPUT(selfDesc, repeatsDesc), OUTPUT(outDesc));
48+ 
49+ uint64_t workspace_size = 0;
50+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
51+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
52+}
53+ 
54+TEST_F(l2_repeat_test, case_02_float16_1d)
55+{
56+ op::SetPlatformSocVersion(op::SocVersion::ASCEND910B);
57+ 
58+ auto selfDesc = TensorDesc({4}, ACL_FLOAT16, ACL_FORMAT_ND).Value(vector<double>{1.0, 2.0, 3.0, 4.0});
59+ auto repeatsDesc = IntArrayDesc(vector<int64_t>{3});
60+ auto outDesc = TensorDesc({12}, ACL_FLOAT16, ACL_FORMAT_ND);
61+ 
62+ auto ut = OP_API_UT(aclnnRepeat, INPUT(selfDesc, repeatsDesc), OUTPUT(outDesc));
63+ 
64+ uint64_t workspace_size = 0;
65+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
66+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
67+}
Aexperimental/math/tile/tests/ut/op_host/test_tile_infershape.cpp+112-0
@@ -0,0 +1,112 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+#include <iostream>
13+#include "infershape_context_faker.h"
14+#include "infershape_case_executor.h"
15+ 
16+class TileInfershape : public testing::Test {
17+protected:
18+ static void SetUpTestCase()
19+ {
20+ std::cout << "TileInfershape SetUp" << std::endl;
21+ }
22+ 
23+ static void TearDownTestCase()
24+ {
25+ std::cout << "TileInfershape TearDown" << std::endl;
26+ }
27+};
28+ 
29+TEST_F(TileInfershape, tile_infershape_2d_dynamic)
30+{
31+ gert::InfershapeContextPara infershapeContextPara(
32+ "Tile",
33+ {
34+ {{{2, 3}, {2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
35+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND},
36+ },
37+ {
38+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
39+ });
40+ std::vector<std::vector<int64_t>> expectOutputShape = {
41+ {-1, -1},
42+ };
43+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
44+}
45+ 
46+TEST_F(TileInfershape, tile_infershape_1d_dynamic)
47+{
48+ gert::InfershapeContextPara infershapeContextPara(
49+ "Tile",
50+ {
51+ {{{128}, {128}}, ge::DT_INT32, ge::FORMAT_ND},
52+ {{{1}, {1}}, ge::DT_INT32, ge::FORMAT_ND},
53+ },
54+ {
55+ {{{}, {}}, ge::DT_INT32, ge::FORMAT_ND},
56+ });
57+ std::vector<std::vector<int64_t>> expectOutputShape = {
58+ {-1},
59+ };
60+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
61+}
62+ 
63+TEST_F(TileInfershape, tile_infershape_4d_dynamic)
64+{
65+ gert::InfershapeContextPara infershapeContextPara(
66+ "Tile",
67+ {
68+ {{{1, 1, 1, 8}, {1, 1, 1, 8}}, ge::DT_FLOAT16, ge::FORMAT_ND},
69+ {{{4}, {4}}, ge::DT_INT32, ge::FORMAT_ND},
70+ },
71+ {
72+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
73+ });
74+ std::vector<std::vector<int64_t>> expectOutputShape = {
75+ {-1, -1, -1, -1},
76+ };
77+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
78+}
79+ 
80+TEST_F(TileInfershape, tile_infershape_broadcast_dims)
81+{
82+ gert::InfershapeContextPara infershapeContextPara(
83+ "Tile",
84+ {
85+ {{{3}, {3}}, ge::DT_FLOAT, ge::FORMAT_ND},
86+ {{{3}, {3}}, ge::DT_INT64, ge::FORMAT_ND},
87+ },
88+ {
89+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
90+ });
91+ std::vector<std::vector<int64_t>> expectOutputShape = {
92+ {-1, -1, -1},
93+ };
94+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
95+}
96+ 
97+TEST_F(TileInfershape, tile_infershape_empty_multiples)
98+{
99+ gert::InfershapeContextPara infershapeContextPara(
100+ "Tile",
101+ {
102+ {{{2, 3}, {2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
103+ {{{0}, {0}}, ge::DT_INT32, ge::FORMAT_ND},
104+ },
105+ {
106+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
107+ });
108+ std::vector<std::vector<int64_t>> expectOutputShape = {
109+ {2, 3},
110+ };
111+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
112+}
Aexperimental/math/tile/tests/ut/op_host/test_tile_tiling.cpp+159-0
@@ -0,0 +1,159 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+#include <iostream>
13+#include <cstring>
14+#include "tile_tiling.h"
15+#include "../../../op_kernel/tile_tiling_data.h"
16+#include "../../../op_kernel/tile_tiling_key.h"
17+#include "tiling_context_faker.h"
18+#include "tiling_case_executor.h"
19+ 
20+class TileTiling : public testing::Test {
21+protected:
22+ static void SetUpTestCase()
23+ {
24+ std::cout << "TileTiling SetUp" << std::endl;
25+ }
26+ 
27+ static void TearDownTestCase()
28+ {
29+ std::cout << "TileTiling TearDown" << std::endl;
30+ }
31+};
32+ 
33+TEST_F(TileTiling, tile_tiling_2d_float32)
34+{
35+ optiling::TileCompileInfo compileInfo;
36+ int32_t multiplesData[] = {3, 2};
37+ gert::TilingContextPara tilingContextPara(
38+ "Tile",
39+ {
40+ {{{2, 3}, {2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
41+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND, true, multiplesData},
42+ },
43+ {
44+ {{{6, 6}, {6, 6}}, ge::DT_FLOAT, ge::FORMAT_ND},
45+ },
46+ &compileInfo);
47+ 
48+ TilingInfo tilingInfo;
49+ bool ret = ExecuteTiling(tilingContextPara, tilingInfo);
50+ ASSERT_TRUE(ret);
51+ 
52+ ASSERT_GE(tilingInfo.tilingDataSize, sizeof(TileTilingData));
53+ const TileTilingData* tiling = reinterpret_cast<const TileTilingData*>(tilingInfo.tilingData.get());
54+ EXPECT_EQ(tiling->numDims, 2);
55+ EXPECT_EQ(tiling->inputShape[0], 2);
56+ EXPECT_EQ(tiling->inputShape[1], 3);
57+ EXPECT_EQ(tiling->multiples[0], 3);
58+ EXPECT_EQ(tiling->multiples[1], 2);
59+ EXPECT_EQ(tiling->outputShape[0], 6);
60+ EXPECT_EQ(tiling->outputShape[1], 6);
61+ EXPECT_EQ(tiling->totalInputElems, 6);
62+ EXPECT_EQ(tiling->totalOutputElems, 36);
63+ EXPECT_EQ(tiling->elemBytes, 4);
64+}
65+ 
66+TEST_F(TileTiling, tile_tiling_1d_int32)
67+{
68+ optiling::TileCompileInfo compileInfo;
69+ int32_t multiplesData[] = {8};
70+ gert::TilingContextPara tilingContextPara(
71+ "Tile",
72+ {
73+ {{{128}, {128}}, ge::DT_INT32, ge::FORMAT_ND},
74+ {{{1}, {1}}, ge::DT_INT32, ge::FORMAT_ND, true, multiplesData},
75+ },
76+ {
77+ {{{1024}, {1024}}, ge::DT_INT32, ge::FORMAT_ND},
78+ },
79+ &compileInfo);
80+ 
81+ TilingInfo tilingInfo;
82+ bool ret = ExecuteTiling(tilingContextPara, tilingInfo);
83+ ASSERT_TRUE(ret);
84+ 
85+ ASSERT_GE(tilingInfo.tilingDataSize, sizeof(TileTilingData));
86+ const TileTilingData* tiling = reinterpret_cast<const TileTilingData*>(tilingInfo.tilingData.get());
87+ EXPECT_EQ(tiling->numDims, 1);
88+ EXPECT_EQ(tiling->inputShape[0], 128);
89+ EXPECT_EQ(tiling->multiples[0], 8);
90+ EXPECT_EQ(tiling->outputShape[0], 1024);
91+ EXPECT_EQ(tiling->totalInputElems, 128);
92+ EXPECT_EQ(tiling->totalOutputElems, 1024);
93+ EXPECT_EQ(tiling->elemBytes, 4);
94+}
95+ 
96+TEST_F(TileTiling, tile_tiling_int64_multiples)
97+{
98+ optiling::TileCompileInfo compileInfo;
99+ int64_t multiplesData[] = {2, 4};
100+ gert::TilingContextPara tilingContextPara(
101+ "Tile",
102+ {
103+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
104+ {{{2}, {2}}, ge::DT_INT64, ge::FORMAT_ND, true, multiplesData},
105+ },
106+ {
107+ {{{6, 20}, {6, 20}}, ge::DT_FLOAT16, ge::FORMAT_ND},
108+ },
109+ &compileInfo);
110+ 
111+ TilingInfo tilingInfo;
112+ bool ret = ExecuteTiling(tilingContextPara, tilingInfo);
113+ ASSERT_TRUE(ret);
114+ 
115+ ASSERT_GE(tilingInfo.tilingDataSize, sizeof(TileTilingData));
116+ const TileTilingData* tiling = reinterpret_cast<const TileTilingData*>(tilingInfo.tilingData.get());
117+ EXPECT_EQ(tiling->numDims, 2);
118+ EXPECT_EQ(tiling->inputShape[0], 3);
119+ EXPECT_EQ(tiling->inputShape[1], 5);
120+ EXPECT_EQ(tiling->multiples[0], 2);
121+ EXPECT_EQ(tiling->multiples[1], 4);
122+ EXPECT_EQ(tiling->outputShape[0], 6);
123+ EXPECT_EQ(tiling->outputShape[1], 20);
124+ EXPECT_EQ(tiling->totalInputElems, 15);
125+ EXPECT_EQ(tiling->totalOutputElems, 120);
126+ EXPECT_EQ(tiling->elemBytes, 2);
127+}
128+ 
129+TEST_F(TileTiling, tile_tiling_dim_merge)
130+{
131+ optiling::TileCompileInfo compileInfo;
132+ int32_t multiplesData[] = {3, 1, 2};
133+ gert::TilingContextPara tilingContextPara(
134+ "Tile",
135+ {
136+ {{{2, 4, 5}, {2, 4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
137+ {{{3}, {3}}, ge::DT_INT32, ge::FORMAT_ND, true, multiplesData},
138+ },
139+ {
140+ {{{6, 4, 10}, {6, 4, 10}}, ge::DT_FLOAT, ge::FORMAT_ND},
141+ },
142+ &compileInfo);
143+ 
144+ TilingInfo tilingInfo;
145+ bool ret = ExecuteTiling(tilingContextPara, tilingInfo);
146+ ASSERT_TRUE(ret);
147+ 
148+ ASSERT_GE(tilingInfo.tilingDataSize, sizeof(TileTilingData));
149+ const TileTilingData* tiling = reinterpret_cast<const TileTilingData*>(tilingInfo.tilingData.get());
150+ EXPECT_EQ(tiling->numDims, 2);
151+ EXPECT_EQ(tiling->inputShape[0], 8);
152+ EXPECT_EQ(tiling->inputShape[1], 5);
153+ EXPECT_EQ(tiling->multiples[0], 3);
154+ EXPECT_EQ(tiling->multiples[1], 2);
155+ EXPECT_EQ(tiling->outputShape[0], 24);
156+ EXPECT_EQ(tiling->outputShape[1], 10);
157+ EXPECT_EQ(tiling->totalInputElems, 40);
158+ EXPECT_EQ(tiling->totalOutputElems, 240);
159+}
Aexperimental/math/tile/tests/ut/op_host/tile_tiling.h+18-0
@@ -0,0 +1,18 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef TILE_TILING_H_
12+#define TILE_TILING_H_
13+ 
14+namespace optiling {
15+struct TileCompileInfo {};
16+} // namespace optiling
17+ 
18+#endif
Aexperimental/math/tile/tests/ut/op_kernel/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+ 
9+if (UT_TEST_ALL OR OP_KERNEL_UT)
10+ set(tile_tiling_files
11+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/tile_tiling.cpp
12+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/tile_infershape.cpp
13+ )
14+ AddOpTestCase(tile "ascend910b" "-DDTYPE_X=float" "${tile_tiling_files}")
15+endif()
Aexperimental/math/tile/tests/ut/op_kernel/test_tile.cpp+189-0
@@ -0,0 +1,189 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <array>
12+#include <vector>
13+#include <iostream>
14+#include <string>
15+#include <cstdint>
16+#include <cstring>
17+#include "gtest/gtest.h"
18+#include "tikicpulib.h"
19+#include "data_utils.h"
20+ 
21+#include "../../../op_kernel/tile.cpp"
22+ 
23+using namespace std;
24+ 
25+class TileTest : public testing::Test {
26+protected:
27+ static void SetUpTestCase()
28+ {
29+ std::cout << "TileTest SetUp" << std::endl;
30+ const string cmd = "cp -rf " + dataPath + " ./";
31+ system(cmd.c_str());
32+ system("chmod -R 755 ./tile_data/");
33+ }
34+ static void TearDownTestCase()
35+ {
36+ std::cout << "TileTest TearDown" << std::endl;
37+ }
38+ 
39+private:
40+ const static std::string rootPath;
41+ const static std::string dataPath;
42+};
43+ 
44+const std::string TileTest::rootPath = "../../../../experimental/";
45+const std::string TileTest::dataPath = rootPath + "math/tile/tests/ut/op_kernel/tile_data";
46+ 
47+template <typename T1, typename T2>
48+inline T1 CeilAlign(T1 a, T2 b)
49+{
50+ if (b == 0) {
51+ return a;
52+ }
53+ return (a + b - 1) / b * b;
54+}
55+ 
56+static void BuildTilingData(
57+ TileTilingData* tilingData, const std::vector<int32_t>& shape, const std::vector<int32_t>& multiples,
58+ int32_t elemBytes, int32_t blockDim, int32_t ubSize)
59+{
60+ memset(tilingData, 0, sizeof(TileTilingData));
61+ int32_t ndim = static_cast<int32_t>(shape.size());
62+ tilingData->numDims = ndim;
63+ tilingData->elemBytes = elemBytes;
64+ tilingData->blockDim = blockDim;
65+ tilingData->ubSize = ubSize;
66+ tilingData->totalInputElems = 1;
67+ tilingData->totalOutputElems = 1;
68+ for (int32_t idx = 0; idx < ndim; idx++) {
69+ tilingData->inputShape[idx] = shape[idx];
70+ tilingData->multiples[idx] = multiples[idx];
71+ tilingData->outputShape[idx] = shape[idx] * multiples[idx];
72+ tilingData->totalInputElems *= shape[idx];
73+ tilingData->totalOutputElems *= tilingData->outputShape[idx];
74+ }
75+ tilingData->inputStrides[ndim - 1] = 1;
76+ tilingData->outputStrides[ndim - 1] = 1;
77+ for (int32_t idx = ndim - 2; idx >= 0; idx--) {
78+ tilingData->inputStrides[idx] = tilingData->inputStrides[idx + 1] * tilingData->inputShape[idx + 1];
79+ tilingData->outputStrides[idx] = tilingData->outputStrides[idx + 1] * tilingData->outputShape[idx + 1];
80+ }
81+}
82+ 
83+TEST_F(TileTest, test_case_float32_2d)
84+{
85+ std::vector<int32_t> shape = {2, 3};
86+ std::vector<int32_t> mult = {3, 2};
87+ 
88+ uint8_t* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(sizeof(TileTilingData), 32)));
89+ TileTilingData* tilingData = reinterpret_cast<TileTilingData*>(tiling);
90+ BuildTilingData(tilingData, shape, mult, 4, 1, 87381);
91+ 
92+ system("cd ./tile_data/ && python3 gen_data.py '2,3' '3,2' float32");
93+ uint32_t inputCount = static_cast<uint32_t>(tilingData->totalInputElems);
94+ uint32_t outputCount = static_cast<uint32_t>(tilingData->totalOutputElems);
95+ size_t inputByteSize = inputCount * sizeof(float);
96+ size_t outputByteSize = outputCount * sizeof(float);
97+ 
98+ uint8_t* x = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(inputByteSize + 32, 32)));
99+ ReadFile("./tile_data/float32_input_t_tile.bin", inputByteSize, x, inputByteSize);
100+ uint8_t* y = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(outputByteSize, 32)));
101+ uint8_t* multGm = static_cast<uint8_t*>(AscendC::GmAlloc(32));
102+ uint8_t* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(32 * 1024 * 1024));
103+ 
104+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
105+ auto func = tile<TILE_TPL_SCH_MODE_DEFAULT>;
106+ ICPU_RUN_KF(func, tilingData->blockDim, x, multGm, y, workspace, tiling);
107+ 
108+ WriteFile("./tile_data/float32_output_t_tile.bin", y, outputByteSize);
109+ 
110+ AscendC::GmFree(static_cast<void*>(x));
111+ AscendC::GmFree(static_cast<void*>(y));
112+ AscendC::GmFree(static_cast<void*>(multGm));
113+ AscendC::GmFree(static_cast<void*>(workspace));
114+ AscendC::GmFree(static_cast<void*>(tiling));
115+ 
116+ system("cd ./tile_data/ && python3 compare_data.py float32_golden_t_tile.bin float32_output_t_tile.bin float32");
117+}
118+ 
119+TEST_F(TileTest, test_case_int32_1d)
120+{
121+ std::vector<int32_t> shape = {128};
122+ std::vector<int32_t> mult = {8};
123+ 
124+ uint8_t* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(sizeof(TileTilingData), 32)));
125+ TileTilingData* tilingData = reinterpret_cast<TileTilingData*>(tiling);
126+ BuildTilingData(tilingData, shape, mult, 4, 1, 87381);
127+ 
128+ system("cd ./tile_data/ && python3 gen_data.py '128' '8' int32");
129+ uint32_t inputCount = static_cast<uint32_t>(tilingData->totalInputElems);
130+ uint32_t outputCount = static_cast<uint32_t>(tilingData->totalOutputElems);
131+ size_t inputByteSize = inputCount * sizeof(int32_t);
132+ size_t outputByteSize = outputCount * sizeof(int32_t);
133+ 
134+ uint8_t* x = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(inputByteSize + 32, 32)));
135+ ReadFile("./tile_data/int32_input_t_tile.bin", inputByteSize, x, inputByteSize);
136+ uint8_t* y = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(outputByteSize, 32)));
137+ uint8_t* multGm = static_cast<uint8_t*>(AscendC::GmAlloc(32));
138+ uint8_t* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(32 * 1024 * 1024));
139+ 
140+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
141+ auto func = tile<TILE_TPL_SCH_MODE_DEFAULT>;
142+ ICPU_RUN_KF(func, tilingData->blockDim, x, multGm, y, workspace, tiling);
143+ 
144+ WriteFile("./tile_data/int32_output_t_tile.bin", y, outputByteSize);
145+ 
146+ AscendC::GmFree(static_cast<void*>(x));
147+ AscendC::GmFree(static_cast<void*>(y));
148+ AscendC::GmFree(static_cast<void*>(multGm));
149+ AscendC::GmFree(static_cast<void*>(workspace));
150+ AscendC::GmFree(static_cast<void*>(tiling));
151+ 
152+ system("cd ./tile_data/ && python3 compare_data.py int32_golden_t_tile.bin int32_output_t_tile.bin int32");
153+}
154+ 
155+TEST_F(TileTest, test_case_float32_large_inner)
156+{
157+ std::vector<int32_t> shape = {1, 1, 25000};
158+ std::vector<int32_t> mult = {3, 3, 12};
159+ 
160+ uint8_t* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(sizeof(TileTilingData), 32)));
161+ TileTilingData* tilingData = reinterpret_cast<TileTilingData*>(tiling);
162+ BuildTilingData(tilingData, shape, mult, 4, 1, 87381);
163+ 
164+ system("cd ./tile_data/ && python3 gen_data.py '1,1,25000' '3,3,12' float32");
165+ uint32_t inputCount = static_cast<uint32_t>(tilingData->totalInputElems);
166+ uint32_t outputCount = static_cast<uint32_t>(tilingData->totalOutputElems);
167+ size_t inputByteSize = inputCount * sizeof(float);
168+ size_t outputByteSize = outputCount * sizeof(float);
169+ 
170+ uint8_t* x = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(inputByteSize + 32, 32)));
171+ ReadFile("./tile_data/float32_input_t_tile.bin", inputByteSize, x, inputByteSize);
172+ uint8_t* y = static_cast<uint8_t*>(AscendC::GmAlloc(CeilAlign(outputByteSize, 32)));
173+ uint8_t* multGm = static_cast<uint8_t*>(AscendC::GmAlloc(32));
174+ uint8_t* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(32 * 1024 * 1024));
175+ 
176+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
177+ auto func = tile<TILE_TPL_SCH_MODE_DEFAULT>;
178+ ICPU_RUN_KF(func, tilingData->blockDim, x, multGm, y, workspace, tiling);
179+ 
180+ WriteFile("./tile_data/float32_output_t_tile.bin", y, outputByteSize);
181+ 
182+ AscendC::GmFree(static_cast<void*>(x));
183+ AscendC::GmFree(static_cast<void*>(y));
184+ AscendC::GmFree(static_cast<void*>(multGm));
185+ AscendC::GmFree(static_cast<void*>(workspace));
186+ AscendC::GmFree(static_cast<void*>(tiling));
187+ 
188+ system("cd ./tile_data/ && python3 compare_data.py float32_golden_t_tile.bin float32_output_t_tile.bin float32");
189+}
Aexperimental/math/tile/tests/ut/op_kernel/tile_data/compare_data.py+46-0
@@ -0,0 +1,46 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
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+import sys
12+import numpy as np
13+ 
14+ 
15+def compare_data(golden_file, output_file, d_type="float32"):
16+ d_type_dict = {
17+ "float32": np.float32,
18+ "float16": np.float16,
19+ "int32": np.int32,
20+ "int8": np.int8,
21+ "uint8": np.uint8,
22+ }
23+ np_type = d_type_dict.get(d_type, np.float32)
24+ golden = np.fromfile(golden_file, dtype=np_type)
25+ output = np.fromfile(output_file, dtype=np_type)
26+ 
27+ if golden.shape != output.shape:
28+ print(f"FAIL: shape mismatch golden={golden.shape} output={output.shape}")
29+ return False
30+ 
31+ if np.issubdtype(np_type, np.integer):
32+ match = np.array_equal(golden, output)
33+ else:
34+ match = np.allclose(golden, output, atol=1e-4, rtol=1e-4)
35+ 
36+ max_diff = np.max(np.abs(golden.astype(np.float64) - output.astype(np.float64)))
37+ print(f"{'PASS' if match else 'FAIL'}: max_diff={max_diff}")
38+ return match
39+ 
40+ 
41+if __name__ == "__main__":
42+ if len(sys.argv) != 4:
43+ print("Usage: compare_data.py <golden_file> <output_file> <dtype>")
44+ exit(1)
45+ ok = compare_data(sys.argv[1], sys.argv[2], sys.argv[3])
46+ exit(0 if ok else 1)
Aexperimental/math/tile/tests/ut/op_kernel/tile_data/gen_data.py+52-0
@@ -0,0 +1,52 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
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+import sys
12+import os
13+import numpy as np
14+ 
15+ 
16+def parse_str_to_shape_list(shape_str):
17+ shape_str = shape_str.strip('(').strip(')')
18+ shape_list = [int(x) for x in shape_str.split(",")]
19+ return np.array(shape_list)
20+ 
21+ 
22+def gen_data_and_golden(shape_str, multiples_str, d_type="float32"):
23+ d_type_dict = {
24+ "float32": np.float32,
25+ "float16": np.float16,
26+ "int32": np.int32,
27+ "int8": np.int8,
28+ "uint8": np.uint8,
29+ }
30+ np_type = d_type_dict.get(d_type, np.float32)
31+ shape = parse_str_to_shape_list(shape_str)
32+ multiples = parse_str_to_shape_list(multiples_str)
33+ 
34+ np.random.seed(42)
35+ if np.issubdtype(np_type, np.integer):
36+ tmp_input = np.random.randint(0, 100, size=shape).astype(np_type)
37+ else:
38+ tmp_input = np.random.randn(*shape).astype(np_type)
39+ 
40+ tmp_golden = np.tile(tmp_input, multiples)
41+ 
42+ tmp_input.tofile(f"{d_type}_input_t_tile.bin")
43+ tmp_golden.tofile(f"{d_type}_golden_t_tile.bin")
44+ 
45+ 
46+if __name__ == "__main__":
47+ if len(sys.argv) != 4:
48+ print("Usage: gen_data.py <shape> <multiples> <dtype>")
49+ print("Example: gen_data.py '2,3' '3,2' float32")
50+ exit(1)
51+ os.system("rm -rf *.bin")
52+ gen_data_and_golden(sys.argv[1], sys.argv[2], sys.argv[3])
Aexperimental/math/tile/tests/ut/op_kernel/tile_tiling.h+56-0
@@ -0,0 +1,56 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef _TILE_TILING_H_
17+#define _TILE_TILING_H_
18+ 
19+#include <cstdint>
20+#include "../../../op_kernel/tile_tiling_data.h"
21+#include "kernel_tiling/kernel_tiling.h"
22+#include "../../../op_kernel/tile.h"
23+ 
24+using namespace TileKernel;
25+ 
26+#ifdef __NPU_TILING__
27+inline [aicore] void InitTilingData(const __gm__ uint8_t* tiling, TileTilingData* constData) {
28+ const __gm__ uint32_t* src = (const __gm__ uint32_t*)tiling;
29+ uint32_t* dst = (uint32_t*)constData;
30+ for (size_t i = 0; i < sizeof(TileTilingData) / 4; i++) {
31+ *(dst + i) = *(src + i);
32+ }
33+}
34+#else
35+inline void InitTilingData(uint8_t* tiling, TileTilingData* constData)
36+{
37+ memcpy(constData, tiling, sizeof(TileTilingData));
38+}
39+#endif
40+ 
41+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
42+ __ubuf__ tilingStruct* tilingDataPointer = \
43+ reinterpret_cast<__ubuf__ tilingStruct*>((__ubuf__ uint8_t*)(tilingPointer))
44+ 
45+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
46+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer)
47+ 
48+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingArg) \
49+ tilingStruct tilingData; \
50+ InitTilingData(tilingArg, &tilingData)
51+ 
52+#define GET_TILING_DATA(tilingData, tilingArg) \
53+ TileTilingData tilingData; \
54+ InitTilingData(tilingArg, &tilingData)
55+ 
56+#endif