Star896
1805
代码
代码
Issues250
Pull Requests311
Actions
项目讨论
Wiki
分析
Star896
1805
已合并
MapIndex支持下一代实现 #1191
Guoqh创建于 1月29日
MapIndex支持下一代实现 #1191
已合并
Guoqh创建于 1月29日
从guoqiuhao/ops-nnmigrate合入到cann/ops-nnmaster
共 15 个文件变更+1664-1
Aindex/map_index/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
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+set(SUPPORT_COMPUTE_UNIT "ascend910_95")
13+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14+set(SUPPORT_TILING_DIR "arch35")
15+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE map_index ACLNNTYPE aclnn_exclude
16+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Aindex/map_index/README.md+3-0
@@ -0,0 +1,3 @@
1+# MapIndex
2+ 
3+本目录仅包含MapIndex算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。
Aindex/map_index/op_graph/map_index_proto.h+55-0
@@ -0,0 +1,55 @@
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 map_index_proto.h
13+ * \brief
14+ */
15+ 
16+#ifndef OPS_OP_PROTO_INC_MAP_INDEX_H_
17+#define OPS_OP_PROTO_INC_MAP_INDEX_H_
18+ 
19+#include "graph/operator_reg.h"
20+#include "graph/types.h"
21+ 
22+namespace ge {
23+ 
24+/**
25+* @brief Returns index of shape in the map.
26+ 
27+* @par Inputs:
CANN-robot
CANN-robotCANN-robot1月29日

注释与命名: 函数/操作符的注释格式存在不一致：1）@brief部分与详细描述之间缺少空行；2）@par Inputs部分使用了项目符号列表（@li），但属性（@par Attributes）和输出（@par Outputs）部分未使用统一格式；3）注释中存在换行不一致（如第31行结尾有换行符\n，但其他地方没有）。

问题类型: 注释与命名 文件路径: index/map_index/op_graph/map_index_proto.h 行号: 27 问题代码:

* @par Inputs:
* Three inputs, including:
* @li x: One dimensional tensor of type int32, specifying queried shape,
* Format support ND, the max dim is 400 (128 on Lhisi(Hi3796CV300CS, SD3403), 24000 on Ascend 910_95 AI Processor).
* @li data_seq: One dimensional tensor of type int32, 
* Format support ND, specifying the mapped table is queried.

修改建议:

统一注释格式：1）在@brief后加空行；2）对Inputs、Attributes、Outputs都使用一致的列表格式（如都使用@li或都使用段落描述）；3）移除不必要的换行符（如第31行的\n），确保注释清晰连贯。遵循项目已有的文档规范。

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

likedislike
28+* Three inputs, including:
29+* @li x: One dimensional tensor of type int32, specifying queried shape,
30+* Format support ND, the max dim is 400 (128 on Lhisi(Hi3796CV300CS, SD3403), 24000 on Ascend 910_95 AI Processor).
CANN-robot
CANN-robotCANN-robot1月29日

代码结构与可维护性: 注释中关于x的最大维度描述存在多个平台特定值（400、128、24000），这些魔法数字直接硬编码在注释中，缺乏明确的常量定义。这会导致代码可维护性差，当平台限制变更时需要多处修改注释，且容易引起开发者混淆。

问题类型: 代码结构与可维护性 文件路径: index/map_index/op_graph/map_index_proto.h 行号: 30 问题代码:

Format support ND, the max dim is 400 (128 on Lhisi(Hi3796CV300CS, SD3403), 24000 on Ascend 910_95 AI Processor).

修改建议:

建议将平台相关的最大维度限制定义为命名常量（如MAX_DIM_DEFAULT、MAX_DIM_LHISI、MAX_DIM_ASCEND_910_95），并在注释中引用这些常量。如果这些限制在运行时由框架或硬件决定，应在注释中说明限制的来源（如框架配置或硬件规格）。

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

likedislike
31+* @li data_seq: One dimensional tensor of type int32,
32+* Format support ND, specifying the mapped table is queried.
33+* The length of data_seq must be multiple of the length of x, and the length of data_seq / x <= 100 (256 on Ascend 910_95 AI Processor).
CANN-robot
CANN-robotCANN-robot1月29日

代码结构与可维护性: 注释中关于data_seq长度与x长度比值的限制（<=100或256）同样是硬编码的魔法数字，且存在平台差异。这种直接嵌入业务逻辑数值的方式不符合代码规范，降低了可读性和可维护性。

问题类型: 代码结构与可维护性 文件路径: index/map_index/op_graph/map_index_proto.h 行号: 33 问题代码:

The length of data_seq must be multiple of the length of x, and the length of data_seq / x <= 100 (256 on Ascend 910_95 AI Processor).

修改建议:

将比值限制定义为命名常量（如MAX_DATA_SEQ_RATIO_DEFAULT、MAX_DATA_SEQ_RATIO_ASCEND），并在注释中引用。同时考虑是否需要在操作符注册时通过属性或约束来显式声明这些限制，以便框架进行静态或运行时验证。

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

likedislike
34+* @li level_index: One dimensional tensor of type int32, the length of level_index must be equal to the length of data_seq divided by the length of x.
CANN-robot
CANN-robotCANN-robot1月29日

错误处理与异常安全: 注释中描述了level_index的长度约束（必须等于data_seq长度除以x长度），但这是一个运行时约束。操作符注册代码（REG_OP）本身没有机制来强制执行这种复杂的动态约束，完全依赖后续算子实现或调用方保证。缺乏显式的约束声明可能导致运行时错误难以调试。

问题类型: 错误处理与异常安全 文件路径: index/map_index/op_graph/map_index_proto.h 行号: 34 问题代码:

the length of level_index must be equal to the length of data_seq divided by the length of x.

修改建议:

虽然操作符原型定义可能无法直接表达动态约束，但应在注释中明确强调该约束必须由算子实现（kernel）在运行时验证，并说明违反约束时的预期行为（如返回错误码）。如果框架支持，考虑添加自定义验证器。

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

likedislike
35+* Format support ND, specifying secondary index. \n
36+ 
37+* @par Attributes:
38+ *@li transpose: An optional bool. specifying the input is transposed on A3 or A5, A3 is true, A5 is false.
CANN-robot
CANN-robotCANN-robot1月29日

代码结构与可维护性: transpose属性的注释描述存在歧义和不准确：1）'A3'和'A5'的含义不明确（可能是硬件架构代号但未说明）；2）描述为'input is transposed on A3 or A5'，但属性是bool类型，无法同时表示两种状态；3）true对应A3、false对应A5的映射关系缺乏上下文解释。

问题类型: 代码结构与可维护性 文件路径: index/map_index/op_graph/map_index_proto.h 行号: 38 问题代码:

@li transpose: An optional bool. specifying the input is transposed on A3 or A5, A3 is true, A5 is false.

修改建议:

重写注释以明确transpose属性的实际含义。例如：'transpose: An optional bool. When true, the input tensor is transposed according to architecture A3 format; when false, it follows architecture A5 format (or no transposition).' 如果A3/A5是内部代号，应考虑使用更通用的描述或添加参考文档链接。

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

likedislike
39+ 
40+* @par Outputs:
41+* y: A scalar of type int32, specifying index of shape in the map.
42+* @par Third-party framework compatibility
43+* It is a custom operator. It has no corresponding operator in Caffe.
44+*/
45+REG_OP(MapIndex)
46+ .INPUT(x, TensorType({DT_INT32}))
47+ .INPUT(data_seq, TensorType({DT_INT32}))
48+ .OPTIONAL_INPUT(level_index, TensorType({DT_INT32}))
49+ .OUTPUT(y, TensorType({DT_INT32}))
50+ .ATTR(transpose, Bool, false)
51+ .OP_END_FACTORY_REG(MapIndex)
52+ 
53+} // namespace ge
54+ 
55+#endif // OPS_OP_PROTO_INC_MAP_INDEX_H_
Aindex/map_index/op_host/config/ascend950/map_index_binary.json+63-0
@@ -0,0 +1,63 @@
1+{
2+ "op_type": "MapIndex",
3+ "op_list": [
4+ {
5+ "bin_filename": "MapIndex_int32_int32_int32",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "int32",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ },
18+ {
19+ "name": "data_seq",
20+ "index": 1,
21+ "dtype": "int32",
22+ "format": "ND",
23+ "paramType": "required",
24+ "shape": [
25+ -2
26+ ],
27+ "format_match_mode": "FormatAgnostic"
28+ },
29+ {
30+ "name": "level_index",
31+ "index": 2,
32+ "dtype": "int32",
33+ "format": "ND",
34+ "paramType": "optional",
35+ "shape": [
36+ -2
37+ ],
38+ "format_match_mode": "FormatAgnostic"
39+ }
40+ ],
41+ "outputs": [
42+ {
43+ "name": "y",
44+ "index": 0,
45+ "dtype": "int32",
46+ "format": "ND",
47+ "paramType": "required",
48+ "shape": [
49+ -2
50+ ],
51+ "format_match_mode": "FormatAgnostic"
52+ }
53+ ],
54+ "attrs": [
55+ {
56+ "name": "transpose",
57+ "dtype": "bool",
58+ "value": false
59+ }
60+ ]
61+ }
62+ ]
63+}
Aindex/map_index/op_host/config/ascend950/map_index_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascnedxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[MapIndex]
13+default=0
Aindex/map_index/op_host/map_index_def.cpp+62-0
@@ -0,0 +1,62 @@
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 map_index_def.cpp
13+ * \brief
14+ */
15+ 
16+#include "register/op_def_registry.h"
17+
18+namespace ops {
19+ class MapIndex : public OpDef {
20+ public:
21+ explicit MapIndex(const char* name) : OpDef(name)
22+ {
23+ this->Input("x")
24+ .ParamType(REQUIRED)
25+ .DataType({ge::DT_INT32})
26+ .Format({ge::FORMAT_ND})
27+ .UnknownShapeFormat({ge::FORMAT_ND})
28+ .AutoContiguous();
29+ this->Input("data_seq")
30+ .ParamType(REQUIRED)
31+ .DataType({ge::DT_INT32})
32+ .Format({ge::FORMAT_ND})
33+ .UnknownShapeFormat({ge::FORMAT_ND})
34+ .AutoContiguous();
35+ this->Input("level_index")
36+ .ParamType(OPTIONAL)
37+ .DataType({ge::DT_INT32})
38+ .Format({ge::FORMAT_ND})
39+ .UnknownShapeFormat({ge::FORMAT_ND})
40+ .AutoContiguous();
41+ this->Output("y")
42+ .ParamType(REQUIRED)
43+ .DataType({ge::DT_INT32})
44+ .Format({ge::FORMAT_ND})
45+ .UnknownShapeFormat({ge::FORMAT_ND});
46+ 
47+ this->Attr("transpose").AttrType(OPTIONAL).Bool(false);
48+
49+ OpAICoreConfig aicoreConfig;
50+ aicoreConfig.DynamicCompileStaticFlag(true)
51+ .DynamicFormatFlag(false)
52+ .DynamicRankSupportFlag(true)
53+ .DynamicShapeSupportFlag(true)
54+ .NeedCheckSupportFlag(false)
55+ .PrecisionReduceFlag(true)
56+ .ExtendCfgInfo("opFile.value", "map_index_apt");
57+ this->AICore().AddConfig("ascend950", aicoreConfig);
58+ }
59+ };
60+
61+ OP_ADD(MapIndex);
62+ } // namespace ops
Aindex/map_index/op_host/map_index_tiling_arch35.cpp+304-0
@@ -0,0 +1,304 @@
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 map_index_tiling_arch35.cpp
13+* \brief
14+*/
15+ 
16+#include "map_index_tiling_arch35.h"
17+#include "util/platform_util.h"
18+#include "util/math_util.h"
19+#include "register/op_impl_registry.h"
20+#include "log/log.h"
21+#include "tiling_base/tiling_util.h"
22+ 
23+using namespace std;
24+using namespace ge;
25+ 
26+namespace optiling {
27+ 
28+constexpr int64_t DIGIT_EIGHT = 8;
29+constexpr int64_t X_MAX_DIM0 = 24000;
30+constexpr int64_t DATA_SEQ_MAX_DIM0 = 256;
31+constexpr int64_t VL_NUMS = 64;
32+constexpr size_t WORKSPACE_SIZE = 32;
CANN-robot
CANN-robotCANN-robot1月29日

代码结构与可维护性: 定义了WORKSPACE_SIZE常量但未使用。可能是残留代码或计划未来使用。

问题类型: 代码结构与可维护性 文件路径: index/map_index/op_host/map_index_tiling_arch35.cpp 行号: 32 问题代码:

constexpr size_t WORKSPACE_SIZE = 32;

修改建议:

如果确实未使用，应删除以避免混淆。如果计划使用，请添加注释说明其用途。

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

likedislike
33+constexpr size_t LEVEL_INDEX_INDEX = 2;
34+constexpr int64_t NUM_TWO_DB = 2;
35+constexpr int64_t RESERVED_UB_SIZE = static_cast<int64_t>(8) * 1024; // 8k
36+constexpr int64_t WORKSPACE_BUFFER = static_cast<int64_t>(20) * 1024 * 1024;
37+constexpr int64_t ATTR_INDEX_TRANSPOSE = 0;
38+const std::set<ge::DataType> INPUT_SUPPORT_DTYPE_SET = { ge::DT_INT32 };
39+ 
40+template <class T> T inline CeilDivide(T num1, T num2)
41+{
42+ if (num2 == 0) {
43+ return 0;
44+ }
45+ return (num1 + num2 - 1) / num2;
46+}
47+ 
48+static ge::graphStatus CheckDtype(const gert::TilingContext *context, MapIndexTilingParam &tilingParam)
49+{
50+ OP_LOGD(context->GetNodeName(), "CheckDtype begin.");
51+ auto inputXPtr = context->GetInputDesc(0);
52+ OP_CHECK_NULL_WITH_CONTEXT(context, inputXPtr);
53+ auto xDtype = inputXPtr->GetDataType();
54+ OP_CHECK_IF(INPUT_SUPPORT_DTYPE_SET.count(xDtype) == 0,
55+ OP_LOGE(context->GetNodeName(),
56+ "Input x's data type is [%s], only supports INT32.",
57+ Ops::Base::ToString(static_cast<ge::DataType>(xDtype)).c_str()),
58+ return ge::GRAPH_FAILED);
59+ 
60+ auto inputDataSeqPtr = context->GetInputDesc(1);
61+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDataSeqPtr);
62+ auto dataSeqDtype = inputDataSeqPtr->GetDataType();
63+ OP_CHECK_IF(INPUT_SUPPORT_DTYPE_SET.count(dataSeqDtype) == 0,
64+ OP_LOGE(context->GetNodeName(),
65+ "Input dataSeq's data type is [%s], only supports INT32.",
66+ Ops::Base::ToString(static_cast<ge::DataType>(dataSeqDtype)).c_str()),
67+ return ge::GRAPH_FAILED);
68+ 
69+ auto levelIndexInput = context->GetOptionalInputDesc(LEVEL_INDEX_INDEX);
70+ if (levelIndexInput == nullptr) {
71+ tilingParam.hasLevelIndex = false;
72+ } else {
73+ auto levelIndexDtype = levelIndexInput->GetDataType();
74+ OP_CHECK_IF(INPUT_SUPPORT_DTYPE_SET.count(levelIndexDtype) == 0,
75+ OP_LOGE(context->GetNodeName(),
76+ "Input levelIndex's data type is [%s], only supports INT32.",
77+ Ops::Base::ToString(static_cast<ge::DataType>(levelIndexDtype)).c_str()),
78+ return ge::GRAPH_FAILED);
79+ }
80+ 
81+ auto outputYPtr = context->GetOutputDesc(0);
82+ OP_CHECK_NULL_WITH_CONTEXT(context, outputYPtr);
83+ auto yDtype = outputYPtr->GetDataType();
84+ OP_CHECK_IF(INPUT_SUPPORT_DTYPE_SET.count(yDtype) == 0,
85+ OP_LOGE(context->GetNodeName(),
86+ "Output y's data type is [%s], only supports INT32.",
87+ Ops::Base::ToString(static_cast<ge::DataType>(yDtype)).c_str()),
88+ return ge::GRAPH_FAILED);
89+ 
90+ return ge::GRAPH_SUCCESS;
91+}
92+ 
93+static ge::graphStatus CheckShape(const gert::TilingContext *context, MapIndexTilingParam &tilingParam)
94+{
95+ OP_LOGD(context->GetNodeName(), "CheckShape begin.");
96+ auto xShapePtr = context->GetInputShape(0);
97+ OP_CHECK_NULL_WITH_CONTEXT(context, xShapePtr);
98+ auto xShape = xShapePtr->GetStorageShape();
99+ 
100+ auto dataSeqShapePtr = context->GetInputShape(1);
101+ OP_CHECK_NULL_WITH_CONTEXT(context, dataSeqShapePtr);
102+ auto dataSeqShape = dataSeqShapePtr->GetStorageShape();
103+ 
104+ auto yShapePtr = context->GetOutputShape(0);
105+ OP_CHECK_NULL_WITH_CONTEXT(context, yShapePtr);
106+ auto yShape = Ops::NN::OpTiling::EnsureNotScalar(yShapePtr->GetStorageShape());
107+ OP_CHECK_IF(yShape.GetDimNum() != 1,
108+ OP_LOGE(context->GetNodeName(),
109+ "The shape of output y must be 1D."),
110+ return ge::GRAPH_FAILED);
111+
112+ OP_CHECK_IF(yShape.GetDim(0) != 1,
113+ OP_LOGE(context->GetNodeName(),
114+ "The shape of output y must be [1]."),
115+ return ge::GRAPH_FAILED);
116+ 
117+ OP_CHECK_IF(xShape.GetDimNum() != 1,
118+ OP_LOGE(context->GetNodeName(),
119+ "The shape of input x must be 1D."),
120+ return ge::GRAPH_FAILED);
121+ OP_CHECK_IF(xShape.GetDim(0) > X_MAX_DIM0,
122+ OP_LOGE(context->GetNodeName(),
123+ "The shape of input x must be less than 24000."),
124+ return ge::GRAPH_FAILED);
125+ OP_CHECK_IF(dataSeqShape.GetDimNum() != 1,
126+ OP_LOGE(context->GetNodeName(),
127+ "The shape of input data_seq must be 1D."),
128+ return ge::GRAPH_FAILED);
129+ 
130+ tilingParam.Dim1Size = xShape.GetDim(0);
131+ OP_CHECK_IF(dataSeqShape.GetDim(0) % xShape.GetDim(0) != 0,
132+ OP_LOGE(context->GetNodeName(),
133+ "the length of data_seq must be multiple of the length of x"),
134+ return ge::GRAPH_FAILED);
135+ 
136+ tilingParam.Dim0Size = dataSeqShape.GetDim(0) / xShape.GetDim(0);
137+ 
138+ OP_CHECK_IF(tilingParam.Dim0Size > DATA_SEQ_MAX_DIM0,
139+ OP_LOGE(context->GetNodeName(),
140+ "The input length of dataseq, which is a multiple of x, should be less than 256."),
141+ return ge::GRAPH_FAILED);
142+
143+ if(tilingParam.hasLevelIndex) {
144+ auto levelIndexShapePtr = context->GetOptionalInputShape(LEVEL_INDEX_INDEX);
145+ OP_CHECK_NULL_WITH_CONTEXT(context, levelIndexShapePtr);
146+ auto levelIndexShape = levelIndexShapePtr->GetStorageShape();
147+ OP_CHECK_IF(levelIndexShape.GetDimNum() != 1,
148+ OP_LOGE(context->GetNodeName(),
149+ "The shape of input level_index must be 1D."),
150+ return ge::GRAPH_FAILED);
151+ OP_CHECK_IF(tilingParam.Dim0Size != levelIndexShape.GetDim(0),
152+ OP_LOGE(context->GetNodeName(),
153+ "The input levelindex shape should be a multiple of dataseq, which should be a multiple of x."),
154+ return ge::GRAPH_FAILED);
155+ }
156+ 
157+ return ge::GRAPH_SUCCESS;
158+}
159+ 
160+static ge::graphStatus CheckAttr(const gert::TilingContext *context)
161+{
162+ OP_LOGD(context->GetNodeName(), "checkAttr begin.");
163+ 
164+ auto attrs = context->GetAttrs();
165+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
166+
167+ auto* attrTransPose = attrs->GetAttrPointer<bool>(ATTR_INDEX_TRANSPOSE);
168+ OP_CHECK_NULL_WITH_CONTEXT(context, attrTransPose);
169+ OP_CHECK_IF((*attrTransPose),
170+ OP_LOGE(
171+ context->GetNodeName(), "The attr transpose should be false on A5, please check"),
172+ return ge::GRAPH_FAILED);
173+
174+ return ge::GRAPH_SUCCESS;
175+}
176+ 
177+static ge::graphStatus GetPlatInfo(const gert::TilingContext *context, MapIndexTilingParam &tilingParam)
178+{
179+ OP_LOGD(context->GetNodeName(), "GetPlatInfo begin.");
180+ auto platformInfo = context->GetPlatformInfo();
181+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
182+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
183+ tilingParam.totalCoreNum = ascendcPlatform.GetCoreNumAiv();
184+ OP_CHECK_IF((tilingParam.totalCoreNum <= 0),
185+ OP_LOGE(context->GetNodeName(), "Failed to get core num."), return ge::GRAPH_FAILED);
186+ uint64_t ubSize;
187+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
188+ tilingParam.ubSize = static_cast<int64_t>(ubSize) - RESERVED_UB_SIZE;
189+ OP_CHECK_IF((tilingParam.ubSize <= 0),
190+ OP_LOGE(context->GetNodeName(), "Failed to get ub size."), return ge::GRAPH_FAILED);
191+ tilingParam.vfLen = Ops::Base::GetVRegSize(context);
192+ tilingParam.workspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
193+ return ge::GRAPH_SUCCESS;
194+}
195+ 
196+static ge::graphStatus DoTiling(const gert::TilingContext *context, MapIndexTilingParam &tilingParam)
197+{
198+ OP_LOGD(context->GetNodeName(), "DoTiling begin.");
199+ tilingParam.normalCoreProcessNum = CeilDivide(tilingParam.Dim0Size, tilingParam.totalCoreNum);
200+ tilingParam.usedCoreNum = CeilDivide(tilingParam.Dim0Size, tilingParam.normalCoreProcessNum);
201+ tilingParam.tailCoreProcessNum = tilingParam.Dim0Size - tilingParam.normalCoreProcessNum * (tilingParam.usedCoreNum - 1);
202+ tilingParam.Dim1SizeAlign = Ops::Base::CeilAlign(tilingParam.Dim1Size, VL_NUMS);
203+ 
204+ int64_t OneRowUB = tilingParam.Dim1SizeAlign * sizeof(int32_t);
205+ int64_t rowsNums = tilingParam.ubSize / OneRowUB;
206+ int64_t dataSeqNums = rowsNums - 1;
207+ 
208+ if (dataSeqNums > tilingParam.normalCoreProcessNum){
209+ tilingParam.CopyInDim0 = tilingParam.normalCoreProcessNum;
210+ tilingParam.CopyInDim0Times = 1;
211+ } else {
212+ tilingParam.CopyInDim0 = dataSeqNums;
213+ tilingParam.CopyInDim0Times = CeilDivide(tilingParam.normalCoreProcessNum, dataSeqNums);
214+ tilingParam.tailCopyInDim0Times = CeilDivide(tilingParam.tailCoreProcessNum, dataSeqNums);
215+ }
216+ 
217+ if (tilingParam.CopyInDim0 < NUM_TWO_DB ){
218+ tilingParam.doubleBuffNum = 1;
219+ } else {
220+ tilingParam.doubleBuffNum = NUM_TWO_DB;
221+ tilingParam.CopyInDim0 = Ops::Base::CeilAlign(tilingParam.CopyInDim0, tilingParam.doubleBuffNum);
222+ }
223+ 
224+ return ge::GRAPH_SUCCESS;
225+}
226+ 
227+inline static ge::graphStatus SetTilingData(gert::TilingContext *context,
228+ const MapIndexTilingParam &tilingParam, MapIndexTilingData &tilingData)
229+{
230+ OP_LOGD(context->GetNodeName(), "SetTilingData begin.");
231+ tilingData.set_totalCoreNum(tilingParam.totalCoreNum);
232+ tilingData.set_usedCoreNum(tilingParam.usedCoreNum);
233+ tilingData.set_normalCoreProcessNum(tilingParam.normalCoreProcessNum);
234+ tilingData.set_tailCoreProcessNum(tilingParam.tailCoreProcessNum);
235+ tilingData.set_Dim1Size(tilingParam.Dim1Size);
236+ tilingData.set_Dim1SizeAlign(tilingParam.Dim1SizeAlign);
237+ tilingData.set_CopyInDim0(tilingParam.CopyInDim0);
238+ tilingData.set_CopyInDim0Times(tilingParam.CopyInDim0Times);
239+ tilingData.set_tailCopyInDim0Times(tilingParam.tailCopyInDim0Times);
240+ tilingData.set_doubleBuffNum(tilingParam.doubleBuffNum);
241+ 
242+ OP_CHECK_IF(tilingData.GetDataSize() > context->GetRawTilingData()->GetCapacity(),
243+ OP_LOGE(context->GetNodeName(), "tiling datasize: %zu is bigger than %zu",
244+ tilingData.GetDataSize(), context->GetRawTilingData()->GetCapacity()),
245+ return ge::GRAPH_FAILED);
246+ tilingData.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
247+ context->GetRawTilingData()->SetDataSize(tilingData.GetDataSize());
248+ context->SetBlockDim(tilingData.get_totalCoreNum());
249+ context->SetTilingKey(1);
250+ size_t *workspaces = context->GetWorkspaceSizes(1);
251+ OP_CHECK_NULL_WITH_CONTEXT(context, workspaces);
252+ workspaces[0] = static_cast<size_t>(WORKSPACE_BUFFER);
253+ return ge::GRAPH_SUCCESS;
254+}
255+ 
256+inline static void PrintTilingData(const gert::TilingContext *context, MapIndexTilingData &tilingData)
257+{
258+ OP_LOGI(context->GetNodeName(), "tilingData is totalCoreNum:%ld, usedCoreNum:%ld, normalCoreProcessNum:%ld, \
259+ tailCoreProcessNum:%ld, Dim1Size:%ld, Dim1SizeAlign:%ld, CopyInDim0:%ld, CopyInDim0Times:%ld, tailCopyInDim0Times:%ld, doubleBuffNum:%ld",
260+ tilingData.get_totalCoreNum(), tilingData.get_usedCoreNum(), tilingData.get_normalCoreProcessNum(),
261+ tilingData.get_tailCoreProcessNum(), tilingData.get_Dim1Size(), tilingData.get_Dim1SizeAlign(),
262+ tilingData.get_CopyInDim0(), tilingData.get_CopyInDim0Times(), tilingData.get_tailCopyInDim0Times(), tilingData.get_doubleBuffNum());
263+}
264+ 
265+ge::graphStatus Tiling4MapIndex(gert::TilingContext *context)
266+{
267+ OP_LOGD(context->GetNodeName(), "Tiling4MapIndex running begin.");
268+ 
269+ MapIndexTilingParam tilingParam;
270+ 
271+ OP_CHECK_IF(CheckDtype(context, tilingParam) != ge::GRAPH_SUCCESS,
272+ OP_LOGE(context->GetNodeName(), "The data type check failed."), return ge::GRAPH_FAILED);
273+ 
274+ OP_CHECK_IF(CheckShape(context, tilingParam) != ge::GRAPH_SUCCESS,
275+ OP_LOGE(context->GetNodeName(), "The shape check failed."), return ge::GRAPH_FAILED);
276+ 
277+ OP_CHECK_IF(GetPlatInfo(context, tilingParam) != ge::GRAPH_SUCCESS,
278+ OP_LOGE(context->GetNodeName(), "GetPlatInfo failed."), return ge::GRAPH_FAILED);
279+ 
280+ OP_CHECK_IF(DoTiling(context, tilingParam) != ge::GRAPH_SUCCESS,
281+ OP_LOGE(context->GetNodeName(), "DoTiling failed."), return ge::GRAPH_FAILED);
282+ 
283+ OP_CHECK_IF(CheckAttr(context) != ge::GRAPH_SUCCESS,
284+ OP_LOGE(context->GetNodeName(), "check attr failed."), return ge::GRAPH_FAILED);
285+ 
286+ MapIndexTilingData tilingData;
287+ OP_CHECK_IF(SetTilingData(context, tilingParam, tilingData) != ge::GRAPH_SUCCESS,
288+ OP_LOGE(context->GetNodeName(), "SetContext fail."),
289+ return ge::GRAPH_FAILED);
290+ 
291+ PrintTilingData(context, tilingData);
292+ return ge::GRAPH_SUCCESS;
293+}
294+ 
295+ge::graphStatus TilingPrepare4MapIndex(gert::TilingParseContext *context)
296+{
297+ OP_LOGD(context->GetNodeName(), "TilingPrepare4MapIndex entering.");
298+ return ge::GRAPH_SUCCESS;
299+}
300+ 
301+IMPL_OP_OPTILING(MapIndex)
302+ .Tiling(Tiling4MapIndex)
303+ .TilingParse<MapIndexCompileInfo>(TilingPrepare4MapIndex);
304+}
Aindex/map_index/op_host/map_index_tiling_arch35.h+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+/* !
13+ * \file map_index_tiling_arch35.h
14+ * \brief
15+ */
16+#ifndef AIR_CXX_RUNTIME_V2_OP_IMPL_MAP_INDEX_TILING_H_
17+#define AIR_CXX_RUNTIME_V2_OP_IMPL_MAP_INDEX_TILING_H_
18+ 
19+#include "register/tilingdata_base.h"
20+ 
21+namespace optiling {
22+ 
23+BEGIN_TILING_DATA_DEF(MapIndexTilingData)
24+ TILING_DATA_FIELD_DEF(int64_t, totalCoreNum);
25+ TILING_DATA_FIELD_DEF(int64_t, usedCoreNum); // 实际使用的核数
26+ TILING_DATA_FIELD_DEF(int64_t, normalCoreProcessNum); // 单核循环次数
27+ TILING_DATA_FIELD_DEF(int64_t, tailCoreProcessNum); // 尾核循环次数
28+ TILING_DATA_FIELD_DEF(int64_t, Dim1Size);
29+ TILING_DATA_FIELD_DEF(int64_t, Dim1SizeAlign);
30+ TILING_DATA_FIELD_DEF(int64_t, CopyInDim0); // 一次最多搬几行
31+ TILING_DATA_FIELD_DEF(int64_t, CopyInDim0Times); // 需要搬运几次
32+ TILING_DATA_FIELD_DEF(int64_t, tailCopyInDim0Times); // 尾块需要搬运几次
33+ TILING_DATA_FIELD_DEF(int64_t, doubleBuffNum);
34+END_TILING_DATA_DEF;
35+ 
36+REGISTER_TILING_DATA_CLASS(MapIndex, MapIndexTilingData)
37+ 
38+struct MapIndexCompileInfo
39+{
40+ int64_t coreNum = 0;
41+ int64_t ubSize = 0;
42+};
43+ 
44+struct MapIndexTilingParam
45+{
46+ int64_t totalCoreNum{ 0 };
47+ int64_t ubSize { 0 };
48+ uint32_t vfLen { 0 };
49+ uint32_t workspaceSize { 0 };
50+ int64_t usedCoreNum { 0 };
51+ int64_t normalCoreProcessNum {0};
52+ int64_t tailCoreProcessNum {0};
53+ int64_t Dim0Size {0};
54+ int64_t Dim1Size {0};
55+ int64_t Dim1SizeAlign {0};
56+ int64_t CopyInDim0 {1};
57+ int64_t CopyInDim0Times {1};
58+ int64_t tailCopyInDim0Times {1};
59+ int64_t doubleBuffNum {0};
60+ bool hasLevelIndex = true;
61+};
62+ 
63+}
64+#endif
Aindex/map_index/op_kernel/arch35/map_index.h+229-0
@@ -0,0 +1,229 @@
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 map_index.h
13+ * \brief
14+ */
15+#ifndef MAP_INDEX_H_
16+#define MAP_INDEX_H_
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_operator_intf.h"
20+#include "kernel_tiling/kernel_tiling.h"
21+ 
22+namespace MapIndexOp {
23+ 
24+using namespace AscendC;
25+constexpr uint32_t BLOCK_INT8 = 8;
26+constexpr uint32_t VF_LEN_INT32 = 64;
27+constexpr uint32_t ONE_BLOCK_NUM = 32;
28+constexpr uint8_t MASK_RESULT = 17; // 00010001
29+class MapIndex {
30+public:
31+ __aicore__ inline MapIndex(){};
32+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR dataSeq, GM_ADDR y, GM_ADDR workspace, const MapIndexTilingData &tilingData, TPipe* inputPipe);
33+ __aicore__ inline void ParseTilingData(const MapIndexTilingData &tilingData);
34+ __aicore__ inline void CopyInX(int64_t calCount);
35+ __aicore__ inline void CopyInDataSeq(int64_t offset, int64_t calCount);
36+ __aicore__ inline void ProcessOneRow(LocalTensor<int32_t> &xLocalRes, int64_t offset, int32_t rowCount, LocalTensor<uint8_t> &compareRes, LocalTensor<uint8_t> &compareResultMask);
37+ __aicore__ inline void Process();
38+ __aicore__ inline void ComputeOneRowMask(LocalTensor<int32_t> &xLocalRes, LocalTensor<int32_t> &dataSeqLocalRes, LocalTensor<uint8_t> &maskRes, int32_t calCount);
39+ __aicore__ inline void CopyOut();
40+ 
41+private:
42+ TPipe *pipe;
43+ /* global memory address */
44+ GlobalTensor<int32_t> xGm_;
45+ GlobalTensor<int32_t> dataSeqGm_;
46+ GlobalTensor<int32_t> yGm_;
47+ 
48+ TQue<QuePosition::VECIN, 1> inXQueue_;
49+ TQue<QuePosition::VECIN, 1> inDataSeqQueue_;
50+ TQue<QuePosition::VECOUT, 1> outYQueue_;
51+ TBuf<TPosition::VECCALC> maskBuf_;
52+ TBuf<TPosition::VECCALC> compareBuf_;
53+ TBuf<TPosition::VECCALC> compareResultBuf_;
54+ 
55+ // 变量区
56+ uint32_t usedCoreNum_ = 0;
57+ int64_t curCoreProcessNum_ = 0;
58+ int64_t normalCoreProcessNum_ = 0;
59+ int64_t tailCoreProcessNum_ = 0;
60+ int64_t Dim1Size_ = 0;
61+ int64_t Dim1SizeAlign_ = 0;
62+ int64_t CopyInDim0_ = 0;
63+ int64_t CopyInDim0Times_ = 0;
64+ int64_t tailCopyInDim0Times_ = 0;
65+ int64_t doubleBuffNum_ = 0;
66+ uint32_t blockIdx_ = 0;
67+ 
68+ /* variable */
69+ int64_t loopNum_ = 0;
70+};
71+ 
72+__aicore__ inline void MapIndex::ParseTilingData(const MapIndexTilingData &tilingData)
73+{
74+ usedCoreNum_ = tilingData.usedCoreNum;
75+ normalCoreProcessNum_ = tilingData.normalCoreProcessNum;
76+ tailCoreProcessNum_ = tilingData.tailCoreProcessNum;
77+ Dim1Size_ = tilingData.Dim1Size;
78+ Dim1SizeAlign_ = tilingData.Dim1SizeAlign;
79+ CopyInDim0_ = tilingData.CopyInDim0;
80+ CopyInDim0Times_ = tilingData.CopyInDim0Times;
81+ tailCopyInDim0Times_ = tilingData.tailCopyInDim0Times;
82+ doubleBuffNum_ = tilingData.doubleBuffNum;
83+}
84+ 
85+__aicore__ inline void MapIndex::Init(GM_ADDR x, GM_ADDR dataSeq, GM_ADDR y, GM_ADDR workspace, const MapIndexTilingData &tilingData, TPipe* inputPipe)
86+{
87+ pipe = inputPipe;
88+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
89+ blockIdx_ = GetBlockIdx();
90+ ParseTilingData(tilingData);
91+ // shield global memory address between different core
92+ uint64_t intraCoreOffset = blockIdx_ * normalCoreProcessNum_ * Dim1Size_;
93+ 
94+ if (blockIdx_ + 1 == usedCoreNum_) {
95+ curCoreProcessNum_ = tailCoreProcessNum_;
96+ } else {
97+ curCoreProcessNum_ = normalCoreProcessNum_;
98+ }
99+ 
100+ xGm_.SetGlobalBuffer((__gm__ int32_t *)x);
101+ dataSeqGm_.SetGlobalBuffer((__gm__ int32_t *)dataSeq);
102+ yGm_.SetGlobalBuffer((__gm__ int32_t *)y);
103+ 
104+ if (blockIdx_ == 0) {
105+ InitOutput<int32_t>(yGm_, 1, -1);
106+ PipeBarrier<PIPE_ALL>();
107+ }
108+ SyncAll();
109+ 
110+ pipe->InitBuffer(inXQueue_, 1, Dim1SizeAlign_ * sizeof(int32_t));
111+ pipe->InitBuffer(inDataSeqQueue_, static_cast<uint8_t>(doubleBuffNum_), Dim1SizeAlign_ * sizeof(int32_t));
112+ pipe->InitBuffer(maskBuf_, BLOCK_INT8 * sizeof(int32_t));
113+ pipe->InitBuffer(compareBuf_, BLOCK_INT8 * sizeof(int32_t));
114+ pipe->InitBuffer(compareResultBuf_, BLOCK_INT8 * sizeof(int32_t));
115+ pipe->InitBuffer(outYQueue_, 1, BLOCK_INT8 * sizeof(int32_t));
116+}
117+ 
118+__aicore__ inline void MapIndex::CopyInX(int64_t calCount)
119+{
120+ LocalTensor<int32_t> xLocal = inXQueue_.AllocTensor<int32_t>();
121+ Duplicate<int32_t>(xLocal[Dim1SizeAlign_ - VF_LEN_INT32], static_cast<int32_t>(0), VF_LEN_INT32);
122+ event_t eventIDVToMTE2 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE2));
123+ SetFlag<HardEvent::V_MTE2>(eventIDVToMTE2);
124+ WaitFlag<HardEvent::V_MTE2>(eventIDVToMTE2);
125+ DataCopyExtParams copyParams{ static_cast<uint16_t>(1), static_cast<uint32_t>(calCount * sizeof(int32_t)),
126+ static_cast<uint32_t>(0), static_cast<uint32_t>(0), static_cast<uint32_t>(0) };
127+ 
128+ DataCopyPadExtParams<int32_t> padParams{ false, static_cast<uint8_t>(0), static_cast<uint8_t>(0),
129+ static_cast<int32_t>(0) };
130+ DataCopyPad(xLocal, xGm_, copyParams, padParams);
131+ inXQueue_.EnQue(xLocal);
132+}
133+ 
134+__aicore__ inline void MapIndex::CopyInDataSeq(int64_t offset, int64_t calCount)
135+{
136+ LocalTensor<int32_t> dataSeqLocal = inDataSeqQueue_.AllocTensor<int32_t>();
137+ Duplicate<int32_t>(dataSeqLocal[Dim1SizeAlign_ - VF_LEN_INT32], static_cast<int32_t>(0), VF_LEN_INT32);
138+ event_t eventIDVToMTE2 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE2));
139+ SetFlag<HardEvent::V_MTE2>(eventIDVToMTE2);
140+ WaitFlag<HardEvent::V_MTE2>(eventIDVToMTE2);
141+ DataCopyExtParams copyParams{ static_cast<uint16_t>(1), static_cast<uint32_t>(calCount * sizeof(int32_t)),
142+ static_cast<uint32_t>(0), static_cast<uint32_t>(0), static_cast<uint32_t>(0) };
143+ uint64_t intraCoreOffset = blockIdx_ * normalCoreProcessNum_ * Dim1Size_;
144+ DataCopyPadExtParams<int32_t> padParams{ false, static_cast<uint8_t>(0), static_cast<uint8_t>(0),
145+ static_cast<int32_t>(0) };
146+ DataCopyPad(dataSeqLocal, dataSeqGm_[offset+intraCoreOffset], copyParams, padParams);
147+ inDataSeqQueue_.EnQue(dataSeqLocal);
148+}
149+ 
150+__aicore__ inline void MapIndex::ProcessOneRow(LocalTensor<int32_t> &xLocalRes, int64_t offset, int32_t rowCount, LocalTensor<uint8_t> &compareRes, LocalTensor<uint8_t> &compareResultMask)
151+{
152+ CopyInDataSeq(offset, rowCount);
153+ LocalTensor<int32_t> dataSeqLocalRes = inDataSeqQueue_.DeQue<int32_t>();
154+ LocalTensor<uint8_t> maskRes = maskBuf_.Get<uint8_t>();
155+ ComputeOneRowMask(xLocalRes, dataSeqLocalRes, maskRes, Dim1SizeAlign_);
156+ inDataSeqQueue_.FreeTensor(dataSeqLocalRes);
157+ Compare(compareResultMask, compareRes, maskRes, CMPMODE::EQ, ONE_BLOCK_NUM);
158+}
159+ 
160+__aicore__ inline void MapIndex::Process()
161+{
162+ if (blockIdx_ >= usedCoreNum_) {
163+ return;
164+ }
165+ CopyInX(Dim1Size_);
166+ LocalTensor<uint8_t> compareRes = compareBuf_.Get<uint8_t>();
167+ LocalTensor<uint8_t> compareResultMask = compareResultBuf_.Get<uint8_t>();
168+ Duplicate<uint8_t>(compareRes, static_cast<uint8_t>(MASK_RESULT), ONE_BLOCK_NUM); // 00010001
169+ LocalTensor<int32_t> xLocalRes = inXQueue_.DeQue<int32_t>();
170+ int64_t loopNum = curCoreProcessNum_;
171+ for (int64_t loopIndex = 0; loopIndex < loopNum; loopIndex++) {
172+ ProcessOneRow(xLocalRes, Dim1Size_ * loopIndex, Dim1Size_, compareRes, compareResultMask);
173+ event_t eventIDVToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_S));
174+ SetFlag<HardEvent::V_S>(eventIDVToS);
175+ WaitFlag<HardEvent::V_S>(eventIDVToS);
176+ uint32_t a = compareResultMask.ReinterpretCast<uint32_t>().GetValue(0);
CANN-robot
CANN-robotCANN-robot1月29日

