adstraw commented on code in PR #13721:
URL: https://github.com/apache/tvm/pull/13721#discussion_r1081717491
##########
include/tvm/meta_schedule/schedule_rule.h:
##########
@@ -210,6 +210,31 @@ class ScheduleRule : public runtime::ObjectRef {
Optional<Array<Integer>> vector_load_lens, Optional<Map<String,
ObjectRef>> reuse_read,
Optional<Map<String, ObjectRef>> reuse_write, bool
use_software_pipeline);
+ /*!
+ * \brief Extension of MultiLevelTiling for auto-tensorization with multiple
groups of candidate
+ * tensor core intrinsics
+ * \param intrin_groups A list of groups of tensor core intrinsics. The map
should contain key
+ * "compute" which represents the tensor intrin for computation. The value
of the map should be
+ * names of tensor intrinsics, must be registered via
+ * TensorIntrin.register(...) beforehand
+ * \param structure The tiling structure. Recommended:
+ * - 'SRSRS' on Hexagon
Review Comment:
You might map `SRSRS` to the layout `NCHWc` in the comment
##########
src/meta_schedule/schedule_rule/multi_level_tiling_hexagon.cc:
##########
@@ -0,0 +1,145 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include "../../tir/schedule/analysis.h"
+#include "../../tir/schedule/transform.h"
+#include "../utils.h"
+#include "multi_level_tiling_with_intrin.h"
+
+namespace tvm {
+namespace meta_schedule {
+
+using tir::BlockRV;
+using tir::LoopRV;
+using tir::Schedule;
+
+class MultiLevelTilingHexagonNode : public MultiLevelTilingWithIntrinNode {
+ private:
+ // Subrule: Add software pipeline
+ inline std::vector<State> AddSoftwarePipeline(State state) const;
+
+ // Override ApplySubRules to apply tensorization-specific sub-rules
+ std::vector<State> ApplySubRules(std::vector<State> states) final;
+
+ // Inherited from ScheduleRuleNode
+ ScheduleRule Clone() const override {
+ ObjectPtr<MultiLevelTilingHexagonNode> n =
make_object<MultiLevelTilingHexagonNode>(*this);
+ return ScheduleRule(n);
+ }
+
+ public:
+ /*! \brief Whether to use software pipeline */
+ bool use_software_pipeline = false;
+ static constexpr const char* _type_key =
"meta_schedule.MultiLevelTilingHexagon";
+ TVM_DECLARE_FINAL_OBJECT_INFO(MultiLevelTilingHexagonNode,
MultiLevelTilingNode);
+};
+
+std::vector<State>
MultiLevelTilingHexagonNode::ApplySubRules(std::vector<State> states) {
+ states = MultiLevelTilingWithIntrinNode::ApplySubRules(states);
+ states = SubRule(std::move(states), [&](State state) { return
AddSoftwarePipeline(state); });
+ return states;
+}
+
+std::vector<State> MultiLevelTilingHexagonNode::AddSoftwarePipeline(State
state) const {
+ if (!use_software_pipeline) {
+ return {state};
+ }
+ // The current config is not suitable for software pipelining.
+ if (r_indices_.size() < 2) {
+ return {state};
+ }
+
+ Schedule& sch = state->sch;
+ // Check reduction length after blockize.
+ int64_t reduction_length = 1;
+ for (int r_index : r_indices_) {
+ const Array<LoopRV>& tiles = state->tiles[r_index];
+ for (const LoopRV& tile : tiles) {
+ const auto* extent = sch->Get(tile)->extent.as<IntImmNode>();
+ ICHECK(extent != nullptr) << "Dynamic extent is not supported.";
+ reduction_length *= extent->value;
+ }
+ }
+ if (reduction_length <= 1) {
Review Comment:
Curious use of `<= 1` here as opposed to `== 1`. Do we support zero or
negative extents?
##########
tests/python/contrib/test_hexagon/test_conv2d_async.py:
##########
@@ -0,0 +1,184 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+# pylint: disable=missing-docstring
+""" Test rpc based launcher for hexagon """
+import tempfile
+
+import numpy as np
+import pytest
+import tvm.testing
+import tvm.topi.testing
+from tvm import meta_schedule as ms
+from tvm import relay
+from tvm.contrib.hexagon.meta_schedule import (
+ get_hexagon_local_builder,
+ get_hexagon_rpc_runner,
+)
+from tvm.meta_schedule import postproc, schedule_rule
+from tvm.tir.tensor_intrin.hexagon import (
+ VRMPY_u8u8i32_VTCM_READS_INTRIN,
+)
+
+from .infrastructure import get_hexagon_target
+
+
+def tune_vrmpy_auto_tensorize(mod, params, hexagon_launcher):
+ sch_rules_async = [
+ schedule_rule.ApplyCustomRule(),
+ schedule_rule.AutoInline(
+ into_producer=False,
+ into_consumer=True,
+ inline_const_tensor=True,
+ disallow_if_then_else=True,
+ require_injective=True,
+ require_ordered=True,
+ disallow_op=["tir.exp"],
+ ),
+ schedule_rule.MultiLevelTilingHexagon(
+ intrin_groups=[
+ {"compute": VRMPY_u8u8i32_VTCM_READS_INTRIN},
+ ],
+ structure="SRSRS",
+ tile_binds=None,
+ max_innermost_factor=64, # 64 // tensor intrin size
+ vector_load_lens=None,
+ reuse_read=ms.schedule_rule.ReuseType(
+ req="must",
+ levels=[2],
+ scope="global.vtcm",
+ ),
+ reuse_write=None,
+ use_software_pipeline=True,
+ ),
+ schedule_rule.ParallelizeVectorizeUnroll(
+ max_jobs_per_core=-1,
+ max_vectorize_extent=-1,
+ unroll_max_steps=[8, 16, 32],
+ unroll_explicit=True,
+ ),
+ ]
+
+ postprocs = [
+ postproc.RewriteParallelVectorizeUnroll(),
+ postproc.RewriteReductionBlock(),
+ postproc.RewriteTensorize(vectorize_init_loop=True),
+ postproc.VerifyVTCMLimit(),
+
postproc.DisallowAsyncStridedMemCopy(merge_async_commit_queue_scope=False),
+ ]
+
+ target = get_hexagon_target("v68")
+ executor = relay.backend.Executor("graph", {"link-params": True})
+ mod = mod.with_attr("executor", executor)
+
+ use_async = True
Review Comment:
This seems strange especially with no `else` case below
##########
python/tvm/tir/tensor_intrin/hexagon.py:
##########
@@ -68,12 +69,12 @@ def sync_dma_load_impl(a: T.handle, c: T.handle) -> None:
return sync_dma_load_desc, sync_dma_load_impl
-def generate_dot_product_32x4_u8u8i32(mem_scope="global"):
+def generate_dot_product_32x4_u8u8i32(read_mem_scope="global",
write_mem_scope="global"):
Review Comment:
Why are the defaults here "global" instead of "global.vtcm"?
##########
python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py:
##########
@@ -197,6 +197,60 @@ def __init__(
)
+@register_object("meta_schedule.MultiLevelTilingHexagon")
+class MultiLevelTilingHexagon(ScheduleRule):
+ """Extension of MultiLevelTiling for auto-tensorizing with multiple groups
of candidate hexagon
+ intrinsics.
+
+ Parameters
+ ----------
+ intrin_groups : List[Mapping[str, str]]
+ A list of groups of tensor core intrinsics. The map should contain key
+ "compute" which represents the tensor intrin for computation. The
value of the map should be
+ names of tensor intrinsics, must be registered via
+ TensorIntrin.register(...) beforehand
+ structure : str
+ The tiling structure. Recommended:
+ - 'SRSRS' on Hexagon
Review Comment:
You might map `SRSRS` to the layout `NCHWc` in the comment
##########
include/tvm/meta_schedule/schedule_rule.h:
##########
@@ -210,6 +210,31 @@ class ScheduleRule : public runtime::ObjectRef {
Optional<Array<Integer>> vector_load_lens, Optional<Map<String,
ObjectRef>> reuse_read,
Optional<Map<String, ObjectRef>> reuse_write, bool
use_software_pipeline);
+ /*!
+ * \brief Extension of MultiLevelTiling for auto-tensorization with multiple
groups of candidate
+ * tensor core intrinsics
+ * \param intrin_groups A list of groups of tensor core intrinsics. The map
should contain key
+ * "compute" which represents the tensor intrin for computation. The value
of the map should be
+ * names of tensor intrinsics, must be registered via
+ * TensorIntrin.register(...) beforehand
+ * \param structure The tiling structure. Recommended:
+ * - 'SRSRS' on Hexagon
+ * \param tile_binds For each level of tiles, which thread axis it is bound
to. These are not
+ * supported on hexagon.
+ * \param max_innermost_factor The maximum size of the innermost factor.
NullOpt means no limit
+ * \param vector_load_lens The length of vector lane in vectorized
cooperative fetching.
Review Comment:
Confused as to why we have both the vector load length and the max innermost
factor. These seem redundant. Aren't we always going to vectorize over the
innermost loop? And, if vectorization is enabled, won't we use
`vector_load_lens` for the size of the innermost loop? Also a "max" for
`max_innermost_factor` seems strange. E.g. if the user says the max is `8` I
imagine that trying `7` is a bad choice whereas a list of [2, 4, 8] are likely
good choices. Anyway... I see that this API is inherited from
`MultiLevelTilingInitCommon` so I won't harp on it. Just an observation.
##########
src/meta_schedule/schedule_rule/multi_level_tiling_hexagon.cc:
##########
@@ -0,0 +1,145 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include "../../tir/schedule/analysis.h"
+#include "../../tir/schedule/transform.h"
+#include "../utils.h"
+#include "multi_level_tiling_with_intrin.h"
+
+namespace tvm {
+namespace meta_schedule {
+
+using tir::BlockRV;
+using tir::LoopRV;
+using tir::Schedule;
+
+class MultiLevelTilingHexagonNode : public MultiLevelTilingWithIntrinNode {
+ private:
+ // Subrule: Add software pipeline
+ inline std::vector<State> AddSoftwarePipeline(State state) const;
+
+ // Override ApplySubRules to apply tensorization-specific sub-rules
+ std::vector<State> ApplySubRules(std::vector<State> states) final;
+
+ // Inherited from ScheduleRuleNode
+ ScheduleRule Clone() const override {
+ ObjectPtr<MultiLevelTilingHexagonNode> n =
make_object<MultiLevelTilingHexagonNode>(*this);
+ return ScheduleRule(n);
+ }
+
+ public:
+ /*! \brief Whether to use software pipeline */
+ bool use_software_pipeline = false;
+ static constexpr const char* _type_key =
"meta_schedule.MultiLevelTilingHexagon";
+ TVM_DECLARE_FINAL_OBJECT_INFO(MultiLevelTilingHexagonNode,
MultiLevelTilingNode);
+};
+
+std::vector<State>
MultiLevelTilingHexagonNode::ApplySubRules(std::vector<State> states) {
+ states = MultiLevelTilingWithIntrinNode::ApplySubRules(states);
+ states = SubRule(std::move(states), [&](State state) { return
AddSoftwarePipeline(state); });
+ return states;
+}
+
+std::vector<State> MultiLevelTilingHexagonNode::AddSoftwarePipeline(State
state) const {
+ if (!use_software_pipeline) {
+ return {state};
+ }
+ // The current config is not suitable for software pipelining.
+ if (r_indices_.size() < 2) {
+ return {state};
+ }
+
+ Schedule& sch = state->sch;
+ // Check reduction length after blockize.
+ int64_t reduction_length = 1;
+ for (int r_index : r_indices_) {
+ const Array<LoopRV>& tiles = state->tiles[r_index];
+ for (const LoopRV& tile : tiles) {
+ const auto* extent = sch->Get(tile)->extent.as<IntImmNode>();
+ ICHECK(extent != nullptr) << "Dynamic extent is not supported.";
+ reduction_length *= extent->value;
+ }
+ }
+ if (reduction_length <= 1) {
+ return {state};
+ }
+
+ // Return if there are more less than 1 or more than 2 cache_reads.
+ size_t cache_read_count = state->read_reuse.size();
+ if (cache_read_count > 2 || cache_read_count == 0) {
+ return {state};
+ }
+
+ // Add annotations for software pipelining at the loop right above the cache
read stages.
+ Array<Integer> software_pipeline_stage;
+ Array<Integer> software_pipeline_order;
+ Array<Integer> software_pipeline_async_stages;
+ if (cache_read_count == 2) {
Review Comment:
This looks correct but this notation is difficult to read for many folks.
Some comments might help.
##########
src/meta_schedule/schedule_rule/multi_level_tiling_hexagon.cc:
##########
@@ -0,0 +1,145 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include "../../tir/schedule/analysis.h"
+#include "../../tir/schedule/transform.h"
+#include "../utils.h"
+#include "multi_level_tiling_with_intrin.h"
+
+namespace tvm {
+namespace meta_schedule {
+
+using tir::BlockRV;
+using tir::LoopRV;
+using tir::Schedule;
+
+class MultiLevelTilingHexagonNode : public MultiLevelTilingWithIntrinNode {
+ private:
+ // Subrule: Add software pipeline
+ inline std::vector<State> AddSoftwarePipeline(State state) const;
+
+ // Override ApplySubRules to apply tensorization-specific sub-rules
+ std::vector<State> ApplySubRules(std::vector<State> states) final;
+
+ // Inherited from ScheduleRuleNode
+ ScheduleRule Clone() const override {
+ ObjectPtr<MultiLevelTilingHexagonNode> n =
make_object<MultiLevelTilingHexagonNode>(*this);
+ return ScheduleRule(n);
+ }
+
+ public:
+ /*! \brief Whether to use software pipeline */
+ bool use_software_pipeline = false;
+ static constexpr const char* _type_key =
"meta_schedule.MultiLevelTilingHexagon";
+ TVM_DECLARE_FINAL_OBJECT_INFO(MultiLevelTilingHexagonNode,
MultiLevelTilingNode);
+};
+
+std::vector<State>
MultiLevelTilingHexagonNode::ApplySubRules(std::vector<State> states) {
+ states = MultiLevelTilingWithIntrinNode::ApplySubRules(states);
+ states = SubRule(std::move(states), [&](State state) { return
AddSoftwarePipeline(state); });
Review Comment:
Seems that `MultiLevelTilingHexagon` could be its own schedule rule as adds
the ability to `AddSoftwarePipeline` but otherwise defers to
`MultiLevelTilingWithIntrin`. I understand the main reason for the inheritance
here is to control the order of the application of the schedule rules, that
tiling must precede software pipelining. Wondering there is some other
solution besides inheritance to solve this problem.
##########
src/meta_schedule/schedule_rule/multi_level_tiling_hexagon.cc:
##########
@@ -0,0 +1,145 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include "../../tir/schedule/analysis.h"
+#include "../../tir/schedule/transform.h"
+#include "../utils.h"
+#include "multi_level_tiling_with_intrin.h"
+
+namespace tvm {
+namespace meta_schedule {
+
+using tir::BlockRV;
+using tir::LoopRV;
+using tir::Schedule;
+
+class MultiLevelTilingHexagonNode : public MultiLevelTilingWithIntrinNode {
+ private:
+ // Subrule: Add software pipeline
+ inline std::vector<State> AddSoftwarePipeline(State state) const;
+
+ // Override ApplySubRules to apply tensorization-specific sub-rules
+ std::vector<State> ApplySubRules(std::vector<State> states) final;
+
+ // Inherited from ScheduleRuleNode
+ ScheduleRule Clone() const override {
+ ObjectPtr<MultiLevelTilingHexagonNode> n =
make_object<MultiLevelTilingHexagonNode>(*this);
+ return ScheduleRule(n);
+ }
+
+ public:
+ /*! \brief Whether to use software pipeline */
+ bool use_software_pipeline = false;
+ static constexpr const char* _type_key =
"meta_schedule.MultiLevelTilingHexagon";
+ TVM_DECLARE_FINAL_OBJECT_INFO(MultiLevelTilingHexagonNode,
MultiLevelTilingNode);
+};
+
+std::vector<State>
MultiLevelTilingHexagonNode::ApplySubRules(std::vector<State> states) {
+ states = MultiLevelTilingWithIntrinNode::ApplySubRules(states);
+ states = SubRule(std::move(states), [&](State state) { return
AddSoftwarePipeline(state); });
+ return states;
+}
+
+std::vector<State> MultiLevelTilingHexagonNode::AddSoftwarePipeline(State
state) const {
+ if (!use_software_pipeline) {
+ return {state};
+ }
+ // The current config is not suitable for software pipelining.
Review Comment:
Update comment to indicate what `r_indices_` represents (reduction axes)
since it's not a member of `MultiLevelTilingHexagonNode` class. And also why
we need at least 2 of them as I can't quite figure out what that's the case.
##########
src/meta_schedule/schedule_rule/multi_level_tiling_hexagon.cc:
##########
@@ -0,0 +1,145 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include "../../tir/schedule/analysis.h"
+#include "../../tir/schedule/transform.h"
+#include "../utils.h"
+#include "multi_level_tiling_with_intrin.h"
+
+namespace tvm {
+namespace meta_schedule {
+
+using tir::BlockRV;
+using tir::LoopRV;
+using tir::Schedule;
+
+class MultiLevelTilingHexagonNode : public MultiLevelTilingWithIntrinNode {
+ private:
+ // Subrule: Add software pipeline
+ inline std::vector<State> AddSoftwarePipeline(State state) const;
+
+ // Override ApplySubRules to apply tensorization-specific sub-rules
+ std::vector<State> ApplySubRules(std::vector<State> states) final;
+
+ // Inherited from ScheduleRuleNode
+ ScheduleRule Clone() const override {
+ ObjectPtr<MultiLevelTilingHexagonNode> n =
make_object<MultiLevelTilingHexagonNode>(*this);
+ return ScheduleRule(n);
+ }
+
+ public:
+ /*! \brief Whether to use software pipeline */
+ bool use_software_pipeline = false;
+ static constexpr const char* _type_key =
"meta_schedule.MultiLevelTilingHexagon";
+ TVM_DECLARE_FINAL_OBJECT_INFO(MultiLevelTilingHexagonNode,
MultiLevelTilingNode);
+};
+
+std::vector<State>
MultiLevelTilingHexagonNode::ApplySubRules(std::vector<State> states) {
+ states = MultiLevelTilingWithIntrinNode::ApplySubRules(states);
+ states = SubRule(std::move(states), [&](State state) { return
AddSoftwarePipeline(state); });
+ return states;
+}
+
+std::vector<State> MultiLevelTilingHexagonNode::AddSoftwarePipeline(State
state) const {
+ if (!use_software_pipeline) {
+ return {state};
+ }
+ // The current config is not suitable for software pipelining.
+ if (r_indices_.size() < 2) {
+ return {state};
+ }
+
+ Schedule& sch = state->sch;
+ // Check reduction length after blockize.
+ int64_t reduction_length = 1;
+ for (int r_index : r_indices_) {
+ const Array<LoopRV>& tiles = state->tiles[r_index];
+ for (const LoopRV& tile : tiles) {
+ const auto* extent = sch->Get(tile)->extent.as<IntImmNode>();
+ ICHECK(extent != nullptr) << "Dynamic extent is not supported.";
+ reduction_length *= extent->value;
+ }
+ }
+ if (reduction_length <= 1) {
+ return {state};
+ }
+
+ // Return if there are more less than 1 or more than 2 cache_reads.
+ size_t cache_read_count = state->read_reuse.size();
+ if (cache_read_count > 2 || cache_read_count == 0) {
+ return {state};
+ }
+
+ // Add annotations for software pipelining at the loop right above the cache
read stages.
+ Array<Integer> software_pipeline_stage;
+ Array<Integer> software_pipeline_order;
+ Array<Integer> software_pipeline_async_stages;
+ if (cache_read_count == 2) {
+ software_pipeline_stage = Array<Integer>{0, 0, 1};
+ software_pipeline_order = Array<Integer>{0, 1, 2};
+ software_pipeline_async_stages = Array<Integer>{0};
+ } else {
+ software_pipeline_stage = Array<Integer>{0, 1};
+ software_pipeline_order = Array<Integer>{0, 1};
+ software_pipeline_async_stages = Array<Integer>{0};
+ }
+
+ tir::BlockRV cache_read_block = state->read_reuse.begin()->second;
+ Array<LoopRV> cache_read_loops = sch->GetLoops(cache_read_block);
+ Array<LoopRV> reduction_loops;
+ for (size_t i = 0; i < cache_read_loops.size() - 1; ++i) {
+ if (tir::GetLoopIterType(sch->GetSRef(cache_read_loops[i])) !=
tir::IterVarType::kDataPar) {
+ reduction_loops.push_back(cache_read_loops[i]);
+ } else if (reduction_loops.size() > 0 &&
+ sch->Get(cache_read_loops[i])->extent.as<IntImmNode>()->value
== 1) {
+ reduction_loops.push_back(cache_read_loops[i]);
+ }
+ }
+ auto fused = sch->Fuse(reduction_loops);
+
+ sch->Annotate(fused, tir::attr::software_pipeline_stage,
software_pipeline_stage);
+ sch->Annotate(fused, tir::attr::software_pipeline_order,
software_pipeline_order);
+ sch->Annotate(fused, tir::attr::software_pipeline_async_stages,
software_pipeline_async_stages);
+
+ // TODO(nverke): Add support for nested async pipelines.
+ // TODO(nverke): Add support for async cache writes.
Review Comment:
Is the lack of cache write support here due to the issue in the
InjectSWPipeline pass where there is no "wait" on cache write stage?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
