獲取CUDA中多個數組的唯一元素

Question

問題： 數組數量很多，例如2000個數組，但每個數組中只有256個整數。整數的範圍非常可觀，例如[0,1000000]。

我想獲得每個數組的唯一元素，換句話說，刪除重複的元素。 我有2個解決方案：

1. 使用推力讓每一個數組的獨特元素，所以我必須做2000次thrust::unique。但是每個陣列都很小，這種方式可能無法獲得良好的性能。

2. 在cuda內核中實現哈希表，使用2000個塊，每個塊中有256個線程。並利用共享內存來實現哈希表，然後每一個塊都會產生一個元素唯一的數組。

以上兩種方法看起來不專業，有沒有優雅的方式來解決CUDA的問題？

Answer 1

您可以使用thrust::unique如果您修改數據類似像它在這太問題做：Segmented Sort with CUDPP/Thrust

爲了簡化起見，我們假設每個數組包含per_array元素，有一個總array_num陣列。每個元素的範圍是[0,max_element]。

演示data與per_array=4，array_num=3和max_element=2可能看起來像這樣：

data = {1,0,1,2},{2,2,0,0},{0,0,0,0} 

爲了表示各元素的成員資格，我們使用相應陣列以下flags：

flags = {0,0,0,0},{1 1 1 1},{2,2,2,2} 

爲了爲了獲得每個分段數據集陣列的獨特元素，我們需要執行以下步驟：

個

1. 變換data所以每個陣列i的元素是獨特的範圍內[i*2*max_element,i*2*max_element+max_element]

   data = data + flags*2*max_element 
   data = {1,0,1,2},{6,6,4,4},{8,8,8,8} 
   

2. 排序變換後的數據：

   data = {0,0,1,2},{4,4,6,6},{8,8,8,8} 
   

3. 使用data作爲鍵和flags作爲應用thrust::unique_by_key值：

   data = {0,1,2}{4,6}{8} 
   flags = {0,0,0}{1,1}{2} 
   

4. 變換data回原始值：

   data = data - flags*2*max_element 
   data = {0,1,2}{0,2}{0} 
   

的max_element最大值由整數的用於表示data大小的限制。如果是與n位無符號整數：

max_max_element(n,array_num) = 2^n/(2*(array_num-1)+1) 

鑑於你array_num=2000，你會得到32位和64位無符號整數以下限制：

max_max_element(32,2000) = 1074010 
max_max_element(64,2000) = 4612839228234447 

下面的代碼實現上述步驟：

unique_per_array.cu

#include <thrust/device_vector.h> 
#include <thrust/extrema.h> 
#include <thrust/transform.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/functional.h> 
#include <thrust/sort.h> 
#include <thrust/unique.h> 
#include <thrust/copy.h> 

#include <iostream> 
#include <cstdint> 

#define PRINTER(name) print(#name, (name)) 
template <template <typename...> class V, typename T, typename ...Args> 
void print(const char* name, const V<T,Args...> & v) 
{ 
    std::cout << name << ":\t"; 
    thrust::copy(v.begin(), v.end(), std::ostream_iterator<T>(std::cout, "\t")); 
    std::cout << std::endl; 
} 

int main() 
{ 
    typedef uint32_t Integer; 

    const std::size_t per_array = 4; 
    const std::size_t array_num = 3; 

    const std::size_t total_count = array_num * per_array; 

    Integer demo_data[] = {1,0,1,2,2,2,0,0,0,0,0,0}; 

    thrust::device_vector<Integer> data(demo_data, demo_data+total_count);  

    PRINTER(data); 

    // if max_element is known for your problem, 
    // you don't need the following operation 
    Integer max_element = *(thrust::max_element(data.begin(), data.end())); 
    std::cout << "max_element=" << max_element << std::endl; 

    using namespace thrust::placeholders; 

    // create the flags 

    // could be a smaller integer type as well 
    thrust::device_vector<uint32_t> flags(total_count); 

    thrust::counting_iterator<uint32_t> flags_cit(0); 

    thrust::transform(flags_cit, 
         flags_cit + total_count, 
         flags.begin(), 
         _1/per_array); 
    PRINTER(flags); 


    // 1. transform data into unique ranges 
    thrust::transform(data.begin(), 
         data.end(), 
         thrust::counting_iterator<Integer>(0), 
         data.begin(), 
         _1 + (_2/per_array)*2*max_element); 
    PRINTER(data); 

    // 2. sort the transformed data 
    thrust::sort(data.begin(), data.end()); 
    PRINTER(data); 

    // 3. eliminate duplicates per array 
    auto new_end = thrust::unique_by_key(data.begin(), 
             data.end(), 
             flags.begin()); 

    uint32_t new_size = new_end.first - data.begin(); 
    data.resize(new_size); 
    flags.resize(new_size); 

    PRINTER(data); 
    PRINTER(flags); 

    // 4. transform data back 
    thrust::transform(data.begin(), 
         data.end(), 
         flags.begin(), 
         data.begin(), 
         _1 - _2*2*max_element); 

    PRINTER(data); 

}  

編譯和運行率：

$ nvcc -std=c++11 unique_per_array.cu -o unique_per_array && ./unique_per_array 

data: 1 0 1 2 2 2 0 0 0 0 0 0 
max_element=2 
flags: 0 0 0 0 1 1 1 1 2 2 2 2 
data: 1 0 1 2 6 6 4 4 8 8 8 8 
data: 0 1 1 2 4 4 6 6 8 8 8 8 
data: 0 1 2 4 6 8 
flags: 0 0 0 1 1 2 
data: 0 1 2 0 2 0 

一兩件事：

在thrust development version存在對thrust::unique*其improves performance by around 25 %實施的改善。如果您希望獲得更好的性能，您可能需要試用此版本。