在推力函子中调用推力算法

Question

我在仿函数内部使用 thrust::reduce，它是 thrust::transform_reduce 中的参数。这种情况看起来像嵌套推力算法。编译成功但运行出错：

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  cudaEventSynchronize in future::wait: an illegal memory access was encountered
Aborted (core dumped)

代码如下：

#include <thrust/inner_product.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>

#include <iostream>
#include <cmath>
#include <boost/concept_check.hpp>


struct aFuntor : public thrust::unary_function<int, int>
{
    aFuntor(int* av__, int* bv__, const int& N__) : av_(av__), bv_(bv__), N_(N__) {};

    __host__ __device__
    int operator()(const int& idx)
    {

    thrust::device_ptr<int> av_dpt = thrust::device_pointer_cast(av_);

    int res = thrust::reduce(av_dpt, av_dpt+N_);

        return res;
    }

    int* av_;
    int* bv_;
    int N_;
};


int main(void)
{
      int N = 5;
      std::vector<int> av = {0,1,3,5};
      std::vector<int> bv = {0,10,20,30};
      thrust::device_vector<int> av_d(N);
      thrust::device_vector<int> bv_d(N);
      av_d = av; bv_d = bv;

      // initial value of the reduction
      int init=0;

      // binary operations
      thrust::plus<int>        bin_op;

      int res =
      thrust::transform_reduce(thrust::counting_iterator<int>(0),
                               thrust::counting_iterator<int>(N-1),
                   aFuntor(thrust::raw_pointer_cast(av_d.data()), 
                      thrust::raw_pointer_cast(bv_d.data()),
                      N),
                init,
                bin_op);    

      std::cout << "result is: " << res << std::endl;
      return 0;
}

推力支持这种嵌套结构吗？或者除了必须重新设计我的算法之外别无他法？据我所知，有些算法很难公开并行性？

提前致谢！

Answer 1

推力允许。但是，有必要确保推力在从设备代码启动推力算法时仅选择设备路径，而在您的情况下，这不会发生。至少在我的系统 (Ubuntu 14.04) 上，当我按原样编译你的代码时，我得到一个指示：

t113.cu(20) (col. 9): warning: calling a __host__ function("thrust::reduce< ::thrust::device_ptr<int> > ") from a __host__ __device__ function("aFuntor::operator ()") is not allowed

所以这显然不是这里想要的。相反，我们可以强制 thrust 使用设备路径（在设备代码中 - 这基本上隐含在你的仿函数定义中，因为你正在传递设备指针）和 thrust 执行策略 thrust::device。当我进行以下更改时，您的代码对我来说编译和运行没有错误：

$ cat t113.cu
#include <thrust/inner_product.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>

#include <iostream>
#include <cmath>
#include <thrust/execution_policy.h>
//#include <boost/concept_check.hpp>


struct aFuntor : public thrust::unary_function<int, int>
{
    aFuntor(int* av__, int* bv__, const int& N__) : av_(av__), bv_(bv__), N_(N__) {};

    __host__ __device__
    int operator()(const int& idx)
    {

    thrust::device_ptr<int> av_dpt = thrust::device_pointer_cast(av_);

    int res = thrust::reduce(thrust::device, av_dpt, av_dpt+N_);

        return res;
    }

    int* av_;
    int* bv_;
    int N_;
};


int main(void)
{
      int N = 5;
      std::vector<int> av = {0,1,3,5};
      std::vector<int> bv = {0,10,20,30};
      thrust::device_vector<int> av_d(N);
      thrust::device_vector<int> bv_d(N);
      av_d = av; bv_d = bv;

      // initial value of the reduction
      int init=0;

      // binary operations
      thrust::plus<int>        bin_op;

      int res =
      thrust::transform_reduce(thrust::counting_iterator<int>(0),
                               thrust::counting_iterator<int>(N-1),
                   aFuntor(thrust::raw_pointer_cast(av_d.data()),
                      thrust::raw_pointer_cast(bv_d.data()),
                      N),
                init,
                bin_op);

      std::cout << "result is: " << res << std::endl;
      return 0;
}
$ nvcc -std=c++11 -arch=sm_61  -o t113 t113.cu
$ ./t113
result is: 36
$

我实际上并没有尝试从代码中解析你的意图，所以我不能肯定地说这是正确的答案，但这似乎不是你要问的问题。（稍后：答案似乎是正确的。您的仿函数只是为每个元素生成 9 的值，并且您在 4 个元素 9x4=36 中减少 9）。

说了这么多，（对我来说）还不完全清楚为什么 thrust 在你原来的情况下选择主机路径。如果您愿意，可以为此提交 thrust issue。但完全有可能我没有足够仔细地考虑推力分配系统。主机代码算法分派 (transform_reduce) 可能有点令人困惑，因为例如，您使用的是主机容器还是设备容器可能并不明显。

在推力函子中调用推力算法

calling thrust algorithms inside a thrust functor

cuda

gpu

thrust