From 29498fd8144a490a61d73797981b3c974a3c94d2 Mon Sep 17 00:00:00 2001 From: Rajeev Rao Date: Thu, 9 Jul 2020 23:43:28 -0700 Subject: [PATCH 1/2] Re-enable GridAnchorRect_TRT plugin with rectangular feature maps Signed-off-by: Rajeev Rao --- plugin/InferPlugin.cpp | 1 + plugin/README.md | 1 + plugin/common/kernels/gridAnchorLayer.cu | 228 +++++++++++-------- plugin/gridAnchorPlugin/gridAnchorPlugin.cpp | 65 ++++-- plugin/gridAnchorPlugin/gridAnchorPlugin.h | 31 ++- 5 files changed, 199 insertions(+), 127 deletions(-) diff --git a/plugin/InferPlugin.cpp b/plugin/InferPlugin.cpp index f973bb8e..bca95b38 100644 --- a/plugin/InferPlugin.cpp +++ b/plugin/InferPlugin.cpp @@ -157,6 +157,7 @@ extern "C" { bool initLibNvInferPlugins(void* logger, const char* libNamespace) { initializePlugin(logger, libNamespace); + initializePlugin(logger, libNamespace); initializePlugin(logger, libNamespace); initializePlugin(logger, libNamespace); initializePlugin(logger, libNamespace); diff --git a/plugin/README.md b/plugin/README.md index 5f9f5c4f..8fbcf532 100644 --- a/plugin/README.md +++ b/plugin/README.md @@ -16,6 +16,7 @@ | [geluPlugin](geluPlugin) | CustomGeluPluginDynamic | 1 | | [generateDetectionPlugin](generateDetectionPlugin) | GenerateDetection_TRT | 1 | | [gridAnchorPlugin](gridAnchorPlugin) | GridAnchor_TRT | 1 | +| [gridAnchorRectPlugin](gridAnchorPlugin) | GridAnchorRect_TRT | 1 | | [groupNormalizationPlugin](groupNormalizationPlugin) | GroupNormalizationPlugin | 1 | | [instanceNormalizationPlugin](instanceNormalizationPlugin) | InstanceNormalization_TRT | 1 | | [multilevelCropAndResizePlugin](multilevelCropAndResizePlugin) | MultilevelCropAndResize_TRT | 1 | diff --git a/plugin/common/kernels/gridAnchorLayer.cu b/plugin/common/kernels/gridAnchorLayer.cu index 6e342647..2a195990 100644 --- a/plugin/common/kernels/gridAnchorLayer.cu +++ b/plugin/common/kernels/gridAnchorLayer.cu @@ -19,105 +19,140 @@ using nvinfer1::plugin::reduced_divisor; template - __launch_bounds__(nthdsPerCTA) - __global__ void gridAnchorKernel( - const GridAnchorParameters param, - const int numAspectRatios, - reduced_divisor divObj, - const float* widths, - const float* heights, - float* outputData - ) - { - // output dims: (H, W, param.numMinSize, (1+haveMaxSize+numAR-1), 4) - const int dim = param.H * param.W * numAspectRatios; - /* - * Parameters used to calculate the bounding box coordinates back to input image scale - * Normally we calculate the anchorStride = image_input_size (in pixel) / feature_map_size - * Here we do not use image_input_size for the moment - * Instead we use 1.0 - * The coordinates calculated are scaled by the input image size. - * Most of the coordinates will be in a range of [0, 1], except for the bounding box coordinates going outside of the image - * Every coordinate will go back to the pixel coordinates in the input image if being multiplied by image_input_size - * Here we implicitly assumes the image input and feature map are square - */ - float anchorStride = (1.0 / param.H); - float anchorOffset = 0.5 * anchorStride; - - int tid = blockIdx.x * blockDim.x + threadIdx.x; - if (tid >= dim) - return; - int arId, currIndex; - divObj.divmod(tid, currIndex, arId); - - const int w = currIndex % param.W; - const int h = currIndex / param.W; - - // Center coordinates - float yC = h * anchorStride + anchorOffset; - float xC = w * anchorStride + anchorOffset; - - // x_min, y_min - float xMin = xC - 0.5 * widths[arId]; - float yMin = yC - 0.5 * heights[arId]; - - // x_max, y_max - float xMax = xC + 0.5 * widths[arId]; - float yMax = yC + 0.5 * heights[arId]; - - outputData[tid * 4] = xMin; - outputData[tid * 4 + 1] = yMin; - outputData[tid * 4 + 2] = xMax; - outputData[tid * 4 + 3] = yMax; - - // Remember to move the output cursor - float* output = outputData + dim * 4; - - // Simply copying the variance - output[tid * 4] = param.variance[0]; - output[tid * 4 + 1] = param.variance[1]; - output[tid * 4 + 2] = param.variance[2]; - output[tid * 4 + 3] = param.variance[3]; - - } - -pluginStatus_t anchorGridInference( - cudaStream_t stream, - const GridAnchorParameters param, - const int numAspectRatios, - const void* widths, - const void* heights, - void* outputData - ) +__launch_bounds__(nthdsPerCTA) +__global__ void gridAnchorKernel( + const GridAnchorParameters param, + const int numAspectRatios, + reduced_divisor divObj, + const float* widths, + const float* heights, + float* outputData + ) { + // output dims: (H, W, param.numMinSize, (1+haveMaxSize+numAR-1), 4) const int dim = param.H * param.W * numAspectRatios; - reduced_divisor divObj(numAspectRatios); - if (dim > 5120) - { - const int BS = 128; - const int GS = (dim + BS - 1) / BS; - gridAnchorKernel<<>>(param, numAspectRatios, divObj, - (const float*) widths, (const float*) heights, - (float*) outputData); - } - else - { - const int BS = 32; - const int GS = (dim + BS - 1) / BS; - gridAnchorKernel<<>>(param, numAspectRatios, divObj, - (const float*) widths, (const float*) heights, - (float*) outputData); + /* + * Parameters used to calculate the bounding box coordinates back to input image scale + * Normally we calculate the anchorStride = image_input_size (in pixel) / feature_map_size + * Here we do not use image_input_size for the moment + * Instead we use 1.0 + * The coordinates calculated are scaled by the input image size. + * Most of the coordinates will be in a range of [0, 1], except for the bounding box coordinates going outside of the image + * Every coordinate will go back to the pixel coordinates in the input image if being multiplied by image_input_size + * Here we implicitly assumes the image input and feature map are square + */ + float anchorStride = (1.0 / param.H); + float anchorOffset = 0.5 * anchorStride; + int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid >= dim) + { + return; } - CSC(cudaGetLastError(), STATUS_FAILURE); - return STATUS_SUCCESS; + + int arId, currIndex; + divObj.divmod(tid, currIndex, arId); + + const int w = currIndex % param.W; + const int h = currIndex / param.W; + + // Center coordinates + float yC = h * anchorStride + anchorOffset; + float xC = w * anchorStride + anchorOffset; + + // x_min, y_min + float xMin = xC - 0.5 * widths[arId]; + float yMin = yC - 0.5 * heights[arId]; + + // x_max, y_max + float xMax = xC + 0.5 * widths[arId]; + float yMax = yC + 0.5 * heights[arId]; + + outputData[tid * 4] = xMin; + outputData[tid * 4 + 1] = yMin; + outputData[tid * 4 + 2] = xMax; + outputData[tid * 4 + 3] = yMax; + + // Remember to move the output cursor + float* output = outputData + dim * 4; + + // Simply copying the variance + output[tid * 4] = param.variance[0]; + output[tid * 4 + 1] = param.variance[1]; + output[tid * 4 + 2] = param.variance[2]; + output[tid * 4 + 3] = param.variance[3]; } -namespace nvinfer1 -{ -namespace plugin +using nvinfer1::plugin::reduced_divisor; +template +__launch_bounds__(nthdsPerCTA) +__global__ void gridAnchorRectKernel( + const GridAnchorParameters param, + const int numAspectRatios, + reduced_divisor divObj, + const float* widths, + const float* heights, + float* outputData + ) { + // output dims: (H, W, param.numMinSize, (1+haveMaxSize+numAR-1), 4) + const int dim = param.H * param.W * numAspectRatios; + + /* + * Parameters used to calculate the bounding box coordinates back to input image scale + * Normally we calculate the anchorStride = image_input_size (in pixel) / feature_map_size + * Here we do not use image_input_size for the moment + * Instead we use 1.0 + * The coordinates calculated are scaled by the input image size. + * Most of the coordinates will be in a range of [0, 1], except for the bounding box coordinates going outside of the image + * Every coordinate will go back to the pixel coordinates in the input image if being multiplied by image_input_size + * Here we implicitly assumes the image input and feature map are square + */ + float anchorStrideH = (1.0 / param.H); + float anchorStrideW = (1.0 / param.W); + float anchorOffsetH = 0.5 * anchorStrideH; + float anchorOffsetW = 0.5 * anchorStrideW; + + int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid >= dim) + { + return; + } + + int arId, currIndex; + divObj.divmod(tid, currIndex, arId); + + const int w = currIndex % param.W; + const int h = currIndex / param.W; + + // Center coordinates + float yC = h * anchorStrideH + anchorOffsetH; + float xC = w * anchorStrideW + anchorOffsetW; + + // x_min, y_min + float xMin = xC - 0.5 * widths[arId]; + float yMin = yC - 0.5 * heights[arId]; + + // x_max, y_max + float xMax = xC + 0.5 * widths[arId]; + float yMax = yC + 0.5 * heights[arId]; + + outputData[tid * 4] = xMin; + outputData[tid * 4 + 1] = yMin; + outputData[tid * 4 + 2] = xMax; + outputData[tid * 4 + 3] = yMax; + + // Remember to move the output cursor + float* output = outputData + dim * 4; + + // Simply copying the variance + output[tid * 4] = param.variance[0]; + output[tid * 4 + 1] = param.variance[1]; + output[tid * 4 + 2] = param.variance[2]; + output[tid * 4 + 3] = param.variance[3]; +} + pluginStatus_t anchorGridInference( cudaStream_t stream, const GridAnchorParameters param, @@ -134,8 +169,8 @@ pluginStatus_t anchorGridInference( const int BS = 128; const int GS = (dim + BS - 1) / BS; gridAnchorKernel<<>>(param, numAspectRatios, divObj, - (const float*) widths, (const float*) heights, - (float*) outputData); + (const float*) widths, (const float*) heights, + (float*) outputData); } else @@ -143,13 +178,10 @@ pluginStatus_t anchorGridInference( const int BS = 32; const int GS = (dim + BS - 1) / BS; gridAnchorKernel<<>>(param, numAspectRatios, divObj, - (const float*) widths, (const float*) heights, - (float*) outputData); - + (const float*) widths, (const float*) heights, + (float*) outputData); } CSC(cudaGetLastError(), STATUS_FAILURE); return STATUS_SUCCESS; } -} -} diff --git a/plugin/gridAnchorPlugin/gridAnchorPlugin.cpp b/plugin/gridAnchorPlugin/gridAnchorPlugin.cpp index 635dbbd4..537d714e 100644 --- a/plugin/gridAnchorPlugin/gridAnchorPlugin.cpp +++ b/plugin/gridAnchorPlugin/gridAnchorPlugin.cpp @@ -22,19 +22,18 @@ #include using namespace nvinfer1; -using nvinfer1::plugin::GridAnchorGenerator; -using nvinfer1::plugin::GridAnchorPluginCreator; namespace { -const char* GRID_ANCHOR_PLUGIN_VERSION{"1"}; -const char* GRID_ANCHOR_PLUGIN_NAME{"GridAnchor_TRT"}; +std::string GRID_ANCHOR_PLUGIN_NAMES[] = {"GridAnchor_TRT", "GridAnchorRect_TRT"}; +const char* GRID_ANCHOR_PLUGIN_VERSION = "1"; } // namespace -PluginFieldCollection GridAnchorPluginCreator::mFC{}; -std::vector GridAnchorPluginCreator::mPluginAttributes; -GridAnchorGenerator::GridAnchorGenerator(const GridAnchorParameters* paramIn, int mNumLayers) - : mNumLayers(mNumLayers) +PluginFieldCollection GridAnchorBasePluginCreator::mFC{}; +std::vector GridAnchorBasePluginCreator::mPluginAttributes; + +GridAnchorGenerator::GridAnchorGenerator(const GridAnchorParameters* paramIn, int numLayers, const char *name) + : mNumLayers(numLayers), mPluginName(name) { CUASSERT(cudaMallocHost((void**) &mNumPriors, mNumLayers * sizeof(int))); CUASSERT(cudaMallocHost((void**) &mDeviceWidths, mNumLayers * sizeof(Weights))); @@ -121,7 +120,8 @@ GridAnchorGenerator::GridAnchorGenerator(const GridAnchorParameters* paramIn, in } } -GridAnchorGenerator::GridAnchorGenerator(const void* data, size_t length) +GridAnchorGenerator::GridAnchorGenerator(const void* data, size_t length, const char *name) : + mPluginName(name) { const char *d = reinterpret_cast(data), *a = d; mNumLayers = read(d); @@ -276,7 +276,7 @@ bool GridAnchorGenerator::supportsFormat(DataType type, PluginFormat format) con const char* GridAnchorGenerator::getPluginType() const { - return GRID_ANCHOR_PLUGIN_NAME; + return mPluginName.c_str(); } const char* GridAnchorGenerator::getPluginVersion() const @@ -284,6 +284,7 @@ const char* GridAnchorGenerator::getPluginVersion() const return GRID_ANCHOR_PLUGIN_VERSION; } + // Set plugin namespace void GridAnchorGenerator::setPluginNamespace(const char* pluginNamespace) { @@ -341,12 +342,12 @@ void GridAnchorGenerator::destroy() IPluginV2Ext* GridAnchorGenerator::clone() const { - IPluginV2Ext* plugin = new GridAnchorGenerator(mParam.data(), mNumLayers); + IPluginV2Ext* plugin = new GridAnchorGenerator(mParam.data(), mNumLayers, mPluginName.c_str()); plugin->setPluginNamespace(mPluginNamespace.c_str()); return plugin; } -GridAnchorPluginCreator::GridAnchorPluginCreator() +GridAnchorBasePluginCreator::GridAnchorBasePluginCreator() { mPluginAttributes.emplace_back(PluginField("minSize", nullptr, PluginFieldType::kFLOAT32, 1)); mPluginAttributes.emplace_back(PluginField("maxSize", nullptr, PluginFieldType::kFLOAT32, 1)); @@ -359,22 +360,22 @@ GridAnchorPluginCreator::GridAnchorPluginCreator() mFC.fields = mPluginAttributes.data(); } -const char* GridAnchorPluginCreator::getPluginName() const +const char* GridAnchorBasePluginCreator::getPluginName() const { - return GRID_ANCHOR_PLUGIN_NAME; + return mPluginName.c_str(); } -const char* GridAnchorPluginCreator::getPluginVersion() const +const char* GridAnchorBasePluginCreator::getPluginVersion() const { return GRID_ANCHOR_PLUGIN_VERSION; } -const PluginFieldCollection* GridAnchorPluginCreator::getFieldNames() +const PluginFieldCollection* GridAnchorBasePluginCreator::getFieldNames() { return &mFC; } -IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const PluginFieldCollection* fc) +IPluginV2Ext* GridAnchorBasePluginCreator::createPlugin(const char* name, const PluginFieldCollection* fc) { float minScale = 0.2F, maxScale = 0.95F; int numLayers = 6; @@ -382,6 +383,8 @@ IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const Plug std::vector fMapShapes; std::vector layerVariances; const PluginField* fields = fc->fields; + + const bool isFMapRect = (GRID_ANCHOR_PLUGIN_NAMES[1] == mPluginName); for (int i = 0; i < fc->nbFields; ++i) { const char* attrName = fields[i].name; @@ -428,6 +431,7 @@ IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const Plug { ASSERT(fields[i].type == PluginFieldType::kINT32); int size = fields[i].length; + ASSERT(!isFMapRect || (size % 2 == 0)); fMapShapes.reserve(size); const int* fMap = static_cast(fields[i].data); for (int j = 0; j < size; j++) @@ -442,7 +446,8 @@ IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const Plug std::vector firstLayerAspectRatios; ASSERT(numLayers > 0); - ASSERT((int) fMapShapes.size() == numLayers); + const int numExpectedLayers = static_cast(fMapShapes.size()) >> (isFMapRect ? 1 : 0); + ASSERT(numExpectedLayers == numLayers); int numFirstLayerARs = 3; // First layer only has the first 3 aspect ratios from aspectRatios @@ -457,30 +462,42 @@ IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const Plug // One set of box parameters for one layer for (int i = 0; i < numLayers; i++) { + int hOffset = (isFMapRect ? i * 2 : i); + int wOffset = (isFMapRect ? i * 2 + 1 : i); // Only the first layer is different if (i == 0) { boxParams[i] = {minScale, maxScale, firstLayerAspectRatios.data(), (int) firstLayerAspectRatios.size(), - fMapShapes[i], fMapShapes[i], + fMapShapes[hOffset], fMapShapes[wOffset], {layerVariances[0], layerVariances[1], layerVariances[2], layerVariances[3]}}; } else { - boxParams[i] = {minScale, maxScale, aspectRatios.data(), (int) aspectRatios.size(), fMapShapes[i], - fMapShapes[i], {layerVariances[0], layerVariances[1], layerVariances[2], layerVariances[3]}}; + boxParams[i] = {minScale, maxScale, aspectRatios.data(), (int) aspectRatios.size(), fMapShapes[hOffset], + fMapShapes[wOffset], {layerVariances[0], layerVariances[1], layerVariances[2], layerVariances[3]}}; } } - GridAnchorGenerator* obj = new GridAnchorGenerator(boxParams.data(), numLayers); + GridAnchorGenerator* obj = new GridAnchorGenerator(boxParams.data(), numLayers, mPluginName.c_str()); obj->setPluginNamespace(mNamespace.c_str()); return obj; } -IPluginV2Ext* GridAnchorPluginCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength) +IPluginV2Ext* GridAnchorBasePluginCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength) { // This object will be deleted when the network is destroyed, which will // call GridAnchor::destroy() - GridAnchorGenerator* obj = new GridAnchorGenerator(serialData, serialLength); + GridAnchorGenerator* obj = new GridAnchorGenerator(serialData, serialLength, mPluginName.c_str()); obj->setPluginNamespace(mNamespace.c_str()); return obj; } + +GridAnchorPluginCreator::GridAnchorPluginCreator() +{ + mPluginName = GRID_ANCHOR_PLUGIN_NAMES[0]; +} + +GridAnchorRectPluginCreator::GridAnchorRectPluginCreator() +{ + mPluginName = GRID_ANCHOR_PLUGIN_NAMES[1]; +} diff --git a/plugin/gridAnchorPlugin/gridAnchorPlugin.h b/plugin/gridAnchorPlugin/gridAnchorPlugin.h index 4ffd903d..1a88b6f9 100644 --- a/plugin/gridAnchorPlugin/gridAnchorPlugin.h +++ b/plugin/gridAnchorPlugin/gridAnchorPlugin.h @@ -29,9 +29,9 @@ namespace plugin class GridAnchorGenerator : public IPluginV2Ext { public: - GridAnchorGenerator(const GridAnchorParameters* param, int numLayers); + GridAnchorGenerator(const GridAnchorParameters* param, int numLayers, const char *version); - GridAnchorGenerator(const void* data, size_t length); + GridAnchorGenerator(const void* data, size_t length, const char *version); ~GridAnchorGenerator() override; @@ -81,6 +81,9 @@ class GridAnchorGenerator : public IPluginV2Ext void detachFromContext() override; +protected: + std::string mPluginName; + private: Weights copyToDevice(const void* hostData, size_t count); @@ -95,12 +98,12 @@ class GridAnchorGenerator : public IPluginV2Ext std::string mPluginNamespace; }; -class GridAnchorPluginCreator : public BaseCreator +class GridAnchorBasePluginCreator : public BaseCreator { public: - GridAnchorPluginCreator(); + GridAnchorBasePluginCreator(); - ~GridAnchorPluginCreator() override = default; + ~GridAnchorBasePluginCreator() override = default; const char* getPluginName() const override; @@ -112,10 +115,28 @@ class GridAnchorPluginCreator : public BaseCreator IPluginV2Ext* deserializePlugin(const char* name, const void* serialData, size_t serialLength) override; +protected: + std::string mPluginName; + private: static PluginFieldCollection mFC; static std::vector mPluginAttributes; }; + +class GridAnchorPluginCreator : public GridAnchorBasePluginCreator +{ +public: + GridAnchorPluginCreator(); + ~GridAnchorPluginCreator() override = default; +}; + +class GridAnchorRectPluginCreator : public GridAnchorBasePluginCreator +{ +public: + GridAnchorRectPluginCreator(); + ~GridAnchorRectPluginCreator() override = default; +}; + } // namespace plugin } // namespace nvinfer1 From c9b4e4dbc127e892988d3937992d07d781ea1c92 Mon Sep 17 00:00:00 2001 From: Rajeev Rao Date: Sat, 11 Jul 2020 02:50:09 -0700 Subject: [PATCH 2/2] Review feedback #1 - remove redundant kernel --- plugin/common/kernels/gridAnchorLayer.cu | 67 ----------------------- samples/opensource/sampleUffSSD/README.md | 2 +- 2 files changed, 1 insertion(+), 68 deletions(-) diff --git a/plugin/common/kernels/gridAnchorLayer.cu b/plugin/common/kernels/gridAnchorLayer.cu index 2a195990..8685d8b4 100644 --- a/plugin/common/kernels/gridAnchorLayer.cu +++ b/plugin/common/kernels/gridAnchorLayer.cu @@ -32,73 +32,6 @@ __global__ void gridAnchorKernel( // output dims: (H, W, param.numMinSize, (1+haveMaxSize+numAR-1), 4) const int dim = param.H * param.W * numAspectRatios; - /* - * Parameters used to calculate the bounding box coordinates back to input image scale - * Normally we calculate the anchorStride = image_input_size (in pixel) / feature_map_size - * Here we do not use image_input_size for the moment - * Instead we use 1.0 - * The coordinates calculated are scaled by the input image size. - * Most of the coordinates will be in a range of [0, 1], except for the bounding box coordinates going outside of the image - * Every coordinate will go back to the pixel coordinates in the input image if being multiplied by image_input_size - * Here we implicitly assumes the image input and feature map are square - */ - float anchorStride = (1.0 / param.H); - float anchorOffset = 0.5 * anchorStride; - - int tid = blockIdx.x * blockDim.x + threadIdx.x; - if (tid >= dim) - { - return; - } - - int arId, currIndex; - divObj.divmod(tid, currIndex, arId); - - const int w = currIndex % param.W; - const int h = currIndex / param.W; - - // Center coordinates - float yC = h * anchorStride + anchorOffset; - float xC = w * anchorStride + anchorOffset; - - // x_min, y_min - float xMin = xC - 0.5 * widths[arId]; - float yMin = yC - 0.5 * heights[arId]; - - // x_max, y_max - float xMax = xC + 0.5 * widths[arId]; - float yMax = yC + 0.5 * heights[arId]; - - outputData[tid * 4] = xMin; - outputData[tid * 4 + 1] = yMin; - outputData[tid * 4 + 2] = xMax; - outputData[tid * 4 + 3] = yMax; - - // Remember to move the output cursor - float* output = outputData + dim * 4; - - // Simply copying the variance - output[tid * 4] = param.variance[0]; - output[tid * 4 + 1] = param.variance[1]; - output[tid * 4 + 2] = param.variance[2]; - output[tid * 4 + 3] = param.variance[3]; -} - -using nvinfer1::plugin::reduced_divisor; -template -__launch_bounds__(nthdsPerCTA) -__global__ void gridAnchorRectKernel( - const GridAnchorParameters param, - const int numAspectRatios, - reduced_divisor divObj, - const float* widths, - const float* heights, - float* outputData - ) -{ - // output dims: (H, W, param.numMinSize, (1+haveMaxSize+numAR-1), 4) - const int dim = param.H * param.W * numAspectRatios; - /* * Parameters used to calculate the bounding box coordinates back to input image scale * Normally we calculate the anchorStride = image_input_size (in pixel) / feature_map_size diff --git a/samples/opensource/sampleUffSSD/README.md b/samples/opensource/sampleUffSSD/README.md index dffe738f..b773805e 100644 --- a/samples/opensource/sampleUffSSD/README.md +++ b/samples/opensource/sampleUffSSD/README.md @@ -217,7 +217,7 @@ The Shuffle layer implements a reshape and transpose operator for tensors. 2. Run the following command for the conversion. `convert-to-uff frozen_inference_graph.pb -O NMS -p config.py` - This saves the converted `.uff` file in the same directory as the input with the name `frozen_inference_graph.pb.uff`. + This saves the converted `.uff` file in the same directory as the input with the name `frozen_inference_graph.uff`. The `config.py` script specifies the preprocessing operations necessary for the SSD TensorFlow graph. The plugin nodes and plugin parameters used in the `config.py` script should match the registered plugins in TensorRT.