[AutoTVM] Find the default optimization configuration of the kernel

sunggg · March 25, 2020, 1:54pm

Hi, I’m trying to get the default schedule configuration for each kernel in tuner modules. e.g., [(‘tile_f’, [-1, 64, 1, 1]), (‘tile_y’, [-1, 1, 1, 7]), …]

I’m bit running out-of-time so I’m asking here to find any luck while looking into the code. Any advice or suggestion would be greatly appreciated. Thanks!

robeastbme · March 26, 2020, 3:09pm

Hi,

I’m not fully sure whether I understood your questions correctly. If you are looking for the default schedule for a selected target, TVM fetches that from this repository during execution time and places a copy under /home/yourusername/.tvm/tophub/

Cheers, Robert

sunggg · March 26, 2020, 7:01pm

Thank you for the advice!

When looking into the autotuners (python/tvm/autotvm/tuner/), they generate a set of optimization configurations that they want to apply to the compilation of each kernel.
(If 1,000 trials are given, they will generate 1,000 configurations and pick the best one)

In this context, I’m wondering what is the default configuration is for each kernel. If default setting is not applying any optimization, I’m not sure how I can trigger compilation and measurement without any optimization inside the tuner module.

Thanks!

robeastbme · March 27, 2020, 9:25am

I see! That’s a very good question. I’m not sure to be honest. I wanted to look into it myself for measuring a performance baseline for each kernel before autotuning. I’ll look into it the next days when I have spare time. Please let me know if you have found something in the meanwhile…

comaniac · March 27, 2020, 4:50pm

It depends on the schedule function. For example:

Directly define a default configuration.

github.com

apache/incubator-tvm/blob/master/topi/python/topi/x86/conv2d.py#L36


from .. import nn
from ..nn.conv2d import conv2d_infer_layout, _get_workload as _get_conv2d_workload
from ..nn.conv2d import unpack_NCHWc_to_nchw
from ..nn.depthwise_conv2d import _get_workload as _get_depthwise_conv2d_workload
from ..nn.util import get_pad_tuple
from ..util import get_const_tuple, traverse_inline
from . import conv2d_avx_1x1, conv2d_avx_common


logger = logging.getLogger('topi')


def _get_default_config(cfg, data, kernel, strides, padding, out_dtype, is_depthwise=False,
                        layout='NCHW'):
    """
    Get default schedule config for the workload
    """
    static_data_shape = []
    for dim in get_const_tuple(data.shape):
        if isinstance(dim, tvm.tir.Var):
            static_data_shape.append(1)
        else:
            static_data_shape.append(dim)

Load the best configurations of similar workloads and make a one.

github.com

apache/incubator-tvm/blob/master/topi/python/topi/cuda/conv2d_direct.py#L45


cfg.define_split("tile_rx", rx, num_outputs=2)
cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])


target = tvm.target.Target.current()
if target.target_name in ['nvptx', 'rocm']:
    cfg.define_knob("unroll_explicit", [1])
else:
    cfg.define_knob("unroll_explicit", [0, 1])


# fallback support
if cfg.is_fallback:
    ref_log = autotvm.tophub.load_reference_log(
        target.target_name, target.model, 'conv2d_nchw.cuda')
    cfg.fallback_with_reference_log(ref_log)
##### space definition end #####


pad_data, kernel = s[conv].op.input_tensors


s[pad_data].compute_inline()
if isinstance(kernel.op, tvm.te.ComputeOp) and 'dilate' in kernel.op.tag:
    s[kernel].compute_inline()

github.com

apache/incubator-tvm/blob/master/python/tvm/autotvm/task/space.py#L1004


            if f > constraints[i]:
                find = j - 1
                break


        if find >= 0:
            entity.size[i] = factors[find]
            now //= factors[find]
        else:
            raise RuntimeError("Cannot find feasible fallback split entity for node: " + name)


