From 66901adc7533e7abcafed50828ee64e6a517a064 Mon Sep 17 00:00:00 2001 From: Tristan Konolige Date: Thu, 15 Oct 2020 16:10:33 -0600 Subject: [PATCH 1/8] Fix multiprocessing with spawn issues --- python/tvm/auto_scheduler/measure.py | 539 ++++++++++-------- python/tvm/auto_scheduler/measure_record.py | 39 +- python/tvm/auto_scheduler/utils.py | 41 +- .../tvm/auto_scheduler/workload_registry.py | 35 ++ python/tvm/autotvm/measure/local_executor.py | 10 +- python/tvm/autotvm/task/task.py | 14 +- .../tvm/autotvm/tuner/xgboost_cost_model.py | 29 +- python/tvm/testing.py | 17 + src/auto_scheduler/measure_record.cc | 73 ++- .../unittest/test_auto_scheduler_measure.py | 6 +- .../test_auto_scheduler_search_policy.py | 4 +- .../unittest/test_autotvm_dispatch_context.py | 11 +- tests/python/unittest/test_runtime_rpc.py | 143 ++--- 13 files changed, 549 insertions(+), 412 deletions(-) mode change 100755 => 100644 src/auto_scheduler/measure_record.cc diff --git a/python/tvm/auto_scheduler/measure.py b/python/tvm/auto_scheduler/measure.py index 8a8b92201d15..1f9c9bf2c157 100644 --- a/python/tvm/auto_scheduler/measure.py +++ b/python/tvm/auto_scheduler/measure.py @@ -51,19 +51,52 @@ from .loop_state import StateObject from .utils import ( get_const_tuple, - NoDaemonPool, call_func_with_timeout, request_remote, check_remote, ) +from .compute_dag import ComputeDAG +from .search_task import SearchTask +from .workload_registry import workload_name, get_workload # The maximum length of error message MAX_ERROR_MSG_LEN = 512 -# We use fork and a global variable to copy arguments between processes. -# This can avoid expensive serialization of TVM IR when using multiprocessing.Pool -GLOBAL_BUILD_ARGUMENTS = None -GLOBAL_RUN_ARGUMENTS = None + +def recover_measure_input(inp, rebuild_state=False): + """ + Recover a deserialized MeasureInput by rebuilding the missing fields. + 1. Rebuid the compute_dag in inp.task + 2. (Optional) Rebuild the stages in inp.state + + Parameters + ---------- + inp: MeasureInput + The deserialized MeasureInput + rebuild_state: bool = False + Whether rebuild the stages in MeasureInput.State + + Returns + ------- + new_input: MeasureInput + The fully recovered MeasureInput with all fields rebuilt. + """ + task = inp.task + print(task.hardware_params) + new_task = SearchTask( + ComputeDAG(task.workload_key), + task.workload_key, + task.target, + task.target_host, + task.hardware_params, + ) + + if rebuild_state: + new_state = new_task.compute_dag.infer_bound_from_state(inp.state) + else: + new_state = inp.state + + return MeasureInput(new_task, new_state) @tvm._ffi.register_object("auto_scheduler.MeasureCallback") @@ -87,6 +120,24 @@ def __init__(self, task, state): state = state if isinstance(state, StateObject) else state.state_object self.__init_handle_by_constructor__(_ffi_api.MeasureInput, task, state) + def serialize(self): + serialize = tvm.get_global_func("auto_scheduler.SerializeMeasureInput", True) + assert serialize + # We serialize the workload function so that it can be used on the deserialized side. + return { + "measureinput": serialize(self), + "name": workload_name(self.task.workload_key), + "func": get_workload(self.task), + } + + @staticmethod + def deserialize(state): + deserialize = tvm.get_global_func("auto_scheduler.DeserializeMeasureInput", True) + assert deserialize + tvm.auto_scheduler.workload_registry.WORKLOAD_FUNC_REGISTRY[state["name"]] = state["func"] + x = deserialize(state["measureinput"]) + return recover_measure_input(x) + @tvm._ffi.register_object("auto_scheduler.BuildResult") class BuildResult(Object): @@ -486,29 +537,63 @@ def make_error_msg(): return error_msg -def local_build_worker(index): +def _timed_func(inp_serialized, build_func, verbose): + tic = time.time() + inp = MeasureInput.deserialize(inp_serialized) + task = inp.task + + error_no = MeasureErrorNo.NO_ERROR + error_msg = None + args = [] + + try: + sch, args = task.compute_dag.apply_steps_from_state(inp.state, layout_rewrite=True) + # pylint: disable=broad-except + except Exception: + error_no = MeasureErrorNo.INSTANTIATION_ERROR + error_msg = make_error_msg() + + if error_no == 0: + dirname = tempfile.mkdtemp() + filename = os.path.join(dirname, "tmp_func." + build_func.output_format) + + try: + # TODO(merrymercy): Port the unroll pass. + with transform.PassContext(): + func = build_module.build( + sch, args, target=task.target, target_host=task.target_host + ) + func.export_library(filename, build_func) + # pylint: disable=broad-except + except Exception: + error_no = MeasureErrorNo.COMPILE_HOST + error_msg = make_error_msg() + else: + filename = "" + + if verbose >= 1: + if error_no == MeasureErrorNo.NO_ERROR: + print(".", end="") + else: + print(".E", end="") # Build error + return filename, args, error_no, error_msg, time.time() - tic + + +def local_build_worker(args): """ Build function of LocalBuilder to be ran in the Builder thread pool. Parameters ---------- - index : int - The MeasureInput index to be processed by the current Builder thread. + args: Tuple[MeasureInput, str, int, int] + inputs, build-func, time, verbose args passed to local_builder_build Returns ------- res : BuildResult The build result of this Builder thread. """ - global GLOBAL_BUILD_ARGUMENTS - - # We use fork and a global variable to copy arguments between processes. - # This can avoid expensive serialization of TVM IR when using multiprocessing.Pool - if not GLOBAL_BUILD_ARGUMENTS: - raise ValueError("GLOBAL_BUILD_ARGUMENTS not found") - measure_inputs, build_func, timeout, verbose = GLOBAL_BUILD_ARGUMENTS - assert isinstance(build_func, str) - + inp, build_func, timeout, verbose = args if build_func == "default": build_func = tar.tar elif build_func == "ndk": @@ -516,48 +601,7 @@ def local_build_worker(index): else: raise ValueError("Invalid build_func" + build_func) - def timed_func(): - tic = time.time() - inp = measure_inputs[index] - task = inp.task - - error_no = MeasureErrorNo.NO_ERROR - error_msg = None - args = [] - - try: - sch, args = task.compute_dag.apply_steps_from_state(inp.state, layout_rewrite=True) - # pylint: disable=broad-except - except Exception: - error_no = MeasureErrorNo.INSTANTIATION_ERROR - error_msg = make_error_msg() - - if error_no == 0: - dirname = tempfile.mkdtemp() - filename = os.path.join(dirname, "tmp_func." + build_func.output_format) - - try: - # TODO(merrymercy): Port the unroll pass. - with transform.PassContext(): - func = build_module.build( - sch, args, target=task.target, target_host=task.target_host - ) - func.export_library(filename, build_func) - # pylint: disable=broad-except - except Exception: - error_no = MeasureErrorNo.COMPILE_HOST - error_msg = make_error_msg() - else: - filename = "" - - if verbose >= 1: - if error_no == MeasureErrorNo.NO_ERROR: - print(".", end="") - else: - print(".E", end="") # Build error - return filename, args, error_no, error_msg, time.time() - tic - - res = call_func_with_timeout(timeout, timed_func) + res = call_func_with_timeout(timeout, _timed_func, args=(inp, build_func, verbose)) if isinstance(res, TimeoutError): if verbose >= 1: print(".T", end="") # Build timeout @@ -590,14 +634,20 @@ def local_builder_build(inputs, timeout, n_parallel, build_func="default", verbo res : List[BuildResult] The build results of these MeasureInputs. """ - # We use fork and a global variable to copy arguments between processes. - # This can avoid expensive serialization of TVM IR when using multiprocessing.Pool - global GLOBAL_BUILD_ARGUMENTS - - GLOBAL_BUILD_ARGUMENTS = (inputs, build_func, timeout, verbose) - - pool = NoDaemonPool(n_parallel) - tuple_res = pool.map(local_build_worker, range(len(inputs))) + # This pool is not doing computationally intensive work, so we can use threads + pool = multiprocessing.pool.ThreadPool(n_parallel) + tuple_res = pool.map( + local_build_worker, + [ + ( + i.serialize(), + build_func, + timeout, + verbose, + ) + for i in inputs + ], + ) pool.terminate() pool.join() del pool @@ -609,6 +659,70 @@ def local_builder_build(inputs, timeout, n_parallel, build_func="default", verbo return results +def _timed_eval_func( + inp_serialized, + build_res, + number, + repeat, + min_repeat_ms, + cooldown_interval, + enable_cpu_cache_flush, + verbose, +): + inp = MeasureInput.deserialize(inp_serialized) + tic = time.time() + error_no = 0 + error_msg = None + try: + func = module.load_module(build_res.filename) + ctx = ndarray.context(str(inp.task.target), 0) + # Limitation: + # We can not get PackFunction directly in the remote mode as it is wrapped + # under the std::function. We could lift the restriction later once we fold + # the PackedFunc as an object. Currently, we pass function name to work + # around it. + f_prepare = "cache_flush_cpu_non_first_arg" if enable_cpu_cache_flush else "" + time_f = func.time_evaluator( + func.entry_name, + ctx, + number=number, + repeat=repeat, + min_repeat_ms=min_repeat_ms, + f_preproc=f_prepare, + ) + # pylint: disable=broad-except + except Exception: + costs = (max_float,) + error_no = MeasureErrorNo.COMPILE_DEVICE + error_msg = make_error_msg() + + if error_no == 0: + try: + args = [ndarray.empty(get_const_tuple(x.shape), x.dtype, ctx) for x in build_res.args] + random_fill = tvm.get_global_func("tvm.contrib.random.random_fill", True) + assert random_fill, "Please make sure USE_RANDOM is ON in the config.cmake" + for arg in args: + random_fill(arg) + ctx.sync() + costs = time_f(*args).results + # pylint: disable=broad-except + except Exception: + costs = (max_float,) + error_no = MeasureErrorNo.RUNTIME_DEVICE + error_msg = make_error_msg() + + shutil.rmtree(os.path.dirname(build_res.filename)) + toc = time.time() + time.sleep(cooldown_interval) + + if verbose >= 1: + if error_no == MeasureErrorNo.NO_ERROR: + print("*", end="") + else: + print("*E", end="") # Run error + return costs, error_no, error_msg, toc - tic + build_res.time_cost, toc + + @tvm._ffi.register_func("auto_scheduler.local_runner.run") def local_run( inputs, @@ -667,61 +781,6 @@ def local_run( """ max_float = 1e10 # We use 1e10 instead of sys.float_info.max for better readability in log - def timed_func(inp, build_res): - tic = time.time() - error_no = 0 - error_msg = None - try: - func = module.load_module(build_res.filename) - ctx = ndarray.context(str(inp.task.target), 0) - # Limitation: - # We can not get PackFunction directly in the remote mode as it is wrapped - # under the std::function. We could lift the restriction later once we fold - # the PackedFunc as an object. Currently, we pass function name to work - # around it. - f_prepare = "cache_flush_cpu_non_first_arg" if enable_cpu_cache_flush else "" - time_f = func.time_evaluator( - func.entry_name, - ctx, - number=number, - repeat=repeat, - min_repeat_ms=min_repeat_ms, - f_preproc=f_prepare, - ) - # pylint: disable=broad-except - except Exception: - costs = (max_float,) - error_no = MeasureErrorNo.COMPILE_DEVICE - error_msg = make_error_msg() - - if error_no == 0: - try: - args = [ - ndarray.empty(get_const_tuple(x.shape), x.dtype, ctx) for x in build_res.args - ] - random_fill = tvm.get_global_func("tvm.contrib.random.random_fill", True) - assert random_fill, "Please make sure USE_RANDOM is ON in the config.cmake" - for arg in args: - random_fill(arg) - ctx.sync() - costs = time_f(*args).results - # pylint: disable=broad-except - except Exception: - costs = (max_float,) - error_no = MeasureErrorNo.RUNTIME_DEVICE - error_msg = make_error_msg() - - shutil.rmtree(os.path.dirname(build_res.filename)) - toc = time.time() - time.sleep(cooldown_interval) - - if verbose >= 1: - if error_no == MeasureErrorNo.NO_ERROR: - print("*", end="") - else: - print("*E", end="") # Run error - return costs, error_no, error_msg, toc - tic + build_res.time_cost, toc - measure_results = [] assert len(inputs) == len(build_results), "Measure input size should be equal to build results" for inp, build_res in zip(inputs, build_results): @@ -734,7 +793,20 @@ def timed_func(inp, build_res): time.time(), ) else: - res = call_func_with_timeout(timeout, timed_func, args=(inp, build_res)) + res = call_func_with_timeout( + timeout, + _timed_eval_func, + args=( + inp.serialize(), + build_res, + number, + repeat, + min_repeat_ms, + cooldown_interval, + enable_cpu_cache_flush, + verbose, + ), + ) if isinstance(res, TimeoutError): if verbose >= 1: print("*T", end="") # Run timeout @@ -753,40 +825,104 @@ def timed_func(inp, build_res): return measure_results -def rpc_run_worker(index): +def _timed_rpc_run( + inp_serialized, + build_res, + key, + host, + port, + priority, + timeout, + number, + repeat, + min_repeat_ms, + cooldown_interval, + enable_cpu_cache_flush, + verbose, +): + inp = MeasureInput.deserialize(inp_serialized) + tic = time.time() + error_no = 0 + error_msg = None + try: + # upload built module + remote = request_remote(key, host, port, priority, timeout) + remote.upload(build_res.filename) + func = remote.load_module(os.path.split(build_res.filename)[1]) + ctx = remote.context(str(inp.task.target), 0) + # Limitation: + # We can not get PackFunction directly in the remote mode as it is wrapped + # under the std::function. We could lift the restriction later once we fold + # the PackedFunc as an object. Currently, we pass function name to work + # around it. + f_prepare = "cache_flush_cpu_non_first_arg" if enable_cpu_cache_flush else "" + time_f = func.time_evaluator( + func.entry_name, + ctx, + number=number, + repeat=repeat, + min_repeat_ms=min_repeat_ms, + f_preproc=f_prepare, + ) + # pylint: disable=broad-except + except Exception: + costs = (max_float,) + error_no = MeasureErrorNo.COMPILE_DEVICE + error_msg = make_error_msg() + + if error_no == 0: + try: + args = [ndarray.empty(get_const_tuple(x.shape), x.dtype, ctx) for x in build_res.args] + try: + random_fill = remote.get_function("tvm.contrib.random.random_fill") + except AttributeError: + raise AttributeError( + "Please make sure USE_RANDOM is ON in the config.cmake " "on the remote devices" + ) + for arg in args: + random_fill(arg) + ctx.sync() + + costs = time_f(*args).results + # clean up remote files + remote.remove(build_res.filename) + remote.remove(os.path.splitext(build_res.filename)[0] + ".so") + remote.remove("") + # pylint: disable=broad-except + except Exception: + costs = (max_float,) + error_no = MeasureErrorNo.RUNTIME_DEVICE + error_msg = make_error_msg() + + shutil.rmtree(os.path.dirname(build_res.filename)) + toc = time.time() + + time.sleep(cooldown_interval) + if verbose >= 1: + if error_no == MeasureErrorNo.NO_ERROR: + print("*", end="") + else: + print("*E", end="") # Run error + + return costs, error_no, error_msg, toc - tic + build_res.time_cost, toc + + +def _rpc_run_worker(args): """Function to be ran in the RPCRunner thread pool. Parameters ---------- - index : int - The MeasureInput and BuildResult index to be processed by the current Runner thread. + args : Tuple[MeasureInput, BuildResult, ...] + Single input and build result plus the rest of the arguments to `rpc_runner_run`. Returns ------- res : MeasureResult The measure result of this Runner thread. """ - global GLOBAL_RUN_ARGUMENTS - ( - inputs, - build_results, - key, - host, - port, - priority, - timeout, - number, - repeat, - min_repeat_ms, - cooldown_interval, - enable_cpu_cache_flush, - verbose, - ) = GLOBAL_RUN_ARGUMENTS - max_float = 1e10 # We use 1e10 instead of sys.float_info.max for better readability in log - inp = inputs[index] - build_res = build_results[index] + _, build_res, _, _, _, _, timeout, _, _, _, _, _, verbose = args if build_res.error_no != MeasureErrorNo.NO_ERROR: return ( (max_float,), @@ -796,76 +932,7 @@ def rpc_run_worker(index): time.time(), ) - def timed_func(): - tic = time.time() - error_no = 0 - error_msg = None - try: - # upload built module - remote = request_remote(key, host, port, priority, timeout) - remote.upload(build_res.filename) - func = remote.load_module(os.path.split(build_res.filename)[1]) - ctx = remote.context(str(inp.task.target), 0) - # Limitation: - # We can not get PackFunction directly in the remote mode as it is wrapped - # under the std::function. We could lift the restriction later once we fold - # the PackedFunc as an object. Currently, we pass function name to work - # around it. - f_prepare = "cache_flush_cpu_non_first_arg" if enable_cpu_cache_flush else "" - time_f = func.time_evaluator( - func.entry_name, - ctx, - number=number, - repeat=repeat, - min_repeat_ms=min_repeat_ms, - f_preproc=f_prepare, - ) - # pylint: disable=broad-except - except Exception: - costs = (max_float,) - error_no = MeasureErrorNo.COMPILE_DEVICE - error_msg = make_error_msg() - - if error_no == 0: - try: - args = [ - ndarray.empty(get_const_tuple(x.shape), x.dtype, ctx) for x in build_res.args - ] - try: - random_fill = remote.get_function("tvm.contrib.random.random_fill") - except AttributeError: - raise AttributeError( - "Please make sure USE_RANDOM is ON in the config.cmake " - "on the remote devices" - ) - for arg in args: - random_fill(arg) - ctx.sync() - - costs = time_f(*args).results - # clean up remote files - remote.remove(build_res.filename) - remote.remove(os.path.splitext(build_res.filename)[0] + ".so") - remote.remove("") - # pylint: disable=broad-except - except Exception: - costs = (max_float,) - error_no = MeasureErrorNo.RUNTIME_DEVICE - error_msg = make_error_msg() - - shutil.rmtree(os.path.dirname(build_res.filename)) - toc = time.time() - - time.sleep(cooldown_interval) - if verbose >= 1: - if error_no == MeasureErrorNo.NO_ERROR: - print("*", end="") - else: - print("*E", end="") # Run error - - return costs, error_no, error_msg, toc - tic + build_res.time_cost, toc - - res = call_func_with_timeout(timeout, timed_func) + res = call_func_with_timeout(timeout, _timed_rpc_run, args=args) if isinstance(res, TimeoutError): if verbose >= 1: @@ -950,26 +1017,30 @@ def rpc_runner_run( res : List[MeasureResult] The measure results of these MeasureInputs. """ - global GLOBAL_RUN_ARGUMENTS - GLOBAL_RUN_ARGUMENTS = ( - inputs, - build_results, - key, - host, - port, - priority, - timeout, - number, - repeat, - min_repeat_ms, - cooldown_interval, - enable_cpu_cache_flush, - verbose, - ) - assert len(inputs) == len(build_results), "Measure input size should be equal to build results" - pool = NoDaemonPool(n_parallel) - tuple_res = pool.map(rpc_run_worker, range(len(build_results))) + # This pool is not doing computationally intensive work, so we can use threads + pool = multiprocessing.pool.ThreadPool(n_parallel) + tuple_res = pool.map( + _rpc_run_worker, + [ + ( + inp.serialize(), + build_res, + key, + host, + port, + priority, + timeout, + number, + repeat, + min_repeat_ms, + cooldown_interval, + enable_cpu_cache_flush, + verbose, + ) + for inp, build_res in zip(inputs, build_results) + ], + ) pool.terminate() pool.join() del pool diff --git a/python/tvm/auto_scheduler/measure_record.py b/python/tvm/auto_scheduler/measure_record.py index 1d0d7650a0f6..f0d930e3257e 100644 --- a/python/tvm/auto_scheduler/measure_record.py +++ b/python/tvm/auto_scheduler/measure_record.py @@ -21,9 +21,7 @@ import tvm._ffi from tvm.runtime import Object -from .compute_dag import ComputeDAG -from .measure import MeasureErrorNo, MeasureInput, MeasureCallback -from .search_task import SearchTask +from .measure import MeasureErrorNo, MeasureCallback from . import _ffi_api @@ -175,38 +173,3 @@ def load_best(filename, workload_key=None, target=None): best_res = res return best_inp, best_res - - -def recover_measure_input(inp, rebuild_state=False): - """ - Recover a deserialized MeasureInput by rebuilding the missing fields. - 1. Rebuid the compute_dag in inp.task - 2. (Optional) Rebuild the stages in inp.state - - Parameters - ---------- - inp: MeasureInput - The deserialized MeasureInput - rebuild_state: bool = False - Whether rebuild the stages in MeasureInput.State - - Returns - ------- - new_input: MeasureInput - The fully recovered MeasureInput with all fields rebuilt. - """ - task = inp.task - new_task = SearchTask( - ComputeDAG(task.workload_key), - task.workload_key, - task.target, - task.target_host, - task.hardware_params, - ) - - if rebuild_state: - new_state = new_task.compute_dag.infer_bound_from_state(inp.state) - else: - new_state = inp.state - - return MeasureInput(new_task, new_state) diff --git a/python/tvm/auto_scheduler/utils.py b/python/tvm/auto_scheduler/utils.py index 75fec9c891e8..49a97d0d6ced 100644 --- a/python/tvm/auto_scheduler/utils.py +++ b/python/tvm/auto_scheduler/utils.py @@ -129,32 +129,6 @@ def deserialize_args(args): return ret -class NoDaemonProcess(multiprocessing.Process): - @property - def daemon(self): - return False - - @daemon.setter - def daemon(self, value): - pass - - -class NoDaemonContext(type(multiprocessing.get_context())): - Process = NoDaemonProcess - - -class NoDaemonPool(multiprocessing.pool.Pool): - """A no daemon pool version of multiprocessing.Pool. - This allows us to start new processes inside the worker function""" - - def __init__(self, *args, **kwargs): - kwargs["context"] = NoDaemonContext() - super().__init__(*args, **kwargs) - - def __reduce__(self): - pass - - def kill_child_processes(parent_pid, sig=signal.SIGTERM): """kill all child processes recursively""" try: @@ -169,17 +143,18 @@ def kill_child_processes(parent_pid, sig=signal.SIGTERM): return +def _func_wrapper(que, func, args, kwargs): + if kwargs: + que.put(func(*args, **kwargs)) + else: + que.put(func(*args)) + + def call_func_with_timeout(timeout, func, args=(), kwargs=None): """Call a function with timeout""" - def func_wrapper(que): - if kwargs: - que.put(func(*args, **kwargs)) - else: - que.put(func(*args)) - que = multiprocessing.Queue(2) - process = multiprocessing.Process(target=func_wrapper, args=(que,)) + process = multiprocessing.Process(target=_func_wrapper, args=(que, func, args, kwargs)) process.start() process.join(timeout) diff --git a/python/tvm/auto_scheduler/workload_registry.py b/python/tvm/auto_scheduler/workload_registry.py index 1d9ee6da4f7a..6bf183b9fd98 100644 --- a/python/tvm/auto_scheduler/workload_registry.py +++ b/python/tvm/auto_scheduler/workload_registry.py @@ -175,6 +175,41 @@ def workload_key_to_tensors(workload_key): return lookup(*args) +def get_workload(task): + """Get the workload function for a given task + + Parameters + ---------- + task : SearchTask + Task to get workload of. + + Returns + ------- + workload : callable + The registered workload function. + """ + name = workload_name(task.workload_key) + lookup = WORKLOAD_FUNC_REGISTRY[name] + assert callable(lookup) + return lookup + + +def workload_name(workload_key): + """Decode a workload key to the registered function name. + + Parameters + ---------- + workload_key : str + The input workload key. + + Returns + ------- + name : str + The function name of this workload key. + """ + return decode_workload_key_to_func_args(workload_key)[0] + + def save_workload_func_registry(filename): """Dump workload function registry to a pickle binary file. diff --git a/python/tvm/autotvm/measure/local_executor.py b/python/tvm/autotvm/measure/local_executor.py index 5dd5cba2b824..af1dc1213292 100644 --- a/python/tvm/autotvm/measure/local_executor.py +++ b/python/tvm/autotvm/measure/local_executor.py @@ -18,7 +18,7 @@ import signal -from multiprocessing import Process, Queue +import multiprocessing try: from queue import Empty @@ -60,7 +60,7 @@ def call_with_timeout(queue, timeout, func, args, kwargs): """A wrapper to support timeout of a function call""" # start a new process for timeout (cannot use thread because we have c function) - p = Process(target=_execute_func, args=(func, queue, args, kwargs)) + p = multiprocessing.Process(target=_execute_func, args=(func, queue, args, kwargs)) p.start() p.join(timeout=timeout) @@ -151,7 +151,9 @@ def submit(self, func, *args, **kwargs): if not self.do_fork: return LocalFutureNoFork(func(*args, **kwargs)) - queue = Queue(2) # Size of 2 to avoid a race condition with size 1. - process = Process(target=call_with_timeout, args=(queue, self.timeout, func, args, kwargs)) + queue = multiprocessing.Queue(2) # Size of 2 to avoid a race condition with size 1. + process = multiprocessing.Process( + target=call_with_timeout, args=(queue, self.timeout, func, args, kwargs) + ) process.start() return LocalFuture(process, queue) diff --git a/python/tvm/autotvm/task/task.py b/python/tvm/autotvm/task/task.py index a7cb9a095765..8822ba971e4c 100644 --- a/python/tvm/autotvm/task/task.py +++ b/python/tvm/autotvm/task/task.py @@ -23,15 +23,14 @@ """ import numpy as np -from tvm.target import Target from tvm import runtime from tvm.ir import container +from tvm.target import Target +from tvm.te import placeholder, tensor from tvm.tir import expr -from tvm.te import tensor, placeholder - from ..util import get_const_int, get_const_tuple -from .dispatcher import DispatchContext, ApplyConfig +from .dispatcher import ApplyConfig, DispatchContext from .space import ConfigSpace @@ -173,6 +172,8 @@ def __getstate__(self): # some unpickable local task functions. # So we only pickle the name of the function # and restore the function by name when unpickling it. + import cloudpickle # pylint: disable=import-outside-toplevel + return { "name": self.name, "args": self.args, @@ -181,14 +182,17 @@ def __getstate__(self): "flop": self.flop, "target": self.target, "target_host": self.target_host, + "func": cloudpickle.dumps(self.func), } def __setstate__(self, state): + import cloudpickle # pylint: disable=import-outside-toplevel + self.name = state["name"] self.args = state["args"] self.kwargs = state["kwargs"] self.config_space = state["config_space"] - self.func = _lookup_task(state["name"]) + self.func = cloudpickle.loads(state["func"]) self.flop = state["flop"] self.target = state["target"] self.target_host = state["target_host"] diff --git a/python/tvm/autotvm/tuner/xgboost_cost_model.py b/python/tvm/autotvm/tuner/xgboost_cost_model.py index 7b9df1c99373..f2e6eb1a2e44 100644 --- a/python/tvm/autotvm/tuner/xgboost_cost_model.py +++ b/python/tvm/autotvm/tuner/xgboost_cost_model.py @@ -153,11 +153,6 @@ def _reset_pool(self, space, target, task): self._close_pool() - # use global variable to pass common arguments - global _extract_space, _extract_target, _extract_task - _extract_space = space - _extract_target = target - _extract_task = task self.pool = multiprocessing.Pool(self.num_threads) def _close_pool(self): @@ -321,10 +316,11 @@ def _get_feature(self, indexes): indexes = np.array(indexes) need_extract = [x for x in indexes if x not in fea_cache] + args = [(self.space.get(x), self.target, self.task) for x in need_extract] if need_extract: pool = self._get_pool() - feas = pool.map(self.feature_extract_func, need_extract) + feas = pool.map(self.feature_extract_func, args) for i, fea in zip(need_extract, feas): fea_cache[i] = fea @@ -344,17 +340,16 @@ def __del__(self): self._close_pool() -_extract_space = None _extract_target = None _extract_task = None -def _extract_itervar_feature_index(index): +def _extract_itervar_feature_index(args): """extract iteration var feature for an index in extract_space""" try: - config = _extract_space.get(index) - with _extract_target: - sch, args = _extract_task.instantiate(config) + config, target, task = args + with target: + sch, args = task.instantiate(config) fea = feature.get_itervar_feature_flatten(sch, args, take_log=True) fea = np.concatenate((fea, list(config.get_other_option().values()))) return fea @@ -381,10 +376,10 @@ def _extract_itervar_feature_log(arg): return None -def _extract_knob_feature_index(index): +def _extract_knob_feature_index(args): """extract knob feature for an index in extract_space""" try: - config = _extract_space.get(index) + config, _, _ = args return config.get_flatten_feature() except Exception: # pylint: disable=broad-except return None @@ -408,12 +403,12 @@ def _extract_knob_feature_log(arg): return None -def _extract_curve_feature_index(index): +def _extract_curve_feature_index(args): """extract sampled curve feature for an index in extract_space""" try: - config = _extract_space.get(index) - with _extract_target: - sch, args = _extract_task.instantiate(config) + config, target, task = args + with target: + sch, args = task.instantiate(config) fea = feature.get_buffer_curve_sample_flatten(sch, args, sample_n=20) fea = np.concatenate((fea, list(config.get_other_option().values()))) return np.array(fea) diff --git a/python/tvm/testing.py b/python/tvm/testing.py index 51fa2d0d7def..e5b17f3d7b53 100644 --- a/python/tvm/testing.py +++ b/python/tvm/testing.py @@ -634,6 +634,23 @@ def requires_micro(*args): return _compose(args, _requires_micro) +def requires_rpc(*args): + """Mark a test as requiring rpc to run. + + Parameters + ---------- + f : function + Function to mark + """ + _requires_rpc = [ + pytest.mark.skipif( + tvm.support.libinfo().get("USE_RPC", "OFF") != "ON", + reason="RPC support not enabled. Set USE_RPC=ON in config.cmake to enable.", + ) + ] + return _compose(args, _requires_rpc) + + def _target_to_requirement(target): # mapping from target to decorator if target.startswith("cuda"): diff --git a/src/auto_scheduler/measure_record.cc b/src/auto_scheduler/measure_record.cc old mode 100755 new mode 100644 index 66f521e17e80..60bf0113f9a6 --- a/src/auto_scheduler/measure_record.cc +++ b/src/auto_scheduler/measure_record.cc @@ -107,6 +107,54 @@ struct Handler<::tvm::auto_scheduler::StateNode> { } }; +template <> +struct Handler<::tvm::auto_scheduler::HardwareParamsNode> { + inline static void Write(dmlc::JSONWriter* writer, + const ::tvm::auto_scheduler::HardwareParamsNode& data) { + writer->BeginArray(false); + writer->WriteArrayItem(data.num_cores); + writer->WriteArrayItem(data.vector_unit_bytes); + writer->WriteArrayItem(data.cache_line_bytes); + writer->WriteArrayItem(data.max_shared_memory_per_block); + writer->WriteArrayItem(data.max_registers_per_block); + writer->WriteArrayItem(data.max_threads_per_block); + writer->WriteArrayItem(data.max_vthread_extent); + writer->WriteArrayItem(data.warp_size); + writer->EndArray(); + } + inline static void Read(dmlc::JSONReader* reader, + ::tvm::auto_scheduler::HardwareParamsNode* data) { + bool s; + reader->BeginArray(); + s = reader->NextArrayItem(); + CHECK(s); + reader->Read(&data->num_cores); + s = reader->NextArrayItem(); + CHECK(s); + reader->Read(&data->vector_unit_bytes); + s = reader->NextArrayItem(); + CHECK(s); + reader->Read(&data->cache_line_bytes); + s = reader->NextArrayItem(); + CHECK(s); + reader->Read(&data->max_shared_memory_per_block); + s = reader->NextArrayItem(); + CHECK(s); + reader->Read(&data->max_registers_per_block); + s = reader->NextArrayItem(); + CHECK(s); + reader->Read(&data->max_threads_per_block); + s = reader->NextArrayItem(); + CHECK(s); + reader->Read(&data->max_vthread_extent); + s = reader->NextArrayItem(); + CHECK(s); + reader->Read(&data->warp_size); + s = reader->NextArrayItem(); + CHECK(!s); + } +}; + template <> struct Handler<::tvm::auto_scheduler::SearchTaskNode> { inline static void Write(dmlc::JSONWriter* writer, @@ -114,11 +162,13 @@ struct Handler<::tvm::auto_scheduler::SearchTaskNode> { writer->BeginArray(false); writer->WriteArrayItem(std::string(data.workload_key)); writer->WriteArrayItem(data.target->str()); + writer->WriteArrayItem(*data.hardware_params.get()); writer->EndArray(); } inline static void Read(dmlc::JSONReader* reader, ::tvm::auto_scheduler::SearchTaskNode* data) { bool s; std::string str_value; + auto hardware_params_node = ::tvm::make_object<::tvm::auto_scheduler::HardwareParamsNode>(); reader->BeginArray(); s = reader->NextArrayItem(); ICHECK(s); @@ -129,7 +179,12 @@ struct Handler<::tvm::auto_scheduler::SearchTaskNode> { reader->Read(&str_value); data->target = ::tvm::Target(str_value); s = reader->NextArrayItem(); - ICHECK(!s); + if (s) { + reader->Read(hardware_params_node.get()); + s = reader->NextArrayItem(); + data->hardware_params = ::tvm::auto_scheduler::HardwareParams(hardware_params_node); + ICHECK(!s); + } } }; @@ -340,5 +395,21 @@ TVM_REGISTER_GLOBAL("auto_scheduler.SaveRecords") std::ofstream ofs(filename, std::ofstream::app); WriteMeasureRecords(&ofs, in, res); }); + +TVM_REGISTER_GLOBAL("auto_scheduler.SerializeMeasureInput") + .set_body_typed([](const MeasureInput& input) { + std::ostringstream os; + dmlc::JSONWriter writer(&os); + writer.Write(*input.get()); + return os.str(); + }); + +TVM_REGISTER_GLOBAL("auto_scheduler.DeserializeMeasureInput").set_body_typed([](String json) { + std::istringstream ss(json); + dmlc::JSONReader reader(&ss); + auto inp = make_object(); + reader.Read(inp.get()); + return ObjectRef(inp); +}); } // namespace auto_scheduler } // namespace tvm diff --git a/tests/python/unittest/test_auto_scheduler_measure.py b/tests/python/unittest/test_auto_scheduler_measure.py index 4369d203b476..f38fb4d74351 100644 --- a/tests/python/unittest/test_auto_scheduler_measure.py +++ b/tests/python/unittest/test_auto_scheduler_measure.py @@ -182,12 +182,10 @@ def test_recover_measure_input(): raw_inp = inputs[0] - correct_inp = auto_scheduler.measure_record.recover_measure_input(raw_inp) + correct_inp = auto_scheduler.measure.recover_measure_input(raw_inp) assert str(correct_inp.task.compute_dag) == str(inp.task.compute_dag) - correct_inp = auto_scheduler.measure_record.recover_measure_input( - raw_inp, rebuild_state=True - ) + correct_inp = auto_scheduler.measure.recover_measure_input(raw_inp, rebuild_state=True) assert str(correct_inp.state) == str(inp.state) diff --git a/tests/python/unittest/test_auto_scheduler_search_policy.py b/tests/python/unittest/test_auto_scheduler_search_policy.py index 07cf4c8141a0..2b5a69a884af 100644 --- a/tests/python/unittest/test_auto_scheduler_search_policy.py +++ b/tests/python/unittest/test_auto_scheduler_search_policy.py @@ -26,6 +26,7 @@ from tvm import auto_scheduler from test_auto_scheduler_common import matmul_auto_scheduler_test, PropagatingThread +import multiprocessing def search_common( @@ -156,9 +157,8 @@ def test_sketch_search_policy_cuda_rpc_runner(): t.join() +@tvm.testing.requires_cuda def test_sketch_search_policy_cuda_xgbmodel_rpc_runner(): - if not tvm.runtime.enabled("cuda"): - return measure_ctx = auto_scheduler.LocalRPCMeasureContext() # wrap the search in a new thread to avoid the conflict # between python's multiprocessing and tvm's thread pool diff --git a/tests/python/unittest/test_autotvm_dispatch_context.py b/tests/python/unittest/test_autotvm_dispatch_context.py index 4064ede3cc06..6ca062047fd7 100644 --- a/tests/python/unittest/test_autotvm_dispatch_context.py +++ b/tests/python/unittest/test_autotvm_dispatch_context.py @@ -21,12 +21,13 @@ from tvm import autotvm -def test_fallback(): - @autotvm.template("testing/dispatch_fallback") - def simple_template(a, b): - cfg = autotvm.get_config() - assert cfg.is_fallback +@autotvm.template("testing/dispatch_fallback") +def simple_template(a, b): + cfg = autotvm.get_config() + assert cfg.is_fallback + +def test_fallback(): simple_template(2, 3) diff --git a/tests/python/unittest/test_runtime_rpc.py b/tests/python/unittest/test_runtime_rpc.py index d25eff23ae76..2739f5637967 100644 --- a/tests/python/unittest/test_runtime_rpc.py +++ b/tests/python/unittest/test_runtime_rpc.py @@ -30,6 +30,7 @@ from tvm.rpc.tracker import Tracker +@tvm.testing.requires_rpc def test_bigendian_rpc(): """Test big endian rpc when there is a PowerPC RPC server available""" host = os.environ.get("TVM_POWERPC_TEST_HOST", None) @@ -61,22 +62,23 @@ def verify_rpc(remote, target, shape, dtype): verify_rpc(remote, target, (10,), dtype) -def test_rpc_simple(): - if not tvm.runtime.enabled("rpc"): - return +@tvm.register_func("rpc.test.addone") +def addone(x): + return x + 1 + + +@tvm.register_func("rpc.test.strcat") +def strcat(name, x): + return "%s:%d" % (name, x) - @tvm.register_func("rpc.test.addone") - def addone(x): - return x + 1 - @tvm.register_func("rpc.test.strcat") - def strcat(name, x): - return "%s:%d" % (name, x) +@tvm.register_func("rpc.test.except") +def remotethrow(name): + raise ValueError("%s" % name) - @tvm.register_func("rpc.test.except") - def remotethrow(name): - raise ValueError("%s" % name) +@tvm.testing.requires_rpc +def test_rpc_simple(): server = rpc.Server("localhost", key="x1") client = rpc.connect(server.host, server.port, key="x1") f1 = client.get_function("rpc.test.addone") @@ -90,14 +92,13 @@ def remotethrow(name): assert f2("abc", 11) == "abc:11" -def test_rpc_runtime_string(): - if not tvm.runtime.enabled("rpc"): - return +@tvm.register_func("rpc.test.runtime_str_concat") +def strcat(x, y): + return x + y - @tvm.register_func("rpc.test.runtime_str_concat") - def strcat(x, y): - return x + y +@tvm.testing.requires_rpc +def test_rpc_runtime_string(): server = rpc.Server("localhost", key="x1") client = rpc.connect(server.host, server.port, key="x1") func = client.get_function("rpc.test.runtime_str_concat") @@ -106,14 +107,15 @@ def strcat(x, y): assert str(func(x, y)) == "abcdef" -def test_rpc_array(): - if not tvm.runtime.enabled("rpc"): - return - x = np.random.randint(0, 10, size=(3, 4)) +@tvm.register_func("rpc.test.remote_array_func") +def remote_array_func(y): + x = np.ones((3, 4)) + np.testing.assert_equal(y.asnumpy(), x) - @tvm.register_func("rpc.test.remote_array_func") - def remote_array_func(y): - np.testing.assert_equal(y.asnumpy(), x) + +@tvm.testing.requires_rpc +def test_rpc_array(): + x = np.ones((3, 4)) server = rpc.Server("localhost") remote = rpc.connect(server.host, server.port) @@ -124,6 +126,7 @@ def remote_array_func(y): fremote(r_cpu) +@tvm.testing.requires_rpc def test_rpc_large_array(): # testcase of large array creation server = rpc.Server("localhost") @@ -137,6 +140,7 @@ def test_rpc_large_array(): np.testing.assert_equal(b.asnumpy(), b_np) +@tvm.testing.requires_rpc def test_rpc_echo(): def check(remote): fecho = remote.get_function("testing.echo") @@ -180,9 +184,8 @@ def check_minrpc(): check_minrpc() +@tvm.testing.requires_rpc def test_rpc_file_exchange(): - if not tvm.runtime.enabled("rpc"): - return server = rpc.Server("localhost") remote = rpc.connect(server.host, server.port) blob = bytearray(np.random.randint(0, 10, size=(10))) @@ -191,10 +194,9 @@ def test_rpc_file_exchange(): assert rev == blob +@tvm.testing.requires_rpc @tvm.testing.requires_llvm def test_rpc_remote_module(): - if not tvm.runtime.enabled("rpc"): - return # graph n = tvm.runtime.convert(102) A = te.placeholder((n,), name="A") @@ -317,11 +319,13 @@ def check_remote_link_cl(remote): check_minrpc() -def test_rpc_return_func(): - @tvm.register_func("rpc.test.remote_func") - def addone(x): - return lambda y: x + y +@tvm.register_func("rpc.test.remote_func") +def addone(x): + return lambda y: x + y + +@tvm.testing.requires_rpc +def test_rpc_return_func(): server = rpc.Server("localhost", key="x1") client = rpc.connect(server.host, server.port, key="x1") f1 = client.get_function("rpc.test.remote_func") @@ -329,6 +333,7 @@ def addone(x): assert fadd(12) == 22 +@tvm.testing.requires_rpc def test_rpc_session_constructor_args(): # start server server0 = rpc.Server("localhost", key="x0") @@ -365,21 +370,23 @@ def check_error_handling(): check_error_handling() -def test_rpc_return_ndarray(): +@tvm.register_func("rpc.test.remote_return_nd") +def my_module(name): # Use closure to check the ref counter correctness nd = tvm.nd.array(np.zeros(10).astype("float32")) - @tvm.register_func("rpc.test.remote_return_nd") - def my_module(name): - if name == "get_arr": - return lambda: nd - elif name == "ref_count": - return lambda: tvm.testing.object_use_count(nd) - elif name == "get_elem": - return lambda idx: nd.asnumpy()[idx] - elif name == "get_arr_elem": - return lambda arr, idx: arr.asnumpy()[idx] + if name == "get_arr": + return lambda: nd + elif name == "ref_count": + return lambda: tvm.testing.object_use_count(nd) + elif name == "get_elem": + return lambda idx: nd.asnumpy()[idx] + elif name == "get_arr_elem": + return lambda arr, idx: arr.asnumpy()[idx] + +@tvm.testing.requires_rpc +def test_rpc_return_ndarray(): # start server server = rpc.Server("localhost", key="x1") client = rpc.connect(server.host, server.port, key="x1") @@ -392,26 +399,19 @@ def my_module(name): # array test def run_arr_test(): arr = get_arr() - assert ref_count() == 2 - arr2 = get_arr() - assert ref_count() == 3 - assert arr.context == client.cpu(0) - arr.copyfrom(np.ones(10).astype(arr.dtype)) - assert arr2.asnumpy()[0] == 1.0 - assert get_elem(0) == 1.0 - assert get_arr_elem(arr2, 0) == 1.0 - - assert ref_count() == 1 + assert get_elem(0) == 0.0 + assert get_arr_elem(arr, 0) == 0.0 + run_arr_test() - # check recycle correctness - assert ref_count() == 1 -def test_local_func(): - @tvm.register_func("rpc.test.remote_func2") - def addone(x): - return lambda y: x + y +@tvm.register_func("rpc.test.remote_func2") +def addone(x): + return lambda y: x + y + +@tvm.testing.requires_rpc +def test_local_func(): client = rpc.LocalSession() f1 = client.get_function("rpc.test.remote_func2") fadd = f1(10) @@ -423,6 +423,7 @@ def addone(x): assert rev == blob +@tvm.testing.requires_rpc def test_rpc_tracker_register(): # test registration tracker = Tracker("localhost", port=9000, port_end=10000) @@ -459,6 +460,15 @@ def test_rpc_tracker_register(): tracker.terminate() +def _target(host, port, device_key, timeout): + client = rpc.connect_tracker(host, port) + remote = client.request(device_key, session_timeout=timeout) + while True: + pass + remote.cpu() + + +@tvm.testing.requires_rpc def test_rpc_tracker_request(): # test concurrent request tracker = Tracker("localhost", port=9000, port_end=10000) @@ -472,16 +482,11 @@ def test_rpc_tracker_request(): ) client = rpc.connect_tracker(tracker.host, tracker.port) - def target(host, port, device_key, timeout): - client = rpc.connect_tracker(host, port) - remote = client.request(device_key, session_timeout=timeout) - while True: - pass - remote.cpu() - - proc1 = multiprocessing.Process(target=target, args=(tracker.host, tracker.port, device_key, 4)) + proc1 = multiprocessing.Process( + target=_target, args=(tracker.host, tracker.port, device_key, 4) + ) proc2 = multiprocessing.Process( - target=target, args=(tracker.host, tracker.port, device_key, 200) + target=_target, args=(tracker.host, tracker.port, device_key, 200) ) proc1.start() time.sleep(0.5) From 50f556202a2dacb09fe58465ddd2761c1c8ce25e Mon Sep 17 00:00:00 2001 From: Tristan Konolige Date: Mon, 26 Oct 2020 09:31:07 -0700 Subject: [PATCH 2/8] address reviewer feedback --- python/tvm/auto_scheduler/measure.py | 14 ++++++++++---- python/tvm/auto_scheduler/utils.py | 1 + python/tvm/auto_scheduler/workload_registry.py | 6 +++--- src/auto_scheduler/measure_record.cc | 2 +- 4 files changed, 15 insertions(+), 8 deletions(-) diff --git a/python/tvm/auto_scheduler/measure.py b/python/tvm/auto_scheduler/measure.py index 1f9c9bf2c157..9f592550cda8 100644 --- a/python/tvm/auto_scheduler/measure.py +++ b/python/tvm/auto_scheduler/measure.py @@ -57,7 +57,7 @@ ) from .compute_dag import ComputeDAG from .search_task import SearchTask -from .workload_registry import workload_name, get_workload +from .workload_registry import workload_func_name, get_workload_func # The maximum length of error message MAX_ERROR_MSG_LEN = 512 @@ -82,7 +82,6 @@ def recover_measure_input(inp, rebuild_state=False): The fully recovered MeasureInput with all fields rebuilt. """ task = inp.task - print(task.hardware_params) new_task = SearchTask( ComputeDAG(task.workload_key), task.workload_key, @@ -121,13 +120,20 @@ def __init__(self, task, state): self.__init_handle_by_constructor__(_ffi_api.MeasureInput, task, state) def serialize(self): + """Custom serialization to workaround MeasureInput not exposing all its + members to the TVM ffi interface. + + Note that we do not implement __getstate__ as it does not seem to work + with initialization of the workload registry (maybe because of + initialization order?). + """ serialize = tvm.get_global_func("auto_scheduler.SerializeMeasureInput", True) assert serialize # We serialize the workload function so that it can be used on the deserialized side. return { "measureinput": serialize(self), - "name": workload_name(self.task.workload_key), - "func": get_workload(self.task), + "name": workload_func_name(self.task.workload_key), + "func": get_workload_func(self.task), } @staticmethod diff --git a/python/tvm/auto_scheduler/utils.py b/python/tvm/auto_scheduler/utils.py index 49a97d0d6ced..2d0ec3efd75d 100644 --- a/python/tvm/auto_scheduler/utils.py +++ b/python/tvm/auto_scheduler/utils.py @@ -144,6 +144,7 @@ def kill_child_processes(parent_pid, sig=signal.SIGTERM): def _func_wrapper(que, func, args, kwargs): + """Call function and return the result over the queue.""" if kwargs: que.put(func(*args, **kwargs)) else: diff --git a/python/tvm/auto_scheduler/workload_registry.py b/python/tvm/auto_scheduler/workload_registry.py index 6bf183b9fd98..c2d7f90771e3 100644 --- a/python/tvm/auto_scheduler/workload_registry.py +++ b/python/tvm/auto_scheduler/workload_registry.py @@ -175,7 +175,7 @@ def workload_key_to_tensors(workload_key): return lookup(*args) -def get_workload(task): +def get_workload_func(task): """Get the workload function for a given task Parameters @@ -188,13 +188,13 @@ def get_workload(task): workload : callable The registered workload function. """ - name = workload_name(task.workload_key) + name = workload_func_name(task.workload_key) lookup = WORKLOAD_FUNC_REGISTRY[name] assert callable(lookup) return lookup -def workload_name(workload_key): +def workload_func_name(workload_key): """Decode a workload key to the registered function name. Parameters diff --git a/src/auto_scheduler/measure_record.cc b/src/auto_scheduler/measure_record.cc index 60bf0113f9a6..1bc2c78a99f0 100644 --- a/src/auto_scheduler/measure_record.cc +++ b/src/auto_scheduler/measure_record.cc @@ -271,7 +271,7 @@ namespace auto_scheduler { TVM_REGISTER_OBJECT_TYPE(RecordToFileNode); TVM_REGISTER_OBJECT_TYPE(RecordReaderNode); -const std::string AUTO_SCHEDULER_LOG_VERSION = "v0.2"; // NOLINT(*) +const std::string AUTO_SCHEDULER_LOG_VERSION = "v0.3"; // NOLINT(*) RecordToFile::RecordToFile(String filename) { auto node = make_object(); From 5a6fd49d0ed21884a0983cd10cce2383a0794ae1 Mon Sep 17 00:00:00 2001 From: Tristan Konolige Date: Tue, 27 Oct 2020 15:16:25 -0600 Subject: [PATCH 3/8] Fix tutorials --- .../auto_scheduler/tune_conv2d_layer_cuda.py | 276 ++++----- tutorials/auto_scheduler/tune_matmul_x86.py | 271 ++++----- tutorials/autotvm/tune_conv2d_cuda.py | 142 ++--- tutorials/autotvm/tune_relay_arm.py | 435 +++++++------- tutorials/autotvm/tune_relay_cuda.py | 533 +++++++++--------- tutorials/autotvm/tune_relay_mobile_gpu.py | 441 +++++++-------- tutorials/autotvm/tune_relay_x86.py | 333 +++++------ tutorials/autotvm/tune_simple_template.py | 232 ++++---- 8 files changed, 1336 insertions(+), 1327 deletions(-) diff --git a/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py b/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py index 10a2d1b44144..157bf9e76728 100644 --- a/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py +++ b/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py @@ -56,141 +56,141 @@ def conv2d_layer(N, H, W, CO, CI, KH, KW, stride, padding): out = topi.nn.relu(conv + bias) return [data, kernel, bias, out] - -###################################################################### -# Create the search task -# ^^^^^^^^^^^^^^^^^^^^^^ -# We then create a search task for the last convolution layer in the resnet. - -target = tvm.target.Target("cuda") - -# Use the last layer in ResNet-50 -N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) -task = auto_scheduler.create_task(conv2d_layer, (N, H, W, CO, CI, KH, KW, strides, padding), target) - -# Inspect the computational graph -print(task.compute_dag) - -###################################################################### -# Next, we set parameters for the auto-scheduler. These parameters -# mainly specify how we do the measurement during the search and auto-tuning. -# -# * :code:`measure_ctx` launches a different process for measurement. This -# provides an isolation. It can protect the master process from GPU crashes -# happended during measurement and avoid other runtime conflicts. -# * :code:`min_repeat_ms` defines the minimum duration of one "repeat" in every measurement. -# This can warmup the GPU, which is necessary to get accurate measurement results. -# Typically, we recommend a value > 300 ms. -# * :code:`num_measure_trials` is the number of measurement trials we can use during the search. -# We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a -# good value for the search to converge. You can do more trials according to your time budget. -# * In addition, we use :code:`RecordToFile` to dump measurement records into a file `conv2d.json`. -# The measurement records can be used to query the history best, resume the search, -# and do more analyses later. -# * see :any:`auto_scheduler.TuningOptions`, -# :any:`auto_scheduler.LocalRPCMeasureContext` for more parameters. - -log_file = "conv2d.json" -measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300) -tune_option = auto_scheduler.TuningOptions( - num_measure_trials=10, - runner=measure_ctx.runner, - measure_callbacks=[auto_scheduler.RecordToFile(log_file)], -) - -###################################################################### -# Run the search -# ^^^^^^^^^^^^^^ -# Now we get all inputs ready. Pretty simple, isn't it? -# We can kick off the search and let the auto-scheduler do its magic. -# After some measurement trials, it will return the best schedule it found. - -sch, args = auto_scheduler.auto_schedule(task, tuning_options=tune_option) - -# Kill the process for measurement -del measure_ctx - -###################################################################### -# We can lower the schedule to see the IR after auto-scheduling. -# The auto-scheduler correctly performs optimizations including multi-level tiling, -# cooperative fetching, unrolling and operator fusion. - -print(tvm.lower(sch, args, simple_mode=True)) - -###################################################################### -# Check correctness and evaluate performance -# ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ -# We build the binary and check its correctness and performance. - -func = tvm.build(sch, args, target) - -# Check correctness -data_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) -weight_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) -bias_np = np.random.uniform(size=(1, CO, 1, 1)).astype(np.float32) -conv_np = conv2d_nchw_python(data_np, weight_np, strides, padding) -out_np = np.maximum(conv_np + bias_np, 0.0) - -ctx = tvm.gpu() -data_tvm = tvm.nd.array(data_np, ctx=ctx) -weight_tvm = tvm.nd.array(weight_np, ctx=ctx) -bias_tvm = tvm.nd.array(bias_np, ctx=ctx) -out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx) -func(data_tvm, weight_tvm, bias_tvm, out_tvm) - -# Check results -np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3) - -# Evaluate execution time -evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500) -print( - "Execution time of this operator: %.3f ms" - % (np.median(evaluator(data_tvm, weight_tvm, bias_tvm, out_tvm).results) * 1000) -) - -###################################################################### -# Using the record file -# ^^^^^^^^^^^^^^^^^^^^^ -# During the search, all measuremnt records are dumpped into the record -# file "conv2d.json". The measurement records can be used to re-apply search results, -# resume the search, and perform other analyses. - -###################################################################### -# Here is an example where we load the best schedule from a file, -# print the equivalent python schedule API, and build the binary again. - -# Load the measuremnt record for the best schedule -inp, res = auto_scheduler.load_best(log_file, task.workload_key) - -# Print equivalent python schedule API. This can be used for debugging and -# learning the behavior of the auto-scheduler. -print("Equivalent python schedule:") -print(task.compute_dag.print_python_code_from_state(inp.state)) - -# Rebuild the binary. This shows how you can apply the best schedule from a -# log file without reruning the search again. -sch, args = task.compute_dag.apply_steps_from_state(inp.state) -func = tvm.build(sch, args, target) - -###################################################################### -# A more complicated example is to resume the search. -# In this case, we need to create the search policy and cost model by ourselves -# and resume the status of search policy and cost model with the log file. -# In the example below we resume the status and do more 5 trials. - - -cost_model = auto_scheduler.XGBModel() -cost_model.update_from_file(log_file) -search_policy = auto_scheduler.SketchPolicy( - task, cost_model, init_search_callbacks=[auto_scheduler.PreloadMeasuredStates(log_file)] -) -measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300) -tune_option = auto_scheduler.TuningOptions( - num_measure_trials=5, - runner=measure_ctx.runner, - measure_callbacks=[auto_scheduler.RecordToFile(log_file)], -) -sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) - -# Kill the measurement process -del measure_ctx +if __name__ == "__main__": + ###################################################################### + # Create the search task + # ^^^^^^^^^^^^^^^^^^^^^^ + # We then create a search task for the last convolution layer in the resnet. + + target = tvm.target.Target("cuda") + + # Use the last layer in ResNet-50 + N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) + task = auto_scheduler.create_task(conv2d_layer, (N, H, W, CO, CI, KH, KW, strides, padding), target) + + # Inspect the computational graph + print(task.compute_dag) + + ###################################################################### + # Next, we set parameters for the auto-scheduler. These parameters + # mainly specify how we do the measurement during the search and auto-tuning. + # + # * :code:`measure_ctx` launches a different process for measurement. This + # provides an isolation. It can protect the master process from GPU crashes + # happended during measurement and avoid other runtime conflicts. + # * :code:`min_repeat_ms` defines the minimum duration of one "repeat" in every measurement. + # This can warmup the GPU, which is necessary to get accurate measurement results. + # Typically, we recommend a value > 300 ms. + # * :code:`num_measure_trials` is the number of measurement trials we can use during the search. + # We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a + # good value for the search to converge. You can do more trials according to your time budget. + # * In addition, we use :code:`RecordToFile` to dump measurement records into a file `conv2d.json`. + # The measurement records can be used to query the history best, resume the search, + # and do more analyses later. + # * see :any:`auto_scheduler.TuningOptions`, + # :any:`auto_scheduler.LocalRPCMeasureContext` for more parameters. + + log_file = "conv2d.json" + measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300) + tune_option = auto_scheduler.TuningOptions( + num_measure_trials=10, + runner=measure_ctx.runner, + measure_callbacks=[auto_scheduler.RecordToFile(log_file)], + ) + + ###################################################################### + # Run the search + # ^^^^^^^^^^^^^^ + # Now we get all inputs ready. Pretty simple, isn't it? + # We can kick off the search and let the auto-scheduler do its magic. + # After some measurement trials, it will return the best schedule it found. + + sch, args = auto_scheduler.auto_schedule(task, tuning_options=tune_option) + + # Kill the process for measurement + del measure_ctx + + ###################################################################### + # We can lower the schedule to see the IR after auto-scheduling. + # The auto-scheduler correctly performs optimizations including multi-level tiling, + # cooperative fetching, unrolling and operator fusion. + + print(tvm.lower(sch, args, simple_mode=True)) + + ###################################################################### + # Check correctness and evaluate performance + # ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + # We build the binary and check its correctness and performance. + + func = tvm.build(sch, args, target) + + # Check correctness + data_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) + weight_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) + bias_np = np.random.uniform(size=(1, CO, 1, 1)).astype(np.float32) + conv_np = conv2d_nchw_python(data_np, weight_np, strides, padding) + out_np = np.maximum(conv_np + bias_np, 0.0) + + ctx = tvm.gpu() + data_tvm = tvm.nd.array(data_np, ctx=ctx) + weight_tvm = tvm.nd.array(weight_np, ctx=ctx) + bias_tvm = tvm.nd.array(bias_np, ctx=ctx) + out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx) + func(data_tvm, weight_tvm, bias_tvm, out_tvm) + + # Check results + np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3) + + # Evaluate execution time + evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500) + print( + "Execution time of this operator: %.3f ms" + % (np.median(evaluator(data_tvm, weight_tvm, bias_tvm, out_tvm).results) * 1000) + ) + + ###################################################################### + # Using the record file + # ^^^^^^^^^^^^^^^^^^^^^ + # During the search, all measuremnt records are dumpped into the record + # file "conv2d.json". The measurement records can be used to re-apply search results, + # resume the search, and perform other analyses. + + ###################################################################### + # Here is an example where we load the best schedule from a file, + # print the equivalent python schedule API, and build the binary again. + + # Load the measuremnt record for the best schedule + inp, res = auto_scheduler.load_best(log_file, task.workload_key) + + # Print equivalent python schedule API. This can be used for debugging and + # learning the behavior of the auto-scheduler. + print("Equivalent python schedule:") + print(task.compute_dag.print_python_code_from_state(inp.state)) + + # Rebuild the binary. This shows how you can apply the best schedule from a + # log file without reruning the search again. + sch, args = task.compute_dag.apply_steps_from_state(inp.state) + func = tvm.build(sch, args, target) + + ###################################################################### + # A more complicated example is to resume the search. + # In this case, we need to create the search policy and cost model by ourselves + # and resume the status of search policy and cost model with the log file. + # In the example below we resume the status and do more 5 trials. + + + cost_model = auto_scheduler.XGBModel() + cost_model.update_from_file(log_file) + search_policy = auto_scheduler.SketchPolicy( + task, cost_model, init_search_callbacks=[auto_scheduler.PreloadMeasuredStates(log_file)] + ) + measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300) + tune_option = auto_scheduler.TuningOptions( + num_measure_trials=5, + runner=measure_ctx.runner, + measure_callbacks=[auto_scheduler.RecordToFile(log_file)], + ) + sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) + + # Kill the measurement process + del measure_ctx diff --git a/tutorials/auto_scheduler/tune_matmul_x86.py b/tutorials/auto_scheduler/tune_matmul_x86.py index 81f2e71ff8f7..327b2805125d 100644 --- a/tutorials/auto_scheduler/tune_matmul_x86.py +++ b/tutorials/auto_scheduler/tune_matmul_x86.py @@ -56,143 +56,144 @@ def matmul_add(N, L, M, dtype): return [A, B, C, out] -###################################################################### -# Create the search task -# ^^^^^^^^^^^^^^^^^^^^^^ -# We then create a search task with N=L=M=128 and dtype="float32" -# If your machine supports avx instructions, you can -# -# - replace "llvm" below with "llvm -mcpu=core-avx2" to enable AVX2 -# - replace "llvm" below with "llvm -mcpu=skylake-avx512" to enable AVX-512 - -target = tvm.target.Target("llvm") -task = tvm.auto_scheduler.create_task(matmul_add, (128, 128, 128, "float32"), target) - -# Inspect the computational graph -print(task.compute_dag) - -###################################################################### -# Next, we set parameters for the auto-scheduler. -# -# * :code:`num_measure_trials` is the number of measurement trials we can use during the search. -# We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a -# good value for the search to converge. You can do more trials according to your time budget. -# * In addition, we use :code:`RecordToFile` to dump measurement records into a file `matmul.json`. -# The measurement records can be used to query the history best, resume the search, -# and do more analyses later. -# * see :any:`auto_scheduler.TuningOptions` for more parameters - -log_file = "matmul.json" -tune_option = auto_scheduler.TuningOptions( - num_measure_trials=10, measure_callbacks=[auto_scheduler.RecordToFile(log_file)] -) - -###################################################################### -# Run the search -# ^^^^^^^^^^^^^^ -# Now we get all inputs ready. Pretty simple, isn't it? -# We can kick off the search and let the auto-scheduler do its magic. -# After some measurement trials, it will return the best schedule it found. - -sch, args = auto_scheduler.auto_schedule(task, tuning_options=tune_option) - -###################################################################### -# We can lower the schedule to see the IR after auto-scheduling. -# The auto-scheduler correctly performs optimizations including multi-level tiling, -# parallelization, vectorization, unrolling and operator fusion. - -print(tvm.lower(sch, args, simple_mode=True)) - -###################################################################### -# Check correctness and evaluate performance -# ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ -# We build the binary and check its correctness and performance. - -func = tvm.build(sch, args) -a_np = np.random.uniform(size=(128, 128)).astype(np.float32) -b_np = np.random.uniform(size=(128, 128)).astype(np.float32) -c_np = np.random.uniform(size=(128, 128)).astype(np.float32) -out_np = a_np.dot(b_np) + c_np - -ctx = tvm.cpu() -a_tvm = tvm.nd.array(a_np, ctx=ctx) -b_tvm = tvm.nd.array(b_np, ctx=ctx) -c_tvm = tvm.nd.array(c_np, ctx=ctx) -out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx) -func(a_tvm, b_tvm, c_tvm, out_tvm) - -# Check results -np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3) - -# Evaluate execution time. -evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500) -print( - "Execution time of this operator: %.3f ms" - % (np.median(evaluator(a_tvm, b_tvm, c_tvm, out_tvm).results) * 1000) -) - - -###################################################################### -# Using the record file -# ^^^^^^^^^^^^^^^^^^^^^ -# During the search, all measuremnt records are dumpped into the record -# file "matmul.json". The measurement records can be used to re-apply search results, -# resume the search, and perform other analyses. - -###################################################################### -# Here is an example where we load the best schedule from a file, -# print the equivalent python schedule API, and build the binary again. - -# Load the measuremnt record for the best schedule -inp, res = auto_scheduler.load_best(log_file, task.workload_key) - -# Print equivalent python schedule API. This can be used for debugging and -# learning the behavior of the auto-scheduler. -print("Equivalent python schedule:") -print(task.compute_dag.print_python_code_from_state(inp.state)) - -# Rebuild the binary. This shows how you can apply the best schedule from a -# log file without reruning the search again. -sch, args = task.compute_dag.apply_steps_from_state(inp.state) -func = tvm.build(sch, args) - -###################################################################### -# A more complicated example is to resume the search. -# In this case, we need to create the search policy and cost model by ourselves -# and resume the status of search policy and cost model with the log file. -# In the example below we resume the status and do more 5 trials. - - -def resume_search(task, log_file_name): - cost_model = auto_scheduler.XGBModel() - cost_model.update_from_file(log_file_name) - search_policy = auto_scheduler.SketchPolicy( - task, - cost_model, - init_search_callbacks=[auto_scheduler.PreloadMeasuredStates(log_file_name)], - ) +if __name__ == "__main__": + ###################################################################### + # Create the search task + # ^^^^^^^^^^^^^^^^^^^^^^ + # We then create a search task with N=L=M=128 and dtype="float32" + # If your machine supports avx instructions, you can + # + # - replace "llvm" below with "llvm -mcpu=core-avx2" to enable AVX2 + # - replace "llvm" below with "llvm -mcpu=skylake-avx512" to enable AVX-512 + + target = tvm.target.Target("llvm") + task = tvm.auto_scheduler.create_task(matmul_add, (128, 128, 128, "float32"), target) + + # Inspect the computational graph + print(task.compute_dag) + + ###################################################################### + # Next, we set parameters for the auto-scheduler. + # + # * :code:`num_measure_trials` is the number of measurement trials we can use during the search. + # We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a + # good value for the search to converge. You can do more trials according to your time budget. + # * In addition, we use :code:`RecordToFile` to dump measurement records into a file `matmul.json`. + # The measurement records can be used to query the history best, resume the search, + # and do more analyses later. + # * see :any:`auto_scheduler.TuningOptions` for more parameters + + log_file = "matmul.json" tune_option = auto_scheduler.TuningOptions( - num_measure_trials=5, measure_callbacks=[auto_scheduler.RecordToFile(log_file_name)] + num_measure_trials=10, measure_callbacks=[auto_scheduler.RecordToFile(log_file)] ) - sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) + ###################################################################### + # Run the search + # ^^^^^^^^^^^^^^ + # Now we get all inputs ready. Pretty simple, isn't it? + # We can kick off the search and let the auto-scheduler do its magic. + # After some measurement trials, it will return the best schedule it found. + + sch, args = auto_scheduler.auto_schedule(task, tuning_options=tune_option) + + ###################################################################### + # We can lower the schedule to see the IR after auto-scheduling. + # The auto-scheduler correctly performs optimizations including multi-level tiling, + # parallelization, vectorization, unrolling and operator fusion. + + print(tvm.lower(sch, args, simple_mode=True)) + + ###################################################################### + # Check correctness and evaluate performance + # ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + # We build the binary and check its correctness and performance. + + func = tvm.build(sch, args) + a_np = np.random.uniform(size=(128, 128)).astype(np.float32) + b_np = np.random.uniform(size=(128, 128)).astype(np.float32) + c_np = np.random.uniform(size=(128, 128)).astype(np.float32) + out_np = a_np.dot(b_np) + c_np + + ctx = tvm.cpu() + a_tvm = tvm.nd.array(a_np, ctx=ctx) + b_tvm = tvm.nd.array(b_np, ctx=ctx) + c_tvm = tvm.nd.array(c_np, ctx=ctx) + out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx) + func(a_tvm, b_tvm, c_tvm, out_tvm) + + # Check results + np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3) + + # Evaluate execution time. + evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500) + print( + "Execution time of this operator: %.3f ms" + % (np.median(evaluator(a_tvm, b_tvm, c_tvm, out_tvm).results) * 1000) + ) -# resume_search(task, log_file) -###################################################################### -# .. note:: -# We cannot run the line above because of the conflict between -# python's multiprocessing and tvm's thread pool. -# After running a tvm generated binary the python's multiprocessing library -# will hang forever. You have to make sure that you don't run any tvm -# generated binaries before calling auot-scheduler's search. -# To run the function above, you should comment out all code in -# "Check correctness and evaluate performance" section. -# -# You should be careful about this problem in your applications. -# There are other workarounds for this problem. -# For example, you can start a new thread/process (with the builtin python library -# threading or multiprocessing) and run the tvm binaries in the new thread/process. -# This provides an isolation and avoids the conflict in the main thread/process. -# You can also use :any:`auto_scheduler.LocalRPCMeasureContext` for auto-scheduler, -# as shown in the GPU tutorial (:ref:`auto-scheduler-conv-gpu`). + ###################################################################### + # Using the record file + # ^^^^^^^^^^^^^^^^^^^^^ + # During the search, all measuremnt records are dumpped into the record + # file "matmul.json". The measurement records can be used to re-apply search results, + # resume the search, and perform other analyses. + + ###################################################################### + # Here is an example where we load the best schedule from a file, + # print the equivalent python schedule API, and build the binary again. + + # Load the measuremnt record for the best schedule + inp, res = auto_scheduler.load_best(log_file, task.workload_key) + + # Print equivalent python schedule API. This can be used for debugging and + # learning the behavior of the auto-scheduler. + print("Equivalent python schedule:") + print(task.compute_dag.print_python_code_from_state(inp.state)) + + # Rebuild the binary. This shows how you can apply the best schedule from a + # log file without reruning the search again. + sch, args = task.compute_dag.apply_steps_from_state(inp.state) + func = tvm.build(sch, args) + + ###################################################################### + # A more complicated example is to resume the search. + # In this case, we need to create the search policy and cost model by ourselves + # and resume the status of search policy and cost model with the log file. + # In the example below we resume the status and do more 5 trials. + + + def resume_search(task, log_file_name): + cost_model = auto_scheduler.XGBModel() + cost_model.update_from_file(log_file_name) + search_policy = auto_scheduler.SketchPolicy( + task, + cost_model, + init_search_callbacks=[auto_scheduler.PreloadMeasuredStates(log_file_name)], + ) + tune_option = auto_scheduler.TuningOptions( + num_measure_trials=5, measure_callbacks=[auto_scheduler.RecordToFile(log_file_name)] + ) + sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) + + + # resume_search(task, log_file) + + ###################################################################### + # .. note:: + # We cannot run the line above because of the conflict between + # python's multiprocessing and tvm's thread pool. + # After running a tvm generated binary the python's multiprocessing library + # will hang forever. You have to make sure that you don't run any tvm + # generated binaries before calling auot-scheduler's search. + # To run the function above, you should comment out all code in + # "Check correctness and evaluate performance" section. + # + # You should be careful about this problem in your applications. + # There are other workarounds for this problem. + # For example, you can start a new thread/process (with the builtin python library + # threading or multiprocessing) and run the tvm binaries in the new thread/process. + # This provides an isolation and avoids the conflict in the main thread/process. + # You can also use :any:`auto_scheduler.LocalRPCMeasureContext` for auto-scheduler, + # as shown in the GPU tutorial (:ref:`auto-scheduler-conv-gpu`). diff --git a/tutorials/autotvm/tune_conv2d_cuda.py b/tutorials/autotvm/tune_conv2d_cuda.py index ce9c19860ff4..5aa3a34d4092 100644 --- a/tutorials/autotvm/tune_conv2d_cuda.py +++ b/tutorials/autotvm/tune_conv2d_cuda.py @@ -50,6 +50,7 @@ import tvm from tvm import te +import tvm.testing from tvm import topi from tvm.topi.testing import conv2d_nchw_python @@ -168,73 +169,74 @@ def conv2d_no_batching(N, H, W, CO, CI, KH, KW, stride, padding): return s, [raw_data, kernel, conv] -###################################################################### -# Step 2: Search through the space -# --------------------------------- -# We pick the last layer on resnet as test case. -# Since our space is very large, :code:`XGBoostTuner` is most suitable -# for our case. Here we only do 20 trials for demonstration. -# In practice, making 1000 trials usually can find some good kernels -# for this template - -# logging config (for printing tuning log to screen) -logging.getLogger("autotvm").setLevel(logging.DEBUG) -logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) - -# the last layer in resnet -N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) -task = autotvm.task.create( - "tutorial/conv2d_no_batching", args=(N, H, W, CO, CI, KH, KW, strides, padding), target="cuda" -) -print(task.config_space) - -# Use local gpu, measure 10 times for every config to reduce variance -# The timeout of compiling a program is 10 seconds, the timeout for running is 4 seconds -measure_option = autotvm.measure_option( - builder=autotvm.LocalBuilder(), - runner=autotvm.LocalRunner(repeat=3, min_repeat_ms=100, timeout=4), -) - -# Begin tuning, log records to file `conv2d.log` -# During tuning we will also try many invalid configs, so you are expected to -# see many error reports. As long as you can see non-zero GFLOPS, it is okay. -tuner = autotvm.tuner.XGBTuner(task) -tuner.tune( - n_trial=20, - measure_option=measure_option, - callbacks=[autotvm.callback.log_to_file("conv2d.log")], -) - -######################################################################### -# Finally we can inspect the best config from log file, check correctness, -# and measure running time. - -# inspect the best config -dispatch_context = autotvm.apply_history_best("conv2d.log") -best_config = dispatch_context.query(task.target, task.workload) -print("\nBest config:") -print(best_config) - -# apply history best from log file -with autotvm.apply_history_best("conv2d.log"): - with tvm.target.Target("cuda"): - s, arg_bufs = conv2d_no_batching(N, H, W, CO, CI, KH, KW, strides, padding) - func = tvm.build(s, arg_bufs) - -# check correctness -a_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) -w_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) -c_np = conv2d_nchw_python(a_np, w_np, strides, padding) - -ctx = tvm.gpu() -a_tvm = tvm.nd.array(a_np, ctx=ctx) -w_tvm = tvm.nd.array(w_np, ctx=ctx) -c_tvm = tvm.nd.empty(c_np.shape, ctx=ctx) -func(a_tvm, w_tvm, c_tvm) - -tvm.testing.assert_allclose(c_np, c_tvm.asnumpy(), rtol=1e-2) - -# Evaluate running time. Here we choose a large repeat number (400) to reduce the noise -# and the overhead of kernel launch. You can also use nvprof to validate the result. -evaluator = func.time_evaluator(func.entry_name, ctx, number=400) -print("Time cost of this operator: %f" % evaluator(a_tvm, w_tvm, c_tvm).mean) +if __name__ == "__main__": + ###################################################################### + # Step 2: Search through the space + # --------------------------------- + # We pick the last layer on resnet as test case. + # Since our space is very large, :code:`XGBoostTuner` is most suitable + # for our case. Here we only do 20 trials for demonstration. + # In practice, making 1000 trials usually can find some good kernels + # for this template + + # logging config (for printing tuning log to screen) + logging.getLogger("autotvm").setLevel(logging.DEBUG) + logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) + + # the last layer in resnet + N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) + task = autotvm.task.create( + "tutorial/conv2d_no_batching", args=(N, H, W, CO, CI, KH, KW, strides, padding), target="cuda" + ) + print(task.config_space) + + # Use local gpu, measure 10 times for every config to reduce variance + # The timeout of compiling a program is 10 seconds, the timeout for running is 4 seconds + measure_option = autotvm.measure_option( + builder=autotvm.LocalBuilder(), + runner=autotvm.LocalRunner(repeat=3, min_repeat_ms=100, timeout=4), + ) + + # Begin tuning, log records to file `conv2d.log` + # During tuning we will also try many invalid configs, so you are expected to + # see many error reports. As long as you can see non-zero GFLOPS, it is okay. + tuner = autotvm.tuner.XGBTuner(task) + tuner.tune( + n_trial=20, + measure_option=measure_option, + callbacks=[autotvm.callback.log_to_file("conv2d.log")], + ) + + ######################################################################### + # Finally we can inspect the best config from log file, check correctness, + # and measure running time. + + # inspect the best config + dispatch_context = autotvm.apply_history_best("conv2d.log") + best_config = dispatch_context.query(task.target, task.workload) + print("\nBest config:") + print(best_config) + + # apply history best from log file + with autotvm.apply_history_best("conv2d.log"): + with tvm.target.Target("cuda"): + s, arg_bufs = conv2d_no_batching(N, H, W, CO, CI, KH, KW, strides, padding) + func = tvm.build(s, arg_bufs) + + # check correctness + a_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) + w_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) + c_np = conv2d_nchw_python(a_np, w_np, strides, padding) + + ctx = tvm.gpu() + a_tvm = tvm.nd.array(a_np, ctx=ctx) + w_tvm = tvm.nd.array(w_np, ctx=ctx) + c_tvm = tvm.nd.empty(c_np.shape, ctx=ctx) + func(a_tvm, w_tvm, c_tvm) + + tvm.testing.assert_allclose(c_np, c_tvm.asnumpy(), rtol=1e-2) + + # Evaluate running time. Here we choose a large repeat number (400) to reduce the noise + # and the overhead of kernel launch. You can also use nvprof to validate the result. + evaluator = func.time_evaluator(func.entry_name, ctx, number=400) + print("Time cost of this operator: %f" % evaluator(a_tvm, w_tvm, c_tvm).mean) diff --git a/tutorials/autotvm/tune_relay_arm.py b/tutorials/autotvm/tune_relay_arm.py index f024ba4f201a..97950d11b669 100644 --- a/tutorials/autotvm/tune_relay_arm.py +++ b/tutorials/autotvm/tune_relay_arm.py @@ -189,227 +189,228 @@ def get_network(name, batch_size): # # You can register multiple devices to the tracker to accelerate the measurement in tuning. -########################################### -# Set Tuning Options -# ------------------ -# Before tuning, we should apply some configurations. Here I use an RK3399 board -# as example. In your setting, you should modify the target and device_key accordingly. -# set :code:`use_android` to True if you use android phone. - -#### DEVICE CONFIG #### - -# Replace "aarch64-linux-gnu" with the correct target of your board. -# This target is used for cross compilation. You can query it by :code:`gcc -v` on your device. -target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu") - -# Also replace this with the device key in your tracker -device_key = "rk3399" - -# Set this to True if you use android phone -use_android = False - -#### TUNING OPTION #### -network = "resnet-18" -log_file = "%s.%s.log" % (device_key, network) -dtype = "float32" - -tuning_option = { - "log_filename": log_file, - "tuner": "xgb", - "n_trial": 1500, - "early_stopping": 800, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(build_func="ndk" if use_android else "default"), - runner=autotvm.RPCRunner( - device_key, - host="0.0.0.0", - port=9190, - number=5, - timeout=10, +if __name__ == "__main__": + ########################################### + # Set Tuning Options + # ------------------ + # Before tuning, we should apply some configurations. Here I use an RK3399 board + # as example. In your setting, you should modify the target and device_key accordingly. + # set :code:`use_android` to True if you use android phone. + + #### DEVICE CONFIG #### + + # Replace "aarch64-linux-gnu" with the correct target of your board. + # This target is used for cross compilation. You can query it by :code:`gcc -v` on your device. + target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu") + + # Also replace this with the device key in your tracker + device_key = "rk3399" + + # Set this to True if you use android phone + use_android = False + + #### TUNING OPTION #### + network = "resnet-18" + log_file = "%s.%s.log" % (device_key, network) + dtype = "float32" + + tuning_option = { + "log_filename": log_file, + "tuner": "xgb", + "n_trial": 1500, + "early_stopping": 800, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(build_func="ndk" if use_android else "default"), + runner=autotvm.RPCRunner( + device_key, + host="0.0.0.0", + port=9190, + number=5, + timeout=10, + ), ), - ), -} - -#################################################################### -# -# .. note:: How to set tuning options -# -# In general, the default values provided here work well. -# If you have enough time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, -# which makes the tuning run longer. -# If your device runs very slow or your conv2d operators have many GFLOPs, considering to -# set timeout larger. -# -# If your model has depthwise convolution, you could consider setting -# :code:`try_spatial_pack_depthwise` be :code:`True`, which perform better than default -# optimization in general. For example, on ARM CPU A53 2.0GHz, we find it could boost 1.6x -# performance of depthwise convolution on Mobilenet V1 model. - -################################################################### -# Begin Tuning -# ------------ -# Now we can extract tuning tasks from the network and begin tuning. -# Here, we provide a simple utility function to tune a list of tasks. -# This function is just an initial implementation which tunes them in sequential order. -# We will introduce a more sophisticated tuning scheduler in the future. - -# You can skip the implementation of this function for this tutorial. -def tune_tasks( - tasks, - measure_option, - tuner="xgb", - n_trial=1000, - early_stopping=None, - log_filename="tuning.log", - use_transfer_learning=True, -): - # create tmp log file - tmp_log_file = log_filename + ".tmp" - if os.path.exists(tmp_log_file): + } + + #################################################################### + # + # .. note:: How to set tuning options + # + # In general, the default values provided here work well. + # If you have enough time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, + # which makes the tuning run longer. + # If your device runs very slow or your conv2d operators have many GFLOPs, considering to + # set timeout larger. + # + # If your model has depthwise convolution, you could consider setting + # :code:`try_spatial_pack_depthwise` be :code:`True`, which perform better than default + # optimization in general. For example, on ARM CPU A53 2.0GHz, we find it could boost 1.6x + # performance of depthwise convolution on Mobilenet V1 model. + + ################################################################### + # Begin Tuning + # ------------ + # Now we can extract tuning tasks from the network and begin tuning. + # Here, we provide a simple utility function to tune a list of tasks. + # This function is just an initial implementation which tunes them in sequential order. + # We will introduce a more sophisticated tuning scheduler in the future. + + # You can skip the implementation of this function for this tutorial. + def tune_tasks( + tasks, + measure_option, + tuner="xgb", + n_trial=1000, + early_stopping=None, + log_filename="tuning.log", + use_transfer_learning=True, + ): + # create tmp log file + tmp_log_file = log_filename + ".tmp" + if os.path.exists(tmp_log_file): + os.remove(tmp_log_file) + + for i, tsk in enumerate(reversed(tasks)): + prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) + + # create tuner + if tuner == "xgb" or tuner == "xgb-rank": + tuner_obj = XGBTuner(tsk, loss_type="rank") + elif tuner == "xgb_knob": + tuner_obj = XGBTuner(tsk, loss_type="rank", feature_type="knob") + elif tuner == "ga": + tuner_obj = GATuner(tsk, pop_size=50) + elif tuner == "random": + tuner_obj = RandomTuner(tsk) + elif tuner == "gridsearch": + tuner_obj = GridSearchTuner(tsk) + else: + raise ValueError("Invalid tuner: " + tuner) + + if use_transfer_learning: + if os.path.isfile(tmp_log_file): + tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) + + # do tuning + tsk_trial = min(n_trial, len(tsk.config_space)) + tuner_obj.tune( + n_trial=tsk_trial, + early_stopping=early_stopping, + measure_option=measure_option, + callbacks=[ + autotvm.callback.progress_bar(tsk_trial, prefix=prefix), + autotvm.callback.log_to_file(tmp_log_file), + ], + ) + + # pick best records to a cache file + autotvm.record.pick_best(tmp_log_file, log_filename) os.remove(tmp_log_file) - for i, tsk in enumerate(reversed(tasks)): - prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) - - # create tuner - if tuner == "xgb" or tuner == "xgb-rank": - tuner_obj = XGBTuner(tsk, loss_type="rank") - elif tuner == "xgb_knob": - tuner_obj = XGBTuner(tsk, loss_type="rank", feature_type="knob") - elif tuner == "ga": - tuner_obj = GATuner(tsk, pop_size=50) - elif tuner == "random": - tuner_obj = RandomTuner(tsk) - elif tuner == "gridsearch": - tuner_obj = GridSearchTuner(tsk) - else: - raise ValueError("Invalid tuner: " + tuner) - - if use_transfer_learning: - if os.path.isfile(tmp_log_file): - tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) - - # do tuning - tsk_trial = min(n_trial, len(tsk.config_space)) - tuner_obj.tune( - n_trial=tsk_trial, - early_stopping=early_stopping, - measure_option=measure_option, - callbacks=[ - autotvm.callback.progress_bar(tsk_trial, prefix=prefix), - autotvm.callback.log_to_file(tmp_log_file), - ], - ) - - # pick best records to a cache file - autotvm.record.pick_best(tmp_log_file, log_filename) - os.remove(tmp_log_file) - - -######################################################################## -# Finally, we launch tuning jobs and evaluate the end-to-end performance. - - -def tune_and_evaluate(tuning_opt): - # extract workloads from relay program - print("Extract tasks...") - mod, params, input_shape, _ = get_network(network, batch_size=1) - tasks = autotvm.task.extract_from_program( - mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) - ) - - # run tuning tasks - print("Tuning...") - tune_tasks(tasks, **tuning_opt) - - # compile kernels with history best records - with autotvm.apply_history_best(log_file): - print("Compile...") - with tvm.transform.PassContext(opt_level=3): - lib = relay.build_module.build(mod, target=target, params=params) - - # export library - tmp = tempdir() - if use_android: - from tvm.contrib import ndk - - filename = "net.so" - lib.export_library(tmp.relpath(filename), ndk.create_shared) - else: - filename = "net.tar" - lib.export_library(tmp.relpath(filename)) - - # upload module to device - print("Upload...") - remote = autotvm.measure.request_remote(device_key, "0.0.0.0", 9190, timeout=10000) - remote.upload(tmp.relpath(filename)) - rlib = remote.load_module(filename) - - # upload parameters to device - ctx = remote.context(str(target), 0) - module = runtime.GraphModule(rlib["default"](ctx)) - data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) - module.set_input("data", data_tvm) - - # evaluate - print("Evaluate inference time cost...") - ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=10) - prof_res = np.array(ftimer().results) * 1000 # convert to millisecond - print( - "Mean inference time (std dev): %.2f ms (%.2f ms)" - % (np.mean(prof_res), np.std(prof_res)) - ) - -# We do not run the tuning in our webpage server since it takes too long. -# Uncomment the following line to run it by yourself. + ######################################################################## + # Finally, we launch tuning jobs and evaluate the end-to-end performance. -# tune_and_evaluate(tuning_option) -###################################################################### -# Sample Output -# ------------- -# The tuning needs to compile many programs and extract feature from them. -# So a high performance CPU is recommended. -# One sample output is listed below. -# It takes about 2 hours on a 32T AMD Ryzen Threadripper. -# -# .. code-block:: bash -# -# Extract tasks... -# Tuning... -# [Task 1/12] Current/Best: 22.37/ 52.19 GFLOPS | Progress: (544/1000) | 406.59 s Done. -# [Task 2/12] Current/Best: 6.51/ 18.77 GFLOPS | Progress: (608/1000) | 325.05 s Done. -# [Task 3/12] Current/Best: 4.67/ 24.87 GFLOPS | Progress: (480/1000) | 372.31 s Done. -# [Task 4/12] Current/Best: 11.35/ 46.83 GFLOPS | Progress: (736/1000) | 602.39 s Done. -# [Task 5/12] Current/Best: 1.01/ 19.80 GFLOPS | Progress: (448/1000) | 262.16 s Done. -# [Task 6/12] Current/Best: 2.47/ 23.76 GFLOPS | Progress: (672/1000) | 563.85 s Done. -# [Task 7/12] Current/Best: 14.57/ 33.97 GFLOPS | Progress: (544/1000) | 465.15 s Done. -# [Task 8/12] Current/Best: 1.13/ 17.65 GFLOPS | Progress: (576/1000) | 365.08 s Done. -# [Task 9/12] Current/Best: 14.45/ 22.66 GFLOPS | Progress: (928/1000) | 724.25 s Done. -# [Task 10/12] Current/Best: 3.22/ 15.36 GFLOPS | Progress: (864/1000) | 564.27 s Done. -# [Task 11/12] Current/Best: 11.03/ 32.23 GFLOPS | Progress: (736/1000) | 635.15 s Done. -# [Task 12/12] Current/Best: 8.00/ 21.65 GFLOPS | Progress: (1000/1000) | 1111.81 s Done. -# Compile... -# Upload... -# Evaluate inference time cost... -# Mean inference time (std dev): 162.59 ms (0.06 ms) + def tune_and_evaluate(tuning_opt): + # extract workloads from relay program + print("Extract tasks...") + mod, params, input_shape, _ = get_network(network, batch_size=1) + tasks = autotvm.task.extract_from_program( + mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) + ) -###################################################################### -# -# .. note:: **Experiencing Difficulties?** -# -# The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", -# then there must be something wrong. -# -# First, make sure you set the correct configuration of your device. -# Then, you can print debug information by adding these lines in the beginning -# of the script. It will print every measurement result, where you can find useful -# error messages. -# -# .. code-block:: python -# -# import logging -# logging.getLogger('autotvm').setLevel(logging.DEBUG) -# -# Finally, always feel free to ask our community for help on https://discuss.tvm.ai + # run tuning tasks + print("Tuning...") + tune_tasks(tasks, **tuning_opt) + + # compile kernels with history best records + with autotvm.apply_history_best(log_file): + print("Compile...") + with tvm.transform.PassContext(opt_level=3): + lib = relay.build_module.build(mod, target=target, params=params) + + # export library + tmp = tempdir() + if use_android: + from tvm.contrib import ndk + + filename = "net.so" + lib.export_library(tmp.relpath(filename), ndk.create_shared) + else: + filename = "net.tar" + lib.export_library(tmp.relpath(filename)) + + # upload module to device + print("Upload...") + remote = autotvm.measure.request_remote(device_key, "0.0.0.0", 9190, timeout=10000) + remote.upload(tmp.relpath(filename)) + rlib = remote.load_module(filename) + + # upload parameters to device + ctx = remote.context(str(target), 0) + module = runtime.GraphModule(rlib["default"](ctx)) + data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) + module.set_input("data", data_tvm) + + # evaluate + print("Evaluate inference time cost...") + ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=10) + prof_res = np.array(ftimer().results) * 1000 # convert to millisecond + print( + "Mean inference time (std dev): %.2f ms (%.2f ms)" + % (np.mean(prof_res), np.std(prof_res)) + ) + + + # We do not run the tuning in our webpage server since it takes too long. + # Uncomment the following line to run it by yourself. + + # tune_and_evaluate(tuning_option) + + ###################################################################### + # Sample Output + # ------------- + # The tuning needs to compile many programs and extract feature from them. + # So a high performance CPU is recommended. + # One sample output is listed below. + # It takes about 2 hours on a 32T AMD Ryzen Threadripper. + # + # .. code-block:: bash + # + # Extract tasks... + # Tuning... + # [Task 1/12] Current/Best: 22.37/ 52.19 GFLOPS | Progress: (544/1000) | 406.59 s Done. + # [Task 2/12] Current/Best: 6.51/ 18.77 GFLOPS | Progress: (608/1000) | 325.05 s Done. + # [Task 3/12] Current/Best: 4.67/ 24.87 GFLOPS | Progress: (480/1000) | 372.31 s Done. + # [Task 4/12] Current/Best: 11.35/ 46.83 GFLOPS | Progress: (736/1000) | 602.39 s Done. + # [Task 5/12] Current/Best: 1.01/ 19.80 GFLOPS | Progress: (448/1000) | 262.16 s Done. + # [Task 6/12] Current/Best: 2.47/ 23.76 GFLOPS | Progress: (672/1000) | 563.85 s Done. + # [Task 7/12] Current/Best: 14.57/ 33.97 GFLOPS | Progress: (544/1000) | 465.15 s Done. + # [Task 8/12] Current/Best: 1.13/ 17.65 GFLOPS | Progress: (576/1000) | 365.08 s Done. + # [Task 9/12] Current/Best: 14.45/ 22.66 GFLOPS | Progress: (928/1000) | 724.25 s Done. + # [Task 10/12] Current/Best: 3.22/ 15.36 GFLOPS | Progress: (864/1000) | 564.27 s Done. + # [Task 11/12] Current/Best: 11.03/ 32.23 GFLOPS | Progress: (736/1000) | 635.15 s Done. + # [Task 12/12] Current/Best: 8.00/ 21.65 GFLOPS | Progress: (1000/1000) | 1111.81 s Done. + # Compile... + # Upload... + # Evaluate inference time cost... + # Mean inference time (std dev): 162.59 ms (0.06 ms) + + ###################################################################### + # + # .. note:: **Experiencing Difficulties?** + # + # The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", + # then there must be something wrong. + # + # First, make sure you set the correct configuration of your device. + # Then, you can print debug information by adding these lines in the beginning + # of the script. It will print every measurement result, where you can find useful + # error messages. + # + # .. code-block:: python + # + # import logging + # logging.getLogger('autotvm').setLevel(logging.DEBUG) + # + # Finally, always feel free to ask our community for help on https://discuss.tvm.ai diff --git a/tutorials/autotvm/tune_relay_cuda.py b/tutorials/autotvm/tune_relay_cuda.py index 4636103a22e2..0bc270816094 100644 --- a/tutorials/autotvm/tune_relay_cuda.py +++ b/tutorials/autotvm/tune_relay_cuda.py @@ -117,276 +117,277 @@ def get_network(name, batch_size): return mod, params, input_shape, output_shape -########################################### -# Set Tuning Options -# ------------------ -# Before tuning, we apply some configurations. - -#### DEVICE CONFIG #### -target = tvm.target.cuda() - -#### TUNING OPTION #### -network = "resnet-18" -log_file = "%s.log" % network -dtype = "float32" - -tuning_option = { - "log_filename": log_file, - "tuner": "xgb", - "n_trial": 2000, - "early_stopping": 600, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(timeout=10), - runner=autotvm.LocalRunner(number=20, repeat=3, timeout=4, min_repeat_ms=150), - ), -} - -#################################################################### -# -# .. note:: How to set tuning options -# -# In general, the default value provided here works well. -# -# If you have large time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, -# which makes the tuning runs longer. -# -# If you have multiple devices, you can use all of them for measurement to -# accelerate the tuning process. (see the 'Scale up measurement` section below). -# - -################################################################### -# Begin Tuning -# ------------ -# Now we can extract tuning tasks from the network and begin tuning. -# Here, we provide a simple utility function to tune a list of tasks. -# This function is just an initial implementation which tunes them in sequential order. -# We will introduce a more sophisticated tuning scheduler in the future. - -# You can skip the implementation of this function for this tutorial. -def tune_tasks( - tasks, - measure_option, - tuner="xgb", - n_trial=1000, - early_stopping=None, - log_filename="tuning.log", - use_transfer_learning=True, -): - # create tmp log file - tmp_log_file = log_filename + ".tmp" - if os.path.exists(tmp_log_file): +if __name__ == "__main__": + ########################################### + # Set Tuning Options + # ------------------ + # Before tuning, we apply some configurations. + + #### DEVICE CONFIG #### + target = tvm.target.cuda() + + #### TUNING OPTION #### + network = "resnet-18" + log_file = "%s.log" % network + dtype = "float32" + + tuning_option = { + "log_filename": log_file, + "tuner": "xgb", + "n_trial": 2000, + "early_stopping": 600, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(timeout=10), + runner=autotvm.LocalRunner(number=20, repeat=3, timeout=4, min_repeat_ms=150), + ), + } + + #################################################################### + # + # .. note:: How to set tuning options + # + # In general, the default value provided here works well. + # + # If you have large time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, + # which makes the tuning runs longer. + # + # If you have multiple devices, you can use all of them for measurement to + # accelerate the tuning process. (see the 'Scale up measurement` section below). + # + + ################################################################### + # Begin Tuning + # ------------ + # Now we can extract tuning tasks from the network and begin tuning. + # Here, we provide a simple utility function to tune a list of tasks. + # This function is just an initial implementation which tunes them in sequential order. + # We will introduce a more sophisticated tuning scheduler in the future. + + # You can skip the implementation of this function for this tutorial. + def tune_tasks( + tasks, + measure_option, + tuner="xgb", + n_trial=1000, + early_stopping=None, + log_filename="tuning.log", + use_transfer_learning=True, + ): + # create tmp log file + tmp_log_file = log_filename + ".tmp" + if os.path.exists(tmp_log_file): + os.remove(tmp_log_file) + + for i, tsk in enumerate(reversed(tasks)): + prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) + + # create tuner + if tuner == "xgb" or tuner == "xgb-rank": + tuner_obj = XGBTuner(tsk, loss_type="rank") + elif tuner == "ga": + tuner_obj = GATuner(tsk, pop_size=100) + elif tuner == "random": + tuner_obj = RandomTuner(tsk) + elif tuner == "gridsearch": + tuner_obj = GridSearchTuner(tsk) + else: + raise ValueError("Invalid tuner: " + tuner) + + if use_transfer_learning: + if os.path.isfile(tmp_log_file): + tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) + + # do tuning + tsk_trial = min(n_trial, len(tsk.config_space)) + tuner_obj.tune( + n_trial=tsk_trial, + early_stopping=early_stopping, + measure_option=measure_option, + callbacks=[ + autotvm.callback.progress_bar(tsk_trial, prefix=prefix), + autotvm.callback.log_to_file(tmp_log_file), + ], + ) + + # pick best records to a cache file + autotvm.record.pick_best(tmp_log_file, log_filename) os.remove(tmp_log_file) - for i, tsk in enumerate(reversed(tasks)): - prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) - - # create tuner - if tuner == "xgb" or tuner == "xgb-rank": - tuner_obj = XGBTuner(tsk, loss_type="rank") - elif tuner == "ga": - tuner_obj = GATuner(tsk, pop_size=100) - elif tuner == "random": - tuner_obj = RandomTuner(tsk) - elif tuner == "gridsearch": - tuner_obj = GridSearchTuner(tsk) - else: - raise ValueError("Invalid tuner: " + tuner) - - if use_transfer_learning: - if os.path.isfile(tmp_log_file): - tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) - - # do tuning - tsk_trial = min(n_trial, len(tsk.config_space)) - tuner_obj.tune( - n_trial=tsk_trial, - early_stopping=early_stopping, - measure_option=measure_option, - callbacks=[ - autotvm.callback.progress_bar(tsk_trial, prefix=prefix), - autotvm.callback.log_to_file(tmp_log_file), - ], - ) - - # pick best records to a cache file - autotvm.record.pick_best(tmp_log_file, log_filename) - os.remove(tmp_log_file) - - -######################################################################## -# Finally, we launch tuning jobs and evaluate the end-to-end performance. - - -def tune_and_evaluate(tuning_opt): - # extract workloads from relay program - print("Extract tasks...") - mod, params, input_shape, out_shape = get_network(network, batch_size=1) - tasks = autotvm.task.extract_from_program( - mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) - ) - - # run tuning tasks - print("Tuning...") - tune_tasks(tasks, **tuning_opt) - - # compile kernels with history best records - with autotvm.apply_history_best(log_file): - print("Compile...") - with tvm.transform.PassContext(opt_level=3): - lib = relay.build_module.build(mod, target=target, params=params) - - # export library - tmp = tempdir() - filename = "net.tar" - lib.export_library(tmp.relpath(filename)) - - # load parameters - ctx = tvm.context(str(target), 0) - module = runtime.GraphModule(lib["default"](ctx)) - data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) - module.set_input("data", data_tvm) - - # evaluate - print("Evaluate inference time cost...") - ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=600) - prof_res = np.array(ftimer().results) * 1000 # convert to millisecond - print( - "Mean inference time (std dev): %.2f ms (%.2f ms)" - % (np.mean(prof_res), np.std(prof_res)) - ) + ######################################################################## + # Finally, we launch tuning jobs and evaluate the end-to-end performance. -# We do not run the tuning in our webpage server since it takes too long. -# Uncomment the following line to run it by yourself. - -# tune_and_evaluate(tuning_option) - -###################################################################### -# Sample Output -# ------------- -# The tuning needs to compile many programs and extract feature from them. -# So a high performance CPU is recommended. One sample output is listed below. -# It takes about 4 hours to get the following output on a 32T AMD Ryzen Threadripper. -# The tuning target is NVIDIA 1080 Ti. -# (You can see some errors during compilation. If the tuning is not stuck, it is okay.) -# -# .. code-block:: bash -# -# Extract tasks... -# Tuning... -# [Task 1/12] Current/Best: 541.83/3570.66 GFLOPS | Progress: (960/2000) | 1001.31 s Done. -# [Task 2/12] Current/Best: 0.56/ 803.33 GFLOPS | Progress: (704/2000) | 608.08 s Done. -# [Task 3/12] Current/Best: 103.69/1141.25 GFLOPS | Progress: (768/2000) | 702.13 s Done. -# [Task 4/12] Current/Best: 2905.03/3925.15 GFLOPS | Progress: (864/2000) | 745.94 sterminate called without an active exception -# [Task 4/12] Current/Best: 2789.36/3925.15 GFLOPS | Progress: (1056/2000) | 929.40 s Done. -# [Task 5/12] Current/Best: 89.06/1076.24 GFLOPS | Progress: (704/2000) | 601.73 s Done. -# [Task 6/12] Current/Best: 40.39/2129.02 GFLOPS | Progress: (1088/2000) | 1125.76 s Done. -# [Task 7/12] Current/Best: 4090.53/5007.02 GFLOPS | Progress: (800/2000) | 903.90 s Done. -# [Task 8/12] Current/Best: 4.78/1272.28 GFLOPS | Progress: (768/2000) | 749.14 s Done. -# [Task 9/12] Current/Best: 1391.45/2325.08 GFLOPS | Progress: (992/2000) | 1084.87 s Done. -# [Task 10/12] Current/Best: 1995.44/2383.59 GFLOPS | Progress: (864/2000) | 862.60 s Done. -# [Task 11/12] Current/Best: 4093.94/4899.80 GFLOPS | Progress: (224/2000) | 240.92 sterminate called without an active exception -# [Task 11/12] Current/Best: 3487.98/4909.91 GFLOPS | Progress: (480/2000) | 534.96 sterminate called without an active exception -# [Task 11/12] Current/Best: 4636.84/4912.17 GFLOPS | Progress: (1184/2000) | 1381.16 sterminate called without an active exception -# [Task 11/12] Current/Best: 50.12/4912.17 GFLOPS | Progress: (1344/2000) | 1602.81 s Done. -# [Task 12/12] Current/Best: 3581.31/4286.30 GFLOPS | Progress: (736/2000) | 943.52 s Done. -# Compile... -# Evaluate inference time cost... -# Mean inference time (std dev): 1.07 ms (0.05 ms) -# -# As a reference baseline, the time cost of MXNet + TensorRT on resnet-18 is 1.30ms. So we are a little faster. - -###################################################################### -# -# .. note:: **Experiencing Difficulties?** -# -# The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", -# then there must be something wrong. -# -# First, make sure you set the correct configuration of your device. -# Then, you can print debug information by adding these lines in the beginning -# of the script. It will print every measurement result, where you can find useful -# error messages. -# -# .. code-block:: python -# -# import logging -# logging.getLogger('autotvm').setLevel(logging.DEBUG) -# -# Finally, always feel free to ask our community for help on https://discuss.tvm.ai + def tune_and_evaluate(tuning_opt): + # extract workloads from relay program + print("Extract tasks...") + mod, params, input_shape, out_shape = get_network(network, batch_size=1) + tasks = autotvm.task.extract_from_program( + mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) + ) -################################################################# -# Scale up measurement by using multiple devices -# ---------------------------------------------- -# -# If you have multiple devices, you can use all of them for measurement. -# TVM uses the RPC Tracker to manage distributed devices. -# The RPC Tracker is a centralized controller node. We can register all devices to -# the tracker. For example, if we have 10 GPU cards, we can register all of them -# to the tracker, and run 10 measurements in parallel, accelerating the tuning process. -# -# To start an RPC tracker, run this command on the host machine. The tracker is -# required during the whole tuning process, so we need to open a new terminal for -# this command: -# -# .. code-block:: bash -# -# python -m tvm.exec.rpc_tracker --host=0.0.0.0 --port=9190 -# -# The expected output is -# -# .. code-block:: bash -# -# INFO:RPCTracker:bind to 0.0.0.0:9190 -# -# Then open another new terminal for the RPC server. We need to start one server -# for each dedicated device. We use a string key to distinguish the types of devices. -# You can pick a name you like. -# (Note: For rocm backend, there are some internal errors with the compiler, -# we need to add `--no-fork` to the argument list.) -# -# .. code-block:: bash -# -# python -m tvm.exec.rpc_server --tracker=0.0.0.0:9190 --key=1080ti -# -# After registering devices, we can confirm it by querying rpc_tracker -# -# .. code-block:: bash -# -# python -m tvm.exec.query_rpc_tracker --host=0.0.0.0 --port=9190 -# -# For example, if we have four 1080ti, two titanx and one gfx900, the output can be -# -# .. code-block:: bash -# -# Queue Status -# ---------------------------------- -# key total free pending -# ---------------------------------- -# 1080ti 4 4 0 -# titanx 2 2 0 -# gfx900 1 1 0 -# ---------------------------------- -# -# Finally, we need to change the tuning option to use RPCRunner. Use the code below -# to replace the corresponding part above. - -tuning_option = { - "log_filename": log_file, - "tuner": "xgb", - "n_trial": 2000, - "early_stopping": 600, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(timeout=10), - runner=autotvm.RPCRunner( - "1080ti", # change the device key to your key - "0.0.0.0", - 9190, - number=20, - repeat=3, - timeout=4, - min_repeat_ms=150, + # run tuning tasks + print("Tuning...") + tune_tasks(tasks, **tuning_opt) + + # compile kernels with history best records + with autotvm.apply_history_best(log_file): + print("Compile...") + with tvm.transform.PassContext(opt_level=3): + lib = relay.build_module.build(mod, target=target, params=params) + + # export library + tmp = tempdir() + filename = "net.tar" + lib.export_library(tmp.relpath(filename)) + + # load parameters + ctx = tvm.context(str(target), 0) + module = runtime.GraphModule(lib["default"](ctx)) + data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) + module.set_input("data", data_tvm) + + # evaluate + print("Evaluate inference time cost...") + ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=600) + prof_res = np.array(ftimer().results) * 1000 # convert to millisecond + print( + "Mean inference time (std dev): %.2f ms (%.2f ms)" + % (np.mean(prof_res), np.std(prof_res)) + ) + + + # We do not run the tuning in our webpage server since it takes too long. + # Uncomment the following line to run it by yourself. + + # tune_and_evaluate(tuning_option) + + ###################################################################### + # Sample Output + # ------------- + # The tuning needs to compile many programs and extract feature from them. + # So a high performance CPU is recommended. One sample output is listed below. + # It takes about 4 hours to get the following output on a 32T AMD Ryzen Threadripper. + # The tuning target is NVIDIA 1080 Ti. + # (You can see some errors during compilation. If the tuning is not stuck, it is okay.) + # + # .. code-block:: bash + # + # Extract tasks... + # Tuning... + # [Task 1/12] Current/Best: 541.83/3570.66 GFLOPS | Progress: (960/2000) | 1001.31 s Done. + # [Task 2/12] Current/Best: 0.56/ 803.33 GFLOPS | Progress: (704/2000) | 608.08 s Done. + # [Task 3/12] Current/Best: 103.69/1141.25 GFLOPS | Progress: (768/2000) | 702.13 s Done. + # [Task 4/12] Current/Best: 2905.03/3925.15 GFLOPS | Progress: (864/2000) | 745.94 sterminate called without an active exception + # [Task 4/12] Current/Best: 2789.36/3925.15 GFLOPS | Progress: (1056/2000) | 929.40 s Done. + # [Task 5/12] Current/Best: 89.06/1076.24 GFLOPS | Progress: (704/2000) | 601.73 s Done. + # [Task 6/12] Current/Best: 40.39/2129.02 GFLOPS | Progress: (1088/2000) | 1125.76 s Done. + # [Task 7/12] Current/Best: 4090.53/5007.02 GFLOPS | Progress: (800/2000) | 903.90 s Done. + # [Task 8/12] Current/Best: 4.78/1272.28 GFLOPS | Progress: (768/2000) | 749.14 s Done. + # [Task 9/12] Current/Best: 1391.45/2325.08 GFLOPS | Progress: (992/2000) | 1084.87 s Done. + # [Task 10/12] Current/Best: 1995.44/2383.59 GFLOPS | Progress: (864/2000) | 862.60 s Done. + # [Task 11/12] Current/Best: 4093.94/4899.80 GFLOPS | Progress: (224/2000) | 240.92 sterminate called without an active exception + # [Task 11/12] Current/Best: 3487.98/4909.91 GFLOPS | Progress: (480/2000) | 534.96 sterminate called without an active exception + # [Task 11/12] Current/Best: 4636.84/4912.17 GFLOPS | Progress: (1184/2000) | 1381.16 sterminate called without an active exception + # [Task 11/12] Current/Best: 50.12/4912.17 GFLOPS | Progress: (1344/2000) | 1602.81 s Done. + # [Task 12/12] Current/Best: 3581.31/4286.30 GFLOPS | Progress: (736/2000) | 943.52 s Done. + # Compile... + # Evaluate inference time cost... + # Mean inference time (std dev): 1.07 ms (0.05 ms) + # + # As a reference baseline, the time cost of MXNet + TensorRT on resnet-18 is 1.30ms. So we are a little faster. + + ###################################################################### + # + # .. note:: **Experiencing Difficulties?** + # + # The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", + # then there must be something wrong. + # + # First, make sure you set the correct configuration of your device. + # Then, you can print debug information by adding these lines in the beginning + # of the script. It will print every measurement result, where you can find useful + # error messages. + # + # .. code-block:: python + # + # import logging + # logging.getLogger('autotvm').setLevel(logging.DEBUG) + # + # Finally, always feel free to ask our community for help on https://discuss.tvm.ai + + + ################################################################# + # Scale up measurement by using multiple devices + # ---------------------------------------------- + # + # If you have multiple devices, you can use all of them for measurement. + # TVM uses the RPC Tracker to manage distributed devices. + # The RPC Tracker is a centralized controller node. We can register all devices to + # the tracker. For example, if we have 10 GPU cards, we can register all of them + # to the tracker, and run 10 measurements in parallel, accelerating the tuning process. + # + # To start an RPC tracker, run this command on the host machine. The tracker is + # required during the whole tuning process, so we need to open a new terminal for + # this command: + # + # .. code-block:: bash + # + # python -m tvm.exec.rpc_tracker --host=0.0.0.0 --port=9190 + # + # The expected output is + # + # .. code-block:: bash + # + # INFO:RPCTracker:bind to 0.0.0.0:9190 + # + # Then open another new terminal for the RPC server. We need to start one server + # for each dedicated device. We use a string key to distinguish the types of devices. + # You can pick a name you like. + # (Note: For rocm backend, there are some internal errors with the compiler, + # we need to add `--no-fork` to the argument list.) + # + # .. code-block:: bash + # + # python -m tvm.exec.rpc_server --tracker=0.0.0.0:9190 --key=1080ti + # + # After registering devices, we can confirm it by querying rpc_tracker + # + # .. code-block:: bash + # + # python -m tvm.exec.query_rpc_tracker --host=0.0.0.0 --port=9190 + # + # For example, if we have four 1080ti, two titanx and one gfx900, the output can be + # + # .. code-block:: bash + # + # Queue Status + # ---------------------------------- + # key total free pending + # ---------------------------------- + # 1080ti 4 4 0 + # titanx 2 2 0 + # gfx900 1 1 0 + # ---------------------------------- + # + # Finally, we need to change the tuning option to use RPCRunner. Use the code below + # to replace the corresponding part above. + + tuning_option = { + "log_filename": log_file, + "tuner": "xgb", + "n_trial": 2000, + "early_stopping": 600, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(timeout=10), + runner=autotvm.RPCRunner( + "1080ti", # change the device key to your key + "0.0.0.0", + 9190, + number=20, + repeat=3, + timeout=4, + min_repeat_ms=150, + ), ), - ), -} + } diff --git a/tutorials/autotvm/tune_relay_mobile_gpu.py b/tutorials/autotvm/tune_relay_mobile_gpu.py index 61254662c463..5e04569e14dc 100644 --- a/tutorials/autotvm/tune_relay_mobile_gpu.py +++ b/tutorials/autotvm/tune_relay_mobile_gpu.py @@ -188,233 +188,234 @@ def get_network(name, batch_size): # # You can register multiple devices to the tracker to accelerate the measurement in tuning. -########################################### -# Set Tuning Options -# ------------------ -# Before tuning, we should apply some configurations. Here I use an RK3399 board -# as example. In your setting, you should modify the target and device_key accordingly. -# set :code:`use_android` to True if you use android phone. - -#### DEVICE CONFIG #### - -target = tvm.target.Target("opencl -device=mali") - -# Replace "aarch64-linux-gnu" with the correct target of your board. -# This target host is used for cross compilation. You can query it by :code:`gcc -v` on your device. -target_host = "llvm -mtriple=aarch64-linux-gnu" - -# Also replace this with the device key in your tracker -device_key = "rk3399" - -# Set this to True if you use android phone -use_android = False - -#### TUNING OPTION #### -network = "resnet-18" -log_file = "%s.%s.log" % (device_key, network) -dtype = "float32" - -tuning_option = { - "log_filename": log_file, - "tuner": "xgb", - "n_trial": 1000, - "early_stopping": 450, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(build_func="ndk" if use_android else "default"), - runner=autotvm.RPCRunner( - device_key, - host="0.0.0.0", - port=9190, - number=10, - timeout=5, +if __name__ == "__main__": + ########################################### + # Set Tuning Options + # ------------------ + # Before tuning, we should apply some configurations. Here I use an RK3399 board + # as example. In your setting, you should modify the target and device_key accordingly. + # set :code:`use_android` to True if you use android phone. + + #### DEVICE CONFIG #### + + target = tvm.target.Target("opencl -device=mali") + + # Replace "aarch64-linux-gnu" with the correct target of your board. + # This target host is used for cross compilation. You can query it by :code:`gcc -v` on your device. + target_host = "llvm -mtriple=aarch64-linux-gnu" + + # Also replace this with the device key in your tracker + device_key = "rk3399" + + # Set this to True if you use android phone + use_android = False + + #### TUNING OPTION #### + network = "resnet-18" + log_file = "%s.%s.log" % (device_key, network) + dtype = "float32" + + tuning_option = { + "log_filename": log_file, + "tuner": "xgb", + "n_trial": 1000, + "early_stopping": 450, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(build_func="ndk" if use_android else "default"), + runner=autotvm.RPCRunner( + device_key, + host="0.0.0.0", + port=9190, + number=10, + timeout=5, + ), ), - ), -} - -#################################################################### -# -# .. note:: How to set tuning options -# -# In general, the default values provided here work well. -# If you have enough time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, -# which makes the tuning run longer. -# If your device runs very slow or your conv2d operators have many GFLOPs, considering to -# set timeout larger. -# + } + + #################################################################### + # + # .. note:: How to set tuning options + # + # In general, the default values provided here work well. + # If you have enough time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, + # which makes the tuning run longer. + # If your device runs very slow or your conv2d operators have many GFLOPs, considering to + # set timeout larger. + # + + ################################################################### + # Begin Tuning + # ------------ + # Now we can extract tuning tasks from the network and begin tuning. + # Here, we provide a simple utility function to tune a list of tasks. + # This function is just an initial implementation which tunes them in sequential order. + # We will introduce a more sophisticated tuning scheduler in the future. + + # You can skip the implementation of this function for this tutorial. + def tune_tasks( + tasks, + measure_option, + tuner="xgb", + n_trial=1000, + early_stopping=None, + log_filename="tuning.log", + use_transfer_learning=True, + ): + # create tmp log file + tmp_log_file = log_filename + ".tmp" + if os.path.exists(tmp_log_file): + os.remove(tmp_log_file) + + for i, tsk in enumerate(reversed(tasks)): + prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) + + # create tuner + if tuner == "xgb" or tuner == "xgb-rank": + tuner_obj = XGBTuner(tsk, loss_type="rank") + elif tuner == "ga": + tuner_obj = GATuner(tsk, pop_size=50) + elif tuner == "random": + tuner_obj = RandomTuner(tsk) + elif tuner == "gridsearch": + tuner_obj = GridSearchTuner(tsk) + else: + raise ValueError("Invalid tuner: " + tuner) + + if use_transfer_learning: + if os.path.isfile(tmp_log_file): + tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) + + # do tuning + tsk_trial = min(n_trial, len(tsk.config_space)) + tuner_obj.tune( + n_trial=tsk_trial, + early_stopping=early_stopping, + measure_option=measure_option, + callbacks=[ + autotvm.callback.progress_bar(tsk_trial, prefix=prefix), + autotvm.callback.log_to_file(tmp_log_file), + ], + ) -################################################################### -# Begin Tuning -# ------------ -# Now we can extract tuning tasks from the network and begin tuning. -# Here, we provide a simple utility function to tune a list of tasks. -# This function is just an initial implementation which tunes them in sequential order. -# We will introduce a more sophisticated tuning scheduler in the future. - -# You can skip the implementation of this function for this tutorial. -def tune_tasks( - tasks, - measure_option, - tuner="xgb", - n_trial=1000, - early_stopping=None, - log_filename="tuning.log", - use_transfer_learning=True, -): - # create tmp log file - tmp_log_file = log_filename + ".tmp" - if os.path.exists(tmp_log_file): + # pick best records to a cache file + autotvm.record.pick_best(tmp_log_file, log_filename) os.remove(tmp_log_file) - for i, tsk in enumerate(reversed(tasks)): - prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) - - # create tuner - if tuner == "xgb" or tuner == "xgb-rank": - tuner_obj = XGBTuner(tsk, loss_type="rank") - elif tuner == "ga": - tuner_obj = GATuner(tsk, pop_size=50) - elif tuner == "random": - tuner_obj = RandomTuner(tsk) - elif tuner == "gridsearch": - tuner_obj = GridSearchTuner(tsk) - else: - raise ValueError("Invalid tuner: " + tuner) - - if use_transfer_learning: - if os.path.isfile(tmp_log_file): - tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) - - # do tuning - tsk_trial = min(n_trial, len(tsk.config_space)) - tuner_obj.tune( - n_trial=tsk_trial, - early_stopping=early_stopping, - measure_option=measure_option, - callbacks=[ - autotvm.callback.progress_bar(tsk_trial, prefix=prefix), - autotvm.callback.log_to_file(tmp_log_file), - ], - ) - # pick best records to a cache file - autotvm.record.pick_best(tmp_log_file, log_filename) - os.remove(tmp_log_file) - - -######################################################################## -# Finally, we launch tuning jobs and evaluate the end-to-end performance. - - -def tune_and_evaluate(tuning_opt): - # extract workloads from relay program - print("Extract tasks...") - mod, params, input_shape, _ = get_network(network, batch_size=1) - tasks = autotvm.task.extract_from_program( - mod["main"], - target=target, - target_host=target_host, - params=params, - ops=(relay.op.get("nn.conv2d"),), - ) - - # run tuning tasks - print("Tuning...") - tune_tasks(tasks, **tuning_opt) - - # compile kernels with history best records - with autotvm.apply_history_best(log_file): - print("Compile...") - with tvm.transform.PassContext(opt_level=3): - lib = relay.build_module.build( - mod, target=target, params=params, target_host=target_host - ) - # export library - tmp = tempdir() - if use_android: - from tvm.contrib import ndk - - filename = "net.so" - lib.export_library(tmp.relpath(filename), ndk.create_shared) - else: - filename = "net.tar" - lib.export_library(tmp.relpath(filename)) - - # upload module to device - print("Upload...") - remote = autotvm.measure.request_remote(device_key, "0.0.0.0", 9190, timeout=10000) - remote.upload(tmp.relpath(filename)) - rlib = remote.load_module(filename) - - # upload parameters to device - ctx = remote.context(str(target), 0) - module = runtime.GraphModule(rlib["default"](ctx)) - data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) - module.set_input("data", data_tvm) - - # evaluate - print("Evaluate inference time cost...") - ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=30) - prof_res = np.array(ftimer().results) * 1000 # convert to millisecond - print( - "Mean inference time (std dev): %.2f ms (%.2f ms)" - % (np.mean(prof_res), np.std(prof_res)) - ) + ######################################################################## + # Finally, we launch tuning jobs and evaluate the end-to-end performance. -# We do not run the tuning in our webpage server since it takes too long. -# Uncomment the following line to run it by yourself. + def tune_and_evaluate(tuning_opt): + # extract workloads from relay program + print("Extract tasks...") + mod, params, input_shape, _ = get_network(network, batch_size=1) + tasks = autotvm.task.extract_from_program( + mod["main"], + target=target, + target_host=target_host, + params=params, + ops=(relay.op.get("nn.conv2d"),), + ) -# tune_and_evaluate(tuning_option) + # run tuning tasks + print("Tuning...") + tune_tasks(tasks, **tuning_opt) + + # compile kernels with history best records + with autotvm.apply_history_best(log_file): + print("Compile...") + with tvm.transform.PassContext(opt_level=3): + lib = relay.build_module.build( + mod, target=target, params=params, target_host=target_host + ) + # export library + tmp = tempdir() + if use_android: + from tvm.contrib import ndk + + filename = "net.so" + lib.export_library(tmp.relpath(filename), ndk.create_shared) + else: + filename = "net.tar" + lib.export_library(tmp.relpath(filename)) + + # upload module to device + print("Upload...") + remote = autotvm.measure.request_remote(device_key, "0.0.0.0", 9190, timeout=10000) + remote.upload(tmp.relpath(filename)) + rlib = remote.load_module(filename) + + # upload parameters to device + ctx = remote.context(str(target), 0) + module = runtime.GraphModule(rlib["default"](ctx)) + data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) + module.set_input("data", data_tvm) + + # evaluate + print("Evaluate inference time cost...") + ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=30) + prof_res = np.array(ftimer().results) * 1000 # convert to millisecond + print( + "Mean inference time (std dev): %.2f ms (%.2f ms)" + % (np.mean(prof_res), np.std(prof_res)) + ) -###################################################################### -# Sample Output -# ------------- -# The tuning needs to compile many programs and extract feature from them. -# So a high performance CPU is recommended. -# One sample output is listed below. It takes about 3 hours on a 32T AMD Ryzen Threadripper. -# -# .. code-block:: bash -# -# Extract tasks... -# Tuning... -# [Task 1/17] Current/Best: 25.30/ 39.12 GFLOPS | Progress: (992/1000) | 751.22 s Done. -# [Task 2/17] Current/Best: 40.70/ 45.50 GFLOPS | Progress: (736/1000) | 545.46 s Done. -# [Task 3/17] Current/Best: 38.83/ 42.35 GFLOPS | Progress: (992/1000) | 1549.85 s Done. -# [Task 4/17] Current/Best: 23.31/ 31.02 GFLOPS | Progress: (640/1000) | 1059.31 s Done. -# [Task 5/17] Current/Best: 0.06/ 2.34 GFLOPS | Progress: (544/1000) | 305.45 s Done. -# [Task 6/17] Current/Best: 10.97/ 17.20 GFLOPS | Progress: (992/1000) | 1050.00 s Done. -# [Task 7/17] Current/Best: 8.98/ 10.94 GFLOPS | Progress: (928/1000) | 421.36 s Done. -# [Task 8/17] Current/Best: 4.48/ 14.86 GFLOPS | Progress: (704/1000) | 582.60 s Done. -# [Task 9/17] Current/Best: 10.30/ 25.99 GFLOPS | Progress: (864/1000) | 899.85 s Done. -# [Task 10/17] Current/Best: 11.73/ 12.52 GFLOPS | Progress: (608/1000) | 304.85 s Done. -# [Task 11/17] Current/Best: 15.26/ 18.68 GFLOPS | Progress: (800/1000) | 747.52 s Done. -# [Task 12/17] Current/Best: 17.48/ 26.71 GFLOPS | Progress: (1000/1000) | 1166.40 s Done. -# [Task 13/17] Current/Best: 0.96/ 11.43 GFLOPS | Progress: (960/1000) | 611.65 s Done. -# [Task 14/17] Current/Best: 17.88/ 20.22 GFLOPS | Progress: (672/1000) | 670.29 s Done. -# [Task 15/17] Current/Best: 11.62/ 13.98 GFLOPS | Progress: (736/1000) | 449.25 s Done. -# [Task 16/17] Current/Best: 19.90/ 23.83 GFLOPS | Progress: (608/1000) | 708.64 s Done. -# [Task 17/17] Current/Best: 17.98/ 22.75 GFLOPS | Progress: (736/1000) | 1122.60 s Done. -# Compile... -# Upload... -# Evaluate inference time cost... -# Mean inference time (std dev): 128.05 ms (7.74 ms) -# -###################################################################### -# -# .. note:: **Experiencing Difficulties?** -# -# The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", -# then there must be something wrong. -# -# First, make sure you set the correct configuration of your device. -# Then, you can print debug information by adding these lines in the beginning -# of the script. It will print every measurement result, where you can find useful -# error messages. -# -# .. code-block:: python -# -# import logging -# logging.getLogger('autotvm').setLevel(logging.DEBUG) -# -# Finally, always feel free to ask our community for help on https://discuss.tvm.ai + # We do not run the tuning in our webpage server since it takes too long. + # Uncomment the following line to run it by yourself. + + # tune_and_evaluate(tuning_option) + + ###################################################################### + # Sample Output + # ------------- + # The tuning needs to compile many programs and extract feature from them. + # So a high performance CPU is recommended. + # One sample output is listed below. It takes about 3 hours on a 32T AMD Ryzen Threadripper. + # + # .. code-block:: bash + # + # Extract tasks... + # Tuning... + # [Task 1/17] Current/Best: 25.30/ 39.12 GFLOPS | Progress: (992/1000) | 751.22 s Done. + # [Task 2/17] Current/Best: 40.70/ 45.50 GFLOPS | Progress: (736/1000) | 545.46 s Done. + # [Task 3/17] Current/Best: 38.83/ 42.35 GFLOPS | Progress: (992/1000) | 1549.85 s Done. + # [Task 4/17] Current/Best: 23.31/ 31.02 GFLOPS | Progress: (640/1000) | 1059.31 s Done. + # [Task 5/17] Current/Best: 0.06/ 2.34 GFLOPS | Progress: (544/1000) | 305.45 s Done. + # [Task 6/17] Current/Best: 10.97/ 17.20 GFLOPS | Progress: (992/1000) | 1050.00 s Done. + # [Task 7/17] Current/Best: 8.98/ 10.94 GFLOPS | Progress: (928/1000) | 421.36 s Done. + # [Task 8/17] Current/Best: 4.48/ 14.86 GFLOPS | Progress: (704/1000) | 582.60 s Done. + # [Task 9/17] Current/Best: 10.30/ 25.99 GFLOPS | Progress: (864/1000) | 899.85 s Done. + # [Task 10/17] Current/Best: 11.73/ 12.52 GFLOPS | Progress: (608/1000) | 304.85 s Done. + # [Task 11/17] Current/Best: 15.26/ 18.68 GFLOPS | Progress: (800/1000) | 747.52 s Done. + # [Task 12/17] Current/Best: 17.48/ 26.71 GFLOPS | Progress: (1000/1000) | 1166.40 s Done. + # [Task 13/17] Current/Best: 0.96/ 11.43 GFLOPS | Progress: (960/1000) | 611.65 s Done. + # [Task 14/17] Current/Best: 17.88/ 20.22 GFLOPS | Progress: (672/1000) | 670.29 s Done. + # [Task 15/17] Current/Best: 11.62/ 13.98 GFLOPS | Progress: (736/1000) | 449.25 s Done. + # [Task 16/17] Current/Best: 19.90/ 23.83 GFLOPS | Progress: (608/1000) | 708.64 s Done. + # [Task 17/17] Current/Best: 17.98/ 22.75 GFLOPS | Progress: (736/1000) | 1122.60 s Done. + # Compile... + # Upload... + # Evaluate inference time cost... + # Mean inference time (std dev): 128.05 ms (7.74 ms) + # + + ###################################################################### + # + # .. note:: **Experiencing Difficulties?** + # + # The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", + # then there must be something wrong. + # + # First, make sure you set the correct configuration of your device. + # Then, you can print debug information by adding these lines in the beginning + # of the script. It will print every measurement result, where you can find useful + # error messages. + # + # .. code-block:: python + # + # import logging + # logging.getLogger('autotvm').setLevel(logging.DEBUG) + # + # Finally, always feel free to ask our community for help on https://discuss.tvm.ai diff --git a/tutorials/autotvm/tune_relay_x86.py b/tutorials/autotvm/tune_relay_x86.py index 1dd947fefd25..2e990efdfa87 100644 --- a/tutorials/autotvm/tune_relay_x86.py +++ b/tutorials/autotvm/tune_relay_x86.py @@ -88,172 +88,173 @@ def get_network(name, batch_size): return mod, params, input_shape, output_shape -# Replace "llvm" with the correct target of your CPU. -# For example, for AWS EC2 c5 instance with Intel Xeon -# Platinum 8000 series, the target should be "llvm -mcpu=skylake-avx512". -# For AWS EC2 c4 instance with Intel Xeon E5-2666 v3, it should be -# "llvm -mcpu=core-avx2". -target = "llvm" - -batch_size = 1 -dtype = "float32" -model_name = "resnet-18" -log_file = "%s.log" % model_name -graph_opt_sch_file = "%s_graph_opt.log" % model_name - -# Set the input name of the graph -# For ONNX models, it is typically "0". -input_name = "data" - -# Set number of threads used for tuning based on the number of -# physical CPU cores on your machine. -num_threads = 1 -os.environ["TVM_NUM_THREADS"] = str(num_threads) - - -################################################################# -# Configure tensor tuning settings and create tasks -# ------------------------------------------------- -# To get better kernel execution performance on x86 CPU, -# we need to change data layout of convolution kernel from -# "NCHW" to "NCHWc". To deal with this situation, we define -# conv2d_NCHWc operator in topi. We will tune this operator -# instead of plain conv2d. -# -# We will use local mode for tuning configuration. RPC tracker -# mode can be setup similarly to the approach in -# :ref:`tune_relay_arm` tutorial. -# -# To perform a precise measurement, we should repeat the measurement several -# times and use the average of results. In addition, we need to flush the cache -# for the weight tensors between repeated measurements. This can make the measured -# latency of one operator closer to its actual latency during end-to-end inference. - -tuning_option = { - "log_filename": log_file, - "tuner": "random", - "early_stopping": None, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(), - runner=autotvm.LocalRunner( - number=1, repeat=10, min_repeat_ms=0, enable_cpu_cache_flush=True +if __name__ == "__main__": + # Replace "llvm" with the correct target of your CPU. + # For example, for AWS EC2 c5 instance with Intel Xeon + # Platinum 8000 series, the target should be "llvm -mcpu=skylake-avx512". + # For AWS EC2 c4 instance with Intel Xeon E5-2666 v3, it should be + # "llvm -mcpu=core-avx2". + target = "llvm" + + batch_size = 1 + dtype = "float32" + model_name = "resnet-18" + log_file = "%s.log" % model_name + graph_opt_sch_file = "%s_graph_opt.log" % model_name + + # Set the input name of the graph + # For ONNX models, it is typically "0". + input_name = "data" + + # Set number of threads used for tuning based on the number of + # physical CPU cores on your machine. + num_threads = 1 + os.environ["TVM_NUM_THREADS"] = str(num_threads) + + + ################################################################# + # Configure tensor tuning settings and create tasks + # ------------------------------------------------- + # To get better kernel execution performance on x86 CPU, + # we need to change data layout of convolution kernel from + # "NCHW" to "NCHWc". To deal with this situation, we define + # conv2d_NCHWc operator in topi. We will tune this operator + # instead of plain conv2d. + # + # We will use local mode for tuning configuration. RPC tracker + # mode can be setup similarly to the approach in + # :ref:`tune_relay_arm` tutorial. + # + # To perform a precise measurement, we should repeat the measurement several + # times and use the average of results. In addition, we need to flush the cache + # for the weight tensors between repeated measurements. This can make the measured + # latency of one operator closer to its actual latency during end-to-end inference. + + tuning_option = { + "log_filename": log_file, + "tuner": "random", + "early_stopping": None, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(), + runner=autotvm.LocalRunner( + number=1, repeat=10, min_repeat_ms=0, enable_cpu_cache_flush=True + ), ), - ), -} - - -# You can skip the implementation of this function for this tutorial. -def tune_kernels( - tasks, measure_option, tuner="gridsearch", early_stopping=None, log_filename="tuning.log" -): - - for i, task in enumerate(tasks): - prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) - - # create tuner - if tuner == "xgb" or tuner == "xgb-rank": - tuner_obj = XGBTuner(task, loss_type="rank") - elif tuner == "ga": - tuner_obj = GATuner(task, pop_size=50) - elif tuner == "random": - tuner_obj = RandomTuner(task) - elif tuner == "gridsearch": - tuner_obj = GridSearchTuner(task) - else: - raise ValueError("Invalid tuner: " + tuner) - - # do tuning - n_trial = len(task.config_space) - tuner_obj.tune( - n_trial=n_trial, - early_stopping=early_stopping, - measure_option=measure_option, - callbacks=[ - autotvm.callback.progress_bar(n_trial, prefix=prefix), - autotvm.callback.log_to_file(log_filename), - ], + } + + + # You can skip the implementation of this function for this tutorial. + def tune_kernels( + tasks, measure_option, tuner="gridsearch", early_stopping=None, log_filename="tuning.log" + ): + + for i, task in enumerate(tasks): + prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) + + # create tuner + if tuner == "xgb" or tuner == "xgb-rank": + tuner_obj = XGBTuner(task, loss_type="rank") + elif tuner == "ga": + tuner_obj = GATuner(task, pop_size=50) + elif tuner == "random": + tuner_obj = RandomTuner(task) + elif tuner == "gridsearch": + tuner_obj = GridSearchTuner(task) + else: + raise ValueError("Invalid tuner: " + tuner) + + # do tuning + n_trial = len(task.config_space) + tuner_obj.tune( + n_trial=n_trial, + early_stopping=early_stopping, + measure_option=measure_option, + callbacks=[ + autotvm.callback.progress_bar(n_trial, prefix=prefix), + autotvm.callback.log_to_file(log_filename), + ], + ) + + + # Use graph tuner to achieve graph level optimal schedules + # Set use_DP=False if it takes too long to finish. + def tune_graph(graph, dshape, records, opt_sch_file, use_DP=True): + target_op = [ + relay.op.get("nn.conv2d"), + ] + Tuner = DPTuner if use_DP else PBQPTuner + executor = Tuner(graph, {input_name: dshape}, records, target_op, target) + executor.benchmark_layout_transform(min_exec_num=2000) + executor.run() + executor.write_opt_sch2record_file(opt_sch_file) + + + ######################################################################## + # Finally, we launch tuning jobs and evaluate the end-to-end performance. + + + def tune_and_evaluate(tuning_opt): + # extract workloads from relay program + print("Extract tasks...") + mod, params, data_shape, out_shape = get_network(model_name, batch_size) + tasks = autotvm.task.extract_from_program( + mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) ) - -# Use graph tuner to achieve graph level optimal schedules -# Set use_DP=False if it takes too long to finish. -def tune_graph(graph, dshape, records, opt_sch_file, use_DP=True): - target_op = [ - relay.op.get("nn.conv2d"), - ] - Tuner = DPTuner if use_DP else PBQPTuner - executor = Tuner(graph, {input_name: dshape}, records, target_op, target) - executor.benchmark_layout_transform(min_exec_num=2000) - executor.run() - executor.write_opt_sch2record_file(opt_sch_file) - - -######################################################################## -# Finally, we launch tuning jobs and evaluate the end-to-end performance. - - -def tune_and_evaluate(tuning_opt): - # extract workloads from relay program - print("Extract tasks...") - mod, params, data_shape, out_shape = get_network(model_name, batch_size) - tasks = autotvm.task.extract_from_program( - mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) - ) - - # run tuning tasks - tune_kernels(tasks, **tuning_opt) - tune_graph(mod["main"], data_shape, log_file, graph_opt_sch_file) - - # compile kernels with graph-level best records - with autotvm.apply_graph_best(graph_opt_sch_file): - print("Compile...") - with tvm.transform.PassContext(opt_level=3): - lib = relay.build_module.build(mod, target=target, params=params) - - # upload parameters to device - ctx = tvm.cpu() - data_tvm = tvm.nd.array((np.random.uniform(size=data_shape)).astype(dtype)) - module = runtime.GraphModule(lib["default"](ctx)) - module.set_input(input_name, data_tvm) - - # evaluate - print("Evaluate inference time cost...") - ftimer = module.module.time_evaluator("run", ctx, number=100, repeat=3) - prof_res = np.array(ftimer().results) * 1000 # convert to millisecond - print( - "Mean inference time (std dev): %.2f ms (%.2f ms)" - % (np.mean(prof_res), np.std(prof_res)) - ) - - -# We do not run the tuning in our webpage server since it takes too long. -# Uncomment the following line to run it by yourself. - -# tune_and_evaluate(tuning_option) - -###################################################################### -# Sample Output -# ------------- -# The tuning needs to compile many programs and extract feature from them. -# So a high performance CPU is recommended. -# One sample output is listed below. -# -# .. code-block:: bash -# -# Extract tasks... -# Tuning... -# [Task 1/12] Current/Best: 598.05/2497.63 GFLOPS | Progress: (252/252) | 1357.95 s Done. -# [Task 2/12] Current/Best: 522.63/2279.24 GFLOPS | Progress: (784/784) | 3989.60 s Done. -# [Task 3/12] Current/Best: 447.33/1927.69 GFLOPS | Progress: (784/784) | 3869.14 s Done. -# [Task 4/12] Current/Best: 481.11/1912.34 GFLOPS | Progress: (672/672) | 3274.25 s Done. -# [Task 5/12] Current/Best: 414.09/1598.45 GFLOPS | Progress: (672/672) | 2720.78 s Done. -# [Task 6/12] Current/Best: 508.96/2273.20 GFLOPS | Progress: (768/768) | 3718.75 s Done. -# [Task 7/12] Current/Best: 469.14/1955.79 GFLOPS | Progress: (576/576) | 2665.67 s Done. -# [Task 8/12] Current/Best: 230.91/1658.97 GFLOPS | Progress: (576/576) | 2435.01 s Done. -# [Task 9/12] Current/Best: 487.75/2295.19 GFLOPS | Progress: (648/648) | 3009.95 s Done. -# [Task 10/12] Current/Best: 182.33/1734.45 GFLOPS | Progress: (360/360) | 1755.06 s Done. -# [Task 11/12] Current/Best: 372.18/1745.15 GFLOPS | Progress: (360/360) | 1684.50 s Done. -# [Task 12/12] Current/Best: 215.34/2271.11 GFLOPS | Progress: (400/400) | 2128.74 s Done. -# Compile... -# Evaluate inference time cost... -# Mean inference time (std dev): 3.16 ms (0.03 ms) + # run tuning tasks + tune_kernels(tasks, **tuning_opt) + tune_graph(mod["main"], data_shape, log_file, graph_opt_sch_file) + + # compile kernels with graph-level best records + with autotvm.apply_graph_best(graph_opt_sch_file): + print("Compile...") + with tvm.transform.PassContext(opt_level=3): + lib = relay.build_module.build(mod, target=target, params=params) + + # upload parameters to device + ctx = tvm.cpu() + data_tvm = tvm.nd.array((np.random.uniform(size=data_shape)).astype(dtype)) + module = runtime.GraphModule(lib["default"](ctx)) + module.set_input(input_name, data_tvm) + + # evaluate + print("Evaluate inference time cost...") + ftimer = module.module.time_evaluator("run", ctx, number=100, repeat=3) + prof_res = np.array(ftimer().results) * 1000 # convert to millisecond + print( + "Mean inference time (std dev): %.2f ms (%.2f ms)" + % (np.mean(prof_res), np.std(prof_res)) + ) + + + # We do not run the tuning in our webpage server since it takes too long. + # Uncomment the following line to run it by yourself. + + # tune_and_evaluate(tuning_option) + + ###################################################################### + # Sample Output + # ------------- + # The tuning needs to compile many programs and extract feature from them. + # So a high performance CPU is recommended. + # One sample output is listed below. + # + # .. code-block:: bash + # + # Extract tasks... + # Tuning... + # [Task 1/12] Current/Best: 598.05/2497.63 GFLOPS | Progress: (252/252) | 1357.95 s Done. + # [Task 2/12] Current/Best: 522.63/2279.24 GFLOPS | Progress: (784/784) | 3989.60 s Done. + # [Task 3/12] Current/Best: 447.33/1927.69 GFLOPS | Progress: (784/784) | 3869.14 s Done. + # [Task 4/12] Current/Best: 481.11/1912.34 GFLOPS | Progress: (672/672) | 3274.25 s Done. + # [Task 5/12] Current/Best: 414.09/1598.45 GFLOPS | Progress: (672/672) | 2720.78 s Done. + # [Task 6/12] Current/Best: 508.96/2273.20 GFLOPS | Progress: (768/768) | 3718.75 s Done. + # [Task 7/12] Current/Best: 469.14/1955.79 GFLOPS | Progress: (576/576) | 2665.67 s Done. + # [Task 8/12] Current/Best: 230.91/1658.97 GFLOPS | Progress: (576/576) | 2435.01 s Done. + # [Task 9/12] Current/Best: 487.75/2295.19 GFLOPS | Progress: (648/648) | 3009.95 s Done. + # [Task 10/12] Current/Best: 182.33/1734.45 GFLOPS | Progress: (360/360) | 1755.06 s Done. + # [Task 11/12] Current/Best: 372.18/1745.15 GFLOPS | Progress: (360/360) | 1684.50 s Done. + # [Task 12/12] Current/Best: 215.34/2271.11 GFLOPS | Progress: (400/400) | 2128.74 s Done. + # Compile... + # Evaluate inference time cost... + # Mean inference time (std dev): 3.16 ms (0.03 ms) diff --git a/tutorials/autotvm/tune_simple_template.py b/tutorials/autotvm/tune_simple_template.py index 357abf19a09c..2243d2df6347 100644 --- a/tutorials/autotvm/tune_simple_template.py +++ b/tutorials/autotvm/tune_simple_template.py @@ -56,6 +56,7 @@ import numpy as np import tvm from tvm import te +import tvm.testing # the module is called `autotvm` from tvm import autotvm @@ -214,118 +215,119 @@ def matmul(N, L, M, dtype): return s, [A, B, C] -###################################################################### -# .. note:: More Explanation on :code:`cfg.defile_split` -# -# In this template, :code:`cfg.define_split("tile_y", y, num_outputs=2)` will enumerate -# all possible combinations that can split axis y into two axes with factors of the length of y. -# For example, if the length of y is 32 and we want to split it into two axes -# using factors of 32, then there are 6 possible values for -# (length of outer axis, length of inner axis) pair, namely -# (32, 1), (16, 2), (8, 4), (4, 8), (2, 16) or (1, 32). -# They are just the 6 possible values of `tile_y`. -# -# During schedule, :code:`cfg["tile_y"]` is a :code:`SplitEntity` object. -# We stores the lengths of outer axes and inner axes in :code:`cfg['tile_y'].size` -# (a tuple with two elements). -# In this template, we apply it by using :code:`yo, yi = cfg['tile_y'].apply(s, C, y)`. -# Actually, this is equivalent to -# :code:`yo, yi = s[C].split(y, cfg["tile_y"].size[1])` -# or :code:`yo, yi = s[C].split(y, nparts=cfg['tile_y"].size[0])` -# -# The advantage of using cfg.apply API is that it makes multi-level split -# (when num_outputs >= 3) easier. - -###################################################################### -# Step 2: Search through the space -# --------------------------------- -# In step 1, we build the search space by extending our old schedule code -# into a template. The next step is to pick a tuner and explore in this space. -# -# Auto-tuners in TVM -# ^^^^^^^^^^^^^^^^^^ -# The job for a tuner can be described by following pseudo code -# -# .. code-block:: c -# -# ct = 0 -# while ct < max_number_of_trials: -# propose a batch of configs -# measure this batch of configs on real hardware and get results -# ct += batch_size -# -# When proposing the next batch of configs, the tuner can take different strategies. We -# provide four tuners with different strategies in autotvm. -# -# * :any:`RandomTuner`: Enumerate the space in a random order -# * :any:`GridSearchTuner`: Enumerate the space in a grid search order -# * :any:`GATuner`: Using genetic algorithm to search through the space -# * :any:`XGBTuner`: Uses a model based method. Train a XGBoost model to predict the speed of lowered IR and pick the next batch according to the prediction. -# -# You can choose the tuner according to the size of your space, your time budget and other factors. -# For example, if your space is very small (less than 1000), a gridsearch tuner or a -# random tuner is good enough. If your space is at the level of 10^9 (this is the space -# size of a conv2d operator on CUDA GPU), XGBoostTuner can explore more efficiently -# and find better configs. - -################################################################ -# Begin tuning -# ^^^^^^^^^^^^ -# Here we continue our matrix multiplication example. -# First we should create a tuning task. -# We can also inspect the initialized search space. -# In this case, for a 512x512 square matrix multiplication, the space size -# is 10x10=100 -N, L, M = 512, 512, 512 -task = autotvm.task.create("tutorial/matmul", args=(N, L, M, "float32"), target="llvm") -print(task.config_space) - -################################################################ -# Then we need to define how to measure the generated code and pick a tuner. -# Since our space is small, a random tuner is just okay. -# -# We only make 10 trials in this tutorial for demonstration. In practice, -# you can do more trials according to your time budget. -# We will log the tuning results into a log file. This file can be -# used to get the best config later. - -# logging config (for printing tuning log to the screen) -logging.getLogger("autotvm").setLevel(logging.DEBUG) -logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) - -# There are two steps for measuring a config: build and run. -# By default, we use all CPU cores to compile program. Then measure them sequentially. -# We measure 5 times and take average to reduce variance. -measure_option = autotvm.measure_option(builder="local", runner=autotvm.LocalRunner(number=5)) - -# Begin tuning with RandomTuner, log records to file `matmul.log` -# You can use alternatives like XGBTuner. -tuner = autotvm.tuner.RandomTuner(task) -tuner.tune( - n_trial=10, - measure_option=measure_option, - callbacks=[autotvm.callback.log_to_file("matmul.log")], -) - -######################################################################### -# Finally we apply history best from the cache file and check its correctness. -# We can call the function :code:`matmul` directly under the -# :any:`autotvm.apply_history_best` context. When we call this function, -# it will query the dispatch context with its argument and get the best config -# with the same argument. - -# apply history best from log file -with autotvm.apply_history_best("matmul.log"): - with tvm.target.Target("llvm"): - s, arg_bufs = matmul(N, L, M, "float32") - func = tvm.build(s, arg_bufs) - -# check correctness -a_np = np.random.uniform(size=(N, L)).astype(np.float32) -b_np = np.random.uniform(size=(L, M)).astype(np.float32) -c_np = a_np.dot(b_np) - -c_tvm = tvm.nd.empty(c_np.shape) -func(tvm.nd.array(a_np), tvm.nd.array(b_np), c_tvm) - -tvm.testing.assert_allclose(c_np, c_tvm.asnumpy(), rtol=1e-2) +if __name__ == "__main__": + ###################################################################### + # .. note:: More Explanation on :code:`cfg.defile_split` + # + # In this template, :code:`cfg.define_split("tile_y", y, num_outputs=2)` will enumerate + # all possible combinations that can split axis y into two axes with factors of the length of y. + # For example, if the length of y is 32 and we want to split it into two axes + # using factors of 32, then there are 6 possible values for + # (length of outer axis, length of inner axis) pair, namely + # (32, 1), (16, 2), (8, 4), (4, 8), (2, 16) or (1, 32). + # They are just the 6 possible values of `tile_y`. + # + # During schedule, :code:`cfg["tile_y"]` is a :code:`SplitEntity` object. + # We stores the lengths of outer axes and inner axes in :code:`cfg['tile_y'].size` + # (a tuple with two elements). + # In this template, we apply it by using :code:`yo, yi = cfg['tile_y'].apply(s, C, y)`. + # Actually, this is equivalent to + # :code:`yo, yi = s[C].split(y, cfg["tile_y"].size[1])` + # or :code:`yo, yi = s[C].split(y, nparts=cfg['tile_y"].size[0])` + # + # The advantage of using cfg.apply API is that it makes multi-level split + # (when num_outputs >= 3) easier. + + ###################################################################### + # Step 2: Search through the space + # --------------------------------- + # In step 1, we build the search space by extending our old schedule code + # into a template. The next step is to pick a tuner and explore in this space. + # + # Auto-tuners in TVM + # ^^^^^^^^^^^^^^^^^^ + # The job for a tuner can be described by following pseudo code + # + # .. code-block:: c + # + # ct = 0 + # while ct < max_number_of_trials: + # propose a batch of configs + # measure this batch of configs on real hardware and get results + # ct += batch_size + # + # When proposing the next batch of configs, the tuner can take different strategies. We + # provide four tuners with different strategies in autotvm. + # + # * :any:`RandomTuner`: Enumerate the space in a random order + # * :any:`GridSearchTuner`: Enumerate the space in a grid search order + # * :any:`GATuner`: Using genetic algorithm to search through the space + # * :any:`XGBTuner`: Uses a model based method. Train a XGBoost model to predict the speed of lowered IR and pick the next batch according to the prediction. + # + # You can choose the tuner according to the size of your space, your time budget and other factors. + # For example, if your space is very small (less than 1000), a gridsearch tuner or a + # random tuner is good enough. If your space is at the level of 10^9 (this is the space + # size of a conv2d operator on CUDA GPU), XGBoostTuner can explore more efficiently + # and find better configs. + + ################################################################ + # Begin tuning + # ^^^^^^^^^^^^ + # Here we continue our matrix multiplication example. + # First we should create a tuning task. + # We can also inspect the initialized search space. + # In this case, for a 512x512 square matrix multiplication, the space size + # is 10x10=100 + N, L, M = 512, 512, 512 + task = autotvm.task.create("tutorial/matmul", args=(N, L, M, "float32"), target="llvm") + print(task.config_space) + + ################################################################ + # Then we need to define how to measure the generated code and pick a tuner. + # Since our space is small, a random tuner is just okay. + # + # We only make 10 trials in this tutorial for demonstration. In practice, + # you can do more trials according to your time budget. + # We will log the tuning results into a log file. This file can be + # used to get the best config later. + + # logging config (for printing tuning log to the screen) + logging.getLogger("autotvm").setLevel(logging.DEBUG) + logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) + + # There are two steps for measuring a config: build and run. + # By default, we use all CPU cores to compile program. Then measure them sequentially. + # We measure 5 times and take average to reduce variance. + measure_option = autotvm.measure_option(builder="local", runner=autotvm.LocalRunner(number=5)) + + # Begin tuning with RandomTuner, log records to file `matmul.log` + # You can use alternatives like XGBTuner. + tuner = autotvm.tuner.RandomTuner(task) + tuner.tune( + n_trial=10, + measure_option=measure_option, + callbacks=[autotvm.callback.log_to_file("matmul.log")], + ) + + ######################################################################### + # Finally we apply history best from the cache file and check its correctness. + # We can call the function :code:`matmul` directly under the + # :any:`autotvm.apply_history_best` context. When we call this function, + # it will query the dispatch context with its argument and get the best config + # with the same argument. + + # apply history best from log file + with autotvm.apply_history_best("matmul.log"): + with tvm.target.Target("llvm"): + s, arg_bufs = matmul(N, L, M, "float32") + func = tvm.build(s, arg_bufs) + + # check correctness + a_np = np.random.uniform(size=(N, L)).astype(np.float32) + b_np = np.random.uniform(size=(L, M)).astype(np.float32) + c_np = a_np.dot(b_np) + + c_tvm = tvm.nd.empty(c_np.shape) + func(tvm.nd.array(a_np), tvm.nd.array(b_np), c_tvm) + + tvm.testing.assert_allclose(c_np, c_tvm.asnumpy(), rtol=1e-2) From 4af275f04a752d778f05dccc75b1efaf92c90f08 Mon Sep 17 00:00:00 2001 From: Tristan Konolige Date: Tue, 27 Oct 2020 15:19:47 -0600 Subject: [PATCH 4/8] formatting --- tutorials/auto_scheduler/tune_conv2d_layer_cuda.py | 6 ++++-- tutorials/auto_scheduler/tune_matmul_x86.py | 3 --- tutorials/autotvm/tune_conv2d_cuda.py | 4 +++- tutorials/autotvm/tune_relay_arm.py | 3 --- tutorials/autotvm/tune_relay_cuda.py | 4 ---- tutorials/autotvm/tune_relay_mobile_gpu.py | 3 --- tutorials/autotvm/tune_relay_x86.py | 6 ------ 7 files changed, 7 insertions(+), 22 deletions(-) diff --git a/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py b/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py index 157bf9e76728..e2db214af388 100644 --- a/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py +++ b/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py @@ -56,6 +56,7 @@ def conv2d_layer(N, H, W, CO, CI, KH, KW, stride, padding): out = topi.nn.relu(conv + bias) return [data, kernel, bias, out] + if __name__ == "__main__": ###################################################################### # Create the search task @@ -66,7 +67,9 @@ def conv2d_layer(N, H, W, CO, CI, KH, KW, stride, padding): # Use the last layer in ResNet-50 N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) - task = auto_scheduler.create_task(conv2d_layer, (N, H, W, CO, CI, KH, KW, strides, padding), target) + task = auto_scheduler.create_task( + conv2d_layer, (N, H, W, CO, CI, KH, KW, strides, padding), target + ) # Inspect the computational graph print(task.compute_dag) @@ -178,7 +181,6 @@ def conv2d_layer(N, H, W, CO, CI, KH, KW, stride, padding): # and resume the status of search policy and cost model with the log file. # In the example below we resume the status and do more 5 trials. - cost_model = auto_scheduler.XGBModel() cost_model.update_from_file(log_file) search_policy = auto_scheduler.SketchPolicy( diff --git a/tutorials/auto_scheduler/tune_matmul_x86.py b/tutorials/auto_scheduler/tune_matmul_x86.py index 327b2805125d..424867d1b7bd 100644 --- a/tutorials/auto_scheduler/tune_matmul_x86.py +++ b/tutorials/auto_scheduler/tune_matmul_x86.py @@ -132,7 +132,6 @@ def matmul_add(N, L, M, dtype): % (np.median(evaluator(a_tvm, b_tvm, c_tvm, out_tvm).results) * 1000) ) - ###################################################################### # Using the record file # ^^^^^^^^^^^^^^^^^^^^^ @@ -163,7 +162,6 @@ def matmul_add(N, L, M, dtype): # and resume the status of search policy and cost model with the log file. # In the example below we resume the status and do more 5 trials. - def resume_search(task, log_file_name): cost_model = auto_scheduler.XGBModel() cost_model.update_from_file(log_file_name) @@ -177,7 +175,6 @@ def resume_search(task, log_file_name): ) sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) - # resume_search(task, log_file) ###################################################################### diff --git a/tutorials/autotvm/tune_conv2d_cuda.py b/tutorials/autotvm/tune_conv2d_cuda.py index 5aa3a34d4092..67809e5074f4 100644 --- a/tutorials/autotvm/tune_conv2d_cuda.py +++ b/tutorials/autotvm/tune_conv2d_cuda.py @@ -186,7 +186,9 @@ def conv2d_no_batching(N, H, W, CO, CI, KH, KW, stride, padding): # the last layer in resnet N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) task = autotvm.task.create( - "tutorial/conv2d_no_batching", args=(N, H, W, CO, CI, KH, KW, strides, padding), target="cuda" + "tutorial/conv2d_no_batching", + args=(N, H, W, CO, CI, KH, KW, strides, padding), + target="cuda", ) print(task.config_space) diff --git a/tutorials/autotvm/tune_relay_arm.py b/tutorials/autotvm/tune_relay_arm.py index 97950d11b669..0197323704ca 100644 --- a/tutorials/autotvm/tune_relay_arm.py +++ b/tutorials/autotvm/tune_relay_arm.py @@ -306,11 +306,9 @@ def tune_tasks( autotvm.record.pick_best(tmp_log_file, log_filename) os.remove(tmp_log_file) - ######################################################################## # Finally, we launch tuning jobs and evaluate the end-to-end performance. - def tune_and_evaluate(tuning_opt): # extract workloads from relay program print("Extract tasks...") @@ -361,7 +359,6 @@ def tune_and_evaluate(tuning_opt): % (np.mean(prof_res), np.std(prof_res)) ) - # We do not run the tuning in our webpage server since it takes too long. # Uncomment the following line to run it by yourself. diff --git a/tutorials/autotvm/tune_relay_cuda.py b/tutorials/autotvm/tune_relay_cuda.py index 0bc270816094..221f286a173a 100644 --- a/tutorials/autotvm/tune_relay_cuda.py +++ b/tutorials/autotvm/tune_relay_cuda.py @@ -213,11 +213,9 @@ def tune_tasks( autotvm.record.pick_best(tmp_log_file, log_filename) os.remove(tmp_log_file) - ######################################################################## # Finally, we launch tuning jobs and evaluate the end-to-end performance. - def tune_and_evaluate(tuning_opt): # extract workloads from relay program print("Extract tasks...") @@ -256,7 +254,6 @@ def tune_and_evaluate(tuning_opt): % (np.mean(prof_res), np.std(prof_res)) ) - # We do not run the tuning in our webpage server since it takes too long. # Uncomment the following line to run it by yourself. @@ -316,7 +313,6 @@ def tune_and_evaluate(tuning_opt): # # Finally, always feel free to ask our community for help on https://discuss.tvm.ai - ################################################################# # Scale up measurement by using multiple devices # ---------------------------------------------- diff --git a/tutorials/autotvm/tune_relay_mobile_gpu.py b/tutorials/autotvm/tune_relay_mobile_gpu.py index 5e04569e14dc..1b521cbfaa0b 100644 --- a/tutorials/autotvm/tune_relay_mobile_gpu.py +++ b/tutorials/autotvm/tune_relay_mobile_gpu.py @@ -301,11 +301,9 @@ def tune_tasks( autotvm.record.pick_best(tmp_log_file, log_filename) os.remove(tmp_log_file) - ######################################################################## # Finally, we launch tuning jobs and evaluate the end-to-end performance. - def tune_and_evaluate(tuning_opt): # extract workloads from relay program print("Extract tasks...") @@ -361,7 +359,6 @@ def tune_and_evaluate(tuning_opt): % (np.mean(prof_res), np.std(prof_res)) ) - # We do not run the tuning in our webpage server since it takes too long. # Uncomment the following line to run it by yourself. diff --git a/tutorials/autotvm/tune_relay_x86.py b/tutorials/autotvm/tune_relay_x86.py index 2e990efdfa87..8c986634caba 100644 --- a/tutorials/autotvm/tune_relay_x86.py +++ b/tutorials/autotvm/tune_relay_x86.py @@ -111,7 +111,6 @@ def get_network(name, batch_size): num_threads = 1 os.environ["TVM_NUM_THREADS"] = str(num_threads) - ################################################################# # Configure tensor tuning settings and create tasks # ------------------------------------------------- @@ -142,7 +141,6 @@ def get_network(name, batch_size): ), } - # You can skip the implementation of this function for this tutorial. def tune_kernels( tasks, measure_option, tuner="gridsearch", early_stopping=None, log_filename="tuning.log" @@ -175,7 +173,6 @@ def tune_kernels( ], ) - # Use graph tuner to achieve graph level optimal schedules # Set use_DP=False if it takes too long to finish. def tune_graph(graph, dshape, records, opt_sch_file, use_DP=True): @@ -188,11 +185,9 @@ def tune_graph(graph, dshape, records, opt_sch_file, use_DP=True): executor.run() executor.write_opt_sch2record_file(opt_sch_file) - ######################################################################## # Finally, we launch tuning jobs and evaluate the end-to-end performance. - def tune_and_evaluate(tuning_opt): # extract workloads from relay program print("Extract tasks...") @@ -226,7 +221,6 @@ def tune_and_evaluate(tuning_opt): % (np.mean(prof_res), np.std(prof_res)) ) - # We do not run the tuning in our webpage server since it takes too long. # Uncomment the following line to run it by yourself. From e9ecb16a10d4f6f932877ab81b65b006868ef940 Mon Sep 17 00:00:00 2001 From: Tristan Konolige Date: Tue, 27 Oct 2020 16:32:11 -0700 Subject: [PATCH 5/8] undo autotvm work --- python/tvm/autotvm/measure/local_executor.py | 10 +++---- python/tvm/autotvm/task/task.py | 14 ++++----- .../tvm/autotvm/tuner/xgboost_cost_model.py | 29 +++++++++++-------- 3 files changed, 26 insertions(+), 27 deletions(-) diff --git a/python/tvm/autotvm/measure/local_executor.py b/python/tvm/autotvm/measure/local_executor.py index af1dc1213292..5dd5cba2b824 100644 --- a/python/tvm/autotvm/measure/local_executor.py +++ b/python/tvm/autotvm/measure/local_executor.py @@ -18,7 +18,7 @@ import signal -import multiprocessing +from multiprocessing import Process, Queue try: from queue import Empty @@ -60,7 +60,7 @@ def call_with_timeout(queue, timeout, func, args, kwargs): """A wrapper to support timeout of a function call""" # start a new process for timeout (cannot use thread because we have c function) - p = multiprocessing.Process(target=_execute_func, args=(func, queue, args, kwargs)) + p = Process(target=_execute_func, args=(func, queue, args, kwargs)) p.start() p.join(timeout=timeout) @@ -151,9 +151,7 @@ def submit(self, func, *args, **kwargs): if not self.do_fork: return LocalFutureNoFork(func(*args, **kwargs)) - queue = multiprocessing.Queue(2) # Size of 2 to avoid a race condition with size 1. - process = multiprocessing.Process( - target=call_with_timeout, args=(queue, self.timeout, func, args, kwargs) - ) + queue = Queue(2) # Size of 2 to avoid a race condition with size 1. + process = Process(target=call_with_timeout, args=(queue, self.timeout, func, args, kwargs)) process.start() return LocalFuture(process, queue) diff --git a/python/tvm/autotvm/task/task.py b/python/tvm/autotvm/task/task.py index 8822ba971e4c..a7cb9a095765 100644 --- a/python/tvm/autotvm/task/task.py +++ b/python/tvm/autotvm/task/task.py @@ -23,14 +23,15 @@ """ import numpy as np +from tvm.target import Target from tvm import runtime from tvm.ir import container -from tvm.target import Target -from tvm.te import placeholder, tensor from tvm.tir import expr +from tvm.te import tensor, placeholder + from ..util import get_const_int, get_const_tuple -from .dispatcher import ApplyConfig, DispatchContext +from .dispatcher import DispatchContext, ApplyConfig from .space import ConfigSpace @@ -172,8 +173,6 @@ def __getstate__(self): # some unpickable local task functions. # So we only pickle the name of the function # and restore the function by name when unpickling it. - import cloudpickle # pylint: disable=import-outside-toplevel - return { "name": self.name, "args": self.args, @@ -182,17 +181,14 @@ def __getstate__(self): "flop": self.flop, "target": self.target, "target_host": self.target_host, - "func": cloudpickle.dumps(self.func), } def __setstate__(self, state): - import cloudpickle # pylint: disable=import-outside-toplevel - self.name = state["name"] self.args = state["args"] self.kwargs = state["kwargs"] self.config_space = state["config_space"] - self.func = cloudpickle.loads(state["func"]) + self.func = _lookup_task(state["name"]) self.flop = state["flop"] self.target = state["target"] self.target_host = state["target_host"] diff --git a/python/tvm/autotvm/tuner/xgboost_cost_model.py b/python/tvm/autotvm/tuner/xgboost_cost_model.py index f2e6eb1a2e44..7b9df1c99373 100644 --- a/python/tvm/autotvm/tuner/xgboost_cost_model.py +++ b/python/tvm/autotvm/tuner/xgboost_cost_model.py @@ -153,6 +153,11 @@ def _reset_pool(self, space, target, task): self._close_pool() + # use global variable to pass common arguments + global _extract_space, _extract_target, _extract_task + _extract_space = space + _extract_target = target + _extract_task = task self.pool = multiprocessing.Pool(self.num_threads) def _close_pool(self): @@ -316,11 +321,10 @@ def _get_feature(self, indexes): indexes = np.array(indexes) need_extract = [x for x in indexes if x not in fea_cache] - args = [(self.space.get(x), self.target, self.task) for x in need_extract] if need_extract: pool = self._get_pool() - feas = pool.map(self.feature_extract_func, args) + feas = pool.map(self.feature_extract_func, need_extract) for i, fea in zip(need_extract, feas): fea_cache[i] = fea @@ -340,16 +344,17 @@ def __del__(self): self._close_pool() +_extract_space = None _extract_target = None _extract_task = None -def _extract_itervar_feature_index(args): +def _extract_itervar_feature_index(index): """extract iteration var feature for an index in extract_space""" try: - config, target, task = args - with target: - sch, args = task.instantiate(config) + config = _extract_space.get(index) + with _extract_target: + sch, args = _extract_task.instantiate(config) fea = feature.get_itervar_feature_flatten(sch, args, take_log=True) fea = np.concatenate((fea, list(config.get_other_option().values()))) return fea @@ -376,10 +381,10 @@ def _extract_itervar_feature_log(arg): return None -def _extract_knob_feature_index(args): +def _extract_knob_feature_index(index): """extract knob feature for an index in extract_space""" try: - config, _, _ = args + config = _extract_space.get(index) return config.get_flatten_feature() except Exception: # pylint: disable=broad-except return None @@ -403,12 +408,12 @@ def _extract_knob_feature_log(arg): return None -def _extract_curve_feature_index(args): +def _extract_curve_feature_index(index): """extract sampled curve feature for an index in extract_space""" try: - config, target, task = args - with target: - sch, args = task.instantiate(config) + config = _extract_space.get(index) + with _extract_target: + sch, args = _extract_task.instantiate(config) fea = feature.get_buffer_curve_sample_flatten(sch, args, sample_n=20) fea = np.concatenate((fea, list(config.get_other_option().values()))) return np.array(fea) From df0b9b8cba108d7376a5d9a8c069a4358e185659 Mon Sep 17 00:00:00 2001 From: Tristan Konolige Date: Wed, 28 Oct 2020 10:11:06 -0700 Subject: [PATCH 6/8] Undo tutorial changes --- .../auto_scheduler/tune_conv2d_layer_cuda.py | 276 +++++---- tutorials/auto_scheduler/tune_matmul_x86.py | 270 ++++----- tutorials/autotvm/tune_conv2d_cuda.py | 144 +++-- tutorials/autotvm/tune_relay_arm.py | 434 +++++++------- tutorials/autotvm/tune_relay_cuda.py | 531 +++++++++--------- tutorials/autotvm/tune_relay_mobile_gpu.py | 442 +++++++-------- tutorials/autotvm/tune_relay_x86.py | 327 +++++------ tutorials/autotvm/tune_simple_template.py | 232 ++++---- 8 files changed, 1331 insertions(+), 1325 deletions(-) diff --git a/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py b/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py index e2db214af388..10a2d1b44144 100644 --- a/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py +++ b/tutorials/auto_scheduler/tune_conv2d_layer_cuda.py @@ -57,142 +57,140 @@ def conv2d_layer(N, H, W, CO, CI, KH, KW, stride, padding): return [data, kernel, bias, out] -if __name__ == "__main__": - ###################################################################### - # Create the search task - # ^^^^^^^^^^^^^^^^^^^^^^ - # We then create a search task for the last convolution layer in the resnet. - - target = tvm.target.Target("cuda") - - # Use the last layer in ResNet-50 - N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) - task = auto_scheduler.create_task( - conv2d_layer, (N, H, W, CO, CI, KH, KW, strides, padding), target - ) - - # Inspect the computational graph - print(task.compute_dag) - - ###################################################################### - # Next, we set parameters for the auto-scheduler. These parameters - # mainly specify how we do the measurement during the search and auto-tuning. - # - # * :code:`measure_ctx` launches a different process for measurement. This - # provides an isolation. It can protect the master process from GPU crashes - # happended during measurement and avoid other runtime conflicts. - # * :code:`min_repeat_ms` defines the minimum duration of one "repeat" in every measurement. - # This can warmup the GPU, which is necessary to get accurate measurement results. - # Typically, we recommend a value > 300 ms. - # * :code:`num_measure_trials` is the number of measurement trials we can use during the search. - # We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a - # good value for the search to converge. You can do more trials according to your time budget. - # * In addition, we use :code:`RecordToFile` to dump measurement records into a file `conv2d.json`. - # The measurement records can be used to query the history best, resume the search, - # and do more analyses later. - # * see :any:`auto_scheduler.TuningOptions`, - # :any:`auto_scheduler.LocalRPCMeasureContext` for more parameters. - - log_file = "conv2d.json" - measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300) - tune_option = auto_scheduler.TuningOptions( - num_measure_trials=10, - runner=measure_ctx.runner, - measure_callbacks=[auto_scheduler.RecordToFile(log_file)], - ) - - ###################################################################### - # Run the search - # ^^^^^^^^^^^^^^ - # Now we get all inputs ready. Pretty simple, isn't it? - # We can kick off the search and let the auto-scheduler do its magic. - # After some measurement trials, it will return the best schedule it found. - - sch, args = auto_scheduler.auto_schedule(task, tuning_options=tune_option) - - # Kill the process for measurement - del measure_ctx - - ###################################################################### - # We can lower the schedule to see the IR after auto-scheduling. - # The auto-scheduler correctly performs optimizations including multi-level tiling, - # cooperative fetching, unrolling and operator fusion. - - print(tvm.lower(sch, args, simple_mode=True)) - - ###################################################################### - # Check correctness and evaluate performance - # ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ - # We build the binary and check its correctness and performance. - - func = tvm.build(sch, args, target) - - # Check correctness - data_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) - weight_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) - bias_np = np.random.uniform(size=(1, CO, 1, 1)).astype(np.float32) - conv_np = conv2d_nchw_python(data_np, weight_np, strides, padding) - out_np = np.maximum(conv_np + bias_np, 0.0) - - ctx = tvm.gpu() - data_tvm = tvm.nd.array(data_np, ctx=ctx) - weight_tvm = tvm.nd.array(weight_np, ctx=ctx) - bias_tvm = tvm.nd.array(bias_np, ctx=ctx) - out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx) - func(data_tvm, weight_tvm, bias_tvm, out_tvm) - - # Check results - np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3) - - # Evaluate execution time - evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500) - print( - "Execution time of this operator: %.3f ms" - % (np.median(evaluator(data_tvm, weight_tvm, bias_tvm, out_tvm).results) * 1000) - ) - - ###################################################################### - # Using the record file - # ^^^^^^^^^^^^^^^^^^^^^ - # During the search, all measuremnt records are dumpped into the record - # file "conv2d.json". The measurement records can be used to re-apply search results, - # resume the search, and perform other analyses. - - ###################################################################### - # Here is an example where we load the best schedule from a file, - # print the equivalent python schedule API, and build the binary again. - - # Load the measuremnt record for the best schedule - inp, res = auto_scheduler.load_best(log_file, task.workload_key) - - # Print equivalent python schedule API. This can be used for debugging and - # learning the behavior of the auto-scheduler. - print("Equivalent python schedule:") - print(task.compute_dag.print_python_code_from_state(inp.state)) - - # Rebuild the binary. This shows how you can apply the best schedule from a - # log file without reruning the search again. - sch, args = task.compute_dag.apply_steps_from_state(inp.state) - func = tvm.build(sch, args, target) - - ###################################################################### - # A more complicated example is to resume the search. - # In this case, we need to create the search policy and cost model by ourselves - # and resume the status of search policy and cost model with the log file. - # In the example below we resume the status and do more 5 trials. - - cost_model = auto_scheduler.XGBModel() - cost_model.update_from_file(log_file) - search_policy = auto_scheduler.SketchPolicy( - task, cost_model, init_search_callbacks=[auto_scheduler.PreloadMeasuredStates(log_file)] - ) - measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300) - tune_option = auto_scheduler.TuningOptions( - num_measure_trials=5, - runner=measure_ctx.runner, - measure_callbacks=[auto_scheduler.RecordToFile(log_file)], - ) - sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) - - # Kill the measurement process - del measure_ctx +###################################################################### +# Create the search task +# ^^^^^^^^^^^^^^^^^^^^^^ +# We then create a search task for the last convolution layer in the resnet. + +target = tvm.target.Target("cuda") + +# Use the last layer in ResNet-50 +N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) +task = auto_scheduler.create_task(conv2d_layer, (N, H, W, CO, CI, KH, KW, strides, padding), target) + +# Inspect the computational graph +print(task.compute_dag) + +###################################################################### +# Next, we set parameters for the auto-scheduler. These parameters +# mainly specify how we do the measurement during the search and auto-tuning. +# +# * :code:`measure_ctx` launches a different process for measurement. This +# provides an isolation. It can protect the master process from GPU crashes +# happended during measurement and avoid other runtime conflicts. +# * :code:`min_repeat_ms` defines the minimum duration of one "repeat" in every measurement. +# This can warmup the GPU, which is necessary to get accurate measurement results. +# Typically, we recommend a value > 300 ms. +# * :code:`num_measure_trials` is the number of measurement trials we can use during the search. +# We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a +# good value for the search to converge. You can do more trials according to your time budget. +# * In addition, we use :code:`RecordToFile` to dump measurement records into a file `conv2d.json`. +# The measurement records can be used to query the history best, resume the search, +# and do more analyses later. +# * see :any:`auto_scheduler.TuningOptions`, +# :any:`auto_scheduler.LocalRPCMeasureContext` for more parameters. + +log_file = "conv2d.json" +measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300) +tune_option = auto_scheduler.TuningOptions( + num_measure_trials=10, + runner=measure_ctx.runner, + measure_callbacks=[auto_scheduler.RecordToFile(log_file)], +) + +###################################################################### +# Run the search +# ^^^^^^^^^^^^^^ +# Now we get all inputs ready. Pretty simple, isn't it? +# We can kick off the search and let the auto-scheduler do its magic. +# After some measurement trials, it will return the best schedule it found. + +sch, args = auto_scheduler.auto_schedule(task, tuning_options=tune_option) + +# Kill the process for measurement +del measure_ctx + +###################################################################### +# We can lower the schedule to see the IR after auto-scheduling. +# The auto-scheduler correctly performs optimizations including multi-level tiling, +# cooperative fetching, unrolling and operator fusion. + +print(tvm.lower(sch, args, simple_mode=True)) + +###################################################################### +# Check correctness and evaluate performance +# ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +# We build the binary and check its correctness and performance. + +func = tvm.build(sch, args, target) + +# Check correctness +data_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) +weight_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) +bias_np = np.random.uniform(size=(1, CO, 1, 1)).astype(np.float32) +conv_np = conv2d_nchw_python(data_np, weight_np, strides, padding) +out_np = np.maximum(conv_np + bias_np, 0.0) + +ctx = tvm.gpu() +data_tvm = tvm.nd.array(data_np, ctx=ctx) +weight_tvm = tvm.nd.array(weight_np, ctx=ctx) +bias_tvm = tvm.nd.array(bias_np, ctx=ctx) +out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx) +func(data_tvm, weight_tvm, bias_tvm, out_tvm) + +# Check results +np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3) + +# Evaluate execution time +evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500) +print( + "Execution time of this operator: %.3f ms" + % (np.median(evaluator(data_tvm, weight_tvm, bias_tvm, out_tvm).results) * 1000) +) + +###################################################################### +# Using the record file +# ^^^^^^^^^^^^^^^^^^^^^ +# During the search, all measuremnt records are dumpped into the record +# file "conv2d.json". The measurement records can be used to re-apply search results, +# resume the search, and perform other analyses. + +###################################################################### +# Here is an example where we load the best schedule from a file, +# print the equivalent python schedule API, and build the binary again. + +# Load the measuremnt record for the best schedule +inp, res = auto_scheduler.load_best(log_file, task.workload_key) + +# Print equivalent python schedule API. This can be used for debugging and +# learning the behavior of the auto-scheduler. +print("Equivalent python schedule:") +print(task.compute_dag.print_python_code_from_state(inp.state)) + +# Rebuild the binary. This shows how you can apply the best schedule from a +# log file without reruning the search again. +sch, args = task.compute_dag.apply_steps_from_state(inp.state) +func = tvm.build(sch, args, target) + +###################################################################### +# A more complicated example is to resume the search. +# In this case, we need to create the search policy and cost model by ourselves +# and resume the status of search policy and cost model with the log file. +# In the example below we resume the status and do more 5 trials. + + +cost_model = auto_scheduler.XGBModel() +cost_model.update_from_file(log_file) +search_policy = auto_scheduler.SketchPolicy( + task, cost_model, init_search_callbacks=[auto_scheduler.PreloadMeasuredStates(log_file)] +) +measure_ctx = auto_scheduler.LocalRPCMeasureContext(min_repeat_ms=300) +tune_option = auto_scheduler.TuningOptions( + num_measure_trials=5, + runner=measure_ctx.runner, + measure_callbacks=[auto_scheduler.RecordToFile(log_file)], +) +sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) + +# Kill the measurement process +del measure_ctx diff --git a/tutorials/auto_scheduler/tune_matmul_x86.py b/tutorials/auto_scheduler/tune_matmul_x86.py index 424867d1b7bd..81f2e71ff8f7 100644 --- a/tutorials/auto_scheduler/tune_matmul_x86.py +++ b/tutorials/auto_scheduler/tune_matmul_x86.py @@ -56,141 +56,143 @@ def matmul_add(N, L, M, dtype): return [A, B, C, out] -if __name__ == "__main__": - ###################################################################### - # Create the search task - # ^^^^^^^^^^^^^^^^^^^^^^ - # We then create a search task with N=L=M=128 and dtype="float32" - # If your machine supports avx instructions, you can - # - # - replace "llvm" below with "llvm -mcpu=core-avx2" to enable AVX2 - # - replace "llvm" below with "llvm -mcpu=skylake-avx512" to enable AVX-512 - - target = tvm.target.Target("llvm") - task = tvm.auto_scheduler.create_task(matmul_add, (128, 128, 128, "float32"), target) - - # Inspect the computational graph - print(task.compute_dag) - - ###################################################################### - # Next, we set parameters for the auto-scheduler. - # - # * :code:`num_measure_trials` is the number of measurement trials we can use during the search. - # We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a - # good value for the search to converge. You can do more trials according to your time budget. - # * In addition, we use :code:`RecordToFile` to dump measurement records into a file `matmul.json`. - # The measurement records can be used to query the history best, resume the search, - # and do more analyses later. - # * see :any:`auto_scheduler.TuningOptions` for more parameters - - log_file = "matmul.json" +###################################################################### +# Create the search task +# ^^^^^^^^^^^^^^^^^^^^^^ +# We then create a search task with N=L=M=128 and dtype="float32" +# If your machine supports avx instructions, you can +# +# - replace "llvm" below with "llvm -mcpu=core-avx2" to enable AVX2 +# - replace "llvm" below with "llvm -mcpu=skylake-avx512" to enable AVX-512 + +target = tvm.target.Target("llvm") +task = tvm.auto_scheduler.create_task(matmul_add, (128, 128, 128, "float32"), target) + +# Inspect the computational graph +print(task.compute_dag) + +###################################################################### +# Next, we set parameters for the auto-scheduler. +# +# * :code:`num_measure_trials` is the number of measurement trials we can use during the search. +# We only make 10 trials in this tutorial for a fast demonstration. In practice, 1000 is a +# good value for the search to converge. You can do more trials according to your time budget. +# * In addition, we use :code:`RecordToFile` to dump measurement records into a file `matmul.json`. +# The measurement records can be used to query the history best, resume the search, +# and do more analyses later. +# * see :any:`auto_scheduler.TuningOptions` for more parameters + +log_file = "matmul.json" +tune_option = auto_scheduler.TuningOptions( + num_measure_trials=10, measure_callbacks=[auto_scheduler.RecordToFile(log_file)] +) + +###################################################################### +# Run the search +# ^^^^^^^^^^^^^^ +# Now we get all inputs ready. Pretty simple, isn't it? +# We can kick off the search and let the auto-scheduler do its magic. +# After some measurement trials, it will return the best schedule it found. + +sch, args = auto_scheduler.auto_schedule(task, tuning_options=tune_option) + +###################################################################### +# We can lower the schedule to see the IR after auto-scheduling. +# The auto-scheduler correctly performs optimizations including multi-level tiling, +# parallelization, vectorization, unrolling and operator fusion. + +print(tvm.lower(sch, args, simple_mode=True)) + +###################################################################### +# Check correctness and evaluate performance +# ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +# We build the binary and check its correctness and performance. + +func = tvm.build(sch, args) +a_np = np.random.uniform(size=(128, 128)).astype(np.float32) +b_np = np.random.uniform(size=(128, 128)).astype(np.float32) +c_np = np.random.uniform(size=(128, 128)).astype(np.float32) +out_np = a_np.dot(b_np) + c_np + +ctx = tvm.cpu() +a_tvm = tvm.nd.array(a_np, ctx=ctx) +b_tvm = tvm.nd.array(b_np, ctx=ctx) +c_tvm = tvm.nd.array(c_np, ctx=ctx) +out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx) +func(a_tvm, b_tvm, c_tvm, out_tvm) + +# Check results +np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3) + +# Evaluate execution time. +evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500) +print( + "Execution time of this operator: %.3f ms" + % (np.median(evaluator(a_tvm, b_tvm, c_tvm, out_tvm).results) * 1000) +) + + +###################################################################### +# Using the record file +# ^^^^^^^^^^^^^^^^^^^^^ +# During the search, all measuremnt records are dumpped into the record +# file "matmul.json". The measurement records can be used to re-apply search results, +# resume the search, and perform other analyses. + +###################################################################### +# Here is an example where we load the best schedule from a file, +# print the equivalent python schedule API, and build the binary again. + +# Load the measuremnt record for the best schedule +inp, res = auto_scheduler.load_best(log_file, task.workload_key) + +# Print equivalent python schedule API. This can be used for debugging and +# learning the behavior of the auto-scheduler. +print("Equivalent python schedule:") +print(task.compute_dag.print_python_code_from_state(inp.state)) + +# Rebuild the binary. This shows how you can apply the best schedule from a +# log file without reruning the search again. +sch, args = task.compute_dag.apply_steps_from_state(inp.state) +func = tvm.build(sch, args) + +###################################################################### +# A more complicated example is to resume the search. +# In this case, we need to create the search policy and cost model by ourselves +# and resume the status of search policy and cost model with the log file. +# In the example below we resume the status and do more 5 trials. + + +def resume_search(task, log_file_name): + cost_model = auto_scheduler.XGBModel() + cost_model.update_from_file(log_file_name) + search_policy = auto_scheduler.SketchPolicy( + task, + cost_model, + init_search_callbacks=[auto_scheduler.PreloadMeasuredStates(log_file_name)], + ) tune_option = auto_scheduler.TuningOptions( - num_measure_trials=10, measure_callbacks=[auto_scheduler.RecordToFile(log_file)] + num_measure_trials=5, measure_callbacks=[auto_scheduler.RecordToFile(log_file_name)] ) + sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) - ###################################################################### - # Run the search - # ^^^^^^^^^^^^^^ - # Now we get all inputs ready. Pretty simple, isn't it? - # We can kick off the search and let the auto-scheduler do its magic. - # After some measurement trials, it will return the best schedule it found. - - sch, args = auto_scheduler.auto_schedule(task, tuning_options=tune_option) - - ###################################################################### - # We can lower the schedule to see the IR after auto-scheduling. - # The auto-scheduler correctly performs optimizations including multi-level tiling, - # parallelization, vectorization, unrolling and operator fusion. - - print(tvm.lower(sch, args, simple_mode=True)) - - ###################################################################### - # Check correctness and evaluate performance - # ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ - # We build the binary and check its correctness and performance. - - func = tvm.build(sch, args) - a_np = np.random.uniform(size=(128, 128)).astype(np.float32) - b_np = np.random.uniform(size=(128, 128)).astype(np.float32) - c_np = np.random.uniform(size=(128, 128)).astype(np.float32) - out_np = a_np.dot(b_np) + c_np - - ctx = tvm.cpu() - a_tvm = tvm.nd.array(a_np, ctx=ctx) - b_tvm = tvm.nd.array(b_np, ctx=ctx) - c_tvm = tvm.nd.array(c_np, ctx=ctx) - out_tvm = tvm.nd.empty(out_np.shape, ctx=ctx) - func(a_tvm, b_tvm, c_tvm, out_tvm) - - # Check results - np.testing.assert_allclose(out_np, out_tvm.asnumpy(), rtol=1e-3) - - # Evaluate execution time. - evaluator = func.time_evaluator(func.entry_name, ctx, min_repeat_ms=500) - print( - "Execution time of this operator: %.3f ms" - % (np.median(evaluator(a_tvm, b_tvm, c_tvm, out_tvm).results) * 1000) - ) - ###################################################################### - # Using the record file - # ^^^^^^^^^^^^^^^^^^^^^ - # During the search, all measuremnt records are dumpped into the record - # file "matmul.json". The measurement records can be used to re-apply search results, - # resume the search, and perform other analyses. - - ###################################################################### - # Here is an example where we load the best schedule from a file, - # print the equivalent python schedule API, and build the binary again. - - # Load the measuremnt record for the best schedule - inp, res = auto_scheduler.load_best(log_file, task.workload_key) - - # Print equivalent python schedule API. This can be used for debugging and - # learning the behavior of the auto-scheduler. - print("Equivalent python schedule:") - print(task.compute_dag.print_python_code_from_state(inp.state)) - - # Rebuild the binary. This shows how you can apply the best schedule from a - # log file without reruning the search again. - sch, args = task.compute_dag.apply_steps_from_state(inp.state) - func = tvm.build(sch, args) - - ###################################################################### - # A more complicated example is to resume the search. - # In this case, we need to create the search policy and cost model by ourselves - # and resume the status of search policy and cost model with the log file. - # In the example below we resume the status and do more 5 trials. - - def resume_search(task, log_file_name): - cost_model = auto_scheduler.XGBModel() - cost_model.update_from_file(log_file_name) - search_policy = auto_scheduler.SketchPolicy( - task, - cost_model, - init_search_callbacks=[auto_scheduler.PreloadMeasuredStates(log_file_name)], - ) - tune_option = auto_scheduler.TuningOptions( - num_measure_trials=5, measure_callbacks=[auto_scheduler.RecordToFile(log_file_name)] - ) - sch, args = auto_scheduler.auto_schedule(task, search_policy, tuning_options=tune_option) - - # resume_search(task, log_file) - - ###################################################################### - # .. note:: - # We cannot run the line above because of the conflict between - # python's multiprocessing and tvm's thread pool. - # After running a tvm generated binary the python's multiprocessing library - # will hang forever. You have to make sure that you don't run any tvm - # generated binaries before calling auot-scheduler's search. - # To run the function above, you should comment out all code in - # "Check correctness and evaluate performance" section. - # - # You should be careful about this problem in your applications. - # There are other workarounds for this problem. - # For example, you can start a new thread/process (with the builtin python library - # threading or multiprocessing) and run the tvm binaries in the new thread/process. - # This provides an isolation and avoids the conflict in the main thread/process. - # You can also use :any:`auto_scheduler.LocalRPCMeasureContext` for auto-scheduler, - # as shown in the GPU tutorial (:ref:`auto-scheduler-conv-gpu`). +# resume_search(task, log_file) + +###################################################################### +# .. note:: +# We cannot run the line above because of the conflict between +# python's multiprocessing and tvm's thread pool. +# After running a tvm generated binary the python's multiprocessing library +# will hang forever. You have to make sure that you don't run any tvm +# generated binaries before calling auot-scheduler's search. +# To run the function above, you should comment out all code in +# "Check correctness and evaluate performance" section. +# +# You should be careful about this problem in your applications. +# There are other workarounds for this problem. +# For example, you can start a new thread/process (with the builtin python library +# threading or multiprocessing) and run the tvm binaries in the new thread/process. +# This provides an isolation and avoids the conflict in the main thread/process. +# You can also use :any:`auto_scheduler.LocalRPCMeasureContext` for auto-scheduler, +# as shown in the GPU tutorial (:ref:`auto-scheduler-conv-gpu`). diff --git a/tutorials/autotvm/tune_conv2d_cuda.py b/tutorials/autotvm/tune_conv2d_cuda.py index 67809e5074f4..ce9c19860ff4 100644 --- a/tutorials/autotvm/tune_conv2d_cuda.py +++ b/tutorials/autotvm/tune_conv2d_cuda.py @@ -50,7 +50,6 @@ import tvm from tvm import te -import tvm.testing from tvm import topi from tvm.topi.testing import conv2d_nchw_python @@ -169,76 +168,73 @@ def conv2d_no_batching(N, H, W, CO, CI, KH, KW, stride, padding): return s, [raw_data, kernel, conv] -if __name__ == "__main__": - ###################################################################### - # Step 2: Search through the space - # --------------------------------- - # We pick the last layer on resnet as test case. - # Since our space is very large, :code:`XGBoostTuner` is most suitable - # for our case. Here we only do 20 trials for demonstration. - # In practice, making 1000 trials usually can find some good kernels - # for this template - - # logging config (for printing tuning log to screen) - logging.getLogger("autotvm").setLevel(logging.DEBUG) - logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) - - # the last layer in resnet - N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) - task = autotvm.task.create( - "tutorial/conv2d_no_batching", - args=(N, H, W, CO, CI, KH, KW, strides, padding), - target="cuda", - ) - print(task.config_space) - - # Use local gpu, measure 10 times for every config to reduce variance - # The timeout of compiling a program is 10 seconds, the timeout for running is 4 seconds - measure_option = autotvm.measure_option( - builder=autotvm.LocalBuilder(), - runner=autotvm.LocalRunner(repeat=3, min_repeat_ms=100, timeout=4), - ) - - # Begin tuning, log records to file `conv2d.log` - # During tuning we will also try many invalid configs, so you are expected to - # see many error reports. As long as you can see non-zero GFLOPS, it is okay. - tuner = autotvm.tuner.XGBTuner(task) - tuner.tune( - n_trial=20, - measure_option=measure_option, - callbacks=[autotvm.callback.log_to_file("conv2d.log")], - ) - - ######################################################################### - # Finally we can inspect the best config from log file, check correctness, - # and measure running time. - - # inspect the best config - dispatch_context = autotvm.apply_history_best("conv2d.log") - best_config = dispatch_context.query(task.target, task.workload) - print("\nBest config:") - print(best_config) - - # apply history best from log file - with autotvm.apply_history_best("conv2d.log"): - with tvm.target.Target("cuda"): - s, arg_bufs = conv2d_no_batching(N, H, W, CO, CI, KH, KW, strides, padding) - func = tvm.build(s, arg_bufs) - - # check correctness - a_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) - w_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) - c_np = conv2d_nchw_python(a_np, w_np, strides, padding) - - ctx = tvm.gpu() - a_tvm = tvm.nd.array(a_np, ctx=ctx) - w_tvm = tvm.nd.array(w_np, ctx=ctx) - c_tvm = tvm.nd.empty(c_np.shape, ctx=ctx) - func(a_tvm, w_tvm, c_tvm) - - tvm.testing.assert_allclose(c_np, c_tvm.asnumpy(), rtol=1e-2) - - # Evaluate running time. Here we choose a large repeat number (400) to reduce the noise - # and the overhead of kernel launch. You can also use nvprof to validate the result. - evaluator = func.time_evaluator(func.entry_name, ctx, number=400) - print("Time cost of this operator: %f" % evaluator(a_tvm, w_tvm, c_tvm).mean) +###################################################################### +# Step 2: Search through the space +# --------------------------------- +# We pick the last layer on resnet as test case. +# Since our space is very large, :code:`XGBoostTuner` is most suitable +# for our case. Here we only do 20 trials for demonstration. +# In practice, making 1000 trials usually can find some good kernels +# for this template + +# logging config (for printing tuning log to screen) +logging.getLogger("autotvm").setLevel(logging.DEBUG) +logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) + +# the last layer in resnet +N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) +task = autotvm.task.create( + "tutorial/conv2d_no_batching", args=(N, H, W, CO, CI, KH, KW, strides, padding), target="cuda" +) +print(task.config_space) + +# Use local gpu, measure 10 times for every config to reduce variance +# The timeout of compiling a program is 10 seconds, the timeout for running is 4 seconds +measure_option = autotvm.measure_option( + builder=autotvm.LocalBuilder(), + runner=autotvm.LocalRunner(repeat=3, min_repeat_ms=100, timeout=4), +) + +# Begin tuning, log records to file `conv2d.log` +# During tuning we will also try many invalid configs, so you are expected to +# see many error reports. As long as you can see non-zero GFLOPS, it is okay. +tuner = autotvm.tuner.XGBTuner(task) +tuner.tune( + n_trial=20, + measure_option=measure_option, + callbacks=[autotvm.callback.log_to_file("conv2d.log")], +) + +######################################################################### +# Finally we can inspect the best config from log file, check correctness, +# and measure running time. + +# inspect the best config +dispatch_context = autotvm.apply_history_best("conv2d.log") +best_config = dispatch_context.query(task.target, task.workload) +print("\nBest config:") +print(best_config) + +# apply history best from log file +with autotvm.apply_history_best("conv2d.log"): + with tvm.target.Target("cuda"): + s, arg_bufs = conv2d_no_batching(N, H, W, CO, CI, KH, KW, strides, padding) + func = tvm.build(s, arg_bufs) + +# check correctness +a_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) +w_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) +c_np = conv2d_nchw_python(a_np, w_np, strides, padding) + +ctx = tvm.gpu() +a_tvm = tvm.nd.array(a_np, ctx=ctx) +w_tvm = tvm.nd.array(w_np, ctx=ctx) +c_tvm = tvm.nd.empty(c_np.shape, ctx=ctx) +func(a_tvm, w_tvm, c_tvm) + +tvm.testing.assert_allclose(c_np, c_tvm.asnumpy(), rtol=1e-2) + +# Evaluate running time. Here we choose a large repeat number (400) to reduce the noise +# and the overhead of kernel launch. You can also use nvprof to validate the result. +evaluator = func.time_evaluator(func.entry_name, ctx, number=400) +print("Time cost of this operator: %f" % evaluator(a_tvm, w_tvm, c_tvm).mean) diff --git a/tutorials/autotvm/tune_relay_arm.py b/tutorials/autotvm/tune_relay_arm.py index 0197323704ca..f024ba4f201a 100644 --- a/tutorials/autotvm/tune_relay_arm.py +++ b/tutorials/autotvm/tune_relay_arm.py @@ -189,225 +189,227 @@ def get_network(name, batch_size): # # You can register multiple devices to the tracker to accelerate the measurement in tuning. -if __name__ == "__main__": - ########################################### - # Set Tuning Options - # ------------------ - # Before tuning, we should apply some configurations. Here I use an RK3399 board - # as example. In your setting, you should modify the target and device_key accordingly. - # set :code:`use_android` to True if you use android phone. - - #### DEVICE CONFIG #### - - # Replace "aarch64-linux-gnu" with the correct target of your board. - # This target is used for cross compilation. You can query it by :code:`gcc -v` on your device. - target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu") - - # Also replace this with the device key in your tracker - device_key = "rk3399" - - # Set this to True if you use android phone - use_android = False - - #### TUNING OPTION #### - network = "resnet-18" - log_file = "%s.%s.log" % (device_key, network) - dtype = "float32" - - tuning_option = { - "log_filename": log_file, - "tuner": "xgb", - "n_trial": 1500, - "early_stopping": 800, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(build_func="ndk" if use_android else "default"), - runner=autotvm.RPCRunner( - device_key, - host="0.0.0.0", - port=9190, - number=5, - timeout=10, - ), +########################################### +# Set Tuning Options +# ------------------ +# Before tuning, we should apply some configurations. Here I use an RK3399 board +# as example. In your setting, you should modify the target and device_key accordingly. +# set :code:`use_android` to True if you use android phone. + +#### DEVICE CONFIG #### + +# Replace "aarch64-linux-gnu" with the correct target of your board. +# This target is used for cross compilation. You can query it by :code:`gcc -v` on your device. +target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu") + +# Also replace this with the device key in your tracker +device_key = "rk3399" + +# Set this to True if you use android phone +use_android = False + +#### TUNING OPTION #### +network = "resnet-18" +log_file = "%s.%s.log" % (device_key, network) +dtype = "float32" + +tuning_option = { + "log_filename": log_file, + "tuner": "xgb", + "n_trial": 1500, + "early_stopping": 800, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(build_func="ndk" if use_android else "default"), + runner=autotvm.RPCRunner( + device_key, + host="0.0.0.0", + port=9190, + number=5, + timeout=10, ), - } - - #################################################################### - # - # .. note:: How to set tuning options - # - # In general, the default values provided here work well. - # If you have enough time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, - # which makes the tuning run longer. - # If your device runs very slow or your conv2d operators have many GFLOPs, considering to - # set timeout larger. - # - # If your model has depthwise convolution, you could consider setting - # :code:`try_spatial_pack_depthwise` be :code:`True`, which perform better than default - # optimization in general. For example, on ARM CPU A53 2.0GHz, we find it could boost 1.6x - # performance of depthwise convolution on Mobilenet V1 model. - - ################################################################### - # Begin Tuning - # ------------ - # Now we can extract tuning tasks from the network and begin tuning. - # Here, we provide a simple utility function to tune a list of tasks. - # This function is just an initial implementation which tunes them in sequential order. - # We will introduce a more sophisticated tuning scheduler in the future. - - # You can skip the implementation of this function for this tutorial. - def tune_tasks( - tasks, - measure_option, - tuner="xgb", - n_trial=1000, - early_stopping=None, - log_filename="tuning.log", - use_transfer_learning=True, - ): - # create tmp log file - tmp_log_file = log_filename + ".tmp" - if os.path.exists(tmp_log_file): - os.remove(tmp_log_file) - - for i, tsk in enumerate(reversed(tasks)): - prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) - - # create tuner - if tuner == "xgb" or tuner == "xgb-rank": - tuner_obj = XGBTuner(tsk, loss_type="rank") - elif tuner == "xgb_knob": - tuner_obj = XGBTuner(tsk, loss_type="rank", feature_type="knob") - elif tuner == "ga": - tuner_obj = GATuner(tsk, pop_size=50) - elif tuner == "random": - tuner_obj = RandomTuner(tsk) - elif tuner == "gridsearch": - tuner_obj = GridSearchTuner(tsk) - else: - raise ValueError("Invalid tuner: " + tuner) - - if use_transfer_learning: - if os.path.isfile(tmp_log_file): - tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) - - # do tuning - tsk_trial = min(n_trial, len(tsk.config_space)) - tuner_obj.tune( - n_trial=tsk_trial, - early_stopping=early_stopping, - measure_option=measure_option, - callbacks=[ - autotvm.callback.progress_bar(tsk_trial, prefix=prefix), - autotvm.callback.log_to_file(tmp_log_file), - ], - ) - - # pick best records to a cache file - autotvm.record.pick_best(tmp_log_file, log_filename) + ), +} + +#################################################################### +# +# .. note:: How to set tuning options +# +# In general, the default values provided here work well. +# If you have enough time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, +# which makes the tuning run longer. +# If your device runs very slow or your conv2d operators have many GFLOPs, considering to +# set timeout larger. +# +# If your model has depthwise convolution, you could consider setting +# :code:`try_spatial_pack_depthwise` be :code:`True`, which perform better than default +# optimization in general. For example, on ARM CPU A53 2.0GHz, we find it could boost 1.6x +# performance of depthwise convolution on Mobilenet V1 model. + +################################################################### +# Begin Tuning +# ------------ +# Now we can extract tuning tasks from the network and begin tuning. +# Here, we provide a simple utility function to tune a list of tasks. +# This function is just an initial implementation which tunes them in sequential order. +# We will introduce a more sophisticated tuning scheduler in the future. + +# You can skip the implementation of this function for this tutorial. +def tune_tasks( + tasks, + measure_option, + tuner="xgb", + n_trial=1000, + early_stopping=None, + log_filename="tuning.log", + use_transfer_learning=True, +): + # create tmp log file + tmp_log_file = log_filename + ".tmp" + if os.path.exists(tmp_log_file): os.remove(tmp_log_file) - ######################################################################## - # Finally, we launch tuning jobs and evaluate the end-to-end performance. + for i, tsk in enumerate(reversed(tasks)): + prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) + + # create tuner + if tuner == "xgb" or tuner == "xgb-rank": + tuner_obj = XGBTuner(tsk, loss_type="rank") + elif tuner == "xgb_knob": + tuner_obj = XGBTuner(tsk, loss_type="rank", feature_type="knob") + elif tuner == "ga": + tuner_obj = GATuner(tsk, pop_size=50) + elif tuner == "random": + tuner_obj = RandomTuner(tsk) + elif tuner == "gridsearch": + tuner_obj = GridSearchTuner(tsk) + else: + raise ValueError("Invalid tuner: " + tuner) + + if use_transfer_learning: + if os.path.isfile(tmp_log_file): + tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) + + # do tuning + tsk_trial = min(n_trial, len(tsk.config_space)) + tuner_obj.tune( + n_trial=tsk_trial, + early_stopping=early_stopping, + measure_option=measure_option, + callbacks=[ + autotvm.callback.progress_bar(tsk_trial, prefix=prefix), + autotvm.callback.log_to_file(tmp_log_file), + ], + ) - def tune_and_evaluate(tuning_opt): - # extract workloads from relay program - print("Extract tasks...") - mod, params, input_shape, _ = get_network(network, batch_size=1) - tasks = autotvm.task.extract_from_program( - mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) + # pick best records to a cache file + autotvm.record.pick_best(tmp_log_file, log_filename) + os.remove(tmp_log_file) + + +######################################################################## +# Finally, we launch tuning jobs and evaluate the end-to-end performance. + + +def tune_and_evaluate(tuning_opt): + # extract workloads from relay program + print("Extract tasks...") + mod, params, input_shape, _ = get_network(network, batch_size=1) + tasks = autotvm.task.extract_from_program( + mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) + ) + + # run tuning tasks + print("Tuning...") + tune_tasks(tasks, **tuning_opt) + + # compile kernels with history best records + with autotvm.apply_history_best(log_file): + print("Compile...") + with tvm.transform.PassContext(opt_level=3): + lib = relay.build_module.build(mod, target=target, params=params) + + # export library + tmp = tempdir() + if use_android: + from tvm.contrib import ndk + + filename = "net.so" + lib.export_library(tmp.relpath(filename), ndk.create_shared) + else: + filename = "net.tar" + lib.export_library(tmp.relpath(filename)) + + # upload module to device + print("Upload...") + remote = autotvm.measure.request_remote(device_key, "0.0.0.0", 9190, timeout=10000) + remote.upload(tmp.relpath(filename)) + rlib = remote.load_module(filename) + + # upload parameters to device + ctx = remote.context(str(target), 0) + module = runtime.GraphModule(rlib["default"](ctx)) + data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) + module.set_input("data", data_tvm) + + # evaluate + print("Evaluate inference time cost...") + ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=10) + prof_res = np.array(ftimer().results) * 1000 # convert to millisecond + print( + "Mean inference time (std dev): %.2f ms (%.2f ms)" + % (np.mean(prof_res), np.std(prof_res)) ) - # run tuning tasks - print("Tuning...") - tune_tasks(tasks, **tuning_opt) - - # compile kernels with history best records - with autotvm.apply_history_best(log_file): - print("Compile...") - with tvm.transform.PassContext(opt_level=3): - lib = relay.build_module.build(mod, target=target, params=params) - - # export library - tmp = tempdir() - if use_android: - from tvm.contrib import ndk - - filename = "net.so" - lib.export_library(tmp.relpath(filename), ndk.create_shared) - else: - filename = "net.tar" - lib.export_library(tmp.relpath(filename)) - - # upload module to device - print("Upload...") - remote = autotvm.measure.request_remote(device_key, "0.0.0.0", 9190, timeout=10000) - remote.upload(tmp.relpath(filename)) - rlib = remote.load_module(filename) - - # upload parameters to device - ctx = remote.context(str(target), 0) - module = runtime.GraphModule(rlib["default"](ctx)) - data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) - module.set_input("data", data_tvm) - - # evaluate - print("Evaluate inference time cost...") - ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=10) - prof_res = np.array(ftimer().results) * 1000 # convert to millisecond - print( - "Mean inference time (std dev): %.2f ms (%.2f ms)" - % (np.mean(prof_res), np.std(prof_res)) - ) - - # We do not run the tuning in our webpage server since it takes too long. - # Uncomment the following line to run it by yourself. - - # tune_and_evaluate(tuning_option) - - ###################################################################### - # Sample Output - # ------------- - # The tuning needs to compile many programs and extract feature from them. - # So a high performance CPU is recommended. - # One sample output is listed below. - # It takes about 2 hours on a 32T AMD Ryzen Threadripper. - # - # .. code-block:: bash - # - # Extract tasks... - # Tuning... - # [Task 1/12] Current/Best: 22.37/ 52.19 GFLOPS | Progress: (544/1000) | 406.59 s Done. - # [Task 2/12] Current/Best: 6.51/ 18.77 GFLOPS | Progress: (608/1000) | 325.05 s Done. - # [Task 3/12] Current/Best: 4.67/ 24.87 GFLOPS | Progress: (480/1000) | 372.31 s Done. - # [Task 4/12] Current/Best: 11.35/ 46.83 GFLOPS | Progress: (736/1000) | 602.39 s Done. - # [Task 5/12] Current/Best: 1.01/ 19.80 GFLOPS | Progress: (448/1000) | 262.16 s Done. - # [Task 6/12] Current/Best: 2.47/ 23.76 GFLOPS | Progress: (672/1000) | 563.85 s Done. - # [Task 7/12] Current/Best: 14.57/ 33.97 GFLOPS | Progress: (544/1000) | 465.15 s Done. - # [Task 8/12] Current/Best: 1.13/ 17.65 GFLOPS | Progress: (576/1000) | 365.08 s Done. - # [Task 9/12] Current/Best: 14.45/ 22.66 GFLOPS | Progress: (928/1000) | 724.25 s Done. - # [Task 10/12] Current/Best: 3.22/ 15.36 GFLOPS | Progress: (864/1000) | 564.27 s Done. - # [Task 11/12] Current/Best: 11.03/ 32.23 GFLOPS | Progress: (736/1000) | 635.15 s Done. - # [Task 12/12] Current/Best: 8.00/ 21.65 GFLOPS | Progress: (1000/1000) | 1111.81 s Done. - # Compile... - # Upload... - # Evaluate inference time cost... - # Mean inference time (std dev): 162.59 ms (0.06 ms) - - ###################################################################### - # - # .. note:: **Experiencing Difficulties?** - # - # The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", - # then there must be something wrong. - # - # First, make sure you set the correct configuration of your device. - # Then, you can print debug information by adding these lines in the beginning - # of the script. It will print every measurement result, where you can find useful - # error messages. - # - # .. code-block:: python - # - # import logging - # logging.getLogger('autotvm').setLevel(logging.DEBUG) - # - # Finally, always feel free to ask our community for help on https://discuss.tvm.ai + +# We do not run the tuning in our webpage server since it takes too long. +# Uncomment the following line to run it by yourself. + +# tune_and_evaluate(tuning_option) + +###################################################################### +# Sample Output +# ------------- +# The tuning needs to compile many programs and extract feature from them. +# So a high performance CPU is recommended. +# One sample output is listed below. +# It takes about 2 hours on a 32T AMD Ryzen Threadripper. +# +# .. code-block:: bash +# +# Extract tasks... +# Tuning... +# [Task 1/12] Current/Best: 22.37/ 52.19 GFLOPS | Progress: (544/1000) | 406.59 s Done. +# [Task 2/12] Current/Best: 6.51/ 18.77 GFLOPS | Progress: (608/1000) | 325.05 s Done. +# [Task 3/12] Current/Best: 4.67/ 24.87 GFLOPS | Progress: (480/1000) | 372.31 s Done. +# [Task 4/12] Current/Best: 11.35/ 46.83 GFLOPS | Progress: (736/1000) | 602.39 s Done. +# [Task 5/12] Current/Best: 1.01/ 19.80 GFLOPS | Progress: (448/1000) | 262.16 s Done. +# [Task 6/12] Current/Best: 2.47/ 23.76 GFLOPS | Progress: (672/1000) | 563.85 s Done. +# [Task 7/12] Current/Best: 14.57/ 33.97 GFLOPS | Progress: (544/1000) | 465.15 s Done. +# [Task 8/12] Current/Best: 1.13/ 17.65 GFLOPS | Progress: (576/1000) | 365.08 s Done. +# [Task 9/12] Current/Best: 14.45/ 22.66 GFLOPS | Progress: (928/1000) | 724.25 s Done. +# [Task 10/12] Current/Best: 3.22/ 15.36 GFLOPS | Progress: (864/1000) | 564.27 s Done. +# [Task 11/12] Current/Best: 11.03/ 32.23 GFLOPS | Progress: (736/1000) | 635.15 s Done. +# [Task 12/12] Current/Best: 8.00/ 21.65 GFLOPS | Progress: (1000/1000) | 1111.81 s Done. +# Compile... +# Upload... +# Evaluate inference time cost... +# Mean inference time (std dev): 162.59 ms (0.06 ms) + +###################################################################### +# +# .. note:: **Experiencing Difficulties?** +# +# The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", +# then there must be something wrong. +# +# First, make sure you set the correct configuration of your device. +# Then, you can print debug information by adding these lines in the beginning +# of the script. It will print every measurement result, where you can find useful +# error messages. +# +# .. code-block:: python +# +# import logging +# logging.getLogger('autotvm').setLevel(logging.DEBUG) +# +# Finally, always feel free to ask our community for help on https://discuss.tvm.ai diff --git a/tutorials/autotvm/tune_relay_cuda.py b/tutorials/autotvm/tune_relay_cuda.py index 221f286a173a..4636103a22e2 100644 --- a/tutorials/autotvm/tune_relay_cuda.py +++ b/tutorials/autotvm/tune_relay_cuda.py @@ -117,273 +117,276 @@ def get_network(name, batch_size): return mod, params, input_shape, output_shape -if __name__ == "__main__": - ########################################### - # Set Tuning Options - # ------------------ - # Before tuning, we apply some configurations. - - #### DEVICE CONFIG #### - target = tvm.target.cuda() - - #### TUNING OPTION #### - network = "resnet-18" - log_file = "%s.log" % network - dtype = "float32" - - tuning_option = { - "log_filename": log_file, - "tuner": "xgb", - "n_trial": 2000, - "early_stopping": 600, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(timeout=10), - runner=autotvm.LocalRunner(number=20, repeat=3, timeout=4, min_repeat_ms=150), - ), - } - - #################################################################### - # - # .. note:: How to set tuning options - # - # In general, the default value provided here works well. - # - # If you have large time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, - # which makes the tuning runs longer. - # - # If you have multiple devices, you can use all of them for measurement to - # accelerate the tuning process. (see the 'Scale up measurement` section below). - # - - ################################################################### - # Begin Tuning - # ------------ - # Now we can extract tuning tasks from the network and begin tuning. - # Here, we provide a simple utility function to tune a list of tasks. - # This function is just an initial implementation which tunes them in sequential order. - # We will introduce a more sophisticated tuning scheduler in the future. - - # You can skip the implementation of this function for this tutorial. - def tune_tasks( - tasks, - measure_option, - tuner="xgb", - n_trial=1000, - early_stopping=None, - log_filename="tuning.log", - use_transfer_learning=True, - ): - # create tmp log file - tmp_log_file = log_filename + ".tmp" - if os.path.exists(tmp_log_file): - os.remove(tmp_log_file) - - for i, tsk in enumerate(reversed(tasks)): - prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) - - # create tuner - if tuner == "xgb" or tuner == "xgb-rank": - tuner_obj = XGBTuner(tsk, loss_type="rank") - elif tuner == "ga": - tuner_obj = GATuner(tsk, pop_size=100) - elif tuner == "random": - tuner_obj = RandomTuner(tsk) - elif tuner == "gridsearch": - tuner_obj = GridSearchTuner(tsk) - else: - raise ValueError("Invalid tuner: " + tuner) - - if use_transfer_learning: - if os.path.isfile(tmp_log_file): - tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) - - # do tuning - tsk_trial = min(n_trial, len(tsk.config_space)) - tuner_obj.tune( - n_trial=tsk_trial, - early_stopping=early_stopping, - measure_option=measure_option, - callbacks=[ - autotvm.callback.progress_bar(tsk_trial, prefix=prefix), - autotvm.callback.log_to_file(tmp_log_file), - ], - ) - - # pick best records to a cache file - autotvm.record.pick_best(tmp_log_file, log_filename) +########################################### +# Set Tuning Options +# ------------------ +# Before tuning, we apply some configurations. + +#### DEVICE CONFIG #### +target = tvm.target.cuda() + +#### TUNING OPTION #### +network = "resnet-18" +log_file = "%s.log" % network +dtype = "float32" + +tuning_option = { + "log_filename": log_file, + "tuner": "xgb", + "n_trial": 2000, + "early_stopping": 600, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(timeout=10), + runner=autotvm.LocalRunner(number=20, repeat=3, timeout=4, min_repeat_ms=150), + ), +} + +#################################################################### +# +# .. note:: How to set tuning options +# +# In general, the default value provided here works well. +# +# If you have large time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, +# which makes the tuning runs longer. +# +# If you have multiple devices, you can use all of them for measurement to +# accelerate the tuning process. (see the 'Scale up measurement` section below). +# + +################################################################### +# Begin Tuning +# ------------ +# Now we can extract tuning tasks from the network and begin tuning. +# Here, we provide a simple utility function to tune a list of tasks. +# This function is just an initial implementation which tunes them in sequential order. +# We will introduce a more sophisticated tuning scheduler in the future. + +# You can skip the implementation of this function for this tutorial. +def tune_tasks( + tasks, + measure_option, + tuner="xgb", + n_trial=1000, + early_stopping=None, + log_filename="tuning.log", + use_transfer_learning=True, +): + # create tmp log file + tmp_log_file = log_filename + ".tmp" + if os.path.exists(tmp_log_file): os.remove(tmp_log_file) - ######################################################################## - # Finally, we launch tuning jobs and evaluate the end-to-end performance. + for i, tsk in enumerate(reversed(tasks)): + prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) + + # create tuner + if tuner == "xgb" or tuner == "xgb-rank": + tuner_obj = XGBTuner(tsk, loss_type="rank") + elif tuner == "ga": + tuner_obj = GATuner(tsk, pop_size=100) + elif tuner == "random": + tuner_obj = RandomTuner(tsk) + elif tuner == "gridsearch": + tuner_obj = GridSearchTuner(tsk) + else: + raise ValueError("Invalid tuner: " + tuner) + + if use_transfer_learning: + if os.path.isfile(tmp_log_file): + tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) + + # do tuning + tsk_trial = min(n_trial, len(tsk.config_space)) + tuner_obj.tune( + n_trial=tsk_trial, + early_stopping=early_stopping, + measure_option=measure_option, + callbacks=[ + autotvm.callback.progress_bar(tsk_trial, prefix=prefix), + autotvm.callback.log_to_file(tmp_log_file), + ], + ) - def tune_and_evaluate(tuning_opt): - # extract workloads from relay program - print("Extract tasks...") - mod, params, input_shape, out_shape = get_network(network, batch_size=1) - tasks = autotvm.task.extract_from_program( - mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) + # pick best records to a cache file + autotvm.record.pick_best(tmp_log_file, log_filename) + os.remove(tmp_log_file) + + +######################################################################## +# Finally, we launch tuning jobs and evaluate the end-to-end performance. + + +def tune_and_evaluate(tuning_opt): + # extract workloads from relay program + print("Extract tasks...") + mod, params, input_shape, out_shape = get_network(network, batch_size=1) + tasks = autotvm.task.extract_from_program( + mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) + ) + + # run tuning tasks + print("Tuning...") + tune_tasks(tasks, **tuning_opt) + + # compile kernels with history best records + with autotvm.apply_history_best(log_file): + print("Compile...") + with tvm.transform.PassContext(opt_level=3): + lib = relay.build_module.build(mod, target=target, params=params) + + # export library + tmp = tempdir() + filename = "net.tar" + lib.export_library(tmp.relpath(filename)) + + # load parameters + ctx = tvm.context(str(target), 0) + module = runtime.GraphModule(lib["default"](ctx)) + data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) + module.set_input("data", data_tvm) + + # evaluate + print("Evaluate inference time cost...") + ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=600) + prof_res = np.array(ftimer().results) * 1000 # convert to millisecond + print( + "Mean inference time (std dev): %.2f ms (%.2f ms)" + % (np.mean(prof_res), np.std(prof_res)) ) - # run tuning tasks - print("Tuning...") - tune_tasks(tasks, **tuning_opt) - - # compile kernels with history best records - with autotvm.apply_history_best(log_file): - print("Compile...") - with tvm.transform.PassContext(opt_level=3): - lib = relay.build_module.build(mod, target=target, params=params) - - # export library - tmp = tempdir() - filename = "net.tar" - lib.export_library(tmp.relpath(filename)) - - # load parameters - ctx = tvm.context(str(target), 0) - module = runtime.GraphModule(lib["default"](ctx)) - data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) - module.set_input("data", data_tvm) - - # evaluate - print("Evaluate inference time cost...") - ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=600) - prof_res = np.array(ftimer().results) * 1000 # convert to millisecond - print( - "Mean inference time (std dev): %.2f ms (%.2f ms)" - % (np.mean(prof_res), np.std(prof_res)) - ) - - # We do not run the tuning in our webpage server since it takes too long. - # Uncomment the following line to run it by yourself. - - # tune_and_evaluate(tuning_option) - - ###################################################################### - # Sample Output - # ------------- - # The tuning needs to compile many programs and extract feature from them. - # So a high performance CPU is recommended. One sample output is listed below. - # It takes about 4 hours to get the following output on a 32T AMD Ryzen Threadripper. - # The tuning target is NVIDIA 1080 Ti. - # (You can see some errors during compilation. If the tuning is not stuck, it is okay.) - # - # .. code-block:: bash - # - # Extract tasks... - # Tuning... - # [Task 1/12] Current/Best: 541.83/3570.66 GFLOPS | Progress: (960/2000) | 1001.31 s Done. - # [Task 2/12] Current/Best: 0.56/ 803.33 GFLOPS | Progress: (704/2000) | 608.08 s Done. - # [Task 3/12] Current/Best: 103.69/1141.25 GFLOPS | Progress: (768/2000) | 702.13 s Done. - # [Task 4/12] Current/Best: 2905.03/3925.15 GFLOPS | Progress: (864/2000) | 745.94 sterminate called without an active exception - # [Task 4/12] Current/Best: 2789.36/3925.15 GFLOPS | Progress: (1056/2000) | 929.40 s Done. - # [Task 5/12] Current/Best: 89.06/1076.24 GFLOPS | Progress: (704/2000) | 601.73 s Done. - # [Task 6/12] Current/Best: 40.39/2129.02 GFLOPS | Progress: (1088/2000) | 1125.76 s Done. - # [Task 7/12] Current/Best: 4090.53/5007.02 GFLOPS | Progress: (800/2000) | 903.90 s Done. - # [Task 8/12] Current/Best: 4.78/1272.28 GFLOPS | Progress: (768/2000) | 749.14 s Done. - # [Task 9/12] Current/Best: 1391.45/2325.08 GFLOPS | Progress: (992/2000) | 1084.87 s Done. - # [Task 10/12] Current/Best: 1995.44/2383.59 GFLOPS | Progress: (864/2000) | 862.60 s Done. - # [Task 11/12] Current/Best: 4093.94/4899.80 GFLOPS | Progress: (224/2000) | 240.92 sterminate called without an active exception - # [Task 11/12] Current/Best: 3487.98/4909.91 GFLOPS | Progress: (480/2000) | 534.96 sterminate called without an active exception - # [Task 11/12] Current/Best: 4636.84/4912.17 GFLOPS | Progress: (1184/2000) | 1381.16 sterminate called without an active exception - # [Task 11/12] Current/Best: 50.12/4912.17 GFLOPS | Progress: (1344/2000) | 1602.81 s Done. - # [Task 12/12] Current/Best: 3581.31/4286.30 GFLOPS | Progress: (736/2000) | 943.52 s Done. - # Compile... - # Evaluate inference time cost... - # Mean inference time (std dev): 1.07 ms (0.05 ms) - # - # As a reference baseline, the time cost of MXNet + TensorRT on resnet-18 is 1.30ms. So we are a little faster. - - ###################################################################### - # - # .. note:: **Experiencing Difficulties?** - # - # The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", - # then there must be something wrong. - # - # First, make sure you set the correct configuration of your device. - # Then, you can print debug information by adding these lines in the beginning - # of the script. It will print every measurement result, where you can find useful - # error messages. - # - # .. code-block:: python - # - # import logging - # logging.getLogger('autotvm').setLevel(logging.DEBUG) - # - # Finally, always feel free to ask our community for help on https://discuss.tvm.ai - - ################################################################# - # Scale up measurement by using multiple devices - # ---------------------------------------------- - # - # If you have multiple devices, you can use all of them for measurement. - # TVM uses the RPC Tracker to manage distributed devices. - # The RPC Tracker is a centralized controller node. We can register all devices to - # the tracker. For example, if we have 10 GPU cards, we can register all of them - # to the tracker, and run 10 measurements in parallel, accelerating the tuning process. - # - # To start an RPC tracker, run this command on the host machine. The tracker is - # required during the whole tuning process, so we need to open a new terminal for - # this command: - # - # .. code-block:: bash - # - # python -m tvm.exec.rpc_tracker --host=0.0.0.0 --port=9190 - # - # The expected output is - # - # .. code-block:: bash - # - # INFO:RPCTracker:bind to 0.0.0.0:9190 - # - # Then open another new terminal for the RPC server. We need to start one server - # for each dedicated device. We use a string key to distinguish the types of devices. - # You can pick a name you like. - # (Note: For rocm backend, there are some internal errors with the compiler, - # we need to add `--no-fork` to the argument list.) - # - # .. code-block:: bash - # - # python -m tvm.exec.rpc_server --tracker=0.0.0.0:9190 --key=1080ti - # - # After registering devices, we can confirm it by querying rpc_tracker - # - # .. code-block:: bash - # - # python -m tvm.exec.query_rpc_tracker --host=0.0.0.0 --port=9190 - # - # For example, if we have four 1080ti, two titanx and one gfx900, the output can be - # - # .. code-block:: bash - # - # Queue Status - # ---------------------------------- - # key total free pending - # ---------------------------------- - # 1080ti 4 4 0 - # titanx 2 2 0 - # gfx900 1 1 0 - # ---------------------------------- - # - # Finally, we need to change the tuning option to use RPCRunner. Use the code below - # to replace the corresponding part above. - - tuning_option = { - "log_filename": log_file, - "tuner": "xgb", - "n_trial": 2000, - "early_stopping": 600, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(timeout=10), - runner=autotvm.RPCRunner( - "1080ti", # change the device key to your key - "0.0.0.0", - 9190, - number=20, - repeat=3, - timeout=4, - min_repeat_ms=150, - ), + +# We do not run the tuning in our webpage server since it takes too long. +# Uncomment the following line to run it by yourself. + +# tune_and_evaluate(tuning_option) + +###################################################################### +# Sample Output +# ------------- +# The tuning needs to compile many programs and extract feature from them. +# So a high performance CPU is recommended. One sample output is listed below. +# It takes about 4 hours to get the following output on a 32T AMD Ryzen Threadripper. +# The tuning target is NVIDIA 1080 Ti. +# (You can see some errors during compilation. If the tuning is not stuck, it is okay.) +# +# .. code-block:: bash +# +# Extract tasks... +# Tuning... +# [Task 1/12] Current/Best: 541.83/3570.66 GFLOPS | Progress: (960/2000) | 1001.31 s Done. +# [Task 2/12] Current/Best: 0.56/ 803.33 GFLOPS | Progress: (704/2000) | 608.08 s Done. +# [Task 3/12] Current/Best: 103.69/1141.25 GFLOPS | Progress: (768/2000) | 702.13 s Done. +# [Task 4/12] Current/Best: 2905.03/3925.15 GFLOPS | Progress: (864/2000) | 745.94 sterminate called without an active exception +# [Task 4/12] Current/Best: 2789.36/3925.15 GFLOPS | Progress: (1056/2000) | 929.40 s Done. +# [Task 5/12] Current/Best: 89.06/1076.24 GFLOPS | Progress: (704/2000) | 601.73 s Done. +# [Task 6/12] Current/Best: 40.39/2129.02 GFLOPS | Progress: (1088/2000) | 1125.76 s Done. +# [Task 7/12] Current/Best: 4090.53/5007.02 GFLOPS | Progress: (800/2000) | 903.90 s Done. +# [Task 8/12] Current/Best: 4.78/1272.28 GFLOPS | Progress: (768/2000) | 749.14 s Done. +# [Task 9/12] Current/Best: 1391.45/2325.08 GFLOPS | Progress: (992/2000) | 1084.87 s Done. +# [Task 10/12] Current/Best: 1995.44/2383.59 GFLOPS | Progress: (864/2000) | 862.60 s Done. +# [Task 11/12] Current/Best: 4093.94/4899.80 GFLOPS | Progress: (224/2000) | 240.92 sterminate called without an active exception +# [Task 11/12] Current/Best: 3487.98/4909.91 GFLOPS | Progress: (480/2000) | 534.96 sterminate called without an active exception +# [Task 11/12] Current/Best: 4636.84/4912.17 GFLOPS | Progress: (1184/2000) | 1381.16 sterminate called without an active exception +# [Task 11/12] Current/Best: 50.12/4912.17 GFLOPS | Progress: (1344/2000) | 1602.81 s Done. +# [Task 12/12] Current/Best: 3581.31/4286.30 GFLOPS | Progress: (736/2000) | 943.52 s Done. +# Compile... +# Evaluate inference time cost... +# Mean inference time (std dev): 1.07 ms (0.05 ms) +# +# As a reference baseline, the time cost of MXNet + TensorRT on resnet-18 is 1.30ms. So we are a little faster. + +###################################################################### +# +# .. note:: **Experiencing Difficulties?** +# +# The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", +# then there must be something wrong. +# +# First, make sure you set the correct configuration of your device. +# Then, you can print debug information by adding these lines in the beginning +# of the script. It will print every measurement result, where you can find useful +# error messages. +# +# .. code-block:: python +# +# import logging +# logging.getLogger('autotvm').setLevel(logging.DEBUG) +# +# Finally, always feel free to ask our community for help on https://discuss.tvm.ai + + +################################################################# +# Scale up measurement by using multiple devices +# ---------------------------------------------- +# +# If you have multiple devices, you can use all of them for measurement. +# TVM uses the RPC Tracker to manage distributed devices. +# The RPC Tracker is a centralized controller node. We can register all devices to +# the tracker. For example, if we have 10 GPU cards, we can register all of them +# to the tracker, and run 10 measurements in parallel, accelerating the tuning process. +# +# To start an RPC tracker, run this command on the host machine. The tracker is +# required during the whole tuning process, so we need to open a new terminal for +# this command: +# +# .. code-block:: bash +# +# python -m tvm.exec.rpc_tracker --host=0.0.0.0 --port=9190 +# +# The expected output is +# +# .. code-block:: bash +# +# INFO:RPCTracker:bind to 0.0.0.0:9190 +# +# Then open another new terminal for the RPC server. We need to start one server +# for each dedicated device. We use a string key to distinguish the types of devices. +# You can pick a name you like. +# (Note: For rocm backend, there are some internal errors with the compiler, +# we need to add `--no-fork` to the argument list.) +# +# .. code-block:: bash +# +# python -m tvm.exec.rpc_server --tracker=0.0.0.0:9190 --key=1080ti +# +# After registering devices, we can confirm it by querying rpc_tracker +# +# .. code-block:: bash +# +# python -m tvm.exec.query_rpc_tracker --host=0.0.0.0 --port=9190 +# +# For example, if we have four 1080ti, two titanx and one gfx900, the output can be +# +# .. code-block:: bash +# +# Queue Status +# ---------------------------------- +# key total free pending +# ---------------------------------- +# 1080ti 4 4 0 +# titanx 2 2 0 +# gfx900 1 1 0 +# ---------------------------------- +# +# Finally, we need to change the tuning option to use RPCRunner. Use the code below +# to replace the corresponding part above. + +tuning_option = { + "log_filename": log_file, + "tuner": "xgb", + "n_trial": 2000, + "early_stopping": 600, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(timeout=10), + runner=autotvm.RPCRunner( + "1080ti", # change the device key to your key + "0.0.0.0", + 9190, + number=20, + repeat=3, + timeout=4, + min_repeat_ms=150, ), - } + ), +} diff --git a/tutorials/autotvm/tune_relay_mobile_gpu.py b/tutorials/autotvm/tune_relay_mobile_gpu.py index 1b521cbfaa0b..61254662c463 100644 --- a/tutorials/autotvm/tune_relay_mobile_gpu.py +++ b/tutorials/autotvm/tune_relay_mobile_gpu.py @@ -188,231 +188,233 @@ def get_network(name, batch_size): # # You can register multiple devices to the tracker to accelerate the measurement in tuning. -if __name__ == "__main__": - ########################################### - # Set Tuning Options - # ------------------ - # Before tuning, we should apply some configurations. Here I use an RK3399 board - # as example. In your setting, you should modify the target and device_key accordingly. - # set :code:`use_android` to True if you use android phone. - - #### DEVICE CONFIG #### - - target = tvm.target.Target("opencl -device=mali") - - # Replace "aarch64-linux-gnu" with the correct target of your board. - # This target host is used for cross compilation. You can query it by :code:`gcc -v` on your device. - target_host = "llvm -mtriple=aarch64-linux-gnu" - - # Also replace this with the device key in your tracker - device_key = "rk3399" - - # Set this to True if you use android phone - use_android = False - - #### TUNING OPTION #### - network = "resnet-18" - log_file = "%s.%s.log" % (device_key, network) - dtype = "float32" - - tuning_option = { - "log_filename": log_file, - "tuner": "xgb", - "n_trial": 1000, - "early_stopping": 450, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(build_func="ndk" if use_android else "default"), - runner=autotvm.RPCRunner( - device_key, - host="0.0.0.0", - port=9190, - number=10, - timeout=5, - ), +########################################### +# Set Tuning Options +# ------------------ +# Before tuning, we should apply some configurations. Here I use an RK3399 board +# as example. In your setting, you should modify the target and device_key accordingly. +# set :code:`use_android` to True if you use android phone. + +#### DEVICE CONFIG #### + +target = tvm.target.Target("opencl -device=mali") + +# Replace "aarch64-linux-gnu" with the correct target of your board. +# This target host is used for cross compilation. You can query it by :code:`gcc -v` on your device. +target_host = "llvm -mtriple=aarch64-linux-gnu" + +# Also replace this with the device key in your tracker +device_key = "rk3399" + +# Set this to True if you use android phone +use_android = False + +#### TUNING OPTION #### +network = "resnet-18" +log_file = "%s.%s.log" % (device_key, network) +dtype = "float32" + +tuning_option = { + "log_filename": log_file, + "tuner": "xgb", + "n_trial": 1000, + "early_stopping": 450, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(build_func="ndk" if use_android else "default"), + runner=autotvm.RPCRunner( + device_key, + host="0.0.0.0", + port=9190, + number=10, + timeout=5, ), - } - - #################################################################### - # - # .. note:: How to set tuning options - # - # In general, the default values provided here work well. - # If you have enough time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, - # which makes the tuning run longer. - # If your device runs very slow or your conv2d operators have many GFLOPs, considering to - # set timeout larger. - # - - ################################################################### - # Begin Tuning - # ------------ - # Now we can extract tuning tasks from the network and begin tuning. - # Here, we provide a simple utility function to tune a list of tasks. - # This function is just an initial implementation which tunes them in sequential order. - # We will introduce a more sophisticated tuning scheduler in the future. - - # You can skip the implementation of this function for this tutorial. - def tune_tasks( - tasks, - measure_option, - tuner="xgb", - n_trial=1000, - early_stopping=None, - log_filename="tuning.log", - use_transfer_learning=True, - ): - # create tmp log file - tmp_log_file = log_filename + ".tmp" - if os.path.exists(tmp_log_file): - os.remove(tmp_log_file) - - for i, tsk in enumerate(reversed(tasks)): - prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) - - # create tuner - if tuner == "xgb" or tuner == "xgb-rank": - tuner_obj = XGBTuner(tsk, loss_type="rank") - elif tuner == "ga": - tuner_obj = GATuner(tsk, pop_size=50) - elif tuner == "random": - tuner_obj = RandomTuner(tsk) - elif tuner == "gridsearch": - tuner_obj = GridSearchTuner(tsk) - else: - raise ValueError("Invalid tuner: " + tuner) - - if use_transfer_learning: - if os.path.isfile(tmp_log_file): - tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) - - # do tuning - tsk_trial = min(n_trial, len(tsk.config_space)) - tuner_obj.tune( - n_trial=tsk_trial, - early_stopping=early_stopping, - measure_option=measure_option, - callbacks=[ - autotvm.callback.progress_bar(tsk_trial, prefix=prefix), - autotvm.callback.log_to_file(tmp_log_file), - ], - ) + ), +} + +#################################################################### +# +# .. note:: How to set tuning options +# +# In general, the default values provided here work well. +# If you have enough time budget, you can set :code:`n_trial`, :code:`early_stopping` larger, +# which makes the tuning run longer. +# If your device runs very slow or your conv2d operators have many GFLOPs, considering to +# set timeout larger. +# - # pick best records to a cache file - autotvm.record.pick_best(tmp_log_file, log_filename) +################################################################### +# Begin Tuning +# ------------ +# Now we can extract tuning tasks from the network and begin tuning. +# Here, we provide a simple utility function to tune a list of tasks. +# This function is just an initial implementation which tunes them in sequential order. +# We will introduce a more sophisticated tuning scheduler in the future. + +# You can skip the implementation of this function for this tutorial. +def tune_tasks( + tasks, + measure_option, + tuner="xgb", + n_trial=1000, + early_stopping=None, + log_filename="tuning.log", + use_transfer_learning=True, +): + # create tmp log file + tmp_log_file = log_filename + ".tmp" + if os.path.exists(tmp_log_file): os.remove(tmp_log_file) - ######################################################################## - # Finally, we launch tuning jobs and evaluate the end-to-end performance. - - def tune_and_evaluate(tuning_opt): - # extract workloads from relay program - print("Extract tasks...") - mod, params, input_shape, _ = get_network(network, batch_size=1) - tasks = autotvm.task.extract_from_program( - mod["main"], - target=target, - target_host=target_host, - params=params, - ops=(relay.op.get("nn.conv2d"),), + for i, tsk in enumerate(reversed(tasks)): + prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) + + # create tuner + if tuner == "xgb" or tuner == "xgb-rank": + tuner_obj = XGBTuner(tsk, loss_type="rank") + elif tuner == "ga": + tuner_obj = GATuner(tsk, pop_size=50) + elif tuner == "random": + tuner_obj = RandomTuner(tsk) + elif tuner == "gridsearch": + tuner_obj = GridSearchTuner(tsk) + else: + raise ValueError("Invalid tuner: " + tuner) + + if use_transfer_learning: + if os.path.isfile(tmp_log_file): + tuner_obj.load_history(autotvm.record.load_from_file(tmp_log_file)) + + # do tuning + tsk_trial = min(n_trial, len(tsk.config_space)) + tuner_obj.tune( + n_trial=tsk_trial, + early_stopping=early_stopping, + measure_option=measure_option, + callbacks=[ + autotvm.callback.progress_bar(tsk_trial, prefix=prefix), + autotvm.callback.log_to_file(tmp_log_file), + ], ) - # run tuning tasks - print("Tuning...") - tune_tasks(tasks, **tuning_opt) - - # compile kernels with history best records - with autotvm.apply_history_best(log_file): - print("Compile...") - with tvm.transform.PassContext(opt_level=3): - lib = relay.build_module.build( - mod, target=target, params=params, target_host=target_host - ) - # export library - tmp = tempdir() - if use_android: - from tvm.contrib import ndk - - filename = "net.so" - lib.export_library(tmp.relpath(filename), ndk.create_shared) - else: - filename = "net.tar" - lib.export_library(tmp.relpath(filename)) - - # upload module to device - print("Upload...") - remote = autotvm.measure.request_remote(device_key, "0.0.0.0", 9190, timeout=10000) - remote.upload(tmp.relpath(filename)) - rlib = remote.load_module(filename) - - # upload parameters to device - ctx = remote.context(str(target), 0) - module = runtime.GraphModule(rlib["default"](ctx)) - data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) - module.set_input("data", data_tvm) - - # evaluate - print("Evaluate inference time cost...") - ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=30) - prof_res = np.array(ftimer().results) * 1000 # convert to millisecond - print( - "Mean inference time (std dev): %.2f ms (%.2f ms)" - % (np.mean(prof_res), np.std(prof_res)) + # pick best records to a cache file + autotvm.record.pick_best(tmp_log_file, log_filename) + os.remove(tmp_log_file) + + +######################################################################## +# Finally, we launch tuning jobs and evaluate the end-to-end performance. + + +def tune_and_evaluate(tuning_opt): + # extract workloads from relay program + print("Extract tasks...") + mod, params, input_shape, _ = get_network(network, batch_size=1) + tasks = autotvm.task.extract_from_program( + mod["main"], + target=target, + target_host=target_host, + params=params, + ops=(relay.op.get("nn.conv2d"),), + ) + + # run tuning tasks + print("Tuning...") + tune_tasks(tasks, **tuning_opt) + + # compile kernels with history best records + with autotvm.apply_history_best(log_file): + print("Compile...") + with tvm.transform.PassContext(opt_level=3): + lib = relay.build_module.build( + mod, target=target, params=params, target_host=target_host ) + # export library + tmp = tempdir() + if use_android: + from tvm.contrib import ndk + + filename = "net.so" + lib.export_library(tmp.relpath(filename), ndk.create_shared) + else: + filename = "net.tar" + lib.export_library(tmp.relpath(filename)) + + # upload module to device + print("Upload...") + remote = autotvm.measure.request_remote(device_key, "0.0.0.0", 9190, timeout=10000) + remote.upload(tmp.relpath(filename)) + rlib = remote.load_module(filename) + + # upload parameters to device + ctx = remote.context(str(target), 0) + module = runtime.GraphModule(rlib["default"](ctx)) + data_tvm = tvm.nd.array((np.random.uniform(size=input_shape)).astype(dtype)) + module.set_input("data", data_tvm) + + # evaluate + print("Evaluate inference time cost...") + ftimer = module.module.time_evaluator("run", ctx, number=1, repeat=30) + prof_res = np.array(ftimer().results) * 1000 # convert to millisecond + print( + "Mean inference time (std dev): %.2f ms (%.2f ms)" + % (np.mean(prof_res), np.std(prof_res)) + ) + + +# We do not run the tuning in our webpage server since it takes too long. +# Uncomment the following line to run it by yourself. + +# tune_and_evaluate(tuning_option) + +###################################################################### +# Sample Output +# ------------- +# The tuning needs to compile many programs and extract feature from them. +# So a high performance CPU is recommended. +# One sample output is listed below. It takes about 3 hours on a 32T AMD Ryzen Threadripper. +# +# .. code-block:: bash +# +# Extract tasks... +# Tuning... +# [Task 1/17] Current/Best: 25.30/ 39.12 GFLOPS | Progress: (992/1000) | 751.22 s Done. +# [Task 2/17] Current/Best: 40.70/ 45.50 GFLOPS | Progress: (736/1000) | 545.46 s Done. +# [Task 3/17] Current/Best: 38.83/ 42.35 GFLOPS | Progress: (992/1000) | 1549.85 s Done. +# [Task 4/17] Current/Best: 23.31/ 31.02 GFLOPS | Progress: (640/1000) | 1059.31 s Done. +# [Task 5/17] Current/Best: 0.06/ 2.34 GFLOPS | Progress: (544/1000) | 305.45 s Done. +# [Task 6/17] Current/Best: 10.97/ 17.20 GFLOPS | Progress: (992/1000) | 1050.00 s Done. +# [Task 7/17] Current/Best: 8.98/ 10.94 GFLOPS | Progress: (928/1000) | 421.36 s Done. +# [Task 8/17] Current/Best: 4.48/ 14.86 GFLOPS | Progress: (704/1000) | 582.60 s Done. +# [Task 9/17] Current/Best: 10.30/ 25.99 GFLOPS | Progress: (864/1000) | 899.85 s Done. +# [Task 10/17] Current/Best: 11.73/ 12.52 GFLOPS | Progress: (608/1000) | 304.85 s Done. +# [Task 11/17] Current/Best: 15.26/ 18.68 GFLOPS | Progress: (800/1000) | 747.52 s Done. +# [Task 12/17] Current/Best: 17.48/ 26.71 GFLOPS | Progress: (1000/1000) | 1166.40 s Done. +# [Task 13/17] Current/Best: 0.96/ 11.43 GFLOPS | Progress: (960/1000) | 611.65 s Done. +# [Task 14/17] Current/Best: 17.88/ 20.22 GFLOPS | Progress: (672/1000) | 670.29 s Done. +# [Task 15/17] Current/Best: 11.62/ 13.98 GFLOPS | Progress: (736/1000) | 449.25 s Done. +# [Task 16/17] Current/Best: 19.90/ 23.83 GFLOPS | Progress: (608/1000) | 708.64 s Done. +# [Task 17/17] Current/Best: 17.98/ 22.75 GFLOPS | Progress: (736/1000) | 1122.60 s Done. +# Compile... +# Upload... +# Evaluate inference time cost... +# Mean inference time (std dev): 128.05 ms (7.74 ms) +# - # We do not run the tuning in our webpage server since it takes too long. - # Uncomment the following line to run it by yourself. - - # tune_and_evaluate(tuning_option) - - ###################################################################### - # Sample Output - # ------------- - # The tuning needs to compile many programs and extract feature from them. - # So a high performance CPU is recommended. - # One sample output is listed below. It takes about 3 hours on a 32T AMD Ryzen Threadripper. - # - # .. code-block:: bash - # - # Extract tasks... - # Tuning... - # [Task 1/17] Current/Best: 25.30/ 39.12 GFLOPS | Progress: (992/1000) | 751.22 s Done. - # [Task 2/17] Current/Best: 40.70/ 45.50 GFLOPS | Progress: (736/1000) | 545.46 s Done. - # [Task 3/17] Current/Best: 38.83/ 42.35 GFLOPS | Progress: (992/1000) | 1549.85 s Done. - # [Task 4/17] Current/Best: 23.31/ 31.02 GFLOPS | Progress: (640/1000) | 1059.31 s Done. - # [Task 5/17] Current/Best: 0.06/ 2.34 GFLOPS | Progress: (544/1000) | 305.45 s Done. - # [Task 6/17] Current/Best: 10.97/ 17.20 GFLOPS | Progress: (992/1000) | 1050.00 s Done. - # [Task 7/17] Current/Best: 8.98/ 10.94 GFLOPS | Progress: (928/1000) | 421.36 s Done. - # [Task 8/17] Current/Best: 4.48/ 14.86 GFLOPS | Progress: (704/1000) | 582.60 s Done. - # [Task 9/17] Current/Best: 10.30/ 25.99 GFLOPS | Progress: (864/1000) | 899.85 s Done. - # [Task 10/17] Current/Best: 11.73/ 12.52 GFLOPS | Progress: (608/1000) | 304.85 s Done. - # [Task 11/17] Current/Best: 15.26/ 18.68 GFLOPS | Progress: (800/1000) | 747.52 s Done. - # [Task 12/17] Current/Best: 17.48/ 26.71 GFLOPS | Progress: (1000/1000) | 1166.40 s Done. - # [Task 13/17] Current/Best: 0.96/ 11.43 GFLOPS | Progress: (960/1000) | 611.65 s Done. - # [Task 14/17] Current/Best: 17.88/ 20.22 GFLOPS | Progress: (672/1000) | 670.29 s Done. - # [Task 15/17] Current/Best: 11.62/ 13.98 GFLOPS | Progress: (736/1000) | 449.25 s Done. - # [Task 16/17] Current/Best: 19.90/ 23.83 GFLOPS | Progress: (608/1000) | 708.64 s Done. - # [Task 17/17] Current/Best: 17.98/ 22.75 GFLOPS | Progress: (736/1000) | 1122.60 s Done. - # Compile... - # Upload... - # Evaluate inference time cost... - # Mean inference time (std dev): 128.05 ms (7.74 ms) - # - - ###################################################################### - # - # .. note:: **Experiencing Difficulties?** - # - # The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", - # then there must be something wrong. - # - # First, make sure you set the correct configuration of your device. - # Then, you can print debug information by adding these lines in the beginning - # of the script. It will print every measurement result, where you can find useful - # error messages. - # - # .. code-block:: python - # - # import logging - # logging.getLogger('autotvm').setLevel(logging.DEBUG) - # - # Finally, always feel free to ask our community for help on https://discuss.tvm.ai +###################################################################### +# +# .. note:: **Experiencing Difficulties?** +# +# The auto tuning module is error-prone. If you always see " 0.00/ 0.00 GFLOPS", +# then there must be something wrong. +# +# First, make sure you set the correct configuration of your device. +# Then, you can print debug information by adding these lines in the beginning +# of the script. It will print every measurement result, where you can find useful +# error messages. +# +# .. code-block:: python +# +# import logging +# logging.getLogger('autotvm').setLevel(logging.DEBUG) +# +# Finally, always feel free to ask our community for help on https://discuss.tvm.ai diff --git a/tutorials/autotvm/tune_relay_x86.py b/tutorials/autotvm/tune_relay_x86.py index 8c986634caba..1dd947fefd25 100644 --- a/tutorials/autotvm/tune_relay_x86.py +++ b/tutorials/autotvm/tune_relay_x86.py @@ -88,167 +88,172 @@ def get_network(name, batch_size): return mod, params, input_shape, output_shape -if __name__ == "__main__": - # Replace "llvm" with the correct target of your CPU. - # For example, for AWS EC2 c5 instance with Intel Xeon - # Platinum 8000 series, the target should be "llvm -mcpu=skylake-avx512". - # For AWS EC2 c4 instance with Intel Xeon E5-2666 v3, it should be - # "llvm -mcpu=core-avx2". - target = "llvm" - - batch_size = 1 - dtype = "float32" - model_name = "resnet-18" - log_file = "%s.log" % model_name - graph_opt_sch_file = "%s_graph_opt.log" % model_name - - # Set the input name of the graph - # For ONNX models, it is typically "0". - input_name = "data" - - # Set number of threads used for tuning based on the number of - # physical CPU cores on your machine. - num_threads = 1 - os.environ["TVM_NUM_THREADS"] = str(num_threads) - - ################################################################# - # Configure tensor tuning settings and create tasks - # ------------------------------------------------- - # To get better kernel execution performance on x86 CPU, - # we need to change data layout of convolution kernel from - # "NCHW" to "NCHWc". To deal with this situation, we define - # conv2d_NCHWc operator in topi. We will tune this operator - # instead of plain conv2d. - # - # We will use local mode for tuning configuration. RPC tracker - # mode can be setup similarly to the approach in - # :ref:`tune_relay_arm` tutorial. - # - # To perform a precise measurement, we should repeat the measurement several - # times and use the average of results. In addition, we need to flush the cache - # for the weight tensors between repeated measurements. This can make the measured - # latency of one operator closer to its actual latency during end-to-end inference. - - tuning_option = { - "log_filename": log_file, - "tuner": "random", - "early_stopping": None, - "measure_option": autotvm.measure_option( - builder=autotvm.LocalBuilder(), - runner=autotvm.LocalRunner( - number=1, repeat=10, min_repeat_ms=0, enable_cpu_cache_flush=True - ), +# Replace "llvm" with the correct target of your CPU. +# For example, for AWS EC2 c5 instance with Intel Xeon +# Platinum 8000 series, the target should be "llvm -mcpu=skylake-avx512". +# For AWS EC2 c4 instance with Intel Xeon E5-2666 v3, it should be +# "llvm -mcpu=core-avx2". +target = "llvm" + +batch_size = 1 +dtype = "float32" +model_name = "resnet-18" +log_file = "%s.log" % model_name +graph_opt_sch_file = "%s_graph_opt.log" % model_name + +# Set the input name of the graph +# For ONNX models, it is typically "0". +input_name = "data" + +# Set number of threads used for tuning based on the number of +# physical CPU cores on your machine. +num_threads = 1 +os.environ["TVM_NUM_THREADS"] = str(num_threads) + + +################################################################# +# Configure tensor tuning settings and create tasks +# ------------------------------------------------- +# To get better kernel execution performance on x86 CPU, +# we need to change data layout of convolution kernel from +# "NCHW" to "NCHWc". To deal with this situation, we define +# conv2d_NCHWc operator in topi. We will tune this operator +# instead of plain conv2d. +# +# We will use local mode for tuning configuration. RPC tracker +# mode can be setup similarly to the approach in +# :ref:`tune_relay_arm` tutorial. +# +# To perform a precise measurement, we should repeat the measurement several +# times and use the average of results. In addition, we need to flush the cache +# for the weight tensors between repeated measurements. This can make the measured +# latency of one operator closer to its actual latency during end-to-end inference. + +tuning_option = { + "log_filename": log_file, + "tuner": "random", + "early_stopping": None, + "measure_option": autotvm.measure_option( + builder=autotvm.LocalBuilder(), + runner=autotvm.LocalRunner( + number=1, repeat=10, min_repeat_ms=0, enable_cpu_cache_flush=True ), - } - - # You can skip the implementation of this function for this tutorial. - def tune_kernels( - tasks, measure_option, tuner="gridsearch", early_stopping=None, log_filename="tuning.log" - ): - - for i, task in enumerate(tasks): - prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) - - # create tuner - if tuner == "xgb" or tuner == "xgb-rank": - tuner_obj = XGBTuner(task, loss_type="rank") - elif tuner == "ga": - tuner_obj = GATuner(task, pop_size=50) - elif tuner == "random": - tuner_obj = RandomTuner(task) - elif tuner == "gridsearch": - tuner_obj = GridSearchTuner(task) - else: - raise ValueError("Invalid tuner: " + tuner) - - # do tuning - n_trial = len(task.config_space) - tuner_obj.tune( - n_trial=n_trial, - early_stopping=early_stopping, - measure_option=measure_option, - callbacks=[ - autotvm.callback.progress_bar(n_trial, prefix=prefix), - autotvm.callback.log_to_file(log_filename), - ], - ) - - # Use graph tuner to achieve graph level optimal schedules - # Set use_DP=False if it takes too long to finish. - def tune_graph(graph, dshape, records, opt_sch_file, use_DP=True): - target_op = [ - relay.op.get("nn.conv2d"), - ] - Tuner = DPTuner if use_DP else PBQPTuner - executor = Tuner(graph, {input_name: dshape}, records, target_op, target) - executor.benchmark_layout_transform(min_exec_num=2000) - executor.run() - executor.write_opt_sch2record_file(opt_sch_file) - - ######################################################################## - # Finally, we launch tuning jobs and evaluate the end-to-end performance. - - def tune_and_evaluate(tuning_opt): - # extract workloads from relay program - print("Extract tasks...") - mod, params, data_shape, out_shape = get_network(model_name, batch_size) - tasks = autotvm.task.extract_from_program( - mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) + ), +} + + +# You can skip the implementation of this function for this tutorial. +def tune_kernels( + tasks, measure_option, tuner="gridsearch", early_stopping=None, log_filename="tuning.log" +): + + for i, task in enumerate(tasks): + prefix = "[Task %2d/%2d] " % (i + 1, len(tasks)) + + # create tuner + if tuner == "xgb" or tuner == "xgb-rank": + tuner_obj = XGBTuner(task, loss_type="rank") + elif tuner == "ga": + tuner_obj = GATuner(task, pop_size=50) + elif tuner == "random": + tuner_obj = RandomTuner(task) + elif tuner == "gridsearch": + tuner_obj = GridSearchTuner(task) + else: + raise ValueError("Invalid tuner: " + tuner) + + # do tuning + n_trial = len(task.config_space) + tuner_obj.tune( + n_trial=n_trial, + early_stopping=early_stopping, + measure_option=measure_option, + callbacks=[ + autotvm.callback.progress_bar(n_trial, prefix=prefix), + autotvm.callback.log_to_file(log_filename), + ], + ) + + +# Use graph tuner to achieve graph level optimal schedules +# Set use_DP=False if it takes too long to finish. +def tune_graph(graph, dshape, records, opt_sch_file, use_DP=True): + target_op = [ + relay.op.get("nn.conv2d"), + ] + Tuner = DPTuner if use_DP else PBQPTuner + executor = Tuner(graph, {input_name: dshape}, records, target_op, target) + executor.benchmark_layout_transform(min_exec_num=2000) + executor.run() + executor.write_opt_sch2record_file(opt_sch_file) + + +######################################################################## +# Finally, we launch tuning jobs and evaluate the end-to-end performance. + + +def tune_and_evaluate(tuning_opt): + # extract workloads from relay program + print("Extract tasks...") + mod, params, data_shape, out_shape = get_network(model_name, batch_size) + tasks = autotvm.task.extract_from_program( + mod["main"], target=target, params=params, ops=(relay.op.get("nn.conv2d"),) + ) + + # run tuning tasks + tune_kernels(tasks, **tuning_opt) + tune_graph(mod["main"], data_shape, log_file, graph_opt_sch_file) + + # compile kernels with graph-level best records + with autotvm.apply_graph_best(graph_opt_sch_file): + print("Compile...") + with tvm.transform.PassContext(opt_level=3): + lib = relay.build_module.build(mod, target=target, params=params) + + # upload parameters to device + ctx = tvm.cpu() + data_tvm = tvm.nd.array((np.random.uniform(size=data_shape)).astype(dtype)) + module = runtime.GraphModule(lib["default"](ctx)) + module.set_input(input_name, data_tvm) + + # evaluate + print("Evaluate inference time cost...") + ftimer = module.module.time_evaluator("run", ctx, number=100, repeat=3) + prof_res = np.array(ftimer().results) * 1000 # convert to millisecond + print( + "Mean inference time (std dev): %.2f ms (%.2f ms)" + % (np.mean(prof_res), np.std(prof_res)) ) - # run tuning tasks - tune_kernels(tasks, **tuning_opt) - tune_graph(mod["main"], data_shape, log_file, graph_opt_sch_file) - - # compile kernels with graph-level best records - with autotvm.apply_graph_best(graph_opt_sch_file): - print("Compile...") - with tvm.transform.PassContext(opt_level=3): - lib = relay.build_module.build(mod, target=target, params=params) - - # upload parameters to device - ctx = tvm.cpu() - data_tvm = tvm.nd.array((np.random.uniform(size=data_shape)).astype(dtype)) - module = runtime.GraphModule(lib["default"](ctx)) - module.set_input(input_name, data_tvm) - - # evaluate - print("Evaluate inference time cost...") - ftimer = module.module.time_evaluator("run", ctx, number=100, repeat=3) - prof_res = np.array(ftimer().results) * 1000 # convert to millisecond - print( - "Mean inference time (std dev): %.2f ms (%.2f ms)" - % (np.mean(prof_res), np.std(prof_res)) - ) - - # We do not run the tuning in our webpage server since it takes too long. - # Uncomment the following line to run it by yourself. - - # tune_and_evaluate(tuning_option) - - ###################################################################### - # Sample Output - # ------------- - # The tuning needs to compile many programs and extract feature from them. - # So a high performance CPU is recommended. - # One sample output is listed below. - # - # .. code-block:: bash - # - # Extract tasks... - # Tuning... - # [Task 1/12] Current/Best: 598.05/2497.63 GFLOPS | Progress: (252/252) | 1357.95 s Done. - # [Task 2/12] Current/Best: 522.63/2279.24 GFLOPS | Progress: (784/784) | 3989.60 s Done. - # [Task 3/12] Current/Best: 447.33/1927.69 GFLOPS | Progress: (784/784) | 3869.14 s Done. - # [Task 4/12] Current/Best: 481.11/1912.34 GFLOPS | Progress: (672/672) | 3274.25 s Done. - # [Task 5/12] Current/Best: 414.09/1598.45 GFLOPS | Progress: (672/672) | 2720.78 s Done. - # [Task 6/12] Current/Best: 508.96/2273.20 GFLOPS | Progress: (768/768) | 3718.75 s Done. - # [Task 7/12] Current/Best: 469.14/1955.79 GFLOPS | Progress: (576/576) | 2665.67 s Done. - # [Task 8/12] Current/Best: 230.91/1658.97 GFLOPS | Progress: (576/576) | 2435.01 s Done. - # [Task 9/12] Current/Best: 487.75/2295.19 GFLOPS | Progress: (648/648) | 3009.95 s Done. - # [Task 10/12] Current/Best: 182.33/1734.45 GFLOPS | Progress: (360/360) | 1755.06 s Done. - # [Task 11/12] Current/Best: 372.18/1745.15 GFLOPS | Progress: (360/360) | 1684.50 s Done. - # [Task 12/12] Current/Best: 215.34/2271.11 GFLOPS | Progress: (400/400) | 2128.74 s Done. - # Compile... - # Evaluate inference time cost... - # Mean inference time (std dev): 3.16 ms (0.03 ms) + +# We do not run the tuning in our webpage server since it takes too long. +# Uncomment the following line to run it by yourself. + +# tune_and_evaluate(tuning_option) + +###################################################################### +# Sample Output +# ------------- +# The tuning needs to compile many programs and extract feature from them. +# So a high performance CPU is recommended. +# One sample output is listed below. +# +# .. code-block:: bash +# +# Extract tasks... +# Tuning... +# [Task 1/12] Current/Best: 598.05/2497.63 GFLOPS | Progress: (252/252) | 1357.95 s Done. +# [Task 2/12] Current/Best: 522.63/2279.24 GFLOPS | Progress: (784/784) | 3989.60 s Done. +# [Task 3/12] Current/Best: 447.33/1927.69 GFLOPS | Progress: (784/784) | 3869.14 s Done. +# [Task 4/12] Current/Best: 481.11/1912.34 GFLOPS | Progress: (672/672) | 3274.25 s Done. +# [Task 5/12] Current/Best: 414.09/1598.45 GFLOPS | Progress: (672/672) | 2720.78 s Done. +# [Task 6/12] Current/Best: 508.96/2273.20 GFLOPS | Progress: (768/768) | 3718.75 s Done. +# [Task 7/12] Current/Best: 469.14/1955.79 GFLOPS | Progress: (576/576) | 2665.67 s Done. +# [Task 8/12] Current/Best: 230.91/1658.97 GFLOPS | Progress: (576/576) | 2435.01 s Done. +# [Task 9/12] Current/Best: 487.75/2295.19 GFLOPS | Progress: (648/648) | 3009.95 s Done. +# [Task 10/12] Current/Best: 182.33/1734.45 GFLOPS | Progress: (360/360) | 1755.06 s Done. +# [Task 11/12] Current/Best: 372.18/1745.15 GFLOPS | Progress: (360/360) | 1684.50 s Done. +# [Task 12/12] Current/Best: 215.34/2271.11 GFLOPS | Progress: (400/400) | 2128.74 s Done. +# Compile... +# Evaluate inference time cost... +# Mean inference time (std dev): 3.16 ms (0.03 ms) diff --git a/tutorials/autotvm/tune_simple_template.py b/tutorials/autotvm/tune_simple_template.py index 2243d2df6347..357abf19a09c 100644 --- a/tutorials/autotvm/tune_simple_template.py +++ b/tutorials/autotvm/tune_simple_template.py @@ -56,7 +56,6 @@ import numpy as np import tvm from tvm import te -import tvm.testing # the module is called `autotvm` from tvm import autotvm @@ -215,119 +214,118 @@ def matmul(N, L, M, dtype): return s, [A, B, C] -if __name__ == "__main__": - ###################################################################### - # .. note:: More Explanation on :code:`cfg.defile_split` - # - # In this template, :code:`cfg.define_split("tile_y", y, num_outputs=2)` will enumerate - # all possible combinations that can split axis y into two axes with factors of the length of y. - # For example, if the length of y is 32 and we want to split it into two axes - # using factors of 32, then there are 6 possible values for - # (length of outer axis, length of inner axis) pair, namely - # (32, 1), (16, 2), (8, 4), (4, 8), (2, 16) or (1, 32). - # They are just the 6 possible values of `tile_y`. - # - # During schedule, :code:`cfg["tile_y"]` is a :code:`SplitEntity` object. - # We stores the lengths of outer axes and inner axes in :code:`cfg['tile_y'].size` - # (a tuple with two elements). - # In this template, we apply it by using :code:`yo, yi = cfg['tile_y'].apply(s, C, y)`. - # Actually, this is equivalent to - # :code:`yo, yi = s[C].split(y, cfg["tile_y"].size[1])` - # or :code:`yo, yi = s[C].split(y, nparts=cfg['tile_y"].size[0])` - # - # The advantage of using cfg.apply API is that it makes multi-level split - # (when num_outputs >= 3) easier. - - ###################################################################### - # Step 2: Search through the space - # --------------------------------- - # In step 1, we build the search space by extending our old schedule code - # into a template. The next step is to pick a tuner and explore in this space. - # - # Auto-tuners in TVM - # ^^^^^^^^^^^^^^^^^^ - # The job for a tuner can be described by following pseudo code - # - # .. code-block:: c - # - # ct = 0 - # while ct < max_number_of_trials: - # propose a batch of configs - # measure this batch of configs on real hardware and get results - # ct += batch_size - # - # When proposing the next batch of configs, the tuner can take different strategies. We - # provide four tuners with different strategies in autotvm. - # - # * :any:`RandomTuner`: Enumerate the space in a random order - # * :any:`GridSearchTuner`: Enumerate the space in a grid search order - # * :any:`GATuner`: Using genetic algorithm to search through the space - # * :any:`XGBTuner`: Uses a model based method. Train a XGBoost model to predict the speed of lowered IR and pick the next batch according to the prediction. - # - # You can choose the tuner according to the size of your space, your time budget and other factors. - # For example, if your space is very small (less than 1000), a gridsearch tuner or a - # random tuner is good enough. If your space is at the level of 10^9 (this is the space - # size of a conv2d operator on CUDA GPU), XGBoostTuner can explore more efficiently - # and find better configs. - - ################################################################ - # Begin tuning - # ^^^^^^^^^^^^ - # Here we continue our matrix multiplication example. - # First we should create a tuning task. - # We can also inspect the initialized search space. - # In this case, for a 512x512 square matrix multiplication, the space size - # is 10x10=100 - N, L, M = 512, 512, 512 - task = autotvm.task.create("tutorial/matmul", args=(N, L, M, "float32"), target="llvm") - print(task.config_space) - - ################################################################ - # Then we need to define how to measure the generated code and pick a tuner. - # Since our space is small, a random tuner is just okay. - # - # We only make 10 trials in this tutorial for demonstration. In practice, - # you can do more trials according to your time budget. - # We will log the tuning results into a log file. This file can be - # used to get the best config later. - - # logging config (for printing tuning log to the screen) - logging.getLogger("autotvm").setLevel(logging.DEBUG) - logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) - - # There are two steps for measuring a config: build and run. - # By default, we use all CPU cores to compile program. Then measure them sequentially. - # We measure 5 times and take average to reduce variance. - measure_option = autotvm.measure_option(builder="local", runner=autotvm.LocalRunner(number=5)) - - # Begin tuning with RandomTuner, log records to file `matmul.log` - # You can use alternatives like XGBTuner. - tuner = autotvm.tuner.RandomTuner(task) - tuner.tune( - n_trial=10, - measure_option=measure_option, - callbacks=[autotvm.callback.log_to_file("matmul.log")], - ) - - ######################################################################### - # Finally we apply history best from the cache file and check its correctness. - # We can call the function :code:`matmul` directly under the - # :any:`autotvm.apply_history_best` context. When we call this function, - # it will query the dispatch context with its argument and get the best config - # with the same argument. - - # apply history best from log file - with autotvm.apply_history_best("matmul.log"): - with tvm.target.Target("llvm"): - s, arg_bufs = matmul(N, L, M, "float32") - func = tvm.build(s, arg_bufs) - - # check correctness - a_np = np.random.uniform(size=(N, L)).astype(np.float32) - b_np = np.random.uniform(size=(L, M)).astype(np.float32) - c_np = a_np.dot(b_np) - - c_tvm = tvm.nd.empty(c_np.shape) - func(tvm.nd.array(a_np), tvm.nd.array(b_np), c_tvm) - - tvm.testing.assert_allclose(c_np, c_tvm.asnumpy(), rtol=1e-2) +###################################################################### +# .. note:: More Explanation on :code:`cfg.defile_split` +# +# In this template, :code:`cfg.define_split("tile_y", y, num_outputs=2)` will enumerate +# all possible combinations that can split axis y into two axes with factors of the length of y. +# For example, if the length of y is 32 and we want to split it into two axes +# using factors of 32, then there are 6 possible values for +# (length of outer axis, length of inner axis) pair, namely +# (32, 1), (16, 2), (8, 4), (4, 8), (2, 16) or (1, 32). +# They are just the 6 possible values of `tile_y`. +# +# During schedule, :code:`cfg["tile_y"]` is a :code:`SplitEntity` object. +# We stores the lengths of outer axes and inner axes in :code:`cfg['tile_y'].size` +# (a tuple with two elements). +# In this template, we apply it by using :code:`yo, yi = cfg['tile_y'].apply(s, C, y)`. +# Actually, this is equivalent to +# :code:`yo, yi = s[C].split(y, cfg["tile_y"].size[1])` +# or :code:`yo, yi = s[C].split(y, nparts=cfg['tile_y"].size[0])` +# +# The advantage of using cfg.apply API is that it makes multi-level split +# (when num_outputs >= 3) easier. + +###################################################################### +# Step 2: Search through the space +# --------------------------------- +# In step 1, we build the search space by extending our old schedule code +# into a template. The next step is to pick a tuner and explore in this space. +# +# Auto-tuners in TVM +# ^^^^^^^^^^^^^^^^^^ +# The job for a tuner can be described by following pseudo code +# +# .. code-block:: c +# +# ct = 0 +# while ct < max_number_of_trials: +# propose a batch of configs +# measure this batch of configs on real hardware and get results +# ct += batch_size +# +# When proposing the next batch of configs, the tuner can take different strategies. We +# provide four tuners with different strategies in autotvm. +# +# * :any:`RandomTuner`: Enumerate the space in a random order +# * :any:`GridSearchTuner`: Enumerate the space in a grid search order +# * :any:`GATuner`: Using genetic algorithm to search through the space +# * :any:`XGBTuner`: Uses a model based method. Train a XGBoost model to predict the speed of lowered IR and pick the next batch according to the prediction. +# +# You can choose the tuner according to the size of your space, your time budget and other factors. +# For example, if your space is very small (less than 1000), a gridsearch tuner or a +# random tuner is good enough. If your space is at the level of 10^9 (this is the space +# size of a conv2d operator on CUDA GPU), XGBoostTuner can explore more efficiently +# and find better configs. + +################################################################ +# Begin tuning +# ^^^^^^^^^^^^ +# Here we continue our matrix multiplication example. +# First we should create a tuning task. +# We can also inspect the initialized search space. +# In this case, for a 512x512 square matrix multiplication, the space size +# is 10x10=100 +N, L, M = 512, 512, 512 +task = autotvm.task.create("tutorial/matmul", args=(N, L, M, "float32"), target="llvm") +print(task.config_space) + +################################################################ +# Then we need to define how to measure the generated code and pick a tuner. +# Since our space is small, a random tuner is just okay. +# +# We only make 10 trials in this tutorial for demonstration. In practice, +# you can do more trials according to your time budget. +# We will log the tuning results into a log file. This file can be +# used to get the best config later. + +# logging config (for printing tuning log to the screen) +logging.getLogger("autotvm").setLevel(logging.DEBUG) +logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) + +# There are two steps for measuring a config: build and run. +# By default, we use all CPU cores to compile program. Then measure them sequentially. +# We measure 5 times and take average to reduce variance. +measure_option = autotvm.measure_option(builder="local", runner=autotvm.LocalRunner(number=5)) + +# Begin tuning with RandomTuner, log records to file `matmul.log` +# You can use alternatives like XGBTuner. +tuner = autotvm.tuner.RandomTuner(task) +tuner.tune( + n_trial=10, + measure_option=measure_option, + callbacks=[autotvm.callback.log_to_file("matmul.log")], +) + +######################################################################### +# Finally we apply history best from the cache file and check its correctness. +# We can call the function :code:`matmul` directly under the +# :any:`autotvm.apply_history_best` context. When we call this function, +# it will query the dispatch context with its argument and get the best config +# with the same argument. + +# apply history best from log file +with autotvm.apply_history_best("matmul.log"): + with tvm.target.Target("llvm"): + s, arg_bufs = matmul(N, L, M, "float32") + func = tvm.build(s, arg_bufs) + +# check correctness +a_np = np.random.uniform(size=(N, L)).astype(np.float32) +b_np = np.random.uniform(size=(L, M)).astype(np.float32) +c_np = a_np.dot(b_np) + +c_tvm = tvm.nd.empty(c_np.shape) +func(tvm.nd.array(a_np), tvm.nd.array(b_np), c_tvm) + +tvm.testing.assert_allclose(c_np, c_tvm.asnumpy(), rtol=1e-2) From 228e382d760c019ce6a29e282ec090223a5528ec Mon Sep 17 00:00:00 2001 From: Tristan Konolige Date: Wed, 28 Oct 2020 11:26:32 -0700 Subject: [PATCH 7/8] Add spawn tests --- .../unittest/test_auto_scheduler_measure.py | 15 +++++++++++++++ .../unittest/test_auto_scheduler_search_policy.py | 15 +++++++++++++++ .../test_auto_scheduler_task_scheduler.py | 14 ++++++++++++++ 3 files changed, 44 insertions(+) diff --git a/tests/python/unittest/test_auto_scheduler_measure.py b/tests/python/unittest/test_auto_scheduler_measure.py index f38fb4d74351..80ce98d0b1c1 100644 --- a/tests/python/unittest/test_auto_scheduler_measure.py +++ b/tests/python/unittest/test_auto_scheduler_measure.py @@ -17,6 +17,7 @@ """ Test measurement and log serialization. """ +import multiprocessing import tvm from tvm import topi from tvm import te, auto_scheduler @@ -230,6 +231,19 @@ def test_measure_local_builder_rpc_runner(): del measure_ctx +def measure_local_builder_rpc_runner_spawn(): + assert multiprocessing.get_start_method(False) == "spawn" + test_measure_local_builder_rpc_runner() + + +@tvm.testing.requires_llvm +def test_measure_local_builder_rpc_runner_spawn(): + ctx = multiprocessing.get_context("spawn") + p = ctx.Process(target=measure_local_builder_rpc_runner_spawn) + p.start() + p.join() + + if __name__ == "__main__": test_record_split_reorder_fuse_annotation() test_record_compute_at_root_inline_cache_read_write() @@ -237,4 +251,5 @@ def test_measure_local_builder_rpc_runner(): test_record_pragma_storage_align_rfactor() test_recover_measure_input() test_measure_local_builder_runner() + test_measure_local_builder_runner_spawn() test_measure_local_builder_rpc_runner() diff --git a/tests/python/unittest/test_auto_scheduler_search_policy.py b/tests/python/unittest/test_auto_scheduler_search_policy.py index 2b5a69a884af..5329f3d50685 100644 --- a/tests/python/unittest/test_auto_scheduler_search_policy.py +++ b/tests/python/unittest/test_auto_scheduler_search_policy.py @@ -18,6 +18,7 @@ """Test search policy""" import random +import multiprocessing import numpy as np import tempfile @@ -123,6 +124,19 @@ def test_sketch_search_policy_basic(): t.join() +def sketch_search_policy_basic_spawn(): + assert multiprocessing.get_start_method(False) == "spawn" + test_sketch_search_policy_basic() + + +@tvm.testing.requires_llvm +def test_sketch_search_policy_basic_spawn(): + ctx = multiprocessing.get_context("spawn") + p = ctx.Process(target=sketch_search_policy_basic_spawn) + p.start() + p.join() + + @tvm.testing.requires_llvm def test_sketch_search_policy_xgbmodel(): # wrap the search in a new thread to avoid the conflict @@ -179,6 +193,7 @@ def test_sketch_search_policy_cuda_xgbmodel_rpc_runner(): if __name__ == "__main__": test_workload_registry_search_basic() test_sketch_search_policy_basic() + test_sketch_search_policy_basic_spawn() test_sketch_search_policy_xgbmodel() test_sketch_search_policy_cuda_rpc_runner() test_sketch_search_policy_cuda_xgbmodel_rpc_runner() diff --git a/tests/python/unittest/test_auto_scheduler_task_scheduler.py b/tests/python/unittest/test_auto_scheduler_task_scheduler.py index 72b998a5a38a..44e315bf40d0 100644 --- a/tests/python/unittest/test_auto_scheduler_task_scheduler.py +++ b/tests/python/unittest/test_auto_scheduler_task_scheduler.py @@ -18,6 +18,7 @@ import tempfile +import multiprocessing import numpy as np from tvm import auto_scheduler @@ -68,6 +69,18 @@ def objective_func(costs): task_scheduler.tune(tune_option, search_policy="sketch.random") +def task_scheduler_round_robin_spawn(): + assert multiprocessing.get_start_method(False) == "spawn" + test_task_scheduler_round_robin() + + +def test_task_scheduler_round_robin_spawn(): + ctx = muliprocessing.get_context("spawn") + p = ctx.Process(target=task_scheduler_round_robin_spawn) + p.start() + p.join() + + def test_task_scheduler_gradient(): tasks = [] for n in [2, 4]: @@ -109,4 +122,5 @@ def objective_func(costs): if __name__ == "__main__": test_task_scheduler_round_robin() + test_task_scheduler_round_robin_spawn() test_task_scheduler_gradient() From 1686131c2f95422636f83658b3f61df23de0ebc6 Mon Sep 17 00:00:00 2001 From: Tristan Konolige Date: Thu, 29 Oct 2020 09:46:16 -0700 Subject: [PATCH 8/8] fix test --- tests/python/unittest/test_auto_scheduler_task_scheduler.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/unittest/test_auto_scheduler_task_scheduler.py b/tests/python/unittest/test_auto_scheduler_task_scheduler.py index 44e315bf40d0..7851d922013d 100644 --- a/tests/python/unittest/test_auto_scheduler_task_scheduler.py +++ b/tests/python/unittest/test_auto_scheduler_task_scheduler.py @@ -75,7 +75,7 @@ def task_scheduler_round_robin_spawn(): def test_task_scheduler_round_robin_spawn(): - ctx = muliprocessing.get_context("spawn") + ctx = multiprocessing.get_context("spawn") p = ctx.Process(target=task_scheduler_round_robin_spawn) p.start() p.join()