深度学习Pytorch框架扩展知识详解

目录

Pytorch目录结构

Pytorch是如何实现C/C++扩展的

Python中的C/C++扩展

使用python的扩展API

使用python的自定义类型

使用pybind11

Pytorch中C/C++的扩展

Pytorch扩展之自动生成代码

Pytorch中op的调用关系

python绑定层

native op的编写

1. 在aten/src/ATen/native/native_functions.yaml声明op

2. 如果需要定义导数,还需要在tools/autograd/derivatives.yaml中声明

3. 在aten/src/ATen/native目录中实现op

如何找到python op对应的C++ op

yaml文件中的函数名,C++ op名以及python op名的关系

如何进行Pytoch性能分析

使用profiler输出每个op的计算时间

编写test case复现op

分析test case的profiler输出结果

总结


Pytorch目录结构

├── android: 用于构建android应用程序
├── aten: A Tensor Library,pytorch的C++ tensor 库,实现了tensor的运算 (不支持自动梯度),大部分op都可以在此找到,Aten的op实现分为两种;1. native op:op的现代C++实现 2. legacy op(TH, THC, THCUNN):原torch遗留下来的op的传统C语言实现,不要在里面添加新的op,这部分正在慢慢转移到native中
│  └── conda:包含了conda构建使用的脚本以及配置文件
│  └── src:op的实现
│     └──ATen:op的现代C++实现
│       └──core:Aten的核心功能,目前正在迁移到c10文件夹中
│       └──native:op的现代实现,在这里添加你想要添加的新的op
│         └──cpu:不是op的实际cpu实现,但是包含了特殊指令集的实现,比如AVX
│         └──cuda:op的cuda实现
│         └──sparse:COO的稀疏tensor的op的cpu和cuda的实现
│         └──mkl mkldnn miopen cudnn:绑定到具体后端的op的实现
│     └──TH:TorcH,CPU tensor的实现。TH,THC,THCUNN是原torch遗留下来的op的C语言实现
│       └──generic:TH,THC,THCUNN三者的generic中包含了操作的实际实现,这些文件会被编译N次
│       └──…
│     └──THC:GPU tensor的实现
│       └──generic
│       └──…
│     └──THCUNN:GPU版底层神经网络实现
│       └──generic
│     └──THH:THC在hip上的实现
│     └──THHUNN:THCUNN在hip上的实现
│  └── tools:包含了一些测试脚本
├── benchmarks:包含了一些pytorch特性的benchmark测试
├── binaries
├── c10:核心Tensor库,从ATen/core迁移过来,只包含最基本的功能,可以运行在服务端和移动端,适用于对二进制文件大小有要求的环境中
│  └── benchmark
│  └── core:CPU的核心库
│  └── cuda:cuda的核心库
│  └── hip:HIP的构建文件,该文件夹不包含实际的文件,这些文件通过tools/amd_build中的HIPIFY脚本拷贝过来
│  └── macros:包含了一些自定义的宏
│  └── test:测试程序
│  └── util:通用的工具库
├── caffe2: caffe2
│  └── core:caffe2的核心文件,包括tensor,workspace,blobs等
│  └── operators:caffe2的op
│  └── python:caffe2的python绑定
│  └── …
├── cmake:用于构建pytorch的cmake文件
├── docker:用于在docker中使用pytorch
├── docs:pytorch相关的说明文档
├── ios:pytroch的ios模块,用于构建ios应用程序
├── modules
├── scripts:包含了一些有用的工具,比如build_android.sh用于创建android上的pytorch库
├── submodules
├── test:pytorch的python前端的python单元测试
│  └── test_torch.py:pytorch功能的基本测试
│  └── test_autograd.py:非神经网络自动微分的测试
│  └── test_nn.py :神经网络op以及自动微分的测试
│  └── test_jit.py: JIT 编译器和TorchScript的测试
│  └──…
│  └──cpp:pytorch的C++前端的C++单元测试
│  └──onnx:onnx输出功能的测试,同时使用pytorch和caffe2
├── third_party:pytorch依赖的第三方库
├── tools:pytorch库的代码生成脚本
│  └── autograd: 生成自微分相关函数的工具
│  └──…
│  └──shared
├── torch:实际的pytorch库。包含了我们平时import的python模块,所有不在csrc文件夹中的是一个python模块,符合pytorch的python前端模块结构
│  └── csrc: pytorch库的C++实现,包含1. C++前端的实现 2.绑定代码(python bindings)。csrc中每个独立文件与python模块对应。python绑定文件的规范列表请查看setup.py,通常他们以python_为前缀
│     └──api:pytorch的C++前端接口
│     └──autograd:自动微分反向模式的实现
│     └──cuda
│     └──distributed:pytorch的分布式训练
│     └──generic
│     └──jit: TorchScript JIT的编译器和前端
│     └──…
│     └──utils
│     └──Module.cpp:构建Python扩展模块
│  └── nn:torch.nn模块
│  └── optim:torch.optim模块
│  └── distributed:torch. distributed模块

总体结构上,C/C++后端用aten库作为上层封装再通过torch/csrc 中的部分胶水代码接入Python(部分为pybind11)。

实际上,Pytorch中重点为下面4个目录:

  1. torch/
  2. torch/csrc/
  3. aten/src/ATen/
  4. c10/

Pytorch是如何实现C/C++扩展的

在了解pytorch是如何实现C++扩展之前,我们先了解一下python是如何实现扩展的。

Python中的C/C++扩展

使用python的扩展API

python有自带的扩展API,可以实现扩展功能,以最简单的加法操作为例展示如何使用python自带的扩展API进行C/C++的扩展

首先是我们需要在python中扩展的原C语言函数:

Add.h:

#ifndef __ADD__
#define __ADD__


#include<stdio.h>
int Add(int a,int b);

#endif

Add.c:

#include"Add.h"
int Add(int a,int b)
{
    return a+b;
}

下面写binding代码,binding代码包含了一组包装器函数,这组函数就是C和python之间的胶水。

Samples.c:

#include"Python.h"
#include"Add.h"
#include"Python_Subtract.h"



