TVM：使用 Schedule 模板和 AutoTVM 来优化算子

TVM：使用 Schedule 模板和 AutoTVM 来优化算子

在本文中，我们将介绍如何使用 TVM 张量表达式（Tensor Expression，TE）语言编写 Schedule 模板，AutoTVM 可以搜索通过这些模板找到最佳 Schedule。这个过程称为自动调整（Auto Tuning），它有助于自动优化张量计算的过程。

本教程需基于之前介绍的如何使用 TE 来写一个矩阵乘法的教程。

Auto Tuning 有两步：

• 第一步是定义一个搜索空间
• 第二步是运行相应的搜索算法来探索这个空间

本教程将展示如何在 TVM 中完成这两步，整个流程将以矩阵乘法为例。

安装依赖

要使用 TVM 中的 AutoTVM 包，需要安装这些额外的依赖：

pip install --user psutil xgboost cloudpickle

为了使 TVM 在 tuning 中运行得更快，建议使用 cython 作为 TVM 的 FFI。在 TVM 的根目录中，执行：

pip3 install --user cython
sudo make cython3

现在我们开始写 Python 代码，先引入包：

import logging
import sys

import numpy as np
import tvm
from tvm import te
import tvm.testing

from tvm import autotvm

TE 实现基本的矩阵乘法

回想一下使用TE实现矩阵乘法的基本方法。我们把它写在这里，稍作修改。我们将把乘法封装在Python 函数定义中。简单起见，我们将把注意力集中在分割优化（split optimization）上，使用一个固定值来定义重排（reordering）的块的大小。

def matmul_basic(N, L, M, dtype):

    A = te.placeholder((N, L), name="A", dtype=dtype)
    B = te.placeholder((L, M), name="B", dtype=dtype)

    k = te.reduce_axis((0, L), name="k")
    C = te.compute((N, M), lambda i, j: te.sum(A[i, k] * B[k, j], axis=k), name="C")
    s = te.create_schedule(C.op)

    # schedule
    y, x = s[C].op.axis
    k = s[C].op.reduce_axis[0]

    yo, yi = s[C].split(y, 8)
    xo, xi = s[C].split(x, 8)

    s[C].reorder(yo, xo, k, yi, xi)

    return s, [A, B, C]

使用 AutoTVM 优化矩阵乘法

在前面的明细表代码中，我们使用常量值 8 作为 tiling 因子。但是，它可能不是最好的，因为最佳的tiling 因子取决于实际的硬件环境和输入的形状。

如果希望 Schedule 代码能够在更大范围的输入形状和目标硬件上移植，最好定义一组候选值，并根据目标硬件上的测量结果选择最佳值。

在 autotvm 中，我们可以定义一个可调参数，或为此类参数定义一个“knob”。

一个基本的矩阵乘法模板

我们的一个示例，介绍如何为 spliting schedule 操作的块大小创建一个可调参数集。

# Matmul V1: List candidate values
@autotvm.template("tutorial/matmul_v1")  # 1. use a decorator
def matmul_v1(N, L, M, dtype):
    A = te.placeholder((N, L), name="A", dtype=dtype)
    B = te.placeholder((L, M), name="B", dtype=dtype)

    k = te.reduce_axis((0, L), name="k")
    C = te.compute((N, M), lambda i, j: te.sum(A[i, k] * B[k, j], axis=k), name="C")
    s = te.create_schedule(C.op)

    # schedule
    y, x = s[C].op.axis
    k = s[C].op.reduce_axis[0]

    # 2. get the config object
    cfg = autotvm.get_config()

    # 3. define search space
    cfg.define_knob("tile_y", [1, 2, 4, 8, 16])
    cfg.define_knob("tile_x", [1, 2, 4, 8, 16])

    # 4. schedule according to config
    yo, yi = s[C].split(y, cfg["tile_y"].val)
    xo, xi = s[C].split(x, cfg["tile_x"].val)

    s[C].reorder(yo, xo, k, yi, xi)

    return s, [A, B, C]

