Ask Your Question
0

GPU SIFT and dealing with GPU pointers

asked 2013-07-01 05:49:50 -0600

proghero gravatar image

updated 2013-07-01 05:51:34 -0600

Hi,

I'm using the GPU libraires of OpenCV to perform a 3D-geometric based image comparison. I use SURF_GPU to get keypoints and descriptors and cv::gpu::BruteForceMatcher_GPU to do the actual matching. I already implemented a serial version which works fine, but I'm having some hard time figuring out how to deal with some OpenCV structures using a GPU without having to convert between data types all the time and saving intermediary results in the CPU. Here is what I already did and what I'm trying to do now:

  1. SIFT extractor stuff and matches in the GPU are OK! I do the matchings in both directions, so I can get only matches which are symmetrical, i.e.


    GpuMat trainIdxMat1, distanceMat1, allDist1;
    GpuMat trainIdxMat2, distanceMat2, allDist2;
    matcher.knnMatchSingle(im1_descriptors_gpu, im2_descriptors_gpu, 
            trainIdxMat1, distanceMat1, allDist1, 2);
    matcher.knnMatchSingle(im2_descriptors_gpu, im1_descriptors_gpu, 
            trainIdxMat2, distanceMat2, allDist2, 2);

  1. Rejecting bad matches - problems dealing with GPU pointers! As I want to filter out bad matches, (in my serial version) I iterate through the found matches and delete matches for which the ratio first_best_match.distance/second_best_match.distance is larger than some threshold (e.g. 0.7). I already wrote a kernel myself and want to pass it pointers to the found matches so it can process them all. I get these pointers like explained here. That's part of my code to do it:


    int2* trainIdx1 = trainIdxMat1.ptr<int2>();
    float2* distance1 = distanceMat1.ptr<float2>();
    int2* trainIdx2 = trainIdxMat2.ptr<int2>();
    float2* distance2 = distanceMat2.ptr<float2>();
Then I pass these pointers to my CUDA-Kernel. My problem here is, that I'm not sure how to best delete the bad matches without having to downloading them and storing them somewhere and so on. It would be great if I could delete them directly from the device, i.e. from my matching array.

  1. Iterating through both match arrays and check which matched points are in matches1 and matches2 as well. Here I have similar questions about deleting matched points from my matches without having to juggle with pointers, arrays and minimizing data transfer between CPU and GPU. I think, I could figure it out, after getting a better understanding about the things in (2).

I hope the question is not too confusing! Please, let me know if I should clarify something better! I would be very thankful for any kind of help!

edit retag flag offensive close merge delete

1 answer

Sort by ยป oldest newest most voted
1

answered 2013-07-01 07:48:41 -0600

Vladislav Vinogradov gravatar image

updated 2013-07-01 07:56:22 -0600

You can try to use Thrust library and thrust::remove_if algorithm (see example). Something like this:

struct is_bad_match
{
    template <typename Tuple>
    __device__ bool operator()(const Tuple& tuple) const
    {
        // unpack the tuple
        const int2 trainIdx1 = thrust::get<0>(tuple);
        const float2 distance1 = thrust::get<1>(tuple);

        // check if it is a bad match
        return (distance1.x / distance1.y) < 0.7;
    }
};

thrust::device_ptr<int2> trainIdx1_ptr(trainIdx1);
thrust::device_ptr<float2> distance1_ptr(distance1);

size_t new_size = thrust::remove_if(thrust::make_zip_iterator(thrust::make_tuple(trainIdx1_ptr, distance1_ptr)),
                                    thrust::make_zip_iterator(thrust::make_tuple(trainIdx1_ptr + size, distance1_ptr + size)),
                                    is_bad_match())
                      - thrust::make_zip_iterator(thrust::make_tuple(trainIdx1_ptr, distance1_ptr));
edit flag offensive delete link more

Comments

I am trying to compile CV3.0 code with Thrust but I get a bunch or errors about undefined

/usr/local/cuda-7.0/include/thrust/system/cuda/detail/error.inl:63: undefined reference to cudaGetErrorString' /usr/local/cuda-7.0/include/thrust/system/cuda/detail/trivial_copy.inl:44: undefined reference tocudaMemcpyAsync'

I did some search and seems that i am missing -lcudaart during compilation, is that needed when compiling thru opencv ? any suggestions on how I can do it ?

abhiguru gravatar imageabhiguru ( 2015-06-10 06:25:21 -0600 )edit

Question Tools

Stats

Asked: 2013-07-01 05:49:50 -0600

Seen: 1,561 times

Last updated: Jul 01 '13