Skip to content

Commit ead9e8b

Browse files
authored
Merge pull request #679 from rajeevsrao/dev/gridAnchorRect
Re-enable GridAnchorRect_TRT plugin with rectangular feature maps
2 parents 9ee8450 + c9b4e4d commit ead9e8b

File tree

6 files changed

+129
-124
lines changed

6 files changed

+129
-124
lines changed

Diff for: plugin/InferPlugin.cpp

+1
Original file line numberDiff line numberDiff line change
@@ -157,6 +157,7 @@ extern "C" {
157157
bool initLibNvInferPlugins(void* logger, const char* libNamespace)
158158
{
159159
initializePlugin<nvinfer1::plugin::GridAnchorPluginCreator>(logger, libNamespace);
160+
initializePlugin<nvinfer1::plugin::GridAnchorRectPluginCreator>(logger, libNamespace);
160161
initializePlugin<nvinfer1::plugin::NMSPluginCreator>(logger, libNamespace);
161162
initializePlugin<nvinfer1::plugin::ReorgPluginCreator>(logger, libNamespace);
162163
initializePlugin<nvinfer1::plugin::RegionPluginCreator>(logger, libNamespace);

Diff for: plugin/README.md

+1
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616
| [geluPlugin](geluPlugin) | CustomGeluPluginDynamic | 1 |
1717
| [generateDetectionPlugin](generateDetectionPlugin) | GenerateDetection_TRT | 1 |
1818
| [gridAnchorPlugin](gridAnchorPlugin) | GridAnchor_TRT | 1 |
19+
| [gridAnchorRectPlugin](gridAnchorPlugin) | GridAnchorRect_TRT | 1 |
1920
| [groupNormalizationPlugin](groupNormalizationPlugin) | GroupNormalizationPlugin | 1 |
2021
| [instanceNormalizationPlugin](instanceNormalizationPlugin) | InstanceNormalization_TRT | 1 |
2122
| [multilevelCropAndResizePlugin](multilevelCropAndResizePlugin) | MultilevelCropAndResize_TRT | 1 |

Diff for: plugin/common/kernels/gridAnchorLayer.cu

+59-94
Original file line numberDiff line numberDiff line change
@@ -19,105 +19,73 @@
1919

2020
using nvinfer1::plugin::reduced_divisor;
2121
template <unsigned nthdsPerCTA>
22-
__launch_bounds__(nthdsPerCTA)
23-
__global__ void gridAnchorKernel(
24-
const GridAnchorParameters param,
25-
const int numAspectRatios,
26-
reduced_divisor divObj,
27-
const float* widths,
28-
const float* heights,
29-
float* outputData
30-
)
31-
{
32-
// output dims: (H, W, param.numMinSize, (1+haveMaxSize+numAR-1), 4)
33-
const int dim = param.H * param.W * numAspectRatios;
34-
/*
35-
* Parameters used to calculate the bounding box coordinates back to input image scale
36-
* Normally we calculate the anchorStride = image_input_size (in pixel) / feature_map_size
37-
* Here we do not use image_input_size for the moment
38-
* Instead we use 1.0
39-
* The coordinates calculated are scaled by the input image size.
40-
* Most of the coordinates will be in a range of [0, 1], except for the bounding box coordinates going outside of the image
41-
* Every coordinate will go back to the pixel coordinates in the input image if being multiplied by image_input_size
42-
* Here we implicitly assumes the image input and feature map are square
43-
*/
44-
float anchorStride = (1.0 / param.H);
45-
float anchorOffset = 0.5 * anchorStride;
22+
__launch_bounds__(nthdsPerCTA)
23+
__global__ void gridAnchorKernel(
24+
const GridAnchorParameters param,
25+
const int numAspectRatios,
26+
reduced_divisor divObj,
27+
const float* widths,
28+
const float* heights,
29+
float* outputData
30+
)
31+
{
32+
// output dims: (H, W, param.numMinSize, (1+haveMaxSize+numAR-1), 4)
33+
const int dim = param.H * param.W * numAspectRatios;
4634

47-
int tid = blockIdx.x * blockDim.x + threadIdx.x;
48-
if (tid >= dim)
49-
return;
50-
int arId, currIndex;
51-
divObj.divmod(tid, currIndex, arId);
35+
/*
36+
* Parameters used to calculate the bounding box coordinates back to input image scale
37+
* Normally we calculate the anchorStride = image_input_size (in pixel) / feature_map_size
38+
* Here we do not use image_input_size for the moment
39+
* Instead we use 1.0
40+
* The coordinates calculated are scaled by the input image size.
41+
* Most of the coordinates will be in a range of [0, 1], except for the bounding box coordinates going outside of the image
42+
* Every coordinate will go back to the pixel coordinates in the input image if being multiplied by image_input_size
43+
* Here we implicitly assumes the image input and feature map are square
44+
*/
45+
float anchorStrideH = (1.0 / param.H);
46+
float anchorStrideW = (1.0 / param.W);
47+
float anchorOffsetH = 0.5 * anchorStrideH;
48+
float anchorOffsetW = 0.5 * anchorStrideW;
5249

53-
const int w = currIndex % param.W;
54-
const int h = currIndex / param.W;
55-
56-
// Center coordinates
57-
float yC = h * anchorStride + anchorOffset;
58-
float xC = w * anchorStride + anchorOffset;
59-
60-
// x_min, y_min
61-
float xMin = xC - 0.5 * widths[arId];
62-
float yMin = yC - 0.5 * heights[arId];
50+
int tid = blockIdx.x * blockDim.x + threadIdx.x;
51+
if (tid >= dim)
52+
{
53+
return;
54+
}
6355

64-
// x_max, y_max
65-
float xMax = xC + 0.5 * widths[arId];
66-
float yMax = yC + 0.5 * heights[arId];
56+
int arId, currIndex;
57+
divObj.divmod(tid, currIndex, arId);
6758

68-
outputData[tid * 4] = xMin;
69-
outputData[tid * 4 + 1] = yMin;
70-
outputData[tid * 4 + 2] = xMax;
71-
outputData[tid * 4 + 3] = yMax;
59+
const int w = currIndex % param.W;
60+
const int h = currIndex / param.W;
7261

73-
// Remember to move the output cursor
74-
float* output = outputData + dim * 4;
62+
// Center coordinates
63+
float yC = h * anchorStrideH + anchorOffsetH;
64+
float xC = w * anchorStrideW + anchorOffsetW;
7565

76-
// Simply copying the variance
77-
output[tid * 4] = param.variance[0];
78-
output[tid * 4 + 1] = param.variance[1];
79-
output[tid * 4 + 2] = param.variance[2];
80-
output[tid * 4 + 3] = param.variance[3];
66+
// x_min, y_min
67+
float xMin = xC - 0.5 * widths[arId];
68+
float yMin = yC - 0.5 * heights[arId];
8169

82-
}
70+
// x_max, y_max
71+
float xMax = xC + 0.5 * widths[arId];
72+
float yMax = yC + 0.5 * heights[arId];
8373

84-
pluginStatus_t anchorGridInference(
85-
cudaStream_t stream,
86-
const GridAnchorParameters param,
87-
const int numAspectRatios,
88-
const void* widths,
89-
const void* heights,
90-
void* outputData
91-
)
92-
{
93-
const int dim = param.H * param.W * numAspectRatios;
94-
reduced_divisor divObj(numAspectRatios);
95-
if (dim > 5120)
96-
{
97-
const int BS = 128;
98-
const int GS = (dim + BS - 1) / BS;
99-
gridAnchorKernel<BS><<<GS, BS, 0, stream>>>(param, numAspectRatios, divObj,
100-
(const float*) widths, (const float*) heights,
101-
(float*) outputData);
74+
outputData[tid * 4] = xMin;
75+
outputData[tid * 4 + 1] = yMin;
76+
outputData[tid * 4 + 2] = xMax;
77+
outputData[tid * 4 + 3] = yMax;
10278

103-
}
104-
else
105-
{
106-
const int BS = 32;
107-
const int GS = (dim + BS - 1) / BS;
108-
gridAnchorKernel<BS><<<GS, BS, 0, stream>>>(param, numAspectRatios, divObj,
109-
(const float*) widths, (const float*) heights,
110-
(float*) outputData);
79+
// Remember to move the output cursor
80+
float* output = outputData + dim * 4;
11181

112-
}
113-
CSC(cudaGetLastError(), STATUS_FAILURE);
114-
return STATUS_SUCCESS;
82+
// Simply copying the variance
83+
output[tid * 4] = param.variance[0];
84+
output[tid * 4 + 1] = param.variance[1];
85+
output[tid * 4 + 2] = param.variance[2];
86+
output[tid * 4 + 3] = param.variance[3];
11587
}
11688

117-
namespace nvinfer1
118-
{
119-
namespace plugin
120-
{
12189
pluginStatus_t anchorGridInference(
12290
cudaStream_t stream,
12391
const GridAnchorParameters param,
@@ -134,22 +102,19 @@ pluginStatus_t anchorGridInference(
134102
const int BS = 128;
135103
const int GS = (dim + BS - 1) / BS;
136104
gridAnchorKernel<BS><<<GS, BS, 0, stream>>>(param, numAspectRatios, divObj,
137-
(const float*) widths, (const float*) heights,
138-
(float*) outputData);
105+
(const float*) widths, (const float*) heights,
106+
(float*) outputData);
139107

140108
}
141109
else
142110
{
143111
const int BS = 32;
144112
const int GS = (dim + BS - 1) / BS;
145113
gridAnchorKernel<BS><<<GS, BS, 0, stream>>>(param, numAspectRatios, divObj,
146-
(const float*) widths, (const float*) heights,
147-
(float*) outputData);
148-
114+
(const float*) widths, (const float*) heights,
115+
(float*) outputData);
149116
}
150117
CSC(cudaGetLastError(), STATUS_FAILURE);
151118
return STATUS_SUCCESS;
152119
}
153120

154-
}
155-
}

Diff for: plugin/gridAnchorPlugin/gridAnchorPlugin.cpp

+41-24
Original file line numberDiff line numberDiff line change
@@ -22,19 +22,18 @@
2222
#include <vector>
2323

2424
using namespace nvinfer1;
25-
using nvinfer1::plugin::GridAnchorGenerator;
26-
using nvinfer1::plugin::GridAnchorPluginCreator;
2725

2826
namespace
2927
{
30-
const char* GRID_ANCHOR_PLUGIN_VERSION{"1"};
31-
const char* GRID_ANCHOR_PLUGIN_NAME{"GridAnchor_TRT"};
28+
std::string GRID_ANCHOR_PLUGIN_NAMES[] = {"GridAnchor_TRT", "GridAnchorRect_TRT"};
29+
const char* GRID_ANCHOR_PLUGIN_VERSION = "1";
3230
} // namespace
33-
PluginFieldCollection GridAnchorPluginCreator::mFC{};
34-
std::vector<PluginField> GridAnchorPluginCreator::mPluginAttributes;
3531

36-
GridAnchorGenerator::GridAnchorGenerator(const GridAnchorParameters* paramIn, int mNumLayers)
37-
: mNumLayers(mNumLayers)
32+
PluginFieldCollection GridAnchorBasePluginCreator::mFC{};
33+
std::vector<PluginField> GridAnchorBasePluginCreator::mPluginAttributes;
34+
35+
GridAnchorGenerator::GridAnchorGenerator(const GridAnchorParameters* paramIn, int numLayers, const char *name)
36+
: mNumLayers(numLayers), mPluginName(name)
3837
{
3938
CUASSERT(cudaMallocHost((void**) &mNumPriors, mNumLayers * sizeof(int)));
4039
CUASSERT(cudaMallocHost((void**) &mDeviceWidths, mNumLayers * sizeof(Weights)));
@@ -121,7 +120,8 @@ GridAnchorGenerator::GridAnchorGenerator(const GridAnchorParameters* paramIn, in
121120
}
122121
}
123122

124-
GridAnchorGenerator::GridAnchorGenerator(const void* data, size_t length)
123+
GridAnchorGenerator::GridAnchorGenerator(const void* data, size_t length, const char *name) :
124+
mPluginName(name)
125125
{
126126
const char *d = reinterpret_cast<const char*>(data), *a = d;
127127
mNumLayers = read<int>(d);
@@ -276,14 +276,15 @@ bool GridAnchorGenerator::supportsFormat(DataType type, PluginFormat format) con
276276

277277
const char* GridAnchorGenerator::getPluginType() const
278278
{
279-
return GRID_ANCHOR_PLUGIN_NAME;
279+
return mPluginName.c_str();
280280
}
281281

282282
const char* GridAnchorGenerator::getPluginVersion() const
283283
{
284284
return GRID_ANCHOR_PLUGIN_VERSION;
285285
}
286286

287+
287288
// Set plugin namespace
288289
void GridAnchorGenerator::setPluginNamespace(const char* pluginNamespace)
289290
{
@@ -341,12 +342,12 @@ void GridAnchorGenerator::destroy()
341342

342343
IPluginV2Ext* GridAnchorGenerator::clone() const
343344
{
344-
IPluginV2Ext* plugin = new GridAnchorGenerator(mParam.data(), mNumLayers);
345+
IPluginV2Ext* plugin = new GridAnchorGenerator(mParam.data(), mNumLayers, mPluginName.c_str());
345346
plugin->setPluginNamespace(mPluginNamespace.c_str());
346347
return plugin;
347348
}
348349

349-
GridAnchorPluginCreator::GridAnchorPluginCreator()
350+
GridAnchorBasePluginCreator::GridAnchorBasePluginCreator()
350351
{
351352
mPluginAttributes.emplace_back(PluginField("minSize", nullptr, PluginFieldType::kFLOAT32, 1));
352353
mPluginAttributes.emplace_back(PluginField("maxSize", nullptr, PluginFieldType::kFLOAT32, 1));
@@ -359,29 +360,31 @@ GridAnchorPluginCreator::GridAnchorPluginCreator()
359360
mFC.fields = mPluginAttributes.data();
360361
}
361362

362-
const char* GridAnchorPluginCreator::getPluginName() const
363+
const char* GridAnchorBasePluginCreator::getPluginName() const
363364
{
364-
return GRID_ANCHOR_PLUGIN_NAME;
365+
return mPluginName.c_str();
365366
}
366367

367-
const char* GridAnchorPluginCreator::getPluginVersion() const
368+
const char* GridAnchorBasePluginCreator::getPluginVersion() const
368369
{
369370
return GRID_ANCHOR_PLUGIN_VERSION;
370371
}
371372

372-
const PluginFieldCollection* GridAnchorPluginCreator::getFieldNames()
373+
const PluginFieldCollection* GridAnchorBasePluginCreator::getFieldNames()
373374
{
374375
return &mFC;
375376
}
376377

377-
IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const PluginFieldCollection* fc)
378+
IPluginV2Ext* GridAnchorBasePluginCreator::createPlugin(const char* name, const PluginFieldCollection* fc)
378379
{
379380
float minScale = 0.2F, maxScale = 0.95F;
380381
int numLayers = 6;
381382
std::vector<float> aspectRatios;
382383
std::vector<int> fMapShapes;
383384
std::vector<float> layerVariances;
384385
const PluginField* fields = fc->fields;
386+
387+
const bool isFMapRect = (GRID_ANCHOR_PLUGIN_NAMES[1] == mPluginName);
385388
for (int i = 0; i < fc->nbFields; ++i)
386389
{
387390
const char* attrName = fields[i].name;
@@ -428,6 +431,7 @@ IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const Plug
428431
{
429432
ASSERT(fields[i].type == PluginFieldType::kINT32);
430433
int size = fields[i].length;
434+
ASSERT(!isFMapRect || (size % 2 == 0));
431435
fMapShapes.reserve(size);
432436
const int* fMap = static_cast<const int*>(fields[i].data);
433437
for (int j = 0; j < size; j++)
@@ -442,7 +446,8 @@ IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const Plug
442446
std::vector<float> firstLayerAspectRatios;
443447

444448
ASSERT(numLayers > 0);
445-
ASSERT((int) fMapShapes.size() == numLayers);
449+
const int numExpectedLayers = static_cast<int>(fMapShapes.size()) >> (isFMapRect ? 1 : 0);
450+
ASSERT(numExpectedLayers == numLayers);
446451

447452
int numFirstLayerARs = 3;
448453
// First layer only has the first 3 aspect ratios from aspectRatios
@@ -457,30 +462,42 @@ IPluginV2Ext* GridAnchorPluginCreator::createPlugin(const char* name, const Plug
457462
// One set of box parameters for one layer
458463
for (int i = 0; i < numLayers; i++)
459464
{
465+
int hOffset = (isFMapRect ? i * 2 : i);
466+
int wOffset = (isFMapRect ? i * 2 + 1 : i);
460467
// Only the first layer is different
461468
if (i == 0)
462469
{
463470
boxParams[i] = {minScale, maxScale, firstLayerAspectRatios.data(), (int) firstLayerAspectRatios.size(),
464-
fMapShapes[i], fMapShapes[i],
471+
fMapShapes[hOffset], fMapShapes[wOffset],
465472
{layerVariances[0], layerVariances[1], layerVariances[2], layerVariances[3]}};
466473
}
467474
else
468475
{
469-
boxParams[i] = {minScale, maxScale, aspectRatios.data(), (int) aspectRatios.size(), fMapShapes[i],
470-
fMapShapes[i], {layerVariances[0], layerVariances[1], layerVariances[2], layerVariances[3]}};
476+
boxParams[i] = {minScale, maxScale, aspectRatios.data(), (int) aspectRatios.size(), fMapShapes[hOffset],
477+
fMapShapes[wOffset], {layerVariances[0], layerVariances[1], layerVariances[2], layerVariances[3]}};
471478
}
472479
}
473480

474-
GridAnchorGenerator* obj = new GridAnchorGenerator(boxParams.data(), numLayers);
481+
GridAnchorGenerator* obj = new GridAnchorGenerator(boxParams.data(), numLayers, mPluginName.c_str());
475482
obj->setPluginNamespace(mNamespace.c_str());
476483
return obj;
477484
}
478485

479-
IPluginV2Ext* GridAnchorPluginCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength)
486+
IPluginV2Ext* GridAnchorBasePluginCreator::deserializePlugin(const char* name, const void* serialData, size_t serialLength)
480487
{
481488
// This object will be deleted when the network is destroyed, which will
482489
// call GridAnchor::destroy()
483-
GridAnchorGenerator* obj = new GridAnchorGenerator(serialData, serialLength);
490+
GridAnchorGenerator* obj = new GridAnchorGenerator(serialData, serialLength, mPluginName.c_str());
484491
obj->setPluginNamespace(mNamespace.c_str());
485492
return obj;
486493
}
494+
495+
GridAnchorPluginCreator::GridAnchorPluginCreator()
496+
{
497+
mPluginName = GRID_ANCHOR_PLUGIN_NAMES[0];
498+
}
499+
500+
GridAnchorRectPluginCreator::GridAnchorRectPluginCreator()
501+
{
502+
mPluginName = GRID_ANCHOR_PLUGIN_NAMES[1];
503+
}

0 commit comments

Comments
 (0)