从 CUDA 数组中删除元素

Question

上下文

我正在尝试使用 CUDA 实现 Boruvka MST 算法，但在这里完全不需要理解该算法来帮助我。

问题

所以，让我来描述一下这个问题：我有一个图形，以边列表格式存储（边数组，每条边用 2 个相邻的顶点 ID 及其权重表示）。

但是，（这非常重要！）为了优化对设备内存的访问，我不是将边缘存储为单个结构数组，而是存储为三个独立的数组：

• int *src_ids
• int *dst_ids
• 浮动*权重

要访问单个边，只需使用相同的索引迭代此数组：

• src_ids[i]
• dst_ids[i]
• 权重[i]

现在，当我描述了数据格式后，问题就来了

我想从这三个数组中删除具有以下条件的元素（边）：

1)如果src_id[i] == dst_id[i]，则删除第i条边

2) if src_id[i] != dst_id[i], 但还有一条边 j 与 src_id[j] 和 dst_id[j] 相同，但权重 [j] 更小, 而不是删除第 i 条边

换句话说，我想：

• 去除边缘
• 连接相同的顶点
• 并复制非最小边

第一个很简单：我可以使用thrust::remove_if 或按照此处parallel removal of elements from an array 的描述进行扫描，以删除具有相同ID 的边缘。 （我已经使用扫描实现了第二个变体）。

但我未能实现第二部分，即删除重复边缘。我有一个想法，但不确定这种方法是否有效。让我描述一下。

首先，我们将按以下方式重新排序（或排序）这三个数组：

• 首先按第一个 (src_id) ID 对边进行排序。
• 然后将第一个 ID 相同的边按第二个 (dst_id) id 排序。
• 最后，按权重对具有相同 id 的边进行排序。

当所有边都按这种方式排序时，删除重复的非最小边相对容易：

• 我们只看前面的边 (i-1)-th，
• 如果它具有相同的 ID，但重量更轻，则标记当前边缘以进行移除。
• 然后我们只需要应用扫描，类似于 (1 src_id[i] == dst_id[i]) 条件。

问题*

但问题是我不知道如何以这种方式有效地对三个数组进行排序。 （可能我可以对转换后的数据、单个结构数组使用thrust::sort，但它似乎会很慢，最好不要删除重复的边）

或者也许有人可以建议更好的方法来删除重复的边缘而不用这种方式排序。

感谢您阅读本文，感谢您的任何建议！

Answer 1

您可以使用thrust::zip_iterator 在一个thrust::sort 调用中轻松排序多个向量。

主要思想是：

auto z = thrust::make_zip_iterator(thrust::make_tuple(d_src_ids.begin(),d_dst_ids.begin(), d_weights.begin()));
thrust::sort(z,z+N);

这将首先按第一个矢量对三个矢量进行排序，然后按第二个，然后按第三个。

以下代码显示了如何在一个完整的示例中使用它。它使用自定义函子（从 thrust::detail 复制）在单个调用中执行 remove_if 步骤，而无需存储中间结果。

#include <thrust/sort.h>
#include <thrust/iterator/zip_iterator.h>
#include <iostream>
#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <thrust/remove.h>

#define PRINTER(name) print(#name, (name))
template <template <typename...> class V, typename T, typename ...Args>
void print(const char* name, const V<T,Args...> & v)
{
    std::cout << name << ":\t";
    thrust::copy(v.begin(), v.end(), std::ostream_iterator<T>(std::cout, "\t"));
    std::cout << std::endl;
}

// copied from https://github.com/thrust/thrust/blob/master/thrust/detail/range/head_flags.h

#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/tuple.h>
#include <thrust/functional.h>

template<typename RandomAccessIterator,
         typename BinaryPredicate = thrust::equal_to<typename thrust::iterator_value<RandomAccessIterator>::type>,
         typename ValueType = bool,
         typename IndexType = typename thrust::iterator_difference<RandomAccessIterator>::type>
  class head_flags
{
  // XXX WAR cudafe issue
  //private:
  public:
    struct head_flag_functor
    {
      BinaryPredicate binary_pred; // this must be the first member for performance reasons
      IndexType n;

      typedef ValueType result_type;

      __host__ __device__
      head_flag_functor(IndexType n)
        : binary_pred(), n(n)
      {}

