目录
- TensorRT的核心在于对模型算子的优化（合并算子、利用GPU特性选择特定核函数等多种策略），通过tensorRT，能够在Nvidia系列GPU上获得最好的性能
- 因此tensorRT的模型，需要在目标GPU上实际运行的方式选择最优算法和配置
- 也因此tensorRT生成的模型只能在特定条件下运行（编译的trt版本、cuda版本、编译时的GPU型号）
- 主要知识点，是模型结构定义方式、编译过程配置、推理过程实现、插件实现、onnx理解
合并算子：
工作流程：
常见的方案：
方案一：基于tensorRT的发布，又有人在之上做了工作https://github.com/wang-xinyu/tensorrtx。为每个模型写硬代码，并已写好了大量的常见模型代码
方案二：onnx路线的模型编译、推理和部署，原因主要有
若使用onnx，则导出或者修改好的onnx模型，可以轻易的移植到其他引擎上、例如ncnn、rknn，这一点硬代码无法做到。并且用于排查错误，修改调整时也非常方便
TensorRT库文件：
学习使用TensorRT-CPP的API构建网络模型，并进行编译的流程
TensorRT工作流程如下图：
- 首先定义网络
- 优化builder参数
- 通过builder生成engine，用于模型保存、推理等
- engine可以通过序列化和逆序列转化模型数据类型（转化为二进制byte文件，加快传输速率），再进一步推动模型由输入张量到输出张量的推理）
// tensorRT include
#include <NvInfer.h>
#include <NvInferRuntime.h>
// cuda include
#include <cuda_runtime.h>
// system include
#include <stdio.h>
class TRTLogger : public nvinfer1::ILogger{
public:
virtual void log(Severity severity, nvinfer1::AsciiChar const* msg) noexcept override{
if(severity <= Severity::kVERBOSE){ //自己判断日志的级别来打印哪些
printf("%d: %s\n", severity, msg);
}
}
};
nvinfer1::Weights make_weights(float* ptr, int n){
nvinfer1::Weights w;
w.count = n;
w.type = nvinfer1::DataType::kFLOAT;
w.values = ptr;
return w;
}int main(){
// 本代码主要实现一个最简单的神经网络 figure/simple_fully_connected_net.png
TRTLogger logger; // logger是必要的，用来捕捉warning和info等
// ----------------------------- 1. 定义 builder, config 和network -----------------------------
// 这是基本需要的组件
//形象的理解是你需要一个builder去build这个网络，网络自身有结构，这个结构可以有不同的配置
nvinfer1::IBuilder* builder = nvinfer1::createInferBuilder(logger);
// 创建一个构建配置，指定TensorRT应该如何优化模型，tensorRT生成的模型只能在特定配置下运行
nvinfer1::IBuilderConfig* config = builder->createBuilderConfig();
// 创建网络定义，其中createNetworkV2(1)表示采用显性batch size，新版tensorRT(>=7.0)时，不建议采用0非显性batch size
// 因此贯穿以后，请都采用createNetworkV2(1)而非createNetworkV2(0)或者createNetwork
nvinfer1::INetworkDefinition* network = builder->createNetworkV2(1);
// 构建一个模型
/*
Network definition:
image
|
linear (fully connected) input = 3, output = 2, bias = True w=[[1.0, 2.0, 0.5], [0.1, 0.2, 0.5]], b=[0.3, 0.8]
|
sigmoid
|
prob
*/
// ----------------------------- 2. 输入，模型结构和输出的基本信息 -----------------------------
const int num_input = 3; // in_channel
const int num_output = 2; // out_channel
float layer1_weight_values[] = {1.0, 2.0, 0.5, 0.1, 0.2, 0.5}; // 前3个给w1的rgb，后3个给w2的rgb
float layer1_bias_values[] = {0.3, 0.8};
//输入指定数据的名称、数据类型和完整维度，将输入层添加到网络
nvinfer1::ITensor* input = network->addInput("image", nvinfer1::DataType::kFLOAT, nvinfer1::Dims4(1, num_input, 1, 1));
nvinfer1::Weights layer1_weight = make_weights(layer1_weight_values, 6);
nvinfer1::Weights layer1_bias = make_weights(layer1_bias_values, 2);
//添加全连接层
auto layer1 = network->addFullyConnected(*input, num_output, layer1_weight, layer1_bias); // 注意对input进行了解引用
//添加激活层
auto prob = network->addActivation(*layer1->getOutput(0), nvinfer1::ActivationType::kSIGMOID); // 注意更严谨的写法是*(layer1->getOutput(0)) 即对getOutput返回的指针进行解引用
// 将我们需要的prob标记为输出
network->markOutput(*prob->getOutput(0));
printf("Workspace Size = %.2f MB\n", (1 << 28) / 1024.0f / 1024.0f); // 256Mib
config->setMaxWorkspaceSize(1 << 28);
builder->setMaxBatchSize(1); // 推理时 batchSize = 1
// ----------------------------- 3. 生成engine模型文件 -----------------------------
//TensorRT 7.1.0版本已弃用buildCudaEngine方法，统一使用buildEngineWithConfig方法
nvinfer1::ICudaEngine* engine = builder->buildEngineWithConfig(*network, *config);
if(engine == nullptr){
printf("Build engine failed.\n");
return -1;
}
// ----------------------------- 4. 序列化模型文件并存储 -----------------------------
// 将模型序列化，并储存为文件
nvinfer1::IHostMemory* model_data = engine->serialize();
FILE* f = fopen("engine.trtmodel", "wb");
fwrite(model_data->data(), 1, model_data->size(), f);
fclose(f);
// 卸载顺序按照构建顺序倒序
model_data->destroy();
engine->destroy();
network->destroy();
config->destroy();
builder->destroy();
printf("Done.\n");
return 0;
}注意：
- 必须使用createNetworkV2，并指定为1（表示显性batch）。createNetwork已经废弃，非显性batch官方不推荐。这个方式直接影响推理时enqueue还是enqueueV2
- builder、config等指针，记得释放，否则会有内存泄漏，使用ptr->destroy()释放
- markOutput表示是该模型的输出节点，mark几次，就有几个输出，addInput几次就有几个输入。这与推理时相呼应
- workspaceSize是工作空间大小，某些layer需要使用额外存储时，不会自己分配空间，而是为了内存复用，直接找tensorRT要workspace空间。指的这个意思
- 一定要记住，保存的模型只能适配编译时的trt版本、编译时指定的设备。也只能保证在这种配置下是最优的。如果用trt跨不同设备执行，有时候可以运行，但不是最优的，也不推荐
编译好的模型进行推理
void inference(){
// ------------------------------ 1. 准备模型并加载 ----------------------------
TRTLogger logger;
auto engine_data = load_file("engine.trtmodel");
// 执行推理前，需要创建一个推理的runtime接口实例。与builer一样，runtime需要logger：
nvinfer1::IRuntime* runtime = nvinfer1::createInferRuntime(logger);
// 将模型从读取到engine_data中，则可以对其进行反序列化以获得engine
nvinfer1::ICudaEngine* engine = runtime->deserializeCudaEngine(engine_data.data(), engine_data.size());
if(engine == nullptr){
printf("Deserialize cuda engine failed.\n");
runtime->destroy();
return;
}
nvinfer1::IExecutionContext* execution_context = engine->createExecutionContext();
cudaStream_t stream = nullptr;
// 创建CUDA流，以确定这个batch的推理是独立的
cudaStreamCreate(&stream);
/*
Network definition:
image
|
linear (fully connected) input = 3, output = 2, bias = True w=[[1.0, 2.0, 0.5], [0.1, 0.2, 0.5]], b=[0.3, 0.8]
|
sigmoid
|
prob
*/
// ------------------------------ 2. 准备好要推理的数据并搬运到GPU ----------------------------
float input_data_host[] = {1, 2, 3};
float* input_data_device = nullptr;
float output_data_host[2];
float* output_data_device = nullptr;
cudaMalloc(&input_data_device, sizeof(input_data_host));
cudaMalloc(&output_data_device, sizeof(output_data_host));
cudaMemcpyAsync(input_data_device, input_data_host, sizeof(input_data_host), cudaMemcpyHostToDevice, stream);
// 用一个指针数组指定input和output在gpu中的指针。
float* bindings[] = {input_data_device, output_data_device};
// ------------------------------ 3. 推理并将结果搬运回CPU ----------------------------
bool success = execution_context->enqueueV2((void**)bindings, stream, nullptr);
cudaMemcpyAsync(output_data_host, output_data_device, sizeof(output_data_host), cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);
printf("output_data_host = %f, %f\n", output_data_host[0], output_data_host[1]);
// ------------------------------ 4. 释放内存 ----------------------------
printf("Clean memory\n");
cudaStreamDestroy(stream);
execution_context->destroy();
engine->destroy();
runtime->destroy();
// ------------------------------ 5. 手动推理进行验证 ----------------------------
const int num_input = 3;
const int num_output = 2;
float layer1_weight_values[] = {1.0, 2.0, 0.5, 0.1, 0.2, 0.5};
float layer1_bias_values[] = {0.3, 0.8};
printf("手动验证计算结果：\n");
for(int io = 0; io < num_output; ++io){
float output_host = layer1_bias_values[io];
for(int ii = 0; ii < num_input; ++ii){
output_host += layer1_weight_values[io * num_input + ii] * input_data_host[ii];
}
// sigmoid
float prob = 1 / (1 + exp(-output_host));
printf("output_prob[%d] = %f\n", io, prob);
}
}
- bindings是tensorRT对输入输出张量的描述，bindings = input-tensor + output-tensor。比如input有a，output有b, c, d，那么bindings = [a, b, c, d]，bindings[0] = a，bindings[2] = c。此时看到engine->getBindingDimensions(0)你得知道获取的是什么
- enqueueV2是异步推理，加入到stream队列等待执行。输入的bindings则是tensors的指针（注意是device pointer）。其shape对应于编译时指定的输入输出的shape（这里只演示全部shape静态）
- createExecutionContext可以执行多次，允许一个引擎具有多个执行上下文，不过看看就好，别当真。
动态shape，即编译时指定可动态的范围[L-H]，推理时可以允许 L <= shape <= H
// --------------------------------- 2.1 关于profile ----------------------------------
// 如果模型有多个输入，则必须多个profile
auto profile = builder->createOptimizationProfile();
// 配置最小允许1 x 1 x 3 x 3
profile->setDimensions(input->getName(), nvinfer1::OptProfileSelector::kMIN, nvinfer1::Dims4(1, num_input, 3, 3));
profile->setDimensions(input->getName(), nvinfer1::OptProfileSelector::kOPT, nvinfer1::Dims4(1, num_input, 3, 3));
// 配置最大允许10 x 1 x 5 x 5
// if networkDims.d[i] != -1, then minDims.d[i] == optDims.d[i] == maxDims.d[i] == networkDims.d[i]
profile->setDimensions(input->getName(), nvinfer1::OptProfileSelector::kMAX, nvinfer1::Dims4(maxBatchSize, num_input, 5, 5));
config->addOptimizationProfile(profile);
- OptimizationProfile是一个优化配置文件，用来指定输入的shape可以变换的范围的，不要被优化两个字蒙蔽了双眼
- 如果onnx的输入某个维度是-1，表示该维度动态，否则表示该维度是明确的，明确维度的minDims, optDims, maxDims一定是一样的
用python将torch转为onnx:
class Model(torch.nn.Module):
def __init__(self):
super().__init__()
self.conv = nn.Conv2d(1, 1, 3, padding=1)
self.relu = nn.ReLU()
self.conv.weight.data.fill_(1)
self.conv.bias.data.fill_(0)
def forward(self, x):
x = self.conv(x)
x = self.relu(x)
return x
# 这个包对应opset11的导出代码，如果想修改导出的细节，可以在这里修改代码
# import torch.onnx.symbolic_opset11
print("对应opset文件夹代码在这里：", os.path.dirname(torch.onnx.__file__))
model = Model()
dummy = torch.zeros(1, 1, 3, 3)
torch.onnx.export(
model,
# 这里的args，是指输入给model的参数，需要传递tuple，因此用括号
(dummy,),
# 储存的文件路径
"demo.onnx",
# 打印详细信息
verbose=True,
# 为输入和输出节点指定名称，方便后面查看或者操作
input_names=["image"],
output_names=["output"],
# 这里的opset，指，各类算子以何种方式导出，对应于symbolic_opset11
opset_version=11,
# 表示他有batch、height、width3个维度是动态的，在onnx中给其赋值为-1
# 通常，我们只设置batch为动态，其他的避免动态
dynamic_axes={
"image": {0: "batch", 2: "height", 3: "width"},
"output": {0: "batch", 2: "height", 3: "width"},
}
)
- ONNX的本质，是一种Protobuf格式文件
- Protobuf则通过onnx-ml.proto编译得到onnx-ml.pb.h和onnx-ml.pb.cc或onnx_ml_pb2.py
- 然后用onnx-ml.pb.cc和代码来操作onnx模型文件，实现增删改
- onnx-ml.proto则是描述onnx文件如何组成的，具有什么结构，他是操作onnx经常参照的东西
Onnx主要结构：日日日日日日日日日日日日日日
model：表示整个onnx的模型，包含图结构和解析器格式、opset版本、导出程序类型
model.graph：表示图结构，通常是我们netron看到的主要结构
model.graph.node：表示图中的所有节点，数组，例如conv、bn等节点就是在这里的，通过input、output表示节点之间的连接关系
model.graph.initializer：权重类的数据大都储存在这里
model.graph.input：整个模型的输入储存在这里，表明哪个节点是输入节点，shape是多少
model.graph.output：整个模型的输出储存在这里，表明哪个节点是输出节点，shape是多少
表示onnx中有节点类型叫node
- input属性，是repeated，即重复类型，数组
- output属性，是repeated，即重复类型，数组
- name属性是string类型
- 对于repeated是数组，对于optional无视他
- 对于input = 1，后面的数字是id，无视他
我们只关心是否数组，类型是什么
查看onnx信息：
model = onnx.load("demo.change.onnx")
#打印信息
print("==============node信息")
# print(helper.printable_graph(model.graph))
print(model)
conv_weight = model.graph.initializer[0]
conv_bias = model.graph.initializer[1]
# 数据是以protobuf的格式存储的，因此当中的数值会以bytes的类型保存，通过np.frombuffer方法还原成类型为float32的ndarray
print(f"===================={conv_weight.name}==========================")
print(conv_weight.name, np.frombuffer(conv_weight.raw_data, dtype=np.float32))
print(f"===================={conv_bias.name}==========================")
print(conv_bias.name, np.frombuffer(conv_bias.raw_data, dtype=np.float32))创建onnx：
import onnx # pip install onnx>=1.10.2
import onnx.helper as helper
import numpy as np
# https://github.com/onnx/onnx/blob/v1.2.1/onnx/onnx-ml.proto
nodes = [
helper.make_node(
name="Conv_0", # 节点名字，不要和op_type搞混了
op_type="Conv", # 节点的算子类型, 比如'Conv'、'Relu'、'Add'这类，详细可以参考onnx给出的算子列表
inputs=["image", "conv.weight", "conv.bias"], # 各个输入的名字，结点的输入包含：输入和算子的权重。必有输入X和权重W，偏置B可以作为可选。
outputs=["3"],
pads=[1, 1, 1, 1], # 其他字符串为节点的属性，attributes在官网被明确的给出了，标注了default的属性具备默认值。
group=1,
dilations=[1, 1],
kernel_shape=[3, 3],
strides=[1, 1]
),
helper.make_node(
name="ReLU_1",
op_type="Relu",
inputs=["3"],
outputs=["output"]
)
]
initializer = [
helper.make_tensor(
name="conv.weight",
data_type=helper.TensorProto.DataType.FLOAT,
dims=[1, 1, 3, 3],
vals=np.array([1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0], dtype=np.float32).tobytes(),
raw=True
),
helper.make_tensor(
name="conv.bias",
data_type=helper.TensorProto.DataType.FLOAT,
dims=[1],
vals=np.array([0.0], dtype=np.float32).tobytes(),
raw=True
)
]
inputs = [
helper.make_value_info(
name="image",
type_proto=helper.make_tensor_type_proto(
elem_type=helper.TensorProto.DataType.FLOAT,
shape=["batch", 1, 3, 3]
)
)
]
outputs = [
helper.make_value_info(
name="output",
type_proto=helper.make_tensor_type_proto(
elem_type=helper.TensorProto.DataType.FLOAT,
shape=["batch", 1, 3, 3]
)
)
]
graph = helper.make_graph(
name="mymodel",
inputs=inputs,
outputs=outputs,
nodes=nodes,
initializer=initializer
)
# 如果名字不是ai.onnx，netron解析就不是太一样了
opset = [
helper.make_operatorsetid("ai.onnx", 11)
]
# producer主要是保持和pytorch一致
model = helper.make_model(graph, opset_imports=opset, producer_name="pytorch", producer_version="1.9")
onnx.save_model(model, "my.onnx")
print(model)
print("Done.!")ONNX重点：
- ONNX的主要结构：graph、graph.node、graph.initializer、graph.input、graph.output
- ONNX的节点构建方式：onnx.helper，各种make函数
- ONNX的proto文件，https://github.com/onnx/onnx/blob/main/onnx/onnx-ml.proto
- 理解模型结构的储存、权重的储存、常量的储存、netron的解读对应到代码中的部分
- ONNX的解析器的理解，包括如何使用nv发布的解析器源代码https://github.com/onnx/onnx-tensorrt
预处理preprocess.onnx
import torch
'''用pytorch写好预处理，生成再加载'''
class Preprocess(torch.nn.modules):
def __init__(self) -> None:
super().__init__()
self.mean = torch.rand(1,1,1,3)
self.std = torch.rand(1,1,1,3)
def forward(self,x):
# x = B*H*W*C Uint8
# y = B*C*H*W F=loart32 减去均值除以标准差
x.float()
x = (x/255.0 - self.mean) /self.std
return x
pre = Preprocess()
torch.onnx.export(
pre,(torch.zeros(1,640,640,3,dtype=torch.uint8),), "preprocess.onnx"
)
pre_onnx = onnx.load("preprocess.onnx")
#0.先把pre_onnx的所有节点以及输入输出名称都加上前缀
#1.yolov5中的image输入节点修改为pre_onnx的输出节点
#2.把pre_onnx的node全部放到yolov5s的node中
#3.把pre_onnx的输入名称作为yolov5s的input名称
for n in pre_onnx.graph.node:
n.name = f"pre/{n.name}"
for n in model.graph.node:
if n.name == "Conv_0":
n.imput[0] = "pre/" + pre_onnx.graph.output[0].name
#2.将pre_onnxpre_onnx的node全部放到yolov5s的node中
for n in pre_onnx.graph.node:
model.graph.node.append(n)
input_name = "pre/" + pre_onnx.graph.input[0].name
model.input[0].CopyFrom(pre_onnx.graph.input[0])
model.input[0].name = input_name正确导出ONNX：
- 对于任何用到shape、size返回值的参数时，例如：tensor.view(tensor.size(0), -1)这类操作，避免直接使用tensor.size的返回值，而是加上int转换，tensor.view(int(tensor.size(0)), -1)，断开跟踪。
- 对于nn.Upsample或nn.functional.interpolate函数，使用scale_factor指定倍率，而不是使用size参数指定大小。
- 对于reshape、view操作时，-1的指定请放到batch维度。其他维度可以计算出来即可。batch维度禁止指定为大于-1的明确数字
- torch.onnx.export指定dynamic_axes参数，并且只指定batch维度，禁止其他动态
- 使用opset_version=11，不要低于11
- 避免使用inplace操作，例如y[…, 0:2] = y[…, 0:2] * 2 - 0.5
- 尽量少的出现5个维度，例如ShuffleNet Module，可以考虑合并wh避免出现5维
- 尽量把让后处理部分在onnx模型中实现，降低后处理复杂度
- 掌握了这些，就可以保证后面各种情况的顺利了
ONNX解析器：
onnx解析器有两个选项：
- libnvonnxparser.so
- https://github.com/onnx/onnx-tensorrt（源代码）。 使用源代码的目的，是为了更好的进行自定义封装，简化插件开发或者模型编译的过程，更加具有定制化，遇到问题可以调试
插件的实现:
重点：
1. 如何在 pytorch里面导出一个插件2. 插件解析时如何对应，在 onnx parser 中如何处理3. 插件的 creator实现4. 插件的具体实现， 继承自IPluginV2DynamicExt5. 插件的序列化与反序列化
MYSELU:
#include "onnx-tensorrt/onnxplugin.hpp"
using namespace ONNXPlugin;
static __device__ float sigmoid(float x){
return 1 / (1 + expf(-x));
}
static __global__ void MYSELU_kernel_fp32(const float* x, float* output, int edge) {
int position = threadIdx.x + blockDim.x * blockIdx.x;
if(position >= edge) return;
output[position] = x[position] * sigmoid(x[position]);
}
class MYSELU : public TRTPlugin {
public:
SetupPlugin(MYSELU); //定义宏
virtual void config_finish() override{
printf("\033[33minit MYSELU config: %s\033[0m\n", config_->info_.c_str());
printf("weights count is %d\n", config_->weights_.size());
}
int enqueue(const std::vector<GTensor>& inputs, std::vector<GTensor>& outputs, const std::vector<GTensor>& weights, void* workspace, cudaStream_t stream) override{
int n = inputs[0].count();
const int nthreads = 512;
int block_size = n < nthreads ? n : nthreads;
int grid_size = (n + block_size - 1) / block_size;
//执行核函数
MYSELU_kernel_fp32 <<<grid_size, block_size, 0, stream>>> (inputs[0].ptr<float>(), outputs[0].ptr<float>(), n);
return 0;
}
};
RegisterPlugin(MYSELU); //注册插件Int8量化:
int8量化是利用int8乘法替换float32乘法实现性能加速的一种方法
1. 对于常规模型有： y = kx + b ，此时 x 、 k 、 b 都是 float32, 对于 kx 的计算使用 float32 的乘法2. 对于 int8 模型有： y = tofp32(toint8(k) * toint8(x)) + b， 其中 int8 * int8 结果为 int163. 因此 int8 模型解决的问题是如何将 float32 合理的转换为 int8 ，使得精度损失最小4. 也因此，经过 int8 量化的精度会受到影响
Int8量化步骤：
1. 配置setFlag nvinfer1::BuilderFlag::kINT8
2. 实现Int8EntropyCalibrator类并继承自IInt8EntropyCalibrator2
3. 实例化Int8EntropyCalibrator并且设置到config.setInt8Calibrator
4. Int8EntropyCalibrator的作用，是读取并预处理图像数据作为输入
- 标定过程的理解：对于输入图像A，使用FP32推理后得到P1再用INT8推理得到 P2，调整int8权重使得P1与P2足够的接近
- 因此标定时需要使用一些图像，正常发布时，使用100张图左右即可
Int8EntropyCalibrator类主要关注：
- getBatchSize，告诉引擎，这次标定的batch是多少
- getBatch，告诉引擎，这次标定的输入数据是什么，把指针赋值给bindings即可，返回false表示没有数据了
- readCalibrationCache，若从缓存文件加载标定信息，则可避免读取文件和预处理，若该函数返回空指针则表示没有缓存，程序会重新通过getBatch重新计算
- writeCalibrationCache，当标定结束后，会调用该函数，我们可以储存标定后的缓存结果，多次标定可以使用该缓存实现加速