Primitives Usage

This section demonstrates matchFrames function implementation in NVX CUDA API.

The first step is to create necessary data objects:

1. Arrays for corners.

       *ref_features = createArrayKeyPoint(1000u);
       *cur_features = createArrayKeyPoint(1000u);

2. CUDA device pointer for the total number of detections.

       uint32_t *num_corners_dev_ptr = NULL;
       cudaMalloc((void **)&num_corners_dev_ptr, sizeof(uint32_t));

3. Pyramid objects.

       const uint32_t num_levels = 3;
       nvxcu_pitch_linear_pyramid_t ref_pyramid = createPyramidHalfScale(ref_image->base.width, ref_image->base.height, num_levels);
       nvxcu_pitch_linear_pyramid_t cur_pyramid = createPyramidHalfScale(ref_image->base.width, ref_image->base.height, num_levels);

4. Execution target (CUDA stream and CUDA device properties).

       nvxcu_stream_exec_target_t exec_target;
       exec_target.base.exec_target_type = NVXCU_STREAM_EXEC_TARGET;
       cudaStreamCreate(&exec_target.stream);
       cudaGetDeviceProperties(&exec_target.dev_prop, 0);

   The second step is to preallocate temporary buffers that are needed for primitives:

1. Gets the required buffer size from all used primitives.

       nvxcu_tmp_buf_size_t harris_buf_size = nvxcuHarrisCorners_GetBufSize(ref_image->base.width, ref_image->base.height,
                                                                            NVXCU_DF_IMAGE_U8,
                                                                            3.0f /*min_distance*/,
                                                                            3 /*gradient_size*/,
                                                                            3 /*block_size*/,
                                                                            NVXCU_TYPE_KEYPOINT, 1000u,
                                                                            &border,
                                                                            &exec_target.dev_prop);
   
       nvxcu_tmp_buf_size_t pyr_buf_size = nvxcuGaussianPyramid_GetBufSize(ref_image->base.width, ref_image->base.height,
                                                                           num_levels, NVXCU_SCALE_PYRAMID_HALF,
                                                                           &border,
                                                                           &exec_target.dev_prop);
   
       nvxcu_tmp_buf_size_t pyrlk_buf_size = nvxcuOpticalFlowPyrLK_GetBufSize(ref_image->base.width, ref_image->base.height,
                                                                              num_levels, NVXCU_SCALE_PYRAMID_HALF,
                                                                              NVXCU_TYPE_KEYPOINT, 1000u,
                                                                              6u /*window_dimension*/,
                                                                              0 /*calcError*/,
                                                                              &border,
                                                                              &exec_target.dev_prop);

2. Calculates total size of temporary buffers and allocates them.

       nvxcu_tmp_buf_size_t final_buf_size;
       final_buf_size.dev_buf_size = harris_buf_size.dev_buf_size;
       final_buf_size.dev_buf_size =
               final_buf_size.dev_buf_size > pyr_buf_size.dev_buf_size ?
                   final_buf_size.dev_buf_size :
                   pyr_buf_size.dev_buf_size;
       final_buf_size.dev_buf_size =
               final_buf_size.dev_buf_size > pyrlk_buf_size.dev_buf_size ?
                   final_buf_size.dev_buf_size :
                   pyrlk_buf_size.dev_buf_size;
       final_buf_size.host_buf_size = harris_buf_size.host_buf_size;
       final_buf_size.host_buf_size =
               final_buf_size.host_buf_size > pyr_buf_size.host_buf_size ?
                   final_buf_size.host_buf_size :
                   pyr_buf_size.host_buf_size;
       final_buf_size.host_buf_size =
               final_buf_size.host_buf_size > pyrlk_buf_size.host_buf_size ?
                   final_buf_size.host_buf_size :
                   pyrlk_buf_size.host_buf_size;
   
       nvxcu_tmp_buf_t tmp_buf = {NULL, NULL};
       if (final_buf_size.dev_buf_size > 0)
       {
           cudaMalloc(&tmp_buf.dev_ptr, final_buf_size.dev_buf_size);
       }
       if (final_buf_size.host_buf_size > 0)
       {
           cudaMallocHost(&tmp_buf.host_ptr, final_buf_size.host_buf_size);
       }

   The third step is to call primitives.

1. Harris corner detector.

       nvxcuHarrisCorners(&ref_image->base,
                          0.00004f /*strength_thresh*/,
                          3.0f /*min_distance*/,
                          0.04f /*sensitivity*/,
                          3 /*gradient_size*/,
                          3 /*block_size*/,
                          &ref_features->base,
                          num_corners_dev_ptr,
                          &tmp_buf,
                          &border,
                          &exec_target.base);