// 定义包装器函数封装C函数(包装器函数就是python解释器和C语言之间的胶水)
static PyObject* Add_Wrapper(PyObject *self,PyObject *args)
{
    // 解析python参数列表args(格式化)
    int x,y;
    if(!PyArg_ParseTuple(args , "ii" , &x , &y))
    {
        return NULL;
    }

    // 实际调用的C函数
    int result=Add(x,y);

    // 从C语言数据创建一个python对象并返回(格式化)
    return Py_BuildValue("i" , result);

}

// 定义函数表
static PyMethodDef sampleMethods[] = 
{
    {"Add", Add_Wrapper, METH_VARARGS, NULL}, // 第二个为包装器函数,第一个参数包装器函数的导出名,最后为函数说明,可以为空
    { NULL, NULL, 0, NULL}
};

// 定义模块
static struct PyModuleDef sampleModule = 
{
        PyModuleDef_HEAD_INIT,
        "sample",           // 模块名
        "A sample module",  // 说明(可以为空)
        -1,                 // 
        sampleMethods       // 函数表
};

// 模块初始化函数,注意命名一定要为PyInit_模块名
PyMODINIT_FUNC PyInit_sample(void) 
{
    return PyModule_Create(&sampleModule);
}

最后需要一个setup.py文件来构建这个扩展模块。

setup.py:

import glob
import os
from setuptools import find_packages
from setuptools import setup
from distutils.core import setup, Extension


currentPath = os.path.dirname(os.path.abspath(__file__))
pathOfExtensionModule = currentPath
fileNameOfSrc = glob.glob(os.path.join(pathOfExtensionModule, "*.c"))
pathOfSrc = [os.path.join(pathOfExtensionModule, s) for s in fileNameOfSrc]

setup(name='sample',# 这里不一定要与模块名一致
      ext_modules=[
        Extension('sample', # 模块名 
                  pathOfSrc,
                  include_dirs = pathOfExtensionModule,
                  define_macros = [('FOO','1')],
                  undef_macros = ['BAR'],
                  library_dirs = ['/usr/local/lib']
                  )
        ]
)

执行python setup.py install,安装成功后,就可以在python中调用了

 

总结一下:

  1. 定义包装器函数
  2. 定义函数表
  3. 定义模块
  4. 定义初始化函数
  5. 执行python setup.py install

注意:

  1. 包装器函数必须是static函数,且返回值必须是PyObject
  2. 包装器函数的参数列表必须是PyObject self, PyObject args
  3. 模块初始化函数命名一定要为PyInit_模块名

我们可以看到,使用python的扩展API实现一个简单的python扩展还是比较容易的。

使用python的自定义类型

除了使用python自带的扩展API,我们还可以通过python的自定义类型PyTypeObject类实现python的扩展,下面我们通过Python自定义类型PyTypeObject向上面的sample模块添加一个subtract类型。

首先还是定义C函数。

Subtract.h:

#ifndef __SUBTRACT__
#define __SUBTRACT__


#include<stdio.h>


int Subtract(int a,int b);

#endif

Subtract.c:

#include"Subtract.h"


int Subtract(int a,int b)
{
    return a-b;
}

下面是binding代码。

Python_Subtract.h:

#ifndef __PYTHON_SUBTRACT__
#define __PYTHON_SUBTRACT__



#include<stdbool.h>
#include"Python.h"

// 用于保存subtract类的实例对象的数据
typedef struct _SubtractObject
{
    PyObject_HEAD
}SubtractObject;

bool AddSubtractModule(PyObject *module);

#endif

Python_Subtract.c:

#include"Python_Subtract.h"
#include"Subtract.h"



// 定义包装器函数封装C函数(包装器函数就是python解释器和C语言之间的胶水)
static PyObject* Subtract_Wrapper(PyObject *self,PyObject *args)
{
    // 解析python参数列表args(格式化)
    int x,y;
    if(!PyArg_ParseTuple(args , "ii" , &x , &y))
    {
        return NULL;
    }

    // 实际调用的C函数
    int result=Subtract(x,y);

    // 从C语言数据创建一个python对象并返回(格式化)
    return Py_BuildValue("i" , result);

}

// 定义函数表
static PyMethodDef subtractMethods[] = 
{
    // 注意:这里要添加METH_STATIC,否则无法直接调用该函数
    {"Subtract", (PyCFunction)Subtract_Wrapper, METH_STATIC | METH_VARARGS | METH_KEYWORDS, NULL}, // 第二个为包装器函数,第一个参数包装器函数的导出名,最后为函数说明,可以为空
    { NULL, NULL, 0, NULL}
};


// 定义了一个subtract类(它保存了该类的元信息,SubtractObject保存subtract类的实例对象的数据)
PyTypeObject subtractClass = 
{
  PyVarObject_HEAD_INIT(0, 0) "sample.subtract",   /* tp_name */
  sizeof(SubtractObject),                         /* tp_basicsize */
  0,                                           /* tp_itemsize */
  0,                                           /* tp_dealloc */
  0,                                           /* tp_vectorcall_offset */
  0,                                     /* tp_getattr */
  0,                                     /* tp_setattr */
  0,                                     /* tp_reserved */
  0,                                     /* tp_repr */
  0,                                     /* tp_as_number */
  0,                                     /* tp_as_sequence */
  0,                     /* tp_as_mapping */
  0,                                     /* tp_hash  */
  0,                                     /* tp_call */
  0,                                     /* tp_str */
  0,                                     /* tp_getattro */
  0,                                     /* tp_setattro */
  0,                                     /* tp_as_buffer */
  0,                                      /* tp_flags */
  0,                                     /* tp_doc */
  0,                                      /* tp_traverse */
  0,                                      /* tp_clear */
  0,                                     /* tp_richcompare */
  0,                                           /* tp_weaklistoffset */
  0,                                     /* tp_iter */
  0,                                     /* tp_iternext */
  subtractMethods,                       /* tp_methods,subtract类的成员函数 */
  0,                                     /* tp_members */
  0,                                      /* tp_getset */
  0,                                     /* tp_base */
  0,                                     /* tp_dict */
  0,                                     /* tp_descr_get */
  0,                                     /* tp_descr_set */
  0,                                           /* tp_dictoffset */
  0,                                     /* tp_init */
  0,                                     /* tp_alloc */
  0                                       /* tp_new */
};

