目录
目录

自定义 C++ 和 CUDA 扩展

创建时间： 2018 年 4 月 26 日 |上次更新时间： 2024 年 7 月 22 日 |上次验证： Nov 05， 2024

作者： Peter Goldsborough

警告

本教程从 PyTorch 2.4 开始已弃用。有关使用自定义 C++/CUDA 扩展扩展 PyTorch 的最新指南，请参阅 PyTorch 自定义运算符。

PyTorch 提供了大量与神经网络相关的操作，任意 张量代数、数据整理和其他目的。但是，您仍然可能会发现 自己需要更加定制的操作。例如，您可能希望 使用您在论文中找到的新颖的激活函数，或实现一个操作 你开发是你研究的一部分。

在 PyTorch 中集成此类自定义操作的最简单方法是编写它 在 Python 中，通过 extend 和 如此处所述。这将为您提供完整的 自动微分的强大功能（让您免于编写导数 函数）以及 Python 通常的表现力。但是，可能会有 当您的操作最好用 C++ 实现时。例如，您的代码 可能需要非常快，因为它在模型中被调用非常频繁 或者即使对于很少的通话也非常昂贵。另一个合理的原因是它 依赖于其他 C 或 C++ 库或与之交互。为了解决此类情况， PyTorch 提供了一种非常简单的方法来编写自定义 C++ 扩展。FunctionModule

C++ 扩展是我们开发的一种机制，允许用户（您）创建 PyTorch 运算符在源外定义，即与 PyTorch 分开 backend 的此方法与本机 PyTorch 操作的方式不同 实现。C++ 扩展旨在为您省去许多样板 与将操作与 PyTorch 的后端集成相关联，同时提供 为基于 PyTorch 的项目提供高度的灵活性。 不过，一旦您将操作定义为 C++ 扩展，将 it 转换为原生 PyTorch 函数在很大程度上是一个代码组织问题， 如果您决定贡献您的运营，您可以在事后解决 上游。

动机和例子

本笔记的其余部分将介绍一个编写和使用 一个 C++（和 CUDA）扩展。如果你被追赶或有人会解雇你，如果 你没有在一天结束时完成那个 OP，你可以跳过这个部分，然后 直接进入下一节中的实施详细信息。

假设您想出了一种新的循环单元，并且您发现 与最先进的技术相比，具有卓越的性能。这个循环单元 与 LSTM 类似，但不同之处在于它没有遗忘门，并使用指数线性单元 （ELU） 作为其内部激活函数。因为 这个单元永远不会忘记，我们称它为 LLTM，或长期记忆单元。

LLTM 与普通 LSTM 的两种不同之处已经足够重要了 我们无法根据自己的目的配置 PyTorch，因此我们必须 创建自定义单元格。第一种方法也是最简单的方法 - 可能在 在所有情况下，一个好的第一步 – 就是在 带有 Python 的普通 PyTorch。为此，我们需要子类化并实现 LLTM 的 forward pass。这将 如下所示：LSTMCell

class LLTM(torch.nn.Module):
    def __init__(self, input_features, state_size):
        super(LLTM, self).__init__()
        self.input_features = input_features
        self.state_size = state_size
        # 3 * state_size for input gate, output gate and candidate cell gate.
        # input_features + state_size because we will multiply with [input, h].
        self.weights = torch.nn.Parameter(
            torch.empty(3 * state_size, input_features + state_size))
        self.bias = torch.nn.Parameter(torch.empty(3 * state_size))
        self.reset_parameters()

    def reset_parameters(self):
        stdv = 1.0 / math.sqrt(self.state_size)
        for weight in self.parameters():
            weight.data.uniform_(-stdv, +stdv)

    def forward(self, input, state):
        old_h, old_cell = state
        X = torch.cat([old_h, input], dim=1)

        # Compute the input, output and candidate cell gates with one MM.
        gate_weights = F.linear(X, self.weights, self.bias)
        # Split the combined gate weight matrix into its components.
        gates = gate_weights.chunk(3, dim=1)

        input_gate = torch.sigmoid(gates[0])
        output_gate = torch.sigmoid(gates[1])
        # Here we use an ELU instead of the usual tanh.
        candidate_cell = F.elu(gates[2])

        # Compute the new cell state.
        new_cell = old_cell + candidate_cell * input_gate
        # Compute the new hidden state and output.
        new_h = torch.tanh(new_cell) * output_gate

        return new_h, new_cell

然后我们可以按预期使用它：

import torch

X = torch.randn(batch_size, input_features)
h = torch.randn(batch_size, state_size)
C = torch.randn(batch_size, state_size)

rnn = LLTM(input_features, state_size)

new_h, new_C = rnn(X, (h, C))

当然，如果可能且合理，您应该使用这种方法来 扩展 PyTorch。由于 PyTorch 对其 CPU 和 GPU 的操作，由 NVIDIA cuDNN、Intel MKL 或 NNPACK 等库提供支持，如上所述的 PyTorch 代码通常是 足够快。然而，我们也可以看到为什么在某些情况下，会有 进一步提高性能的空间。最明显的原因是 PyTorch 不知道您正在实施的算法。它只知道 用于组合算法的各个操作中。因此，PyTorch 必须一个接一个地单独执行您的操作。由于每个 对操作的实现（或内核）的单独调用，这可能会 涉及启动一个 CUDA 内核，有一定的开销，这个 在许多函数调用中，开销可能会变得很大。此外， 运行我们代码的 Python 解释器本身会减慢我们的程序速度。

因此，加快速度的明确方法是用 C++ 重写部分（或 CUDA） 并融合特定的操作组。熔合是指将 将许多函数实现为单个函数，该函数受益于 更少的内核启动以及我们可以执行的其他优化 提高了全球数据流的可见性。

让我们看看如何使用 C++ 扩展来实现 LLTM.我们将首先用纯 C++ 编写它，使用 ATen 库，该库为 PyTorch 的大部分 backend 的 Python 代码，看看它让我们翻译 Python 代码是多么容易。然后 通过将模型的某些部分移动到 CUDA 内核以获得更多好处，从而进一步加快速度 从 GPU 提供的大规模并行性。

