评价此页

Inductor CPU 后端调试和性能分析#

创建日期: 2023年07月01日 | 最后更新: 2025年01月08日 | 最后验证: 2024年11月05日

作者: Xuan Liao, Haozhe Zhu, Jiong Gong, Weihan Wang

概述#

PyTorch 2.0 引入了一个名为 torch.compile 的编译 API。这项新功能通过默认的 Inductor 后端驱动的图级优化,在急切模式执行之上提供了显著的速度提升。

本教程旨在通过深入研究 torch.compile 的细节,提供关于 Inductor CPU 后端调试和性能分析的深度介绍。

同时,您也可以找到关于 torch.compile 的相关教程,包括 基本用法、全面的 故障排除 以及 GPU 特定知识,例如 GPU 性能分析

我们将以一个触发编译问题和精度问题的激励性示例开始调试,演示调试过程以查明问题。

通过启用日志记录并探索底层生成的代码,您可以学会如何逐步缩小故障范围,最终找出根本原因。

在此之后,我们将继续讨论如何分析编译后的代码的性能,并通过与急切模式进行性能比较,阐述 torch.compile 相较于其急切模式的对应项能提供额外性能提升的原因。

调试#

这是一个简单的示例,用于使用 Inductor 运行 torch.compile 并将其结果与急切模式进行比较

import torch

def foo1(x1, x2):
    a = torch.neg(x1)
    b = torch.maximum(x2, a)
    y = torch.cat([b], dim=0)
    return y

x1 = torch.randint(256, (1, 8), dtype=torch.uint8)
x2 = torch.randint(256, (8390, 8), dtype=torch.uint8)

compiled_foo1 = torch.compile(foo1)
result = compiled_foo1(x1, x2)

cpp 代码生成中的 neg 的正确实现如下

def neg1(x):
    return f"decltype({x})(-{x})"

为了演示调试,我们稍后会将其修改为错误的版本。

获取更多日志信息#

默认情况下运行此简单示例不会提供任何调试信息。为了获得更有用的调试和日志信息,我们通常会像下面这样添加一个 TORCH_COMPILE_DEBUG 环境变量

TORCH_COMPILE_DEBUG=1 python xx.py

这将会在输出日志中打印更多调试信息,并且还会转储在代码生成过程中生成的中间 IR。您可以在日志中找到转储的文件路径,如下所示

torch._inductor.debug: [WARNING] model___20 debug trace: /tmp/torchinductor_root/rx/crxfi2ybd7yp5sbj2pnhw33wfhtdw7wumvrobyp5sjvdui5ktjc2.debug

在此目录中,为调试目的保存了以下文件

文件

描述

fx_graph_runnable.py

可执行的 FX 图,分解后,模式匹配前

fx_graph_transformed.py

转换后的 FX 图,模式匹配后

ir_pre_fusion.txt

融合前的 Inductor IR

ir_post_fusion.txt

融合后的 Inductor IR

output_code.py

生成的图的 Python 代码,包含 C++/Triton 内核

请注意,fx_graph_runnable.pyoutput_code.py 都是可执行和可编辑的,以便于调试。以下是提取自文件的代码的主要部分,我们将 C++ 生成的行与 FX 代码行进行关联。

fx_graph_runnable:

def forward1(self, arg0_1, arg1_1):
    neg = torch.ops.aten.neg.default(arg0_1);  arg0_1 = None
    maximum = torch.ops.aten.maximum.default(arg1_1, neg);  arg1_1 = neg = None
    clone = torch.ops.aten.clone.default(maximum);  maximum = None
    return (clone,)

output_code 中的 C++ 内核

import torch
from torch._inductor.async_compile import AsyncCompile
async_compile = AsyncCompile()

