向量的 N 个最大元素及其索引

N largest elements of a vector along with their indices

我有一个thrust::device_vector <float> vec。假设 vec.size() = LN < L。我想在 vec 中找到最大的 N 元素及其索引。我们如何使用原始 CUDA 或 thrust 有效地做到这一点?

你应该结帐this question.

我喜欢 Ricky Bobby 的回答(如果 N 远小于 L)。

我还建议查看以下论文。 Fast K-selection Algorithm for Graphics Processing Units 作者:Alabi T 等人

它为K-selection提供了3种不同的并行算法。 他们描述的 bucket-select 表现最好。 该算法有两个步骤:

第1步(预处理将原始数据分割成大小小于2^21个元素的向量)

第 2 步:

1. Choose bucket containing Kth element
2. split the bucket again.
3. Repeat until the kth element is found (the min and max of the new bucket are equal).

现在您已经将其他 k-1 个最大元素划分到桶中。

此方法也称为分布式分区。

一个简单的解决方案是先对值进行排序,然后 select 最后 N 个元素。

以下示例 select 是 N=5 个最大元素及其来自 L=18 个值的原始索引。

编译使用

nvcc -std=c++11 nlargest.cu -o nlargest


输出时运行./nlargest

d_values:   1   2   3   4   5   6   7   8   9   4   5   6   7   8   9   0   1   2   
d_indices:  0   1   2   3   4   5   6   7   8   9   10  11  12  13  14  15  16  17  
d_values:   0   1   1   2   2   3   4   4   5   5   6   6   7   7   8   8   9   9   
d_indices:  15  0   16  1   17  2   3   9   4   10  5   11  6   12  7   13  8   14  
d_values_s: 7   8   8   9   9   
d_indices_s:12  7   13  8   14  

nlargest.cu

#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <thrust/sequence.h>
#include <iostream>

#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;
}

template<typename... Iterators>
__host__ __device__
thrust::zip_iterator<thrust::tuple<Iterators...>> zip(Iterators... its)
{
    return thrust::make_zip_iterator(thrust::make_tuple(its...));
}


int main()
{

    const int size = 18;
    const int select_size = 5;

    float values[size] = {1,2,3,
                          4,5,6,
                          7,8,9,
                          4,5,6,
                          7,8,9,
                          0,1,2
    };

    thrust::host_vector<float> h_values (values, values+size);
    thrust::device_vector<float> d_values = h_values;
    thrust::device_vector<int> d_indices(size);
    thrust::sequence(d_indices.begin(), d_indices.end());

    PRINTER(d_values);
    PRINTER(d_indices);
    thrust::sort(zip(d_values.begin(), d_indices.begin()),zip(d_values.end(), d_indices.end()));
    PRINTER(d_values);
    PRINTER(d_indices);

    thrust::device_vector<float> d_values_s(select_size);
    thrust::device_vector<int> d_indices_s(select_size);

    thrust::copy(zip(d_values.end()-select_size, d_indices.end()-select_size),
                zip(d_values.end(), d_indices.end()),
                zip(d_values_s.begin(), d_indices_s.begin())
                );
    PRINTER(d_values_s);
    PRINTER(d_indices_s);

    return 0;
}