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: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