2. Gaussian pyramids. Note that pyramids are built from the first level, so the code copies input images to the first level of the corresponding pydamid.

       cudaMemcpy2DAsync(ref_pyramid.levels[0].planes[0].dev_ptr, ref_pyramid.levels[0].planes[0].pitch_in_bytes,
                         ref_image->planes[0].dev_ptr, ref_image->planes[0].pitch_in_bytes,
                         ref_image->base.width * sizeof(uint8_t), ref_image->base.height,
                         cudaMemcpyDeviceToDevice,
                         exec_target.stream);
   
       nvxcuGaussianPyramid(&ref_pyramid.base,
                            &tmp_buf,
                            &border,
                            &exec_target.base);
   
       cudaMemcpy2DAsync(cur_pyramid.levels[0].planes[0].dev_ptr, cur_pyramid.levels[0].planes[0].pitch_in_bytes,
                         cur_image->planes[0].dev_ptr, cur_image->planes[0].pitch_in_bytes,
                         cur_image->base.width * sizeof(uint8_t), cur_image->base.height,
                         cudaMemcpyDeviceToDevice,
                         exec_target.stream);
   
       nvxcuGaussianPyramid(&cur_pyramid.base,
                            &tmp_buf,
                            &border,
                            &exec_target.base);

3. Optical flow. Note that the cur_features parameter must be initialized with initial estimation. Points from the first image (ref_features) are used.

       cudaMemcpyAsync(cur_features->dev_ptr, ref_features->dev_ptr,
                       ref_features->base.capacity * sizeof(nvxcu_keypoint_t),
                       cudaMemcpyDeviceToDevice,
                       exec_target.stream);
   
       cudaMemcpyAsync(cur_features->num_items_dev_ptr, ref_features->num_items_dev_ptr,
                       sizeof(uint32_t),
                       cudaMemcpyDeviceToDevice,
                       exec_target.stream);
   
       nvxcuOpticalFlowPyrLK(&ref_pyramid.base,
                             &cur_pyramid.base,
                             &ref_features->base,
                             &cur_features->base,
                             NVXCU_TERM_CRITERIA_BOTH,
                             0.01f /*epsilon*/,
                             30 /*num_iterations*/,
                             6u /*window_dimension*/,
                             0 /*calcError*/,
                             &tmp_buf,
                             &border,
                             &exec_target.base);

4. Synchronize with CUDA stream to be sure that all pending operations are finished.

       cudaStreamSynchronize(exec_target.stream);

   The last step is to clean up all the temporarily allocated buffers and data objects:

    cudaFree(num_corners_dev_ptr);

    for (uint32_t i = 0; i < num_levels; ++i)
    {
        cudaFree(ref_pyramid.levels[i].planes[0].dev_ptr);
        cudaFree(cur_pyramid.levels[i].planes[0].dev_ptr);
    }
    free(ref_pyramid.levels);
    free(cur_pyramid.levels);

    cudaStreamDestroy(exec_target.stream);

    cudaFree(tmp_buf.dev_ptr);
    cudaFreeHost(tmp_buf.host_ptr);

The full code of the matchFrames function:

