c++ - 如何在 CUDA 中从稀疏数组表示到密集数组

Question

我是 CUDA 的初学者，正在尝试找出最有效的方法来做某事。

我有一组值。我想构建一个数组，它本质上是每个值在数组中出现的次数。是否有使用 CUDA 执行此操作的有效算法？

例如，我们假设值的范围是 0-10。实际上我也有负值，但我需要忽略这些。 (编辑:我试过 thrust::remove_if，然后是 thrust::reduce_by_key，但我正在寻找在忽略我不关心的元素方面更有效的东西。几乎像 thrust::reduce_by_key_if)。该列表比值的范围小得多(即，绝大多数值都在我关心的范围之外)。我可能有:

int32_t values[5] = {3, 5, 2, 5, 1, -1};

我想构建数组:

int32_t result[10] = {0, 1, 1, 1, 0, 2, 0, 0, 0, 0};

现在我主要在 CPU 上完成。我尝试使用 thrust 对索引列表进行排序以提高内存缓存性能，但性能改进充其量是微不足道的。

有什么想法吗？有没有一种优雅的方式来做到这一点？

Answer 1

您可以修改 thrust histogram example在排序后构建直方图时只考虑非负值:

#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <thrust/adjacent_difference.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/find.h>
#include <thrust/functional.h>
#include <thrust/binary_search.h>

#include <iostream>
#include <iomanip>
#include <iterator>

#include <cstdint>

template <typename Vector>
void print_vector(const std::string& name, const Vector& v)
{
  typedef typename Vector::value_type T;
  std::cout << "  " << std::setw(20) << name << "  ";
  thrust::copy(v.begin(), v.end(), std::ostream_iterator<T>(std::cout, " "));
  std::cout << std::endl;
}

int main(void)
{

  const std::int32_t N = 6;

  std::int32_t values[6] = {3, 5, 2, 5, 1, -1};

  thrust::device_vector<std::int32_t> d_values(values, values + N);

  print_vector("initial data", d_values);

  // sort values to bring equal elements together
  thrust::sort(d_values.begin(), d_values.end());

  print_vector("sorted data", d_values);

  using thrust::placeholders::_1;
  auto first_non_negative = thrust::find_if(d_values.begin(), d_values.end(), _1>=0);

  // number of histogram bins is equal to the maximum value plus one
  std::int32_t num_bins = d_values.back() + 1;

  thrust::device_vector<std::int32_t> histogram(num_bins);

  thrust::counting_iterator<std::int32_t> search_begin(0);
  thrust::upper_bound(first_non_negative, d_values.end(),
                      search_begin, search_begin + num_bins,
                      histogram.begin());

  thrust::adjacent_difference(histogram.begin(), histogram.end(),
                              histogram.begin());

  print_vector("histogram", histogram);

  return 0;
}

输出

initial data  3 5 2 5 1 -1 
sorted data  -1 1 2 3 5 5 
histogram     0 1 1 1 0 2