// 初始化模块
bool AddSubtractModule(PyObject *module)
{
  if (PyType_Ready(&subtractClass) < 0) 
  {
    return false;
  }
  Py_INCREF(&subtractClass);

  // 向模块添加subtract类
  if(PyModule_AddObject(module, "subtract",   (PyObject *)&subtractClass)!=0)
  {
      return false;
  }

  return true;
}

注意命名规范:python binding代码文件以Python_为前缀

最后修改Sample.c:

#include"Python.h"
#include"Add.h"
#include"Python_Subtract.h"



// 定义包装器函数封装C函数(包装器函数就是python解释器和C语言之间的胶水)
static PyObject* Add_Wrapper(PyObject *self,PyObject *args)
{
    // 解析python参数列表args(格式化)
    int x,y;
    if(!PyArg_ParseTuple(args , "ii" , &x , &y))
    {
        return NULL;
    }

    // 实际调用的C函数
    int result=Add(x,y);

    // 从C语言数据创建一个python对象并返回(格式化)
    return Py_BuildValue("i" , result);

}

// 定义函数表
static PyMethodDef sampleMethods[] = 
{
    {"Add", Add_Wrapper, METH_VARARGS, NULL}, // 第二个为包装器函数,第一个参数包装器函数的导出名,最后为函数说明,可以为空
    { NULL, NULL, 0, NULL}
};

// 定义模块
static struct PyModuleDef sampleModule = 
{
        PyModuleDef_HEAD_INIT,
        "sample",           // 模块名
        "A sample module",  // 说明(可以为空)
        -1,                 // 
        sampleMethods       // 函数表
};

// 模块初始化函数,注意命名一定要为PyInit_模块名
// PyMODINIT_FUNC PyInit_sample(void) 
// {
//     return PyModule_Create(&sampleModule);
// }

PyMODINIT_FUNC PyInit_sample(void) 
{
    // 添加Add操作
    PyObject* module=PyModule_Create(&sampleModule);



  // 添加Subtract操作
  AddSubtractModule(module);


    return module;

}

这里我们修改了原来的PyInit_sample()函数,通过两种方式实现了C扩展

下面执行:

python setup.py install

编程成功后,在python中调用

编译后的目录结构:

使用pybind11

pybind11是一个轻量级的只包含头文件的库,主要用于实现python的C++扩展。

下面我们先来看一下如何使用pybind11实现一个简单的C++扩展。我们来实现一个简单的乘法函数。

Python_Multiply.c:

#include<pybind11/pybind11.h>


int Multiply(int i, int j) 
{
    return i * j;
}

/* multiply是模块名,注意不能加引号,m是创建绑定代码的接口,m.def生成了实际的绑定代码 
编译命令:c++ -O3 -Wall -shared -std=c++11 -fPIC `python3 -m pybind11 --includes` Python_Multiply.c -o multiply`python3-config --extension-suffix`
*/
PYBIND11_MODULE(multiply, m) 
{
    m.doc() = "multiply"; 
    m.def("Multiply", &Multiply, "A function which multiplies two numbers");
}

编译:

c++ -O3 -Wall -shared -std=c++11 -fPIC `python3 -m pybind11 --includes` Python_Multiply.c -o multiply`python3-config --extension-suffix`

编译成功后在python中调用

下面我们将上面的Multiply操作添加到上面的sample模块中,添加:

Python_Multiply.h:

#ifndef __PYTHON_MULTIPLY__
#define __PYTHON_MULTIPLY__


#include<stdbool.h>
#include"Python.h"

void AddMultiplyModule(PyObject *module);

#endif

修改Python_Multiply.c:

#include"Python_Multiply.h"
#include<pybind11/pybind11.h>


int Multiply(int i, int j) 
{
    return i * j;
}

/* multiply是模块名,注意不能加引号,m是创建绑定代码的接口,m.def生成了实际的绑定代码 
编译命令:c++ -O3 -Wall -shared -std=c++11 -fPIC `python3 -m pybind11 --includes` Python_Multiply.c -o multiply`python3-config --extension-suffix`
*/
// PYBIND11_MODULE(multiply, m) 
// {
//     m.doc() = "multiply"; 
//     m.def("Multiply", &Multiply, "A function which multiplies two numbers");
// }

void AddMultiplyModule(PyObject *module)
{
    // 通过pybind11实现C/C++扩展
    auto py_module = pybind11::reinterpret_borrow<pybind11::module>(module);
    py_module.def("Multiply", &Multiply, "A function which multiplies two numbers");
}

修改Sample.c:

#include"Python.h"
#include"Add.h"
#include"Python_Subtract.h"
#include"Python_Multiply.h"



// 定义包装器函数封装C函数(包装器函数就是python解释器和C语言之间的胶水)
static PyObject* Add_Wrapper(PyObject *self,PyObject *args)
{
    // 解析python参数列表args(格式化)
    int x,y;
    if(!PyArg_ParseTuple(args , "ii" , &x , &y))
    {
        return NULL;
    }

    // 实际调用的C函数
    int result=Add(x,y);

    // 从C语言数据创建一个python对象并返回(格式化)
    return Py_BuildValue("i" , result);

}

// 定义函数表
static PyMethodDef sampleMethods[] = 
{
    {"Add", Add_Wrapper, METH_VARARGS, NULL}, // 第二个为包装器函数,第一个参数包装器函数的导出名,最后为函数说明,可以为空
    { NULL, NULL, 0, NULL}
};

// 定义模块
static struct PyModuleDef sampleModule = 
{
        PyModuleDef_HEAD_INIT,
        "sample",           // 模块名
        "A sample module",  // 说明(可以为空)
        -1,                 // 
        sampleMethods       // 函数表
};

// 模块初始化函数,注意命名一定要为PyInit_模块名
// PyMODINIT_FUNC PyInit_sample(void) 
// {
//     return PyModule_Create(&sampleModule);
// }

