查找 CUDA 数组中每个不同值的第一个索引答案

【问题标题】：Finding the first index of every distinct value in CUDA array查找 CUDA 数组中每个不同值的第一个索引
【发布时间】：2017-12-14 14:24:37
【问题描述】：

假设我们有一个这样的数组：

0, 0, 0, 1, 2, 2, 2, 3, 3, 4, ...

我想要每个值第一次出现的索引，所以在这个例子中是 [0, 3, 4, 7, 9]。数组已排序，所有可能的值都是已知且连续的。

我的可能解决方案是为这个数组中的每个元素使用一个内核，并使用 atomicmin 来保存最低索引。但我认为可能有更好的方法。

【问题讨论】：

将每个值与其前面的值进行比较，以确定它是否是第一次出现。
您的数组如何作为输入显示？如果它是从外部源（例如文件）读取的，那么您需要一个单独的内核。但是，如果它是另一个内核的输出，比如排序算法，那么最好在那个时候获得所需的输出，而不是在单独的步骤中。
这是多个推力调用的输出，因此需要单独的步骤

标签： cuda thrust

【解决方案1】：

如果您提供索引向量，例如，您可以通过一次调用 thrust::unique_by_key()通过thrust::sequence()。这是一个有效的例子：

$ cat t3.cu
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/unique.h>
#include <thrust/sequence.h>
#include <iostream>

int main(){

  int keys[] = {0, 0, 0, 1, 2, 2, 2, 3, 3, 4};
  int ks = sizeof(keys)/sizeof(keys[0]);
  thrust::device_vector<int> d_keys(keys, keys+ks);
  thrust::device_vector<int> d_result(ks);
  thrust::sequence(d_result.begin(), d_result.end());
  int rs  = (thrust::unique_by_key(d_keys.begin(), d_keys.end(), d_result.begin())).first - d_keys.begin();
  thrust::copy_n(d_result.begin(), rs, std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
}

$ nvcc -arch=sm_35 -o t3 t3.cu
$ ./t3
0,3,4,7,9,
$

这里发生的重要活动是流压缩，thrust 为各种用例提供了nice set of routines。例如，此操作也可以使用 thrust::unique_copy() 完成，在这种情况下，由于代码复杂性增加，您可以消除对 thrust::sequence() 调用的需要（它将替换为与您的数据一起压缩的 thrust::counting_iterator，和适当的选择函子），但它仍然需要相同长度的输出向量。

【讨论】：

不错！优雅的解决方案。
我们尝试了这两种解决方案，这个解决方案的速度更快

【解决方案2】：

正如@tera 所指出的，您可以将一个数字与前一个数字进行比较，以确定它是否是唯一数字序列中的第一个出现。您可以编写一个内核来为此标准生成掩码，以便掩码数组包含第一次出现的数字的索引和负数（如 -1，因为它不能是索引）。之后，使用推力通过谓词计算非 -1 值。然后使用与上述相同的谓词从掩码中复制这些值。最后，将结果复制回主机。

这是上述方法的示例实现。

#include <iostream>
#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/count.h>
#include <thrust/copy.h>

using namespace std;


//Copy index
__global__ void is_first_occurence(int* input, int* is, int count)
{
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;

    if(tid<count)
    {
        if(tid == 0)
        {
            is[0] = 0;
        }
        else if(input[tid] != input[tid-1])
        {
            is[tid] = tid;
        }
        else
        {
            is[tid] = -1;
        }
    }
}


struct isFirst
{
  __host__ __device__ bool operator()(const int x)
  {
    return (x != -1);
  }
};

int main(int argc, char** argv)
{
    const int count = 13;

    std::vector<int> arr = { 0, 0, 0, 1, 1, 2, 2, 2, 3, 3, 4, 4 ,4 };

    thrust::device_vector<int> arr_d = arr;
    thrust::device_vector<int> mask_d(arr_d.size());

    int* pArr = thrust::raw_pointer_cast(arr_d.data() );
    int* pMask = thrust::raw_pointer_cast(mask_d.data() );

    dim3 block(16);
    dim3 grid((count + block.x -1)/block.x);

    is_first_occurence<<<grid,block>>>(pArr, pMask, count);
    cudaDeviceSynchronize();

    int num_unique = thrust::count_if(mask_d.begin(), mask_d.end(), isFirst());

    thrust::copy_if(mask_d.begin(), mask_d.end(), arr_d.begin(), isFirst());

    std::vector<int> unique_indices(num_unique);

    thrust::copy(arr_d.begin(), arr_d.begin() + num_unique, unique_indices.begin());

    for(auto i:unique_indices)
    {
        cout<<i<<endl;
    }

    return 0;
}

使用以下命令编译和测试：

nvcc -o get_unique get_unique.cu -std=c++11 -arch=sm_61

【讨论】：