• 教程 >
  • Inductor CPU 後端除錯與效能分析
快捷方式

Inductor CPU 後端除錯與效能分析

建立日期:2023 年 7 月 1 日 | 最後更新:2025 年 1 月 8 日 | 最後驗證:2024 年 11 月 5 日

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

概述

PyTorch 2.0 引入了名為 torch.compile 的編譯 API。這項新功能透過由預設 Inductor 後端驅動的圖級別最佳化,提供了比 eager 模式執行顯著的加速。

本教程旨在深入探討 torch.compile 的細節,從而深入介紹 Inductor CPU 後端的除錯和效能分析。

同時,你可能還會找到關於 torch.compile 的相關教程,涵蓋基本用法、全面的故障排除以及 GPU 特定的知識,例如GPU 效能分析

我們將從一個引發編譯問題和精度問題的激勵性示例開始除錯,透過演示除錯過程來精確定位問題。

透過啟用日誌記錄和探索底層生成的程式碼,你可以學習如何一步步縮小故障範圍,最終找出根本原因。

接下來,我們將討論如何對編譯後的程式碼進行效能分析,並透過與 eager 模式的效能比較,闡述 torch.compile 相較於其 eager 對應物為何能提供額外的效能提升。

除錯

這是一個使用 Inductor 執行 torch.compile 並將其結果與 eager 模式進行比較的簡單示例

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。編譯錯誤的根本原因通常會顯示在回溯日誌中。

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

例如,neg 函式被修改如下

 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 節點。

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;
            }
        }
    }
}

C++ 核心

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

IR 節點

根據回溯日誌,編譯錯誤是由 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

例如,我們像這樣修改示例

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

並同時修改 neg 函式

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

將會引發一個精度問題,如下所示

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

要使用 Minifier 除錯精度問題,需要兩個環境變數

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

這會給我們提供日誌資訊,展示最小化的步驟

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

執行後,我們得到最終最小化的圖,其中包含目標節點 neg

有關 Minifier 的更多使用詳情,請參閱故障排除

效能分析

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}")

在本節中,我們將演示如何對使用 Inductor CPU 後端編譯的模型進行效能分析。在下面的示例中,我們使用 eager 模式和 Inductor 圖模式對 Hugging Face Transformer 模型 MobileBertForQuestionAnswering 進行基準測試。基準測試後會列印 Inductor 的執行時間和加速比。我們使用 英特爾(R) Xeon(R) Platinum 8358 CPU @ 2.60GHz 並在第一個 socket 上執行基準測試,以演示本節中的最佳化。我們將以下環境變數設定為在 英特爾(R) CPU 上進行基準測試的最佳實踐。

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

輸出

在我們自己的測試中,我們發現 Inductor CPU 後端將模型加速了約 2.355 倍。

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

接下來,讓我們深入研究操作層面的效能,以瞭解加速來自何處。PyTorch Profiler 是一個很好的工具來幫助我們。Inductor CPU 後端支援使用 enable_kernel_profile 配置選項向 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()

按照PyTorch Profiler 中的步驟,我們可以獲得性能分析表和跟蹤檔案。

-------------------------  ------------  ------------  ------------
                     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

我們得到了 eager 模式模型的以下效能分析表(省略了一些列)

-----------------------------------------------  ------------  ------------  ------------
                                           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

類似地,我們也獲得了使用 Inductor 編譯的模型的效能分析表(省略了一些列)

從 eager 模型的效能分析表中,我們可以看到最耗時的操作是 [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 次,這與 eager 模型效能分析表中的 aten::linear 完全相同。aten::linear 的 CPU 總時間為 376.888ms,而 mkl::_mkl_linear 的時間為 231.573ms。這表明“linear”部分的加速約為 1.63 倍。加速主要來自將權重張量打包成塊狀記憶體格式,並在 Inductor CPU 後端中呼叫cblas_sgemm_compute,以在 GEMM 計算期間獲得更好的快取行為。

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)));
                }
            }
        }
    }
}''')

(2) 關於其他記憶體密集型操作:在我們自己的測試中,eager/inductor 模型的端到端延遲分別為 802ms/339ms。因此,我們可以大致推斷其他記憶體密集型操作的加速約為 3.94 倍。讓我們閱讀生成的程式碼,瞭解 inductor 是如何實現這一令人印象深刻的最佳化的。你可以在 output_code.py 中搜索 cpp_fused__mkl_linear_add_mul_relu_151 來找到生成的程式碼。

# 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}")

在本節中,我們將演示如何對使用 Inductor CPU 後端編譯的模型進行效能分析。在下面的示例中,我們使用 eager 模式和 Inductor 圖模式對 Hugging Face Transformer 模型 MobileBertForQuestionAnswering 進行基準測試。基準測試後會列印 Inductor 的執行時間和加速比。我們使用 英特爾(R) Xeon(R) Platinum 8358 CPU @ 2.60GHz 並在第一個 socket 上執行基準測試,以演示本節中的最佳化。我們將以下環境變數設定為在 英特爾(R) CPU 上進行基準測試的最佳實踐。

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

從上面的生成程式碼中,我們可以看到該核心對 [add, add, mul, add] 執行了典型的迴圈融合。這是一個記憶體限制的瓶頸,阻礙了良好的效能。為了更直觀地瞭解這種最佳化,我們可以推斷輸入的大小和步長,並進一步對 [add, add, mul, add] 模式進行基準測試。

這只是一個示例。效能分析表顯示,在該模型中,Inductor 自動融合了所有逐元素操作。你可以在 output_code.py 中閱讀更多核心。

結論

本文件提供了關於 Inductor CPU 後端的深入教程。

透過激勵性示例,我們逐步介紹了除錯和效能分析的過程。核心思想是縮小問題範圍。

我們逐步演示瞭如何深入探究問題並找到故障的根本原因,這得益於除錯日誌和 Minifier 工具的幫助。首先確定故障發生在哪個元件中,然後嘗試生成能夠重現故障的最小程式碼片段。

當 Inductor 的效能優於 eager 模式時,我們提供了一種可靠的效能分析方法。我們展示瞭如何使用 PyTorch Profiler 找到耗時的熱點,並確定操作級別或核心級別的原因來解釋這種現象。

下載 Jupyter notebook: inductor_debug_cpu.ipynb



評價本教程

© 版權所有 2024, PyTorch。