void matchFrames(const nvxcu_pitch_linear_image_t* ref_image,
                 const nvxcu_pitch_linear_image_t* cur_image,
                 nvxcu_plain_array_t* ref_features,
                 nvxcu_plain_array_t* cur_features)
{
    assert(ref_image->base.format == NVXCU_DF_IMAGE_U8);
    assert(ref_image->base.width == cur_image->base.width);
    assert(ref_image->base.height == cur_image->base.height);
    assert(ref_image->base.format == cur_image->base.format);

    // create data objects

    *ref_features = createArrayKeyPoint(1000u);
    *cur_features = createArrayKeyPoint(1000u);

    uint32_t *num_corners_dev_ptr = NULL;
    cudaMalloc((void **)&num_corners_dev_ptr, sizeof(uint32_t));

    const uint32_t num_levels = 3;
    nvxcu_pitch_linear_pyramid_t ref_pyramid = createPyramidHalfScale(ref_image->base.width, ref_image->base.height, num_levels);
    nvxcu_pitch_linear_pyramid_t cur_pyramid = createPyramidHalfScale(ref_image->base.width, ref_image->base.height, num_levels);

    nvxcu_stream_exec_target_t exec_target;
    exec_target.base.exec_target_type = NVXCU_STREAM_EXEC_TARGET;
    cudaStreamCreate(&exec_target.stream);
    cudaGetDeviceProperties(&exec_target.dev_prop, 0);

    nvxcu_border_t border;
    border.mode = NVXCU_BORDER_MODE_REPLICATE;

    // allocate temporary buffers

    nvxcu_tmp_buf_size_t harris_buf_size = nvxcuHarrisCorners_GetBufSize(ref_image->base.width, ref_image->base.height,
                                                                         NVXCU_DF_IMAGE_U8,
                                                                         3.0f /*min_distance*/,
                                                                         3 /*gradient_size*/,
                                                                         3 /*block_size*/,
                                                                         NVXCU_TYPE_KEYPOINT, 1000u,
                                                                         &border,
                                                                         &exec_target.dev_prop);

    nvxcu_tmp_buf_size_t pyr_buf_size = nvxcuGaussianPyramid_GetBufSize(ref_image->base.width, ref_image->base.height,
                                                                        num_levels, NVXCU_SCALE_PYRAMID_HALF,
                                                                        &border,
                                                                        &exec_target.dev_prop);

    nvxcu_tmp_buf_size_t pyrlk_buf_size = nvxcuOpticalFlowPyrLK_GetBufSize(ref_image->base.width, ref_image->base.height,
                                                                           num_levels, NVXCU_SCALE_PYRAMID_HALF,
                                                                           NVXCU_TYPE_KEYPOINT, 1000u,
                                                                           6u /*window_dimension*/,
                                                                           0 /*calcError*/,
                                                                           &border,
                                                                           &exec_target.dev_prop);

    nvxcu_tmp_buf_size_t final_buf_size;
    final_buf_size.dev_buf_size = harris_buf_size.dev_buf_size;
    final_buf_size.dev_buf_size =
            final_buf_size.dev_buf_size > pyr_buf_size.dev_buf_size ?
                final_buf_size.dev_buf_size :
                pyr_buf_size.dev_buf_size;
    final_buf_size.dev_buf_size =
            final_buf_size.dev_buf_size > pyrlk_buf_size.dev_buf_size ?
                final_buf_size.dev_buf_size :
                pyrlk_buf_size.dev_buf_size;
    final_buf_size.host_buf_size = harris_buf_size.host_buf_size;
    final_buf_size.host_buf_size =
            final_buf_size.host_buf_size > pyr_buf_size.host_buf_size ?
                final_buf_size.host_buf_size :
                pyr_buf_size.host_buf_size;
    final_buf_size.host_buf_size =
            final_buf_size.host_buf_size > pyrlk_buf_size.host_buf_size ?
                final_buf_size.host_buf_size :
                pyrlk_buf_size.host_buf_size;

    nvxcu_tmp_buf_t tmp_buf = {NULL, NULL};
    if (final_buf_size.dev_buf_size > 0)
    {
        cudaMalloc(&tmp_buf.dev_ptr, final_buf_size.dev_buf_size);
    }
    if (final_buf_size.host_buf_size > 0)
    {
        cudaMallocHost(&tmp_buf.host_ptr, final_buf_size.host_buf_size);
    }

    // call primitives

    nvxcuHarrisCorners(&ref_image->base,
                       0.00004f /*strength_thresh*/,
                       3.0f /*min_distance*/,
                       0.04f /*sensitivity*/,
                       3 /*gradient_size*/,
                       3 /*block_size*/,
                       &ref_features->base,
                       num_corners_dev_ptr,
                       &tmp_buf,
                       &border,
                       &exec_target.base);

    cudaMemcpy2DAsync(ref_pyramid.levels[0].planes[0].dev_ptr, ref_pyramid.levels[0].planes[0].pitch_in_bytes,
                      ref_image->planes[0].dev_ptr, ref_image->planes[0].pitch_in_bytes,
                      ref_image->base.width * sizeof(uint8_t), ref_image->base.height,
                      cudaMemcpyDeviceToDevice,
                      exec_target.stream);

    nvxcuGaussianPyramid(&ref_pyramid.base,
                         &tmp_buf,
                         &border,
                         &exec_target.base);

    cudaMemcpy2DAsync(cur_pyramid.levels[0].planes[0].dev_ptr, cur_pyramid.levels[0].planes[0].pitch_in_bytes,
                      cur_image->planes[0].dev_ptr, cur_image->planes[0].pitch_in_bytes,
                      cur_image->base.width * sizeof(uint8_t), cur_image->base.height,
                      cudaMemcpyDeviceToDevice,
                      exec_target.stream);

    nvxcuGaussianPyramid(&cur_pyramid.base,
                         &tmp_buf,
                         &border,
                         &exec_target.base);

    cudaMemcpyAsync(cur_features->dev_ptr, ref_features->dev_ptr,
                    ref_features->base.capacity * sizeof(nvxcu_keypoint_t),
                    cudaMemcpyDeviceToDevice,
                    exec_target.stream);

    cudaMemcpyAsync(cur_features->num_items_dev_ptr, ref_features->num_items_dev_ptr,
                    sizeof(uint32_t),
                    cudaMemcpyDeviceToDevice,
                    exec_target.stream);

    nvxcuOpticalFlowPyrLK(&ref_pyramid.base,
                          &cur_pyramid.base,
                          &ref_features->base,
                          &cur_features->base,
                          NVXCU_TERM_CRITERIA_BOTH,
                          0.01f /*epsilon*/,
                          30 /*num_iterations*/,
                          6u /*window_dimension*/,
                          0 /*calcError*/,
                          &tmp_buf,
                          &border,
                          &exec_target.base);

    cudaStreamSynchronize(exec_target.stream);

    // clean up

    cudaFree(num_corners_dev_ptr);

    for (uint32_t i = 0; i < num_levels; ++i)
    {
        cudaFree(ref_pyramid.levels[i].planes[0].dev_ptr);
        cudaFree(cur_pyramid.levels[i].planes[0].dev_ptr);
    }
    free(ref_pyramid.levels);
    free(cur_pyramid.levels);

    cudaStreamDestroy(exec_target.stream);

    cudaFree(tmp_buf.dev_ptr);
    cudaFreeHost(tmp_buf.host_ptr);
}