编写 C++ 扩展

C++ 扩展有两种形式：可以使用 “提前”构建 ，也可以通过 .我们将从第一种方法开始，然后 稍后讨论后者。setuptools

构建setuptools

对于 “提前” 风格，我们通过编写一个脚本来构建我们的 C++ 扩展，该脚本使用 setuptools 来编译我们的 C++ 代码。对于 LLTM，它 看起来就这么简单：setup.py

from setuptools import setup, Extension
from torch.utils import cpp_extension

setup(name='lltm_cpp',
      ext_modules=[cpp_extension.CppExtension('lltm_cpp', ['lltm.cpp'])],
      cmdclass={'build_ext': cpp_extension.BuildExtension})

在此代码中，是一个方便的包装器，它传递正确的 include 路径和集 C++ 扩展的语言。等效的原版代码就是：CppExtensionsetuptools.Extensionsetuptools

Extension(
   name='lltm_cpp',
   sources=['lltm.cpp'],
   include_dirs=cpp_extension.include_paths(),
   language='c++')

BuildExtension执行许多必需的配置步骤，并 检查并管理混合 C++/CUDA 情况下的混合编译 扩展。这就是我们真正需要了解的有关构建 C++ 扩展的全部内容 暂时！现在让我们看一下 C++ 扩展的实现 它进入 。lltm.cpp

编写 C++ 运算

让我们开始用 C++ 实现 LLTM！我们需要的一个函数 backward pass 是 sigmoid 的导数。这是一个足够小的部分 讨论编写 C++ 时可用的整体环境的代码 扩展：

#include <torch/extension.h>

#include <iostream>

torch::Tensor d_sigmoid(torch::Tensor z) {
  auto s = torch::sigmoid(z);
  return (1 - s) * s;
}

<torch/extension.h>是包含所有必需 PyTorch 的一站式标头 位来写入 C++ 扩展。它包括：

  • ATen 库是我们用于张量计算的主要 API，

  • pybind11，这就是我们为 C++ 代码创建 Python 绑定的方式，

  • 管理 ATen 和 pybind11 之间交互细节的 Headers。

的实现展示了如何使用 ATen API。 PyTorch 的 tensor 和 variable 接口是从 ATen 库，因此我们可以或多或少地将我们的 Python 实现 1：1 翻译 转换为 C++。我们所有计算的主要数据类型将是 。可以在此处检查其完整的 API。通知 此外，我们可以包含或任何其他 C 或 C++ 头文件 —— 我们有 C++11 的全部功能可供我们使用。d_sigmoid()torch::Tensor<iostream>

请注意，CUDA-11.5 nvcc 在 Windows 上解析 torch/extension.h 时会遇到内部编译器错误。 要解决此问题，请将 python 绑定逻辑移动到纯 C++ 文件。 示例使用：

#include <ATen/ATen.h>
at::Tensor SigmoidAlphaBlendForwardCuda(....)

而不是：

#include <torch/extension.h>
torch::Tensor SigmoidAlphaBlendForwardCuda(...)

当前未解决的 nvcc 错误问题在这里。 此处提供了完整的解决方法代码示例。

前进传球

接下来，我们可以将整个前向传递移植到 C++：

#include <vector>

std::vector<at::Tensor> lltm_forward(
    torch::Tensor input,
    torch::Tensor weights,
    torch::Tensor bias,
    torch::Tensor old_h,
    torch::Tensor old_cell) {
  auto X = torch::cat({old_h, input}, /*dim=*/1);

  auto gate_weights = torch::addmm(bias, X, weights.transpose(0, 1));
  auto gates = gate_weights.chunk(3, /*dim=*/1);

  auto input_gate = torch::sigmoid(gates[0]);
  auto output_gate = torch::sigmoid(gates[1]);
  auto candidate_cell = torch::elu(gates[2], /*alpha=*/1.0);

  auto new_cell = old_cell + candidate_cell * input_gate;
  auto new_h = torch::tanh(new_cell) * output_gate;

  return {new_h,
          new_cell,
          input_gate,
          output_gate,
          candidate_cell,
          X,
          gate_weights};
}

Backward Pass

C++ 扩展 API 目前不提供自动 为我们生成一个 backwards 函数。因此，我们还必须实现 我们的 LLTM 的反向传递，它计算损失的导数 尊重 forward pass 的每个 input。最终，我们将把 forward 和 backward 函数转换为 a 来创建 一个很好的 Python 绑定。backward 函数稍微复杂一些，因此 我们不会更深入地研究代码（如果您有兴趣，Alex Graves 的论文是一篇很好的阅读 有关此内容的信息）：

// tanh'(z) = 1 - tanh^2(z)
torch::Tensor d_tanh(torch::Tensor z) {
  return 1 - z.tanh().pow(2);
}

// elu'(z) = relu'(z) + { alpha * exp(z) if (alpha * (exp(z) - 1)) < 0, else 0}
torch::Tensor d_elu(torch::Tensor z, torch::Scalar alpha = 1.0) {
  auto e = z.exp();
  auto mask = (alpha * (e - 1)) < 0;
  return (z > 0).type_as(z) + mask.type_as(z) * (alpha * e);
}