这里对于之前的 schedule 代码，我们有四处调整，从而可以得到一个可调节的 “模板”：

1. 使用装饰器 @autotvm.template() 来将此函数标记为一个模板

2. 获取配置项： cfg 可以视为本函数的一个参数，但是我们是通过一种不同的方式得到它的。有了这个参数，这个函数就不再是一个确定的 schedule。而是我们可以为这个函数传入不同的配置来得到不同的 schedules。一个像这样含有配置项的函数称为 “模板”。

   为了使得模板函数更加 compact，我们通过完成以下两件事情来在一个函数中定义参数搜索空间：

   • 通过一组值定义搜索空间。这是通过将 cfg 设置为一个 ConfigSpace 对象完成的。它会收集本函数中所有的可以调节的 knobs ，并从中建立一个搜索空间。
   • 根据此空间中的实体进行 schedule。这是通过将 cfg 设置为一个 ConfigEntity 对象完成的。当它是一个 ConfigEntity 时，它会忽略所有空间定义的API（即 cfg.define_XXX(...)）。相反，它会所有可调节的 knobs 保存确定的值，然后我们根据这些值进行 schedule。

   在 auto-tuning 时，我们将首先使用 ConfigSpace 对象调用此模板以构建搜索空间。然后，我们在构建空间中使用不同的 ConfigEntity 调用此模板，以获得不同的 schedule。最后，我们将度量由不同计划生成的代码，并选择最优的。

3. 定义两个可调节的 knobs。第一个是具有 5 个可能值的 tile_y。第二个是 tile_x，具有相同的可能值列表。这两个 knob 是独立的，因此它们会有大小为 25=5x5 的搜索空间。

4. 配置 knobs 被传递到 split schedule 操作，允许我们根据之前在 cfg 中定义的 5x5 确定值进行 schedule。

具有高级参数API的矩阵乘法模板

在前面的模板中，我们手动列出了 knobs 的所有可能值。这是定义空间的最低级别API，并提供要搜索的参数空间的显式枚举。但是，我们还提供了另一组API，可以使搜索空间的定义更简单、更智能。在可能的情况下，我们建议您使用此高级API。

在下面的示例中，我们使用 ConfigSpace.define_split 来定义拆分 knob。它将列举所有分割轴和构建空间的可能方法。

我们还有 ConfigSpace.define_reorder用于 reorder knob，ConfigSpace.define_annotation 用于展开、矢量化、线程绑定等注释。当高级API不能满足您的需求时，您可以随时使用低级API。

@autotvm.template("tutorial/matmul")
def matmul(N, L, M, dtype):
    A = te.placeholder((N, L), name="A", dtype=dtype)
    B = te.placeholder((L, M), name="B", dtype=dtype)

    k = te.reduce_axis((0, L), name="k")
    C = te.compute((N, M), lambda i, j: te.sum(A[i, k] * B[k, j], axis=k), name="C")
    s = te.create_schedule(C.op)

    # schedule
    y, x = s[C].op.axis
    k = s[C].op.reduce_axis[0]

    ##### define space begin #####
    cfg = autotvm.get_config()
    cfg.define_split("tile_y", y, num_outputs=2)
    cfg.define_split("tile_x", x, num_outputs=2)
    ##### define space end #####

    # schedule according to config
    yo, yi = cfg["tile_y"].apply(s, C, y)
    xo, xi = cfg["tile_x"].apply(s, C, x)

    s[C].reorder(yo, xo, k, yi, xi)

    return s, [A, B, C]

注意：More Explanation on cfg.define_split

更多关于 cfg.define_split 的解释

在此模板中，cfg.define_split(“tile_y”，y，num_outputs=2) 将枚举所有可能的组合，这些组合可以将y轴拆分为两个具有y长度因子的轴。例如，如果y的长度为32，我们希望使用32的因子将其拆分为两个轴，那么（外轴长度、内轴长度）对有6个可能的值，即（32,1）、（16,2）、（8,4）、（4,8）、（2,16）或（1,32）。这些都是 tile_y 的6个可能值。

