cuda - 3种不同大小向量的推力复数变换

您好，我在C +中有此循环，我试图将其转换为推力，但没有得到相同的结果...
有任何想法吗？
谢谢你

C++代码

for (i=0;i<n;i++)
    for (j=0;j<n;j++)
      values[i]=values[i]+(binv[i*n+j]*d[j]);

推力代码

thrust::fill(values.begin(), values.end(), 0);
thrust::transform(make_zip_iterator(make_tuple(
                thrust::make_permutation_iterator(values.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexDivFunctor(n))),
                binv.begin(),
                thrust::make_permutation_iterator(d.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexModFunctor(n))))),
                make_zip_iterator(make_tuple(
                thrust::make_permutation_iterator(values.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexDivFunctor(n))) + n,
                binv.end(),
                thrust::make_permutation_iterator(d.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexModFunctor(n))) + n)),
                thrust::make_permutation_iterator(values.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexDivFunctor(n))),
                function1()
                );

推力函数

struct IndexDivFunctor: thrust::unary_function<int, int>
{
  int n;

  IndexDivFunctor(int n_) : n(n_) {}

  __host__ __device__
  int operator()(int idx)
  {
    return idx / n;
  }
};

struct IndexModFunctor: thrust::unary_function<int, int>
{
  int n;

  IndexModFunctor(int n_) : n(n_) {}

  __host__ __device__
  int operator()(int idx)
  {
    return idx % n;
  }
};


struct function1
{
  template <typename Tuple>
  __host__ __device__
  double operator()(Tuple v)
  {
    return thrust::get<0>(v) + thrust::get<1>(v) * thrust::get<2>(v);
  }
};

最佳答案

首先，一些一般性评论。你的循环

for (i=0;i<n;i++)
    for (j=0;j<n;j++)
      v[i]=v[i]+(B[i*n+j]*d[j]);

与标准BLAS gemv操作等效

矩阵以行主要顺序存储的位置。在设备上执行此操作的最佳方法是使用CUBLAS，而不是由推力图元构造的东西。

话虽这么说，您发布的推力代码绝对不可能做您的串行代码所能做的。您看到的错误不是浮点关联性的结果。根本上thrust::transform将提供的仿函数应用于输入迭代器的每个元素，并将结果存储在输出迭代器上。为了产生与您发布的循环相同的结果，thrust::transform调用将需要对您发布的fmad仿函数执行(n * n)个操作。显然不是。此外，不能保证thrust::transform将以对内存竞争无害的方式执行求和/归约操作。

正确的解决方案可能是这样的:

使用推力::变换来计算 B 和 d

元素的(n * n)乘积

使用推力::: reduce_by_key将乘积简化为部分和，产生 Bd

使用推力::变换将结果矩阵向量乘积添加到 v 中以产生最终结果。

在代码中，首先定义一个仿函数，如下所示:

struct functor
{
  template <typename Tuple>
  __host__ __device__
  double operator()(Tuple v)
  {
    return thrust::get<0>(v) * thrust::get<1>(v);
  }
};

然后执行以下操作以计算矩阵向量乘法

  typedef thrust::device_vector<int> iVec;
  typedef thrust::device_vector<double> dVec;

  typedef thrust::counting_iterator<int> countIt;
  typedef thrust::transform_iterator<IndexDivFunctor, countIt> columnIt;
  typedef thrust::transform_iterator<IndexModFunctor, countIt> rowIt;

  // Assuming the following allocations on the device
  dVec B(n*n), v(n), d(n);

  // transformation iterators mapping to vector rows and columns
  columnIt cv_begin = thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexDivFunctor(n));
  columnIt cv_end   = cv_begin + (n*n);

  rowIt rv_begin = thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexModFunctor(n));
  rowIt rv_end   = rv_begin + (n*n);

  dVec temp(n*n);
  thrust::transform(make_zip_iterator(
                      make_tuple(
                        B.begin(),
                        thrust::make_permutation_iterator(d.begin(),rv_begin) ) ),
                    make_zip_iterator(
                      make_tuple(
                        B.end(),
                        thrust::make_permutation_iterator(d.end(),rv_end) ) ),
                    temp.begin(),
                    functor());

  iVec outkey(n);
  dVec Bd(n);
  thrust::reduce_by_key(cv_begin, cv_end, temp.begin(), outkey.begin(), Bd.begin());
  thrust::transform(v.begin(), v.end(), Bd.begin(), v.begin(), thrust::plus<double>());

当然，与使用专门设计的矩阵矢量乘法代码(例如来自CUBLAS的dgemv)相比，这是一种极其低效的计算方法。

关于cuda - 3种不同大小向量的推力复数变换，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/7617160/