如何在 CUDA 中从稀疏数组表示形式变为密集数组表示形式

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

本文关键字：数组 表示 CUDA      更新时间：2023-10-16

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

我有一个值数组。我想构建一个数组，它本质上是每个值在数组中出现的次数。有没有一种有效的算法来使用 CUDA 来做到这一点？

例如，假设值的范围为 0-10。实际上，我也有负值，但我需要忽略这些。（编辑：我已经尝试了thrust：：remove_if，然后是thrust：：reduce_by_key，但我正在寻找在忽略我不关心的元素方面更有效的东西。几乎就像推力：：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 上做这件事。我尝试使用推力对索引列表进行排序以提高内存缓存性能，但那里的性能改进充其量是微不足道的。

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

您可以修改推力直方图示例，以便在排序后构建直方图时仅考虑非负值：

#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

thrust：：remove_if 后跟 thrust：：reduce_by_key 最终成为为我的应用程序执行此操作的最佳方式。

相关文章：