PyMODINIT_FUNC PyInit_sample(void) 
{
    // 添加Add操作
    PyObject* module=PyModule_Create(&sampleModule);

    // 添加Subtract操作
    AddSubtractModule(module);

    // 添加Multiply操作
    AddMultiplyModule(module);

    return module;

}

执行python setup.py install报错:

发现是由于之前用的c语言方式编译的,需要使用C++的方式编译,将所有.c文件修改为.cpp文件,并且修改setup.py:

import glob
import os
from setuptools import find_packages
from setuptools import setup
from distutils.core import setup, Extension


currentPath = os.path.dirname(os.path.abspath(__file__))
pathOfExtensionModule = currentPath
fileNameOfSrc = glob.glob(os.path.join(pathOfExtensionModule, "*.cpp"))
pathOfSrc = [os.path.join(pathOfExtensionModule, s) for s in fileNameOfSrc]

extraCompileArgs = ['-std=c++11','-Wall','-O3']

setup(name='sample',# 这里不一定要与模块名一致
      ext_modules=[
        Extension('sample', # 模块名 
                  sources=pathOfSrc,
                  language='c++',
                  extra_compile_args=extraCompileArgs,
                  include_dirs = [pathOfExtensionModule],
                  define_macros = [('FOO','1')],
                  undef_macros = ['BAR'],
                  library_dirs = ['/usr/local/lib']
                  )
        ]
)

执行python setup.py install安装成功后就可以在python中使用了

附编译成功后的目录结构:

学完上面的内容后,就很容易理解Pytorch中的扩展了。

Pytorch中C/C++的扩展

从setup.py中我们可以发现,C的扩展的编译过程在函数def configure_extension_build()中:

setpu.py:

...

def configure_extension_build():
    ...
    main_sources = ["torch/csrc/stub.cpp"]

   ...

    ################################################################################
    # Declare extensions and package
    ################################################################################

    extensions = []
    packages = find_packages(exclude=('tools', 'tools.*'))
    C = Extension("torch._C",
                  libraries=main_libraries,
                  sources=main_sources,
                  language='c++',
                  extra_compile_args=main_compile_args + extra_compile_args,
                  include_dirs=[],
                  library_dirs=library_dirs,
                  extra_link_args=extra_link_args + main_link_args + [make_relative_rpath('lib')])
    extensions.append(C)

    ...

    return extensions, cmdclass, packages, entry_points, extra_install_requires

...

我们可以看到torch._C的接口由torch\csrc\stub.cpp来定义,我们来看一下"torch/csrc/stub.cpp"文件里是什么,查看stub.cpp我们可以发现,模块初始化是通过PyInitC()函数函数完成的, PyInitC()函数实际上调用的是torch\csrc\Module.cpp中的PyObject* initModule()函数,这个函数是什么作用呢?这个函数其实就是构建了Python的C/C++扩展模块。

torch/csrc/Module.cpp:

…

//NOLINTNEXTLINE(cppcoreguidelines-avoid-c-arrays, modernize-avoid-c-arrays)
// TorchMethods函数表的定义
static PyMethodDef TorchMethods[] = {
  {"_initExtension",  (PyCFunction)THPModule_initExtension,   METH_O,       nullptr},
  {"_autograd_init",  (PyCFunction)THPAutograd_initExtension, METH_NOARGS,  nullptr},
  {"_add_docstr",     (PyCFunction)THPModule_addDocStr,       METH_VARARGS, nullptr},
  {"_init_names",     (PyCFunction)THPModule_initNames,       METH_O,       nullptr},
  {"_has_distributed",(PyCFunction)THPModule_hasDistributed,  METH_NOARGS,  nullptr},
  {"_safe_call",      (PyCFunction)(void(*)())THPModule_safeCall, METH_VARARGS | METH_KEYWORDS, nullptr},
  {"_set_default_tensor_type", (PyCFunction)THPModule_setDefaultTensorType, METH_O, nullptr},
  {"_set_default_dtype", (PyCFunction)THPModule_setDefaultDtype, METH_O, nullptr},
  {"_infer_size",     (PyCFunction)THPModule_inferSize,         METH_VARARGS, nullptr},

…

{nullptr, nullptr, 0, nullptr}
};

…

PyObject* initModule() 
{

…

// 添加函数表的定义
  THPUtils_addPyMethodDefs(methods, TorchMethods);
  THPUtils_addPyMethodDefs(methods, DataLoaderMethods);
  THPUtils_addPyMethodDefs(methods, torch::autograd::python_functions());
  THPUtils_addPyMethodDefs(methods, torch::multiprocessing::python_functions());

…

// 定义模块
  static struct PyModuleDef torchmodule = {
     PyModuleDef_HEAD_INIT,
     "torch._C",
     nullptr,
     -1,
     methods.data()
  };
  // 创建模块
  ASSERT_TRUE(module = PyModule_Create(&torchmodule));



// 初始化其他模块(采用了PyTypeObject自定义类型的方式)
  ASSERT_TRUE(THPWrapper_init(module));
  ASSERT_TRUE(THPGenerator_init(module));
  ASSERT_TRUE(THPException_init(module));
  THPSize_init(module);
  THPDtype_init(module);
  THPDTypeInfo_init(module);
  THPLayout_init(module);
  THPMemoryFormat_init(module);
  THPQScheme_init(module);
  THPDevice_init(module);
  ASSERT_TRUE(THPVariable_initModule(module));
  ASSERT_TRUE(THPFunction_initModule(module));
  ASSERT_TRUE(THPEngine_initModule(module));

…

// 通过pybind11实现C/C++扩展
  auto py_module = py::reinterpret_borrow<py::module>(module);
  py_module.def("_demangle", &c10::demangle);
  py_module.def("_log_api_usage_once", &LogAPIUsageOnceFromPython);

…

return module;
  END_HANDLE_TH_ERRORS
}

我们看一下torch._C的定义:

// 定义模块
  static struct PyModuleDef torchmodule = {
     PyModuleDef_HEAD_INIT,
     "torch._C",
     nullptr,
     -1,
     methods.data()
  };

其中torch._C模块就是采用了python扩展API的方法。

