CUDA / CUDA Thrust 中的多态性和派生类

Question

这是我在 Stack Overflow 上的第一个问题，这是一个很长的问题。thrust::device_vector<BaseClass>tl;dr 版本是：如果我希望 a 存储不同类型的对象，我该如何使用 aDerivedClass1，DerivedClass2a ？

\n\n

我想利用 CUDA Thrust 的多态性。我正在编译一个-arch=sm_30GPU (GeForce GTX 670) 进行编译。

\n\n

让我们看一下下面的问题：假设镇上有 80 户家庭。其中60户是已婚夫妇，20户是单亲家庭。因此，每个家庭都有不同数量的成员。现在是人口普查时间，家庭必须说明父母的年龄和孩子的数量。Family因此，政府构建了一个对象数组，thrust::device_vector<Family> familiesInTown(80)即家庭信息familiesInTown[0]对应familiesInTown[59]为已婚夫妇，其余信息（familiesInTown[60]为 ）familiesInTown[79] ) 为单亲家庭。

\n\n

\n
• Family是基类 - 家庭中父母的数量（单亲父母数量为 1，夫妻数量为 2）以及他们拥有的孩子数量作为成员存储在此处。
\n
• SingleParent，派生自Family，包括一个新成员 - 单亲父母的年龄，unsigned int ageOfParent。
\n

• MarriedCouple，也源自Family，但是，引入了两个新成员 - 父母的年龄，unsigned int ageOfParent1以及unsigned int ageOfParent2。

  \n\n

  #include <iostream>\n#include <stdio.h>\n#include <thrust/device_vector.h>\n\nclass Family\n{\nprotected:\n  unsigned int numParents;\n  unsigned int numChildren;\npublic:\n  __host__ __device__ Family() {};\n  __host__ __device__ Family(const unsigned int& nPars, const unsigned int& nChil) : numParents(nPars), numChildren(nChil) {};\n  __host__ __device__ virtual ~Family() {};\n\n  __host__ __device__ unsigned int showNumOfParents() {return numParents;}\n  __host__ __device__ unsigned int showNumOfChildren() {return numChildren;}\n};\n\nclass SingleParent : public Family\n{\nprotected:\n  unsigned int ageOfParent;\npublic:\n  __host__ __device__ SingleParent() {};\n  __host__ __device__ SingleParent(const unsigned int& nChil, const unsigned int& age) : Family(1, nChil), ageOfParent(age) {};\n\n  __host__ __device__ unsigned int showAgeOfParent() {return ageOfParent;}\n};\n\nclass MarriedCouple : public Family\n{\nprotected:\n  unsigned int ageOfParent1;\n  unsigned int ageOfParent2;\npublic:\n  __host__ __device__ MarriedCouple() {};\n  __host__ __device__ MarriedCouple(const unsigned int& nChil, const unsigned int& age1, const unsigned int& age2) : Family(2, nChil), ageOfParent1(age1), ageOfParent2(age2) {};\n\n  __host__ __device__ unsigned int showAgeOfParent1() {return ageOfParent1;}\n  __host__ __device__ unsigned int showAgeOfParent2() {return ageOfParent2;}\n};\n

\n

\n\n

如果我要 na\xc3\xafvely 启动我的对象中的对象thrust::device_vector<Family>使用以下函子启动 my 中的对象：

\n\n

