Star1150
290
代码
代码
Issues55
Pull Requests215
Actions
项目讨论
Wiki
分析
Star1150
290
已合并
feat: 支持 IndirectLoad Load-Transpose 布局并完善 SIMD/SIMT codegen #2029
xiebangrui2025创建于 9月9日
feat: 支持 IndirectLoad Load-Transpose 布局并完善 SIMD/SIMT codegen #2029
已合并
xiebangrui2025创建于 9月9日
从xiebangrui2025/graph-autofusionIL合入到cann/graph-autofusiondevelop
共 19 个文件变更+2862-1324
Mautofuse/common/indirect_load_utils.cpp+638-0
@@ -11,13 +11,16 @@
11#include "indirect_load_utils.h"11#include "indirect_load_utils.h"
12 12 
13#include <algorithm>13#include <algorithm>
14+#include <limits>
14#include <string>15#include <string>
16+#include <unordered_map>
15#include <unordered_set>17#include <unordered_set>
16 18 
17#include "ascir_ops.h"19#include "ascir_ops.h"
18#include "ascir_ops_utils.h"20#include "ascir_ops_utils.h"
19#include "graph/symbolizer/symbolic_utils.h"21#include "graph/symbolizer/symbolic_utils.h"
20#include "schedule_result.h"22#include "schedule_result.h"
23+#include "utils/extern_math_util.h"
21 24 
22namespace ascgen_utils::indirect_load {25namespace ascgen_utils::indirect_load {
23namespace {26namespace {
Mautofuse/common/indirect_load_utils.h+96-0
@@ -24,6 +24,7 @@ constexpr size_t kInputTensorIndex = 0UL;
24constexpr size_t kIndexTensorIndex = 1UL;24constexpr size_t kIndexTensorIndex = 1UL;
25constexpr char kTemplateLogicalViewAttr[] = "af.internal.indirect_load.logical_view";25constexpr char kTemplateLogicalViewAttr[] = "af.internal.indirect_load.logical_view";
26constexpr char kAccessInfoAttr[] = "af.internal.indirect_load.access_info";26constexpr char kAccessInfoAttr[] = "af.internal.indirect_load.access_info";
27+constexpr char kLoweringMetadataAttr[] = "af.internal.indirect_load.lowering_metadata";
27 28 
28enum class TemplateRole : int64_t {29enum class TemplateRole : int64_t {
29 kNone,30 kNone,
@@ -103,6 +104,97 @@ struct IndirectLoadAccessInfo {
103 bool can_use_simd_embedding = false;104 bool can_use_simd_embedding = false;
104};105};
105 106 
107+enum class SimdFallback : int64_t {
108+ kRegisterGather = 0,
109+ kGatherApi = 1,
110+ kStrided = 2,
111+};
112+ 
113+struct SimdLoweringMetadata {
114+ SimdFallback fallback = SimdFallback::kRegisterGather;
115+ bool try_embedding = false;
116+};
117+ 
118+enum class SimtAddressPolicy : int64_t {
119+ kStaticPowerOfTwo = 0,
120+ kStaticInner = 1,
121+ kStructuredMagic = 2,
122+ kEmbedding = 3,
123+ kRecursive = 4,
124+ kStrided = 5,
125+};
126+ 
127+enum class SimtOffsetWidth : int64_t {
128+ kUint32,
129+ kUint64,
130+};
131+ 
132+enum class SimtLoadAddressSource : int64_t {
133+ kZeroOffset,
134+ kOutputOffset,
135+ kIndexOffset,
136+};
137+ 
138+struct SimtPolicyMetadata {
139+ SimtAddressPolicy policy = SimtAddressPolicy::kRecursive;
140+ SimtOffsetWidth offset_width = SimtOffsetWidth::kUint64;
141+ uint64_t inner_span = 0U;
142+ uint64_t output_axis_span = 0U;
143+ uint64_t input_axis_stride = 0U;
144+ uint64_t input_axis_span = 0U;
145+ uint64_t input_stride_mask = 0U;
146+ uint64_t index_stride_mask = 0U;
147+ std::vector<af::Expression> runtime_params;
148+};
149+ 
150+struct SimtGmTensorMetadata {
151+ ascir::TensorId value_tensor_id = af::kIdNone;
152+ ascir::TensorId gm_tensor_id = af::kIdNone;
153+ af::DataType dtype = af::DT_UNDEFINED;
154+ bool is_scalar = false;
155+};
156+ 
157+struct SimtLoadMetadata {
158+ std::string node_name;
159+ SimtLoadAddressSource address_source = SimtLoadAddressSource::kOutputOffset;
160+ bool use_logical_offset = false;
161+ LogicalTensorView physical_view;
162+};
163+ 
164+struct SimtOutputChainMetadata {
165+ std::vector<std::string> node_names;
166+ ascir::TensorId result_tensor_id = af::kIdNone;
167+ ascir::TensorId target_tensor_id = af::kIdNone;
168+ af::DataType dtype = af::DT_UNDEFINED;
169+ bool local_target = false;
170+};
171+ 
172+struct SimtLoweringMetadata {
173+ bool has_post_reduce = false;
174+ std::vector<std::string> index_node_names;
175+ std::vector<std::string> output_node_names;
176+ std::vector<SimtLoadMetadata> index_loads;
177+ std::vector<SimtLoadMetadata> output_loads;
178+ std::vector<SimtOutputChainMetadata> output_chains;
179+ std::vector<SimtGmTensorMetadata> gm_tensors;
180+ SimtPolicyMetadata policy;
181+ ascir::TensorId index_result_tensor_id = af::kIdNone;
182+ ascir::TensorId value_tensor_id = af::kIdNone;
183+ ascir::TensorId output_result_tensor_id = af::kIdNone;
184+ af::DataType index_dtype = af::DT_UNDEFINED;
185+ af::DataType output_dtype = af::DT_UNDEFINED;
186+};
187+ 
188+struct IndirectLoadLoweringMetadata {
189+ uint32_t version = 1U;
190+ int64_t axis = -1L;
191+ af::AxisId outer_axis = af::kIdNone;
192+ TemplateLogicalView logical_view;
193+ IndirectLoadAccessInfo access_info;
194+ SimdLoweringMetadata simd;
195+ SimtLoweringMetadata simt;
196+};
197+ 
106TemplateBehavior GetTemplateBehavior(const af::AscNodePtr &node);198TemplateBehavior GetTemplateBehavior(const af::AscNodePtr &node);
107TemplateRole GetTemplateRole(const af::AscNodePtr &node);199TemplateRole GetTemplateRole(const af::AscNodePtr &node);
108af::AscNodePtr GetPostReduceConsumer(const af::AscNodePtr &node);200af::AscNodePtr GetPostReduceConsumer(const af::AscNodePtr &node);
@@ -118,6 +210,10 @@ af::Status SetTemplateLogicalView(const af::AscNodePtr &node, const TemplateLogi
118af::Status GetTemplateLogicalView(const af::AscNodePtr &node, TemplateLogicalView &view);210af::Status GetTemplateLogicalView(const af::AscNodePtr &node, TemplateLogicalView &view);
119af::Status SetIndirectLoadAccessInfo(const af::AscNodePtr &node, const IndirectLoadAccessInfo &info);211af::Status SetIndirectLoadAccessInfo(const af::AscNodePtr &node, const IndirectLoadAccessInfo &info);
120af::Status GetIndirectLoadAccessInfo(const af::AscNodePtr &node, IndirectLoadAccessInfo &info);212af::Status GetIndirectLoadAccessInfo(const af::AscNodePtr &node, IndirectLoadAccessInfo &info);
213+af::Status SetLoweringMetadata(const af::AscNodePtr &node, const IndirectLoadLoweringMetadata &metadata);
214+af::Status GetLoweringMetadata(const af::AscNodePtr &node, IndirectLoadLoweringMetadata &metadata);
215+bool HasLoweringMetadata(const af::AscNodePtr &node);
216+af::Status FinalizeLoweringMetadata(const af::AscNodePtr &node, bool &is_supported);
121af::Status AnalyzeIndirectLoadAccess(const af::AscNodePtr &node, const TemplateLogicalView &logical_view,217af::Status AnalyzeIndirectLoadAccess(const af::AscNodePtr &node, const TemplateLogicalView &logical_view,
122 IndirectLoadAccessInfo &info);218 IndirectLoadAccessInfo &info);
123af::Status SetImplementation(const af::AscNodePtr &node, Implementation implementation);219af::Status SetImplementation(const af::AscNodePtr &node, Implementation implementation);
Mautofuse/optimize/autoschedule/schedule.cpp+0-15
@@ -629,13 +629,6 @@ Status Scheduler::InitIndirectLoadScheduleCase() {
629 return af::SUCCESS;629 return af::SUCCESS;
630 }630 }
631 631 
632- for (const af::AscNodePtr &node : graph_.GetAllNodes()) {
633- const auto role = ascgen_utils::indirect_load::GetTemplateRole(node);
634- if (role == ascgen_utils::indirect_load::TemplateRole::kStridedUbPath ||
635- role == ascgen_utils::indirect_load::TemplateRole::kSimdInputPreStridedUbPath) {
636- indirect_load_info_.aligned_strided_path.emplace_back(node);
637- }
638- }
639 GE_ASSERT_NOTNULL(tiling_case_.ub_tiling_y.first);632 GE_ASSERT_NOTNULL(tiling_case_.ub_tiling_y.first);
640 GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, indirect_load_info_.axes));633 GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::GetTemplateAxes(indirect_load, indirect_load_info_.axes));
641 indirect_load_info_.active = true;634 indirect_load_info_.active = true;
@@ -772,14 +765,6 @@ Status Scheduler::DoScheduler() {
772 }765 }
773 GE_ASSERT_SUCCESS(NodeCacheMarker(graph_).MarkIfNodeNeedsCache());766 GE_ASSERT_SUCCESS(NodeCacheMarker(graph_).MarkIfNodeNeedsCache());
774 GE_ASSERT_SUCCESS(AlignmentHandler::ModifyVectorizedStrides(graph_));767 GE_ASSERT_SUCCESS(AlignmentHandler::ModifyVectorizedStrides(graph_));
775- for (const af::AscNodePtr &node : indirect_load_info_.aligned_strided_path) {
776- for (const auto &output : node->outputs()) {
777- if (!output->attr.vectorized_axis.empty()) {
778- GE_ASSERT_SUCCESS(
779- BaseAlignmentStrategy::SetVectorizedStridesForTensor(node, output->attr, AlignmentType::kAligned));
780- }
781- }
782- }
783 ascir::utils::DumpGraph(graph_, "AfterDoTiling");768 ascir::utils::DumpGraph(graph_, "AfterDoTiling");
784 return af::SUCCESS;769 return af::SUCCESS;
785}770}
Mautofuse/optimize/autoschedule/schedule.h+0-1
@@ -43,7 +43,6 @@ struct TilingCase {
43 43 
44struct IndirectLoadInfo {44struct IndirectLoadInfo {
45 bool active = false;45 bool active = false;
46- std::vector<af::AscNodePtr> aligned_strided_path;
47 ascgen_utils::indirect_load::TemplateAxes axes;46 ascgen_utils::indirect_load::TemplateAxes axes;
48};47};
49 48 
Mautofuse/optimize/task_generator/indirect_load_schedule_case_generator.cpp+206-49
@@ -56,6 +56,7 @@ struct RewrittenGraphAnalysis {
56 NodePath input_path;56 NodePath input_path;
57 NodePath index_path;57 NodePath index_path;
58 af::AscNodePtr input_root;58 af::AscNodePtr input_root;
59+ af::AscNodePtr input_boundary;
59 af::AscNodePtr index_root;60 af::AscNodePtr index_root;
60 af::AscNodePtr output_store;61 af::AscNodePtr output_store;
61 af::AscNodePtr post_reduce;62 af::AscNodePtr post_reduce;
@@ -66,6 +67,7 @@ struct RewrittenGraphAnalysis {
66 67 
67struct InputViewPlan {68struct InputViewPlan {
68 NodePath path;69 NodePath path;
70+ af::AscNodePtr load_transpose;
69 ascgen_utils::indirect_load::IndirectLoadTensorLayout layout;71 ascgen_utils::indirect_load::IndirectLoadTensorLayout layout;
70 int64_t path_broadcast_index = kInvalidBroadcastIndex;72 int64_t path_broadcast_index = kInvalidBroadcastIndex;
71 bool simd_index_uses_output_inner_axis = false;73 bool simd_index_uses_output_inner_axis = false;
@@ -85,10 +87,24 @@ struct PostReduceLayout {
85 size_t first_reduce;87 size_t first_reduce;
86};88};
87 89 
88-bool IsInputRegionBoundary(const af::AscNodePtr &node) {90+bool IsInputDataSource(const af::AscNodePtr &node) {
89 return ScheduleUtils::IsDataInput(node) || af::ops::IsOps<af::ascir_op::Scalar>(node);91 return ScheduleUtils::IsDataInput(node) || af::ops::IsOps<af::ascir_op::Scalar>(node);
90}92}
91 93 
94+// A retained Transpose separates two physical coordinate systems. Do not overwrite its source GM view.
95+bool IsInputRegionBoundary(const af::AscNodePtr &node) {
96+ if (IsInputDataSource(node)) {
97+ return true;
98+ }
99+ if (!af::ops::IsOps<af::ascir_op::Load>(node)) {
100+ return false;
101+ }
102+ const auto consumers = node->GetOutDataNodes();
103+ return std::any_of(consumers.begin(), consumers.end(), [](const af::NodePtr &consumer) {
104+ return af::ops::IsOps<af::ascir_op::Transpose>(std::dynamic_pointer_cast<af::AscNode>(consumer));
105+ });
106+}
107+ 
92bool HasControlEdge(const af::AscNodePtr &node) {108bool HasControlEdge(const af::AscNodePtr &node) {
93 return node->GetInControlNodesSize() != 0UL || node->GetOutControlNodesSize() != 0UL;109 return node->GetInControlNodesSize() != 0UL || node->GetOutControlNodesSize() != 0UL;
94}110}
@@ -153,7 +169,10 @@ af::Status GetBroadcastPhysicalAttr(const af::AscNodePtr &broadcast, ascir::Temp
153 static_cast<int32_t>(template_id), broadcast->GetNamePtr());169 static_cast<int32_t>(template_id), broadcast->GetNamePtr());
154 GE_ASSERT_NOTNULL(producer_outputs[output_idx], "IndirectLoad template[%d] Broadcast node[%s] source output is null.",170 GE_ASSERT_NOTNULL(producer_outputs[output_idx], "IndirectLoad template[%d] Broadcast node[%s] source output is null.",
155 static_cast<int32_t>(template_id), broadcast->GetNamePtr());171 static_cast<int32_t>(template_id), broadcast->GetNamePtr());
156- physical_attr = producer_outputs[output_idx]->attr;172+ const auto &producer_attr = producer_outputs[output_idx]->attr;
173+ physical_attr.axis = producer_attr.axis;
174+ physical_attr.repeats = producer_attr.repeats;
175+ physical_attr.strides = producer_attr.strides;
157 return af::SUCCESS;176 return af::SUCCESS;
158}177}
159 178 
@@ -241,7 +260,7 @@ bool NeedsAlignedUbWindow(const ascgen_utils::indirect_load::IndirectLoadTensorL
241 260 
242af::Status AnnotateStridedUbPath(const NodePath &path) {261af::Status AnnotateStridedUbPath(const NodePath &path) {
243 for (const af::AscNodePtr &node : path) {262 for (const af::AscNodePtr &node : path) {
244- if (IsInputRegionBoundary(node) || ScheduleUtils::IsBuffer(node)) {263+ if (IsInputDataSource(node) || ScheduleUtils::IsBuffer(node)) {
245 continue;264 continue;
246 }265 }
247 const auto role = ascgen_utils::indirect_load::GetTemplateRole(node);266 const auto role = ascgen_utils::indirect_load::GetTemplateRole(node);
@@ -315,6 +334,19 @@ af::Status TraverseOutputConsumers(const NodePath &roots, NodeSet &visited, cons
315 return af::SUCCESS;334 return af::SUCCESS;
316}335}
317 336 
337+af::AscNodePtr GetLoadTransposeSource(const af::AscNodePtr &node) {
338+ if (node == nullptr || !af::ops::IsOps<af::ascir_op::Transpose>(node)) {
339+ return nullptr;
340+ }
341+ const auto producer = ascgen_utils::indirect_load::GetInputProducer(node, 0UL);
342+ return producer != nullptr && af::ops::IsOps<af::ascir_op::Load>(producer) ? producer : nullptr;
343+}
344+ 
345+af::AscNodePtr GetFoldableLoadTransposeSource(const af::AscNodePtr &node) {
346+ const auto load = GetLoadTransposeSource(node);
347+ return IsSingleConsumerWithoutControlEdge(node) && IsSingleConsumerWithoutControlEdge(load) ? load : nullptr;
348+}
349+ 
318bool IsSupportedIndirectLoadTopologyNode(const af::AscNodePtr &node) {350bool IsSupportedIndirectLoadTopologyNode(const af::AscNodePtr &node) {
319 if (node == nullptr) {351 if (node == nullptr) {
320 return false;352 return false;
@@ -323,6 +355,9 @@ bool IsSupportedIndirectLoadTopologyNode(const af::AscNodePtr &node) {
323 af::ops::IsOps<af::ascir_op::IndirectLoad>(node)) {355 af::ops::IsOps<af::ascir_op::IndirectLoad>(node)) {
324 return true;356 return true;
325 }357 }
358+ if (GetLoadTransposeSource(node) != nullptr) {
359+ return true;
360+ }
326 return ScheduleUtils::IsLoad(node) || ScheduleUtils::IsStore(node) || ScheduleUtils::IsElewise(node) ||361 return ScheduleUtils::IsLoad(node) || ScheduleUtils::IsStore(node) || ScheduleUtils::IsElewise(node) ||
327 ScheduleUtils::IsBroadcast(node) || ScheduleUtils::IsReduce(node);362 ScheduleUtils::IsBroadcast(node) || ScheduleUtils::IsReduce(node);
328}363}
@@ -366,7 +401,7 @@ bool CollectSimtBackwardRegion(const NodePath &roots, const af::AscNodePtr &indi
366 if (node == nullptr) {401 if (node == nullptr) {
367 return false;402 return false;
368 }403 }
369- if (node == indirect_load || IsInputRegionBoundary(node)) {404+ if (node == indirect_load || IsInputDataSource(node)) {
370 continue;405 continue;
371 }406 }
372 if (!region.emplace(node.get()).second || af::ops::IsOps<af::ascir_op::Load>(node)) {407 if (!region.emplace(node.get()).second || af::ops::IsOps<af::ascir_op::Load>(node)) {
@@ -921,7 +956,7 @@ af::Status ReplaceAxisPrefix(std::vector<af::AxisId> &target_axes, const std::ve
921af::Status ReplaceRegionAxisPrefix(const NodePath &region, const std::vector<af::AxisId> &output_axes,956af::Status ReplaceRegionAxisPrefix(const NodePath &region, const std::vector<af::AxisId> &output_axes,
922 size_t axis_index, bool is_input_region) {957 size_t axis_index, bool is_input_region) {
923 for (const af::AscNodePtr &node : region) {958 for (const af::AscNodePtr &node : region) {
924- if (!is_input_region && IsInputRegionBoundary(node)) {959+ if (!is_input_region && IsInputDataSource(node)) {
925 continue;960 continue;
926 }961 }
927 const auto role = ascgen_utils::indirect_load::GetTemplateRole(node);962 const auto role = ascgen_utils::indirect_load::GetTemplateRole(node);
@@ -1012,6 +1047,27 @@ bool CanEmitSimtScalar(const af::AscNodePtr &node) {
1012 return v2_impl != nullptr && v2_impl->IsSimtScalarSupported(*node);1047 return v2_impl != nullptr && v2_impl->IsSimtScalarSupported(*node);
1013}1048}
1014 1049 
1050+// The IndirectLoad logical view is expressed in the Transpose output order,
1051+// while the direct-GM boundary still reads the original Load buffer. Map the
1052+// source Load strides into that output order for the strided address policy.
1053+af::Status BuildLoadTransposeSourceView(const af::AscNodePtr &transpose, const af::AscTensorAttr &logical_attr,
1054+ ascgen_utils::indirect_load::LogicalTensorView &view) {
1055+ const auto load = ascgen_utils::indirect_load::GetInputProducer(transpose, 0UL);
1056+ const auto &source = load->outputs[0].attr;
1057+ GE_ASSERT_TRUE(source.axis.size() == source.strides.size() &&
1058+ logical_attr.axis.size() == logical_attr.repeats.size() &&
1059+ logical_attr.axis.size() == logical_attr.strides.size(),
1060+ "IndirectLoad Load->Transpose tensor metadata rank mismatch.");
1061+ view = {logical_attr.axis, logical_attr.repeats, logical_attr.strides};
1062+ for (size_t output_dim = 0UL; output_dim < logical_attr.axis.size(); ++output_dim) {
1063+ const auto source_axis = std::find(source.axis.begin(), source.axis.end(), logical_attr.axis[output_dim]);
1064+ GE_ASSERT_TRUE(source_axis != source.axis.end(), "IndirectLoad Transpose output axis is absent from Load source.");
1065+ const size_t source_dim = static_cast<size_t>(std::distance(source.axis.begin(), source_axis));
1066+ view.strides[output_dim] = source.strides[source_dim];
1067+ }
1068+ return af::SUCCESS;
1069+}
1070+ 
1015af::Status ValidateReduceOutput(const af::AscNodePtr &reduce) {1071af::Status ValidateReduceOutput(const af::AscNodePtr &reduce) {
1016 const auto reduce_outputs = reduce->GetOutDataNodes();1072 const auto reduce_outputs = reduce->GetOutDataNodes();
1017 GE_ASSERT_TRUE(reduce_outputs.size() == 1UL, "[IndirectLoad] Reduce[%s] must have one Store output, got %zu.",1073 GE_ASSERT_TRUE(reduce_outputs.size() == 1UL, "[IndirectLoad] Reduce[%s] must have one Store output, got %zu.",
@@ -1067,6 +1123,10 @@ af::Status CollectOutputBoundaries(const af::AscNodePtr &indirect_load, Rewritte
1067af::Status CollectRewrittenBoundaries(const af::AscNodePtr &indirect_load, RewrittenGraphAnalysis &analysis) {1123af::Status CollectRewrittenBoundaries(const af::AscNodePtr &indirect_load, RewrittenGraphAnalysis &analysis) {
1068 analysis.input_root =1124 analysis.input_root =
1069 ascgen_utils::indirect_load::GetInputProducer(indirect_load, ascgen_utils::indirect_load::kInputTensorIndex);1125 ascgen_utils::indirect_load::GetInputProducer(indirect_load, ascgen_utils::indirect_load::kInputTensorIndex);
1126+ analysis.input_boundary = GetLoadTransposeSource(analysis.input_root);
1127+ if (analysis.input_boundary == nullptr) {
1128+ analysis.input_boundary = analysis.input_root;
1129+ }
1070 analysis.index_root =1130 analysis.index_root =
1071 ascgen_utils::indirect_load::GetInputProducer(indirect_load, ascgen_utils::indirect_load::kIndexTensorIndex);1131 ascgen_utils::indirect_load::GetInputProducer(indirect_load, ascgen_utils::indirect_load::kIndexTensorIndex);
1072 GE_ASSERT_SUCCESS(CollectOutputBoundaries(indirect_load, analysis));1132 GE_ASSERT_SUCCESS(CollectOutputBoundaries(indirect_load, analysis));
@@ -1160,6 +1220,14 @@ af::Status RewriteInputPreNodes(af::AscGraph &graph, const af::AscNodePtr &indir
1160 indirect_load->GetNamePtr(), af::ToString(input_shape).c_str(), af::ToString(output_shape).c_str());1220 indirect_load->GetNamePtr(), af::ToString(input_shape).c_str(), af::ToString(output_shape).c_str());
1161 return af::SUCCESS;1221 return af::SUCCESS;
1162 }1222 }
1223+ // Keep a direct Load->Transpose boundary on the SIMT input side. Its
1224+ // permutation is represented by the SIMT address policy; moving it across
1225+ // IndirectLoad would discard the source physical view.
1226+ const auto input_producer =
1227+ ascgen_utils::indirect_load::GetInputProducer(indirect_load, ascgen_utils::indirect_load::kInputTensorIndex);
1228+ if (template_id == ascir::TemplateId::kIndirectLoadSimt && GetLoadTransposeSource(input_producer) != nullptr) {
1229+ return af::SUCCESS;
1230+ }
1163 NodePath movable_nodes;1231 NodePath movable_nodes;
1164 af::AscNodePtr consumer = indirect_load;1232 af::AscNodePtr consumer = indirect_load;
1165 for (af::AscNodePtr node =1233 for (af::AscNodePtr node =
@@ -1262,47 +1330,71 @@ void CollectInputPaths(const af::AscNodePtr &indirect_load, ascir::TemplateId te
1262 }1330 }
1263}1331}
1264 1332 
1265-af::Status AnalyzeInputPath(const af::AscNodePtr &indirect_load, size_t input_idx, ascir::TemplateId template_id,1333+af::Status ResolveInputPhysicalView(const af::AscTensorAttr &logical_attr, ascir::TemplateId template_id,
1266- InputViewPlan *const plans[], bool &is_path_supported) {1334+ InputViewPlan &plan, af::AscTensorAttr &physical_attr, bool &is_supported) {
1267- // 返回 SUCCESS 表示分析流程正常；is_path_supported 为 false 表示当前模板淘汰，不中止 Generate。1335+ physical_attr = logical_attr;
1268- is_path_supported = true;1336+ plan.load_transpose = nullptr;
1269- const auto inputs = indirect_load->inputs();
1270- const size_t axis_index = GetIndirectLoadAxisIndex(indirect_load);
1271- GE_ASSERT_TRUE(axis_index != kIndirectLoadInvalidAxisIndex, "IndirectLoad axis index is invalid.");
1272- InputViewPlan &plan = *plans[input_idx];
1273- const auto &logical_attr = inputs[input_idx]->attr;
1274- plan.simd_index_uses_output_inner_axis = false;
1275- ascgen_utils::indirect_load::LogicalTensorView view;
1276 const bool has_broadcast = plan.path_broadcast_index != kInvalidBroadcastIndex;1337 const bool has_broadcast = plan.path_broadcast_index != kInvalidBroadcastIndex;
1277- if (!has_broadcast) {1338+ af::AscNodePtr source = plan.path.empty() ? nullptr : plan.path.front();
1278- // 无直连广播（SIMD/SIMT）或无广播（SK）：按逻辑视图直接分类，非直连广播由通用逻辑承载。1339+ if (has_broadcast) {
1279- view = {logical_attr.axis, logical_attr.repeats, logical_attr.strides};1340+ const auto broadcast = plan.path[static_cast<size_t>(plan.path_broadcast_index)];
1280- } else {1341+ source = ascgen_utils::indirect_load::GetInputProducer(broadcast, 0UL);
1281- const size_t broadcast_index = static_cast<size_t>(plan.path_broadcast_index);1342+ if (source == nullptr) {
1282- const af::AscNodePtr &broadcast = plan.path[broadcast_index];
1283- // 取源物理属性前保留源存在性检查。
1284- if (ascgen_utils::indirect_load::GetInputProducer(broadcast, 0UL) == nullptr) {
1285 GELOGI("[IndirectLoad] Reject candidate[%d]: Broadcast node[%s] source is invalid.",1343 GELOGI("[IndirectLoad] Reject candidate[%d]: Broadcast node[%s] source is invalid.",
1286 static_cast<int32_t>(template_id), broadcast->GetNamePtr());1344 static_cast<int32_t>(template_id), broadcast->GetNamePtr());
1287- is_path_supported = false;1345+ is_supported = false;
1288 return af::SUCCESS;1346 return af::SUCCESS;
1289 }1347 }
1290- // 基于广播源物理属性构造执行视图：源形状为 1 而逻辑视图非 1 的维度置 stride 0。
1291- af::AscTensorAttr physical_attr;
1292 GE_ASSERT_SUCCESS(GetBroadcastPhysicalAttr(broadcast, template_id, physical_attr));1348 GE_ASSERT_SUCCESS(GetBroadcastPhysicalAttr(broadcast, template_id, physical_attr));
1349+ }
1350+ if (template_id == ascir::TemplateId::kIndirectLoadSK || GetLoadTransposeSource(source) == nullptr) {
1351+ return af::SUCCESS;
1352+ }
1353+ // SIMD may use the source Load only when the intervening Transpose can be folded safely.
1354+ // SIMT already reads GM directly, so it only needs the axis-mapped view, not a graph rewrite.
1355+ if (has_broadcast && template_id == ascir::TemplateId::kIndirectLoadSimd &&
1356+ GetFoldableLoadTransposeSource(source) == nullptr) {
1357+ return af::SUCCESS;
1358+ }
1359+ plan.load_transpose = source;
1360+ ascgen_utils::indirect_load::LogicalTensorView source_view;
1361+ GE_ASSERT_SUCCESS(BuildLoadTransposeSourceView(source, physical_attr, source_view));
1362+ physical_attr.strides = std::move(source_view.strides);
1363+ return af::SUCCESS;
1364+}
1365+ 
1366+af::Status AnalyzeInputPath(const af::AscNodePtr &indirect_load, size_t input_idx, ascir::TemplateId template_id,
1367+ InputViewPlan &plan, bool &is_path_supported) {
1368+ is_path_supported = true;
1369+ const auto inputs = indirect_load->inputs();
1370+ const auto &logical_attr = inputs[input_idx]->attr;
1371+ af::AscTensorAttr physical_attr;
1372+ GE_ASSERT_SUCCESS(ResolveInputPhysicalView(logical_attr, template_id, plan, physical_attr, is_path_supported));
1373+ if (!is_path_supported) {
1374+ return af::SUCCESS;
1375+ }
1376+ const bool has_broadcast = plan.path_broadcast_index != kInvalidBroadcastIndex;
1377+ ascgen_utils::indirect_load::LogicalTensorView view{logical_attr.axis, logical_attr.repeats, physical_attr.strides};
1378+ if (has_broadcast) {
1293 GE_ASSERT_SUCCESS(BuildBroadcastLogicalView(logical_attr, physical_attr, template_id, view));1379 GE_ASSERT_SUCCESS(BuildBroadcastLogicalView(logical_attr, physical_attr, template_id, view));
1294 }1380 }
1295- 1381+ plan.simd_index_uses_output_inner_axis = template_id == ascir::TemplateId::kIndirectLoadSimd &&
1296- const bool allow_simt_strided_broadcast = template_id == ascir::TemplateId::kIndirectLoadSimt && has_broadcast;1382+ input_idx == ascgen_utils::indirect_load::kIndexTensorIndex &&
1297- GE_ASSERT_SUCCESS(1383+ (has_broadcast || plan.load_transpose != nullptr);
1298- ascgen_utils::indirect_load::ClassifyIndirectLoadLayout(view, plan.layout, allow_simt_strided_broadcast));1384+ if (plan.load_transpose != nullptr) {
1385+ // Permutations are valid non-overlapping strided views even when strides are not monotonically decreasing.
1386+ plan.layout = {{view.axis_ids, view.sizes, view.strides},
1387+ ascgen_utils::indirect_load::IndirectLoadLayoutKind::kStrided,
1388+ physical_attr.repeats};
1389+ } else {
1390+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::ClassifyIndirectLoadLayout(
1391+ view, plan.layout, template_id == ascir::TemplateId::kIndirectLoadSimt && has_broadcast));
1392+ }
1299 if (plan.layout.kind == ascgen_utils::indirect_load::IndirectLoadLayoutKind::kUnsupported) {1393 if (plan.layout.kind == ascgen_utils::indirect_load::IndirectLoadLayoutKind::kUnsupported) {
1300 GELOGI("[IndirectLoad] Reject candidate[%d]: input path layout%s is unsupported.",1394 GELOGI("[IndirectLoad] Reject candidate[%d]: input path layout%s is unsupported.",
1301 static_cast<int32_t>(template_id), has_broadcast ? " with Broadcast source view" : "");1395 static_cast<int32_t>(template_id), has_broadcast ? " with Broadcast source view" : "");
1302 is_path_supported = false;1396 is_path_supported = false;
1303 }1397 }
1304- plan.simd_index_uses_output_inner_axis = template_id == ascir::TemplateId::kIndirectLoadSimd &&
1305- input_idx == ascgen_utils::indirect_load::kIndexTensorIndex && has_broadcast;
1306 return af::SUCCESS;1398 return af::SUCCESS;
1307}1399}
1308 1400 
@@ -1319,6 +1411,24 @@ af::Status RewriteSkInputPaths(PhysicalViewPreparation &preparation, RewrittenGr
1319 return af::SUCCESS;1411 return af::SUCCESS;
1320}1412}
1321 1413 
1414+af::Status RewriteSimdInputLayouts(PhysicalViewPreparation &preparation) {
1415+ for (InputViewPlan *plan : {&preparation.input, &preparation.index}) {
1416+ if (plan->path_broadcast_index == kInvalidBroadcastIndex || plan->load_transpose == nullptr) {
1417+ continue;
1418+ }
1419+ const auto broadcast = plan->path[static_cast<size_t>(plan->path_broadcast_index)];
1420+ const auto &transpose = plan->load_transpose;
1421+ const auto load = ascgen_utils::indirect_load::GetInputProducer(transpose, 0UL);
1422+ GE_ASSERT_NOTNULL(load);
1423+ const auto owner = transpose->GetOwnerComputeGraph();
1424+ GE_ASSERT_GRAPH_SUCCESS(af::GraphUtils::IsolateNodeOneIO(transpose));
1425+ GE_ASSERT_GRAPH_SUCCESS(af::GraphUtils::RemoveNodeWithoutRelink(owner, transpose));
1426+ // Once the Transpose is removed, the Load and Broadcast share the resolved physical view.
1427+ GE_ASSERT_SUCCESS(ApplyPhysicalView(NodePath{load, broadcast}, plan->layout));
1428+ }
1429+ return af::SUCCESS;
1430+}
1431+ 
1322af::Status RewriteBroadcastPaths(const af::AscNodePtr &indirect_load, ascir::TemplateId template_id,1432af::Status RewriteBroadcastPaths(const af::AscNodePtr &indirect_load, ascir::TemplateId template_id,
1323 PhysicalViewPreparation &preparation, RewrittenGraphAnalysis &analysis,1433 PhysicalViewPreparation &preparation, RewrittenGraphAnalysis &analysis,
1324 bool &is_candidate_legal) {1434 bool &is_candidate_legal) {
@@ -1328,7 +1438,7 @@ af::Status RewriteBroadcastPaths(const af::AscNodePtr &indirect_load, ascir::Tem
1328 if (template_id == ascir::TemplateId::kIndirectLoadSimt) {1438 if (template_id == ascir::TemplateId::kIndirectLoadSimt) {
1329 return RewriteSimtInputBroadcast(indirect_load, preparation.input, is_candidate_legal);1439 return RewriteSimtInputBroadcast(indirect_load, preparation.input, is_candidate_legal);
1330 }1440 }
1331- return af::SUCCESS;1441+ return RewriteSimdInputLayouts(preparation);
1332}1442}
1333 1443 
1334af::Status PreparePhysicalViews(const af::AscNodePtr &indirect_load, ascir::TemplateId template_id,1444af::Status PreparePhysicalViews(const af::AscNodePtr &indirect_load, ascir::TemplateId template_id,
@@ -1336,10 +1446,11 @@ af::Status PreparePhysicalViews(const af::AscNodePtr &indirect_load, ascir::Temp
1336 // 输入/输出个数已在 Generate 入口统一校验，axis 由 AnalyzeRewrittenGraph 统一获取，此处不再重复获取。1446 // 输入/输出个数已在 Generate 入口统一校验，axis 由 AnalyzeRewrittenGraph 统一获取，此处不再重复获取。
1337 is_candidate_legal = true;1447 is_candidate_legal = true;
1338 CollectInputPaths(indirect_load, template_id, preparation);1448 CollectInputPaths(indirect_load, template_id, preparation);
1339- InputViewPlan *const plans[] = {&preparation.input, &preparation.index};
1340 for (size_t input_idx = 0UL; input_idx < kIndirectLoadInputCount; ++input_idx) {1449 for (size_t input_idx = 0UL; input_idx < kIndirectLoadInputCount; ++input_idx) {
1450+ InputViewPlan &plan =
1451+ input_idx == ascgen_utils::indirect_load::kInputTensorIndex ? preparation.input : preparation.index;
1341 bool is_path_supported = true;1452 bool is_path_supported = true;
1342- GE_ASSERT_SUCCESS(AnalyzeInputPath(indirect_load, input_idx, template_id, plans, is_path_supported));1453+ GE_ASSERT_SUCCESS(AnalyzeInputPath(indirect_load, input_idx, template_id, plan, is_path_supported));
1343 if (!is_path_supported) {1454 if (!is_path_supported) {
1344 is_candidate_legal = false;1455 is_candidate_legal = false;
1345 const char *const path_name = input_idx == ascgen_utils::indirect_load::kIndexTensorIndex ? "index" : "input";1456 const char *const path_name = input_idx == ascgen_utils::indirect_load::kIndexTensorIndex ? "index" : "input";
@@ -1450,7 +1561,7 @@ af::Status AnalyzeRewrittenGraph(af::AscGraph &graph, const af::AscNodePtr &indi
1450 1561 
1451af::Status AnnotateSimdTemplateRoles(const RewrittenGraphAnalysis &analysis) {1562af::Status AnnotateSimdTemplateRoles(const RewrittenGraphAnalysis &analysis) {
1452 for (const af::AscNodePtr &node : analysis.input_region) {1563 for (const af::AscNodePtr &node : analysis.input_region) {
1453- if (IsInputRegionBoundary(node)) {1564+ if (IsInputDataSource(node)) {
1454 continue;1565 continue;
1455 }1566 }
1456 const auto role = ascgen_utils::indirect_load::GetTemplateRole(node);1567 const auto role = ascgen_utils::indirect_load::GetTemplateRole(node);
@@ -1461,7 +1572,7 @@ af::Status AnnotateSimdTemplateRoles(const RewrittenGraphAnalysis &analysis) {
1461 }1572 }
1462 if (analysis.simd_index_uses_output_inner_axis) {1573 if (analysis.simd_index_uses_output_inner_axis) {
1463 for (const af::AscNodePtr &node : analysis.index_region) {1574 for (const af::AscNodePtr &node : analysis.index_region) {
1464- if (IsInputRegionBoundary(node)) {1575+ if (IsInputDataSource(node)) {
1465 continue;1576 continue;
1466 }1577 }
1467 GE_ASSERT_SUCCESS(1578 GE_ASSERT_SUCCESS(
@@ -1473,10 +1584,13 @@ af::Status AnnotateSimdTemplateRoles(const RewrittenGraphAnalysis &analysis) {
1473 1584 
1474af::Status ValidateSimtTemplateRegion(const RewrittenGraphAnalysis &analysis, bool &is_candidate_legal) {1585af::Status ValidateSimtTemplateRegion(const RewrittenGraphAnalysis &analysis, bool &is_candidate_legal) {
1475 is_candidate_legal = false;1586 is_candidate_legal = false;
1476- if (!af::ops::IsOps<af::ascir_op::Load>(analysis.input_root) || analysis.index_region.empty()) {1587+ if (!af::ops::IsOps<af::ascir_op::Load>(analysis.input_boundary) || analysis.index_region.empty()) {
1477 return af::SUCCESS;1588 return af::SUCCESS;
1478 }1589 }
1479 for (const af::AscNodePtr &node : analysis.index_region) {1590 for (const af::AscNodePtr &node : analysis.index_region) {
1591+ if (GetLoadTransposeSource(node) != nullptr) {
1592+ continue;
1593+ }
1480 // Compile-time Scalar values are emitted as local constants by the SIMT evaluator.1594 // Compile-time Scalar values are emitted as local constants by the SIMT evaluator.
1481 // ScalarData is runtime input and still requires an explicit context/GM binding.1595 // ScalarData is runtime input and still requires an explicit context/GM binding.
1482 if (af::ops::IsOps<af::ascir_op::ScalarData>(node) || HasControlEdge(node)) {1596 if (af::ops::IsOps<af::ascir_op::ScalarData>(node) || HasControlEdge(node)) {
@@ -1494,8 +1608,13 @@ af::Status ValidateSimtTemplateRegion(const RewrittenGraphAnalysis &analysis, bo
1494}1608}
1495 1609 
1496af::Status AnnotateSimtTemplateRoles(const af::AscNodePtr &indirect_load, const RewrittenGraphAnalysis &analysis) {1610af::Status AnnotateSimtTemplateRoles(const af::AscNodePtr &indirect_load, const RewrittenGraphAnalysis &analysis) {
1611+ GE_ASSERT_NOTNULL(analysis.input_boundary, "IndirectLoad SIMT input boundary is missing.");
1497 GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(1612 GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(
1498- analysis.input_root, ascgen_utils::indirect_load::TemplateRole::kSimtInputBoundary));1613+ analysis.input_boundary, ascgen_utils::indirect_load::TemplateRole::kSimtInputBoundary));
1614+ if (analysis.input_boundary != analysis.input_root) {
1615+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(
1616+ analysis.input_root, ascgen_utils::indirect_load::TemplateRole::kSimtInlineTransform));
1617+ }
1499 GE_ASSERT_SUCCESS(1618 GE_ASSERT_SUCCESS(
1500 ascgen_utils::indirect_load::SetTemplateRole(indirect_load, ascgen_utils::indirect_load::TemplateRole::kSimtOp));1619 ascgen_utils::indirect_load::SetTemplateRole(indirect_load, ascgen_utils::indirect_load::TemplateRole::kSimtOp));
1501 for (const af::AscNodePtr &node : analysis.index_region) {1620 for (const af::AscNodePtr &node : analysis.index_region) {
@@ -1526,7 +1645,7 @@ af::Status AnnotateSimtMainOutputPath(const af::AscNodePtr &indirect_load, const
1526 if (node == nullptr || node == indirect_load || !visited.emplace(node.get()).second) {1645 if (node == nullptr || node == indirect_load || !visited.emplace(node.get()).second) {
1527 continue;1646 continue;
1528 }1647 }
1529- if (!IsInputRegionBoundary(node) && !ScheduleUtils::IsReduce(node) && !ScheduleUtils::IsStore(node)) {1648+ if (!IsInputDataSource(node) && !ScheduleUtils::IsReduce(node) && !ScheduleUtils::IsStore(node)) {
1530 GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(1649 GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(
1531 node, ascgen_utils::indirect_load::TemplateRole::kSimtInlineTransform));1650 node, ascgen_utils::indirect_load::TemplateRole::kSimtInlineTransform));
1532 }1651 }
@@ -1555,7 +1674,7 @@ af::Status AnnotateSimtSideInputClosure(const af::AscNodePtr &branch_node, const
1555 ScheduleUtils::IsReduce(producer) || !visited.emplace(producer.get()).second) {1674 ScheduleUtils::IsReduce(producer) || !visited.emplace(producer.get()).second) {
1556 continue;1675 continue;
1557 }1676 }
1558- if (IsInputRegionBoundary(producer) && !af::ops::IsOps<af::ascir_op::Load>(producer)) {1677+ if (IsInputDataSource(producer) && !af::ops::IsOps<af::ascir_op::Load>(producer)) {
1559 continue;1678 continue;
1560 }1679 }
1561 const auto role = ascgen_utils::indirect_load::GetTemplateRole(producer);1680 const auto role = ascgen_utils::indirect_load::GetTemplateRole(producer);
@@ -1601,7 +1720,7 @@ af::Status AnnotateSimtFanoutBranches(const af::AscNodePtr &indirect_load, const
1601 }1720 }
1602 const auto visit = [&](const af::AscNodePtr &node, bool &stop) -> af::Status {1721 const auto visit = [&](const af::AscNodePtr &node, bool &stop) -> af::Status {
1603 const bool is_fanout_branch =1722 const bool is_fanout_branch =
1604- selected.count(node.get()) == 0UL && !IsInputRegionBoundary(node) && !ScheduleUtils::IsReduce(node);1723+ selected.count(node.get()) == 0UL && !IsInputDataSource(node) && !ScheduleUtils::IsReduce(node);
1605 if (is_fanout_branch) {1724 if (is_fanout_branch) {
1606 GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(1725 GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateRole(
1607 node, ascgen_utils::indirect_load::TemplateRole::kSimtFanoutBranch));1726 node, ascgen_utils::indirect_load::TemplateRole::kSimtFanoutBranch));
@@ -1661,8 +1780,20 @@ af::Status FinalizeTemplate(const af::AscNodePtr &indirect_load, ascir::Template
1661 GE_ASSERT_TRUE(false, "IndirectLoad template id %d is invalid.", static_cast<int32_t>(template_id));1780 GE_ASSERT_TRUE(false, "IndirectLoad template id %d is invalid.", static_cast<int32_t>(template_id));
1662}1781}
1663 1782 
1664-af::Status ApplySkGraphPass(af::AscGraph &graph, const af::AscNodePtr &indirect_load,1783+af::Status ApplySkGraphPass(af::AscGraph &graph, const af::AscNodePtr &indirect_load, bool &is_candidate_legal) {
1665- const RewrittenGraphAnalysis &analysis, bool &is_candidate_legal) {1784+ is_candidate_legal = false;
1785+ for (const auto &node : graph.GetAllNodes()) {
1786+ if (af::ops::IsOps<af::ascir_op::Transpose>(node)) {
1787+ GELOGI("[IndirectLoad] Skip SK candidate: Transpose is only supported by SIMD/SIMT.");
1788+ return af::SUCCESS;
1789+ }
1790+ }
1791+ RewrittenGraphAnalysis analysis;
1792+ GE_ASSERT_SUCCESS(
1793+ AnalyzeRewrittenGraph(graph, indirect_load, ascir::TemplateId::kIndirectLoadSK, is_candidate_legal, analysis));
1794+ if (!is_candidate_legal) {
1795+ return af::SUCCESS;
1796+ }
1666 is_candidate_legal = false;1797 is_candidate_legal = false;
1667 if (!IsSkTemplateCandidateLegal(indirect_load)) {1798 if (!IsSkTemplateCandidateLegal(indirect_load)) {
1668 GELOGI("[IndirectLoad] Reject SK candidate for node[%s]: candidate legality check failed.",1799 GELOGI("[IndirectLoad] Reject SK candidate for node[%s]: candidate legality check failed.",
@@ -1691,14 +1822,14 @@ af::Status ApplyGraphPass(af::AscGraph &graph, const af::AscNodePtr &indirect_lo
1691 GELOGD("[IndirectLoad] Apply graph pass for node[%s], template_id[%d].", indirect_load->GetNamePtr(),1822 GELOGD("[IndirectLoad] Apply graph pass for node[%s], template_id[%d].", indirect_load->GetNamePtr(),
1692 static_cast<int32_t>(template_id));1823 static_cast<int32_t>(template_id));
1693 is_candidate_legal = true;1824 is_candidate_legal = true;
1825+ if (template_id == ascir::TemplateId::kIndirectLoadSK) {
1826+ return ApplySkGraphPass(graph, indirect_load, is_candidate_legal);
1827+ }
1694 RewrittenGraphAnalysis analysis;1828 RewrittenGraphAnalysis analysis;
1695 GE_ASSERT_SUCCESS(AnalyzeRewrittenGraph(graph, indirect_load, template_id, is_candidate_legal, analysis));1829 GE_ASSERT_SUCCESS(AnalyzeRewrittenGraph(graph, indirect_load, template_id, is_candidate_legal, analysis));
1696 if (!is_candidate_legal) {1830 if (!is_candidate_legal) {
1697 return af::SUCCESS;1831 return af::SUCCESS;
1698 }1832 }
1699- if (template_id == ascir::TemplateId::kIndirectLoadSK) {
1700- return ApplySkGraphPass(graph, indirect_load, analysis, is_candidate_legal);
1701- }
1702 GE_ASSERT_SUCCESS(CompleteInputDataTensorAttrs(analysis));1833 GE_ASSERT_SUCCESS(CompleteInputDataTensorAttrs(analysis));
1703 1834 
1704 // The SIMT template overwrites this boundary from its post-reduce layout; without a post Reduce the1835 // The SIMT template overwrites this boundary from its post-reduce layout; without a post Reduce the
@@ -1716,7 +1847,13 @@ af::Status ApplyGraphPass(af::AscGraph &graph, const af::AscNodePtr &indirect_lo
1716 }1847 }
1717 GE_ASSERT_SUCCESS(NormalizeTemplateAxes(graph, indirect_load, template_id, analysis, boundary));1848 GE_ASSERT_SUCCESS(NormalizeTemplateAxes(graph, indirect_load, template_id, analysis, boundary));
1718 GE_ASSERT_SUCCESS(CompletePreservedVectorizedViews(graph, indirect_load));1849 GE_ASSERT_SUCCESS(CompletePreservedVectorizedViews(graph, indirect_load));
1719- return FinalizeTemplate(indirect_load, template_id);1850+ GE_ASSERT_SUCCESS(FinalizeTemplate(indirect_load, template_id));
1851+ bool metadata_supported = false;
1852+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::FinalizeLoweringMetadata(indirect_load, metadata_supported));
1853+ if (!metadata_supported) {
1854+ is_candidate_legal = false;
1855+ }
1856+ return af::SUCCESS;
1720}1857}
1721 1858 
1722bool IsEmbeddingFastPathCapable(const TemplateCase &template_case,1859bool IsEmbeddingFastPathCapable(const TemplateCase &template_case,
@@ -1803,6 +1940,25 @@ af::Status ReduceTaskGraphCount(std::vector<ascir::ImplGraph> &grouped_graphs, c
1803 GE_ASSERT_NOTNULL(backend_spec);1940 GE_ASSERT_NOTNULL(backend_spec);
1804 return ScheduleGroupGraphPartitioner::ReduceGraphCount(grouped_graphs, backend_spec->max_group_num_per_compile_unit);1941 return ScheduleGroupGraphPartitioner::ReduceGraphCount(grouped_graphs, backend_spec->max_group_num_per_compile_unit);
1805}1942}
1943+ 
1944+af::Status FinalizeGroupedGraphLoweringMetadata(std::vector<ascir::ImplGraph> &grouped_graphs) {
1945+ for (auto &graph : grouped_graphs) {
1946+ af::AscNodePtr indirect_load;
1947+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::ValidateSingleIndirectLoadNode(graph, indirect_load));
1948+ if (indirect_load == nullptr) {
1949+ continue;
1950+ }
1951+ const auto template_id = ascir::GetTemplateIdOrDefault(*indirect_load);
1952+ if (template_id != ascir::TemplateId::kIndirectLoadSimd && template_id != ascir::TemplateId::kIndirectLoadSimt) {
1953+ continue;
1954+ }
1955+ bool metadata_supported = false;
1956+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::FinalizeLoweringMetadata(indirect_load, metadata_supported));
1957+ GE_ASSERT_TRUE(metadata_supported, "IndirectLoad lowering metadata is unsupported in grouped graph[%s], node[%s].",
1958+ graph.GetName().c_str(), indirect_load->GetNamePtr());
1959+ }
1960+ return af::SUCCESS;
1961+}
1806} // namespace1962} // namespace
1807 1963 
1808Status IndirectLoadScheduleCaseGenerator::Generate(ascir::HintGraph &graph, std::vector<ascir::ImplGraph> &graphs,1964Status IndirectLoadScheduleCaseGenerator::Generate(ascir::HintGraph &graph, std::vector<ascir::ImplGraph> &graphs,
@@ -1867,6 +2023,7 @@ Status IndirectLoadScheduleCaseGenerator::GeneratorTask(ascir::HintGraph &optimi
1867 "Failed to partition graph");2023 "Failed to partition graph");
1868 GE_ASSERT_SUCCESS(RefreshTaskAxisSizes(task.grouped_graphs, need_update_axis, is_sk_template));2024 GE_ASSERT_SUCCESS(RefreshTaskAxisSizes(task.grouped_graphs, need_update_axis, is_sk_template));
1869 GE_CHK_STATUS_RET(ReduceTaskGraphCount(task.grouped_graphs, options), "Failed to reduce graph count");2025 GE_CHK_STATUS_RET(ReduceTaskGraphCount(task.grouped_graphs, options), "Failed to reduce graph count");
2026+ GE_ASSERT_SUCCESS(FinalizeGroupedGraphLoweringMetadata(task.grouped_graphs));
1870 tasks.emplace_back(std::move(task));2027 tasks.emplace_back(std::move(task));
1871 }2028 }
1872 return af::SUCCESS;2029 return af::SUCCESS;
Mautofuse/optimize/task_generator/transpose_schedule_case_generator.cpp+7-0
@@ -17,6 +17,7 @@
17#include "ascir_utils.h"17#include "ascir_utils.h"
18#include "ascir_ops_utils.h"18#include "ascir_ops_utils.h"
19#include "ascir_ops.h"19#include "ascir_ops.h"
20+#include "indirect_load_utils.h"
20#include "schedule_utils.h"21#include "schedule_utils.h"
21#include "graph_properties_cache.h"22#include "graph_properties_cache.h"
22 23 
@@ -97,6 +98,12 @@ Status TransposeFusionCaseGenerator::TransposeConvertProcess(ascir::HintGraph &g
97 98 
98Status TransposeFusionCaseGenerator::Generate(ascir::HintGraph &graph, std::vector<ascir::ImplGraph> &graphs,99Status TransposeFusionCaseGenerator::Generate(ascir::HintGraph &graph, std::vector<ascir::ImplGraph> &graphs,
99 std::vector<std::string> &score_functions) {100 std::vector<std::string> &score_functions) {
101+ const auto indirect_load = ascgen_utils::indirect_load::FindIndirectLoadNode(graph);
102+ if (indirect_load != nullptr) {
103+ GELOGI("Graph %s has indirect load, skip transpose task generation; handled by IndirectLoad tasks.",
104+ graph.GetName().c_str());
105+ return af::SUCCESS;
106+ }
100 /*107 /*
101 单个Transpose场景：108 单个Transpose场景：
102 场景1： 尾轴转置， 需要UB重排，Transpose节点保留；109 场景1： 尾轴转置， 需要UB重排，Transpose节点保留；
Mautofuse/tests/ut/codegen/api_call/test_codegen_indirect_load_api_call.cpp+180-7
@@ -16,6 +16,8 @@
16#include "codegen_kernel.h"16#include "codegen_kernel.h"
17#include "common_utils.h"17#include "common_utils.h"
18#include "graph/ascendc_ir/utils/asc_tensor_utils.h"18#include "graph/ascendc_ir/utils/asc_tensor_utils.h"
19+#include "graph/ascendc_ir/utils/asc_graph_utils.h"
20+#include "graph/utils/graph_utils.h"
19#include "indirect_load_utils.h"21#include "indirect_load_utils.h"
20#include "platform_context.h"22#include "platform_context.h"
21#include "runtime_stub.h"23#include "runtime_stub.h"
@@ -1013,7 +1015,9 @@ TEST(IndirectLoadApiCallTest, GenerateSimdDenseSelectsDenseApi) {
1013 1015 
1014 EXPECT_NE(result.find("IndirectLoadSimd<half, int32_t, 2, 1>"), std::string::npos);1016 EXPECT_NE(result.find("IndirectLoadSimd<half, int32_t, 2, 1>"), std::string::npos);
1015 EXPECT_EQ(result.find("IndirectLoadSimdStrided"), std::string::npos);1017 EXPECT_EQ(result.find("IndirectLoadSimdStrided"), std::string::npos);
1016- EXPECT_EQ(result.find("indirect_load_simd_params"), std::string::npos);1018+ EXPECT_NE(result.find("IndirectLoadSimdFallback>(0), false>"), std::string::npos);
1019+ EXPECT_NE(result.find("IndirectLoadSimdParams<IndirectLoadSimdCase, 2> indirect_load_simd_params"),
1020+ std::string::npos);
1017}1021}
1018 1022 
1019TEST(IndirectLoadApiCallTest, GenerateSimdDenseGatherSelectsGatherApi) {1023TEST(IndirectLoadApiCallTest, GenerateSimdDenseGatherSelectsGatherApi) {
@@ -1089,7 +1093,7 @@ TEST(IndirectLoadApiCallTest, GenerateFuncDefinitionSimtProducesBodyStruct) {
1089 EXPECT_NE(def.find("using Context = IndirectLoadSimtContext_"), std::string::npos);1093 EXPECT_NE(def.find("using Context = IndirectLoadSimtContext_"), std::string::npos);
1090 EXPECT_NE(def.find("__simt_callee__ __aicore__ inline static"), std::string::npos);1094 EXPECT_NE(def.find("__simt_callee__ __aicore__ inline static"), std::string::npos);
1091 EXPECT_NE(def.find(" Index(uint64_t output_index, const Context &context)"), std::string::npos);1095 EXPECT_NE(def.find(" Index(uint64_t output_index, const Context &context)"), std::string::npos);
1092- EXPECT_NE(def.find(" Output("), std::string::npos);1096+ EXPECT_NE(def.find(" Outputs("), std::string::npos);
1093}1097}
1094 1098 
1095TEST(IndirectLoadApiCallTest, GenerateFuncDefinitionSimtEmitsMultiOutputEvaluator) {1099TEST(IndirectLoadApiCallTest, GenerateFuncDefinitionSimtEmitsMultiOutputEvaluator) {
@@ -1117,6 +1121,79 @@ TEST(IndirectLoadApiCallTest, GenerateFuncDefinitionUsesUint32OffsetsForStaticSh
1117 EXPECT_NE(def.find("uint32_t output_index"), std::string::npos);1121 EXPECT_NE(def.find("uint32_t output_index"), std::string::npos);
1118}1122}
1119 1123 
1124+void InsertSimtIdentityBroadcast(ILTestGraph &g, const char *name, const char *producer_name, const char *consumer_name,
1125+ int32_t input_index, ascir::TensorId tensor_id) {
1126+ const auto producer = g.graph.FindNode(producer_name);
1127+ const auto consumer = g.graph.FindNode(consumer_name);
1128+ ASSERT_NE(producer, nullptr);
1129+ ASSERT_NE(consumer, nullptr);
1130+ Broadcast broadcast(name);
1131+ g.graph.AddNode(broadcast);
1132+ const auto node = g.graph.FindNode(name);
1133+ ASSERT_NE(node, nullptr);
1134+ node->attr.sched.axis = producer->attr.sched.axis;
1135+ node->attr.api.compute_type = af::ComputeType::kComputeBroadcast;
1136+ node->outputs[0].attr.dtype = producer->outputs[0].attr.dtype;
1137+ node->outputs[0].attr.axis = producer->outputs[0].attr.axis;
1138+ node->outputs[0].attr.repeats = producer->outputs[0].attr.repeats;
1139+ node->outputs[0].attr.strides = producer->outputs[0].attr.strides;
1140+ node->outputs[0].attr.mem.tensor_id = tensor_id;
1141+ ASSERT_EQ(SetTemplateRole(node, TemplateRole::kSimtInlineTransform), af::SUCCESS);
1142+ ASSERT_EQ(af::GraphUtils::ReplaceEdgeSrc(producer->GetOutDataAnchor(0), consumer->GetInDataAnchor(input_index),
1143+ node->GetOutDataAnchor(0)),
1144+ af::SUCCESS);
1145+ ASSERT_EQ(af::GraphUtils::AddEdge(producer->GetOutDataAnchor(0), node->GetInDataAnchor(0)), af::SUCCESS);
1146+ ASSERT_EQ(af::AscGraphUtils::GetComputeGraph(g.graph)->TopologicalSorting(), af::SUCCESS);
1147+}
1148+ 
1149+TEST_F(IndirectLoadConstructFromNodesTest, OmitsIdentityScalarTemporaries) {
1150+ ILTestGraph g("simt_identity_chain", 2, 8, 2, 8);
1151+ BuildSimtGraph(g);
1152+ ASSERT_NO_FATAL_FAILURE(InsertSimtIdentityBroadcast(g, "identity0", "indirect_load", "store", 0, 20));
1153+ ASSERT_NO_FATAL_FAILURE(InsertSimtIdentityBroadcast(g, "identity1", "identity0", "store", 0, 21));
1154+ AnnotateSimtTemplate(g.graph);
1155+ std::string definition;
1156+ GenerateSimtFuncDefinitionFromGraph(g, definition);
1157+ EXPECT_EQ(definition.find("v_20 ="), std::string::npos);
1158+ EXPECT_EQ(definition.find("v_21 ="), std::string::npos);
1159+ EXPECT_NE(definition.find("outputs.output0 = value;"), std::string::npos);
1160+}
1161+ 
1162+TEST_F(IndirectLoadConstructFromNodesTest, KeepsGmLoadMaterializedThroughIdentityBroadcast) {
1163+ ILTestGraph g("simt_identity_gm_load");
1164+ BuildSimtPostReduceGraph(g);
1165+ ASSERT_NO_FATAL_FAILURE(InsertSimtIdentityBroadcast(g, "load_identity", "addend_load", "output_transform", 1, 50));
1166+ std::string definition;
1167+ GenerateSimtPostReduceFuncDefinitionFromGraph(g, definition);
1168+ EXPECT_NE(definition.find("half v_50 = context.gm_9["), std::string::npos);
1169+ const auto load = definition.find("context.gm_9[");
1170+ ASSERT_NE(load, std::string::npos);
1171+ EXPECT_EQ(definition.find("context.gm_9[", load + 1), std::string::npos);
1172+}
1173+ 
1174+TEST_F(IndirectLoadConstructFromNodesTest, InlinesCompileTimeScalarWithoutTemporary) {
1175+ ILTestGraph g("simt_inline_scalar");
1176+ BuildSimtPostReduceGraph(g);
1177+ Scalar scalar("constant_bias", g.graph);
1178+ scalar.ir_attr.SetValue("2.5");
1179+ scalar.y.dtype = af::DT_FLOAT16;
1180+ const auto scalar_node = g.graph.FindNode("constant_bias");
1181+ const auto load = g.graph.FindNode("addend_load");
1182+ const auto add = g.graph.FindNode("output_transform");
1183+ ASSERT_NE(scalar_node, nullptr);
1184+ ASSERT_NE(load, nullptr);
1185+ ASSERT_NE(add, nullptr);
1186+ scalar_node->outputs[0].attr.mem.tensor_id = 40;
1187+ ASSERT_EQ(af::GraphUtils::ReplaceEdgeSrc(load->GetOutDataAnchor(0), add->GetInDataAnchor(1),
1188+ scalar_node->GetOutDataAnchor(0)),
1189+ af::SUCCESS);
1190+ ASSERT_EQ(af::AscGraphUtils::GetComputeGraph(g.graph)->TopologicalSorting(), af::SUCCESS);
1191+ std::string definition;
1192+ GenerateSimtPostReduceFuncDefinitionFromGraph(g, definition);
1193+ EXPECT_EQ(definition.find("v_40 ="), std::string::npos);
1194+ EXPECT_NE(definition.find("static_cast<half>(2.5"), std::string::npos);
1195+}
1196+ 
1120TEST(IndirectLoadApiCallTest, GenerateSimtOutputUsesIndexOffsetForMatchingLoadView) {1197TEST(IndirectLoadApiCallTest, GenerateSimtOutputUsesIndexOffsetForMatchingLoadView) {
1121 ILTestGraph g("simt_post_offsets");1198 ILTestGraph g("simt_post_offsets");
1122 std::string definition;1199 std::string definition;
@@ -1139,6 +1216,101 @@ TEST(IndirectLoadApiCallTest, GenerateSimtOutputUsesOutputOffsetForDistinctLoadV
1139 EXPECT_NE(definition.find("context.gm_9[output_index]"), std::string::npos);1216 EXPECT_NE(definition.find("context.gm_9[output_index]"), std::string::npos);
1140}1217}
1141 1218 
1219+TEST(IndirectLoadApiCallTest, GenerateSimtOutputMapsBroadcastLoadToColumn) {
1220+ ILTestGraph g("simt_post_broadcast_offset", 2, 8, 2, 8);
1221+ BuildSimtPostReduceGraph(g);
1222+ // Reuse the side-input fixture with a direct Store, matching the position-bias output chain.
1223+ const auto reduce = g.graph.FindNode("reduce");
1224+ ASSERT_NE(reduce, nullptr);
1225+ ASSERT_EQ(af::GraphUtils::IsolateNodeOneIO(reduce), af::SUCCESS);
1226+ ASSERT_EQ(af::GraphUtils::RemoveNodeWithoutRelink(af::AscGraphUtils::GetComputeGraph(g.graph), reduce), af::SUCCESS);
1227+ const auto addend_load = g.graph.FindNode("addend_load");
1228+ ASSERT_NE(addend_load, nullptr);
1229+ addend_load->outputs[0].attr.repeats = {af::sym::kSymbolOne, g.s3};
1230+ addend_load->outputs[0].attr.strides = {af::sym::kSymbolZero, af::sym::kSymbolOne};
1231+ 
1232+ std::string definition;
1233+ GenerateSimtPostReduceFuncDefinitionFromGraph(g, definition);
1234+ EXPECT_NE(definition.find("context.gm_9[index_coord_0]"), std::string::npos);
1235+ EXPECT_NE(definition.find("index_coord_0 = static_cast<uint32_t>(output_index % 8);"), std::string::npos);
1236+ EXPECT_EQ(definition.find("context.gm_9[output_index]"), std::string::npos);
1237+ EXPECT_EQ(definition.find("context.gm_9[index_offset]"), std::string::npos);
1238+}
1239+ 
1240+void GenerateSimtBroadcastPairDefinition(int32_t second_extent, std::string &definition) {
1241+ SimtCodegenContext context;
1242+ ILTestGraph g("simt_broadcast_pair", 2, 8, 2, 8);
1243+ BuildSimtPostReduceGraph(g);
1244+ const auto reduce = g.graph.FindNode("reduce");
1245+ ASSERT_NE(reduce, nullptr);
1246+ ASSERT_EQ(af::GraphUtils::IsolateNodeOneIO(reduce), af::SUCCESS);
1247+ ASSERT_EQ(af::GraphUtils::RemoveNodeWithoutRelink(af::AscGraphUtils::GetComputeGraph(g.graph), reduce), af::SUCCESS);
1248+ const auto first_load = g.graph.FindNode("addend_load");
1249+ ASSERT_NE(first_load, nullptr);
1250+ first_load->outputs[0].attr.repeats = {af::ops::One, g.s3};
1251+ first_load->outputs[0].attr.strides = {af::ops::Zero, af::ops::One};
1252+ 
1253+ Data extra("extra", g.graph);
1254+ extra.ir_attr.SetIndex(3);
1255+ extra.y.dtype = af::DT_FLOAT16;
1256+ g.graph.FindNode("extra")->outputs[0].attr.mem.tensor_id = 41;
1257+ Load second_load("second_bias_load");
1258+ g.graph.AddNode(second_load);
1259+ second_load.x = extra.y;
1260+ second_load.y.dtype = af::DT_FLOAT16;
1261+ second_load.attr.sched.axis = {g.z2.id, g.z3.id};
1262+ *second_load.y.axis = {g.z2.id, g.z3.id};
1263+ *second_load.y.repeats = {af::ops::One, af::Symbol(second_extent)};
1264+ *second_load.y.strides = {af::ops::Zero, af::ops::One};
1265+ const auto second_node = g.graph.FindNode("second_bias_load");
1266+ ASSERT_NE(second_node, nullptr);
1267+ second_node->outputs[0].attr.mem.tensor_id = 51;
1268+ ASSERT_EQ(SetTemplateRole(second_node, TemplateRole::kSimtDirectGmBoundary), af::SUCCESS);
1269+ 
1270+ Add add("add_second_bias");
1271+ g.graph.AddNode(add);
1272+ const auto add_node = g.graph.FindNode("add_second_bias");
1273+ const auto first_add = g.graph.FindNode("output_transform");
1274+ const auto store = g.graph.FindNode("store");
1275+ ASSERT_NE(add_node, nullptr);
1276+ ASSERT_NE(first_add, nullptr);
1277+ ASSERT_NE(store, nullptr);
1278+ add.y.dtype = af::DT_FLOAT16;
1279+ add.attr.sched.axis = first_add->attr.sched.axis;
1280+ *add.y.axis = first_add->outputs[0].attr.axis;
1281+ *add.y.repeats = first_add->outputs[0].attr.repeats;
1282+ *add.y.strides = first_add->outputs[0].attr.strides;
1283+ add_node->outputs[0].attr.mem.tensor_id = 52;
1284+ ASSERT_EQ(SetTemplateRole(add_node, TemplateRole::kSimtInlineTransform), af::SUCCESS);
1285+ ASSERT_EQ(af::GraphUtils::ReplaceEdgeSrc(first_add->GetOutDataAnchor(0), store->GetInDataAnchor(0),
1286+ add_node->GetOutDataAnchor(0)),
1287+ af::SUCCESS);
1288+ ASSERT_EQ(af::GraphUtils::AddEdge(first_add->GetOutDataAnchor(0), add_node->GetInDataAnchor(0)), af::SUCCESS);
1289+ ASSERT_EQ(af::GraphUtils::AddEdge(second_node->GetOutDataAnchor(0), add_node->GetInDataAnchor(1)), af::SUCCESS);
1290+ ASSERT_EQ(af::AscGraphUtils::GetComputeGraph(g.graph)->TopologicalSorting(), af::SUCCESS);
1291+ GenerateSimtPostReduceFuncDefinitionFromGraph(g, definition);
1292+}
1293+ 
1294+TEST(IndirectLoadApiCallTest, ReusesBroadcastCoordinatesAcrossLoads) {
1295+ std::string definition;
1296+ GenerateSimtBroadcastPairDefinition(8, definition);
1297+ EXPECT_NE(definition.find("context.gm_9[index_coord_0]"), std::string::npos);
1298+ EXPECT_NE(definition.find("context.gm_51[index_coord_0]"), std::string::npos);
1299+ const auto coordinate = definition.find("output_index % 8");
1300+ ASSERT_NE(coordinate, std::string::npos);
1301+ EXPECT_EQ(definition.find("output_index % 8", coordinate + 1), std::string::npos);
1302+ EXPECT_EQ(definition.find("index_coord_1"), std::string::npos);
1303+}
1304+ 
1305+TEST(IndirectLoadApiCallTest, KeepsDistinctPhysicalExtentsInBroadcastCoordinates) {
1306+ std::string definition;
1307+ GenerateSimtBroadcastPairDefinition(4, definition);
1308+ EXPECT_NE(definition.find("output_index % 8"), std::string::npos);
1309+ EXPECT_NE(definition.find("output_index % 4"), std::string::npos);
1310+ EXPECT_NE(definition.find("context.gm_9[index_coord_0]"), std::string::npos);
1311+ EXPECT_NE(definition.find("context.gm_51[index_coord_1]"), std::string::npos);
1312+}
1313+ 
1142TEST(IndirectLoadApiCallTest, GenerateSimtOutputUsesZeroOffsetForScalarLoadView) {1314TEST(IndirectLoadApiCallTest, GenerateSimtOutputUsesZeroOffsetForScalarLoadView) {
1143 ILTestGraph g("simt_post_zero_offset");1315 ILTestGraph g("simt_post_zero_offset");
1144 BuildSimtPostReduceGraph(g);1316 BuildSimtPostReduceGraph(g);
@@ -1263,7 +1435,7 @@ TEST_F(IndirectLoadConstructFromNodesTest, ConnectsRedirectedSimtPostReduceWithP
1263 1435 
1264 std::string generated;1436 std::string generated;
1265 ASSERT_EQ(kernel.root_loop.Generate(kernel.tiler, kernel.tpipe, generated), af::SUCCESS);1437 ASSERT_EQ(kernel.root_loop.Generate(kernel.tiler, kernel.tpipe, generated), af::SUCCESS);
1266- EXPECT_NE(generated.find("IndirectLoadSimtMulti"), std::string::npos);1438+ EXPECT_NE(generated.find("AscendC::IndirectLoadSimt<"), std::string::npos);
1267 EXPECT_NE(generated.find("global_10.GetPhyAddr()"), std::string::npos);1439 EXPECT_NE(generated.find("global_10.GetPhyAddr()"), std::string::npos);
1268}1440}
1269 1441 
@@ -1319,7 +1491,7 @@ TEST_F(IndirectLoadConstructFromNodesTest, GeneratesSimtMultiOutputCall) {
1319 Kernel kernel("simt_multi_output_call");1491 Kernel kernel("simt_multi_output_call");
1320 std::string generated;1492 std::string generated;
1321 ASSERT_EQ(ParseGraphAndGenerateLoop(g, {"x", "idx"}, kernel, generated, {"y", "second_y"}), af::SUCCESS);1493 ASSERT_EQ(ParseGraphAndGenerateLoop(g, {"x", "idx"}, kernel, generated, {"y", "second_y"}), af::SUCCESS);
1322- EXPECT_NE(generated.find("IndirectLoadSimtMulti"), std::string::npos);1494+ EXPECT_NE(generated.find("AscendC::IndirectLoadSimt<"), std::string::npos);
1323 EXPECT_NE(generated.find("OutputTargets"), std::string::npos);1495 EXPECT_NE(generated.find("OutputTargets"), std::string::npos);
1324 EXPECT_NE(generated.find("global_2.GetPhyAddr()"), std::string::npos);1496 EXPECT_NE(generated.find("global_2.GetPhyAddr()"), std::string::npos);
1325 EXPECT_NE(generated.find("global_10.GetPhyAddr()"), std::string::npos);1497 EXPECT_NE(generated.find("global_10.GetPhyAddr()"), std::string::npos);
@@ -1336,7 +1508,8 @@ TEST_F(IndirectLoadConstructFromNodesTest, GeneratesStaticPowerOfTwoPolicyWithUi
1336 Kernel kernel("simt_static_power_of_two");1508 Kernel kernel("simt_static_power_of_two");
1337 std::string generated;1509 std::string generated;
1338 ASSERT_EQ(ParseGraphAndGenerateLoop(g, {"x", "idx"}, kernel, generated), af::SUCCESS);1510 ASSERT_EQ(ParseGraphAndGenerateLoop(g, {"x", "idx"}, kernel, generated), af::SUCCESS);
1339- EXPECT_NE(generated.find("IndirectLoadSimtStaticPowerOfTwoPolicy<uint32_t, 1ULL, 8ULL, 1ULL, 8ULL>"),1511+ EXPECT_NE(generated.find("IndirectLoadSimtCase>(0), uint32_t, 2, 1, 1ULL, 8ULL, 1ULL, 8ULL"), std::string::npos);
1512+ EXPECT_NE(generated.find("IndirectLoadSimtParams<IndirectLoadSimtCase> indirect_load_simt_params{}"),
1340 std::string::npos);1513 std::string::npos);
1341}1514}
1342 1515 
@@ -1351,8 +1524,8 @@ TEST_F(IndirectLoadConstructFromNodesTest, GeneratesStaticInnerPolicyForLastAxis
1351 Kernel kernel("simt_static_inner");1524 Kernel kernel("simt_static_inner");
1352 std::string generated;1525 std::string generated;
1353 ASSERT_EQ(ParseGraphAndGenerateLoop(g, {"x", "idx"}, kernel, generated), af::SUCCESS);1526 ASSERT_EQ(ParseGraphAndGenerateLoop(g, {"x", "idx"}, kernel, generated), af::SUCCESS);
1354- EXPECT_NE(generated.find("IndirectLoadSimtStaticInnerPolicy<uint32_t, 1ULL, 1ULL, 7ULL>"), std::string::npos);1527+ EXPECT_NE(generated.find("IndirectLoadSimtCase>(1), uint32_t, 2, 1, 1ULL, 7ULL, 1ULL, 7ULL"), std::string::npos);
1355- EXPECT_NE(generated.find(", 7);"), std::string::npos);1528+ EXPECT_NE(generated.find("indirect_load_simt_params{static_cast<uint32_t>(7)}"), std::string::npos);
1356}1529}
1357 1530 
1358TEST_F(IndirectLoadConstructFromNodesTest, GeneratesSimtPostReduceCallWithUbOutput) {1531TEST_F(IndirectLoadConstructFromNodesTest, GeneratesSimtPostReduceCallWithUbOutput) {
Mautofuse/tests/ut/optimize/task_generator/test_indirect_load_schedule_case_generator.cpp+98-2
@@ -309,7 +309,7 @@ void BuildBroadcastInputPreChain(const af::AscOpOutput &broadcast_output, const
309}309}
310 310 
311af::AscGraph BuildIndirectLoadBroadcastGraph(bool with_input_element = true, bool with_index_broadcast = false,311af::AscGraph BuildIndirectLoadBroadcastGraph(bool with_input_element = true, bool with_index_broadcast = false,
312- int64_t axis = 1L) {312+ int64_t axis = 1L, bool transpose_input = false) {
313 af::AscGraph graph("indirect_load_broadcast_ut_graph");313 af::AscGraph graph("indirect_load_broadcast_ut_graph");
314 const BroadcastGraphView view = CreateBroadcastGraphView(graph);314 const BroadcastGraphView view = CreateBroadcastGraphView(graph);
315 const std::vector<af::Expression> source_strides = {view.source_repeats[2], view.source_repeats[2], af::ops::One};315 const std::vector<af::Expression> source_strides = {view.source_repeats[2], view.source_repeats[2], af::ops::One};
@@ -325,9 +325,20 @@ af::AscGraph BuildIndirectLoadBroadcastGraph(bool with_input_element = true, boo
325 af::ascir_op::Load input_load("broadcast_input_load");325 af::ascir_op::Load input_load("broadcast_input_load");
326 input_load.x = x.y;326 input_load.x = x.y;
327 SetNodeView(input_load, af::DT_FLOAT16, view.source_axes, view.source_repeats, source_strides);327 SetNodeView(input_load, af::DT_FLOAT16, view.source_axes, view.source_repeats, source_strides);
328+ af::ascir_op::Transpose transpose("broadcast_input_transpose");
329+ if (transpose_input) {
330+ auto physical_axes = view.source_axes;
331+ auto physical_repeats = view.source_repeats;
332+ std::swap(physical_axes[0], physical_axes[2]);
333+ std::swap(physical_repeats[0], physical_repeats[2]);
334+ SetNodeView(input_load, af::DT_FLOAT16, physical_axes, physical_repeats,
335+ {af::Symbol(2), af::Symbol(2), af::ops::One});
336+ transpose.x = input_load.y;
337+ SetNodeView(transpose, af::DT_FLOAT16, view.source_axes, view.source_repeats, source_strides);
338+ }
328 af::ascir_op::Broadcast broadcast("input_broadcast");339 af::ascir_op::Broadcast broadcast("input_broadcast");
329 broadcast.attr.api.compute_type = af::ComputeType::kComputeBroadcast;340 broadcast.attr.api.compute_type = af::ComputeType::kComputeBroadcast;
330- broadcast.x = input_load.y;341+ broadcast.x = transpose_input ? transpose.y : input_load.y;
331 SetNodeView(broadcast, af::DT_FLOAT16, view.input_axes, view.input_repeats, input_strides);342 SetNodeView(broadcast, af::DT_FLOAT16, view.input_axes, view.input_repeats, input_strides);
332 auto index_source_repeats = view.output_repeats;343 auto index_source_repeats = view.output_repeats;
333 auto index_source_strides = index_strides;344 auto index_source_strides = index_strides;
@@ -1139,6 +1150,91 @@ TEST(IndirectLoadScheduleCaseGeneratorTest, SimtRewritesInputBroadcastAndSetsTem
1139 EXPECT_TRUE(indirect_load->outputs()[0]->attr.vectorized_strides.empty());1150 EXPECT_TRUE(indirect_load->outputs()[0]->attr.vectorized_strides.empty());
1140}1151}
1141 1152 
1153+TEST(IndirectLoadScheduleCaseGeneratorTest, BroadcastAfterTransposePreservesSourcePhysicalView) {
1154+ PlatformContextReset platform_reset;
1155+ ge::PlatformContext::GetInstance().Reset();
1156+ ge::PlatformInfo platform_info;
1157+ platform_info.soc_ver = "3510";
1158+ platform_info.ub_size = 256 * 1024;
1159+ platform_info.aiv_num = 48;
1160+ ge::PlatformContext::GetInstance().SetPlatformInfo(platform_info);
1161+ auto graph = BuildIndirectLoadBroadcastGraph(false, false, 2L, true);
1162+ const auto source = graph.FindNode("broadcast_input_load")->outputs[0].attr;
1163+ optimize::IndirectLoadScheduleCaseGenerator generator;
1164+ std::vector<af::AscGraph> graphs;
1165+ std::vector<std::string> score_functions;
1166+ ASSERT_EQ(generator.Generate(graph, graphs, score_functions), af::SUCCESS);
1167+ ASSERT_EQ(graphs.size(), 3UL);
1168+ for (auto &candidate : graphs) {
1169+ const auto il = candidate.FindNode("indirect_load");
1170+ ASSERT_NE(il, nullptr);
1171+ EXPECT_NE(ascir::GetTemplateIdOrDefault(*il), ascir::TemplateId::kIndirectLoadSK);
1172+ const auto load = candidate.FindNode("broadcast_input_load");
1173+ ASSERT_NE(load, nullptr);
1174+ const auto transpose = candidate.FindNode("broadcast_input_transpose");
1175+ ascgen_utils::indirect_load::TemplateLogicalView view;
1176+ ASSERT_EQ(ascgen_utils::indirect_load::GetTemplateLogicalView(il, view), af::SUCCESS);
1177+ // The scheduler must publish the final GM address strides for SIMT as well as SIMD.
1178+ EXPECT_EQ(view.input.strides, (std::vector<af::Expression>{af::ops::One, af::ops::Zero, af::Symbol(2)}));
1179+ if (ascir::GetTemplateIdOrDefault(*il) == ascir::TemplateId::kIndirectLoadSimd) {
1180+ EXPECT_EQ(transpose, nullptr);
1181+ EXPECT_EQ(load->outputs[0].attr.strides, view.input.strides);
1182+ continue;
1183+ }
1184+ ASSERT_NE(transpose, nullptr);
1185+ EXPECT_EQ(load->outputs[0].attr.axis.front(), transpose->outputs[0].attr.axis.back());
1186+ EXPECT_EQ(load->outputs[0].attr.axis.back(), transpose->outputs[0].attr.axis.front());
1187+ EXPECT_EQ(load->outputs[0].attr.repeats, source.repeats);
1188+ EXPECT_EQ(load->outputs[0].attr.strides, source.strides);
1189+ }
1190+}
1191+ 
1192+TEST(IndirectLoadScheduleCaseGeneratorTest, PreservesSharedOrControlledTransposeLoadBoundary) {
1193+ PlatformContextReset platform_reset;
1194+ ge::PlatformContext::GetInstance().Reset();
1195+ ge::PlatformInfo platform_info;
1196+ platform_info.soc_ver = "3510";
1197+ platform_info.ub_size = 256 * 1024;
1198+ platform_info.aiv_num = 48;
1199+ ge::PlatformContext::GetInstance().SetPlatformInfo(platform_info);
1200+ for (const bool control_edge : {false, true}) {
1201+ SCOPED_TRACE(control_edge);
1202+ auto graph = BuildIndirectLoadBroadcastGraph(false, false, 2L, true);
1203+ const auto load = graph.FindNode("broadcast_input_load");
1204+ ASSERT_NE(load, nullptr);
1205+ const auto source = load->outputs[0].attr;
1206+ if (control_edge) {
1207+ const auto index = graph.FindNode("broadcast_index_load");
1208+ ASSERT_NE(index, nullptr);
1209+ ASSERT_EQ(af::GraphUtils::AddEdge(load->GetOutControlAnchor(), index->GetInControlAnchor()), af::SUCCESS);
1210+ } else {
1211+ af::ascir_op::Store shared_store("shared_store");
1212+ graph.AddNode(shared_store);
1213+ SetNodeView(shared_store, source.dtype, source.axis, source.repeats, source.strides);
1214+ ASSERT_EQ(af::GraphUtils::AddEdge(load->GetOutDataAnchor(0), graph.FindNode("shared_store")->GetInDataAnchor(0)),
1215+ af::SUCCESS);
1216+ af::ascir_op::Output shared_output("shared_output");
1217+ graph.AddNode(shared_output);
1218+ shared_output.ir_attr.SetIndex(1);
1219+ shared_output.x = shared_store.y;
1220+ SetNodeView(shared_output, source.dtype, source.axis, source.repeats, source.strides);
1221+ }
1222+ optimize::IndirectLoadScheduleCaseGenerator generator;
1223+ std::vector<af::AscGraph> graphs;
1224+ std::vector<std::string> score_functions;
1225+ ASSERT_EQ(generator.Generate(graph, graphs, score_functions), af::SUCCESS);
1226+ ASSERT_EQ(graphs.size(), 2UL);
1227+ for (auto &candidate : graphs) {
1228+ const auto input_load = candidate.FindNode("broadcast_input_load");
1229+ ASSERT_NE(input_load, nullptr);
1230+ EXPECT_NE(candidate.FindNode("broadcast_input_transpose"), nullptr);
1231+ EXPECT_EQ(input_load->outputs[0].attr.repeats, source.repeats);
1232+ EXPECT_EQ(input_load->outputs[0].attr.strides, source.strides);
1233+ EXPECT_TRUE(ascgen_utils::indirect_load::ShouldApplyInputInnerVectorization(input_load));
1234+ }
1235+ }
1236+}
1237+ 
1142TEST(IndirectLoadScheduleCaseGeneratorTest, BroadcastDirectPathUsesPhysicalViewForAllTemplates) {1238TEST(IndirectLoadScheduleCaseGeneratorTest, BroadcastDirectPathUsesPhysicalViewForAllTemplates) {
1143 PlatformContextReset platform_reset;1239 PlatformContextReset platform_reset;
1144 ge::PlatformContext::GetInstance().Reset();1240 ge::PlatformContext::GetInstance().Reset();
Mautofuse/tests/v35/st/backend_e2e_v2/indirect_load_store_test/CMakeLists.txt+38-0
@@ -343,6 +343,35 @@ mark_indirect_load_codegen_and_e2e(indirect_load_broadcast_retained_simd_test IL
343add_indirect_load_broadcast_test(indirect_load_broadcast_retained_simt_test simt 0 0 10 0 0 0 0)343add_indirect_load_broadcast_test(indirect_load_broadcast_retained_simt_test simt 0 0 10 0 0 0 0)
344mark_indirect_load_codegen_and_e2e(indirect_load_broadcast_retained_simt_test IL_RETAIN_BROADCAST=1)344mark_indirect_load_codegen_and_e2e(indirect_load_broadcast_retained_simt_test IL_RETAIN_BROADCAST=1)
345 345 
346+# Retain the two-input Transpose baseline. The other existing targets below compose input Transpose,
347+# Broadcast, a plain index, and an output-side broadcast Load without increasing the target count.
348+add_indirect_load_broadcast_test(indirect_load_both_transpose_simd_test simd 0 0 0 0 0 0 0)
349+mark_indirect_load_static_shape(indirect_load_both_transpose_simd_test)
350+mark_indirect_load_codegen_and_e2e(indirect_load_both_transpose_simd_test
351+ IL_BOTH_TRANSPOSE=1 IL_OUTPUT_S0=2 IL_OUTPUT_S1=2 IL_OUTPUT_S2=1 IL_OUTPUT_S3=4)
352+ 
353+# Exercise both SIMD tiling cases explicitly. Mixed cases transpose across the gather axis, broadcast
354+# the singleton suffix, and add a column-only GM bias; non-square shapes and non-symmetric data expose
355+# both permutation mistakes and output-linear offsets incorrectly used for the broadcast side input.
356+add_indirect_load_broadcast_test(indirect_load_both_transpose_simd_fallback_test simd 0 0 8 0 0 1 0)
357+mark_indirect_load_static_shape(indirect_load_both_transpose_simd_fallback_test)
358+mark_indirect_load_codegen_and_e2e(indirect_load_both_transpose_simd_fallback_test
359+ IL_INPUT_TRANSPOSE=1 IL_OUTPUT_BROADCAST_ADD=1 IL_FORCE_TILING_CASE=0
360+ IL_OUTPUT_S0=2 IL_OUTPUT_S1=3 IL_OUTPUT_S2=4 IL_OUTPUT_S3=32)
361+mark_indirect_load_red(indirect_load_both_transpose_simd_fallback_test)
362+ 
363+add_indirect_load_broadcast_test(indirect_load_both_transpose_simd_nddma_test simd 0 0 8 0 0 1 0)
364+mark_indirect_load_static_shape(indirect_load_both_transpose_simd_nddma_test)
365+mark_indirect_load_codegen_and_e2e(indirect_load_both_transpose_simd_nddma_test
366+ IL_INPUT_TRANSPOSE=1 IL_OUTPUT_BROADCAST_ADD=1 IL_FORCE_TILING_CASE=1
367+ IL_OUTPUT_S0=2 IL_OUTPUT_S1=3 IL_OUTPUT_S2=4 IL_OUTPUT_S3=32)
368+ 
369+add_indirect_load_broadcast_test(indirect_load_both_transpose_simt_test simt 0 0 8 0 0 1 0)
370+mark_indirect_load_static_shape(indirect_load_both_transpose_simt_test)
371+mark_indirect_load_codegen_and_e2e(indirect_load_both_transpose_simt_test
372+ IL_INPUT_TRANSPOSE=1 IL_OUTPUT_BROADCAST_ADD=1
373+ IL_OUTPUT_S0=2 IL_OUTPUT_S1=3 IL_OUTPUT_S2=4 IL_OUTPUT_S3=32)
374+ 
346# Keep one SK Broadcast regression outside the SIMD/SIMT matrix.375# Keep one SK Broadcast regression outside the SIMD/SIMT matrix.
347add_indirect_load_broadcast_test(indirect_load_broadcast_elements_sk_test sk 2 2 3 0 0)376add_indirect_load_broadcast_test(indirect_load_broadcast_elements_sk_test sk 2 2 3 0 0)
348add_indirect_load_broadcast_test(indirect_load_broadcast_index_physical_view_simt_test simt 2 2 3 0 0)377add_indirect_load_broadcast_test(indirect_load_broadcast_index_physical_view_simt_test simt 2 2 3 0 0)
@@ -417,6 +446,11 @@ do_backend_e2e_st_test(indirect_load_graph_hint_simd_repro
417mark_indirect_load_codegen_and_e2e(indirect_load_graph_hint_simd_repro446mark_indirect_load_codegen_and_e2e(indirect_load_graph_hint_simd_repro
418 IL_CASE_BROADCAST_WHERE IL_GRAPH_HINT_SIMD_REPRO)447 IL_CASE_BROADCAST_WHERE IL_GRAPH_HINT_SIMD_REPRO)
419 448 
449+# Exact reproduction of the user-provided embedding slice GraphHint.
450+add_indirect_load_graph_case(indirect_load_graph_hint_embedding_slice
451+ indirect_load_graph_hint_embedding_slice 0
452+ IL_CASE_BROADCAST_WHERE IL_GRAPH_HINT_EMBEDDING_SLICE)
453+ 
420# Minimal SIMT reproducer: a partially zero-stride mask view is enough to trigger the bad evaluator offset.454# Minimal SIMT reproducer: a partially zero-stride mask view is enough to trigger the bad evaluator offset.
421add_indirect_load_graph_case(indirect_load_user_masked_embedding_minimal455add_indirect_load_graph_case(indirect_load_user_masked_embedding_minimal
422 indirect_load_user_masked_embedding_minimal 1456 indirect_load_user_masked_embedding_minimal 1
@@ -492,6 +526,10 @@ add_indirect_load_user_fanout_case(indirect_load_user_fanout_post_reduce_simd
492add_indirect_load_user_fanout_case(indirect_load_user_fanout_post_reduce_simt526add_indirect_load_user_fanout_case(indirect_load_user_fanout_post_reduce_simt
493 IL_USER_FANOUT_POST IL_USER_FANOUT_REDUCE)527 IL_USER_FANOUT_POST IL_USER_FANOUT_REDUCE)
494 528 
529+# Regression: two GM output chains, with one chain reading a broadcast GM side input.
530+add_indirect_load_graph_case(indirect_load_user_fanout_side_input_simt
531+ indirect_load_user_fanout_side_input 1 IL_USER_FANOUT_SIDE_INPUT)
532+ 
495# Minimal red regression: a post-IndirectLoad fanout Add consumes an independent Load side input.533# Minimal red regression: a post-IndirectLoad fanout Add consumes an independent Load side input.
496add_indirect_load_graph_case(indirect_load_user_side_input_fanout_simd534add_indirect_load_graph_case(indirect_load_user_side_input_fanout_simd
497 indirect_load_user_side_input_fanout_simd 0535 indirect_load_user_side_input_fanout_simd 0
Mautofuse/tests/v35/st/backend_e2e_v2/indirect_load_store_test/test_e2e_indirect_load_store_kernel.cpp+206-11
@@ -11,6 +11,7 @@
11#ifndef AUTOFUSE_TESTS_V35_ST_BACKEND_E2E_V2_INDIRECT_LOAD_STORE_TEST_INDIRECT_LOAD_KERNEL_TEST_COMMON_H_11#ifndef AUTOFUSE_TESTS_V35_ST_BACKEND_E2E_V2_INDIRECT_LOAD_STORE_TEST_INDIRECT_LOAD_KERNEL_TEST_COMMON_H_
12#define AUTOFUSE_TESTS_V35_ST_BACKEND_E2E_V2_INDIRECT_LOAD_STORE_TEST_INDIRECT_LOAD_KERNEL_TEST_COMMON_H_12#define AUTOFUSE_TESTS_V35_ST_BACKEND_E2E_V2_INDIRECT_LOAD_STORE_TEST_INDIRECT_LOAD_KERNEL_TEST_COMMON_H_
13 13 
14+#include <algorithm>
14#include <cstdint>15#include <cstdint>
15#include <memory>16#include <memory>
16#include <vector>17#include <vector>
@@ -21,12 +22,19 @@
21#include "autofuse_tiling_data.h"22#include "autofuse_tiling_data.h"
22 23 
23#if !defined(IL_CASE_STORE) && !defined(IL_CASE_MIXED)24#if !defined(IL_CASE_STORE) && !defined(IL_CASE_MIXED)
25+struct ResLimit;
24extern "C" int64_t AutofuseTiling(AutofuseTilingData *, uint32_t *, uint32_t *, uint32_t, uint32_t);26extern "C" int64_t AutofuseTiling(AutofuseTilingData *, uint32_t *, uint32_t *, uint32_t, uint32_t);
27+extern "C" int64_t AutofuseTilingWithConfig(const char *, AutofuseTilingData *, uint32_t *, uint32_t *, ResLimit *,
28+ int32_t);
25#endif29#endif
26 30 
27#if defined(IL_USER_FANOUT)31#if defined(IL_USER_FANOUT)
28extern "C" __global__ __aicore__ void user_fanout(GM_ADDR indices, GM_ADDR embedding, GM_ADDR weight, GM_ADDR output0,32extern "C" __global__ __aicore__ void user_fanout(GM_ADDR indices, GM_ADDR embedding, GM_ADDR weight, GM_ADDR output0,
29 GM_ADDR output1, GM_ADDR workspace, GM_ADDR gm_tiling_data);33 GM_ADDR output1, GM_ADDR workspace, GM_ADDR gm_tiling_data);
34+#elif defined(IL_USER_FANOUT_SIDE_INPUT)
35+extern "C" __global__ __aicore__ void user_fanout_side_input(GM_ADDR indices, GM_ADDR embedding, GM_ADDR side_input,
36+ GM_ADDR output0, GM_ADDR output1, GM_ADDR workspace,
37+ GM_ADDR gm_tiling_data);
30#elif defined(IL_USER_SIDE_INPUT_FANOUT)38#elif defined(IL_USER_SIDE_INPUT_FANOUT)
31extern "C" __global__ __aicore__ void user_side_input_fanout(GM_ADDR input0, GM_ADDR input1, GM_ADDR input2,39extern "C" __global__ __aicore__ void user_side_input_fanout(GM_ADDR input0, GM_ADDR input1, GM_ADDR input2,
32 GM_ADDR input3, GM_ADDR input4, GM_ADDR input5,40 GM_ADDR input3, GM_ADDR input4, GM_ADDR input5,
@@ -83,6 +91,10 @@ struct KernelData {
83struct KernelTiling {91struct KernelTiling {
84 explicit KernelTiling(uint32_t core_num = 48U) : workspace(nullptr, GmFree) {92 explicit KernelTiling(uint32_t core_num = 48U) : workspace(nullptr, GmFree) {
85 EXPECT_EQ(AutofuseTiling(&data, &workspace_size, &block_dim, core_num, 192U * 1024U), 0);93 EXPECT_EQ(AutofuseTiling(&data, &workspace_size, &block_dim, core_num, 192U * 1024U), 0);
94+#ifdef IL_FORCE_TILING_CASE
95+ EXPECT_EQ(AutofuseTilingWithConfig(nullptr, &data, &workspace_size, &block_dim, nullptr, IL_FORCE_TILING_CASE), 0);
96+ EXPECT_EQ(data.graph0_result0_g0_tiling_data.tiling_key, static_cast<uint32_t>(IL_FORCE_TILING_CASE));
97+#endif
86 EXPECT_GT(data.block_dim, 0U);98 EXPECT_GT(data.block_dim, 0U);
87 if (workspace_size != 0U) {99 if (workspace_size != 0U) {
88 workspace.reset(reinterpret_cast<uint8_t *>(AscendC::GmAlloc(workspace_size)));100 workspace.reset(reinterpret_cast<uint8_t *>(AscendC::GmAlloc(workspace_size)));
@@ -189,7 +201,9 @@ using OutputType = float;
189#else201#else
190using OutputType = DataType;202using OutputType = DataType;
191#endif203#endif
192-#ifdef IL_INDEX_INT64204+#if IL_BOTH_TRANSPOSE
205+using IndexType = int32_t;
206+#elif defined(IL_INDEX_INT64)
193using IndexType = int64_t;207using IndexType = int64_t;
194#else208#else
195using IndexType = int32_t;209using IndexType = int32_t;
@@ -212,6 +226,12 @@ float BesselK0Reference(float x) {
212}226}
213 227 
214void RunTiling(AutofuseTilingData &tiling, uint32_t &workspace_size, uint32_t &block_dim) {228void RunTiling(AutofuseTilingData &tiling, uint32_t &workspace_size, uint32_t &block_dim) {
229+#ifdef IL_FORCE_TILING_CASE
230+ ASSERT_EQ(AutofuseTilingWithConfig(nullptr, &tiling, &workspace_size, &block_dim, nullptr, IL_FORCE_TILING_CASE), 0);
231+ ASSERT_EQ(tiling.graph0_tiling_key, 0U);
232+ ASSERT_EQ(tiling.graph0_result0_g0_tiling_data.tiling_key, static_cast<uint32_t>(IL_FORCE_TILING_CASE));
233+ return;
234+#endif
215#ifdef IL_STATIC_SHAPE235#ifdef IL_STATIC_SHAPE
216 AutofuseTiling(&tiling, &workspace_size, &block_dim, 48, 192 * 1024);236 AutofuseTiling(&tiling, &workspace_size, &block_dim, 48, 192 * 1024);
217#elif IL_RANK == 2237#elif IL_RANK == 2
@@ -648,6 +668,12 @@ TEST(E2EIndirectLoadStore, GeneratedKernelMatchesReference) {
648#ifndef IL_INPUT_ABS_BEFORE_BROADCAST668#ifndef IL_INPUT_ABS_BEFORE_BROADCAST
649#define IL_INPUT_ABS_BEFORE_BROADCAST 0669#define IL_INPUT_ABS_BEFORE_BROADCAST 0
650#endif670#endif
671+#ifndef IL_BOTH_TRANSPOSE
672+#define IL_BOTH_TRANSPOSE 0
673+#endif
674+#ifndef IL_INPUT_TRANSPOSE
675+#define IL_INPUT_TRANSPOSE 0
676+#endif
651#ifndef IL_OUTPUT_S0677#ifndef IL_OUTPUT_S0
652#define IL_OUTPUT_S0 4678#define IL_OUTPUT_S0 4
653#endif679#endif
@@ -664,6 +690,9 @@ TEST(E2EIndirectLoadStore, GeneratedKernelMatchesReference) {
664#if IL_AIC_REPRO690#if IL_AIC_REPRO
665extern "C" __global__ __aicore__ void indirect_load_aic_repro(GM_ADDR input, GM_ADDR index, GM_ADDR output,691extern "C" __global__ __aicore__ void indirect_load_aic_repro(GM_ADDR input, GM_ADDR index, GM_ADDR output,
666 GM_ADDR workspace, GM_ADDR tiling);692 GM_ADDR workspace, GM_ADDR tiling);
693+#elif IL_OUTPUT_BROADCAST_ADD
694+extern "C" __global__ __aicore__ void indirect_load_broadcast_test(GM_ADDR x, GM_ADDR index, GM_ADDR bias, GM_ADDR y,
695+ GM_ADDR workspace, GM_ADDR tiling);
667#else696#else
668extern "C" __global__ __aicore__ void indirect_load_broadcast_test(GM_ADDR x, GM_ADDR index, GM_ADDR y,697extern "C" __global__ __aicore__ void indirect_load_broadcast_test(GM_ADDR x, GM_ADDR index, GM_ADDR y,
669 GM_ADDR workspace, GM_ADDR tiling);698 GM_ADDR workspace, GM_ADDR tiling);
@@ -694,7 +723,11 @@ void InitializeAicReproData(DataType *input, IndexType *index, DataType *expecte
694}723}
695#else724#else
696using DataType = half;725using DataType = half;
726+#if IL_BOTH_TRANSPOSE
727+using IndexType = int32_t;
728+#else
697using IndexType = int64_t;729using IndexType = int64_t;
730+#endif
698constexpr std::array<int32_t, 4> kOutputShape = {IL_OUTPUT_S0, IL_OUTPUT_S1, IL_OUTPUT_S2, IL_OUTPUT_S3};731constexpr std::array<int32_t, 4> kOutputShape = {IL_OUTPUT_S0, IL_OUTPUT_S1, IL_OUTPUT_S2, IL_OUTPUT_S3};
699constexpr bool kInputBroadcast = IL_INPUT_BROADCAST;732constexpr bool kInputBroadcast = IL_INPUT_BROADCAST;
700constexpr bool kIndexBroadcast = IL_INDEX_BROADCAST;733constexpr bool kIndexBroadcast = IL_INDEX_BROADCAST;
@@ -707,6 +740,8 @@ constexpr bool kRetainBroadcast = IL_RETAIN_BROADCAST;
707constexpr bool kBroadcastPostReduce = IL_BROADCAST_POST_REDUCE;740constexpr bool kBroadcastPostReduce = IL_BROADCAST_POST_REDUCE;
708constexpr bool kIndexMixedView = IL_INDEX_MIXED_VIEW;741constexpr bool kIndexMixedView = IL_INDEX_MIXED_VIEW;
709constexpr bool kInputAbsBeforeBroadcast = IL_INPUT_ABS_BEFORE_BROADCAST;742constexpr bool kInputAbsBeforeBroadcast = IL_INPUT_ABS_BEFORE_BROADCAST;
743+constexpr bool kBothTranspose = IL_BOTH_TRANSPOSE;
744+constexpr bool kInputTranspose = IL_INPUT_TRANSPOSE;
710constexpr uint32_t kBroadcastAxesMask = IL_BROADCAST_AXES_MASK;745constexpr uint32_t kBroadcastAxesMask = IL_BROADCAST_AXES_MASK;
711 746 
712constexpr std::array<int32_t, 4> MakeBroadcastSourceShape() {747constexpr std::array<int32_t, 4> MakeBroadcastSourceShape() {
@@ -720,8 +755,19 @@ constexpr std::array<int32_t, 4> MakeBroadcastSourceShape() {
720}755}
721 756 
722constexpr std::array<int32_t, 4> kBroadcastSourceShape = MakeBroadcastSourceShape();757constexpr std::array<int32_t, 4> kBroadcastSourceShape = MakeBroadcastSourceShape();
723-constexpr std::array<int32_t, 4> kInputShape = kInputBroadcast && !kComplexSimt ? kBroadcastSourceShape : kOutputShape;758+constexpr std::array<int32_t, 4> MakeTransposeSourceShape(const std::array<int32_t, 4> &shape) {
724-constexpr std::array<int32_t, 4> kIndexShape = kIndexBroadcast ? kBroadcastSourceShape : kOutputShape;759+ return {shape[1], shape[0], shape[2], shape[3]};
760+}
761+constexpr std::array<int32_t, 4> kTransposeInputShape = MakeTransposeSourceShape(kOutputShape);
762+constexpr std::array<int32_t, 4> MakeGatherTransposeSourceShape(std::array<int32_t, 4> shape) {
763+ return {shape[2], shape[1], shape[0], shape[3]};
764+}
765+constexpr std::array<int32_t, 4> kInputShape =
766+ kBothTranspose ? kTransposeInputShape
767+ : kInputTranspose ? MakeGatherTransposeSourceShape(kInputBroadcast ? kBroadcastSourceShape : kOutputShape)
768+ : (kInputBroadcast && !kComplexSimt ? kBroadcastSourceShape : kOutputShape);
769+constexpr std::array<int32_t, 4> kIndexShape =
770+ kBothTranspose ? kTransposeInputShape : (kIndexBroadcast ? kBroadcastSourceShape : kOutputShape);
725constexpr int32_t kInputElementCount = IL_HAS_INPUT_ELEMENT;771constexpr int32_t kInputElementCount = IL_HAS_INPUT_ELEMENT;
726constexpr int32_t kIndexElementCount = IL_HAS_INDEX_ELEMENT;772constexpr int32_t kIndexElementCount = IL_HAS_INDEX_ELEMENT;
727constexpr bool kHasOutputRelu = IL_HAS_OUTPUT_RELU;773constexpr bool kHasOutputRelu = IL_HAS_OUTPUT_RELU;
@@ -751,12 +797,38 @@ int32_t DenseOffset(const std::array<int32_t, 4> &coordinate, const std::array<i
751}797}
752 798 
753void InitializeData(DataType *x, IndexType *index, DataType *expected) {799void InitializeData(DataType *x, IndexType *index, DataType *expected) {
754- for (int32_t i = 0; i < ElementCount(kInputShape); ++i) {800+ if constexpr (kBothTranspose) {
755- x[i] = static_cast<DataType>(static_cast<float>((i % 29) - 14) * 0.25F);801+ for (int32_t a = 0; a < kInputShape[0]; ++a) {
756- }802+ for (int32_t b = 0; b < kInputShape[1]; ++b) {
757- for (int32_t i = 0; i < ElementCount(kIndexShape); ++i) {803+ for (int32_t c = 0; c < kInputShape[2]; ++c) {
758- const int32_t gathered_axis = (i * 3 + 1) % kInputShape[2];804+ for (int32_t d = 0; d < kInputShape[3]; ++d) {
759- index[i] = static_cast<IndexType>(kIndexElementCount == 0 || i % 2 == 0 ? gathered_axis : -gathered_axis);805+ const std::array<int32_t, 4> coordinate = {a, b, c, d};
806+ const int32_t value =
807+ kOutputShape[0] == kOutputShape[1] ? (a + b) * 7 + c * 3 + d : a * 11 + b * 3 + c * 5 + d;
808+ x[DenseOffset(coordinate, kInputShape)] = static_cast<DataType>(static_cast<float>(value) * 0.125F);
809+ }
810+ }
811+ }
812+ }
813+ for (int32_t a = 0; a < kIndexShape[0]; ++a) {
814+ for (int32_t b = 0; b < kIndexShape[1]; ++b) {
815+ for (int32_t c = 0; c < kIndexShape[2]; ++c) {
816+ for (int32_t d = 0; d < kIndexShape[3]; ++d) {
817+ const std::array<int32_t, 4> coordinate = {a, b, c, d};
818+ const int32_t value = kOutputShape[0] == kOutputShape[1] ? a + b + c + d : a * 5 + b * 2 + c + d;
819+ index[DenseOffset(coordinate, kIndexShape)] = static_cast<IndexType>(value % kInputShape[2]);
820+ }
821+ }
822+ }
823+ }
824+ } else {
825+ for (int32_t i = 0; i < ElementCount(kInputShape); ++i) {
826+ x[i] = static_cast<DataType>(static_cast<float>((i % 29) - 14) * 0.25F);
827+ }
828+ for (int32_t i = 0; i < ElementCount(kIndexShape); ++i) {
829+ const int32_t gathered_axis = (i * 3 + 1) % kInputShape[2];
830+ index[i] = static_cast<IndexType>(kIndexElementCount == 0 || i % 2 == 0 ? gathered_axis : -gathered_axis);
831+ }
760 }832 }
761 for (int32_t i = 0; i < ElementCount(kOutputShape); ++i) {833 for (int32_t i = 0; i < ElementCount(kOutputShape); ++i) {
762 int32_t coordinate = i;834 int32_t coordinate = i;
@@ -767,7 +839,8 @@ void InitializeData(DataType *x, IndexType *index, DataType *expected) {
767 const int32_t b = coordinate % kOutputShape[1];839 const int32_t b = coordinate % kOutputShape[1];
768 coordinate /= kOutputShape[1];840 coordinate /= kOutputShape[1];
769 const int32_t a = coordinate;841 const int32_t a = coordinate;
770- std::array<int32_t, 4> index_coordinate = {a, b, c, d};842+ std::array<int32_t, 4> index_coordinate =
843+ kBothTranspose ? std::array<int32_t, 4>{b, a, c, d} : std::array<int32_t, 4>{a, b, c, d};
771 if (kIndexBroadcast) {844 if (kIndexBroadcast) {
772 for (size_t dim = 0UL; dim < index_coordinate.size(); ++dim) {845 for (size_t dim = 0UL; dim < index_coordinate.size(); ++dim) {
773 if ((kBroadcastAxesMask & (1U << dim)) != 0U) {846 if ((kBroadcastAxesMask & (1U << dim)) != 0U) {
@@ -788,7 +861,8 @@ void InitializeData(DataType *x, IndexType *index, DataType *expected) {
788 gathered_index = std::abs(gathered_index);861 gathered_index = std::abs(gathered_index);
789 }862 }
790 const int32_t gathered_axis = static_cast<int32_t>(gathered_index);863 const int32_t gathered_axis = static_cast<int32_t>(gathered_index);
791- std::array<int32_t, 4> input_coordinate = {a, b, gathered_axis, d};864+ std::array<int32_t, 4> input_coordinate = kBothTranspose ? std::array<int32_t, 4>{b, a, gathered_axis, d}
865+ : std::array<int32_t, 4>{a, b, gathered_axis, d};
792 if (kInputBroadcast && !kComplexSimt) {866 if (kInputBroadcast && !kComplexSimt) {
793 for (size_t dim = 0UL; dim < input_coordinate.size(); ++dim) {867 for (size_t dim = 0UL; dim < input_coordinate.size(); ++dim) {
794 if ((kBroadcastAxesMask & (1U << dim)) != 0U) {868 if ((kBroadcastAxesMask & (1U << dim)) != 0U) {
@@ -796,6 +870,9 @@ void InitializeData(DataType *x, IndexType *index, DataType *expected) {
796 }870 }
797 }871 }
798 }872 }
873+ if constexpr (kInputTranspose) {
874+ std::swap(input_coordinate[0], input_coordinate[2]);
875+ }
799 const int32_t input_offset = DenseOffset(input_coordinate, kInputShape);876 const int32_t input_offset = DenseOffset(input_coordinate, kInputShape);
800 float value = static_cast<float>(x[input_offset]);877 float value = static_cast<float>(x[input_offset]);
801 if constexpr (kInputAbsBeforeBroadcast) {878 if constexpr (kInputAbsBeforeBroadcast) {
@@ -863,9 +940,25 @@ TEST(E2EIndirectLoadBroadcast, GeneratedKernelMatchesReference) {
863 ASSERT_TRUE(tiling.IsValid());940 ASSERT_TRUE(tiling.IsValid());
864 941 
865 AscendC::SetKernelMode(KernelMode::AIV_MODE);942 AscendC::SetKernelMode(KernelMode::AIV_MODE);
943+#if IL_OUTPUT_BROADCAST_ADD
944+ auto bias = indirect_load_test::AllocGmBuffer<DataType>(kOutputShape[3]);
945+ ASSERT_NE(bias, nullptr);
946+ for (int32_t d = 0; d < kOutputShape[3]; ++d) {
947+ bias.get()[d] = static_cast<DataType>(static_cast<float>((d * 7) % 19 - 9) * 0.125F);
948+ }
949+ for (int32_t i = 0; i < output_count; ++i) {
950+ buffers.expected[i] = static_cast<DataType>(static_cast<float>(buffers.expected[i]) +
951+ static_cast<float>(bias.get()[i % kOutputShape[3]]));
952+ }
953+ ICPU_RUN_KF(indirect_load_broadcast_test, tiling.data.block_dim, reinterpret_cast<uint8_t *>(buffers.input.get()),
954+ reinterpret_cast<uint8_t *>(buffers.index.get()), reinterpret_cast<uint8_t *>(bias.get()),
955+ reinterpret_cast<uint8_t *>(buffers.output.get()), tiling.workspace.get(),
956+ reinterpret_cast<uint8_t *>(&tiling.data));
957+#else
866 ICPU_RUN_KF(indirect_load_broadcast_test, tiling.data.block_dim, reinterpret_cast<uint8_t *>(buffers.input.get()),958 ICPU_RUN_KF(indirect_load_broadcast_test, tiling.data.block_dim, reinterpret_cast<uint8_t *>(buffers.input.get()),
867 reinterpret_cast<uint8_t *>(buffers.index.get()), reinterpret_cast<uint8_t *>(buffers.output.get()),959 reinterpret_cast<uint8_t *>(buffers.index.get()), reinterpret_cast<uint8_t *>(buffers.output.get()),
868 tiling.workspace.get(), reinterpret_cast<uint8_t *>(&tiling.data));960 tiling.workspace.get(), reinterpret_cast<uint8_t *>(&tiling.data));
961+#endif
869 for (int32_t i = 0; i < output_count; ++i) {962 for (int32_t i = 0; i < output_count; ++i) {
870 EXPECT_NEAR(static_cast<float>(buffers.output.get()[i]),963 EXPECT_NEAR(static_cast<float>(buffers.output.get()[i]),
871 static_cast<float>(buffers.expected[static_cast<size_t>(i)]), 0.0625F)964 static_cast<float>(buffers.expected[static_cast<size_t>(i)]), 0.0625F)
@@ -941,6 +1034,51 @@ TEST(UserGraphConstruction, GeneratedKernelMatchesReference) {
941 }1034 }
942}1035}
943 1036 
1037+#elif defined(IL_USER_FANOUT_SIDE_INPUT)
1038+TEST(UserGraphConstruction, GeneratedKernelMatchesReference) {
1039+ constexpr int32_t kRows = 2;
1040+ constexpr int32_t kDim = 16;
1041+ constexpr int32_t kTableRows = 32;
1042+ auto indices = indirect_load_test::AllocGmBuffer<int64_t>(kRows);
1043+ auto embedding = indirect_load_test::AllocGmBuffer<bfloat16_t>(kTableRows * kDim);
1044+ auto side_input = indirect_load_test::AllocGmBuffer<bfloat16_t>(kDim);
1045+ auto output0 = indirect_load_test::AllocGmBuffer<bfloat16_t>(kRows * kDim);
1046+ auto output1 = indirect_load_test::AllocGmBuffer<bfloat16_t>(kRows * kDim);
1047+ ASSERT_TRUE(indices && embedding && side_input && output0 && output1);
1048+ for (int32_t row = 0; row < kRows; ++row) {
1049+ indices.get()[row] = row + 1;
1050+ }
1051+ for (int32_t row = 0; row < kTableRows; ++row) {
1052+ for (int32_t col = 0; col < kDim; ++col) {
1053+ embedding.get()[row * kDim + col] = static_cast<bfloat16_t>((row + 1) * 0.01F + col * 0.001F);
1054+ }
1055+ }
1056+ for (int32_t col = 0; col < kDim; ++col) {
1057+ side_input.get()[col] = static_cast<bfloat16_t>((col - 3) * 0.02F);
1058+ }
1059+ std::fill_n(output0.get(), kRows * kDim, static_cast<bfloat16_t>(0.0F));
1060+ std::fill_n(output1.get(), kRows * kDim, static_cast<bfloat16_t>(0.0F));
1061+ 
1062+ indirect_load_test::KernelTiling tiling;
1063+ ASSERT_TRUE(tiling.IsValid());
1064+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
1065+ ICPU_RUN_KF(user_fanout_side_input, tiling.block_dim, reinterpret_cast<uint8_t *>(indices.get()),
1066+ reinterpret_cast<uint8_t *>(embedding.get()), reinterpret_cast<uint8_t *>(side_input.get()),
1067+ reinterpret_cast<uint8_t *>(output0.get()), reinterpret_cast<uint8_t *>(output1.get()),
1068+ reinterpret_cast<uint8_t *>(tiling.workspace.get()), reinterpret_cast<uint8_t *>(&tiling.data));
1069+ 
1070+ for (int32_t row = 0; row < kRows; ++row) {
1071+ for (int32_t col = 0; col < kDim; ++col) {
1072+ const float value = static_cast<float>(embedding.get()[indices.get()[row] * kDim + col]);
1073+ const float bias = static_cast<float>(side_input.get()[col]);
1074+ EXPECT_NEAR(static_cast<float>(output0.get()[row * kDim + col]), value, 0.125F)
1075+ << "output0 row=" << row << ", col=" << col;
1076+ EXPECT_NEAR(static_cast<float>(output1.get()[row * kDim + col]), value + bias, 0.125F)
1077+ << "output1 row=" << row << ", col=" << col;
1078+ }
1079+ }
1080+}
1081+ 
944#elif defined(IL_USER_SIDE_INPUT_FANOUT)1082#elif defined(IL_USER_SIDE_INPUT_FANOUT)
945TEST(UserGraphConstruction, GeneratedKernelMatchesReference) {1083TEST(UserGraphConstruction, GeneratedKernelMatchesReference) {
946 constexpr int32_t kElements = 64;1084 constexpr int32_t kElements = 64;
@@ -1290,6 +1428,9 @@ extern "C" int64_t AutofuseTiling(AutofuseTilingData *, uint32_t *, uint32_t *,
1290#if defined(IL_USER_MASKED_EMBEDDING_MINIMAL) || defined(IL_USER_MASKED_EMBEDDING_SUM_FULL)1428#if defined(IL_USER_MASKED_EMBEDDING_MINIMAL) || defined(IL_USER_MASKED_EMBEDDING_SUM_FULL)
1291extern "C" __global__ __aicore__ void user_masked_embedding_sum(GM_ADDR input0, GM_ADDR input1, GM_ADDR input2,1429extern "C" __global__ __aicore__ void user_masked_embedding_sum(GM_ADDR input0, GM_ADDR input1, GM_ADDR input2,
1292 GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling);1430 GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling);
1431+#elif defined(IL_GRAPH_HINT_EMBEDDING_SLICE)
1432+extern "C" __global__ __aicore__ void autofused_to_copy_embedding_slice_161da55df6362966862381c5d053b7e7(
1433+ GM_ADDR input0, GM_ADDR input1, GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling);
1293#elif defined(IL_GRAPH_HINT_SIMD_REPRO)1434#elif defined(IL_GRAPH_HINT_SIMD_REPRO)
1294extern "C" __global__ __aicore__ void indirect_load_graph_hint_simd_repro(GM_ADDR input0, GM_ADDR input1,1435extern "C" __global__ __aicore__ void indirect_load_graph_hint_simd_repro(GM_ADDR input0, GM_ADDR input1,
1295 GM_ADDR output, GM_ADDR workspace,1436 GM_ADDR output, GM_ADDR workspace,
@@ -1390,6 +1531,60 @@ TEST(E2EUserMaskedEmbeddingSum, GeneratedKernelMatchesReference) {
1390 }1531 }
1391 }1532 }
1392}1533}
1534+#elif defined(IL_GRAPH_HINT_EMBEDDING_SLICE)
1535+constexpr int32_t kGraphHintEmbeddingSliceRows = 128;
1536+constexpr int32_t kGraphHintEmbeddingSliceColumns = 512;
1537+constexpr int32_t kGraphHintEmbeddingSliceTableRows = 65536;
1538+constexpr int32_t kGraphHintEmbeddingSliceIndexStride = 39;
1539+constexpr int32_t kGraphHintEmbeddingSliceIndexOffset = 13;
1540+ 
1541+TEST(E2EIndirectLoadGraphHintEmbeddingSlice, GeneratedKernelMatchesReference) {
1542+ constexpr int64_t table_count =
1543+ static_cast<int64_t>(kGraphHintEmbeddingSliceTableRows) * kGraphHintEmbeddingSliceColumns;
1544+ constexpr int64_t index_count =
1545+ static_cast<int64_t>(kGraphHintEmbeddingSliceIndexOffset) +
1546+ static_cast<int64_t>(kGraphHintEmbeddingSliceRows - 1) * kGraphHintEmbeddingSliceIndexStride + 1;
1547+ constexpr int64_t output_count = static_cast<int64_t>(kGraphHintEmbeddingSliceRows) * kGraphHintEmbeddingSliceColumns;
1548+ auto table = indirect_load_test::AllocGmBuffer<float>(table_count);
1549+ auto index = indirect_load_test::AllocGmBuffer<float>(index_count);
1550+ auto output = indirect_load_test::AllocGmBuffer<float>(output_count);
1551+ ASSERT_TRUE(table && index && output);
1552+ 
1553+ for (int32_t row = 0; row < kGraphHintEmbeddingSliceTableRows; ++row) {
1554+ for (int32_t column = 0; column < kGraphHintEmbeddingSliceColumns; ++column) {
1555+ table.get()[static_cast<int64_t>(row) * kGraphHintEmbeddingSliceColumns + column] =
1556+ static_cast<float>((row % 97) * 0.25F + (column % 31) * 0.001F);
1557+ }
1558+ }
1559+ std::fill_n(index.get(), index_count, 0.0F);
1560+ for (int32_t row = 0; row < kGraphHintEmbeddingSliceRows; ++row) {
1561+ const int32_t positive_index = (row * 997 + 17) % kGraphHintEmbeddingSliceTableRows;
1562+ const int32_t index_value = positive_index;
1563+ index.get()[static_cast<int64_t>(kGraphHintEmbeddingSliceIndexOffset) +
1564+ static_cast<int64_t>(row) * kGraphHintEmbeddingSliceIndexStride] = static_cast<float>(index_value);
1565+ }
1566+ std::fill_n(output.get(), output_count, 0.0F);
1567+ 
1568+ indirect_load_test::KernelTiling tiling;
1569+ ASSERT_TRUE(tiling.IsValid());
1570+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
1571+ ICPU_RUN_KF(autofused_to_copy_embedding_slice_161da55df6362966862381c5d053b7e7, tiling.block_dim,
1572+ reinterpret_cast<uint8_t *>(table.get()), reinterpret_cast<uint8_t *>(index.get()),
1573+ reinterpret_cast<uint8_t *>(output.get()), reinterpret_cast<uint8_t *>(tiling.workspace.get()),
1574+ reinterpret_cast<uint8_t *>(&tiling.data));
1575+ 
1576+ for (int32_t row = 0; row < kGraphHintEmbeddingSliceRows; ++row) {
1577+ const int32_t raw_index =
1578+ static_cast<int32_t>(index.get()[static_cast<int64_t>(kGraphHintEmbeddingSliceIndexOffset) +
1579+ static_cast<int64_t>(row) * kGraphHintEmbeddingSliceIndexStride]);
1580+ const int32_t selected_index = raw_index;
1581+ for (int32_t column = 0; column < kGraphHintEmbeddingSliceColumns; ++column) {
1582+ const int64_t offset = static_cast<int64_t>(row) * kGraphHintEmbeddingSliceColumns + column;
1583+ const int64_t table_offset = static_cast<int64_t>(selected_index) * kGraphHintEmbeddingSliceColumns + column;
1584+ EXPECT_FLOAT_EQ(output.get()[offset], table.get()[table_offset]) << "row=" << row << ", column=" << column;
1585+ }
1586+ }
1587+}
1393#elif defined(IL_GRAPH_HINT_SIMD_REPRO)1588#elif defined(IL_GRAPH_HINT_SIMD_REPRO)
1394constexpr int32_t kGraphHintSimdRows = 30;1589constexpr int32_t kGraphHintSimdRows = 30;
1395constexpr int32_t kGraphHintSimdIndexColumns = 3;1590constexpr int32_t kGraphHintSimdIndexColumns = 3;
Mautofuse/v35/ascendc/api_regbase/indirect_load_simd.h+151-3
@@ -12,12 +12,22 @@
12#define AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMD_H_12#define AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMD_H_
13 13 
14#include <type_traits>14#include <type_traits>
15+#include <utility>
15 16 
16#ifndef AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMD_POLICY_H_17#ifndef AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMD_POLICY_H_
17#include "indirect_load_simd_policy.h"18#include "indirect_load_simd_policy.h"
18#endif19#endif
19 20 
20namespace AscendC {21namespace AscendC {
22+enum class IndirectLoadSimdFallback : uint8_t {
23+ kRegisterGather = 0,
24+ kGatherApi = 1,
25+ kStrided = 2,
26+};
27+ 
28+template <IndirectLoadSimdFallback Fallback, bool TryEmbedding>
29+struct IndirectLoadSimdCaseTag {};
30+ 
21template <int32_t Rank>31template <int32_t Rank>
22struct IndirectLoadSimdStridedParams {32struct IndirectLoadSimdStridedParams {
23 uint32_t logical_size;33 uint32_t logical_size;
@@ -29,6 +39,29 @@ struct IndirectLoadSimdStridedParams {
29 int64_t output_strides[Rank];39 int64_t output_strides[Rank];
30};40};
31 41 
42+template <int32_t Rank>
43+struct IndirectLoadSimdDenseParams {
44+ uint32_t physical_size;
45+ uint32_t input_actual_size;
46+ int64_t output_offset;
47+ int64_t input_axis;
48+ int64_t index_sizes[Rank];
49+ int64_t input_strides[Rank];
50+};
51+ 
52+template <int32_t Rank>
53+struct IndirectLoadSimdEmbeddingParams {
54+ uint32_t logical_size;
55+ uint32_t physical_size;
56+ uint32_t input_actual_size;
57+ int64_t output_offset;
58+ int64_t input_axis;
59+ int64_t index_sizes[Rank];
60+ int64_t input_strides[Rank];
61+ int64_t index_strides[Rank];
62+ int64_t output_strides[Rank];
63+};
64+ 
32namespace Internal {65namespace Internal {
33template <int32_t Dim, int32_t Axis, int32_t StrideBase>66template <int32_t Dim, int32_t Axis, int32_t StrideBase>
34struct IndirectLoadSimdInnerOffset {67struct IndirectLoadSimdInnerOffset {
@@ -396,14 +429,19 @@ __aicore__ inline bool TryIndirectLoadSimdEmbedding(const LocalTensor<X> &x, con
396 int64_t expected_input_stride = 1;429 int64_t expected_input_stride = 1;
397 for (int32_t dim = Rank - 1; dim > Axis; --dim) {430 for (int32_t dim = Rank - 1; dim > Axis; --dim) {
398 if (shape[dim] <= 0 || shape[index_stride_base + dim] != 0 ||431 if (shape[dim] <= 0 || shape[index_stride_base + dim] != 0 ||
399- shape[input_stride_base + dim] != expected_input_stride) {432+ shape[input_stride_base + dim] != expected_input_stride || shape[output_stride_base + dim] != payload_span) {
400 return false;433 return false;
401 }434 }
402 payload_span *= shape[dim];435 payload_span *= shape[dim];
403 expected_input_stride *= shape[dim];436 expected_input_stride *= shape[dim];
404 }437 }
438+ const int64_t output_axis_stride = shape[output_stride_base + Axis];
405 if (payload_span <= 0 || shape[input_stride_base + Axis] != expected_input_stride || output_offset < 0 ||439 if (payload_span <= 0 || shape[input_stride_base + Axis] != expected_input_stride || output_offset < 0 ||
406- shape[output_stride_base + Axis] < payload_span) {440+ output_axis_stride <= 0 || output_axis_stride < payload_span) {
441+ return false;
442+ }
443+ const int64_t index_axis_stride = shape[index_stride_base + Axis];
444+ if (index_axis_stride <= 0) {
407 return false;445 return false;
408 }446 }
409 const int64_t output_slice_span = shape[Axis] * payload_span;447 const int64_t output_slice_span = shape[Axis] * payload_span;
@@ -448,7 +486,7 @@ __aicore__ inline bool TryIndirectLoadSimdEmbedding(const LocalTensor<X> &x, con
448 input_offset += coordinate * shape[input_stride_base + dim];486 input_offset += coordinate * shape[input_stride_base + dim];
449 }487 }
450 }488 }
451- const int64_t index_value = static_cast<int64_t>(index.GetValue(index_axis_coordinate));489+ const int64_t index_value = static_cast<int64_t>(index.GetValue(index_axis_coordinate * index_axis_stride));
452 const int64_t source_offset = index_value * shape[input_stride_base + Axis] + input_offset + payload_offset;490 const int64_t source_offset = index_value * shape[input_stride_base + Axis] + input_offset + payload_offset;
453 const int64_t payload_remaining = payload_span - payload_offset;491 const int64_t payload_remaining = payload_span - payload_offset;
454 const uint32_t copy_count = static_cast<uint32_t>(492 const uint32_t copy_count = static_cast<uint32_t>(
@@ -656,6 +694,116 @@ __aicore__ inline void IndirectLoadSimdGatherApi(const LocalTensor<X> &x, const
656 Gather(y, x, offsets, 0U, actual_size);694 Gather(y, x, offsets, 0U, actual_size);
657}695}
658 696 
697+namespace Internal {
698+template <typename CaseTag, int32_t Rank>
699+struct IndirectLoadSimdCaseSelector;
700+ 
701+template <IndirectLoadSimdFallback Fallback, bool TryEmbedding, int32_t Rank>
702+struct IndirectLoadSimdCaseSelector<IndirectLoadSimdCaseTag<Fallback, TryEmbedding>, Rank> {
703+ using Params =
704+ std::conditional_t<TryEmbedding, IndirectLoadSimdEmbeddingParams<Rank>,
705+ std::conditional_t<Fallback == IndirectLoadSimdFallback::kStrided,
706+ IndirectLoadSimdStridedParams<Rank>, IndirectLoadSimdDenseParams<Rank>>>;
707+ 
708+ template <typename X, typename Index, int32_t Axis>
709+ __aicore__ inline static void RunDense(const LocalTensor<X> &x, const LocalTensor<Index> &index,
710+ const LocalTensor<X> &y, const Params &params) {
711+ static_assert(Fallback != IndirectLoadSimdFallback::kStrided,
712+ "IndirectLoad SIMD dense facade requires a dense fallback.");
713+ if constexpr (TryEmbedding) {
714+ if (TryEmbeddingImpl<X, Index, Axis>(x, index, y, params, std::make_index_sequence<Rank>{})) {
715+ return;
716+ }
717+ }
718+ RunDenseFallback<X, Index, Axis>(x, index, y, params, std::make_index_sequence<Rank>{});
719+ }
720+ 
721+ template <typename X, typename Index, int32_t Axis>
722+ __aicore__ inline static void RunStrided(const LocalTensor<X> &x, const LocalTensor<Index> &index,
723+ const LocalTensor<X> &y, const LocalTensor<uint8_t> &tmp,
724+ const Params &params) {
725+ static_assert(Fallback == IndirectLoadSimdFallback::kStrided,
726+ "IndirectLoad SIMD strided facade requires a strided fallback.");
727+ if constexpr (TryEmbedding) {
728+ if (TryEmbeddingImpl<X, Index, Axis>(x, index, y, params, std::make_index_sequence<Rank>{})) {
729+ return;
730+ }
731+ }
732+ if constexpr (TryEmbedding) {
733+ IndirectLoadSimdStridedParams<Rank> strided_params{
734+ params.logical_size, params.physical_size, params.output_offset, {}, {}, {}, {}};
735+ for (int32_t dim = 0; dim < Rank; ++dim) {
736+ strided_params.index_sizes[dim] = params.index_sizes[dim];
737+ strided_params.input_strides[dim] = params.input_strides[dim];
738+ strided_params.index_strides[dim] = params.index_strides[dim];
739+ strided_params.output_strides[dim] = params.output_strides[dim];
740+ }
741+ AscendC::IndirectLoadSimdStrided<X, Index, Rank, Axis>(x, index, y, tmp, strided_params);
742+ } else {
743+ AscendC::IndirectLoadSimdStrided<X, Index, Rank, Axis>(x, index, y, tmp, params);
744+ }
745+ }
746+ 
747+ private:
748+ template <typename X, typename Index, int32_t Axis, size_t... Dims>
749+ __aicore__ inline static bool TryEmbeddingImpl(const LocalTensor<X> &x, const LocalTensor<Index> &index,
750+ const LocalTensor<X> &y, const Params &params,
751+ std::index_sequence<Dims...>) {
752+ return TryIndirectLoadSimdEmbedding<X, Index, Rank, Axis>(
753+ x, index, y, params.logical_size, params.output_offset, params.index_sizes[Dims]...,
754+ params.input_strides[Dims]..., params.index_strides[Dims]..., params.output_strides[Dims]...);
755+ }
756+ 
757+ template <typename X, typename Index, int32_t Axis, size_t... Dims>
758+ __aicore__ inline static void RunDenseFallback(const LocalTensor<X> &x, const LocalTensor<Index> &index,
759+ const LocalTensor<X> &y, const Params &params,
760+ std::index_sequence<Dims...>) {
761+ if constexpr (Fallback == IndirectLoadSimdFallback::kRegisterGather) {
762+ AscendC::IndirectLoadSimd<X, Index, Rank, Axis>(x, index, y, params.physical_size, params.output_offset,
763+ params.input_actual_size, params.input_axis,
764+ params.index_sizes[Dims]..., params.input_strides[Dims]...);
765+ } else {
766+ AscendC::IndirectLoadSimdGatherApi<X, Index, Rank, Axis>(
767+ x, index, y, params.physical_size, params.output_offset, params.input_actual_size, params.input_axis,
768+ params.index_sizes[Dims]..., params.input_strides[Dims]...);
769+ }
770+ }
771+};
772+} // namespace Internal
773+ 
774+template <typename CaseTag, int32_t Rank>
775+using IndirectLoadSimdParams = typename Internal::IndirectLoadSimdCaseSelector<CaseTag, Rank>::Params;
776+ 
777+template <typename X, typename Index, int32_t Rank, int32_t Axis, bool TryEmbedding>
778+__aicore__ inline void IndirectLoadSimd(
779+ const LocalTensor<X> &x, const LocalTensor<Index> &index, const LocalTensor<X> &y,
780+ IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kRegisterGather, TryEmbedding>,
781+ const IndirectLoadSimdParams<IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kRegisterGather, TryEmbedding>, Rank>
782+ &params) {
783+ using CaseTag = IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kRegisterGather, TryEmbedding>;
784+ Internal::IndirectLoadSimdCaseSelector<CaseTag, Rank>::template RunDense<X, Index, Axis>(x, index, y, params);
785+}
786+ 
787+template <typename X, typename Index, int32_t Rank, int32_t Axis, bool TryEmbedding>
788+__aicore__ inline void IndirectLoadSimdGatherApi(
789+ const LocalTensor<X> &x, const LocalTensor<Index> &index, const LocalTensor<X> &y,
790+ IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kGatherApi, TryEmbedding>,
791+ const IndirectLoadSimdParams<IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kGatherApi, TryEmbedding>, Rank>
792+ &params) {
793+ using CaseTag = IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kGatherApi, TryEmbedding>;
794+ Internal::IndirectLoadSimdCaseSelector<CaseTag, Rank>::template RunDense<X, Index, Axis>(x, index, y, params);
795+}
796+ 
797+template <typename X, typename Index, int32_t Rank, int32_t Axis, bool TryEmbedding>
798+__aicore__ inline void IndirectLoadSimdStrided(
799+ const LocalTensor<X> &x, const LocalTensor<Index> &index, const LocalTensor<X> &y, const LocalTensor<uint8_t> &tmp,
800+ IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kStrided, TryEmbedding>,
801+ const IndirectLoadSimdParams<IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kStrided, TryEmbedding>, Rank>
802+ &params) {
803+ using CaseTag = IndirectLoadSimdCaseTag<IndirectLoadSimdFallback::kStrided, TryEmbedding>;
804+ Internal::IndirectLoadSimdCaseSelector<CaseTag, Rank>::template RunStrided<X, Index, Axis>(x, index, y, tmp, params);
805+}
806+ 
659} // namespace AscendC807} // namespace AscendC
660 808 
661#endif // AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMD_H_809#endif // AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMD_H_
Mautofuse/v35/ascendc/api_regbase/indirect_load_simt.h+239-138
@@ -11,7 +11,61 @@
11#ifndef AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMT_H_11#ifndef AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMT_H_
12#define AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMT_H_12#define AUTOFUSE_V35_ASCENDC_API_REGBASE_INDIRECT_LOAD_SIMT_H_
13 13 
14+#include <type_traits>
15+ 
14namespace AscendC {16namespace AscendC {
17+enum class IndirectLoadSimtCase : uint8_t {
18+ kStaticPowerOfTwo = 0,
19+ kStaticInner = 1,
20+ kStructuredMagic = 2,
21+ kEmbedding = 3,
22+ kRecursive = 4,
23+ kStrided = 5,
24+};
25+ 
26+template <IndirectLoadSimtCase Case, typename OffsetT, int32_t Rank, int32_t Axis, uint64_t InnerSpan = 0U,
27+ uint64_t OutputAxisSpan = 0U, uint64_t InputAxisStride = 0U, uint64_t InputAxisSpan = 0U,
28+ uint64_t InputStrideMask = 0U, uint64_t IndexStrideMask = 0U>
29+struct IndirectLoadSimtCaseTag {};
30+ 
31+template <int32_t Rank>
32+struct IndirectLoadSimtShapeLayout {
33+ static constexpr int32_t kInputStrideBase = Rank;
34+ static constexpr int32_t kIndexStrideBase = kInputStrideBase + Rank;
35+ static constexpr int32_t kRecursiveShapeSize = kIndexStrideBase;
36+ static constexpr int32_t kStridedShapeSize = kIndexStrideBase + Rank;
37+};
38+ 
39+struct IndirectLoadSimtStaticPowerOfTwoParams {};
40+ 
41+template <typename OffsetT>
42+struct IndirectLoadSimtStaticInnerParams {
43+ OffsetT output_axis_span;
44+};
45+ 
46+template <typename OffsetT>
47+struct IndirectLoadSimtStructuredMagicParams {
48+ OffsetT inner_span;
49+ OffsetT output_axis_span;
50+ OffsetT input_axis_stride;
51+ OffsetT input_axis_span;
52+};
53+ 
54+template <typename OffsetT, int32_t Rank>
55+struct IndirectLoadSimtEmbeddingParams {
56+ OffsetT shape[IndirectLoadSimtShapeLayout<Rank>::kStridedShapeSize];
57+};
58+ 
59+template <typename OffsetT, int32_t Rank>
60+struct IndirectLoadSimtRecursiveParams {
61+ OffsetT shape[IndirectLoadSimtShapeLayout<Rank>::kRecursiveShapeSize];
62+};
63+ 
64+template <typename OffsetT, int32_t Rank>
65+struct IndirectLoadSimtStridedParams {
66+ OffsetT shape[IndirectLoadSimtShapeLayout<Rank>::kStridedShapeSize];
67+};
68+ 
15namespace Internal {69namespace Internal {
16__aicore__ inline uint64_t IndirectLoadGetUintDivMagic(uint64_t dividend, uint64_t divisor) {70__aicore__ inline uint64_t IndirectLoadGetUintDivMagic(uint64_t dividend, uint64_t divisor) {
17 uint64_t quotient = 0U;71 uint64_t quotient = 0U;
Mautofuse/v35/codegen/reg_api_call/reg_indirect_load_api_call.h+11-50
@@ -14,20 +14,12 @@
14#include <cstdint>14#include <cstdint>
15#include <string>15#include <string>
16#include <unordered_map>16#include <unordered_map>
17-#include <unordered_set>
18#include <vector>17#include <vector>
19 18 
20#include "codegen_kernel.h"19#include "codegen_kernel.h"
21#include "indirect_load_utils.h"20#include "indirect_load_utils.h"
22 21 
23namespace codegen {22namespace codegen {
24-struct SimtGmTensor {
25- ascir::TensorId value_tensor_id;
26- ascir::TensorId gm_tensor_id;
27- af::DataType dtype;
28- bool is_scalar{false};
29-};
30- 
31struct SimtOutputChain {23struct SimtOutputChain {
32 std::vector<af::AscNodePtr> nodes;24 std::vector<af::AscNodePtr> nodes;
33 ascir::TensorId result_tensor_id{af::kIdNone};25 ascir::TensorId result_tensor_id{af::kIdNone};
@@ -47,39 +39,9 @@ struct LogicalTensorInfo {
47 std::vector<ascir::SizeExpr> strides;39 std::vector<ascir::SizeExpr> strides;
48};40};
49 41 
50-enum class SimtAddressPolicy {42+using SimtLoadMetadataMap = std::unordered_map<std::string, ascgen_utils::indirect_load::SimtLoadMetadata>;
51- kStaticPowerOfTwo,43+using SimtLoadIndexOffsetExpressions = std::unordered_map<std::string, std::string>;
52- kStaticInner,44+using SimtGraphNodeMap = std::unordered_map<std::string, af::AscNodePtr>;
53- kStructuredMagic,
54- kEmbedding,
55- kRecursive,
56- kStrided,
57-};
58- 
59-struct SimtCodegenPlan {
60- SimtAddressPolicy policy = SimtAddressPolicy::kRecursive;
61- std::string offset_type = "uint64_t";
62- af::Expression inner_span = af::sym::kSymbolOne;
63- af::Expression output_axis_span = af::sym::kSymbolOne;
64- af::Expression input_axis_stride = af::sym::kSymbolOne;
65- af::Expression input_axis_span = af::sym::kSymbolOne;
66- uint64_t inner_span_value = 0U;
67- uint64_t output_axis_span_value = 0U;
68- uint64_t input_axis_stride_value = 0U;
69- uint64_t input_axis_span_value = 0U;
70- uint64_t input_stride_mask = 0U;
71- uint64_t index_stride_mask = 0U;
72- bool use_per_load_index_offsets = false;
73-};
74- 
75-enum class SimtLoadAddressSource {
76- kZeroOffset,
77- kOutputOffset,
78- kIndexOffset,
79-};
80- 
81-using SimtLoadAddressSources = std::unordered_map<const af::AscNode *, SimtLoadAddressSource>;
82-using SimtLoadIndexOffsetExpressions = std::unordered_map<const af::AscNode *, std::string>;
83 45 
84class IndirectLoadRegApiCall final : public ApiCall {46class IndirectLoadRegApiCall final : public ApiCall {
85 public:47 public:
@@ -96,7 +58,9 @@ class IndirectLoadRegApiCall final : public ApiCall {
96 Status ParseAttr(const ascir::NodeView &node) override;58 Status ParseAttr(const ascir::NodeView &node) override;
97 59 
98 private:60 private:
99- Status ParseSimtAttr(const ascir::NodeView &node);61+ Status ResolveSimtMetadata(const ascir::NodeView &node,
62+ const ascgen_utils::indirect_load::SimtLoweringMetadata &metadata,
63+ const SimtGraphNodeMap &node_map);
100 Status GenerateSk(const TPipe &tpipe, const std::vector<ascir::AxisId> &current_axis,64 Status GenerateSk(const TPipe &tpipe, const std::vector<ascir::AxisId> &current_axis,
101 const std::vector<std::reference_wrapper<const Tensor>> &inputs,65 const std::vector<std::reference_wrapper<const Tensor>> &inputs,
102 const std::vector<std::reference_wrapper<const Tensor>> &outputs, std::string &result) const;66 const std::vector<std::reference_wrapper<const Tensor>> &outputs, std::string &result) const;
@@ -113,22 +77,19 @@ class IndirectLoadRegApiCall final : public ApiCall {
113 ascir::TemplateId template_id_{ascir::TemplateId::kDefault};77 ascir::TemplateId template_id_{ascir::TemplateId::kDefault};
114 ascgen_utils::indirect_load::TemplateLogicalView logical_view_;78 ascgen_utils::indirect_load::TemplateLogicalView logical_view_;
115 ascgen_utils::indirect_load::IndirectLoadAccessInfo access_info_;79 ascgen_utils::indirect_load::IndirectLoadAccessInfo access_info_;
80+ ascgen_utils::indirect_load::SimdLoweringMetadata simd_metadata_;
116 bool has_post_reduce_{false};81 bool has_post_reduce_{false};
117 std::vector<af::AscNodePtr> index_nodes_;82 std::vector<af::AscNodePtr> index_nodes_;
118 std::vector<af::AscNodePtr> output_nodes_;83 std::vector<af::AscNodePtr> output_nodes_;
119 std::vector<SimtOutputChain> output_chains_;84 std::vector<SimtOutputChain> output_chains_;
120- // SIMT layout metadata and policy are built once during ParseSimtAttr and reused by both emitters.85+ ascgen_utils::indirect_load::SimtPolicyMetadata simt_policy_;
121- LogicalTensorInfo simt_input_info_;86+ SimtLoadMetadataMap simt_index_loads_;
122- LogicalTensorInfo simt_index_info_;87+ SimtLoadMetadataMap simt_output_loads_;
123- SimtCodegenPlan simt_plan_;88+ std::vector<ascgen_utils::indirect_load::SimtGmTensorMetadata> simt_gm_tensors_;
124- SimtLoadAddressSources simt_load_address_sources_;
125- std::vector<SimtGmTensor> simt_gm_tensors_;
126- std::unordered_set<ascir::TensorId> simt_gm_tensor_ids_;
127 ascir::TensorId index_result_tensor_id_{af::kIdNone};89 ascir::TensorId index_result_tensor_id_{af::kIdNone};
128 ascir::TensorId simt_value_tensor_id_{af::kIdNone};90 ascir::TensorId simt_value_tensor_id_{af::kIdNone};
129 ascir::TensorId output_result_tensor_id_{af::kIdNone};91 ascir::TensorId output_result_tensor_id_{af::kIdNone};
130 af::DataType index_dtype_{af::DT_UNDEFINED};92 af::DataType index_dtype_{af::DT_UNDEFINED};
131- af::DataType output_dtype_{af::DT_UNDEFINED};
132 ascgen_utils::indirect_load::Implementation implementation_{ascgen_utils::indirect_load::Implementation::kDefault};93 ascgen_utils::indirect_load::Implementation implementation_{ascgen_utils::indirect_load::Implementation::kDefault};
133};94};
134} // namespace codegen95} // namespace codegen
Mautofuse/v35/optimize/template/nddma_template.cpp+29-1
@@ -28,6 +28,25 @@ constexpr char kTransposePrefix[] = "transpose_";
28bool IsNddma(const af::AscNodePtr &node) {28bool IsNddma(const af::AscNodePtr &node) {
29 return ScheduleUtils::IsLoad(node) && node->attr.type == "Nddma";29 return ScheduleUtils::IsLoad(node) && node->attr.type == "Nddma";
30}30}
31+ 
32+Status UpdateSimdIndirectLoadDenseInputLayouts(const af::OutDataAnchorPtr &producer_anchor) {
33+ for (const auto &consumer_anchor : producer_anchor->GetPeerInDataAnchors()) {
34+ const auto consumer = std::dynamic_pointer_cast<af::AscNode>(consumer_anchor->GetOwnerNode());
35+ if (!af::ops::IsOps<af::ascir_op::IndirectLoad>(consumer) ||
36+ ascir::GetTemplateIdOrDefault(*consumer) != ascir::TemplateId::kIndirectLoadSimd) {
37+ continue;
38+ }
39+ ascgen_utils::indirect_load::TemplateLogicalView logical_view;
40+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::GetTemplateLogicalView(consumer, logical_view));
41+ const auto input_index = static_cast<size_t>(consumer_anchor->GetIdx());
42+ auto &layout =
43+ input_index == ascgen_utils::indirect_load::kInputTensorIndex ? logical_view.input : logical_view.index;
44+ GE_ASSERT_SUCCESS(ScheduleUtils::RecalculateStridesFromRepeats(layout.sizes, layout.strides));
45+ layout.kind = ascgen_utils::indirect_load::IndirectLoadLayoutKind::kDense;
46+ GE_ASSERT_SUCCESS(ascgen_utils::indirect_load::SetTemplateLogicalView(consumer, logical_view));
47+ }
48+ return af::SUCCESS;
49+}
31} // namespace50} // namespace
32 51 
33std::string NddmaTemplate::GenName(const std::string &general_case_name) {52std::string NddmaTemplate::GenName(const std::string &general_case_name) {
@@ -143,6 +162,7 @@ Status NddmaTemplate::MergeLoadAndTranspose(const af::AscNodePtr &load_node, af:
143 load_node->attr.type = "Nddma";162 load_node->attr.type = "Nddma";
144 // 删除transpose节点163 // 删除transpose节点
145 GE_ASSERT_SUCCESS(ScheduleUtils::RemoveNode(new_case, out_node, load_node->GetOutDataAnchor(0)));164 GE_ASSERT_SUCCESS(ScheduleUtils::RemoveNode(new_case, out_node, load_node->GetOutDataAnchor(0)));
165+ GE_ASSERT_SUCCESS(UpdateSimdIndirectLoadDenseInputLayouts(load_out_anchor));
146 return af::SUCCESS;166 return af::SUCCESS;
147}167}
148 168 
@@ -241,6 +261,14 @@ af::Status NddmaTemplate::ProcessTransposeNodes(af::AscGraph &new_case, bool &is
241 261 
242af::Status NddmaTemplate::Generate([[maybe_unused]] const af::AscGraph &origin_graph,262af::Status NddmaTemplate::Generate([[maybe_unused]] const af::AscGraph &origin_graph,
243 [[maybe_unused]] const af::AscGraph &based_case, af::AscGraph &new_case) {263 [[maybe_unused]] const af::AscGraph &based_case, af::AscGraph &new_case) {
264+ const auto indirect_load = ascgen_utils::indirect_load::FindIndirectLoadNode(new_case);
265+ const bool is_simd_indirect_load =
266+ indirect_load != nullptr && ascir::GetTemplateIdOrDefault(*indirect_load) == ascir::TemplateId::kIndirectLoadSimd;
267+ if (indirect_load != nullptr && !is_simd_indirect_load &&
268+ ScheduleUtils::HasComputeType(origin_graph, af::ComputeType::kComputeTranspose)) {
269+ GELOGD("Skip NDDMA transpose template for IndirectLoad graph[%s].", new_case.GetName().c_str());
270+ return af::FAILED;
271+ }
244 bool is_nddma_generated = false;272 bool is_nddma_generated = false;
245 GE_ASSERT_SUCCESS(ProcessTransposeNodes(new_case, is_nddma_generated));273 GE_ASSERT_SUCCESS(ProcessTransposeNodes(new_case, is_nddma_generated));
246 bool is_transpose_nddma_generated = is_nddma_generated;274 bool is_transpose_nddma_generated = is_nddma_generated;
@@ -283,7 +311,7 @@ af::Status NddmaTemplate::Generate([[maybe_unused]] const af::AscGraph &origin_g
283 GELOGD("No nddma template generated.");311 GELOGD("No nddma template generated.");
284 return af::FAILED;312 return af::FAILED;
285 }313 }
286- if (is_transpose_nddma_generated) {314+ if (is_transpose_nddma_generated && indirect_load == nullptr) {
287 GE_ASSERT_SUCCESS(315 GE_ASSERT_SUCCESS(
288 UnAlignmentStrategy::ModifyTransposeFusionVectorizedStrides(new_case, BaseAlignmentStrategy::GetAlignWidth()));316 UnAlignmentStrategy::ModifyTransposeFusionVectorizedStrides(new_case, BaseAlignmentStrategy::GetAlignWidth()));
289 }317 }
Mautofuse/v35/optimize/un_alignment_strategy.cpp+64-0
@@ -13,10 +13,36 @@
13#include <stack>13#include <stack>
14#include "ascir_utils.h"14#include "ascir_utils.h"
15#include "common_utils.h"15#include "common_utils.h"
16+#include "indirect_load_utils.h"
16#include "tensor_layout_utils.h"17#include "tensor_layout_utils.h"
17#include "platform/v1/alignment_strategy.h"18#include "platform/v1/alignment_strategy.h"
18 19 
19namespace optimize {20namespace optimize {
21+namespace {
22+af::Status NormalizeStoreSingletonStrides(const af::AscNodePtr &node) {
23+ GE_ASSERT_TRUE(node->outputs().size() == 1UL);
24+ auto &attr = node->outputs[0].attr;
25+ GE_ASSERT_TRUE(attr.vectorized_axis.size() == attr.vectorized_strides.size());
26+ GE_ASSERT_TRUE(attr.axis.size() == attr.repeats.size());
27+ af::Expression inner_span = af::sym::kSymbolOne;
28+ for (size_t index = attr.vectorized_axis.size(); index > 0UL; --index) {
29+ const size_t vector_index = index - 1UL;
30+ const auto axis_iter = std::find(attr.axis.begin(), attr.axis.end(), attr.vectorized_axis[vector_index]);
31+ GE_ASSERT_TRUE(axis_iter != attr.axis.end());
32+ const size_t axis_index = static_cast<size_t>(std::distance(attr.axis.begin(), axis_iter));
33+ auto &stride = attr.vectorized_strides[vector_index];
34+ if (af::SymbolicUtils::StaticCheckEq(attr.repeats[axis_index], af::sym::kSymbolOne) == af::TriBool::kTrue &&
35+ af::SymbolicUtils::StaticCheckEq(stride, af::sym::kSymbolZero) == af::TriBool::kTrue) {
36+ stride = inner_span;
37+ }
38+ if (af::SymbolicUtils::StaticCheckEq(stride, af::sym::kSymbolZero) != af::TriBool::kTrue) {
39+ inner_span = af::sym::Mul(stride, attr.repeats[axis_index]);
40+ }
41+ }
42+ return af::SUCCESS;
43+}
44+} // namespace
45+ 
20AlignmentType UnAlignmentStrategy::GetDefaultAlignmentType() {46AlignmentType UnAlignmentStrategy::GetDefaultAlignmentType() {
21 return AlignmentType::kNotAligned;47 return AlignmentType::kNotAligned;
22}48}
@@ -440,7 +466,45 @@ Status UnAlignmentStrategy::ModifyTransposeFusionVectorizedStrides(af::AscGraph
440 return af::SUCCESS;466 return af::SUCCESS;
441}467}
442 468 
469+Status UnAlignmentStrategy::ModifyIndirectLoadVectorizedStrides(ascir::ImplGraph &impl_graph) {
470+ bool has_transpose = false;
471+ bool has_nddma_load = false;
472+ for (const auto &node : impl_graph.GetAllNodes()) {
473+ has_transpose = has_transpose || node->attr.api.compute_type == af::ComputeType::kComputeTranspose;
474+ has_nddma_load = has_nddma_load || (ScheduleUtils::IsLoad(node) && node->attr.type == "Nddma");
475+ const auto role = ascgen_utils::indirect_load::GetTemplateRole(node);
476+ if (role != ascgen_utils::indirect_load::TemplateRole::kStridedUbPath &&
477+ role != ascgen_utils::indirect_load::TemplateRole::kSimdInputPreStridedUbPath) {
478+ continue;
479+ }
480+ for (const auto &output : node->outputs()) {
481+ if (!output->attr.vectorized_axis.empty()) {
482+ GE_ASSERT_SUCCESS(
483+ BaseAlignmentStrategy::SetVectorizedStridesForTensor(node, output->attr, AlignmentType::kAligned));
484+ }
485+ }
486+ }
487+ // Only the direct IL->Store boundary inherits the transposed/NDDMA layout.
488+ // Other IL stores (for example, stores after elementwise or reduce nodes) use
489+ // their own layout and must not be normalized here.
490+ if (has_transpose || has_nddma_load) {
491+ for (const auto &node : impl_graph.GetAllNodes()) {
492+ if (!af::ops::IsOps<af::ascir_op::Store>(node) || node->inputs.Size() != 1UL || node->outputs().size() != 1UL) {
493+ continue;
494+ }
495+ const auto producer = ascgen_utils::indirect_load::GetInputProducer(node, 0UL);
496+ if (af::ops::IsOps<af::ascir_op::IndirectLoad>(producer)) {
497+ GE_ASSERT_SUCCESS(NormalizeStoreSingletonStrides(node));
498+ }
499+ }
500+ }
501+ return af::SUCCESS;
502+}
503+ 
443Status UnAlignmentStrategy::ModifyVectorizedStrides(af::AscGraph &impl_graph) {504Status UnAlignmentStrategy::ModifyVectorizedStrides(af::AscGraph &impl_graph) {
505+ if (ascgen_utils::indirect_load::FindIndirectLoadNode(impl_graph) != nullptr) {
506+ return ModifyIndirectLoadVectorizedStrides(impl_graph);
507+ }
444 bool has_transpose = false;508 bool has_transpose = false;
445 for (auto node : impl_graph.GetAllNodes()) {509 for (auto node : impl_graph.GetAllNodes()) {
446 if (af::ops::IsOps<af::ascir_op::Transpose>(node)) {510 if (af::ops::IsOps<af::ascir_op::Transpose>(node)) {
Mautofuse/v35/optimize/un_alignment_strategy.h+1-0
@@ -37,6 +37,7 @@ class UnAlignmentStrategy : public BaseAlignmentStrategy {
37 static af::Status CollectTransposePreNodes(const af::AscGraph &graph, std::set<af::AscNodePtr> &transpose_pre_nodes);37 static af::Status CollectTransposePreNodes(const af::AscGraph &graph, std::set<af::AscNodePtr> &transpose_pre_nodes);
38 static af::Status UpdateOutputVectorizedStrides(const af::AscNodePtr &node, uint32_t continuous_tail_axis_num,38 static af::Status UpdateOutputVectorizedStrides(const af::AscNodePtr &node, uint32_t continuous_tail_axis_num,
39 uint32_t align_width);39 uint32_t align_width);
40+ static af::Status ModifyIndirectLoadVectorizedStrides(ascir::ImplGraph &impl_graph);
40};41};
41 42 
42Status GenLoadToGenNddmaNode(const af::AscNodePtr &node_load);43Status GenLoadToGenNddmaNode(const af::AscNodePtr &node_load);
Mscripts/test/run_autofuse_test.sh+9-2
@@ -733,7 +733,6 @@ build_backend() {
733 cmake $CMAKE_ARGS ../733 cmake $CMAKE_ARGS ../
734 734 
735 # st用例可执行文件的列表，inductor split_compile 仅保留 presubmit 代表用例，其余放 nightly。735 # st用例可执行文件的列表，inductor split_compile 仅保留 presubmit 代表用例，其余放 nightly。
736- # Temporarily exclude embedding_test and graph_hint_simd_repro because the SIMD embedding fast path has an index-stride issue.
737 MAKE_TARGET_LIST="add_abs_test_e2e \736 MAKE_TARGET_LIST="add_abs_test_e2e \
738 axpy_abs_test_e2e \737 axpy_abs_test_e2e \
739 sub_abs_test_e2e \738 sub_abs_test_e2e \
@@ -845,6 +844,7 @@ build_backend() {
845 indirect_load_broadcast_inner_adjacent_simd_test_e2e_v2 \844 indirect_load_broadcast_inner_adjacent_simd_test_e2e_v2 \
846 indirect_load_broadcast_continuous_simd_test_e2e_v2 \845 indirect_load_broadcast_continuous_simd_test_e2e_v2 \
847 indirect_load_broadcast_continuous_index_simt_test_e2e_v2 \846 indirect_load_broadcast_continuous_index_simt_test_e2e_v2 \
847+ indirect_load_embedding_test_e2e_v2 \
848 indirect_load_embedding_tail_simd_e2e_v2 \848 indirect_load_embedding_tail_simd_e2e_v2 \
849 indirect_load_embedding_tail_simt_e2e_v2 \849 indirect_load_embedding_tail_simt_e2e_v2 \
850 indirect_load_embedding_aligned_simd_e2e_v2 \850 indirect_load_embedding_aligned_simd_e2e_v2 \
@@ -855,8 +855,12 @@ build_backend() {
855 indirect_load_complex_input_broadcast_simd_test_e2e_v2 \855 indirect_load_complex_input_broadcast_simd_test_e2e_v2 \
856 indirect_load_broadcast_cross_boundary_simt_test_e2e_v2 \856 indirect_load_broadcast_cross_boundary_simt_test_e2e_v2 \
857 indirect_load_broadcast_axis_simt_test_e2e_v2 \857 indirect_load_broadcast_axis_simt_test_e2e_v2 \
858- indirect_load_broadcast_reduce_simt_fallback_test_e2e_v2 \858+ indirect_load_broadcast_reduce_simt_fallback_test_e2e_v2 \
859 indirect_load_broadcast_retained_simt_test_e2e_v2 \859 indirect_load_broadcast_retained_simt_test_e2e_v2 \
860+ indirect_load_both_transpose_simd_test_e2e_v2 \
861+ indirect_load_both_transpose_simd_fallback_test_e2e_v2 \
862+ indirect_load_both_transpose_simd_nddma_test_e2e_v2 \
863+ indirect_load_both_transpose_simt_test_e2e_v2 \
860 indirect_load_broadcast_elements_sk_test_e2e_v2 \864 indirect_load_broadcast_elements_sk_test_e2e_v2 \
861 indirect_load_broadcast_index_physical_view_simt_test_e2e_v2 \865 indirect_load_broadcast_index_physical_view_simt_test_e2e_v2 \
862 indirect_load_broadcast_index_mixed_view_simt_test_e2e_v2 \866 indirect_load_broadcast_index_mixed_view_simt_test_e2e_v2 \
@@ -872,6 +876,8 @@ build_backend() {
872 indirect_load_rank3_axis1_input_index_outer_gap_simt_test_e2e_v2 \876 indirect_load_rank3_axis1_input_index_outer_gap_simt_test_e2e_v2 \
873 indirect_load_rank3_axis1_torch_gather_frontend_e2e_v2 \877 indirect_load_rank3_axis1_torch_gather_frontend_e2e_v2 \
874 indirect_load_graph_hint_reduce_simt_test_e2e_v2 \878 indirect_load_graph_hint_reduce_simt_test_e2e_v2 \
879+ indirect_load_graph_hint_simd_repro_e2e_v2 \
880+ indirect_load_graph_hint_embedding_slice_e2e_v2 \
875 indirect_load_embedding_reduce_simt_test_e2e_v2 \881 indirect_load_embedding_reduce_simt_test_e2e_v2 \
876 indirect_load_add_il_reduce_test_e2e_v2 \882 indirect_load_add_il_reduce_test_e2e_v2 \
877 indirect_load_user_masked_embedding_minimal_e2e_v2 \883 indirect_load_user_masked_embedding_minimal_e2e_v2 \
@@ -893,6 +899,7 @@ build_backend() {
893 indirect_load_user_fanout_post_stores_simt_e2e_v2 \899 indirect_load_user_fanout_post_stores_simt_e2e_v2 \
894 indirect_load_user_fanout_post_reduce_simd_e2e_v2 \900 indirect_load_user_fanout_post_reduce_simd_e2e_v2 \
895 indirect_load_user_fanout_post_reduce_simt_e2e_v2 \901 indirect_load_user_fanout_post_reduce_simt_e2e_v2 \
902+ indirect_load_user_fanout_side_input_simt_e2e_v2 \
896 indirect_load_user_side_input_fanout_simd_e2e_v2 \903 indirect_load_user_side_input_fanout_simd_e2e_v2 \
897 indirect_load_user_side_input_fanout_simt_e2e_v2 \904 indirect_load_user_side_input_fanout_simt_e2e_v2 \
898 load_where_x2_x3_is_ubscalar_store_test_e2e_v2 \905 load_where_x2_x3_is_ubscalar_store_test_e2e_v2 \