环境: 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=0FC层只有一个输出,可能像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++)之后,再填.