Is there any way to use thrust with cufftComplex data type when the data is in device?

Question

I'm making program where I have 4D data. In a loop I iterate through the 4th dimension, and take the 3D data which is originally float, I convert that data to cufftComplex and then copy it to device and do several device operations like fourier transform(from cufft library), image denoising(my own cuda kernel). While being in that loop in the between these operations after the fourier transform I want find the median of the 3D data in that particular iteration.

I know thrust allows me to do sorting from which I can easily find the median value like in the following example with a normal float data type I do

int row, column, slice;
row = 5;
column = 5;
slice = 5;
int alloc_size = row*column*slice;
float * theMainData = new float[alloc_size];
srand(time(NULL));
for(int k=0; k<slice; k++)
{
    for(int j=0; j<column; j++)
    {
       for(int i=0; i<row; i++)
       {
          theMainData[i+(j*column)+(k*row*column)] = rand()%1000;;
       }
    }
}
thrust::sort(theMainData, theMainData + alloc_size);
float median = theMainData[(alloc_size/2)];
cout<<endl<<"The median value is "<<median<<endl;

So, now my question is, is there any way to make thrust::sort() function work with cufftComplex when the data already copied to the device? Or is there any other function similar to thrust which can deal with cufftComplex data type when the data is already copied to the device?

Answer 1

So, now my question is, is there any way to make thrust::sort() function work with cufftComplex when the data already copied to the device?

Yes.

Thrust provides versions of the sort operation that accept user defined comparison operators, to order the data any way you wish.

An example usage with a custom functor is given in the thrust sorting example (using the evens_before_odds functor).

So you'll need a comparison functor. I don't know how you sort complex values, but as an example if you wanted to sort them based on the magnitude you could do:

#include <cuComplex.h>
#include <cufft.h>


struct my_complex_sort
{
  __host__ __device__
  bool operator()(cufftComplex x, cufftComplex y)
  {
    return (cuCabs(x) > cuCabs(y));
  }
};

If you already have cufftComplex data on the device, you could create a thrust pointer for it. Let's suppose the pointer to your device data is d_data and the data size is N:

thrust::device_ptr<cufftComplex> dp_data = thrust::device_pointer_cast(d_data);

Then you could sort the data with:

thrust::sort(dp_data, dp_data+N, my_complex_sort());

(disclaimer: coded in browser, not tested)