通过构造函数将 device_vector 存储在仿函数中？

Question

我正在尝试将 thrust::device_vector 存储在仿函数中。简单解读如下：

struct StructOperator : public thrust::unary_function<float, int>  {
  int num_;
  thrust::device_vector<int> v_test;

  explicit StructOperator(thrust::device_vector<int> const& input_v) :
    v_test(input_v), num_(input_v.size()) {};

  __host__ __device__
   float operator()(int index) {
      // magic happens
   }
};

无法编译 - nvcc 一直说从 __host__ __device__ 调用 __host__ 是不允许的。我看到 this 问题 - 这是实现此目标的唯一方法吗？

Answer 1

当您将 __device__ 装饰器放在仿函数运算符上时，您现在已将在该运算符主体中可以做的事情限制为与 CUDA 设备代码兼容的事情。

A thrust::device_vector 是一个 class 定义，旨在促进推力的 expression/computation 模型（大致类似于 STL containers/algorithms）。因此它包括主机和设备代码。 thrust::device_vector 中的主机代码未修饰以供在设备上使用，普通主机代码在 CUDA 设备代码中不可用。

thrust::device_vector 并非旨在也无意直接在设备代码中使用。它不能像你建议的那样使用。与可能的推测相反，它并非设计为可在 CUDA 设备代码中使用的 std::vector 的模拟。它被设计为可用于推力算法的 std::vector 的模拟（根据设计，来自主机代码的 callable/usable）。这就是您在编译时收到消息的原因，并且没有简单的方法 (*) 可以解决这个问题。

据推测，thrust::device_vector 的主要目的是作为一个容器来保存设备上 usable/accessible 的数据。 CUDA 设备代码已经支持的 POD 类型数据中最直接的等价物是数组或指向数据的指针。

因此我认为用 "yes" 来回答您的问题是合理的 - 这是实现此目的的唯一方法。

我将各种类似的方法归为一类，例如传递推力指针而不是裸指针。
(*)我忽略了这样的想法，例如编写允许在设备上使用的自己的容器 class，或者进行大量修改以推动自身以某种方式允许这种行为。

这是一个完整的示例，围绕您所展示的内容：

$ cat t1385.cu
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/copy.h>


struct StructOperator : public thrust::unary_function<float, int>  {
  int num_;
  int *v_test;

  explicit StructOperator(int *input_v, int input_v_size) :
    v_test(input_v), num_(input_v_size) {};

  __host__ __device__
   float operator()(int index) {
      if (index < num_)  return v_test[index] + 0.5;
      return 0.0f;
   }
};

const int ds = 3;
int main(){

  thrust::device_vector<int> d(ds);
  thrust::sequence(d.begin(), d.end());
  thrust::device_vector<float> r(ds);
  thrust::transform(d.begin(), d.end(), r.begin(), StructOperator(thrust::raw_pointer_cast(d.data()), d.size()));
  thrust::copy(r.begin(), r.end(), std::ostream_iterator<float>(std::cout, ","));
  std::cout << std::endl;
}
$ nvcc t1385.cu -o t1385
$ ./t1385
0.5,1.5,2.5,
$

通过构造函数将 device_vector 存储在仿函数中？

Storing a device_vector inside a functor through the constructor?

cuda

functor

thrust