cpp_fused_cat_maximum_neg_0 = async_compile.cpp('''
#include "/tmp/torchinductor_root/gv/cgv6n5aotqjo5w4vknjibhengeycuattfto532hkxpozszcgxr3x.h"
extern "C" void kernel(const unsigned char* in_ptr0,
                       const unsigned char* in_ptr1,
                       unsigned char* out_ptr0)
{
    {
        #pragma GCC ivdep
        for(long i0=static_cast<long>(0L); i0<static_cast<long>(8390L); i0+=static_cast<long>(1L))
        {
            #pragma GCC ivdep
            for(long i1=static_cast<long>(0L); i1<static_cast<long>(8L); i1+=static_cast<long>(1L))
            {
                auto tmp0 = in_ptr0[static_cast<long>(i1 + (8L*i0))];
                auto tmp1 = in_ptr1[static_cast<long>(i1)];
                // Corresponding FX code line: neg = torch.ops.aten.neg.default(arg0_1);  arg0_1 = None
                auto tmp2 = decltype(tmp1)(-tmp1);
                // Corresponding FX code line: maximum = torch.ops.aten.maximum.default(arg1_1, neg);  arg1_1 = neg = None
                auto tmp3 = max_propagate_nan(tmp0, tmp2);
                // Corresponding FX code line: clone = torch.ops.aten.clone.default(maximum);  maximum = None
                out_ptr0[static_cast<long>(i1 + (8L*i0))] = tmp3;
            }
        }
    }
}''')

确定错误组件#

当遇到错误或精度问题时,查找 bug 的一个简单方法是缩小问题的范围。首先要做的是确定错误发生的组件。幸运的是,这可以通过更改 torch.compile 的后端来实现。

代码

描述

torch.compile(fn, backend="eager")

启用 Dynamo

torch.compile(fn, backend="aot_eager")

启用 Dynamo + AOT Autograd

torch.compile(fn, backend="inductor")

启用 Dynamo + AOT Autograd + Inductor

如果模型在后端设置为 eageraot_eager 时可以成功运行,而在 inductor 下失败,那么我们可以将失败范围缩小到 Inductor。

编译错误#

我们知道,图级优化的演变链是这样的

torch.neg (Python) -> torch.ops.aten.neg.default (within FX graph) -> ops.neg (within IR node) -> tmp2 = -tmp1 (within C++ kernel)

如果您遇到编译错误,说明在编译输出代码中的 C++ 内核时出了问题。此类错误表明在将 IR 节点降低到输出代码时引入了 bug。编译错误的根本原因通常会显示在回溯日志中。

例如,neg 函数修改如下

def neg2(x):
    return f"-{x}"