代码结构与可维护性: 变量名'a'过于简单，无法表达其含义。从上下文看，它存储的是compareResultMask转换为uint32_t后的值，用于判断是否全为1（0xFFFFFFFF）。

问题类型: 代码结构与可维护性 文件路径: index/map_index/op_kernel/arch35/map_index.h 行号: 176 问题代码:

uint32_t a = compareResultMask.ReinterpretCast<uint32_t>().GetValue(0);

修改建议:

使用更有意义的变量名，如：uint32_t maskValue = compareResultMask.ReinterpretCast<uint32_t>().GetValue(0);

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

likedislike
177+ if(a == static_cast<uint32_t>(0xFFFFFFFF)) {
178+ LocalTensor<int32_t> outYLocal = outYQueue_.AllocTensor<int32_t>();
179+ Duplicate<int32_t>(outYLocal, static_cast<int32_t>(loopIndex + blockIdx_ * normalCoreProcessNum_), BLOCK_INT8);
180+ outYQueue_.EnQue(outYLocal);
181+ CopyOut();
182+ break;
183+ }
184+ }
185+ inXQueue_.FreeTensor(xLocalRes);
186+}
187+ 
188+__aicore__ inline void MapIndex::ComputeOneRowMask(LocalTensor<int32_t> &xLocalRes, LocalTensor<int32_t> &dataSeqLocalRes, LocalTensor<uint8_t> &maskRes, int32_t calCount)
189+{
190+ uint32_t dtypeSize = sizeof(int32_t);
191+ uint32_t vl = VECTOR_REG_WIDTH / dtypeSize;
192+ uint16_t loopNum = (calCount + vl - 1) / vl;
193+ uint32_t vlSize = vl;
CANN-robot
CANN-robotCANN-robot1月29日

未使用的变量: ComputeOneRowMask函数第194行定义了vlSize变量，但在后续代码中从未使用。这可能是调试遗留的代码或复制粘贴错误，应该移除未使用的变量以避免混淆。

问题类型: 未使用的变量 文件路径: index/map_index/op_kernel/arch35/map_index.h 行号: 193 问题代码:

uint32_t vlSize = vl;

修改建议:

删除未使用的vlSize变量定义，或者如果确实需要，添加注释说明其用途。

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

likedislike
194+ __ubuf__ int32_t* xAddr = (__ubuf__ int32_t*)xLocalRes.GetPhyAddr();
195+ __ubuf__ int32_t* dataSeqAddr = (__ubuf__ int32_t*)dataSeqLocalRes.GetPhyAddr();
196+ __ubuf__ uint8_t* maskAddr = (__ubuf__ uint8_t*)maskRes.GetPhyAddr();
197+ 
198+ __VEC_SCOPE__
199+ {
200+ AscendC::MicroAPI::RegTensor<int32_t> xReg;
201+ AscendC::MicroAPI::RegTensor<int32_t> dataSeqReg;
202+ AscendC::MicroAPI::MaskReg preg0;
203+ AscendC::MicroAPI::MaskReg preg1 = AscendC::MicroAPI::CreateMask<int32_t, AscendC::MicroAPI::MaskPattern::ALL>();
CANN-robot
CANN-robotCANN-robot1月29日

代码结构与可维护性: ComputeOneRowMask函数中，preg1变量被创建但未使用。创建了AscendC::MicroAPI::MaskReg preg1 = AscendC::MicroAPI::CreateMask<int32_t, AscendC::MicroAPI::MaskPattern::ALL>()，但在后续代码中从未使用。

问题类型: 代码结构与可维护性 文件路径: index/map_index/op_kernel/arch35/map_index.h 行号: 203 问题代码:

AscendC::MicroAPI::MaskReg preg1 = AscendC::MicroAPI::CreateMask<int32_t, AscendC::MicroAPI::MaskPattern::ALL>();

修改建议:

删除未使用的preg1变量，除非它是为未来的功能预留（此时应添加注释说明）。

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

likedislike
204+ AscendC::MicroAPI::MaskReg preg2;
205+ AscendC::MicroAPI::MaskReg pregResultMask = AscendC::MicroAPI::CreateMask<int32_t, AscendC::MicroAPI::MaskPattern::ALL>();
206+
207+ uint32_t sreg0 = calCount;
208+ for (uint16_t i = 0; i < loopNum; i++) { // 256B
209+ preg0 = AscendC::MicroAPI::UpdateMask<int32_t>(sreg0);
210+ AscendC::MicroAPI::DataCopy(xReg, xAddr + i * vl);
211+ AscendC::MicroAPI::DataCopy(dataSeqReg, dataSeqAddr + i * vl);
212+ AscendC::MicroAPI::Compare<int32_t, CMPMODE::EQ>(preg2, xReg, dataSeqReg, preg0);
213+ AscendC::MicroAPI::MaskAnd(pregResultMask, pregResultMask, preg2, preg0);
214+ }
215+ AscendC::MicroAPI::DataCopy(maskAddr, pregResultMask);
216+ }
217+}
218+ 
219+__aicore__ inline void MapIndex::CopyOut()
220+{
221+ LocalTensor<int32_t> yOutLocal = outYQueue_.DeQue<int32_t>();
222+ DataCopyExtParams copyParams{ static_cast<uint16_t>(1), static_cast<uint32_t>(1 * sizeof(int32_t)),
223+ static_cast<uint32_t>(0), static_cast<uint32_t>(0), static_cast<uint32_t>(0) };
224+ DataCopyPad(yGm_, yOutLocal, copyParams);
225+ outYQueue_.FreeTensor(yOutLocal);
226+}
227+ 
228+}
CANN-robot
CANN-robotCANN-robot1月29日

