自動為自定義核生成轉換器

我們將演示如何利用 TensorRT 10.8 中的全新基於 Python 的外掛系統，使用 Torch-TensorRT 自動為自定義核生成轉換器。

如果 Torch-TensorRT 不知道如何在 TensorRT 中編譯某個操作，它支援回退到 PyTorch 實現。但是，這會帶來圖中斷的開銷，並會降低模型的效能。解決缺乏對運算元支援的最簡單方法是新增一個分解（參見：為 Dynamo 前端編寫降級過程）——它將運算元定義為 Torch-TensorRT 中支援的 PyTorch 運算元；或者一個轉換器（參見：為 Dynamo 前端編寫轉換器）——它將運算元定義為 TensorRT 運算元。

在某些情況下，這兩種方法都不太適用，可能是因為該運算元是一個自定義核，不是標準 PyTorch 的一部分，或者 TensorRT 無法原生支援它。

對於這些情況，可以使用 TensorRT 外掛在 TensorRT 引擎內部替換該運算元，從而避免圖中斷帶來的效能和資源開銷。

以前這涉及一個複雜的過程，不僅要構建一個高效能核，還要將其設定為在 TensorRT 中執行（參見：在 Torch-TensorRT 中使用自定義核處理 TensorRT 引擎）。藉助 TensorRT 10.8，有一個全新的 Python 原生外掛系統，極大地簡化了這一過程。該外掛系統還允許 Torch-TensorRT 自動生成必要的轉換程式碼，將 PyTorch 中的操作轉換為 TensorRT。

在 PyTorch 中編寫自定義運算元

先前的教程已經涵蓋了在 PyTorch 中建立自定義運算元，這些運算元後續將與 Torch-TensorRT 一起使用。

在這裡，我們在 Triton 中定義一個簡單的逐元素乘法運算元。然後，此運算元在 PyTorch 中註冊為自定義運算元，包括其主機啟動程式碼以及一個“Meta-核”。Meta-核是一個描述運算元將執行的形狀和資料型別轉換的函式。Dynamo 和 Torch-TensorRT 都會使用此 Meta-核，因此必須對其進行定義。

from typing import Tuple

import tensorrt.plugin as trtp
import torch
import torch_tensorrt
import triton
import triton.language as tl


@triton.jit
def elementwise_mul_kernel(X, Y, Z, BLOCK_SIZE: tl.constexpr):
    # Program ID determines the block of data each thread will process
    pid = tl.program_id(0)
    # Compute the range of elements that this thread block will work on
    block_start = pid * BLOCK_SIZE
    # Range of indices this thread will handle
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    # Load elements from the X and Y tensors
    x_vals = tl.load(X + offsets)
    y_vals = tl.load(Y + offsets)
    # Perform the element-wise multiplication
    z_vals = x_vals * y_vals
    # Store the result in Z
    tl.store(Z + offsets, z_vals)


@torch.library.custom_op("torchtrt_ex::elementwise_mul", mutates_args=())  # type: ignore[misc]
def elementwise_mul(
    X: torch.Tensor, Y: torch.Tensor, b: float = 0.2, a: int = 2
) -> torch.Tensor:
    # Ensure the tensors are on the GPU
    assert X.is_cuda and Y.is_cuda, "Tensors must be on CUDA device."
    assert X.shape == Y.shape, "Tensors must have the same shape."

    # Create output tensor
    Z = torch.empty_like(X)

    # Define block size
    BLOCK_SIZE = 1024

    # Grid of programs
    grid = lambda meta: (X.numel() // meta["BLOCK_SIZE"],)

    # Launch the kernel
    elementwise_mul_kernel[grid](X, Y, Z, BLOCK_SIZE=BLOCK_SIZE)

    return Z

逐元素操作的 Meta-核就是其中一個輸入的形狀和資料型別，因為我們在操作過程中不會改變形狀。

@torch.library.register_fake("torchtrt_ex::elementwise_mul")
def _(x: torch.Tensor, y: torch.Tensor, b: float = 0.2, a: int = 2) -> torch.Tensor:
    return x

使用快速部署外掛系統為 TensorRT 編寫外掛

TensorRT 10.8 中的快速部署外掛系統允許在 Python 中建立自定義外掛，且樣板程式碼顯著減少。它使用一個與 PyTorch 類似的系統，您可以在其中定義一個描述運算元將執行的形狀和資料型別轉換的函式，然後定義在給定 GPU 視訊記憶體控制代碼的情況下啟動核的程式碼。

正如 PyTorch Meta-核一樣，輸入和輸出之間沒有形狀或資料型別的轉換，因此我們可以直接告訴 TensorRT 預期與輸入相同的形狀。

@trtp.register("torchtrt_ex::elementwise_mul")
def _(
    x: trtp.TensorDesc, y: trtp.TensorDesc, b: float, a: int
) -> Tuple[trtp.TensorDesc]:
    return x.like()

在這裡，我們重用與 PyTorch 相似的主機啟動程式碼，但我們需要在啟動核之前將 TensorRT 張量轉換為 PyTorch 張量。這些操作也是原地操作，因此結果必須放在 TensorRT 提供的輸出張量中。

@trtp.impl("torchtrt_ex::elementwise_mul")
def _(
    x: trtp.Tensor,
    y: trtp.Tensor,
    b: float,
    a: int,
    outputs: Tuple[trtp.Tensor],
    stream: int,
):
    # Define block size
    BLOCK_SIZE = 1024

    # Grid of programs
    grid = lambda meta: (x.numel() // meta["BLOCK_SIZE"],)

    x_t = torch.as_tensor(x, device="cuda")
    y_t = torch.as_tensor(y, device="cuda")
    z_t = torch.as_tensor(outputs[0], device="cuda")
    # Launch the kernel
    elementwise_mul_kernel[grid](x_t, y_t, z_t, BLOCK_SIZE=BLOCK_SIZE)

生成轉換器

鑑於我們已經在 PyTorch 和 TensorRT 中定義了自定義運算元，現在我們可以為該操作生成轉換器了。只要名稱空間和名稱匹配，以下函式將自動生成該操作的轉換器。

torch_tensorrt.dynamo.conversion.plugins.generate_plugin_converter(
    "torchtrt_ex::elementwise_mul", supports_dynamic_shapes=True
)

將我們的轉換器與模型一起使用

現在我們可以在模型中使用我們的自定義運算元，並使用 Torch-TensorRT 對其進行編譯。我們可以看到，自定義運算元被用作模型前向傳播中的一個操作。此時編譯模型的過程與標準 Torch-TensorRT 用法相同。

class MyModel(torch.nn.Module):  # type: ignore[misc]
    def __init__(self):
        super().__init__()

    def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
        z = torch.add(x, y)
        res = torch.ops.torchtrt_ex.elementwise_mul.default(x, z, a=1)

        return res


my_model = MyModel().to("cuda")
m = torch.full((64, 64), 2, device="cuda", dtype=torch.float)
n = torch.full((64, 64), 3, device="cuda", dtype=torch.float)

with torch_tensorrt.logging.errors():
    model_trt = torch_tensorrt.compile(
        my_model, inputs=[m, n], debug=True, min_block_size=1
    )
    for i in range(300):
        res = model_trt(m, n)
        assert torch.allclose(res, my_model(m, n))

print("Ran with custom plugin!")