def fallback_with_reference_log(self, ref_log):
    """A data driven fallback mechanism.
    We use tuned parameters from TopHub as reference data.
    For an unseen shape, we find the most similar tuned one from TopHub and
    mimic its parameters.
    Note that we are not matching by workload (e.g., input size, kernel size),
    but instead matching by configuration space. The idea is that if two workloads have
    similar configuration space, their optimal configurations are also likely to be similar.


    Parameters
    ----------

sunggg · March 27, 2020, 7:53pm

Thanks for the advice. This seems like a good start.

I’ve been playing around with auto-tuner module and thinking that we may utilize the performance of the default schedule during the tuning process.

2nd approach seems more relevant for my purpose. Rather than defining a default configuration by myself as the 1st approach, I want to measure the performance of the current default configuration in the tuner model.

Is fallback_with_reference_log the one I can use? I’m not sure since it does not return any configuration. (Seems like it applies directly to the kernel?)

As far as I know, I need a set of optimization configurations (config in the following code) for the default setting to trigger the measurement inside the tuner.

github.com

apache/incubator-tvm/blob/master/python/tvm/autotvm/tuner/tuner.py#L122


self.early_stopping = early_stopping


old_level = logger.level


GLOBAL_SCOPE.in_tuning = True
i = error_ct = 0
while i < n_trial:
    if not self.has_next():
        break


    configs = self.next_batch(min(n_parallel, n_trial - i))


    inputs = [MeasureInput(self.task.target, self.task, config) for config in configs]
    results = measure_batch(inputs)


    # keep best config
    for k, (inp, res) in enumerate(zip(inputs, results)):
        config = inp.config
        if res.error_no == 0:
            flops = inp.task.flop / np.mean(res.costs)
            error_ct = 0

Thank you so much for the great advices!

comaniac · March 27, 2020, 8:02pm

Is fallback_with_reference_log the one I can use? I’m not sure since it does not return any configuration. (Seems like it applies directly to the kernel?)

Yes this is the one. It didn’t return because it is a method of a config entity. What it did was copying the config map from other config enrities to itself, such as https://github.com/apache/incubator-tvm/blob/master/python/tvm/autotvm/task/space.py#L1045.

In this case, you don’t need to run AutoTVM for the fallback config. When you build the model without calling apply_history_best, the fallback config will be used automatically.

sunggg · March 28, 2020, 12:02am

Thank you for the help. Here’s what I’ve learned and found issues.

This is what I tried inside the my custom tuner.

  cfg = FallbackConfigEntity()
  func_name = self.task.name          // conv2d_nchw_winograd.cuda
  ref_log = autotvm.tophub.load_referernce_log( target.target_name, 
                                                '1080ti', func_name)
  print(cfg)                          // Returns 'None'

Automatic device model selection in line 229 did not work for my case: when selecting alternative model, current code does not check whether it contains data for the workload and pick the wrong model. So, currently, I picked the valid model manually.

github.com

apache/incubator-tvm/blob/master/python/tvm/autotvm/tophub.py#L229


        download_package(tophub_location, package_name)
    if os.path.isfile(filename): # in case download failed
        find = False
        inp = None
        counts = {}
        for inp, res in load_from_file(filename):
            counts[inp.target.model] = counts.get(inp.target.model, 0) + 1
            if model == inp.target.model:
                find = True
                break
        # if device model is not find, use the device model with the most tuned workloads
        if not find and counts:
            model = max(counts.items(), key=lambda k: k[1])[0]


        for inp, res in load_from_file(filename):
            if model == inp.target.model and inp.task.workload[0] == workload_name:
                tmp.append((inp, res))
    REFERENCE_LOG_CACHE[key] = tmp


return REFERENCE_LOG_CACHE[key]

Is there any reason for FallbackConfigEntity to inherit ConfigSpace, not the ConfigEntity? In my code, cfg.fallback_with_reference_log updates cfg with reference configuration but since it is not a ConfigEntity object, tuner cannot handle this object. For example, index field does not exist with FallbackConfigEntity although it loads the best knobs from the reference. When I pass this object to the tuner, it complains the runtime error with error_no=4.

github.com

apache/incubator-tvm/blob/master/python/tvm/autotvm/task/space.py#L954


            else:
                raise RuntimeError("Invalid config knob type: " + knob_type)
            entity_map[str(key)] = entity


        return ConfigEntity(index, code_hash, entity_map, constraints)


    def __repr__(self):
        return "%s,%s,%d" % (str(self._entity_map)[12:-1], self.code_hash, self.index)




class FallbackConfigEntity(ConfigSpace):
    """The config entity created to support fallback"""


    def __init__(self):
        super(FallbackConfigEntity, self).__init__()
        self.is_fallback = True


    def fallback_split(self, name, constraints):
        """Fallback a split knob


        Parameters