Sortieren von Arrays von Strukturen in CUDA

Question

Ich habe einen Laptop mit einer NVIDIA GT750M 4Gb (Compute Capability 3.0) Grafikkarte. Ich muss ein Array von Strukturen auf CUDA sortieren (ca. 3 × 10^7 Elemente). Aber ich kann nicht herausfinden, wie, da ich nicht genug Erfahrung in CUDA habe. Wenn ich thrust::sort benutze, erhalte ich seltsame Ergebnisse (es dauert einige zehn Minuten, während std::sort 1 Minute dauert).

struct MyStruct
{
    float key;
    float a;
    float b;
    int c;
    int d;
    int e;
    int f;
    bool flag;
}
bool minCompare(const MyStruct lhs, const MyStruct rhs)
{
    return lhs.key < rhs.key;
}

Answer 1

Wie Robert Crovella in seinem Kommentar bereits angemerkt hat, bedeutet „Tents of minutes“ höchstwahrscheinlich, dass Sie etwas falsch machen. Im Folgenden finden Sie ein Beispiel, in dem ich die Leistung beim Sortieren eines Arrays von Strukturen (AoS) und einer Struktur von Arrays (SoA) mit thrust::sort und thrust::sort_by_key vergleiche. Ich arbeite mit einer Laptop GeForce GT 540M und kompiliere mit CUDA 5.5, daher haben Sie eine leistungsstärkere Karte als ich. Für 100000 Elemente beträgt die Ausführungszeit in beiden Fällen nur wenige Sekunden. Wie ich bereits in meinem Kommentar erwähnte, ist der erste Fall in Bezug auf die Rechenzeit anspruchsvoller (1675ms) als der zweite (668.9ms).

#include 
#include 

struct MyStruct1
{
    int key;
    int value1;
    int value2;
};

struct MyStruct2
{
    int N;
    int* key;
    int* value1;
    int* value2;

    MyStruct2(int N_) {
        N = N_;
        cudaMalloc((void**)&key,N*sizeof(int));
        cudaMalloc((void**)&value1,N*sizeof(int));
        cudaMalloc((void**)&value2,N*sizeof(int));
    }

};

__host__ __device__ bool operator<(const MyStruct1 &lhs, const MyStruct1 &rhs) { return (lhs.key < rhs.key); };

void main(void)
{
    const int N = 10000;

    float time;
    cudaEvent_t start, stop;

    /*******************************/
    /* SORTING ARRAY OF STRUCTURES */
    /*******************************/
    thrust::host_vector h_struct1(N);
    for (int i = 0; i d_struct(h_struct1);

cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);

thrust::sort(d_struct.begin(), d_struct.end());

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
printf("Sortieren von Strukturen - vergangene Zeit:  %3.1f ms \n", time);  

h_struct1 = d_struct;

//for (int i = 0; i h_temp_key(N);
thrust::host_vector h_temp_value1(N);
thrust::host_vector h_temp_value2(N);

//for (int i = 0; i dev_ptr_key     = thrust::device_pointer_cast(d_struct2.key);
thrust::device_ptr dev_ptr_value1  = thrust::device_pointer_cast(d_struct2.value1);
thrust::device_ptr dev_ptr_value2  = thrust::device_pointer_cast(d_struct2.value2);

thrust::device_vector d_indices(N);
thrust::sequence(d_indices.begin(), d_indices.end(), 0, 1);

// first sort the keys and indices by the keys
thrust::sort_by_key(dev_ptr_key, dev_ptr_key + N, d_indices.begin());

// Now reorder the ID arrays using the sorted indices
thrust::gather(d_indices.begin(), d_indices.end(), dev_ptr_value1, dev_ptr_value1);
thrust::gather(d_indices.begin(), d_indices.end(), dev_ptr_value2, dev_ptr_value2);

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
printf("Sortieren der Struktur von Arrays - vergangene Zeit:  %3.1f ms \n", time);  

cudaMemcpy(h_temp_key.data(),d_struct2.key,N*sizeof(int),cudaMemcpyDeviceToHost);
cudaMemcpy(h_temp_value1.data(),d_struct2.value1,N*sizeof(int),cudaMemcpyDeviceToHost);
cudaMemcpy(h_temp_value2.data(),d_struct2.value2,N*sizeof(int),cudaMemcpyDeviceToHost);

//for (int i = 0; i

`

Zur Vereinfachung habe ich darauf verzichtet, eine ordnungsgemäße CUDA-Fehlerprüfung hinzuzufügen, wie sie unter Was ist der canonische Weg, um Fehler mit der CUDA-Laufzeit-API zu überprüfen? beschrieben ist.

`

Answer 2

Falls noch jemand sucht, wenn Sie sich entscheiden, die Thrust-Bibliotheken zu verwenden, können Sie einen Zip-Iterator von Tupeln aus der Struktur von Arrays erstellen, die Sie sortieren möchten, und diesen Iterator an thrust::sort_by_key() übergeben, wie folgt:

void sort() {
   const int N = 6;
   int  *keys = new int[N]; // = { 1,   4,   2,   8,   5,   7};
   char *vals = new char[N]; // = {'a', 'b', 'c', 'd', 'e', 'f'};
   int  *addr = new int[N]; // = { 1,   2,   3,   4,   5,   6};
   keys[0]=1; keys[1]=4; keys[2]=2; keys[3]=8; keys[4]=5; keys[5]=7;
   vals[0]='a'; vals[1]='b'; vals[2]='c'; vals[3]='d'; vals[4]='e'; vals[5]='f';
   addr[0]=1; addr[1]=2; addr[2]=3; addr[3]=4; addr[4]=5; addr[5]=6;

   int *d_keys, *d_addr;
   char *d_vals;
   CUDA_SAFE(cudaMalloc((void **)&d_keys, N * sizeof(int)));
   CUDA_SAFE(cudaMalloc((void **)&d_addr, N * sizeof(int)));
   CUDA_SAFE(cudaMalloc((void **)&d_vals, N * sizeof(char)));
   CUDA_SAFE(cudaMemcpy(d_keys, keys, N * sizeof(int), cudaMemcpyHostToDevice));
   CUDA_SAFE(cudaMemcpy(d_vals, vals, N * sizeof(char), cudaMemcpyHostToDevice));
   CUDA_SAFE(cudaMemcpy(d_addr, addr, N * sizeof(int), cudaMemcpyHostToDevice));

   auto it = thrust::make_zip_iterator(thrust::make_tuple(d_vals, d_addr));

   thrust::sort_by_key(thrust::device, d_keys, d_keys+N, it);

   CUDA_SAFE(cudaMemcpy(keys, d_keys, N * sizeof(int), cudaMemcpyDeviceToHost));
   CUDA_SAFE(cudaMemcpy(vals, d_vals, N * sizeof(char), cudaMemcpyDeviceToHost));
   CUDA_SAFE(cudaMemcpy(addr, d_addr, N * sizeof(int), cudaMemcpyDeviceToHost));

   printf("Keys: "); for (int i = 0; i < N; i++) printf("%d ", keys[i]); printf("\n");
   printf("Vals: "); for (int i = 0; i < N; i++) printf("%c ", vals[i]); printf("\n");
   printf("Addr: "); for (int i = 0; i < N; i++) printf("%d ", addr[i]); printf("\n");
}

Zu Zwecken der Kürze habe ich die Kopfzeilen und Überprüfungen ausgelassen.