Using thrust with cv::cuda::GpuMat

asked 2015-09-01 10:46:19 -0500

dtmoodie gravatar image

updated 2015-09-02 10:21:45 -0500

I just struggled with this for a bit so I wanted to post it somewhere where it may be helpful for someone else.

This function can be used for creating a thrust iterator that correctly indexes a cv::cuda::GpuMat.

struct step_functor : public thrust::unary_function<int, int>
{
    int columns;
    int step;
    step_functor(int columns_, int step_) : columns(columns_), step(step_)  {   };
    __host__ __device__
    int operator()(int x) const
    {
        int row = x / columns;
        int idx = (row * step) + x % columns;
        return idx;
    }
};


template<typename T>
thrust::permutation_iterator<thrust::device_ptr<T>, thrust::transform_iterator<step_functor,     thrust::counting_iterator<int>>>  GpuMatBeginItr(cv::cuda::GpuMat mat)
{
    return thrust::make_permutation_iterator(thrust::device_pointer_cast(mat.ptr<T>(0)),
        thrust::make_transform_iterator(thrust::make_counting_iterator(0), 
           step_functor(mat.cols, mat.step / sizeof(T))));
}

template<typename T>
thrust::permutation_iterator<thrust::device_ptr<T>, thrust::transform_iterator<step_functor, thrust::counting_iterator<int>>>  GpuMatEndItr(cv::cuda::GpuMat mat)
{
    return thrust::make_permutation_iterator(thrust::device_pointer_cast(mat.ptr<T>(0)),
        thrust::make_transform_iterator(thrust::make_counting_iterator(mat.rows),
          step_functor(mat.cols*mat.rows, mat.step / sizeof(T))));
}

Thus performing thrust operations on rows / columns is as easy as:

cv::cuda::GpuMat d_test(h_test);

auto keyBegin = GpuMatBeginItr<int>(d_test.col(4));
auto keyEnd = GpuMatEndItr<int>(d_test.col(4));
auto valueBegin = GpuMatBeginItr<int>(d_test.col(5));

thrust::sort_by_key(keyBegin, keyEnd, valueBegin);
edit retag flag offensive close merge delete

Comments

Would you be up to creating a guide for using Thrust and OpenCV combined?

StevenPuttemans gravatar imageStevenPuttemans ( 2015-09-02 07:45:37 -0500 )edit
1

Sure I'll look into it.

dtmoodie gravatar imagedtmoodie ( 2015-09-02 10:22:24 -0500 )edit

I assume you are the one who write this tutorial in OpenCV

http://docs.opencv.org/3.1.0/d5/d8f/m...

Do you have any idea, if I want to pass the type with 3 channels at the same time to thrust kernel (e.g. CV_8UC3, CV_32FC3)? I tried uchar3 and float3, but it seems their size are not the same as cv::Vec3b and cv::Vec3f.

Surely, I tried to send them each separately and they worked fine. Just want the code to be more compact. Thank you.

milLII3yway gravatar imagemilLII3yway ( 2016-04-05 01:57:08 -0500 )edit

I've used GpuMatBeginItr<cv::vec3b> on CV_8UC3 images without a problem, similarly with CV_32FC3.

dtmoodie gravatar imagedtmoodie ( 2016-11-30 10:01:44 -0500 )edit