std::vector<torch::Tensor> lltm_backward(
    torch::Tensor grad_h,
    torch::Tensor grad_cell,
    torch::Tensor new_cell,
    torch::Tensor input_gate,
    torch::Tensor output_gate,
    torch::Tensor candidate_cell,
    torch::Tensor X,
    torch::Tensor gate_weights,
    torch::Tensor weights) {
  auto d_output_gate = torch::tanh(new_cell) * grad_h;
  auto d_tanh_new_cell = output_gate * grad_h;
  auto d_new_cell = d_tanh(new_cell) * d_tanh_new_cell + grad_cell;

  auto d_old_cell = d_new_cell;
  auto d_candidate_cell = input_gate * d_new_cell;
  auto d_input_gate = candidate_cell * d_new_cell;

  auto gates = gate_weights.chunk(3, /*dim=*/1);
  d_input_gate *= d_sigmoid(gates[0]);
  d_output_gate *= d_sigmoid(gates[1]);
  d_candidate_cell *= d_elu(gates[2]);

  auto d_gates =
      torch::cat({d_input_gate, d_output_gate, d_candidate_cell}, /*dim=*/1);

  auto d_weights = d_gates.t().mm(X);
  auto d_bias = d_gates.sum(/*dim=*/0, /*keepdim=*/true);

  auto d_X = d_gates.mm(weights);
  const auto state_size = grad_h.size(1);
  auto d_old_h = d_X.slice(/*dim=*/1, 0, state_size);
  auto d_input = d_X.slice(/*dim=*/1, state_size);

  return {d_old_h, d_input, d_weights, d_bias, d_old_cell};
}

绑定到 Python

一旦你用 C++ 和 ATen 编写了你的操作，你就可以使用 pybind11 来 以非常简单的方式将 C++ 函数或类绑定到 Python 中。 您对 PyTorch C++ 扩展的这一部分有疑问或问题将 主要由 Pybind11 文档解决。

对于我们的扩展，必要的绑定代码仅跨越四行：

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
  m.def("forward", &lltm_forward, "LLTM forward");
  m.def("backward", &lltm_backward, "LLTM backward");
}

这里需要注意的一点是 宏 .割torch扩展 build 会将其定义为你在脚本中为扩展指定的名称。在这种情况下，的值将为 “lltm_cpp”。 这是为了避免在两个地方维护扩展的名称 （构建脚本和您的 C++ 代码），因为两者之间的不匹配可能会导致 令人讨厌且难以跟踪的问题。TORCH_EXTENSION_NAMEsetup.pyTORCH_EXTENSION_NAME

使用您的扩展

现在，我们可以在 PyTorch 中导入我们的扩展。此时，您的目录 结构可能如下所示：

pytorch/
  lltm-extension/
    lltm.cpp
    setup.py

现在，运行以构建和安装您的扩展。这 应如下所示：python setup.py install

running install
running bdist_egg
running egg_info
creating lltm_cpp.egg-info
writing lltm_cpp.egg-info/PKG-INFO
writing dependency_links to lltm_cpp.egg-info/dependency_links.txt
writing top-level names to lltm_cpp.egg-info/top_level.txt
writing manifest file 'lltm_cpp.egg-info/SOURCES.txt'
reading manifest file 'lltm_cpp.egg-info/SOURCES.txt'
writing manifest file 'lltm_cpp.egg-info/SOURCES.txt'
installing library code to build/bdist.linux-x86_64/egg
running install_lib
running build_ext
building 'lltm_cpp' extension
creating build
creating build/temp.linux-x86_64-3.7
gcc -pthread -B ~/local/miniconda/compiler_compat -Wl,--sysroot=/ -Wsign-compare -DNDEBUG -g -fwrapv -O3 -Wall -Wstrict-prototypes -fPIC -I~/local/miniconda/lib/python3.7/site-packages/torch/include -I~/local/miniconda/lib/python3.7/site-packages/torch/include/torch/csrc/api/include -I~/local/miniconda/lib/python3.7/site-packages/torch/include/TH -I~/local/miniconda/lib/python3.7/site-packages/torch/include/THC -I~/local/miniconda/include/python3.7m -c lltm.cpp -o build/temp.linux-x86_64-3.7/lltm.o -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=lltm_cpp -D_GLIBCXX_USE_CXX11_ABI=1 -std=c++11
cc1plus: warning: command line option ‘-Wstrict-prototypes’ is valid for C/ObjC but not for C++
creating build/lib.linux-x86_64-3.7
g++ -pthread -shared -B ~/local/miniconda/compiler_compat -L~/local/miniconda/lib -Wl,-rpath=~/local/miniconda/lib -Wl,--no-as-needed -Wl,--sysroot=/ build/temp.linux-x86_64-3.7/lltm.o -o build/lib.linux-x86_64-3.7/lltm_cpp.cpython-37m-x86_64-linux-gnu.so
creating build/bdist.linux-x86_64
creating build/bdist.linux-x86_64/egg
copying build/lib.linux-x86_64-3.7/lltm_cpp.cpython-37m-x86_64-linux-gnu.so -> build/bdist.linux-x86_64/egg
creating stub loader for lltm_cpp.cpython-37m-x86_64-linux-gnu.so
byte-compiling build/bdist.linux-x86_64/egg/lltm_cpp.py to lltm_cpp.cpython-37.pyc
creating build/bdist.linux-x86_64/egg/EGG-INFO
copying lltm_cpp.egg-info/PKG-INFO -> build/bdist.linux-x86_64/egg/EGG-INFO
copying lltm_cpp.egg-info/SOURCES.txt -> build/bdist.linux-x86_64/egg/EGG-INFO
copying lltm_cpp.egg-info/dependency_links.txt -> build/bdist.linux-x86_64/egg/EGG-INFO
copying lltm_cpp.egg-info/top_level.txt -> build/bdist.linux-x86_64/egg/EGG-INFO
writing build/bdist.linux-x86_64/egg/EGG-INFO/native_libs.txt
zip_safe flag not set; analyzing archive contents...
__pycache__.lltm_cpp.cpython-37: module references __file__
creating 'dist/lltm_cpp-0.0.0-py3.7-linux-x86_64.egg' and adding 'build/bdist.linux-x86_64/egg' to it
removing 'build/bdist.linux-x86_64/egg' (and everything under it)
Processing lltm_cpp-0.0.0-py3.7-linux-x86_64.egg
removing '~/local/miniconda/lib/python3.7/site-packages/lltm_cpp-0.0.0-py3.7-linux-x86_64.egg' (and everything under it)
creating ~/local/miniconda/lib/python3.7/site-packages/lltm_cpp-0.0.0-py3.7-linux-x86_64.egg
Extracting lltm_cpp-0.0.0-py3.7-linux-x86_64.egg to ~/local/miniconda/lib/python3.7/site-packages
lltm-cpp 0.0.0 is already the active version in easy-install.pth