struct initSlicedCouples : public thrust::unary_function<unsigned int, MarriedCouple>\n{\n  __device__ MarriedCouple operator()(const unsigned int& idx) const\n  // I use a thrust::counting_iterator to get idx\n  {\n    return MarriedCouple(idx % 3, 20 + idx, 19 + idx); \n    // Couple 0: Ages 20 and 19, no children\n    // Couple 1: Ages 21 and 20, 1 child\n    // Couple 2: Ages 22 and 21, 2 children\n    // Couple 3: Ages 23 and 22, no children\n    // etc\n  }\n};\n\nstruct initSlicedSingles : public thrust::unary_function<unsigned int, SingleParent>\n{\n  __device__ SingleParent operator()(const unsigned int& idx) const\n  {\n    return SingleParent(idx % 3, 25 + idx);\n  }\n};\n\nint main()\n{\n  unsigned int Num_couples = 60;\n  unsigned int Num_single_parents = 20;\n\n  thrust::device_vector<Family> familiesInTown(Num_couples + Num_single_parents);\n  // Families [0] to [59] are couples. Families [60] to [79] are single-parent households.\n  thrust::transform(thrust::counting_iterator<unsigned int>(0),\n                    thrust::counting_iterator<unsigned int>(Num_couples),\n                    familiesInTown.begin(),\n                    initSlicedCouples());\n  thrust::transform(thrust::counting_iterator<unsigned int>(Num_couples),\n                    thrust::counting_iterator<unsigned int>(Num_couples + Num_single_parents),\n                    familiesInTown.begin() + Num_couples,\n                    initSlicedSingles());\n  return 0;\n}\n

\n\n

我肯定会对一些经典的对象切片感到内疚感到内疚......

\n\n

所以，我问自己，一个可以给我一些甜蜜的多态性的指针向量怎么样？C++ 中的智能指针是一个东西，并且thrust迭代器可以做一些非常令人印象深刻的事情，所以让我们尝试一下，我想。以下代码可编译。

\n\n

struct initCouples : public thrust::unary_function<unsigned int, MarriedCouple*>\n{\n  __device__ MarriedCouple* operator()(const unsigned int& idx) const\n  {\n    return new MarriedCouple(idx % 3, 20 + idx, 19 + idx); // Memory issues?\n  }\n};\nstruct initSingles : public thrust::unary_function<unsigned int, SingleParent*>\n{\n  __device__ SingleParent* operator()(const unsigned int& idx) const\n  {\n    return new SingleParent(idx % 3, 25 + idx);\n  }\n};\n\nint main()\n{\n  unsigned int Num_couples = 60;\n  unsigned int Num_single_parents = 20;\n\n  thrust::device_vector<Family*> familiesInTown(Num_couples + Num_single_parents);\n  // Families [0] to [59] are couples. Families [60] to [79] are single-parent households.\n  thrust::transform(thrust::counting_iterator<unsigned int>(0),\n                    thrust::counting_iterator<unsigned int>(Num_couples),\n                    familiesInTown.begin(),\n                    initCouples()); \n  thrust::transform(thrust::counting_iterator<unsigned int>(Num_couples),\n                    thrust::counting_iterator<unsigned int>(Num_couples + Num_single_parents),\n                    familiesInTown.begin() + Num_couples,\n                    initSingles());\n\n  Family A = *(familiesInTown[2]); // Compiles, but object slicing takes place (in theory)\n  std::cout << A.showNumOfParents() << "\\n"; // Segmentation fault\n return 0;\n}\n

\n\n

看来我在这里碰壁了。我对内存管理的理解正确吗？(虚拟表等）。我的对象是否正在实例化并填充到设备上？我是否正在泄漏记忆，就像没有明天一样？

\n\n

无论如何，为了避免对象切片，我尝试使用dynamic_cast<DerivedPointer*>(basePointer). 这就是我创建Family析构函数的原因virtual。

\n\n

Family *pA = familiesInTown[2];\nMarriedCouple *pB = dynamic_cast<MarriedCouple*>(pA);\n

\n\n

以下几行可以编译，但不幸的是，再次抛出段错误。CUDA-Memcheck 不会告诉我原因。

\n\n

  std::cout << "Ages " << (pB -> showAgeOfParent1()) << ", " << (pB -> showAgeOfParent2()) << "\\n";\n

\n\n

和

\n\n

  MarriedCouple B = *pB;\n  std::cout << "Ages " << B.showAgeOfParent1() << ", " << B.showAgeOfParent2() << "\\n";\n

\n\n

简而言之，我需要的是一个类接口，用于具有不同属性、彼此之间具有不同数量成员的对象，但我可以将其存储在一个公共向量中（这就是我想要一个基类的原因）可以在GPU上进行操作。thrust我的目的是通过 ing 在转换和 CUDA 内核中使用它们thrust::raw_pointer_cast，这对我来说一直完美地工作，直到我需要将我的类分支为一个基础类和几个派生类。标准程序是什么？

\n\n

提前致谢！

\n

Answer 1

我完全同意@talonmies 的回答。（例如，我不知道推力已经通过多态性进行了广泛的测试。）此外，我还没有完全解析您的代码。我发布这个答案是为了添加额外的信息，特别是我相信某种程度的多态性可以与推力一起工作。

我要做的一个关键观察是，不允许将__global__具有虚函数的类的对象作为参数传递给函数。 这意味着在主机上创建的多态对象无法传递到设备（通过推力，或在普通 CUDA C++ 中）。（此限制的一个基础是对象中对虚拟函数表的要求，主机和设备之间的虚拟函数表必然有所不同，再加上在主机代码中直接获取设备函数的地址是非法的）。

然而，多态性可以在设备代码中发挥作用，包括推力设备函数。

下面的示例演示了这个想法，将我们限制为在设备上创建的对象，尽管我们当然可以使用主机数据初始化它们。我创建了两个类Triangle和，派生自包含虚函数 的Rectangle基类。 并从基类继承该函数但替换虚函数。PolygonareaTriangleRectangleset_valuesarea

然后我们可以多态地操作这些类的对象，如下所示：

#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/for_each.h>
#include <thrust/sequence.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/copy.h>
#define N 4


class Polygon {
  protected:
    int width, height;
  public:
  __host__ __device__  void set_values (int a, int b)
      { width=a; height=b; }
  __host__ __device__  virtual int area ()
      { return 0; }
};

class Rectangle: public Polygon {
  public:
  __host__ __device__  int area ()
      { return width * height; }
};

class Triangle: public Polygon {
  public:
  __host__ __device__   int area ()
      { return (width * height / 2); }
};


struct init_f {
  template <typename Tuple>
  __host__ __device__ void operator()(const Tuple &arg) {
    (thrust::get<0>(arg)).set_values(thrust::get<1>(arg), thrust::get<2>(arg));}
};

struct setup_f {
  template <typename Tuple>
  __host__ __device__ void operator()(const Tuple &arg) {
    if (thrust::get<0>(arg) == 0)
      thrust::get<1>(arg) = &(thrust::get<2>(arg));
    else
      thrust::get<1>(arg) = &(thrust::get<3>(arg));}
};

struct area_f {
  template <typename Tuple>
  __host__ __device__ void operator()(const Tuple &arg) {
    thrust::get<1>(arg) = (thrust::get<0>(arg))->area();}
};


int main () {

  thrust::device_vector<int>  widths(N);
  thrust::device_vector<int> heights(N);
  thrust::sequence( widths.begin(),  widths.end(), 2);
  thrust::sequence(heights.begin(), heights.end(), 3);
  thrust::device_vector<Rectangle> rects(N);
  thrust::device_vector<Triangle>  trgls(N);
  thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(rects.begin(), widths.begin(), heights.begin())), thrust::make_zip_iterator(thrust::make_tuple(rects.end(), widths.end(), heights.end())), init_f());
  thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(trgls.begin(), widths.begin(), heights.begin())), thrust::make_zip_iterator(thrust::make_tuple(trgls.end(), widths.end(), heights.end())), init_f());
  thrust::device_vector<Polygon *> polys(N);
  thrust::device_vector<int> selector(N);
  for (int i = 0; i<N; i++) selector[i] = i%2;
  thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(selector.begin(), polys.begin(), rects.begin(), trgls.begin())), thrust::make_zip_iterator(thrust::make_tuple(selector.end(), polys.end(), rects.end(), trgls.end())), setup_f());
  thrust::device_vector<int> areas(N);
  thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(polys.begin(), areas.begin())), thrust::make_zip_iterator(thrust::make_tuple(polys.end(), areas.end())), area_f());
  thrust::copy(areas.begin(), areas.end(), std::ostream_iterator<int>(std::cout, "\n"));
  return 0;
}

