Thrust sort on tuples is very slow

Question

I need to sort an array of tuples, so I'm defining an operator for tuples and sorting using thrust::sort.

So what I found is that sorting an array of tuples is significantly slower the sorting an array of numbers. Here is my code:

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/set_operations.h>
#include <thrust/reduce.h>
#include <thrust/unique.h>
#include <thrust/binary_search.h>
#include <thrust/gather.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>
#include <iostream>

static const int size = 100000;

#define mzi(x) thrust::make_zip_iterator(x)

#define mt(...) thrust::make_tuple(__VA_ARGS__)

typedef thrust::tuple<int, int> IntTuple;
typedef thrust::device_vector<IntTuple>::iterator TupleIterator;
typedef thrust::device_vector<int>::iterator IntIterator;
typedef thrust::tuple<IntIterator, IntIterator> IteratorTuple;
typedef thrust::zip_iterator<IteratorTuple> ZipIterator;

struct TupleComp
{
    __host__ __device__
    bool operator()(const IntTuple& t1, const IntTuple& t2)
    {
        return t1.get<0>() != t2.get<0>() ? t1.get<0>() < t2.get<0>() : t1.get<1>() > t2.get<1>();
    }
};

int main()
{
    timespec start;
    clock_gettime(0, &start);
    thrust::device_vector<int> dataA1(size);
    thrust::device_vector<int> dataA2(size);
    thrust::device_vector<int> dataB1(size);
    thrust::device_vector<int> dataB2(size);

    srand(time(NULL));
    for (int i = 0; i < size; i++)
    {
        //dataA[i] = dataA[i - 1] + (rand() % 100);
        dataA1[i] = (rand() % 100);
        dataA2[i] = (rand() % 100);
        dataB1[i] = (rand() % 100);
        dataB2[i] = (rand() % 100);
        std::cout << dataA1[i] << "\t" << dataA2[i] << "\t" << dataB1[i] << "\t" << dataB2[i];
        std::cout << std::endl;
    }
    timespec end;
    clock_gettime(0, &end);
    std::cout << "gendb took: " << end.tv_sec - start.tv_sec << "s" << end.tv_nsec - start.tv_nsec << "ns" << std::endl;
    ZipIterator beginA = mzi(mt(dataA1.begin(), dataA2.begin()));
    ZipIterator beginB = mzi(mt(dataB1.begin(), dataB2.begin()));
    ZipIterator endA = mzi(mt(dataA1.end(), dataA2.end()));
    ZipIterator endB = mzi(mt(dataB1.end(), dataB2.end()));
    thrust::device_vector<IntTuple> A(size);
    thrust::device_vector<IntTuple> B(size);

    clock_gettime(0, &start);
    thrust::copy(beginA, endA, A.begin());
    thrust::copy(beginB, endB, B.begin());
    clock_gettime(0, &end);
    std::cout << "thrust::copy took: " << end.tv_sec - start.tv_sec << "s" << end.tv_nsec - start.tv_nsec << "ns" << std::endl;

    clock_gettime(0, &start);
    thrust::sort(A.begin(), A.end());
    clock_gettime(0, &end);
    std::cout << "A thrust::sort took: " << end.tv_sec - start.tv_sec << "s" << end.tv_nsec - start.tv_nsec << "ns" << std::endl;
    clock_gettime(0, &start);
    thrust::sort(B.begin(), B.end(), TupleComp());
    clock_gettime(0, &end);
    std::cout << "B thrust::sort took: " << end.tv_sec - start.tv_sec << "s" << end.tv_nsec - start.tv_nsec << "ns" << std::endl;

    clock_gettime(0, &start);
    thrust::sort(dataA1.begin(), dataA1.end());
    clock_gettime(0, &end);
    std::cout << "regular thrust::sort took: " << end.tv_sec - start.tv_sec << "s" << end.tv_nsec - start.tv_nsec << "ns" << std::endl;
    clock_gettime(0, &start);
    thrust::sort(beginA, endA, TupleComp());
    thrust::sort(beginB, endB, TupleComp());
    clock_gettime(0, &end);
    std::cout << "thrust::sort took: " << end.tv_sec - start.tv_sec << "s" << end.tv_nsec - start.tv_nsec << "ns" << std::endl;
}

I'm getting that tuple sort is ~10X times slower than regular sort.

I don't understand why. Is the complexity of the sort in thrust is directly affected by the operator? Even though, my operator is not 10x times slower than a regular comparator.

Note: It's not just 10x times slower: for 100000 it's ~10x slower for 1000000 it's ~20x slower

I also found that coping two arrays into an array of tuples and sorting that array instead is about 150% faster while the thrust::copy take almost nothing (0.3 for 1M).

Note2:

I changed my operator to this:

struct TupleComp
{
    __host__ __device__
    bool operator()(const IntTuple& t1, const IntTuple& t2)
    {
        if(t1.get<0>() < t2.get<0>())
            return true;
        if(t1.get<0>() > t2.get<0>())
            return false;
        return t1.get<1>() > t2.get<1>();
    }
};

and now the sort is about 112.5% faster, which is probably because equals on the first value is rarely happens, this way there is less if's to check in the operator in general.

Answer 1

Sorry but Nsight totally confused me, All this time I believed I was in release mode, but the run configuration it self was set to run the debug mode.

Now I've made sure everything is set to release, and it works much better.

The difference is only ~150% between int sort and tuple sort, which makes much more sense. Not sure what else can I do to improve the performance but it's good enough already.

The conclusion is: be careful with eclipse preferences, especially on linux.

Answer 2

Just a followup on this that the use of functors gets slower if their call cannot be inlined by the compiler (which could happen due to various reasons like separate compilation units) . A good resource is here