Installed ~/local/miniconda/lib/python3.7/site-packages/lltm_cpp-0.0.0-py3.7-linux-x86_64.egg
Processing dependencies for lltm-cpp==0.0.0
Finished processing dependencies for lltm-cpp==0.0.0

关于编译器的小说明：由于 ABI 版本控制问题，您用于 build 您的 C++ 扩展必须与编译器 PyTorch 的 ABI 兼容 构建与。实际上，这意味着您必须在 Linux 上使用 GCC 版本 4.9 及更高版本。 对于 Ubuntu 16.04 和其他更新的 Linux 发行版，这应该是 default 编译器。在 MacOS 上，您必须使用 clang（它没有任何 ABI 版本控制问题）。在最坏的情况下 的情况下，您可以使用编译器从源代码构建 PyTorch，然后构建 扩展。

构建扩展后，您只需在 Python 中使用 您在脚本中指定的名称。请务必先这样做，因为这将解析动态链接器必须解析的一些符号 看：setup.pyimport torch

In [1]: import torch
In [2]: import lltm_cpp
In [3]: lltm_cpp.forward
Out[3]: <function lltm.PyCapsule.forward>

如果我们调用函数或模块，我们可以看到它的签名 匹配我们的 C++ 代码：help()

In[4] help(lltm_cpp.forward)
forward(...) method of builtins.PyCapsule instance
    forward(arg0: torch::Tensor, arg1: torch::Tensor, arg2: torch::Tensor, arg3: torch::Tensor, arg4: torch::Tensor) -> List[torch::Tensor]

    LLTM forward

由于我们现在能够从 Python 调用 C++ 函数，因此我们可以包装它们 with 和 to make them first PyTorch 的类公民：

import math
import torch

# Our module!
import lltm_cpp

class LLTMFunction(torch.autograd.Function):
    @staticmethod
    def forward(ctx, input, weights, bias, old_h, old_cell):
        outputs = lltm_cpp.forward(input, weights, bias, old_h, old_cell)
        new_h, new_cell = outputs[:2]
        variables = outputs[1:] + [weights]
        ctx.save_for_backward(*variables)

        return new_h, new_cell

    @staticmethod
    def backward(ctx, grad_h, grad_cell):
        outputs = lltm_cpp.backward(
            grad_h.contiguous(), grad_cell.contiguous(), *ctx.saved_tensors)
        d_old_h, d_input, d_weights, d_bias, d_old_cell = outputs
        return d_input, d_weights, d_bias, d_old_h, d_old_cell


class LLTM(torch.nn.Module):
    def __init__(self, input_features, state_size):
        super(LLTM, self).__init__()
        self.input_features = input_features
        self.state_size = state_size
        self.weights = torch.nn.Parameter(
            torch.empty(3 * state_size, input_features + state_size))
        self.bias = torch.nn.Parameter(torch.empty(3 * state_size))
        self.reset_parameters()

    def reset_parameters(self):
        stdv = 1.0 / math.sqrt(self.state_size)
        for weight in self.parameters():
            weight.data.uniform_(-stdv, +stdv)

    def forward(self, input, state):
        return LLTMFunction.apply(input, self.weights, self.bias, *state)

性能比较

现在，我们能够从 PyTorch 使用和调用我们的 C++ 代码，我们可以运行 小基准测试，看看我们在 C++。我们将向前和向后运行 LLTM 几次，并测量 期间：

import time

import torch

batch_size = 16
input_features = 32
state_size = 128

X = torch.randn(batch_size, input_features)
h = torch.randn(batch_size, state_size)
C = torch.randn(batch_size, state_size)

rnn = LLTM(input_features, state_size)

forward = 0
backward = 0
for _ in range(100000):
    start = time.time()
    new_h, new_C = rnn(X, (h, C))
    forward += time.time() - start

    start = time.time()
    (new_h.sum() + new_C.sum()).backward()
    backward += time.time() - start

print('Forward: {:.3f} s | Backward {:.3f} s'.format(forward, backward))

如果我们使用原始 LLTM 运行此代码，则我们一开始就用纯 Python 编写 在这篇文章中，我们得到以下数字（在我的机器上）：

Forward: 506.480 us | Backward 444.694 us

以及我们的新 C++ 版本：

Forward: 349.335 us | Backward 443.523 us

我们已经可以看到 forward 函数的显著加速（超过 30%).对于向后功能，可以看到加速，尽管不是主要加速。 我上面写的向后传递没有特别优化，可以 肯定会得到改进。此外，PyTorch 的自动微分引擎可以 自动并行计算图，可以使用更高效的 操作，并且也是用 C++ 实现的，因此它应该是 快。尽管如此，这是一个好的开始。

GPU 设备上的性能

关于 PyTorch 的 ATen 后端的一个绝妙事实是，它将 运行设备的 Computing 设备。这意味着我们为 CPU 编写的代码相同 也可以在 GPU 上运行，并且各个操作将相应地调度 到 GPU 优化的实现。对于某些运算，如矩阵乘法 （喜欢 or ），这是一个巨大的胜利。我们来看看有多少 使用 CUDA 张量运行 C++ 代码的性能。没有变化 我们的实现是必需的，我们只需要将我们的张量放入 GPU 中 memory 的 Python 中，在 创建时间或使用 after creation：mmaddmmdevice=cuda_device.to(cuda_device)

import torch

assert torch.cuda.is_available()
cuda_device = torch.device("cuda")  # device object representing GPU

batch_size = 16
input_features = 32
state_size = 128

# Note the device=cuda_device arguments here
X = torch.randn(batch_size, input_features, device=cuda_device)
h = torch.randn(batch_size, state_size, device=cuda_device)
C = torch.randn(batch_size, state_size, device=cuda_device)

rnn = LLTM(input_features, state_size).to(cuda_device)

