TensorRT5.1入门 samplePlugin

最新推荐文章于 2024-02-23 16:13:14 发布

Nonentity_neu

最新推荐文章于 2024-02-23 16:13:14 发布

阅读量2k

点赞数 1

分类专栏： TensorRT 文章标签： TensorRT

本文链接：https://blog.csdn.net/qq_33120609/article/details/93905105

版权

TensorRT 专栏收录该内容

19 篇文章 3 订阅

订阅专栏

环境: Ubuntu16.04 , gcc5.4.0 , TensorRT 5.1.5.0 .

1. Introduction

为了尽早做好准确的网络优化测试,我开始学习如何使用自定义层.但因为绝大多数的实践博文都是如何在TensorRT上跑通自己的网络,而我的需求是要得到测试网络在TensorRT的帮助下前馈时间的缩减程度,即和不用TensorRT情况下的对比数据.
在不需要自定义层的网络中,观察到效果确实不错.尤其是在亲儿子GoogleNet上,跑的很好,如下图所示.
在这里插入图片描述
但是,实际应用起来,就必须要学会把自定义层放进

2. samplePlugin 结构

samplePlugin是自定义层的构建demo
结合工程的README,最为核心的部分是 IPluginExt类和IPluginFactory类.如这个sample中展示的(C++水平非常有限)FCPlugin是由IPluginExt继承而来,PluginFactroy是由IPluginFactory继承而来.
我的理解是IPluginFactory派生的类作为一个plugin产出工厂,而IPluginExt是plugin的图纸.如这个sample中展示的,FCPlugin是用来定义"自己定义的全连接层",而PluginFactory更主要的是把创建的FCPlugin实例化.

下面,按照教程上的顺序一一记录.

Define Network

getNbOutput()用来返回该层output的个数
getOutputDimensions()用来获取Output tensor的信息

index : output的index
inputs: Dims类型,内含input tensor的信息
nbInputDims: input tensor的个数

在这个sample中,上图是我Debug的结果,对应的
index=0

FC层只有一个输出,可能像Slice层就有两个输出,所以可能Slice自定义时,前后两次调用这个函数,index就会分别为0和1,这里只有一个,序号也就只是0

inputs

MAX_DIMS是一个tensor支持的最大维度,这里设置为8
nbDims是输出的维度
d 和 type 是对应关系. eg:500 指的是channel type,1指的 spatial type

nbInputDims

input tensor 的数量,这里是1

所以,这个方法的目的就是返回输出tensor的信息（以Dims这种类型的方式），这个函数所做的就是通过输入的Dims* inputs来求得Dims 数据类型的输出信息。

Enabling custom layers in NvCaffeParser

这一步是利用PluginFactory上为实例化FCPlugin做准备

isPlugin/isPluginExt都是为了下文assert(layername)准备的.
tensorRT中的自定义层,定义的是层的名称,不是层的type(还没看自定义层的注册机制,按理说是可以定义type的),所以这里定义name就很重要了,比如sample里面将这个层定义为"ip2",那么在prototxt中的写法就应该如下图所示.

就是只要把名字定义成ip2,type是随意的(以我现阶段的test,改变type是不影响的)
virtual nvinfer1::IPlugin* createPlugin
这个函数是用来实例化FCPlugin的,官方的解释是从serialized data中实例化(create)plugin.

layername
weight:　Weights指针,内部存储weight的type,values和count,比如这个全连接层的weight的count为5000(50010)
nbWeights:　int类型,看FCPlugin的构造函数，这里的2,指的是filter和bias两种参数.

所以这个函数可能在只是验证了部分参数信息，然后利用类中定义的智能指针mPlugin来实例化一个新的FCPlugin,最后返回这个指针(不是很懂智能指针).

Building the Engine

这里介绍的是FCPlugin中定义的两个配置函数,主要定义支持的Format(NCHW…)和type(Float32,Float16…)

Serializing and deserializing

这里介绍的是串并行化的过程

getSerializationSize()
得到总的需要串行化的长度
serialize()
执行串行化
IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength) override
在部署的时候,会先调用IPluginFactory中的这个函数,再次实例化FCPlugin.
FCPlugin(const void* data, size_t length)
这里和在串行化之前不同,串行化(serialized)之前是为了构造Engine然后串行化存到设备中.这里是为了从设备中读出串行化的数据,将其并行化做Inferrence.
所以其实,这第二个FCPlugin的构造函数大致就是读数据,并行化;最后的assert确保数据读取正确.

Resource management and execution

initialize()和terminate()
在自定义层被执行之前,plugin就已经被初始化了.以下的部分介绍了在plugin的生命周期中,资源初始化和获取的手段和位置.在这个example中,weight先被存在CPU memory中,所以在building阶段,每个配置的测试都可以被设置成固定格式,然后在initialize的时候被存到device中.
initialize方法创建了需要的cuBLAS和cuDNN句柄,设置了tensor描述符,分配了内存,还将weight复制到了device memory上.相反的,terminate方法负责销毁.
enqueue()
plugin的核心就是enqueue(),这个方法用来在runtime的时候执行plugin.
这个函数的输入提供了batch size,inputs,outputs,cuBLAS和cuDNN提供了操作方法(卷积的偏移,和卷积的点乘),在building the Engine步骤又设置了type和format,因此Plugin可以据此执行.
clone()

3.多个plugin层的使用学习记录

这里仅仅是把两个FC层尝试放到代码中.

设置pluginImplement的头文件和源文件(也就是Factory的文件)
(1) isPluginExt里面加上需要判断名字的两个层,

