注意
单击此处下载完整的示例代码
电感器 CPU 后端调试和分析¶
创建时间: Jul 01, 2023 |上次更新时间: 2024 年 7 月 23 日 |上次验证: Nov 05, 2024
作者: Xuan Liao, Haozhe Zhu, Jiong Gong, Weihan Wang
概述¶
PyTorch 2.0 引入了名为 .
这项新功能通过由默认 Inductor 后端提供支持的图形级优化,显著提高了 Eager 模式的执行速度。torch.compile
本教程旨在对调试进行深入介绍
以及 Inductor CPU 后端的性能分析,深入研究 .torch.compile
同时,您还可以找到有关基本用法的相关教程,
全面的故障排除和 GPU 特定知识,例如 GPU 性能分析。torch.compile
我们将从一个触发编译问题和准确性问题的激励性示例开始调试 通过演示调试过程来查明问题。
通过启用日志记录并探索底层生成的代码, 您可以学习如何逐步缩小失败范围,并最终找出路由原因。
之后,我们将继续讨论如何分析编译后的代码,以及
通过与 Eager 模式的性能比较,
详细说明为什么 Can 提供额外的性能提升 与 Eager 对应物相比。torch.compile
调试¶
这是一个运行 using Inductor 并将其结果与 eager 模式进行比较的简单示例: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)
/usr/local/lib/python3.10/dist-packages/onnxscript/converter.py:820: FutureWarning:
'onnxscript.values.Op.param_schemas' is deprecated in version 0.1 and will be removed in the future. Please use '.op_signature' instead.
/usr/local/lib/python3.10/dist-packages/onnxscript/converter.py:820: FutureWarning:
'onnxscript.values.OnnxFunction.param_schemas' is deprecated in version 0.1 and will be removed in the future. Please use '.op_signature' instead.
codegen 中 的正确实现如下:neg
cpp
def neg1(x):
return f"decltype({x})(-{x})"
为了演示调试,我们稍后会将函数修改为错误的函数。
获取更多日志记录信息¶
如果默认运行此简单示例,则不会提供调试信息。为了获得更有用的调试和日志信息,我们通常会添加一个环境变量,如下所示:TORCH_COMPILE_DEBUG
TORCH_COMPILE_DEBUG=1 python xx.py
这将在输出日志中打印更多调试信息,并转储在 codegen 过程中生成的中间 IR。您可以在日志中找到转储的文件路径,如下所示:
torch._inductor.debug: [WARNING] model___20 debug trace: /tmp/torchinductor_root/rx/crxfi2ybd7yp5sbj2pnhw33wfhtdw7wumvrobyp5sjvdui5ktjc2.debug
在此目录中,保存了以下文件以供调试:
文件 |
描述 |
---|---|
|
可执行 FX 图,分解后,模式匹配前 |
|
形态匹配后的转换 FX 图 |
|
熔合前的电感器 IR |
|
熔合后的电感器 IR |
|
为图形生成的 Python 代码,使用 C++/Triton 内核 |
请注意,和 都是可运行和可编辑的,以便于调试。
以下是从文件中提取的代码的主要部分,我们将 C++ 生成的行与 FX 代码行相关联。fx_graph_runnable.py
output_code.py
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,)
C++ 内核中 :output_code
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;
}
}
}
}''')
确定错误的组成部分¶
当遇到错误或准确性问题时,查找错误的直接解决方案是缩小问题范围。首先要做的是确定发生错误的组件。幸运的是,它可以通过更改 的后端来实现。torch.compile
法典 |
描述 |
---|---|
|
启用 Dynamo |
|
启用 Dynamo + AOT Autograd |
|
启用 Dynamo + AOT Autograd + 电感器 |
如果模型可以在 backend 设置为 或 failing时成功运行,我们可以将失败范围缩小到 Inductor。eager
aot_eager
inductor
编译错误¶
正如我们所知,图级优化的进化链是这样的:
torch.neg (Python) -> torch.ops.aten.neg.default (within FX graph) -> ops.neg (within IR node) -> tmp2 = -tmp1 (within C++ kernel)
如果遇到编译错误,则表示在输出代码中编译 C++ 内核时出现问题。 这种类型的错误表示在降低 IR 节点以输出代码时引入了错误。 编译错误的根本原因通常显示在回溯日志中。
例如,该函数的修改方式如下: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);
| ^
我们还在 output code 和 IR 节点中查看相应的 C++ 内核。
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
根据回溯日志记录,编译错误是由 输入的数据类型不一致引起的。
通过检查 C++ 内核,我们知道 is 不再是 按原样执行。
我们可以轻松地将 和 在 C++ 内核中分别与 IR 节点和 IR 节点匹配。max_propagate_nan
tmp2
long
-
tmp0
long
-
max_propagate_nan
ops.neg
ops.maximum
现在我们成功发现,根本原因是 codegen 中的 implementation ,它在执行 .ops.neg
cpp
neg
精度调试¶
否则,如果模型运行时出现其他错误或准确性问题,您可以使用名为 Minifier 的 PyTorch 调试工具。
其核心思想是不断删除 graph 的节点和输入,直到找到有问题的最小 graph。
它通过 4 种策略帮助自动生成缩小的问题图:截断后缀、增量调试、消除死代码和删除未使用的输入。Minifier
现在,我们将在 的帮助下展示精度问题的调试过程。
精度问题是指 backends eager 和 inductor 的输出不同的情况。Minifer
例如,我们像这样修改示例:
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 后端编译的模型进行性能分析的过程。
在下面的示例中,我们使用 Eager 模式和 Inductor graph 模式对 Hugging Face Transformer 模型进行基准测试。
Inductor 的执行时间和加速比是在 benchmark 之后打印的。
我们使用 Intel(R) Xeon(R) Platinum 8358 CPU @ 2.60GHz,并在第一个插槽上运行基准测试,以演示本节中的优化。
我们将以下环境变量设置为在 Intel(R) CPU 上进行基准测试的最佳实践。MobileBertForQuestionAnswering
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}")
/usr/local/lib/python3.10/dist-packages/numpy/core/getlimits.py:518: UserWarning:
The value of the smallest subnormal for <class 'numpy.float32'> type is zero.
/usr/local/lib/python3.10/dist-packages/numpy/core/getlimits.py:89: UserWarning:
The value of the smallest subnormal for <class 'numpy.float32'> type is zero.
输出:
eager use: 802.1023553796113 ms/iter
inductor use: 339.95180135127157 ms/iter
speed up ratio: 2.359459053287382
在我们自己的测试中,我们发现 Inductor CPU 后端将模型速度提高了约 2.355 倍。
接下来,让我们深入研究操作级别的性能,以了解加速的来源。Pytorch Profiler 是一个很好的工具来帮助我们。
电感器 CPU 后端支持使用 configuration 选项将 fusion kernel 的时间报告给 profiler: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()
我们得到以下 eager-mode 模型的性能分析表(省略了一些列):
------------------------- ------------ ------------ ------------
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
从 Eager 模型的分析表中,我们可以看到最耗时的运算是 [, , ]。
与电感器模型分析表进行比较,我们注意到一个条目和多个 fused kernel 的形式为 。他们是主要的
inductor 模型正在进行的优化。让我们分别讨论它们。aten::addmm
aten::add
aten::copy_
aten::mul
aten::clamp_min
aten::bmm
mkl::_mkl_linear
graph_0_cpp_fused_*
(1) 关于 :您可能会注意到对此内核的调用次数为 362,这与 Eager Model 性能分析表中的调用次数完全相同。
CPU 总值为 376.888 毫秒,而 为 231.573 毫秒。这表明 “linear” 部分为 ~1.63x。
加速主要来自将权重张量打包为 block memory 格式并在 Inductor CPU 后端调用 cblas_sgemm_compute
在 GEMM 计算期间具有更好的缓存行为。mkl::_mkl_linear
aten::linear
aten::linear
mkl::_mkl_linear
(2) 关于其他内存密集型操作:在我们的测试中,eager/inductor 模型的端到端延迟为 802/339ms。因此,我们可以大致推断出其他内存密集型运算的速度约为 3.94 倍。
让我们阅读生成的代码,以了解 inductor 如何实现这一令人印象深刻的优化。您可以通过以下方式找到生成的代码
搜索范围cpp_fused__mkl_linear_add_mul_relu_151
output_code.py
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)));
}
}
}
}
}''')
从上面生成的代码中,我们可以看到这个内核在 上做了一个典型的 Loop Fusion。
这是一个受内存限制的瓶颈,无法获得良好的性能。为了更直观地了解此优化,
我们可以推断 inputs 的大小和步幅,并进一步对这个模式进行基准测试。[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
这只是一个例子。分析表显示,在此模型中,所有元件级运算都自动熔接在电感器内。您可以在 output_code.py 中阅读更多内核
结论¶
本文档提供了 Inductor CPU 后端的深入教程。
通过激励性示例,我们将逐步完成调试和分析的过程。 主要思想是缩小问题范围。
我们逐步演示了在调试日志记录和工具 Minifier 的帮助下,深入研究问题并找到故障根本原因的方法。 首先确定故障发生在哪个组件中,然后尝试生成可以重现故障的最小代码片段。
当 Inductor 的性能优于 Eager 模式时,我们提供了一种可靠的性能分析方法。 我们将展示如何使用 PyTorch Profiler 找到耗时的热点,并找出运算符级别或内核级别的原因来解释这种现象。
脚本总运行时间:(8 分 48.727 秒)