forward = 0
backward = 0
for _ in range(100000):
    start = time.time()
    new_h, new_C = rnn(X, (h, C))
    torch.cuda.synchronize()
    forward += time.time() - start

    start = time.time()
    (new_h.sum() + new_C.sum()).backward()
    torch.cuda.synchronize()
    backward += time.time() - start

print('Forward: {:.3f} us | Backward {:.3f} us'.format(forward * 1e6/1e5, backward * 1e6/1e5))

再次将我们的普通 PyTorch 代码与我们的 C++ 版本进行比较，现在两者都是 在 CUDA 设备上运行，我们再次看到性能提升。对于 Python/PyTorch：

Forward: 187.719 us | Backward 410.815 us

和 C++/ATen：

Forward: 149.802 us | Backward 393.458 us

与非 CUDA 代码相比，这是一个很大的整体加速。但是，我们可以拉取 通过编写自定义 CUDA 内核，我们的 C++ 代码性能更高，这 我们很快就会深入研究。在此之前，让我们讨论另一种构建 C++ 的方法 扩展。

JIT 编译扩展

之前，我提到过构建 C++ 扩展有两种方法：使用或即时 （JIT）。介绍完前者后，让我们 详细说明后者。JIT 编译机制为您提供了一种方法 通过调用简单的 在 PyTorch 的 API 中名为 .为 LLTM，这看起来就像这样简单：setuptools

from torch.utils.cpp_extension import load

lltm_cpp = load(name="lltm_cpp", sources=["lltm.cpp"])

在这里，我们为函数提供与 相同的信息。在后台，这将执行以下操作：setuptools

  1. 创建临时目录 ，/tmp/torch_extensions/lltm

  2. 将 Ninja 构建文件发送到该临时目录

  3. 将源文件编译到共享库中，

  4. 将此共享库作为 Python 模块导入。

事实上，如果传递给 ，则会 了解该过程：verbose=Truecpp_extension.load()

Using /tmp/torch_extensions as PyTorch extensions root...
Emitting ninja build file /tmp/torch_extensions/lltm_cpp/build.ninja...
Building extension module lltm_cpp...
Loading extension module lltm_cpp...

生成的 Python 模块将与 setuptools 生成的完全相同。 但消除了必须维护单独构建的要求 文件。如果您的设置更复杂，并且您确实需要 的全部功能，您可以编写自己的 - 但在许多情况下 这种 JIT 技术就可以了。第一次跑过这条线时， 这将需要一些时间，因为扩展正在后台编译。因为 我们使用 Ninja 构建系统来构建您的源代码，重新编译是 incremental 的 incremental 函数，从而在运行 Python 模块 a second time 速度快，如果您没有更改扩展的 源文件。setup.pysetuptoolssetup.py

编写混合 C++/CUDA 扩展

为了真正将我们的实现提升到一个新的水平，我们可以手写 我们使用自定义 CUDA 内核的正向和反向传递。对于 LLTM，这具有 特别有效的前景，因为有大量的 按顺序进行逐点操作，这些操作都可以在 单个 CUDA 内核。让我们看看如何编写这样的 CUDA 内核和 使用此扩展机制将其与 PyTorch 集成。

编写 CUDA 扩展的一般策略是首先编写一个 C++ 文件 它定义了将从 Python 调用的函数，并将这些函数绑定 函数添加到 Python。此外，此文件还将声明 CUDA （） 文件中定义的函数。然后，C++ 函数将 进行一些检查，并最终将其调用转发到 CUDA 函数。在 CUDA 文件，我们编写实际的 CUDA 内核。包装 然后，将负责使用 C++ 编译器（如 C++ 编译器）编译 C++ 源，并使用 NVIDIA 的编译器编译 CUDA 源。这可确保 每个编译器都会处理它最擅长编译的文件。最终，他们 将链接到一个共享库，该库可从 Python 获得 法典。.cucpp_extensiongccnvcc

我们将从 C++ 文件开始，我们将其称为 ，例如：lltm_cuda.cpp

#include <torch/extension.h>

#include <vector>

// CUDA forward declarations

std::vector<torch::Tensor> lltm_cuda_forward(
    torch::Tensor input,
    torch::Tensor weights,
    torch::Tensor bias,
    torch::Tensor old_h,
    torch::Tensor old_cell);

std::vector<torch::Tensor> lltm_cuda_backward(
    torch::Tensor grad_h,
    torch::Tensor grad_cell,
    torch::Tensor new_cell,
    torch::Tensor input_gate,
    torch::Tensor output_gate,
    torch::Tensor candidate_cell,
    torch::Tensor X,
    torch::Tensor gate_weights,
    torch::Tensor weights);

// C++ interface

#define CHECK_CUDA(x) TORCH_CHECK(x.device().is_cuda(), #x " must be a CUDA tensor")
#define CHECK_CONTIGUOUS(x) TORCH_CHECK(x.is_contiguous(), #x " must be contiguous")
#define CHECK_INPUT(x) CHECK_CUDA(x); CHECK_CONTIGUOUS(x)

std::vector<torch::Tensor> lltm_forward(
    torch::Tensor input,
    torch::Tensor weights,
    torch::Tensor bias,
    torch::Tensor old_h,
    torch::Tensor old_cell) {
  CHECK_INPUT(input);
  CHECK_INPUT(weights);
  CHECK_INPUT(bias);
  CHECK_INPUT(old_h);
  CHECK_INPUT(old_cell);

  return lltm_cuda_forward(input, weights, bias, old_h, old_cell);
}

std::vector<torch::Tensor> lltm_backward(
    torch::Tensor grad_h,
    torch::Tensor grad_cell,
    torch::Tensor new_cell,
    torch::Tensor input_gate,
    torch::Tensor output_gate,
    torch::Tensor candidate_cell,
    torch::Tensor X,
    torch::Tensor gate_weights,
    torch::Tensor weights) {
  CHECK_INPUT(grad_h);
  CHECK_INPUT(grad_cell);
  CHECK_INPUT(input_gate);
  CHECK_INPUT(output_gate);
  CHECK_INPUT(candidate_cell);
  CHECK_INPUT(X);
  CHECK_INPUT(gate_weights);
  CHECK_INPUT(weights);

  return lltm_cuda_backward(
      grad_h,
      grad_cell,
      new_cell,
      input_gate,
      output_gate,
      candidate_cell,
      X,
      gate_weights,
      weights);
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
  m.def("forward", &lltm_forward, "LLTM forward (CUDA)");
  m.def("backward", &lltm_backward, "LLTM backward (CUDA)");
}

