循环旋转GPU矢量？

Question

我有一个想要实现的算法，其中涉及

1. 坐标加法，
2. 坐标乘法，以及
3. 坐标循环旋转。

我的加法和乘法有点烦人（它们是算术模 n），但我想我可以移植标准算法，例如使用蒙哥马利算术，所以我现在不关注这个话题。

对于最后一步，我的意思是映射一个向量

[a, b, c, d] -> [b, c, d, a] -> [c, d, a, b] -> [d, a, b, c] -> [a,b,c,d]

，尽管一般来说我的向量将是~1000维，所以一般的循环旋转会更复杂一些。

完成我要求的第三步的最佳方法是什么？查看 Thrust 库，似乎最简单的方法就是使用

permutation_iterator

。这应该足以提供正确的实现，但考虑到我想要计算的排列非常简单，我想知道这是否是最快的选择。

另外，虽然这个问题被标记为

thrust

，但如果有一些明显更好的库/框架可供使用，我将不胜感激。

Answer 1

如果您使用变换迭代器来即时构建排列，它应该非常高效；当然比实际移动数据更快。

这应该有效：

#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/transform_iterator.h>

struct rotate_right : public thrust::unary_function<int,int>
{
  int rotate_by, end;

  rotate_right(int rotate_by, int end) noexcept
  : rotate_by(rotate_by), end(end)
  { assert(rotate_by >= 0 && rotate_by <= end); }

  __host__ __device__
  int operator()(int idx) const noexcept
  {
      idx += rotate_by;
      if(idx >= end)
          idx -= end;
      return idx;
  }
};

template<class Iterator>
auto make_rotate_right(Iterator unrotated, int rotate_by, int len)
{
  using counting_iterator = thrust::counting_iterator<int>;
  counting_iterator base_index(0);
  using index_iterator = thrust::transform_iterator<rotate_right, counting_iterator>;
  index_iterator permutation{base_index, rotate_right{rotate_by, len}};
  return thrust::make_permutation_iterator(unrotated, permutation);
}

thrust::device_vector<float> foo()
{
  int n = 4;
  thrust::device_vector<float> v(n);
  v[0] = 1.0f;
  v[1] = 4.0f;
  v[2] = 9.0f;
  v[3] = 16.0f;
  auto rotated = make_rotate_right(v.begin(), 2, n);
  return {rotated, rotated + n};
}

int main()
{
  auto f = foo();
  for(auto fi: f)
    std::cout << fi << '\n';
}

打印

9
16
1
4