文章/答案/技术大牛

发布

社区首页 >问答首页 >CUDA / CUDA推力中的多态性和派生类

问CUDA / CUDA推力中的多态性和派生类
EN

Stack Overflow用户

提问于 2014-04-10 12:35:01

回答 2查看 3.4K关注 0票数 5

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

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

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

Family是基类--家庭中父母的数量(单亲父母1人，夫妇2人)，以及他们作为成员存储在这里的子女数量。
SingleParent来源于Family，包括一个新成员--单亲父母的年龄，unsigned int ageOfParent。
然而，MarriedCouple也来自于Family，它引入了两个新成员--父母的年龄，unsigned int ageOfParent1和unsigned int ageOfParent2。 #include # #include 类族{受保护:无符号int numParents；无符号int numChildren；public：__host__ __device__族() {}；__host__ __device__系列( const unsigned int& nPars，const unsigned int& nChil)：numParents(nPars)，numChildren(nChil) {}；__host__ __device__虚拟~系列() {}；__host__ __device__无符号int showNumOfParents() { int；} __host__ __device__ unsigned int showNumOfChildren() { class；}；类SingleParent :公共家庭{受保护:无符号int ageOfParent；public：__host__ __device__ SingleParent() {}；__host__ __device__ SingleParent(const unsigned int& nChil，const unsigned int& age)：家庭(1，__device__)，__device__(Age) {}；__host__ __device__无符号int showAgeOfParent() {返回ageOfParent; }；类MarriedCouple :公共家庭{受保护:无符号int ageOfParent1；无符号int ageOfParent2；public：__host__ __device__ MarriedCouple() {}；__host__ __device__ MarriedCouple(无符号int& nChil，const unsigned int& age1，const无签名int& age2)：家庭(2，nChil)，ageOfParent1(NChil)，(en20 20#) {}；__host__ __device__无符号int showAgeOfParent1() {返回ageOfParent1;} __host__ __device__无符号int showAgeOfParent2() {返回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;
}

我肯定会犯一些经典的对象切片.

所以，我问自己，一个指针向量会给我一些甜蜜的多态吗？智能指针 in 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等)。我的对象是否被实例化并在设备上填充？我是不是漏了记忆，好像没有明天一样？

为了避免对象切片，我尝试了一个dynamic_cast<DerivedPointer*>(basePointer)。这就是为什么我做了我的Family析构函数virtual。

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

下面的行会编译，但不幸的是，会再次引发分段错误。库达-梅切克不告诉我原因。

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

和

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

简而言之，我需要的是一个对象的类接口，这些对象具有不同的属性，成员数目不同，但我可以存储在一个公共向量中(这就是我需要一个基类的原因)，我可以在GPU上操作它。我的目的是在thrust转换和CUDA内核中通过thrust::raw_pointer_casting与它们一起工作，在我需要将类扩展到一个基本类和几个派生类之前，这一点对我来说是完美的。标准的程序是什么？

提前感谢！

cuda

polymorphism

thrust

回答 2

Stack Overflow用户

回答已采纳

发布于 2014-04-12 11:11:52

我不打算回答这个问题的所有问题，因为它太大了。话虽如此，以下是关于您发布的代码的一些意见，这些意见可能会有所帮助：

GPU端new操作符从私有运行时堆中分配内存。从CUDA 6开始，主机端的CUDA API无法访问该内存。您可以从内核和设备函数中访问内存，但主机不能访问该内存。因此，在推力装置函子内使用new是一个无法工作的坏设计。这就是为什么您的“指针向量”模型失败的原因。
推力的基本意图是允许将典型STL算法的数据并行版本应用于POD类型。使用复杂的多态对象构建代码库，并尝试通过推力容器和算法填充代码，这可能是可行的，但这不是推力设计的目的，我不推荐它。如果你以意想不到的方式打破推力，那就不要感到惊讶。
CUDA支持许多C++特性，但是编译和对象模型比它们所基于的C++98标准要简单得多。CUDA缺乏几个关键特性(例如RTTI)，使得复杂的多态对象设计在C++中是可行的。我的建议是节约使用C++特性。仅仅因为你可以在数据自动化系统中做点什么并不意味着你应该。GPU是一个简单的体系结构，简单的数据结构和代码几乎总是比功能相似的复杂对象具有更高的性能。

在浏览了你发布的代码后，我的总体建议是回到画板上。如果您想看一些非常优雅的CUDA/C++设计，请花一些时间阅读幼崽和尖尖的代码库。它们都是非常不同的，但两者都有很多值得学习的地方(而且CUSP构建在推力之上，这使得它与您的用例更加相关，我怀疑这一点)。

票数 4

Stack Overflow用户

发布于 2014-05-05 15:42:08

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

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

然而，多态可以在设备代码中工作，包括推力装置功能。

下面的示例演示了这种想法，将自己限制在设备上创建的对象上，尽管我们肯定可以使用主机数据初始化它们。我创建了两个类，Triangle和Rectangle，它们来自一个基类Polygon，其中包含一个虚拟函数area。Triangle和Rectangle从基类继承函数set_values，但替换虚拟area函数。

然后，我们可以对这些类的对象进行多形性操作，如下所示：

#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进行了测试。

(多态示例思想和部分代码取自这里。)

票数 4

页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持

原文链接：

https://stackoverflow.com/questions/22988244

复制

相似问题

问CUDA / CUDA推力中的多态性和派生类
EN

回答 2

Stack Overflow用户

Stack Overflow用户

社区

活动

圈层

关于

腾讯云开发者

热门产品

热门推荐

更多推荐

问CUDA / CUDA推力中的多态性和派生类EN

回答 2

Stack Overflow用户

Stack Overflow用户

社区

活动

圈层

关于

腾讯云开发者

热门产品

热门推荐

更多推荐

问CUDA / CUDA推力中的多态性和派生类
EN