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

Question

这是我关于 Stack Overflow 的第一个问题，这是一个很长的问题。 tl;dr 版本是:我如何使用 thrust::device_vector<BaseClass>如果我希望它存储不同类型的对象 DerivedClass1 , DerivedClass2等，同时？

我想利用 CUDA Thrust 的多态性。我正在编译 -arch=sm_30 GPU(GeForce GTX 670)。

让我们看看下面的问题:假设镇上有 80 个家庭。其中60人为已婚夫妇，20人为单亲家庭。因此，每个家庭都有不同数量的成员。现在是人口普查时间，家庭必须说明 parent 的年龄和他们有多少 child 。因此，Family 的数组对象由政府 build ，即thrust::device_vector<Family> familiesInTown(80) , 这样的家庭信息 familiesInTown[0]至familiesInTown[59]对应已婚夫妇，其余(familiesInTown[60] 至 familiesInTown[79])为单亲家庭。

Family是基类 - 家庭中 parent 的数量(单亲 parent 为 1，夫妻为 2)和他们拥有的 child 的数量作为成员存储在这里。

SingleParent , 源自 Family ，包括一个新成员 - 单亲的年龄，unsigned int ageOfParent .

MarriedCouple , 也源自 Family然而，引入了两个新成员 - 双方 parent 的年龄，unsigned int ageOfParent1和 unsigned int ageOfParent2 .

#include <iostream>
#include <stdio.h>
#include <thrust/device_vector.h>

class Family
{
protected:
  unsigned int numParents;
  unsigned int numChildren;
public:
  __host__ __device__ Family() {};
  __host__ __device__ Family(const unsigned int& nPars, const unsigned int& nChil) : numParents(nPars), numChildren(nChil) {};
  __host__ __device__ virtual ~Family() {};

  __host__ __device__ unsigned int showNumOfParents() {return numParents;}
  __host__ __device__ unsigned int showNumOfChildren() {return numChildren;}
};

class SingleParent : public Family
{
protected:
  unsigned int ageOfParent;
public:
  __host__ __device__ SingleParent() {};
  __host__ __device__ SingleParent(const unsigned int& nChil, const unsigned int& age) : Family(1, nChil), ageOfParent(age) {};

  __host__ __device__ unsigned int showAgeOfParent() {return ageOfParent;}
};

class MarriedCouple : public Family
{
protected:
  unsigned int ageOfParent1;
  unsigned int ageOfParent2;
public:
  __host__ __device__ MarriedCouple() {};
  __host__ __device__ MarriedCouple(const unsigned int& nChil, const unsigned int& age1, const unsigned int& age2) : Family(2, nChil), ageOfParent1(age1), ageOfParent2(age2) {};

  __host__ __device__ unsigned int showAgeOfParent1() {return ageOfParent1;}
  __host__ __device__ unsigned int showAgeOfParent2() {return ageOfParent2;}
};

如果我天真地启动我的 thrust::device_vector<Family> 中的对象使用以下仿函数:

struct initSlicedCouples : public thrust::unary_function<unsigned int, MarriedCouple>
{
  __device__ MarriedCouple operator()(const unsigned int& idx) const
  // I use a thrust::counting_iterator to get idx
  {
    return MarriedCouple(idx % 3, 20 + idx, 19 + idx); 
    // Couple 0: Ages 20 and 19, no children
    // Couple 1: Ages 21 and 20, 1 child
    // Couple 2: Ages 22 and 21, 2 children
    // Couple 3: Ages 23 and 22, no children
    // etc
  }
};

struct initSlicedSingles : public thrust::unary_function<unsigned int, SingleParent>
{
  __device__ SingleParent operator()(const unsigned int& idx) const
  {
    return SingleParent(idx % 3, 25 + idx);
  }
};

int main()
{
  unsigned int Num_couples = 60;
  unsigned int Num_single_parents = 20;

  thrust::device_vector<Family> familiesInTown(Num_couples + Num_single_parents);
  // Families [0] to [59] are couples. Families [60] to [79] are single-parent households.
  thrust::transform(thrust::counting_iterator<unsigned int>(0),
                    thrust::counting_iterator<unsigned int>(Num_couples),
                    familiesInTown.begin(),
                    initSlicedCouples());
  thrust::transform(thrust::counting_iterator<unsigned int>(Num_couples),
                    thrust::counting_iterator<unsigned int>(Num_couples + Num_single_parents),
                    familiesInTown.begin() + Num_couples,
                    initSlicedSingles());
  return 0;
}

