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.

Answer 2

I try to use thrust::sort to sort thrust::host_vector<struct> and thrust::device_vector<struct>. They are both running well in my device.

This is my code:

#include <bits/stdc++.h>
#include <stdio.h>
#include <cuda_runtime.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <time.h>
#include <thrust/functional.h>
const int MAXN = 1e2;
typedef struct Node {
    int x, y;
    __host__ __device__ bool operator < (const Node& a) const {
        if (x == a.x) {
            return x ? y<a.y : y>a.y;
        }else 
            return x < a.x;
    }
}Node;
int main() {
    thrust::host_vector<Node> h(MAXN),ne(MAXN);
    for (int i = 0; i < MAXN; i++) {
        h[i].x = rand() % 2;
        h[i].y = rand() % 997;
    }
    thrust::device_vector<Node> d = h;
    thrust::sort(h.begin(), h.end());
    thrust::sort(d.begin(), d.end());
    ne = d;
    for (int i = 0; i < MAXN; i++) {
        std::cout << h[i].x << " " << h[i].y << "  |  " << ne[i].x << " " << ne[i].y << std::endl;
    }
}