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

Question

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

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

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

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

Answer 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::unique_copy(),在这种情况下,由于一些额外的代码复杂性,你可以消除对thrust::sequence()调用的需要(它将被替换为thrust::counting_iterator与你的数据一起压缩,以及一个合适的选择函子),但它仍然需要相同长度的输出向量.