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

Question

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

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

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

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 以在排序后构建直方图时仅考虑 non-negative 值：

#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

Answer 2

thrust::remove_if 后跟 thrust::reduce_by_key 最终成为我的应用程序的最佳方式。

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

How to go from a sparse array representation to a dense one in CUDA

c++

arrays

cuda

thrust