You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "zxybazh (via GitHub)" <gi...@apache.org> on 2023/04/14 07:00:43 UTC
[GitHub] [tvm] zxybazh opened a new pull request, #14624: [Unity] Introduce FewShotTuning Pass
zxybazh opened a new pull request, #14624:
URL: https://github.com/apache/tvm/pull/14624
Based on feedback for `DefaultGPUSchedule` pass, we noticed that while PrimFuncs can be supported without tuning on GPU devices, the performance is not competitive. Therefore we introduce a new FewShotTuning pass for static shape PrimFuncs that directly samples from the search space and make sure the returned candidate can pass post-processors.
By design, this pass enables most of the performance optimizations according to the target, and avoided tuning by:
- Disable benchmarking by default.
- Only do random sampling.
- Limit valid trials to user specified small number, by default 1.
- Skip tuning apis and directly do sampling and compilation.
- Still guarantee result is compilable.
This PR is a follow-up for @psrivas2's work and based on discussion with @sunggg and @psrivas2.
CC: @jwfromm @junrushao
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] sunggg commented on a diff in pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "sunggg (via GitHub)" <gi...@apache.org>.
sunggg commented on code in PR #14624:
URL: https://github.com/apache/tvm/pull/14624#discussion_r1185618857
##########
tests/python/relax/test_transform_few_shot_tuning.py:
##########
@@ -0,0 +1,436 @@
+# 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=invalid-name,,missing-function-docstring
+from typing import List, Tuple
+import numpy as np
+
+import tvm
+from tvm.script import tir as T
+from tvm.tir.tensor_intrin.cuda import * # pylint: disable=wildcard-import,unused-wildcard-import
+from tvm.tir.tensor_intrin.x86 import * # pylint: disable=wildcard-import,unused-wildcard-import
+from tvm.meta_schedule.testing.tune_utils import generate_input_data
+from tvm.meta_schedule.arg_info import ArgInfo
+from tvm.relax.transform import FewShotTuning
+import tvm.testing
+
+# pylint: disable=no-self-argument,missing-class-docstring,line-too-long
+# fmt: off
+@tvm.script.ir_module
+class MatMul:
+ @T.prim_func
+ def matmul(
+ A: T.Buffer((32, 32), "float16"),
+ B: T.Buffer((32, 32), "float16"),
+ C: T.Buffer((32, 32), "float16"),
+ ):
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ # with T.block("root"):
+ for i, j, k in T.grid(32, 32, 32):
+ with T.block("C"):
+ v_i, v_j, v_k = T.axis.remap("SSR", [i, j, k])
+ T.reads(A[v_i, v_k], B[v_k, v_j])
+ T.writes(C[v_i, v_j])
+ with T.init():
+ C[v_i, v_j] = T.float16(0)
+ C[v_i, v_j] = C[v_i, v_j] + A[v_i, v_k] * B[v_k, v_j]
+# fmt: on
+# pylint: enable=no-self-argument,missing-class-docstring,line-too-long
+
+
+def _target():
+ return tvm.target.Target("llvm -num-cores=4")
+ # for local testing only
+ # return tvm.target.Target("nvidia/geforce-rtx-3070")
+
+
+def _get_input_output_info(func: tvm.tir.PrimFunc):
+ args = ArgInfo.from_prim_func(func)
+ inputs = [generate_input_data(x.shape, x.dtype) for x in args[:-1]]
+ output_shape = args[-1].shape
+ output_dtype = args[-1].dtype
+ return inputs, output_shape, output_dtype
+
+
+def _expected_results(
+ mod: tvm.ir.IRModule, inputs: List[np.ndarray], output_shape: Tuple, output_dtype: str
+):
+ rt_mod = tvm.build(mod, target="llvm")
+ data = [
+ tvm.nd.array(x)
+ for x in [
+ *inputs,
+ np.zeros(output_shape, dtype=output_dtype),
+ ]
+ ]
+ rt_mod(*data)
+ return data[-1].numpy()
+
+
+def _actual_results(
+ actual: tvm.ir.IRModule, inputs: List[np.ndarray], output_shape: Tuple, output_dtype: str
+):
+ target = _target()
+ actual_rt_mod = tvm.build(actual, target=target)
+ actual_data = [
+ tvm.nd.array(x, device=tvm.cuda() if target.kind.name == "cuda" else tvm.cpu())
+ for x in [
+ *inputs,
+ np.zeros(output_shape, dtype=output_dtype),
+ ]
+ ]
+ actual_rt_mod(*actual_data)
+ return actual_data[-1].numpy()
+
+
+def _assert_allclose(mod: tvm.ir.IRModule, actual: tvm.ir.IRModule, func: tvm.tir.PrimFunc):
+ inputs, output_shape, output_dtype = _get_input_output_info(func)
+ expected_output = _expected_results(mod, inputs, output_shape, output_dtype)
+ actual_output = _actual_results(actual, inputs, output_shape, output_dtype)
+ tvm.testing.assert_allclose(expected_output, actual_output, rtol=1e-2, atol=1e-2)
Review Comment:
We usually use `rtol=1e-7, atol=1e-7`. Shall we change to this?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] zxybazh commented on pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on PR #14624:
URL: https://github.com/apache/tvm/pull/14624#issuecomment-1533888568
@sunggg Now all the comments are addressed, and the tests are working. For testing, note that we cannot know which search space and what decision is picked becaused it’s randomly sampling from the search space. Therefore, I only made sure that the result is correct and schedule is applied (so that it can run on cuda). Please take another look.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] sunggg commented on a diff in pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "sunggg (via GitHub)" <gi...@apache.org>.
sunggg commented on code in PR #14624:
URL: https://github.com/apache/tvm/pull/14624#discussion_r1167102287
##########
tests/python/relax/test_transform_few_shot_tuning.py:
##########
@@ -0,0 +1,56 @@
+# 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=invalid-name,,missing-function-docstring
+import tvm
+from tvm.tir.tensor_intrin.cuda import *
+from tvm.tir.tensor_intrin.x86 import *
+from tvm.relax.transform import FewShotTuning
+from tvm.script import tir as T
+import tvm.testing
+
+
+def test_matmul():
+ # pylint: disable=no-self-argument,missing-class-docstring,line-too-long
+ # fmt: off
+ @tvm.script.ir_module
+ class Before:
+ @T.prim_func
+ def matmul(
+ A: T.Buffer((32, 32), "float16"),
+ B: T.Buffer((32, 32), "float16"),
+ C: T.Buffer((32, 32), "float16"),
+ ):
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ # with T.block("root"):
+ for i, j, k in T.grid(32, 32, 32):
+ with T.block("C"):
+ v_i, v_j, v_k = T.axis.remap("SSR", [i, j, k])
+ T.reads(A[v_i, v_k], B[v_k, v_j])
+ T.writes(C[v_i, v_j])
+ with T.init():
+ C[v_i, v_j] = T.float16(0)
+ C[v_i, v_j] = C[v_i, v_j] + A[v_i, v_k] * B[v_k, v_j]
+ # fmt: on
+ # pylint: enable=no-self-argument,missing-class-docstring,line-too-long
+ target = tvm.target.Target("nvidia/geforce-rtx-3070")
+ with target, tvm.transform.PassContext(opt_level=3):
+ After = FewShotTuning()(Before)
+ After.show()
Review Comment:
Let's use the structural equality checks e.g. https://github.com/tlc-pack/relax/blob/relax/tests/python/relax/test_transform_legalize_ops_binary.py#L195
##########
src/relax/transform/few_shot_tuning.cc:
##########
@@ -0,0 +1,169 @@
+/*
+ * 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 <tvm/relax/transform.h>
+
+#include "../../meta_schedule/utils.h"
+
+namespace tvm {
+namespace relax {
+namespace transform {
+
+tir::PrimFunc FewShotTunePrimFunc(const tir::PrimFunc& prim_func, const Target& target,
+ int64_t valid_count, Optional<meta_schedule::Runner> runner) {
+ // fetch a local builder
+ static const auto* f_get_local_builder =
+ runtime::Registry::Get("meta_schedule.builder.get_local_builder");
+ ICHECK(f_get_local_builder)
+ << "ValueError: Cannot find the packed function \"meta_schedule.builder.get_local_builder\"";
+ meta_schedule::Builder builder = (*f_get_local_builder)();
+ // create an IRModule
+ IRModule mod = IRModule(Map<GlobalVar, BaseFunc>({{GlobalVar("main"), prim_func}}));
+ // fetch the number of physical cores
+ static const auto* f_cpu_count = runtime::Registry::Get("meta_schedule.cpu_count");
+ ICHECK(f_cpu_count) << "ValueError: Cannot find the packed function \"meta_schedule._cpu_count\"";
+ int num_threads = (*f_cpu_count)(false);
+ // store the results
+ Array<IRModule> results;
+ std::vector<double> costs;
+ // create a TuneContext
+ meta_schedule::TuneContext task = meta_schedule::TuneContext(
+ /*mod=*/mod,
+ /*target=*/target,
+ /*space_generator=*/
+ meta_schedule::SpaceGenerator::PostOrderApply(/*f_block_filter=*/nullptr,
+ /*sch_rules=*/NullOpt,
+ /*postprocs=*/NullOpt,
+ /*mutator_probs=*/NullOpt),
+ /*search_strategy=*/meta_schedule::SearchStrategy::ReplayTrace(/*max_fail_count=*/100),
+ /*task_name=*/NullOpt,
+ /*num_threads=*/num_threads, // use all available local threads
+ /*rand_state=*/-1, // -1 means use random seed
+ /*logger=*/nullptr);
+ task->Initialize();
+ task->search_strategy.value()->PreTuning(
+ /*max_trials=*/valid_count, /*num_trials_per_iter=*/valid_count,
+ /*design_spaces=*/task->space_generator.value()->GenerateDesignSpace(mod),
+ /*database=*/NullOpt,
+ /*cost_model=*/NullOpt);
+ while (valid_count > 0) {
+ Optional<Array<meta_schedule::MeasureCandidate>> candidates =
+ task->search_strategy.value()->GenerateMeasureCandidates();
+ if (!candidates.defined()) break;
+ Array<meta_schedule::BuilderInput> builder_inputs;
+ for (const meta_schedule::MeasureCandidate& candidate : candidates.value()) {
+ builder_inputs.push_back(meta_schedule::BuilderInput(
+ /*mod=*/candidate->sch->mod(),
+ /*target=*/target));
+ }
+ Array<meta_schedule::BuilderResult> builder_results = builder->Build(builder_inputs);
+ ICHECK_EQ(builder_results.size(), candidates.value().size());
+ int idx = 0;
+ for (const meta_schedule::BuilderResult& builder_result : builder_results) {
+ if (!builder_result->error_msg.defined()) {
+ results.push_back(candidates.value()[idx]->sch->mod());
+ valid_count--;
+ }
+ idx++;
+ }
+ if (runner.defined()) {
+ Array<meta_schedule::RunnerInput> runner_inputs;
+ int idx = 0;
+ for (const meta_schedule::BuilderResult& builder_result : builder_results) {
+ if (!builder_result->error_msg.defined()) {
+ runner_inputs.push_back(meta_schedule::RunnerInput(
+ /*artifact_path=*/builder_result->artifact_path.value(),
+ /*device_type=*/target->kind->name,
+ /*args_info=*/candidates.value()[idx]->args_info));
+ }
+ idx++;
+ }
+ Array<meta_schedule::RunnerFuture> runner_futures = runner.value()->Run(runner_inputs);
+ for (const meta_schedule::RunnerFuture& runner_future : runner_futures) {
+ meta_schedule::RunnerResult runner_result = runner_future->Result();
+ if (runner_result->error_msg.defined()) {
+ costs.push_back(1e10);
+ } else {
+ double sum = 0;
+ for (const FloatImm& cost : runner_result->run_secs.value()) {
+ sum += cost->value;
+ }
+ costs.push_back(sum / runner_result->run_secs.value().size());
+ }
+ }
+ ICHECK_EQ(costs.size(), results.size());
+ }
+ }
+ if (results.size() == 0) {
+ LOG(WARNING) << "No valid schedule found";
+ return prim_func;
+ }
+ int best_idx = 0;
+ if (runner.defined()) {
+ for (int i = 1; i < costs.size(); ++i) {
+ if (costs[i] < costs[best_idx]) {
+ best_idx = i;
+ }
+ }
+ } else {
+ best_idx = results.size() - 1;
Review Comment:
So, if we don't define runner, we pick the last one among the buildable candidates.
Is there any benefit over using `valid_count` of 1 in this case?
##########
include/tvm/relax/transform.h:
##########
@@ -481,6 +481,18 @@ TVM_DLL Pass DeadCodeElimination(Array<runtime::String> entry_functions);
*/
TVM_DLL Pass ToMixedPrecision(const DataType& out_dtype);
+/*!
+ * \brief The pass is designed for few shot tuning for static shape PrimFuncs. It examines all the
+ * blocks within the PrimFunc and conducts loop fusion, splitting, and other transformations based
+ * on MetaSchedule schedule rules but directly samples from the search space instead of using the
+ * tuning algorithm. User can specify the number of valid counts to try and whether to use runner
+ * for evaluation.
+ * \param valid_count The number of valid counts to try.
+ * \param runner The runner to evaluate the generated schedules.
+ * \return The Pass.
+ */
+TVM_DLL Pass FewShotTuning(Integer valid_count, ObjectRef runner);
Review Comment:
Passing `runner` might be too much overhead for most of the cases.
Can we use the default runner and use the boolean argument like `enable_runner`?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] zxybazh commented on a diff in pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on code in PR #14624:
URL: https://github.com/apache/tvm/pull/14624#discussion_r1175577652
##########
src/relax/transform/few_shot_tuning.cc:
##########
@@ -0,0 +1,169 @@
+/*
+ * 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 <tvm/relax/transform.h>
+
+#include "../../meta_schedule/utils.h"
+
+namespace tvm {
+namespace relax {
+namespace transform {
+
+tir::PrimFunc FewShotTunePrimFunc(const tir::PrimFunc& prim_func, const Target& target,
+ int64_t valid_count, Optional<meta_schedule::Runner> runner) {
+ // fetch a local builder
+ static const auto* f_get_local_builder =
+ runtime::Registry::Get("meta_schedule.builder.get_local_builder");
+ ICHECK(f_get_local_builder)
+ << "ValueError: Cannot find the packed function \"meta_schedule.builder.get_local_builder\"";
+ meta_schedule::Builder builder = (*f_get_local_builder)();
+ // create an IRModule
+ IRModule mod = IRModule(Map<GlobalVar, BaseFunc>({{GlobalVar("main"), prim_func}}));
+ // fetch the number of physical cores
+ static const auto* f_cpu_count = runtime::Registry::Get("meta_schedule.cpu_count");
+ ICHECK(f_cpu_count) << "ValueError: Cannot find the packed function \"meta_schedule._cpu_count\"";
+ int num_threads = (*f_cpu_count)(false);
+ // store the results
+ Array<IRModule> results;
+ std::vector<double> costs;
+ // create a TuneContext
+ meta_schedule::TuneContext task = meta_schedule::TuneContext(
+ /*mod=*/mod,
+ /*target=*/target,
+ /*space_generator=*/
+ meta_schedule::SpaceGenerator::PostOrderApply(/*f_block_filter=*/nullptr,
+ /*sch_rules=*/NullOpt,
+ /*postprocs=*/NullOpt,
+ /*mutator_probs=*/NullOpt),
+ /*search_strategy=*/meta_schedule::SearchStrategy::ReplayTrace(/*max_fail_count=*/100),
+ /*task_name=*/NullOpt,
+ /*num_threads=*/num_threads, // use all available local threads
+ /*rand_state=*/-1, // -1 means use random seed
+ /*logger=*/nullptr);
+ task->Initialize();
+ task->search_strategy.value()->PreTuning(
+ /*max_trials=*/valid_count, /*num_trials_per_iter=*/valid_count,
+ /*design_spaces=*/task->space_generator.value()->GenerateDesignSpace(mod),
+ /*database=*/NullOpt,
+ /*cost_model=*/NullOpt);
+ while (valid_count > 0) {
+ Optional<Array<meta_schedule::MeasureCandidate>> candidates =
+ task->search_strategy.value()->GenerateMeasureCandidates();
+ if (!candidates.defined()) break;
+ Array<meta_schedule::BuilderInput> builder_inputs;
+ for (const meta_schedule::MeasureCandidate& candidate : candidates.value()) {
+ builder_inputs.push_back(meta_schedule::BuilderInput(
+ /*mod=*/candidate->sch->mod(),
+ /*target=*/target));
+ }
+ Array<meta_schedule::BuilderResult> builder_results = builder->Build(builder_inputs);
+ ICHECK_EQ(builder_results.size(), candidates.value().size());
+ int idx = 0;
+ for (const meta_schedule::BuilderResult& builder_result : builder_results) {
+ if (!builder_result->error_msg.defined()) {
+ results.push_back(candidates.value()[idx]->sch->mod());
+ valid_count--;
+ }
+ idx++;
+ }
+ if (runner.defined()) {
+ Array<meta_schedule::RunnerInput> runner_inputs;
+ int idx = 0;
+ for (const meta_schedule::BuilderResult& builder_result : builder_results) {
+ if (!builder_result->error_msg.defined()) {
+ runner_inputs.push_back(meta_schedule::RunnerInput(
+ /*artifact_path=*/builder_result->artifact_path.value(),
+ /*device_type=*/target->kind->name,
+ /*args_info=*/candidates.value()[idx]->args_info));
+ }
+ idx++;
+ }
+ Array<meta_schedule::RunnerFuture> runner_futures = runner.value()->Run(runner_inputs);
+ for (const meta_schedule::RunnerFuture& runner_future : runner_futures) {
+ meta_schedule::RunnerResult runner_result = runner_future->Result();
+ if (runner_result->error_msg.defined()) {
+ costs.push_back(1e10);
+ } else {
+ double sum = 0;
+ for (const FloatImm& cost : runner_result->run_secs.value()) {
+ sum += cost->value;
+ }
+ costs.push_back(sum / runner_result->run_secs.value().size());
+ }
+ }
+ ICHECK_EQ(costs.size(), results.size());
+ }
+ }
+ if (results.size() == 0) {
+ LOG(WARNING) << "No valid schedule found";
+ return prim_func;
+ }
+ int best_idx = 0;
+ if (runner.defined()) {
+ for (int i = 1; i < costs.size(); ++i) {
+ if (costs[i] < costs[best_idx]) {
+ best_idx = i;
+ }
+ }
+ } else {
+ best_idx = results.size() - 1;
Review Comment:
Good point, if `valid_count` is `1`, I think there's no gain / benefit to generate more candidate without benchmarking. We can fix the option to default to 1 when benchmarking is not enabled.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tvm-bot commented on pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #14624:
URL: https://github.com/apache/tvm/pull/14624#issuecomment-1508021487
<!---bot-comment-->
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] zxybazh commented on a diff in pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on code in PR #14624:
URL: https://github.com/apache/tvm/pull/14624#discussion_r1175576516
##########
tests/python/relax/test_transform_few_shot_tuning.py:
##########
@@ -0,0 +1,56 @@
+# 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=invalid-name,,missing-function-docstring
+import tvm
+from tvm.tir.tensor_intrin.cuda import *
+from tvm.tir.tensor_intrin.x86 import *
+from tvm.relax.transform import FewShotTuning
+from tvm.script import tir as T
+import tvm.testing
+
+
+def test_matmul():
+ # pylint: disable=no-self-argument,missing-class-docstring,line-too-long
+ # fmt: off
+ @tvm.script.ir_module
+ class Before:
+ @T.prim_func
+ def matmul(
+ A: T.Buffer((32, 32), "float16"),
+ B: T.Buffer((32, 32), "float16"),
+ C: T.Buffer((32, 32), "float16"),
+ ):
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ # with T.block("root"):
+ for i, j, k in T.grid(32, 32, 32):
+ with T.block("C"):
+ v_i, v_j, v_k = T.axis.remap("SSR", [i, j, k])
+ T.reads(A[v_i, v_k], B[v_k, v_j])
+ T.writes(C[v_i, v_j])
+ with T.init():
+ C[v_i, v_j] = T.float16(0)
+ C[v_i, v_j] = C[v_i, v_j] + A[v_i, v_k] * B[v_k, v_j]
+ # fmt: on
+ # pylint: enable=no-self-argument,missing-class-docstring,line-too-long
+ target = tvm.target.Target("nvidia/geforce-rtx-3070")
+ with target, tvm.transform.PassContext(opt_level=3):
+ After = FewShotTuning()(Before)
+ After.show()
Review Comment:
Thanks for the suggestion! It seems we cannot know for sure which search space we will be using and what decision will be used in this case. Therefore, `structural_equal` might be unfit here, what about we check for the trace and see if the trace is one of our search space?
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] zxybazh commented on a diff in pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on code in PR #14624:
URL: https://github.com/apache/tvm/pull/14624#discussion_r1185621213
##########
tests/python/relax/test_transform_few_shot_tuning.py:
##########
@@ -0,0 +1,436 @@
+# 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=invalid-name,,missing-function-docstring
+from typing import List, Tuple
+import numpy as np
+
+import tvm
+from tvm.script import tir as T
+from tvm.tir.tensor_intrin.cuda import * # pylint: disable=wildcard-import,unused-wildcard-import
+from tvm.tir.tensor_intrin.x86 import * # pylint: disable=wildcard-import,unused-wildcard-import
+from tvm.meta_schedule.testing.tune_utils import generate_input_data
+from tvm.meta_schedule.arg_info import ArgInfo
+from tvm.relax.transform import FewShotTuning
+import tvm.testing
+
+# pylint: disable=no-self-argument,missing-class-docstring,line-too-long
+# fmt: off
+@tvm.script.ir_module
+class MatMul:
+ @T.prim_func
+ def matmul(
+ A: T.Buffer((32, 32), "float16"),
+ B: T.Buffer((32, 32), "float16"),
+ C: T.Buffer((32, 32), "float16"),
+ ):
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ # with T.block("root"):
+ for i, j, k in T.grid(32, 32, 32):
+ with T.block("C"):
+ v_i, v_j, v_k = T.axis.remap("SSR", [i, j, k])
+ T.reads(A[v_i, v_k], B[v_k, v_j])
+ T.writes(C[v_i, v_j])
+ with T.init():
+ C[v_i, v_j] = T.float16(0)
+ C[v_i, v_j] = C[v_i, v_j] + A[v_i, v_k] * B[v_k, v_j]
+# fmt: on
+# pylint: enable=no-self-argument,missing-class-docstring,line-too-long
+
+
+def _target():
+ return tvm.target.Target("llvm -num-cores=4")
+ # for local testing only
+ # return tvm.target.Target("nvidia/geforce-rtx-3070")
+
+
+def _get_input_output_info(func: tvm.tir.PrimFunc):
+ args = ArgInfo.from_prim_func(func)
+ inputs = [generate_input_data(x.shape, x.dtype) for x in args[:-1]]
+ output_shape = args[-1].shape
+ output_dtype = args[-1].dtype
+ return inputs, output_shape, output_dtype
+
+
+def _expected_results(
+ mod: tvm.ir.IRModule, inputs: List[np.ndarray], output_shape: Tuple, output_dtype: str
+):
+ rt_mod = tvm.build(mod, target="llvm")
+ data = [
+ tvm.nd.array(x)
+ for x in [
+ *inputs,
+ np.zeros(output_shape, dtype=output_dtype),
+ ]
+ ]
+ rt_mod(*data)
+ return data[-1].numpy()
+
+
+def _actual_results(
+ actual: tvm.ir.IRModule, inputs: List[np.ndarray], output_shape: Tuple, output_dtype: str
+):
+ target = _target()
+ actual_rt_mod = tvm.build(actual, target=target)
+ actual_data = [
+ tvm.nd.array(x, device=tvm.cuda() if target.kind.name == "cuda" else tvm.cpu())
+ for x in [
+ *inputs,
+ np.zeros(output_shape, dtype=output_dtype),
+ ]
+ ]
+ actual_rt_mod(*actual_data)
+ return actual_data[-1].numpy()
+
+
+def _assert_allclose(mod: tvm.ir.IRModule, actual: tvm.ir.IRModule, func: tvm.tir.PrimFunc):
+ inputs, output_shape, output_dtype = _get_input_output_info(func)
+ expected_output = _expected_results(mod, inputs, output_shape, output_dtype)
+ actual_output = _actual_results(actual, inputs, output_shape, output_dtype)
+ tvm.testing.assert_allclose(expected_output, actual_output, rtol=1e-2, atol=1e-2)
Review Comment:
Thanks Sung, the reason we were using 1e-2 is because the unit test was running on cuda, which has slight difference compare to on our baseline on llvm. Now that the default target is changed to llvm I'll change the threshold accordingly!
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] zxybazh commented on a diff in pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on code in PR #14624:
URL: https://github.com/apache/tvm/pull/14624#discussion_r1186746272
##########
tests/python/relax/test_transform_few_shot_tuning.py:
##########
@@ -0,0 +1,436 @@
+# 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=invalid-name,,missing-function-docstring
+from typing import List, Tuple
+import numpy as np
+
+import tvm
+from tvm.script import tir as T
+from tvm.tir.tensor_intrin.cuda import * # pylint: disable=wildcard-import,unused-wildcard-import
+from tvm.tir.tensor_intrin.x86 import * # pylint: disable=wildcard-import,unused-wildcard-import
+from tvm.meta_schedule.testing.tune_utils import generate_input_data
+from tvm.meta_schedule.arg_info import ArgInfo
+from tvm.relax.transform import FewShotTuning
+import tvm.testing
+
+# pylint: disable=no-self-argument,missing-class-docstring,line-too-long
+# fmt: off
+@tvm.script.ir_module
+class MatMul:
+ @T.prim_func
+ def matmul(
+ A: T.Buffer((32, 32), "float16"),
+ B: T.Buffer((32, 32), "float16"),
+ C: T.Buffer((32, 32), "float16"),
+ ):
+ T.func_attr({"global_symbol": "main", "tir.noalias": True})
+ # with T.block("root"):
+ for i, j, k in T.grid(32, 32, 32):
+ with T.block("C"):
+ v_i, v_j, v_k = T.axis.remap("SSR", [i, j, k])
+ T.reads(A[v_i, v_k], B[v_k, v_j])
+ T.writes(C[v_i, v_j])
+ with T.init():
+ C[v_i, v_j] = T.float16(0)
+ C[v_i, v_j] = C[v_i, v_j] + A[v_i, v_k] * B[v_k, v_j]
+# fmt: on
+# pylint: enable=no-self-argument,missing-class-docstring,line-too-long
+
+
+def _target():
+ return tvm.target.Target("llvm -num-cores=4")
+ # for local testing only
+ # return tvm.target.Target("nvidia/geforce-rtx-3070")
+
+
+def _get_input_output_info(func: tvm.tir.PrimFunc):
+ args = ArgInfo.from_prim_func(func)
+ inputs = [generate_input_data(x.shape, x.dtype) for x in args[:-1]]
+ output_shape = args[-1].shape
+ output_dtype = args[-1].dtype
+ return inputs, output_shape, output_dtype
+
+
+def _expected_results(
+ mod: tvm.ir.IRModule, inputs: List[np.ndarray], output_shape: Tuple, output_dtype: str
+):
+ rt_mod = tvm.build(mod, target="llvm")
+ data = [
+ tvm.nd.array(x)
+ for x in [
+ *inputs,
+ np.zeros(output_shape, dtype=output_dtype),
+ ]
+ ]
+ rt_mod(*data)
+ return data[-1].numpy()
+
+
+def _actual_results(
+ actual: tvm.ir.IRModule, inputs: List[np.ndarray], output_shape: Tuple, output_dtype: str
+):
+ target = _target()
+ actual_rt_mod = tvm.build(actual, target=target)
+ actual_data = [
+ tvm.nd.array(x, device=tvm.cuda() if target.kind.name == "cuda" else tvm.cpu())
+ for x in [
+ *inputs,
+ np.zeros(output_shape, dtype=output_dtype),
+ ]
+ ]
+ actual_rt_mod(*actual_data)
+ return actual_data[-1].numpy()
+
+
+def _assert_allclose(mod: tvm.ir.IRModule, actual: tvm.ir.IRModule, func: tvm.tir.PrimFunc):
+ inputs, output_shape, output_dtype = _get_input_output_info(func)
+ expected_output = _expected_results(mod, inputs, output_shape, output_dtype)
+ actual_output = _actual_results(actual, inputs, output_shape, output_dtype)
+ tvm.testing.assert_allclose(expected_output, actual_output, rtol=1e-2, atol=1e-2)
Review Comment:
The accuracy has been fixed. For softmax test, fp16 doesn't satisfy 1e-7 accuracy but fp32 can work so I changed the PrimFunc to work in fp32. For Fused Variance Test I found that there's an issue with RewriteReduction Postproc so I've reported this issue with a reproducible script in #14791.
Other than that, the tests are all fixed and working.
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen merged pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "tqchen (via GitHub)" <gi...@apache.org>.
tqchen merged PR #14624:
URL: https://github.com/apache/tvm/pull/14624
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] zxybazh commented on a diff in pull request #14624: [Unity] Introduce FewShotTuning Pass
Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on code in PR #14624:
URL: https://github.com/apache/tvm/pull/14624#discussion_r1175599526
##########
include/tvm/relax/transform.h:
##########
@@ -481,6 +481,18 @@ TVM_DLL Pass DeadCodeElimination(Array<runtime::String> entry_functions);
*/
TVM_DLL Pass ToMixedPrecision(const DataType& out_dtype);
+/*!
+ * \brief The pass is designed for few shot tuning for static shape PrimFuncs. It examines all the
+ * blocks within the PrimFunc and conducts loop fusion, splitting, and other transformations based
+ * on MetaSchedule schedule rules but directly samples from the search space instead of using the
+ * tuning algorithm. User can specify the number of valid counts to try and whether to use runner
+ * for evaluation.
+ * \param valid_count The number of valid counts to try.
+ * \param runner The runner to evaluate the generated schedules.
+ * \return The Pass.
+ */
+TVM_DLL Pass FewShotTuning(Integer valid_count, ObjectRef runner);
Review Comment:
Yeah sure, will do!
--
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: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org