How to use a struct array in thrust::sort?

Question

I'm fairly new to CUDA development and I'm trying to sort a struct array using thrust library's sort method. My struct is like this:

#define N 307200    
struct distanceVector {
   Point3D a, b;
   float distance;
};

I want to sort the array on "distance", however, the sort function requires two random access iterators and since I'm not using vectors I don't have any. I've tried doing something like this:

bool distance_sort(distanceVector A, distanceVector B){
   return (A.distance > B.distance);
}

distanceVector * MyStructArray;
cudaMalloc((void**)&MyStructArray, sizeof(distanceVector) * N);
//LAUNCH KERNEL WHICH FILLS MYSTRUCTARRAY AND THEN...
thrust::sort(MyStructArray, MyStructArray + N, distance_sort);

... which I saw as an example in [thrust's guide][1]:

#include <thrust/sort.h>
#include <thrust/functional.h>
...
const int N = 6;
int A[N] = {1, 4, 2, 8, 5, 7};
thrust::stable_sort(A, A + N, thrust::greater<int>());
// A is now {8, 7, 5, 4, 2, 1}

Although it compiles, during execution I get a "Access violation reading location 0x405e041c." error. While debugging the application stops at this section in the insertion_sort.h file:

for(RandomAccessIterator i = first + 1; i != last; ++i)
  {
    value_type tmp = *i;

    if (wrapped_comp(tmp, *first)).......

Is there a way to solve this without using thrust's vectors?

Answer 1

Ok, so I found my error. The problem is I'm trying to use thrust on memory that is allocated on the device. I tried copying MyStructArray to the host device first and then using thrust's sort and it worked perfectly. In order to work on the device's memory I have to use thrust::device_ptr<MyStruct> pointer(MyStructArray). Hope this helps someone else.