代码结构与可维护性: 文件末尾缺少换行符（No newline at end of file），这虽然不是功能性问题，但可能在某些工具中引起警告，且不符合代码规范。

问题类型: 代码结构与可维护性 文件路径: index/map_index/op_kernel/arch35/map_index.h 行号: 228 问题代码:

}
#endif

修改建议:

在文件末尾添加一个空行（换行符）。

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

likedislike
229+#endif
Aindex/map_index/op_kernel/map_index_apt.cpp+39-0
@@ -0,0 +1,39 @@
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 map_index_apt.cpp
13+ * \brief kernel file of map_index
14+ */
15+ 
16+#include "kernel_operator.h"
17+#include "arch35/map_index.h"
18+
19+using namespace MapIndexOp;
20+ 
21+extern "C" __global__ __aicore__ void map_index(
22+ GM_ADDR x, GM_ADDR data_seq, GM_ADDR level_index, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
23+ if (workspace == nullptr) {
24+ return;
25+ }
26+ SetSysWorkspace(workspace);
27+ GM_ADDR userWs = AscendC::GetUserWorkspace(workspace);
28+ if (userWs == nullptr) {
29+ return;
30+ }
31+ GET_TILING_DATA(tilingData, tiling);
32+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
33+ TPipe pipe;
34+ if (TILING_KEY_IS(1)) {
35+ MapIndex op;
36+ op.Init(x, data_seq, y, userWs, tilingData, &pipe);
37+ op.Process();
38+ }
39+}
Aindex/map_index/tests/CMakeLists.txt+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+message(STATUS "=== Debug: start ops.index.map_index.tests.CMakeLists.txt ")
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
14+foreach(SUB_DIR ${CURRENT_DIRS})
15+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16+ add_subdirectory(${SUB_DIR})
17+ endif()
18+endforeach()
Aindex/map_index/tests/ut/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
13+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aindex/map_index/tests/ut/op_host/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
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+file(GLOB CURRENT_DIR RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_HOST_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
15+endif()
Aindex/map_index/tests/ut/op_host/arch35/test_map_index_tiling.cpp+765-0
@@ -0,0 +1,765 @@
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_map_index_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include <iostream>
17+#include <fstream>
18+#include <vector>
19+#include "log/log.h"
20+#include <gtest/gtest.h>
21+#include "register/op_impl_registry.h"
22+#include "platform/platform_infos_def.h"
23+#include "ut_op_common.h"
24+#include "ut_op_util.h"
25+#include "index/map_index/op_host/map_index_tiling_arch35.h"
26+#include "kernel_run_context_facker.h"
27+#include "test_cube_util.h"
28+#include "exe_graph/runtime/storage_format.h"
29+#include "exe_graph/runtime/storage_shape.h"
30+ 
31+ using namespace std;
32+
33+class MapIndexTiling : public testing::Test
34+{
35+protected:
36+ static void SetUpTestCase()
37+ {
38+ std::cout << "MapIndexTiling SetUp" << std::endl;
39+ }
40+ 
41+ static void TearDownTestCase()
42+ {
43+ std::cout << "MapIndexTiling TearDown" << std::endl;
44+ }
45+};
46+ 
47+TEST_F(MapIndexTiling, test_tiling_pass_01) {
48+ std::string opType("MapIndex");
49+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
50+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
51+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
52+ 
53+ string compileInfoString = R"({
54+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
55+ "Intrinsic_fix_pipe_l0c2out": false,
56+ "Intrinsic_data_move_l12ub": true,
57+ "Intrinsic_data_move_l0c2ub": true,
58+ "Intrinsic_data_move_out2l1_nd2nz": false,
59+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
60+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
61+ "CORE_NUM": 64}
62+ })";
63+ map<string, string> socInfos;
64+ map<string, string> aicoreSpec;
65+ map<string, string> intrinsics;
66+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
67+ 
68+ // platform info
69+ fe::PlatFormInfos platformInfo;
70+ platformInfo.Init();
71+ // compile info
72+ optiling::MapIndexCompileInfo compileInfo;
73+ compileInfo.coreNum = 64;
74+ compileInfo.ubSize = 253952;
75+ 
76+ // tilingFunc simulate
77+ auto param = gert::TilingData::CreateCap(4096);
78+ ASSERT_NE(param, nullptr);
79+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
80+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
81+ gert::StorageShape xShape = {{64}, {64}};
82+ gert::StorageShape dataSeqShape = {{640}, {640}};
83+ gert::StorageShape levelIndexShape = {{10}, {10}};
84+ gert::StorageShape yShape = {{1}, {1}};
85+ auto holder = gert::TilingContextFaker()
86+ .NodeIoNum(3, 1)
87+ .IrInstanceNum({1, 1, 1})
88+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
89+ .OutputShapes({&yShape})
90+ .CompileInfo(&compileInfo)
91+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
92+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
93+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
94+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
95+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
96+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
97+ .TilingData(param.get())
98+ .Workspace(wsSize)
99+ .Build();
100+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
101+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
102+ 
103+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
104+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
105+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
106+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
107+ 
108+ // workspaces nullptr return failed
109+ std::cout << "test>> holder.GetContext end" << std::endl;
110+ 
111+ if (tilingFunc == nullptr) {
112+ std::cout << "test>> tilingFunc is invalid" << std::endl;
113+ } else {
114+ std::cout << "test>> tilingFunc is valid" << std::endl;
115+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_SUCCESS);
116+ }
117+}
118+
119+TEST_F(MapIndexTiling, test_tiling_dtype_dataSeq_dtype_fail_02) {
120+ std::string opType("MapIndex");
121+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
122+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
123+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
124+ 
125+ string compileInfoString = R"({
126+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
127+ "Intrinsic_fix_pipe_l0c2out": false,
128+ "Intrinsic_data_move_l12ub": true,
129+ "Intrinsic_data_move_l0c2ub": true,
130+ "Intrinsic_data_move_out2l1_nd2nz": false,
131+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
132+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
133+ "CORE_NUM": 64}
134+ })";
135+ map<string, string> socInfos;
136+ map<string, string> aicoreSpec;
137+ map<string, string> intrinsics;
138+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
139+ 
140+ // platform info
141+ fe::PlatFormInfos platformInfo;
142+ platformInfo.Init();
143+ // compile info
144+ optiling::MapIndexCompileInfo compileInfo;
145+ compileInfo.coreNum = 64;
146+ compileInfo.ubSize = 253952;
147+ 
148+ // tilingFunc simulate
149+ auto param = gert::TilingData::CreateCap(4096);
150+ ASSERT_NE(param, nullptr);
151+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
152+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
153+ gert::StorageShape xShape = {{64}, {64}};
154+ gert::StorageShape dataSeqShape = {{640}, {640}};
155+ gert::StorageShape levelIndexShape = {{10}, {10}};
156+ gert::StorageShape yShape = {{1}, {1}};
157+ auto holder = gert::TilingContextFaker()
158+ .NodeIoNum(3, 1)
159+ .IrInstanceNum({1, 1, 1})
160+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
161+ .OutputShapes({&yShape})
162+ .CompileInfo(&compileInfo)
163+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
164+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
165+ .NodeInputTd(1, ge::DT_INT64, ge::FORMAT_ND, ge::FORMAT_ND)
166+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
167+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
168+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
169+ .TilingData(param.get())
170+ .Workspace(wsSize)
171+ .Build();
172+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
173+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
174+ 
175+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
176+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
177+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
178+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
179+ 
180+ // workspaces nullptr return failed
181+ std::cout << "test>> holder.GetContext end" << std::endl;
182+ 
183+ if (tilingFunc == nullptr) {
184+ std::cout << "test>> tilingFunc is invalid" << std::endl;
185+ } else {
186+ std::cout << "test>> tilingFunc is valid" << std::endl;
187+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
188+ }
189+}
190+ 
191+TEST_F(MapIndexTiling, test_tiling_dtype_y_dtype_fail_03) {
192+ std::string opType("MapIndex");
193+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
194+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
195+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
196+ 
197+ string compileInfoString = R"({
198+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
199+ "Intrinsic_fix_pipe_l0c2out": false,
200+ "Intrinsic_data_move_l12ub": true,
201+ "Intrinsic_data_move_l0c2ub": true,
202+ "Intrinsic_data_move_out2l1_nd2nz": false,
203+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
204+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
205+ "CORE_NUM": 64}
206+ })";
207+ map<string, string> socInfos;
208+ map<string, string> aicoreSpec;
209+ map<string, string> intrinsics;
210+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
211+ 
212+ // platform info
213+ fe::PlatFormInfos platformInfo;
214+ platformInfo.Init();
215+ // compile info
216+ optiling::MapIndexCompileInfo compileInfo;
217+ compileInfo.coreNum = 64;
218+ compileInfo.ubSize = 253952;
219+ 
220+ // tilingFunc simulate
221+ auto param = gert::TilingData::CreateCap(4096);
222+ ASSERT_NE(param, nullptr);
223+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
224+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
225+ gert::StorageShape xShape = {{64}, {64}};
226+ gert::StorageShape dataSeqShape = {{640}, {640}};
227+ gert::StorageShape levelIndexShape = {{10}, {10}};
228+ gert::StorageShape yShape = {{1}, {1}};
229+ auto holder = gert::TilingContextFaker()
230+ .NodeIoNum(3, 1)
231+ .IrInstanceNum({1, 1, 1})
232+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
233+ .OutputShapes({&yShape})
234+ .CompileInfo(&compileInfo)
235+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
236+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
237+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
238+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
239+ .NodeOutputTd(0, ge::DT_INT64, ge::FORMAT_ND, ge::FORMAT_ND)
240+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
241+ .TilingData(param.get())
242+ .Workspace(wsSize)
243+ .Build();
244+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
245+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
246+ 
247+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
248+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
249+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
250+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
251+ 
252+ // workspaces nullptr return failed
253+ std::cout << "test>> holder.GetContext end" << std::endl;
254+ 
255+ if (tilingFunc == nullptr) {
256+ std::cout << "test>> tilingFunc is invalid" << std::endl;
257+ } else {
258+ std::cout << "test>> tilingFunc is valid" << std::endl;
259+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
260+ }
261+}
262+ 
263+TEST_F(MapIndexTiling, test_tiling_dtype_dataSeq_shape_fail_03) {
264+ std::string opType("MapIndex");
265+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
266+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
267+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
268+ 
269+ string compileInfoString = R"({
270+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
271+ "Intrinsic_fix_pipe_l0c2out": false,
272+ "Intrinsic_data_move_l12ub": true,
273+ "Intrinsic_data_move_l0c2ub": true,
274+ "Intrinsic_data_move_out2l1_nd2nz": false,
275+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
276+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
277+ "CORE_NUM": 64}
278+ })";
279+ map<string, string> socInfos;
280+ map<string, string> aicoreSpec;
281+ map<string, string> intrinsics;
282+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
283+ 
284+ // platform info
285+ fe::PlatFormInfos platformInfo;
286+ platformInfo.Init();
287+ // compile info
288+ optiling::MapIndexCompileInfo compileInfo;
289+ compileInfo.coreNum = 64;
290+ compileInfo.ubSize = 253952;
291+ 
292+ // tilingFunc simulate
293+ auto param = gert::TilingData::CreateCap(4096);
294+ ASSERT_NE(param, nullptr);
295+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
296+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
297+ gert::StorageShape xShape = {{64}, {64}};
298+ gert::StorageShape dataSeqShape = {{641}, {641}};
299+ gert::StorageShape levelIndexShape = {{10}, {10}};
300+ gert::StorageShape yShape = {{1}, {1}};
301+ auto holder = gert::TilingContextFaker()
302+ .NodeIoNum(3, 1)
303+ .IrInstanceNum({1, 1, 1})
304+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
305+ .OutputShapes({&yShape})
306+ .CompileInfo(&compileInfo)
307+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
308+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
309+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
310+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
311+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
312+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
313+ .TilingData(param.get())
314+ .Workspace(wsSize)
315+ .Build();
316+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
317+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
318+ 
319+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
320+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
321+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
322+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
323+ 
324+ // workspaces nullptr return failed
325+ std::cout << "test>> holder.GetContext end" << std::endl;
326+ 
327+ if (tilingFunc == nullptr) {
328+ std::cout << "test>> tilingFunc is invalid" << std::endl;
329+ } else {
330+ std::cout << "test>> tilingFunc is valid" << std::endl;
331+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
332+ }
333+}
334+ 
335+TEST_F(MapIndexTiling, test_tiling_dtype_y_shape_fail_04) {
336+ std::string opType("MapIndex");
337+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
338+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
339+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
340+ 
341+ string compileInfoString = R"({
342+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
343+ "Intrinsic_fix_pipe_l0c2out": false,
344+ "Intrinsic_data_move_l12ub": true,
345+ "Intrinsic_data_move_l0c2ub": true,
346+ "Intrinsic_data_move_out2l1_nd2nz": false,
347+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
348+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
349+ "CORE_NUM": 64}
350+ })";
351+ map<string, string> socInfos;
352+ map<string, string> aicoreSpec;
353+ map<string, string> intrinsics;
354+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
355+ 
356+ // platform info
357+ fe::PlatFormInfos platformInfo;
358+ platformInfo.Init();
359+ // compile info
360+ optiling::MapIndexCompileInfo compileInfo;
361+ compileInfo.coreNum = 64;
362+ compileInfo.ubSize = 253952;
363+ 
364+ // tilingFunc simulate
365+ auto param = gert::TilingData::CreateCap(4096);
366+ ASSERT_NE(param, nullptr);
367+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
368+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
369+ gert::StorageShape xShape = {{64}, {64}};
370+ gert::StorageShape dataSeqShape = {{640}, {640}};
371+ gert::StorageShape levelIndexShape = {{10}, {10}};
372+ gert::StorageShape yShape = {{8}, {8}};
373+ auto holder = gert::TilingContextFaker()
374+ .NodeIoNum(3, 1)
375+ .IrInstanceNum({1, 1, 1})
376+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
377+ .OutputShapes({&yShape})
378+ .CompileInfo(&compileInfo)
379+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
380+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
381+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
382+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
383+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
384+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
385+ .TilingData(param.get())
386+ .Workspace(wsSize)
387+ .Build();
388+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
389+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
390+ 
391+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
392+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
393+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
394+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
395+ 
396+ // workspaces nullptr return failed
397+ std::cout << "test>> holder.GetContext end" << std::endl;
398+ 
399+ if (tilingFunc == nullptr) {
400+ std::cout << "test>> tilingFunc is invalid" << std::endl;
401+ } else {
402+ std::cout << "test>> tilingFunc is valid" << std::endl;
403+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
404+ }
405+}
406+ 
407+TEST_F(MapIndexTiling, test_tiling_dtype_x_shape_fail_05) {
408+ std::string opType("MapIndex");
409+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
410+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
411+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
412+ 
413+ string compileInfoString = R"({
414+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
415+ "Intrinsic_fix_pipe_l0c2out": false,
416+ "Intrinsic_data_move_l12ub": true,
417+ "Intrinsic_data_move_l0c2ub": true,
418+ "Intrinsic_data_move_out2l1_nd2nz": false,
419+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
420+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
421+ "CORE_NUM": 64}
422+ })";
423+ map<string, string> socInfos;
424+ map<string, string> aicoreSpec;
425+ map<string, string> intrinsics;
426+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
427+ 
428+ // platform info
429+ fe::PlatFormInfos platformInfo;
430+ platformInfo.Init();
431+ // compile info
432+ optiling::MapIndexCompileInfo compileInfo;
433+ compileInfo.coreNum = 64;
434+ compileInfo.ubSize = 253952;
435+ 
436+ // tilingFunc simulate
437+ auto param = gert::TilingData::CreateCap(4096);
438+ ASSERT_NE(param, nullptr);
439+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
440+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
441+ gert::StorageShape xShape = {{64,1}, {64,1}};
442+ gert::StorageShape dataSeqShape = {{640}, {640}};
443+ gert::StorageShape levelIndexShape = {{10}, {10}};
444+ gert::StorageShape yShape = {{1}, {1}};
445+ auto holder = gert::TilingContextFaker()
446+ .NodeIoNum(3, 1)
447+ .IrInstanceNum({1, 1, 1})
448+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
449+ .OutputShapes({&yShape})
450+ .CompileInfo(&compileInfo)
451+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
452+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
453+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
454+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
455+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
456+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
457+ .TilingData(param.get())
458+ .Workspace(wsSize)
459+ .Build();
460+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
461+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
462+ 
463+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
464+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
465+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
466+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
467+ 
468+ // workspaces nullptr return failed
469+ std::cout << "test>> holder.GetContext end" << std::endl;
470+ 
471+ if (tilingFunc == nullptr) {
472+ std::cout << "test>> tilingFunc is invalid" << std::endl;
473+ } else {
474+ std::cout << "test>> tilingFunc is valid" << std::endl;
475+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
476+ }
477+}
478+ 
479+TEST_F(MapIndexTiling, test_tiling_dtype_x_shape_fail_06) {
480+ std::string opType("MapIndex");
481+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
482+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
483+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
484+ 
485+ string compileInfoString = R"({
486+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
487+ "Intrinsic_fix_pipe_l0c2out": false,
488+ "Intrinsic_data_move_l12ub": true,
489+ "Intrinsic_data_move_l0c2ub": true,
490+ "Intrinsic_data_move_out2l1_nd2nz": false,
491+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
492+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
493+ "CORE_NUM": 64}
494+ })";
495+ map<string, string> socInfos;
496+ map<string, string> aicoreSpec;
497+ map<string, string> intrinsics;
498+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
499+ 
500+ // platform info
501+ fe::PlatFormInfos platformInfo;
502+ platformInfo.Init();
503+ // compile info
504+ optiling::MapIndexCompileInfo compileInfo;
505+ compileInfo.coreNum = 64;
506+ compileInfo.ubSize = 253952;
507+ 
508+ // tilingFunc simulate
509+ auto param = gert::TilingData::CreateCap(4096);
510+ ASSERT_NE(param, nullptr);
511+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
512+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
513+ gert::StorageShape xShape = {{25000}, {25000}};
514+ gert::StorageShape dataSeqShape = {{640}, {640}};
515+ gert::StorageShape levelIndexShape = {{10}, {10}};
516+ gert::StorageShape yShape = {{1}, {1}};
517+ auto holder = gert::TilingContextFaker()
518+ .NodeIoNum(3, 1)
519+ .IrInstanceNum({1, 1, 1})
520+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
521+ .OutputShapes({&yShape})
522+ .CompileInfo(&compileInfo)
523+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
524+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
525+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
526+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
527+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
528+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
529+ .TilingData(param.get())
530+ .Workspace(wsSize)
531+ .Build();
532+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
533+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
534+ 
535+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
536+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
537+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
538+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
539+ 
540+ // workspaces nullptr return failed
541+ std::cout << "test>> holder.GetContext end" << std::endl;
542+ 
543+ if (tilingFunc == nullptr) {
544+ std::cout << "test>> tilingFunc is invalid" << std::endl;
545+ } else {
546+ std::cout << "test>> tilingFunc is valid" << std::endl;
547+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
548+ }
549+}
550+ 
551+TEST_F(MapIndexTiling, test_tiling_dtype_dataSeq_shape_fail_07) {
552+ std::string opType("MapIndex");
553+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
554+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
555+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
556+ 
557+ string compileInfoString = R"({
558+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
559+ "Intrinsic_fix_pipe_l0c2out": false,
560+ "Intrinsic_data_move_l12ub": true,
561+ "Intrinsic_data_move_l0c2ub": true,
562+ "Intrinsic_data_move_out2l1_nd2nz": false,
563+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
564+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
565+ "CORE_NUM": 64}
566+ })";
567+ map<string, string> socInfos;
568+ map<string, string> aicoreSpec;
569+ map<string, string> intrinsics;
570+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
571+ 
572+ // platform info
573+ fe::PlatFormInfos platformInfo;
574+ platformInfo.Init();
575+ // compile info
576+ optiling::MapIndexCompileInfo compileInfo;
577+ compileInfo.coreNum = 64;
578+ compileInfo.ubSize = 253952;
579+ 
580+ // tilingFunc simulate
581+ auto param = gert::TilingData::CreateCap(4096);
582+ ASSERT_NE(param, nullptr);
583+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
584+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
585+ gert::StorageShape xShape = {{64}, {64}};
586+ gert::StorageShape dataSeqShape = {{19200}, {19200}};
587+ gert::StorageShape levelIndexShape = {{10}, {10}};
588+ gert::StorageShape yShape = {{1}, {1}};
589+ auto holder = gert::TilingContextFaker()
590+ .NodeIoNum(3, 1)
591+ .IrInstanceNum({1, 1, 1})
592+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
593+ .OutputShapes({&yShape})
594+ .CompileInfo(&compileInfo)
595+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
596+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
597+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
598+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
599+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
600+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
601+ .TilingData(param.get())
602+ .Workspace(wsSize)
603+ .Build();
604+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
605+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
606+ 
607+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
608+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
609+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
610+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
611+ 
612+ // workspaces nullptr return failed
613+ std::cout << "test>> holder.GetContext end" << std::endl;
614+ 
615+ if (tilingFunc == nullptr) {
616+ std::cout << "test>> tilingFunc is invalid" << std::endl;
617+ } else {
618+ std::cout << "test>> tilingFunc is valid" << std::endl;
619+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
620+ }
621+}
622+ 
623+TEST_F(MapIndexTiling, test_tiling_dtype_dataSeq_shape_fail_08) {
624+ std::string opType("MapIndex");
625+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
626+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
627+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
628+ 
629+ string compileInfoString = R"({
630+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
631+ "Intrinsic_fix_pipe_l0c2out": false,
632+ "Intrinsic_data_move_l12ub": true,
633+ "Intrinsic_data_move_l0c2ub": true,
634+ "Intrinsic_data_move_out2l1_nd2nz": false,
635+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
636+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
637+ "CORE_NUM": 64}
638+ })";
639+ map<string, string> socInfos;
640+ map<string, string> aicoreSpec;
641+ map<string, string> intrinsics;
642+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
643+ 
644+ // platform info
645+ fe::PlatFormInfos platformInfo;
646+ platformInfo.Init();
647+ // compile info
648+ optiling::MapIndexCompileInfo compileInfo;
649+ compileInfo.coreNum = 64;
650+ compileInfo.ubSize = 253952;
651+ 
652+ // tilingFunc simulate
653+ auto param = gert::TilingData::CreateCap(4096);
654+ ASSERT_NE(param, nullptr);
655+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
656+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
657+ gert::StorageShape xShape = {{64}, {64}};
658+ gert::StorageShape dataSeqShape = {{240, 64}, {240, 64}};
659+ gert::StorageShape levelIndexShape = {{10}, {10}};
660+ gert::StorageShape yShape = {{1}, {1}};
661+ auto holder = gert::TilingContextFaker()
662+ .NodeIoNum(3, 1)
663+ .IrInstanceNum({1, 1, 1})
664+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
665+ .OutputShapes({&yShape})
666+ .CompileInfo(&compileInfo)
667+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
668+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
669+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
670+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
671+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
672+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
673+ .TilingData(param.get())
674+ .Workspace(wsSize)
675+ .Build();
676+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
677+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
678+ 
679+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
680+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
681+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
682+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
683+ 
684+ // workspaces nullptr return failed
685+ std::cout << "test>> holder.GetContext end" << std::endl;
686+ 
687+ if (tilingFunc == nullptr) {
688+ std::cout << "test>> tilingFunc is invalid" << std::endl;
689+ } else {
690+ std::cout << "test>> tilingFunc is valid" << std::endl;
691+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
692+ }
693+}
694+ 
695+TEST_F(MapIndexTiling, test_tiling_dtype_dataSeq_shape_fail_09) {
696+ std::string opType("MapIndex");
697+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
698+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
699+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
700+ 
701+ string compileInfoString = R"({
702+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
703+ "Intrinsic_fix_pipe_l0c2out": false,
704+ "Intrinsic_data_move_l12ub": true,
705+ "Intrinsic_data_move_l0c2ub": true,
706+ "Intrinsic_data_move_out2l1_nd2nz": false,
707+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
708+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
709+ "CORE_NUM": 64}
710+ })";
711+ map<string, string> socInfos;
712+ map<string, string> aicoreSpec;
713+ map<string, string> intrinsics;
714+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
715+ 
716+ // platform info
717+ fe::PlatFormInfos platformInfo;
718+ platformInfo.Init();
719+ // compile info
720+ optiling::MapIndexCompileInfo compileInfo;
721+ compileInfo.coreNum = 64;
722+ compileInfo.ubSize = 253952;
723+ 
724+ // tilingFunc simulate
725+ auto param = gert::TilingData::CreateCap(4096);
726+ ASSERT_NE(param, nullptr);
727+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
728+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
729+ gert::StorageShape xShape = {{64}, {64}};
730+ gert::StorageShape dataSeqShape = {{640}, {640}};
731+ gert::StorageShape levelIndexShape = {{12}, {12}};
732+ gert::StorageShape yShape = {{1}, {1}};
733+ auto holder = gert::TilingContextFaker()
734+ .NodeIoNum(3, 1)
735+ .IrInstanceNum({1, 1, 1})
736+ .InputShapes({&xShape, &dataSeqShape, &levelIndexShape})
737+ .OutputShapes({&yShape})
738+ .CompileInfo(&compileInfo)
739+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
740+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
741+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
742+ .NodeInputTd(2, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
743+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
744+ .NodeAttrs({{"transpose", Ops::NN::AnyValue::CreateFrom<bool>(false)}})
745+ .TilingData(param.get())
746+ .Workspace(wsSize)
747+ .Build();
748+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
749+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
750+ 
751+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
752+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
753+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
754+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
755+ 
756+ // workspaces nullptr return failed
757+ std::cout << "test>> holder.GetContext end" << std::endl;
758+ 
759+ if (tilingFunc == nullptr) {
760+ std::cout << "test>> tilingFunc is invalid" << std::endl;
761+ } else {
762+ std::cout << "test>> tilingFunc is valid" << std::endl;
763+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_FAILED);
764+ }
765+}
Mscripts/kernel/binary_config/ascendc_config.json+1-1
@@ -236,11 +236,11 @@
236 {"name":"DynamicMxQuantWithDualAxis", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},236 {"name":"DynamicMxQuantWithDualAxis", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
237 {"name":"DynamicDualLevelMxQuant", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},237 {"name":"DynamicDualLevelMxQuant", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
238 {"name":"GroupedDynamicBlockQuant", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},238 {"name":"GroupedDynamicBlockQuant", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
239+ {"name":"MapIndex", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
239 {"name":"GeluGradV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},240 {"name":"GeluGradV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
240 {"name":"GeluV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},241 {"name":"GeluV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
241 {"name":"PRelu", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},242 {"name":"PRelu", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
242 {"name":"DynamicQuantV3", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},243 {"name":"DynamicQuantV3", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
243- {"name":"MapIndex", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
244 {"name":"DynamicBlockQuant", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},244 {"name":"DynamicBlockQuant", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
245 {"name":"GroupedDynamicMxQuant", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},245 {"name":"GroupedDynamicMxQuant", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
246 {"name":"SiluGrad", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},246 {"name":"SiluGrad", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},