我建议为 cc2.0 或更新的架构编译上述代码。我在 RHEL 5.5 上使用 CUDA 6 进行了测试。

（多态示例想法和一些代码取自此处。）

Answer 2

我不会尝试回答这个问题的所有内容，它太大了。话虽如此，这里有一些关于您发布的代码的观察结果可能会有所帮助：

• GPU 端new运算符从私有运行时堆中分配内存。从 CUDA 6 开始，主机端 CUDA API 无法访问该内存。您可以从内核和设备函数内部访问内存，但主机无法访问该内存。因此，new在推力装置函子内部使用是一种损坏的设计，永远无法工作。这就是您的“指针向量”模型失败的原因。
• Thrust 的根本目的是允许典型 STL 算法的数据并行版本应用于 POD 类型。使用复杂的多态对象构建代码库并尝试通过 Thrust 容器和算法填充这些代码库可能会起作用，但这不是 Thrust 的设计目的，我不会推荐它。如果您以意想不到的方式破坏了推力，请不要感到惊讶。
• CUDA 支持许多 C++ 功能，但编译和对象模型甚至比它们所基于的 C++98 标准简单得多。CUDA 缺乏几个关键功能（例如 RTTI），而这些功能使得复杂的多态对象设计可以在 C++ 中使用。我的建议是谨慎使用 C++ 功能。仅仅因为您可以在 CUDA 中做某事并不意味着您应该这样做。GPU 是一个简单的架构，简单的数据结构和代码几乎总是比功能相似的复杂对象具有更高的性能。

浏览完您发布的代码后，我的总体建议是回到绘图板。如果您想了解一些非常优雅的 CUDA/C++ 设计，请花一些时间阅读CUB和CUSP的代码库。它们都有很大不同，但都有很多值得学习的地方（我怀疑 CUSP 是建立在 Thrust 之上的，这使得它与您的用例更加相关）。