我肯定会犯一些经典的object slicing ...

所以，我问自己，一个指针向量可能会给我一些甜蜜的多态性呢？ Smart pointers在 C++ 中是一个东西，而 thrust迭代器可以做一些非常令人印象深刻的事情，所以让我们试一试吧，我想。以下代码编译。

struct initCouples : public thrust::unary_function<unsigned int, MarriedCouple*>
{
  __device__ MarriedCouple* operator()(const unsigned int& idx) const
  {
    return new MarriedCouple(idx % 3, 20 + idx, 19 + idx); // Memory issues?
  }
};
struct initSingles : public thrust::unary_function<unsigned int, SingleParent*>
{
  __device__ SingleParent* operator()(const unsigned int& idx) const
  {
    return new SingleParent(idx % 3, 25 + idx);
  }
};

int main()
{
  unsigned int Num_couples = 60;
  unsigned int Num_single_parents = 20;

  thrust::device_vector<Family*> familiesInTown(Num_couples + Num_single_parents);
  // Families [0] to [59] are couples. Families [60] to [79] are single-parent households.
  thrust::transform(thrust::counting_iterator<unsigned int>(0),
                    thrust::counting_iterator<unsigned int>(Num_couples),
                    familiesInTown.begin(),
                    initCouples()); 
  thrust::transform(thrust::counting_iterator<unsigned int>(Num_couples),
                    thrust::counting_iterator<unsigned int>(Num_couples + Num_single_parents),
                    familiesInTown.begin() + Num_couples,
                    initSingles());

  Family A = *(familiesInTown[2]); // Compiles, but object slicing takes place (in theory)
  std::cout << A.showNumOfParents() << "\n"; // Segmentation fault
 return 0;
}

好像我在这里碰壁了。我是否正确理解内存管理？ ( VTables 等)。我的对象是否在设备上被实例化和填充？我是否像没有明天一样泄漏内存？

对于它的值(value)，为了避免对象切片，我尝试使用 dynamic_cast<DerivedPointer*>(basePointer) .这就是为什么我做了我的Family析构函数 virtual .

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

以下行编译，但不幸的是，再次引发了段错误。 CUDA-Memcheck 不会告诉我原因。

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

和

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

简而言之，我需要一个对象的类接口(interface)，这些对象将具有不同的属性，彼此之间具有不同数量的成员，但我可以将其存储在一个可以操作的公共(public)向量中(这就是我想要一个基类的原因)显卡。我打算在 thrust 中与他们一起工作。通过 thrust::raw_pointer_cast 转换和在 CUDA 内核中ing，这对我来说完美无缺，直到我需要将我的类扩展到一个基础类和几个派生类。这样做的标准程序是什么？

提前致谢!

Answer 1

我不会试图回答这个问题的所有内容，它太大了。话虽如此，以下是对您发布的代码的一些观察，可能会有所帮助:

GPU方面new运算符从私有(private)运行时堆中分配内存。从 CUDA 6 开始，主机端 CUDA API 无法访问该内存。您可以从内核和设备函数中访问内存，但主机无法访问该内存。所以使用 new推力装置仿函数内部是一个永远无法工作的损坏设计。这就是您的“指针向量”模型失败的原因。

Thrust 的根本目的是允许将典型 STL 算法的数据并行版本应用于 POD 类型。使用复杂的多态对象构建代码库并尝试通过 Thrust 容器和算法填充它们可能会起作用，但这不是 Thrust 的设计目的，我不会推荐它。如果您以意想不到的方式打破推力，请不要感到惊讶。

CUDA 支持许多 C++ 特性，但编译和对象模型甚至比它们所基于的 C++98 标准要简单得多。 CUDA 缺少一些使复杂的多态对象设计在 C++ 中可行的关键特性(例如 RTTI)。我的建议是谨慎使用 C++ 功能。仅仅因为您可以在 CUDA 中做某事并不意味着您应该这样做。 GPU 是一个简单的架构，简单的数据结构和代码几乎总是比功能相似的复杂对象具有更高的性能。

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