      __host__ __device__
      head_flag_functor(IndexType n, BinaryPredicate binary_pred)
        : binary_pred(binary_pred), n(n)
      {}

      template<typename Tuple>
      __host__ __device__ __thrust_forceinline__
      result_type operator()(const Tuple &t)
      {
        const IndexType i = thrust::get<0>(t);

        // note that we do not dereference the tuple's 2nd element when i <= 0
        // and therefore do not dereference a bad location at the boundary
        return (i == 0 || !binary_pred(thrust::get<1>(t), thrust::get<2>(t)));
      }
    };

    typedef thrust::counting_iterator<IndexType> counting_iterator;

  public:
    typedef thrust::transform_iterator<
      head_flag_functor,
      thrust::zip_iterator<thrust::tuple<counting_iterator,RandomAccessIterator,RandomAccessIterator> >
    > iterator;

    __host__ __device__
    head_flags(RandomAccessIterator first, RandomAccessIterator last)
      : m_begin(thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator<IndexType>(0), first, first - 1)),
                                                head_flag_functor(last - first))),
        m_end(m_begin + (last - first))
    {}

    __host__ __device__
    head_flags(RandomAccessIterator first, RandomAccessIterator last, BinaryPredicate binary_pred)
      : m_begin(thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator<IndexType>(0), first, first - 1)),
                                                head_flag_functor(last - first, binary_pred))),
        m_end(m_begin + (last - first))
    {}

    __host__ __device__
    iterator begin() const
    {
      return m_begin;
    }

    __host__ __device__
    iterator end() const
    {
      return m_end;
    }

    template<typename OtherIndex>
    __host__ __device__
    typename iterator::reference operator[](OtherIndex i)
    {
      return *(begin() + i);
    }

  private:
    iterator m_begin, m_end;
};

template<typename RandomAccessIterator>
__host__ __device__
head_flags<RandomAccessIterator>
  make_head_flags(RandomAccessIterator first, RandomAccessIterator last)
{
  return head_flags<RandomAccessIterator>(first, last);
}

int main()
{
  const int N = 6;
  int src_ids[] = {3,1,2,2,3,3};
  int dst_ids[] = {2,2,3,3,1,1};
  float weights[] = {1,2,8,4,5,6};

  thrust::device_vector<int> d_src_ids(src_ids,src_ids+N);
  thrust::device_vector<int> d_dst_ids(dst_ids,dst_ids+N);
  thrust::device_vector<float> d_weights(weights,weights+N); 

  std::cout << "--- initial values ---" << std::endl;
  PRINTER(d_src_ids);
  PRINTER(d_dst_ids);
  PRINTER(d_weights);


  auto z = thrust::make_zip_iterator(thrust::make_tuple(d_src_ids.begin(),d_dst_ids.begin(), d_weights.begin()));

  thrust::sort(z,z+N);

  std::cout << "--- after sort ---" << std::endl;
  PRINTER(d_src_ids);
  PRINTER(d_dst_ids);
  PRINTER(d_weights);

  auto z2 = thrust::make_zip_iterator(thrust::make_tuple(d_src_ids.begin(),d_dst_ids.begin()));


  auto t = make_head_flags(z2,z2+N);
  using namespace thrust::placeholders;
  auto end = thrust::remove_if(z,z+N, t.begin(), !_1);
  int new_size = thrust::get<0>(end.get_iterator_tuple()) - d_src_ids.begin();

  d_src_ids.resize(new_size);
  d_dst_ids.resize(new_size);
  d_weights.resize(new_size);

  std::cout << "--- after remove_if ---" << std::endl;
  PRINTER(d_src_ids);
  PRINTER(d_dst_ids);
  PRINTER(d_weights);


  return 0;
}

输出：

--- initial values ---
d_src_ids:  3   1   2   2   3   3   
d_dst_ids:  2   2   3   3   1   1   
d_weights:  1   2   8   4   5   6   
--- after sort ---
d_src_ids:  1   2   2   3   3   3   
d_dst_ids:  2   3   3   1   1   2   
d_weights:  2   4   8   5   6   1   
--- after remove_if ---
d_src_ids:  1   2   3   3   
d_dst_ids:  2   3   1   2   
d_weights:  2   4   5   1