基于TensorRT API实现模型网络时用到CUDA核函数时需要注意的一个问题

TensorRT提供了一些常用的基本API，例如2d和3d卷积、池化、上采样、反卷积(转置卷积)、ReLU和全连接等等，但是更新比较慢，对于新的算子和激活函数没有及时提供实现，至于类似不同输出层多种维度特征数据融合这些复杂一点的组合功能是不可能提供现成的API的，所以这些在自己基于TensorRT API实现某个比较新的模型的网络时是需要自己去实现的。

对于这种需要自己定制的部分，TensorRT提供了Plugin机制，也就是提供了几个Plugin接口和超级类分别用于支持输入数据维度是静态的或者动态的，我们自己实现Plugin时只需去集成这些超级类后实现其中必须实现的函数，然后使用时按规定注册，然后就可以在网络中像调用TensorRT的API一样调用来创建定制的layer。Plugin的实现可以参考TensorRT/plugin at master · NVIDIA/TensorRT · GitHub 这里提供了几十个定制实现的算子，包括leakly ReLU、instanceNormalization、groupNormalization等非常常用的网络节点以及nms之类常在后处理中实现的功能(这样做的好处是：如果能把后处理放在网络内部实现，就能作为网络的图(或会话)的一部分一起导出到onnx或者生成到engine中，这样在推理时就不用再去专门做nms之类的后处理了，另外改用GPU计算实现比CPU计算性能好)。

像涉及到多个不同维度的特征数据的输出融合处理，为了性能考虑肯定是需要进行并行计算的，就需要用到CUDA的核函数形式来实现这样的并行运算功能，这样的功能一般放在Plugin参与推理时被调用的enqueue函数里实现:

int enqueue(const nvinfer1::PluginTensorDesc* inputDesc, const nvinfer1::PluginTensorDesc* outputDesc, const void* const* inputs, void* const* outputs, void* workspace, cudaStream_t stream)

CUDA核函数在调用时具有这样的形式：

kernel<<>>(parameters...)

gird参数类型为int或者dim3 (x,y,z)类型，用于指定一个grid中有多少个block

block参数类型为int或者dim3(x,y,z)类型，用于指定一个block中有多少个thread，在GPU中一个thread分配使用一个core执行，一个block的线程都分配在同一个SM(stream multiprocessor)中执行，一个SM中可以分配给一个block或者多个block。

后面的两个参数memory size和stream都有默认值0，前面的用于指定动态内存大小，0表示不需要动态内存，后面的stream是用于指定cuda stream，对于在网络里，这里的stream参数提供了个默认值0就是个坑，当调用核函数时，如果这个参数不指定为全网络统一使用的cuda stream(一般是推理前通过调用cudaStreamCreate()创建的自己的stream)，而是不传入参数让它使用默认的0的话，运行是不会报错的，但是模型推理的精度会下降较多！也就是说一个基于TensorRT API实现的网络，如果网络的不同部分没有使用同一个cuda stream，会导致网络的推理精度下降，我当时的直观的感受是，当全网络使用一个cuda stream时，视频推理的结果视频中bbox框能每帧稳定准确画出来，当网络里有核函数使用了默认的cuda stream 0时，视频推理的结果就变得很差了，经常有帧没检测出目标来，造成结果视频看起来画面一闪一闪的，体验很差！找到问题将使用默认值0的地方全部改为整个网络使用的cuda stream后，结果视频就回复正常的推理效果了，而且这个问题是在Deepstream里跑TensorRT API实现的网络模型时表现明显，如果是使用的自己写的独立的单线程程应用程序调用这种网络模型时推理结果几乎看不出差异来！

创作打卡挑战赛 赢取流量/现金/CSDN周边激励大奖