下面我们以_TensorBase类为例来看一下是如何使用python的自定义类型实现扩展的。_TensorBase类是通过THPVariable_initModule()函数添加到torch._C模块中。我们来看一下该函数的实现。

torch/csrc/autograd/python_variable.cpp:

...

PyTypeObject THPVariableType = {
  PyVarObject_HEAD_INIT(nullptr, 0)
  "torch._C._TensorBase",                      /* tp_name */
  sizeof(THPVariable),                         /* tp_basicsize */
  0,                                           /* tp_itemsize */
  (destructor)THPVariable_dealloc,             /* tp_dealloc */
  0,                                           /* tp_vectorcall_offset */
  nullptr,                                     /* tp_getattr */
  nullptr,                                     /* tp_setattr */
  nullptr,                                     /* tp_reserved */
  nullptr,                                     /* tp_repr */
  nullptr,                                     /* tp_as_number */
  nullptr,                                     /* tp_as_sequence */
  &THPVariable_as_mapping,                     /* tp_as_mapping */
  nullptr,                                     /* tp_hash  */
  nullptr,                                     /* tp_call */
  nullptr,                                     /* tp_str */
  nullptr,                                     /* tp_getattro */
  nullptr,                                     /* tp_setattro */
  nullptr,                                     /* tp_as_buffer */
  Py_TPFLAGS_DEFAULT | Py_TPFLAGS_BASETYPE | Py_TPFLAGS_HAVE_GC, /* tp_flags */
  nullptr,                                     /* tp_doc */
  (traverseproc)THPVariable_traverse,          /* tp_traverse */
  (inquiry)THPVariable_clear,                  /* tp_clear */
  nullptr,                                     /* tp_richcompare */
  0,                                           /* tp_weaklistoffset */
  nullptr,                                     /* tp_iter */
  nullptr,                                     /* tp_iternext */
  nullptr,                                     /* tp_methods */
  nullptr,                                     /* tp_members */
  THPVariable_properties,                      /* tp_getset */
  nullptr,                                     /* tp_base */
  nullptr,                                     /* tp_dict */
  nullptr,                                     /* tp_descr_get */
  nullptr,                                     /* tp_descr_set */
  0,                                           /* tp_dictoffset */
  nullptr,                                     /* tp_init */
  nullptr,                                     /* tp_alloc */
  THPVariable_pynew                            /* tp_new */
};

...

bool THPVariable_initModule(PyObject *module)
{
  // 添加多个函数定义PyMethodDef
  static std::vector<PyMethodDef> methods;
  THPUtils_addPyMethodDefs(methods, torch::autograd::variable_methods);
  THPUtils_addPyMethodDefs(methods, extra_methods);
  THPVariableType.tp_methods = methods.data();
  if (PyType_Ready(&THPVariableType) < 0)
    return false;
  Py_INCREF(&THPVariableType);



  // 向模块添加类型
  PyModule_AddObject(module, "_TensorBase",   (PyObject *)&THPVariableType);
  torch::autograd::initTorchFunctions(module);
  torch::autograd::initTensorImplConversion(module);
  return true;
}

我们可以看到THPVariable_initModule()函数通过python的自定义类型PyTypeObject添加了torch._C._TensorBase类,从而实现了python的扩展。

通过上面的分析可以看出pytorch实现C/C++扩展采用的方式就是前面提到python实现扩展的三种方式。

Pytorch扩展之自动生成代码

前面提到的是一般情况下采用手动添加函数表定义和模块定义的方式实现扩展的,但是pytorch的一个特别之处在于pytorch中大部分op采用的是自动生成代码的方式实现扩展的。下面我们来看一下是如何实现的。

注:这里只简单分析python的op到C++的op的调用流程,通过这个调用流程了解pytorch的自动生成代码机制,不深入分析是如何自动生成python绑定代码的,这部分后面会详细分析。

这里以平时最常见的卷积操作为例。

import torch 
input = torch.rand(1,3,3,3) 
conv=torch.nn.Conv2d(3,5,(3,3)) 
result=conv(input)
print(result)

上面执行了一个最简单的卷积操作。这段代码的背后执行了什么呢?我们看一下卷积层的定义。

torch/nn/modules/conv.py:

class Conv2d(_ConvNd):

...

def _conv_forward(self, input, weight):
        if self.padding_mode != 'zeros':
            return F.conv2d(F.pad(input, self._padding_repeated_twice, mode=self.padding_mode),
                            weight, self.bias, self.stride,
                            _pair(0), self.dilation, self.groups)
        return F.conv2d(input, weight, self.bias, self.stride,
                        self.padding, self.dilation, self.groups)



    def forward(self, input):
        return self._conv_forward(input, self.weight)

我们看到其实调用了F.conv2d这个函数,定位到这个函数:

torch/nn/functional.py:

...