(2) 尝试修改create方法
(3)修改另一个create
(4)把对应的destroy方法也修改了就行了
添加对应的layer的头文件和源文件
这里加了两个FC层,分别定义为ip1.h,ip2.h,ip1.cpp,ip2.cpp.如上面所说的,在PluginFactory(我命名的是pluginImplement)完成改动后,直接链接这些层文件即可.
下面记录一下自己看的层的具体定义,拿一个FC层做例子:

class FCPlugin : public IPluginExt
{
public:
FC2Plugin(const Weights* weights, int nbWeights, int nbOutputChannels);
FC2Plugin(const void* data, size_t length);
~FC2Plugin();
virtual int getNbOutputs() const override;
virtual Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override;
bool supportsFormat(DataType type, PluginFormat format) const override ;
void configureWithFormat(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, DataType type, PluginFormat format, int maxBatchSize) override;
virtual int initialize() override;
virtual void terminate() override;
virtual size_t getWorkspaceSize(int maxBatchSize) const override;
virtual int enqueue(int batchSize, const void* const* inputs, void** outputs, void* workspace, cudaStream_t stream) override;
virtual size_t getSerializationSize() override;
virtual void serialize(void* buffer) override;
private:
size_t type2size(DataType type) { return type == DataType::kFLOAT ? sizeof(float) : sizeof(__half); }
template < typename T>
void write(char*& buffer, const T& val)
{
reinterpret_cast<T>(buffer) = val;
buffer += sizeof(T);
}
template < typename T>
void read(const char*& buffer, T& val)
{
val = reinterpret_cast<const T>(buffer);
buffer += sizeof(T);
}
void* copyToDevice(const void* data, size_t count);
void convertAndCopyToDevice(void*& deviceWeights, const Weights& weights);
void convertAndCopyToBuffer(char*& buffer, const Weights& weights);
void deserializeToDevice(const char*& hostBuffer, void*& deviceWeights, size_t size);
int mNbOutputChannels, mNbInputChannels;
Weights mKernelWeights, mBiasWeights;
DataType mDataType{DataType::kFLOAT};
void* mDeviceKernel{nullptr};
void* mDeviceBias{nullptr};
cudnnHandle_t mCudnn;
cublasHandle_t mCublas;
cudnnTensorDescriptor_t mSrcDescriptor, mDstDescriptor;
};

随笔记录 :

FC2Plugin(weight,nbWeight,nbOutputChannels)
验证nbWeight,即需要的weight种类的数量
验证两种weight:mKernelWeights和mBiasWeights的各个属性
为两种weight分配内存
定义mNbInputChannels

FC2Plugin(data length)  //用来处理序列化的engine
	a指针指向d开头的地方
	利用read函数,读取长度为..和..的device上的数据
	重新定义count和value
	利用read函数,读取..长度的数据
	重新定义mBiasWeights
	利用read函数,读取mDataType长度的数据
	利用deserializeToDevice将mDeviceKernel和mDeviceBias指针移动到固定的长度位置处

~FC2Plugin()
	释放mKernelWeights和mBiasWeights的value指针
	
getNbOutputs()
	返回输出的个数
	
getOutputDimensions(index,inputs,nbInputDims)
	检查这三个输入参数的值
	检查mNbOutputChannels的大小
	返回一个应该返回的结构
	
supportsFormat和configureWithFormat是检测format和type的

getSerializationSize是计算mNbInputChannels和mNbOutputChannels和mBiasWeighys....的长度,为并行化做准备.
serialize(buffer)   //随着engine被串行化,这些变量也被串行化,需要转换weight,写到缓冲区中
	同样的设置a和d两个指针
	利用write函数,不断写入数据
	利用convertAndCopyToBuffer来将mKernelWeights和mBiasWeight这两个weight类型的数据结构转换之火存到buffer中(利用memcpy函数)
	最后验证输入数据长度
	
copyToDevice(const void* data,size_t count)
	声明一个deviceData指针
	CHECK+ cudaMalloc和cudaMemcpy
	
convertAndCopyToDevice(void*& deviceWeights, const Weights& weights)
	利用上面介绍的copyToDevice函数实现保存到device中
	
convertAndCopyToBuffer(char*& buffer,const Weights& weights)
	利用memcpy将weight的信息存进去,然后更改buffer指针的位置
	
deserializeToDevice(const char*& hostBuffer,void*& deviceWeights,size_t size)
	并行化的时候同样也利用copyToDevice来实现.

三个火枪手系列:

initialize()
	这里引用了cudnn和cublas,并且创造了描述符.
	当检测mKernelWeights和mBiasWeights的value都有值的时候,将其利用convertAndCopyToDevice函数存到设备上.
	
terminate()
	检测两个描述符
	检测cudnn和cublas是否被销毁
	当mDeviceKernel和mDeviceBias有值的事或,释放他们
	
enqueue(int batchSize,const void* const* inputs ,void** outputs, void* workspace,cudaStream_t stream)
cudaStream_t是对CU_st* 结构的typedef(不开源),这是个工作流?
	float了onef和zerof
	__half了oneh和zeroh
	进行cudnn和cublas流的设置
	先进行了基于cublas的矩阵相乘的运算
	又检测mBiasWeights是否存在,来判断是否进行cudnn的矩阵偏移运算

具体的FC层的实现原理在这里FC层.

这里坑还很多没填,今天只是勉强看懂了,记录一下.具体每个函数的用法以后有时间(会C++)之后,再填.

Nonentity_neu

关注

1
点赞
踩
8

收藏

觉得还不错? 一键收藏
6
评论
TensorRT5.1入门 samplePlugin

FCPlugin类:继承IPluginExt类.第一个构造函数FCPlugin// 用于build,对mKernelWeights和mBiasWeights的check和内存分配(传入的Weight类型的weights存储了这两个值,其中Weight类包含权重的type,value和count)FCPlugin(const Weights weights, int nbWeights, ...
复制链接

扫一扫