DriveWorks SDK Reference
3.0.4260 Release
For Test and Development only

Single Camera Feature Tracking
Note
SW Release Applicability: This tutorial is applicable to modules in both NVIDIA DriveWorks and NVIDIA DRIVE Software releases.

This tutorial shows how dwFeature2DDetector and dwFeature2DTracker modules are typically used. Note that error handling is left out for clarity.

Initialize feature history array on both CPU and GPU

dwFeatureHistoryArray_create(&featureHistoryCPU, maxFeatureCount, maxHistoryCapacity,
DW_MEMORY_TYPE_CPU, contextHandle);
dwFeatureHistoryArray_create(&featureHistoryGPU, maxFeatureCount, maxHistoryCapacity,
DW_MEMORY_TYPE_CUDA, contextHandle);

Initialize feature array on GPU for detector

dwFeatureArray_create(&featuresDetected, maxFeatureCount, DW_MEMORY_TYPE_CUDA, contextHandle);

Initialize feature detector and tracker parameters with default values

Modify parameters according to input image and detection requirements, create detector. Detector will first split the given WxH input image into cellSize x cellSize cells, detect Harris corner whose score is higher than scoreThreshold. Features with score >= scoreThreshold will be considered as high frequency details (red crosses). Detail features will be added to candidate list first, detector will then sort the rest features with score between (scoreThreshold, detailThreshold) from highest to lowest by cell, select the first numEvenDistributionPerCell features in each cell (blue crosses, here numEvenDistributionPerCell = 5, so there're 5 second-highest-score features in each cell). These second-highest-score features are distributed evenly across cells.

feature_tracker.png

Setting scoreThreshold >= detailThreshold means all features are high frequency details, numEvenDistributionPerCell will be invalid. Setting numEvenDistributionPerCell = 0 turns off the even distribution component. Both of the 2 cases will make detected features be clustered in high frequency regions, i.e. only red crosses will be left in above illustration. contextHandle is assumed to be previously initialized dwContextHandle_t.

detectorConfig.imageWidth = inputImageProps.width;
detectorConfig.imageHeight = inputImageProps.height;
detectorConfig.scoreThreshold = scoreThreshold;
detectorConfig.detailThreshold = detailThreshold;
detectorConfig.cellSize = cellSize;
detectorConfig.maxFeatureCount = maxFeatureCount;
detectorConfig.numEvenDistributionPerCell = numEvenDistributionPerCell;
dwFeature2DDetector_initialize(&detector, &detectorConfig, 0, contextHandle);

Modify parameters to use PVA Harris feature detector, create detector. The standard Harris corner detection process is applied. First detector computes the spatial gradient, using a separable gradient filter. Gradient filter size is 3x3, 5x5 or 7x7, correspond to gradientSize value 3,5,7. Then Computes a gradients covariance matrix (structure tensor) for each pixel within a block window, the block window size is 3x3, 5x5, or 7x7, correspond to blockSize value 3,5,7. Later Computes a Harris response score using a sensitivity factor, setting harrisK to modify the factor. Last applies a threshold-strength criterion, the threshold value is setted by scoreThreshold, it should be a uint32 value. contextHandle is assumed to be previously initialized dwContextHandle_t.

detectorConfig.imageWidth = inputImageProps.width;
detectorConfig.imageHeight = inputImageProps.height;
detectorConfig.scoreThreshold = scoreThreshold;
detectorConfig.blockSize = blockSize;
detectorConfig.gradientSize = gradientSize;
detectorConfig.maxFeatureCount = maxFeatureCount;
detectorConfig.harrisK = harrisK;
detectorConfig.processorType = DW_PROCESSOR_TYPE_PVA_0;//or DW_PROCESSOR_TYPE_PVA_1
dwFeature2DDetector_initialize(&detector, &detectorConfig, 0, contextHandle);

Modify parameters to use Fast9 feature detector, create detector. Fast9 Corner Detector will detect corner whose pixel vaule is larger or smaller than at least 9 continous pixels on the Bresenham circle around it by at least scoreThreshold. And non-maximum suppression can be choosen to apply to filter detected features. contextHandle is assumed to be previously initialized dwContextHandle_t.

detectorConfig.imageWidth = inputImageProps.width;
detectorConfig.imageHeight = inputImageProps.height;
detectorConfig.scoreThreshold = scoreThreshold;
detectorConfig.maxFeatureCount = maxFeatureCount;
detectorConfig.NMSRadius = nmsRadius;
detectorConfig.processorType = DW_PROCESSOR_TYPE_GPU;//or DW_PROCESSOR_TYPE_PVA_0, DW_PROCESSOR_TYPE_PVA_1
dwFeature2DDetector_initialize(&detector, &detectorConfig, 0, contextHandle);

Modify parameters according to input image and tracking requirements, create tracker. trackerConfig.isInputCompact must be read from detector. contextHandle is assumed to be previously initialized dwContextHandle_t.

trackerConfig.imageWidth = inputImageProps.width;
trackerConfig.imageHeight = inputImageProps.height;
trackerConfig.windowSizeLK = 10;
trackerConfig.iterationsLK = 10;
trackerConfig.detectorType = detectorConfig.type;
dwFeature2DTracker_initialize(&tracker, &trackerConfig, 0, contextHandle);

Create 2 image pyramids as detection/tracking input: one is current, one is previous, the 2 pyramids work as a double buffer (or ping-pong buffer). They swap each other after each detection/tracking cycle. So new frame will always be read into pyrmaidCurrent and the previous frame will be stored to pyramidPrevious.

dwPyramid_create only creates pyramid image and allocates memory, pyramid will be filled in dwImageFilter_computePyramid. Top level (level 0) in pyramid is always the same size as original input image.

inputImageProps.width, inputImageProps.height are the same in tracker/detector initialization

imagePxlType should be got from dwImage_getPixelType(inputImageProps.format).

contextHandle is assumed to be previously initialized dwContextHandle_t.

dwPyramid_create(&pyramidPrevious, levelCount,
inputImageProps.width, inputImageProps.height, pxlType, context);
dwPyramid_create(&pyramidCurrent, levelCount,
inputImageProps.width, inputImageProps.height, pxlType,context);

#### Create nccScore array on GPU

cudaMalloc(&d_nccScores, maxFeatureCount * sizeof(float32_t));

Start tracking the loop

while(true)
{
// ping-pong current/previous pyramid
// so new frame will always be built in pyramidCurrent
// and old pyramid will be stored in pyramidPrevious
std::swap(pyramidCurrent, pyramidPrevious);
// CODE: Read image frame code
dwImageHandle_t image = ...;
// compute pyramid based on the new read frame image
dwImageFilter_computePyramid(&pyramidCurrent, image, 0, contextHandle);
// do tracking
dwFeatureArray featurePredicted;
dwFeature2DTracker_trackFeatures(&featureHistoryGPU, &featurePredicted,
d_nccScores, &featuresDetected,
&pyramidPrevious, &pyramidCurrent, tracker)
// detect new features in the current frame
dwFeature2DDetector_detectFromImage(&featuresDetected, pyramidCurrent.levelImages[0],
&featurePredicted, d_nccScores, detector)
// copy detected results to CPU
dwFeatureHistoryArray_copy(&featureHistoryCPU, &featureHistoryGPU, 0);
// do some stuff in CPU
featureHistoryCPU.locationHistory[...];
featureHistoryCPU.statuses[...];
}

Finally, free previously allocated memory.

cudaFree(d_nccScores);
dwPyramid_destroy(pyramidPrevious);
dwPyramid_destroy(pyramidCurrent);
dwFeatureHistoryArray_destroy(featureHistoryCPU);
dwFeatureHistoryArray_destroy(featureHistoryGPU);
dwFeatureArray_destroy(featuresDetected);

For the full implementation refer to Feature Tracker Sample.