在 schedule 期间，cfg[“tile_y”] 是一个 SplitEntity 对象。我们将外轴和内轴的长度存储在 cfg['tile_y'].size（包含两个元素的元组）中。在这个模板中，我们使用yo, yi=cfg['tile_y']] 来应用它。实际上，这相当于yo, yi=s[C].split（y，cfg[“tile\u y”].size[1]）或 yo, yi=s[C].split（y，npart=cfg[“tile\u y”].size[0]）

使用 cfg.apply API 的优点是，它使多级拆分（即当num_outputs>=3时）更加容易。

第二步：使用AutoTVM优化矩阵乘法

在第一步中，我们编写了一个矩阵乘法模板，该模板允许我们对分割 schedule 中使用的块大小进行参数化。现在我们可以在这个参数空间上进行搜索。下一步是选择一个调谐器（tuner）来指导这个空间的探索。

TVM 中的 Auto-tuners

调谐器的工作可以通过以下伪代码来描述：

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

提出下一批配置时，调谐器可以采取不同的策略。TVM提供的一些调谐器策略包括：

• tvm.autotvm.tuner.RandomTuner：按随机顺序枚举空间
• tvm.autotvm.tuner.GridSearchTuner：按网格搜索顺序枚举空间
• tvm.autotvm.tuner.GATuner：利用遗传算法进行空间搜索
• tvm.autotvm.tuner.XGBTuner：使用基于模型的方法。训练 XGBoost 模型预测降低 IR 的速度，并根据预测选择下一批。

我们可以根据空间大小、时间预算和其他因素选择调谐器。例如，如果您的空间非常小（小于1000），gridsearch 调谐器或随机调谐器就足够了。如果您的空间级别为 10^9（这是CUDA GPU上conv2d 算子的空间大小），XGBoostTuner 可以更高效地探索并找到更好的配置。

开始 tuning

这里我们继续我们的矩阵乘法示例。首先，我们创建一个 tuning 任务。我们还可以检查初始化的搜索空间。在这种情况下，对于 512x512 平方矩阵乘法，空间大小为 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)

此处输出：

N, L, M = 512, 512, 512
task = autotvm.task.create("tutorial/matmul", args=(N, L, M, "float32"), target="llvm")
print(task.config_space)

然后我们需要定义如何测量生成的代码并选择调谐器。因为我们的空间很小，随机调谐器也可以。

在本教程中，我们只做了 10 次试验来演示。实际上，你可以根据你的时间预算做更多的试验。我们将把调优结果记录到日志文件中。此文件可用于选择调谐器稍后发现的最佳配置。

# logging config (for printing tuning log to the screen)
logging.getLogger("autotvm").setLevel(logging.DEBUG)
logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout))

测量配置有两个步骤：构建和运行。默认情况下，我们使用所有CPU核来编译程序。然后我们依次测量它们。为了减少方差，我们进行了5次测量并取平均值。

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")],
)

调优 tuning 完成后，我们可以从日志文件中选择性能最好的配置，并使用相应的参数编译 schedule。我们还快速验证了 schedule 是否正确。我们可以直接在 autotvm.apply_history_best 上下文下调用函数 matmul。当我们调用此函数时，它将使用其参数查询分派上下文，并使用相同的参数获得最佳配置。

# 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.numpy(), rtol=1e-4)

最后的注意事项与总结

在本教程中，我们展示了如何构建算子模板，以使得 TVM 能够对参数空间进行搜索并选择最优的的 schedule 配置。为了更深入地了解其工作原理，我们建议在此示例上进行扩展，可以根据使用张量表达式（TE）入门教程中演示的 schedule 操作向计划添加新的搜索参数。在接下来的部分中，我们将演示 AutoScheduler，这是一种TVM优化常用算子的方法，用户无需提供用户定义的模板。

Ref：
https://tvm.apache.org/docs/tutorial/autotvm_matmul_x86.html