日志给出了以下编译错误,原因相当清楚。

 torch._dynamo.exc.BackendCompilerFailed: backend='inductor' raised:
 CppCompileError: C++ compile error
 /tmp/torchinductor_root/xg/cxga5tk3b4lkwoxyigrtocjp5s7vc5cg2ikuscf6bk6pjqip2bhx.cpp: In function ‘void kernel(const unsigned char*, const unsigned char*, unsigned char*)’:
 /tmp/torchinductor_root/xg/cxga5tk3b4lkwoxyigrtocjp5s7vc5cg2ikuscf6bk6pjqip2bhx.cpp:17:57: error: no matching function for call to ‘max_propagate_nan(unsigned char&, int&)’
   17 |                 auto tmp3 = max_propagate_nan(tmp0, tmp2);
        |                                                         ^
 In file included from /tmp/torchinductor_root/xg/cxga5tk3b4lkwoxyigrtocjp5s7vc5cg2ikuscf6bk6pjqip2bhx.cpp:2:
 /tmp/torchinductor_root/gv/cgv6n5aotqjo5w4vknjibhengeycuattfto532hkxpozszcgxr3x.h:27:17: note: candidate: ‘template<class scalar_t> scalar_t max_propagate_nan(scalar_t, scalar_t)’
 27 | inline scalar_t max_propagate_nan(scalar_t a, scalar_t b) {
      |                 ^~~~~~~~~~~~~~~~~
 /tmp/torchinductor_root/gv/cgv6n5aotqjo5w4vknjibhengeycuattfto532hkxpozszcgxr3x.h:27:17: note:   template argument deduction/substitution failed:
/tmp/torchinductor_root/xg/cxga5tk3b4lkwoxyigrtocjp5s7vc5cg2ikuscf6bk6pjqip2bhx.cpp:17:57: note:   deduced conflicting types for parameter ‘scalar_t’ (‘unsigned char’ and ‘int’)
 17 |                 auto tmp3 = max_propagate_nan(tmp0, tmp2);
      |                                                         ^

我们也可以看看输出代码中相应的 C++ 内核和 IR 节点。

C++ 内核

include "/tmp/torchinductor_root/gv/cgv6n5aotqjo5w4vknjibhengeycuattfto532hkxpozszcgxr3x.h"
extern "C" void kernel(const unsigned char* in_ptr0,
                    const unsigned char* in_ptr1,
                    unsigned char* out_ptr0)
{
    {
        #pragma GCC ivdep
        for(long i0=static_cast<long>(0L); i0<static_cast<long>(8390L); i0+=static_cast<long>(1L))
        {
            #pragma GCC ivdep
            for(long i1=static_cast<long>(0L); i1<static_cast<long>(8L); i1+=static_cast<long>(1L))
            {
                auto tmp0 = in_ptr0[static_cast<long>(i1 + (8L*i0))];
                auto tmp1 = in_ptr1[static_cast<long>(i1)];
                auto tmp2 = -tmp1;
                auto tmp3 = max_propagate_nan(tmp0, tmp2);
                out_ptr0[static_cast<long>(i1 + (8L*i0))] = tmp3;
            }
        }
    }
}

IR 节点

buf0: SchedulerNode(ComputedBuffer)
buf0.writes = [MemoryDep('buf0', c0, {c0: 67120})]
buf0.unmet_dependencies = []
buf0.met_dependencies =
    [   MemoryDep('arg0_1', c1, {c0: 8390, c1: 8}),
        MemoryDep('arg1_1', c0, {c0: 67120})]
buf0.users = [NodeUser(node=OUTPUT, can_inplace=False)]
buf0.group.device = cpu
buf0.group.iteration = ((8390, 8), ())
buf0.sizes = ([8390, 8], [])
class buf0_loop_body:
    var_ranges = {z0: 8390, z1: 8}
    index0 = 8*z0 + z1
    index1 = z1
    def body(self, ops):
        get_index = self.get_index('index0')
        load = ops.load('arg1_1', get_index)
        get_index_1 = self.get_index('index1')
        load_1 = ops.load('arg0_1', get_index_1)
        neg = ops.neg(load_1)
        maximum = ops.maximum(load, neg)
        get_index_2 = self.get_index('index0')
        store = ops.store('buf0', get_index_2, maximum, None)
        return store

根据回溯日志,编译错误是由 max_propagate_nan 输入的数据类型不一致引起的。通过检查 C++ 内核,我们知道在执行 - 操作后,tmp2 不再是 long 类型,因为 tmp0long 类型。我们可以轻松地将 C++ 内核中的 -max_propagate_nan 分别与 IR 节点中的 ops.negops.maximum 匹配起来。

现在我们成功找到了根本原因,即 cpp 代码生成中 ops.neg 的实现,它在执行 neg 操作时会默默地改变数据类型。

精度调试#

否则,如果模型遇到其他错误或精度问题,您可以使用 PyTorch 调试工具 Minifier

Minifier 的核心思想是不断删除图的节点和输入,直到找到存在问题的最小图。它通过 4 种策略自动生成最小化的问题图:截断后缀、增量调试、消除死代码和移除未使用的输入。

现在,我们将借助 Minifer 来展示精度问题的调试过程。精度问题是指后端 eager 和 inductor 的输出不同的情况。

例如,我们修改示例如下

from torch._dynamo.utils import same

def foo2(x1, x2):
    a = torch.neg(x1)
    b = torch.maximum(x2, a)
    y = torch.cat([b], dim=0)
    return y

x1 = torch.randn((1, 8), dtype=torch.float32)
x2 = torch.randn((8390, 8), dtype=torch.float32)

expected_result = foo2(x1, x2)

compiled_foo2 = torch.compile(foo2)
actual_result = compiled_foo2(x1, x2)

assert same(expected_result, actual_result) == True

并且还修改了 neg 函数

def neg3(x):
    return f"decltype({x})(2 * {x})"

会引发如下精度问题

torch._dynamo.utils: [ERROR] Accuracy failed: allclose not within tol=0.0001
Traceback (most recent call last):
  File "test_script.py", line 18, in <module>
    assert same(expected_result, actual_result) == True
AssertionError

要使用 Minifier 调试精度问题,需要两个环境变量

TORCHDYNAMO_REPRO_AFTER="aot" TORCHDYNAMO_REPRO_LEVEL=4 python xx.py

这为我们提供了展示最小化过程的日志信息

Started off with 6 nodes

Trying granularity 2
Strategy: Truncate suffix (G: 2) (6 nodes, 2 inputs)
SUCCESS: Went from 6 to 4 nodes

Trying granularity 4
Strategy: Remove unused inputs (G: 4) (4 nodes, 2 inputs)
SUCCESS: Went from 4 to 3 nodes

运行后,我们得到最终最小化后的图,目标节点为 neg

def forward2(self, arg0_1):
    neg = torch.ops.aten.neg.default(arg0_1);  arg0_1 = None
    return (neg,)

有关 Minifier 的更多使用细节,请参阅 故障排除

性能分析#

在本节中,我们将演示对使用 Inductor CPU 后端编译的模型进行性能分析的过程。在下面的示例中,我们将对 Hugging Face Transformer 模型 MobileBertForQuestionAnswering 进行急切模式和 Inductor 图模式的基准测试。在基准测试后,将打印 Inductor 的执行时间和加速比。我们使用 Intel(R) Xeon(R) Platinum 8358 CPU @ 2.60GHz 并在第一个套接字上运行基准测试,以演示此部分的优化。我们设置了以下环境变量作为在 Intel(R) CPU 上进行基准测试的最佳实践。

export KMP_BLOCKTIME=1
export KMP_SETTINGS=1
export KMP_AFFINITY=granularity=fine,compact,1,0
export LD_PRELOAD=${CONDA_PREFIX:-"$(dirname $(which conda))/../"}/lib/libiomp5.so:${CONDA_PREFIX:-"$(dirname $(which conda))/../"}/lib/libjemalloc.so
export MALLOC_CONF="oversize_threshold:1,background_thread:true,metadata_thp:auto,dirty_decay_ms:-1,muzzy_decay_ms:-1"
numactl -C 0-31 -m 0 python bench.py
# bench.py
from transformers import MobileBertForQuestionAnswering
# Initialize an eager model
model = MobileBertForQuestionAnswering.from_pretrained("csarron/mobilebert-uncased-squad-v2")
seq_length = 128
bs = 128
vocab_size = model.config.vocab_size
input = torch.randint(0, vocab_size, (bs, seq_length), dtype=torch.int64)
input_dict = {"input_ids": input}

# Initialize the inductor model
compiled_model = torch.compile(model)
with torch.no_grad():
    compiled_model(**input_dict)

NUM_ITERS=50
import timeit
with torch.no_grad():
    # warmup
    for _ in range(10):
        model(**input_dict)
    eager_t = timeit.timeit("model(**input_dict)", number=NUM_ITERS, globals=globals())

with torch.no_grad():
    # warmup
    for _ in range(10):
        compiled_model(**input_dict)
    inductor_t = timeit.timeit("compiled_model(**input_dict)", number=NUM_ITERS, globals=globals())
# print(f"eager use: {eager_t * 1000 / NUM_ITERS} ms/iter")
# print(f"inductor use: {inductor_t * 1000 / NUM_ITERS} ms/iter")
# print(f"speed up ratio: {eager_t / inductor_t}")

输出

eager use: 802.1023553796113 ms/iter
inductor use: 339.95180135127157 ms/iter
speed up ratio: 2.359459053287382

在我们的测试中,发现 Inductor CPU 后端将模型加速了约 2.355 倍。

接下来,让我们深入了解操作级别的性能,以了解加速的来源。Pytorch Profiler 是一个很好的工具来帮助我们。Inductor CPU 后端支持通过 enable_kernel_profile 配置选项将融合内核的时间报告给分析器。

from torch._inductor import config
config.cpp.enable_kernel_profile = True

遵循 Pytorch Profiler 中的步骤,我们能够获得性能分析表和跟踪文件。

# bench.py
from torch.profiler import profile, schedule, ProfilerActivity
RESULT_DIR = "./prof_trace"
my_schedule = schedule(
    skip_first=10,
    wait=5,
    warmup=5,
    active=1,
    repeat=5)

def trace_handler(p):
    output = p.key_averages().table(sort_by="self_cpu_time_total", row_limit=20)
    # print(output)
    p.export_chrome_trace(f"{RESULT_DIR}/{p.step_num}.json")

for _ in range(10):
    model(**input_dict)  # compiled_model(**input_dict) to get inductor model profiling

total = 0
with profile(
    activities=[ProfilerActivity.CPU],
    schedule=my_schedule,
    on_trace_ready=trace_handler
) as p:
    for _ in range(50):
        model(**input_dict)  # compiled_model(**input_dict) to get inductor model profiling
        p.step()

我们为急切模式模型获得了以下性能分析表(省略了部分列)

-------------------------  ------------  ------------  ------------
                     Name   CPU total %     CPU total    # of Calls
-------------------------  ------------  ------------  ------------
              aten::addmm        45.73%     370.814ms           362
                aten::add        19.89%     161.276ms           363
              aten::copy_        14.97%     121.416ms           488
                aten::mul         9.02%      73.154ms           194
          aten::clamp_min         8.81%      71.444ms            96
                aten::bmm         5.46%      44.258ms            48
            ProfilerStep*       100.00%     810.920ms             1
                aten::div         2.89%      23.447ms            24
           aten::_softmax         1.00%       8.087ms            24
             aten::linear        46.48%     376.888ms           362
              aten::clone         2.77%      22.430ms            98
                  aten::t         0.31%       2.502ms           362
               aten::view         0.14%       1.161ms           850
          aten::transpose         0.17%       1.377ms           386
       aten::index_select         0.12%     952.000us             3
             aten::expand         0.12%     986.000us           458
             aten::matmul         8.31%      67.420ms            48
                aten::cat         0.09%     703.000us             1
         aten::as_strided         0.08%     656.000us           963
               aten::relu         8.86%      71.864ms            96
-------------------------  ------------  ------------  ------------
Self CPU time total: 810.920ms

同样,我们也为使用 Inductor 编译的模型获得了表格(省略了部分列)

-----------------------------------------------  ------------  ------------  ------------
                                           Name   CPU total %     CPU total    # of Calls
-----------------------------------------------  ------------  ------------  ------------
                               mkl::_mkl_linear        68.79%     231.573ms           362
                                      aten::bmm         8.02%      26.992ms            48
                                  ProfilerStep*       100.00%     336.642ms             1
  graph_0_cpp_fused_constant_pad_nd_embedding_0         0.27%     915.000us             1
                                    aten::empty         0.27%     911.000us           362
 graph_0_cpp_fused__mkl_linear_add_mul_relu_151         0.27%     901.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_226         0.27%     899.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_361         0.27%     898.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_121         0.27%     895.000us             1
  graph_0_cpp_fused__mkl_linear_add_mul_relu_31         0.27%     893.000us             1
  graph_0_cpp_fused__mkl_linear_add_mul_relu_76         0.26%     892.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_256         0.26%     892.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_346         0.26%     892.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_241         0.26%     891.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_316         0.26%     891.000us             1
  graph_0_cpp_fused__mkl_linear_add_mul_relu_91         0.26%     890.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_106         0.26%     890.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_211         0.26%     890.000us             1
  graph_0_cpp_fused__mkl_linear_add_mul_relu_61         0.26%     889.000us             1
 graph_0_cpp_fused__mkl_linear_add_mul_relu_286         0.26%     889.000us             1
-----------------------------------------------  ------------  ------------  ------------
Self CPU time total: 336.642ms

从急切模式模型的分析表中,我们可以看到耗时最多的算子是 [aten::addmm, aten::add, aten::copy_, aten::mul, aten::clamp_min, aten::bmm]。与 Inductor 模型分析表相比,我们注意到一个 mkl::_mkl_linear 条目和多个形式为 graph_0_cpp_fused_* 的融合内核。这些是 Inductor 模型进行的主要优化。让我们分别讨论它们。

(1) 关于 mkl::_mkl_linear:您可能会注意到此内核的调用次数为 362,这与急切模式分析表中的 aten::linear 调用次数完全相同。急切模式 aten::linear 的 CPU 总计时间为 376.888ms,而 mkl::_mkl_linear 的总计时间为 231.573ms。这表明“线性”部分的加速约为 1.63 倍。加速主要来自于将权重张量 打包成块状内存格式 并调用 cblas_sgemm_compute,这是 Inductor CPU 后端在 GEMM 计算中实现更好的缓存行为。

(2) 关于其他内存密集型算子:在我们测试中,急切/Inductor 模型的端到端延迟分别为 802/339ms。因此,我们可以大致推断其他内存密集型算子的加速约为 3.94 倍。让我们阅读生成的代码,了解 Inductor 如何实现这种令人印象深刻的优化。您可以在 output_code.py 中搜索 cpp_fused__mkl_linear_add_mul_relu_151 来找到生成的代码。

cpp_fused__mkl_linear_add_mul_relu_151 = async_compile.cpp('''
#include <ATen/record_function.h>
#include "/tmp/torchinductor_root/lr/clrlgu27q4ggd472umdzwsu6qcpqxcuusjxqvx2hwitjbujiiz7z.h"
extern "C" void kernel(float* in_out_ptr0,
                       const float* in_ptr0,
                       const float* in_ptr1,
                       const float* in_ptr2,
                       const float* in_ptr3)
{
    RECORD_FUNCTION("graph_0_cpp_fused__mkl_linear_add_mul_relu_151", c10::ArrayRef<c10::IValue>({}));
    #pragma omp parallel num_threads(32)
    {
        {
            #pragma omp for
            for(long i0=static_cast<long>(0L); i0<static_cast<long>(16384L); i0+=static_cast<long>(1L))
            {
                for(long i1=static_cast<long>(0L); i1<static_cast<long>(512L); i1+=static_cast<long>(8L))
                {
                    auto tmp0 = at::vec::Vectorized<float>::loadu(in_ptr0 + static_cast<long>(i1 + (512L*i0)));
                    auto tmp1 = at::vec::Vectorized<float>::loadu(in_ptr1 + static_cast<long>(i1));
                    auto tmp3 = at::vec::Vectorized<float>::loadu(in_out_ptr0 + static_cast<long>(i1 + (512L*i0)));
                    auto tmp5 = at::vec::Vectorized<float>::loadu(in_ptr2 + static_cast<long>(i1));
                    auto tmp7 = at::vec::Vectorized<float>::loadu(in_ptr3 + static_cast<long>(i1));
                    auto tmp2 = tmp0 + tmp1;
                    auto tmp4 = tmp2 + tmp3;
                    auto tmp6 = tmp4 * tmp5;
                    auto tmp8 = tmp6 + tmp7;
                    tmp8.store(in_out_ptr0 + static_cast<long>(i1 + (512L*i0)));
                }
            }
        }
    }
}''')

从上面的生成的代码可以看出,该内核对 [add, add, mul, add] 进行了典型的 循环融合。这是阻碍良好性能的内存瓶颈。为了更直观地感受这种优化,我们可以推断输入的大小和步幅,并对这种 [add, add, mul, add] 模式进行进一步的基准测试。

# bench.py
def func(arg_0, arg_1, arg_2, arg_3, arg_4):
    add_0 = arg_0 + arg_1
    add_1 = add_0 + arg_2
    mul_1 = add_1 * arg_3
    add_2 = mul_1 + arg_4
    arg_2 = add_2
    return arg_2

arg_0 = torch.rand(16384, 512)
arg_1 = torch.rand(1, 512)
arg_2 = torch.zeros(16384, 512)
arg_3 = torch.rand(1, 512)
arg_4 = torch.rand(1, 512)

input = (arg_0, arg_1, arg_2, arg_3, arg_4)
inductor_func = torch.compile(func)
with torch.no_grad():
    inductor_func(*input)

import timeit
NUM_ITERS=100
with torch.no_grad():
    # warmup
    for _ in range(10):
        func(*input)
    eager_t = timeit.timeit("func(*input)", number=NUM_ITERS, globals=globals())

with torch.no_grad():
    # warmup
    for _ in range(10):
        inductor_func(*input)
    inductor_t = timeit.timeit("inductor_func(*input)", number=NUM_ITERS, globals=globals())
# print(f"eager use: {eager_t * 1000 / NUM_ITERS} ms/iter")
# print(f"inductor use: {inductor_t * 1000 / NUM_ITERS} ms/iter")
# print(f"speed up ratio: {eager_t / inductor_t}")

输出

eager use: 5.780875144992024 ms/iter
inductor use: 0.9588955780491233 ms/iter
speed up ratio: 6.0286805751604735

这只是一个示例。分析表显示,在此模型中,所有元素级算子都已在 Inductor 中自动融合。您可以在 output_code.py 中阅读更多内核。

结论#

该文档为 Inductor CPU 后端提供了深度教程。

通过激励性的示例,我们逐步介绍了调试和性能分析的过程。主要思想是缩小问题的范围。

我们通过调试日志和 Minifier 工具,一步一步地演示了如何深入研究问题并找到失败的根本原因。首先确定失败发生在哪个组件,然后尝试生成能够重现失败的最小代码片段。

当 Inductor 的性能优于急切模式时,我们提供了一个扎实的分析方法来进行性能分析。我们展示了如何使用 PyTorch Profiler 找到耗时的热点,并找出算子级别或内核级别的根本原因来解释这种现象。

脚本总运行时间: (10 分钟 40.868 秒)