Star632
236
代码
代码
Issues33
Pull Requests90
Actions
项目讨论
Wiki
分析
Star632
236
已合并
feat: add IndirectLoad fusion support #1463
xiebangrui2025创建于 7月24日
feat: add IndirectLoad fusion support #1463
已合并
xiebangrui2025创建于 7月24日
从xiebangrui2025/graph-autofusionIL合入到cann/graph-autofusiondevelop
共 52 个文件变更+3552-100
Mautofuse/att/base/att_const_values.h+1-0
@@ -205,6 +205,7 @@ inline const std::string kXor = "Xor";
205inline const std::string kVshrs = "Vshrs";205inline const std::string kVshrs = "Vshrs";
206// 下面均为ASCIR定义（已有性能评估）206// 下面均为ASCIR定义（已有性能评估）
207inline const std::string kGather = "Gather";207inline const std::string kGather = "Gather";
208+inline const std::string kIndirectLoad = "IndirectLoad";
208inline const std::string kAbs = "Abs";209inline const std::string kAbs = "Abs";
209inline const std::string kAdd = "Add";210inline const std::string kAdd = "Add";
210inline const std::string kBroadcast = "Broadcast";211inline const std::string kBroadcast = "Broadcast";
Mautofuse/att/gen_model_info/parser/ascend_graph_parser.cpp+12-3
@@ -24,6 +24,8 @@
24#include "ascir_node_param/ascir_node_param.h"24#include "ascir_node_param/ascir_node_param.h"
25#include "base_types_printer.h"25#include "base_types_printer.h"
26#include "common_utils.h"26#include "common_utils.h"
27+#include "indirect_load_utils.h"
28+#include "schedule_result.h"
27#include "specific_params_builder.h"29#include "specific_params_builder.h"
28#include "vector_function_graph_parser.h"30#include "vector_function_graph_parser.h"
29 31 
@@ -853,11 +855,11 @@ void AscendGraphParser::ParserOptionalInfos(const af::AscGraph &graph) const {
853}855}
854 856 
855af::Status AscendGraphParser::CalculateReservedUbSize(const af::AscGraph &graph) {857af::Status AscendGraphParser::CalculateReservedUbSize(const af::AscGraph &graph) {
856- constexpr int32_t kSimtDcacheSize = 32 * 1024;
857 tuning_space_->reserve_ub["ascendc"] = ascgen_utils::CalcReservedTmpBufSizeForAscGraph(graph);858 tuning_space_->reserve_ub["ascendc"] = ascgen_utils::CalcReservedTmpBufSizeForAscGraph(graph);
858 for (const auto &node : graph.GetAllNodes()) {859 for (const auto &node : graph.GetAllNodes()) {
859- if (node->GetType() == kGather) {860+ const int64_t simt_dcache_size = ::ascir::GetDcacheSize(*node);
860- tuning_space_->reserve_ub["simt_dcache"] = kSimtDcacheSize;861+ if (simt_dcache_size > 0) {
862+ tuning_space_->reserve_ub["simt_dcache"] = static_cast<uint32_t>(simt_dcache_size);
861 break;863 break;
862 }864 }
863 }865 }
@@ -875,6 +877,13 @@ af::Status AscendGraphParser::ConvertToTuningSpace(const af::AscGraph &graph) {
875 if (node_iter == graph_sched_info_.end()) {877 if (node_iter == graph_sched_info_.end()) {
876 continue;878 continue;
877 }879 }
880+ const auto indirect_load_behavior = ascgen_utils::indirect_load::GetTemplateBehavior(node);
881+ if (indirect_load_behavior.uses_direct_gm_pipeline || indirect_load_behavior.skips_api_emit) {
882+ GELOGD("[IndirectLoad] Skip tuning-space node info for node[%s], direct_gm[%d], skips_emit[%d].",
883+ node->GetNamePtr(), static_cast<int32_t>(indirect_load_behavior.uses_direct_gm_pipeline),
884+ static_cast<int32_t>(indirect_load_behavior.skips_api_emit));
885+ continue;
886+ }
878 const auto &sched_attrs = node_iter->second;887 const auto &sched_attrs = node_iter->second;
879 GE_ASSERT_SUCCESS(ConvertNodeInfos(node, sched_attrs, graph, use_cache_flag), "Parse node info failed.");888 GE_ASSERT_SUCCESS(ConvertNodeInfos(node, sched_attrs, graph, use_cache_flag), "Parse node info failed.");
880 }889 }
Mautofuse/att/generator/tiling_code_gen_impl.cpp+5-0
@@ -2384,6 +2384,11 @@ af::Status TilingCodeGenImpl::GenSelectBetterTilingBasedOnObjAndUbRatio() {
2384 tiling_func_.AddLine(" tilingCaseImplPtr->GetTilingData(tmp_tiling, tiling_data);");2384 tiling_func_.AddLine(" tilingCaseImplPtr->GetTilingData(tmp_tiling, tiling_data);");
2385 tiling_func_.AddLine(" }");2385 tiling_func_.AddLine(" }");
2386 tiling_func_.AddLine(" }");2386 tiling_func_.AddLine(" }");
2387+ } else {
2388+ tiling_func_.AddLine(GenCallUpdateBetterTiling(is_uniq_group_));
2389+ tiling_func_.AddLine(" sub_case_flag = is_sub_case;");
2390+ tiling_func_.AddLine(" obj = cur_obj;");
2391+ tiling_func_.AddLine(" return true;");
2387 }2392 }
2388 return af::SUCCESS;2393 return af::SUCCESS;
2389}2394}
Mautofuse/codegen/codegen_graph_check.cpp+4-0
@@ -13,6 +13,7 @@
13#include "ascir_ops.h"13#include "ascir_ops.h"
14#include "common/platform_context.h"14#include "common/platform_context.h"
15#include "common_utils.h"15#include "common_utils.h"
16+#include "indirect_load_utils.h"
16 17 
17namespace codegen {18namespace codegen {
18 19 
@@ -214,6 +215,9 @@ Status IsRepeatStrideValid(const ascir::ImplGraph &graph) {
214 215 
215Status IsGraphNodeValid(const ascir::ImplGraph &graph) {216Status IsGraphNodeValid(const ascir::ImplGraph &graph) {
216 for (const auto &node : graph.GetAllNodes()) {217 for (const auto &node : graph.GetAllNodes()) {
218+ if (ascgen_utils::indirect_load::GetTemplateBehavior(node).skips_api_emit) {
219+ continue;
220+ }
217 auto impl = ascgen_utils::GetAscIrCodegenImpl(node->GetType());221 auto impl = ascgen_utils::GetAscIrCodegenImpl(node->GetType());
218 GE_ASSERT_NOTNULL(impl, "GetAscIrCodegenImpl of node %s[%s] is null", node->GetTypePtr(), node->GetNamePtr());222 GE_ASSERT_NOTNULL(impl, "GetAscIrCodegenImpl of node %s[%s] is null", node->GetTypePtr(), node->GetNamePtr());
219 GE_ASSERT_TRUE(impl->IsNodeValid(*node), "Node %s[%s] is invalid", node->GetTypePtr(), node->GetNamePtr());223 GE_ASSERT_TRUE(impl->IsNodeValid(*node), "Node %s[%s] is invalid", node->GetTypePtr(), node->GetNamePtr());
Mautofuse/codegen/codegen_kernel.cpp+70-43
@@ -25,7 +25,10 @@
25#include "backend/backend_spec.h"25#include "backend/backend_spec.h"
26#include "graph/symbolizer/symbolic_utils.h"26#include "graph/symbolizer/symbolic_utils.h"
27#include "ascendc_api_registry.h"27#include "ascendc_api_registry.h"
28+#include "indirect_load_utils.h"
28#include "optimize/platform/platform_factory.h"29#include "optimize/platform/platform_factory.h"
30+#include "optimize/schedule_utils.h"
31+#include "v35/codegen/simt_scalar_call/simt_scalar_emitter.h"
29#include "common/platform_context.h"32#include "common/platform_context.h"
30#include "codegen_graph_check.h"33#include "codegen_graph_check.h"
31 34 
@@ -40,6 +43,43 @@ constexpr uint32_t kFuncIdBegin = 20000000U;
40constexpr const char kInputTensorDescName[] = "input_tensor_desc";43constexpr const char kInputTensorDescName[] = "input_tensor_desc";
41constexpr const char kOutputTensorDescName[] = "output_tensor_desc";44constexpr const char kOutputTensorDescName[] = "output_tensor_desc";
42 45 
46+struct QueCollection {
47+ std::unordered_map<ascir::QueId, af::Position> que_id_to_src_position;
48+ std::set<ascir::QueId> need_bind_que_id;
49+ std::map<ascir::QueId, bool> direct_gm_ques;
50+};
51+ 
52+Status AllocateQues(TPipe &tpipe, const QueCollection &collection) {
53+ for (const auto &iter : collection.que_id_to_src_position) {
54+ if (tpipe.ques.count(iter.first) > 0UL) {
55+ continue;
56+ }
57+ std::string position;
58+ GE_CHK_STATUS_RET(PositionValue(iter.second, position), "Codegen get position value failed");
59+ const bool need_que_bind = collection.need_bind_que_id.count(iter.first) > 0UL;
60+ if (need_que_bind) {
61+ std::string dst_position;
62+ GE_CHK_STATUS_RET(PositionValue(af::Position::kPositionVecOut, dst_position),
63+ "Codegen get position value failed");
64+ const auto new_que = tpipe.ques.emplace(iter.first, TQue{iter.first, iter.second, position, dst_position});
65+ GE_CHK_BOOL_RET_STATUS(new_que.second, af::FAILED, "Codegen emplace que [%ld] failed", iter.first);
66+ } else {
67+ const auto new_que = tpipe.ques.emplace(iter.first, TQue{iter.first, iter.second, position});
68+ GE_CHK_BOOL_RET_STATUS(new_que.second, af::FAILED, "Codegen emplace que [%ld] failed", iter.first);
69+ }
70+ const auto skip_iter = collection.direct_gm_ques.find(iter.first);
71+ if (skip_iter != collection.direct_gm_ques.end()) {
72+ tpipe.ques.at(iter.first).skip_init_for_simt_direct_gm = skip_iter->second;
73+ }
74+ }
75+ for (auto &[id, que] : tpipe.ques) {
76+ if (id != tpipe.cube_output_que_id) {
77+ que.is_cv_ub_fusion = (tpipe.cv_fusion_type == ascir::CubeTemplateType::kUBFuse);
78+ }
79+ }
80+ return af::SUCCESS;
81+}
82+ 
43std::string GetTensorName(const ascir::TensorAttr &tensor) {83std::string GetTensorName(const ascir::TensorAttr &tensor) {
44 const auto node = tensor.anchor.GetOwnerNodeBarePtr();84 const auto node = tensor.anchor.GetOwnerNodeBarePtr();
45 if (node != nullptr && node->GetType() == ScalarData::Type) {85 if (node != nullptr && node->GetType() == ScalarData::Type) {
@@ -1288,7 +1328,8 @@ Status Kernel::ParseUbScalarOptimizationInfo(const ascir::NodeView &node, Tensor
1288Status Kernel::JudgeIsLoadLinkStoreAndVec(const ascir::NodeView &node, Tensor &t, ascir::TensorId id) const {1328Status Kernel::JudgeIsLoadLinkStoreAndVec(const ascir::NodeView &node, Tensor &t, ascir::TensorId id) const {
1289 // todo: 解决load多引用场景，被store, vec 同时引用的缺少mte3到mte2的同步的问题,1329 // todo: 解决load多引用场景，被store, vec 同时引用的缺少mte3到mte2的同步的问题,
1290 // 临时方案，从这里解析下是否该场景1330 // 临时方案，从这里解析下是否该场景
1291- if ((node->attr.api.compute_type == ascir::ComputeType::kComputeLoad) && (!IsOps<Gather>(node))) {1331+ if ((node->attr.api.compute_type == ascir::ComputeType::kComputeLoad) &&
1332+ (!optimize::ScheduleUtils::IsGatherLikeLoad(std::dynamic_pointer_cast<af::AscNode>(node)))) {
1292 bool link_to_store = false;1333 bool link_to_store = false;
1293 bool link_to_vec = false;1334 bool link_to_vec = false;
1294 for (auto &out : node->outputs()) {1335 for (auto &out : node->outputs()) {
@@ -1416,6 +1457,10 @@ Status TPipe::TensorAlloc(const Tensor &tensor, std::string &result) const {
1416}1457}
1417 1458 
1418Status TPipe::InitTQueBuffers(const TQue &que, std::string &result) const {1459Status TPipe::InitTQueBuffers(const TQue &que, std::string &result) const {
1460+ if (que.skip_init_for_simt_direct_gm) {
1461+ result.clear();
1462+ return af::SUCCESS;
1463+ }
1419 stringstream ss;1464 stringstream ss;
1420 std::string blk_align;1465 std::string blk_align;
1421 GE_CHK_STATUS_RET(KernelUtils::BlkAlign(ge::DT_UINT8, blk_align), "Codegen blk align failed");1466 GE_CHK_STATUS_RET(KernelUtils::BlkAlign(ge::DT_UINT8, blk_align), "Codegen blk align failed");
@@ -1465,6 +1510,10 @@ std::string TPipe::TensorSizeCalc() const {
1465 for (const auto &pair : this->tensors) {1510 for (const auto &pair : this->tensors) {
1466 const auto &t = pair.second;1511 const auto &t = pair.second;
1467 if (t.alloc_type == af::AllocType::kAllocTypeQueue) {1512 if (t.alloc_type == af::AllocType::kAllocTypeQueue) {
1513+ const TQue *que = GetQue(t.que_id);
1514+ if (que != nullptr && que->skip_init_for_simt_direct_gm) {
1515+ continue;
1516+ }
1468 ss << t.size.DefineConst(this->tiler.TensorVectorizedSize(t)) << std::endl;1517 ss << t.size.DefineConst(this->tiler.TensorVectorizedSize(t)) << std::endl;
1469 ss << t.que_buf_num.DefineConst(to_string(t.que_buf_num_value)) << std::endl;1518 ss << t.que_buf_num.DefineConst(to_string(t.que_buf_num_value)) << std::endl;
1470 } else if (t.alloc_type == af::AllocType::kAllocTypeBuffer) {1519 } else if (t.alloc_type == af::AllocType::kAllocTypeBuffer) {
@@ -1547,7 +1596,7 @@ Status TPipe::LocalTQueAlloc(std::string &result) const {
1547 stringstream ss;1596 stringstream ss;
1548 1597 
1549 for (auto &[id, que] : this->ques) {1598 for (auto &[id, que] : this->ques) {
1550- if (id == this->cube_output_que_id) {1599+ if (id == this->cube_output_que_id || que.skip_init_for_simt_direct_gm) {
1551 continue;1600 continue;
1552 }1601 }
1553 stringstream tensor_size_max;1602 stringstream tensor_size_max;
@@ -1560,10 +1609,7 @@ Status TPipe::LocalTQueAlloc(std::string &result) const {
1560 1609 
1561 for (auto mid : que.merge_scopes) {1610 for (auto mid : que.merge_scopes) {
1562 auto merge_scope = this->merge_scopes.find(mid);1611 auto merge_scope = this->merge_scopes.find(mid);
1563- if (merge_scope == this->merge_scopes.end()) {1612+ GE_ASSERT_TRUE(merge_scope != this->merge_scopes.end(), "Codegen merge scope not found:%ld", mid);
1564- GELOGE(af::FAILED, "Codegen merge scope not found:%ld", mid);
1565- return af::FAILED;
1566- }
1567 1613 
1568 if (is_first) {1614 if (is_first) {
1569 is_first = false;1615 is_first = false;
@@ -1579,10 +1625,7 @@ Status TPipe::LocalTQueAlloc(std::string &result) const {
1579 uint32_t tensor_buf_num_max_val = 0;1625 uint32_t tensor_buf_num_max_val = 0;
1580 for (auto tid : que.not_merge_tensors) {1626 for (auto tid : que.not_merge_tensors) {
1581 auto tensor = this->tensors.find(tid);1627 auto tensor = this->tensors.find(tid);
1582- if (tensor == this->tensors.end()) {1628+ GE_ASSERT_TRUE(tensor != this->tensors.end(), "Codegen tensor not found:%ld", tid);
1583- GELOGE(af::FAILED, "Codegen tensor not found:%ld", tid);
1584- return af::FAILED;
1585- }
1586 1629 
1587 if (is_first) {1630 if (is_first) {
1588 is_first = false;1631 is_first = false;
@@ -1815,8 +1858,7 @@ std::string TPipe::SyncMte2ToMte3(const Tensor in_tensor) const {
1815}1858}
1816 1859 
1817Status TPipe::CollectQues(const ascir::ImplGraph &graph) {1860Status TPipe::CollectQues(const ascir::ImplGraph &graph) {
1818- std::unordered_map<ascir::QueId, af::Position> que_id_to_src_position;1861+ QueCollection collection;
1819- std::set<ascir::QueId> need_bind_que_id;
1820 for (auto node : graph.GetAllNodes()) {1862 for (auto node : graph.GetAllNodes()) {
1821 if (node->attr.api.type == ge::ApiType::kAPITypeBuffer) {1863 if (node->attr.api.type == ge::ApiType::kAPITypeBuffer) {
1822 continue;1864 continue;
@@ -1826,8 +1868,13 @@ Status TPipe::CollectQues(const ascir::ImplGraph &graph) {
1826 continue;1868 continue;
1827 }1869 }
1828 const int64_t tensor_que_id = out_tensor->attr.que.id;1870 const int64_t tensor_que_id = out_tensor->attr.que.id;
1871+ const bool uses_direct_gm = ascgen_utils::indirect_load::GetTemplateBehavior(node).uses_direct_gm_pipeline;
1872+ auto [direct_gm_iter, inserted] = collection.direct_gm_ques.emplace(tensor_que_id, uses_direct_gm);
1873+ if (!inserted) {
1874+ direct_gm_iter->second = direct_gm_iter->second && uses_direct_gm;
1875+ }
1829 if (out_tensor->attr.mem.position == af::Position::kPositionVecIn) {1876 if (out_tensor->attr.mem.position == af::Position::kPositionVecIn) {
1830- que_id_to_src_position.emplace(tensor_que_id, out_tensor->attr.mem.position);1877+ collection.que_id_to_src_position.emplace(tensor_que_id, out_tensor->attr.mem.position);
1831 1878 
1832 std::set<std::string> peer_node_types;1879 std::set<std::string> peer_node_types;
1833 for (const auto &peer_in_anchor : out_tensor->anchor.GetPeerInDataAnchorsPtr()) {1880 for (const auto &peer_in_anchor : out_tensor->anchor.GetPeerInDataAnchorsPtr()) {
@@ -1836,40 +1883,14 @@ Status TPipe::CollectQues(const ascir::ImplGraph &graph) {
1836 }1883 }
1837 }1884 }
1838 if ((peer_node_types.size() == 1U) && *peer_node_types.begin() == Store::Type) {1885 if ((peer_node_types.size() == 1U) && *peer_node_types.begin() == Store::Type) {
1839- need_bind_que_id.emplace(tensor_que_id);1886+ collection.need_bind_que_id.emplace(tensor_que_id);
1840 }1887 }
1841 } else if (out_tensor->attr.mem.position == af::Position::kPositionVecOut) {1888 } else if (out_tensor->attr.mem.position == af::Position::kPositionVecOut) {
1842- que_id_to_src_position.emplace(tensor_que_id, out_tensor->attr.mem.position);1889+ collection.que_id_to_src_position.emplace(tensor_que_id, out_tensor->attr.mem.position);
1843 }1890 }
1844 }1891 }
1845 }1892 }
1846- 1893+ return AllocateQues(*this, collection);
1847- // Allocate
1848- for (const auto &iter : que_id_to_src_position) {
1849- if (this->ques.count(iter.first) > 0UL) {
1850- continue;
1851- }
1852- std::string position;
1853- GE_CHK_STATUS_RET(PositionValue(iter.second, position), "Codegen get position value failed");
1854- // 如果qid被非load->store的load复用，则不能使用TQueBind
1855- bool need_que_bind = need_bind_que_id.count(iter.first) > 0UL;
1856- if (need_que_bind) {
1857- std::string dst_position;
1858- GE_CHK_STATUS_RET(PositionValue(af::Position::kPositionVecOut, dst_position),
1859- "Codegen get position value failed");
1860- auto new_que = this->ques.emplace(iter.first, TQue{iter.first, iter.second, position, dst_position});
1861- GE_CHK_BOOL_RET_STATUS(new_que.second, af::FAILED, "Codegen emplace que [%ld] failed", iter.first);
1862- } else {
1863- auto new_que = this->ques.emplace(iter.first, TQue{iter.first, iter.second, position});
1864- GE_CHK_BOOL_RET_STATUS(new_que.second, af::FAILED, "Codegen emplace que [%ld] failed", iter.first);
1865- }
1866- }
1867- for (auto &[id, que] : this->ques) {
1868- if (id != this->cube_output_que_id) {
1869- que.is_cv_ub_fusion = (this->cv_fusion_type == ascir::CubeTemplateType::kUBFuse);
1870- }
1871- }
1872- return af::SUCCESS;
1873}1894}
1874 1895 
1875void TPipe::SetUsingAttCalcQBTSizeConfig(bool using_att_calc_qbt_size) {1896void TPipe::SetUsingAttCalcQBTSizeConfig(bool using_att_calc_qbt_size) {
@@ -2155,7 +2176,8 @@ Status Kernel::ParseGraph(const ascir::ImplGraph &graph, const ascir::FusedSched
2155 }2176 }
2156 continue;2177 continue;
2157 }2178 }
2158- has_gather = (has_gather || IsOps<Gather>(node));2179+ has_gather =
2180+ (has_gather || optimize::ScheduleUtils::IsGatherLikeLoad(std::dynamic_pointer_cast<af::AscNode>(node)));
2159 }2181 }
2160 for (const auto &pair : kernel_outputs) {2182 for (const auto &pair : kernel_outputs) {
2161 kernel.outputs.emplace_back(GM_ADDR(GenValidName(pair.second.first)));2183 kernel.outputs.emplace_back(GM_ADDR(GenValidName(pair.second.first)));
@@ -2192,6 +2214,11 @@ Status Kernel::ParseGraph(const ascir::ImplGraph &graph, const ascir::FusedSched
2192 if (IsOps<Output>(node) || IsOps<Data>(node) || IsOps<ScalarData>(node)) {2214 if (IsOps<Output>(node) || IsOps<Data>(node) || IsOps<ScalarData>(node)) {
2193 continue;2215 continue;
2194 }2216 }
2217+ const auto indirect_load_behavior = ascgen_utils::indirect_load::GetTemplateBehavior(node);
2218+ if (indirect_load_behavior.skips_api_emit && indirect_load_behavior.uses_direct_gm_pipeline &&
2219+ !IsOps<Store>(node)) {
2220+ continue;
2221+ }
2195 2222 
2196 auto desc = node->GetOpDesc();2223 auto desc = node->GetOpDesc();
2197 for (auto output : node->outputs()) {2224 for (auto output : node->outputs()) {
Mautofuse/codegen/codegen_kernel.h+1-0
@@ -203,6 +203,7 @@ class TQue : public Variable {
203 203 
204 Variable buf;204 Variable buf;
205 bool is_cv_ub_fusion{false};205 bool is_cv_ub_fusion{false};
206+ bool skip_init_for_simt_direct_gm{false};
206 207 
207 TQue(ascir::QueId que_id, ascir::Position pos, std::string &position_name);208 TQue(ascir::QueId que_id, ascir::Position pos, std::string &position_name);
208 TQue(ascir::QueId que_id, ascir::Position src_position, const std::string &src_position_name,209 TQue(ascir::QueId que_id, ascir::Position src_position, const std::string &src_position_name,
Mautofuse/codegen/codegen_kernel_loop.cpp+97-11
@@ -18,6 +18,8 @@
18#include "optimize/platform/platform_factory.h"18#include "optimize/platform/platform_factory.h"
19#include "api_call/utils/api_call_factory.h"19#include "api_call/utils/api_call_factory.h"
20#include "graph/ascendc_ir/utils/asc_tensor_utils.h"20#include "graph/ascendc_ir/utils/asc_tensor_utils.h"
21+#include "indirect_load_utils.h"
22+#include "v35/codegen/simt_scalar_call/simt_scalar_emitter.h"
21 23 
22using namespace std;24using namespace std;
23using namespace af::ops;25using namespace af::ops;
@@ -32,6 +34,59 @@ const std::string kEnCacheOriginBroadcastAxis = "enable_cache_origin_brc_axis";
32const std::string kEnCacheFusedBroadcastAxis = "enable_cache_fused_brc_axis";34const std::string kEnCacheFusedBroadcastAxis = "enable_cache_fused_brc_axis";
33const std::string kEnCacheA = "dis_enable_cache_a";35const std::string kEnCacheA = "dis_enable_cache_a";
34const std::string kEnCacheR = "dis_enable_cache_r";36const std::string kEnCacheR = "dis_enable_cache_r";
37+ 
38+ApiCall *FindIndirectLoadOpCall(Loop *loop, bool require_direct_gm) {
39+ if (loop == nullptr) {
40+ return nullptr;
41+ }
42+ for (const auto &body : loop->bodys) {
43+ if (body.type == LoopType::CALL) {
44+ if (body.call == nullptr || !af::ops::IsOps<af::ascir_op::IndirectLoad>(body.call->node)) {
45+ continue;
46+ }
47+ const auto behavior =
48+ ascgen_utils::indirect_load::GetTemplateBehavior(std::dynamic_pointer_cast<af::AscNode>(body.call->node));
49+ if (!require_direct_gm || (behavior.uses_direct_gm_pipeline && behavior.skips_ub_lifecycle)) {
50+ return body.call;
51+ }
52+ } else if (ApiCall *call = FindIndirectLoadOpCall(body.loop, require_direct_gm); call != nullptr) {
53+ return call;
54+ }
55+ }
56+ return nullptr;
57+}
58+ 
59+Status AddSkippedApiEmitProcessCall(const ascir::NodeView &node, Loop *current_loop,
60+ const std::vector<ascir::AxisId> &current_axis,
61+ std::map<ascir::TensorId, ApiCall *> &tensor_calls) {
62+ auto call = CreateApiCallObject(node);
63+ GE_ASSERT_NOTNULL(call, "Create api call object failed, ascir type:%s", node->GetTypePtr());
64+ current_loop->AddCall(call);
65+ GE_CHK_STATUS_RET(call->Init(node), "ApiCall Init failed, ascir type:%s", node->GetTypePtr());
66+ call->skip_api_emit = true;
67+ call->exec_condition = node->attr.sched.exec_condition;
68+ call->axis = current_loop->axis_id;
69+ call->depth = current_axis.size();
70+ if (IsOps<Store>(node)) {
71+ for (auto out : node->outputs()) {
72+ tensor_calls.insert({out->attr.mem.tensor_id, call});
73+ }
74+ return af::SUCCESS;
75+ }
76+ GE_ASSERT_TRUE(node->inputs.Size() == 1UL,
77+ "Skipped api emit process node only supports single-input tensor mapping, node[%s].",
78+ node->GetNamePtr());
79+ auto input = node->inputs()[0];
80+ GE_ASSERT_NOTNULL(input, "Skipped api emit process node[%s] input is null", node->GetNamePtr());
81+ auto in_call = tensor_calls.find(input->attr.mem.tensor_id);
82+ GE_CHK_BOOL_RET_STATUS(in_call != tensor_calls.end(), af::FAILED,
83+ "Codegen node[%s] no API call found for process input tensor id[%ld]", node->GetNamePtr(),
84+ input->attr.mem.tensor_id);
85+ for (auto out : node->outputs()) {
86+ tensor_calls.insert({out->attr.mem.tensor_id, in_call->second});
87+ }
88+ return af::SUCCESS;
89+}
35} // namespace90} // namespace
36 91 
37Loop::Loop(const ascir::AxisId axis) : axis_id(axis), parent(nullptr) {}92Loop::Loop(const ascir::AxisId axis) : axis_id(axis), parent(nullptr) {}
@@ -250,6 +305,11 @@ Status Loop::ConstructFromNodes(ascir::NodeViewVisitorConst nodes, const Tiler &
250 TraverseGraphForReduceNodes(nodes, current_loop->is_graph_has_reduce_node, current_loop->is_ar);305 TraverseGraphForReduceNodes(nodes, current_loop->is_graph_has_reduce_node, current_loop->is_ar);
251 auto lifecycle_edge = GetLifecycleEdge(nodes, tpipe);306 auto lifecycle_edge = GetLifecycleEdge(nodes, tpipe);
252 for (auto node : nodes) {307 for (auto node : nodes) {
308+ if (IsSkippedApiEmitProcessNode(node)) {
309+ GE_CHK_STATUS_RET(AddSkippedApiEmitProcessCall(node, current_loop, current_axis, tensor_calls));
310+ continue;
311+ }
312+ 
253 // Loop enter or create313 // Loop enter or create
254 GELOGI("node:%s, ComputeUnit:%u\r\n", node->GetNamePtr(), static_cast<uint32_t>(node->attr.api.unit));314 GELOGI("node:%s, ComputeUnit:%u\r\n", node->GetNamePtr(), static_cast<uint32_t>(node->attr.api.unit));
255 if (node->attr.api.unit != af::ComputeUnit::kUnitNone) {315 if (node->attr.api.unit != af::ComputeUnit::kUnitNone) {
@@ -454,21 +514,30 @@ Status Loop::GenerateBody(const Tiler &tiler, const TPipe &tpipe, std::vector<as
454 }514 }
455 if (body.type == LoopType::LOOP) {515 if (body.type == LoopType::LOOP) {
456 for (auto call : target_calls) {516 for (auto call : target_calls) {
457- GE_CHK_STATUS_RET(call->AllocOutputs(tpipe, ss), "Codegen alloc outputs failed");517+ if (!ascgen_utils::indirect_load::GetTemplateBehavior(std::dynamic_pointer_cast<af::AscNode>(call->node))
518+ .skips_ub_lifecycle) {
519+ GE_CHK_STATUS_RET(call->AllocOutputs(tpipe, ss), "Codegen alloc outputs failed");
520+ }
458 used_calls.insert(call);521 used_calls.insert(call);
459 }522 }
460 body.loop->compute_stage = this->compute_stage;523 body.loop->compute_stage = this->compute_stage;
461 GE_CHK_STATUS_RET(body.loop->GenerateLoop(tiler, tpipe, current_axis, ss), "Generate loop for body failed");524 GE_CHK_STATUS_RET(body.loop->GenerateLoop(tiler, tpipe, current_axis, ss), "Generate loop for body failed");
462 for (auto call : target_calls) {525 for (auto call : target_calls) {
463- GE_CHK_BOOL_RET_STATUS(call->SyncOutputs(tpipe, ss), af::FAILED, "Func SyncOutputs return false");526+ if (!ascgen_utils::indirect_load::GetTemplateBehavior(std::dynamic_pointer_cast<af::AscNode>(call->node))
527+ .skips_ub_lifecycle) {
528+ GE_CHK_BOOL_RET_STATUS(call->SyncOutputs(tpipe, ss), af::FAILED, "Func SyncOutputs return false");
529+ }
464 }530 }
465 used_calls.clear();531 used_calls.clear();
466 } else {532 } else {
467- if (body.call->unit == af::ComputeUnit::kUnitNone) {533+ if (body.call->unit == af::ComputeUnit::kUnitNone || body.call->skip_api_emit) {
468 continue;534 continue;
469 }535 }
536+ const bool skips_ub_lifecycle =
537+ ascgen_utils::indirect_load::GetTemplateBehavior(std::dynamic_pointer_cast<af::AscNode>(body.call->node))
538+ .skips_ub_lifecycle;
470 GE_CHK_BOOL_RET_STATUS(body.call->WaitInputs(tpipe, ss), af::FAILED, "Func WaitInputs return false");539 GE_CHK_BOOL_RET_STATUS(body.call->WaitInputs(tpipe, ss), af::FAILED, "Func WaitInputs return false");
471- if (!IsFindInUsedCalls(body.call)) {540+ if (!IsFindInUsedCalls(body.call) && !skips_ub_lifecycle) {
472 GE_CHK_STATUS_RET(body.call->AllocOutputs(tpipe, ss), "Codegen alloc outputs failed");541 GE_CHK_STATUS_RET(body.call->AllocOutputs(tpipe, ss), "Codegen alloc outputs failed");
473 }542 }
474 std::string call;543 std::string call;
@@ -501,12 +570,14 @@ Status Loop::GenerateBody(const Tiler &tiler, const TPipe &tpipe, std::vector<as
501 }570 }
502 }571 }
503 572 
504- if (!IsFindInUsedCalls(body.call)) {573+ if (!skips_ub_lifecycle) {
505- GE_CHK_BOOL_RET_STATUS(body.call->SyncOutputs(tpipe, ss), af::FAILED, "Func SyncOutputs return false");574+ if (!IsFindInUsedCalls(body.call)) {
575+ GE_CHK_BOOL_RET_STATUS(body.call->SyncOutputs(tpipe, ss), af::FAILED, "Func SyncOutputs return false");
576+ }
577+ GE_CHK_BOOL_RET_STATUS(body.call->FreeInputs(tpipe, ss), af::FAILED, "Func FreeInputs return false");
578+ GE_CHK_BOOL_RET_STATUS(body.call->FreeUnusedOutputs(tpipe, ss), af::FAILED,
579+ "Func FreeUnusedOutputs return false");
506 }580 }
507- GE_CHK_BOOL_RET_STATUS(body.call->FreeInputs(tpipe, ss), af::FAILED, "Func FreeInputs return false");
508- GE_CHK_BOOL_RET_STATUS(body.call->FreeUnusedOutputs(tpipe, ss), af::FAILED,
509- "Func FreeUnusedOutputs return false");
510 ss << std::endl;581 ss << std::endl;
511 }582 }
512 }583 }
@@ -627,13 +698,27 @@ Status Loop::GenerateLoop(const Tiler &tiler, const TPipe &tpipe, std::vector<as
627 if (axis.type == Axis::Type::kAxisTypeBlockInner) {698 if (axis.type == Axis::Type::kAxisTypeBlockInner) {
628 auto peer = tiler.GetAxis(axis.split_pair_other_id);699 auto peer = tiler.GetAxis(axis.split_pair_other_id);
629 ss << "int32_t block_dim_offset = " << peer.Str() << " * " << tiler.Size(axis.size) << ";" << std::endl;700 ss << "int32_t block_dim_offset = " << peer.Str() << " * " << tiler.Size(axis.size) << ";" << std::endl;
701+ ApiCall *call = FindIndirectLoadOpCall(this, true);
702+ if (call != nullptr) {
703+ std::string call_code;
704+ GE_CHK_STATUS_RET(call->Generate(tpipe, current_axis, call_code), "Codegen generate hoisted call failed");
705+ ss << call_code << std::endl;
706+ call->skip_api_emit = true;
707+ current_axis.pop_back();
708+ return af::SUCCESS;
709+ }
710+ }
711+ if (axis.type == Axis::Type::kAxisTypeTileInner && FindIndirectLoadOpCall(this, false) != nullptr) {
712+ ss << tiler.GenInnerLoopSizeAndActualSize(axis.id, axis.split_pair_other_id);
630 }713 }
631 if (tpipe.cv_fusion_type == ascir::CubeTemplateType::kUBFuse && axis.type == Axis::Type::kAxisTypeTileOuter) {714 if (tpipe.cv_fusion_type == ascir::CubeTemplateType::kUBFuse && axis.type == Axis::Type::kAxisTypeTileOuter) {
632 ss << axis.loop_size.AsArg() << " = 1;" << std::endl;715 ss << axis.loop_size.AsArg() << " = 1;" << std::endl;
633 }716 }
634 ss << "for (" << axis.AsArg() << " = 0; " << axis << " < " << axis.loop_size.Str() << "; " << axis << "++) "717 ss << "for (" << axis.AsArg() << " = 0; " << axis << " < " << axis.loop_size.Str() << "; " << axis << "++) "
635 << "{" << std::endl;718 << "{" << std::endl;
636- if (tpipe.cv_fusion_type != ascir::CubeTemplateType::kUBFuse) {719+ const bool skip_calc_from_axis_for_indirect_load_simt =
720+ axis.type == Axis::Type::kAxisTypeTileInner && FindIndirectLoadOpCall(this, true) != nullptr;
721+ if (tpipe.cv_fusion_type != ascir::CubeTemplateType::kUBFuse && !skip_calc_from_axis_for_indirect_load_simt) {
637 ss << tiler.CalcFromAxis(axis.id);722 ss << tiler.CalcFromAxis(axis.id);
638 }723 }
639 GenerateEnCacheCondition(tiler, tpipe, axis, ss);724 GenerateEnCacheCondition(tiler, tpipe, axis, ss);
@@ -1187,7 +1272,8 @@ bool ApiCall::WaitInputVector(const TPipe &tpipe, const ApiTensor *in, const Ten
1187bool ApiCall::WaitInputMte(const TPipe &tpipe, const ApiTensor *in, const Tensor &t, std::stringstream &ss) const {1272bool ApiCall::WaitInputMte(const TPipe &tpipe, const ApiTensor *in, const Tensor &t, std::stringstream &ss) const {
1188 // 1. load->store 2. load->store store 3. load->vec store store1273 // 1. load->store 2. load->store store 3. load->vec store store
1189 if (this->type == Store::Type &&1274 if (this->type == Store::Type &&
1190- ((in->write->compute_type == ascir::ComputeType::kComputeLoad) && (in->write->type != Gather::Type)) &&1275+ ((in->write->compute_type == ascir::ComputeType::kComputeLoad) && (in->write->type != Gather::Type) &&
1276+ (in->write->type != IndirectLoad::Type)) &&
1191 IsUnitFirstRead(*this, *in)) {1277 IsUnitFirstRead(*this, *in)) {
1192 ss << tpipe.SyncMte2ToMte3(t) << std::endl;1278 ss << tpipe.SyncMte2ToMte3(t) << std::endl;
1193 }1279 }
Mautofuse/codegen/codegen_kernel_loop.h+1-0
@@ -136,6 +136,7 @@ class ApiCall {
136 std::vector<const ApiTensor *> inputs;136 std::vector<const ApiTensor *> inputs;
137 bool enable_cache{false};137 bool enable_cache{false};
138 bool is_input_tbuf_contiguous = false;138 bool is_input_tbuf_contiguous = false;
139+ bool skip_api_emit{false};
139 std::string enable_cache_with_condition;140 std::string enable_cache_with_condition;
140 // 用于标记Call节点执行状态141 // 用于标记Call节点执行状态
141 // broadcast cache场景：在Call节点外生成控制条件142 // broadcast cache场景：在Call节点外生成控制条件
Mautofuse/common/ascgraph_info_complete.cpp+5-0
@@ -16,11 +16,14 @@
16#include "graph/symbolizer/symbolic.h"16#include "graph/symbolizer/symbolic.h"
17#include "graph/attribute_group/attr_group_shape_env.h"17#include "graph/attribute_group/attr_group_shape_env.h"
18#include "ascir_ops_utils.h"18#include "ascir_ops_utils.h"
19+#include "schedule_result.h"
19 20 
20using namespace af::ascir_op;21using namespace af::ascir_op;
21 22 
22namespace optimize {23namespace optimize {
23namespace {24namespace {
25+constexpr int64_t kGatherSimtDcacheSize = 32 * 1024;
26+ 
24static Status GetNodeIrAttrOffset(const af::NodePtr &node, af::Expression &offset) {27static Status GetNodeIrAttrOffset(const af::NodePtr &node, af::Expression &offset) {
25 auto asc_node = std::dynamic_pointer_cast<af::AscNode>(node);28 auto asc_node = std::dynamic_pointer_cast<af::AscNode>(node);
26 GE_ASSERT_NOTNULL(asc_node);29 GE_ASSERT_NOTNULL(asc_node);
@@ -76,6 +79,7 @@ void CompleteSplitApiInfo(af::AscNodePtr &node) {
76void CompleteGatherApiInfo(af::AscNodePtr &node) {79void CompleteGatherApiInfo(af::AscNodePtr &node) {
77 node->attr.api.type = af::ApiType::kAPITypeCompute;80 node->attr.api.type = af::ApiType::kAPITypeCompute;
78 node->attr.api.unit = af::ComputeUnit::kUnitMTE2;81 node->attr.api.unit = af::ComputeUnit::kUnitMTE2;
82+ (void)::ascir::SetDcacheSize(node, kGatherSimtDcacheSize);
79}83}
80 84 
81void CompleteCubeApiInfo(af::AscNodePtr &node) {85void CompleteCubeApiInfo(af::AscNodePtr &node) {
@@ -154,6 +158,7 @@ static const std::map<std::string, af::ComputeType> kOpTypeToComputeType = {
154 {Sigmoid::Type, af::ComputeType::kComputeElewise},158 {Sigmoid::Type, af::ComputeType::kComputeElewise},
155 {Concat::Type, af::ComputeType::kComputeConcat},159 {Concat::Type, af::ComputeType::kComputeConcat},
156 {Gather::Type, af::ComputeType::kComputeGather},160 {Gather::Type, af::ComputeType::kComputeGather},
161+ {IndirectLoad::Type, af::ComputeType::kComputeLoad},
157 162 
158 {Where::Type, af::ComputeType::kComputeElewise},163 {Where::Type, af::ComputeType::kComputeElewise},
159 {Select::Type, af::ComputeType::kComputeElewise},164 {Select::Type, af::ComputeType::kComputeElewise},
Aautofuse/common/indirect_load_utils.cpp+243-0
@@ -0,0 +1,243 @@
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 "indirect_load_utils.h"
12+ 
13+#include <string>
14+ 
15+#include "ascir_ops.h"
16+#include "ascir_ops_utils.h"
17+#include "common_utils.h"
18+#include "schedule_result.h"
19+ 
20+namespace ascgen_utils::indirect_load {
21+namespace {
22+constexpr char kTemplateOuterAxisAttr[] = "af.internal.indirect_load.outer_axis";
23+constexpr char kTemplateInnerAxisAttr[] = "af.internal.indirect_load.inner_axis";
24+constexpr char kTemplateInputInnerAxisAttr[] = "af.internal.indirect_load.input_inner_axis";
25+constexpr char kTemplateLogicalViewAttr[] = "af.internal.indirect_load.logical_view";
26+ 
27+bool IsValidLogicalTensorView(const LogicalTensorView &view) {
28+ return !view.axis_ids.empty() && view.axis_ids.size() == view.strides.size();
29+}
30+ 
31+TemplateRole GetAnnotatedTemplateRole(const af::AscNodePtr &node) {
32+ if (node == nullptr) {
33+ return TemplateRole::kNone;
34+ }
35+ return static_cast<TemplateRole>(::ascir::GetTemplateRoleOrDefault(*node, static_cast<int64_t>(TemplateRole::kNone)));
36+}
37+ 
38+TemplateBehavior GetBehavior(TemplateRole role) {
39+ TemplateBehavior behavior;
40+ switch (role) {
41+ case TemplateRole::kSimdInputPre:
42+ behavior.skips_main_schedule_tiling = true;
43+ behavior.skips_api_emit = true;
44+ behavior.preserves_vectorized_axis = true;
45+ break;
46+ case TemplateRole::kSimdOp:
47+ break;
48+ case TemplateRole::kSimtInputBoundary:
49+ behavior.skips_main_schedule_tiling = true;
50+ behavior.skips_api_emit = true;
51+ behavior.uses_direct_gm_pipeline = true;
52+ behavior.skips_ub_expr = true;
53+ behavior.preserves_vectorized_axis = true;
54+ break;
55+ case TemplateRole::kSimtDirectGmBoundary:
56+ case TemplateRole::kSimtInlineTransform:
57+ behavior.skips_api_emit = true;
58+ behavior.uses_direct_gm_pipeline = true;
59+ behavior.skips_ub_expr = true;
60+ behavior.preserves_vectorized_axis = true;
61+ break;
62+ case TemplateRole::kSimtOp:
63+ behavior.uses_direct_gm_pipeline = true;
64+ behavior.skips_ub_lifecycle = true;
65+ behavior.skips_ub_expr = true;
66+ behavior.preserves_vectorized_axis = true;
67+ break;
68+ case TemplateRole::kNone:
69+ break;
70+ }
71+ return behavior;
72+}
73+} // namespace
74+ 
75+TemplateRole GetTemplateRole(const af::AscNodePtr &node) {
76+ return GetAnnotatedTemplateRole(node);
77+}
78+ 
79+TemplateBehavior GetTemplateBehavior(const af::AscNodePtr &node) {
80+ return GetBehavior(GetTemplateRole(node));
81+}
82+ 
83+af::Status InheritTemplateRoleIfIL(af::AscGraph &graph, const std::string &vf_node_name, const af::AscNodePtr &src) {
84+ GE_ASSERT_NOTNULL(src);
85+ const auto &behavior = GetTemplateBehavior(src);
86+ if (!behavior.skips_api_emit && !behavior.uses_direct_gm_pipeline) {
87+ return af::SUCCESS;
88+ }
89+ auto vf_node = graph.FindNode(vf_node_name.c_str());
90+ GE_ASSERT_NOTNULL(vf_node, "IndirectLoad: cannot find new VectorFunc node %s.", vf_node_name.c_str());
91+ GE_ASSERT_SUCCESS(::ascir::SetTemplateRole(vf_node, ::ascir::GetTemplateRoleOrDefault(*src)));
92+ return af::SUCCESS;
93+}
94+ 
95+af::Status SetTemplateRole(const af::AscNodePtr &node, TemplateRole role) {
96+ return ::ascir::SetTemplateRole(node, static_cast<int64_t>(role));
97+}
98+ 
99+af::Status SetTemplateAxes(const af::AscNodePtr &node, const TemplateAxes &axes) {
100+ GE_ASSERT_NOTNULL(node);
101+ auto op_desc = node->GetOpDesc();
102+ GE_ASSERT_NOTNULL(op_desc);
103+ GE_ASSERT_TRUE(op_desc->SetExtAttr(kTemplateOuterAxisAttr, static_cast<int64_t>(axes.outer_axis)),
104+ "Set IndirectLoad outer axis failed, node = %s", node->GetNamePtr());
105+ GE_ASSERT_TRUE(op_desc->SetExtAttr(kTemplateInnerAxisAttr, static_cast<int64_t>(axes.inner_axis)),
106+ "Set IndirectLoad inner axis failed, node = %s", node->GetNamePtr());
107+ GE_ASSERT_TRUE(op_desc->SetExtAttr(kTemplateInputInnerAxisAttr, static_cast<int64_t>(axes.input_inner_axis)),
108+ "Set IndirectLoad input inner axis failed, node = %s", node->GetNamePtr());
109+ return af::SUCCESS;
110+}
111+ 
112+af::Status GetTemplateAxes(const af::AscNodePtr &node, TemplateAxes &axes) {
113+ GE_ASSERT_NOTNULL(node);
114+ auto op_desc = node->GetOpDesc();
115+ GE_ASSERT_NOTNULL(op_desc);
116+ axes.outer_axis = op_desc->TryGetExtAttr(kTemplateOuterAxisAttr, static_cast<int64_t>(af::kIdNone));
117+ axes.inner_axis = op_desc->TryGetExtAttr(kTemplateInnerAxisAttr, static_cast<int64_t>(af::kIdNone));
118+ axes.input_inner_axis = op_desc->TryGetExtAttr(kTemplateInputInnerAxisAttr, static_cast<int64_t>(af::kIdNone));
119+ GE_ASSERT_TRUE(axes.outer_axis != af::kIdNone, "IndirectLoad template axes are missing, node = %s",
120+ node->GetNamePtr());
121+ return af::SUCCESS;
122+}
123+ 
124+af::Status SetTemplateLogicalView(const af::AscNodePtr &node, const TemplateLogicalView &view) {
125+ GE_ASSERT_NOTNULL(node);
126+ GE_ASSERT_TRUE(IsValidLogicalTensorView(view.data) && IsValidLogicalTensorView(view.index) &&
127+ IsValidLogicalTensorView(view.output),
128+ "IndirectLoad logical view is invalid, node = %s", node->GetNamePtr());
129+ auto op_desc = node->GetOpDesc();
130+ GE_ASSERT_NOTNULL(op_desc);
131+ GE_ASSERT_TRUE(op_desc->SetExtAttr(kTemplateLogicalViewAttr, view), "Set IndirectLoad logical view failed, node = %s",
132+ node->GetNamePtr());
133+ return af::SUCCESS;
134+}
135+ 
136+af::Status GetTemplateLogicalView(const af::AscNodePtr &node, TemplateLogicalView &view) {
137+ GE_ASSERT_NOTNULL(node);
138+ auto op_desc = node->GetOpDesc();
139+ GE_ASSERT_NOTNULL(op_desc);
140+ view = op_desc->TryGetExtAttr(kTemplateLogicalViewAttr, TemplateLogicalView{});
141+ GE_ASSERT_TRUE(IsValidLogicalTensorView(view.data) && IsValidLogicalTensorView(view.index) &&
142+ IsValidLogicalTensorView(view.output),
143+ "IndirectLoad logical view is missing or invalid, node = %s", node->GetNamePtr());
144+ return af::SUCCESS;
145+}
146+ 
147+bool ShouldSkipMainScheduleTiling(const af::AscNodePtr &node) {
148+ return GetTemplateBehavior(node).skips_main_schedule_tiling;
149+}
150+ 
151+bool ShouldPreserveVectorizedAxis(const af::AscNodePtr &node) {
152+ return GetTemplateBehavior(node).preserves_vectorized_axis;
153+}
154+ 
155+bool ShouldApplyInputInnerVectorization(const af::AscNodePtr &node) {
156+ return GetTemplateRole(node) == TemplateRole::kSimdInputPre;
157+}
158+ 
159+bool ShouldDisableRegularVectorFunc(const af::AscNodePtr &node) {
160+ const TemplateRole role = GetTemplateRole(node);
161+ return role == TemplateRole::kSimtInputBoundary || role == TemplateRole::kSimtDirectGmBoundary ||
162+ role == TemplateRole::kSimtInlineTransform || role == TemplateRole::kSimtOp;
163+}
164+ 
165+af::AscNodePtr GetInputProducer(const af::AscNodePtr &node, size_t input_index) {
166+ auto input_anchor = node == nullptr ? nullptr : node->GetInDataAnchor(input_index);
167+ if (input_anchor == nullptr || input_anchor->GetPeerOutAnchor() == nullptr) {
168+ return nullptr;
169+ }
170+ return std::dynamic_pointer_cast<af::AscNode>(input_anchor->GetPeerOutAnchor()->GetOwnerNode());
171+}
172+ 
173+af::AscNodePtr GetOnlyOutputConsumer(const af::AscNodePtr &node) {
174+ if (node == nullptr || node->GetOutDataNodesSize() != 1UL) {
175+ return nullptr;
176+ }
177+ return std::dynamic_pointer_cast<af::AscNode>(*node->GetOutDataNodes().begin());
178+}
179+ 
180+af::AscNodePtr FindIndirectLoadNode(const af::AscGraph &graph) {
181+ for (const af::AscNodePtr &node : graph.GetAllNodes()) {
182+ if (af::ops::IsOps<af::ascir_op::IndirectLoad>(node)) {
183+ return node;
184+ }
185+ }
186+ return nullptr;
187+}
188+ 
189+af::Status ValidateSingleIndirectLoadNode(const af::AscGraph &graph, af::AscNodePtr &node) {
190+ node = nullptr;
191+ for (const af::AscNodePtr &candidate : graph.GetAllNodes()) {
192+ if (!af::ops::IsOps<af::ascir_op::IndirectLoad>(candidate)) {
193+ continue;
194+ }
195+ if (node != nullptr) {
196+ GELOGE(af::FAILED, "[IndirectLoad] Graph[%s] contains multiple IndirectLoad nodes, first[%s], next[%s].",
197+ graph.GetName().c_str(), node->GetNamePtr(), candidate->GetNamePtr());
198+ }
199+ GE_ASSERT_TRUE(node == nullptr, "Graph contains multiple IndirectLoad nodes, only one is supported.");
200+ node = candidate;
201+ }
202+ if (node != nullptr) {
203+ GELOGD("[IndirectLoad] Graph[%s] found IndirectLoad node[%s].", graph.GetName().c_str(), node->GetNamePtr());
204+ }
205+ return af::SUCCESS;
206+}
207+ 
208+af::Status GetPrebuiltYTilingCase(const af::AscGraph &graph, bool &has_case, af::AxisId &tile_id,
209+ std::pair<af::AxisPtr, af::AxisPtr> &tiling) {
210+ has_case = false;
211+ tile_id = af::kIdNone;
212+ tiling = {nullptr, nullptr};
213+ 
214+ const af::AscNodePtr indirect_load = FindIndirectLoadNode(graph);
215+ if (indirect_load == nullptr) {
216+ return af::SUCCESS;
217+ }
218+ 
219+ TemplateAxes axes;
220+ GE_ASSERT_SUCCESS(GetTemplateAxes(indirect_load, axes), "[IndirectLoad] Failed to get template axes for node[%s].",
221+ indirect_load->GetNamePtr());
222+ 
223+ for (const auto &axis : graph.GetAllAxis()) {
224+ if (axis == nullptr || axis->type != af::Axis::Type::kAxisTypeTileOuter || axis->from.size() != 1UL ||
225+ axis->from[0] != axes.outer_axis) {
226+ continue;
227+ }
228+ GE_ASSERT_TRUE(
229+ axis->split_pair_other_id >= 0L && static_cast<size_t>(axis->split_pair_other_id) < graph.GetAllAxis().size(),
230+ "[IndirectLoad] Invalid split pair axis[%ld] for graph[%s].", axis->split_pair_other_id,
231+ graph.GetName().c_str());
232+ tiling.first = axis;
233+ tiling.second = graph.GetAllAxis()[static_cast<size_t>(axis->split_pair_other_id)];
234+ has_case = true;
235+ tile_id = axes.outer_axis;
236+ return af::SUCCESS;
237+ }
238+ GELOGE(af::FAILED, "[IndirectLoad] Template fixed split is missing for axis[%ld] in graph[%s].", axes.outer_axis,
239+ graph.GetName().c_str());
240+ return af::FAILED;
241+}
242+ 
243+} // namespace ascgen_utils::indirect_load
Aautofuse/common/indirect_load_utils.h+77-0
@@ -0,0 +1,77 @@
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 __INDIRECT_LOAD_UTILS_H__
12+#define __INDIRECT_LOAD_UTILS_H__
13+ 
14+#include <cstddef>
15+#include <string>
16+#include <utility>
17+#include <vector>
18+#include "graph/ascendc_ir/ascendc_ir_core/ascendc_ir.h"
19+ 
20+namespace ascgen_utils::indirect_load {
21+enum class TemplateRole : int64_t {
22+ kNone,
23+ kSimdInputPre,
24+ kSimdOp,
25+ kSimtInputBoundary,
26+ kSimtDirectGmBoundary,
27+ kSimtInlineTransform,
28+ kSimtOp,
29+};
30+ 
31+struct TemplateBehavior {
32+ bool skips_main_schedule_tiling = false;
33+ bool skips_api_emit = false;
34+ bool uses_direct_gm_pipeline = false;
35+ bool skips_ub_lifecycle = false;
36+ bool skips_ub_expr = false;
37+ bool preserves_vectorized_axis = false;
38+};
39+ 
40+struct TemplateAxes {
41+ af::AxisId outer_axis = af::kIdNone;
42+ af::AxisId inner_axis = af::kIdNone;
43+ af::AxisId input_inner_axis = af::kIdNone;
44+};
45+ 
46+struct LogicalTensorView {
47+ std::vector<af::AxisId> axis_ids;
48+ std::vector<af::Expression> strides;
49+};
50+ 
51+struct TemplateLogicalView {
52+ LogicalTensorView data;
53+ LogicalTensorView index;
54+ LogicalTensorView output;
55+};
56+ 
57+TemplateBehavior GetTemplateBehavior(const af::AscNodePtr &node);
58+TemplateRole GetTemplateRole(const af::AscNodePtr &node);
59+af::Status InheritTemplateRoleIfIL(af::AscGraph &graph, const std::string &vf_node_name, const af::AscNodePtr &src);
60+af::Status SetTemplateRole(const af::AscNodePtr &node, TemplateRole role);
61+af::Status SetTemplateAxes(const af::AscNodePtr &node, const TemplateAxes &axes);
62+af::Status GetTemplateAxes(const af::AscNodePtr &node, TemplateAxes &axes);
63+af::Status SetTemplateLogicalView(const af::AscNodePtr &node, const TemplateLogicalView &view);
64+af::Status GetTemplateLogicalView(const af::AscNodePtr &node, TemplateLogicalView &view);
65+bool ShouldSkipMainScheduleTiling(const af::AscNodePtr &node);
66+bool ShouldPreserveVectorizedAxis(const af::AscNodePtr &node);
67+bool ShouldApplyInputInnerVectorization(const af::AscNodePtr &node);
68+bool ShouldDisableRegularVectorFunc(const af::AscNodePtr &node);
69+af::AscNodePtr GetInputProducer(const af::AscNodePtr &node, size_t input_index);
70+af::AscNodePtr GetOnlyOutputConsumer(const af::AscNodePtr &node);
71+af::AscNodePtr FindIndirectLoadNode(const af::AscGraph &graph);
72+af::Status ValidateSingleIndirectLoadNode(const af::AscGraph &graph, af::AscNodePtr &node);
73+af::Status GetPrebuiltYTilingCase(const af::AscGraph &graph, bool &has_case, af::AxisId &tile_id,
74+ std::pair<af::AxisPtr, af::AxisPtr> &tiling);
75+} // namespace ascgen_utils::indirect_load
76+ 
77+#endif // __INDIRECT_LOAD_UTILS_H__
Mautofuse/common/schedule_result.h+61-0
@@ -12,6 +12,13 @@
12#define ASCGEN_DEV_BASE_COMMON_SCHEDULE_RESULT_H_12#define ASCGEN_DEV_BASE_COMMON_SCHEDULE_RESULT_H_
13 13 
14#include "ascendc_ir/ascendc_ir_core/ascendc_ir.h"14#include "ascendc_ir/ascendc_ir_core/ascendc_ir.h"
15+#include "common/checker.h"
16+ 
17+namespace {
18+constexpr char kTemplateIdAttr[] = "af.internal.template.id";
19+constexpr char kTemplateRoleAttr[] = "af.internal.indirect_load.role";
20+constexpr char kDcacheSizeAttr[] = "af.internal.template.dcache_size";
21+} // namespace
15 22 
16namespace ascir {23namespace ascir {
17struct ScheduleGroup {24struct ScheduleGroup {
@@ -46,6 +53,60 @@ struct FusedScheduledResult {
46 std::vector<af::Expression> origin_vars;53 std::vector<af::Expression> origin_vars;
47 std::vector<std::vector<ScheduledResult>> node_idx_to_scheduled_results;54 std::vector<std::vector<ScheduledResult>> node_idx_to_scheduled_results;
48};55};
56+ 
57+enum class TemplateId : int64_t {
58+ kDefault = -1,
59+ kIndirectLoadSimd = 0,
60+ kIndirectLoadSimt = 1,
61+};
62+ 
63+inline af::Status SetTemplateId(const af::AscNodePtr &node, TemplateId template_id) {
64+ GE_ASSERT_NOTNULL(node);
65+ auto op_desc = node->GetOpDesc();
66+ GE_ASSERT_NOTNULL(op_desc);
67+ GE_ASSERT_TRUE(op_desc->SetExtAttr(kTemplateIdAttr, static_cast<int64_t>(template_id)),
68+ "Set internal template id failed, node = %s", node->GetNamePtr());
69+ return af::SUCCESS;
70+}
71+ 
72+inline af::Status SetTemplateRole(const af::AscNodePtr &node, int64_t role) {
73+ GE_ASSERT_NOTNULL(node);
74+ auto op_desc = node->GetOpDesc();
75+ GE_ASSERT_NOTNULL(op_desc);
76+ GE_ASSERT_TRUE(op_desc->SetExtAttr(kTemplateRoleAttr, role), "Set internal template role failed, node = %s",
77+ node->GetNamePtr());
78+ return af::SUCCESS;
79+}
80+ 
81+inline int64_t GetTemplateRoleOrDefault(const af::AscNode &node, int64_t default_role = -1) {
82+ if (node.GetOpDesc() == nullptr) {
83+ return default_role;
84+ }
85+ return node.GetOpDesc()->TryGetExtAttr(kTemplateRoleAttr, default_role);
86+}
87+ 
88+inline TemplateId GetTemplateIdOrDefault(const af::AscNode &node, TemplateId default_id = TemplateId::kDefault) {
89+ if (node.GetOpDesc() == nullptr) {
90+ return default_id;
91+ }
92+ return static_cast<TemplateId>(node.GetOpDesc()->TryGetExtAttr(kTemplateIdAttr, static_cast<int64_t>(default_id)));
93+}
94+ 
95+inline af::Status SetDcacheSize(const af::AscNodePtr &node, int64_t dcache_size) {
96+ GE_ASSERT_NOTNULL(node);
97+ auto op_desc = node->GetOpDesc();
98+ GE_ASSERT_NOTNULL(op_desc);
99+ GE_ASSERT_TRUE(op_desc->SetExtAttr(kDcacheSizeAttr, dcache_size),
100+ "Set internal template dcache size failed, node = %s", node->GetNamePtr());
101+ return af::SUCCESS;
102+}
103+ 
104+inline int64_t GetDcacheSize(const af::AscNode &node) {
105+ if (node.GetOpDesc() == nullptr) {
106+ return 0;
107+ }
108+ return node.GetOpDesc()->TryGetExtAttr(kDcacheSizeAttr, int64_t{0});
109+}
49} // namespace ascir110} // namespace ascir
50 111 
51#endif // ASCGEN_DEV_BASE_COMMON_SCHEDULE_RESULT_H_112#endif // ASCGEN_DEV_BASE_COMMON_SCHEDULE_RESULT_H_
Mautofuse/common/ub_expr/asc_graph_ub_expr_builder.cpp+8-3
@@ -17,11 +17,12 @@
17#include "common/checker.h"17#include "common/checker.h"
18#include "common_utils.h"18#include "common_utils.h"
19#include "graph/types_af.h"19#include "graph/types_af.h"
20+#include "indirect_load_utils.h"
21+#include "schedule_result.h"
20 22 
21namespace ascir {23namespace ascir {
22namespace {24namespace {
23constexpr int64_t kMinTmpBufferSize = 8 * 1024;25constexpr int64_t kMinTmpBufferSize = 8 * 1024;
24-constexpr int64_t kSimtDcacheSize = 32 * 1024;
25constexpr int64_t kBlockAlignBytes = 32;26constexpr int64_t kBlockAlignBytes = 32;
26 27 
27std::string MakeQueueName(int64_t id) {28std::string MakeQueueName(int64_t id) {
@@ -220,8 +221,9 @@ UbExpr CalcReservedUbSize(const af::AscGraph &graph) {
220 UbExpr reserved_ub_size = af::Symbol(ascgen_utils::CalcReservedTmpBufSizeForAscGraph(graph));221 UbExpr reserved_ub_size = af::Symbol(ascgen_utils::CalcReservedTmpBufSizeForAscGraph(graph));
221 for (const auto &node : graph.GetAllNodes()) {222 for (const auto &node : graph.GetAllNodes()) {
222 GE_ASSERT_NOTNULL(node);223 GE_ASSERT_NOTNULL(node);
223- if (node->GetType() == af::ascir_op::Gather::Type) {224+ const int64_t simt_dcache_size = ::ascir::GetDcacheSize(*node);
224- reserved_ub_size = reserved_ub_size + af::Symbol(kSimtDcacheSize);225+ if (simt_dcache_size > 0) {
226+ reserved_ub_size = reserved_ub_size + af::Symbol(simt_dcache_size);
225 break;227 break;
226 }228 }
227 }229 }
@@ -305,6 +307,9 @@ af::Status AscGraphUbExprBuilder::Build(const af::AscGraph &graph, UbExprContext
305 std::map<int64_t, ContainerState> buffer_bytes;307 std::map<int64_t, ContainerState> buffer_bytes;
306 for (const auto &node : graph.GetAllNodes()) {308 for (const auto &node : graph.GetAllNodes()) {
307 GE_ASSERT_NOTNULL(node);309 GE_ASSERT_NOTNULL(node);
310+ if (ascgen_utils::indirect_load::GetTemplateBehavior(node).skips_ub_expr) {
311+ continue;
312+ }
308 for (const auto &output : node->outputs()) {313 for (const auto &output : node->outputs()) {
309 AddTensor(graph, output->attr, queue_bytes, buffer_bytes);314 AddTensor(graph, output->attr, queue_bytes, buffer_bytes);
310 }315 }
Mautofuse/compiler/py_module/pyascir.cpp+12-0
@@ -52,6 +52,8 @@ inline constexpr char kScalarDataOpType[] = "ScalarData";
52inline constexpr char kAscGraphAttr[] = "ascgraph";52inline constexpr char kAscGraphAttr[] = "ascgraph";
53inline constexpr char kNegativeSlopeAttr[] = "negative_slope";53inline constexpr char kNegativeSlopeAttr[] = "negative_slope";
54inline constexpr char kNegativeIndexSupportAttr[] = "negative_index_support";54inline constexpr char kNegativeIndexSupportAttr[] = "negative_index_support";
55+inline constexpr char kNeedCheckBoundAttr[] = "need_check_bound";
56+inline constexpr char kMaxAttr[] = "max";
55inline constexpr char kAlphaAttr[] = "alpha";57inline constexpr char kAlphaAttr[] = "alpha";
56inline constexpr char kNAttr[] = "n";58inline constexpr char kNAttr[] = "n";
57inline constexpr char kErrorMsgAttr[] = "error_msg";59inline constexpr char kErrorMsgAttr[] = "error_msg";
@@ -1150,6 +1152,14 @@ DEFINE_IR_ATTR_ACCESSORS(IndexExpr, AscIndexExprIrAttrDef, kExprAttr, int64_t, P
1150 PyLong_AsLong, SetExpr, GetExpr)1152 PyLong_AsLong, SetExpr, GetExpr)
1151DEFINE_IR_ATTR_ACCESSORS(Gather, AscGatherIrAttrDef, kAxisAttr, int64_t, PyLong_Check, PyLong_FromLong, PyLong_AsLong,1153DEFINE_IR_ATTR_ACCESSORS(Gather, AscGatherIrAttrDef, kAxisAttr, int64_t, PyLong_Check, PyLong_FromLong, PyLong_AsLong,
1152 SetAxis, GetAxis)1154 SetAxis, GetAxis)
1155+DEFINE_IR_ATTR_ACCESSORS(IndirectLoad, AscIndirectLoadIrAttrDef, kAxisAttr, int64_t, PyLong_Check, PyLong_FromLong,
1156+ PyLong_AsLong, SetAxis, GetAxis)
1157+DEFINE_IR_ATTR_ACCESSORS(IndirectLoad, AscIndirectLoadIrAttrDef, kNegativeIndexSupportAttr, bool, PyBool_Check,
1158+ PyBool_FromLong, PyObject_IsTrue, SetNegative_index_support, GetNegative_index_support)
1159+DEFINE_IR_ATTR_ACCESSORS(IndirectLoad, AscIndirectLoadIrAttrDef, kNeedCheckBoundAttr, bool, PyBool_Check,
1160+ PyBool_FromLong, PyObject_IsTrue, SetNeed_check_bound, GetNeed_check_bound)
1161+DEFINE_IR_ATTR_ACCESSORS(IndirectLoad, AscIndirectLoadIrAttrDef, kMaxAttr, int64_t, PyLong_Check, PyLong_FromLong,
1162+ PyLong_AsLong, SetMax, GetMax)
1153DEFINE_IR_ATTR_ACCESSORS(MatMul, AscMatMulIrAttrDef, kHasRelu, int64_t, PyLong_Check, PyLong_FromLong, PyLong_AsLong,1163DEFINE_IR_ATTR_ACCESSORS(MatMul, AscMatMulIrAttrDef, kHasRelu, int64_t, PyLong_Check, PyLong_FromLong, PyLong_AsLong,
1154 SetHas_relu, GetHas_relu)1164 SetHas_relu, GetHas_relu)
1155DEFINE_IR_ATTR_ACCESSORS(MatMul, AscMatMulIrAttrDef, kTransposeX1, int64_t, PyLong_Check, PyLong_FromLong,1165DEFINE_IR_ATTR_ACCESSORS(MatMul, AscMatMulIrAttrDef, kTransposeX1, int64_t, PyLong_Check, PyLong_FromLong,
@@ -1287,6 +1297,8 @@ const std::map<std::string, typename IrAttr<OpType>::handler> IrAttr<OpType>::at
1287 {"Load", AutoRegAttrHandle<af::ascir_op::Load, kOffsetAttr>::RegHandle},1297 {"Load", AutoRegAttrHandle<af::ascir_op::Load, kOffsetAttr>::RegHandle},
1288 {"Store", AutoRegAttrHandle<af::ascir_op::Store, kOffsetAttr>::RegHandle},1298 {"Store", AutoRegAttrHandle<af::ascir_op::Store, kOffsetAttr>::RegHandle},
1289 {"Gather", AutoRegAttrHandle<af::ascir_op::Gather, kAxisAttr>::RegHandle},1299 {"Gather", AutoRegAttrHandle<af::ascir_op::Gather, kAxisAttr>::RegHandle},
1300+ {"IndirectLoad", AutoRegAttrHandle<af::ascir_op::IndirectLoad, kAxisAttr, kNegativeIndexSupportAttr,
1301+ kNeedCheckBoundAttr, kMaxAttr>::RegHandle},
1290 {"MatMul",1302 {"MatMul",
1291 AutoRegAttrHandle<af::ascir_op::MatMul, kHasRelu, kOffsetX, kTransposeX1, kTransposeX2, kEnableHf32>::RegHandle},1303 AutoRegAttrHandle<af::ascir_op::MatMul, kHasRelu, kOffsetX, kTransposeX1, kTransposeX2, kEnableHf32>::RegHandle},
1292 {"MatMulBias", AutoRegAttrHandle<af::ascir_op::MatMulBias, kHasRelu, kOffsetX, kTransposeX1, kTransposeX2,1304 {"MatMulBias", AutoRegAttrHandle<af::ascir_op::MatMulBias, kHasRelu, kOffsetX, kTransposeX1, kTransposeX2,
Mautofuse/compiler/py_module/pyascir.h+1-0
@@ -116,6 +116,7 @@ PyMODINIT_FUNC PyInit_pyautofuse(void);
116 OP(BatchMatMulBias) \116 OP(BatchMatMulBias) \
117 OP(Where) \117 OP(Where) \
118 OP(Gather) \118 OP(Gather) \
119+ OP(IndirectLoad) \
119 OP(Transpose) \120 OP(Transpose) \
120 OP(BitwiseAnd) \121 OP(BitwiseAnd) \
121 OP(Ln) \122 OP(Ln) \
Mautofuse/compiler/python/ascir_api.py+24-0
@@ -1837,6 +1837,30 @@ def Gather(
1837 return op.y1837 return op.y
1838 1838 
1839 1839 
1840+def IndirectLoad(
1841+ owner_graph: ascir.HintGraph,
1842+ x1: ascir.OpsOperatorOutput,
1843+ x2: ascir.OpsOperatorOutput,
1844+ *,
1845+ axis: int,
1846+ sched_axis: List[ascir.Axis],
1847+ size: Optional[List[ascir.SizeExpr]] = None,
1848+ stride: Optional[List[ascir.SizeExpr]] = None,
1849+) -> ascir.OpsOperatorOutput:
1850+ name = _generate_op_name(owner_graph, "indirectload")
1851+ op = ascir.ops.IndirectLoad(name)
1852+ meta = _get_metadata(owner_graph)
1853+ meta.ops.append(op)
1854+ 
1855+ op.attr.ir_attr.axis = axis
1856+ op.attr.sched.axis = sched_axis
1857+ op.x1 = x1
1858+ op.x2 = x2
1859+ _infer_or_set_view(op.y, sched_axis, size, stride)
1860+ op.infer_dtype()
1861+ return op.y
1862+ 
1863+ 
1840def BitwiseAnd(1864def BitwiseAnd(
1841 owner_graph: ascir.HintGraph,1865 owner_graph: ascir.HintGraph,
1842 x1: ascir.OpsOperatorOutput,1866 x1: ascir.OpsOperatorOutput,
Mautofuse/inc/graph_metadef/graph/ascendc_ir/ascir_registry.h+13-0
@@ -188,6 +188,19 @@ class AscIrCodegen {
188 (void)node;188 (void)node;
189 return true;189 return true;
190 }190 }
191+ 
192+ virtual bool IsSimtScalarSupported(const AscNode &node) const {
193+ (void)node;
194+ return false;
195+ }
196+ 
197+ virtual ge::graphStatus GenerateSimtScalarExpr(const AscNode &node, const std::vector<std::string> &inputs,
198+ std::string &expr) const {
199+ (void)node;
200+ (void)inputs;
201+ (void)expr;
202+ return af::FAILED;
203+ }
191};204};
192 205 
193class AscIrAtt {206class AscIrAtt {
Mautofuse/optimize/autoschedule/autoschedule.cpp+61-2
@@ -10,6 +10,7 @@
10 10 
11#include "autoschedule.h"11#include "autoschedule.h"
12#include <algorithm>12#include <algorithm>
13+#include <numeric>
13#include <sstream>14#include <sstream>
14#include <string>15#include <string>
15#include <queue>16#include <queue>
@@ -18,6 +19,7 @@
18#include "ascir_utils.h"19#include "ascir_utils.h"
19#include "schedule_utils.h"20#include "schedule_utils.h"
20#include "common_utils.h"21#include "common_utils.h"
22+#include "indirect_load_utils.h"
21#include "node_utils.h"23#include "node_utils.h"
22#include "ascendc_ir/core/ascendc_ir_impl.h"24#include "ascendc_ir/core/ascendc_ir_impl.h"
23#include "graph/symbolizer/symbolic_utils.h"25#include "graph/symbolizer/symbolic_utils.h"
@@ -179,6 +181,33 @@ void AppendIdIfNotDefault(std::stringstream &ss, const std::string &prefix, int6
179 ss << prefix << id;181 ss << prefix << id;
180 }182 }
181}183}
184+ 
185+// Returns true if an IndirectLoad prebuilt tiling case was generated, and the caller should return early.
186+bool TryGenIndirectLoadTilingCase(ascir::ImplGraph &graph,
187+ std::vector<optimize::autoschedule::TilingCase> &tiling_cases) {
188+ bool has_indirect_load_case = false;
189+ af::AxisId indirect_load_tile_id = kDefaultAxisId;
190+ std::pair<af::AxisPtr, af::AxisPtr> indirect_load_tiling;
191+ if (ascgen_utils::indirect_load::GetPrebuiltYTilingCase(graph, has_indirect_load_case, indirect_load_tile_id,
192+ indirect_load_tiling) != af::SUCCESS) {
193+ GELOGE(af::FAILED, "[IndirectLoad] Failed to generate prebuilt tiling case for graph[%s].",
194+ graph.GetName().c_str());
195+ return true;
196+ }
197+ if (!has_indirect_load_case) {
198+ return false;
199+ }
200+ optimize::autoschedule::TilingCase tiling_case;
201+ if (indirect_load_tile_id != kDefaultAxisId) {
202+ tiling_case.ub_tiling_id_y = indirect_load_tile_id;
203+ }
204+ tiling_case.ub_tiling_y = indirect_load_tiling;
205+ tiling_case.block_tiling_id = 0;
206+ tiling_cases.push_back(tiling_case);
207+ GELOGD("[IndirectLoad] Graph[%s] generate prebuilt outer tiling case for axis[%ld].", graph.GetName().c_str(),
208+ indirect_load_tile_id);
209+ return true;
210+}
182} // namespace211} // namespace
183 212 
184namespace optimize::autoschedule {213namespace optimize::autoschedule {
@@ -187,7 +216,8 @@ Status AutoSchedule::SelectLoopAxis(ascir::ImplGraph &impl_graph, bool is_reduce
187 for (auto node : impl_graph.GetAllNodes()) {216 for (auto node : impl_graph.GetAllNodes()) {
188 GE_ASSERT_NOTNULL(node);217 GE_ASSERT_NOTNULL(node);
189 node->attr.sched.loop_axis = af::kIdNone;218 node->attr.sched.loop_axis = af::kIdNone;
190- if (node->attr.api.type != af::ApiType::kAPITypeCompute) {219+ const auto behavior = ascgen_utils::indirect_load::GetTemplateBehavior(node);
220+ if (node->attr.api.type != af::ApiType::kAPITypeCompute || behavior.skips_main_schedule_tiling) {
191 continue;221 continue;
192 }222 }
193 if (ScheduleUtils::IsReduce(node) && !is_reduce_fullload) {223 if (ScheduleUtils::IsReduce(node) && !is_reduce_fullload) {
@@ -220,6 +250,10 @@ void AutoSchedule::GenTilingCase(std::vector<TilingCase> &tiling_cases) {
220 }250 }
221 };251 };
222 252 
253+ if (TryGenIndirectLoadTilingCase(graph_, tiling_cases)) {
254+ return;
255+ }
256+ 
223 if (cube_template_ != ascir::CubeTemplateType::kDefault) {257 if (cube_template_ != ascir::CubeTemplateType::kDefault) {
224 for (const auto &y_id : axes_group_.y_group) {258 for (const auto &y_id : axes_group_.y_group) {
225 TilingCase tiling_case;259 TilingCase tiling_case;
@@ -305,6 +339,27 @@ static std::string GetTilingCaseStr(const std::string &graph_name, const TilingC
305 return ss.str();339 return ss.str();
306}340}
307 341 
342+Status BuildIndirectLoadAxisGroup(const ascir::ImplGraph &graph, const AxisGroup &base_group, AxisGroup &axis_group) {
343+ const af::AscNodePtr indirect_load = ascgen_utils::indirect_load::FindIndirectLoadNode(graph);
344+ if (indirect_load == nullptr) {
345+ axis_group = base_group;
346+ return af::SUCCESS;
347+ }
348+ 
349+ ascgen_utils::indirect_load::TemplateAxes axes;
350+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, axes));
351+ axis_group = {};
352+ axis_group.y_group.push_back(axes.outer_axis);
353+ if (axes.inner_axis != af::kIdNone) {
354+ axis_group.y_group.push_back(axes.inner_axis);
355+ }
356+ axis_group.axes_order.resize(axis_group.y_group.size());
357+ std::iota(axis_group.axes_order.begin(), axis_group.axes_order.end(), 0UL);
358+ GELOGD("[IndirectLoad] Graph[%s] use template axis group[%s].", graph.GetName().c_str(),
359+ axis_group.ToString().c_str());
360+ return af::SUCCESS;
361+}
362+ 
308Status AutoSchedule::DoAutoSchedule() {363Status AutoSchedule::DoAutoSchedule() {
309 graph_.SetGraphType(af::AscGraphType::kImplGraph);364 graph_.SetGraphType(af::AscGraphType::kImplGraph);
310 ReorderBroadcastAxesInner(graph_);365 ReorderBroadcastAxesInner(graph_);
@@ -356,7 +411,11 @@ Status AutoSchedule::ProcessOneTilingCase(TilingCase &tiling_case, size_t index,
356 GE_ASSERT_TRUE(output.scheduled_graph.CopyFrom(graph_), "Failed to copy graph for tiling case %zu in graph: [%s]",411 GE_ASSERT_TRUE(output.scheduled_graph.CopyFrom(graph_), "Failed to copy graph for tiling case %zu in graph: [%s]",
357 index, graph_.GetName().c_str());412 index, graph_.GetName().c_str());
358 413 
359- Scheduler scheduler(output.scheduled_graph, axes_group_, tiling_case, is_last_axis_reduce, reduce_template_,414+ AxisGroup axis_group;
415+ GE_CHK_STATUS_RET(BuildIndirectLoadAxisGroup(output.scheduled_graph, axes_group_, axis_group),
416+ "Failed to build scheduler axis group for tiling case %zu in graph: [%s]", index,
417+ graph_name.c_str());
418+ Scheduler scheduler(output.scheduled_graph, axis_group, tiling_case, is_last_axis_reduce, reduce_template_,
360 cube_template_);419 cube_template_);
361 420 
362 auto ret = scheduler.DoScheduler();421 auto ret = scheduler.DoScheduler();
Mautofuse/optimize/autoschedule/schedule.cpp+165-27
@@ -11,6 +11,7 @@
11#include "schedule.h"11#include "schedule.h"
12#include <numeric>12#include <numeric>
13#include "alignment_handler.h"13#include "alignment_handler.h"
14+#include "indirect_load_utils.h"
14#include "schedule_utils.h"15#include "schedule_utils.h"
15#include "node_cache_marker.h"16#include "node_cache_marker.h"
16 17 
@@ -179,12 +180,6 @@ void AppendAxisOrder(const std::vector<size_t> &axes_order, size_t group_axis_si
179 axes_order.begin() + inner_axis_end_idx);180 axes_order.begin() + inner_axis_end_idx);
180}181}
181 182 
182-void AdjustAxisOrderOffsets(std::vector<size_t> &axes_order, size_t start_idx, size_t end_idx, size_t offset) {
183- for (size_t i = start_idx; i < end_idx; ++i) {
184- axes_order[i] += offset;
185- }
186-}
187- 
188void GetOuterAxes(const std::vector<ascir::AxisId> &axes_group, const ascir::AxisId &ub_tiling_id,183void GetOuterAxes(const std::vector<ascir::AxisId> &axes_group, const ascir::AxisId &ub_tiling_id,
189 const ascir::Axis &ub_tiling_outer_axis, const std::vector<size_t> &axes_order,184 const ascir::Axis &ub_tiling_outer_axis, const std::vector<size_t> &axes_order,
190 std::vector<ascir::AxisId> &outer_axes, std::vector<size_t> &outer_axes_index,185 std::vector<ascir::AxisId> &outer_axes, std::vector<size_t> &outer_axes_index,
@@ -405,8 +400,25 @@ Status Scheduler::ApplyBlockSplitToNode(ascir::NodeView &node, bool is_store_aft
405 return af::SUCCESS;400 return af::SUCCESS;
406}401}
407 402 
403+void Scheduler::AdjustVectorizedAxesOrderOffsets(std::vector<size_t> &vectorized_axes_order, size_t split_point,
404+ size_t end, size_t offset) const {
405+ const size_t start_idx = is_last_axis_reduce_ ? split_point : 0;
406+ const size_t end_idx = is_last_axis_reduce_ ? end : split_point;
407+ for (size_t i = start_idx; i < end_idx; ++i) {
408+ vectorized_axes_order[i] += offset;
409+ }
410+}
411+ 
408void Scheduler::FindVectorizedAxes(std::vector<ascir::AxisId> &vectorized_axes,412void Scheduler::FindVectorizedAxes(std::vector<ascir::AxisId> &vectorized_axes,
409 std::vector<size_t> &vectorized_axes_order) {413 std::vector<size_t> &vectorized_axes_order) {
414+ if (is_indirect_load_schedule_case_) {
415+ if (indirect_load_axes_.inner_axis != af::kIdNone) {
416+ vectorized_axes.push_back(indirect_load_axes_.inner_axis);
417+ vectorized_axes_order.push_back(0UL);
418+ }
419+ return;
420+ }
421+ 
410 size_t last_ub_size = 0UL;422 size_t last_ub_size = 0UL;
411 size_t group_axis_size = 0UL;423 size_t group_axis_size = 0UL;
412 const auto &axes_order = axes_group_.axes_order;424 const auto &axes_order = axes_group_.axes_order;
@@ -436,11 +448,7 @@ void Scheduler::FindVectorizedAxes(std::vector<ascir::AxisId> &vectorized_axes,
436 448 
437 // 带reduce 需要在ub内确保向量化轴是RA或者AR排布449 // 带reduce 需要在ub内确保向量化轴是RA或者AR排布
438 const size_t offset = axes_order.size() + axes_group_.n_group.size();450 const size_t offset = axes_order.size() + axes_group_.n_group.size();
439- if (is_last_axis_reduce_) {451+ AdjustVectorizedAxesOrderOffsets(vectorized_axes_order, prev_ub_size, current_ub_size, offset);
440- AdjustAxisOrderOffsets(vectorized_axes_order, prev_ub_size, current_ub_size, offset);
441- } else {
442- AdjustAxisOrderOffsets(vectorized_axes_order, 0, prev_ub_size, offset);
443- }
444 last_ub_size = current_ub_size;452 last_ub_size = current_ub_size;
445 }453 }
446 454 
@@ -453,11 +461,7 @@ void Scheduler::FindVectorizedAxes(std::vector<ascir::AxisId> &vectorized_axes,
453 }461 }
454 462 
455 const size_t offset = axes_order.size() + vectorized_axes.size();463 const size_t offset = axes_order.size() + vectorized_axes.size();
456- if (is_last_axis_reduce_) {464+ AdjustVectorizedAxesOrderOffsets(vectorized_axes_order, non_reduce_axis_size, vectorized_axes.size(), offset);
457- AdjustAxisOrderOffsets(vectorized_axes_order, non_reduce_axis_size, vectorized_axes.size(), offset);
458- } else {
459- AdjustAxisOrderOffsets(vectorized_axes_order, 0, non_reduce_axis_size, offset);
460- }
461 }465 }
462}466}
463 467 
@@ -496,15 +500,123 @@ Status Scheduler::RemoveRedundantBroadcastNode(const ascir::ImplGraph &impl_grap
496 return af::SUCCESS;500 return af::SUCCESS;
497}501}
498 502 
499-Status Scheduler::TileSplit() {503+namespace {
500- // split ub504+Status AddIndirectLoadSyntheticOuterAxis(const af::AscNodePtr &node, ascir::AxisId outer_axis_id,
501- TileTiling(tiling_case_.ub_tiling_id_x, tiling_case_.ub_tiling_x);505+ bool has_synthetic_outer_axis) {
502- TileTiling(tiling_case_.ub_tiling_id_y, tiling_case_.ub_tiling_y);506+ if (!has_synthetic_outer_axis) {
503- TileTiling(tiling_case_.ub_tiling_id_r, tiling_case_.ub_tiling_r);507+ return af::SUCCESS;
508+ }
504 509 
510+ node->attr.sched.axis.insert(node->attr.sched.axis.begin(), outer_axis_id);
511+ for (const auto &output : node->outputs()) {
512+ ascir::SizeExpr stride = af::sym::kSymbolOne;
513+ if (!output->attr.axis.empty()) {
514+ stride = af::sym::Mul(output->attr.repeats.front(), output->attr.strides.front());
515+ }
516+ output->attr.axis.insert(output->attr.axis.begin(), outer_axis_id);
517+ output->attr.repeats.insert(output->attr.repeats.begin(), af::sym::kSymbolOne);
518+ output->attr.strides.insert(output->attr.strides.begin(), stride);
519+ }
520+ return af::SUCCESS;
521+}
522+ 
523+Status ApplyIndirectLoadTemplateMerge(ascir::ImplGraph &graph, const af::AscNodePtr &node, ascir::AxisId axis_id) {
524+ if (axis_id == af::kIdNone) {
525+ return af::SUCCESS;
526+ }
527+ const auto axis = graph.FindAxis(axis_id);
528+ GE_ASSERT_NOTNULL(axis, "IndirectLoad template axis[%ld] is not found.", axis_id);
529+ if (axis->type == ascir::Axis::Type::kAxisTypeMerged) {
530+ GELOGD("[IndirectLoad] Graph[%s] apply template merge axis[%ld] for node[%s].", graph.GetName().c_str(), axis_id,
531+ node->GetNamePtr());
532+ GE_ASSERT_TRUE(graph.ApplySchedAxisMerge(node, axis_id), "Failed to merge schedule axis[%ld] for node[%s].",
533+ axis_id, node->GetNamePtr());
534+ GE_ASSERT_TRUE(graph.ApplyTensorAxisMerge(node, axis_id), "Failed to merge tensor axis[%ld] for node[%s].", axis_id,
535+ node->GetNamePtr());
536+ } else {
537+ GELOGD("[IndirectLoad] Graph[%s] apply single-axis template merge axis[%ld] for node[%s].", graph.GetName().c_str(),
538+ axis_id, node->GetNamePtr());
539+ GE_ASSERT_TRUE(graph.ApplySchedAxisMerge(node, axis_id, {axis_id}),
540+ "Failed to merge schedule axis[%ld] for node[%s].", axis_id, node->GetNamePtr());
541+ GE_ASSERT_TRUE(graph.ApplyTensorAxisMerge(node, axis_id, {axis_id}),
542+ "Failed to merge tensor axis[%ld] for node[%s].", axis_id, node->GetNamePtr());
543+ }
544+ return af::SUCCESS;
545+}
546+ 
547+Status ApplyInputInnerVectorizedAxis(ascir::ImplGraph &graph, const af::AscNodePtr &node,
548+ ascir::AxisId input_inner_axis_id) {
549+ if (input_inner_axis_id == af::kIdNone) {
550+ return af::SUCCESS;
551+ }
552+ GE_ASSERT_TRUE(!node->outputs().empty(), "IndirectLoad input-pre node[%s] has no output.", node->GetNamePtr());
553+ for (const auto &output : node->outputs()) {
554+ output->attr.vectorized_axis = {input_inner_axis_id};
555+ output->attr.vectorized_strides = {af::sym::kSymbolOne};
556+ }
557+ 
558+ const auto input_inner_axis = graph.FindAxis(input_inner_axis_id);
559+ GE_ASSERT_NOTNULL(input_inner_axis, "IndirectLoad input inner axis[%ld] is not found.", input_inner_axis_id);
560+ if (input_inner_axis->type != ascir::Axis::Type::kAxisTypeMerged) {
561+ return af::SUCCESS;
562+ }
563+ GELOGD("[IndirectLoad] Graph[%s] apply input inner vectorized axis[%ld] for node[%s].", graph.GetName().c_str(),
564+ input_inner_axis_id, node->GetNamePtr());
565+ GE_ASSERT_TRUE(graph.ApplySchedAxisMerge(node, input_inner_axis_id),
566+ "Failed to merge input inner schedule axis[%ld] for node[%s].", input_inner_axis_id,
567+ node->GetNamePtr());
568+ GE_ASSERT_TRUE(graph.ApplyTensorAxisMerge(node, input_inner_axis_id),
569+ "Failed to merge input inner tensor axis[%ld] for node[%s].", input_inner_axis_id, node->GetNamePtr());
570+ return af::SUCCESS;
571+}
572+} // namespace
573+ 
574+Status Scheduler::InitIndirectLoadScheduleCase() {
575+ is_indirect_load_schedule_case_ = false;
576+ has_indirect_load_synthetic_outer_axis_ = false;
577+ if (tiling_case_.ub_tiling_y.first == nullptr) {
578+ return af::SUCCESS;
579+ }
580+ const af::AscNodePtr indirect_load = ascgen_utils::indirect_load::FindIndirectLoadNode(graph_);
581+ if (indirect_load == nullptr) {
582+ return af::SUCCESS;
583+ }
584+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, indirect_load_axes_));
585+ const auto outer_axis = graph_.FindAxis(indirect_load_axes_.outer_axis);
586+ GE_ASSERT_NOTNULL(outer_axis, "IndirectLoad outer axis[%ld] is not found.", indirect_load_axes_.outer_axis);
587+ has_indirect_load_synthetic_outer_axis_ =
588+ outer_axis->from.empty() &&
589+ std::find(indirect_load->attr.sched.axis.begin(), indirect_load->attr.sched.axis.end(),
590+ indirect_load_axes_.outer_axis) == indirect_load->attr.sched.axis.end();
591+ is_indirect_load_schedule_case_ = true;
592+ return af::SUCCESS;
593+}
594+ 
595+Status Scheduler::ApplyIndirectLoadNodeAxes(const af::AscNodePtr &node, bool &skip_main_tiling) const {
596+ skip_main_tiling = false;
597+ if (!is_indirect_load_schedule_case_) {
598+ return af::SUCCESS;
599+ }
600+ if (ascgen_utils::indirect_load::ShouldApplyInputInnerVectorization(node)) {
601+ skip_main_tiling = true;
602+ GE_ASSERT_SUCCESS(ApplyInputInnerVectorizedAxis(graph_, node, indirect_load_axes_.input_inner_axis));
603+ return af::SUCCESS;
604+ }
605+ if (ascgen_utils::indirect_load::ShouldSkipMainScheduleTiling(node)) {
606+ skip_main_tiling = true;
607+ return af::SUCCESS;
608+ }
609+ GE_ASSERT_SUCCESS(
610+ AddIndirectLoadSyntheticOuterAxis(node, indirect_load_axes_.outer_axis, has_indirect_load_synthetic_outer_axis_));
611+ GE_ASSERT_SUCCESS(ApplyIndirectLoadTemplateMerge(graph_, node, indirect_load_axes_.outer_axis));
612+ GE_ASSERT_SUCCESS(ApplyIndirectLoadTemplateMerge(graph_, node, indirect_load_axes_.inner_axis));
613+ return af::SUCCESS;
614+}
615+ 
616+std::vector<ascir::AxisId> Scheduler::GetSortedNodeVectorizedAxes(Scheduler &scheduler) {
505 std::vector<ascir::AxisId> vectorized_axes;617 std::vector<ascir::AxisId> vectorized_axes;
506 std::vector<size_t> vectorized_axes_order;618 std::vector<size_t> vectorized_axes_order;
507- FindVectorizedAxes(vectorized_axes, vectorized_axes_order);619+ scheduler.FindVectorizedAxes(vectorized_axes, vectorized_axes_order);
508 620 
509 // reorder vectorized axis by original node axis order621 // reorder vectorized axis by original node axis order
510 std::vector<size_t> base_order(vectorized_axes.size());622 std::vector<size_t> base_order(vectorized_axes.size());
@@ -513,17 +625,32 @@ Status Scheduler::TileSplit() {
513 return vectorized_axes_order[a] < vectorized_axes_order[b];625 return vectorized_axes_order[a] < vectorized_axes_order[b];
514 });626 });
515 627 
516- std::vector<ascir::AxisId> sorted_node_vectorized_axes;628+ std::vector<ascir::AxisId> sorted;
517- sorted_node_vectorized_axes.reserve(base_order.size());629+ sorted.reserve(base_order.size());
518 for (const size_t index : base_order) {630 for (const size_t index : base_order) {
519- sorted_node_vectorized_axes.push_back(vectorized_axes[index]);631+ sorted.push_back(vectorized_axes[index]);
520 }632 }
633+ return sorted;
634+}
635+ 
636+Status Scheduler::TileSplit() {
637+ // split ub
638+ TileTiling(tiling_case_.ub_tiling_id_x, tiling_case_.ub_tiling_x);
639+ TileTiling(tiling_case_.ub_tiling_id_y, tiling_case_.ub_tiling_y);
640+ TileTiling(tiling_case_.ub_tiling_id_r, tiling_case_.ub_tiling_r);
641+ 
642+ auto sorted_node_vectorized_axes = GetSortedNodeVectorizedAxes(*this);
521 643 
522 bool has_reduce = graph_cache_.HasComputeType(af::ComputeType::kComputeReduce);644 bool has_reduce = graph_cache_.HasComputeType(af::ComputeType::kComputeReduce);
523 for (auto node : graph_.GetAllNodes()) {645 for (auto node : graph_.GetAllNodes()) {
524 if (ScheduleUtils::IsBuffer(node)) {646 if (ScheduleUtils::IsBuffer(node)) {
525 continue;647 continue;
526 }648 }
649+ bool skip_main_tiling = false;
650+ GE_CHK_STATUS_RET(ApplyIndirectLoadNodeAxes(node, skip_main_tiling));
651+ if (skip_main_tiling) {
652+ continue;
653+ }
527 ApplyTiling(node, tiling_case_.ub_tiling_id_x, tiling_case_.ub_tiling_x);654 ApplyTiling(node, tiling_case_.ub_tiling_id_x, tiling_case_.ub_tiling_x);
528 ApplyTiling(node, tiling_case_.ub_tiling_id_y, tiling_case_.ub_tiling_y);655 ApplyTiling(node, tiling_case_.ub_tiling_id_y, tiling_case_.ub_tiling_y);
529 ApplyTiling(node, tiling_case_.ub_tiling_id_r, tiling_case_.ub_tiling_r);656 ApplyTiling(node, tiling_case_.ub_tiling_id_r, tiling_case_.ub_tiling_r);
@@ -543,6 +670,9 @@ Status Scheduler::TileSplit() {
543 // 非reduce场景应该将向量化轴调整为tensor中的相对顺序, 带reduce场景由于tiling策略已经做了特别的reorder,需要跳过670 // 非reduce场景应该将向量化轴调整为tensor中的相对顺序, 带reduce场景由于tiling策略已经做了特别的reorder,需要跳过
544 // tiling策略暂时无法支持具有reduce和transpose融合的场景671 // tiling策略暂时无法支持具有reduce和transpose融合的场景
545 for (auto &output : node->outputs()) {672 for (auto &output : node->outputs()) {
673+ if (is_indirect_load_schedule_case_ && ascgen_utils::indirect_load::ShouldPreserveVectorizedAxis(node)) {
674+ continue;
675+ }
546 output->attr.vectorized_axis = node_vectorized_axes;676 output->attr.vectorized_axis = node_vectorized_axes;
547 if (!has_reduce) {677 if (!has_reduce) {
548 auto tensor_axis = output->attr.axis;678 auto tensor_axis = output->attr.axis;
@@ -561,6 +691,8 @@ Status Scheduler::DoScheduler() {
561 ascir::utils::DumpGraph(graph_, "AfterDoTiling");691 ascir::utils::DumpGraph(graph_, "AfterDoTiling");
562 return af::SUCCESS;692 return af::SUCCESS;
563 }693 }
694+ GE_CHK_STATUS_RET(InitIndirectLoadScheduleCase(), "Failed to init IndirectLoad schedule case for graph[%s].",
695+ graph_.GetName().c_str());
564 RemoveDuplicatedAxisFromGroup();696 RemoveDuplicatedAxisFromGroup();
565 // Tile Split697 // Tile Split
566 TileSplit();698 TileSplit();
@@ -586,7 +718,8 @@ Status Scheduler::DoScheduler() {
586Status Scheduler::ApplyBlockSplit(const std::vector<ascir::AxisId> &new_sched_axes) {718Status Scheduler::ApplyBlockSplit(const std::vector<ascir::AxisId> &new_sched_axes) {
587 bool is_reduce_after = false;719 bool is_reduce_after = false;
588 for (auto node : graph_.GetAllNodes()) {720 for (auto node : graph_.GetAllNodes()) {
589- if (ScheduleUtils::IsBuffer(node)) {721+ if (ScheduleUtils::IsBuffer(node) ||
722+ (is_indirect_load_schedule_case_ && ascgen_utils::indirect_load::ShouldSkipMainScheduleTiling(node))) {
590 continue;723 continue;
591 }724 }
592 if ((!is_reduce_after) && ScheduleUtils::IsReduce(node)) {725 if ((!is_reduce_after) && ScheduleUtils::IsReduce(node)) {
@@ -596,9 +729,14 @@ Status Scheduler::ApplyBlockSplit(const std::vector<ascir::AxisId> &new_sched_ax
596 std::vector<ascir::AxisId> node_new_sched_axes = new_sched_axes;729 std::vector<ascir::AxisId> node_new_sched_axes = new_sched_axes;
597 GE_ASSERT_TRUE(!node->outputs.operator()().empty());730 GE_ASSERT_TRUE(!node->outputs.operator()().empty());
598 auto &vectorized_axis = node->outputs[0].attr.vectorized_axis;731 auto &vectorized_axis = node->outputs[0].attr.vectorized_axis;
599- node_new_sched_axes.insert(node_new_sched_axes.end(), vectorized_axis.begin(), vectorized_axis.end());732+ if (!is_indirect_load_schedule_case_) {
733+ node_new_sched_axes.insert(node_new_sched_axes.end(), vectorized_axis.begin(), vectorized_axis.end());
734+ }
600 bool is_store_after_reduce = is_reduce_after && ScheduleUtils::IsStore(node);735 bool is_store_after_reduce = is_reduce_after && ScheduleUtils::IsStore(node);
601 GE_ASSERT_SUCCESS(ApplyBlockSplitToNode(node, is_store_after_reduce));736 GE_ASSERT_SUCCESS(ApplyBlockSplitToNode(node, is_store_after_reduce));
737+ if (is_indirect_load_schedule_case_) {
738+ continue;
739+ }
602 graph_.ApplySchedAxisReorder(node, node_new_sched_axes);740 graph_.ApplySchedAxisReorder(node, node_new_sched_axes);
603 }741 }
604 return af::SUCCESS;742 return af::SUCCESS;
Mautofuse/optimize/autoschedule/schedule.h+11-1
@@ -14,6 +14,7 @@
14#include <utility>14#include <utility>
15 15 
16#include "ascir.h"16#include "ascir.h"
17+#include "indirect_load_utils.h"
17#include "optimize.h"18#include "optimize.h"
18#include "tiling_group.h"19#include "tiling_group.h"
19#include "graph_properties_cache.h"20#include "graph_properties_cache.h"
@@ -66,6 +67,7 @@ class Scheduler {
66 private:67 private:
67 // ub 切分68 // ub 切分
68 Status TileSplit();69 Status TileSplit();
70+ static std::vector<ascir::AxisId> GetSortedNodeVectorizedAxes(Scheduler &scheduler);
69 // block 切分71 // block 切分
70 Status BlockSplit(std::vector<ascir::AxisId> &tile_out_axes);72 Status BlockSplit(std::vector<ascir::AxisId> &tile_out_axes);
71 void FuseTileOutAxes(const std::vector<ascir::AxisId> &non_reduce_outer_axes,73 void FuseTileOutAxes(const std::vector<ascir::AxisId> &non_reduce_outer_axes,
@@ -74,10 +76,12 @@ class Scheduler {
74 const std::vector<ascir::AxisId> &non_reduce_outer_axes,76 const std::vector<ascir::AxisId> &non_reduce_outer_axes,
75 const std::vector<ascir::AxisId> &reduce_outer_axes);77 const std::vector<ascir::AxisId> &reduce_outer_axes);
76 void RemoveDuplicatedAxisFromGroup();78 void RemoveDuplicatedAxisFromGroup();
79+ Status InitIndirectLoadScheduleCase();
80+ Status ApplyIndirectLoadNodeAxes(const af::AscNodePtr &node, bool &skip_main_tiling) const;
77 Status ModifyStoreAfterReduce(ascir::NodeView &node, ascir::AxisId reduce_block_id);81 Status ModifyStoreAfterReduce(ascir::NodeView &node, ascir::AxisId reduce_block_id);
78 Status ApplyBlockSplitToNode(ascir::NodeView &node, bool is_store_after_reduce);82 Status ApplyBlockSplitToNode(ascir::NodeView &node, bool is_store_after_reduce);
79 void TileTiling(ascir::AxisId tile_id, std::pair<af::AxisPtr, af::AxisPtr> &tiled_axes) const {83 void TileTiling(ascir::AxisId tile_id, std::pair<af::AxisPtr, af::AxisPtr> &tiled_axes) const {
80- if (tile_id != kDefaultAxisId) {84+ if (tile_id != kDefaultAxisId && tiled_axes.first == nullptr) {
81 tiled_axes = graph_.TileSplit(tile_id);85 tiled_axes = graph_.TileSplit(tile_id);
82 }86 }
83 }87 }
@@ -96,12 +100,18 @@ class Scheduler {
96 bool HasRGroup() const {100 bool HasRGroup() const {
97 return tiling_case_.ub_tiling_id_r != kDefaultAxisId;101 return tiling_case_.ub_tiling_id_r != kDefaultAxisId;
98 }102 }
103+ 
104+ void AdjustVectorizedAxesOrderOffsets(std::vector<size_t> &vectorized_axes_order, size_t split_point, size_t end,
105+ size_t offset) const;
99 ascir::ImplGraph &graph_;106 ascir::ImplGraph &graph_;
100 AxisGroup axes_group_;107 AxisGroup axes_group_;
101 TilingCase &tiling_case_;108 TilingCase &tiling_case_;
102 bool is_last_axis_reduce_;109 bool is_last_axis_reduce_;
103 optimize::ReduceTemplateType reduce_template_;110 optimize::ReduceTemplateType reduce_template_;
104 ascir::CubeTemplateType cube_template_;111 ascir::CubeTemplateType cube_template_;
112+ bool is_indirect_load_schedule_case_ = false;
113+ bool has_indirect_load_synthetic_outer_axis_ = false;
114+ ascgen_utils::indirect_load::TemplateAxes indirect_load_axes_;
105 GraphPropertiesCache graph_cache_; // 图属性缓存，避免重复遍历115 GraphPropertiesCache graph_cache_; // 图属性缓存，避免重复遍历
106};116};
107} // namespace optimize::autoschedule117} // namespace optimize::autoschedule
Mautofuse/optimize/autoschedule/tiling_group.cpp+3-1
@@ -12,6 +12,7 @@
12#include "ascir_ops_utils.h"12#include "ascir_ops_utils.h"
13#include "ascgraph_info_complete.h"13#include "ascgraph_info_complete.h"
14#include "ascir_utils.h"14#include "ascir_utils.h"
15+#include "indirect_load_utils.h"
15#include "schedule_utils.h"16#include "schedule_utils.h"
16#include "graph/symbolizer/symbolic_utils.h"17#include "graph/symbolizer/symbolic_utils.h"
17 18 
@@ -483,7 +484,8 @@ Status TilingGroup::GenTilingGroup(const ascir::ImplGraph &impl_graph, AxisGroup
483 std::vector<std::pair<std::string, AxisGroup>> node_name_to_tiling_group;484 std::vector<std::pair<std::string, AxisGroup>> node_name_to_tiling_group;
484 std::set<af::AxisId> n_groupset;485 std::set<af::AxisId> n_groupset;
485 for (const auto &node : impl_graph.GetAllNodes()) {486 for (const auto &node : impl_graph.GetAllNodes()) {
486- if (ScheduleUtils::IsBuffer(node)) {487+ if (ScheduleUtils::IsBuffer(node) ||
488+ ascgen_utils::indirect_load::GetTemplateBehavior(node).skips_main_schedule_tiling) {
487 continue;489 continue;
488 }490 }
489 AxisGroup single_node_axes_group;491 AxisGroup single_node_axes_group;
Mautofuse/optimize/buffer_allocate/buf_que_allocator.cpp+6-1
@@ -16,6 +16,7 @@
16#include "schedule_utils.h"16#include "schedule_utils.h"
17#include "graph_utils.h"17#include "graph_utils.h"
18#include "common_utils.h"18#include "common_utils.h"
19+#include "indirect_load_utils.h"
19#include "attribute_group/attr_group_symbolic_desc.h"20#include "attribute_group/attr_group_symbolic_desc.h"
20#include "platform/platform_factory.h"21#include "platform/platform_factory.h"
21#include "mem_reuse_manager.h"22#include "mem_reuse_manager.h"
@@ -382,6 +383,10 @@ Status BufQueAllocator::SetOutputTensorAttr(const af::AscGraph &impl_graph) cons
382}383}
383 384 
384Status BufQueAllocator::GetAndSetNodeTempBuffer(const af::AscNodePtr &node) {385Status BufQueAllocator::GetAndSetNodeTempBuffer(const af::AscNodePtr &node) {
386+ if (ascgen_utils::indirect_load::GetTemplateBehavior(node).skips_api_emit) {
387+ node->attr.tmp_buffers.clear();
388+ return af::SUCCESS;
389+ }
385 auto impl = ascgen_utils::GetAscIrCodegenImpl(node->GetType());390 auto impl = ascgen_utils::GetAscIrCodegenImpl(node->GetType());
386 GE_ASSERT_NOTNULL(impl, "GetAscIrCodegenImpl of node %s[%s] is null", node->GetTypePtr(), node->GetNamePtr());391 GE_ASSERT_NOTNULL(impl, "GetAscIrCodegenImpl of node %s[%s] is null", node->GetTypePtr(), node->GetNamePtr());
387 std::vector<std::unique_ptr<af::TmpBufDesc>> buffers =392 std::vector<std::unique_ptr<af::TmpBufDesc>> buffers =
@@ -401,7 +406,7 @@ Status BufQueAllocator::GetAndSetNodeTempBuffer(const af::AscNodePtr &node) {
401}406}
402 407 
403bool BufQueAllocator::IsTensorUsedByOtherUnit(const af::AscNodePtr &node, const af::AscTensor *output) {408bool BufQueAllocator::IsTensorUsedByOtherUnit(const af::AscNodePtr &node, const af::AscTensor *output) {
404- if (ScheduleUtils::IsLoad(node) || IsOps<Gather>(node)) {409+ if (ScheduleUtils::IsLoad(node) || ScheduleUtils::IsGatherLikeLoad(node)) {
405 return true;410 return true;
406 }411 }
407 for (const auto &input : output->anchor.GetPeerInDataAnchorsPtr()) {412 for (const auto &input : output->anchor.GetPeerInDataAnchorsPtr()) {
Mautofuse/optimize/optimize.cpp+2-1
@@ -24,6 +24,7 @@
24#include "fusion/autofuse_attrs.h"24#include "fusion/autofuse_attrs.h"
25#include "buffer_allocate/buf_que_allocator.h"25#include "buffer_allocate/buf_que_allocator.h"
26#include "ascgraph_info_complete.h"26#include "ascgraph_info_complete.h"
27+#include "indirect_load_utils.h"
27#include "schedule_utils.h"28#include "schedule_utils.h"
28#include "common_utils.h"29#include "common_utils.h"
29#include "node_utils.h"30#include "node_utils.h"
@@ -802,7 +803,7 @@ Status Optimizer::RemoveAllZeroStrideLoopAxis(ascir::ImplGraph &owner_graph) {
802 803 
803Status Optimizer::MergeContinuousAxis(ascir::ImplGraph &impl_graph, ascir::CubeTemplateType cube_type) {804Status Optimizer::MergeContinuousAxis(ascir::ImplGraph &impl_graph, ascir::CubeTemplateType cube_type) {
804 auto all_axis = impl_graph.GetAllAxis();805 auto all_axis = impl_graph.GetAllAxis();
805- if (all_axis.size() <= 1UL) {806+ if (all_axis.size() <= 1UL || ascgen_utils::indirect_load::FindIndirectLoadNode(impl_graph) != nullptr) {
806 return af::SUCCESS;807 return af::SUCCESS;
807 }808 }
808 // concat等场景,会有多套轴, 只能先用循环轴的index来生成潜在连续组, 后续根据连续组会找到多个连续轴809 // concat等场景,会有多套轴, 只能先用循环轴的index来生成潜在连续组, 后续根据连续组会找到多个连续轴
Mautofuse/optimize/platform/common/base_alignment_strategy.cpp+6-2
@@ -11,6 +11,7 @@
11#include "base_alignment_strategy.h"11#include "base_alignment_strategy.h"
12#include "common_utils.h"12#include "common_utils.h"
13#include "graph/symbolizer/symbolic_utils.h"13#include "graph/symbolizer/symbolic_utils.h"
14+#include "indirect_load_utils.h"
14#include "platform/platform_factory.h"15#include "platform/platform_factory.h"
15 16 
16namespace optimize {17namespace optimize {
@@ -301,7 +302,8 @@ af::Status BaseAlignmentStrategy::AddPadForAlignmentConflictNode(ascir::ImplGrap
301 bool inserted = false;302 bool inserted = false;
302 for (const auto &node : impl_graph.GetAllNodes()) {303 for (const auto &node : impl_graph.GetAllNodes()) {
303 GE_ASSERT_NOTNULL(node);304 GE_ASSERT_NOTNULL(node);
304- if (ScheduleUtils::IsBuffer(node)) {305+ const auto indirect_load_behavior = ascgen_utils::indirect_load::GetTemplateBehavior(node);
306+ if (ScheduleUtils::IsBuffer(node) || indirect_load_behavior.uses_direct_gm_pipeline) {
305 continue;307 continue;
306 }308 }
307 309 
@@ -383,7 +385,9 @@ af::Status BaseAlignmentStrategy::AlignVectorizedStrides(ascir::ImplGraph &impl_
383 385 
384 for (const auto &node : impl_graph.GetAllNodes()) {386 for (const auto &node : impl_graph.GetAllNodes()) {
385 GE_ASSERT_NOTNULL(node);387 GE_ASSERT_NOTNULL(node);
386- if (ScheduleUtils::IsBuffer(node)) {388+ const auto indirect_load_behavior = ascgen_utils::indirect_load::GetTemplateBehavior(node);
389+ if (ScheduleUtils::IsBuffer(node) || indirect_load_behavior.uses_direct_gm_pipeline ||
390+ indirect_load_behavior.skips_api_emit) {
387 continue;391 continue;
388 }392 }
389 GE_ASSERT_SUCCESS(SetVectorizedStridesForOneNode(node));393 GE_ASSERT_SUCCESS(SetVectorizedStridesForOneNode(node));
Mautofuse/optimize/schedule_utils.h+6-0
@@ -19,6 +19,7 @@
19#include "ascir.h"19#include "ascir.h"
20#include "ascir_ops.h"20#include "ascir_ops.h"
21#include "common/platform_context.h"21#include "common/platform_context.h"
22+#include "schedule_result.h"
22 23 
23namespace optimize {24namespace optimize {
24class ScheduleUtils {25class ScheduleUtils {
@@ -66,6 +67,11 @@ class ScheduleUtils {
66 return node->attr.api.compute_type == af::ComputeType::kComputeGather;67 return node->attr.api.compute_type == af::ComputeType::kComputeGather;
67 }68 }
68 69 
70+ static bool IsGatherLikeLoad(const af::AscNodePtr &node) {
71+ return af::ops::IsOps<af::ascir_op::Gather>(node) ||
72+ ::ascir::GetTemplateIdOrDefault(*node) == ascir::TemplateId::kIndirectLoadSimd;
73+ }
74+ 
69 static bool IsBuffer(const af::AscNodePtr &node) {75 static bool IsBuffer(const af::AscNodePtr &node) {
70 return node->attr.api.type == af::ApiType::kAPITypeBuffer;76 return node->attr.api.type == af::ApiType::kAPITypeBuffer;
71 }77 }
Aautofuse/optimize/task_generator/indirect_load_schedule_case_generator.cpp+526-0
@@ -0,0 +1,526 @@
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 "task_generator/indirect_load_schedule_case_generator.h"
12+ 
13+#include "ascir_ops.h"
14+#include "ascir_ops_utils.h"
15+#include "common_utils.h"
16+#include "graph_utils.h"
17+#include "graph/symbolizer/symbolic.h"
18+#include "indirect_load_utils.h"
19+#include "schedule_result.h"
20+#include "schedule_utils.h"
21+ 
22+#include <algorithm>
23+#include <sstream>
24+#include <vector>
25+ 
26+namespace optimize {
27+namespace {
28+constexpr int64_t kIndirectLoadSimtDcacheSize = 32 * 1024;
29+using NodePath = std::vector<af::AscNodePtr>;
30+ 
31+struct IndirectLoadGraphPaths {
32+ NodePath data_input;
33+ NodePath index_input;
34+ NodePath output;
35+};
36+ 
37+NodePath CollectInputPath(const af::AscNodePtr &indirect_load, size_t input_index) {
38+ NodePath path;
39+ for (auto current = ascgen_utils::indirect_load::GetInputProducer(indirect_load, input_index); current != nullptr;
40+ current = current->inputs.Size() == 1UL ? ascgen_utils::indirect_load::GetInputProducer(current, 0UL)
41+ : nullptr) {
42+ path.emplace_back(current);
43+ }
44+ return path;
45+}
46+ 
47+NodePath CollectOutputPath(const af::AscNodePtr &indirect_load) {
48+ NodePath path;
49+ for (auto current = ascgen_utils::indirect_load::GetOnlyOutputConsumer(indirect_load); current != nullptr;
50+ current = ascgen_utils::indirect_load::GetOnlyOutputConsumer(current)) {
51+ path.emplace_back(current);
52+ }
53+ return path;
54+}
55+ 
56+IndirectLoadGraphPaths CollectGraphPaths(const af::AscNodePtr &indirect_load) {
57+ return {CollectInputPath(indirect_load, 0UL), CollectInputPath(indirect_load, 1UL), CollectOutputPath(indirect_load)};
58+}
59+ 
60+std::string GenerateScoreFunc(ascir::TemplateId template_id, bool prefer_simd) {
61+ const bool is_simd = template_id == ascir::TemplateId::kIndirectLoadSimd;
62+ std::stringstream ss;
63+ ss << "int32_t CalcScore(AutofuseTilingData &tiling_data) {" << std::endl;
64+ ss << " (void)tiling_data;" << std::endl;
65+ ss << " return " << (is_simd == prefer_simd ? 1 : 0) << ";" << std::endl;
66+ ss << "}" << std::endl;
67+ return ss.str();
68+}
69+ 
70+bool HasExp2(const af::AscGraph &graph) {
71+ for (const auto &node : graph.GetAllNodes()) {
72+ if (af::ops::IsOps<af::ascir_op::Exp2>(node)) {
73+ return true;
74+ }
75+ }
76+ return false;
77+}
78+ 
79+bool IsTemplateCandidateLegal(ascir::TemplateId template_id, const IndirectLoadGraphPaths &paths) {
80+ if (template_id != ascir::TemplateId::kIndirectLoadSimt) {
81+ return true;
82+ }
83+ if (paths.data_input.empty()) {
84+ return false;
85+ }
86+ const auto &input_node = paths.data_input.front();
87+ return af::ops::IsOps<af::ascir_op::Data>(input_node) || af::ops::IsOps<af::ascir_op::Load>(input_node);
88+}
89+ 
90+af::Status GetIndirectLoadAxis(const af::AscNodePtr &node, int64_t &axis) {
91+ GE_ASSERT_NOTNULL(node, "IndirectLoad node is null.");
92+ GE_ASSERT_NOTNULL(node->attr.ir_attr, "IndirectLoad ir attr is null, node = %s", node->GetNamePtr());
93+ const auto *ir_attr = node->attr.ir_attr->DownCastTo<af::ascir_op::IndirectLoad::AscIndirectLoadIrAttrDef>();
94+ GE_ASSERT_NOTNULL(ir_attr, "IndirectLoad ir attr type is invalid, node = %s", node->GetNamePtr());
95+ GE_ASSERT_GRAPH_SUCCESS(ir_attr->GetAxis(axis), "Failed to get IndirectLoad axis, node = %s", node->GetNamePtr());
96+ return af::SUCCESS;
97+}
98+ 
99+af::Status ValidateIndirectLoadInputRank(const af::AscNodePtr &indirect_load, size_t input_idx, size_t output_rank) {
100+ auto input_anchor = indirect_load->GetInDataAnchor(input_idx);
101+ GE_ASSERT_NOTNULL(input_anchor, "IndirectLoad input%zu anchor is null.", input_idx);
102+ auto peer_out_anchor = input_anchor->GetPeerOutAnchor();
103+ GE_ASSERT_NOTNULL(peer_out_anchor, "IndirectLoad input%zu peer anchor is null.", input_idx);
104+ 
105+ auto input_node = std::dynamic_pointer_cast<af::AscNode>(peer_out_anchor->GetOwnerNode());
106+ GE_ASSERT_NOTNULL(input_node, "IndirectLoad input%zu node is invalid.", input_idx);
107+ const size_t output_idx = static_cast<size_t>(peer_out_anchor->GetIdx());
108+ const auto input_outputs = input_node->outputs();
109+ GE_ASSERT_TRUE(output_idx < input_outputs.size(),
110+ "IndirectLoad input%zu output index %zu is out of range, output num:%zu.", input_idx, output_idx,
111+ input_outputs.size());
112+ const auto input_output = input_outputs[output_idx];
113+ GE_ASSERT_NOTNULL(input_output, "IndirectLoad input%zu output tensor is null.", input_idx);
114+ const size_t input_rank = input_output->attr.axis.size();
115+ GE_ASSERT_TRUE(input_rank == output_rank,
116+ "IndirectLoad input%zu rank must equal output rank, input rank:%zu, output rank:%zu.", input_idx,
117+ input_rank, output_rank);
118+ return af::SUCCESS;
119+}
120+ 
121+af::Status ValidateIndirectLoadNode(const af::AscNodePtr &indirect_load) {
122+ GE_ASSERT_NOTNULL(indirect_load, "IndirectLoad node is null.");
123+ const auto outputs = indirect_load->outputs();
124+ GE_ASSERT_TRUE(!outputs.empty(), "IndirectLoad graph is invalid.");
125+ const auto output = outputs[0];
126+ GE_ASSERT_NOTNULL(output, "IndirectLoad output tensor is null.");
127+ const size_t output_rank = output->attr.axis.size();
128+ int64_t axis = 0L;
129+ GE_ASSERT_SUCCESS(GetIndirectLoadAxis(indirect_load, axis));
130+ const int64_t rank = static_cast<int64_t>(output_rank);
131+ GE_ASSERT_TRUE(axis >= -rank && axis < rank, "IndirectLoad axis %ld is out of range for output rank %zu.", axis,
132+ output_rank);
133+ GE_ASSERT_SUCCESS(ValidateIndirectLoadInputRank(indirect_load, 0UL, output_rank));
134+ GE_ASSERT_SUCCESS(ValidateIndirectLoadInputRank(indirect_load, 1UL, output_rank));
135+ return af::SUCCESS;
136+}
137+ 
138+af::Status MergeAxesForTemplate(af::AscGraph &graph, const std::vector<af::AxisId> &axes, const std::string &name,
139+ af::AxisId &merged_axis) {
140+ GE_ASSERT_TRUE(!axes.empty(), "IndirectLoad merge axis source is empty, name:%s.", name.c_str());
141+ merged_axis = axes.size() == 1UL ? axes.front() : graph.MergeAxis(axes, name)->id;
142+ return af::SUCCESS;
143+}
144+ 
145+af::Status CreateFixedTileSplit(af::AscGraph &graph, af::AxisId axis_id, af::AxisId &outer_id, af::AxisId &inner_id) {
146+ const auto *axis = graph.FindAxis(axis_id);
147+ GE_ASSERT_NOTNULL(axis, "IndirectLoad fixed tile axis %ld is not found.", axis_id);
148+ outer_id =
149+ graph.CreateAxis(axis->name + "T", ascir::Axis::Type::kAxisTypeTileOuter, axis->size, {axis_id}, af::kIdNone).id;
150+ inner_id =
151+ graph
152+ .CreateAxis(axis->name + "t", ascir::Axis::Type::kAxisTypeTileInner, af::sym::kSymbolOne, {axis_id}, outer_id)
153+ .id;
154+ auto *outer_axis = graph.FindAxis(outer_id);
155+ GE_ASSERT_NOTNULL(outer_axis, "IndirectLoad fixed tile outer axis %ld is not found.", outer_id);
156+ outer_axis->split_pair_other_id = inner_id;
157+ return af::SUCCESS;
158+}
159+ 
160+af::Status BuildSimdInputInnerAxis(af::AscGraph &graph, const af::AscNodePtr &input_producer, size_t axis_index,
161+ ascir::AxisId &input_inner_axis) {
162+ GE_ASSERT_NOTNULL(input_producer, "IndirectLoad SIMD input0 producer is null.");
163+ GE_ASSERT_TRUE(!input_producer->outputs().empty(), "IndirectLoad SIMD input0 producer has no output.");
164+ const auto input_axes = input_producer->outputs()[0]->attr.axis;
165+ GE_ASSERT_TRUE(axis_index < input_axes.size(), "IndirectLoad SIMD input axis index is out of range.");
166+ std::vector<ascir::AxisId> input_inner_axes(input_axes.begin() + static_cast<int64_t>(axis_index), input_axes.end());
167+ GE_ASSERT_SUCCESS(MergeAxesForTemplate(graph, input_inner_axes, "indirect_load_input_inner", input_inner_axis));
168+ return af::SUCCESS;
169+}
170+af::Status BuildAxisViewByBoundary(af::AscGraph &graph, const std::vector<af::AxisId> &axes, size_t boundary,
171+ af::AxisId &outer_axis, af::AxisId &inner_axis, std::vector<af::AxisId> &outer_axes,
172+ std::vector<af::AxisId> &inner_axes) {
173+ GE_ASSERT_TRUE(!axes.empty(), "IndirectLoad output axis is empty.");
174+ const size_t split = std::min(boundary, axes.size());
175+ outer_axes.assign(axes.begin(), axes.begin() + static_cast<int64_t>(split));
176+ inner_axes.assign(axes.begin() + static_cast<int64_t>(split), axes.end());
177+ if (outer_axes.empty()) {
178+ outer_axis = graph.CreateAxis("indirect_load_single_outer", af::sym::kSymbolOne).id;
179+ } else {
180+ GE_ASSERT_SUCCESS(MergeAxesForTemplate(graph, outer_axes, "indirect_load_outer", outer_axis));
181+ }
182+ if (!inner_axes.empty()) {
183+ GE_ASSERT_SUCCESS(MergeAxesForTemplate(graph, inner_axes, "indirect_load_inner", inner_axis));
184+ }
185+ return af::SUCCESS;
186+}
187+ 
188+af::Status NormalizeAxesForTemplate(af::AscGraph &graph, const af::AscNodePtr &indirect_load, int64_t axis,
189+ int64_t rank, ascir::AxisId input_inner_axis) {
190+ const auto output_axes = indirect_load->outputs()[0]->attr.axis;
191+ GE_ASSERT_TRUE(!output_axes.empty(), "IndirectLoad output axis is empty.");
192+ const size_t boundary = static_cast<size_t>(axis < 0L ? axis + rank : axis);
193+ ascir::AxisId outer_axis = af::kIdNone;
194+ ascir::AxisId inner_axis = af::kIdNone;
195+ std::vector<ascir::AxisId> outer_axes;
196+ std::vector<ascir::AxisId> inner_axes;
197+ GE_ASSERT_SUCCESS(
198+ BuildAxisViewByBoundary(graph, output_axes, boundary, outer_axis, inner_axis, outer_axes, inner_axes));
199+ ascir::AxisId fixed_tile_outer_axis = af::kIdNone;
200+ ascir::AxisId fixed_tile_inner_axis = af::kIdNone;
201+ GE_ASSERT_SUCCESS(CreateFixedTileSplit(graph, outer_axis, fixed_tile_outer_axis, fixed_tile_inner_axis));
202+ GE_ASSERT_SUCCESS(
203+ ascgen_utils::indirect_load::SetTemplateAxes(indirect_load, {outer_axis, inner_axis, input_inner_axis}));
204+ return af::SUCCESS;
205+}
206+ 
207+af::Status NormalizeSimdAxesForTemplate(af::AscGraph &graph, const af::AscNodePtr &indirect_load,
208+ const IndirectLoadGraphPaths &paths) {
209+ const auto output_axes = indirect_load->outputs()[0]->attr.axis;
210+ GE_ASSERT_TRUE(!output_axes.empty(), "IndirectLoad SIMD output axis is empty.");
211+ int64_t axis = 0L;
212+ GE_ASSERT_SUCCESS(GetIndirectLoadAxis(indirect_load, axis));
213+ const int64_t rank = static_cast<int64_t>(output_axes.size());
214+ const size_t axis_index = static_cast<size_t>(axis < 0L ? axis + rank : axis);
215+ ascir::AxisId input_inner_axis = af::kIdNone;
216+ const af::AscNodePtr input_producer = paths.data_input.empty() ? nullptr : paths.data_input.front();
217+ if (input_producer != nullptr && !af::ops::IsOps<af::ascir_op::Data>(input_producer)) {
218+ GE_ASSERT_SUCCESS(BuildSimdInputInnerAxis(graph, input_producer, axis_index, input_inner_axis));
219+ }
220+ return NormalizeAxesForTemplate(graph, indirect_load, axis, rank, input_inner_axis);
221+}
222+ 
223+af::Status NormalizeSimtAxesForTemplate(af::AscGraph &graph, const af::AscNodePtr &indirect_load) {
224+ const auto output_axes = indirect_load->outputs()[0]->attr.axis;
225+ GE_ASSERT_TRUE(!output_axes.empty(), "IndirectLoad SIMT output axis is empty.");
226+ const int64_t rank = static_cast<int64_t>(output_axes.size());
227+ return NormalizeAxesForTemplate(graph, indirect_load, rank, rank, af::kIdNone);
228+}
229+ 
230+bool CanEmitSimtScalar(const af::AscNodePtr &node) {
231+ if (node == nullptr) {
232+ return false;
233+ }
234+ const auto impl = ascgen_utils::GetAscIrCodegenImpl(node->GetType());
235+ return impl != nullptr && impl->IsSimtScalarSupported(*node);
236+}
237+ 
238+bool CollectMovableSimtInputPreNodes(const af::AscNodePtr &indirect_load, const NodePath &data_input, NodePath &nodes) {
239+ af::AscNodePtr consumer = indirect_load;
240+ for (const auto &current : data_input) {
241+ if (af::ops::IsOps<af::ascir_op::Data>(current) || af::ops::IsOps<af::ascir_op::Load>(current)) {
242+ return true;
243+ }
244+ if (current->inputs.Size() != 1UL || current->outputs().empty() || !CanEmitSimtScalar(current) ||
245+ current->GetInControlNodesSize() != 0UL || current->GetOutControlNodesSize() != 0UL ||
246+ ascgen_utils::indirect_load::GetOnlyOutputConsumer(current) != consumer) {
247+ return false;
248+ }
249+ nodes.emplace_back(current);
250+ consumer = current;
251+ }
252+ return false;
253+}
254+ 
255+af::Status MoveSimtInputPreNode(const af::AscNodePtr &node, const af::AscNodePtr &indirect_load) {
256+ const auto producer = ascgen_utils::indirect_load::GetInputProducer(node, 0UL);
257+ GE_ASSERT_NOTNULL(producer);
258+ const auto producer_out = producer->GetOutDataAnchor(0UL);
259+ const auto node_in = node->GetInDataAnchor(0UL);
260+ const auto node_out = node->GetOutDataAnchor(0UL);
261+ const auto indirect_in = indirect_load->GetInDataAnchor(0UL);
262+ const auto indirect_out = indirect_load->GetOutDataAnchor(0UL);
263+ GE_ASSERT_NOTNULL(producer_out);
264+ GE_ASSERT_NOTNULL(node_in);
265+ GE_ASSERT_NOTNULL(node_out);
266+ GE_ASSERT_NOTNULL(indirect_in);
267+ GE_ASSERT_NOTNULL(indirect_out);
268+ GE_ASSERT_GRAPH_SUCCESS(af::GraphUtils::ReplaceEdgeSrc(node_out, indirect_in, producer_out));
269+ GE_ASSERT_GRAPH_SUCCESS(af::GraphUtils::ReplaceEdgeSrc(producer_out, node_in, indirect_out));
270+ const auto peer_inputs = indirect_out->GetPeerInDataAnchors();
271+ for (const auto &peer_in : peer_inputs) {
272+ if (peer_in != node_in) {
273+ GE_ASSERT_GRAPH_SUCCESS(af::GraphUtils::ReplaceEdgeSrc(indirect_out, peer_in, node_out));
274+ }
275+ }
276+ 
277+ const auto indirect_output = indirect_load->outputs()[0];
278+ const auto node_output = node->outputs()[0];
279+ node->attr.sched.axis = indirect_load->attr.sched.axis;
280+ node_output->attr.axis = indirect_output->attr.axis;
281+ node_output->attr.repeats = indirect_output->attr.repeats;
282+ node_output->attr.strides = indirect_output->attr.strides;
283+ indirect_output->attr.dtype = producer->outputs()[0]->attr.dtype;
284+ GELOGD("[IndirectLoad] Move SIMT input pre node[%s] after IndirectLoad[%s].", node->GetNamePtr(),
285+ indirect_load->GetNamePtr());
286+ return af::SUCCESS;
287+}
288+ 
289+af::Status MoveSimtInputPreNodes(const af::AscNodePtr &indirect_load, IndirectLoadGraphPaths &paths) {
290+ NodePath nodes;
291+ if (!CollectMovableSimtInputPreNodes(indirect_load, paths.data_input, nodes) ||
292+ (!nodes.empty() &&
293+ (indirect_load->GetInControlNodesSize() != 0UL || indirect_load->GetOutControlNodesSize() != 0UL)) ||
294+ ascgen_utils::indirect_load::GetOnlyOutputConsumer(indirect_load) == nullptr) {
295+ return af::SUCCESS;
296+ }
297+ for (const auto &node : nodes) {
298+ GE_ASSERT_SUCCESS(MoveSimtInputPreNode(node, indirect_load));
299+ }
300+ paths.data_input.erase(paths.data_input.begin(), paths.data_input.begin() + static_cast<int64_t>(nodes.size()));
301+ paths.output.insert(paths.output.begin(), nodes.rbegin(), nodes.rend());
302+ return af::SUCCESS;
303+}
304+ 
305+af::Status PropagateInputTensorAttrs(const IndirectLoadGraphPaths &paths) {
306+ for (const NodePath *path : {&paths.data_input, &paths.index_input}) {
307+ for (size_t i = 0UL; i < path->size(); ++i) {
308+ const auto &load = (*path)[i];
309+ if (!af::ops::IsOps<af::ascir_op::Load>(load)) {
310+ continue;
311+ }
312+ const af::AscNodePtr data = i + 1UL < path->size() ? (*path)[i + 1UL] : nullptr;
313+ if (data == nullptr || !af::ops::IsOps<af::ascir_op::Data>(data) || load->outputs().empty() ||
314+ data->outputs().empty()) {
315+ break;
316+ }
317+ const auto load_out = load->outputs()[0];
318+ const auto data_out = data->outputs()[0];
319+ if (data_out->attr.axis.empty() && !load_out->attr.axis.empty()) {
320+ GELOGD("[IndirectLoad] Propagate tensor attrs from Load[%s] to Data[%s].", load->GetNamePtr(),
321+ data->GetNamePtr());
322+ data_out->attr.axis = load_out->attr.axis;
323+ data_out->attr.repeats = load_out->attr.repeats;
324+ data_out->attr.strides = load_out->attr.strides;
325+ data_out->attr.dtype = load_out->attr.dtype;
326+ }
327+ break;
328+ }
329+ }
330+ return af::SUCCESS;
331+}
332+ 
333+af::Status RecordTemplateLogicalView(const af::AscNodePtr &indirect_load) {
334+ GE_ASSERT_TRUE(indirect_load->inputs.Size() == 2UL && indirect_load->outputs().size() == 1UL,
335+ "IndirectLoad expects 2 inputs and 1 output.");
336+ ascgen_utils::indirect_load::TemplateLogicalView view;
337+ view.data.axis_ids = indirect_load->inputs()[0]->attr.axis;
338+ view.data.strides = indirect_load->inputs()[0]->attr.strides;
339+ view.index.axis_ids = indirect_load->inputs()[1]->attr.axis;
340+ view.index.strides = indirect_load->inputs()[1]->attr.strides;
341+ view.output.axis_ids = indirect_load->outputs()[0]->attr.axis;
342+ view.output.strides = indirect_load->outputs()[0]->attr.strides;
343+ return ascgen_utils::indirect_load::SetTemplateLogicalView(indirect_load, view);
344+}
345+ 
346+af::Status PrepareCandidateGraph(const af::AscNodePtr &indirect_load, ascir::TemplateId template_id,
347+ IndirectLoadGraphPaths &paths, bool &is_candidate_legal) {
348+ paths = CollectGraphPaths(indirect_load);
349+ if (template_id == ascir::TemplateId::kIndirectLoadSimt) {
350+ GE_ASSERT_SUCCESS(MoveSimtInputPreNodes(indirect_load, paths));
351+ }
352+ is_candidate_legal = IsTemplateCandidateLegal(template_id, paths);
353+ if (!is_candidate_legal) {
354+ return af::SUCCESS;
355+ }
356+ GE_ASSERT_SUCCESS(PropagateInputTensorAttrs(paths));
357+ GE_ASSERT_SUCCESS(RecordTemplateLogicalView(indirect_load));
358+ GE_ASSERT_SUCCESS(::ascir::SetTemplateId(indirect_load, template_id));
359+ return af::SUCCESS;
360+}
361+ 
362+af::Status AnnotateSimdTemplateRoles(const af::AscNodePtr &indirect_load, const IndirectLoadGraphPaths &paths) {
363+ GE_ASSERT_NOTNULL(indirect_load);
364+ GE_ASSERT_SUCCESS(
365+ ascgen_utils::indirect_load::SetTemplateRole(indirect_load, ascgen_utils::indirect_load::TemplateRole::kSimdOp));
366+ 
367+ for (const auto &node : paths.data_input) {
368+ if (af::ops::IsOps<af::ascir_op::Data>(node)) {
369+ break;
370+ }
371+ if (af::ops::IsOps<af::ascir_op::VectorFunc>(node)) {
372+ if (ascgen_utils::indirect_load::GetOnlyOutputConsumer(node).get() == indirect_load.get()) {
373+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(
374+ node, ascgen_utils::indirect_load::TemplateRole::kSimdInputPre));
375+ }
376+ break;
377+ }
378+ if (!ascgen_utils::IsNodeSupportsVectorFunction(node) && !af::ops::IsOps<af::ascir_op::Load>(node)) {
379+ break;
380+ }
381+ GE_ASSERT_SUCCESS(
382+ ascgen_utils::indirect_load::SetTemplateRole(node, ascgen_utils::indirect_load::TemplateRole::kSimdInputPre));
383+ }
384+ return af::SUCCESS;
385+}
386+ 
387+af::Status AnnotateSimtInputRoles(const NodePath &path, bool &is_candidate_legal) {
388+ is_candidate_legal = false;
389+ for (const auto &node : path) {
390+ if (af::ops::IsOps<af::ascir_op::Data>(node)) {
391+ is_candidate_legal = true;
392+ return af::SUCCESS;
393+ }
394+ if (af::ops::IsOps<af::ascir_op::Load>(node)) {
395+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(
396+ node, ascgen_utils::indirect_load::TemplateRole::kSimtInputBoundary));
397+ is_candidate_legal = true;
398+ return af::SUCCESS;
399+ }
400+ return af::SUCCESS;
401+ }
402+ return af::SUCCESS;
403+}
404+ 
405+af::Status AnnotateSimtTransformRoles(const NodePath &path, const std::string &terminal_type,
406+ const std::string &direct_gm_boundary_type, bool &is_candidate_legal) {
407+ is_candidate_legal = false;
408+ for (const auto &node : path) {
409+ if (node->GetType() == direct_gm_boundary_type) {
410+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(
411+ node, ascgen_utils::indirect_load::TemplateRole::kSimtDirectGmBoundary));
412+ is_candidate_legal = true;
413+ return af::SUCCESS;
414+ }
415+ if (node->GetType() == terminal_type) {
416+ is_candidate_legal = true;
417+ return af::SUCCESS;
418+ }
419+ if (af::ops::IsOps<af::ascir_op::VectorFunc>(node) || !CanEmitSimtScalar(node)) {
420+ return af::SUCCESS;
421+ }
422+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(
423+ node, ascgen_utils::indirect_load::TemplateRole::kSimtInlineTransform));
424+ }
425+ return af::SUCCESS;
426+}
427+ 
428+af::Status AnnotateSimtTemplateRoles(const af::AscNodePtr &indirect_load, const IndirectLoadGraphPaths &paths,
429+ bool &is_candidate_legal) {
430+ GE_ASSERT_NOTNULL(indirect_load);
431+ is_candidate_legal = false;
432+ GE_ASSERT_SUCCESS(
433+ ascgen_utils::indirect_load::SetTemplateRole(indirect_load, ascgen_utils::indirect_load::TemplateRole::kSimtOp));
434+ GE_ASSERT_SUCCESS(AnnotateSimtInputRoles(paths.data_input, is_candidate_legal));
435+ if (!is_candidate_legal) {
436+ return af::SUCCESS;
437+ }
438+ GE_ASSERT_SUCCESS(AnnotateSimtTransformRoles(paths.index_input, af::ascir_op::Data::Type, af::ascir_op::Load::Type,
439+ is_candidate_legal));
440+ if (!is_candidate_legal) {
441+ return af::SUCCESS;
442+ }
443+ GE_ASSERT_SUCCESS(AnnotateSimtTransformRoles(paths.output, af::ascir_op::Output::Type, af::ascir_op::Store::Type,
444+ is_candidate_legal));
445+ return af::SUCCESS;
446+}
447+ 
448+af::Status ApplySimdGraphPass(af::AscGraph &graph, const af::AscNodePtr &indirect_load, bool &is_candidate_legal) {
449+ IndirectLoadGraphPaths paths;
450+ GE_ASSERT_SUCCESS(
451+ PrepareCandidateGraph(indirect_load, ascir::TemplateId::kIndirectLoadSimd, paths, is_candidate_legal));
452+ if (!is_candidate_legal) {
453+ return af::SUCCESS;
454+ }
455+ GE_ASSERT_SUCCESS(NormalizeSimdAxesForTemplate(graph, indirect_load, paths));
456+ GE_ASSERT_SUCCESS(AnnotateSimdTemplateRoles(indirect_load, paths));
457+ return af::SUCCESS;
458+}
459+ 
460+af::Status ApplySimtGraphPass(af::AscGraph &graph, const af::AscNodePtr &indirect_load, bool &is_candidate_legal) {
461+ GELOGD("[IndirectLoad] Apply SIMT graph pass for node[%s], dcache_size[%ld].", indirect_load->GetNamePtr(),
462+ kIndirectLoadSimtDcacheSize);
463+ IndirectLoadGraphPaths paths;
464+ GE_ASSERT_SUCCESS(
465+ PrepareCandidateGraph(indirect_load, ascir::TemplateId::kIndirectLoadSimt, paths, is_candidate_legal));
466+ if (!is_candidate_legal) {
467+ return af::SUCCESS;
468+ }
469+ GE_ASSERT_SUCCESS(AnnotateSimtTemplateRoles(indirect_load, paths, is_candidate_legal));
470+ if (!is_candidate_legal) {
471+ return af::SUCCESS;
472+ }
473+ GE_ASSERT_SUCCESS(::ascir::SetDcacheSize(indirect_load, kIndirectLoadSimtDcacheSize));
474+ GE_ASSERT_SUCCESS(NormalizeSimtAxesForTemplate(graph, indirect_load));
475+ return af::SUCCESS;
476+}
477+ 
478+af::Status ApplyGraphPass(af::AscGraph &graph, const af::AscNodePtr &indirect_load, ascir::TemplateId template_id,
479+ bool &is_candidate_legal) {
480+ GELOGD("[IndirectLoad] Apply graph pass for node[%s], template_id[%d].", indirect_load->GetNamePtr(),
481+ static_cast<int32_t>(template_id));
482+ is_candidate_legal = false;
483+ if (template_id == ascir::TemplateId::kIndirectLoadSimd) {
484+ return ApplySimdGraphPass(graph, indirect_load, is_candidate_legal);
485+ }
486+ if (template_id == ascir::TemplateId::kIndirectLoadSimt) {
487+ return ApplySimtGraphPass(graph, indirect_load, is_candidate_legal);
488+ }
489+ return af::SUCCESS;
490+}
491+} // namespace
492+ 
493+Status IndirectLoadScheduleCaseGenerator::Generate(ascir::HintGraph &graph, std::vector<ascir::ImplGraph> &graphs,
494+ std::vector<std::string> &score_functions) {
495+ af::AscNodePtr indirect_load;
496+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::ValidateSingleIndirectLoadNode(graph, indirect_load));
497+ if (indirect_load == nullptr) {
498+ return af::SUCCESS;
499+ }
500+ GE_ASSERT_SUCCESS(ValidateIndirectLoadNode(indirect_load));
501+ GELOGI("[IndirectLoad] Generate schedule candidates for graph[%s], node[%s].", graph.GetName().c_str(),
502+ indirect_load->GetNamePtr());
503+ const std::string indirect_load_name = indirect_load->GetName();
504+ const bool prefer_simd = HasExp2(graph);
505+ for (ascir::TemplateId template_id : {ascir::TemplateId::kIndirectLoadSimd, ascir::TemplateId::kIndirectLoadSimt}) {
506+ ascir::ImplGraph candidate_graph(graph.GetName().c_str());
507+ GE_ASSERT_TRUE(candidate_graph.CopyFrom(graph), "Failed to copy graph [%s].", graph.GetName().c_str());
508+ const af::AscNodePtr candidate_indirect_load = candidate_graph.FindNode(indirect_load_name.c_str());
509+ GE_ASSERT_NOTNULL(candidate_indirect_load, "Failed to find copied IndirectLoad node[%s].",
510+ indirect_load_name.c_str());
511+ bool is_candidate_legal = false;
512+ GE_ASSERT_SUCCESS(ApplyGraphPass(candidate_graph, candidate_indirect_load, template_id, is_candidate_legal));
513+ if (!is_candidate_legal) {
514+ GELOGW("[IndirectLoad] Skip illegal template candidate[%d] for node[%s].", static_cast<int32_t>(template_id),
515+ candidate_indirect_load->GetNamePtr());
516+ continue;
517+ }
518+ graphs.emplace_back(std::move(candidate_graph));
519+ score_functions.emplace_back(GenerateScoreFunc(template_id, prefer_simd));
520+ GELOGI("[IndirectLoad] Add schedule candidate[%d] for node[%s].", static_cast<int32_t>(template_id),
521+ candidate_indirect_load->GetNamePtr());
522+ }
523+ return af::SUCCESS;
524+}
525+ 
526+} // namespace optimize
Aautofuse/optimize/task_generator/indirect_load_schedule_case_generator.h+24-0
@@ -0,0 +1,24 @@
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 ASCGEN_DEV_SRC_OPTIMIZE_TASK_GENERATOR_INDIRECT_LOAD_SCHEDULE_CASE_GENERATOR_H_
12+#define ASCGEN_DEV_SRC_OPTIMIZE_TASK_GENERATOR_INDIRECT_LOAD_SCHEDULE_CASE_GENERATOR_H_
13+ 
14+#include "task_generator/schedule_case_generator.h"
15+ 
16+namespace optimize {
17+class IndirectLoadScheduleCaseGenerator : public FusionCaseGenerator {
18+ public:
19+ Status Generate(ascir::HintGraph &graph, std::vector<ascir::ImplGraph> &graphs,
20+ std::vector<std::string> &score_functions) override;
21+};
22+} // namespace optimize
23+ 
24+#endif // ASCGEN_DEV_SRC_OPTIMIZE_TASK_GENERATOR_INDIRECT_LOAD_SCHEDULE_CASE_GENERATOR_H_
Mautofuse/tests/framework/share_graph/include/share_graph.h+2-0
@@ -112,6 +112,8 @@ struct ShareGraph {
112 static af::ComputeGraphPtr LoadGatherAbsStore(int64_t gather_axis, af::DataType data_type);112 static af::ComputeGraphPtr LoadGatherAbsStore(int64_t gather_axis, af::DataType data_type);
113 static af::ComputeGraphPtr LoadGatherTailAbsStore(int64_t gather_axis, af::DataType data_type);113 static af::ComputeGraphPtr LoadGatherTailAbsStore(int64_t gather_axis, af::DataType data_type);
114 static af::ComputeGraphPtr LoadGatherOneAxisAbsStore(int64_t gather_axis, af::DataType data_type);114 static af::ComputeGraphPtr LoadGatherOneAxisAbsStore(int64_t gather_axis, af::DataType data_type);
115+ static af::ComputeGraphPtr IndirectLoadStoreFusedGraph(size_t rank, int64_t axis, af::DataType data_type,
116+ bool has_input_pre, bool use_exp2);
115 static af::ComputeGraphPtr MatMulFusedGraph(size_t dims_size);117 static af::ComputeGraphPtr MatMulFusedGraph(size_t dims_size);
116 static af::ComputeGraphPtr GatherReduceStore(int64_t gather_axis, af::DataType data_type);118 static af::ComputeGraphPtr GatherReduceStore(int64_t gather_axis, af::DataType data_type);
117 static af::ComputeGraphPtr LoadWhereReduceStoreFusedGraph(size_t dims_size, bool x2_scalar, bool x3_scalar);119 static af::ComputeGraphPtr LoadWhereReduceStoreFusedGraph(size_t dims_size, bool x2_scalar, bool x3_scalar);
Mautofuse/tests/framework/share_graph/src/share_graph.cc+171-0
@@ -16,6 +16,7 @@
16#include "ascendc_ir/utils/asc_graph_utils.h"16#include "ascendc_ir/utils/asc_graph_utils.h"
17#include "ascir_ops.h"17#include "ascir_ops.h"
18#include "ascir_ops_utils.h"18#include "ascir_ops_utils.h"
19+#include "fusion/autofuse_attrs.h"
19 20 
20using namespace af;21using namespace af;
21 22 
@@ -8287,6 +8288,176 @@ af::ComputeGraphPtr ShareGraph::LoadGatherTailAbsStore(int64_t gather_axis, af::
8287 return compute_graph;8288 return compute_graph;
8288}8289}
8289 8290 
8291+template <typename Op>
8292+void SetIndirectLoadNodeView(Op &op, const std::vector<int64_t> &axes, const std::vector<af::Expression> &repeats,
8293+ const std::vector<af::Expression> &strides, af::DataType data_type) {
8294+ op.attr.sched.axis = axes;
8295+ op.y.dtype = data_type;
8296+ *op.y.axis = axes;
8297+ *op.y.repeats = repeats;
8298+ *op.y.strides = strides;
8299+}
8300+ 
8301+static void CreateIndirectLoadAxes(af::AscGraph &graph, size_t rank, size_t offset, std::vector<int64_t> &axes,
8302+ std::vector<af::Expression> &repeats, std::vector<af::Expression> &strides) {
8303+ for (size_t i = 0UL; i < rank; ++i) {
8304+ const size_t dim = offset + i;
8305+ repeats.emplace_back(graph.CreateSizeVar("s" + std::to_string(dim)));
8306+ axes.emplace_back(graph.CreateAxis("z" + std::to_string(dim), repeats.back()).id);
8307+ }
8308+ strides.assign(rank, af::ops::One);
8309+ for (size_t i = rank - 1UL; i > 0UL; --i) {
8310+ strides[i - 1UL] = repeats[i] * strides[i];
8311+ }
8312+}
8313+ 
8314+static void AddIndirectLoadDataChain(af::AscGraph &graph, const std::vector<int64_t> &axes,
8315+ const std::vector<af::Expression> &repeats,
8316+ const std::vector<af::Expression> &strides, af::DataType data_type,
8317+ bool has_input_pre, af::ascir_op::IndirectLoad &indirect_load) {
8318+ af::ascir_op::Data x("x");
8319+ graph.AddNode(x);
8320+ x.y.dtype = data_type;
8321+ x.ir_attr.SetIndex(0);
8322+ af::ascir_op::Load input_load("input_load");
8323+ graph.AddNode(input_load);
8324+ input_load.x = x.y;
8325+ SetIndirectLoadNodeView(input_load, axes, repeats, strides, data_type);
8326+ if (!has_input_pre) {
8327+ indirect_load.x1 = input_load.y;
8328+ return;
8329+ }
8330+ af::ascir_op::Relu input_relu("input_relu");
8331+ graph.AddNode(input_relu);
8332+ input_relu.x = input_load.y;
8333+ SetIndirectLoadNodeView(input_relu, axes, repeats, strides, data_type);
8334+ indirect_load.x1 = input_relu.y;
8335+}
8336+ 
8337+static void AddIndirectLoadIndexChain(af::AscGraph &graph, const std::vector<int64_t> &axes,
8338+ const std::vector<af::Expression> &repeats,
8339+ const std::vector<af::Expression> &strides,
8340+ af::ascir_op::IndirectLoad &indirect_load) {
8341+ af::ascir_op::Data index("index");
8342+ graph.AddNode(index);
8343+ index.y.dtype = af::DT_INT32;
8344+ index.ir_attr.SetIndex(1);
8345+ af::ascir_op::Load index_load("index_load");
8346+ graph.AddNode(index_load);
8347+ index_load.x = index.y;
8348+ SetIndirectLoadNodeView(index_load, axes, repeats, strides, af::DT_INT32);
8349+ af::ascir_op::Abs index_abs("index_abs");
8350+ graph.AddNode(index_abs);
8351+ index_abs.x = index_load.y;
8352+ SetIndirectLoadNodeView(index_abs, axes, repeats, strides, af::DT_INT32);
8353+ indirect_load.x2 = index_abs.y;
8354+}
8355+ 
8356+static void AddIndirectLoadOutputChain(af::AscGraph &graph, const std::vector<int64_t> &axes,
8357+ const std::vector<af::Expression> &repeats,
8358+ const std::vector<af::Expression> &strides, af::DataType data_type,
8359+ bool use_exp2, af::ascir_op::IndirectLoad &indirect_load) {
8360+ af::ascir_op::Neg output_neg("output_neg");
8361+ graph.AddNode(output_neg);
8362+ if (use_exp2) {
8363+ af::ascir_op::Exp2 exp2("output_exp");
8364+ graph.AddNode(exp2);
8365+ exp2.x = indirect_load.y;
8366+ SetIndirectLoadNodeView(exp2, axes, repeats, strides, data_type);
8367+ output_neg.x = exp2.y;
8368+ } else {
8369+ af::ascir_op::Exp exp("output_exp");
8370+ graph.AddNode(exp);
8371+ exp.x = indirect_load.y;
8372+ SetIndirectLoadNodeView(exp, axes, repeats, strides, data_type);
8373+ output_neg.x = exp.y;
8374+ }
8375+ SetIndirectLoadNodeView(output_neg, axes, repeats, strides, data_type);
8376+ 
8377+ af::ascir_op::Store store("store");
8378+ graph.AddNode(store);
8379+ store.x = output_neg.y;
8380+ SetIndirectLoadNodeView(store, axes, repeats, strides, data_type);
8381+ 
8382+ af::ascir_op::Output y("y");
8383+ graph.AddNode(y);
8384+ y.x = store.y;
8385+ SetIndirectLoadNodeView(y, axes, repeats, strides, data_type);
8386+ y.ir_attr.SetIndex(0);
8387+}
8388+ 
8389+static void IndirectLoadStore_BeforeAutofuse(af::AscGraph &graph, size_t rank, int64_t axis, af::DataType data_type,
8390+ bool has_input_pre, bool use_exp2) {
8391+ std::vector<int64_t> input_axes;
8392+ std::vector<af::Expression> input_repeats;
8393+ std::vector<af::Expression> input_strides;
8394+ CreateIndirectLoadAxes(graph, rank, 0UL, input_axes, input_repeats, input_strides);
8395+ std::vector<int64_t> output_axes;
8396+ std::vector<af::Expression> output_repeats;
8397+ std::vector<af::Expression> output_strides;
8398+ CreateIndirectLoadAxes(graph, rank, rank, output_axes, output_repeats, output_strides);
8399+ af::ascir_op::IndirectLoad indirect_load("indirect_load");
8400+ graph.AddNode(indirect_load);
8401+ AddIndirectLoadDataChain(graph, input_axes, input_repeats, input_strides, data_type, has_input_pre, indirect_load);
8402+ AddIndirectLoadIndexChain(graph, output_axes, output_repeats, output_strides, indirect_load);
8403+ indirect_load.ir_attr.SetAxis(axis);
8404+ SetIndirectLoadNodeView(indirect_load, output_axes, output_repeats, output_strides, data_type);
8405+ AddIndirectLoadOutputChain(graph, output_axes, output_repeats, output_strides, data_type, use_exp2, indirect_load);
8406+}
8407+ 
8408+af::ComputeGraphPtr ShareGraph::IndirectLoadStoreFusedGraph(size_t rank, int64_t axis, af::DataType data_type,
8409+ bool has_input_pre, bool use_exp2) {
8410+ if (rank < 2UL || rank > 4UL || axis < -static_cast<int64_t>(rank) || axis >= static_cast<int64_t>(rank)) {
8411+ return nullptr;
8412+ }
8413+ af::AscGraph fused_graph("indirect_load_store_test");
8414+ af::ascir_op::Data data0("data0", fused_graph);
8415+ data0.ir_attr.SetIndex(0);
8416+ af::ascir_op::Data data1("data1", fused_graph);
8417+ data1.ir_attr.SetIndex(1);
8418+ ComputeGraphPtr compute_graph = af::AscGraphUtils::GetComputeGraph(fused_graph);
8419+ if (compute_graph == nullptr) {
8420+ return nullptr;
8421+ }
8422+ 
8423+ auto data_desc = std::make_shared<GeTensorDesc>();
8424+ data_desc->SetDataType(data_type);
8425+ auto index_desc = std::make_shared<GeTensorDesc>();
8426+ index_desc->SetDataType(af::DT_INT32);
8427+ auto backend_desc = std::make_shared<OpDesc>("asc_backend", "AscBackend");
8428+ backend_desc->AddInputDesc(data_desc->Clone());
8429+ backend_desc->AddInputDesc(index_desc->Clone());
8430+ backend_desc->AddOutputDesc(data_desc->Clone());
8431+ auto backend = compute_graph->AddNode(backend_desc);
8432+ if (backend == nullptr) {
8433+ return nullptr;
8434+ }
8435+ 
8436+ auto sub_graph = std::make_shared<af::AscGraph>("indirect_load_store_test");
8437+ IndirectLoadStore_BeforeAutofuse(*sub_graph, rank, axis, data_type, has_input_pre, use_exp2);
8438+ auto fuse_attrs = backend->GetOpDesc()->GetOrCreateAttrsGroup<af::AutoFuseAttrs>();
8439+ if (fuse_attrs == nullptr) {
8440+ return nullptr;
8441+ }
8442+ fuse_attrs->SetAscGraph(sub_graph);
8443+ 
8444+ af::ascir_op::Output output("output");
8445+ output.ir_attr.SetIndex(0);
8446+ auto output_node = compute_graph->AddNode(af::OpDescUtils::GetOpDescFromOperator(output));
8447+ auto data0_node = fused_graph.FindNode("data0");
8448+ auto data1_node = fused_graph.FindNode("data1");
8449+ if (data0_node == nullptr || data1_node == nullptr || output_node == nullptr) {
8450+ return nullptr;
8451+ }
8452+ if (af::GraphUtils::AddEdge(data0_node->GetOutDataAnchor(0), backend->GetInDataAnchor(0)) != ge::GRAPH_SUCCESS ||
8453+ af::GraphUtils::AddEdge(data1_node->GetOutDataAnchor(0), backend->GetInDataAnchor(1)) != ge::GRAPH_SUCCESS ||
8454+ af::GraphUtils::AddEdge(backend->GetOutDataAnchor(0), output_node->GetInDataAnchor(0)) != ge::GRAPH_SUCCESS ||
8455+ compute_graph->TopologicalSorting() != ge::GRAPH_SUCCESS) {
8456+ return nullptr;
8457+ }
8458+ return compute_graph;
8459+}
8460+ 
8290static void ConstructVVAscGraphAxisInfoForOneAxisGather(af::AscGraph &graph, size_t dims_size) {8461static void ConstructVVAscGraphAxisInfoForOneAxisGather(af::AscGraph &graph, size_t dims_size) {
8291 std::vector<int64_t> axis;8462 std::vector<int64_t> axis;
8292 std::vector<af::Expression> repeats;8463 std::vector<af::Expression> repeats;
Mautofuse/tests/st/python/test_python_ascir.py+25-0
@@ -161,6 +161,31 @@ class TestAscir:
161 == "Check dtype failed for cast_0 Cast; input_dtypes: [DT_INT8], output_dytpes: [DT_INT4]"161 == "Check dtype failed for cast_0 Cast; input_dtypes: [DT_INT8], output_dytpes: [DT_INT4]"
162 )162 )
163 163 
164+ @staticmethod
165+ def test_graph_create_node_with_indirect_load_api():
166+ ascir.utils.set_platform("3510", 1, 245760)
167+ try:
168+ graph = ascir.HintGraph("test_indirect_load_api")
169+ s0 = graph.create_size("s0")
170+ s1 = graph.create_size("s1")
171+ s2 = graph.create_size("s2")
172+ z0 = graph.create_axis("z0", s0)
173+ z1 = graph.create_axis("z1", s1)
174+ z2 = graph.create_axis("z2", s2)
175+ 
176+ x = ascir_api.Data(graph, dtype=ascir.dtypes.float32)
177+ x.axis = [z0, z1, z2]
178+ index = ascir_api.Data(graph, dtype=ascir.dtypes.int32)
179+ index.axis = [z0, z1, z2]
180+ y = ascir_api.IndirectLoad(graph, x, index, axis=1, sched_axis=[z0, z1, z2])
181+ 
182+ assert y.dtype == ascir.dtypes.float32
183+ debug_str = ascir.utils.debug_str(graph)
184+ assert "IndirectLoad" in debug_str
185+ assert "axis" in debug_str
186+ finally:
187+ ascir.utils.set_platform("2201", 1, 245760)
188+ 
164 @staticmethod189 @staticmethod
165 def test_graph_create_const_node_with_value_str_attr():190 def test_graph_create_const_node_with_value_str_attr():
166 graph = ascir.HintGraph("test")191 graph = ascir.HintGraph("test")
Mautofuse/tests/ut/att/testcase/gen_model_info/parser/test_ascend_graph_parser.cpp+2-0
@@ -9,6 +9,7 @@
9 */9 */
10#include <iostream>10#include <iostream>
11#include "gtest/gtest.h"11#include "gtest/gtest.h"
12+#include "ascgraph_info_complete.h"
12#include "gen_model_info.h"13#include "gen_model_info.h"
13#include "test_fa_ascir_graph.h"14#include "test_fa_ascir_graph.h"
14#include "base/att_const_values.h"15#include "base/att_const_values.h"
@@ -247,6 +248,7 @@ TEST_F(TestAscendGraphParser, case1) {
247TEST_F(TestAscendGraphParser, test_gather_graph_parse) {248TEST_F(TestAscendGraphParser, test_gather_graph_parse) {
248 af::AscGraph graph1("gather_graph");249 af::AscGraph graph1("gather_graph");
249 ASSERT_EQ(af::ascir::cg::BuildGatherAscendGraphND(graph1), af::SUCCESS);250 ASSERT_EQ(af::ascir::cg::BuildGatherAscendGraphND(graph1), af::SUCCESS);
251+ ASSERT_EQ(optimize::AscGraphInfoComplete::CompleteApiInfo(graph1), af::SUCCESS);
250 att::TuningSpacePtr tuning_space = std::make_shared<att::TuningSpace>();252 att::TuningSpacePtr tuning_space = std::make_shared<att::TuningSpace>();
251 EXPECT_NE(tuning_space, nullptr);253 EXPECT_NE(tuning_space, nullptr);
252 att::AscendGraphParser ascend_graph_parser(tuning_space);254 att::AscendGraphParser ascend_graph_parser(tuning_space);
Mautofuse/tests/ut/common/test_ub_expr_utils.cpp+2-0
@@ -11,6 +11,7 @@
11#include <gtest/gtest.h>11#include <gtest/gtest.h>
12 12 
13#include "ascir_ops.h"13#include "ascir_ops.h"
14+#include "ascgraph_info_complete.h"
14#include "common/ub_expr/asc_graph_ub_expr_builder.h"15#include "common/ub_expr/asc_graph_ub_expr_builder.h"
15#include "common/ub_expr/ub_expr_utils.h"16#include "common/ub_expr/ub_expr_utils.h"
16#include "graph/utils/graph_utils.h"17#include "graph/utils/graph_utils.h"
@@ -331,6 +332,7 @@ TEST(AscGraphUbExprBuilderTest, BuildIncludesReservedUbForGather) {
331 af::Operator op("Gather", "Gather");332 af::Operator op("Gather", "Gather");
332 auto node = graph.AddNode(op);333 auto node = graph.AddNode(op);
333 ASSERT_NE(node, nullptr);334 ASSERT_NE(node, nullptr);
335+ ASSERT_EQ(optimize::AscGraphInfoComplete::CompleteApiInfo(graph), af::SUCCESS);
334 336 
335 ascir::UbExprContext context;337 ascir::UbExprContext context;
336 EXPECT_EQ(ascir::AscGraphUbExprBuilder().Build(graph, context), af::SUCCESS);338 EXPECT_EQ(ascir::AscGraphUbExprBuilder().Build(graph, context), af::SUCCESS);
Aautofuse/tests/ut/optimize/task_generator/test_indirect_load_schedule_case_generator.cpp+442-0
@@ -0,0 +1,442 @@
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+ 
13+#include <algorithm>
14+#include <string>
15+#include <vector>
16+ 
17+#include "ascir_ops.h"
18+#include "common/platform_context.h"
19+#include "graph/debug/ge_attr_define.h"
20+#include "graph/ascendc_ir/utils/asc_graph_utils.h"
21+#include "indirect_load_utils.h"
22+#include "schedule_result.h"
23+#include "task_generator/indirect_load_schedule_case_generator.h"
24+ 
25+namespace {
26+constexpr int64_t kSimtDcacheSize = 32 * 1024;
27+ 
28+std::vector<std::string> AxisNames(af::AscGraph &graph, const std::vector<af::AxisId> &axis_ids) {
29+ std::vector<std::string> names;
30+ names.reserve(axis_ids.size());
31+ for (af::AxisId axis_id : axis_ids) {
32+ const auto *axis = graph.FindAxis(axis_id);
33+ names.push_back(axis == nullptr ? "" : axis->name);
34+ }
35+ return names;
36+}
37+ 
38+af::AxisId FindAxisByName(af::AscGraph &graph, const std::string &name) {
39+ for (const auto &axis : graph.GetAllAxis()) {
40+ if (axis != nullptr && axis->name == name) {
41+ return axis->id;
42+ }
43+ }
44+ return af::kIdNone;
45+}
46+ 
47+void ExpectAxisNames(af::AscGraph &graph, const std::vector<af::AxisId> &axis_ids,
48+ const std::vector<std::string> &expected) {
49+ EXPECT_EQ(AxisNames(graph, axis_ids), expected);
50+}
51+ 
52+void ExpectMergedFrom(af::AscGraph &graph, const std::string &merged_name,
53+ const std::vector<std::string> &expected_from) {
54+ const af::AxisId merged_axis = FindAxisByName(graph, merged_name);
55+ ASSERT_NE(merged_axis, af::kIdNone);
56+ const auto *axis = graph.FindAxis(merged_axis);
57+ ASSERT_NE(axis, nullptr);
58+ EXPECT_EQ(axis->type, af::Axis::Type::kAxisTypeMerged);
59+ ExpectAxisNames(graph, axis->from, expected_from);
60+}
61+ 
62+void ExpectFixedTileSplit(af::AscGraph &graph, const std::string &axis_name) {
63+ const af::AxisId outer_axis_id = FindAxisByName(graph, axis_name);
64+ const af::AxisId tile_outer_axis_id = FindAxisByName(graph, axis_name + "T");
65+ const af::AxisId tile_inner_axis_id = FindAxisByName(graph, axis_name + "t");
66+ ASSERT_NE(outer_axis_id, af::kIdNone);
67+ ASSERT_NE(tile_outer_axis_id, af::kIdNone);
68+ ASSERT_NE(tile_inner_axis_id, af::kIdNone);
69+ 
70+ const auto *tile_outer_axis = graph.FindAxis(tile_outer_axis_id);
71+ const auto *tile_inner_axis = graph.FindAxis(tile_inner_axis_id);
72+ ASSERT_NE(tile_outer_axis, nullptr);
73+ ASSERT_NE(tile_inner_axis, nullptr);
74+ EXPECT_EQ(tile_outer_axis->type, af::Axis::Type::kAxisTypeTileOuter);
75+ EXPECT_EQ(tile_inner_axis->type, af::Axis::Type::kAxisTypeTileInner);
76+ ExpectAxisNames(graph, tile_outer_axis->from, {axis_name});
77+ ExpectAxisNames(graph, tile_inner_axis->from, {axis_name});
78+ EXPECT_EQ(tile_outer_axis->split_pair_other_id, tile_inner_axis_id);
79+ EXPECT_EQ(tile_inner_axis->split_pair_other_id, tile_outer_axis_id);
80+}
81+ 
82+af::AscGraph BuildIndirectLoadGraph(int64_t axis, bool has_input_pre_node = false) {
83+ af::AscGraph graph("indirect_load_ut_graph");
84+ const af::Expression s0 = graph.CreateSizeVar("s0");
85+ const af::Expression s1 = graph.CreateSizeVar("s1");
86+ const af::Expression s2 = graph.CreateSizeVar("s2");
87+ const af::Expression s3 = graph.CreateSizeVar("s3");
88+ const af::Expression s4 = graph.CreateSizeVar("s4");
89+ const af::Expression s5 = graph.CreateSizeVar("s5");
90+ const af::Expression s6 = graph.CreateSizeVar("s6");
91+ const af::Expression s7 = graph.CreateSizeVar("s7");
92+ const auto z0 = graph.CreateAxis("z0", s0);
93+ const auto z1 = graph.CreateAxis("z1", s1);
94+ const auto z2 = graph.CreateAxis("z2", s2);
95+ const auto z3 = graph.CreateAxis("z3", s3);
96+ const auto z4 = graph.CreateAxis("z4", s4);
97+ const auto z5 = graph.CreateAxis("z5", s5);
98+ const auto z6 = graph.CreateAxis("z6", s6);
99+ const auto z7 = graph.CreateAxis("z7", s7);
100+ const std::vector<af::AxisId> input_axes = {z0.id, z1.id, z2.id, z3.id};
101+ const std::vector<af::AxisId> output_axes = {z4.id, z5.id, z6.id, z7.id};
102+ const std::vector<af::Expression> input_repeats = {s0, s1, s2, s3};
103+ const std::vector<af::Expression> output_repeats = {s4, s5, s6, s7};
104+ const std::vector<af::Expression> input_strides = {s1 * s2 * s3, s2 * s3, s3, af::sym::kSymbolOne};
105+ const std::vector<af::Expression> output_strides = {s5 * s6 * s7, s6 * s7, s7, af::sym::kSymbolOne};
106+ 
107+ af::ascir_op::Data x("x", graph);
108+ x.y.dtype = ge::DT_FLOAT16;
109+ x.attr.api.compute_type = af::ComputeType::kComputeInvalid;
110+ x.attr.api.type = af::ApiType::kAPITypeBuffer;
111+ x.ir_attr.SetIndex(0);
112+ x.attr.sched.axis = input_axes;
113+ *x.y.axis = input_axes;
114+ *x.y.repeats = input_repeats;
115+ *x.y.strides = input_strides;
116+ 
117+ af::ascir_op::Data index("index", graph);
118+ index.y.dtype = ge::DT_INT32;
119+ index.attr.api.compute_type = af::ComputeType::kComputeInvalid;
120+ index.attr.api.type = af::ApiType::kAPITypeBuffer;
121+ index.ir_attr.SetIndex(1);
122+ index.attr.sched.axis = output_axes;
123+ *index.y.axis = output_axes;
124+ *index.y.repeats = output_repeats;
125+ *index.y.strides = output_strides;
126+ 
127+ af::ascir_op::IndirectLoad indirect_load("indirect_load");
128+ if (has_input_pre_node) {
129+ af::ascir_op::Abs pre_abs("pre_abs");
130+ pre_abs.x = x.y;
131+ pre_abs.y.dtype = ge::DT_FLOAT16;
132+ pre_abs.attr.sched.axis = input_axes;
133+ *pre_abs.y.axis = input_axes;
134+ *pre_abs.y.repeats = input_repeats;
135+ *pre_abs.y.strides = input_strides;
136+ indirect_load.x1 = pre_abs.y;
137+ } else {
138+ indirect_load.x1 = x.y;
139+ }
140+ 
141+ indirect_load.x2 = index.y;
142+ indirect_load.y.dtype = ge::DT_FLOAT16;
143+ indirect_load.attr.sched.axis = output_axes;
144+ indirect_load.ir_attr.SetAxis(axis);
145+ *indirect_load.y.axis = output_axes;
146+ *indirect_load.y.repeats = output_repeats;
147+ *indirect_load.y.strides = output_strides;
148+ 
149+ af::ascir_op::Output y("y");
150+ y.x = indirect_load.y;
151+ y.y.dtype = ge::DT_FLOAT16;
152+ y.attr.api.compute_type = af::ComputeType::kComputeInvalid;
153+ y.attr.api.type = af::ApiType::kAPITypeBuffer;
154+ y.ir_attr.SetIndex(0);
155+ y.attr.sched.axis = output_axes;
156+ *y.y.axis = output_axes;
157+ *y.y.repeats = output_repeats;
158+ *y.y.strides = output_strides;
159+ 
160+ return graph;
161+}
162+ 
163+template <typename Op>
164+void SetNodeView(Op &op, af::DataType dtype, const std::vector<af::AxisId> &axes,
165+ const std::vector<af::Expression> &repeats, const std::vector<af::Expression> &strides) {
166+ op.y.dtype = dtype;
167+ op.attr.sched.axis = axes;
168+ *op.y.axis = axes;
169+ *op.y.repeats = repeats;
170+ *op.y.strides = strides;
171+}
172+ 
173+af::AscGraph BuildIndirectLoadPrecisionCastGraph() {
174+ af::AscGraph graph("indirect_load_precision_cast_ut_graph");
175+ const af::Expression s0 = graph.CreateSizeVar("s0");
176+ const af::Expression s1 = graph.CreateSizeVar("s1");
177+ const af::Expression s2 = graph.CreateSizeVar("s2");
178+ const af::Expression s3 = graph.CreateSizeVar("s3");
179+ const auto z0 = graph.CreateAxis("z0", s0);
180+ const auto z1 = graph.CreateAxis("z1", s1);
181+ const auto z2 = graph.CreateAxis("z2", s2);
182+ const auto z3 = graph.CreateAxis("z3", s3);
183+ const std::vector<af::AxisId> input_axes = {z0.id, z1.id};
184+ const std::vector<af::AxisId> output_axes = {z2.id, z3.id};
185+ const std::vector<af::Expression> input_repeats = {s0, s1};
186+ const std::vector<af::Expression> output_repeats = {s2, s3};
187+ const std::vector<af::Expression> input_strides = {s1, af::sym::kSymbolOne};
188+ const std::vector<af::Expression> output_strides = {s3, af::sym::kSymbolOne};
189+ 
190+ af::ascir_op::Data x("x", graph);
191+ x.ir_attr.SetIndex(0);
192+ SetNodeView(x, af::DT_FLOAT16, input_axes, input_repeats, input_strides);
193+ af::ascir_op::Load input_load("input_load");
194+ input_load.x = x.y;
195+ SetNodeView(input_load, af::DT_FLOAT16, input_axes, input_repeats, input_strides);
196+ af::ascir_op::Cast input_cast("input_cast");
197+ input_cast.x = input_load.y;
198+ SetNodeView(input_cast, af::DT_FLOAT, input_axes, input_repeats, input_strides);
199+ 
200+ af::ascir_op::Data index("index", graph);
201+ index.ir_attr.SetIndex(1);
202+ SetNodeView(index, af::DT_INT32, output_axes, output_repeats, output_strides);
203+ 
204+ af::ascir_op::IndirectLoad indirect_load("indirect_load");
205+ indirect_load.x1 = input_cast.y;
206+ indirect_load.x2 = index.y;
207+ indirect_load.ir_attr.SetAxis(1);
208+ SetNodeView(indirect_load, af::DT_FLOAT, output_axes, output_repeats, output_strides);
209+ af::ascir_op::Exp output_exp("output_exp");
210+ output_exp.x = indirect_load.y;
211+ SetNodeView(output_exp, af::DT_FLOAT, output_axes, output_repeats, output_strides);
212+ af::ascir_op::Cast output_cast("output_cast");
213+ output_cast.x = output_exp.y;
214+ SetNodeView(output_cast, af::DT_FLOAT16, output_axes, output_repeats, output_strides);
215+ af::ascir_op::Store store("store");
216+ store.x = output_cast.y;
217+ SetNodeView(store, af::DT_FLOAT16, output_axes, output_repeats, output_strides);
218+ af::ascir_op::Output y("y");
219+ y.x = store.y;
220+ y.ir_attr.SetIndex(0);
221+ SetNodeView(y, af::DT_FLOAT16, output_axes, output_repeats, output_strides);
222+ return graph;
223+}
224+ 
225+std::vector<af::AscGraph> GenerateIndirectLoadCases(int64_t axis) {
226+ auto graph = BuildIndirectLoadGraph(axis);
227+ optimize::IndirectLoadScheduleCaseGenerator generator;
228+ std::vector<af::AscGraph> graphs;
229+ std::vector<std::string> score_functions;
230+ EXPECT_EQ(generator.Generate(graph, graphs, score_functions), af::SUCCESS);
231+ EXPECT_EQ(graphs.size(), 2UL);
232+ EXPECT_EQ(score_functions.size(), 2UL);
233+ return graphs;
234+}
235+ 
236+af::AscGraph &FindGeneratedGraphByTemplate(std::vector<af::AscGraph> &graphs, ascir::TemplateId template_id) {
237+ const auto iter = std::find_if(graphs.begin(), graphs.end(), [template_id](const af::AscGraph &graph) {
238+ const auto node = graph.FindNode("indirect_load");
239+ return node != nullptr && ascir::GetTemplateIdOrDefault(*node) == template_id;
240+ });
241+ EXPECT_NE(iter, graphs.end());
242+ return *iter;
243+}
244+ 
245+class IndirectLoadScheduleCaseGeneratorTest : public ::testing::TestWithParam<int64_t> {};
246+ 
247+TEST(IndirectLoadScheduleCaseGeneratorTest, SimtSetsDcacheAndUsesUnifiedVectorizedAxisWithoutReduce) {
248+ auto graphs = GenerateIndirectLoadCases(2);
249+ auto &simt_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimt);
250+ const auto indirect_load = simt_graph.FindNode("indirect_load");
251+ ASSERT_NE(indirect_load, nullptr);
252+ 
253+ EXPECT_EQ(ascir::GetDcacheSize(*indirect_load), kSimtDcacheSize);
254+ ExpectMergedFrom(simt_graph, "indirect_load_outer", {"z4", "z5", "z6", "z7"});
255+ ascgen_utils::indirect_load::TemplateAxes axes;
256+ ASSERT_EQ(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, axes), af::SUCCESS);
257+ ExpectAxisNames(simt_graph, indirect_load->attr.sched.axis, {"z4", "z5", "z6", "z7"});
258+ ExpectAxisNames(simt_graph, {axes.outer_axis}, {"indirect_load_outer"});
259+ EXPECT_TRUE(indirect_load->outputs()[0]->attr.vectorized_axis.empty());
260+ EXPECT_TRUE(indirect_load->outputs()[0]->attr.vectorized_strides.empty());
261+}
262+ 
263+TEST(IndirectLoadScheduleCaseGeneratorTest, StoresTemplateAxesWithoutOverwritingSchedAxis) {
264+ auto graphs = GenerateIndirectLoadCases(2);
265+ auto &simd_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimd);
266+ const auto indirect_load = simd_graph.FindNode("indirect_load");
267+ ASSERT_NE(indirect_load, nullptr);
268+ 
269+ ascgen_utils::indirect_load::TemplateAxes axes;
270+ ASSERT_EQ(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, axes), af::SUCCESS);
271+ ExpectAxisNames(simd_graph, indirect_load->attr.sched.axis, {"z4", "z5", "z6", "z7"});
272+ ExpectAxisNames(simd_graph, {axes.outer_axis}, {"indirect_load_outer"});
273+ ExpectAxisNames(simd_graph, {axes.inner_axis}, {"indirect_load_inner"});
274+ ExpectFixedTileSplit(simd_graph, "indirect_load_outer");
275+ 
276+ ascgen_utils::indirect_load::TemplateLogicalView logical_view;
277+ ASSERT_EQ(ascgen_utils::indirect_load::GetTemplateLogicalView(indirect_load, logical_view), af::SUCCESS);
278+ ExpectAxisNames(simd_graph, logical_view.data.axis_ids, {"z0", "z1", "z2", "z3"});
279+ ExpectAxisNames(simd_graph, logical_view.index.axis_ids, {"z4", "z5", "z6", "z7"});
280+ ExpectAxisNames(simd_graph, logical_view.output.axis_ids, {"z4", "z5", "z6", "z7"});
281+ EXPECT_EQ(logical_view.data.strides.size(), 4UL);
282+ EXPECT_EQ(logical_view.index.strides.size(), 4UL);
283+ EXPECT_EQ(logical_view.output.strides.size(), 4UL);
284+}
285+ 
286+TEST(IndirectLoadScheduleCaseGeneratorTest, DoesNotStoreFixedTileAxesAsTemplateMetadata) {
287+ auto graphs = GenerateIndirectLoadCases(2);
288+ auto &simd_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimd);
289+ const auto indirect_load = simd_graph.FindNode("indirect_load");
290+ ASSERT_NE(indirect_load, nullptr);
291+ const auto op_desc = indirect_load->GetOpDesc();
292+ ASSERT_NE(op_desc, nullptr);
293+ 
294+ EXPECT_EQ(op_desc->TryGetExtAttr("af.internal.indirect_load.tile_outer_axis", static_cast<int64_t>(af::kIdNone)),
295+ af::kIdNone);
296+ EXPECT_EQ(op_desc->TryGetExtAttr("af.internal.indirect_load.tile_inner_axis", static_cast<int64_t>(af::kIdNone)),
297+ af::kIdNone);
298+}
299+ 
300+TEST(IndirectLoadScheduleCaseGeneratorTest, StoresInputVectorizedAxisAsTemplateMetadata) {
301+ auto graph = BuildIndirectLoadGraph(2, true);
302+ optimize::IndirectLoadScheduleCaseGenerator generator;
303+ std::vector<af::AscGraph> graphs;
304+ std::vector<std::string> score_functions;
305+ ASSERT_EQ(generator.Generate(graph, graphs, score_functions), af::SUCCESS);
306+ auto &simd_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimd);
307+ const auto indirect_load = simd_graph.FindNode("indirect_load");
308+ const auto pre_abs = simd_graph.FindNode("pre_abs");
309+ ASSERT_NE(indirect_load, nullptr);
310+ ASSERT_NE(pre_abs, nullptr);
311+ 
312+ ascgen_utils::indirect_load::TemplateAxes axes;
313+ ASSERT_EQ(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, axes), af::SUCCESS);
314+ ExpectAxisNames(simd_graph, pre_abs->attr.sched.axis, {"z0", "z1", "z2", "z3"});
315+ ExpectAxisNames(simd_graph, pre_abs->outputs()[0]->attr.axis, {"z0", "z1", "z2", "z3"});
316+ EXPECT_TRUE(pre_abs->outputs()[0]->attr.vectorized_axis.empty());
317+ ExpectAxisNames(simd_graph, {axes.input_inner_axis}, {"indirect_load_input_inner"});
318+}
319+ 
320+TEST(IndirectLoadScheduleCaseGeneratorTest, GeneratedSimdCandidateKeepsPublicBehavior) {
321+ auto graph = BuildIndirectLoadGraph(2, true);
322+ optimize::IndirectLoadScheduleCaseGenerator generator;
323+ std::vector<af::AscGraph> graphs;
324+ std::vector<std::string> score_functions;
325+ ASSERT_EQ(generator.Generate(graph, graphs, score_functions), af::SUCCESS);
326+ auto &simd_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimd);
327+ const auto indirect_load = simd_graph.FindNode("indirect_load");
328+ ASSERT_NE(indirect_load, nullptr);
329+ 
330+ const auto il_behavior = ascgen_utils::indirect_load::GetTemplateBehavior(indirect_load);
331+ EXPECT_FALSE(il_behavior.skips_main_schedule_tiling);
332+ EXPECT_FALSE(il_behavior.skips_api_emit);
333+ EXPECT_FALSE(il_behavior.uses_direct_gm_pipeline);
334+ EXPECT_FALSE(il_behavior.skips_ub_lifecycle);
335+ EXPECT_FALSE(il_behavior.skips_ub_expr);
336+ EXPECT_FALSE(il_behavior.preserves_vectorized_axis);
337+ EXPECT_FALSE(ascgen_utils::indirect_load::ShouldDisableRegularVectorFunc(indirect_load));
338+}
339+ 
340+TEST(IndirectLoadScheduleCaseGeneratorTest, GeneratedSimtCandidateKeepsPublicBehavior) {
341+ auto graphs = GenerateIndirectLoadCases(2);
342+ auto &simt_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimt);
343+ const auto indirect_load = simt_graph.FindNode("indirect_load");
344+ ASSERT_NE(indirect_load, nullptr);
345+ 
346+ const auto behavior = ascgen_utils::indirect_load::GetTemplateBehavior(indirect_load);
347+ EXPECT_FALSE(behavior.skips_main_schedule_tiling);
348+ EXPECT_FALSE(behavior.skips_api_emit);
349+ EXPECT_TRUE(behavior.uses_direct_gm_pipeline);
350+ EXPECT_TRUE(behavior.skips_ub_lifecycle);
351+ EXPECT_TRUE(behavior.skips_ub_expr);
352+ EXPECT_TRUE(behavior.preserves_vectorized_axis);
353+ EXPECT_FALSE(ascgen_utils::indirect_load::ShouldApplyInputInnerVectorization(indirect_load));
354+ EXPECT_FALSE(ascgen_utils::indirect_load::ShouldSkipMainScheduleTiling(indirect_load));
355+ EXPECT_TRUE(ascgen_utils::indirect_load::ShouldPreserveVectorizedAxis(indirect_load));
356+ EXPECT_TRUE(ascgen_utils::indirect_load::ShouldDisableRegularVectorFunc(indirect_load));
357+}
358+ 
359+TEST(IndirectLoadScheduleCaseGeneratorTest, SimtMovesInputPrecisionCastAfterIndirectLoad) {
360+ auto graph = BuildIndirectLoadPrecisionCastGraph();
361+ optimize::IndirectLoadScheduleCaseGenerator generator;
362+ std::vector<af::AscGraph> graphs;
363+ std::vector<std::string> score_functions;
364+ ge::PlatformContext::GetInstance().SetPlatform("3510");
365+ const auto status = generator.Generate(graph, graphs, score_functions);
366+ ge::PlatformContext::GetInstance().Reset();
367+ ASSERT_EQ(status, af::SUCCESS);
368+ ASSERT_EQ(graphs.size(), 2UL);
369+ 
370+ auto &simd_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimd);
371+ EXPECT_NE(simd_graph.FindNode("input_cast"), nullptr);
372+ EXPECT_NE(simd_graph.FindNode("output_cast"), nullptr);
373+ ASSERT_NE(simd_graph.FindNode("indirect_load"), nullptr);
374+ ASSERT_NE(simd_graph.FindNode("output_exp"), nullptr);
375+ EXPECT_EQ(simd_graph.FindNode("indirect_load")->outputs()[0]->attr.dtype, af::DT_FLOAT);
376+ EXPECT_EQ(simd_graph.FindNode("output_exp")->outputs()[0]->attr.dtype, af::DT_FLOAT);
377+ 
378+ auto &simt_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimt);
379+ const auto input_cast = simt_graph.FindNode("input_cast");
380+ EXPECT_NE(input_cast, nullptr);
381+ EXPECT_NE(simt_graph.FindNode("output_cast"), nullptr);
382+ const auto simt_indirect_load = simt_graph.FindNode("indirect_load");
383+ const auto output_exp = simt_graph.FindNode("output_exp");
384+ ASSERT_NE(simt_indirect_load, nullptr);
385+ ASSERT_NE(output_exp, nullptr);
386+ EXPECT_EQ(simt_indirect_load->outputs()[0]->attr.dtype, af::DT_FLOAT16);
387+ EXPECT_EQ(output_exp->outputs()[0]->attr.dtype, af::DT_FLOAT);
388+ const auto input_producer = ascgen_utils::indirect_load::GetInputProducer(simt_indirect_load, 0UL);
389+ const auto cast_producer = ascgen_utils::indirect_load::GetInputProducer(input_cast, 0UL);
390+ const auto cast_consumer = ascgen_utils::indirect_load::GetOnlyOutputConsumer(input_cast);
391+ ASSERT_NE(input_producer, nullptr);
392+ EXPECT_EQ(input_producer->GetName(), "input_load");
393+ EXPECT_EQ(cast_producer, simt_indirect_load);
394+ EXPECT_EQ(cast_consumer, output_exp);
395+}
396+ 
397+TEST_P(IndirectLoadScheduleCaseGeneratorTest, SimdSplitsOuterAndInnerAxesByNormalizedAxis) {
398+ const int64_t axis = GetParam();
399+ auto graphs = GenerateIndirectLoadCases(axis);
400+ auto &simd_graph = FindGeneratedGraphByTemplate(graphs, ascir::TemplateId::kIndirectLoadSimd);
401+ const auto indirect_load = simd_graph.FindNode("indirect_load");
402+ ASSERT_NE(indirect_load, nullptr);
403+ 
404+ const int64_t output_rank = 4L;
405+ const size_t expected_axis_index = static_cast<size_t>(axis < 0L ? axis + output_rank : axis);
406+ const std::vector<std::string> output_axis_names = {"z4", "z5", "z6", "z7"};
407+ std::vector<std::string> expected_outer;
408+ std::vector<std::string> expected_inner;
409+ expected_outer.assign(output_axis_names.begin(),
410+ output_axis_names.begin() + static_cast<int64_t>(expected_axis_index));
411+ expected_inner.assign(output_axis_names.begin() + static_cast<int64_t>(expected_axis_index), output_axis_names.end());
412+ 
413+ ascgen_utils::indirect_load::TemplateAxes axes;
414+ ASSERT_EQ(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, axes), af::SUCCESS);
415+ ExpectAxisNames(simd_graph, indirect_load->attr.sched.axis, output_axis_names);
416+ if (expected_outer.empty()) {
417+ EXPECT_NE(FindAxisByName(simd_graph, "indirect_load_single_outer"), af::kIdNone);
418+ ExpectAxisNames(simd_graph, {axes.outer_axis}, {"indirect_load_single_outer"});
419+ } else {
420+ ExpectMergedFrom(simd_graph, "indirect_load_outer", expected_outer);
421+ ExpectAxisNames(simd_graph, {axes.outer_axis}, {"indirect_load_outer"});
422+ }
423+ 
424+ if (expected_inner.size() > 1UL) {
425+ ExpectMergedFrom(simd_graph, "indirect_load_inner", expected_inner);
426+ }
427+}
428+ 
429+TEST(IndirectLoadScheduleCaseGeneratorTest, GenerateFailsWhenAxisOutOfRange) {
430+ for (int64_t axis : {-5L, 8L}) {
431+ auto graph = BuildIndirectLoadGraph(axis);
432+ optimize::IndirectLoadScheduleCaseGenerator generator;
433+ std::vector<af::AscGraph> graphs;
434+ std::vector<std::string> score_functions;
435+ EXPECT_NE(generator.Generate(graph, graphs, score_functions), af::SUCCESS) << "axis=" << axis;
436+ EXPECT_TRUE(graphs.empty());
437+ EXPECT_TRUE(score_functions.empty());
438+ }
439+}
440+ 
441+INSTANTIATE_TEST_SUITE_P(AxisBoundary, IndirectLoadScheduleCaseGeneratorTest, ::testing::Values(-1, 0, 2, 3));
442+} // namespace
Mautofuse/tests/ut/optimize/test_optimizer.cpp+32-0
@@ -2200,6 +2200,38 @@ TEST_F(TestOptimizer, MergeAxesGatherOnlyOneDim) {
2200 EXPECT_EQ(new_axis[3]->size, s1 * s2);2200 EXPECT_EQ(new_axis[3]->size, s1 * s2);
2201}2201}
2202 2202 
2203+TEST_F(TestOptimizer, CompleteGatherApiInfoSetsDcacheSize) {
2204+ constexpr int64_t kSimtDcacheSize = 32 * 1024;
2205+ af::AscGraph graph("GatherDcacheGraph");
2206+ 
2207+ af::ascir_op::Data data0("data0", graph);
2208+ data0.y.dtype = ge::DT_FLOAT16;
2209+ data0.ir_attr.SetIndex(0);
2210+ 
2211+ af::ascir_op::Data data1("data1", graph);
2212+ data1.y.dtype = ge::DT_INT32;
2213+ data1.ir_attr.SetIndex(1);
2214+ 
2215+ af::ascir_op::Gather gather("gather");
2216+ gather.x1 = data0.y;
2217+ gather.x2 = data1.y;
2218+ gather.ir_attr.SetAxis(0);
2219+ gather.y.dtype = ge::DT_FLOAT16;
2220+ 
2221+ af::ascir_op::Output output("output");
2222+ output.x = gather.y;
2223+ output.y.dtype = ge::DT_FLOAT16;
2224+ output.ir_attr.SetIndex(0);
2225+ 
2226+ auto gather_node = graph.FindNode("gather");
2227+ ASSERT_NE(gather_node, nullptr);
2228+ gather_node->attr.api.compute_type = af::ComputeType::kComputeGather;
2229+ EXPECT_EQ(::ascir::GetDcacheSize(*gather_node), 0);
2230+ 
2231+ ASSERT_EQ(optimize::AscGraphInfoComplete::CompleteApiInfo(graph), af::SUCCESS);
2232+ EXPECT_EQ(::ascir::GetDcacheSize(*gather_node), kSimtDcacheSize);
2233+}
2234+ 
2203TEST_F(TestOptimizer, MergeAxesReduce) {2235TEST_F(TestOptimizer, MergeAxesReduce) {
2204 af::AscGraph graph("LoadAbsStore");2236 af::AscGraph graph("LoadAbsStore");
2205 auto s0 = graph.CreateSizeVar("s0");2237 auto s0 = graph.CreateSizeVar("s0");
Mautofuse/tests/ut/python/test_python_ascir.py+49-0
@@ -10,10 +10,12 @@
10# -----------------------------------------------------------------------------------------------------------10# -----------------------------------------------------------------------------------------------------------
11 11 
12import pytest12import pytest
13+import importlib.util
13import json14import json
14import time15import time
15import os16import os
16import shutil17import shutil
18+import sys
17from autofuse.pyautofuse import ascir, Autofuser, AutofuserOptions, Schedule, CodeGen19from autofuse.pyautofuse import ascir, Autofuser, AutofuserOptions, Schedule, CodeGen
18 20 
19try:21try:
@@ -161,6 +163,53 @@ class TestAscir:
161 == "Check dtype failed for cast_0 Cast; input_dtypes: [DT_INT8], output_dytpes: [DT_INT4]"163 == "Check dtype failed for cast_0 Cast; input_dtypes: [DT_INT8], output_dytpes: [DT_INT4]"
162 )164 )
163 165 
166+ @staticmethod
167+ def test_graph_create_node_with_indirect_load_api():
168+ import pyautofuse
169+ 
170+ origin_pyautofuse = sys.modules.get("autofuse.pyautofuse")
171+ sys.modules["autofuse.pyautofuse"] = pyautofuse
172+ ascir_mod = pyautofuse.ascir
173+ ascir_mod.utils.set_platform("3510", 1, 245760)
174+ try:
175+ source_path = os.path.abspath(
176+ os.path.join(PYF_PATH, "../../..", "compiler", "python", "ascir_api.py")
177+ )
178+ spec = importlib.util.spec_from_file_location(
179+ "source_ascir_api", source_path
180+ )
181+ source_ascir_api = importlib.util.module_from_spec(spec)
182+ try:
183+ spec.loader.exec_module(source_ascir_api)
184+ finally:
185+ if origin_pyautofuse is None:
186+ sys.modules.pop("autofuse.pyautofuse", None)
187+ else:
188+ sys.modules["autofuse.pyautofuse"] = origin_pyautofuse
189+ 
190+ graph = ascir_mod.HintGraph("test_indirect_load_api")
191+ s0 = graph.create_size("s0")
192+ s1 = graph.create_size("s1")
193+ s2 = graph.create_size("s2")
194+ z0 = graph.create_axis("z0", s0)
195+ z1 = graph.create_axis("z1", s1)
196+ z2 = graph.create_axis("z2", s2)
197+ 
198+ x = source_ascir_api.Data(graph, dtype=ascir_mod.dtypes.float32)
199+ x.axis = [z0, z1, z2]
200+ index = source_ascir_api.Data(graph, dtype=ascir_mod.dtypes.int32)
201+ index.axis = [z0, z1, z2]
202+ y = source_ascir_api.IndirectLoad(
203+ graph, x, index, axis=1, sched_axis=[z0, z1, z2]
204+ )
205+ 
206+ assert y.dtype == ascir_mod.dtypes.float32
207+ debug_str = ascir_mod.utils.debug_str(graph)
208+ assert "IndirectLoad" in debug_str
209+ assert "axis" in debug_str
210+ finally:
211+ ascir_mod.utils.set_platform("2201", 1, 245760)
212+ 
164 @staticmethod213 @staticmethod
165 def test_graph_create_const_node_with_value_str_attr():214 def test_graph_create_const_node_with_value_str_attr():
166 graph = ascir.HintGraph("test")215 graph = ascir.HintGraph("test")
Mautofuse/tests/v35/st/backend_e2e_v2/CMakeLists.txt+1-0
@@ -134,6 +134,7 @@ add_subdirectory(load_logicaland_store_test)
134add_subdirectory(load_gather_split_b_t_abs_store_test)134add_subdirectory(load_gather_split_b_t_abs_store_test)
135add_subdirectory(load_gather_tail_split_b_t_abs_store_test)135add_subdirectory(load_gather_tail_split_b_t_abs_store_test)
136add_subdirectory(load_gather_one_axis_split_b_t_abs_store_test)136add_subdirectory(load_gather_one_axis_split_b_t_abs_store_test)
137+add_subdirectory(indirect_load_store_test)
137# add_subdirectory(split_test)138# add_subdirectory(split_test)
138add_subdirectory(load_where_x2_x3_is_ubscalar_store_test)139add_subdirectory(load_where_x2_x3_is_ubscalar_store_test)
139add_subdirectory(gather_reduce_store_test)140add_subdirectory(gather_reduce_store_test)
Aautofuse/tests/v35/st/backend_e2e_v2/indirect_load_store_test/CMakeLists.txt+41-0
@@ -0,0 +1,41 @@
1+function(add_indirect_load_e2e_case test_name rank axis use_exp2 tiling_key x0 x1 x2 x3 index0 index1 index2 index3)
2+ set(case_workdir ${CMAKE_CURRENT_BINARY_DIR}/${test_name})
3+ file(MAKE_DIRECTORY ${case_workdir})
4+ do_backend_e2e_st_test(${test_name}
5+ WORKDIR ${case_workdir}
6+ CODEGEN indirect_load_store_backend_generator.cpp
7+ TILING_KEY ${tiling_key}
8+ KERNEL_SRC
9+ indirect_load_store_test_kernel.cpp
10+ indirect_load_store_test_tiling.cpp
11+ autofuse_tiling_data.h
12+ TEST_SRC test_e2e_indirect_load_store_kernel.cpp)
13+ 
14+ set(case_definitions
15+ IL_RANK=${rank}
16+ IL_AXIS=${axis}
17+ IL_HAS_INPUT_PRE=1
18+ IL_USE_EXP2=${use_exp2}
19+ IL_TILING_KEY=${tiling_key}
20+ IL_X_S0=${x0}
21+ IL_X_S1=${x1}
22+ IL_X_S2=${x2}
23+ IL_X_S3=${x3}
24+ IL_INDEX_S0=${index0}
25+ IL_INDEX_S1=${index1}
26+ IL_INDEX_S2=${index2}
27+ IL_INDEX_S3=${index3})
28+ target_compile_definitions(${test_name}_codegen_v2 PRIVATE ${case_definitions})
29+ target_compile_definitions(${test_name}_e2e_v2 PRIVATE ${case_definitions})
30+endfunction()
31+ 
32+add_indirect_load_e2e_case(indirect_load_rank2_axis1_simt 2 1 0 1 5 8 1 1 4 6 1 1)
33+add_indirect_load_e2e_case(indirect_load_rank2_axis1_simd 2 1 1 0 5 8 1 1 4 6 1 1)
34+add_indirect_load_e2e_case(indirect_load_rank2_axis0_simt 2 0 0 1 4 9 1 1 6 8 1 1)
35+add_indirect_load_e2e_case(indirect_load_rank2_axis0_simd 2 0 1 0 4 9 1 1 6 8 1 1)
36+add_indirect_load_e2e_case(indirect_load_rank2_axisneg2_simt 2 -2 0 1 4 9 1 1 4 8 1 1)
37+add_indirect_load_e2e_case(indirect_load_rank2_axisneg2_simd 2 -2 1 0 4 9 1 1 4 8 1 1)
38+add_indirect_load_e2e_case(indirect_load_rank3_axis1_simt 3 1 0 1 3 5 8 1 2 7 6 1)
39+add_indirect_load_e2e_case(indirect_load_rank3_axis1_simd 3 1 1 0 3 5 8 1 2 7 6 1)
40+add_indirect_load_e2e_case(indirect_load_rank4_axis2_simt 4 2 0 1 3 3 5 8 2 2 5 6)
41+add_indirect_load_e2e_case(indirect_load_rank4_axis2_simd 4 2 1 0 3 3 5 8 2 2 5 6)
Aautofuse/tests/v35/st/backend_e2e_v2/indirect_load_store_test/indirect_load_store_backend_generator.cpp+205-0
@@ -0,0 +1,205 @@
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 <exception>
12+#include <fstream>
13+#include <gtest/gtest.h>
14+#include <map>
15+#include <string>
16+#include <vector>
17+ 
18+#include "backend_common.h"
19+#include "ascir_ops.h"
20+#include "ascir_ops_utils.h"
21+#include "codegen.h"
22+#include "common/platform_context.h"
23+#include "graph/symbolizer/symbolic_utils.h"
24+#include "indirect_load_utils.h"
25+#include "optimize.h"
26+#include "runtime_stub.h"
27+#include "share_graph.h"
28+ 
29+namespace {
30+constexpr size_t kRank = IL_RANK;
31+constexpr int64_t kAxis = IL_AXIS;
32+constexpr bool kHasInputPre = IL_HAS_INPUT_PRE;
33+constexpr bool kUseExp2 = IL_USE_EXP2;
34+ 
35+const af::Axis *FindDerivedAxis(const af::AscGraph &graph, af::Axis::Type type, af::AxisId from) {
36+ for (const auto &axis : graph.GetAllAxis()) {
37+ if (axis != nullptr && axis->type == type && axis->from == std::vector<af::AxisId>{from}) {
38+ return axis.get();
39+ }
40+ }
41+ return nullptr;
42+}
43+ 
44+void CollectOriginAxes(af::AscGraph &graph, af::AxisId axis_id, std::vector<af::AxisId> &origins) {
45+ const auto *axis = graph.FindAxis(axis_id);
46+ ASSERT_NE(axis, nullptr);
47+ if (axis->from.empty()) {
48+ origins.emplace_back(axis_id);
49+ return;
50+ }
51+ for (af::AxisId from : axis->from) {
52+ CollectOriginAxes(graph, from, origins);
53+ }
54+}
55+ 
56+void ExpectAxisOrigins(af::AscGraph &graph, af::AxisId axis_id, const std::vector<af::AxisId> &expected) {
57+ std::vector<af::AxisId> origins;
58+ CollectOriginAxes(graph, axis_id, origins);
59+ EXPECT_EQ(origins, expected);
60+}
61+ 
62+void ExpectLoopFramework(af::AscGraph &graph, const af::AscNodePtr &indirect_load) {
63+ ascgen_utils::indirect_load::TemplateAxes axes;
64+ ascgen_utils::indirect_load::TemplateLogicalView logical_view;
65+ ASSERT_EQ(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, axes), af::SUCCESS);
66+ ASSERT_EQ(ascgen_utils::indirect_load::GetTemplateLogicalView(indirect_load, logical_view), af::SUCCESS);
67+ 
68+ const auto *outer = graph.FindAxis(axes.outer_axis);
69+ ASSERT_NE(outer, nullptr);
70+ const auto *tile_outer = FindDerivedAxis(graph, af::Axis::Type::kAxisTypeTileOuter, outer->id);
71+ const auto *tile_inner = FindDerivedAxis(graph, af::Axis::Type::kAxisTypeTileInner, outer->id);
72+ ASSERT_NE(tile_outer, nullptr);
73+ ASSERT_NE(tile_inner, nullptr);
74+ EXPECT_EQ(tile_outer->split_pair_other_id, tile_inner->id);
75+ EXPECT_EQ(tile_inner->split_pair_other_id, tile_outer->id);
76+ 
77+ const auto *block_outer = FindDerivedAxis(graph, af::Axis::Type::kAxisTypeBlockOuter, tile_outer->id);
78+ const auto *block_inner = FindDerivedAxis(graph, af::Axis::Type::kAxisTypeBlockInner, tile_outer->id);
79+ ASSERT_NE(block_outer, nullptr);
80+ ASSERT_NE(block_inner, nullptr);
81+ EXPECT_EQ(block_outer->split_pair_other_id, block_inner->id);
82+ EXPECT_EQ(block_inner->split_pair_other_id, block_outer->id);
83+ 
84+ const int64_t normalized_axis = kAxis < 0L ? kAxis + static_cast<int64_t>(kRank) : kAxis;
85+ const auto template_id = ascir::GetTemplateIdOrDefault(*indirect_load);
86+ if (template_id == ascir::TemplateId::kIndirectLoadSimt) {
87+ ExpectAxisOrigins(graph, axes.outer_axis, logical_view.output.axis_ids);
88+ EXPECT_EQ(axes.inner_axis, af::kIdNone);
89+ return;
90+ }
91+ 
92+ ASSERT_EQ(template_id, ascir::TemplateId::kIndirectLoadSimd);
93+ const size_t split = static_cast<size_t>(normalized_axis);
94+ if (split == 0UL) {
95+ EXPECT_TRUE(outer->from.empty());
96+ EXPECT_EQ(af::SymbolicUtils::StaticCheckEq(outer->size, af::ops::One), af::TriBool::kTrue);
97+ } else {
98+ ExpectAxisOrigins(
99+ graph, axes.outer_axis,
100+ std::vector<af::AxisId>(logical_view.output.axis_ids.begin(), logical_view.output.axis_ids.begin() + split));
101+ }
102+ ExpectAxisOrigins(
103+ graph, axes.inner_axis,
104+ std::vector<af::AxisId>(logical_view.output.axis_ids.begin() + split, logical_view.output.axis_ids.end()));
105+ if (kHasInputPre) {
106+ ExpectAxisOrigins(
107+ graph, axes.input_inner_axis,
108+ std::vector<af::AxisId>(logical_view.data.axis_ids.begin() + split, logical_view.data.axis_ids.end()));
109+ }
110+}
111+ 
112+void CheckScheduledLoopFramework(ascir::FusedScheduledResult &result) {
113+ size_t simd_count = 0UL;
114+ size_t simt_count = 0UL;
115+ for (auto &candidates : result.node_idx_to_scheduled_results) {
116+ for (auto &candidate : candidates) {
117+ for (auto &group : candidate.schedule_groups) {
118+ for (auto &graph : group.impl_graphs) {
119+ const af::AscNodePtr indirect_load = ascgen_utils::indirect_load::FindIndirectLoadNode(graph);
120+ if (indirect_load == nullptr) {
121+ continue;
122+ }
123+ ExpectLoopFramework(graph, indirect_load);
124+ if (ascir::GetTemplateIdOrDefault(*indirect_load) == ascir::TemplateId::kIndirectLoadSimd) {
125+ ++simd_count;
126+ } else {
127+ ++simt_count;
128+ }
129+ }
130+ }
131+ }
132+ }
133+ EXPECT_GT(simd_count, 0UL);
134+ EXPECT_GT(simt_count, 0UL);
135+}
136+ 
137+void CheckGeneratedKernel(const std::string &kernel) {
138+ EXPECT_NE(kernel.find("// IndirectLoad SIMD"), std::string::npos);
139+ EXPECT_NE(kernel.find("// IndirectLoad SIMT"), std::string::npos);
140+}
141+} // namespace
142+ 
143+class TestBackendIndirectLoadStoreE2e : public testing::Test {
144+ protected:
145+ void SetUp() override {
146+ dlog_setlevel(ASCGEN_MODULE_NAME, DLOG_ERROR, 0);
147+ ge::PlatformContext::GetInstance().Reset();
148+ ge::RuntimeStub::SetInstance(std::make_shared<af::RuntimeStubV2>());
149+ }
150+ 
151+ void TearDown() override {
152+ dlog_setlevel(ASCGEN_MODULE_NAME, DLOG_ERROR, 0);
153+ ge::RuntimeStub::Reset();
154+ }
155+};
156+ 
157+TEST_F(TestBackendIndirectLoadStoreE2e, IndirectLoadStoreCodegen) {
158+ const std::string tiling_stub = R"(
159+#define REGISTER_TILING_DEFAULT(tiling)
160+#define GET_TILING_DATA(t, tiling) AutofuseTilingData t = *(AutofuseTilingData*)tiling;
161+)";
162+ auto graph = ascir::ShareGraph::IndirectLoadStoreFusedGraph(kRank, kAxis, af::DT_FLOAT16, kHasInputPre, kUseExp2);
163+ ASSERT_NE(graph, nullptr);
164+ std::map<std::string, std::string> shape_info;
165+ for (size_t i = 0UL; i < 2UL * kRank; ++i) {
166+ shape_info.emplace("s" + std::to_string(i), "stub_s" + std::to_string(i));
167+ }
168+ 
169+ const std::vector<std::string> parts = splitString(KERNEL_SRC_LIST, ':');
170+ ASSERT_EQ(parts.size(), 3U);
171+ try {
172+ optimize::Optimizer optimizer(optimize::OptimizerOptions{.graph_type = optimize::GraphType::kFusedAscBackend});
173+ codegen::Codegen codegen(codegen::CodegenOptions{});
174+ ascir::FusedScheduledResult fused_schedule_result;
175+ codegen::CodegenResult result;
176+ testing::internal::CaptureStdout();
177+ const auto optimize_status = optimizer.Optimize(graph, fused_schedule_result);
178+ if (optimize_status == 0) {
179+ CheckScheduledLoopFramework(fused_schedule_result);
180+ }
181+ const auto codegen_status = optimize_status == 0 ? codegen.Generate(shape_info, fused_schedule_result, result) : -1;
182+ const std::string logs = testing::internal::GetCapturedStdout();
183+ EXPECT_EQ(logs.find("[ERROR]"), std::string::npos) << logs;
184+ ASSERT_EQ(optimize_status, 0) << logs;
185+ ASSERT_EQ(codegen_status, 0) << logs;
186+ CheckGeneratedKernel(result.kernel);
187+ 
188+ std::fstream kernel_file(parts[0], std::ios::out);
189+ std::fstream tiling_file(parts[1], std::ios::out);
190+ std::fstream tiling_data_file(parts[2], std::ios::out);
191+ ASSERT_TRUE(kernel_file.is_open());
192+ ASSERT_TRUE(tiling_file.is_open());
193+ ASSERT_TRUE(tiling_data_file.is_open());
194+ kernel_file << tiling_stub << RemoveSubDirInclude(result.kernel);
195+ tiling_file << result.tiling;
196+ tiling_data_file << result.tiling_data;
197+ EXPECT_TRUE(kernel_file.good());
198+ EXPECT_TRUE(tiling_file.good());
199+ EXPECT_TRUE(tiling_data_file.good());
200+ } catch (const std::exception &e) {
201+ FAIL() << e.what();
202+ } catch (...) {
203+ FAIL() << "Unknown exception";
204+ }
205+}
Aautofuse/tests/v35/st/backend_e2e_v2/indirect_load_store_test/test_e2e_indirect_load_store_kernel.cpp+132-0
@@ -0,0 +1,132 @@
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 <algorithm>
12+#include <array>
13+#include <cmath>
14+#include <cstdint>
15+#include <memory>
16+#include <vector>
17+ 
18+#include <gtest/gtest.h>
19+#include "tikicpulib.h"
20+ 
21+#include "autofuse_tiling_data.h"
22+ 
23+extern "C" __global__ __aicore__ void indirect_load_store_test(GM_ADDR x, GM_ADDR index, GM_ADDR y, GM_ADDR workspace,
24+ GM_ADDR tiling);
25+#if IL_RANK == 2
26+extern "C" int64_t AutofuseTiling(uint32_t s0, uint32_t s1, uint32_t s2, uint32_t s3, AutofuseTilingData *, uint32_t *,
27+ uint32_t *, uint32_t, uint32_t);
28+constexpr std::array<int32_t, 2> kInputShape = {IL_X_S0, IL_X_S1};
29+constexpr std::array<int32_t, 2> kIndexShape = {IL_INDEX_S0, IL_INDEX_S1};
30+#elif IL_RANK == 3
31+extern "C" int64_t AutofuseTiling(uint32_t s0, uint32_t s1, uint32_t s2, uint32_t s3, uint32_t s4, uint32_t s5,
32+ AutofuseTilingData *, uint32_t *, uint32_t *, uint32_t, uint32_t);
33+constexpr std::array<int32_t, 3> kInputShape = {IL_X_S0, IL_X_S1, IL_X_S2};
34+constexpr std::array<int32_t, 3> kIndexShape = {IL_INDEX_S0, IL_INDEX_S1, IL_INDEX_S2};
35+#elif IL_RANK == 4
36+extern "C" int64_t AutofuseTiling(uint32_t s0, uint32_t s1, uint32_t s2, uint32_t s3, uint32_t s4, uint32_t s5,
37+ uint32_t s6, uint32_t s7, AutofuseTilingData *, uint32_t *, uint32_t *, uint32_t,
38+ uint32_t);
39+constexpr std::array<int32_t, 4> kInputShape = {IL_X_S0, IL_X_S1, IL_X_S2, IL_X_S3};
40+constexpr std::array<int32_t, 4> kIndexShape = {IL_INDEX_S0, IL_INDEX_S1, IL_INDEX_S2, IL_INDEX_S3};
41+#endif
42+ 
43+namespace {
44+constexpr int32_t kAxis = IL_AXIS < 0 ? IL_AXIS + IL_RANK : IL_AXIS;
45+ 
46+void RunTiling(AutofuseTilingData &tiling, uint32_t &workspace_size, uint32_t &block_dim) {
47+#if IL_RANK == 2
48+ AutofuseTiling(kInputShape[0], kInputShape[1], kIndexShape[0], kIndexShape[1], &tiling, &workspace_size, &block_dim,
49+ 48, 192 * 1024);
50+#elif IL_RANK == 3
51+ AutofuseTiling(kInputShape[0], kInputShape[1], kInputShape[2], kIndexShape[0], kIndexShape[1], kIndexShape[2],
52+ &tiling, &workspace_size, &block_dim, 48, 192 * 1024);
53+#elif IL_RANK == 4
54+ AutofuseTiling(kInputShape[0], kInputShape[1], kInputShape[2], kInputShape[3], kIndexShape[0], kIndexShape[1],
55+ kIndexShape[2], kIndexShape[3], &tiling, &workspace_size, &block_dim, 48, 192 * 1024);
56+#endif
57+}
58+ 
59+template <size_t N>
60+int32_t ElementCount(const std::array<int32_t, N> &shape) {
61+ int32_t count = 1;
62+ for (int32_t dim : shape) {
63+ count *= dim;
64+ }
65+ return count;
66+}
67+ 
68+void InitializeData(half *x, int32_t *index, half *expected, int32_t input_count, int32_t output_count) {
69+ std::array<int32_t, IL_RANK> input_strides{};
70+ input_strides.back() = 1;
71+ for (size_t i = input_strides.size() - 1UL; i > 0UL; --i) {
72+ input_strides[i - 1UL] = kInputShape[i] * input_strides[i];
73+ }
74+ for (int32_t i = 0; i < input_count; ++i) {
75+ x[i] = static_cast<half>(static_cast<float>((i % 29) - 14) * 0.25F);
76+ }
77+ for (int32_t i = 0; i < output_count; ++i) {
78+ const int32_t index_value = (i * 3 + 1) % kInputShape[kAxis];
79+ index[i] = i % 2 == 0 ? index_value : -index_value;
80+ int32_t remaining = i;
81+ int32_t input_offset = 0;
82+ for (size_t dim = kIndexShape.size(); dim-- > 0UL;) {
83+ const int32_t coord = remaining % kIndexShape[dim];
84+ remaining /= kIndexShape[dim];
85+ input_offset += (static_cast<int32_t>(dim) == kAxis ? index_value : coord) * input_strides[dim];
86+ }
87+ float value = static_cast<float>(x[input_offset]);
88+#if IL_HAS_INPUT_PRE
89+ value = std::max(value, 0.0F);
90+#endif
91+#if IL_USE_EXP2
92+ expected[i] = static_cast<half>(-std::exp2(value));
93+#else
94+ expected[i] = static_cast<half>(-std::exp(value));
95+#endif
96+ }
97+}
98+} // namespace
99+ 
100+TEST(E2EIndirectLoadStore, GeneratedKernelMatchesReference) {
101+ const int32_t input_count = ElementCount(kInputShape);
102+ const int32_t output_count = ElementCount(kIndexShape);
103+ const auto gm_free = [](void *ptr) { AscendC::GmFree(ptr); };
104+ std::unique_ptr<half, decltype(gm_free)> x(reinterpret_cast<half *>(AscendC::GmAlloc(input_count * sizeof(half))),
105+ gm_free);
106+ std::unique_ptr<int32_t, decltype(gm_free)> index(
107+ reinterpret_cast<int32_t *>(AscendC::GmAlloc(output_count * sizeof(int32_t))), gm_free);
108+ std::unique_ptr<half, decltype(gm_free)> output(
109+ reinterpret_cast<half *>(AscendC::GmAlloc(output_count * sizeof(half))), gm_free);
110+ ASSERT_NE(x, nullptr);
111+ ASSERT_NE(index, nullptr);
112+ ASSERT_NE(output, nullptr);
113+ std::vector<half> expected(static_cast<size_t>(output_count));
114+ InitializeData(x.get(), index.get(), expected.data(), input_count, output_count);
115+ std::fill_n(output.get(), output_count, static_cast<half>(0.0F));
116+ 
117+ AutofuseTilingData tiling_data{};
118+ uint32_t workspace_size = 0U;
119+ uint32_t block_dim = 48U;
120+ RunTiling(tiling_data, workspace_size, block_dim);
121+ ASSERT_EQ(tiling_data.graph0_tiling_key, IL_TILING_KEY);
122+ ASSERT_GT(tiling_data.block_dim, 0U);
123+ 
124+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
125+ ICPU_RUN_KF(indirect_load_store_test, tiling_data.block_dim, reinterpret_cast<uint8_t *>(x.get()),
126+ reinterpret_cast<uint8_t *>(index.get()), reinterpret_cast<uint8_t *>(output.get()), nullptr,
127+ reinterpret_cast<uint8_t *>(&tiling_data));
128+ for (int32_t i = 0; i < output_count; ++i) {
129+ EXPECT_NEAR(static_cast<float>(output.get()[i]), static_cast<float>(expected[static_cast<size_t>(i)]), 0.0625F)
130+ << "offset=" << i;
131+ }
132+}
Mautofuse/v35/ascir/generator/ascir_builtin_ops_v2.cpp+15-0
@@ -929,6 +929,21 @@ REG_ASC_IR(Gather).Impl(v2_soc_versions,
929 {{"T1", TensorType{DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_FLOAT16, DT_BF16, DT_FLOAT}},929 {{"T1", TensorType{DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_FLOAT16, DT_BF16, DT_FLOAT}},
930 {"T2", TensorType{DT_INT32, DT_INT64}}}});930 {"T2", TensorType{DT_INT32, DT_INT64}}}});
931 931 
932+REG_ASC_IR(IndirectLoad)
933+ .Input("x1", "T1")
934+ .Input("x2", "T2")
935+ .Output("y", "T1")
936+ .Attr<int64_t>("axis")
937+ .Attr<bool>("negative_index_support")
938+ .Attr<bool>("need_check_bound")
939+ .Attr<int64_t>("max")
940+ .ComputeType(ComputeType::kComputeLoad)
941+ .Impl(v2_soc_versions,
942+ {af::ascir::AscIrImplCreator<af::ascir::IndirectLoadAscIrAttImplV2>(),
943+ af::ascir::AscIrImplCreator<af::ascir::IndirectLoadAscIrCodegenImplV2>(),
944+ {{"T1", TensorType{DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_FLOAT16, DT_BF16, DT_FLOAT}},
945+ {"T2", TensorType{DT_INT32, DT_INT64}}}});
Z
Zzhang_shengjie7月27日

注册了： negative_index_support need_check_bound max 但：

  • ascir_api.py:1840-1861 的高层 Python API 只接受 axis。
  • reg_indirect_load_api_call.cpp:386-397 只解析 axis。
  • SIMT 固定对负索引和越界索引写 0。
  • SIMD 不做负索引归一化和越界检查。
  • max 完全未使用。 例如设置 negative_index_support=true，索引 -1 应访问最后一个元素：
  • SIMT 当前直接写 0。
  • SIMD 会把负 src_idx 转成 uint32_t，形成巨大 offset，存在越界访问风险。 设置 need_check_bound=true 时，SIMD 同样没有边界检查。 建议二选一：
  • 完整解析并实现三个属性，同时补充 Python API 参数和两条后端的一致语义；
  • 当前版本不准备支持时，删除属性注册和访问器，避免形成虚假的公开接口契约。
likedislike
946+ 
932REG_ASC_IR(Transpose).Impl(v2_soc_versions, {af::ascir::AscIrImplCreator<af::ascir::TransposeAscIrAttImplV2>(),947REG_ASC_IR(Transpose).Impl(v2_soc_versions, {af::ascir::AscIrImplCreator<af::ascir::TransposeAscIrAttImplV2>(),
933 af::ascir::AscIrImplCreator<af::ascir::TransposeAscIrCodegenImplV2>(),948 af::ascir::AscIrImplCreator<af::ascir::TransposeAscIrCodegenImplV2>(),
934 {{"T", TensorType{DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_FLOAT16,949 {{"T", TensorType{DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_FLOAT16,
Mautofuse/v35/ascir/generator/v2_ascir_att_impl.h+1-0
@@ -35,6 +35,7 @@ namespace ascir {
35 }35 }
36REG_ASC_IR_ATT_V2_CLASS_DEFINE(Add);36REG_ASC_IR_ATT_V2_CLASS_DEFINE(Add);
37REG_ASC_IR_ATT_V2_CLASS_DEFINE(Gather);37REG_ASC_IR_ATT_V2_CLASS_DEFINE(Gather);
38+REG_ASC_IR_ATT_V2_CLASS_DEFINE(IndirectLoad);
38REG_ASC_IR_ATT_V2_CLASS_DEFINE(Abs);39REG_ASC_IR_ATT_V2_CLASS_DEFINE(Abs);
39REG_ASC_IR_ATT_V2_CLASS_DEFINE(Broadcast);40REG_ASC_IR_ATT_V2_CLASS_DEFINE(Broadcast);
40REG_ASC_IR_ATT_V2_CLASS_DEFINE(Cast);41REG_ASC_IR_ATT_V2_CLASS_DEFINE(Cast);
Mautofuse/v35/ascir/generator/v2_ascir_codegen_impl.h+224-0
@@ -18,6 +18,7 @@
18#include "reg_func/default_reg_func_v2.h"18#include "reg_func/default_reg_func_v2.h"
19#include "symbolizer/symbolic_utils.h"19#include "symbolizer/symbolic_utils.h"
20#include "ascir_codegen_v2.h"20#include "ascir_codegen_v2.h"
21+#include "schedule_result.h"
21#include "generator/ascir_common.h"22#include "generator/ascir_common.h"
22 23 
23namespace af {24namespace af {
@@ -281,6 +282,20 @@ class CastAscIrCodegenImplV2 : public AscIrCodegenV2 {
281 }282 }
282 return false;283 return false;
283 }284 }
285+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
286+ return IsVectorFunctionSupported(node);
287+ }
288+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr(const AscNode &node, const std::vector<std::string> &inputs,
289+ std::string &expr) const override {
290+ auto outputs = const_cast<AscNode &>(node).outputs();
291+ GE_ASSERT_TRUE(outputs.size() == 1UL, "Cast node %s[%s] must have one output", node.GetTypePtr(),
292+ node.GetNamePtr());
293+ const char *output_dtype = GetSimtDtypeName(outputs[0]->attr.dtype);
294+ GE_ASSERT_NOTNULL(output_dtype, "Cast node %s[%s] has unsupported SIMT output dtype %d", node.GetTypePtr(),
295+ node.GetNamePtr(), static_cast<int32_t>(outputs[0]->attr.dtype));
296+ expr = std::string("static_cast<") + output_dtype + ">(" + inputs[0] + ")";
297+ return ge::GRAPH_SUCCESS;
298+ }
284 [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {299 [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {
285 return {300 return {
286 "basic_api/reg_compute/kernel_reg_compute_intf.h",301 "basic_api/reg_compute/kernel_reg_compute_intf.h",
@@ -291,6 +306,20 @@ class CastAscIrCodegenImplV2 : public AscIrCodegenV2 {
291 node.GetTypePtr(), node.GetNamePtr());306 node.GetTypePtr(), node.GetNamePtr());
292 return true;307 return true;
293 }308 }
309+ 
310+ private:
311+ static const char *GetSimtDtypeName(ge::DataType dtype) {
312+ static constexpr std::pair<ge::DataType, const char *> kDtypeNames[] = {
313+ {DT_FLOAT, "float"}, {DT_FLOAT16, "half"}, {DT_BF16, "bfloat16_t"}, {DT_INT8, "int8_t"},
314+ {DT_UINT8, "uint8_t"}, {DT_INT16, "int16_t"}, {DT_INT32, "int32_t"}, {DT_INT64, "int64_t"},
315+ };
316+ for (const auto &entry : kDtypeNames) {
317+ if (entry.first == dtype) {
318+ return entry.second;
319+ }
320+ }
321+ return nullptr;
322+ }
294};323};
295 324 
296class AbsAscIrCodegenImplV2 : public AscIrCodegenV2 {325class AbsAscIrCodegenImplV2 : public AscIrCodegenV2 {
@@ -315,6 +344,18 @@ class AbsAscIrCodegenImplV2 : public AscIrCodegenV2 {
315 return true;344 return true;
316 }345 }
317 346 
347+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
348+ (void)node;
349+ return true;
350+ }
351+ 
352+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr([[maybe_unused]] const AscNode &node,
353+ const std::vector<std::string> &inputs,
354+ std::string &expr) const override {
355+ expr = "AscendC::Simt::Abs(" + inputs[0] + ")";
356+ return ge::GRAPH_SUCCESS;
357+ }
358+ 
318 [[nodiscard]] bool IsInplaceSupported(const AscNode &abs_node) const override {359 [[nodiscard]] bool IsInplaceSupported(const AscNode &abs_node) const override {
319 (void)abs_node;360 (void)abs_node;
320 return true;361 return true;
@@ -384,6 +425,23 @@ class ExpAscIrCodegenImplV2 : public AscIrCodegenV2 {
384 node.GetTypePtr(), node.GetNamePtr());425 node.GetTypePtr(), node.GetNamePtr());
385 return true;426 return true;
386 }427 }
428+ 
429+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
430+ auto &outputs = const_cast<AscNode &>(node).outputs;
431+ auto dt = static_cast<ge::DataType>(outputs[0].attr.dtype);
432+ return dt == DT_FLOAT16 || dt == DT_FLOAT || dt == DT_BF16;
433+ }
434+ 
435+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr(const AscNode &node, const std::vector<std::string> &inputs,
436+ std::string &expr) const override {
437+ auto &outputs = const_cast<AscNode &>(node).outputs;
438+ if (outputs[0].attr.dtype == DT_FLOAT) {
439+ expr = "AscendC::Simt::Exp(" + inputs[0] + ")";
440+ } else {
441+ expr = "AscendC::Simt::Exp(static_cast<float>(" + inputs[0] + "))";
442+ }
443+ return ge::GRAPH_SUCCESS;
444+ }
387};445};
388 446 
389class Exp2AscIrCodegenImplV2 : public AscIrCodegenV2 {447class Exp2AscIrCodegenImplV2 : public AscIrCodegenV2 {
@@ -412,6 +470,23 @@ class Exp2AscIrCodegenImplV2 : public AscIrCodegenV2 {
412 node.GetTypePtr(), node.GetNamePtr());470 node.GetTypePtr(), node.GetNamePtr());
413 return true;471 return true;
414 }472 }
473+ 
474+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
475+ auto &outputs = const_cast<AscNode &>(node).outputs;
476+ auto dt = static_cast<ge::DataType>(outputs[0].attr.dtype);
477+ return dt == DT_FLOAT16 || dt == DT_FLOAT || dt == DT_BF16;
478+ }
479+ 
480+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr(const AscNode &node, const std::vector<std::string> &inputs,
481+ std::string &expr) const override {
482+ auto &outputs = const_cast<AscNode &>(node).outputs;
483+ if (outputs[0].attr.dtype == DT_FLOAT) {
484+ expr = "AscendC::Simt::Exp2(" + inputs[0] + ")";
485+ } else {
486+ expr = "AscendC::Simt::Exp2(static_cast<float>(" + inputs[0] + "))";
487+ }
488+ return ge::GRAPH_SUCCESS;
489+ }
415};490};
416 491 
417class FloorAscIrCodegenImplV2 : public AscIrCodegenV2 {492class FloorAscIrCodegenImplV2 : public AscIrCodegenV2 {
@@ -1223,6 +1298,23 @@ class LnAscIrCodegenImplV2 : public AscIrCodegenV2 {
1223 node.GetTypePtr(), node.GetNamePtr());1298 node.GetTypePtr(), node.GetNamePtr());
1224 return true;1299 return true;
1225 }1300 }
1301+ 
1302+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
1303+ auto &outputs = const_cast<AscNode &>(node).outputs;
1304+ auto dt = static_cast<ge::DataType>(outputs[0].attr.dtype);
1305+ return dt == DT_FLOAT16 || dt == DT_FLOAT || dt == DT_BF16;
1306+ }
1307+ 
1308+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr(const AscNode &node, const std::vector<std::string> &inputs,
1309+ std::string &expr) const override {
1310+ auto &outputs = const_cast<AscNode &>(node).outputs;
1311+ if (outputs[0].attr.dtype == DT_FLOAT) {
1312+ expr = "AscendC::Simt::Log(" + inputs[0] + ")";
1313+ } else {
1314+ expr = "AscendC::Simt::Log(static_cast<float>(" + inputs[0] + "))";
1315+ }
1316+ return ge::GRAPH_SUCCESS;
1317+ }
1226};1318};
1227 1319 
1228class ExpmAscIrCodegenImplV2 : public AscIrCodegenV2 {1320class ExpmAscIrCodegenImplV2 : public AscIrCodegenV2 {
@@ -1390,6 +1482,23 @@ class SqrtAscIrCodegenImplV2 : public AscIrCodegenV2 {
1390 node.GetTypePtr(), node.GetNamePtr());1482 node.GetTypePtr(), node.GetNamePtr());
1391 return true;1483 return true;
1392 }1484 }
1485+ 
1486+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
1487+ auto &outputs = const_cast<AscNode &>(node).outputs;
1488+ auto dt = static_cast<ge::DataType>(outputs[0].attr.dtype);
1489+ return dt == DT_FLOAT16 || dt == DT_FLOAT || dt == DT_BF16;
1490+ }
1491+ 
1492+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr(const AscNode &node, const std::vector<std::string> &inputs,
1493+ std::string &expr) const override {
1494+ auto &outputs = const_cast<AscNode &>(node).outputs;
1495+ if (outputs[0].attr.dtype == DT_FLOAT) {
1496+ expr = "AscendC::Simt::Sqrt(" + inputs[0] + ")";
1497+ } else {
1498+ expr = "AscendC::Simt::Sqrt(static_cast<float>(" + inputs[0] + "))";
1499+ }
1500+ return ge::GRAPH_SUCCESS;
1501+ }
1393};1502};
1394 1503 
1395class RsqrtAscIrCodegenImplV2 : public AscIrCodegenV2 {1504class RsqrtAscIrCodegenImplV2 : public AscIrCodegenV2 {
@@ -1452,6 +1561,22 @@ class NegAscIrCodegenImplV2 : public AscIrCodegenV2 {
1452 return true;1561 return true;
1453 }1562 }
1454 1563 
1564+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
1565+ (void)node;
1566+ return true;
1567+ }
1568+ 
1569+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr(const AscNode &node, const std::vector<std::string> &inputs,
1570+ std::string &expr) const override {
1571+ auto &outputs = const_cast<AscNode &>(node).outputs;
1572+ if (outputs[0].attr.dtype == DT_FLOAT16 || outputs[0].attr.dtype == DT_BF16) {
1573+ expr = "-static_cast<float>(" + inputs[0] + ")";
1574+ } else {
1575+ expr = "-(" + inputs[0] + ")";
1576+ }
1577+ return ge::GRAPH_SUCCESS;
1578+ }
1579+ 
1455 [[nodiscard]] bool IsInplaceSupported(const AscNode &neg_node) const override {1580 [[nodiscard]] bool IsInplaceSupported(const AscNode &neg_node) const override {
1456 (void)neg_node;1581 (void)neg_node;
1457 return true;1582 return true;
@@ -1500,6 +1625,18 @@ class ReluAscIrCodegenImplV2 : public AscIrCodegenV2 {
1500 return true;1625 return true;
1501 }1626 }
1502 1627 
1628+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
1629+ (void)node;
1630+ return true;
1631+ }
1632+ 
1633+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr([[maybe_unused]] const AscNode &node,
1634+ const std::vector<std::string> &inputs,
1635+ std::string &expr) const override {
1636+ expr = "AscendC::Simt::Max(" + inputs[0] + ", static_cast<decltype(" + inputs[0] + ")>(0))";
1637+ return ge::GRAPH_SUCCESS;
1638+ }
1639+ 
1503 [[nodiscard]] bool IsInplaceSupported(const AscNode &relu_node) const override {1640 [[nodiscard]] bool IsInplaceSupported(const AscNode &relu_node) const override {
1504 (void)relu_node;1641 (void)relu_node;
1505 return true;1642 return true;
@@ -2571,7 +2708,25 @@ class LeakyReluAscIrCodegenImplV2 : public AscIrCodegenV2 {
2571 node.GetTypePtr(), node.GetNamePtr());2708 node.GetTypePtr(), node.GetNamePtr());
2572 return true;2709 return true;
2573 }2710 }
2711+ 
2712+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
2713+ (void)node;
2714+ return true;
2715+ }
2716+ 
2717+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr(const AscNode &node, const std::vector<std::string> &inputs,
2718+ std::string &expr) const override {
2719+ float negative_slope = 0.0f;
2720+ GE_ASSERT_NOTNULL(node.attr.ir_attr, "LeakyRelu node %s has no ir attr", node.GetNamePtr());
2721+ GE_ASSERT_GRAPH_SUCCESS(
2722+ const_cast<AscIrAttrDefBase *>(node.attr.ir_attr.get())->GetAttrValue("negative_slope", negative_slope));
2723+ const std::string &input = inputs[0];
2724+ expr = "(" + input + " > static_cast<decltype(" + input + ")>(0)) ? " + input + " : static_cast<decltype(" + input +
2725+ ")>(" + std::to_string(negative_slope) + ") * " + input;
2726+ return ge::GRAPH_SUCCESS;
2727+ }
2574};2728};
2729+ 
2575/*********************************************************************************/2730/*********************************************************************************/
2576class ClipByValueAscIrCodegenImplV2 : public AscIrCodegenV2 {2731class ClipByValueAscIrCodegenImplV2 : public AscIrCodegenV2 {
2577 public:2732 public:
@@ -2711,6 +2866,61 @@ class GatherAscIrCodegenImplV2 : public AscIrCodegenV2 {
2711 }2866 }
2712};2867};
2713/*********************************************************************************/2868/*********************************************************************************/
2869+class IndirectLoadAscIrCodegenImplV2 : public AscIrCodegenV2 {
2870+ public:
2871+ [[nodiscard]] std::vector<std::unique_ptr<TmpBufDesc>> CalcTmpBufSize(const AscNode &node) override {
2872+ if (::ascir::GetTemplateIdOrDefault(node) != ::ascir::TemplateId::kIndirectLoadSimd) {
2873+ return {};
2874+ }
2875+ auto node_inputs = node.inputs;
2876+ auto node_outputs = node.outputs;
2877+ const auto &x = node_inputs[0].attr;
2878+ const auto &index = node_inputs[1].attr;
2879+ const auto &y = node_outputs[0].attr;
2880+ Expression x_size = Symbol(GetSizeByDataType(x.dtype));
2881+ for (const auto &repeat : x.repeats) {
2882+ x_size = x_size * repeat;
2883+ }
2884+ Expression index_size = Symbol(GetSizeByDataType(index.dtype));
2885+ for (const auto &repeat : y.repeats) {
2886+ index_size = index_size * repeat;
2887+ }
2888+ Expression offset_size = Symbol(sizeof(uint32_t));
2889+ for (const auto &repeat : y.repeats) {
2890+ offset_size = offset_size * repeat;
2891+ }
2892+ const size_t index_dtype_size = static_cast<size_t>(GetSizeByDataType(index.dtype));
2893+ const Expression aligned_x_size = af::sym::Align(x_size, 32);
2894+ const Expression offset_or_index_size = index_dtype_size > sizeof(uint32_t) ? index_size : offset_size;
2895+ TmpBufDesc desc = {Symbol(2) * aligned_x_size + af::sym::Align(offset_or_index_size, 32), -1};
Z
Zzhang_shengjie7月27日

当前分配： 2 * aligned_x_size + aligned_offset_size 但实际 SIMD buffer 布局是： 原始输入窗口 + 预处理后输入窗口 + offset buffer 例如： GM float32 -> Cast float16 -> IndirectLoad 元素数 N 临时空间计算看到的是 Cast 后的 IndirectLoad 输入： 当前分配：2 * 2N = 4N bytes 实际需要：4N + 2N = 6N bytes 不含 offset buffer 就已经少了 2N，后续 offset_buf 会落到已分配空间之外，可能覆盖相邻 UB。 当前测试只覆盖了同 dtype 的 Relu，以及 float16→float32 的放大 Cast；没有覆盖 float32→float16、int32→int16 等降精度场景。 建议分别计算： aligned(source_dtype_size * input_elements)

  • aligned(preprocess_output_dtype_size * input_elements)
  • aligned(offset_or_index_size)
likedislike
2896+ std::vector<std::unique_ptr<TmpBufDesc>> tmp_buf_descs;
2897+ tmp_buf_descs.emplace_back(std::make_unique<TmpBufDesc>(desc));
2898+ return tmp_buf_descs;
2899+ }
2900+ 
2901+ [[nodiscard]] std::string GetApiCallName() const override {
2902+ return "IndirectLoadRegApiCall";
2903+ }
2904+ [[nodiscard]] std::string GetApiName() const override {
2905+ return "IndirectLoad";
2906+ }
2907+ [[nodiscard]] bool IsVectorFunctionSupported(const AscNode &node) const override {
2908+ (void)node;
2909+ return false;
2910+ }
2911+ [[nodiscard]] std::vector<std::string> LoadApiHeaderFiles([[maybe_unused]] bool is_dynamic) const override {
2912+ return {"datacopy_reg_base.h"};
2913+ }
2914+ [[nodiscard]] std::vector<std::string> IncludeApiHeaderFiles() const override {
2915+ return {"basic_api/kernel_operator_vec_gather_intf.h", "simt_api/cpp/kernel_simt_intf.h"};
2916+ }
2917+ [[nodiscard]] bool IsNodeValid(const AscNode &node) const override {
2918+ GE_ASSERT_TRUE(!IsNodeHasScalarInput(node), "Node %s[%s] not support scalar input", node.GetTypePtr(),
2919+ node.GetNamePtr());
2920+ return true;
2921+ }
2922+};
2923+/*********************************************************************************/
2714class TransposeAscIrCodegenImplV2 : public AscIrCodegenV2 {2924class TransposeAscIrCodegenImplV2 : public AscIrCodegenV2 {
2715 public:2925 public:
2716 [[nodiscard]] std::vector<std::unique_ptr<TmpBufDesc>> CalcTmpBufSize(const AscNode &node) override {2926 [[nodiscard]] std::vector<std::unique_ptr<TmpBufDesc>> CalcTmpBufSize(const AscNode &node) override {
@@ -3447,6 +3657,20 @@ class BitwiseNotAscIrCodegenImplV2 : public AscIrCodegenV2 {
3447 node.GetTypePtr(), node.GetNamePtr());3657 node.GetTypePtr(), node.GetNamePtr());
3448 return true;3658 return true;
3449 }3659 }
3660+ 
3661+ [[nodiscard]] bool IsSimtScalarSupported(const AscNode &node) const override {
3662+ auto &outputs = const_cast<AscNode &>(node).outputs;
3663+ auto dt = static_cast<ge::DataType>(outputs[0].attr.dtype);
3664+ return dt == DT_INT8 || dt == DT_INT16 || dt == DT_INT32 || dt == DT_INT64 || dt == DT_UINT8 || dt == DT_UINT16 ||
3665+ dt == DT_UINT32 || dt == DT_UINT64;
3666+ }
3667+ 
3668+ [[nodiscard]] ge::graphStatus GenerateSimtScalarExpr([[maybe_unused]] const AscNode &node,
3669+ const std::vector<std::string> &inputs,
3670+ std::string &expr) const override {
3671+ expr = "~(" + inputs[0] + ")";
3672+ return ge::GRAPH_SUCCESS;
3673+ }
3450};3674};
3451 3675 
3452class BitwiseOrAscIrCodegenImplV2 : public AscIrCodegenV2 {3676class BitwiseOrAscIrCodegenImplV2 : public AscIrCodegenV2 {
Mautofuse/v35/att/api_perf_register/ascir_api_perf_v2.cpp+3-0
@@ -503,6 +503,9 @@ ApiPerfRegister<ApiPerf> add_api_perf_v2(ApiPerfRegisterV2(kAdd, GetPerfFunc(kAd
503 &perf_param_table_v2, &tiling_schedule_config_table_v2));503 &perf_param_table_v2, &tiling_schedule_config_table_v2));
504ApiPerfRegister<ApiPerf> gather_api_perf_v2(ApiPerfRegisterV2(kGather, GetPerfFunc(kGather), nullptr,504ApiPerfRegister<ApiPerf> gather_api_perf_v2(ApiPerfRegisterV2(kGather, GetPerfFunc(kGather), nullptr,
505 &perf_param_table_v2, &tiling_schedule_config_table_v2));505 &perf_param_table_v2, &tiling_schedule_config_table_v2));
506+ApiPerfRegister<ApiPerf> indirect_load_api_perf_v2(ApiPerfRegisterV2(kIndirectLoad, GetPerfFunc(kUnitVector), nullptr,
507+ &perf_param_table_v2,
508+ &tiling_schedule_config_table_v2));
506ApiPerfRegister<ApiPerf> abs_api_perf_v2(ApiPerfRegisterV2(kAbs, GetPerfFunc(kAbs + "V2"), nullptr,509ApiPerfRegister<ApiPerf> abs_api_perf_v2(ApiPerfRegisterV2(kAbs, GetPerfFunc(kAbs + "V2"), nullptr,
507 &perf_param_table_v2, &tiling_schedule_config_table_v2));510 &perf_param_table_v2, &tiling_schedule_config_table_v2));
508ApiPerfRegister<ApiPerf> broadcast_api_perf_v2(ApiPerfRegisterV2(kBroadcast, GetPerfFunc(kBroadcast), nullptr,511ApiPerfRegister<ApiPerf> broadcast_api_perf_v2(ApiPerfRegisterV2(kBroadcast, GetPerfFunc(kBroadcast), nullptr,