必须用 ipluginv2 而非旧版 iplugin，因 tensorrt 7.0+ 已彻底废弃 iplugin 接口；ipluginv2 及其子类（如 ipluginv2dynamicext）统一要求插件自主管理内存布局、显式推导输出维度、在 enqueue() 中封装 cuda 执行逻辑，并通过序列化/反序列化传递非 tensor 参数，确保动态 shape 和多 batch 下正确运行。
为什么必须用 IPluginV2 而不是旧版 IPlugin
TensorRT 7.0+ 彻底废弃了 IPlugin 接口，所有新插件必须继承 IPluginV2 或其子类（如 IPluginV2Ext、IPluginV2IOExt）。如果你在代码里还看到 createPlugin 返回 IPlugin*，说明你正踩在兼容性断层上——编译可能通过，但运行时会触发 Invalid plugin type 或直接段错误。
关键区别在于内存管理模型：IPluginV2 要求插件自己负责输入/输出 tensor 的内存布局、尺寸推导和执行调度，不再依赖 TensorRT 内部的隐式 shape 推导逻辑。
-
getOutputDimensions()必须显式返回每个输出的Dims，不能返回Dims{-1}模糊值 -
configurePlugin()中拿到的inputDims是运行时确定的 shape（含 batch 维），不是网络定义时的 symbolic shape -
enqueue()的workspace指针是唯一可申请临时显存的地方，不能在initialize()里调cudMalloc
enqueue() 里怎么安全传参并调用 CUDA kernel
插件的 CUDA 执行逻辑必须封装在 enqueue() 中，且所有参数都要从 const void* const* inputs 和 void* const* outputs 解包。常见错误是把 host 端参数（比如缩放系数 alpha）硬编码进 kernel launch，导致多 batch 或动态 shape 下结果错乱。
正确做法是：在 serialize() 里把所有非 tensor 参数（float alpha、int mode 等）写入序列化 buffer；在 deserializePlugin() 里读出并存在成员变量中；enqueue() 中再传给 kernel。
__global__ void my_custom_kernel(const float* input, float* output, int N, float alpha) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >>(
static_cast<const float>(inputs[0]),
static_cast<float>(outputs[0]),
N,
mAlpha // 来自 deserializePlugin() 存储的成员变量
);
</float></const>
怎么让插件支持 dynamic shape 和 implicit batch mode
默认插件只支持 implicit batch（即 batch 维被剥离，inputDims[0] 对应 HWC）。若需支持 explicit batch（NCHW）或 dynamic shape（如 -1x3x224x224），必须在 supportsFormatCombination() 中显式声明，并在 getOutputDimensions() 中处理 DimensionType::kDYNAMIC。
- 对 dynamic 维，不能直接用
d[i]取值，得用inputDims[0].d[i] == -1 ? -1 : actual_value判断 - 若插件内部有固定 size 的 workspace（比如预分配 1MB buffer），必须在
getWorkspaceSize()中按最大可能 shape 计算，否则executeV2()会因 workspace 不足静默失败 - 支持 dynamic 的插件必须重载
configurePlugin()，不能只靠构造函数初始化
调试 segmentation fault 时优先查哪几个点
90% 的插件段错误发生在 enqueue() 首行或 getOutputDimensions() 返回非法 Dims。不要急着加 printf，先确认：
- 是否在
enqueue()中用了未初始化的成员变量（比如mAlpha没在deserializePlugin()里赋值） -
getOutputDimensions(0, inputDims, nbInputs)是否返回了Dims4(-1, 3, -1, -1)这类非法组合（TensorRT 不允许输出 shape 含多个 -1） - CUDA kernel 中是否越界访问：
input[idx]的idx是否超了N，尤其注意inputDims在 dynamic 模式下可能含 -1 - 是否在
destroy()里重复cudFree同一块 memory（workspace由 TensorRT 管理，插件不能 free）
最有效的调试手段是：在 enqueue() 开头加 cudaError_t err = cudaGetLastError(); assert(err == cudaSuccess);，能立刻暴露 kernel launch 失败点。
C++免费学习笔记（深入）：立即使用
在学习笔记中，你将探索 C++ 的入门与实战技巧！
