Un bon algorithme de tri parallèle implémenté avec CUDA?

Question

Est-ce que quelqu'un connaît un bon algorithme de tri pour les GPU, implémenté dans CUDA?

J'ai lu peu de papiers et quelques exemples de code NVIDIA, mais rien qui explique l'algorithme de tri utilisé.

Answer 1

De nombreuses implémentations de tri GPU sont des variantes du bitonic sort, ce qui est assez bien connu et décrit dans la plupart des textes raisonnables sur les algorithmes publiés au cours des 25 ou 30 dernières années.

La « référence » pour la mise en œuvre de tri CUDA fait par Nadathur Satish de Berkeley et Mark Harris et Michael Garland de NVIDIA (papier here) est un radix sort, et constitue la base de ce qui est dans NPP et Thrust.

Answer 2

Depuis mai 2014, je voudrais mentionner deux possibilités pour effectuer le tri dans CUDA en utilisant des bibliothèques. La première utilise Thrust, comme mentionné par @talonmies, et la seconde est CUB, qui a plusieurs primitives "à l'échelle du périphérique". Thrust utilise le tri par base pour les types de données fondamentaux lors de l'utilisation des opérations de comparaison déjà définies, voir Thrust: A Productivity-Oriented Library for CUDA. À l'opposé de cela, lorsque vous devez trier des objets en définissant des opérateurs de comparaison de structure personnalisés, Thrust passe à fusionner-trier. Le tri radix est remarquablement plus rapide que le tri par fusion, de sorte que la recommandation utilise des opérateurs de comparaison intégrés lorsque cela est possible.

Outre Thrust, je voudrais également mentionner CUB qui a DeviceRadixSort dispositif qui fournit, à l'échelle des opérations en parallèle pour le calcul d'un tri de base sur une séquence d'éléments de données résidant dans la mémoire globale. Ci-dessous, je donne un exemple montrant l'utilisation des deux bibliothèques pour effectuer le tri. Il convient de noter que cub::DeviceRadixSort effectue une "tri par clé", tandis que Thrust a à la fois thrust::sort et thrust::sort_by_key disponibles. Bien sûr, CUB cub::DeviceRadixSort peut être utilisé pour un tri simple en faisant coïncider le tableau de clés avec le tableau de valeurs. Les lignes commentées du code ci-dessous fournissent un tri simple pour CUB ou un tri par clé pour 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(); 

}