Một thuật toán phân loại song song tốt được thực hiện với CUDA?

Question

Có ai biết thuật toán sắp xếp tốt cho GPU, được triển khai trong CUDA không?

Tôi đã đọc vài bài báo và một số mã mẫu của NVIDIA, nhưng không có gì giải thích thuật toán sắp xếp được sử dụng.

Answer 1

Nhiều triển khai phân loại GPU là các biến thể của bitonic sort, được biết đến và được mô tả trong hầu hết các văn bản hợp lý về thuật toán được xuất bản trong 25 hoặc 30 năm qua.

Các "tham khảo" sắp xếp thực hiện cho CUDA thực hiện bằng cách Nadathur Satish từ Berkeley và Mark Harris và Michael Garland từ NVIDIA (giấy here) là một radix sort, và hình thành nền tảng của những gì có trong NPP và Thrust.

Answer 2

Tính đến tháng 5 năm 2014, tôi muốn đề cập đến hai khả năng để thực hiện sắp xếp trong CUDA bằng cách sử dụng thư viện. Đầu tiên là sử dụng Thrust, như đã đề cập bởi @talonmies và thứ hai là CUB, trong đó có một số nguyên thủy "toàn bộ thiết bị".

Lực đẩy sử dụng phân loại gốc cho các loại dữ liệu cơ bản khi sử dụng các hoạt động so sánh đã xác định, xem Thrust: A Productivity-Oriented Library for CUDA. Ngược lại với điều đó, khi bạn cần sắp xếp các đối tượng bằng cách xác định các toán tử so sánh struct tùy chỉnh, Thrust sẽ chuyển sang phối hợp sắp xếp. Radix sắp xếp nhanh hơn đáng kể so với sắp xếp hợp nhất, do đó đề xuất sử dụng toán tử so sánh được tích hợp khi có thể.

Bên cạnh, tôi cũng muốn đề cập đến CUB có DeviceRadixSort cung cấp thiết bị rộng, hoạt động song song để tính toán sắp xếp trên một chuỗi các mục dữ liệu nằm trong bộ nhớ chung.

Dưới đây, tôi đang cung cấp ví dụ cho thấy việc sử dụng cả hai thư viện để thực hiện sắp xếp. Cần lưu ý rằng cub::DeviceRadixSort thực hiện "sắp xếp theo khóa", trong khi Thrust có cả hai thrust::sort và thrust::sort_by_key khả dụng. Tất nhiên, CUB cub::DeviceRadixSort có thể được sử dụng để phân loại đơn giản bằng cách tạo mảng khóa trùng với mảng giá trị. Các dòng nhận xét của mã bên dưới cung cấp phân loại đơn giản cho CUB hoặc sắp xếp theo khóa cho Lực đẩy.

// 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(); 

}