conv2d = _add_docstr(torch.conv2d,

...

我们看到其实调用了torch.conv2d()这个函数,那么这个函数哪里来的呢?这个函数不在python前端定义,那一定就是使用了扩展来定义的。既然使用了扩展定义,那么一定会定义一个函数表导出这个conv2d的函数,在pytorch的github工程中搜索”conv2d”,结果发现没有函数定义定义了conv2d这个函数,这是怎么回事呢?通过一番查找资料发现,这个函数表的定义是在编译pytorch过程中自动生成的!在github工程中是招不到这个函数表的。编译下载下来的pytorch工程,然后再搜索”conv2d”:

终于找到这个函数表的定义了!我们来看一下这个函数表的定义。

torch/csrc/autograd/generated/python_torch_functions.cpp:

static PyMethodDef torch_functions[] = {

...

{"conv1d", (PyCFunction)(void(*)(void))THPVariable_conv1d, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"conv2d", (PyCFunction)(void(*)(void))THPVariable_conv2d, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"conv3d", (PyCFunction)(void(*)(void))THPVariable_conv3d, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"conv_tbc", (PyCFunction)(void(*)(void))THPVariable_conv_tbc, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"conv_transpose1d", (PyCFunction)(void(*)(void))THPVariable_conv_transpose1d, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"conv_transpose2d", (PyCFunction)(void(*)(void))THPVariable_conv_transpose2d, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"conv_transpose3d", (PyCFunction)(void(*)(void))THPVariable_conv_transpose3d, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"convolution", (PyCFunction)(void(*)(void))THPVariable_convolution, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"copy_imag", (PyCFunction)(void(*)(void))THPVariable_copy_imag, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"copy_real", (PyCFunction)(void(*)(void))THPVariable_copy_real, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cos", (PyCFunction)(void(*)(void))THPVariable_cos, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cos_", (PyCFunction)(void(*)(void))THPVariable_cos_, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cosh", (PyCFunction)(void(*)(void))THPVariable_cosh, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cosh_", (PyCFunction)(void(*)(void))THPVariable_cosh_, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cosine_embedding_loss", (PyCFunction)(void(*)(void))THPVariable_cosine_embedding_loss, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cosine_similarity", (PyCFunction)(void(*)(void))THPVariable_cosine_similarity, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cross", (PyCFunction)(void(*)(void))THPVariable_cross, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"ctc_loss", (PyCFunction)(void(*)(void))THPVariable_ctc_loss, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cudnn_affine_grid_generator", (PyCFunction)(void(*)(void))THPVariable_cudnn_affine_grid_generator, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cudnn_batch_norm", (PyCFunction)(void(*)(void))THPVariable_cudnn_batch_norm, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cudnn_convolution", (PyCFunction)(void(*)(void))THPVariable_cudnn_convolution, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cudnn_convolution_transpose", (PyCFunction)(void(*)(void))THPVariable_cudnn_convolution_transpose, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},
  {"cudnn_grid_sampler", (PyCFunction)(void(*)(void))THPVariable_cudnn_grid_sampler, METH_VARARGS | METH_KEYWORDS | METH_STATIC, NULL},

...

  {NULL}

};

这个torch/csrc/autograd/generated/python_torch_functions.cpp就是在编译pytorch过程中自动生成的python绑定代码文件,里面包含了torch模块的函数,比如常用的conv2d, relu, log_softmax。

通过上面的分析我们看到python中的很多op是通过自动生成代码的方式实现扩展的。

Pytorch中op的调用关系

了解了Pytorch中自动生成代码机制之后,就很容易理解Pytorch中op的调用关系了。为了了解python到C++的调用关系,我们需要编译debug版本的pytorch,然后在自动生成的python绑定代码处设置断点,最后使用gdb调试跟踪代码的执行流程就可以看到调用关系了。下面我们以torch.abs(x)的执行过程为例看一下调用关系。

test.py:

import torch

x=torch.Tensor([-1,-2,-3])

torch.abs(x)

在torch/csrc/autograd/generated/python_torch_functions.cpp 中的THPVariable_abs()函数里打个断点,然后使用gdb调试,为了方便观察调用关系,自己写了一个简单的日志工具, 然后在每个调用函数里输出日志:

下面我们结合日志看一下调用关系:

我们看一下调用关系:

调用层次文件函数备注
python层test.pytorch.abs(x)
python绑定层torch/csrc/autograd/generated/python_torch_functions.cppstatic PyObject * THPVariable_abs()
第一个调度(dispatch)层build/aten/src/ATen/core/TensorMethods.h,build/aten/src/ATen/TypeDefault.cppinline Tensor Tensor::abs(),Tensor abs(const Tensor & self)T由于abs在所有设备上的实现都是一样的,所以会调度到TypeDefault::abs()
native op(kernel的实现)aten/src/ATen/native/UnaryOps.cppTensor abs(const Tensor& self),Tensor& abs_out(Tensor& result, const Tensor& self);native op部分会执行第二次数据类型的调度

我们可以看到最后会调用到abs_out这个函数,下面我们会说到abs和abs_out之间的关系。最后总结一下这个调用关系:python层->python绑定层->调度层->最后的native op。

python绑定层

我们看一下python绑定层的代码

torch/csrc/autograd/generated/python_torch_functions.cpp:

static PyObject * THPVariable_abs(PyObject* self_, PyObject* args, PyObject* kwargs)
{
  HANDLE_TH_ERRORS



  // 将Python 参数args 和 kwargs 转换为C++参数.
  static PythonArgParser parser({
    "abs(Tensor input, *, Tensor out=None)",
  }, /*traceable=*/true);

  ParsedArgs<2> parsed_args;
  auto _r = parser.parse(args, kwargs, parsed_args);
  if(_r.has_torch_function()) {
    return handle_torch_function(_r, args, kwargs, THPVariableFunctionsModule, "torch");
  }
  if (_r.isNone(1)) { // 调用形式a.abs()
    // aten::abs(Tensor self) -> Tensor
    auto dispatch_abs = [](const Tensor & self) -> Tensor {
      pybind11::gil_scoped_release no_gil;
      return self.abs();
    };
    // dispatch_()释放全局解释器锁,然后调用在 C++ 张量自身上的一个普通的旧方法
    // wrap()将返回的 Tensor 重新包装进 PyObject
    return wrap(dispatch_abs(_r.tensor(0)));
  } else { // 调用形式:torch.abs(a)
    // aten::abs.out(Tensor self, *, Tensor(a!) out) -> Tensor(a!)
    auto dispatch_abs_out = [](Tensor out, const Tensor & self) -> Tensor {
      pybind11::gil_scoped_release no_gil;
      return at::abs_out(out, self);
    };
    return wrap(dispatch_abs_out(_r.tensor(1), _r.tensor(0)));
  }
  Py_RETURN_NONE;
  END_HANDLE_TH_ERRORS
}

PythonArgParser 类的作用是将Python 参数args 和 kwargs 转换为C++参数。

dispatch_():释放全局解释器锁,然后调用在 C++ 张量自身上的一个普通的旧方法。

wrap():将返回的 Tensor 重新包装进 PyObject。

native op的编写

op调用关系的最后就是op的实现了,下面我们看一下实现一个op的基本步骤:

1. 在aten/src/ATen/native/native_functions.yaml声明op

- func: abs(Tensor self) -> Tensor
  use_c10_dispatcher: full
  variants: function, method
  supports_named_tensor: True

native_functions.yaml的作用:

  1. 生成代码(如中间的调度代码)
  2. 生成python绑定代码

2. 如果需要定义导数,还需要在tools/autograd/derivatives.yaml中声明

- name: abs(Tensor self) -> Tensor
  self: grad * self.sign()

3. 在aten/src/ATen/native目录中实现op

实现一个op大致划分为下面三个步骤

  1. 类型检查
  2. 为输出结果分配空间
  3. 数据类型调度

下面是一个模板op的实现示例:

Tensor my_op_out_cpu(Tensor &result,const Tensor &self,const Tensor &other)
{
  // 类型检查
  TORCH_CHECK(result.is_cpu()&&self.is_cpu()&&other.is_cpu());
  TORCH_CHECK(self.dim()==1);
  TORCH_CHECK(self.size()==other.sizes());



  // 为输出结果分配空间
  result.resize_(self.sizes());

  // 数据类型调度
  AT_DISPATCH_ALL_TYPES(self.scalar_type(),"my_op_cpu",[&]
  {
    my_op_cpu_kernel<scalar_t>(result,self,other);

  });
}

写op的时候,你经常会写三个op,以abs为例:

  1. abs: 算子的普通形式

    Tensor abs(const Tensor& self)
    
    {
    
             Tensor result = at::empty({0}, self.options());
    
             return abs_out(result,self);
    
    }
  1. abs: 执行就地操作(in place), 后缀为””的函数执行的是就地操作

    Tensor&  abs_(Tensor& self)
    
    {
    
         return abs_out(self,self);
    
    }
  1. absout: abs_out是真正执行的函数,abs和abs只是对abs_out的简单封装

    Tensor& abs_out(Tensor& result, const Tensor& self)
    
    {
    
        result.resize_(self.sizes());
    
        // …the real implementation…
    
        return result;
    
    }

    实现一个op的第一步就是在native_functions.yaml中声明这个Op,所以了解native_functions.yaml的编码规范是非常重要的。

如何找到python op对应的C++ op

yaml文件中的函数名,C++ op名以及python op名的关系

要找到python op对应的C++ op,首先我们需要知道yaml文件中的函数名,C++ op名以及python op名之间的关系。通过上面的学习,我们可以发现,三者的命名是一样的。这样就大大方便了我们查找相关的代码。

我们知道python中的大部分op是在aten/src/ATen/native下定义的,而定义之前需要在aten/src/ATen/native/native_functions.yaml文件中声明,所以我们首先可以通过aten/src/ATen/native/native_functions.yaml文件来定位。

具体查找方法如下:

  1. 查找aten/src/ATen/native/native_functions.yaml。

  2. 如果能查到,说明这个op是native op, python绑定代码是自动生成的,这种情况下有两种方式:

    1. 直接找op定义的源码

    到aten/src/ATen/native目录下查找,由于C++文件命名是按照对应的python功能模块命名的,所以很容易找到对应的.cpp文件。比如torch.conv2d(),是卷积操作,则找到aten/src/ATen/native/Convolution.cpp文件,然后从native_functions.yaml的编码规范可以知道函数声明,就可以找到对应的at::Tensor conv2d()函数了。注意,对于有dispatch的,需要查找dispatch中的函数,比如_log_softmax,在yaml中的声明为:

    - func: _log_softmax(Tensor self, int dim, bool half_to_float) -> Tensor
    
      use_c10_dispatcher: full
    
      dispatch:
    
        CPU: log_softmax_cpu
    
        CUDA: log_softmax_cuda

    所以如果在cpu上执行,我们需要查找log_softmax_cpu,如果在GPU上执行,需要查找log_softmax_cuda

    2.通过op的调用关系找 首先要找这个op的python绑定代码。native代码的python绑定代码是在编译的时候自动生成的,生成的代码文件在 torch/csrc/autograd/generated/下,生成的绑定代码文件有一定的命名规则:python模块名_functions.cpp或者python类名_methods.cpp。这样根据python模块或者类就可以找到对应的代码文件,然后找函数表,其中函数表中导出的函数名跟python前端命名一样。比如要找torch.conv2d(),这个函数是torch模块的,则找到python_torch_functions.cpp,然后查找”conv2d”,就可以找到函数表的定义了,然后可以使用gdb调试一步一步跟踪代码的执行流程就可以找到对应的C++代码了。或者最直接的是搜索整个工程(使用grep),搜索”conv2d”就可以找到绑定代码,再使用gdb调试

  3. 如果没有查到,说明不是一个native op,python绑定代码是手动添加的,这部分op基本是通过torch/csrc/Module.cpp中的PyObject initModule()添加的,由于这部分op不像native op那样绑定代码文件有一定的命名规范,文件组织有点复杂,一个模块中的函数表可以在多个文件中定义,比如torch.C模块,函数表的定义分布在torch/csrc/Module.cpp, torch/csrc/autograd/init.cpp, torch/csrc/multiprocessing/init.cpp等文件中,所以一般直接搜索github工程查找函数表的定义。比如torch.C._add_docstr()这个函数,直接搜索” _add_docstr”。

如何进行Pytoch性能分析

在使用Pytorch训练网络的时候,有的时候会发现有些网络在某些情况下速度会很慢,此时需要对网络进行性能分析。pytorch自带了性能分析工具torch.autograd.profiler,下面我们就利用该工具进行pytorch的性能分析。

使用profiler输出每个op的计算时间

首先使用torch.autograd.profiler打印出每个Op的计算时间。

profiler的基本使用方法如下:

with torch.autograd.profiler.profile(enabled=True, use_cuda=True, record_shapes=False) as prof:



   需要进行性能分析的代码



print(prof.key_averages().table(sort_by="self_cpu_time_total"))

这里以SSD_MobileNet和SSD_VGG为例,下表为两个网络在cuda和rocm上的profiler输出的结果。

SSD_mobilenet:

cudarocm
cudnn_convolution 206.071usmiopen_convolution 189.372us
cudnn_convolution_backward 652.616usmiopen_convolution_backward 579.941us
batch_norm 261.369usbatch_norm 272.146us
cudnn_batch_norm 231.619usmiopen_batch_norm 194.284us
cudnn_batch_norm_backward 414.876usmiopen_batch_norm_backward 227.372us
thnn_conv_depthwise2d 393.443usmiopen_depthwise_convolution 175.188us
thnn_conv_depthwise2d_backward 1.589msmiopen_depthwise_convolution_backward 330.349us
smooth_l1_loss 62.250ussmooth_l1_loss 36.469us
smooth_l1_loss_backward 173.438ussmooth_l1_loss_backward 60.812us
_log_softmax 50.236us_log_softmax 110.985us
log_softmax 56.094uslog_softmax 121.406us
nll_loss 134.344usnll_loss 352.625us
nll_loss_forward 125.156usnll_loss_forward 343.344us
nll_loss_backward 196.547usnll_loss_backward 241.469us

SSD_VGG16:

cudarocm
cudnn_convolution 1.857msmiopen_convolution 1.943ms
cudnn_convolution_backward 4.131msmiopen_convolution_backward 5.293ms
cudnn_convolution_transpose 1.005msmiopen_convolution_transpose 787.521us
conv_transpose2d 756.807usconv_transpose2d 820.156us
cudnn_convolution_transpose_backward 1.005msmiopen_convolution_transpose_backward 1.649ms
smooth_l1_loss 69.781ussmooth_l1_loss 98.719us
smooth_l1_loss_backward 94.328ussmooth_l1_loss_backward 124.969us
_log_softmax 23.905us_log_softmax 83.309us
log_softmax 29.440uslog_softmax 98.094us
nll_loss 55.680usnll_loss 162.547us
nll_loss_forward 49.844usnll_loss_forward 153.125us
nll_loss_backward 64.781usnll_loss_backward 144.312us

我们发现log_softmax和nll_loss两个op在rocm上运行速度较慢,但是输出结果中有多个op与这两个op有关,比如与log_softmax有关的还有_log_softmax,与nll_loss有关的还有nll_loss_forward,这些op之间有什么关系呢?只有了解这些op之间的关系,我们才能进一步确定是哪个op慢。在对输出结果做进一步分析之前,我们需要写一个test case,复现一下这两个op。编写test case的目的主要是为了方便后面对该op进行优化和调试。如果直接在网络中调试某一个op效率会比较低。

编写test case复现op

要复现op,我们需要知道:

  1. 该op对应的python op
  2. 该op在实际计算过程中的参数

我们知道log_softmax和nll_loss是cross_entropy()函数用到的,所以python端的op为cross_entropy,我们看一下这部分代码:

        pos_idx = pos.unsqueeze(2).expand_as(conf_data)
        neg_idx = neg.unsqueeze(2).expand_as(conf_data)
        conf_p = conf_data[(pos_idx + neg_idx).gt(0)].view(-1, self.num_classes)
        targets_weighted = conf_t[(pos + neg).gt(0)]
        ipdb.set_trace()
        loss_c = F.cross_entropy(conf_p, targets_weighted, reduction='sum')

然后我们需要知道conf_p和targets_weighted这两个tensor的大小,我们将这两个tensor的大小打印出来。这里以SSD_VGG为例,输出结果如下:

下面我们就可以写test case了。

import torch
import os
import torch.nn.functional as F
os.environ['CUDA_VISIBLE_DEVICES']="0"



for i in range(100):
    with torch.autograd.profiler.profile(enabled=True, use_cuda=True, record_shapes=False) as prof:
        conf_p=torch.randn(2896,2).cuda()
        targets_weighted=torch.randint(0,1,(2896,)).cuda()
        loss_c = F.cross_entropy(conf_p, targets_weighted, reduction='sum')
    print(prof.key_averages().table(sort_by="self_cpu_time_total"))

由于rocm上第一次计算较慢,这里我们需要写一个循环。

我们来看一下中间的一组输出结果:

上面的结果可以看出test case复现了rocm上的op的计算。

下面在cuda上测试一下,输出结果如下:

我们可以看到在cuda上也复现了op的计算。

复现了op之后,就可以借助该test case对输出结果做进一步的分析了。

分析test case的profiler输出结果

下面我们在cuda上以log_softmax为例分析一下上面test case的profiler输出结果。前面提到了与log_softmax有关的还有_log_softmax,那这两个op有什么关系呢?我们看一下log_softmax这个op的调用关系,还是采用打印日志的方式:

我们可以看到log_softmax调用了_log_softmax,这样我们就可以定位到我们需要的那个op了。通过日志,我们发现最终调用的native op为log_softmax_cuda()。下面在rocm平台采用上述方法看一下op的调用关系,这里省略了中间调度层日志:

我们看到在rocm平台上,最终调用的native op为ATen/native/hip/SoftMax.hip中的log_softmax_cuda()。这就是我们需要优化的op了。

下面我们进一步分析一下这个log_softmax_cuda()函数,我们看到是用到了host_softmax()这个模板函数。下面我们将这个程序分段打印日志,通过日志分析每一段的耗时:

通过日志我们看到第690行到底750行最耗时。下面继续对这两段程序进行分析,还是采用打印日志的方式:

通过日志观察到有三处地方比较耗时:

693~695:114us

710~726:61us

726~729:79us

下面以693~695为例来说明,我们将这段代码与cuda进行对比:

通过与cuda对比我们发现,693~695这段代码,cuda上只需要约67us。

我们来看一下这段代码是什么:

Tensor output = half_to_float ? at::empty_like(input, input.options().dtype(ScalarType::Float), LEGACY_CONTIGUOUS_MEMORY_FORMAT) : at::empty_like(input, LEGACY_CONTIGUOUS_MEMORY_FORMAT);

最后需要优化的代码就是这段代码。

总结

pytorch进行性能分析基本步骤:

  1. 使用profiler输出每个op的计算时间
  2. 编写test case复现op,方便后面对op的调试和优化
  3. 分析test case的profiler输出结果,主要是了解op之间的调用关系找到我们最终需要优化的native op
  4. 通过分析op每段代码的耗时,找到最终要优化的代码
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

技术瘾君子1573

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值