如您所见，它主要是样板、检查和转发到函数 我们将在 CUDA 文件中定义。我们将为此文件命名（注意扩展名！NVCC 可以合理地 编译 C++11，因此我们仍然有 ATen 和 C++ 标准库可用 对我们来说（但不是 ）。请注意，无法处理文件 具有相同的名称但不同的扩展名，因此如果您使用该方法而不是 JIT 方法，则必须为您的 CUDA 文件指定一个不同的名称 比您的 C++ 文件（对于 JIT 方法）和 很好）。让我们稍微看一下这个文件会是什么样子：lltm_cuda_kernel.cu.cutorch.hsetuptoolssetup.pylltm.cpplltm.cu

#include <torch/extension.h>

#include <cuda.h>
#include <cuda_runtime.h>

#include <vector>

template <typename scalar_t>
__device__ __forceinline__ scalar_t sigmoid(scalar_t z) {
  return 1.0 / (1.0 + exp(-z));
}

在这里，我们看到了我刚才描述的标头，以及我们正在使用 特定于 CUDA 的声明，如 and 和 函数（如 .让我们继续介绍另外几个帮助程序函数，这些函数 我们需要：__device____forceinline__exp

template <typename scalar_t>
__device__ __forceinline__ scalar_t d_sigmoid(scalar_t z) {
  const auto s = sigmoid(z);
  return (1.0 - s) * s;
}

template <typename scalar_t>
__device__ __forceinline__ scalar_t d_tanh(scalar_t z) {
  const auto t = tanh(z);
  return 1 - (t * t);
}

template <typename scalar_t>
__device__ __forceinline__ scalar_t elu(scalar_t z, scalar_t alpha = 1.0) {
  return fmax(0.0, z) + fmin(0.0, alpha * (exp(z) - 1.0));
}

template <typename scalar_t>
__device__ __forceinline__ scalar_t d_elu(scalar_t z, scalar_t alpha = 1.0) {
  const auto e = exp(z);
  const auto d_relu = z < 0.0 ? 0.0 : 1.0;
  return d_relu + (((alpha * (e - 1.0)) < 0.0) ? (alpha * e) : 0.0);
}

要实际实现一个函数，我们再次需要两样东西：一个函数 执行我们不希望显式手动写入的操作，并调用 转换为 CUDA 内核，然后是我们想要的部分的实际 CUDA 内核 加速。对于前向传递，第一个函数应如下所示：

std::vector<torch::Tensor> lltm_cuda_forward(
    torch::Tensor input,
    torch::Tensor weights,
    torch::Tensor bias,
    torch::Tensor old_h,
    torch::Tensor old_cell) {
  auto X = torch::cat({old_h, input}, /*dim=*/1);
  auto gates = torch::addmm(bias, X, weights.transpose(0, 1));

  const auto batch_size = old_cell.size(0);
  const auto state_size = old_cell.size(1);

  auto new_h = torch::zeros_like(old_cell);
  auto new_cell = torch::zeros_like(old_cell);
  auto input_gate = torch::zeros_like(old_cell);
  auto output_gate = torch::zeros_like(old_cell);
  auto candidate_cell = torch::zeros_like(old_cell);

  const int threads = 1024;
  const dim3 blocks((state_size + threads - 1) / threads, batch_size);

  AT_DISPATCH_FLOATING_TYPES(gates.type(), "lltm_forward_cuda", ([&] {
    lltm_cuda_forward_kernel<scalar_t><<<blocks, threads>>>(
        gates.data<scalar_t>(),
        old_cell.data<scalar_t>(),
        new_h.data<scalar_t>(),
        new_cell.data<scalar_t>(),
        input_gate.data<scalar_t>(),
        output_gate.data<scalar_t>(),
        candidate_cell.data<scalar_t>(),
        state_size);
  }));

  return {new_h, new_cell, input_gate, output_gate, candidate_cell, X, gates};
}

这里的主要兴趣点是宏和 内核启动（由 表示）。虽然 ATen 抽象出 我们处理的张量的设备和数据类型，张量将在运行时， 仍然由 Concrete 设备上的 Concrete 类型的内存提供支持。因此，我们 需要一种方法在运行时确定 Tensor 是什么类型，然后有选择地确定 调用具有相应正确类型签名的函数。手动完成， 这（概念上）看起来像这样：AT_DISPATCH_FLOATING_TYPES<<<...>>>

switch (tensor.type().scalarType()) {
  case torch::ScalarType::Double:
    return function<double>(tensor.data<double>());
  case torch::ScalarType::Float:
    return function<float>(tensor.data<float>());
  ...
}

目的是照顾这个调度 对我们来说。它需要一个类型（在我们的例子中为），一个名称（用于错误 messages） 和 Lambda 函数。在此 lambda 函数中，类型别名可用，并被定义为张量实际 在运行时。因此，如果我们有一个模板函数（其中 我们的 CUDA 内核将是），我们可以使用这个别名来实例化它， ，并且将调用正确的函数。在这种情况下，我们还希望检索 张量的数据指针作为该类型的指针。如果你 想要分派所有类型，而不仅仅是浮点类型（ 和 ），你可以使用 。AT_DISPATCH_FLOATING_TYPESgates.type()scalar_tscalar_tscalar_tFloatDoubleAT_DISPATCH_ALL_TYPES

请注意，我们使用普通的 ATen 执行一些操作。这些操作将 仍然在 GPU 上运行，但使用 ATen 的默认实现。这使得 sense 的原因，因为 ATen 将对 matrix 等内容使用高度优化的例程 乘法（例如 ）或卷积，这将更难 实施和改进自己。addmm

