Dobry algorytm równoległego sortowania zaimplementowany z CUDA?

Question

Czy ktoś zna dobry algorytm sortowania dla GPU, zaimplementowany w CUDA?

Przeczytałem kilka artykułów i przykładowy kod NVIDIA, ale nic, co wyjaśnia zastosowany algorytm sortowania.

Answer 1

Wiele implementacji sortowania GPU to warianty bitonic sort, które są dość dobrze znane i opisane w najrozsądniejszych tekstach na temat algorytmów opublikowanych w ciągu ostatnich 25 lub 30 lat.

W „odniesienia” sortowanie realizację CUDA wykonaną przez Nadathur Satish z Berkeley i Mark Harris i Michael Garland firmy NVIDIA (papier here) jest radix sort, i stanowi podstawę tego, co jest w NPP i Thrust.

Answer 2

Od maja 2014 r. Chciałbym wymienić dwie możliwości sortowania w CUDA przy użyciu bibliotek. Pierwszy z nich to Thrust, jak wspomniano przez @talonmies, a drugi to CUB, który ma kilka prymitywów "w całym urządzeniu".

Metoda Thrust używa sortowania radix dla podstawowych typów danych podczas korzystania z już zdefiniowanych operacji porównania, patrz Thrust: A Productivity-Oriented Library for CUDA. W przeciwieństwie do tego, kiedy trzeba posortować obiekty przez zdefiniowanie spersonalizowanych operatorów porównania struktur, Thrust przełączy się na sortowanie scalone. Radix-sort jest zadziwiająco szybszy niż sortowanie scalone, więc jeśli to możliwe, zalecane jest używanie wbudowanych operatorów porównania.

Poza Thrust, chciałbym również wspomnieć CUB który ma DeviceRadixSort która zapewnia urządzenie szerokości, równoległe operacje obliczania sortowania radix całej sekwencji elementów danych przebywających w obrębie pamięci globalnej.

Poniżej zamieszczam przykład pokazujący użycie obu bibliotek do sortowania. Należy zauważyć, że cub::DeviceRadixSort wykonuje "sort-by-key", podczas gdy Thrust ma dostępne zarówno thrust::sort i thrust::sort_by_key. Oczywiście, CUB cub::DeviceRadixSort może być użyty do prostego sortowania, ponieważ tablica kluczy jest zbieżna z tablicą wartości. Skomentowane wiersze poniższego kodu zapewniają proste sortowanie dla CUB lub sortowanie według klucza dla Thrust.

// Ensure printing of CUDA runtime errors to console 
#define CUB_STDERR 

#include <stdio.h> 
#include <algorithm> // std::generate 

#include <cub/cub.cuh> // or equivalently <cub/device/device_scan.cuh> 
#include <thrust\device_vector.h> 
#include <thrust\host_vector.h> 
#include <thrust\sort.h> 

/********************/ 
/* CUDA ERROR CHECK */ 
/********************/ 
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

void main(void) 
{ 

    // Declare, allocate, and initialize device pointers for input and output 
    int num_items = 7; 

    thrust::device_vector<float> d_key_buf(num_items); 
    thrust::device_vector<float> d_key_alt_buf(num_items); 
    thrust::device_vector<float> d_value_buf(num_items); 
    thrust::device_vector<float> d_value_alt_buf(num_items); 

    thrust::host_vector<float> h_key(num_items); 
    thrust::host_vector<float> h_value(num_items); 

    std::generate(h_key.begin(), h_key.end(), rand); 
    std::generate(h_value.begin(), h_value.end(), rand); 

    d_key_buf = h_key; 
    d_key_alt_buf = d_key_buf; 
    d_value_buf = h_value; 
    d_value_alt_buf = d_value_buf; 
    //d_value_buf = h_key; 
    //d_value_alt_buf = d_key_buf; 

    /*******/ 
    /* CUB */ 
    /*******/ 

    // Create a set of DoubleBuffers to wrap pairs of device pointers 
    cub::DoubleBuffer<float> d_keys(thrust::raw_pointer_cast(d_key_buf.data()), thrust::raw_pointer_cast(d_key_alt_buf.data())); 
    cub::DoubleBuffer<float> d_values(thrust::raw_pointer_cast(d_value_buf.data()), thrust::raw_pointer_cast(d_value_alt_buf.data())); 

    // Determine temporary device storage requirements 
    void  *d_temp_storage = NULL; 
    size_t temp_storage_bytes = 0; 
    cub::DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items); 

    // Allocate temporary storage 
    gpuErrchk(cudaMalloc(&d_temp_storage, temp_storage_bytes)); 

    // Run sorting operation 
    cub::DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items); 

    printf("*** CUB ***\n"); 
    for (int i=0; i<num_items; i++) { 
     float val1 = d_key_buf[i]; 
     float val2 = d_value_buf[i]; 
     printf("%i %f %f\n",i,val1,val2); 
    } 
    printf("\n\n"); 

    /**********/ 
    /* THRUST */ 
    /**********/ 

    d_key_buf = h_key; 
    d_value_buf = h_value; 
    //thrust::sort_by_key(d_key_buf.begin(), d_key_buf.end(), d_value_buf.begin()); 
    thrust::sort(d_key_buf.begin(), d_key_buf.end()); 

    printf("*** THRUST ***\n"); 
    for (int i=0; i<num_items; i++) { 
     float val1 = d_key_buf[i]; 
    printf("%i %f\n",i,val1); 
     //float val2 = d_value_buf[i]; 
     //printf("%i %f %f\n",i,val1,val2); 
    } 
    printf("\n\n"); 

    cudaFree(d_temp_storage); 

    getchar(); 

}