至于内核启动本身，我们在这里指定每个 CUDA 块 将有 1024 个线程，并且整个 GPU 网格被分割成尽可能多的 线程块，以将我们的矩阵填充一个 每个组件的线程数。例如，如果我们的 state size 为 2048，并且我们的 batch 大小 4，则每 1024 个线程总共启动一个区块。如果 您以前从未听说过 CUDA “块” 或 “网格” ，介绍性读物 关于 CUDA may 帮助。1 x 10244 x 2 = 8

实际的 CUDA 内核相当简单（如果您以前曾对 GPU 进行过编程）：

template <typename scalar_t>
__global__ void lltm_cuda_forward_kernel(
    const scalar_t* __restrict__ gates,
    const scalar_t* __restrict__ old_cell,
    scalar_t* __restrict__ new_h,
    scalar_t* __restrict__ new_cell,
    scalar_t* __restrict__ input_gate,
    scalar_t* __restrict__ output_gate,
    scalar_t* __restrict__ candidate_cell,
    size_t state_size) {
  const int column = blockIdx.x * blockDim.x + threadIdx.x;
  const int index = blockIdx.y * state_size + column;
  const int gates_row = blockIdx.y * (state_size * 3);
  if (column < state_size) {
    input_gate[index] = sigmoid(gates[gates_row + column]);
    output_gate[index] = sigmoid(gates[gates_row + state_size + column]);
    candidate_cell[index] = elu(gates[gates_row + 2 * state_size + column]);
    new_cell[index] =
        old_cell[index] + candidate_cell[index] * input_gate[index];
    new_h[index] = tanh(new_cell[index]) * output_gate[index];
  }
}

这里最有趣的是，我们能够计算所有这些 Pointwise 操作完全并行地执行 门矩阵。如果你想象必须用一个巨大的循环来做这件事 100 万个元素，你可以明白为什么这会快得多。for

使用访问器

您可以在 CUDA 内核中看到，我们直接在具有右侧 类型。事实上，直接在 cuda 中使用高级类型不可知的张量 kernels 将非常低效。

然而，这是以易用性和可读性为代价的，特别是对于 高维数据。在我们的示例中，我们知道连续张量有 3 个维度：gates

  1. 批次、大小和步幅batch_size3*state_size

  2. 行、大小和步幅3state_size

  3. 指数、大小和步幅state_size1

那么我们如何访问内核内部的元素呢？ 事实证明，您需要大步通过一些简单的 算术。gates[n][row][column]

gates.data<scalar_t>()[n*3*state_size + row*state_size + column]

除了冗长之外，此表达式还需要显式 stride known，从而在其参数中传递给内核函数。你可以看到 在核函数接受具有不同 sizes 的参数，则最终会得到一个很长的参数列表。

幸运的是，ATen 提供了使用单个 dynamic 检查 Tensor 是否为维度的类型和数量。 然后，访问器公开一个 API 以高效访问 Tensor 元素 而无需转换为单个指针：

torch::Tensor foo = torch::rand({12, 12});

// assert foo is 2-dimensional and holds floats.
auto foo_a = foo.accessor<float,2>();
float trace = 0;

for(int i = 0; i < foo_a.size(0); i++) {
  // use the accessor foo_a to get tensor data.
  trace += foo_a[i][i];
}

访问器对象具有相对高级的接口，包括 和 方法和多维索引。该接口旨在有效地访问 cpu 张量上的数据。等效的 对于 CUDA 张量是 和 ，其中 生成具有 64 位或 32 位整数索引的 Packed Accessor。.size().stride().accessor<>packed_accessor64<>packed_accessor32<>

与 Accessor 的根本区别在于 Packed Accessor 复制大小 和 stride 数据，而不是指向它。它允许我们 将其传递给 CUDA 内核函数并在其中使用其接口。

我们可以设计一个采用 Packed Accessors 而不是指针的函数。

__global__ void lltm_cuda_forward_kernel(
    const torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> gates,
    const torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> old_cell,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> new_h,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> new_cell,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> input_gate,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> output_gate,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> candidate_cell)

让我们分解此处使用的模板。前两个参数 和 与常规 Accessor 相同。该参数指示关键字必须是 使用。另请注意，我们使用了 variant 来存储 sizes 和 strides 在 .这在使用 64 位 variant （） 会使内核变慢。scalar_t2torch::RestrictPtrTraits__restrict__PackedAccessor32int32_tPackedAccessor64

函数声明变为

template <typename scalar_t>
__global__ void lltm_cuda_forward_kernel(
    const torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> gates,
    const torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> old_cell,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> new_h,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> new_cell,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> input_gate,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> output_gate,
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> candidate_cell) {
  //batch index
  const int n = blockIdx.y;
  // column index
  const int c = blockIdx.x * blockDim.x + threadIdx.x;
  if (c < gates.size(2)){
    input_gate[n][c] = sigmoid(gates[n][0][c]);
    output_gate[n][c] = sigmoid(gates[n][1][c]);
    candidate_cell[n][c] = elu(gates[n][2][c]);
    new_cell[n][c] =
        old_cell[n][c] + candidate_cell[n][c] * input_gate[n][c];
    new_h[n][c] = tanh(new_cell[n][c]) * output_gate[n][c];
  }
}

实现更具可读性！然后，此函数由 使用 host 函数。.packed_accessor32<>

std::vector<torch::Tensor> lltm_cuda_forward(
    torch::Tensor input,
    torch::Tensor weights,
    torch::Tensor bias,
    torch::Tensor old_h,
    torch::Tensor old_cell) {
  auto X = torch::cat({old_h, input}, /*dim=*/1);
  auto gate_weights = torch::addmm(bias, X, weights.transpose(0, 1));

  const auto batch_size = old_cell.size(0);
  const auto state_size = old_cell.size(1);

  auto gates = gate_weights.reshape({batch_size, 3, state_size});
  auto new_h = torch::zeros_like(old_cell);
  auto new_cell = torch::zeros_like(old_cell);
  auto input_gate = torch::zeros_like(old_cell);
  auto output_gate = torch::zeros_like(old_cell);
  auto candidate_cell = torch::zeros_like(old_cell);

  const int threads = 1024;
  const dim3 blocks((state_size + threads - 1) / threads, batch_size);

  AT_DISPATCH_FLOATING_TYPES(gates.type(), "lltm_forward_cuda", ([&] {
    lltm_cuda_forward_kernel<scalar_t><<<blocks, threads>>>(
        gates.packed_accessor32<scalar_t,3,torch::RestrictPtrTraits>(),
        old_cell.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        new_h.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        new_cell.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        input_gate.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        output_gate.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        candidate_cell.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>());
  }));

  return {new_h, new_cell, input_gate, output_gate, candidate_cell, X, gates};
}

向后传递遵循大致相同的模式，我不会进一步详细说明 在它上面：

template <typename scalar_t>
__global__ void lltm_cuda_backward_kernel(
    torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> d_old_cell,
    torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> d_gates,
    const torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> grad_h,
    const torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> grad_cell,
    const torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> new_cell,
    const torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> input_gate,
    const torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> output_gate,
    const torch::PackedTensorAccessor32<scalar_t,2,torch::RestrictPtrTraits> candidate_cell,
    const torch::PackedTensorAccessor32<scalar_t,3,torch::RestrictPtrTraits> gate_weights) {
  //batch index
  const int n = blockIdx.y;
  // column index
  const int c = blockIdx.x * blockDim.x + threadIdx.x;
  if (c < d_gates.size(2)){
    const auto d_output_gate = tanh(new_cell[n][c]) * grad_h[n][c];
    const auto d_tanh_new_cell = output_gate[n][c] * grad_h[n][c];
    const auto d_new_cell =
        d_tanh(new_cell[n][c]) * d_tanh_new_cell + grad_cell[n][c];


    d_old_cell[n][c] = d_new_cell;
    const auto d_candidate_cell = input_gate[n][c] * d_new_cell;
    const auto d_input_gate = candidate_cell[n][c] * d_new_cell;

    d_gates[n][0][c] =
        d_input_gate * d_sigmoid(gate_weights[n][0][c]);
    d_gates[n][1][c] =
        d_output_gate * d_sigmoid(gate_weights[n][1][c]);
    d_gates[n][2][c] =
        d_candidate_cell * d_elu(gate_weights[n][2][c]);
  }
}

std::vector<torch::Tensor> lltm_cuda_backward(
    torch::Tensor grad_h,
    torch::Tensor grad_cell,
    torch::Tensor new_cell,
    torch::Tensor input_gate,
    torch::Tensor output_gate,
    torch::Tensor candidate_cell,
    torch::Tensor X,
    torch::Tensor gates,
    torch::Tensor weights) {
  auto d_old_cell = torch::zeros_like(new_cell);
  auto d_gates = torch::zeros_like(gates);

  const auto batch_size = new_cell.size(0);
  const auto state_size = new_cell.size(1);

  const int threads = 1024;
  const dim3 blocks((state_size + threads - 1) / threads, batch_size);

  AT_DISPATCH_FLOATING_TYPES(X.type(), "lltm_backward_cuda", ([&] {
    lltm_cuda_backward_kernel<scalar_t><<<blocks, threads>>>(
        d_old_cell.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        d_gates.packed_accessor32<scalar_t,3,torch::RestrictPtrTraits>(),
        grad_h.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        grad_cell.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        new_cell.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        input_gate.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        output_gate.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        candidate_cell.packed_accessor32<scalar_t,2,torch::RestrictPtrTraits>(),
        gates.packed_accessor32<scalar_t,3,torch::RestrictPtrTraits>());
  }));

  auto d_gate_weights = d_gates.reshape({batch_size, 3*state_size});
  auto d_weights = d_gate_weights.t().mm(X);
  auto d_bias = d_gate_weights.sum(/*dim=*/0, /*keepdim=*/true);

  auto d_X = d_gate_weights.mm(weights);
  auto d_old_h = d_X.slice(/*dim=*/1, 0, state_size);
  auto d_input = d_X.slice(/*dim=*/1, state_size);

  return {d_old_h, d_input, d_weights, d_bias, d_old_cell, d_gates};
}

将 C++/CUDA 操作与 PyTorch 集成

将支持 CUDA 的操作与 PyTorch 集成也非常简单。 如果要编写脚本，它可能如下所示：setup.py

from setuptools import setup
from torch.utils.cpp_extension import BuildExtension, CUDAExtension

setup(
    name='lltm',
    ext_modules=[
        CUDAExtension('lltm_cuda', [
            'lltm_cuda.cpp',
            'lltm_cuda_kernel.cu',
        ])
    ],
    cmdclass={
        'build_ext': BuildExtension
    })

我们现在使用 ，而不是 。我们可以 指定文件和文件 – 库采用 照顾这给您带来的所有麻烦。JIT 机制是偶数 简单：CppExtension()CUDAExtension().cu.cpp

from torch.utils.cpp_extension import load

lltm = load(name='lltm', sources=['lltm_cuda.cpp', 'lltm_cuda_kernel.cu'])

性能比较

我们希望并行化和融合代码的逐点操作 使用 CUDA 将提高我们 LLTM 的性能。让我们看看这是否成立 真。我们可以运行我之前列出的代码来运行基准测试。我们最快的 版本之前的版本是基于 CUDA 的 C++ 代码：

Forward: 149.802 us | Backward 393.458 us

现在使用我们的自定义 CUDA 内核：

Forward: 129.431 us | Backward 304.641 us

更多性能提升！

结论

现在，您应该对 PyTorch 的 C++ 扩展有了很好的了解 机制以及使用它们的动机。您可以找到代码 此处为本说明中显示的示例。如果您有任何疑问，请使用论坛。另外，请务必查看我们的常见问题解答，以防您遇到任何问题。

文档

访问 PyTorch 的全面开发人员文档

查看文档

教程

获取面向初学者和高级开发人员的深入教程

查看教程

资源

查找开发资源并解答您的问题

查看资源